[SYCL] Add support for half type

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;
     }

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

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"

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

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);
   }