From 87141decd8837ba889add1dfd4a92cde7ce84f1a Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Thu, 5 Jun 2025 01:41:36 -0700 Subject: [PATCH 1/3] [SYCL] do not reassign boolean flag if there is more than one attribute --- clang/lib/Sema/SemaSYCL.cpp | 1 + .../FreeFunctionKernels/properties.cpp | 89 +++++++++++++++++++ 2 files changed, 90 insertions(+) create mode 100644 sycl/test-e2e/FreeFunctionKernels/properties.cpp diff --git a/clang/lib/Sema/SemaSYCL.cpp b/clang/lib/Sema/SemaSYCL.cpp index cf64331198c91..3e19efe765c0e 100644 --- a/clang/lib/Sema/SemaSYCL.cpp +++ b/clang/lib/Sema/SemaSYCL.cpp @@ -1176,6 +1176,7 @@ bool SemaSYCL::isFreeFunction(const FunctionDecl *FD) { NameValuePair.first == "sycl-single-task-kernel"; }); IsFreeFunctionAttr = it != NameValuePairs.end(); + if (IsFreeFunctionAttr) break; } if (Redecl->isFirstDecl()) { if (IsFreeFunctionAttr) diff --git a/sycl/test-e2e/FreeFunctionKernels/properties.cpp b/sycl/test-e2e/FreeFunctionKernels/properties.cpp new file mode 100644 index 0000000000000..e75685a7badfc --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/properties.cpp @@ -0,0 +1,89 @@ +// REQUIRES: aspect-usm_shared_allocations +// RUN: %{build} %cxx_std_optionc++20 -o %t.out +// RUN: %{run} %t.out + +#include +#include +#include +#include +#include + +namespace syclext = sycl::ext::oneapi; +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +inline void kernel_code(float start, float *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void func_range_wg_1dsize_before(float start, float *ptr) { + kernel_code(start, ptr); +} + + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +void func_range_wg_1dsize_after(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) +void func_range_wg_1dsize_hint_after(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void func_range_wg_1dsize_hint_before(float start, float *ptr) { + kernel_code(start, ptr); +} + +template +bool check_result(T *ptr) { + for (size_t i = 0; i < NUM; ++i) { + const float expected = static_cast(3.14f) + static_cast(i); + if (ptr[i] != expected) + return true; + } + return false; +} + +template +static bool call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { + T *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + cgh.set_args(3.14f, ptr); + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, kernel); + }).wait(); + const bool ret = check_result(ptr); + sycl::free(ptr, q); + return ret; +} + +template +bool test_function(sycl::queue &q, sycl::context &ctxt) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + bool ret = call_kernel_code(q, k_func); + auto attrs_info_kernel = sycl::ext::oneapi::experimental::get_kernel_info(ctxt); +} + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + + bool ret = 0; + ret |= test_function(q, ctxt); + ret |= test_function(q, ctxt); + ret |= test_function(q, ctxt); + ret |= test_function(q, ctxt); + return ret; +} From fdf6a0df3c6aa3874881dc870df94ea73cf47647 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 30 Jun 2025 13:46:47 +0200 Subject: [PATCH 2/3] [SYCL][E2E] add more tests to check kernel free function properties --- .../FreeFunctionKernels/properties.cpp | 120 ++++++++++++++---- 1 file changed, 96 insertions(+), 24 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/properties.cpp b/sycl/test-e2e/FreeFunctionKernels/properties.cpp index e75685a7badfc..3c178611ceb05 100644 --- a/sycl/test-e2e/FreeFunctionKernels/properties.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/properties.cpp @@ -1,18 +1,19 @@ // REQUIRES: aspect-usm_shared_allocations -// RUN: %{build} %cxx_std_optionc++20 -o %t.out +// RUN: %{build} -o %t.out // RUN: %{run} %t.out -#include #include +#include #include + #include -#include namespace syclext = sycl::ext::oneapi; namespace syclexp = sycl::ext::oneapi::experimental; static constexpr size_t NUM = 1024; -static constexpr size_t WGSIZE = 16; +static constexpr size_t WGSIZE = 32; +static constexpr size_t SGSIZE = 16; inline void kernel_code(float start, float *ptr) { size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); @@ -21,35 +22,59 @@ inline void kernel_code(float start, float *ptr) { SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void func_range_wg_1dsize_before(float start, float *ptr) { +void range_wg_1dsize_before(float start, float *ptr) { kernel_code(start, ptr); } - SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) -void func_range_wg_1dsize_after(float start, float *ptr) { - kernel_code(start, ptr); -} +void range_wg_1dsize_after(float start, float *ptr) { kernel_code(start, ptr); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) -void func_range_wg_1dsize_hint_after(float start, float *ptr) { +void range_wg_1dsize_hint_after(float start, float *ptr) { kernel_code(start, ptr); } SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size_hint)) SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) -void func_range_wg_1dsize_hint_before(float start, float *ptr) { +void range_wg_1dsize_hint_before(float start, float *ptr) { kernel_code(start, ptr); } -template -bool check_result(T *ptr) { +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_sg_1dsize_before(float start, float *ptr) { + kernel_code(start, ptr); +} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +void range_sg_1dsize_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::device_has)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_has_before(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::device_has)) +void range_has_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +void range_several_after(float start, float *ptr) { kernel_code(start, ptr); } + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::work_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::sub_group_size)) +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) +void range_several_before(float start, float *ptr) { kernel_code(start, ptr); } + +template bool check_result(T *ptr) { for (size_t i = 0; i < NUM; ++i) { - const float expected = static_cast(3.14f) + static_cast(i); + const T expected = static_cast(3.14f) + static_cast(i); if (ptr[i] != expected) - return true; + return true; } return false; } @@ -67,23 +92,70 @@ static bool call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { return ret; } -template -bool test_function(sycl::queue &q, sycl::context &ctxt) { +template +bool test(sycl::queue &q, sycl::context &ctxt, const size_t cmp_value) { auto exe_bndl = syclexp::get_kernel_bundle(ctxt); sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); bool ret = call_kernel_code(q, k_func); - auto attrs_info_kernel = sycl::ext::oneapi::experimental::get_kernel_info(ctxt); + auto value = syclexp::get_kernel_info(q); + ret &= (cmp_value >= value); + + const auto kernel_ids = exe_bndl.get_kernel_ids(); + if (kernel_ids.empty()) + return true; + sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); + const size_t kernel_value = k.get_info(q.get_device()); + ret &= (value != kernel_value); + return ret; +} + +template +bool test_has_desc(sycl::queue &q, sycl::context &ctxt) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + return call_kernel_code(q, k_func); +} + +using wg_size_desc = sycl::info::kernel_device_specific::work_group_size; +using sg_size_desc = sycl::info::kernel_device_specific::compile_sub_group_size; + +template +bool test_several_properties(sycl::queue &q, sycl::context &ctxt) { + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); + bool ret = call_kernel_code(q, k_func); + const size_t value_wg_size = syclexp::get_kernel_info(q); + const size_t value_sg_size = syclexp::get_kernel_info(q); + + const auto kernel_ids = exe_bndl.get_kernel_ids(); + if (kernel_ids.empty()) + return true; + sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); + const size_t kernel_value_wg = k.get_info(q.get_device()); + const size_t kernel_value_sg = k.get_info(q.get_device()); + ret &= (value_wg_size != kernel_value_wg); + ret &= (value_sg_size != kernel_value_sg); + return ret; } int main() { + sycl::queue q; sycl::context ctxt = q.get_context(); - - bool ret = 0; - ret |= test_function(q, ctxt); - ret |= test_function(q, ctxt); - ret |= test_function(q, ctxt); - ret |= test_function(q, ctxt); + + int ret = 0; + ret |= test(q, ctxt, WGSIZE); + ret |= test(q, ctxt, WGSIZE); + ret |= + test(q, ctxt, WGSIZE); + ret |= test(q, ctxt, WGSIZE); + ret |= test(q, ctxt, SGSIZE); + ret |= test(q, ctxt, SGSIZE); + ret |= test_has_desc(q, ctxt); + ret |= test_has_desc(q, ctxt); + ret |= test_several_properties(q, ctxt); return ret; } From 5fbaf8a65d94e06e51c2c83463fcc8ca956d2799 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Tue, 1 Jul 2025 16:22:32 +0200 Subject: [PATCH 3/3] [SYCL][E2E] add message if test failed --- .../FreeFunctionKernels/properties.cpp | 50 ++++++++++++------- 1 file changed, 33 insertions(+), 17 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/properties.cpp b/sycl/test-e2e/FreeFunctionKernels/properties.cpp index 3c178611ceb05..46124b577dbd9 100644 --- a/sycl/test-e2e/FreeFunctionKernels/properties.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/properties.cpp @@ -2,6 +2,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out +#include #include #include #include @@ -93,20 +94,23 @@ static bool call_kernel_code(sycl::queue &q, sycl::kernel &kernel) { } template -bool test(sycl::queue &q, sycl::context &ctxt, const size_t cmp_value) { +bool test(sycl::queue &q, sycl::context &ctxt, std::string_view name) { auto exe_bndl = syclexp::get_kernel_bundle(ctxt); sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); bool ret = call_kernel_code(q, k_func); auto value = syclexp::get_kernel_info(q); - ret &= (cmp_value >= value); - const auto kernel_ids = exe_bndl.get_kernel_ids(); if (kernel_ids.empty()) return true; sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); const size_t kernel_value = k.get_info(q.get_device()); - ret &= (value != kernel_value); + ret |= (value != kernel_value); + if (ret) + std::cout << "Test " << name + << " did not pass: value got from get_kernel_info " << value + << ", value got from kernel get_info " << kernel_value + << std::endl; return ret; } @@ -122,40 +126,52 @@ using wg_size_desc = sycl::info::kernel_device_specific::work_group_size; using sg_size_desc = sycl::info::kernel_device_specific::compile_sub_group_size; template -bool test_several_properties(sycl::queue &q, sycl::context &ctxt) { +bool test_several_properties(sycl::queue &q, sycl::context &ctxt, + std::string_view name) { auto exe_bndl = syclexp::get_kernel_bundle(ctxt); sycl::kernel k_func = exe_bndl.template ext_oneapi_get_kernel(); bool ret = call_kernel_code(q, k_func); const size_t value_wg_size = syclexp::get_kernel_info(q); const size_t value_sg_size = syclexp::get_kernel_info(q); - const auto kernel_ids = exe_bndl.get_kernel_ids(); if (kernel_ids.empty()) return true; sycl::kernel k = exe_bndl.get_kernel(kernel_ids[0]); const size_t kernel_value_wg = k.get_info(q.get_device()); const size_t kernel_value_sg = k.get_info(q.get_device()); - ret &= (value_wg_size != kernel_value_wg); - ret &= (value_sg_size != kernel_value_sg); + ret |= (value_wg_size != kernel_value_wg); + ret |= (value_sg_size != kernel_value_sg); + if (ret) + std::cout << "Test " << name << " did not pass: value_wg_size " + << value_wg_size << ", value_sg_size " << value_sg_size + << ", kernel_value_wg " << kernel_value_wg << ", kernel_value_sg " + << kernel_value_sg << std::endl; return ret; } int main() { - sycl::queue q; sycl::context ctxt = q.get_context(); int ret = 0; - ret |= test(q, ctxt, WGSIZE); - ret |= test(q, ctxt, WGSIZE); - ret |= - test(q, ctxt, WGSIZE); - ret |= test(q, ctxt, WGSIZE); - ret |= test(q, ctxt, SGSIZE); - ret |= test(q, ctxt, SGSIZE); + ret |= test( + q, ctxt, "range_wg_1dsize_before"); + ret |= test( + q, ctxt, "range_wg_1dsize_after"); + ret |= test( + q, ctxt, "range_wg_1dsize_hint_before"); + ret |= test( + q, ctxt, "range_wg_1dsize_hint_after"); + ret |= test( + q, ctxt, "range_sg_1dsize_before"); + ret |= test( + q, ctxt, "range_sg_1dsize_after"); ret |= test_has_desc(q, ctxt); ret |= test_has_desc(q, ctxt); - ret |= test_several_properties(q, ctxt); + ret |= test_several_properties( + q, ctxt, "range_several_before"); + ret |= test_several_properties( + q, ctxt, "range_several_after"); return ret; }