diff --git a/clang/lib/Driver/Driver.cpp b/clang/lib/Driver/Driver.cpp index 6546e611d11e6..913466a0e0b6f 100644 --- a/clang/lib/Driver/Driver.cpp +++ b/clang/lib/Driver/Driver.cpp @@ -7919,18 +7919,11 @@ Driver::getOffloadArchs(Compilation &C, const llvm::opt::DerivedArgList &Args, ArgStringList TargetArgs; DeviceTC->TranslateBackendTargetArgs(DeviceTC->getTriple(), C.getInputArgs(), TargetArgs); - // Look for -device and use that as the known - // arch to be associated with the current spir64_gen entry. Grab - // the right most entry. - for (int i = TargetArgs.size() - 2; i >= 0; --i) { - if (StringRef(TargetArgs[i]) == "-device") { - StringRef Arch; - Arch = TargetArgs[i + 1]; - if (!Arch.empty()) - Archs.insert(Arch); - break; - } - } + // Use the rightmost embedded "-device " as the arch bound to + // the raw spir64_gen entry. + if (StringRef Arch = tools::SYCL::gen::getEmbeddedDeviceArch(TargetArgs); + !Arch.empty()) + Archs.insert(Arch); } } diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index 48fe5a1fc809f..5d308526b4c2a 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -12331,6 +12331,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 key is qualified with "/" and emitted per + // (triple, arch) to keep per-arch tokens from leaking between archs. const toolchains::SYCLToolChain &SYCLTC = static_cast(getToolChain()); for (auto &ToolChainMember : @@ -12338,20 +12340,41 @@ 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 (/). + 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 5df5a8d60bed0..7ce99b1c3bd2e 100644 --- a/clang/lib/Driver/ToolChains/SYCL.cpp +++ b/clang/lib/Driver/ToolChains/SYCL.cpp @@ -1097,6 +1097,13 @@ StringRef SYCL::gen::getGenGRFFlag(StringRef GRFMode) { return GRFModeFlagMap[GRFMode]; } +StringRef SYCL::gen::getEmbeddedDeviceArch(ArrayRef Tokens) { + for (int I = static_cast(Tokens.size()) - 2; I >= 0; --I) + if (StringRef(Tokens[I]) == "-device") + return Tokens[I + 1]; + return {}; +} + void SYCL::gen::BackendCompiler::ConstructJob(Compilation &C, const JobAction &JA, const InputInfo &Output, @@ -1663,14 +1670,18 @@ void SYCLToolChain::TranslateTargetOpt(const llvm::Triple &Triple, bool IsGenTriple = Triple.isSPIR() && Triple.getSubArch() == llvm::Triple::SPIRSubArch_gen; if (IsGenTriple) { - if (Device != GenDevice && !Device.empty()) + // "aot_generic" is a new-offload-model pseudo-arch bucket for raw + // -Xsycl-target-backend=spir64_gen entries; inert in old model. + StringRef AOTGenericArch = "aot_generic"; + if (Device != GenDevice && !Device.empty() && + !(GenDevice.empty() && Device == AOTGenericArch)) 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. + if (OptTargetTriple == Triple && !Device.empty() && + Device != AOTGenericArch) continue; } else if (OptTargetTriple != Triple) continue; diff --git a/clang/lib/Driver/ToolChains/SYCL.h b/clang/lib/Driver/ToolChains/SYCL.h index 7b38ba8c61e6f..7057267362220 100644 --- a/clang/lib/Driver/ToolChains/SYCL.h +++ b/clang/lib/Driver/ToolChains/SYCL.h @@ -82,6 +82,9 @@ StringRef resolveGenDevice(StringRef DeviceName); SmallString<64> getGenDeviceMacro(StringRef DeviceName); StringRef getGenGRFFlag(StringRef GRFMode); +// Returns the rightmost "-device " value in Tokens, or empty if none. +StringRef getEmbeddedDeviceArch(ArrayRef Tokens); + // Returns the full path of the ocloc tool to be used for AOT compilation and // for emitting the ocloc help information. A user provided --ocloc-path= is // honored above all other lookup locations. If not found, the tool (ocloc) is diff --git a/clang/test/Driver/clang-linker-wrapper.cpp b/clang/test/Driver/clang-linker-wrapper.cpp index 642e5201b0937..2bd7f4d04b6f8 100644 --- a/clang/test/Driver/clang-linker-wrapper.cpp +++ b/clang/test/Driver/clang-linker-wrapper.cpp @@ -153,6 +153,19 @@ // RUN: not clang-linker-wrapper --ocloc-path= --linker-path=/usr/bin/ld -o /dev/null %t1.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-OCLOC-PATH-NOARG %s // CHK-OCLOC-PATH-NOARG: no directory given for '--ocloc-path=' +// "/" 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..8cd155c88ccbf 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=intel_gpu_pvc,intel_gpu_skl \ +// RUN: -Xsycl-target-backend=intel_gpu_pvc "-options -extraopt_pvc" \ +// RUN: -Xsycl-target-backend=intel_gpu_skl "-options -extraopt_skl" \ +// RUN: -### %s 2>&1 \ +// RUN: | FileCheck --implicit-check-not='/pvc=-extraopt_skl' \ +// RUN: --implicit-check-not='/skl=-extraopt_pvc' \ +// RUN: -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" /// 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 a336dabc47a5c..41934a01dc348 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2569,12 +2569,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)) { @@ -2587,12 +2586,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/doc/design/OffloadDesign.md b/sycl/doc/design/OffloadDesign.md index 0485e6ad48e63..b8097265bd854 100644 --- a/sycl/doc/design/OffloadDesign.md +++ b/sycl/doc/design/OffloadDesign.md @@ -256,27 +256,33 @@ model's usage pattern. This will be implemented by invoking `clang-linker-wrappe TU's device code independently, embedding the result directly into the host object. #### Format of the --device-compiler Option -The `--device-compiler` option uses the format `--device-compiler=[:][=]` where: +The `--device-compiler` option uses the format `--device-compiler=[:][[/]=]` where: - `` : specifies the offloading kind (e.g., sycl, hip, openmp) and is optional. - `` : specifies the target triple (e.g., `spir64_gen-unknown-unknown`, `spir64_x86_64-unknown-unknown`) and is optional. -- `` : contains the arguments to be passed to the backend compiler. +- `` : optional architecture qualifier appended to the triple after a `/`. Used for `spir64_gen`, where a single triple may back several GPU architectures. +- `` : one option token to be passed to the backend compiler. Each `--device-compiler` occurrence carries a single token; multi-token option strings are emitted as multiple `--device-compiler` occurrences with the same key. -In clang-linker-wrapper, the `` and `` are matched against the current compilation target. Only arguments that match both the offloading kind and target triple will be passed to the backend compiler. If `` is not specified, the arguments will match any offloading kind; if `` is not specified, the arguments will match any target triple; and if neither is specified, the arguments will be applied to all targets. +In clang-linker-wrapper, the ``, ``, and `` are matched against the current compilation target. Only arguments that match all specified filters are forwarded to the backend compiler. If `` is not specified, the arguments will match any offloading kind; if `` is not specified, the arguments will match any target triple; if `` is not specified, the arguments will match every architecture of the matching triple. -To support multiple device architectures, a new `--device-compiler` option must be specified for each device. For example, to compile for Ponte Vecchio (PVC) and Skylake (SKL) architectures and put them in a fat binary, the user must add the following two `--device-compiler` options: +To supply per-architecture backend options, emit a separate `--device-compiler` occurrence for each `(triple, arch)` pair and for each option token. For example, to build for Ponte Vecchio (PVC) and Skylake (SKL) architectures and put them in a fat binary, the driver emits one `--device-compiler` occurrence per token per arch: -`--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -options ...` - -`--device-compiler=sycl:spir64_gen-unknown-unknown=-device skl -options ...` - -Device specific options for each of the device architectures should be specified after `-device `. +``` +--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-options +--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-cl-mad-enable +--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-options +--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-cl-unsafe-math-optimizations +``` -Here is an example of a clang-linker-wrapper invocation where ther user wants to create a fat binary with PVC and SKL architectures to be run on a x86_64 Linux host. In addition, they would like to enable aggressive mathematical optimizations and are tolerant for slightly imprecise floating-point values just for SKL, that is, use the `-cl-unsafe-math-optimizations` flag. For PVC, they would like to enable the multiply and add instruction usage (`-cl-mad-enable`). The source binaries are called host.o and kernel.o and the output should be called out.exe. +Here is an example of a clang-linker-wrapper invocation where the user wants to create a fat binary with PVC and SKL architectures to run on an x86_64 Linux host. For SKL they want aggressive floating-point relaxation (`-cl-unsafe-math-optimizations`); for PVC they want multiply-and-add fusion (`-cl-mad-enable`). The source binaries are called `host.o` and `kernel.o` and the output should be called `out.exe`. -`clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu - --device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -options "-cl-mad-enable" - --device-compiler=sycl:spir64_gen-unknown-unknown=-device skl -options "-cl-unsafe-math-optimizations" -- /usr/bin/ld -host.o kernel.o -o out.exe` +``` +clang-linker-wrapper --host-triple=x86_64-unknown-linux-gnu \ + --device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-options \ + --device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-cl-mad-enable \ + --device-compiler=sycl:spir64_gen-unknown-unknown/skl=-options \ + --device-compiler=sycl:spir64_gen-unknown-unknown/skl=-cl-unsafe-math-optimizations \ + -- /usr/bin/ld host.o kernel.o -o out.exe +``` #### Other Supported Options To complete the support needed for the various targets using the @@ -316,19 +322,21 @@ list to be passed along. *Example: spir64_gen enabling options* -> "--device-compiler=sycl:spir64_gen-unknown-unknown=-device pvc -options extraopt_pvc" -"--device-compiler=sycl:spir64_gen-unknown-unknown=-options -extraopt_skl" +> --device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-device +--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=pvc +--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-options +--device-compiler=sycl:spir64_gen-unknown-unknown/pvc=-extraopt_pvc +--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-options +--device-compiler=sycl:spir64_gen-unknown-unknown/skl=-extraopt_skl *Example: clang-linker-wrapper options* -Each OCLOC call will be represented as a separate device binary that is -individually wrapped and linked into the final executable. - -Additionally, the syntax can be expanded to enable the ability to pass specific -options to a specific device GPU target for spir64_gen. The syntax will -resemble `--device-compiler=sycl:spir64_gen-unknown-unknown= `. This corresponds to the existing -option syntax of `-fsycl-targets=intel_gpu_arch` where `arch` can be a fixed -set of targets. +Each `(triple, arch)` pair produces its own OCLOC call and its own device +binary that is individually wrapped and linked into the final executable. The +`/` qualifier on the key routes each option token to the ocloc +invocation for that arch; a `--device-compiler` occurrence with no +`/` (or from a non-gen triple) applies to every arch of the matching +triple. #### --offload-arch 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..8cbad50dfda17 --- /dev/null +++ b/sycl/test-e2e/NewOffloadDriver/per-arch-backend-options.cpp @@ -0,0 +1,36 @@ +//==-- 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 +// +//===----------------------------------------------------------------------===// +// Build-only test that -Xsycl-target-backend options for two AOT Intel GPU +// archs (pvc + dg2_g10) reach only their own arch's ocloc invocation. + +// REQUIRES: ocloc + +// 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 \ +// RUN: --implicit-check-not='ocloc{{.*}} -device pvc {{.*}}-cl-unsafe-math-optimizations' %s +// RUN: FileCheck --input-file=%t.log --check-prefix=CHECK-ACM \ +// RUN: --implicit-check-not='-device acm_g10 {{.*}}-cl-mad-enable' %s + +// CHECK-PVC: ocloc{{.*}} -device pvc {{.*}}-cl-mad-enable +// dg2_g10's canonical ocloc device name is acm_g10. +// CHECK-ACM: ocloc{{.*}} -device acm_g10 {{.*}}-cl-unsafe-math-optimizations + +#include + +int main() { + sycl::queue q; + q.submit([&](sycl::handler &h) { + h.parallel_for(sycl::range<1>{1}, [=](sycl::id<1>) {}); + }); + return 0; +}