From 2ee12b057c29807bcff2bdb341aaecba44ab867e Mon Sep 17 00:00:00 2001 From: y Date: Tue, 18 Nov 2025 12:18:04 -0800 Subject: [PATCH 01/11] enable new offloading model by default --- clang/include/clang/Driver/Driver.h | 7 +++++ clang/lib/Driver/Driver.cpp | 37 +++++++++++++++++++++++---- clang/lib/Driver/ToolChains/Clang.cpp | 4 ++- 3 files changed, 42 insertions(+), 6 deletions(-) diff --git a/clang/include/clang/Driver/Driver.h b/clang/include/clang/Driver/Driver.h index 3dd8822c4844b..126f127de8674 100644 --- a/clang/include/clang/Driver/Driver.h +++ b/clang/include/clang/Driver/Driver.h @@ -626,6 +626,13 @@ class Driver { /// @name Helper Methods /// @{ + /// Utility function to parse all devices passed via -fsycl-targets. + /// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets. + /// Otherwise return 'false'. + bool + GetUseNewOffloadDriverForSYCLOffload(Compilation &C, + const llvm::opt::ArgList &Args) const; + /// getSYCLDeviceTriple - Returns the SYCL device triple for the /// specified subarch // TODO: Additional Arg input parameter is for diagnostic output information diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 475c1fddb0363..9e039ef3724c5 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1565,6 +1565,32 @@ static void appendOneArg(InputArgList &Args, const Arg *Opt) { } } +// Utility function to parse all devices passed via -fsycl-targets. +// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets. +// Otherwise return 'false'. +bool Driver::GetUseNewOffloadDriverForSYCLOffload(Compilation &C, + const ArgList &Args) const { + // Check only if enabled with -fsycl + if (!Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false)) + return false; + + if (Args.hasFlag(options::OPT_no_offload_new_driver, + options::OPT_offload_new_driver, false)) + return false; + + if (Args.hasArg(options::OPT_fintelfpga)) + return false; + + if (const Arg *A = Args.getLastArg(options::OPT_fsycl_targets_EQ)) { + for (const char *Val : A->getValues()) { + llvm::Triple TT(C.getDriver().getSYCLDeviceTriple(Val)); + if ((!TT.isSPIROrSPIRV()) || TT.isSPIRAOT()) + return false; + } + } + return true; +} + bool Driver::readConfigFile(StringRef FileName, llvm::cl::ExpansionContext &ExpCtx) { // Try opening the given file. @@ -2195,12 +2221,12 @@ Compilation *Driver::BuildCompilation(ArrayRef ArgList) { // Use new offloading path for OpenMP. This is disabled as the SYCL // offloading path is not properly setup to use the updated device linking // scheme. - if ((C->isOffloadingHostKind(Action::OFK_OpenMP) && - TranslatedArgs->hasFlag(options::OPT_fopenmp_new_driver, - options::OPT_no_offload_new_driver, true)) || + if (C->isOffloadingHostKind(Action::OFK_OpenMP) || TranslatedArgs->hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false)) + options::OPT_no_offload_new_driver, false) || + GetUseNewOffloadDriverForSYCLOffload(*C, *TranslatedArgs)) { setUseNewOffloadingDriver(); + } // Construct the list of abstract actions to perform for this compilation. On // MachO targets this uses the driver-driver and universal actions. @@ -7095,7 +7121,8 @@ void Driver::BuildDefaultActions(Compilation &C, DerivedArgList &Args, options::OPT_fno_offload_via_llvm, false) || Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, - C.isOffloadingHostKind(Action::OFK_Cuda)); + C.isOffloadingHostKind(Action::OFK_Cuda)) || + GetUseNewOffloadDriverForSYCLOffload(C, Args); bool HIPNoRDC = C.isOffloadingHostKind(Action::OFK_HIP) && diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 2e857f13990be..1c485708caa10 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5232,7 +5232,9 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, (JA.isHostOffloading(C.getActiveOffloadKinds()) && Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, - C.isOffloadingHostKind(Action::OFK_Cuda))); + C.isOffloadingHostKind(Action::OFK_Cuda))) || + (JA.isHostOffloading(Action::OFK_SYCL) && + C.getDriver().GetUseNewOffloadDriverForSYCLOffload(C, Args)); bool IsRDCMode = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, IsSYCL); From 148845461f3f9237af460c1c0fe55f0fd795beac Mon Sep 17 00:00:00 2001 From: y Date: Wed, 19 Nov 2025 10:58:42 -0800 Subject: [PATCH 02/11] resolve the issue of backend compile option got passed into ld --- .../ClangLinkerWrapper.cpp | 18 ++++++++++++++++++ 1 file changed, 18 insertions(+) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 204d6375fd782..2cfa43d8dc487 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -427,6 +427,24 @@ Error runLinker(ArrayRef Files, const ArgList &Args) { return createStringError("linker path missing, must pass 'linker-path'"); ArgStringList NewLinkerArgs; for (const opt::Arg *Arg : Args) { + // DEBUG: Print argument information before checking WrapperOnlyOption + llvm::errs() << "DEBUG: Processing argument: " << Arg->getSpelling(); + + // Safely check if the argument has values before accessing them + if (Arg->getNumValues() > 0) { + llvm::errs() << " with value: " << Arg->getValue(); + } else { + llvm::errs() << " (no value)"; + } + + llvm::errs() << " hasFlag(WrapperOnlyOption): " << Arg->getOption().hasFlag(WrapperOnlyOption) + << "\n"; + + StringRef Spelling = Arg->getSpelling(); + if (Spelling.starts_with("--sycl-target-link-options")) { + continue; + } + // Do not forward arguments only intended for the linker wrapper. if (Arg->getOption().hasFlag(WrapperOnlyOption)) continue; From da1c270dbcf16f1c2af41a9f238cf7b831ccdf2f Mon Sep 17 00:00:00 2001 From: y Date: Wed, 19 Nov 2025 19:49:08 -0800 Subject: [PATCH 03/11] remove linker and backend option to be passed into clang-linker-wrapper --- clang/lib/Driver/ToolChains/Clang.cpp | 52 +++++++++++-------- .../ClangLinkerWrapper.cpp | 25 --------- 2 files changed, 29 insertions(+), 48 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 1c485708caa10..958009ee6d301 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11487,29 +11487,35 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, WrapperOption = "--cpu-tool-arg="; } else continue; - ArgStringList BuildArgs; - SmallString<128> BackendOptString; - SmallString<128> LinkOptString; - SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); - for (const auto &A : BuildArgs) - appendOption(BackendOptString, A); - - BuildArgs.clear(); - SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); - for (const auto &A : BuildArgs) { - if (IsJIT) - appendOption(LinkOptString, A); - else - // For AOT, combine the Backend and Linker strings into one. - appendOption(BackendOptString, A); - } - if (!BackendOptString.empty()) - CmdArgs.push_back( - Args.MakeArgString(Twine(WrapperOption) + BackendOptString)); - if (!LinkOptString.empty()) - CmdArgs.push_back( - Args.MakeArgString("--sycl-target-link-options=" + LinkOptString)); - } + // ArgStringList BuildArgs; + // SmallString<128> BackendOptString; + // SmallString<128> LinkOptString; + // SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); + // for (const auto &A : BuildArgs) + // appendOption(BackendOptString, A); + + // BuildArgs.clear(); + // SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); + // for (const auto &A : BuildArgs) { + // if (IsJIT) + // int x = 1; + // // appendOption(LinkOptString, A); + // else + // // For AOT, combine the Backend and Linker strings into one. + // appendOption(BackendOptString, A); + // } + // if (!BackendOptString.empty()) { + // llvm::errs() << "DEBUG: Adding argument: '" << BackendOptString << "'\n"; + // CmdArgs.push_back( + // Args.MakeArgString(Twine(WrapperOption) + BackendOptString)); + // } + // if (!LinkOptString.empty()){ + // llvm::errs() << "DEBUG: Adding argument: '" << LinkOptString << "'\n"; + // CmdArgs.push_back( + // Args.MakeArgString("--sycl-target-link-options=" + LinkOptString)); + // } + } + // Add option to enable creating of the .syclbin file. if (Arg *A = Args.getLastArg(options::OPT_fsyclbin_EQ)) CmdArgs.push_back( diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 2cfa43d8dc487..16b11e494b26a 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -224,11 +224,6 @@ void printCommands(ArrayRef CmdArgs) { if (CmdArgs.empty()) return; - llvm::errs() << " \"" << CmdArgs.front() << "\" "; - for (auto IC = std::next(CmdArgs.begin()), IE = CmdArgs.end(); IC != IE; ++IC) - llvm::errs() << *IC << (std::next(IC) != IE ? " " : "\n"); -} - [[noreturn]] void reportError(Error E) { outs().flush(); logAllUnhandledErrors(std::move(E), @@ -427,24 +422,6 @@ Error runLinker(ArrayRef Files, const ArgList &Args) { return createStringError("linker path missing, must pass 'linker-path'"); ArgStringList NewLinkerArgs; for (const opt::Arg *Arg : Args) { - // DEBUG: Print argument information before checking WrapperOnlyOption - llvm::errs() << "DEBUG: Processing argument: " << Arg->getSpelling(); - - // Safely check if the argument has values before accessing them - if (Arg->getNumValues() > 0) { - llvm::errs() << " with value: " << Arg->getValue(); - } else { - llvm::errs() << " (no value)"; - } - - llvm::errs() << " hasFlag(WrapperOnlyOption): " << Arg->getOption().hasFlag(WrapperOnlyOption) - << "\n"; - - StringRef Spelling = Arg->getSpelling(); - if (Spelling.starts_with("--sycl-target-link-options")) { - continue; - } - // Do not forward arguments only intended for the linker wrapper. if (Arg->getOption().hasFlag(WrapperOnlyOption)) continue; @@ -1284,8 +1261,6 @@ Error copyFileToFinalExecutable(StringRef File, const ArgList &Args) { llvm::Triple Triple(Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple())); StringRef CopyCommand = Triple.isOSWindows() ? "copy" : "cp"; - llvm::errs() << "\"" << CopyCommand << "\" " << File << " " - << ExecutableName << "\n"; } // TODO: check if copy can be replaced by rename. if (std::error_code EC = sys::fs::copy_file(File, ExecutableName)) From 0452eb7ea88f9ecb15c056a16e167d37d249a5c5 Mon Sep 17 00:00:00 2001 From: y Date: Wed, 19 Nov 2025 20:10:17 -0800 Subject: [PATCH 04/11] code clean up --- clang/lib/Driver/ToolChains/Clang.cpp | 34 +------------------ .../ClangLinkerWrapper.cpp | 7 ++++ 2 files changed, 8 insertions(+), 33 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 958009ee6d301..cf581f58bd252 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11473,47 +11473,15 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, for (auto &ToolChainMember : llvm::make_range(ToolChainRange.first, ToolChainRange.second)) { const ToolChain *TC = ToolChainMember.second; - bool IsJIT = false; StringRef WrapperOption; - StringRef WrapperLinkOption; if (TC->getTriple().isSPIROrSPIRV()) { - if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) { - IsJIT = true; - WrapperOption = "--sycl-backend-compile-options="; - } if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) WrapperOption = "--gpu-tool-arg="; if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64) WrapperOption = "--cpu-tool-arg="; + Args.MakeArgString(Twine(WrapperOption)); } else continue; - // ArgStringList BuildArgs; - // SmallString<128> BackendOptString; - // SmallString<128> LinkOptString; - // SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); - // for (const auto &A : BuildArgs) - // appendOption(BackendOptString, A); - - // BuildArgs.clear(); - // SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); - // for (const auto &A : BuildArgs) { - // if (IsJIT) - // int x = 1; - // // appendOption(LinkOptString, A); - // else - // // For AOT, combine the Backend and Linker strings into one. - // appendOption(BackendOptString, A); - // } - // if (!BackendOptString.empty()) { - // llvm::errs() << "DEBUG: Adding argument: '" << BackendOptString << "'\n"; - // CmdArgs.push_back( - // Args.MakeArgString(Twine(WrapperOption) + BackendOptString)); - // } - // if (!LinkOptString.empty()){ - // llvm::errs() << "DEBUG: Adding argument: '" << LinkOptString << "'\n"; - // CmdArgs.push_back( - // Args.MakeArgString("--sycl-target-link-options=" + LinkOptString)); - // } } // Add option to enable creating of the .syclbin file. diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 16b11e494b26a..204d6375fd782 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -224,6 +224,11 @@ void printCommands(ArrayRef CmdArgs) { if (CmdArgs.empty()) return; + llvm::errs() << " \"" << CmdArgs.front() << "\" "; + for (auto IC = std::next(CmdArgs.begin()), IE = CmdArgs.end(); IC != IE; ++IC) + llvm::errs() << *IC << (std::next(IC) != IE ? " " : "\n"); +} + [[noreturn]] void reportError(Error E) { outs().flush(); logAllUnhandledErrors(std::move(E), @@ -1261,6 +1266,8 @@ Error copyFileToFinalExecutable(StringRef File, const ArgList &Args) { llvm::Triple Triple(Args.getLastArgValue(OPT_host_triple_EQ, sys::getDefaultTargetTriple())); StringRef CopyCommand = Triple.isOSWindows() ? "copy" : "cp"; + llvm::errs() << "\"" << CopyCommand << "\" " << File << " " + << ExecutableName << "\n"; } // TODO: check if copy can be replaced by rename. if (std::error_code EC = sys::fs::copy_file(File, ExecutableName)) From 71c475bc9fecb6c35f85569f12f1066f9325f601 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 25 Nov 2025 07:05:11 -0800 Subject: [PATCH 05/11] fix the CI failure --- clang/lib/Driver/ToolChains/Clang.cpp | 16 ++++++++++++++++ .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 4 +--- sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp | 2 +- 3 files changed, 18 insertions(+), 4 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index cf581f58bd252..bca722e07bcd2 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11473,8 +11473,12 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, for (auto &ToolChainMember : llvm::make_range(ToolChainRange.first, ToolChainRange.second)) { const ToolChain *TC = ToolChainMember.second; + bool IsJIT = false; StringRef WrapperOption; if (TC->getTriple().isSPIROrSPIRV()) { + if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) { + IsJIT = true; + } if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) WrapperOption = "--gpu-tool-arg="; if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64) @@ -11482,6 +11486,18 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, Args.MakeArgString(Twine(WrapperOption)); } else continue; + + ArgStringList BuildArgs; + SmallString<128> BackendOptString; + SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); + for (const auto &A : BuildArgs){ + appendOption(BackendOptString, A); + } + + if (!BackendOptString.empty() && !IsJIT){ + StringRef ArgString = Args.MakeArgString(Twine(WrapperOption) + BackendOptString); + CmdArgs.push_back(ArgString.data()); + } } // Add option to enable creating of the .syclbin file. diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 204d6375fd782..247c4ae3a222a 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -954,6 +954,7 @@ static Expected runLLVMToSPIRVTranslation(StringRef File, static void addSYCLBackendOptions(const ArgList &Args, SmallVector &CmdArgs, bool IsCPU, StringRef BackendOptions) { + llvm::errs() << "[DEBUG] ClangLinkerWrapper.cpp: BackendOptions " << BackendOptions << "\n"; if (IsCPU) { BackendOptions.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); } else { @@ -983,9 +984,6 @@ static void addSYCLBackendOptions(const ArgList &Args, } } - StringRef OptTool = (IsCPU) ? Args.getLastArgValue(OPT_cpu_tool_arg_EQ) - : Args.getLastArgValue(OPT_gpu_tool_arg_EQ); - OptTool.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); return; } diff --git a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp index 99dab3ed39d52..bfb31c7f7239b 100644 --- a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp +++ b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp @@ -1,7 +1,7 @@ // REQUIRES: ocloc, gpu, target-spir, !gpu-intel-gen12 // Test with `--offload-new-driver` // -// RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ +// RUN: %clangxx --no-offload-new-driver -fsycl -fsycl-device-code-split=per_source \ // RUN: -fsycl-targets=spir64_gen \ // RUN: -Xsycl-target-backend=spir64_gen \ // RUN: "-device dg2" -I %S/Inputs -o %t.out \ From 29931529c148574746352551d2a0e5fc92310464 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 25 Nov 2025 11:22:43 -0800 Subject: [PATCH 06/11] remove the backend option to be passed into clanglinkerwrapper, and add the link time option to device_compiler_args_EQ --- clang/lib/Driver/ToolChains/Clang.cpp | 42 +++++-------------- clang/test/Driver/clang-linker-wrapper.cpp | 3 +- clang/test/Driver/sycl-offload-new-driver.cpp | 17 +------- clang/test/Driver/sycl-offload.c | 9 ---- clang/test/Driver/sycl-offload.cpp | 10 ++--- .../ClangLinkerWrapper.cpp | 27 ++++++++---- .../clang-linker-wrapper/LinkerWrapperOpts.td | 10 ----- sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp | 2 +- 8 files changed, 40 insertions(+), 80 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index bca722e07bcd2..27195fa61bba1 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11460,47 +11460,27 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, CmdArgs.push_back( Args.MakeArgString("-sycl-allow-device-image-dependencies")); - // Formulate and add any offload-wrapper and AOT specific options. These - // are additional options passed in via -Xsycl-target-linker and - // -Xsycl-target-backend. + // For AOT, pass along backend target args via --device-compiler options + // to the clang-linker-wrapper. const toolchains::SYCLToolChain &SYCLTC = static_cast(getToolChain()); - // Only store compile/link opts in the image descriptor for the SPIR-V - // target. For AOT, pass along the addition options via GPU or CPU - // specific clang-linker-wrapper options. - const ArgList &Args = - C.getArgsForToolChain(nullptr, StringRef(), Action::OFK_SYCL); for (auto &ToolChainMember : llvm::make_range(ToolChainRange.first, ToolChainRange.second)) { const ToolChain *TC = ToolChainMember.second; - bool IsJIT = false; - StringRef WrapperOption; - if (TC->getTriple().isSPIROrSPIRV()) { - if (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) { - IsJIT = true; - } - if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) - WrapperOption = "--gpu-tool-arg="; - if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64) - WrapperOption = "--cpu-tool-arg="; - Args.MakeArgString(Twine(WrapperOption)); - } else - continue; - ArgStringList BuildArgs; - SmallString<128> BackendOptString; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); - for (const auto &A : BuildArgs){ - appendOption(BackendOptString, A); - } - - if (!BackendOptString.empty() && !IsJIT){ - StringRef ArgString = Args.MakeArgString(Twine(WrapperOption) + BackendOptString); - CmdArgs.push_back(ArgString.data()); + if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) { + SmallString<128> BackendOptString; + for (const auto &A : BuildArgs) + appendOption(BackendOptString, A); + CmdArgs.push_back(Args.MakeArgString( + "--device-compiler=" + TC->getTripleString() + "=" + BackendOptString)); } } - + // Add option to enable creating of the .syclbin file. + const ArgList &Args = + C.getArgsForToolChain(nullptr, StringRef(), Action::OFK_SYCL); if (Arg *A = Args.getLastArg(options::OPT_fsyclbin_EQ)) CmdArgs.push_back( Args.MakeArgString("--syclbin=" + StringRef{A->getValue()})); diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 96181605b3c7b..aa6b4b4be640b 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -74,7 +74,8 @@ // Check that when --gpu-tool-arg is specified in clang-linker-wrapper // (happen when AOT device is specified via -Xsycl-target-backend '-device pvc' in clang), // the target is not passed to sycl-post-link for filtering. -// RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--gpu-tool-arg=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s +// RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--device-compiler=spir64_gen-unknown-unknown=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s +// // CHK-NO-CMDS-AOT-GEN: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o {{[^,]*}}.table {{.*}}.bc /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU) diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index f449b3839602d..1f719722aaada 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -141,27 +141,14 @@ // MULT_TARG_PHASES: 15: backend, {14}, assembler, (host-sycl) // MULT_TARG_PHASES: 16: assembler, {15}, object, (host-sycl) -/// Test option passing behavior for clang-offload-wrapper options. -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ -// RUN: -Xsycl-target-backend -backend-opt -### %s 2>&1 \ -// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND %s -// WRAPPER_OPTIONS_BACKEND: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options={{.*}}-backend-opt{{.*}}" - -// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ -// RUN: -Xsycl-target-linker -link-opt -### %s 2>&1 \ -// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_LINK %s -// WRAPPER_OPTIONS_LINK: clang-linker-wrapper{{.*}} "--sycl-target-link-options={{.*}}-link-opt{{.*}}" - -/// Test option passing behavior for clang-offload-wrapper options for AOT. +/// Test option passing behavior for clang-offload-wrapper options for AOT GPU targets. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ // RUN: -fsycl-targets=spir64_gen,spir64_x86_64 \ // RUN: -Xsycl-target-backend=spir64_gen -backend-gen-opt \ // RUN: -Xsycl-target-backend=spir64_x86_64 -backend-cpu-opt \ // RUN: -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND_AOT %s -// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu" -// WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--gpu-tool-arg=-backend-gen-opt" -// WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--cpu-tool-arg=-backend-cpu-opt" +// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu" {{.*}} "--device-compiler=spir64_gen-unknown-unknown=-backend-gen-opt" /// Verify arch settings for nvptx and amdgcn targets // RUN: %clangxx -fsycl -### -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv \ diff --git a/clang/test/Driver/sycl-offload.c b/clang/test/Driver/sycl-offload.c index ce045accaced1..d51a97332fa30 100644 --- a/clang/test/Driver/sycl-offload.c +++ b/clang/test/Driver/sycl-offload.c @@ -325,11 +325,6 @@ // CHK-NO-FSYCL-TARGET-ERROR-NOT: clang{{.*}} error: cannot deduce implicit triple value for '-Xsycl-target-frontend', specify triple using '-Xsycl-target-frontend=' /// ########################################################################### - -// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-TOOLS-OPTS %s -// CHK-TOOLS-OPTS: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options=-DFOO1 -DFOO2" - /// Check for implied options (-g -O0) // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -g -O0 -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-OPTS %s @@ -346,10 +341,6 @@ // RUN: | FileCheck -check-prefix=CHK-TOOLS-IMPLIED-OPTS-O0 %s // CHK-TOOLS-IMPLIED-OPTS-O0-NOT: llvm-offload-binary{{.*}} {{.*}}compile-opts={{.*}}-cl-opt-disable" -// RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64-unknown-unknown -Xsycl-target-linker "-DFOO1 -DFOO2" %s 2>&1 \ -// RUN: | FileCheck -check-prefix=CHK-TOOLS-OPTS2 %s -// CHK-TOOLS-OPTS2: clang-linker-wrapper{{.*}} "--sycl-target-link-options=-DFOO1 -DFOO2" - /// -fsycl-range-rounding settings /// /// // Check that driver flag is passed to cc1 diff --git a/clang/test/Driver/sycl-offload.cpp b/clang/test/Driver/sycl-offload.cpp index ca3a1a1d9fb91..a1479b1a10c15 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -57,16 +57,16 @@ // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64 -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \ // RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT %s // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \ -// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT %s -// SYCL_TARGET_OPT: clang-linker-wrapper{{.*}} "--sycl-backend-compile-options=-DFOO" "--sycl-target-link-options=-DFOO2" +// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT --implicit-check-not="--sycl-backend-compile-options" --implicit-check-not="--sycl-target-link-options" %s +// SYCL_TARGET_OPT: clang-linker-wrapper{{.*}} // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64 -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ -// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU %s +// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU --implicit-check-not="--device-compiler" %s // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_gen -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ // RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_GPU %s // SYCL_TARGET_OPT_AOT-NOT: error: cannot deduce implicit triple value for '-Xsycl-target-backend' -// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} "--cpu-tool-arg=-DFOO" -// SYCL_TARGET_OPT_GPU: clang-linker-wrapper{{.*}} "--gpu-tool-arg=-DFOO" +// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} +// SYCL_TARGET_OPT_GPU: clang-linker-wrapper{{.*}} "--device-compiler=spir64_gen-unknown-unknown=-DFOO" /// Check -fsycl-targets=spir64 enables addition of -ffine-grained-bitfield-accesses option // RUN: %clangxx -### -fsycl-device-only --offload-new-driver %s 2>&1 | FileCheck -check-prefixes=CHECK_BITFIELD_OPTION %s diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 247c4ae3a222a..94bb8fd5d1873 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -716,7 +716,7 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args, /// code and will be parsed to generate options required to be passed into the /// sycl-post-link tool. static Expected> -runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { +runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, bool HasGPUTool) { Expected SYCLPostLinkPath = findProgram("sycl-post-link", {getMainExecutable("sycl-post-link")}); if (!SYCLPostLinkPath) @@ -731,14 +731,13 @@ runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args) { // Enable the driver to invoke sycl-post-link with the device architecture // when Intel GPU targets are passed in -fsycl-targets. - // OPT_gpu_tool_arg_EQ is checked to ensure the device architecture is not + // HasGPUTool is checked to ensure the device architecture is not // passed through -Xsycl-target-backend=spir64_gen "-device " format const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); StringRef Arch = Args.getLastArgValue(OPT_arch_EQ); - StringRef IsGPUTool = Args.getLastArgValue(OPT_gpu_tool_arg_EQ); if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && !Arch.empty() && - IsGPUTool.empty() && Arch != "*") + !HasGPUTool && Arch != "*") OutputPathWithArch = "intel_gpu_" + Arch.str() + "," + OutputPathWithArch; else if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64) OutputPathWithArch = "spir64_x86_64," + OutputPathWithArch; @@ -954,7 +953,6 @@ static Expected runLLVMToSPIRVTranslation(StringRef File, static void addSYCLBackendOptions(const ArgList &Args, SmallVector &CmdArgs, bool IsCPU, StringRef BackendOptions) { - llvm::errs() << "[DEBUG] ClangLinkerWrapper.cpp: BackendOptions " << BackendOptions << "\n"; if (IsCPU) { BackendOptions.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); } else { @@ -983,7 +981,6 @@ static void addSYCLBackendOptions(const ArgList &Args, CmdArgs.push_back(Args.MakeArgString(JoinedOptions)); } } - return; } @@ -2220,6 +2217,17 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, extractSYCLCompileLinkOptions(Input); if (!CompileLinkOptionsOrErr) return CompileLinkOptionsOrErr.takeError(); + + // Append any additional backend compiler options specified at link time for GPU + const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + for (StringRef Arg : LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) { + auto [ArgTriple, ArgValue] = Arg.split('='); + if (ArgTriple == LinkerArgs.getLastArgValue(OPT_triple_EQ) && + Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && + !ArgValue.empty()) { + CompileLinkOptionsOrErr->first += Twine(" ", ArgValue).str(); + } + } SmallVector InputFiles; // Write device inputs to an output file for the linker. @@ -2236,16 +2244,19 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, return TmpOutputOrErr.takeError(); SmallVector InputFilesSYCL; InputFilesSYCL.emplace_back(*TmpOutputOrErr); + + SmallVector Args; + StringRef(CompileLinkOptionsOrErr->first).split(Args, ' '); + bool HasGPUTool = std::find(Args.begin(), Args.end(), "-device") != Args.end(); auto SplitModulesOrErr = UseSYCLPostLinkTool - ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs) + ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs, HasGPUTool) : sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, *SYCLModuleSplitMode); if (!SplitModulesOrErr) return SplitModulesOrErr.takeError(); auto &SplitModules = *SplitModulesOrErr; - const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); bool IsJIT = Triple.isSPIROrSPIRV() && Triple.getSubArch() == llvm::Triple::NoSubArch; if ((Triple.isNVPTX() || Triple.isAMDGCN()) && diff --git a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td index 5d3f2eb19c8e8..2780ddd326bf2 100644 --- a/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td +++ b/clang/tools/clang-linker-wrapper/LinkerWrapperOpts.td @@ -199,16 +199,6 @@ def sycl_add_default_spec_consts_image : Flag<["--", "-"], "sycl-add-default-spe def no_sycl_add_default_spec_consts_image : Flag<["--", "-"], "no-sycl-add-default-spec-consts-image">, Flags<[WrapperOnlyOption, HelpHidden]>; -// Special options to pass backend options required for AOT compilation -def gpu_tool_arg_EQ : - Joined<["--", "-"], "gpu-tool-arg=">, - Flags<[WrapperOnlyOption]>, - HelpText<"Options that are passed to the backend of target device compiler for Intel GPU during AOT compilation">; -def cpu_tool_arg_EQ : - Joined<["--", "-"], "cpu-tool-arg=">, - Flags<[WrapperOnlyOption]>, - HelpText<"Options that are passed to the backend of target device compiler for Intel CPU during AOT compilation">; - def sycl_thin_lto : Flag<["--", "-"], "sycl-thin-lto">, Flags<[WrapperOnlyOption]>, HelpText<"Link SYCL device code using thinLTO">; diff --git a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp index bfb31c7f7239b..99dab3ed39d52 100644 --- a/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp +++ b/sycl/test-e2e/NewOffloadDriver/aot-gpu.cpp @@ -1,7 +1,7 @@ // REQUIRES: ocloc, gpu, target-spir, !gpu-intel-gen12 // Test with `--offload-new-driver` // -// RUN: %clangxx --no-offload-new-driver -fsycl -fsycl-device-code-split=per_source \ +// RUN: %clangxx -fsycl -fsycl-device-code-split=per_source \ // RUN: -fsycl-targets=spir64_gen \ // RUN: -Xsycl-target-backend=spir64_gen \ // RUN: "-device dg2" -I %S/Inputs -o %t.out \ From 7d0ee38c0a101635863732241dcdeacd86c9543c Mon Sep 17 00:00:00 2001 From: Yixing Zhang Date: Mon, 1 Dec 2025 22:55:16 -0800 Subject: [PATCH 07/11] add passing the link time compiler option for cpu as well --- clang/lib/Driver/ToolChains/Clang.cpp | 12 ++---------- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 4 +--- 2 files changed, 3 insertions(+), 13 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 27195fa61bba1..4534b42614fa4 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11248,14 +11248,6 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX())) LinkerArgs.emplace_back("-lompdevice"); - // Forward all of these to the appropriate toolchain. - for (StringRef Arg : CompilerArgs) - CmdArgs.push_back(Args.MakeArgString( - "--device-compiler=" + TC->getTripleString() + "=" + Arg)); - for (StringRef Arg : LinkerArgs) - CmdArgs.push_back(Args.MakeArgString( - "--device-linker=" + TC->getTripleString() + "=" + Arg)); - // Forward the LTO mode relying on the Driver's parsing. if (C.getDriver().getOffloadLTOMode() == LTOK_Full) CmdArgs.push_back(Args.MakeArgString( @@ -11469,8 +11461,8 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, const ToolChain *TC = ToolChainMember.second; ArgStringList BuildArgs; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); - if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) { - SmallString<128> BackendOptString; + SmallString<128> BackendOptString; + if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen || (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64)) { for (const auto &A : BuildArgs) appendOption(BackendOptString, A); CmdArgs.push_back(Args.MakeArgString( diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 94bb8fd5d1873..3f42d6f81b4fb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2222,9 +2222,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); for (StringRef Arg : LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) { auto [ArgTriple, ArgValue] = Arg.split('='); - if (ArgTriple == LinkerArgs.getLastArgValue(OPT_triple_EQ) && - Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen && - !ArgValue.empty()) { + if (ArgTriple == Triple.getTriple() && !ArgValue.empty()) { CompileLinkOptionsOrErr->first += Twine(" ", ArgValue).str(); } } From e88a3fb1a72503bc57860fa36f5977ff6cd0df5a Mon Sep 17 00:00:00 2001 From: y Date: Tue, 2 Dec 2025 07:42:59 -0800 Subject: [PATCH 08/11] fix driver test failure --- clang/lib/Driver/ToolChains/Clang.cpp | 10 ++++++++++ clang/test/Driver/sycl-offload-new-driver.cpp | 4 ++-- clang/test/Driver/sycl-offload.cpp | 12 ++---------- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 2 +- 4 files changed, 15 insertions(+), 13 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 4534b42614fa4..2406452f24beb 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11247,6 +11247,16 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Kind == Action::OFK_OpenMP && !Args.hasArg(OPT_no_offloadlib) && (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX())) LinkerArgs.emplace_back("-lompdevice"); + + // Forward all of these to the appropriate toolchain. + if (!C.hasOffloadToolChain()) { + for (StringRef Arg : CompilerArgs) + CmdArgs.push_back(Args.MakeArgString( + "--device-compiler=" + TC->getTripleString() + "=" + Arg)); + for (StringRef Arg : LinkerArgs) + CmdArgs.push_back(Args.MakeArgString( + "--device-linker=" + TC->getTripleString() + "=" + Arg)); + } // Forward the LTO mode relying on the Driver's parsing. if (C.getDriver().getOffloadLTOMode() == LTOK_Full) diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 1f719722aaada..3ba3cbd49f9be 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -141,14 +141,14 @@ // MULT_TARG_PHASES: 15: backend, {14}, assembler, (host-sycl) // MULT_TARG_PHASES: 16: assembler, {15}, object, (host-sycl) -/// Test option passing behavior for clang-offload-wrapper options for AOT GPU targets. +/// Test option passing behavior for clang-offload-wrapper options for AOT // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ // RUN: -fsycl-targets=spir64_gen,spir64_x86_64 \ // RUN: -Xsycl-target-backend=spir64_gen -backend-gen-opt \ // RUN: -Xsycl-target-backend=spir64_x86_64 -backend-cpu-opt \ // RUN: -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_BACKEND_AOT %s -// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu" {{.*}} "--device-compiler=spir64_gen-unknown-unknown=-backend-gen-opt" +// WRAPPER_OPTIONS_BACKEND_AOT: clang-linker-wrapper{{.*}} "--host-triple=x86_64-unknown-linux-gnu" {{.*}} "--device-compiler=spir64_gen-unknown-unknown=-backend-gen-opt" "--device-compiler=spir64_x86_64-unknown-unknown=-backend-cpu-opt" /// Verify arch settings for nvptx and amdgcn targets // RUN: %clangxx -fsycl -### -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv \ diff --git a/clang/test/Driver/sycl-offload.cpp b/clang/test/Driver/sycl-offload.cpp index a1479b1a10c15..07ce0caed80a5 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -52,20 +52,12 @@ // RUN: | FileCheck -check-prefix DEFAULT_LINK %s // DEFAULT_LINK: clang-linker-wrapper{{.*}} -/// Passing in the default triple should allow for -Xsycl-target options, both the -/// "=" and the default spelling -// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64 -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \ -// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT %s -// RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -Xsycl-target-backend=spir64 -DFOO -Xsycl-target-linker=spir64 -DFOO2 %S/Inputs/SYCL/objlin64.o 2>&1 \ -// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT --implicit-check-not="--sycl-backend-compile-options" --implicit-check-not="--sycl-target-link-options" %s -// SYCL_TARGET_OPT: clang-linker-wrapper{{.*}} - // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_x86_64 -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ -// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU --implicit-check-not="--device-compiler" %s +// RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_CPU %s // RUN: %clangxx -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver -fsycl-targets=spir64_gen -Xsycl-target-backend -DFOO %S/Inputs/SYCL/objlin64.o 2>&1 \ // RUN: | FileCheck -check-prefixes=SYCL_TARGET_OPT_AOT,SYCL_TARGET_OPT_GPU %s // SYCL_TARGET_OPT_AOT-NOT: error: cannot deduce implicit triple value for '-Xsycl-target-backend' -// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} +// SYCL_TARGET_OPT_CPU: clang-linker-wrapper{{.*}} "--device-compiler=spir64_x86_64-unknown-unknown=-DFOO" // SYCL_TARGET_OPT_GPU: clang-linker-wrapper{{.*}} "--device-compiler=spir64_gen-unknown-unknown=-DFOO" /// Check -fsycl-targets=spir64 enables addition of -ffine-grained-bitfield-accesses option diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 3f42d6f81b4fb..5b8a49759f0a6 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2218,7 +2218,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, if (!CompileLinkOptionsOrErr) return CompileLinkOptionsOrErr.takeError(); - // Append any additional backend compiler options specified at link time for GPU + // Append any additional backend compiler options specified at link time. const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); for (StringRef Arg : LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) { auto [ArgTriple, ArgValue] = Arg.split('='); From f4730f1ee76eb66c1bbc6f8a19b12aa6b5e177ad Mon Sep 17 00:00:00 2001 From: y Date: Tue, 2 Dec 2025 08:39:31 -0800 Subject: [PATCH 09/11] code clean up --- clang/include/clang/Driver/Driver.h | 7 ---- clang/lib/Driver/Driver.cpp | 39 +++---------------- clang/lib/Driver/ToolChains/Clang.cpp | 6 +-- clang/test/Driver/sycl-offload-new-driver.cpp | 2 +- 4 files changed, 9 insertions(+), 45 deletions(-) diff --git a/clang/include/clang/Driver/Driver.h b/clang/include/clang/Driver/Driver.h index 126f127de8674..3dd8822c4844b 100644 --- a/clang/include/clang/Driver/Driver.h +++ b/clang/include/clang/Driver/Driver.h @@ -626,13 +626,6 @@ class Driver { /// @name Helper Methods /// @{ - /// Utility function to parse all devices passed via -fsycl-targets. - /// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets. - /// Otherwise return 'false'. - bool - GetUseNewOffloadDriverForSYCLOffload(Compilation &C, - const llvm::opt::ArgList &Args) const; - /// getSYCLDeviceTriple - Returns the SYCL device triple for the /// specified subarch // TODO: Additional Arg input parameter is for diagnostic output information diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 9e039ef3724c5..bfa6893dbd9b0 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1563,33 +1563,7 @@ static void appendOneArg(InputArgList &Args, const Arg *Opt) { AliasCopy->claim(); Copy->setAlias(std::move(AliasCopy)); } -} - -// Utility function to parse all devices passed via -fsycl-targets. -// Return 'true' for JIT, AOT Intel CPU/GPUs and NVidia/AMD targets. -// Otherwise return 'false'. -bool Driver::GetUseNewOffloadDriverForSYCLOffload(Compilation &C, - const ArgList &Args) const { - // Check only if enabled with -fsycl - if (!Args.hasFlag(options::OPT_fsycl, options::OPT_fno_sycl, false)) - return false; - - if (Args.hasFlag(options::OPT_no_offload_new_driver, - options::OPT_offload_new_driver, false)) - return false; - - if (Args.hasArg(options::OPT_fintelfpga)) - return false; - - if (const Arg *A = Args.getLastArg(options::OPT_fsycl_targets_EQ)) { - for (const char *Val : A->getValues()) { - llvm::Triple TT(C.getDriver().getSYCLDeviceTriple(Val)); - if ((!TT.isSPIROrSPIRV()) || TT.isSPIRAOT()) - return false; - } - } - return true; -} +} bool Driver::readConfigFile(StringRef FileName, llvm::cl::ExpansionContext &ExpCtx) { @@ -2221,12 +2195,12 @@ Compilation *Driver::BuildCompilation(ArrayRef ArgList) { // Use new offloading path for OpenMP. This is disabled as the SYCL // offloading path is not properly setup to use the updated device linking // scheme. - if (C->isOffloadingHostKind(Action::OFK_OpenMP) || + if ((C->isOffloadingHostKind(Action::OFK_OpenMP) && + TranslatedArgs->hasFlag(options::OPT_fopenmp_new_driver, + options::OPT_no_offload_new_driver, true)) || TranslatedArgs->hasFlag(options::OPT_offload_new_driver, - options::OPT_no_offload_new_driver, false) || - GetUseNewOffloadDriverForSYCLOffload(*C, *TranslatedArgs)) { + options::OPT_no_offload_new_driver, false)) setUseNewOffloadingDriver(); - } // Construct the list of abstract actions to perform for this compilation. On // MachO targets this uses the driver-driver and universal actions. @@ -7121,8 +7095,7 @@ void Driver::BuildDefaultActions(Compilation &C, DerivedArgList &Args, options::OPT_fno_offload_via_llvm, false) || Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, - C.isOffloadingHostKind(Action::OFK_Cuda)) || - GetUseNewOffloadDriverForSYCLOffload(C, Args); + C.isOffloadingHostKind(Action::OFK_Cuda)); bool HIPNoRDC = C.isOffloadingHostKind(Action::OFK_HIP) && diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 2406452f24beb..b899022d564ff 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -5232,9 +5232,7 @@ void Clang::ConstructJob(Compilation &C, const JobAction &JA, (JA.isHostOffloading(C.getActiveOffloadKinds()) && Args.hasFlag(options::OPT_offload_new_driver, options::OPT_no_offload_new_driver, - C.isOffloadingHostKind(Action::OFK_Cuda))) || - (JA.isHostOffloading(Action::OFK_SYCL) && - C.getDriver().GetUseNewOffloadDriverForSYCLOffload(C, Args)); + C.isOffloadingHostKind(Action::OFK_Cuda))); bool IsRDCMode = Args.hasFlag(options::OPT_fgpu_rdc, options::OPT_fno_gpu_rdc, IsSYCL); @@ -11247,7 +11245,7 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Kind == Action::OFK_OpenMP && !Args.hasArg(OPT_no_offloadlib) && (TC->getTriple().isAMDGPU() || TC->getTriple().isNVPTX())) LinkerArgs.emplace_back("-lompdevice"); - + // Forward all of these to the appropriate toolchain. if (!C.hasOffloadToolChain()) { for (StringRef Arg : CompilerArgs) diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 3ba3cbd49f9be..2be5f3966ec7f 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -141,7 +141,7 @@ // MULT_TARG_PHASES: 15: backend, {14}, assembler, (host-sycl) // MULT_TARG_PHASES: 16: assembler, {15}, object, (host-sycl) -/// Test option passing behavior for clang-offload-wrapper options for AOT +/// Test option passing behavior for clang-offload-wrapper options for AOT. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver \ // RUN: -fsycl-targets=spir64_gen,spir64_x86_64 \ // RUN: -Xsycl-target-backend=spir64_gen -backend-gen-opt \ From 2f11978e31ab1bfb3fd901c03472d68f4d312c63 Mon Sep 17 00:00:00 2001 From: y Date: Tue, 2 Dec 2025 09:11:34 -0800 Subject: [PATCH 10/11] code clean up --- clang/lib/Driver/Driver.cpp | 2 +- clang/lib/Driver/ToolChains/Clang.cpp | 10 +++++++--- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 14 +++++++++----- 3 files changed, 17 insertions(+), 9 deletions(-) diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index bfa6893dbd9b0..475c1fddb0363 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -1563,7 +1563,7 @@ static void appendOneArg(InputArgList &Args, const Arg *Opt) { AliasCopy->claim(); Copy->setAlias(std::move(AliasCopy)); } -} +} bool Driver::readConfigFile(StringRef FileName, llvm::cl::ExpansionContext &ExpCtx) { diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index b899022d564ff..7512a87502dc8 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11247,6 +11247,8 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, LinkerArgs.emplace_back("-lompdevice"); // Forward all of these to the appropriate toolchain. + // SYCL offload toolchains handle device compiler and linker argument forwarding using SYCLImage. + // Therefore, we skip forwarding these arguments here when a SYCL offload toolchain is present. if (!C.hasOffloadToolChain()) { for (StringRef Arg : CompilerArgs) CmdArgs.push_back(Args.MakeArgString( @@ -11470,11 +11472,13 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, ArgStringList BuildArgs; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); SmallString<128> BackendOptString; - if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen || (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64)) { + if (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen || + (TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_x86_64)) { for (const auto &A : BuildArgs) appendOption(BackendOptString, A); - CmdArgs.push_back(Args.MakeArgString( - "--device-compiler=" + TC->getTripleString() + "=" + BackendOptString)); + CmdArgs.push_back( + Args.MakeArgString("--device-compiler=" + TC->getTripleString() + + "=" + BackendOptString)); } } diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 5b8a49759f0a6..b7d69747dee01 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -716,7 +716,8 @@ getTripleBasedSYCLPostLinkOpts(const ArgList &Args, /// code and will be parsed to generate options required to be passed into the /// sycl-post-link tool. static Expected> -runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, bool HasGPUTool) { +runSYCLPostLinkTool(ArrayRef InputFiles, const ArgList &Args, + bool HasGPUTool) { Expected SYCLPostLinkPath = findProgram("sycl-post-link", {getMainExecutable("sycl-post-link")}); if (!SYCLPostLinkPath) @@ -2217,10 +2218,11 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, extractSYCLCompileLinkOptions(Input); if (!CompileLinkOptionsOrErr) return CompileLinkOptionsOrErr.takeError(); - + // Append any additional backend compiler options specified at link time. const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); - for (StringRef Arg : LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) { + for (StringRef Arg : + LinkerArgs.getAllArgValues(OPT_device_compiler_args_EQ)) { auto [ArgTriple, ArgValue] = Arg.split('='); if (ArgTriple == Triple.getTriple() && !ArgValue.empty()) { CompileLinkOptionsOrErr->first += Twine(" ", ArgValue).str(); @@ -2245,10 +2247,12 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, SmallVector Args; StringRef(CompileLinkOptionsOrErr->first).split(Args, ' '); - bool HasGPUTool = std::find(Args.begin(), Args.end(), "-device") != Args.end(); + bool HasGPUTool = + std::find(Args.begin(), Args.end(), "-device") != Args.end(); auto SplitModulesOrErr = UseSYCLPostLinkTool - ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs, HasGPUTool) + ? sycl::runSYCLPostLinkTool(InputFilesSYCL, LinkerArgs, + HasGPUTool) : sycl::runSYCLSplitLibrary(InputFilesSYCL, LinkerArgs, *SYCLModuleSplitMode); if (!SplitModulesOrErr) From fe390ac353c65ec85ccd2b3bf7883a1efca55a5e Mon Sep 17 00:00:00 2001 From: y Date: Tue, 2 Dec 2025 09:17:45 -0800 Subject: [PATCH 11/11] code clean up --- clang/test/Driver/clang-linker-wrapper.cpp | 1 - 1 file changed, 1 deletion(-) diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index aa6b4b4be640b..de1c2c60f75ea 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -75,7 +75,6 @@ // (happen when AOT device is specified via -Xsycl-target-backend '-device pvc' in clang), // the target is not passed to sycl-post-link for filtering. // RUN: clang-linker-wrapper -sycl-embed-ir -sycl-device-libraries=%t1.devicelib.o -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" -llvm-spirv-options="LLVM_SPIRV_OPTIONS" "--host-triple=x86_64-unknown-linux-gnu" "--device-compiler=spir64_gen-unknown-unknown=-device pvc" "--linker-path=/usr/bin/ld" "--" HOST_LINKER_FLAGS "-dynamic-linker" HOST_DYN_LIB "-o" "a.out" HOST_LIB_PATH HOST_STAT_LIB %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN %s -// // CHK-NO-CMDS-AOT-GEN: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o {{[^,]*}}.table {{.*}}.bc /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU)