[SYCL] Perform SubgroupShuffle emulation for vectors of long long and half (#9102)

Author: maksimsab

sycl/include/sycl/detail/spirv.hpp

@@ -13,6 +13,7 @@
 #include <cstring>
 #include <sycl/detail/generic_type_traits.hpp>
 #include <sycl/detail/helpers.hpp>
+#include <sycl/detail/type_list.hpp>
 #include <sycl/detail/type_traits.hpp>
 #include <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
 #include <sycl/id.hpp>
@@ -696,14 +697,19 @@ AtomicMax(multi_ptr<T, AddressSpace, IsDecorated> MPtr, memory_scope Scope,
 // - The Intel SPIR-V extension natively supports all arithmetic types.
 //   However, OpenCL extension natively supports float vectors,
 //   integer vectors, half scalar and double scalar.
-//   For double vectors we perform emulation with scalar version.
+//   For double, long, long long, unsigned long, unsigned long long
+//   and half vectors we perform emulation with scalar version.
 // - The CUDA shfl intrinsics do not support vectors, and we use the _i32
 //   variants for all scalar types
 #ifndef __NVPTX__
 
+using ProhibitedTypesForShuffleEmulation =
+    type_list<double, long, long long, unsigned long, unsigned long long, half>;
+
 template <typename T>
 struct TypeIsProhibitedForShuffleEmulation
-    : std::bool_constant<std::is_same_v<vector_element_t<T>, double>> {};
+    : std::bool_constant<is_contained<
+          vector_element_t<T>, ProhibitedTypesForShuffleEmulation>::value> {};
 
 template <typename T>
 struct VecTypeIsProhibitedForShuffleEmulation
@@ -790,6 +796,12 @@ EnableIfBitcastShuffle<T> SubgroupShuffle(T x, id<1> local_id);
 template <typename T>
 EnableIfBitcastShuffle<T> SubgroupShuffleXor(T x, id<1> local_id);
 
+template <typename T>
+EnableIfBitcastShuffle<T> SubgroupShuffleDown(T x, uint32_t delta);
+
+template <typename T>
+EnableIfBitcastShuffle<T> SubgroupShuffleUp(T x, uint32_t delta);
+
 template <typename T>
 EnableIfGenericShuffle<T> SubgroupShuffle(T x, id<1> local_id);
 

sycl/test-e2e/SubGroup/shuffle.cpp

@@ -33,8 +33,28 @@ int main() {
   check<unsigned int, 8>(Queue);
   check<unsigned int, 16>(Queue);
   check<long>(Queue);
+  check<long, 2>(Queue);
+  check<long, 4>(Queue);
+  check<long, 8>(Queue);
+  check<long, 16>(Queue);
   check<unsigned long>(Queue);
+  check<unsigned long, 2>(Queue);
+  check<unsigned long, 4>(Queue);
+  check<unsigned long, 8>(Queue);
+  check<unsigned long, 16>(Queue);
   check<float>(Queue);
+  check<float, 2>(Queue);
+  check<float, 4>(Queue);
+  check<float, 8>(Queue);
+  check<float, 16>(Queue);
+
+  // Check long long and unsigned long long because they differ from
+  // long and unsigned long according to C++ rules even if they have the same
+  // size at some system.
+  check<long long>(Queue);
+  check<long long, 16>(Queue);
+  check<unsigned long long>(Queue);
+  check<unsigned long long, 16>(Queue);
   std::cout << "Test passed." << std::endl;
   return 0;
 }

sycl/test-e2e/SubGroup/shuffle_fp16.cpp

@@ -24,6 +24,10 @@ int main() {
   queue Queue;
   if (Queue.get_device().has(sycl::aspect::fp16)) {
     check<half>(Queue);
+    check<half, 2>(Queue);
+    check<half, 4>(Queue);
+    check<half, 8>(Queue);
+    check<half, 16>(Queue);
     std::cout << "Test passed." << std::endl;
   } else {
     std::cout << "Test skipped because device doesn't support aspect::fp16"