From 7c6305977d39a815abac600f07c5fb814d4c8e11 Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Wed, 1 Sep 2021 15:05:47 +0300 Subject: [PATCH 1/4] [SYCL] Test for Group Mask feature --- .github/CODEOWNERS | 1 + SYCL/GroupMask/Basic.cpp | 59 ++++++++++++++++++++++++++++++++++++++++ 2 files changed, 60 insertions(+) create mode 100644 SYCL/GroupMask/Basic.cpp diff --git a/.github/CODEOWNERS b/.github/CODEOWNERS index f8847ead4e..64f97b80b5 100644 --- a/.github/CODEOWNERS +++ b/.github/CODEOWNERS @@ -38,6 +38,7 @@ SYCL/Functor @AlexeySachkov # Group algorithms SYCL/GroupAlgorithm @Pennycook @AlexeySachkov SYCL/SubGroup @Pennycook @AlexeySachkov +SYCL/GroupMask @Pennycook @vladimilaz # Group local memory SYCL/GroupLocalMemory @sergey-semenov @Pennycook diff --git a/SYCL/GroupMask/Basic.cpp b/SYCL/GroupMask/Basic.cpp new file mode 100644 index 0000000000..3d5b56495a --- /dev/null +++ b/SYCL/GroupMask/Basic.cpp @@ -0,0 +1,59 @@ +// RUN: %clangxx -fsycl -fsycl-targets=%sycl_triple %s -o %t.out +// REQUIRES: gpu +// UNSUPPORTED: cuda, hip +// GroupNonUniformBallot capability is supported on Intel GPU only +// RUN: %GPU_RUN_PLACEHOLDER %t.out + +//==---------- barrier.cpp - SYCL sub_group barrier test -------*- C++ -*---==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +#include +#include +#include +using namespace sycl; +constexpr int global_size = 128; +constexpr int local_size = 64; +int main() { + queue Queue; + + try { + nd_range<1> NdRange(global_size, local_size); + std::vector> data(global_size); + std::iota(data.begin(), data.end(), sizeof(int)); + { + buffer andbuf(data.data(), range<1>(global_size)); + buffer orbuf(data.data(), range<1>(global_size)); + buffer xorbuf(data.data(), range<1>(global_size)); + + buffer sgsizebuf(1); + Queue.submit([&](handler &cgh) { + auto andacc = andbuf.get_access(cgh); + auto oracc = orbuf.get_access(cgh); + auto xoracc = xorbuf.get_access(cgh); + + cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { + size_t GID = NdItem.get_global_linear_id(); + size_t LGID = NdItem.get_group(0); + auto gmask_gid = + ext::oneapi::group_ballot(NdItem.get_group(), GID % 2); + auto gmask_lgid = + ext::oneapi::group_ballot(NdItem.get_group(), LGID % 2); + andacc[GID] = (gmask_gid & gmask_lgid).extract_bits(); + oracc[GID] = (gmask_gid | gmask_lgid).extract_bits(); + xoracc[GID] = (gmask_gid ^ gmask_lgid).extract_bits(); + }); + }); + } + } catch (exception e) { + std::cout << "SYCL exception caught: " << e.what(); + exit(1); + } + + std::cout << "Test passed." << std::endl; + return 0; +} From 7c50f62c2fa2ee65276b08b3524678c4c37928eb Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Mon, 6 Sep 2021 17:07:01 +0300 Subject: [PATCH 2/4] Extend test and bugfix --- SYCL/GroupMask/Basic.cpp | 61 ++++++++++++++++++++++++++++------------ 1 file changed, 43 insertions(+), 18 deletions(-) diff --git a/SYCL/GroupMask/Basic.cpp b/SYCL/GroupMask/Basic.cpp index 3d5b56495a..9480aada94 100644 --- a/SYCL/GroupMask/Basic.cpp +++ b/SYCL/GroupMask/Basic.cpp @@ -17,38 +17,63 @@ #include using namespace sycl; constexpr int global_size = 128; -constexpr int local_size = 64; +constexpr int local_size = 32; int main() { queue Queue; try { nd_range<1> NdRange(global_size, local_size); - std::vector> data(global_size); - std::iota(data.begin(), data.end(), sizeof(int)); + int Res = 0; { - buffer andbuf(data.data(), range<1>(global_size)); - buffer orbuf(data.data(), range<1>(global_size)); - buffer xorbuf(data.data(), range<1>(global_size)); + buffer resbuf(&Res, range<1>(1)); - buffer sgsizebuf(1); Queue.submit([&](handler &cgh) { - auto andacc = andbuf.get_access(cgh); - auto oracc = orbuf.get_access(cgh); - auto xoracc = xorbuf.get_access(cgh); + auto resacc = resbuf.get_access(cgh); cgh.parallel_for(NdRange, [=](nd_item<1> NdItem) { size_t GID = NdItem.get_global_linear_id(); - size_t LGID = NdItem.get_group(0); - auto gmask_gid = - ext::oneapi::group_ballot(NdItem.get_group(), GID % 2); - auto gmask_lgid = - ext::oneapi::group_ballot(NdItem.get_group(), LGID % 2); - andacc[GID] = (gmask_gid & gmask_lgid).extract_bits(); - oracc[GID] = (gmask_gid | gmask_lgid).extract_bits(); - xoracc[GID] = (gmask_gid ^ gmask_lgid).extract_bits(); + auto SG = NdItem.get_sub_group(); + auto gmask_gid2 = + ext::oneapi::group_ballot(NdItem.get_sub_group(), GID % 2); + auto gmask_gid3 = + ext::oneapi::group_ballot(NdItem.get_sub_group(), GID % 3); + NdItem.barrier(); + + if (!GID) { + int res = 0; + + for (size_t i = 0; i < SG.get_max_local_range()[0]; i++) { + /* res |= !((gmask_gid2 | gmask_gid3)[i] == (i % 2 || + i % 3)) << 1; res |= !((gmask_gid2 & gmask_gid3)[i] == (i % 2 + && i % 3)) << 2; res |= !((gmask_gid2 ^ gmask_gid3)[i] == + ((bool)(i % 2) ^ (bool)(i % 3))) + << 3;*/ + } + gmask_gid2 <<= 32; + res |= (gmask_gid2.extract_bits()[2] != 0xaaaaaaaa) << 4; + res |= ((gmask_gid2 >> 8).extract_bits()[3] != 0xaa000000) << 5; + res |= ((gmask_gid3 >> 8).extract_bits()[3] != 0xb6db6d) << 6; + res |= (!gmask_gid2[32] && gmask_gid2[31]) << 7; + gmask_gid3[0] = gmask_gid3[3] = gmask_gid3[6] = true; + res |= (gmask_gid3.extract_bits()[3] != 0xb6db6dff) << 7; + gmask_gid3.reset(); + res |= !(gmask_gid3.none() && gmask_gid2.any() && !gmask_gid2.all()) + << 8; + gmask_gid2.set(); + res |= !(gmask_gid3.none() && gmask_gid2.any() && gmask_gid2.all()) + << 9; + gmask_gid3.flip(); + res |= (gmask_gid3 != gmask_gid2) << 10; + resacc[0] = res; + } }); }); } + if (Res) { + std::cout << "Unexpected result for group_mask operation: " << Res + << std::endl; + exit(1); + } } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); exit(1); From 09343a9dbc2fac780fd7b4956cb8eacf15a7229f Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Tue, 7 Sep 2021 07:45:31 +0300 Subject: [PATCH 3/4] Apply comments --- SYCL/GroupMask/Basic.cpp | 14 ++++++-------- 1 file changed, 6 insertions(+), 8 deletions(-) diff --git a/SYCL/GroupMask/Basic.cpp b/SYCL/GroupMask/Basic.cpp index 9480aada94..a9ad66561b 100644 --- a/SYCL/GroupMask/Basic.cpp +++ b/SYCL/GroupMask/Basic.cpp @@ -4,7 +4,7 @@ // GroupNonUniformBallot capability is supported on Intel GPU only // RUN: %GPU_RUN_PLACEHOLDER %t.out -//==---------- barrier.cpp - SYCL sub_group barrier test -------*- C++ -*---==// +//==---------- Basic.cpp - SYCL Group Mask basic test ----------*- C++ -*---==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -13,8 +13,6 @@ //===----------------------------------------------------------------------===// #include -#include -#include using namespace sycl; constexpr int global_size = 128; constexpr int local_size = 32; @@ -43,11 +41,11 @@ int main() { int res = 0; for (size_t i = 0; i < SG.get_max_local_range()[0]; i++) { - /* res |= !((gmask_gid2 | gmask_gid3)[i] == (i % 2 || - i % 3)) << 1; res |= !((gmask_gid2 & gmask_gid3)[i] == (i % 2 - && i % 3)) << 2; res |= !((gmask_gid2 ^ gmask_gid3)[i] == - ((bool)(i % 2) ^ (bool)(i % 3))) - << 3;*/ + res |= !((gmask_gid2 | gmask_gid3)[i] == (i % 2 || i % 3)) << 1; + res |= !((gmask_gid2 & gmask_gid3)[i] == (i % 2 && i % 3)) << 2; + res |= !((gmask_gid2 ^ gmask_gid3)[i] == + ((bool)(i % 2) ^ (bool)(i % 3))) + << 3; } gmask_gid2 <<= 32; res |= (gmask_gid2.extract_bits()[2] != 0xaaaaaaaa) << 4; From efcc6ea746da656882f7b5e2a20babd9289f52fa Mon Sep 17 00:00:00 2001 From: Vladimir Lazarev Date: Tue, 7 Sep 2021 08:55:20 +0300 Subject: [PATCH 4/4] Enable test only when extension is present --- SYCL/GroupMask/Basic.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/SYCL/GroupMask/Basic.cpp b/SYCL/GroupMask/Basic.cpp index a9ad66561b..6fe6956e75 100644 --- a/SYCL/GroupMask/Basic.cpp +++ b/SYCL/GroupMask/Basic.cpp @@ -17,6 +17,7 @@ using namespace sycl; constexpr int global_size = 128; constexpr int local_size = 32; int main() { +#ifdef SYCL_EXT_ONEAPI_GROUP_MASK queue Queue; try { @@ -78,5 +79,8 @@ int main() { } std::cout << "Test passed." << std::endl; +#else + std::cout << "Test skipped due to missing extension." << std::endl; +#endif return 0; }