diff --git a/sycl/source/detail/jit_compiler.cpp b/sycl/source/detail/jit_compiler.cpp index 8e466b97d75cf..c366996aadc9b 100644 --- a/sycl/source/detail/jit_compiler.cpp +++ b/sycl/source/detail/jit_compiler.cpp @@ -1211,17 +1211,21 @@ sycl_device_binaries 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 << 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 = 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 450674e876d3e..563f75d313e95 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, int *unused) { 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, T *unused) { + + 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) { @@ -126,19 +137,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_S1_"); + // 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; }