[SYCL][RTC] Disable dead argument elimination (#16139)

Unused arguments need to remain in place for free-function kernels, in order not to perturb the argument indices used with `handler::set_arg`. Signed-off-by: Julian Oppermann <[email protected]>
intel · Nov 26, 2024 · 5afa292 · 5afa292
1 parent 3c274a8
commit 5afa292
Show file tree

Hide file tree

Showing 2 changed files with 11 additions and 1 deletion.
diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp
@@ -175,6 +175,8 @@ jit_compiler::compileDeviceCode(InMemoryFile SourceFile,
   DerivedArgList DAL{UserArgList};
   const auto &OptTable = getDriverOptTable();
   DAL.AddFlagArg(nullptr, OptTable.getOption(OPT_fsycl_device_only));
+  DAL.AddFlagArg(nullptr,
+                 OptTable.getOption(OPT_fno_sycl_dead_args_optimization));
   DAL.AddJoinedArg(
       nullptr, OptTable.getOption(OPT_resource_dir_EQ),
       (DPCPPRoot + "/lib/clang/" + Twine(CLANG_VERSION_MAJOR)).str());
@@ -518,5 +520,11 @@ jit_compiler::parseUserArgs(View<const char *> UserArgs) {
         "Runtime compilation of ESIMD kernels is not yet supported");
   }
 
+  if (AL.hasFlag(OPT_fsycl_dead_args_optimization,
+                 OPT_fno_sycl_dead_args_optimization, false)) {
+    return createStringError(
+        "Dead argument optimization must be disabled for runtime compilation");
+  }
+
   return std::move(AL);
 }
diff --git a/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp b/sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp
@@ -56,7 +56,7 @@ auto constexpr SYCLSource = R"===(
 
 // use extern "C" to avoid name mangling
 extern "C" SYCL_EXTERNAL SYCL_EXT_ONEAPI_FUNCTION_PROPERTY((sycl::ext::oneapi::experimental::nd_range_kernel<1>))
-void ff_cp(int *ptr) {
+void ff_cp(int *ptr, int *unused) {
 
   // intentionally using deprecated routine, as opposed to this_work_item::get_nd_item<1>()
   sycl::nd_item<1> Item = sycl::ext::oneapi::experimental::this_nd_item<1>();
@@ -78,6 +78,7 @@ void test_1(sycl::queue &Queue, sycl::kernel &Kernel, int seed) {
   memset(usmPtr, 0, Range * sizeof(int));
   Queue.submit([&](sycl::handler &Handler) {
     Handler.set_arg(0, usmPtr);
+    Handler.set_arg(1, usmPtr);
     Handler.parallel_for(R1, Kernel);
   });
   Queue.wait();
@@ -181,6 +182,7 @@ int test_unsupported_options() {
   CheckUnsupported({"-Xarch_device", "-fsanitize=address"});
   CheckUnsupported({"-fsycl-device-code-split=kernel"});
   CheckUnsupported({"-fsycl-device-code-split-esimd"});
+  CheckUnsupported({"-fsycl-dead-args-optimization"});
 
   return 0;
 }