diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 0aa88c190ad0f..522bdee6a4c4e 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="); } @@ -12312,6 +12314,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 +12323,42 @@ 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 --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); + 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( diff --git a/clang/lib/Driver/ToolChains/SYCL.cpp b/clang/lib/Driver/ToolChains/SYCL.cpp index 084cf7614158b..fae65d200a3ac 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/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 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) diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 9c20a19601bee..3103d4cd1adcc 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 +// "/" 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. diff --git a/clang/test/Driver/sycl-offload-new-driver.cpp b/clang/test/Driver/sycl-offload-new-driver.cpp index c823bd6ab6484..60a3b7a4c7c89 100644 --- a/clang/test/Driver/sycl-offload-new-driver.cpp +++ b/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 "/" 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 "/" 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 \ diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 59aa81ed94a74..e22180c43c807 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2541,12 +2541,11 @@ 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, 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)) { @@ -2559,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); } } 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..05e494eede29f --- /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: %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 " 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 + +#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) { + 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]; + }); + }).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; +}