From 17bf4b68973a2902d07e03791aa3322de8182e78 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Fri, 5 Apr 2019 13:53:01 +0300
Subject: [PATCH 1/2] [SYCL] Add support for half type

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
---
 sycl/CMakeLists.txt                      |   1 +
 sycl/include/CL/sycl/half_type.hpp       |  84 +++++++
 sycl/include/CL/sycl/intel/sub_group.hpp |  21 +-
 sycl/include/CL/sycl/types.hpp           | 285 +++++++++++++++++------
 sycl/source/half_type.cpp                | 189 +++++++++++++++
 sycl/test/basic_tests/half_type.cpp      | 184 +++++++++++++++
 sycl/test/lit.cfg                        |   1 +
 sycl/test/sub_group/helper.hpp           |   8 +
 sycl/test/sub_group/load_store.cpp       |  15 +-
 sycl/test/sub_group/shuffle.cpp          |  13 +-
 10 files changed, 718 insertions(+), 83 deletions(-)
 create mode 100644 sycl/include/CL/sycl/half_type.hpp
 create mode 100644 sycl/source/half_type.cpp
 create mode 100644 sycl/test/basic_tests/half_type.cpp

diff --git a/sycl/CMakeLists.txt b/sycl/CMakeLists.txt
index 14ac967d233d8..57b66b2185227 100644
--- a/sycl/CMakeLists.txt
+++ b/sycl/CMakeLists.txt
@@ -99,6 +99,7 @@ add_library("${SYCLLibrary}" SHARED
   "${sourceRootPath}/device_selector.cpp"
   "${sourceRootPath}/event.cpp"
   "${sourceRootPath}/exception.cpp"
+  "${sourceRootPath}/half_type.cpp"
   "${sourceRootPath}/kernel.cpp"
   "${sourceRootPath}/platform.cpp"
   "${sourceRootPath}/queue.cpp"
diff --git a/sycl/include/CL/sycl/half_type.hpp b/sycl/include/CL/sycl/half_type.hpp
new file mode 100644
index 0000000000000..f3a6e57b5a66c
--- /dev/null
+++ b/sycl/include/CL/sycl/half_type.hpp
@@ -0,0 +1,84 @@
+//==-------------- half_type.hpp --- SYCL half type ------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#pragma once
+
+#include <cstdint>
+#include <functional>
+
+namespace cl {
+namespace sycl {
+namespace detail {
+namespace half_impl {
+
+class half {
+public:
+  half() = default;
+  half(const half &) = default;
+  half(half &&) = default;
+
+  half(const float &rhs);
+
+  half &operator=(const half &rhs) = default;
+
+  // Operator +=, -=, *=, /=
+  half &operator+=(const half &rhs);
+
+  half &operator-=(const half &rhs);
+
+  half &operator*=(const half &rhs);
+
+  half &operator/=(const half &rhs);
+
+  // Operator ++, --
+  half &operator++() {
+    *this += 1;
+    return *this;
+  }
+
+  half operator++(int) {
+    half ret(*this);
+    operator++();
+    return ret;
+  }
+
+  half &operator--() {
+    *this -= 1;
+    return *this;
+  }
+
+  half operator--(int) {
+    half ret(*this);
+    operator--();
+    return ret;
+  }
+
+  // Operator float
+  operator float() const;
+
+  template <typename Key> friend struct std::hash;
+
+private:
+  uint16_t Buf;
+};
+} // namespace half_impl
+} // namespace detail
+
+} // namespace sycl
+} // namespace cl
+
+namespace std {
+
+template <> struct hash<cl::sycl::detail::half_impl::half> {
+  size_t operator()(cl::sycl::detail::half_impl::half const &key) const
+      noexcept {
+    return hash<uint16_t>()(key.Buf);
+  }
+};
+
+} // namespace std
diff --git a/sycl/include/CL/sycl/intel/sub_group.hpp b/sycl/include/CL/sycl/intel/sub_group.hpp
index 6701a1d64e92b..e78e3d0af3bb0 100644
--- a/sycl/include/CL/sycl/intel/sub_group.hpp
+++ b/sycl/include/CL/sycl/intel/sub_group.hpp
@@ -139,11 +139,18 @@ struct sub_group {
     return BinaryOperation::template calc<T, cl::__spirv::InclusiveScan>(x);
   }
 
+  template <typename T>
+  using EnableIfIsArithmeticOrHalf = typename std::enable_if<
+      (std::is_arithmetic<T>::value ||
+       std::is_same<typename std::remove_const<T>::type, half>::value),
+      T>::type;
+
+
   /* --- one - input shuffles --- */
   /* indices in [0 , sub - group size ) */
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle(T x, id<1> local_id) {
     return cl::__spirv::OpSubgroupShuffleINTEL(x, local_id.get(0));
   }
@@ -156,7 +163,7 @@ struct sub_group {
   }
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle_down(T x, uint32_t delta) {
     return shuffle_down(x, x, delta);
   }
@@ -168,7 +175,7 @@ struct sub_group {
   }
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle_up(T x, uint32_t delta) {
     return shuffle_up(x, x, delta);
   }
@@ -180,7 +187,7 @@ struct sub_group {
   }
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle_xor(T x, id<1> value) {
     return cl::__spirv::OpSubgroupShuffleXorINTEL(x, (uint32_t)value.get(0));
   }
@@ -195,7 +202,7 @@ struct sub_group {
   /* --- two - input shuffles --- */
   /* indices in [0 , 2* sub - group size ) */
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle(T x, T y, id<1> local_id) {
     return cl::__spirv::OpSubgroupShuffleDownINTEL(
         x, y, local_id.get(0) - get_local_id().get(0));
@@ -210,7 +217,7 @@ struct sub_group {
   }
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle_down(T current, T next, uint32_t delta) {
     return cl::__spirv::OpSubgroupShuffleDownINTEL(current, next, delta);
   }
@@ -223,7 +230,7 @@ struct sub_group {
   }
 
   template <typename T>
-  typename std::enable_if<std::is_arithmetic<T>::value, T>::type
+  EnableIfIsArithmeticOrHalf<T>
   shuffle_up(T previous, T current, uint32_t delta) {
     return cl::__spirv::OpSubgroupShuffleUpINTEL(previous, current, delta);
   }
diff --git a/sycl/include/CL/sycl/types.hpp b/sycl/include/CL/sycl/types.hpp
index 64f93b8e6f5ea..da5d278c1ede9 100644
--- a/sycl/include/CL/sycl/types.hpp
+++ b/sycl/include/CL/sycl/types.hpp
@@ -45,11 +45,18 @@
 #endif // __HAS_EXT_VECTOR_TYPE__
 
 #include <CL/sycl/detail/common.hpp>
+#include <CL/sycl/half_type.hpp>
 #include <CL/sycl/multi_ptr.hpp>
 
 // 4.10.1: Scalar data types
 // 4.10.2: SYCL vector types
 
+#ifdef __SYCL_DEVICE_ONLY__
+using half = _Float16;
+#else
+using half = cl::sycl::detail::half_impl::half;
+#endif
+
 namespace cl {
 namespace sycl {
 
@@ -373,16 +380,45 @@ template <typename Type, int NumElements> class vec {
   }
 
 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
-  explicit vec(const DataT &arg) {
+  template <typename T = void>
+  using EnableIfNotHostHalf = typename std::enable_if<
+      !std::is_same<DataT, cl::sycl::detail::half_impl::half>::value, T>::type;
+  template <typename T = void>
+  using EnableIfHostHalf = typename std::enable_if<
+      std::is_same<DataT, cl::sycl::detail::half_impl::half>::value, T>::type;
+
+  template <typename Ty = DataT>
+  explicit vec(const EnableIfNotHostHalf<Ty> &arg) {
     m_Data = (DataType)arg;
   }
 
   template <typename Ty = DataT>
-  typename std::enable_if<std::is_fundamental<Ty>::value, vec &>::type
-  operator=(const DataT &Rhs) {
+  typename std::enable_if<
+      std::is_fundamental<Ty>::value ||
+          std::is_same<typename std::remove_const<Ty>::type, half>::value,
+      vec &>::type
+  operator=(const EnableIfNotHostHalf<Ty> &Rhs) {
     m_Data = (DataType)Rhs;
     return *this;
   }
+
+  template <typename Ty = DataT> explicit vec(const EnableIfHostHalf<Ty> &arg) {
+    for (int i = 0; i < NumElements; ++i) {
+      setValue(i, arg);
+    }
+  }
+
+  template <typename Ty = DataT>
+  typename std::enable_if<
+      std::is_fundamental<Ty>::value ||
+          std::is_same<typename std::remove_const<Ty>::type, half>::value,
+      vec &>::type
+  operator=(const EnableIfHostHalf<Ty> &Rhs) {
+    for (int i = 0; i < NumElements; ++i) {
+      setValue(i, Rhs);
+    }
+    return *this;
+  }
 #else
   explicit vec(const DataT &arg) {
     for (int i = 0; i < NumElements; ++i) {
@@ -391,7 +427,10 @@ template <typename Type, int NumElements> class vec {
   }
 
   template <typename Ty = DataT>
-  typename std::enable_if<std::is_fundamental<Ty>::value, vec &>::type
+  typename std::enable_if<
+      std::is_fundamental<Ty>::value ||
+          std::is_same<typename std::remove_const<Ty>::type, half>::value,
+      vec &>::type
   operator=(const DataT &Rhs) {
     for (int i = 0; i < NumElements; ++i) {
       setValue(i, Rhs);
@@ -407,32 +446,32 @@ template <typename Type, int NumElements> class vec {
   // Helper type to make specific constructors available only for specific
   // number of elements.
   template <int IdxNum, typename T = void>
-  using EnableIfMultipleElems =
-      typename std::enable_if<std::is_convertible<T, DataT>::value &&
-                                  NumElements == IdxNum,
-                              DataT>::type;
+  using EnableIfMultipleElems = typename std::enable_if<
+      std::is_convertible<T, DataT>::value && NumElements == IdxNum,
+      DataT>::type;
   template <typename Ty = DataT>
-  vec(const EnableIfMultipleElems<2, Ty> Arg0, const DataT Arg1)
+  vec(const EnableIfMultipleElems<2, Ty> Arg0,
+      const EnableIfNotHostHalf<Ty> Arg1)
       : m_Data{Arg0, Arg1} {}
   template <typename Ty = DataT>
-  vec(const EnableIfMultipleElems<3, Ty> Arg0, const DataT Arg1,
-      const DataT Arg2)
+  vec(const EnableIfMultipleElems<3, Ty> Arg0,
+      const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2)
       : m_Data{Arg0, Arg1, Arg2} {}
   template <typename Ty = DataT>
-  vec(const EnableIfMultipleElems<4, Ty> Arg0, const DataT Arg1,
-      const DataT Arg2, const Ty Arg3)
+  vec(const EnableIfMultipleElems<4, Ty> Arg0,
+      const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2, const Ty Arg3)
       : m_Data{Arg0, Arg1, Arg2, Arg3} {}
   template <typename Ty = DataT>
-  vec(const EnableIfMultipleElems<8, Ty> Arg0, const DataT Arg1,
-      const DataT Arg2, const DataT Arg3, const DataT Arg4, const DataT Arg5,
-      const DataT Arg6, const DataT Arg7)
+  vec(const EnableIfMultipleElems<8, Ty> Arg0,
+      const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2, const DataT Arg3,
+      const DataT Arg4, const DataT Arg5, const DataT Arg6, const DataT Arg7)
       : m_Data{Arg0, Arg1, Arg2, Arg3, Arg4, Arg5, Arg6, Arg7} {}
   template <typename Ty = DataT>
-  vec(const EnableIfMultipleElems<16, Ty> Arg0, const DataT Arg1,
-      const DataT Arg2, const DataT Arg3, const DataT Arg4, const DataT Arg5,
-      const DataT Arg6, const DataT Arg7, const DataT Arg8, const DataT Arg9,
-      const DataT ArgA, const DataT ArgB, const DataT ArgC, const DataT ArgD,
-      const DataT ArgE, const DataT ArgF)
+  vec(const EnableIfMultipleElems<16, Ty> Arg0,
+      const EnableIfNotHostHalf<Ty> Arg1, const DataT Arg2, const DataT Arg3,
+      const DataT Arg4, const DataT Arg5, const DataT Arg6, const DataT Arg7,
+      const DataT Arg8, const DataT Arg9, const DataT ArgA, const DataT ArgB,
+      const DataT ArgC, const DataT ArgD, const DataT ArgE, const DataT ArgF)
       : m_Data{Arg0, Arg1, Arg2, Arg3, Arg4, Arg5, Arg6, Arg7,
                Arg8, Arg9, ArgA, ArgB, ArgC, ArgD, ArgE, ArgF} {}
 #endif
@@ -580,15 +619,26 @@ template <typename Type, int NumElements> class vec {
 
 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
 #define __SYCL_BINOP(BINOP, OPASSIGN)                                          \
-  vec operator BINOP(const vec &Rhs) const {                                   \
+  template <typename Ty = vec>                                                 \
+  vec operator BINOP(const EnableIfNotHostHalf<Ty> &Rhs) const {               \
     vec Ret;                                                                   \
     Ret.m_Data = m_Data BINOP Rhs.m_Data;                                      \
     return Ret;                                                                \
   }                                                                            \
+  template <typename Ty = vec>                                                 \
+  vec operator BINOP(const EnableIfHostHalf<Ty> &Rhs) const {                  \
+    vec Ret;                                                                   \
+    for (size_t I = 0; I < NumElements; ++I) {                                 \
+      Ret.setValue(I, (getValue(I) BINOP Rhs.getValue(I)));                    \
+    }                                                                          \
+    return Ret;                                                                \
+  }                                                                            \
   template <typename T>                                                        \
-  typename std::enable_if<std::is_convertible<T, DataT>::value &&              \
-                              std::is_fundamental<T>::value,                   \
-                          vec>::type                                           \
+  typename std::enable_if<                                                     \
+      std::is_convertible<DataT, T>::value &&                                  \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec>::type                                                               \
   operator BINOP(const T &Rhs) const {                                         \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
   }                                                                            \
@@ -597,8 +647,8 @@ template <typename Type, int NumElements> class vec {
     return *this;                                                              \
   }                                                                            \
   template <int Num = NumElements>                                             \
-  typename std::enable_if<Num != 1, vec &>::type operator OPASSIGN(            \
-      const DataT &Rhs) {                                                      \
+  typename std::enable_if<Num != 1, vec &>::type                               \
+  operator OPASSIGN(const DataT &Rhs) {                                        \
     *this = *this BINOP vec(Rhs);                                              \
     return *this;                                                              \
   }
@@ -612,9 +662,11 @@ template <typename Type, int NumElements> class vec {
     return Ret;                                                                \
   }                                                                            \
   template <typename T>                                                        \
-  typename std::enable_if<std::is_convertible<T, DataT>::value &&              \
-                              std::is_fundamental<T>::value,                   \
-                          vec>::type                                           \
+  typename std::enable_if<                                                     \
+      std::is_convertible<DataT, T>::value &&                                  \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec>::type                                                               \
   operator BINOP(const T &Rhs) const {                                         \
     return *this BINOP vec(static_cast<const DataT &>(Rhs));                   \
   }                                                                            \
@@ -666,9 +718,11 @@ template <typename Type, int NumElements> class vec {
     return Ret;                                                                \
   }                                                                            \
   template <typename T>                                                        \
-  typename std::enable_if<std::is_convertible<T, DataT>::value &&              \
-                              std::is_fundamental<T>::value,                   \
-                          vec<rel_t, NumElements>>::type                       \
+  typename std::enable_if<                                                     \
+      std::is_convertible<T, DataT>::value &&                                  \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec<rel_t, NumElements>>::type                                           \
   operator RELLOGOP(const T &Rhs) const {                                      \
     return *this RELLOGOP vec(static_cast<const DataT &>(Rhs));                \
   }
@@ -729,36 +783,67 @@ template <typename Type, int NumElements> class vec {
   // vec<RET, NumElements> operatorOP(const DataT &Rhs) const;
 private:
   // Generic method that execute "Operation" on underlying values.
-  template <template <typename> class Operation>
+#ifdef __SYCL_USE_EXT_VECTOR_TYPE__
+  template <template <typename> class Operation,
+            typename Ty = vec<DataT, NumElements>>
   vec<DataT, NumElements>
-  operatorHelper(const vec<DataT, NumElements> &Rhs) const {
+  operatorHelper(const EnableIfNotHostHalf<Ty> &Rhs) const {
     vec<DataT, NumElements> Result;
-#ifdef __SYCL_USE_EXT_VECTOR_TYPE__
     Operation<DataType> Op;
     Result.m_Data = Op(m_Data, Rhs.m_Data);
+    return Result;
+  }
+
+  template <template <typename> class Operation,
+            typename Ty = vec<DataT, NumElements>>
+  vec<DataT, NumElements>
+  operatorHelper(const EnableIfHostHalf<Ty> &Rhs) const {
+    vec<DataT, NumElements> Result;
+    Operation<DataT> Op;
+    for (size_t I = 0; I < NumElements; ++I) {
+      Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
+    }
+    return Result;
+  }
 #else  // __SYCL_USE_EXT_VECTOR_TYPE__
+  template <template <typename> class Operation>
+  vec<DataT, NumElements>
+  operatorHelper(const vec<DataT, NumElements> &Rhs) const {
+    vec<DataT, NumElements> Result;
     Operation<DataT> Op;
     for (size_t I = 0; I < NumElements; ++I) {
       Result.setValue(I, Op(Rhs.getValue(I), getValue(I)));
     }
-#endif // __SYCL_USE_EXT_VECTOR_TYPE__
     return Result;
   }
+#endif // __SYCL_USE_EXT_VECTOR_TYPE__
 
 // setValue and getValue should be able to operate on different underlying
 // types: enum cl_float#N , builtin vector float#N, builtin type float.
 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
-  template <int Num = NumElements,
+  template <int Num = NumElements, typename Ty = int,
             typename = typename std::enable_if<1 != Num>::type>
-  void setValue(int Index, const DataT &Value, int) {
+  void setValue(EnableIfNotHostHalf<Ty> Index, const DataT &Value, int) {
     m_Data[Index] = Value;
   }
 
-  template <int Num = NumElements,
+  template <int Num = NumElements, typename Ty = int,
             typename = typename std::enable_if<1 != Num>::type>
-  DataT getValue(int Index, int) const {
+  DataT getValue(EnableIfNotHostHalf<Ty> Index, int) const {
     return m_Data[Index];
   }
+
+  template <int Num = NumElements, typename Ty = int,
+            typename = typename std::enable_if<1 != Num>::type>
+  void setValue(EnableIfHostHalf<Ty> Index, const DataT &Value, int) {
+    m_Data.s[Index] = Value;
+  }
+
+  template <int Num = NumElements, typename Ty = int,
+            typename = typename std::enable_if<1 != Num>::type>
+  DataT getValue(EnableIfHostHalf<Ty> Index, int) const {
+    return m_Data.s[Index];
+  }
 #else  // __SYCL_USE_EXT_VECTOR_TYPE__
   template <int Num = NumElements,
             typename = typename std::enable_if<1 != Num>::type>
@@ -908,14 +993,16 @@ class SwizzleOp {
       1 != IdxNum && SwizzleOp::getNumElements() == IdxNum, T>::type;
 
   template <typename T>
-  using EnableIfScalarType =
-      typename std::enable_if<std::is_convertible<DataT, T>::value &&
-                              std::is_fundamental<T>::value>::type;
+  using EnableIfScalarType = typename std::enable_if<
+      std::is_convertible<DataT, T>::value &&
+      (std::is_fundamental<T>::value ||
+       std::is_same<typename std::remove_const<T>::type, half>::value)>::type;
 
   template <typename T>
-  using EnableIfNoScalarType =
-      typename std::enable_if<!std::is_convertible<DataT, T>::value ||
-                              !std::is_fundamental<T>::value>::type;
+  using EnableIfNoScalarType = typename std::enable_if<
+      !std::is_convertible<DataT, T>::value ||
+      !(std::is_fundamental<T>::value ||
+        std::is_same<typename std::remove_const<T>::type, half>::value)>::type;
 
   template <int... Indices>
   using Swizzle =
@@ -1437,7 +1524,10 @@ class SwizzleOp {
 #endif
 #define __SYCL_BINOP(BINOP)                                                    \
   template <typename T, int Num>                                               \
-  typename std::enable_if<std::is_fundamental<T>::value, vec<T, Num>>::type    \
+  typename std::enable_if<                                                     \
+      std::is_fundamental<T>::value ||                                         \
+          std::is_same<typename std::remove_const<T>::type, half>::value,      \
+      vec<T, Num>>::type                                                       \
   operator BINOP(const T &Lhs, const vec<T, Num> &Rhs) {                       \
     return vec<T, Num>(Lhs) BINOP Rhs;                                         \
   }                                                                            \
@@ -1445,13 +1535,14 @@ class SwizzleOp {
             template <typename> class OperationCurrentT, int... Indexes,       \
             typename T, typename T1 = typename VecT::element_type,             \
             int Num = sizeof...(Indexes)>                                      \
-  typename std::enable_if<std::is_convertible<T, T1>::value &&                 \
-                              std::is_fundamental<T>::value,                   \
-                          vec<T1, Num>>::type                                  \
-  operator BINOP(                                                              \
-      const T &Lhs,                                                            \
-      const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT,           \
-                              OperationCurrentT, Indexes...> &Rhs) {           \
+  typename std::enable_if<                                                     \
+      std::is_convertible<T, T1>::value &&                                     \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec<T1, Num>>::type                                                      \
+  operator BINOP(const T &Lhs,                                                 \
+                 const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT,\
+                                         OperationCurrentT, Indexes...> &Rhs) {\
     vec<T1, Num> Tmp = Rhs;                                                    \
     return Lhs BINOP Tmp;                                                      \
   }                                                                            \
@@ -1486,9 +1577,11 @@ __SYCL_BINOP(<<)
 #endif
 #define __SYCL_RELLOGOP(RELLOGOP)                                              \
   template <typename T, typename DataT, int Num>                               \
-  typename std::enable_if<std::is_convertible<T, DataT>::value &&              \
-                              std::is_fundamental<T>::value,                   \
-                          vec<detail::rel_t<DataT>, Num>>::type                \
+  typename std::enable_if<                                                     \
+      std::is_convertible<T, DataT>::value &&                                  \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec<detail::rel_t<DataT>, Num>>::type                                    \
   operator RELLOGOP(const T &Lhs, const vec<DataT, Num> &Rhs) {                \
     return vec<T, Num>(static_cast<T>(Lhs)) RELLOGOP Rhs;                      \
   }                                                                            \
@@ -1496,9 +1589,11 @@ __SYCL_BINOP(<<)
             template <typename> class OperationCurrentT, int... Indexes,       \
             typename T, typename T1 = typename VecT::element_type,             \
             int Num = sizeof...(Indexes)>                                      \
-  typename std::enable_if<std::is_convertible<T, T1>::value &&                 \
-                              std::is_fundamental<T>::value,                   \
-                          vec<detail::rel_t<T1>, Num>>::type                   \
+  typename std::enable_if<                                                     \
+      std::is_convertible<T, T1>::value &&                                     \
+          (std::is_fundamental<T>::value ||                                    \
+           std::is_same<typename std::remove_const<T>::type, half>::value),    \
+      vec<detail::rel_t<T1>, Num>>::type                                       \
   operator RELLOGOP(                                                           \
       const T &Lhs,                                                            \
       const detail::SwizzleOp<VecT, OperationLeftT, OperationRightT,           \
@@ -1531,6 +1626,7 @@ __SYCL_RELLOGOP(||)
 } // namespace sycl
 } // namespace cl
 
+
 #ifdef __SYCL_USE_EXT_VECTOR_TYPE__
 typedef char __char_t;
 typedef char __char2_vec_t __attribute__((ext_vector_type(2)));
@@ -1609,13 +1705,18 @@ typedef float __float3_vec_t __attribute__((ext_vector_type(3)));
 typedef float __float4_vec_t __attribute__((ext_vector_type(4)));
 typedef float __float8_vec_t __attribute__((ext_vector_type(8)));
 typedef float __float16_vec_t __attribute__((ext_vector_type(16)));
-// TODO: Add support for half builtin type in clang for sycl target.
-// typedef half __half_t;
-// typedef half __half2_vec_t __attribute__((ext_vector_type(2)));
-// typedef half __half3_vec_t __attribute__((ext_vector_type(3)));
-// typedef half __half4_vec_t __attribute__((ext_vector_type(4)));
-// typedef half __half8_vec_t __attribute__((ext_vector_type(8)));
-// typedef half __half16_vec_t __attribute__((ext_vector_type(16)));
+// Half type is defined as custom class for host and _Float16 for device.
+// The ext_vector_type attribute is only applicable to integral and float
+// scalars so it's not possible to use attribute ext_vector_type for half on
+// host.
+#ifdef __SYCL_DEVICE_ONLY__
+typedef half __half_t;
+typedef half __half2_vec_t __attribute__((ext_vector_type(2)));
+typedef half __half3_vec_t __attribute__((ext_vector_type(3)));
+typedef half __half4_vec_t __attribute__((ext_vector_type(4)));
+typedef half __half8_vec_t __attribute__((ext_vector_type(8)));
+typedef half __half16_vec_t __attribute__((ext_vector_type(16)));
+#endif
 typedef double __double_t;
 typedef double __double2_vec_t __attribute__((ext_vector_type(2)));
 typedef double __double3_vec_t __attribute__((ext_vector_type(3)));
@@ -1625,6 +1726,7 @@ typedef double __double16_vec_t __attribute__((ext_vector_type(16)));
 
 #define GET_CL_TYPE(target, num) __##target##num##_vec_t
 #define GET_SCALAR_CL_TYPE(target) target
+
 #else // __SYCL_USE_EXT_VECTOR_TYPE__
 // For signed char. OpenCL doesn't have any type about `signed char`, therefore
 // we use type alias of cl_char instead.
@@ -1639,6 +1741,12 @@ using cl_schar16 = cl_char16;
 #define GET_SCALAR_CL_TYPE(target) cl_##target
 #endif // __SYCL_USE_EXT_VECTOR_TYPE__
 
+#ifndef __SYCL_DEVICE_ONLY__
+#define GET_CL_HALF_TYPE(target, num) cl_##target##num
+#else
+#define GET_CL_HALF_TYPE(target, num) __##target##num##_vec_t
+#endif
+
 namespace cl {
 namespace sycl {
 
@@ -1654,6 +1762,12 @@ namespace sycl {
     using DataType = ::GET_CL_TYPE(base, num);                                 \
   };
 
+#define DECLARE_HALF_CONVERTER(base, num)                                      \
+  template <> class BaseCLTypeConverter<base, num> {                           \
+  public:                                                                      \
+    using DataType = GET_CL_HALF_TYPE(base, num);                              \
+  };
+
 #define DECLARE_VECTOR_CONVERTERS(base)                                        \
   namespace detail {                                                           \
   DECLARE_CONVERTER(base, 2)                                                   \
@@ -1667,6 +1781,19 @@ namespace sycl {
   };                                                                           \
   } // namespace detail
 
+#define DECLARE_HALF_VECTOR_CONVERTERS(base)                                   \
+  namespace detail {                                                           \
+  DECLARE_HALF_CONVERTER(base, 2)                                              \
+  DECLARE_HALF_CONVERTER(base, 3)                                              \
+  DECLARE_HALF_CONVERTER(base, 4)                                              \
+  DECLARE_HALF_CONVERTER(base, 8)                                              \
+  DECLARE_HALF_CONVERTER(base, 16)                                             \
+  template <> class BaseCLTypeConverter<base, 1> {                             \
+  public:                                                                      \
+    using DataType = half;                                                     \
+  };                                                                           \
+  } // namespace detail
+
 #define DECLARE_VECTOR_LONGLONG_CONVERTERS(base)                               \
   namespace detail {                                                           \
   DECLARE_LONGLONG_CONVERTER(base, 2)                                          \
@@ -1706,6 +1833,19 @@ namespace sycl {
   using char3 = vec<char, 3>;                                                  \
   using char2 = vec<char, 2>;
 
+#define DECLARE_SYCL_VEC_HALF_WO_CONVERTERS                                    \
+  using cl_half16 = vec<half, 16>;                                             \
+  using cl_half8 = vec<half, 8>;                                               \
+  using cl_half4 = vec<half, 4>;                                               \
+  using cl_half3 = vec<half, 3>;                                               \
+  using cl_half2 = vec<half, 2>;                                               \
+  using cl_half = half;                                                        \
+  using half16 = vec<half, 16>;                                                \
+  using half8 = vec<half, 8>;                                                  \
+  using half4 = vec<half, 4>;                                                  \
+  using half3 = vec<half, 3>;                                                  \
+  using half2 = vec<half, 2>;
+
 // cl_longlong/cl_ulonglong are not supported in SYCL
 #define DECLARE_SYCL_VEC_LONGLONG_WO_CONVERTERS(base)                          \
   using base##long16 = vec<base##long, 16>;                                    \
@@ -1726,6 +1866,10 @@ namespace sycl {
   DECLARE_VECTOR_LONGLONG_CONVERTERS(base)                                     \
   DECLARE_SYCL_VEC_LONGLONG_WO_CONVERTERS(base)
 
+#define DECLARE_SYCL_VEC_HALF(base)                                            \
+  DECLARE_HALF_VECTOR_CONVERTERS(base)                                         \
+  DECLARE_SYCL_VEC_HALF_WO_CONVERTERS
+
 DECLARE_SYCL_VEC_CHAR
 DECLARE_SYCL_VEC(schar)
 DECLARE_SYCL_VEC(uchar)
@@ -1737,14 +1881,13 @@ DECLARE_SYCL_VEC(long)
 DECLARE_SYCL_VEC(ulong)
 DECLARE_SYCL_VEC_LONGLONG(long)
 DECLARE_SYCL_VEC_LONGLONG(ulong)
+DECLARE_SYCL_VEC_HALF(half)
 DECLARE_SYCL_VEC(float)
 DECLARE_SYCL_VEC(double)
-// TODO: Fix half.
-typedef ::cl_half half;
-DECLARE_SYCL_VEC_WO_CONVERTERS(half)
 
 using cl_bool = ::cl_bool;
 using byte = uchar;
+using half = ::half;
 
 #undef GET_CL_TYPE
 #undef GET_SCALAR_CL_TYPE
diff --git a/sycl/source/half_type.cpp b/sycl/source/half_type.cpp
new file mode 100644
index 0000000000000..3b11ef1498266
--- /dev/null
+++ b/sycl/source/half_type.cpp
@@ -0,0 +1,189 @@
+//==-------------- half_type.cpp --- SYCL half type ------------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl/half_type.hpp>
+#include <iostream>
+
+namespace cl {
+namespace sycl {
+namespace detail {
+
+static uint16_t float2Half(const float &Val) {
+  const uint32_t Bits = *reinterpret_cast<const uint32_t *>(&Val);
+
+  // Extract the sign from the float value
+  const uint16_t Sign = (Bits & 0x80000000) >> 16;
+  // Extract the fraction from the float value
+  const uint32_t Frac32 = Bits & 0x7fffff;
+  // Extract the exponent from the float value
+  const uint8_t Exp32 = (Bits & 0x7f800000) >> 23;
+  const int8_t Exp32Diff = Exp32 - 127;
+
+  uint16_t Exp16 = 0;
+  uint16_t Frac16 = Frac32 >> 13;
+
+  if (__builtin_expect(Exp32 == 0xff || Exp32Diff > 15, 0)) {
+    Exp16 = 0x1f;
+  } else if (__builtin_expect(Exp32 == 0 || Exp32Diff < -14, 0)) {
+    Exp16 = 0;
+  } else {
+    Exp16 = Exp32Diff + 15;
+  }
+
+  if (__builtin_expect(Exp32 == 0xff && Frac32 != 0 && Frac16 == 0, 0)) {
+    // corner case 1: NaN
+    // This case happens when FP32 value is NaN whose the fraction part
+    // transformed to FP16 counterpart is truncated to 0. We need to flip the
+    // high bit to 1 to make it distinguished from inf.
+    Frac16 = 0x200;
+  } else if (__builtin_expect(Exp32 == 0 || (Exp16 == 0x1f && Exp32 != 0xff),
+                              0)) {
+    // corner case 2: subnormal
+    // All FP32 subnormal values are under the range of FP16 so the fraction
+    // part is set to 0.
+    // corner case 3: overflow
+    Frac16 = 0;
+  } else if (__builtin_expect(Exp16 == 0 && Exp32 != 0, 0)) {
+    // corner case 4: underflow
+    // We use `truncate` mode here.
+    Frac16 = 0x100 | (Frac16 >> 2);
+  }
+
+  // Compose the final FP16 binary
+  uint16_t Ret = 0;
+  Ret |= Sign;
+  Ret |= Exp16 << 10;
+  Ret |= Frac16;
+
+  return Ret;
+}
+
+static float half2Float(const uint16_t &Val) {
+  // Extract the sign from the bits
+  const uint32_t Sign = static_cast<uint32_t>(Val & 0x8000) << 16;
+  // Extract the exponent from the bits
+  const uint8_t Exp16 = (Val & 0x7c00) >> 10;
+  // Extract the fraction from the bits
+  uint16_t Frac16 = Val & 0x3ff;
+
+  uint32_t Exp32 = 0;
+  if (__builtin_expect(Exp16 == 0x1f, 0)) {
+    Exp32 = 0xff;
+  } else if (__builtin_expect(Exp16 == 0, 0)) {
+    Exp32 = 0;
+  } else {
+    Exp32 = static_cast<uint32_t>(Exp16) + 112;
+  }
+
+  // corner case: subnormal -> normal
+  // The denormal number of FP16 can be represented by FP32, therefore we need
+  // to recover the exponent and recalculate the fration.
+  if (__builtin_expect(Exp16 == 0 && Frac16 != 0, 0)) {
+    uint8_t OffSet = 0;
+    do {
+      ++OffSet;
+      Frac16 <<= 1;
+    } while ((Frac16 & 0x400) != 0x400);
+    // mask the 9th bit
+    Frac16 &= 0x3ff;
+    Exp32 = 113 - OffSet;
+  }
+
+  uint32_t Frac32 = Frac16 << 13;
+
+  // Compose the final FP32 binary
+  uint32_t Bits = 0;
+
+  Bits |= Sign;
+  Bits |= (Exp32 << 23);
+  Bits |= Frac32;
+
+  return *reinterpret_cast<float *>(&Bits);
+}
+
+std::ostream &operator<<(std::ostream &O, const half_impl::half &Val) {
+  O << static_cast<float>(Val);
+  return O;
+}
+
+std::istream &operator>>(std::istream &I, half_impl::half &ValHalf) {
+  float ValFloat = 0.0f;
+  I >> ValFloat;
+  ValHalf = ValFloat;
+  return I;
+}
+
+namespace half_impl {
+
+half::half(const float &RHS) : Buf(float2Half(RHS)) {}
+
+half &half::operator+=(const half &RHS) {
+  *this = operator float() + static_cast<float>(RHS);
+  return *this;
+}
+
+half &half::operator-=(const half &RHS) {
+  *this = operator float() - static_cast<float>(RHS);
+  return *this;
+}
+
+half &half::operator*=(const half &RHS) {
+  *this = operator float() * static_cast<float>(RHS);
+  return *this;
+}
+
+half &half::operator/=(const half &RHS) {
+  *this = operator float() / static_cast<float>(RHS);
+  return *this;
+}
+
+half::operator float() const { return half2Float(Buf); }
+
+// Operator +, -, *, /
+half operator+(half LHS, const half &RHS) {
+  LHS += RHS;
+  return LHS;
+}
+
+half operator-(half LHS, const half &RHS) {
+  LHS -= RHS;
+  return LHS;
+}
+
+half operator*(half LHS, const half &RHS) {
+  LHS *= RHS;
+  return LHS;
+}
+
+half operator/(half LHS, const half &RHS) {
+  LHS /= RHS;
+  return LHS;
+}
+
+// Operator <, >, <=, >=
+bool operator<(const half &LHS, const half &RHS) {
+  return static_cast<float>(LHS) < static_cast<float>(RHS);
+}
+
+bool operator>(const half &LHS, const half &RHS) { return RHS < LHS; }
+
+bool operator<=(const half &LHS, const half &RHS) { return !(LHS > RHS); }
+
+bool operator>=(const half &LHS, const half &RHS) { return !(LHS < RHS); }
+
+// Operator ==, !=
+bool operator==(const half &LHS, const half &RHS) {
+  return static_cast<float>(LHS) == static_cast<float>(RHS);
+}
+
+bool operator!=(const half &LHS, const half &RHS) { return !(LHS == RHS); }
+} // namespace half_impl
+
+} // namespace detail
+} // namespace sycl
+} // namespace cl
diff --git a/sycl/test/basic_tests/half_type.cpp b/sycl/test/basic_tests/half_type.cpp
new file mode 100644
index 0000000000000..fe30a4aa5beaf
--- /dev/null
+++ b/sycl/test/basic_tests/half_type.cpp
@@ -0,0 +1,184 @@
+// RUN: %clang -std=c++11 -fsycl %s -o %t.out -lstdc++ -lOpenCL
+// RUN: env SYCL_DEVICE_TYPE=HOST %t.out
+// RUN: %CPU_RUN_PLACEHOLDER %t.out
+// RUN: %GPU_RUN_PLACEHOLDER %t.out
+// RUN: %ACC_RUN_PLACEHOLDER %t.out
+
+//==--------------- half_type.cpp - SYCL half type test --------------------==//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+
+#include <CL/sycl.hpp>
+
+using namespace cl::sycl;
+
+constexpr float FLT_EPSILON = 9.77e-4;
+
+constexpr size_t N = 100;
+
+template <typename T> void assert_close(const T &C, const float ref) {
+  for (size_t i = 0; i < N; i++) {
+    float diff = C[i] - ref;
+    assert(std::fabs(diff) < FLT_EPSILON);
+  }
+}
+
+void verify_add(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
+                const float ref) {
+  buffer<half, 1> c{r};
+
+  q.submit([&](handler &cgh) {
+    auto A = a.get_access<access::mode::read>(cgh);
+    auto B = b.get_access<access::mode::read>(cgh);
+    auto C = c.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<class calc_add>(
+        r, [=](id<1> index) { C[index] = A[index] + B[index]; });
+  });
+
+  assert_close(c.get_access<access::mode::read>(), ref);
+}
+
+void verify_min(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
+                const float ref) {
+  buffer<half, 1> c{r};
+
+  q.submit([&](handler &cgh) {
+    auto A = a.get_access<access::mode::read>(cgh);
+    auto B = b.get_access<access::mode::read>(cgh);
+    auto C = c.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<class calc_min>(
+        r, [=](id<1> index) { C[index] = A[index] - B[index]; });
+  });
+
+  assert_close(c.get_access<access::mode::read>(), ref);
+}
+
+void verify_mul(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
+                const float ref) {
+  buffer<half, 1> c{r};
+
+  q.submit([&](handler &cgh) {
+    auto A = a.get_access<access::mode::read>(cgh);
+    auto B = b.get_access<access::mode::read>(cgh);
+    auto C = c.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<class calc_mul>(
+        r, [=](id<1> index) { C[index] = A[index] * B[index]; });
+  });
+
+  assert_close(c.get_access<access::mode::read>(), ref);
+}
+
+void verify_div(queue &q, buffer<half, 1> &a, buffer<half, 1> &b, range<1> &r,
+                const float ref) {
+  buffer<half, 1> c{r};
+
+  q.submit([&](handler &cgh) {
+    auto A = a.get_access<access::mode::read>(cgh);
+    auto B = b.get_access<access::mode::read>(cgh);
+    auto C = c.get_access<access::mode::write>(cgh);
+    cgh.parallel_for<class calc_div>(
+        r, [=](id<1> index) { C[index] = A[index] / B[index]; });
+  });
+
+  assert_close(c.get_access<access::mode::read>(), ref);
+}
+
+inline bool bitwise_comparison_fp16(const half val, const uint16_t exp) {
+  return reinterpret_cast<const uint16_t&>(val) == exp;
+}
+
+inline bool bitwise_comparison_fp32(const half val, const uint32_t exp) {
+  const float fp32 = static_cast<float>(val);
+  return reinterpret_cast<const uint32_t&>(fp32) == exp;
+}
+
+int main() {
+  // We assert that the length is 1 because we use macro to select the device
+  assert(device::get_devices().size() == 1);
+
+  auto dev = device::get_devices()[0];
+  if (!dev.is_host() && !dev.has_extension("cl_khr_fp16")) {
+    std::cout << "This device doesn't support the extension cl_khr_fp16"
+              << std::endl;
+    return 0;
+  }
+
+  std::vector<half> vec_a(N, 5.0);
+  std::vector<half> vec_b(N, 2.0);
+
+  range<1> r(N);
+  buffer<half, 1> a{vec_a.data(), r};
+  buffer<half, 1> b{vec_b.data(), r};
+
+  queue q;
+
+  verify_add(q, a, b, r, 7.0);
+  verify_min(q, a, b, r, 3.0);
+  verify_mul(q, a, b, r, 10.0);
+  verify_div(q, a, b, r, 2.5);
+
+  if (!dev.is_host()) {
+    return 0;
+  }
+
+  // Basic tests: fp32->fp16
+  // The following references are from `_cvtss_sh` with truncate mode.
+  // +inf
+  assert(bitwise_comparison_fp16(75514, 31744));
+  // -inf
+  assert(bitwise_comparison_fp16(-75514, 64512));
+  // +0
+  assert(bitwise_comparison_fp16(0.0, 0));
+  // -0
+  assert(bitwise_comparison_fp16(-0.0, 32768));
+  // nan
+  assert(bitwise_comparison_fp16(0.0 / 0.0, 32256));
+  assert(bitwise_comparison_fp16(-0.0 / 0.0, 32256));
+  // special nan
+  uint32_t special_nan = 0x7f800001;
+  assert(
+      bitwise_comparison_fp16(reinterpret_cast<float &>(special_nan), 32256));
+  special_nan = 0xff800001;
+  assert(
+      bitwise_comparison_fp16(reinterpret_cast<float &>(special_nan), 65024));
+  // subnormal
+  assert(bitwise_comparison_fp16(9.8E-45, 0));
+  assert(bitwise_comparison_fp16(-9.8E-45, 32768));
+  // overflow
+  assert(bitwise_comparison_fp16(half(55504) * 3, 31744));
+  assert(bitwise_comparison_fp16(half(-55504) * 3, 64512));
+  // underflow
+  assert(bitwise_comparison_fp16(half(8.1035e-05) / half(3), 453));
+
+  // Basic tests: fp16->fp32
+  // The following references are from `_cvtsh_ss`.
+  // +inf
+  const uint16_t pinf = 0x7a00;
+  assert(bitwise_comparison_fp32(reinterpret_cast<const half &>(pinf),
+                                 1195376640));
+  // -inf
+  const uint16_t ninf = 0xfa00;
+  assert(bitwise_comparison_fp32(reinterpret_cast<const half &>(ninf),
+                                 3342860288));
+  // +0
+  const uint16_t p0 = 0x0;
+  assert(bitwise_comparison_fp32(reinterpret_cast<const half &>(p0), 0));
+  // -0
+  const uint16_t n0 = 0x8000;
+  assert(
+      bitwise_comparison_fp32(reinterpret_cast<const half &>(n0), 2147483648));
+  // nan
+  const uint16_t nan16 = 0x7a03;
+  assert(bitwise_comparison_fp32(reinterpret_cast<const half &>(nan16),
+                                 1195401216));
+  // subnormal
+  const uint16_t subnormal = 0x0005;
+  assert(bitwise_comparison_fp32(reinterpret_cast<const half &>(subnormal),
+                                 882900992));
+
+  return 0;
+}
diff --git a/sycl/test/lit.cfg b/sycl/test/lit.cfg
index 6b8a469163cc9..137d518e40065 100644
--- a/sycl/test/lit.cfg
+++ b/sycl/test/lit.cfg
@@ -85,6 +85,7 @@ gpu_run_substitute = "echo"
 if getDeviceCount("gpu"):
     print("Found available GPU device")
     gpu_run_substitute = " env SYCL_DEVICE_TYPE=GPU "
+    config.available_features.add('gpu')
 config.substitutions.append( ('%GPU_RUN_PLACEHOLDER',  gpu_run_substitute) )
 
 acc_run_substitute = "echo"
diff --git a/sycl/test/sub_group/helper.hpp b/sycl/test/sub_group/helper.hpp
index 091822dfadaf9..9ee0db87c83f7 100644
--- a/sycl/test/sub_group/helper.hpp
+++ b/sycl/test/sub_group/helper.hpp
@@ -111,6 +111,14 @@ template <typename T> void exit_if_not_equal(T val, T ref, const char *name) {
   }
 }
 
+template <> void exit_if_not_equal(half val, half ref, const char *name) {
+  if (std::fabs((float)val - (float)ref) > 0.01) {
+    std::cout << "Unexpected result for " << name << ": " << (float)val
+              << " expected value: " << (float)ref << std::endl;
+    exit(1);
+  }
+}
+
 template <typename T, int N>
 void exit_if_not_equal_vec(vec<T, N> val, vec<T, N> ref, const char *name) {
   if (!utils<T, N>::cmp_vec(ref, val)) {
diff --git a/sycl/test/sub_group/load_store.cpp b/sycl/test/sub_group/load_store.cpp
index 19cb8ef10183c..3f1e4748173fb 100644
--- a/sycl/test/sub_group/load_store.cpp
+++ b/sycl/test/sub_group/load_store.cpp
@@ -11,11 +11,17 @@
 //
 //===----------------------------------------------------------------------===//
 
+// XFAIL: gpu
+
 #include "helper.hpp"
 #include <CL/sycl.hpp>
 template <typename T, int N> class sycl_subgr;
 
 using namespace cl::sycl;
+// TODO remove this workaround when integration header will support correct
+// half generation
+struct wa_half;
+typedef half aligned_half __attribute__((aligned(16)));
 
 template <typename T, int N> void check(queue &Queue) {
   const int G = 1024, L = 64;
@@ -30,10 +36,12 @@ template <typename T, int N> void check(queue &Queue) {
         acc[i] += 0.1; // Check that floating point types are not casted to int
       }
     }
+    using TT = typename std::conditional<std::is_same<T, aligned_half>::value,
+                                         wa_half, T>::type;
     Queue.submit([&](handler &cgh) {
       auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
-      cgh.parallel_for<sycl_subgr<T, N>>(NdRange, [=](nd_item<1> NdItem) {
+      cgh.parallel_for<sycl_subgr<TT, N>>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         if (SG.get_group_id().get(0) % N == 0) {
           size_t WGSGoffset =
@@ -97,10 +105,12 @@ template <typename T> void check(queue &Queue) {
       }
     }
 
+    using TT = typename std::conditional<std::is_same<T, aligned_half>::value,
+                                         wa_half, T>::type;
     Queue.submit([&](handler &cgh) {
       auto acc = syclbuf.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
-      cgh.parallel_for<sycl_subgr<T, 0>>(NdRange, [=](nd_item<1> NdItem) {
+      cgh.parallel_for<sycl_subgr<TT, 0>>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         if (NdItem.get_global_id(0) == 0)
           sgsizeacc[0] = SG.get_max_local_range()[0];
@@ -172,7 +182,6 @@ int main() {
     check<aligned_short, 4>(Queue);
     check<aligned_short, 8>(Queue);
     if (Queue.get_device().has_extension("cl_khr_fp16")) {
-      typedef half aligned_half __attribute__((aligned(16)));
       check<aligned_half>(Queue);
       check<aligned_half, 1>(Queue);
       check<aligned_half, 2>(Queue);
diff --git a/sycl/test/sub_group/shuffle.cpp b/sycl/test/sub_group/shuffle.cpp
index 7f246e2161d8c..66b74767710e9 100644
--- a/sycl/test/sub_group/shuffle.cpp
+++ b/sycl/test/sub_group/shuffle.cpp
@@ -16,6 +16,11 @@
 template <typename T, int N> class sycl_subgr;
 
 using namespace cl::sycl;
+
+// TODO remove this workaround when clang will support correct generation of
+// half typename in integration header
+struct wa_half;
+
 template <typename T, int N>
 void check(queue &Queue, size_t G = 240, size_t L = 60) {
   try {
@@ -28,6 +33,8 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
     buffer<vec<T, N>> buf_down(G);
     buffer<vec<T, N>> buf_xor(G);
     buffer<size_t> sgsizebuf(1);
+    using TT = typename std::conditional<std::is_same<T, half>::value, wa_half,
+                                         T>::type;
     Queue.submit([&](handler &cgh) {
       auto acc2 = buf2.template get_access<access::mode::read_write>(cgh);
       auto acc2_up = buf2_up.template get_access<access::mode::read_write>(cgh);
@@ -41,7 +48,7 @@ void check(queue &Queue, size_t G = 240, size_t L = 60) {
       auto acc_xor = buf_xor.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
 
-      cgh.parallel_for<sycl_subgr<T, N>>(NdRange, [=](nd_item<1> NdItem) {
+      cgh.parallel_for<sycl_subgr<TT, N>>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         uint32_t wggid = NdItem.get_global_id(0);
         uint32_t sgid = SG.get_group_id().get(0);
@@ -134,6 +141,8 @@ template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
     buffer<T> buf_down(G);
     buffer<T> buf_xor(G);
     buffer<size_t> sgsizebuf(1);
+    using TT = typename std::conditional<std::is_same<T, half>::value, wa_half,
+                                         T>::type;
     Queue.submit([&](handler &cgh) {
       auto acc2 = buf2.template get_access<access::mode::read_write>(cgh);
       auto acc2_up = buf2_up.template get_access<access::mode::read_write>(cgh);
@@ -147,7 +156,7 @@ template <typename T> void check(queue &Queue, size_t G = 240, size_t L = 60) {
       auto acc_xor = buf_xor.template get_access<access::mode::read_write>(cgh);
       auto sgsizeacc = sgsizebuf.get_access<access::mode::read_write>(cgh);
 
-      cgh.parallel_for<sycl_subgr<T, 0>>(NdRange, [=](nd_item<1> NdItem) {
+      cgh.parallel_for<sycl_subgr<TT, 0>>(NdRange, [=](nd_item<1> NdItem) {
         intel::sub_group SG = NdItem.get_sub_group();
         uint32_t wggid = NdItem.get_global_id(0);
         uint32_t sgid = SG.get_group_id().get(0);

From 1e6ab9236e4ee00c7c4182bd5b45ab6a49dc4c82 Mon Sep 17 00:00:00 2001
From: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
Date: Fri, 5 Apr 2019 13:55:29 +0300
Subject: [PATCH 2/2] [SYCL] Add support for _Float16 type mangling to
 OclCxxRewrite

Half type is defined on SYCL device as _Float16. This type is mangled by
clang as "DF16_". OclCxxRewrite doesn't support this mangling and skips
all OpenCL built-ins with _Float16 type. So these built-ins are not
translated properly to SPIRV without this change.

Signed-off-by: Mariya Podchishchaeva <mariya.podchishchaeva@intel.com>
---
 clang/lib/CodeGen/OclCxxRewrite/OclCxxDemangler.cpp       | 2 ++
 clang/lib/CodeGen/OclCxxRewrite/OclCxxMangleEncodings.inc | 2 +-
 2 files changed, 3 insertions(+), 1 deletion(-)

diff --git a/clang/lib/CodeGen/OclCxxRewrite/OclCxxDemangler.cpp b/clang/lib/CodeGen/OclCxxRewrite/OclCxxDemangler.cpp
index a4bde63c18928..6a61b1781d23d 100644
--- a/clang/lib/CodeGen/OclCxxRewrite/OclCxxDemangler.cpp
+++ b/clang/lib/CodeGen/OclCxxRewrite/OclCxxDemangler.cpp
@@ -353,6 +353,8 @@ static bool matchFixedBiTypeAndAdvance(const std::string &ParsedText,
       break;
     case 'n': MatchedType = DBT_NullPtr;
       break;
+    case 'F': MatchedType = DBT_Half; Pos += 3; // _Float16 -> DF16_
+      break;
     default:
       return false;
     }
diff --git a/clang/lib/CodeGen/OclCxxRewrite/OclCxxMangleEncodings.inc b/clang/lib/CodeGen/OclCxxRewrite/OclCxxMangleEncodings.inc
index 9badeb9be4bdf..021bb8c3a9f58 100644
--- a/clang/lib/CodeGen/OclCxxRewrite/OclCxxMangleEncodings.inc
+++ b/clang/lib/CodeGen/OclCxxRewrite/OclCxxMangleEncodings.inc
@@ -115,7 +115,7 @@ OCLCXX_MENC_BITYPE_FIXED(Float128R,    "De", "__fp128r")
                              ///< IEEE 754r decimal floating point (128 bits)
 OCLCXX_MENC_BITYPE_FIXED(Float32R,     "Df", "__fp32r")
                              ///< IEEE 754r decimal floating point (32 bits)
-OCLCXX_MENC_BITYPE_FIXED(Float16R,     "Dh", "half")
+OCLCXX_MENC_BITYPE_FIXED(Float16R,     "DF16_", "_Float16")
                              ///< IEEE 754r half-prec floating point (16 bits)
 OCLCXX_MENC_BITYPE_ALIAS(Half,         Float16R)
 OCLCXX_MENC_BITYPE_FIXED(Char32,       "Di", "char32_t")       ///< char32_t