Skip to content
Draft
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
56 changes: 33 additions & 23 deletions clang/lib/Driver/OffloadBundler.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -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<StringRef, 4> 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<char, 0> 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);
Expand Down
30 changes: 30 additions & 0 deletions clang/test/OffloadTools/clang-offload-bundler/sycl-target-id.c
Original file line number Diff line number Diff line change
Expand Up @@ -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 \
Expand Down
53 changes: 53 additions & 0 deletions sycl/test-e2e/AOT/spir64_gen_multi_arch_link.cpp
Original file line number Diff line number Diff line change
@@ -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 <iostream>
#include <sycl/detail/core.hpp>

SYCL_EXTERNAL int get_value();

#ifdef SOURCE1
int main() {
sycl::queue q;

int data = 0;
{
sycl::buffer<int, 1> buf(&data, sycl::range<1>(1));

q.submit([&](sycl::handler &h) {
auto acc = buf.get_access<sycl::access_mode::write>(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
Loading