diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 5c2b183337e43..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; @@ -12008,7 +12030,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 +12049,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 +12205,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); + CmdArgs.push_back( + Args.MakeArgString("--device-compiler=" + + Action::GetOffloadKindName(Action::OFK_SYCL) + + ":" + TC->getTripleString() + "=" + 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()) { + 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..c8b63c10bde91 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -126,11 +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 (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) // ------- @@ -329,9 +332,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 +347,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..acabc3576d17f 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1000,9 +1000,25 @@ 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. +// 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) { auto [BeforeOptions, AfterOptions] = BackendOptions.split("-options "); BeforeOptions.split(CmdArgs, " ", /*MaxSplit=*/-1, /*KeepEmpty=*/false); @@ -1011,6 +1027,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 +1035,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 +1054,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 +1073,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 +1098,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 +1117,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 +1852,8 @@ namespace sycl { /// compilation (Intel CPU/GPU). Expected invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, - StringRef SYCLBackendOptions = StringRef()) { + StringRef SYCLBackendOptions, + ArrayRef AOTDeviceArgs) { const llvm::Triple Triple(Args.getLastArgValue(OPT_triple_EQ)); switch (Triple.getArch()) { case Triple::nvptx: @@ -1849,7 +1882,8 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, Triple.getSubArch() == llvm::Triple::SPIRSubArch_x86_64); Expected AOTFile = (NeedAOTCompile) - ? sycl::runAOTCompile(*SPVFile, Args, SYCLBackendOptions) + ? sycl::runAOTCompile(*SPVFile, Args, SYCLBackendOptions, + AOTDeviceArgs) : *SPVFile; if (!AOTFile) return AOTFile.takeError(); @@ -1865,12 +1899,13 @@ invokeBackendForSYCLDevice(StringRef InputFile, const ArgList &Args, } } -Expected compileDeviceAndBundle(StringRef ModuleFilePath, - const ArgList &LinkerArgs, - const llvm::Triple &Triple, - StringRef AdditionalCompileOptions) { +Expected +compileDeviceAndBundle(StringRef ModuleFilePath, const ArgList &LinkerArgs, + const llvm::Triple &Triple, + StringRef AdditionalCompileOptions, + ArrayRef AOTDeviceArgs) { Expected OutputOrErr = invokeBackendForSYCLDevice( - ModuleFilePath, LinkerArgs, AdditionalCompileOptions); + ModuleFilePath, LinkerArgs, AdditionalCompileOptions, AOTDeviceArgs); if (!OutputOrErr) return OutputOrErr.takeError(); @@ -1895,17 +1930,19 @@ Expected compileDeviceAndBundle(StringRef ModuleFilePath, /// * 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, StringRef BackendOptions, + bool IsDevicePassedWithSyclTargetBackend, function_ref WrappedOutputCallback) { - SmallVector CompileArgsSplit; - StringRef(CompileLinkOptions.first).split(CompileArgsSplit, ' '); - bool IsDevicePassedWithSyclTargetBackend = - std::find(CompileArgsSplit.begin(), CompileArgsSplit.end(), "-device") != - CompileArgsSplit.end(); - SmallVector InputFilesSYCL = {ModuleFilePath}; Expected> SplitModulesOrErr = UseSYCLPostLinkTool @@ -1917,7 +1954,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 +1984,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,7 +2012,72 @@ Expected> postLinkProcessModule( Expected> runSYCLOffloadingPipeline( ArrayRef InputModules, 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. + // 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; + // 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()) { + 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; @@ -1999,7 +2100,9 @@ Expected> runSYCLOffloadingPipeline( // Note: sycl-post-link can produce more modules than incoming due to module // split. Expected> ModulesOrErr = - postLinkProcessModule(Module, LinkerArgs, CompileLinkOptions, + postLinkProcessModule(Module, LinkerArgs, Triple, CompileLinkOptions, + AOTDeviceArgs, BackendOptions, + IsDevicePassedWithSyclTargetBackend, WrappedOutputCallback); if (!ModulesOrErr) return ModulesOrErr.takeError(); @@ -2465,19 +2568,57 @@ 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; + // 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)) + 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 +2632,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, Expected> ModulesOrErr = sycl::runSYCLOffloadingPipeline(InputFiles, LinkerArgs, - CompileLinkOptions, + CompileLinkOptions, AOTDeviceArgs, AppendImageToWrapperOutput); if (!ModulesOrErr) return ModulesOrErr.takeError();