[SYCL] Add half constexpr default constructor (#4518)

Add default constexpr constructor for half and host half classes, and also for half_vec struct.
Due to bitselect built-in uses union with half specialization, but object of a class without trivial constructor cannot be a member of union, the builtins_helper struct was added. This struct allow to use host half class instead of main half class in the builtin.

Signed-off-by: mdimakov [email protected]

Author: maximdimakov

sycl/include/CL/sycl/half_type.hpp

@@ -37,6 +37,11 @@ __SYCL_INLINE_NAMESPACE(cl) {
 namespace sycl {
 namespace detail {
 
+// bitselect builtin uses union with half specialization, but half has no
+// trivial constructor. This struct allow to use host half class instead of main
+// half class in the bitselect builtin.
+template <typename T> struct builtins_helper;
+
 inline __SYCL_CONSTEXPR_HALF uint16_t float2Half(const float &Val) {
   const uint32_t Bits = sycl::bit_cast<uint32_t>(Val);
 
@@ -189,7 +194,7 @@ class __SYCL_EXPORT half {
 // The main host half class
 class __SYCL_EXPORT half_v2 {
 public:
-  half_v2() = default;
+  __SYCL_CONSTEXPR_HALF half_v2() : Buf(float2Half(0.0f)) {}
   constexpr half_v2(const half_v2 &) = default;
   constexpr half_v2(half_v2 &&) = default;
 
@@ -255,6 +260,8 @@ class __SYCL_EXPORT half_v2 {
   // Initialize underlying data
   constexpr explicit half_v2(uint16_t x) : Buf(x) {}
 
+  template <typename T> friend struct cl::sycl::detail::builtins_helper;
+
 private:
   uint16_t Buf;
 };
@@ -299,7 +306,12 @@ using BIsRepresentationT = half;
 // as a kernel argument which is expected to be floating point number.
 template <int NumElements> struct half_vec {
   alignas(detail::vector_alignment<StorageT, NumElements>::value)
-      std::array<StorageT, NumElements> s;
+      StorageT s[NumElements];
+
+  __SYCL_CONSTEXPR_HALF half_vec() {
+    for (int i = 0; i < NumElements; i++)
+      s[i] = StorageT(0.0f);
+  }
 };
 
   using Vec2StorageT = half_vec<2>;
@@ -311,7 +323,7 @@ template <int NumElements> struct half_vec {
 
 class half {
 public:
-  half() = default;
+  __SYCL_CONSTEXPR_HALF half() : Data(0.0f){};
   constexpr half(const half &) = default;
   constexpr half(half &&) = default;
 
@@ -383,6 +395,8 @@ class half {
   }
 
   template <typename Key> friend struct std::hash;
+  template <typename T> friend struct cl::sycl::detail::builtins_helper;
+
 private:
   StorageT Data;
 };

sycl/source/detail/builtins_relational.cpp

@@ -18,6 +18,21 @@ namespace s = cl::sycl;
 namespace d = s::detail;
 
 __SYCL_INLINE_NAMESPACE(cl) {
+namespace sycl {
+namespace detail {
+
+// The declaration of this struct is in the CL/sycl/half_type.hpp
+template <typename T> struct builtins_helper {
+  using RetType = T;
+  static constexpr RetType get(T value) { return value; }
+};
+
+template <> struct builtins_helper<s::cl_half> {
+  using RetType = uint16_t;
+  static constexpr RetType get(s::cl_half value) { return value.Data.Buf; }
+};
+} // namespace detail
+} // namespace sycl
 namespace __host_std {
 namespace {
 
@@ -103,19 +118,20 @@ template <> union databitset<s::cl_double> {
 template <> union databitset<s::cl_half> {
   static_assert(sizeof(s::cl_short) == sizeof(s::cl_half),
                 "size of cl_half is not equal to 16 bits(cl_short).");
-  s::cl_half f;
+  uint16_t f;
   s::cl_short i;
 };
 
 template <typename T>
 typename sycl::detail::enable_if_t<d::is_sgenfloat<T>::value,
                                    T> inline __bitselect(T a, T b, T c) {
+  d::builtins_helper<T> helper;
   databitset<T> ba;
-  ba.f = a;
+  ba.f = helper.get(a);
   databitset<T> bb;
-  bb.f = b;
+  bb.f = helper.get(b);
   databitset<T> bc;
-  bc.f = c;
+  bc.f = helper.get(c);
   databitset<T> br;
   br.f = 0;
   br.i = ((ba.i & ~bc.i) | (bb.i & bc.i));

sycl/test/abi/sycl_symbols_windows.dump

@@ -340,6 +340,7 @@
 ??0half_v2@host_half_impl@detail@sycl@cl@@QEAA@AEBM@Z
 ??0half_v2@host_half_impl@detail@sycl@cl@@QEAA@AEBV01234@@Z
 ??0half_v2@host_half_impl@detail@sycl@cl@@QEAA@G@Z
+??0half_v2@host_half_impl@detail@sycl@cl@@QEAA@XZ
 ??0handler@sycl@cl@@AEAA@V?$shared_ptr@Vqueue_impl@detail@sycl@cl@@@std@@_N@Z
 ??0host_selector@sycl@cl@@QEAA@$$QEAV012@@Z
 ??0host_selector@sycl@cl@@QEAA@AEBV012@@Z

sycl/test/basic_tests/types.cpp

@@ -90,6 +90,9 @@ template <> inline void checkSizeForFloatingPoint<s::half, sizeof(int16_t)>() {
 }
 
 int main() {
+  // Test for half constexpr default constructors
+  constexpr sycl::specialization_id<sycl::vec<sycl::half, 2>> id(1.0);
+  constexpr sycl::marray<sycl::half, 2> MH(3);
   // Check the size and alignment of the SYCL vectors.
   checkVectors();