diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 910bfef64cd0f..1d75136160e99 100644 --- 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 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 index 3402691453287..450674e876d3e 100644 --- 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; }