From 2018109089fa25ecaddd90f246012e2a56d4549b Mon Sep 17 00:00:00 2001 From: Julian Oppermann Date: Wed, 20 Nov 2024 17:59:41 +0000 Subject: [PATCH] [SYCL][RTC] Disable dead argument elimination Signed-off-by: Julian Oppermann --- sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp | 8 ++++++++ sycl/test-e2e/KernelCompiler/kernel_compiler_sycl_jit.cpp | 4 +++- 2 files changed, 11 insertions(+), 1 deletion(-) diff --git a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp index 1bdfe7d63b641..6d96e4c35ed16 100644 --- a/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp +++ b/sycl-jit/jit-compiler/lib/rtc/DeviceCompilation.cpp @@ -167,6 +167,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()); @@ -410,5 +412,11 @@ jit_compiler::parseUserArgs(View UserArgs) { } } + 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 Expected{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 01f25f813b826..e89009ded0282 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(); @@ -179,6 +180,7 @@ int test_unsupported_options() { CheckUnsupported({"-Xsycl-target-frontend", "-fsanitize=address"}); CheckUnsupported({"-Xsycl-target-frontend=spir64", "-fsanitize=address"}); CheckUnsupported({"-Xarch_device", "-fsanitize=address"}); + CheckUnsupported({"-fsycl-dead-args-optimization"}); return 0; }