From 5846c09c22b57e1422f58e40815fa2b6641137b8 Mon Sep 17 00:00:00 2001 From: Nick Sarnie Date: Thu, 29 Oct 2026 12:23:49 -0700 Subject: [PATCH] [Clang][OffloadBundler] Extract SYCL code without a target ID for each target ID When unbundling, a bundle entry was only extracted for the first requested target it matched. A SYCL entry without a target ID, such as from -fsycl-targets=spir64_gen, is compatible with any target ID of the same triple, so with e.g. -fsycl-targets=intel_gpu_tgllp,intel_gpu_pvc it was only extracted for one of them, picked by hash order, and the others got no device code. Linking such an object or static archive for several architectures then failed at runtime with "Kernel ... has no image for the selected device". Extract such an entry for every target ID that has no exact entry. The entry is read once and copied to each output, as reading it may advance the file handler. Co-Authored-By: Claude Opus 5.5 --- clang/lib/Driver/OffloadBundler.cpp | 56 +++++++++++-------- .../clang-offload-bundler/sycl-target-id.c | 30 ++++++++++ .../AOT/spir64_gen_multi_arch_link.cpp | 53 ++++++++++++++++++ 3 files changed, 116 insertions(+), 23 deletions(-) create mode 100644 sycl/test-e2e/AOT/spir64_gen_multi_arch_link.cpp diff --git a/clang/lib/Driver/OffloadBundler.cpp b/clang/lib/Driver/OffloadBundler.cpp index 2caafcdd4a657..bfee398c15c25 100644 --- a/clang/lib/Driver/OffloadBundler.cpp +++ b/clang/lib/Driver/OffloadBundler.cpp @@ -2031,39 +2031,49 @@ Error OffloadBundler::UnbundleFiles() { "invalid bundle id read from the bundle"); OffloadTargetInfo CodeObject(CurTriple, BundlerConfig); - // First, look for a requested -targets ID (Target) that this bundle - // entry (CodeObject) matches exactly, e.g. entry - // sycl-spir64_gen-unknown-unknown-pvc for target + // Collect the requested -targets IDs (Target) to extract this bundle + // entry (CodeObject) for. First, a target the entry matches exactly, + // e.g. entry sycl-spir64_gen-unknown-unknown-pvc for target // sycl-spir64_gen-unknown-unknown-pvc. - auto Output = llvm::find_if(Worklist, [&](const auto &W) { + SmallVector Outputs; + auto Exact = llvm::find_if(Worklist, [&](const auto &W) { OffloadTargetInfo Target(W.first(), BundlerConfig); return TargetIDPresenceMatches(CodeObject, Target) && isCodeObjectCompatible(CodeObject, Target); }); - // Otherwise, an ID-less SYCL entry (e.g. sycl-spir64_gen-unknown-unknown) - // may be used for a target with a target ID, or the reverse, but only if - // the blob has no exact entry for that target. - if (Output == Worklist.end()) - Output = llvm::find_if(Worklist, [&](const auto &W) { - OffloadTargetInfo Target(W.first(), BundlerConfig); - return !TargetIDPresenceMatches(CodeObject, Target) && - isCodeObjectCompatible(CodeObject, Target) && - !HasExactEntry(Target); - }); - - if (Output == Worklist.end()) + if (Exact != Worklist.end()) + Outputs.push_back(Exact->first()); + // Then, an ID-less SYCL entry (e.g. sycl-spir64_gen-unknown-unknown) is + // used for every target with a target ID, or the reverse, for which the + // blob has no exact entry. + for (const auto &W : Worklist) { + OffloadTargetInfo Target(W.first(), BundlerConfig); + if (!TargetIDPresenceMatches(CodeObject, Target) && + isCodeObjectCompatible(CodeObject, Target) && + !HasExactEntry(Target)) + Outputs.push_back(W.first()); + } + + if (Outputs.empty()) continue; - // Check if the output file can be opened and copy the bundle to it. - std::error_code EC; - raw_fd_ostream OutputFile((*Output).second, EC, sys::fs::OF_None); - if (EC) - return createFileError((*Output).second, EC); - if (Error Err = FH->ReadBundle(OutputFile, Input)) + // Read the bundle once, as reading it may advance the file handler, and + // copy it to each output file. + SmallVector Bundle; + raw_svector_ostream BundleOS(Bundle); + if (Error Err = FH->ReadBundle(BundleOS, Input)) return Err; if (Error Err = FH->ReadBundleEnd(Input)) return Err; - Worklist.erase(Output); + for (StringRef Target : Outputs) { + StringRef OutputFileName = Worklist.lookup(Target); + std::error_code EC; + raw_fd_ostream OutputFile(OutputFileName, EC, sys::fs::OF_None); + if (EC) + return createFileError(OutputFileName, EC); + OutputFile << BundleOS.str(); + Worklist.erase(Target); + } // Record if we found the host bundle. auto OffloadInfo = OffloadTargetInfo(CurTriple, BundlerConfig); diff --git a/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c b/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c index 18fe06195817e..35c9fba8f440b 100644 --- a/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c +++ b/clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c @@ -69,6 +69,36 @@ // RUN: clang-offload-bundler -type=ii -input=%t.bundle.mixed.ii -output=%t.res -unbundle \ // RUN: -targets=sycl-spir64_gen-unknown-unknown-acm_g10 // RUN: diff %t.noid %t.res +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.mixed.ii -output=%t.res1 -output=%t.res2 -output=%t.res3 -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-pvc,sycl-spir64_gen-unknown-unknown-acm_g10,sycl-spir64_gen-unknown-unknown +// RUN: diff %t.tgt2 %t.res1 +// RUN: diff %t.noid %t.res2 +// RUN: diff %t.noid %t.res3 + +// The entry without a target ID is used for each of several target IDs, as +// for -fsycl-targets=intel_gpu_tgllp,intel_gpu_pvc. +// RUN: clang-offload-bundler -type=ii -input=%t.bundle.noid.ii -output=%t.res1 -output=%t.res2 -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt1 %t.res1 +// RUN: diff %t.tgt1 %t.res2 +// RUN: clang-offload-bundler -type=o -input=%t.o -input=%t.tgt1 -output=%t.bundle.noid.o \ +// RUN: -targets=host-x86_64-unknown-linux-gnu,sycl-spir64_gen-unknown-unknown +// RUN: clang-offload-bundler -type=o -input=%t.bundle.noid.o -output=%t.res1 -output=%t.res2 -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt1 %t.res1 +// RUN: diff %t.tgt1 %t.res2 +// RUN: rm -f %t.noid.a +// RUN: llvm-ar cr %t.noid.a %t.bundle.noid.o +// RUN: clang-offload-bundler -type=ao -input=%t.noid.a -output=%t.res1 -output=%t.res2 -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: diff %t.tgt1 %t.res1 +// RUN: diff %t.tgt1 %t.res2 +// RUN: clang-offload-bundler -type=aoo -input=%t.noid.a -output=%t.res1.lst -output=%t.res2.lst -unbundle \ +// RUN: -targets=sycl-spir64_gen-unknown-unknown-tgllp,sycl-spir64_gen-unknown-unknown-pvc +// RUN: tr -d '\n' < %t.res1.lst > %t.res1.path +// RUN: diff %t.tgt1 %{readfile:%t.res1.path} +// RUN: tr -d '\n' < %t.res2.lst > %t.res2.path +// RUN: diff %t.tgt1 %{readfile:%t.res2.path} // 4-component triples are unaffected. // RUN: clang-offload-bundler -type=ii -input=%t.host -input=%t.tgt1 -input=%t.tgt2 -output=%t.bundle.hip.ii \ diff --git a/sycl/test-e2e/AOT/spir64_gen_multi_arch_link.cpp b/sycl/test-e2e/AOT/spir64_gen_multi_arch_link.cpp new file mode 100644 index 0000000000000..b21738bea5f6c --- /dev/null +++ b/sycl/test-e2e/AOT/spir64_gen_multi_arch_link.cpp @@ -0,0 +1,53 @@ +// Check that device code compiled for spir64_gen, which has no target ID, is +// used for each of several GPU architectures when linking objects and a static +// archive, and runs on any of these devices. + +// REQUIRES: ocloc, target-spir +// REQUIRES: arch-intel_gpu_pvc || gpu-intel-dg2 || arch-intel_gpu_bmg_g21 +// UNSUPPORTED: new-offload-model +// UNSUPPORTED-INTENDED: The new offloading model does not support +// -fsycl-targets=spir64_gen. + +// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Wno-deprecated -DSOURCE1 -c %s -o %t1.o +// RUN: %clangxx -fsycl -fsycl-targets=spir64_gen -Wno-deprecated -DSOURCE2 -c %s -o %t2.o +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2,intel_gpu_bmg_g21 -Wno-sycl-target %t1.o %t2.o -o %t.out +// RUN: %{run} %t.out + +// The kernel is in an object so that it pulls the device function from the +// archive. +// RUN: rm -f %t.a +// RUN: llvm-ar cr %t.a %t2.o +// RUN: %clangxx -fsycl -fsycl-targets=intel_gpu_pvc,intel_gpu_dg2,intel_gpu_bmg_g21 -Wno-sycl-target %t1.o %t.a -o %t.a.out +// RUN: %{run} %t.a.out + +#include +#include + +SYCL_EXTERNAL int get_value(); + +#ifdef SOURCE1 +int main() { + sycl::queue q; + + int data = 0; + { + sycl::buffer buf(&data, sycl::range<1>(1)); + + q.submit([&](sycl::handler &h) { + auto acc = buf.get_access(h); + h.single_task([=]() { acc[0] = get_value(); }); + }).wait(); + } + + if (data != 42) { + std::cerr << "Error: expected 42, got " << data << std::endl; + return 1; + } + + return 0; +} +#endif // SOURCE1 + +#ifdef SOURCE2 +SYCL_EXTERNAL int get_value() { return 42; } +#endif // SOURCE2