From ba0e6f8bdfcd2afb2635bcc7ac295c8f854337cd Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 20 Nov 2024 17:32:21 +0000 Subject: [PATCH 1/3] [SYCL][RTC] Ensure template kernel instantiation Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.cpp | 17 ++++++---- .../kernel_compiler_sycl_jit.cpp | 32 ++++++++++++++++--- 2 files changed, 38 insertions(+), 11 deletions(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 19f1915943f05..6812fc584e8b5 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1173,17 +1173,20 @@ std::vector jit_compiler::compileSYCL( const std::vector &UserArgs, std::string *LogPtr, const std::vector &RegisteredKernelNames) { - // TODO: Handle template instantiation. - if (!RegisteredKernelNames.empty()) { - throw sycl::exception( - sycl::errc::build, - "Property `sycl::ext::oneapi::experimental::registered_kernel_names` " - "is not yet supported for the `sycl_jit` source language"); + // RegisteredKernelNames may contain template specializations, so we just put + // them in main() which ensures they are instantiated. + std::ostringstream ss; + ss << "int main() {\n"; + for (const std::string &KernelName : RegisteredKernelNames) { + ss << " (void)" << KernelName << ";\n"; } + ss << " return 0;\n}\n" << std::endl; + + std::string FinalSource = SYCLSource + ss.str(); std::string SYCLFileName = Id + ".cpp"; ::jit_compiler::InMemoryFile SourceFile{SYCLFileName.c_str(), - SYCLSource.c_str()}; + FinalSource.c_str()}; std::vector<::jit_compiler::InMemoryFile> IncludeFilesView; IncludeFilesView.reserve(IncludePairs.size()); diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index 01f25f813b826..72bdfcd9d36bb 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -64,6 +64,17 @@ void ff_cp(int *ptr) { sycl::id<1> GId = Item.get_global_id(); ptr[GId.get(0)] = AddEm(GId.get(0), 37); } + +// this name will be mangled +template +SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) +void ff_templated(T *ptr) { + + sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); + + sycl::id<1> GId = Item.get_global_id(); + ptr[GId.get(0)] = PlusEm(GId.get(0), 38); +} )==="; void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) { @@ -125,19 +136,32 @@ int test_build_and_run() { // Compilation of empty prop list, no devices. exe_kb kbExe1 = syclex::build(kbSrc); - // // Compilation with props and devices + // Compilation with props and devices std::string log; std::vector flags{"-g", "-fno-fast-math", "-fsycl-instrument-device-code"}; std::vector devs = kbSrc.get_devices(); exe_kb kbExe2 = syclex::build( - kbSrc, devs, syclex::properties{syclex::build_options{flags}}); + kbSrc, devs, + syclex::properties{syclex::build_options{flags}, syclex::save_log{&log}, + syclex::registered_kernel_names{"ff_templated"}}); - // extern "C" was used, so the name "ff_cp" is not mangled. + // extern "C" was used, so the name "ff_cp" is not mangled and can be used + // directly. sycl::kernel k = kbExe2.ext_oneapi_get_kernel("ff_cp"); + // The templated function name will have been mangled. Mapping from original + // name to mangled is not yet supported. So we cannot yet do this: + // sycl::kernel k2 = kbExe2.ext_oneapi_get_kernel("ff_templated"); + + // Instead, we can TEMPORARILY use the mangled name. Once demangling is + // supported this might no longer work. + sycl::kernel k2 = + kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); + // Test the kernels. - test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. + test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more. + test_1(q, k2, 38 + 6); // ff_templated seeds 38. PlusEm adds 6 more. return 0; } From 7585efb3329fca01b94574276b1adb9bd9a9b5f0 Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Thu, 21 Nov 2024 10:32:40 +0000 Subject: [PATCH 2/3] Review feedback Signed-off-by: Julian Oppermann --- sycl/source/detail/jit_compiler.cpp | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 6812fc584e8b5..9c84507a2905b 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1176,13 +1176,14 @@ std::vector jit_compiler::compileSYCL( // RegisteredKernelNames may contain template specializations, so we just put // them in main() which ensures they are instantiated. std::ostringstream ss; + ss << SYCLSource << '\n'; ss << "int main() {\n"; for (const std::string &KernelName : RegisteredKernelNames) { ss << " (void)" << KernelName << ";\n"; } ss << " return 0;\n}\n" << std::endl; - std::string FinalSource = SYCLSource + ss.str(); + std::string FinalSource = ss.str(); std::string SYCLFileName = Id + ".cpp"; ::jit_compiler::InMemoryFile SourceFile{SYCLFileName.c_str(), From 22c4b0c00eafcc37cc3597f2147e257a1114205f Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Tue, 26 Nov 2024 11:26:29 +0000 Subject: [PATCH 3/3] Add unused arg to template kernel in E2E test Signed-off-by: Julian Oppermann --- sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp index bd1b9ca0f760b..563f75d313e95 100644 --- a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp +++ b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp @@ -68,7 +68,7 @@ void ff_cp(int *ptr, int *unused) { // this name will be mangled template SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>)) -void ff_templated(T *ptr) { +void ff_templated(T *ptr, T *unused) { sycl::nd_item<1> Item = sycl::ext::oneapi::this_work_item::get_nd_item<1>(); @@ -158,7 +158,7 @@ int test_build_and_run() { // Instead, we can TEMPORARILY use the mangled name. Once demangling is // supported this might no longer work. sycl::kernel k2 = - kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_"); + kbExe2.ext_oneapi_get_kernel("_Z26__sycl_kernel_ff_templatedIiEvPT_S1_"); // Test the kernels. test_1(q, k, 37 + 5); // ff_cp seeds 37. AddEm will add 5 more.