[SYCL] Add opportunistic_group implementation

Signed-off-by: John Pennycook <[email protected]>

Author: Pennycook

sycl/include/sycl/ext/oneapi/experimental/opportunistic_group.hpp

@@ -0,0 +1,147 @@
+//==--- opportunistic_group.hpp --- SYCL extension for non-uniform groups --==//
+//
+// 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 <sycl/ext/oneapi/experimental/non_uniform_groups.hpp>
+#include <sycl/ext/oneapi/sub_group_mask.hpp>
+
+namespace sycl {
+__SYCL_INLINE_VER_NAMESPACE(_V1) {
+namespace ext::oneapi::experimental {
+
+class opportunistic_group;
+
+namespace this_kernel {
+inline opportunistic_group get_opportunistic_group();
+}
+
+class opportunistic_group {
+public:
+  using id_type = id<1>;
+  using range_type = range<1>;
+  using linear_id_type = uint32_t;
+  static constexpr int dimensions = 1;
+  static constexpr sycl::memory_scope fence_scope =
+      sycl::memory_scope::sub_group;
+
+  id_type get_group_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<id_type>(0);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  id_type get_local_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return detail::CallerPositionInMask(Mask);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  range_type get_group_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return 1;
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  range_type get_local_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return Mask.count();
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  linear_id_type get_group_linear_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_group_id()[0]);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  linear_id_type get_local_linear_id() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_local_id()[0]);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  linear_id_type get_group_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_group_range()[0]);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  linear_id_type get_local_linear_range() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return static_cast<linear_id_type>(get_local_range()[0]);
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+  bool leader() const {
+#ifdef __SYCL_DEVICE_ONLY__
+    return __spirv_SubgroupLocalInvocationId() == Mask.find_low();
+#else
+    throw runtime_error("Non-uniform groups are not supported on host device.",
+                        PI_ERROR_INVALID_DEVICE);
+#endif
+  }
+
+private:
+  sub_group_mask Mask;
+
+protected:
+  opportunistic_group(sub_group_mask m) : Mask(m) {}
+
+  friend opportunistic_group this_kernel::get_opportunistic_group();
+};
+
+namespace this_kernel {
+
+inline opportunistic_group get_opportunistic_group() {
+#ifdef __SYCL_DEVICE_ONLY__
+#if defined(__SPIR__)
+  // TODO: It may be wiser to call the intrinsic than rely on this_group()
+  sycl::sub_group sg = sycl::ext::oneapi::this_sub_group();
+  sub_group_mask mask = sycl::ext::oneapi::group_ballot(sg, true);
+  return opportunistic_group(mask);
+#elif defined(__NVPTX__)
+  // TODO: Construct from __activemask
+#endif
+#else
+  throw runtime_error("Non-uniform groups are not supported on host device.",
+                      PI_ERROR_INVALID_DEVICE);
+#endif
+}
+
+} // namespace this_kernel
+
+template <>
+struct is_user_constructed_group<opportunistic_group> : std::true_type {};
+
+} // namespace ext::oneapi::experimental
+} // __SYCL_INLINE_VER_NAMESPACE(_V1)
+} // namespace sycl

sycl/include/sycl/sycl.hpp

@@ -71,6 +71,7 @@
 #include <sycl/ext/oneapi/experimental/builtins.hpp>
 #include <sycl/ext/oneapi/experimental/cluster_group.hpp>
 #include <sycl/ext/oneapi/experimental/cuda/barrier.hpp>
+#include <sycl/ext/oneapi/experimental/opportunistic_group.hpp>
 #include <sycl/ext/oneapi/experimental/tangle_group.hpp>
 #include <sycl/ext/oneapi/filter_selector.hpp>
 #include <sycl/ext/oneapi/group_algorithm.hpp>