From 3c0307b4614effe1e24509ec993be1b6b6a79f2d Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 18 Jul 2025 20:08:05 +0200 Subject: [PATCH 1/4] [SYCL][E2E] add tests of kernel bundle --- .../FreeFunctionKernels/kernel_bundle.cpp | 91 +++++++++++++++++++ .../KernelAndProgram/free_function_apis.cpp | 38 +++++++- 2 files changed, 127 insertions(+), 2 deletions(-) create mode 100644 sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp diff --git a/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp b/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp new file mode 100644 index 0000000000000..dc61807f0d26d --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp @@ -0,0 +1,91 @@ +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out + +#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 auto FFTestMark = "Free function Kernel Test:"; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>)) +void func_range(float start, float *ptr) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void func_single(float start, float *ptr) {} + +template +int test_free_function_kernel_id(sycl::context &ctxt) { + sycl::kernel_id id = syclexp::get_kernel_id(); + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + const auto kernel_ids = exe_bndl.get_kernel_ids(); + const bool res = + std::find(kernel_ids.begin(), kernel_ids.end(), id) == kernel_ids.end(); + if (res) + std::cout << FFTestMark << "test_kernel_id failed: kernel id is not found" + << std::endl; + return res; +} + +template +int test_kernel_bundle_ctxt(sycl::context &ctxt, std::string_view fname) { + sycl::kernel_id id = syclexp::get_kernel_id(); + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt); + const bool res = + exe_bndl.has_kernel(id) && + exe_bndl.get_kernel(id) + .template get_info() == fname; + if (!res) + std::cout + << FFTestMark + << "test_kernel_bundle_ctxt failed: bundle does not contain kernel id " + "or function name " + << fname << std::endl; + return res; +} + +template +int test_kernel_bundle_ctxt_dev(sycl::context &ctxt, sycl::device &dev, + std::string_view fname) { + sycl::kernel_id id = syclexp::get_kernel_id(); + auto exe_bndl = + syclexp::get_kernel_bundle(ctxt, + {dev}); + const bool res = + exe_bndl.has_kernel(id) && + exe_bndl.get_kernel(id) + .template get_info() == fname; + if (!res) + std::cout << FFTestMark + << "test_kernel_bundle_ctxt_dev failed: bundle does not contain " + "kernel id " + "or function name " + << fname << std::endl; + return res; +} + + + + +int main() { + sycl::queue q; + sycl::context ctxt = q.get_context(); + sycl::device dev = q.get_device(); + + int ret = test_free_function_kernel_id(ctxt); + ret |= test_free_function_kernel_id(ctxt); + ret |= test_kernel_bundle_ctxt(ctxt, "func_range"); + ret |= test_kernel_bundle_ctxt(ctxt, "func_single"); + ret |= test_kernel_bundle_ctxt_dev(ctxt, dev, "func_range"); + ret |= test_kernel_bundle_ctxt_dev(ctxt, dev, "func_single"); + ret |= test_get_kernel_exception(); + return ret; +} + diff --git a/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp b/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp index 3293ece0441b0..690454e45e448 100644 --- a/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp +++ b/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp @@ -22,6 +22,10 @@ void ff_2(int *ptr, int start) { ptr2D[GId.get(0)][GId.get(1)] = LId.get(0) + LId.get(1) + start; } +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( + (ext::oneapi::experimental::nd_range_kernel<2>)) +void ff_b(int *ptr, int start) {} + // Templated free function definition. template SYCL_EXT_ONEAPI_FUNCTION_PROPERTY( @@ -97,6 +101,21 @@ bool test_kernel_apis(queue Queue) { return Pass; } +template +bool check_kernel_id(const sycl::kernel k, const sycl::context &ctx) { + namespace exp = ext::oneapi::experimental; + const auto id = exp::get_kernel_id(); + auto exe_bndl = + exp::get_kernel_bundle(ctx); + if (!exe_bndl.has_kernel(id)) + return false; + const auto kb = sycl::get_kernel_bundle(ctx); + const auto kernel_ids = kb.get_kernel_ids(); + bool ret = + std::find(kernel_ids.begin(), kernel_ids.end(), id) != kernel_ids.end(); + return ret; +} + bool test_bundle_apis(queue Queue) { bool Pass = true; @@ -134,6 +153,10 @@ bool test_bundle_apis(queue Queue) { std::cout << "PassE=" << PassE << std::endl; Pass &= PassE; + bool PassE2 = ext::oneapi::experimental::is_compatible(Device); + std::cout << "PassE2=" << PassE2 << std::endl; + Pass &= PassE2; + // Check that ff_2 is found in bundle. kernel_bundle Bundle2 = ext::oneapi::experimental::get_kernel_bundle< ff_2, bundle_state::executable>(Context); @@ -145,7 +168,7 @@ bool test_bundle_apis(queue Queue) { std::cout << "PassG=" << PassG << std::endl; Pass &= PassG; kernel Kernel2 = Bundle2.ext_oneapi_get_kernel(); - bool PassH = true; + bool PassH = check_kernel_id(Kernel2, Context); std::cout << "PassH=" << PassH << std::endl; Pass &= PassH; @@ -162,7 +185,8 @@ bool test_bundle_apis(queue Queue) { Pass &= PassJ; kernel Kernel3 = Bundle3.ext_oneapi_get_kernel<(void (*)(int *, int))ff_3>(); - bool PassK = true; + bool PassK = + check_kernel_id<(void (*)(int *, int))ff_3>(Kernel3, Context); std::cout << "PassK=" << PassK << std::endl; Pass &= PassK; @@ -197,6 +221,16 @@ bool test_bundle_apis(queue Queue) { std::cout << "PassP=" << PassP << std::endl; Pass &= PassP; + bool PassO1 = false; + try { + kernel Kernel51 = Bundle5.ext_oneapi_get_kernel(); + std::cout << "Wrong PATH of Kernel51" << std::endl; + } catch (const sycl::exception &e) { + PassO1 = e.code() == sycl::errc::invalid; + } + std::cout << "PassO1=" << PassO1 << std::endl; + Pass &= PassO1; + bool PassQ = Bundle6.ext_oneapi_has_kernel<(void (*)(int *, int))ff_3>(Device); std::cout << "PassQ=" << PassQ << std::endl; From e268d8dbf5b460a5367d5e02c81bed960c7bd3dc Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Fri, 18 Jul 2025 20:12:02 +0200 Subject: [PATCH 2/4] [SYCL][E2E] add test of separate binaries --- .../FreeFunctionKernels/separate_binaries.cpp | 38 +++++++++++++++++++ 1 file changed, 38 insertions(+) create mode 100644 sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp diff --git a/sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp b/sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp new file mode 100644 index 0000000000000..b40abbb2d00d5 --- /dev/null +++ b/sycl/test-e2e/FreeFunctionKernels/separate_binaries.cpp @@ -0,0 +1,38 @@ +// RUN: %{build} -fsycl-device-code-split=per_kernel -o %t.out +// RUN: %{run} %t.out + +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<2>)) +void func_range(float start, float *ptr) {} + +SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) +void func_single(float start, float *ptr) {} + +int main() { + sycl::queue q; + auto bundle_range = + syclexp::get_kernel_bundle( + q.get_context()); + try { + sycl::kernel k = bundle_range.ext_oneapi_get_kernel(); + } catch (const sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + return 0; + } + auto bundle_single = + syclexp::get_kernel_bundle( + q.get_context()); + try { + sycl::kernel k = bundle_single.ext_oneapi_get_kernel(); + } catch (const sycl::exception &e) { + assert(e.code() == sycl::errc::invalid); + return 0; + } + return 1; +} From 90d1f428d8459d5a3977a49a4b26413be774e8bb Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 21 Jul 2025 10:34:16 +0200 Subject: [PATCH 3/4] [SYCL][E2E] fix formatting [SYCL][E2E] remove unused function --- sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp | 8 +------- 1 file changed, 1 insertion(+), 7 deletions(-) diff --git a/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp b/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp index dc61807f0d26d..067f3f9b8c429 100644 --- a/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp +++ b/sycl/test-e2e/FreeFunctionKernels/kernel_bundle.cpp @@ -19,8 +19,7 @@ void func_range(float start, float *ptr) {} SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::single_task_kernel)) void func_single(float start, float *ptr) {} -template -int test_free_function_kernel_id(sycl::context &ctxt) { +template int test_free_function_kernel_id(sycl::context &ctxt) { sycl::kernel_id id = syclexp::get_kernel_id(); auto exe_bndl = syclexp::get_kernel_bundle(ctxt); @@ -71,9 +70,6 @@ int test_kernel_bundle_ctxt_dev(sycl::context &ctxt, sycl::device &dev, return res; } - - - int main() { sycl::queue q; sycl::context ctxt = q.get_context(); @@ -85,7 +81,5 @@ int main() { ret |= test_kernel_bundle_ctxt(ctxt, "func_single"); ret |= test_kernel_bundle_ctxt_dev(ctxt, dev, "func_range"); ret |= test_kernel_bundle_ctxt_dev(ctxt, dev, "func_single"); - ret |= test_get_kernel_exception(); return ret; } - From 3bf00ec41add7a40f3df044c24f9f43742644315 Mon Sep 17 00:00:00 2001 From: "Klochkov, Denis" Date: Mon, 21 Jul 2025 10:37:27 +0200 Subject: [PATCH 4/4] [SYCL][E2E] check exception in a separate test --- sycl/test-e2e/KernelAndProgram/free_function_apis.cpp | 10 ---------- 1 file changed, 10 deletions(-) diff --git a/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp b/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp index 690454e45e448..53737712c5bd9 100644 --- a/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp +++ b/sycl/test-e2e/KernelAndProgram/free_function_apis.cpp @@ -221,16 +221,6 @@ bool test_bundle_apis(queue Queue) { std::cout << "PassP=" << PassP << std::endl; Pass &= PassP; - bool PassO1 = false; - try { - kernel Kernel51 = Bundle5.ext_oneapi_get_kernel(); - std::cout << "Wrong PATH of Kernel51" << std::endl; - } catch (const sycl::exception &e) { - PassO1 = e.code() == sycl::errc::invalid; - } - std::cout << "PassO1=" << PassO1 << std::endl; - Pass &= PassO1; - bool PassQ = Bundle6.ext_oneapi_has_kernel<(void (*)(int *, int))ff_3>(Device); std::cout << "PassQ=" << PassQ << std::endl;