[SYCL][clang-linker-wrapper] Wrap all architectures in one module for -fno-sycl-rdc

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

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

@@ -1215,13 +1215,17 @@ static Error compressImages(SmallVectorImpl<offloading::SYCLImage> &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<StringRef>
 wrapSYCLBinariesFromFile(ArrayRef<module_split::SplitModule> 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<module_split::SplitModule> 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<llvm::offloading::SYCLImage> 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<StringRef>
 runWrapperAndCompile(ArrayRef<module_split::SplitModule> 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<SmallVector<OffloadFile>> LinkerInputFiles,
   std::mutex SYCLBINModulesMtx;
   SmallVector<SYCLBIN::SYCLBINModuleDesc> 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<module_split::SplitModule>;
+  std::vector<SplitModulesVec> 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<std::string> NoRDCTriple;
+  if (!NeedsWrapping && !OutputSYCLBIN) {
+    for (const SmallVector<OffloadFile> &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<SmallVector<OffloadFile>> LinkerInputFiles,
         std::scoped_lock<std::mutex> 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<StringRef> 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<StringRef> OutputFile =
@@ -3140,6 +3168,20 @@ linkAndWrapDeviceFiles(ArrayRef<SmallVector<OffloadFile>> 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<StringRef> 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)

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