[SYCL][COMPAT] Support for experimental masked shifts (#12972)

This PR adds an experimental approach to support
`select_from_sub_group`,
`shift_sub_group_left`, `shift_sub_group_right`, and
`permute_sub_group_by_xor` functions on SPIRV devices.
 
As these are based on still to be supported intrinsics, tests have been
left out.

Author: Alcpz

sycl/doc/syclcompat/README.md

@@ -1088,6 +1088,10 @@ functionality to `sycl::select_from_group`, `sycl::shift_group_left`,
 However, they provide an optional argument to represent the `logical_group` size
 (default 32).
 
+Experimental support for masked versions of `select_from_sub_group`,
+`shift_sub_group_left`, `shift_sub_group_right` and `permute_sub_group_by_xor` is
+provided only for SPIRV devices.
+
 ```c++
 namespace syclcompat {
 
@@ -1116,6 +1120,25 @@ template <typename ValueT>
 ValueT permute_sub_group_by_xor(sycl::sub_group g, ValueT x, unsigned int mask,
                            int logical_sub_group_size = 32);
 
+namespace experimental {
+
+template <typename ValueT>
+ValueT select_from_sub_group(unsigned int member_mask, sycl::sub_group g, ValueT x,
+                             int remote_local_id, int logical_sub_group_size = 32);
+
+template <typename ValueT>
+ValueT shift_sub_group_left(unsigned int member_mask, sycl::sub_group g, ValueT x,
+                            unsigned int delta, int logical_sub_group_size = 32);
+
+template <typename ValueT>
+ValueT shift_sub_group_right(unsigned int member_mask, sycl::sub_group g, ValueT x,
+                             unsigned int delta, int logical_sub_group_size = 32);
+
+template <typename ValueT>
+ValueT permute_sub_group_by_xor(unsigned int member_mask, sycql::sub_group g, ValueT x,
+                                unsigned int mask, int logical_sub_group_size = 32);
+
+} // namespace experimental
 } // namespace syclcompat
 ```
 

sycl/include/syclcompat/util.hpp

@@ -39,6 +39,25 @@
 #include <syclcompat/math.hpp>
 #include <syclcompat/memory.hpp>
 
+// TODO: Remove these function definitions once they exist in the DPC++ compiler
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
+template <typename T>
+__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT
+    __attribute__((noduplicate)) T
+    __spirv_GroupNonUniformShuffle(__spv::Scope::Flag, T, unsigned) noexcept;
+
+template <typename T>
+__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT
+    __attribute__((noduplicate)) T
+    __spirv_GroupNonUniformShuffleDown(__spv::Scope::Flag, T,
+                                       unsigned) noexcept;
+
+template <typename T>
+__SYCL_CONVERGENT__ extern SYCL_EXTERNAL __SYCL_EXPORT
+    __attribute__((noduplicate)) T
+    __spirv_GroupNonUniformShuffleUp(__spv::Scope::Flag, T, unsigned) noexcept;
+#endif
+
 namespace syclcompat {
 
 namespace detail {
@@ -269,6 +288,196 @@ T permute_sub_group_by_xor(sycl::sub_group g, T x, unsigned int mask,
                                      : id);
 }
 
+namespace experimental {
+/// Masked version of select_from_sub_group, which execute masked sub-group
+/// operation. The parameter member_mask indicating the work-items participating
+/// the call. Whether the n-th bit is set to 1 representing whether the
+/// work-item with id n is participating the call. All work-items named in
+/// member_mask must be executed with the same member_mask, or the result is
+/// undefined.
+/// \tparam T Input value type
+/// \param [in] member_mask Input mask
+/// \param [in] g Input sub_group
+/// \param [in] x Input value
+/// \param [in] remote_local_id Input source work item id
+/// \param [in] logical_sub_group_size Input logical sub_group size
+/// \returns The result
+template <typename T>
+T select_from_sub_group(unsigned int member_mask, sycl::sub_group g, T x,
+                        int remote_local_id, int logical_sub_group_size = 32) {
+  unsigned int start_index =
+      g.get_local_linear_id() / logical_sub_group_size * logical_sub_group_size;
+  unsigned logical_remote_id =
+      start_index + remote_local_id % logical_sub_group_size;
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
+#if defined(__SPIR__)
+  return __spirv_GroupNonUniformShuffle(__spv::Scope::Subgroup, x,
+                                        logical_remote_id);
+#else
+  // TODO: Check
+  // https://github.com/intel/llvm/pull/12972#issuecomment-2034218881 for the
+  // __NVPTX__ implementation
+  throw sycl::exception(sycl::errc::runtime,
+                        "[SYCLcompat] Masked version of select_from_sub_group "
+                        "only supports SPIR-V backends.");
+#endif // __SPIR__
+#else
+  (void)g;
+  (void)x;
+  (void)remote_local_id;
+  (void)logical_sub_group_size;
+  (void)member_mask;
+  throw sycl::exception(
+      sycl::errc::runtime,
+      "[SYCLcompat] Masked version of select_from_sub_group not "
+      "supported on host device and none intel compiler.");
+#endif // __SYCL_DEVICE_ONLY__ && __INTEL_LLVM_COMPILER
+}
+
+/// Masked version of shift_sub_group_left, which execute masked sub-group
+/// operation. The parameter member_mask indicating the work-items participating
+/// the call. Whether the n-th bit is set to 1 representing whether the
+/// work-item with id n is participating the call. All work-items named in
+/// member_mask must be executed with the same member_mask, or the result is
+/// undefined.
+/// \tparam T Input value type
+/// \param [in] member_mask Input mask
+/// \param [in] g Input sub_group
+/// \param [in] x Input value
+/// \param [in] delta Input delta
+/// \param [in] logical_sub_group_size Input logical sub_group size
+/// \returns The result
+template <typename T>
+T shift_sub_group_left(unsigned int member_mask, sycl::sub_group g, T x,
+                       unsigned int delta, int logical_sub_group_size = 32) {
+  unsigned int id = g.get_local_linear_id();
+  unsigned int end_index =
+      (id / logical_sub_group_size + 1) * logical_sub_group_size;
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
+#if defined(__SPIR__)
+  T result =
+      __spirv_GroupNonUniformShuffleDown(__spv::Scope::Subgroup, x, delta);
+  if ((id + delta) >= end_index) {
+    result = x;
+  }
+  return result;
+#else
+  // TODO: Check
+  // https://github.com/intel/llvm/pull/12972#issuecomment-2034218881 for the
+  // __NVPTX__ implementation
+  throw sycl::exception(sycl::errc::runtime,
+                        "[SYCLcompat] Masked version of shift_sub_group_left "
+                        "only supports SPIR-V backends.");
+#endif // __SPIR__
+#else
+  (void)g;
+  (void)x;
+  (void)delta;
+  (void)logical_sub_group_size;
+  (void)member_mask;
+  throw sycl::exception(
+      sycl::errc::runtime,
+      "[SYCLcompat] Masked version of select_from_sub_group not "
+      "supported on host device and none intel compiler.");
+#endif // __SYCL_DEVICE_ONLY__ && __INTEL_LLVM_COMPILER
+}
+
+/// Masked version of shift_sub_group_right, which execute masked sub-group
+/// operation. The parameter member_mask indicating the work-items participating
+/// the call. Whether the n-th bit is set to 1 representing whether the
+/// work-item with id n is participating the call. All work-items named in
+/// member_mask must be executed with the same member_mask, or the result is
+/// undefined.
+/// \tparam T Input value type
+/// \param [in] member_mask Input mask
+/// \param [in] g Input sub_group
+/// \param [in] x Input value
+/// \param [in] delta Input delta
+/// \param [in] logical_sub_group_size Input logical sub_group size
+/// \returns The result
+template <typename T>
+T shift_sub_group_right(unsigned int member_mask, sycl::sub_group g, T x,
+                        unsigned int delta, int logical_sub_group_size = 32) {
+  unsigned int id = g.get_local_linear_id();
+  unsigned int start_index =
+      id / logical_sub_group_size * logical_sub_group_size;
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
+#if defined(__SPIR__)
+  T result = __spirv_GroupNonUniformShuffleUp(__spv::Scope::Subgroup, x, delta);
+  if ((id - start_index) < delta) {
+    result = x;
+  }
+  return result;
+#else
+  // TODO: Check
+  // https://github.com/intel/llvm/pull/12972#issuecomment-2034218881 for the
+  // __NVPTX__ implementation
+  throw sycl::exception(sycl::errc::runtime,
+                        "Masked version of shift_sub_group_right "
+                        "only supports SPIR-V backends.");
+#endif // __SPIR__
+#else
+  (void)g;
+  (void)x;
+  (void)delta;
+  (void)logical_sub_group_size;
+  (void)member_mask;
+  throw sycl::exception(sycl::errc::runtime,
+                        "Masked version of select_from_sub_group not "
+                        "supported on host device and none intel compiler.");
+#endif // __SYCL_DEVICE_ONLY && __INTEL_LLVM_COMPILER
+}
+
+/// Masked version of permute_sub_group_by_xor, which execute masked sub-group
+/// operation. The parameter member_mask indicating the work-items participating
+/// the call. Whether the n-th bit is set to 1 representing whether the
+/// work-item with id n is participating the call. All work-items named in
+/// member_mask must be executed with the same member_mask, or the result is
+/// undefined.
+/// \tparam T Input value type
+/// \param [in] member_mask Input mask
+/// \param [in] g Input sub_group
+/// \param [in] x Input value
+/// \param [in] mask Input mask
+/// \param [in] logical_sub_group_size Input logical sub_group size
+/// \returns The result
+template <typename T>
+T permute_sub_group_by_xor(unsigned int member_mask, sycl::sub_group g, T x,
+                           unsigned int mask, int logical_sub_group_size = 32) {
+  unsigned int id = g.get_local_linear_id();
+  unsigned int start_index =
+      id / logical_sub_group_size * logical_sub_group_size;
+  unsigned int target_offset = (id % logical_sub_group_size) ^ mask;
+  unsigned logical_remote_id = (target_offset < logical_sub_group_size)
+                                   ? start_index + target_offset
+                                   : id;
+#if defined(__SYCL_DEVICE_ONLY__) && defined(__INTEL_LLVM_COMPILER)
+#if defined(__SPIR__)
+  return __spirv_GroupNonUniformShuffle(__spv::Scope::Subgroup, x,
+                                        logical_remote_id);
+#else
+  // TODO: Check
+  // https://github.com/intel/llvm/pull/12972#issuecomment-2034218881 for the
+  // __NVPTX__ implementation
+  throw sycl::exception(
+      sycl::errc::runtime,
+      "[SYCLcompat] Masked version of permute_sub_group_by_xor "
+      "only supports SPIR-V backends.");
+#endif // __SPIR__
+#else
+  (void)g;
+  (void)x;
+  (void)mask;
+  (void)logical_sub_group_size;
+  (void)member_mask;
+  throw sycl::exception(
+      sycl::errc::runtime,
+      "[SYCLcompat]Masked version of select_from_sub_group not "
+      "supported on host device and none intel compiler.");
+#endif // __SYCL_DEVICE_ONLY__ && __INTEL_LLVM_COMPILER
+}
+} // namespace experimental
+
 /// Inherited from the original SYCLomatic compatibility headers.
 /// @return compiler's SYCL version if defined, 202000 otherwise.
 inline int get_sycl_language_version() {