[SYCL] Added changes to support multiple archs in command line

clang/lib/Driver/ToolChains/Clang.cpp

@@ -11001,10 +11001,12 @@ void OffloadPackager::ConstructJob(Compilation &C, const JobAction &JA,
           static_cast<const toolchains::SYCLToolChain &>(*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=");
     }
 
@@ -12312,27 +12314,51 @@ 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 "/<arch>" and emitted per
+    // (triple, arch) to keep per-arch tokens from crossing across archs.
     const toolchains::SYCLToolChain &SYCLTC =
         static_cast<const toolchains::SYCLToolChain &>(getToolChain());
     for (auto &ToolChainMember :
          llvm::make_range(ToolChainRange.first, ToolChainRange.second)) {
       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<StringRef, 4> 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 --device-compiler/--device-linker per token; per-arch routing
+      // rides on the key (<triple>/<arch>). Preserves dd9abc1's per-token
+      // AOT forwarding invariant. Wrapper filters by key, no reparse.
+      StringRef KindPrefix = Action::GetOffloadKindName(Action::OFK_SYCL);
+      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);
+        for (const char *T : BuildArgs)
+          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));
+      }
 
       BuildArgs.clear();
       SYCLTC.TranslateTargetOpt(

clang/lib/Driver/ToolChains/SYCL.cpp

@@ -1654,15 +1654,30 @@ 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())
-          // Triples match, but we are expecting a specific device to be set.
-          continue;
+        if (OptTargetTriple == Triple && !Device.empty()) {
+          // 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<const char *, 8> Tokens;
+          llvm::BumpPtrAllocator Alloc;
+          llvm::StringSaver S(Alloc);
+          llvm::cl::TokenizeGNUCommandLine(Value, S, Tokens);
+          bool EmbDeviceNoMatch = false;
+          for (size_t I = 0; I + 1 < Tokens.size(); ++I) {
+            if (StringRef(Tokens[I]) == "-device") {
+              EmbDeviceNoMatch = StringRef(Tokens[I + 1]) != Device;
+              break;
+            }
+          }
+          if (EmbDeviceNoMatch)
+            continue;
+        }
       } else if (OptTargetTriple != Triple)
         continue;
     } else if (!OptNoTriple)

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
 
+// "/<arch>" 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/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
+// 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.

clang/test/Driver/sycl-offload-new-driver.cpp

@@ -155,17 +155,30 @@
 // 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= one token per occurrence; per-arch
+/// routing rides on a "/<arch>" 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=-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: each arch's tokens
+/// carry their own "/<arch>" 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" \
+// 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/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 \

clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp

@@ -2541,12 +2541,11 @@ DerivedArgList getLinkerArgs(ArrayRef<OffloadFile> 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 [<kind>:][<triple>=]<value>.
-  // An example of passing such an option to clang-linker-wrapper is:
-  // --device-compiler=sycl:spir64_gen-unknown-unknown=opt_val.
+  // Filter by kind, triple, and optional /<arch> qualifier on the key.
+  // Format: [<kind>:][<triple>[/<arch>]=]<value>. Entries without /<arch>
+  // 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)) {
@@ -2559,12 +2558,15 @@ DerivedArgList getLinkerArgs(ArrayRef<OffloadFile> 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);
         }
       }

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: %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, 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 <arch>" in the
+// backend option value must also route per-arch without leakage.
+// 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" \
+// 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 <sycl/sycl.hpp>
+
+#include <array>
+#include <cstddef>
+#include <iostream>
+
+constexpr std::size_t N = 16;
+
+class VecAdd;
+
+int main() {
+  std::array<int, N> a{}, b{}, c{};
+  for (std::size_t i = 0; i < N; ++i) {
+    a[i] = static_cast<int>(i);
+    b[i] = static_cast<int>(2 * i);
+  }
+
+  {
+    sycl::queue q;
+    sycl::buffer<int, 1> bufA{a.data(), sycl::range<1>{N}};
+    sycl::buffer<int, 1> bufB{b.data(), sycl::range<1>{N}};
+    sycl::buffer<int, 1> bufC{c.data(), sycl::range<1>{N}};
+
+    q.submit([&](sycl::handler &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<VecAdd>(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<int>(i) + static_cast<int>(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;
+}