From f301dc6c275ad26563756d43211ff5839f183aa7 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 10:11:28 -0700 Subject: [PATCH 1/9] [SYCL] Forward AOT backend/linker options as individual tokens Route -Xsycl-target-backend and -Xsycl-target-linker options through --device-compiler=/--device-linker= uniformly for JIT and AOT SYCL targets, one flag per token instead of joining tokens into a single string. clang-linker-wrapper now keeps CLI-supplied AOT options as a token list end-to-end and appends them to ocloc/opencl-aot argv individually, avoiding a lossy re-split of values with embedded spaces. JIT targets keep the existing string-joining behavior for the image's compile-opts/link-opts properties. --- clang/lib/Driver/ToolChains/Clang.cpp | 33 +--- clang/test/Driver/clang-linker-wrapper.cpp | 20 +- clang/test/Driver/sycl-offload-new-driver.cpp | 12 ++ clang/test/Driver/sycl-offload.cpp | 4 +- .../ClangLinkerWrapper.cpp | 174 +++++++++++++----- 5 files changed, 168 insertions(+), 75 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 5c2b183337e43..69d71e9722c69 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12008,7 +12008,6 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, Args); SYCLInstallation.getSYCLDeviceLibPath(LibLocCandidates); SmallString<128> LibName("libsycl-crt.bc"); - bool IsNewOffload = D.getUseNewOffloadingDriver(); for (const auto &LibLoc : LibLocCandidates) { SmallString<128> FullLibName(LibLoc); llvm::sys::path::append(FullLibName, LibName); @@ -12028,12 +12027,6 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // targets (NVPTX, AMD) pass device libraries to clang-linker-wrapper. SmallVector BCLibList; - auto appendToList = [](SmallString<256> &List, const Twine &Arg) { - if (List.size() > 0) - List += ","; - List += Arg.str(); - }; - auto ToolChainRange = C.getOffloadToolChains(); for (const auto &[Kind, TC] : llvm::make_range(ToolChainRange.first, ToolChainRange.second)) { @@ -12190,33 +12183,19 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (!TC->getTriple().isSPIROrSPIRV()) 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 (TC->getTriple().getSubArch() == llvm::Triple::NoSubArch) - 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( "--device-compiler=" + Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + - TC->getTripleString() + "=" + BackendOptString)); - } - if (!LinkOptString.empty()) { + TC->getTripleString() + "=" + A)); + + BuildArgs.clear(); + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); + for (const auto &A : BuildArgs) CmdArgs.push_back(Args.MakeArgString( "--device-linker=" + Action::GetOffloadKindName(Action::OFK_SYCL) + - ":" + TC->getTripleString() + "=" + LinkOptString)); - } + ":" + TC->getTripleString() + "=" + A)); } // Add option to enable creating of the .syclbin file. diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index bb0600c8d6d49..0cd30caae9a35 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -132,6 +132,17 @@ // RUN: clang-linker-wrapper -sycl-embed-ir --bitcode-library=spir64_gen-unknown-unknown=%t1.devicelib.bc -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 that when "-device pvc" specification, split across two separate +// --device-linker= arguments, is reconstructed correctly and still suppresses +// the sycl-post-link target filtering prefix. +// RUN: clang-linker-wrapper -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" --device-linker=spir64_gen-unknown-unknown=-device --device-linker=spir64_gen-unknown-unknown=pvc --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN-LINKERARG %s +// CHK-NO-CMDS-AOT-GEN-LINKERARG: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o {{[^,]*}}.table {{.*}}.bc + +// Check that multiple distinct --device-linker= values for the same AOT +// target all reach the ocloc invocation. +// RUN: clang-linker-wrapper --device-linker=spir64_gen-unknown-unknown=-foo-opt --device-linker=spir64_gen-unknown-unknown=-bar-opt --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-MULTI-DEVICE-LINKER-AOT %s +// CHK-MULTI-DEVICE-LINKER-AOT: ocloc{{.*}} -foo-opt -bar-opt{{.*}} + /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU) // ------- // Generate .o file as linker wrapper input. @@ -329,9 +340,10 @@ // RUN: --bitcode-library=spir64_gen-unknown-unknown=%t1.devicelib.bc \ // RUN: %t.aot.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s --check-prefix=CHECK-COMPILE-LINK-OPTS-AOT // -// Check that in AOT case backend options are passed to ocloc and are not passed to offload wrapper -// because SYCL Runtime can't make any use of it in AOT case. -// CHECK-COMPILE-LINK-OPTS-AOT: ocloc{{.*}} -device pvc ccc ccc -output +// Check that in AOT case backend and linker options are passed to ocloc and +// are not passed to offload wrapper because SYCL runtime can't make any use of +// it in AOT case. +// CHECK-COMPILE-LINK-OPTS-AOT: ocloc{{.*}} -device pvc ccc ccc ddd ddd -output // CHECK-COMPILE-LINK-OPTS-AOT: offload-wrapper: {{.*}} compile-opts: , link-opts: // TODO: The following check is turned off because clang-linker-wrapper doesn't support @@ -343,7 +355,7 @@ // rUN: %t.jit_and_aot.o -o %t.out 2>&1 --linker-path="/usr/bin/ld" | FileCheck %s --check-prefix=CHECK-COMPILE-LINK-OPTS-JIT-AND-AOT // cHECK-COMPILE-LINK-OPTS-JIT-AND-AOT: offload-wrapper: {{.*}} compile-opts: aaa aaa, link-opts: bbb bbb -// cHECK-COMPILE-LINK-OPTS-JIT-AND-AOT: ocloc{{.*}} -device pvc ccc ccc +// cHECK-COMPILE-LINK-OPTS-JIT-AND-AOT: ocloc{{.*}} -device pvc ccc ccc ddd ddd // cHECK-COMPILE-LINK-OPTS-JIT-AND-AOT: offload-wrapper: {{.*}} compile-opts: , link-opts: // Check that missed triple in image causes an error. diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 628221630313c..b617653bb2742 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -155,6 +155,18 @@ // WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-backend-gen-opt" // WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-backend-cpu-opt" +/// Test that -Xsycl-target-backend and -Xsycl-target-linker options for an +/// AOT (ocloc) target are forwarded via --device-compiler=/--device-linker= +/// respectively, each token as its own argument, the same as for JIT +/// targets. +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ +// RUN: -fsycl-targets=intel_gpu_pvc \ +// RUN: -Xsycl-target-backend -opt1 -Xsycl-target-linker -opt2 \ +// RUN: -### %s 2>&1 \ +// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_AOT_SEPARATE %s +// WRAPPER_OPTIONS_AOT_SEPARATE: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-opt1" +// WRAPPER_OPTIONS_AOT_SEPARATE-SAME: "--device-linker=sycl:spir64_gen-unknown-unknown=-opt2" + /// Verify arch settings for nvptx and amdgcn targets // RUN: %clangxx -fsycl -### -fsycl-targets=amdgcn-amd-amdhsa -fno-sycl-libspirv \ // RUN: -nocudalib --offload-new-driver --sysroot=%S/Inputs/SYCL \ diff --git a/clang/test/Driver/sycl-offload.cpp b/clang/test/Driver/sycl-offload.cpp index ddf623ddbf37a..240a62e7de905 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -328,7 +328,7 @@ // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -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{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-DFOO1 -DFOO2" +// CHK-TOOLS-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-DFOO1" "--device-compiler=sycl:spir64-unknown-unknown=-DFOO2" /// Check for implied options (-g -O0) // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -fsycl-targets=spir64-unknown-unknown -g -O0 -Xsycl-target-backend "-DFOO1 -DFOO2" %s 2>&1 \ @@ -348,7 +348,7 @@ // RUN: %clang -### -target x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL -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{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-DFOO1 -DFOO2" +// CHK-TOOLS-OPTS2: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-DFOO1" "--device-linker=sycl:spir64-unknown-unknown=-DFOO2" /// -fsycl-range-rounding settings /// diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index ca1de7603b10e..1837e03b684f9 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1000,9 +1000,15 @@ static Expected runLLVMToSPIRVTranslation(StringRef File, /// ("-options") in the first member of the pair, and everything after the /// separator in the second part of the pair. The separator is not included in /// any of them. -/// \p BackendOptions is a string containing backend compilation options. For -/// example, "-options -cl-opt-disable". +/// \p BackendOptions is a string containing backend compilation options +/// extracted from the device image (e.g. "-options -cl-opt-disable"). +/// \p AOTDeviceArgs are additional options supplied on the clang-linker- +/// wrapper command line via --device-compiler=/--device-linker=; each is +/// already an individual token and is appended to \p CmdArgs verbatim, +/// without being merged into \p BackendOptions and re-split, so that tokens +/// containing embedded spaces are preserved intact. static void addOclocOptions(StringRef BackendOptions, + ArrayRef AOTDeviceArgs, SmallVector &CmdArgs) { auto [BeforeOptions, AfterOptions] = BackendOptions.split("-options "); BeforeOptions.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); @@ -1011,6 +1017,7 @@ static void addOclocOptions(StringRef BackendOptions, CmdArgs.push_back("-options"); CmdArgs.push_back(AfterOptions); } + llvm::append_range(CmdArgs, AOTDeviceArgs); } /// Run AOT compilation for Intel CPU. @@ -1018,11 +1025,15 @@ static void addOclocOptions(StringRef BackendOptions, /// \p InputFile is the input SPIR-V file. /// \p Args encompasses all arguments required for linking and wrapping device /// code. -/// \p BackendOptions is a string containing backend compilation options. For -/// example, "-options -cl-opt-disable". -static Expected runAOTCompileIntelCPU(StringRef InputFile, - const ArgList &Args, - StringRef BackendOptions) { +/// \p BackendOptions is a string containing backend compilation options +/// extracted from the device image. For example, "-options -cl-opt-disable". +/// \p AOTDeviceArgs are additional individual option tokens supplied on the +/// clang-linker-wrapper command line, appended verbatim as separate argv +/// entries. +static Expected +runAOTCompileIntelCPU(StringRef InputFile, const ArgList &Args, + StringRef BackendOptions, + ArrayRef AOTDeviceArgs) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); SmallVector CmdArgs; Expected OpenCLAOTPath = @@ -1033,6 +1044,7 @@ static Expected runAOTCompileIntelCPU(StringRef InputFile, CmdArgs.push_back(*OpenCLAOTPath); CmdArgs.push_back("--device=cpu"); BackendOptions.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); + llvm::append_range(CmdArgs, AOTDeviceArgs); // Create a new file to write the translated file to. auto TempFileOrErr = createOutputFile(sys::path::filename(ExecutableName), "out"); @@ -1051,11 +1063,15 @@ static Expected runAOTCompileIntelCPU(StringRef InputFile, /// \p InputFile is the input SPIR-V file. /// \p Args encompasses all arguments required for linking and wrapping device /// code. -/// \p BackendOptions is a string containing backend compilation options. For -/// example, "-options -cl-opt-disable". -static Expected runAOTCompileIntelGPU(StringRef InputFile, - const ArgList &Args, - StringRef BackendOptions) { +/// \p BackendOptions is a string containing backend compilation options +/// extracted from the device image. For example, "-options -cl-opt-disable". +/// \p AOTDeviceArgs are additional individual option tokens supplied on the +/// clang-linker-wrapper command line, appended verbatim as separate argv +/// entries (see addOclocOptions). +static Expected +runAOTCompileIntelGPU(StringRef InputFile, const ArgList &Args, + StringRef BackendOptions, + ArrayRef AOTDeviceArgs) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); StringRef Arch(Args.getLastArgValue(OPT_arch_EQ)); SmallVector CmdArgs; @@ -1072,7 +1088,7 @@ static Expected runAOTCompileIntelGPU(StringRef InputFile, CmdArgs.push_back("-device"); CmdArgs.push_back(Arch); } - addOclocOptions(BackendOptions, CmdArgs); + addOclocOptions(BackendOptions, AOTDeviceArgs, CmdArgs); // Create a new file to write the translated file to. auto TempFileOrErr = createOutputFile(sys::path::filename(ExecutableName), "out"); @@ -1091,17 +1107,23 @@ static Expected runAOTCompileIntelGPU(StringRef InputFile, /// \p InputFile is the input SPIR-V file. /// \p Args encompasses all arguments required for linking and wrapping device /// code. -/// \p BackendOptions is a string containing backend compilation options. For -/// example, "-options -cl-opt-disable". +/// \p BackendOptions is a string containing backend compilation options +/// extracted from the device image. For example, "-options -cl-opt-disable". +/// \p AOTDeviceArgs are additional individual option tokens supplied on the +/// clang-linker-wrapper command line via --device-compiler=/ +/// --device-linker=. static Expected runAOTCompile(StringRef InputFile, const ArgList &Args, - StringRef BackendOptions) { + StringRef BackendOptions, + ArrayRef AOTDeviceArgs) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); if (Triple.isSPIRAOT()) { if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen) - return runAOTCompileIntelGPU(InputFile, Args, BackendOptions); + return runAOTCompileIntelGPU(InputFile, Args, BackendOptions, + AOTDeviceArgs); if (Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64) - return runAOTCompileIntelCPU(InputFile, Args, BackendOptions); + return runAOTCompileIntelCPU(InputFile, Args, BackendOptions, + AOTDeviceArgs); } return createStringError(inconvertibleErrorCode(), "Unsupported SYCL Triple and Arch"); @@ -1820,7 +1842,8 @@ namespace sycl { /// compilation (Intel CPU/GPU). Expected invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, - StringRef SYCLBackendOptions = StringRef()) { + StringRef SYCLBackendOptions = StringRef(), + ArrayRef AOTDeviceArgs = {}) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); switch (Triple.getArch()) { case Triple::nvptx: @@ -1848,9 +1871,10 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen || Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64); Expected AOTFile = - (NeedAOTCompile) - ? sycl::runAOTCompile(*SPVFile, Args, SYCLBackendOptions) - : *SPVFile; + (NeedAOTCompile) ? sycl::runAOTCompile(*SPVFile, Args, + SYCLBackendOptions, + AOTDeviceArgs) + : *SPVFile; if (!AOTFile) return AOTFile.takeError(); return NeedAOTCompile ? *AOTFile : *SPVFile; @@ -1868,9 +1892,10 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, Expected compileDeviceAndBundle(StringRef ModuleFilePath, const ArgList &LinkerArgs, const llvm::Triple &Triple, - StringRef AdditionalCompileOptions) { + StringRef AdditionalCompileOptions, + ArrayRef AOTDeviceArgs = {}) { Expected OutputOrErr = invokeBackendForSYCLDevice( - ModuleFilePath, LinkerArgs, AdditionalCompileOptions); + ModuleFilePath, LinkerArgs, AdditionalCompileOptions, AOTDeviceArgs); if (!OutputOrErr) return OutputOrErr.takeError(); @@ -1899,9 +1924,44 @@ Expected compileDeviceAndBundle(StringRef ModuleFilePath, Expected> postLinkProcessModule( StringRef ModuleFilePath, const ArgList &LinkerArgs, const std::pair &CompileLinkOptions, + ArrayRef AOTDeviceArgs, function_ref WrappedOutputCallback) { + const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + + // AOT tools (ocloc/opencl-aot) don't distinguish compile vs. link options, + // so combine both here regardless of whether an option arrived via + // --device-compiler= or --device-linker=. CompileLinkOptions holds only + // the options extracted from the image for AOT triples (a flat, already + // space-joined string); options supplied on the CLI for this invocation + // live separately in AOTDeviceArgs, as individual tokens, and are appended + // later (in runAOTCompileIntelGPU/CPU) without being folded into this + // string, so that values with embedded spaces survive intact. + std::string AOTOptions; + if (Triple.isSPIRAOT()) { + AOTOptions = CompileLinkOptions.first; + if (!CompileLinkOptions.second.empty()) { + if (!AOTOptions.empty()) + AOTOptions += ' '; + AOTOptions += CompileLinkOptions.second; + } + } + StringRef BackendOptions = Triple.isSPIRAOT() + ? StringRef(AOTOptions) + : StringRef(CompileLinkOptions.first); + + // Detect whether the user already specified "-device " via + // -Xsycl-target-backend, either embedded in the image string above or + // supplied on the CLI (AOTDeviceArgs). Split every source on spaces for + // this search (a single CLI token can still contain an embedded space, + // e.g. when clang-linker-wrapper is invoked directly rather than through + // the driver's tokenizing -Xsycl-target-backend handling); this does not + // affect how AOTDeviceArgs are forwarded to the backend tool's argv below, + // which keeps each token intact. StringRef::split() appends to its output + // vector, so accumulate tokens from both sources before searching. SmallVector CompileArgsSplit; - StringRef(CompileLinkOptions.first).split(CompileArgsSplit, ' '); + BackendOptions.split(CompileArgsSplit, ' '); + for (StringRef Arg : AOTDeviceArgs) + Arg.split(CompileArgsSplit, ' '); bool IsDevicePassedWithSyclTargetBackend = std::find(CompileArgsSplit.begin(), CompileArgsSplit.end(), "-device") != CompileArgsSplit.end(); @@ -1917,7 +1977,6 @@ Expected> postLinkProcessModule( return SplitModulesOrErr.takeError(); std::vector &SplitModules = *SplitModulesOrErr; - const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); if ((Triple.isNVPTX() || Triple.isAMDGCN()) && LinkerArgs.hasArg(OPT_sycl_embed_ir)) { // When compiling for Nvidia/AMD devices and the user requested the @@ -1948,7 +2007,7 @@ Expected> postLinkProcessModule( for (size_t I = 0, E = SplitModules.size(); I != E; ++I) { Expected OutputOrErr = compileDeviceAndBundle(SplitModules[I].ModuleFilePath, LinkerArgs, - Triple, CompileLinkOptions.first); + Triple, BackendOptions, AOTDeviceArgs); if (!OutputOrErr) return OutputOrErr.takeError(); @@ -1976,6 +2035,7 @@ Expected> postLinkProcessModule( Expected> runSYCLOffloadingPipeline( ArrayRef InputModules, const ArgList &LinkerArgs, const std::pair &CompileLinkOptions, + ArrayRef AOTDeviceArgs, function_ref WrappedOutputCallback) { // Note: pipeline can skip linking due to -fno-sycl-rdc option. // In that case, we apply sycl processing to several modules. @@ -2000,7 +2060,7 @@ Expected> runSYCLOffloadingPipeline( // split. Expected> ModulesOrErr = postLinkProcessModule(Module, LinkerArgs, CompileLinkOptions, - WrappedOutputCallback); + AOTDeviceArgs, WrappedOutputCallback); if (!ModulesOrErr) return ModulesOrErr.takeError(); @@ -2465,19 +2525,49 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, *CompileLinkOptionsOrErr; // Append device compiler and linker options passed via - // -device-compiler= and -device-linker= to clang-linker-warpper, - // together with options extracted from the image. - StringRef DeviceCompilerArgs = - LinkerArgs.getLastArgValue(OPT_compiler_arg_EQ); - if (!DeviceCompilerArgs.empty()) { - CompileLinkOptions.first += " "; - CompileLinkOptions.first += DeviceCompilerArgs; - } - StringRef DeviceLinkerArgs = - LinkerArgs.getLastArgValue(OPT_linker_arg_EQ); - if (!DeviceLinkerArgs.empty()) { - CompileLinkOptions.second += " "; - CompileLinkOptions.second += DeviceLinkerArgs; + // --device-compiler= and --device-linker= to clang-linker-wrapper. + // Each occurrence of --device-compiler=/--device-linker= that matched + // this triple/kind was forwarded as its own compiler-arg=/linker-arg= + // by getLinkerArgs(). + // + // JIT targets: the SYCL runtime consumes these as flat strings, so + // join them (after the options already extracted from the image) into + // CompileLinkOptions. + // + // AOT targets (ocloc/opencl-aot): keep the CLI-supplied tokens as a + // separate list (AOTDeviceArgs) rather than folding them into + // CompileLinkOptions, so that a token containing an embedded space + // isn't re-split downstream. They are appended to the AOT tool's argv + // as individual entries, alongside the image-embedded options (which + // remain a flat, space-tokenized string, unrelated to this list). + const llvm::Triple TargetTriple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + std::vector AOTDeviceArgs; + if (TargetTriple.isSPIRAOT()) { + for (std::string &DeviceCompilerArg : + LinkerArgs.getAllArgValues(OPT_compiler_arg_EQ)) + if (!DeviceCompilerArg.empty()) + AOTDeviceArgs.push_back(std::move(DeviceCompilerArg)); + for (std::string &DeviceLinkerArg : + LinkerArgs.getAllArgValues(OPT_linker_arg_EQ)) + if (!DeviceLinkerArg.empty()) + AOTDeviceArgs.push_back(std::move(DeviceLinkerArg)); + } else { + for (const std::string &DeviceCompilerArg : + LinkerArgs.getAllArgValues(OPT_compiler_arg_EQ)) { + if (DeviceCompilerArg.empty()) + continue; + if (!CompileLinkOptions.first.empty()) + CompileLinkOptions.first += " "; + CompileLinkOptions.first += DeviceCompilerArg; + } + for (const std::string &DeviceLinkerArg : + LinkerArgs.getAllArgValues(OPT_linker_arg_EQ)) { + if (DeviceLinkerArg.empty()) + continue; + if (!CompileLinkOptions.second.empty()) + CompileLinkOptions.second += " "; + CompileLinkOptions.second += DeviceLinkerArg; + } } SmallVector InputFiles; @@ -2491,7 +2581,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, Expected> ModulesOrErr = sycl::runSYCLOffloadingPipeline(InputFiles, LinkerArgs, - CompileLinkOptions, + CompileLinkOptions, AOTDeviceArgs, AppendImageToWrapperOutput); if (!ModulesOrErr) return ModulesOrErr.takeError(); From 131a6716a9a29c51db3a76b8259b89474682d4b5 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 10:24:04 -0700 Subject: [PATCH 2/9] clang-format --- clang/lib/Driver/ToolChains/Clang.cpp | 8 +++---- .../ClangLinkerWrapper.cpp | 21 ++++++++++--------- 2 files changed, 15 insertions(+), 14 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 69d71e9722c69..528292ad523b9 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12185,10 +12185,10 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, ArgStringList BuildArgs; SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); for (const auto &A : BuildArgs) - CmdArgs.push_back(Args.MakeArgString( - "--device-compiler=" + - Action::GetOffloadKindName(Action::OFK_SYCL) + ":" + - TC->getTripleString() + "=" + A)); + CmdArgs.push_back( + Args.MakeArgString("--device-compiler=" + + Action::GetOffloadKindName(Action::OFK_SYCL) + + ":" + TC->getTripleString() + "=" + A)); BuildArgs.clear(); SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 1837e03b684f9..6f1fb68a957a2 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1871,10 +1871,10 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, (Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen || Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64); Expected AOTFile = - (NeedAOTCompile) ? sycl::runAOTCompile(*SPVFile, Args, - SYCLBackendOptions, - AOTDeviceArgs) - : *SPVFile; + (NeedAOTCompile) + ? sycl::runAOTCompile(*SPVFile, Args, SYCLBackendOptions, + AOTDeviceArgs) + : *SPVFile; if (!AOTFile) return AOTFile.takeError(); return NeedAOTCompile ? *AOTFile : *SPVFile; @@ -1889,11 +1889,11 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, } } -Expected compileDeviceAndBundle(StringRef ModuleFilePath, - const ArgList &LinkerArgs, - const llvm::Triple &Triple, - StringRef AdditionalCompileOptions, - ArrayRef AOTDeviceArgs = {}) { +Expected +compileDeviceAndBundle(StringRef ModuleFilePath, const ArgList &LinkerArgs, + const llvm::Triple &Triple, + StringRef AdditionalCompileOptions, + ArrayRef AOTDeviceArgs = {}) { Expected OutputOrErr = invokeBackendForSYCLDevice( ModuleFilePath, LinkerArgs, AdditionalCompileOptions, AOTDeviceArgs); if (!OutputOrErr) @@ -2540,7 +2540,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // isn't re-split downstream. They are appended to the AOT tool's argv // as individual entries, alongside the image-embedded options (which // remain a flat, space-tokenized string, unrelated to this list). - const llvm::Triple TargetTriple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + const llvm::Triple TargetTriple( + LinkerArgs.getLastArgValue(OPT_triple_EQ)); std::vector AOTDeviceArgs; if (TargetTriple.isSPIRAOT()) { for (std::string &DeviceCompilerArg : From f2b5c59f9f7e88754b7698fe29f3fd5d3ab4c69b Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 15:46:18 -0700 Subject: [PATCH 3/9] [SYCL] Fix stale AOT device-token tests in clang-linker-wrapper.cpp Update the CHK-NO-CMDS-AOT-GEN-LINKERARG test to assert on the actual ocloc invocation instead of an unrelated sycl-post-link line, matching the new one-CLI-token-equals-one-argv-entry contract for --device-compiler=/--device-linker=. Remove CHK-NO-CMDS-AOT-GEN (relied on a single-token "-device pvc" CLI value that is no longer split, and is unreachable via the real driver since -Xsycl-target-backend already tokenizes it into separate flags) and CHK-MULTI-DEVICE-LINKER-AOT (duplicated the same reconstruction behavior as the fixed test). --- clang/test/Driver/clang-linker-wrapper.cpp | 22 +++++++--------------- 1 file changed, 7 insertions(+), 15 deletions(-) diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 0cd30caae9a35..c8b63c10bde91 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -126,22 +126,14 @@ // CHK-CMDS-AOT-GEN-NEXT: clang{{.*}} -c -o [[LLCOUT:.*]].o [[WRAPPEROUT]].bc // CHK-CMDS-AOT-GEN-NEXT: "{{.*}}/ld" -- HOST_LINKER_FLAGS -dynamic-linker HOST_DYN_LIB -o a.out [[LLCOUT]].o HOST_LIB_PATH HOST_STAT_LIB {{.*}}.o -// Check that when "--device-compiler=triple=-device pvc" 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 --bitcode-library=spir64_gen-unknown-unknown=%t1.devicelib.bc -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 that when "-device pvc" specification, split across two separate -// --device-linker= arguments, is reconstructed correctly and still suppresses -// the sycl-post-link target filtering prefix. -// RUN: clang-linker-wrapper -sycl-post-link-options="SYCL_POST_LINK_OPTIONS" --device-linker=spir64_gen-unknown-unknown=-device --device-linker=spir64_gen-unknown-unknown=pvc --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN-LINKERARG %s -// CHK-NO-CMDS-AOT-GEN-LINKERARG: sycl-post-link{{.*}} SYCL_POST_LINK_OPTIONS -o {{[^,]*}}.table {{.*}}.bc - -// Check that multiple distinct --device-linker= values for the same AOT -// target all reach the ocloc invocation. -// RUN: clang-linker-wrapper --device-linker=spir64_gen-unknown-unknown=-foo-opt --device-linker=spir64_gen-unknown-unknown=-bar-opt --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-MULTI-DEVICE-LINKER-AOT %s -// CHK-MULTI-DEVICE-LINKER-AOT: ocloc{{.*}} -foo-opt -bar-opt{{.*}} +// --device-linker= arguments (each --device-compiler=/--device-linker= CLI +// occurrence becomes exactly one ocloc argv entry, so a multi-token value +// like "-device pvc" must be supplied as separate occurrences, not as one +// value containing an embedded space), is reconstructed correctly in the +// ocloc invocation. +// RUN: clang-linker-wrapper --device-linker=spir64_gen-unknown-unknown=-device --device-linker=spir64_gen-unknown-unknown=pvc --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-NO-CMDS-AOT-GEN-LINKERARG %s +// CHK-NO-CMDS-AOT-GEN-LINKERARG: ocloc{{.*}} -device pvc -output /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU) // ------- From 7f190ca26ad1cee69ca2d1ac40f53e30b3049d8b Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 15:54:40 -0700 Subject: [PATCH 4/9] [SYCL] Add FIXME on ocloc option string parsing in clang-linker-wrapper MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit -Xsycl-target-backend/-Xsycl-target-linker require the driver and clang-linker-wrapper to parse and re-serialize ocloc's own option syntax (e.g. the literal "-options " split in addOclocOptions), which is fragile — appended link-opts can be silently absorbed into the -options value. Note the deprecation/removal of these options as a follow-up to avoid this class of parsing entirely. --- .../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 6f1fb68a957a2..87ebc8c329841 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1007,6 +1007,16 @@ static Expected runLLVMToSPIRVTranslation(StringRef File, /// already an individual token and is appended to \p CmdArgs verbatim, /// without being merged into \p BackendOptions and re-split, so that tokens /// containing embedded spaces are preserved intact. +// FIXME: This literal-substring split on "-options " is inherently fragile +// (e.g. link-opts appended after a compile-opts "-options ..." blob get +// silently absorbed into the -options value). The root issue is that +// -Xsycl-target-backend/-Xsycl-target-linker require the driver and this +// tool to parse and re-serialize ocloc's own option syntax. Consider +// deprecating and removing -Xsycl-target-backend/-Xsycl-target-linker (and +// the compile-opts=/link-opts= image encoding they feed) in favor of a +// mechanism that forwards backend/linker options as opaque tokens end to +// end, so neither the driver nor clang-linker-wrapper needs to understand +// ocloc's option grammar. static void addOclocOptions(StringRef BackendOptions, ArrayRef AOTDeviceArgs, SmallVector &CmdArgs) { @@ -1936,6 +1946,14 @@ Expected> postLinkProcessModule( // live separately in AOTDeviceArgs, as individual tokens, and are appended // later (in runAOTCompileIntelGPU/CPU) without being folded into this // string, so that values with embedded spaces survive intact. + // FIXME: Concatenating compile-opts and link-opts into one flat string + // here means any "-options ..." wrapper already present in compile-opts + // (see SYCLToolChain::AddSPIRVImpliedTargetArgs) will swallow the + // appended link-opts into ocloc's -options value in addOclocOptions + // below. This is a symptom of -Xsycl-target-backend/-Xsycl-target-linker + // requiring us to parse and re-serialize ocloc's option syntax; consider + // deprecating and removing those options so backend/linker options can be + // forwarded as opaque tokens without this kind of string surgery. std::string AOTOptions; if (Triple.isSPIRAOT()) { AOTOptions = CompileLinkOptions.first; From 9b434effda32aa416b619f6fa3e93f7cb6f1eff6 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 16:06:13 -0700 Subject: [PATCH 5/9] [SYCL] Note driver-side counterpart to linker-wrapper -device detection Add a comment cross-referencing Driver::getOffloadArchs, which performs a similar "-device" token detection on the driver side for spir64_gen Arch selection, so a future edit to either detector's handling of -Xsycl-target-backend syntax prompts a check of the other. --- .../tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 87ebc8c329841..08b8f584b6c2f 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1976,6 +1976,17 @@ Expected> postLinkProcessModule( // affect how AOTDeviceArgs are forwarded to the backend tool's argv below, // which keeps each token intact. StringRef::split() appends to its output // vector, so accumulate tokens from both sources before searching. + // + // NOTE: this mirrors a separate "-device" detector on the driver side + // (Driver::getOffloadArchs in clang/lib/Driver/Driver.cpp, which scans the + // tokenized -Xsycl-target-backend argv for spir64_gen to pick the Arch + // forwarded to this tool). The two detectors operate on different data + // shapes (pre- vs. post-serialization) and aren't required to literally + // match today, but if -Xsycl-target-backend's "-device" syntax ever grows + // (e.g. comma-separated multi-device lists, already anticipated in a + // comment there) only one of the two may get updated, which can desync + // the Arch this tool receives from whether it thinks "-device" was + // explicit. SmallVector CompileArgsSplit; BackendOptions.split(CompileArgsSplit, ' '); for (StringRef Arg : AOTDeviceArgs) From bd41c5b63826dd296e76bf62c89cdb83d2dd28cb Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 16:14:56 -0700 Subject: [PATCH 6/9] [SYCL] Hoist per-triple AOT option computation out of postLinkProcessModule Triple parsing, AOTOptions/BackendOptions derivation, and the -device detection in IsDevicePassedWithSyclTargetBackend were recomputed on every call to postLinkProcessModule, even though the value is identical across all split modules sharing a triple. Moved the computation up to runSYCLOffloadingPipeline, which now computes it once before looping over Modules and passes the results down. --- .../ClangLinkerWrapper.cpp | 133 ++++++++++-------- 1 file changed, 71 insertions(+), 62 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 08b8f584b6c2f..315ec2b58046b 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1930,71 +1930,19 @@ compileDeviceAndBundle(StringRef ModuleFilePath, const ArgList &LinkerArgs, /// * Set Compile/Link options to the output Modules. /// * Invokes device backend compilation + bundling. /// +/// \p Triple, \p BackendOptions and \p IsDevicePassedWithSyclTargetBackend +/// are derived from \p CompileLinkOptions/\p AOTDeviceArgs once per triple by +/// the caller (they are identical across every module for a given triple), +/// rather than being recomputed here on each call. +/// /// \returns The list of the processed Modules. Expected> postLinkProcessModule( StringRef ModuleFilePath, const ArgList &LinkerArgs, + const llvm::Triple &Triple, const std::pair &CompileLinkOptions, - ArrayRef AOTDeviceArgs, + ArrayRef AOTDeviceArgs, StringRef BackendOptions, + bool IsDevicePassedWithSyclTargetBackend, function_ref WrappedOutputCallback) { - const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); - - // AOT tools (ocloc/opencl-aot) don't distinguish compile vs. link options, - // so combine both here regardless of whether an option arrived via - // --device-compiler= or --device-linker=. CompileLinkOptions holds only - // the options extracted from the image for AOT triples (a flat, already - // space-joined string); options supplied on the CLI for this invocation - // live separately in AOTDeviceArgs, as individual tokens, and are appended - // later (in runAOTCompileIntelGPU/CPU) without being folded into this - // string, so that values with embedded spaces survive intact. - // FIXME: Concatenating compile-opts and link-opts into one flat string - // here means any "-options ..." wrapper already present in compile-opts - // (see SYCLToolChain::AddSPIRVImpliedTargetArgs) will swallow the - // appended link-opts into ocloc's -options value in addOclocOptions - // below. This is a symptom of -Xsycl-target-backend/-Xsycl-target-linker - // requiring us to parse and re-serialize ocloc's option syntax; consider - // deprecating and removing those options so backend/linker options can be - // forwarded as opaque tokens without this kind of string surgery. - std::string AOTOptions; - if (Triple.isSPIRAOT()) { - AOTOptions = CompileLinkOptions.first; - if (!CompileLinkOptions.second.empty()) { - if (!AOTOptions.empty()) - AOTOptions += ' '; - AOTOptions += CompileLinkOptions.second; - } - } - StringRef BackendOptions = Triple.isSPIRAOT() - ? StringRef(AOTOptions) - : StringRef(CompileLinkOptions.first); - - // Detect whether the user already specified "-device " via - // -Xsycl-target-backend, either embedded in the image string above or - // supplied on the CLI (AOTDeviceArgs). Split every source on spaces for - // this search (a single CLI token can still contain an embedded space, - // e.g. when clang-linker-wrapper is invoked directly rather than through - // the driver's tokenizing -Xsycl-target-backend handling); this does not - // affect how AOTDeviceArgs are forwarded to the backend tool's argv below, - // which keeps each token intact. StringRef::split() appends to its output - // vector, so accumulate tokens from both sources before searching. - // - // NOTE: this mirrors a separate "-device" detector on the driver side - // (Driver::getOffloadArchs in clang/lib/Driver/Driver.cpp, which scans the - // tokenized -Xsycl-target-backend argv for spir64_gen to pick the Arch - // forwarded to this tool). The two detectors operate on different data - // shapes (pre- vs. post-serialization) and aren't required to literally - // match today, but if -Xsycl-target-backend's "-device" syntax ever grows - // (e.g. comma-separated multi-device lists, already anticipated in a - // comment there) only one of the two may get updated, which can desync - // the Arch this tool receives from whether it thinks "-device" was - // explicit. - SmallVector CompileArgsSplit; - BackendOptions.split(CompileArgsSplit, ' '); - for (StringRef Arg : AOTDeviceArgs) - Arg.split(CompileArgsSplit, ' '); - bool IsDevicePassedWithSyclTargetBackend = - std::find(CompileArgsSplit.begin(), CompileArgsSplit.end(), "-device") != - CompileArgsSplit.end(); - SmallVector InputFilesSYCL = {ModuleFilePath}; Expected> SplitModulesOrErr = UseSYCLPostLinkTool @@ -2066,6 +2014,65 @@ Expected> runSYCLOffloadingPipeline( const std::pair &CompileLinkOptions, ArrayRef AOTDeviceArgs, function_ref WrappedOutputCallback) { + const llvm::Triple Triple(LinkerArgs.getLastArgValue(OPT_triple_EQ)); + + // AOT tools (ocloc/opencl-aot) don't distinguish compile vs. link options, + // so combine both here regardless of whether an option arrived via + // --device-compiler= or --device-linker=. CompileLinkOptions holds only + // the options extracted from the image for AOT triples (a flat, already + // space-joined string); options supplied on the CLI for this invocation + // live separately in AOTDeviceArgs, as individual tokens, and are appended + // later (in runAOTCompileIntelGPU/CPU) without being folded into this + // string, so that values with embedded spaces survive intact. + // FIXME: Concatenating compile-opts and link-opts into one flat string + // here means any "-options ..." wrapper already present in compile-opts + // (see SYCLToolChain::AddSPIRVImpliedTargetArgs) will swallow the + // appended link-opts into ocloc's -options value in addOclocOptions + // below. This is a symptom of -Xsycl-target-backend/-Xsycl-target-linker + // requiring us to parse and re-serialize ocloc's option syntax; consider + // deprecating and removing those options so backend/linker options can be + // forwarded as opaque tokens without this kind of string surgery. + std::string AOTOptions; + if (Triple.isSPIRAOT()) { + AOTOptions = CompileLinkOptions.first; + if (!CompileLinkOptions.second.empty()) { + if (!AOTOptions.empty()) + AOTOptions += ' '; + AOTOptions += CompileLinkOptions.second; + } + } + StringRef BackendOptions = Triple.isSPIRAOT() + ? StringRef(AOTOptions) + : StringRef(CompileLinkOptions.first); + + // Detect whether the user already specified "-device " via + // -Xsycl-target-backend, either embedded in the image string above or + // supplied on the CLI (AOTDeviceArgs). Split every source on spaces for + // this search (a single CLI token can still contain an embedded space, + // e.g. when clang-linker-wrapper is invoked directly rather than through + // the driver's tokenizing -Xsycl-target-backend handling); this does not + // affect how AOTDeviceArgs are forwarded to the backend tool's argv below, + // which keeps each token intact. StringRef::split() appends to its output + // vector, so accumulate tokens from both sources before searching. + // + // NOTE: this mirrors a separate "-device" detector on the driver side + // (Driver::getOffloadArchs in clang/lib/Driver/Driver.cpp, which scans the + // tokenized -Xsycl-target-backend argv for spir64_gen to pick the Arch + // forwarded to this tool). The two detectors operate on different data + // shapes (pre- vs. post-serialization) and aren't required to literally + // match today, but if -Xsycl-target-backend's "-device" syntax ever grows + // (e.g. comma-separated multi-device lists, already anticipated in a + // comment there) only one of the two may get updated, which can desync + // the Arch this tool receives from whether it thinks "-device" was + // explicit. + SmallVector CompileArgsSplit; + BackendOptions.split(CompileArgsSplit, ' '); + for (StringRef Arg : AOTDeviceArgs) + Arg.split(CompileArgsSplit, ' '); + bool IsDevicePassedWithSyclTargetBackend = + std::find(CompileArgsSplit.begin(), CompileArgsSplit.end(), "-device") != + CompileArgsSplit.end(); + // Note: pipeline can skip linking due to -fno-sycl-rdc option. // In that case, we apply sycl processing to several modules. std::vector Modules; @@ -2088,8 +2095,10 @@ Expected> runSYCLOffloadingPipeline( // Note: sycl-post-link can produce more modules than incoming due to module // split. Expected> ModulesOrErr = - postLinkProcessModule(Module, LinkerArgs, CompileLinkOptions, - AOTDeviceArgs, WrappedOutputCallback); + postLinkProcessModule(Module, LinkerArgs, Triple, CompileLinkOptions, + AOTDeviceArgs, BackendOptions, + IsDevicePassedWithSyclTargetBackend, + WrappedOutputCallback); if (!ModulesOrErr) return ModulesOrErr.takeError(); From 6bbda3f4dccf69a4b9a4013ebf759be873ce0d9c Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 16:18:41 -0700 Subject: [PATCH 7/9] [SYCL] Drop dead default args on AOTDeviceArgs parameters invokeBackendForSYCLDevice and compileDeviceAndBundle are single-caller helpers whose only call sites already pass AOTDeviceArgs explicitly, so the `= {}` defaults were unreachable and misleadingly implied some caller omits the argument. --- clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 315ec2b58046b..7e9eea2e19aeb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1852,8 +1852,8 @@ namespace sycl { /// compilation (Intel CPU/GPU). Expected invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, - StringRef SYCLBackendOptions = StringRef(), - ArrayRef AOTDeviceArgs = {}) { + StringRef SYCLBackendOptions, + ArrayRef AOTDeviceArgs) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); switch (Triple.getArch()) { case Triple::nvptx: @@ -1903,7 +1903,7 @@ Expected compileDeviceAndBundle(StringRef ModuleFilePath, const ArgList &LinkerArgs, const llvm::Triple &Triple, StringRef AdditionalCompileOptions, - ArrayRef AOTDeviceArgs = {}) { + ArrayRef AOTDeviceArgs) { Expected OutputOrErr = invokeBackendForSYCLDevice( ModuleFilePath, LinkerArgs, AdditionalCompileOptions, AOTDeviceArgs); if (!OutputOrErr) From ba12722273d3273a1fee8c276eaa03671d614419 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 16:58:32 -0700 Subject: [PATCH 8/9] [SYCL] Add FIXME on image-embedded compile/link option space-joining BuildArgs is joined into a single space-separated compile-opts=/link-opts= string before being embedded in the offload image, then later re-split on spaces by clang-linker-wrapper for AOT triples. This corrupts any token containing an embedded space unless it happens to fall inside the pre-existing "-options \"...\"" wrapper convention. The CLI-supplied counterpart (--device-compiler=/--device-linker=) already avoids this by forwarding individual tokens via AOTDeviceArgs. --- clang/lib/Driver/ToolChains/Clang.cpp | 22 ++++++++++++++++++++++ 1 file changed, 22 insertions(+) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 528292ad523b9..8768d20519491 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10522,6 +10522,17 @@ static void addRunTimeWrapperOpts(Compilation &C, // Grab any Target specific options that need to be added to the wrapper // information. ArgStringList BuildArgs; + // FIXME: BuildArgs is joined with spaces below into a single + // -compile-opts=/-link-opts= string embedded in the image descriptor, then + // later re-split on spaces by clang-linker-wrapper for AOT triples (see + // addOclocOptions/runAOTCompileIntelCPU in + // clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp). Any individual + // token containing an embedded space is corrupted by this round trip + // unless it happens to fall inside the pre-existing "-options \"...\"" + // wrapper convention. Forwarding these as individual tokens instead of a + // joined string (as is already done for the CLI-supplied + // --device-compiler=/--device-linker= counterpart via AOTDeviceArgs) would + // avoid this. auto createArgString = [&](const char *Opt) { if (BuildArgs.empty()) return; @@ -10840,6 +10851,17 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, // compilers and the clang-offload-wrapper in the case of SYCL offloading. if (OffloadAction->getOffloadingDeviceKind() == Action::OFK_SYCL) { ArgStringList BuildArgs; + // FIXME: BuildArgs is joined with spaces into a single compile-opts=/ + // link-opts= string below, then later re-split on spaces by + // clang-linker-wrapper (see addOclocOptions/runAOTCompileIntelCPU in + // clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp). Any + // individual token containing an embedded space (e.g. from + // -Xsycl-target-backend "-abc 'multi word'") is corrupted by this + // round trip unless it happens to fall inside the pre-existing + // "-options \"...\"" wrapper convention. The CLI-supplied counterpart + // of these options (--device-compiler=/--device-linker=) avoids this + // by forwarding individual tokens (AOTDeviceArgs); the image-embedded + // path here still does not. auto createArgString = [&](const char *Opt) { if (BuildArgs.empty()) return; From e82f1eab402746b3dab954d5fa241da9136811e1 Mon Sep 17 00:00:00 2001 From: Alexey Bader Date: Wed, 15 Jul 2026 17:03:12 -0700 Subject: [PATCH 9/9] [SYCL] Note duplicated isSPIRAOT() branch between AOT option channels linkAndWrapDeviceFiles() and runSYCLOffloadingPipeline() each decide independently whether CLI-supplied AOT options go through the token vector or the flat-string path, with nothing enforcing the two isSPIRAOT() checks stay in agreement. --- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 12 ++++++++++++ 1 file changed, 12 insertions(+) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 7e9eea2e19aeb..acabc3576d17f 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2033,6 +2033,11 @@ Expected> runSYCLOffloadingPipeline( // deprecating and removing those options so backend/linker options can be // forwarded as opaque tokens without this kind of string surgery. std::string AOTOptions; + // TODO: This isSPIRAOT() check must stay in sync with the equivalent check + // in linkAndWrapDeviceFiles() above, which decides whether CLI-supplied + // options for this triple go into AOTDeviceArgs (token vector) or get + // folded into CompileLinkOptions (flat string). Nothing enforces + // agreement between the two; see the TODO there for the suggested fix. if (Triple.isSPIRAOT()) { AOTOptions = CompileLinkOptions.first; if (!CompileLinkOptions.second.empty()) { @@ -2581,6 +2586,13 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, const llvm::Triple TargetTriple( LinkerArgs.getLastArgValue(OPT_triple_EQ)); std::vector AOTDeviceArgs; + // TODO: This isSPIRAOT() check must stay in sync with the equivalent + // check in runSYCLOffloadingPipeline() below, which decides whether to + // fold CompileLinkOptions into AOTOptions for the same triple. Nothing + // enforces agreement between the two; consider unifying the option + // representation (e.g. making CompileLinkOptions itself a token + // vector) so both AOT option origins share one code path instead of + // two independently-maintained branches. if (TargetTriple.isSPIRAOT()) { for (std::string &DeviceCompilerArg : LinkerArgs.getAllArgValues(OPT_compiler_arg_EQ))