From b4c399f5901cf1f0697435bb7a89adccad07460d Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Mon, 24 Jun 2024 12:36:51 -0700 Subject: [PATCH 1/9] Remove `core_sg_supported` function --- sycl/test-e2e/Basic/linear-sub_group.cpp | 5 -- sycl/test-e2e/SubGroup/attributes.cpp | 15 +---- sycl/test-e2e/SubGroup/helper.hpp | 21 ------- sycl/test-e2e/SubGroup/info.cpp | 58 ++++++------------- sycl/test-e2e/SubGroup/reduce.cpp | 4 -- sycl/test-e2e/SubGroup/reduce_fp16.cpp | 4 -- sycl/test-e2e/SubGroup/reduce_fp64.cpp | 4 -- sycl/test-e2e/SubGroup/reduce_spirv13.cpp | 4 -- .../test-e2e/SubGroup/reduce_spirv13_fp16.cpp | 4 -- .../test-e2e/SubGroup/reduce_spirv13_fp64.cpp | 4 -- sycl/test-e2e/SubGroup/scan.cpp | 4 -- sycl/test-e2e/SubGroup/scan_fp16.cpp | 4 -- sycl/test-e2e/SubGroup/scan_fp64.cpp | 4 -- sycl/test-e2e/SubGroup/scan_spirv13.cpp | 4 -- sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp | 4 -- sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp | 4 -- sycl/test-e2e/SubGroup/vote.cpp | 4 -- 17 files changed, 19 insertions(+), 132 deletions(-) diff --git a/sycl/test-e2e/Basic/linear-sub_group.cpp b/sycl/test-e2e/Basic/linear-sub_group.cpp index df99664c8f9b9..a7b5f8c720732 100644 --- a/sycl/test-e2e/Basic/linear-sub_group.cpp +++ b/sycl/test-e2e/Basic/linear-sub_group.cpp @@ -9,7 +9,6 @@ // //===----------------------------------------------------------------------===// -#include "../SubGroup/helper.hpp" #include #include #include @@ -20,10 +19,6 @@ using namespace sycl; int main(int argc, char *argv[]) { queue q; - if (!core_sg_supported(q.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } // Fill output array with sub-group IDs const uint32_t outer = 2; diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index 44de853e07f91..1107e0554626f 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -13,7 +13,7 @@ #define KERNEL_FUNCTOR_WITH_SIZE(SIZE) \ class KernelFunctor##SIZE { \ public: \ - [[intel::reqd_sub_group_size(SIZE)]] void \ + [[sycl::reqd_sub_group_size(SIZE)]] void \ operator()(sycl::nd_item<1> Item) const { \ const auto GID = Item.get_global_id(); \ } \ @@ -49,19 +49,6 @@ int main() { queue Queue; device Device = Queue.get_device(); - // According to specification, this kernel query requires `cl_khr_subgroups` - // or `cl_intel_subgroups`, and also `cl_intel_required_subgroup_size` - auto Vec = Device.get_info(); - if (std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") == - std::end(Vec) && - std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") == - std::end(Vec) || - std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") == - std::end(Vec)) { - std::cout << "Skipping test\n"; - return 0; - } - try { const auto SGSizes = Device.get_info(); diff --git a/sycl/test-e2e/SubGroup/helper.hpp b/sycl/test-e2e/SubGroup/helper.hpp index dfb47988d85e0..2a88bb0f5652e 100644 --- a/sycl/test-e2e/SubGroup/helper.hpp +++ b/sycl/test-e2e/SubGroup/helper.hpp @@ -164,24 +164,3 @@ void exit_if_not_equal_vec(vec val, vec ref, const char *name) { exit(1); } } - -bool core_sg_supported(const device &Device) { - auto Vec = Device.get_info(); - if (std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") != std::end(Vec)) - return true; - - if (std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") != std::end(Vec)) - return true; - - if (Device.get_backend() == sycl::backend::opencl) { - // Extract the numerical version from the version string, OpenCL version - // string have the format "OpenCL . ". - std::string ver = Device.get_info().substr(7, 3); - - // cl_khr_subgroups was core in OpenCL 2.1 and 2.2, but went back to - // optional in 3.0 - return ver >= "2.1" && ver < "3.0"; - } - - return false; -} diff --git a/sycl/test-e2e/SubGroup/info.cpp b/sycl/test-e2e/SubGroup/info.cpp index 83e9fdd5a64bf..e6b4ad65b6c96 100644 --- a/sycl/test-e2e/SubGroup/info.cpp +++ b/sycl/test-e2e/SubGroup/info.cpp @@ -17,12 +17,6 @@ int main() { queue Queue; device Device = Queue.get_device(); - /* Basic sub-group functionality is supported as part of cl_khr_subgroups - * extension or as core OpenCL 2.1 feature. */ - if (!core_sg_supported(Device)) { - std::cout << "Skipping test\n"; - return 0; - } /* Check info::device parameters. */ Device.get_info(); Device.get_info(); @@ -49,30 +43,24 @@ int main() { }); uint32_t Res = 0; - /* sub_group_sizes can be queried only if cl_intel_required_subgroup_size - * extension is supported by device*/ - auto Vec = Device.get_info(); - if (std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") != - std::end(Vec)) { - auto sg_sizes = Device.get_info(); + auto sg_sizes = Device.get_info(); + + // Max sub-group size for a particular kernel might not be the max + // supported size on the device in general. Can only check that it is + // contained in list of valid sizes. + Res = Kernel.get_info( + Device); + bool Expected = + std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end(); + exit_if_not_equal(Expected, true, "max_sub_group_size"); - // Max sub-group size for a particular kernel might not be the max - // supported size on the device in general. Can only check that it is - // contained in list of valid sizes. + for (auto r : {range<3>(3, 4, 5), range<3>(1, 1, 1), range<3>(4, 2, 1), + range<3>(32, 3, 4), range<3>(7, 9, 11)}) { Res = Kernel.get_info( Device); - bool Expected = + Expected = std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end(); exit_if_not_equal(Expected, true, "max_sub_group_size"); - - for (auto r : {range<3>(3, 4, 5), range<3>(1, 1, 1), range<3>(4, 2, 1), - range<3>(32, 3, 4), range<3>(7, 9, 11)}) { - Res = Kernel.get_info( - Device); - Expected = - std::find(sg_sizes.begin(), sg_sizes.end(), Res) != sg_sizes.end(); - exit_if_not_equal(Expected, true, "max_sub_group_size"); - } } Res = Kernel.get_info( @@ -81,21 +69,11 @@ int main() { /* Sub-group size is not specified in kernel or IL*/ exit_if_not_equal(Res, 0, "compile_num_sub_groups"); - // According to specification, this kernel query requires `cl_khr_subgroups` - // or `cl_intel_subgroups` - if ((std::find(Vec.begin(), Vec.end(), "cl_khr_subgroups") != - std::end(Vec)) || - std::find(Vec.begin(), Vec.end(), "cl_intel_subgroups") != - std::end(Vec) && - std::find(Vec.begin(), Vec.end(), - "cl_intel_required_subgroup_size") != std::end(Vec)) { - Res = - Kernel.get_info( - Device); - - /* Required sub-group size is not specified in kernel or IL*/ - exit_if_not_equal(Res, 0, "compile_sub_group_size"); - } + Res = Kernel.get_info( + Device); + + /* Required sub-group size is not specified in kernel or IL*/ + exit_if_not_equal(Res, 0, "compile_sub_group_size"); } catch (exception e) { std::cout << "SYCL exception caught: " << e.what(); diff --git a/sycl/test-e2e/SubGroup/reduce.cpp b/sycl/test-e2e/SubGroup/reduce.cpp index 9422db73a599a..165556df711b1 100644 --- a/sycl/test-e2e/SubGroup/reduce.cpp +++ b/sycl/test-e2e/SubGroup/reduce.cpp @@ -13,10 +13,6 @@ #include int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); check(Queue); check(Queue); diff --git a/sycl/test-e2e/SubGroup/reduce_fp16.cpp b/sycl/test-e2e/SubGroup/reduce_fp16.cpp index 1140ab26677f9..ea40b0897e7c1 100644 --- a/sycl/test-e2e/SubGroup/reduce_fp16.cpp +++ b/sycl/test-e2e/SubGroup/reduce_fp16.cpp @@ -10,10 +10,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/reduce_fp64.cpp b/sycl/test-e2e/SubGroup/reduce_fp64.cpp index 0fd801deda7ec..f00b0000a6351 100644 --- a/sycl/test-e2e/SubGroup/reduce_fp64.cpp +++ b/sycl/test-e2e/SubGroup/reduce_fp64.cpp @@ -8,10 +8,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/reduce_spirv13.cpp b/sycl/test-e2e/SubGroup/reduce_spirv13.cpp index 43fbdb3b215c5..aaee44963d54a 100644 --- a/sycl/test-e2e/SubGroup/reduce_spirv13.cpp +++ b/sycl/test-e2e/SubGroup/reduce_spirv13.cpp @@ -8,10 +8,6 @@ #include int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check_mul(Queue); check_mul(Queue); diff --git a/sycl/test-e2e/SubGroup/reduce_spirv13_fp16.cpp b/sycl/test-e2e/SubGroup/reduce_spirv13_fp16.cpp index d7e074551d9f9..d74fbd9a856a8 100644 --- a/sycl/test-e2e/SubGroup/reduce_spirv13_fp16.cpp +++ b/sycl/test-e2e/SubGroup/reduce_spirv13_fp16.cpp @@ -11,10 +11,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check_mul(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/reduce_spirv13_fp64.cpp b/sycl/test-e2e/SubGroup/reduce_spirv13_fp64.cpp index 6106ad2ba86a1..543dc32e24fc3 100644 --- a/sycl/test-e2e/SubGroup/reduce_spirv13_fp64.cpp +++ b/sycl/test-e2e/SubGroup/reduce_spirv13_fp64.cpp @@ -10,10 +10,6 @@ #include int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check_mul(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/scan.cpp b/sycl/test-e2e/SubGroup/scan.cpp index 989fcf7f588a7..74db27c7f5208 100644 --- a/sycl/test-e2e/SubGroup/scan.cpp +++ b/sycl/test-e2e/SubGroup/scan.cpp @@ -14,10 +14,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); check(Queue); check(Queue); diff --git a/sycl/test-e2e/SubGroup/scan_fp16.cpp b/sycl/test-e2e/SubGroup/scan_fp16.cpp index d49b960f4148d..a9d26f4a396cd 100644 --- a/sycl/test-e2e/SubGroup/scan_fp16.cpp +++ b/sycl/test-e2e/SubGroup/scan_fp16.cpp @@ -11,10 +11,6 @@ #include int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/scan_fp64.cpp b/sycl/test-e2e/SubGroup/scan_fp64.cpp index 6e32d278386ab..abd29566e9d77 100644 --- a/sycl/test-e2e/SubGroup/scan_fp64.cpp +++ b/sycl/test-e2e/SubGroup/scan_fp64.cpp @@ -9,10 +9,6 @@ #include int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/scan_spirv13.cpp b/sycl/test-e2e/SubGroup/scan_spirv13.cpp index 0966161844d1f..774283d0d6e13 100644 --- a/sycl/test-e2e/SubGroup/scan_spirv13.cpp +++ b/sycl/test-e2e/SubGroup/scan_spirv13.cpp @@ -9,10 +9,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check_mul(Queue); check_mul(Queue); check_mul(Queue); diff --git a/sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp b/sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp index 9a213bd885b1b..136c3210d122a 100644 --- a/sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp +++ b/sycl/test-e2e/SubGroup/scan_spirv13_fp16.cpp @@ -12,10 +12,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check_mul(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp b/sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp index f78a86c9a8e2d..8ea36613ed613 100644 --- a/sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp +++ b/sycl/test-e2e/SubGroup/scan_spirv13_fp64.cpp @@ -11,10 +11,6 @@ int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue); std::cout << "Test passed." << std::endl; return 0; diff --git a/sycl/test-e2e/SubGroup/vote.cpp b/sycl/test-e2e/SubGroup/vote.cpp index d835e2d2d9ec1..5fb9ac0066194 100644 --- a/sycl/test-e2e/SubGroup/vote.cpp +++ b/sycl/test-e2e/SubGroup/vote.cpp @@ -69,10 +69,6 @@ void check(queue Queue, const int G, const int L, const int D, const int R) { } int main() { queue Queue; - if (!core_sg_supported(Queue.get_device())) { - std::cout << "Skipping test\n"; - return 0; - } check(Queue, 240, 80, 3, 1); check(Queue, 24, 12, 3, 4); check(Queue, 1024, 256, 3, 1); From c533717ffc6ad17fea789fd313a19e7846f740b9 Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Mon, 24 Jun 2024 14:32:51 -0700 Subject: [PATCH 2/9] Mark attributes test as unsupported due to bug --- sycl/test-e2e/SubGroup/attributes.cpp | 4 ++++ 1 file changed, 4 insertions(+) diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index 1107e0554626f..e812a5435e6ef 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -1,3 +1,7 @@ +// UNSUPPORTED: accelerator +// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing +// this test to fail + // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out // RUN: %{run} %t.out //==------- attributes.cpp - SYCL sub_group attributes test ----*- C++ -*---==// From f36aaad73155b1beb08b3a1e40ea708f9d2a7f5b Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Wed, 26 Jun 2024 07:56:18 -0700 Subject: [PATCH 3/9] Mark info test unsupported due to bug --- sycl/test-e2e/SubGroup/info.cpp | 18 +++++++++++++++++- 1 file changed, 17 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/SubGroup/info.cpp b/sycl/test-e2e/SubGroup/info.cpp index e6b4ad65b6c96..c18b11362b390 100644 --- a/sycl/test-e2e/SubGroup/info.cpp +++ b/sycl/test-e2e/SubGroup/info.cpp @@ -1,3 +1,7 @@ +// UNSUPPORTED: accelerator +// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing +// this test to fail. Additionally, the kernel max_sub_group_size checks +// crash on FPGAs // RUN: %{build} -o %t.out // RUN: %{run} %t.out @@ -17,8 +21,20 @@ int main() { queue Queue; device Device = Queue.get_device(); + bool old_opencl = false; + if (Device.get_backend() == sycl::backend::opencl) { + // Extract the numerical version from the version string, OpenCL version + // string have the format "OpenCL . ". + std::string ver = Device.get_info().substr(7, 3); + old_opencl = (ver < "2.1"); + } + /* Check info::device parameters. */ - Device.get_info(); + if (!old_opencl) { + // Independent forward progress is missing on OpenCL backend prior to + // version 2.1 + Device.get_info(); + } Device.get_info(); try { From ccf3d48d67d1a141f1eef9a7f11f99b24eba90be Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Wed, 26 Jun 2024 08:13:12 -0700 Subject: [PATCH 4/9] Update supgroup size regression test Now checks if the subgroup sizes list should be empty --- sycl/test-e2e/Regression/get_subgroup_sizes.cpp | 15 ++++++++++----- 1 file changed, 10 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp index eb910425ea8d4..248d4cbb0b96e 100644 --- a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp +++ b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp @@ -1,3 +1,6 @@ +// UNSUPPORTED: accelerator +// TODO: FPGAs currently report `sub_group_sizes` as non-empty list, +// despite not having extension `cl_intel_required_subgroup_size` // RUN: %{build} -o %t.out // RUN: %{run} %t.out @@ -18,13 +21,15 @@ int main() { queue Q; auto Dev = Q.get_device(); auto Vec = Dev.get_info(); + std::vector SubGroupSizes = + Dev.get_info(); if (std::find(Vec.begin(), Vec.end(), "cl_intel_required_subgroup_size") != std::end(Vec)) { - std::vector SubGroupSizes = - Dev.get_info(); - std::vector::const_iterator MaxIter = - std::max_element(SubGroupSizes.begin(), SubGroupSizes.end()); - int MaxSubGroup_size = *MaxIter; + assert(!SubGroupSizes.empty() && + "Required sub-group size list should not be empty"); + } else { + assert(SubGroupSizes.empty() && + "Required sub-group size list should be empty"); } return 0; } From 4f62e3083122475ac8b70d060094ca0e8ba5dc63 Mon Sep 17 00:00:00 2001 From: David Garcia Date: Wed, 26 Jun 2024 15:21:48 -0400 Subject: [PATCH 5/9] clang format --- sycl/test-e2e/Regression/get_subgroup_sizes.cpp | 2 +- sycl/test-e2e/SubGroup/attributes.cpp | 15 +++++++-------- sycl/test-e2e/SubGroup/info.cpp | 4 ++-- 3 files changed, 10 insertions(+), 11 deletions(-) diff --git a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp index 248d4cbb0b96e..86ba7ed796163 100644 --- a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp +++ b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp @@ -1,5 +1,5 @@ // UNSUPPORTED: accelerator -// TODO: FPGAs currently report `sub_group_sizes` as non-empty list, +// TODO: FPGAs currently report `sub_group_sizes` as non-empty list, // despite not having extension `cl_intel_required_subgroup_size` // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index e812a5435e6ef..6c2b5372b2c00 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -1,5 +1,5 @@ // UNSUPPORTED: accelerator -// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing +// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing // this test to fail // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out @@ -15,13 +15,12 @@ #include "helper.hpp" #define KERNEL_FUNCTOR_WITH_SIZE(SIZE) \ - class KernelFunctor##SIZE { \ - public: \ - [[sycl::reqd_sub_group_size(SIZE)]] void \ - operator()(sycl::nd_item<1> Item) const { \ - const auto GID = Item.get_global_id(); \ - } \ - }; + class KernelFunctor## \ + SIZE{public : [[sycl::reqd_sub_group_size(SIZE)]] void operator()( \ + sycl::nd_item<1> Item) const {const auto GID = Item.get_global_id(); \ + } \ + } \ + ; KERNEL_FUNCTOR_WITH_SIZE(1); KERNEL_FUNCTOR_WITH_SIZE(2); diff --git a/sycl/test-e2e/SubGroup/info.cpp b/sycl/test-e2e/SubGroup/info.cpp index c18b11362b390..51e5f760b8de0 100644 --- a/sycl/test-e2e/SubGroup/info.cpp +++ b/sycl/test-e2e/SubGroup/info.cpp @@ -1,5 +1,5 @@ // UNSUPPORTED: accelerator -// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing +// TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing // this test to fail. Additionally, the kernel max_sub_group_size checks // crash on FPGAs // RUN: %{build} -o %t.out @@ -31,7 +31,7 @@ int main() { /* Check info::device parameters. */ if (!old_opencl) { - // Independent forward progress is missing on OpenCL backend prior to + // Independent forward progress is missing on OpenCL backend prior to // version 2.1 Device.get_info(); } From 5eb429afef611ef9620135ef5df63db6386127a7 Mon Sep 17 00:00:00 2001 From: David Garcia Date: Wed, 26 Jun 2024 15:37:30 -0400 Subject: [PATCH 6/9] clang format macro --- sycl/test-e2e/SubGroup/attributes.cpp | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index 6c2b5372b2c00..c9a25aa3b0ea1 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -15,12 +15,13 @@ #include "helper.hpp" #define KERNEL_FUNCTOR_WITH_SIZE(SIZE) \ - class KernelFunctor## \ - SIZE{public : [[sycl::reqd_sub_group_size(SIZE)]] void operator()( \ - sycl::nd_item<1> Item) const {const auto GID = Item.get_global_id(); \ +class KernelFunctor##SIZE { \ +public: \ + [[sycl::reqd_sub_group_size(SIZE)]] void \ + operator()(sycl::nd_item<1> Item) const { \ + const auto GID = Item.get_global_id(); \ } \ - } \ - ; +}; KERNEL_FUNCTOR_WITH_SIZE(1); KERNEL_FUNCTOR_WITH_SIZE(2); From 1fc8f1e825fec81ff71d2903e43f33f82c9131d6 Mon Sep 17 00:00:00 2001 From: David Garcia Date: Wed, 26 Jun 2024 15:44:00 -0400 Subject: [PATCH 7/9] format macro again --- sycl/test-e2e/SubGroup/attributes.cpp | 14 +++++++------- 1 file changed, 7 insertions(+), 7 deletions(-) diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index c9a25aa3b0ea1..75c1e7acf1384 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -15,13 +15,13 @@ #include "helper.hpp" #define KERNEL_FUNCTOR_WITH_SIZE(SIZE) \ -class KernelFunctor##SIZE { \ -public: \ - [[sycl::reqd_sub_group_size(SIZE)]] void \ - operator()(sycl::nd_item<1> Item) const { \ - const auto GID = Item.get_global_id(); \ - } \ -}; + class KernelFunctor##SIZE { \ + public: \ + [[sycl::reqd_sub_group_size(SIZE)]] void \ + operator()(sycl::nd_item<1> Item) const { \ + const auto GID = Item.get_global_id(); \ + } \ + }; KERNEL_FUNCTOR_WITH_SIZE(1); KERNEL_FUNCTOR_WITH_SIZE(2); From 48bd9c15d736f5b99411a7a663ac48dab02db069 Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Thu, 27 Jun 2024 12:20:57 -0700 Subject: [PATCH 8/9] Add unsupported hip/cuda for failing tests --- sycl/test-e2e/Regression/get_subgroup_sizes.cpp | 3 +++ sycl/test-e2e/SubGroup/attributes.cpp | 4 ++++ 2 files changed, 7 insertions(+) diff --git a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp index 86ba7ed796163..38a07ce20ef79 100644 --- a/sycl/test-e2e/Regression/get_subgroup_sizes.cpp +++ b/sycl/test-e2e/Regression/get_subgroup_sizes.cpp @@ -1,6 +1,9 @@ // UNSUPPORTED: accelerator // TODO: FPGAs currently report `sub_group_sizes` as non-empty list, // despite not having extension `cl_intel_required_subgroup_size` +// UNSUPPORTED: cuda || hip +// TODO: Similar issue to FPGAs + // RUN: %{build} -o %t.out // RUN: %{run} %t.out diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index 75c1e7acf1384..4dec73c144343 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -1,6 +1,10 @@ // UNSUPPORTED: accelerator // TODO: FPGAs currently report supported subgroups as {4,8,16,32,64}, causing // this test to fail +// UNSUPPORTED: cuda || hip +// TODO: Device subgroup sizes reports {32}, but when we try to use it with a +// kernel attribute and check it, we get a subgroup size of 0. + // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out // RUN: %{run} %t.out From a75c13eb3b39f0827cc8458a1b12541febe9ed1c Mon Sep 17 00:00:00 2001 From: "Garcia Orozco, David" Date: Thu, 27 Jun 2024 12:28:32 -0700 Subject: [PATCH 9/9] clang format --- sycl/test-e2e/SubGroup/attributes.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/sycl/test-e2e/SubGroup/attributes.cpp b/sycl/test-e2e/SubGroup/attributes.cpp index 4dec73c144343..a4503726c1dec 100644 --- a/sycl/test-e2e/SubGroup/attributes.cpp +++ b/sycl/test-e2e/SubGroup/attributes.cpp @@ -5,7 +5,6 @@ // TODO: Device subgroup sizes reports {32}, but when we try to use it with a // kernel attribute and check it, we get a subgroup size of 0. - // RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out // RUN: %{run} %t.out //==------- attributes.cpp - SYCL sub_group attributes test ----*- C++ -*---==//