From c8d8e15e026aabfcbc65b553bd0b23bc896cdc24 Mon Sep 17 00:00:00 2001 From: Haroon Rasheed Date: Mon, 6 Feb 2023 20:14:03 +0500 Subject: [PATCH 1/6] [SYCL][FIX] Fixed method definition of sub_group_mask::group_ballot --- sycl/include/sycl/ext/oneapi/sub_group_mask.hpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp index 16a10b2304312..e7f44f5d71707 100644 --- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp +++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp @@ -267,7 +267,7 @@ struct sub_group_mask { template detail::enable_if_t, sub_group>::value, sub_group_mask> -group_ballot(Group g, bool predicate) { +group_ballot(Group g, bool predicate = true) { (void)g; #ifdef __SYCL_DEVICE_ONLY__ auto res = __spirv_GroupNonUniformBallot( From 28a604488d53f0f0ceed2c24beb72cb54ac879e1 Mon Sep 17 00:00:00 2001 From: Haroon Rasheed Date: Tue, 7 Feb 2023 19:11:50 +0500 Subject: [PATCH 2/6] [SYCL][FIX] Fixed group_ballot default argument for predicate --- sycl/include/sycl/ext/oneapi/sub_group_mask.hpp | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp index e7f44f5d71707..a63fe2eb1f3db 100644 --- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp +++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp @@ -29,6 +29,11 @@ namespace ext::oneapi { #define BITS_TYPE uint32_t #endif +// defining `group_ballot` here to make predicate default `true` +template +detail::enable_if_t, 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; @@ -267,7 +272,7 @@ struct sub_group_mask { template detail::enable_if_t, sub_group>::value, sub_group_mask> -group_ballot(Group g, bool predicate = true) { +group_ballot(Group g, bool predicate) { (void)g; #ifdef __SYCL_DEVICE_ONLY__ auto res = __spirv_GroupNonUniformBallot( From 1eadcf9eaf35ebe35bcdadd5dda18325e92c570b Mon Sep 17 00:00:00 2001 From: Haroon Rasheed Date: Tue, 7 Feb 2023 19:21:27 +0500 Subject: [PATCH 3/6] [SYCL][FIX] Fixed linting issue for group_ballot definition --- sycl/include/sycl/ext/oneapi/sub_group_mask.hpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp index a63fe2eb1f3db..e67f1a3c0e5b9 100644 --- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp +++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp @@ -31,7 +31,8 @@ namespace ext::oneapi { // defining `group_ballot` here to make predicate default `true` template -detail::enable_if_t, sub_group>::value, sub_group_mask> +detail::enable_if_t, sub_group>::value, + sub_group_mask> group_ballot(Group g, bool predicate = true); struct sub_group_mask { From f02b3849632795419f1baadce73a81bee7f58c8e Mon Sep 17 00:00:00 2001 From: Haroon Rasheed Date: Wed, 8 Feb 2023 19:41:12 +0500 Subject: [PATCH 4/6] [SYCL][FIX] Forward declared sub_group_mask for group_ballot declaration --- sycl/include/sycl/ext/oneapi/sub_group_mask.hpp | 2 ++ 1 file changed, 2 insertions(+) diff --git a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp index e67f1a3c0e5b9..e268f20982c98 100644 --- a/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp +++ b/sycl/include/sycl/ext/oneapi/sub_group_mask.hpp @@ -30,6 +30,8 @@ namespace ext::oneapi { #endif // defining `group_ballot` here to make predicate default `true` +// need to forward declare sub_group_mask first +struct sub_group_mask; template detail::enable_if_t, sub_group>::value, sub_group_mask> From a329e4f3ae5e7969b9401a23d9c3b67f67b43654 Mon Sep 17 00:00:00 2001 From: Haroon Rasheed Date: Wed, 8 Feb 2023 20:56:49 +0500 Subject: [PATCH 5/6] [SYCL][FIX] Test for default predicate in group_ballot --- sycl/test/check_device_code/group_ballot.cpp | 10 ++++++++++ 1 file changed, 10 insertions(+) create mode 100644 sycl/test/check_device_code/group_ballot.cpp diff --git a/sycl/test/check_device_code/group_ballot.cpp b/sycl/test/check_device_code/group_ballot.cpp new file mode 100644 index 0000000000000..a5ecf520dc53f --- /dev/null +++ b/sycl/test/check_device_code/group_ballot.cpp @@ -0,0 +1,10 @@ +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s + +#include + +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()); + }); +} From 8d4e7699b7e397aa15cde27a8b3443bc4a8e0b19 Mon Sep 17 00:00:00 2001 From: Alexey Sachkov Date: Wed, 8 Mar 2023 10:30:33 +0100 Subject: [PATCH 6/6] Update sycl/test/check_device_code/group_ballot.cpp --- sycl/test/check_device_code/group_ballot.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/test/check_device_code/group_ballot.cpp b/sycl/test/check_device_code/group_ballot.cpp index a5ecf520dc53f..58f700ea15d77 100644 --- a/sycl/test/check_device_code/group_ballot.cpp +++ b/sycl/test/check_device_code/group_ballot.cpp @@ -1,4 +1,5 @@ -// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify -Xclang -verify-ignore-unexpected=note,warning %s +// RUN: %clangxx -fsycl -fsycl-device-only -fsyntax-only -Xclang -verify %s +// expected-no-diagnostics #include