[SYCL] Add 64-bit type support to load/store sub-group functions (#1999)

Also, correct select_cl_scalar_integral_xxx type transformation naming typo.

Signed-off-by: Dmitri Mokhov <[email protected]>

Author: dnmokhov

sycl/include/CL/__spirv/spirv_ops.hpp

@@ -205,6 +205,15 @@ __SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
 __spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint32_t *Ptr,
                                 dataT Data) noexcept;
 
+template <typename dataT>
+__SYCL_CONVERGENT__ extern SYCL_EXTERNAL dataT __spirv_SubgroupBlockReadINTEL(
+    const __attribute__((opencl_global)) uint64_t *Ptr) noexcept;
+
+template <typename dataT>
+__SYCL_CONVERGENT__ extern SYCL_EXTERNAL void
+__spirv_SubgroupBlockWriteINTEL(__attribute__((opencl_global)) uint64_t *Ptr,
+                                dataT Data) noexcept;
+
 template <typename dataT>
 extern SYCL_EXTERNAL int32_t __spirv_ReadPipe(RPipeTy<dataT> Pipe, dataT *Data,
                                               int32_t Size,

sycl/include/CL/sycl/detail/generic_type_traits.hpp

@@ -367,12 +367,12 @@ using select_apply_cl_scalar_t =
 
 // Shortcuts for selecting scalar int/unsigned int/fp type.
 template <typename T>
-using select_cl_scalar_intergal_signed_t =
+using select_cl_scalar_integral_signed_t =
     select_apply_cl_scalar_t<T, sycl::cl_char, sycl::cl_short, sycl::cl_int,
                              sycl::cl_long>;
 
 template <typename T>
-using select_cl_scalar_intergal_unsigned_t =
+using select_cl_scalar_integral_unsigned_t =
     select_apply_cl_scalar_t<T, sycl::cl_uchar, sycl::cl_ushort, sycl::cl_uint,
                              sycl::cl_ulong>;
 
@@ -382,16 +382,16 @@ using select_cl_scalar_float_t =
                              sycl::cl_double>;
 
 template <typename T>
-using select_cl_scalar_intergal_t =
+using select_cl_scalar_integral_t =
     conditional_t<std::is_signed<T>::value,
-                  select_cl_scalar_intergal_signed_t<T>,
-                  select_cl_scalar_intergal_unsigned_t<T>>;
+                  select_cl_scalar_integral_signed_t<T>,
+                  select_cl_scalar_integral_unsigned_t<T>>;
 
 // select_cl_scalar_t picks corresponding cl_* type for input
 // scalar T or returns T if T is not scalar.
 template <typename T>
 using select_cl_scalar_t = conditional_t<
-    std::is_integral<T>::value, select_cl_scalar_intergal_t<T>,
+    std::is_integral<T>::value, select_cl_scalar_integral_t<T>,
     conditional_t<
         std::is_floating_point<T>::value, select_cl_scalar_float_t<T>,
         // half is a special case: it is implemented differently on host and

sycl/include/CL/sycl/intel/sub_group.hpp

@@ -55,11 +55,9 @@ __SYCL_SG_GENERATE_BODY_2ARG(shuffle_up, SubgroupShuffleUpINTEL)
 
 #undef __SYCL_SG_GENERATE_BODY_2ARG
 
-// Selects 8-bit, 16-bit or 32-bit type depending on size of T. If T doesn't
-// maps to mentioned types, then void is returned
+// Selects 8, 16, 32, or 64-bit type depending on size of scalar type T.
 template <typename T>
-using SelectBlockT =
-    select_apply_cl_scalar_t<T, uint8_t, uint16_t, uint32_t, void>;
+using SelectBlockT = select_cl_scalar_integral_unsigned_t<T>;
 
 template <typename T, access::address_space Space>
 using AcceptableForGlobalLoadStore =

sycl/test/sub_group/load_store.cpp

@@ -158,7 +158,8 @@ template <typename T> void check(queue &Queue) {
 int main() {
   queue Queue;
   if (!Queue.get_device().has_extension("cl_intel_subgroups") &&
-      !Queue.get_device().has_extension("cl_intel_subgroups_short")) {
+      !Queue.get_device().has_extension("cl_intel_subgroups_short") &&
+      !Queue.get_device().has_extension("cl_intel_subgroups_long")) {
     std::cout << "Skipping test\n";
     return 0;
   }
@@ -200,6 +201,26 @@ int main() {
       check<aligned_half, 8>(Queue);
     }
   }
+  if (Queue.get_device().has_extension("cl_intel_subgroups_long")) {
+    typedef long aligned_long __attribute__((aligned(16)));
+    check<aligned_long>(Queue);
+    check<aligned_long, 1>(Queue);
+    check<aligned_long, 2>(Queue);
+    check<aligned_long, 4>(Queue);
+    check<aligned_long, 8>(Queue);
+    typedef unsigned long aligned_ulong __attribute__((aligned(16)));
+    check<aligned_ulong>(Queue);
+    check<aligned_ulong, 1>(Queue);
+    check<aligned_ulong, 2>(Queue);
+    check<aligned_ulong, 4>(Queue);
+    check<aligned_ulong, 8>(Queue);
+    typedef double aligned_double __attribute__((aligned(16)));
+    check<aligned_double>(Queue);
+    check<aligned_double, 1>(Queue);
+    check<aligned_double, 2>(Queue);
+    check<aligned_double, 4>(Queue);
+    check<aligned_double, 8>(Queue);
+  }
   std::cout << "Test passed." << std::endl;
   return 0;
 }