From 0df930db0242333f113513ac5c5ae82521d29ff6 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Tue, 25 Feb 2025 16:57:17 +0000 Subject: [PATCH 1/5] Add tests for jit RTC --- .../KernelCompiler/kernel_compiler_basic.cpp | 79 +++++++++++++++++ .../kernel_compiler_namespaces.cpp | 82 ++++++++++++++++++ .../kernel_compiler_overload.cpp | 84 +++++++++++++++++++ .../no_sycl_hpp_in_e2e_tests.cpp | 2 +- 4 files changed, 246 insertions(+), 1 deletion(-) create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp create mode 100644 sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp new file mode 100644 index 0000000000000..49d758205c483 --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp @@ -0,0 +1,79 @@ +//==- kernel_compiler_sycl_jit.cpp --- kernel_compiler extension tests -----==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) +// REQUIRES: aspect-usm_device_allocations + +// UNSUPPORTED: accelerator +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out 1 + +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +int main() { + sycl::queue q; + + // The source code for a kernel, defined as a SYCL "free function kernel". + std::string source = R"""( + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + extern "C" + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(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); + } + )"""; + + // Create a kernel bundle in "source" state. + sycl::kernel_bundle kb_src = + syclexp::create_kernel_bundle_from_source( + q.get_context(), syclexp::source_language::sycl_jit, source); + + // Compile the kernel. There is no need to use the "registered_kernel_names" + // property because the kernel is declared extern "C". + sycl::kernel_bundle kb_exe = + syclexp::build(kb_src); + + // Get the kernel via its compiler-generated name. + sycl::kernel iota = kb_exe.ext_oneapi_get_kernel("iota"); + + float *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + // Set the values of the kernel arguments. + cgh.set_args(3.14f, ptr); + + // Launch the kernel according to its type, in this case an nd-range + // kernel. + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, iota); + }).wait(); + + constexpr float eps = 0.001; + for (int i = 0; i < NUM; i++) { + const float truth = 3.14f + static_cast(i); + if (std::abs(ptr[i] - truth) > eps) { + std::cout << "Result: " << ptr[i] << " expected " << i << "\n"; + sycl::free(ptr, q); + exit(1); + } + } + sycl::free(ptr, q); +} diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp new file mode 100644 index 0000000000000..f024ce29ea14c --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp @@ -0,0 +1,82 @@ +//==- kernel_compiler_overload.cpp --- kernel_compiler extension tests -----==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) +// REQUIRES: aspect-usm_device_allocations + +// UNSUPPORTED: accelerator +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out 1 + +#include +#include +#include +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +int main() { + sycl::queue q; + + // The source code for two kernels defined in different namespaces + std::string source = R"""( + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(int start, int *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } + + namespace mykernels { + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(int start, int *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } + } // namespace mykernels + )"""; + + // Create a kernel bundle in "source" state. + sycl::kernel_bundle kb_src = + syclexp::create_kernel_bundle_from_source( + q.get_context(), syclexp::source_language::sycl_jit, source); + + // Compile the kernel. Select kernel from the mykernels namespace + std::string iota_name{"mykernels::iota"}; + sycl::kernel_bundle kb_exe = syclexp::build( + kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + + sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); + + int *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + // Set the values of the kernel arguments. + cgh.set_args(3, ptr); + + // Launch the kernel according to its type, in this case an nd-range + // kernel. + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, iota); + }).wait(); + + for (int i = 0; i < NUM; i++) { + if (ptr[i] != i + 3) { + std::cout << "Result: " << ptr[i] << " expected " << i << "\n"; + sycl::free(ptr, q); + exit(1); + } + } + sycl::free(ptr, q); +} diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp new file mode 100644 index 0000000000000..381da8039f29e --- /dev/null +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp @@ -0,0 +1,84 @@ +//==- kernel_compiler_overload.cpp --- kernel_compiler extension tests -----==// +// +// 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 +// +//===----------------------------------------------------------------------===// + +// REQUIRES: (opencl || level_zero) +// REQUIRES: aspect-usm_device_allocations + +// UNSUPPORTED: accelerator +// UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. + +// RUN: %{build} -o %t.out +// RUN: %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out 1 + +#include +#include +#include +namespace syclexp = sycl::ext::oneapi::experimental; + +static constexpr size_t NUM = 1024; +static constexpr size_t WGSIZE = 16; + +int main() { + sycl::queue q; + + // The source code for two kernels defined as overloaded functions. + std::string source = R"""( + #include + namespace syclext = sycl::ext::oneapi; + namespace syclexp = sycl::ext::oneapi::experimental; + + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) + void iota(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::nd_range_kernel<1>)) + void iota(int start, int *ptr) { + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); + ptr[id] = start + static_cast(id); + } + )"""; + + // Create a kernel bundle in "source" state. + sycl::kernel_bundle kb_src = + syclexp::create_kernel_bundle_from_source( + q.get_context(), syclexp::source_language::sycl_jit, source); + + // Compile the kernel. Because there are two overloads of "iota", we need to + // use a C++ cast to disambiguate between them. Here, we are selecting the + // "int" overload. + std::string iota_name{"(void(*)(int, int*))iota"}; + sycl::kernel_bundle kb_exe = syclexp::build( + kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + + // Get the kernel by passing the same string we used to construct the + // "registered_kernel_names" property. + sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); + + int *ptr = sycl::malloc_shared(NUM, q); + q.submit([&](sycl::handler &cgh) { + // Set the values of the kernel arguments. + cgh.set_args(3, ptr); + + // Launch the kernel according to its type, in this case an nd-range + // kernel. + sycl::nd_range ndr{{NUM}, {WGSIZE}}; + cgh.parallel_for(ndr, iota); + }).wait(); + + for (int i = 0; i < NUM; i++) { + if (ptr[i] != i + 3) { + std::cout << "Result: " << ptr[i] << " expected " << i << "\n"; + sycl::free(ptr, q); + exit(1); + } + } + sycl::free(ptr, q); +} diff --git a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp index dd068fb40752a..23ea51be5d7c4 100644 --- a/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp +++ b/sycl/test/e2e_test_requirements/no_sycl_hpp_in_e2e_tests.cpp @@ -6,7 +6,7 @@ // CHECK-DAG: README.md // CHECK-DAG: lit.cfg.py // -// CHECK-NUM-MATCHES: 7 +// CHECK-NUM-MATCHES: 10 // // This test verifies that `` isn't used in E2E tests. Instead, // fine-grained includes should used, see From ea1008a8825b3ba67dde763ec2ffffe8a6376823 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Mon, 3 Mar 2025 08:58:56 +0000 Subject: [PATCH 2/5] Small fixes in spec doc examples --- .../sycl_ext_oneapi_kernel_compiler.asciidoc | 10 ++++++---- 1 file changed, 6 insertions(+), 4 deletions(-) diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc index d796e49b628a7..43794ea0c06f1 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_kernel_compiler.asciidoc @@ -818,7 +818,7 @@ int main() { extern "C" SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void iota(float start, float *ptr) { - size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id] = start + static_cast(id); } )"""; @@ -847,6 +847,7 @@ int main() { sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, iota); }).wait(); + sycl::free(ptr, q); } ---- @@ -874,13 +875,13 @@ int main() { SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void iota(float start, float *ptr) { - size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + 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::range_kernel<1>)) + SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((syclexp::nd_range_kernel<1>)) void iota(int start, int *ptr) { - size_t id = syclext::this_work_item::get_nd_item().get_global_linear_id(); + size_t id = syclext::this_work_item::get_nd_item<1>().get_global_linear_id(); ptr[id] = start + static_cast(id); } )"""; @@ -912,6 +913,7 @@ int main() { sycl::nd_range ndr{{NUM}, {WGSIZE}}; cgh.parallel_for(ndr, iota); }).wait(); + sycl::free(ptr, q); } ---- From 22a8dddc990c01e533adbe52f8556a366ac9b0f1 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Mon, 3 Mar 2025 09:02:40 +0000 Subject: [PATCH 3/5] Typo in comment --- sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp index f024ce29ea14c..0977119e19166 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp @@ -1,4 +1,4 @@ -//==- kernel_compiler_overload.cpp --- kernel_compiler extension tests -----==// +//==- kernel_compiler_namespaces.cpp --- kernel_compiler extension tests -----==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From 771b65da4add6a4d962547268b9231cca0879208 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Mon, 3 Mar 2025 09:14:35 +0000 Subject: [PATCH 4/5] Formatting --- sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp index 0977119e19166..452a929760cff 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp @@ -1,4 +1,4 @@ -//==- kernel_compiler_namespaces.cpp --- kernel_compiler extension tests -----==// +//==- kernel_compiler_namespaces.cpp --- kernel_compiler extension tests ---==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. From eac9671607c87807db769fc8fff74a90ae33e4e7 Mon Sep 17 00:00:00 2001 From: PietroGhg Date: Tue, 4 Mar 2025 10:21:49 +0000 Subject: [PATCH 5/5] Review comments --- sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp | 5 ++--- .../test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp | 5 ++--- sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp | 7 +++---- 3 files changed, 7 insertions(+), 10 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp index 49d758205c483..a2352977f0c2c 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_basic.cpp @@ -13,8 +13,7 @@ // UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. // RUN: %{build} -o %t.out -// RUN: %{run} %t.out 1 -// RUN: %{l0_leak_check} %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out #include #include @@ -47,7 +46,7 @@ int main() { syclexp::create_kernel_bundle_from_source( q.get_context(), syclexp::source_language::sycl_jit, source); - // Compile the kernel. There is no need to use the "registered_kernel_names" + // Compile the kernel. There is no need to use the "registered_names" // property because the kernel is declared extern "C". sycl::kernel_bundle kb_exe = syclexp::build(kb_src); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp index 452a929760cff..5bd1fb66f68ed 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_namespaces.cpp @@ -13,8 +13,7 @@ // UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. // RUN: %{build} -o %t.out -// RUN: %{run} %t.out 1 -// RUN: %{l0_leak_check} %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out #include #include @@ -56,7 +55,7 @@ int main() { // Compile the kernel. Select kernel from the mykernels namespace std::string iota_name{"mykernels::iota"}; sycl::kernel_bundle kb_exe = syclexp::build( - kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + kb_src, syclexp::properties{syclexp::registered_names{iota_name}}); sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp index 381da8039f29e..51e142f1939db 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_overload.cpp @@ -13,8 +13,7 @@ // UNSUPPORTED-INTENDED: while accelerator is AoT only, this cannot run there. // RUN: %{build} -o %t.out -// RUN: %{run} %t.out 1 -// RUN: %{l0_leak_check} %{run} %t.out 1 +// RUN: %{l0_leak_check} %{run} %t.out #include #include @@ -56,10 +55,10 @@ int main() { // "int" overload. std::string iota_name{"(void(*)(int, int*))iota"}; sycl::kernel_bundle kb_exe = syclexp::build( - kb_src, syclexp::properties{syclexp::registered_kernel_names{iota_name}}); + kb_src, syclexp::properties{syclexp::registered_names{iota_name}}); // Get the kernel by passing the same string we used to construct the - // "registered_kernel_names" property. + // "registered_names" property. sycl::kernel iota = kb_exe.ext_oneapi_get_kernel(iota_name); int *ptr = sycl::malloc_shared(NUM, q);