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