From 4ff6a3c1ada5e2a180b641b2654622f17ef5149d Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Tue, 13 Oct 2026 15:10:18 -0700 Subject: [PATCH 1/8] [SYCL][clang-linker-wrapper] Wrap all architectures in one module for -fno-sycl-rdc The -fno-sycl-rdc compile step (--emit-fatbin-only) wrapped the images of each target group separately and then expected a single output, so a compile for several architectures of one triple, e.g. -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2, failed with "Expect single output from the SYCL device linker". Collect the images of all groups and wrap them into a single wrapper module, as is done for the other offload kinds. The per-architecture information, such as the compile_target, is carried by each image. Inputs with different target triples, which the driver already rejects with -fno-sycl-rdc, are diagnosed rather than labelled with one triple. --- .../linker-wrapper-sycl.cpp | 11 +++++ .../ClangLinkerWrapper.cpp | 46 ++++++++++++++----- 2 files changed, 46 insertions(+), 11 deletions(-) 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..4c9642072cd3e 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -6,6 +6,7 @@ // // 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 +349,16 @@ // 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: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}}, input: {{.*}} +// 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 + // 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..9948bb5abcecb 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1216,12 +1216,14 @@ static Error compressImages(SmallVectorImpl &Images, } /// Reads device images from the given \p InputFile and wraps them -/// in one LLVM IR Module as a constant data. +/// in one LLVM IR Module as a constant data. \p Target overrides the target +/// triple given by \p Args. /// /// \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, + StringRef Target = "") { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); if (!OutputFileOrErr) @@ -1246,7 +1248,8 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, return OutputFilePath; } - StringRef Target = Args.getLastArgValue(OPT_triple_EQ); + if (Target.empty()) + Target = Args.getLastArgValue(OPT_triple_EQ); if (Target.empty()) return createStringError( inconvertibleErrorCode(), @@ -2955,6 +2958,12 @@ 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. Keep them per input group for a stable order. + std::vector> NoRDCModules( + LinkerInputFiles.size()); + std::string NoRDCTriple; + // Initialize the images with any overriding inputs. if (Args.hasArg(OPT_override_image)) if (Error Err = handleOverrideImages(Args, Images)) @@ -3060,14 +3069,15 @@ 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. + StringRef Triple = LinkerArgs.getLastArgValue(OPT_triple_EQ); + std::scoped_lock Guard(ImageMtx); + if (!NoRDCTriple.empty() && NoRDCTriple != Triple) + return createStringError( + "--no-sycl-rdc does not support multiple SYCL target triples"); + NoRDCModules[&Input - LinkerInputFiles.data()] = std::move(Modules); + NoRDCTriple = Triple.str(); } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap Expected OutputFile = @@ -3140,6 +3150,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.empty()) { + std::vector Modules; + for (std::vector &GroupModules : NoRDCModules) + std::move(GroupModules.begin(), GroupModules.end(), + std::back_inserter(Modules)); + Expected OutputFile = sycl::wrapSYCLBinariesFromFile( + Modules, Args, /*IsEmbeddedIR=*/false, NoRDCTriple); + if (!OutputFile) + return OutputFile.takeError(); + WrappedOutput.push_back(*OutputFile); + } + if (OutputSYCLBIN) { auto OutputOrErr = sycl::packageSYCLBIN(SYCLBINState, SYCLBINModules); if (!OutputOrErr) From c56a019678b8d4f4d15214c23388d33c28e6c082 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 14:06:35 -0700 Subject: [PATCH 2/8] Trigger CI with the new-offload-model label From aa39cb9c13e0fa3945cbe9b0772c827402c1eee1 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Wed, 14 Oct 2026 15:37:42 -0700 Subject: [PATCH 3/8] Name the triples in the --no-sycl-rdc mixed-triple error Also diagnose a missing triple, use a dedicated mutex and explain why the global arguments are used for wrapping. --- .../clang-linker-wrapper/linker-wrapper-sycl.cpp | 2 +- .../clang-linker-wrapper/ClangLinkerWrapper.cpp | 15 +++++++++++---- 2 files changed, 12 insertions(+), 5 deletions(-) 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 4c9642072cd3e..afe8937acd392 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -357,7 +357,7 @@ // 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 +// CHK-FATBIN-ONLY-MIXED: error: --no-sycl-rdc does not support multiple SYCL target triples ('{{.*}}' and '{{.*}}') // 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 diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 9948bb5abcecb..820223b504680 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -2963,6 +2963,7 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, std::vector> NoRDCModules( LinkerInputFiles.size()); std::string NoRDCTriple; + std::mutex NoRDCMtx; // Initialize the images with any overriding inputs. if (Args.hasArg(OPT_override_image)) @@ -3072,11 +3073,16 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // -fno-sycl-rdc compile step: the modules are wrapped once all inputs // are processed. StringRef Triple = LinkerArgs.getLastArgValue(OPT_triple_EQ); - std::scoped_lock Guard(ImageMtx); + if (Triple.empty()) + return createStringError( + "-triple argument is missing for --no-sycl-rdc"); + size_t GroupIdx = &Input - LinkerInputFiles.data(); + std::scoped_lock Guard(NoRDCMtx); if (!NoRDCTriple.empty() && NoRDCTriple != Triple) return createStringError( - "--no-sycl-rdc does not support multiple SYCL target triples"); - NoRDCModules[&Input - LinkerInputFiles.data()] = std::move(Modules); + "--no-sycl-rdc does not support multiple SYCL target triples ('" + + NoRDCTriple + "' and '" + Triple + "')"); + NoRDCModules[GroupIdx] = std::move(Modules); NoRDCTriple = Triple.str(); } else { // TODO(NOM7): Remove this call and use community flow for bundle/wrap @@ -3151,7 +3157,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, 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. + // compilation links it in via -foffload-include-binary. The per-input + // arguments only differ from Args in the triple, which is passed explicitly. if (!NoRDCTriple.empty()) { std::vector Modules; for (std::vector &GroupModules : NoRDCModules) From 240a124d9ddb055f4be8b7f7759df2be4dbcb049 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 16 Oct 2026 15:15:25 -0700 Subject: [PATCH 4/8] Check the --no-sycl-rdc target triples before device compilation Collect the per-group modules without a lock, make the target triple a required argument of wrapSYCLBinariesFromFile, and check in the test that the images of both architectures reach the one wrapper module. --- .../linker-wrapper-sycl.cpp | 6 +- .../ClangLinkerWrapper.cpp | 88 +++++++++++-------- 2 files changed, 54 insertions(+), 40 deletions(-) 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 afe8937acd392..7987ee300d61c 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -3,6 +3,8 @@ // 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 @@ -352,7 +354,9 @@ // 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: offload-wrapper: output: [[WRAPPEROUT:.*]].bc, input: {{.*}}, input: {{.*}} +// 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" diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 820223b504680..6f50902672f18 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1215,15 +1215,16 @@ 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. \p Target overrides the target -/// triple given by \p Args. +/// 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: image compression and host triple. /// /// \returns A path to the LLVM Module that contains wrapped images. Expected wrapSYCLBinariesFromFile(ArrayRef SplitModules, const ArgList &Args, bool IsEmbeddedIR, - StringRef Target = "") { + const llvm::Triple &Target) { auto OutputFileOrErr = createOutputFile( sys::path::filename(ExecutableName) + ".sycl.image.wrapper", "bc"); if (!OutputFileOrErr) @@ -1248,21 +1249,13 @@ wrapSYCLBinariesFromFile(ArrayRef SplitModules, return OutputFilePath; } - if (Target.empty()) - 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"; @@ -1648,8 +1641,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 @@ -2959,11 +2957,34 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, SmallVector SYCLBINModules; // The -fno-sycl-rdc compile step wraps the modules of all architectures into - // a single wrapper module. Keep them per input group for a stable order. - std::vector> NoRDCModules( - LinkerInputFiles.size()); - std::string NoRDCTriple; - std::mutex NoRDCMtx; + // 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 SplitModuleList = 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)) @@ -3072,18 +3093,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, } else if (!NeedsWrapping) { // -fno-sycl-rdc compile step: the modules are wrapped once all inputs // are processed. - StringRef Triple = LinkerArgs.getLastArgValue(OPT_triple_EQ); - if (Triple.empty()) - return createStringError( - "-triple argument is missing for --no-sycl-rdc"); - size_t GroupIdx = &Input - LinkerInputFiles.data(); - std::scoped_lock Guard(NoRDCMtx); - if (!NoRDCTriple.empty() && NoRDCTriple != Triple) - return createStringError( - "--no-sycl-rdc does not support multiple SYCL target triples ('" + - NoRDCTriple + "' and '" + Triple + "')"); - NoRDCModules[GroupIdx] = std::move(Modules); - NoRDCTriple = Triple.str(); + 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 = @@ -3157,15 +3168,14 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, 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. The per-input - // arguments only differ from Args in the triple, which is passed explicitly. - if (!NoRDCTriple.empty()) { - std::vector Modules; - for (std::vector &GroupModules : NoRDCModules) + // compilation links it in via -foffload-include-binary. + if (NoRDCTriple) { + SplitModuleList Modules; + for (SplitModuleList &GroupModules : NoRDCModulesPerGroup) std::move(GroupModules.begin(), GroupModules.end(), std::back_inserter(Modules)); Expected OutputFile = sycl::wrapSYCLBinariesFromFile( - Modules, Args, /*IsEmbeddedIR=*/false, NoRDCTriple); + Modules, Args, /*IsEmbeddedIR=*/false, llvm::Triple(*NoRDCTriple)); if (!OutputFile) return OutputFile.takeError(); WrappedOutput.push_back(*OutputFile); From 33a8f57c9ae7089bf4089e80721c7305f88c4819 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Fri, 16 Oct 2026 15:15:26 -0700 Subject: [PATCH 5/8] [SYCL][E2E] Run a -fno-sycl-rdc object built for two GPU architectures --- sycl/test-e2e/AOT/early_aot_multi_arch.cpp | 9 +++++++++ 1 file changed, 9 insertions(+) create mode 100644 sycl/test-e2e/AOT/early_aot_multi_arch.cpp 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..8d5009d3f5fed --- /dev/null +++ b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp @@ -0,0 +1,9 @@ +// 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 +// RUN: %clangxx -fsycl %t.o -o %t.out +// RUN: %{run} %t.out From 8b8ff56274b434ee51e23593b320ec184071ddc1 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 19 Oct 2026 13:45:17 -0700 Subject: [PATCH 6/8] Rename SplitModuleList to SplitModulesVec and test the missing triple error --- .../clang-linker-wrapper/linker-wrapper-sycl.cpp | 6 ++++++ .../tools/clang-linker-wrapper/ClangLinkerWrapper.cpp | 11 ++++++----- 2 files changed, 12 insertions(+), 5 deletions(-) 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 7987ee300d61c..eff9831b8ffbd 100644 --- a/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp +++ b/clang/test/OffloadTools/clang-linker-wrapper/linker-wrapper-sycl.cpp @@ -362,6 +362,12 @@ // 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 diff --git a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp index 6f50902672f18..147fcc102ba24 100644 --- a/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp +++ b/clang/tools/clang-linker-wrapper/ClangLinkerWrapper.cpp @@ -1218,7 +1218,8 @@ static Error compressImages(SmallVectorImpl &Images, /// 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: image compression and host triple. +/// \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 @@ -2960,8 +2961,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // 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 SplitModuleList = std::vector; - std::vector NoRDCModulesPerGroup(LinkerInputFiles.size()); + 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. @@ -3170,8 +3171,8 @@ linkAndWrapDeviceFiles(ArrayRef> LinkerInputFiles, // -fno-sycl-rdc compile step: emit the wrapper module as bitcode; the host // compilation links it in via -foffload-include-binary. if (NoRDCTriple) { - SplitModuleList Modules; - for (SplitModuleList &GroupModules : NoRDCModulesPerGroup) + SplitModulesVec Modules; + for (SplitModulesVec &GroupModules : NoRDCModulesPerGroup) std::move(GroupModules.begin(), GroupModules.end(), std::back_inserter(Modules)); Expected OutputFile = sycl::wrapSYCLBinariesFromFile( From c99f528f842c46dee76dd3be3a06d25adb32e213 Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 19 Oct 2026 14:18:38 -0700 Subject: [PATCH 7/8] [SYCL][E2E] Link the -fno-sycl-rdc object without -fsycl in the new offloading model --- sycl/test-e2e/AOT/early_aot_multi_arch.cpp | 5 ++++- 1 file changed, 4 insertions(+), 1 deletion(-) diff --git a/sycl/test-e2e/AOT/early_aot_multi_arch.cpp b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp index 8d5009d3f5fed..cb1e478534a4b 100644 --- a/sycl/test-e2e/AOT/early_aot_multi_arch.cpp +++ b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp @@ -5,5 +5,8 @@ // 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 -// RUN: %clangxx -fsycl %t.o -o %t.out +// 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. +// RUN: %clangxx %if new-offload-model %{ -Wno-unused-command-line-argument %t.o -o %t.out %sycl_options %} %else %{ -fsycl %t.o -o %t.out %} // RUN: %{run} %t.out From 50b44e3a106463d09398a4020e59d477662af71d Mon Sep 17 00:00:00 2001 From: srividya sundaram Date: Mon, 19 Oct 2026 16:35:57 -0700 Subject: [PATCH 8/8] [SYCL][E2E] Keep the -fsycl link of early_aot_multi_arch.cpp on Windows Without -fsycl, the Windows link uses the static C runtime, but the object was compiled for the dynamic one (__imp__wassert is unresolved). --- sycl/test-e2e/AOT/early_aot_multi_arch.cpp | 5 +++-- 1 file changed, 3 insertions(+), 2 deletions(-) diff --git a/sycl/test-e2e/AOT/early_aot_multi_arch.cpp b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp index cb1e478534a4b..5b3cdce617f38 100644 --- a/sycl/test-e2e/AOT/early_aot_multi_arch.cpp +++ b/sycl/test-e2e/AOT/early_aot_multi_arch.cpp @@ -7,6 +7,7 @@ // 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. -// RUN: %clangxx %if new-offload-model %{ -Wno-unused-command-line-argument %t.o -o %t.out %sycl_options %} %else %{ -fsycl %t.o -o %t.out %} +// 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