From ba8c1fc64854b2c6b3c76244bd228d9e2c6f5682 Mon Sep 17 00:00:00 2001
From: johnlu <chengjun.lu@intel.com>
Date: Wed, 8 Sep 2021 08:24:46 -0700
Subject: [PATCH] Add Half conversion of bit cast for SYCL kernel (#64340)

Summary:
## Motivation
Enhance the performance of Half/float conversion in SYCL kernels.

## Solution
Add the native SYCL half type to help convert the half from/to float in the kernel code.

## Additional Context
`__SYCL_DEVICE_ONLY__` is a MACRO only valid when compiling the kernel code for SYCL backend.

Pull Request resolved: https://github.com/pytorch/pytorch/pull/64340

Reviewed By: gchanan

Differential Revision: D30720823

Pulled By: ezyang

fbshipit-source-id: e7e770d02df5b2d45da61d2fed3ba59383b3dc3a
---
 c10/util/Half-inl.h | 8 ++++++++
 1 file changed, 8 insertions(+)

diff --git a/c10/util/Half-inl.h b/c10/util/Half-inl.h
index 5547874..1982dc7 100644
--- a/c10/util/Half-inl.h
+++ b/c10/util/Half-inl.h
@@ -12,6 +12,10 @@
 #include <hip/hip_fp16.h>
 #endif
 
+#ifdef __SYCL_DEVICE_ONLY__
+#include <CL/sycl.hpp>
+#endif
+
 namespace c10 {
 
 /// Constructors
@@ -19,6 +23,8 @@ namespace c10 {
 inline C10_HOST_DEVICE Half::Half(float value) {
 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
   x = __half_as_short(__float2half(value));
+#elif defined(__SYCL_DEVICE_ONLY__)
+  x = sycl::bit_cast<uint16_t>(sycl::half(value));
 #else
   x = detail::fp16_ieee_from_fp32_value(value);
 #endif
@@ -29,6 +35,8 @@ inline C10_HOST_DEVICE Half::Half(float value) {
 inline C10_HOST_DEVICE Half::operator float() const {
 #if defined(__CUDA_ARCH__) || defined(__HIP_DEVICE_COMPILE__)
   return __half2float(*reinterpret_cast<const __half*>(&x));
+#elif defined(__SYCL_DEVICE_ONLY__)
+  return float(sycl::bit_cast<sycl::half>(x));
 #else
   return detail::fp16_ieee_to_fp32_value(x);
 #endif
-- 
2.7.4