From 3850c239de3269f3f474b45c1266130818faaa98 Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Fri, 14 Aug 2026 13:21:12 -0700 Subject: [PATCH 1/6] Added changes to support multiple archs in command line --- clang/lib/Driver/ToolChains/Clang.cpp | 64 +++++++++++++++---- clang/lib/Driver/ToolChains/SYCL.cpp | 30 ++++++++- clang/test/Driver/clang-linker-wrapper.cpp | 13 ++++ clang/test/Driver/sycl-offload-new-driver.cpp | 23 +++++-- clang/test/Driver/sycl-offload.cpp | 4 +- .../ClangLinkerWrapper.cpp | 38 +++++++---- 6 files changed, 136 insertions(+), 36 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0aa88c190ad0f..bdd7f1ec1bdc1 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12312,6 +12312,8 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, // -Xdevice-post-link -> --sycl-post-link-options // -Xspirv-translator -> --llvm-spirv-options // -Xspirv-to-ir-wrapper -> --spirv-to-ir-wrapper-options. + // For spir64_gen the value is qualified with "/" and emitted per + // (triple, arch) to keep per-arch tokens from crossing across archs. const toolchains::SYCLToolChain &SYCLTC = static_cast(getToolChain()); for (auto &ToolChainMember : @@ -12319,20 +12321,56 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, const ToolChain *TC = ToolChainMember.second; if (!TC->getTriple().isSPIROrSPIRV()) continue; - 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)); - 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() + "=" + A)); + SmallVector Devices; + if (TC->getTriple().isSPIR() && + TC->getTriple().getSubArch() == llvm::Triple::SPIRSubArch_gen) { + for (BoundArch BA : C.getDriver().getOffloadArchs( + C, C.getArgs(), Action::OFK_SYCL, *TC)) + if (!BA.ArchName.empty()) + Devices.push_back(BA.ArchName); + } + if (Devices.empty()) + Devices.push_back(StringRef()); + + // One entry per (triple, arch); value is space-joined with a leading + // "-device " for gen. Wrapper tokenizes and routes by -device. + ArgStringList BuildArgs; + auto EmitPerArch = [&](StringRef Flag, StringRef Device, + bool PrependDevice) { + SmallString<128> Joined; + // Prepend "-device " only if the user's tokens don't already + // start with one (raw -Xsycl-target-backend=spir64_gen may embed it). + if (PrependDevice && !Device.empty() && + !(BuildArgs.size() >= 2 && StringRef(BuildArgs[0]) == "-device")) { + Joined = "-device "; + Joined += Device; + } + for (const char *T : BuildArgs) { + if (!Joined.empty()) + Joined += ' '; + Joined += T; + } + if (!Joined.empty()) + CmdArgs.push_back(Args.MakeArgString( + Flag + ":" + TC->getTripleString() + "=" + Joined)); + }; + StringRef KindPrefix = Action::GetOffloadKindName(Action::OFK_SYCL); + SmallString<32> CompilerFlag("--device-compiler="); + CompilerFlag += KindPrefix; + SmallString<32> LinkerFlag("--device-linker="); + LinkerFlag += KindPrefix; + for (StringRef Device : Devices) { + BuildArgs.clear(); + SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, + Device); + EmitPerArch(CompilerFlag, Device, /*PrependDevice=*/true); + + BuildArgs.clear(); + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, + Device); + EmitPerArch(LinkerFlag, Device, /*PrependDevice=*/false); + } BuildArgs.clear(); SYCLTC.TranslateTargetOpt( diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 084cf7614158b..737056c06b509 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1654,15 +1654,39 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; if (IsGenTriple) { - if (Device != GenDevice && !Device.empty()) + if (!GenDevice.empty() && Device != GenDevice && !Device.empty()) continue; if (OptTargetTriple != Triple && GenDevice.empty()) // Triples do not match, but only skip when we know we are not // comparing against intel_gpu_* continue; - if (OptTargetTriple == Triple && !Device.empty()) + if (OptTargetTriple == Triple && !Device.empty()) { // Triples match, but we are expecting a specific device to be set. - continue; + // The raw spir64_gen form of -Xsycl-target-backend may still be + // intended for a specific arch when the caller embedded a + // "-device " inside the value (e.g. + // -Xsycl-target-backend=spir64_gen "-device pvc -options ..."). + // Peek at the value and route it only to the matching arch so + // per-arch options do not cross-contaminate other spir64_gen + // sub-targets. If no "-device" is embedded, the entry is treated + // as shared and only surfaces in the unqualified pass. + StringRef Value = A->getValue(1); + SmallVector Tokens; + llvm::BumpPtrAllocator Alloc; + llvm::StringSaver S(Alloc); + llvm::cl::TokenizeGNUCommandLine(Value, S, Tokens); + bool HasEmbeddedDevice = false; + bool EmbeddedDeviceMatches = false; + for (size_t I = 0; I + 1 < Tokens.size(); ++I) { + if (StringRef(Tokens[I]) == "-device") { + HasEmbeddedDevice = true; + EmbeddedDeviceMatches = StringRef(Tokens[I + 1]) == Device; + break; + } + } + if (!HasEmbeddedDevice || !EmbeddedDeviceMatches) + continue; + } } else if (OptTargetTriple != Triple) continue; } else if (!OptNoTriple) diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 9c20a19601bee..3afafddd94d65 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -138,6 +138,19 @@ // CHK-NO-CMDS-AOT-GEN-LINKERARG: sycl-post-link{{.*}} -o {{[^,]*}}.table {{.*}}.bc // CHK-NO-CMDS-AOT-GEN-LINKERARG: ocloc{{.*}} -device pvc -output +// Leading "-device " in each --device-compiler=/--device-linker= +// value routes it to that arch's ocloc call and per-arch sycl-post-link. +// RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_skl -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t1_skl.o +// RUN: clang-linker-wrapper \ +// RUN: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -extraopt_pvc" \ +// RUN: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl -extraopt_skl" \ +// RUN: --linker-path=/usr/bin/ld -o /dev/null %t1.o %t1_skl.o --dry-run 2>&1 \ +// RUN: | FileCheck -check-prefix=CHK-PER-ARCH-DC %s +// CHK-PER-ARCH-DC-DAG: sycl-post-link{{.*}} -o intel_gpu_pvc,{{.*}}.table +// CHK-PER-ARCH-DC-DAG: ocloc{{.*}} -device pvc{{.*}}-extraopt_pvc +// CHK-PER-ARCH-DC-DAG: sycl-post-link{{.*}} -o intel_gpu_skl,{{.*}}.table +// CHK-PER-ARCH-DC-DAG: ocloc{{.*}} -device skl{{.*}}-extraopt_skl + /// Check for list of commands for standalone clang-linker-wrapper run for sycl (AOT for Intel CPU) // ------- // Generate .o file as linker wrapper input. diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index c823bd6ab6484..40e088f5e81a4 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -155,18 +155,31 @@ // 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. +/// -Xsycl-target-backend/-Xsycl-target-linker forward to +/// --device-compiler=/--device-linker= as a space-joined value per +/// (triple, arch); for gen the value has a leading "-device ". // 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: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -opt1" // WRAPPER_OPTIONS_AOT_SEPARATE-SAME: "--device-linker=sycl:spir64_gen-unknown-unknown=-opt2" +/// Two spir64_gen sub-targets on the same triple: one --device-compiler +/// per arch, each carrying its own leading "-device ". +// RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ +// RUN: -fsycl-targets=spir64_gen,intel_gpu_skl \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc -options -extraopt_pvc" \ +// RUN: -Xsycl-target-backend=intel_gpu_skl "-options -extraopt_skl" \ +// RUN: -### %s 2>&1 \ +// RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_MULTI_GEN %s +// WRAPPER_OPTIONS_MULTI_GEN: clang-linker-wrapper +// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc {{[^"]*}}-extraopt_pvc" +// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl {{[^"]*}}-extraopt_skl" +// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc {{[^"]*}}extraopt_skl +// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl {{[^"]*}}extraopt_pvc + /// 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 c7fa5e879e255..e05dfa7915f60 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -329,7 +329,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" "--device-compiler=sycl:spir64-unknown-unknown=-DFOO2" +// CHK-TOOLS-OPTS: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64-unknown-unknown=-DFOO1 -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 \ @@ -349,7 +349,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" "--device-linker=sycl:spir64-unknown-unknown=-DFOO2" +// CHK-TOOLS-OPTS2: clang-linker-wrapper{{.*}} "--device-linker=sycl:spir64-unknown-unknown=-DFOO1 -DFOO2" /// -fsycl-range-rounding settings /// diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 59aa81ed94a74..94f3422ff0077 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2541,11 +2541,8 @@ DerivedArgList getLinkerArgs(ArrayRef Input, if (llvm::all_of(Input, ContainsBitcode)) DAL.AddFlagArg(nullptr, Tbl.getOption(OPT_whole_program)); - // This function filters the SYCL device compiler, linker, sycl-post-link, - // llvm-spirv and spirv-to-ir-wrapper options by target triple and offload - // kind. The options accept values in the form [:][=]. - // An example of passing such an option to clang-linker-wrapper is: - // --device-compiler=sycl:spir64_gen-unknown-unknown=opt_val. + // Filter by kind and triple. Format: [:][=]. + // Arch routing happens on the AOT path below (via embedded "-device"). const StringRef TripleStr = DAL.getLastArgValue(OPT_triple_EQ); auto ProcessDeviceArgs = [&](llvm::opt::OptSpecifier DeviceArgsOptionID, llvm::opt::OptSpecifier ForwardedOptionID) { @@ -2738,14 +2735,29 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // 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)); + // Value is space-joined; a leading "-device " routes it to + // that arch only. No -device prefix -> apply to every arch. + const StringRef Arch = LinkerArgs.getLastArgValue(OPT_arch_EQ); + auto RouteAOT = [&](llvm::opt::OptSpecifier OptID) { + for (StringRef Value : LinkerArgs.getAllArgValues(OptID)) { + if (Value.empty()) + continue; + SmallVector Tokens; + llvm::BumpPtrAllocator Alloc; + llvm::StringSaver S(Alloc); + llvm::cl::TokenizeGNUCommandLine(Value, S, Tokens); + size_t StartIdx = 0; + if (Tokens.size() >= 2 && StringRef(Tokens[0]) == "-device") { + if (StringRef(Tokens[1]) != Arch) + continue; + StartIdx = 2; + } + for (size_t I = StartIdx; I < Tokens.size(); ++I) + AOTDeviceArgs.emplace_back(Tokens[I]); + } + }; + RouteAOT(OPT_compiler_arg_EQ); + RouteAOT(OPT_linker_arg_EQ); } else { for (const std::string &DeviceCompilerArg : LinkerArgs.getAllArgValues(OPT_compiler_arg_EQ)) { From 5246d874e015874ebe24e8bacf6b998e29a0886b Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Mon, 17 Aug 2026 12:48:30 -0700 Subject: [PATCH 2/6] Route each arch. backend options to their own ocloc call --- clang/lib/Driver/ToolChains/Clang.cpp | 6 ++++-- clang/lib/Driver/ToolChains/SYCL.cpp | 19 +++++-------------- .../ClangLinkerWrapper.cpp | 6 ++++-- 3 files changed, 13 insertions(+), 18 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index bdd7f1ec1bdc1..7e022e585a863 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -11001,10 +11001,12 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA, static_cast(*TC); SYCLTC.AddSPIRVImpliedTargetArgs(TC->getTriple(), Args, BuildArgs, JA, *HostTC, Arch.ArchName); - SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs); + SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, + Arch.ArchName); createArgString("compile-opts="); BuildArgs.clear(); - SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs); + SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, + Arch.ArchName); createArgString("link-opts="); } diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 737056c06b509..fae65d200a3ac 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1661,30 +1661,21 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, // comparing against intel_gpu_* continue; if (OptTargetTriple == Triple && !Device.empty()) { - // Triples match, but we are expecting a specific device to be set. - // The raw spir64_gen form of -Xsycl-target-backend may still be - // intended for a specific arch when the caller embedded a - // "-device " inside the value (e.g. - // -Xsycl-target-backend=spir64_gen "-device pvc -options ..."). - // Peek at the value and route it only to the matching arch so - // per-arch options do not cross-contaminate other spir64_gen - // sub-targets. If no "-device" is embedded, the entry is treated - // as shared and only surfaces in the unqualified pass. + // Raw spir64_gen entry: if the value embeds "-device X", route + // only to arch X. Absent -> shared, applies to every arch. StringRef Value = A->getValue(1); SmallVector Tokens; llvm::BumpPtrAllocator Alloc; llvm::StringSaver S(Alloc); llvm::cl::TokenizeGNUCommandLine(Value, S, Tokens); - bool HasEmbeddedDevice = false; - bool EmbeddedDeviceMatches = false; + bool EmbDeviceNoMatch = false; for (size_t I = 0; I + 1 < Tokens.size(); ++I) { if (StringRef(Tokens[I]) == "-device") { - HasEmbeddedDevice = true; - EmbeddedDeviceMatches = StringRef(Tokens[I + 1]) == Device; + EmbDeviceNoMatch = StringRef(Tokens[I + 1]) != Device; break; } } - if (!HasEmbeddedDevice || !EmbeddedDeviceMatches) + if (EmbDeviceNoMatch) continue; } } else if (OptTargetTriple != Triple) diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 94f3422ff0077..fad79a64b9632 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2735,8 +2735,10 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // vector) so both AOT option origins share one code path instead of // two independently-maintained branches. if (TargetTriple.isSPIRAOT()) { - // Value is space-joined; a leading "-device " routes it to - // that arch only. No -device prefix -> apply to every arch. + // Leading "-device " is a wrapper-side routing hint, stripped + // before forwarding tokens to ocloc. ocloc runs once per (triple, + // arch) with a single -device, so it never sees the hint or another + // arch's options. No -device prefix -> apply to every arch. const StringRef Arch = LinkerArgs.getLastArgValue(OPT_arch_EQ); auto RouteAOT = [&](llvm::opt::OptSpecifier OptID) { for (StringRef Value : LinkerArgs.getAllArgValues(OptID)) { From d6bd1ed4fe1a3aea13cb710bcf6c83944808a900 Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Wed, 19 Aug 2026 13:17:45 -0700 Subject: [PATCH 3/6] Route per-arch backend options to their own ocloc call --- clang/lib/Driver/ToolChains/Clang.cpp | 46 +++----- clang/test/Driver/clang-linker-wrapper.cpp | 8 +- clang/test/Driver/sycl-offload-new-driver.cpp | 20 ++-- clang/test/Driver/sycl-offload.cpp | 4 +- .../ClangLinkerWrapper.cpp | 44 +++----- .../per-arch-backend-options.cpp | 102 ++++++++++++++++++ 6 files changed, 151 insertions(+), 73 deletions(-) create mode 100644 sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 7e022e585a863..28252dc1b8c30 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12335,43 +12335,31 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, if (Devices.empty()) Devices.push_back(StringRef()); - // One entry per (triple, arch); value is space-joined with a leading - // "-device " for gen. Wrapper tokenizes and routes by -device. - ArgStringList BuildArgs; - auto EmitPerArch = [&](StringRef Flag, StringRef Device, - bool PrependDevice) { - SmallString<128> Joined; - // Prepend "-device " only if the user's tokens don't already - // start with one (raw -Xsycl-target-backend=spir64_gen may embed it). - if (PrependDevice && !Device.empty() && - !(BuildArgs.size() >= 2 && StringRef(BuildArgs[0]) == "-device")) { - Joined = "-device "; - Joined += Device; - } - for (const char *T : BuildArgs) { - if (!Joined.empty()) - Joined += ' '; - Joined += T; - } - if (!Joined.empty()) - CmdArgs.push_back(Args.MakeArgString( - Flag + ":" + TC->getTripleString() + "=" + Joined)); - }; + // One --device-compiler/--device-linker per token; per-arch routing + // rides on the key (/). Preserves dd9abc1's per-token + // AOT forwarding invariant. Wrapper filters by key, no reparse. StringRef KindPrefix = Action::GetOffloadKindName(Action::OFK_SYCL); - SmallString<32> CompilerFlag("--device-compiler="); - CompilerFlag += KindPrefix; - SmallString<32> LinkerFlag("--device-linker="); - LinkerFlag += KindPrefix; + ArgStringList BuildArgs; for (StringRef Device : Devices) { + SmallString<64> Key(TC->getTripleString()); + if (!Device.empty()) { + Key += '/'; + Key += Device; + } BuildArgs.clear(); SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); - EmitPerArch(CompilerFlag, Device, /*PrependDevice=*/true); - + for (const char *T : BuildArgs) + CmdArgs.push_back(Args.MakeArgString("--device-compiler=" + + KindPrefix + ":" + Key + "=" + + T)); BuildArgs.clear(); SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, Device); - EmitPerArch(LinkerFlag, Device, /*PrependDevice=*/false); + for (const char *T : BuildArgs) + CmdArgs.push_back(Args.MakeArgString("--device-linker=" + + KindPrefix + ":" + Key + "=" + + T)); } BuildArgs.clear(); diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 3afafddd94d65..3103d4cd1adcc 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -138,12 +138,12 @@ // CHK-NO-CMDS-AOT-GEN-LINKERARG: sycl-post-link{{.*}} -o {{[^,]*}}.table {{.*}}.bc // CHK-NO-CMDS-AOT-GEN-LINKERARG: ocloc{{.*}} -device pvc -output -// Leading "-device " in each --device-compiler=/--device-linker= -// value routes it to that arch's ocloc call and per-arch sycl-post-link. +// "/" qualifier on --device-compiler=/--device-linker= routes each +// value to that arch's ocloc call and per-arch sycl-post-link output name. // RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_skl -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t1_skl.o // RUN: clang-linker-wrapper \ -// RUN: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -extraopt_pvc" \ -// RUN: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl -extraopt_skl" \ +// RUN: --device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-extraopt_pvc \ +// RUN: --device-compiler=sycl:spir64_gen-unknown-unknown/skl=-extraopt_skl \ // RUN: --linker-path=/usr/bin/ld -o /dev/null %t1.o %t1_skl.o --dry-run 2>&1 \ // RUN: | FileCheck -check-prefix=CHK-PER-ARCH-DC %s // CHK-PER-ARCH-DC-DAG: sycl-post-link{{.*}} -o intel_gpu_pvc,{{.*}}.table diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index 40e088f5e81a4..60a3b7a4c7c89 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/clang/test/Driver/sycl-offload-new-driver.cpp @@ -156,18 +156,18 @@ // WRAPPER_OPTIONS_BACKEND_AOT-SAME: "--device-compiler=sycl:spir64_x86_64-unknown-unknown=-backend-cpu-opt" /// -Xsycl-target-backend/-Xsycl-target-linker forward to -/// --device-compiler=/--device-linker= as a space-joined value per -/// (triple, arch); for gen the value has a leading "-device ". +/// --device-compiler=/--device-linker= one token per occurrence; per-arch +/// routing rides on a "/" qualifier appended to the triple key. // 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=-device pvc -opt1" -// WRAPPER_OPTIONS_AOT_SEPARATE-SAME: "--device-linker=sycl:spir64_gen-unknown-unknown=-opt2" +// WRAPPER_OPTIONS_AOT_SEPARATE: clang-linker-wrapper{{.*}} "--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-opt1" +// WRAPPER_OPTIONS_AOT_SEPARATE-SAME: "--device-linker=sycl:spir64_gen-unknown-unknown/pvc=-opt2" -/// Two spir64_gen sub-targets on the same triple: one --device-compiler -/// per arch, each carrying its own leading "-device ". +/// Two spir64_gen sub-targets on the same triple: each arch's tokens +/// carry their own "/" qualifier so options don't cross-contaminate. // RUN: %clangxx --target=x86_64-unknown-linux-gnu -fsycl --offload-new-driver --sysroot=%S/Inputs/SYCL \ // RUN: -fsycl-targets=spir64_gen,intel_gpu_skl \ // RUN: -Xsycl-target-backend=spir64_gen "-device pvc -options -extraopt_pvc" \ @@ -175,10 +175,10 @@ // RUN: -### %s 2>&1 \ // RUN: | FileCheck -check-prefix WRAPPER_OPTIONS_MULTI_GEN %s // WRAPPER_OPTIONS_MULTI_GEN: clang-linker-wrapper -// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc {{[^"]*}}-extraopt_pvc" -// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl {{[^"]*}}-extraopt_skl" -// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc {{[^"]*}}extraopt_skl -// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl {{[^"]*}}extraopt_pvc +// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-extraopt_pvc" +// WRAPPER_OPTIONS_MULTI_GEN-SAME: "--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-extraopt_skl" +// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-extraopt_skl" +// WRAPPER_OPTIONS_MULTI_GEN-NOT: "--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-extraopt_pvc" /// 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 e05dfa7915f60..c7fa5e879e255 100644 --- a/clang/test/Driver/sycl-offload.cpp +++ b/clang/test/Driver/sycl-offload.cpp @@ -329,7 +329,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 \ @@ -349,7 +349,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 fad79a64b9632..e22180c43c807 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2541,9 +2541,11 @@ DerivedArgList getLinkerArgs(ArrayRef Input, if (llvm::all_of(Input, ContainsBitcode)) DAL.AddFlagArg(nullptr, Tbl.getOption(OPT_whole_program)); - // Filter by kind and triple. Format: [:][=]. - // Arch routing happens on the AOT path below (via embedded "-device"). + // Filter by kind, triple, and optional / qualifier on the key. + // Format: [:][[/]=]. Entries without / + // apply to every arch of the matching triple. const StringRef TripleStr = DAL.getLastArgValue(OPT_triple_EQ); + const StringRef ArchStr = DAL.getLastArgValue(OPT_arch_EQ); auto ProcessDeviceArgs = [&](llvm::opt::OptSpecifier DeviceArgsOptionID, llvm::opt::OptSpecifier ForwardedOptionID) { for (StringRef DeviceArgValue : Args.getAllArgValues(DeviceArgsOptionID)) { @@ -2556,12 +2558,15 @@ DerivedArgList getLinkerArgs(ArrayRef Input, } size_t EqPos = DeviceArgValue.find('='); if (EqPos != StringRef::npos) { - StringRef ArgTargetTripleStr = DeviceArgValue.take_front(EqPos); + StringRef Key = DeviceArgValue.take_front(EqPos); + auto [ArgTargetTripleStr, ArgArchStr] = Key.split('/'); llvm::Triple ArgTargetTriple(ArgTargetTripleStr); // If this isn't a recognized triple then it's an `arg=value` option. if (ArgTargetTriple.getArch() != Triple::ArchType::UnknownArch) { if (ArgTargetTripleStr != TripleStr) continue; + if (!ArgArchStr.empty() && ArgArchStr != ArchStr) + continue; DeviceArgValue = DeviceArgValue.drop_front(EqPos + 1); } } @@ -2735,31 +2740,14 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // vector) so both AOT option origins share one code path instead of // two independently-maintained branches. if (TargetTriple.isSPIRAOT()) { - // Leading "-device " is a wrapper-side routing hint, stripped - // before forwarding tokens to ocloc. ocloc runs once per (triple, - // arch) with a single -device, so it never sees the hint or another - // arch's options. No -device prefix -> apply to every arch. - const StringRef Arch = LinkerArgs.getLastArgValue(OPT_arch_EQ); - auto RouteAOT = [&](llvm::opt::OptSpecifier OptID) { - for (StringRef Value : LinkerArgs.getAllArgValues(OptID)) { - if (Value.empty()) - continue; - SmallVector Tokens; - llvm::BumpPtrAllocator Alloc; - llvm::StringSaver S(Alloc); - llvm::cl::TokenizeGNUCommandLine(Value, S, Tokens); - size_t StartIdx = 0; - if (Tokens.size() >= 2 && StringRef(Tokens[0]) == "-device") { - if (StringRef(Tokens[1]) != Arch) - continue; - StartIdx = 2; - } - for (size_t I = StartIdx; I < Tokens.size(); ++I) - AOTDeviceArgs.emplace_back(Tokens[I]); - } - }; - RouteAOT(OPT_compiler_arg_EQ); - RouteAOT(OPT_linker_arg_EQ); + 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)) { diff --git a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp new file mode 100644 index 0000000000000..8fba4838d890f --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp @@ -0,0 +1,102 @@ +//==-- per-arch-backend-options.cpp -----------------------------------------==// +// +// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// +//===----------------------------------------------------------------------===// +// End-to-end test: build a SYCL program AOT-compiled for two Intel GPU archs +// (pvc + dg2_g10) with distinct -Xsycl-target-backend options for each, and +// verify that: +// 1. The driver emits per-(triple, arch) --device-compiler entries and +// the wrapper routes each arch's tokens to its own ocloc invocation +// (compile-time check on -v output). +// 2. The resulting fat binary runs correctly on a matching device +// (runtime check). + +// REQUIRES: ocloc, target-spir +// REQUIRES: arch-intel_gpu_pvc +// One physical arch (pvc) is required so %{run} has an AOT image matching +// the local device. The second arch (dg2_g10) is a build-only target that +// exercises the routing logic; ocloc builds its image but the runtime never +// executes it. + +// RUN: %{run-aux} %clangxx -Wno-error=unused-command-line-argument \ +// RUN: --offload-new-driver -fsycl \ +// RUN: -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2_g10 \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc "-options -cl-mad-enable" \ +// RUN: -Xsycl-target-backend=intel_gpu_dg2_g10 "-options -cl-unsafe-math-optimizations" \ +// RUN: -v %s -o %t.out > %t.log 2>&1 +// RUN: FileCheck --input-file=%t.log --check-prefix=CHECK-PVC %s +// RUN: FileCheck --input-file=%t.log --check-prefix=CHECK-ACM %s +// RUN: %{run} %t.out + +// pvc's ocloc call carries -cl-mad-enable and NOT -cl-unsafe-math-optimizations. +// CHECK-PVC: ocloc{{.*}} -device pvc {{.*}}-cl-mad-enable +// CHECK-PVC-NOT: ocloc{{.*}} -device pvc {{.*}}-cl-unsafe-math-optimizations + +// dg2_g10's canonical ocloc device name is acm_g10. +// CHECK-ACM: ocloc{{.*}} -device acm_g10 {{.*}}-cl-unsafe-math-optimizations +// CHECK-ACM-NOT: ocloc{{.*}} -device acm_g10 {{.*}}-cl-mad-enable + +// Regression: raw spir64_gen with an embedded "-device " in the +// backend option value must also route per-arch without leakage. +// RUN: %{run-aux} %clangxx -Wno-error=unused-command-line-argument \ +// RUN: --offload-new-driver -fsycl \ +// RUN: -fsycl-targets=intel_gpu_dg2_g10,spir64_gen \ +// RUN: -Xsycl-target-backend=spir64_gen "-device pvc -options -cl-mad-enable" \ +// RUN: -Xsycl-target-backend=intel_gpu_dg2_g10 "-options -cl-unsafe-math-optimizations" \ +// RUN: -v %s -o %t_raw.out > %t_raw.log 2>&1 +// RUN: FileCheck --input-file=%t_raw.log --check-prefix=CHECK-RAW-PVC %s +// RUN: FileCheck --input-file=%t_raw.log --check-prefix=CHECK-RAW-ACM %s + +// CHECK-RAW-PVC: ocloc{{.*}} -device pvc {{.*}}-cl-mad-enable +// CHECK-RAW-PVC-NOT: ocloc{{.*}} -device pvc {{.*}}-cl-unsafe-math-optimizations +// CHECK-RAW-ACM: ocloc{{.*}} -device acm_g10 {{.*}}-cl-unsafe-math-optimizations +// CHECK-RAW-ACM-NOT: ocloc{{.*}} -device acm_g10 {{.*}}-cl-mad-enable + +#include + +#include +#include +#include + +constexpr std::size_t N = 16; + +class VecAdd; + +int main() { + std::array a{}, b{}, c{}; + for (std::size_t i = 0; i < N; ++i) { + a[i] = static_cast(i); + b[i] = static_cast(2 * i); + } + + { + sycl::queue q; + sycl::buffer bufA{a.data(), sycl::range<1>{N}}; + sycl::buffer bufB{b.data(), sycl::range<1>{N}}; + sycl::buffer bufC{c.data(), sycl::range<1>{N}}; + + q.submit([&](sycl::handler &h) { + auto accA = bufA.get_access(h); + auto accB = bufB.get_access(h); + auto accC = bufC.get_access(h); + h.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> i) { + accC[i] = accA[i] + accB[i]; + }); + }).wait(); + } + + for (std::size_t i = 0; i < N; ++i) { + int expected = static_cast(i) + static_cast(2 * i); + if (c[i] != expected) { + std::cerr << "FAIL at i=" << i << ": got " << c[i] << ", expected " + << expected << "\n"; + return 1; + } + } + + std::cout << "PASS\n"; + return 0; +} From f58455ff18d46b9d12648ca90d80dd42c7bae5e7 Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Wed, 19 Aug 2026 14:21:25 -0700 Subject: [PATCH 4/6] Fixed an issue in e2e test --- sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp index 8fba4838d890f..33566f8cd1f10 100644 --- a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp +++ b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp @@ -21,7 +21,7 @@ // exercises the routing logic; ocloc builds its image but the runtime never // executes it. -// RUN: %{run-aux} %clangxx -Wno-error=unused-command-line-argument \ +// RUN: %clangxx -Wno-error=unused-command-line-argument \ // RUN: --offload-new-driver -fsycl \ // RUN: -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2_g10 \ // RUN: -Xsycl-target-backend=intel_gpu_pvc "-options -cl-mad-enable" \ @@ -41,7 +41,7 @@ // Regression: raw spir64_gen with an embedded "-device " in the // backend option value must also route per-arch without leakage. -// RUN: %{run-aux} %clangxx -Wno-error=unused-command-line-argument \ +// RUN: %clangxx -Wno-error=unused-command-line-argument \ // RUN: --offload-new-driver -fsycl \ // RUN: -fsycl-targets=intel_gpu_dg2_g10,spir64_gen \ // RUN: -Xsycl-target-backend=spir64_gen "-device pvc -options -cl-mad-enable" \ From 4bc16a42422c7e9bce160f19d895fbc0755d8f03 Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Wed, 19 Aug 2026 15:51:55 -0700 Subject: [PATCH 5/6] Fixed an issue in test --- sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp index 33566f8cd1f10..f69b5c7bc3023 100644 --- a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp +++ b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp @@ -79,9 +79,9 @@ int main() { sycl::buffer bufC{c.data(), sycl::range<1>{N}}; q.submit([&](sycl::handler &h) { - auto accA = bufA.get_access(h); - auto accB = bufB.get_access(h); - auto accC = bufC.get_access(h); + sycl::accessor accA{bufA, h, sycl::read_only}; + sycl::accessor accB{bufB, h, sycl::read_only}; + sycl::accessor accC{bufC, h, sycl::write_only}; h.parallel_for(sycl::range<1>{N}, [=](sycl::id<1> i) { accC[i] = accA[i] + accB[i]; }); From 94fba2f5efbde2e3bd5c4a96c5e1c28ba9079e8b Mon Sep 17 00:00:00 2001 From: "Balaji V. Iyer" Date: Wed, 19 Aug 2026 16:15:11 -0700 Subject: [PATCH 6/6] Fixed clang format --- clang/lib/Driver/ToolChains/Clang.cpp | 10 ++++------ .../NewOffloadDriver/per-arch-backend-options.cpp | 4 ++-- 2 files changed, 6 insertions(+), 8 deletions(-) diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 28252dc1b8c30..522bdee6a4c4e 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12350,16 +12350,14 @@ void LinkerWrapper::ConstructJob(Compilation &C, const JobAction &JA, SYCLTC.TranslateBackendTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const char *T : BuildArgs) - CmdArgs.push_back(Args.MakeArgString("--device-compiler=" + - KindPrefix + ":" + Key + "=" + - T)); + CmdArgs.push_back(Args.MakeArgString( + "--device-compiler=" + KindPrefix + ":" + Key + "=" + T)); BuildArgs.clear(); SYCLTC.TranslateLinkerTargetArgs(TC->getTriple(), Args, BuildArgs, Device); for (const char *T : BuildArgs) - CmdArgs.push_back(Args.MakeArgString("--device-linker=" + - KindPrefix + ":" + Key + "=" + - T)); + CmdArgs.push_back(Args.MakeArgString("--device-linker=" + KindPrefix + + ":" + Key + "=" + T)); } BuildArgs.clear(); diff --git a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp index f69b5c7bc3023..05e494eede29f 100644 --- a/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp +++ b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp @@ -1,4 +1,4 @@ -//==-- per-arch-backend-options.cpp -----------------------------------------==// +//==-- per-arch-backend-options.cpp ---------------------------------------==// // // Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions. // See https://llvm.org/LICENSE.txt for license information. @@ -31,7 +31,7 @@ // RUN: FileCheck --input-file=%t.log --check-prefix=CHECK-ACM %s // RUN: %{run} %t.out -// pvc's ocloc call carries -cl-mad-enable and NOT -cl-unsafe-math-optimizations. +// pvc's ocloc call carries -cl-mad-enable, NOT -cl-unsafe-math-optimizations. // CHECK-PVC: ocloc{{.*}} -device pvc {{.*}}-cl-mad-enable // CHECK-PVC-NOT: ocloc{{.*}} -device pvc {{.*}}-cl-unsafe-math-optimizations