diff --git a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp index 1b380896381cf..eff9831b8ffbd 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -3,9 +3,12 @@ // REQUIRES: spirv-to-ir-wrapper, sycl-post-link // ------- // Generate .o file as linker wrapper input. +// These objects are built in RDC mode on purpose: they carry the packaged +// device code that the clang-linker-wrapper runs below consume. // // RUN: %clang %s -fsycl -fsycl-targets=spir64-unknown-unknown -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t.o // RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_pvc -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu.o +// RUN: %clang %s -fsycl -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_gpu_multi.o // RUN: %clang %s -fsycl -fsycl-targets=spir64_x86_64 -c --offload-new-driver --no-offloadlib -fno-sycl-instrument-device-code -o %t_aot_cpu.o // RUN: %clang %s -fsycl -fsycl-targets=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_50 -nocudalib -fno-sycl-libspirv -c --offload-new-driver -o %t_nvptx.o // RUN: %clang %s -fsycl -fsycl-targets=amdgcn-amd-amdhsa -Xsycl-target-backend=amdgcn-amd-amdhsa --offload-arch=gfx803 -fgpu-rdc -nogpulib -fno-sycl-libspirv -c --offload-new-driver -o %t_amdgcn.o @@ -348,6 +351,24 @@ // CHK-FATBIN-ONLY-NEXT: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.syclfb // CHK-FATBIN-ONLY-NOT: "{{.*}}/ld" +// With several GPU architectures, all their images are wrapped into a single +// wrapper module. +// RUN: clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.multi.syclfb %t_aot_gpu_multi.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MULTI %s +// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device pvc{{.*}} -output [[PVC:[^ ]+]] +// CHK-FATBIN-ONLY-MULTI-DAG: ocloc{{.*}} -device dg2{{.*}} -output [[DG2:[^ ]+]] +// CHK-FATBIN-ONLY-MULTI: offload-wrapper: output: [[WRAPPEROUT:.*]].bc,{{.*}} input: [[DG2]],{{.*}} input: [[PVC]], +// CHK-FATBIN-ONLY-MULTI-NOT: offload-wrapper: +// CHK-FATBIN-ONLY-MULTI: "{{.*cp|copy}}" [[WRAPPEROUT]].bc {{.*}}.multi.syclfb +// CHK-FATBIN-ONLY-MULTI-NOT: "{{.*}}/ld" +// RUN: not clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.mixed.syclfb %t_aot_gpu.o %t_aot_cpu.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-MIXED %s +// CHK-FATBIN-ONLY-MIXED: error: --no-sycl-rdc does not support multiple SYCL target triples ('{{.*}}' and '{{.*}}') +// A SYCL input without a target triple is diagnosed. +// RUN: %clang -cc1 -triple spir64-unknown-unknown -emit-llvm-bc -x c /dev/null -o %t_notriple.bc +// RUN: llvm-offload-binary -o %t_notriple.fat --image=file=%t_notriple.bc,kind=sycl,triple= +// RUN: %clang -cc1 -triple x86_64-unknown-linux-gnu -emit-obj -x c /dev/null -o %t_notriple.o -fembed-offload-object=%t_notriple.fat +// RUN: not clang-linker-wrapper --emit-fatbin-only --no-sycl-rdc --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o %t.notriple.syclfb %t_notriple.o --dry-run 2>&1 | FileCheck -check-prefix=CHK-FATBIN-ONLY-NOTRIPLE %s +// CHK-FATBIN-ONLY-NOTRIPLE: error: --no-sycl-rdc requires a target triple for every SYCL input + // Verify list of commands when syclbin is used // RUN: clang-linker-wrapper --bitcode-library=spir64-unknown-unknown=%t.devicelib.bc -sycl-post-link-options=SYCL_POST_LINK_OPTIONS --host-triple=x86_64-unknown-linux-gnu --linker-path=/usr/bin/ld -o /dev/null %t.o --dry-run -syclbin=executable 2>&1 | FileCheck --check-prefix CHK-SYCLBIN-CMDS %s // CHK-SYCLBIN-CMDS: spirv-to-ir-wrapper{{.*}} -o [[FIRSTLLVMLINKIN:.*]].bc --llvm-spirv-opts --spirv-preserve-auxdata --spirv-target-env=SPV-IR --spirv-builtin-format=global diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index d2e66ccf44dbf..147fcc102ba24 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1215,13 +1215,17 @@ static Error compressImages(SmallVectorImpl &Images, return Error::success(); } -/// Reads device images from the given \p InputFile and wraps them -/// in one LLVM IR Module as a constant data. +/// Wraps the device images of \p SplitModules, built for the device target +/// triple \p Target, in one LLVM IR module as constant data. If +/// \p IsEmbeddedIR is true, the images are LLVM IR to be compiled at runtime. +/// \p Args provides the global options, such as image compression and the +/// host triple. /// /// \returns A path to the LLVM Module that contains wrapped images. Expected wrapSYCLBinariesFromFile(ArrayRef SplitModules, - const ArgList &Args, bool IsEmbeddedIR) { + const ArgList &Args, bool IsEmbeddedIR, + const llvm::Triple &Target) { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); if (!OutputFileOrErr) @@ -1246,20 +1250,13 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, return OutputFilePath; } - StringRef Target = Args.getLastArgValue(OPT_triple_EQ); - if (Target.empty()) - return createStringError( - inconvertibleErrorCode(), - "can't wrap SYCL image. -triple argument is missing."); - SmallVector Images; // SYCL runtime currently works for spir64 target triple and not for // spir64-unknown-unknown/spirv64-unknown-unknown/spirv64. // TODO: Fix SYCL runtime to accept other triples - llvm::Triple T(Target); std::string EmbeddedIRTarget("llvm_"); - EmbeddedIRTarget.append(T.getArchName()); - StringRef RegularTarget(T.getArchName()); + EmbeddedIRTarget.append(Target.getArchName()); + StringRef RegularTarget(Target.getArchName()); if (RegularTarget == "spirv64") RegularTarget = "spir64"; @@ -1645,8 +1642,13 @@ static Error checkForUndefinedSYCLBINSymbols(const SYCLBINSymbols &Symbols) { static Expected runWrapperAndCompile(ArrayRef SplitModules, const ArgList &Args, bool IsEmbeddedIR = false) { - auto OutputFile = - sycl::wrapSYCLBinariesFromFile(SplitModules, Args, IsEmbeddedIR); + StringRef Target = Args.getLastArgValue(OPT_triple_EQ); + if (Target.empty()) + return createStringError( + inconvertibleErrorCode(), + "can't wrap SYCL image. -triple argument is missing."); + auto OutputFile = sycl::wrapSYCLBinariesFromFile( + SplitModules, Args, IsEmbeddedIR, llvm::Triple(Target)); if (!OutputFile) return OutputFile.takeError(); // call to clang @@ -2955,6 +2957,36 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, std::mutex SYCLBINModulesMtx; SmallVector SYCLBINModules; + // The -fno-sycl-rdc compile step wraps the modules of all architectures into + // a single wrapper module. One vector of split modules per input group, so + // the wrap order doesn't depend on thread scheduling. Each group only writes + // its own slot, so no lock is needed. + using SplitModulesVec = std::vector; + std::vector NoRDCModulesPerGroup(LinkerInputFiles.size()); + + // The single wrapper module has one target triple, so check before any + // device compilation that all SYCL input groups share it. + std::optional NoRDCTriple; + if (!NeedsWrapping && !OutputSYCLBIN) { + for (const SmallVector &Input : LinkerInputFiles) { + const auto *It = llvm::find_if(Input, [](const OffloadFile &F) { + return F.getBinary()->getOffloadKind() == OFK_SYCL; + }); + if (It == Input.end()) + continue; + StringRef Triple = It->getBinary()->getTriple(); + if (Triple.empty()) + return createStringError( + "--no-sycl-rdc requires a target triple for every SYCL input"); + if (!NoRDCTriple) + NoRDCTriple = Triple.str(); + else if (*NoRDCTriple != Triple) + return createStringError( + "--no-sycl-rdc does not support multiple SYCL target triples ('" + + *NoRDCTriple + "' and '" + Triple + "')"); + } + } + // Initialize the images with any overriding inputs. if (Args.hasArg(OPT_override_image)) if (Error Err = handleOverrideImages(Args, Images)) @@ -3060,14 +3092,10 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, std::scoped_lock Guard(SYCLBINModulesMtx); SYCLBINModules.emplace_back(std::move(MD)); } else if (!NeedsWrapping) { - // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the - // host compilation links it in via -foffload-include-binary. - Expected OutputFile = - sycl::wrapSYCLBinariesFromFile(Modules, LinkerArgs, - /*IsEmbeddedIR=*/false); - if (!OutputFile) - return OutputFile.takeError(); - AppendImageToWrapperOutput(*OutputFile); + // -fno-sycl-rdc compile step: the modules are wrapped once all inputs + // are processed. + size_t GroupIdx = std::distance(LinkerInputFiles.begin(), &Input); + NoRDCModulesPerGroup[GroupIdx] = std::move(Modules); } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap Expected OutputFile = @@ -3140,6 +3168,20 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, if (Err) return std::move(Err); + // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the host + // compilation links it in via -foffload-include-binary. + if (NoRDCTriple) { + SplitModulesVec Modules; + for (SplitModulesVec &GroupModules : NoRDCModulesPerGroup) + std::move(GroupModules.begin(), GroupModules.end(), + std::back_inserter(Modules)); + Expected OutputFile = sycl::wrapSYCLBinariesFromFile( + Modules, Args, /*IsEmbeddedIR=*/false, llvm::Triple(*NoRDCTriple)); + if (!OutputFile) + return OutputFile.takeError(); + WrappedOutput.push_back(*OutputFile); + } + if (OutputSYCLBIN) { auto OutputOrErr = sycl::packageSYCLBIN(SYCLBINState, SYCLBINModules); if (!OutputOrErr) diff --git a/sycl/test-e2e/AOT/early_aot_multi_arch.cpp b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp new file mode 100644 index 0000000000000..5b3cdce617f38 --- /dev/null +++ b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp @@ -0,0 +1,13 @@ +// Check that a -fno-sycl-rdc object built for two GPU architectures carries +// the device images of both and runs on either device. + +// REQUIRES: ocloc, target-spir +// REQUIRES: arch-intel_gpu_acm_g10 || arch-intel_gpu_bmg_g21 + +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_acm_g10,intel_gpu_bmg_g21 -fno-sycl-rdc -c %S/Inputs/simple.cpp -o %t.o +// With the new offloading model the object already contains the wrapped device +// images, so a plain host link is enough. The old model extracts them from the +// object's offload bundle at a -fsycl link. On Windows, -fsycl also selects the +// dynamic C runtime that the object was compiled for. +// RUN: %clangxx %if new-offload-model && linux %{ -Wno-unused-command-line-argument %t.o -o %t.out %sycl_options %} %else %{ -fsycl %t.o -o %t.out %} +// RUN: %{run} %t.out