[SYCL][RTC] Ensure template kernel instantiation

Signed-off-by: Julian Oppermann <[email protected]>

Author: jopperm

sycl/source/detail/jit_compiler.cpp

@@ -1173,17 +1173,20 @@ std::vector<uint8_t> jit_compiler::compileSYCL(
     const std::vector<std::string> &UserArgs, std::string *LogPtr,
     const std::vector<std::string> &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());

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 <typename T>
+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<std::string> flags{"-g", "-fno-fast-math",
                                  "-fsycl-instrument-device-code"};
   std::vector<sycl::device> 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<int>"}});
 
-  // 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<int>");
+
+  // 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;
 }