[SYCL] Fix method definition of sub_group_mask::group_ballot (intel#8212

) Closes intel#8201 provided default value for `predicate` argument in `sub_group_mask::group_ballot` definition --------- Co-authored-by: Alexey Sachkov <alexey.sachkov@intel.com>
YuriPlyakhin · Mar 28, 2023 · e4024d7 · e4024d7
1 parent 4bf87d2
commit e4024d7
Show file tree

Hide file tree

Showing 2 changed files with 19 additions and 0 deletions.
diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp
@@ -29,6 +29,14 @@ namespace ext::oneapi {
 #define BITS_TYPE uint32_t
 #endif
 
+// defining `group_ballot` here to make predicate default `true`
+// need to forward declare sub_group_mask first
+struct sub_group_mask;
+template <typename Group>
+detail::enable_if_t<std::is_same<std::decay_t<Group>, sub_group>::value,
+                    sub_group_mask>
+group_ballot(Group g, bool predicate = true);
+
 struct sub_group_mask {
   friend class detail::Builder;
   using BitsType = BITS_TYPE;

diff --git a/sycl/test/check_device_code/group_ballot.cpp b/sycl/test/check_device_code/group_ballot.cpp
@@ -0,0 +1,11 @@
+// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s
+// expected-no-diagnostics
+
+#include <sycl/sycl.hpp>
+
+int main() {
+  sycl::queue Q;
+  Q.parallel_for(sycl::nd_range<1>{32, 32}, [=](sycl::nd_item<1> item) {
+    auto Mask = sycl::ext::oneapi::group_ballot(item.get_sub_group());
+  });
+}