[clang] [llvm] [SYCL] Shorten triple form in device binary. (PR #231446)
Kseniya Tikhomirova via cfe-commits
cfe-commits at lists.llvm.org
Thu Oct 15 03:30:02 PDT 2026
https://github.com/KseniyaTikhomirova updated https://github.com/llvm/llvm-project/pull/231446
>From 78c74be3632f7a17e452d90a387eb6ff6f78ef84 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Tue, 13 Oct 2026 03:14:00 -0700
Subject: [PATCH 1/2] [SYCL] Shorten triple form in device binary
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
.../OffloadTools/clang-sycl-linker/triple.ll | 18 ++++++++++++++++++
.../clang-sycl-linker/ClangSYCLLinker.cpp | 6 +++++-
.../src/detail/device_binary_structures.hpp | 4 +++-
3 files changed, 26 insertions(+), 2 deletions(-)
diff --git a/clang/test/OffloadTools/clang-sycl-linker/triple.ll b/clang/test/OffloadTools/clang-sycl-linker/triple.ll
index 7f432296acb8c4..683941bab0e7c8 100644
--- a/clang/test/OffloadTools/clang-sycl-linker/triple.ll
+++ b/clang/test/OffloadTools/clang-sycl-linker/triple.ll
@@ -10,6 +10,15 @@
; RUN: | FileCheck %s --check-prefix=NO-TRIPLE-INPUT
; NO-TRIPLE-INPUT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
;
+; Test that only the architecture is recorded in the offload image, both when
+; the full triple comes from -triple= and when it is inferred from the input.
+; RUN: clang-sycl-linker --dry-run -v -triple=spirv64-unknown-unknown %t/no-triple.bc -o %t/short-explicit.out 2>&1 \
+; RUN: | FileCheck %s --check-prefix=SHORT-TRIPLE
+; RUN: llvm-as %t/full-triple.ll -o %t/full-triple.bc
+; RUN: clang-sycl-linker --dry-run -v %t/full-triple.bc -o %t/short-inferred.out 2>&1 \
+; RUN: | FileCheck %s --check-prefix=SHORT-TRIPLE
+; SHORT-TRIPLE: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+;
; Test that triple was inferred from inputs and recorded in the offload image.
; RUN: llvm-as %t/input1.ll -o %t/input1.bc
; RUN: llvm-as %t/input2.ll -o %t/input2.bc
@@ -81,6 +90,15 @@ define spir_kernel void @kernel_d() #0 {
attributes #0 = { "sycl-module-id"="TU4.cpp" }
+;--- full-triple.ll
+target triple = "spirv64-unknown-unknown"
+
+define spir_kernel void @kernel_e() #0 {
+ ret void
+}
+
+attributes #0 = { "sycl-module-id"="TU5.cpp" }
+
;--- archive-full-triple.ll
target triple = "spirv64-unknown-unknown"
diff --git a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
index 60ab7dc4e2a7d1..e65ea5eb29b61b 100644
--- a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
+++ b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
@@ -1043,8 +1043,12 @@ static Error runSYCLLink(ArrayRef<std::unique_ptr<MemoryBuffer>> Inputs,
OffloadingImage TheImage{};
TheImage.TheImageKind = IsAOTCompileNeeded ? IMG_Object : IMG_SPIRV;
TheImage.TheOffloadKind = OFK_SYCL;
+ // Pack the short form of the triple. The runtime only needs to know which
+ // binary format the image holds, and the architecture is the single
+ // canonical spelling of it: inputs may spell the same target as 'spirv64'
+ // or 'spirv64-unknown-unknown'.
TheImage.StringData["triple"] =
- Args.MakeArgString(Result.TargetTriple.str());
+ Args.MakeArgString(Result.TargetTriple.getArchName());
TheImage.StringData["arch"] =
Args.MakeArgString(Args.getLastArgValue(OPT_arch_EQ));
TheImage.StringData["symbols"] = SI.Symbols;
diff --git a/libsycl/src/detail/device_binary_structures.hpp b/libsycl/src/detail/device_binary_structures.hpp
index 29e71b45b33542..93aaafa8565f1c 100644
--- a/libsycl/src/detail/device_binary_structures.hpp
+++ b/libsycl/src/detail/device_binary_structures.hpp
@@ -26,9 +26,11 @@ namespace detail {
/// A device type represented by a particular target
/// triple requires specific binary images. We need
/// to map the image type onto the device target triple.
+/// clang-sycl-linker packs the short form of the triple, i.e. the
+/// architecture only, so these strings carry no vendor/OS components.
/// SPIR-V with 64-bit pointers.
-static constexpr char DeviceBinaryTripleSPIRV64[] = "spirv64-unknown-unknown";
+static constexpr char DeviceBinaryTripleSPIRV64[] = "spirv64";
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
>From 2c4d3a761cdd438024e2a1ee281ef467a9a02b69 Mon Sep 17 00:00:00 2001
From: "Tikhomirova, Kseniya" <kseniya.tikhomirova at intel.com>
Date: Thu, 15 Oct 2026 03:23:31 -0700
Subject: [PATCH 2/2] redesign
Signed-off-by: Tikhomirova, Kseniya <kseniya.tikhomirova at intel.com>
---
.../OffloadTools/clang-sycl-linker/basic.ll | 12 +--
.../clang-sycl-linker/split-mode.ll | 12 +--
.../OffloadTools/clang-sycl-linker/triple.ll | 18 ++--
.../clang-sycl-linker/ClangSYCLLinker.cpp | 10 +-
.../src/detail/device_binary_structures.hpp | 8 +-
libsycl/src/detail/device_image_wrapper.hpp | 9 +-
libsycl/src/detail/program_manager.cpp | 5 +-
libsycl/unittests/common/device_images.hpp | 19 ++--
.../common/scoped_binary_registration.hpp | 11 +-
.../unittests/program_manager/CMakeLists.txt | 1 +
.../program_manager/image_compatibility.cpp | 100 ++++++++++++++++++
11 files changed, 160 insertions(+), 45 deletions(-)
create mode 100644 libsycl/unittests/program_manager/image_compatibility.cpp
diff --git a/clang/test/OffloadTools/clang-sycl-linker/basic.ll b/clang/test/OffloadTools/clang-sycl-linker/basic.ll
index 0ba81f61005a43..edd17f0c90217b 100644
--- a/clang/test/OffloadTools/clang-sycl-linker/basic.ll
+++ b/clang/test/OffloadTools/clang-sycl-linker/basic.ll
@@ -30,7 +30,7 @@
; RUN: | FileCheck %s --check-prefix=SIMPLE-FO
; SIMPLE-FO: link: inputs: {{.*}}.bc, {{.*}}.bc output: [[LLVMLINKOUT:.*]].bc
; SIMPLE-FO-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}.spv
-; SIMPLE-FO-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; SIMPLE-FO-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
; SIMPLE-FO-NOT: {{.+}}
;
; Test the dry run of a simple case with device library archive specified using --whole-archive.
@@ -42,7 +42,7 @@
; RUN: | FileCheck %s --check-prefix=DEVLIBS
; DEVLIBS: link: inputs: {{.*}}.bc, {{.*}}.bc, {{.*}}libdevice.a(lib1.bc), {{.*}}libdevice.a(lib2.bc) output: [[LLVMLINKOUT:.*]].bc
; DEVLIBS-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}.spv
-; DEVLIBS-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; DEVLIBS-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
; DEVLIBS-NOT: {{.+}}
;
; Test -L short form (joined) and -l with archive using --whole-archive.
@@ -115,7 +115,7 @@
; AOT-INTEL-GPU: link: inputs: {{.*}}.bc, {{.*}}.bc output: [[LLVMLINKOUT:.*]].bc
; AOT-INTEL-GPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPVSTEM:.*]].spv
; AOT-INTEL-GPU-NEXT: "{{.*}}ocloc{{.*}}" {{.*}}-device bmg_g21 -a -b {{.*}}-output [[SPVSTEM]].out -file [[SPVSTEM]].spv
-; AOT-INTEL-GPU-NEXT: sycl-bundle: image kind: o, triple: spirv64, arch: bmg_g21
+; AOT-INTEL-GPU-NEXT: sycl-bundle: image kind: o, triple: spirv64-unknown-unknown, arch: bmg_g21
; AOT-INTEL-GPU-NOT: {{.+}}
;
; Test that all --ocloc-options are passed to ocloc, even if they contain spaces or quotes.
@@ -132,7 +132,7 @@
; AOT-INTEL-CPU: link: inputs: {{.*}}.bc, {{.*}}.bc output: [[LLVMLINKOUT:.*]].bc
; AOT-INTEL-CPU-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: [[SPVSTEM:.*]].spv
; AOT-INTEL-CPU-NEXT: "{{.*}}opencl-aot{{.*}}" {{.*}}--device=cpu -a -b {{.*}}-o [[SPVSTEM]].out [[SPVSTEM]].spv
-; AOT-INTEL-CPU-NEXT: sycl-bundle: image kind: o, triple: spirv64, arch: graniterapids
+; AOT-INTEL-CPU-NEXT: sycl-bundle: image kind: o, triple: spirv64-unknown-unknown, arch: graniterapids
; AOT-INTEL-CPU-NOT: {{.+}}
;
; Test that AOT temp file names ("<stem>-<unique>.spv"/"<stem>-<unique>.out")
@@ -169,10 +169,10 @@
; RUN: llvm-objdump --offloading %t/srcsplit.out | FileCheck %s --check-prefix=SRCSPLIT
; SRCSPLIT: OFFLOADING IMAGE [0]:
; SRCSPLIT: kind spir-v
-; SRCSPLIT: triple spirv64
+; SRCSPLIT: triple spirv64-unknown-unknown
; SRCSPLIT: OFFLOADING IMAGE [1]:
; SRCSPLIT: kind spir-v
-; SRCSPLIT: triple spirv64
+; SRCSPLIT: triple spirv64-unknown-unknown
;--- input1.ll
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
diff --git a/clang/test/OffloadTools/clang-sycl-linker/split-mode.ll b/clang/test/OffloadTools/clang-sycl-linker/split-mode.ll
index 7bc227aafc0079..1116b7863f2bdf 100644
--- a/clang/test/OffloadTools/clang-sycl-linker/split-mode.ll
+++ b/clang/test/OffloadTools/clang-sycl-linker/split-mode.ll
@@ -14,7 +14,7 @@
; RUN: | FileCheck %s --check-prefix=SPLIT-NONE
; SPLIT-NONE: link: inputs: {{.*}}.bc output: [[LLVMLINKOUT:.*]].bc
; SPLIT-NONE-NEXT: LLVM backend: input: [[LLVMLINKOUT]].bc, output: {{.*}}.spv
-; SPLIT-NONE-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; SPLIT-NONE-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
; SPLIT-NONE-NOT: {{.+}}
;
; Test the split mode ("kernel"): each SPIR_KERNEL function produces its own device image.
@@ -28,9 +28,9 @@
; SPLIT-KERNEL-NEXT: LLVM backend: input: [[SPLIT0]].bc, output: {{.*}}.spv
; SPLIT-KERNEL-NEXT: LLVM backend: input: [[SPLIT1]].bc, output: {{.*}}.spv
; SPLIT-KERNEL-NEXT: LLVM backend: input: [[SPLIT2]].bc, output: {{.*}}.spv
-; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
-; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
-; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
+; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
+; SPLIT-KERNEL-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
; SPLIT-KERNEL-NOT: {{.+}}
;
; Test default split mode ('translation_unit'): no --module-split-mode flag needed.
@@ -49,8 +49,8 @@
; SPLIT-SRC-NEXT: [[S1:.*]].bc [kernel_a ]
; SPLIT-SRC-NEXT: LLVM backend: input: [[S0]].bc, output: {{.*}}.spv
; SPLIT-SRC-NEXT: LLVM backend: input: [[S1]].bc, output: {{.*}}.spv
-; SPLIT-SRC-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
-; SPLIT-SRC-NEXT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; SPLIT-SRC-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
+; SPLIT-SRC-NEXT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
; SPLIT-SRC-NOT: {{.+}}
target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64-G1"
diff --git a/clang/test/OffloadTools/clang-sycl-linker/triple.ll b/clang/test/OffloadTools/clang-sycl-linker/triple.ll
index 683941bab0e7c8..43bc0d0a8b22bf 100644
--- a/clang/test/OffloadTools/clang-sycl-linker/triple.ll
+++ b/clang/test/OffloadTools/clang-sycl-linker/triple.ll
@@ -8,23 +8,23 @@
; RUN: llvm-as %t/no-triple.ll -o %t/no-triple.bc
; RUN: clang-sycl-linker --dry-run -v -triple=spirv64 %t/no-triple.bc -o %t/no-triple-input.out 2>&1 \
; RUN: | FileCheck %s --check-prefix=NO-TRIPLE-INPUT
-; NO-TRIPLE-INPUT: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; NO-TRIPLE-INPUT: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
;
-; Test that only the architecture is recorded in the offload image, both when
-; the full triple comes from -triple= and when it is inferred from the input.
-; RUN: clang-sycl-linker --dry-run -v -triple=spirv64-unknown-unknown %t/no-triple.bc -o %t/short-explicit.out 2>&1 \
-; RUN: | FileCheck %s --check-prefix=SHORT-TRIPLE
+; Test that an already three-part triple is recorded unchanged, whether it comes
+; from -triple= or is inferred from the input.
+; RUN: clang-sycl-linker --dry-run -v -triple=spirv64-unknown-unknown %t/no-triple.bc -o %t/full-explicit.out 2>&1 \
+; RUN: | FileCheck %s --check-prefix=FULL-TRIPLE
; RUN: llvm-as %t/full-triple.ll -o %t/full-triple.bc
-; RUN: clang-sycl-linker --dry-run -v %t/full-triple.bc -o %t/short-inferred.out 2>&1 \
-; RUN: | FileCheck %s --check-prefix=SHORT-TRIPLE
-; SHORT-TRIPLE: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; RUN: clang-sycl-linker --dry-run -v %t/full-triple.bc -o %t/full-inferred.out 2>&1 \
+; RUN: | FileCheck %s --check-prefix=FULL-TRIPLE
+; FULL-TRIPLE: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
;
; Test that triple was inferred from inputs and recorded in the offload image.
; RUN: llvm-as %t/input1.ll -o %t/input1.bc
; RUN: llvm-as %t/input2.ll -o %t/input2.bc
; RUN: clang-sycl-linker --dry-run -v --module-split-mode=link_unit %t/input1.bc %t/input2.bc -o %t/spirv.out 2>&1 \
; RUN: | FileCheck %s --check-prefix=TRIPLE-INFERENCE
-; TRIPLE-INFERENCE: sycl-bundle: image kind: spv, triple: spirv64, arch: {{$}}
+; TRIPLE-INFERENCE: sycl-bundle: image kind: spv, triple: spirv64-unknown-unknown, arch: {{$}}
;
; Test error on mismatched triple between inputs.
; RUN: llvm-as %t/input-mismatch.ll -o %t/input-mismatch.bc
diff --git a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
index e65ea5eb29b61b..53ef6375039ac3 100644
--- a/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
+++ b/clang/tools/clang-sycl-linker/ClangSYCLLinker.cpp
@@ -1043,12 +1043,10 @@ static Error runSYCLLink(ArrayRef<std::unique_ptr<MemoryBuffer>> Inputs,
OffloadingImage TheImage{};
TheImage.TheImageKind = IsAOTCompileNeeded ? IMG_Object : IMG_SPIRV;
TheImage.TheOffloadKind = OFK_SYCL;
- // Pack the short form of the triple. The runtime only needs to know which
- // binary format the image holds, and the architecture is the single
- // canonical spelling of it: inputs may spell the same target as 'spirv64'
- // or 'spirv64-unknown-unknown'.
- TheImage.StringData["triple"] =
- Args.MakeArgString(Result.TargetTriple.getArchName());
+ // Normalize to the three-part form so that the recorded string does not
+ // depend on how the inputs spelled the target.
+ TheImage.StringData["triple"] = Args.MakeArgString(
+ Result.TargetTriple.normalize(Triple::CanonicalForm::THREE_IDENT));
TheImage.StringData["arch"] =
Args.MakeArgString(Args.getLastArgValue(OPT_arch_EQ));
TheImage.StringData["symbols"] = SI.Symbols;
diff --git a/libsycl/src/detail/device_binary_structures.hpp b/libsycl/src/detail/device_binary_structures.hpp
index 93aaafa8565f1c..e9a9d997e1a225 100644
--- a/libsycl/src/detail/device_binary_structures.hpp
+++ b/libsycl/src/detail/device_binary_structures.hpp
@@ -17,6 +17,8 @@
#include <sycl/__impl/detail/config.hpp>
+#include <string_view>
+
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace detail {
@@ -26,11 +28,11 @@ namespace detail {
/// A device type represented by a particular target
/// triple requires specific binary images. We need
/// to map the image type onto the device target triple.
-/// clang-sycl-linker packs the short form of the triple, i.e. the
-/// architecture only, so these strings carry no vendor/OS components.
+/// Offload images record the triple in its normalized three-part form, of which
+/// the architecture component alone identifies the binary format.
/// SPIR-V with 64-bit pointers.
-static constexpr char DeviceBinaryTripleSPIRV64[] = "spirv64";
+static constexpr std::string_view DeviceBinaryArchSPIRV64 = "spirv64";
} // namespace detail
_LIBSYCL_END_NAMESPACE_SYCL
diff --git a/libsycl/src/detail/device_image_wrapper.hpp b/libsycl/src/detail/device_image_wrapper.hpp
index d66e2b1e623e95..da4023d795ee2d 100644
--- a/libsycl/src/detail/device_image_wrapper.hpp
+++ b/libsycl/src/detail/device_image_wrapper.hpp
@@ -92,7 +92,7 @@ class ProgramWrapper {
class DeviceImageManager {
public:
explicit DeviceImageManager(std::unique_ptr<llvm::object::OffloadBinary> Bin)
- : MBin(std::move(Bin)) {}
+ : MBin(std::move(Bin)), MTargetArch(MBin->getTriple().split('-').first) {}
// Explicitly delete copy constructor/operator= to avoid unintentional copies.
DeviceImageManager(const DeviceImageManager &) = delete;
DeviceImageManager &operator=(const DeviceImageManager &) = delete;
@@ -105,8 +105,15 @@ class DeviceImageManager {
/// \return a reference to the corresponding parsed OffloadBinary object.
const llvm::object::OffloadBinary &getOffloadBinary() const { return *MBin; }
+ /// \return the architecture component of this image's target triple.
+ std::string_view getTargetArch() const { return MTargetArch; }
+
private:
std::unique_ptr<llvm::object::OffloadBinary> MBin;
+
+ // Points into the image's triple string, so it stays valid exactly as long as
+ // the fat binary this image was parsed from is registered.
+ std::string_view MTargetArch;
};
} // namespace detail
diff --git a/libsycl/src/detail/program_manager.cpp b/libsycl/src/detail/program_manager.cpp
index e3b646c053ab48..aeec65aebed8f1 100644
--- a/libsycl/src/detail/program_manager.cpp
+++ b/libsycl/src/detail/program_manager.cpp
@@ -162,13 +162,12 @@ void ProgramAndKernelManager::unregisterFatBin(const void *BinaryStart,
static bool isImageCompatible(const DeviceImageManager &Image,
const DeviceImpl &Device) {
- const llvm::object::OffloadBinary &OB = Image.getOffloadBinary();
- if (!(OB.getTriple() == DeviceBinaryTripleSPIRV64 &&
+ if (!(Image.getTargetArch() == DeviceBinaryArchSPIRV64 &&
Device.getBackend() == sycl::backend::level_zero))
return false;
bool IsValid{};
- llvm::StringRef ImageBytes = OB.getImage();
+ llvm::StringRef ImageBytes = Image.getOffloadBinary().getImage();
callAndThrow(olIsValidBinary, Device.getOLHandle(), ImageBytes.data(),
ImageBytes.size(), &IsValid);
return IsValid;
diff --git a/libsycl/unittests/common/device_images.hpp b/libsycl/unittests/common/device_images.hpp
index 7402214d8f5191..d5fdd2f2b54ffc 100644
--- a/libsycl/unittests/common/device_images.hpp
+++ b/libsycl/unittests/common/device_images.hpp
@@ -23,14 +23,18 @@
_LIBSYCL_BEGIN_NAMESPACE_SYCL
namespace unittests {
-inline llvm::object::OffloadBinary::OffloadingImage createSYCLImage(
- llvm::StringRef SymbolsBlob,
- llvm::object::ImageKind ImageKind = llvm::object::IMG_SPIRV,
- llvm::object::OffloadKind OffloadKind = llvm::object::OFK_SYCL) {
+/// The triple clang-sycl-linker records in SYCL offload images.
+inline constexpr llvm::StringRef DefaultImageTriple = "spirv64-unknown-unknown";
+
+inline llvm::object::OffloadBinary::OffloadingImage
+createSYCLImage(llvm::StringRef SymbolsBlob,
+ llvm::object::ImageKind ImageKind = llvm::object::IMG_SPIRV,
+ llvm::object::OffloadKind OffloadKind = llvm::object::OFK_SYCL,
+ llvm::StringRef Triple = DefaultImageTriple) {
llvm::object::OffloadBinary::OffloadingImage Image;
Image.TheImageKind = ImageKind;
Image.TheOffloadKind = OffloadKind;
- Image.StringData["triple"] = sycl::detail::DeviceBinaryTripleSPIRV64;
+ Image.StringData["triple"] = Triple;
Image.StringData["symbols"] = SymbolsBlob;
static constexpr char DummyImageData[] = "dummy image data";
Image.Image = llvm::MemoryBuffer::getMemBufferCopy(
@@ -41,12 +45,13 @@ inline llvm::object::OffloadBinary::OffloadingImage createSYCLImage(
inline llvm::SmallString<0> createSYCLDeviceBinary(
llvm::ArrayRef<llvm::StringRef> KernelNames,
llvm::object::ImageKind ImageKind = llvm::object::IMG_SPIRV,
- llvm::object::OffloadKind OffloadKind = llvm::object::OFK_SYCL) {
+ llvm::object::OffloadKind OffloadKind = llvm::object::OFK_SYCL,
+ llvm::StringRef Triple = DefaultImageTriple) {
llvm::SmallString<0> SymbolsBlob;
llvm::offloading::sycl::writeSymbolTable(KernelNames, SymbolsBlob);
llvm::object::OffloadBinary::OffloadingImage Image =
- createSYCLImage(SymbolsBlob, ImageKind, OffloadKind);
+ createSYCLImage(SymbolsBlob, ImageKind, OffloadKind, Triple);
return llvm::object::OffloadBinary::write(Image);
}
diff --git a/libsycl/unittests/common/scoped_binary_registration.hpp b/libsycl/unittests/common/scoped_binary_registration.hpp
index bd97d6a5825037..2ab2b25d5e97ee 100644
--- a/libsycl/unittests/common/scoped_binary_registration.hpp
+++ b/libsycl/unittests/common/scoped_binary_registration.hpp
@@ -27,8 +27,10 @@ namespace unittests {
class ScopedBinaryRegistration {
public:
- explicit ScopedBinaryRegistration(llvm::ArrayRef<llvm::StringRef> KernelNames)
- : MBinary(createSYCLDeviceBinary(KernelNames)) {
+ explicit ScopedBinaryRegistration(llvm::ArrayRef<llvm::StringRef> KernelNames,
+ llvm::StringRef Triple = DefaultImageTriple)
+ : MBinary(createSYCLDeviceBinary(KernelNames, llvm::object::IMG_SPIRV,
+ llvm::object::OFK_SYCL, Triple)) {
sycl::detail::ProgramAndKernelManager::getInstance().registerFatBin(
MBinary.data(), MBinary.size());
}
@@ -44,9 +46,10 @@ class ScopedBinaryRegistration {
class ScopedKernelRegistration {
public:
- explicit ScopedKernelRegistration(std::string KernelName)
+ explicit ScopedKernelRegistration(std::string KernelName,
+ llvm::StringRef Triple = DefaultImageTriple)
: MKernelName(std::move(KernelName)),
- MRegistration(createKernelNames(MKernelName)) {}
+ MRegistration(createKernelNames(MKernelName), Triple) {}
private:
static std::array<llvm::StringRef, 1>
diff --git a/libsycl/unittests/program_manager/CMakeLists.txt b/libsycl/unittests/program_manager/CMakeLists.txt
index 3604b86bf6ed2a..8aaefb21e00174 100644
--- a/libsycl/unittests/program_manager/CMakeLists.txt
+++ b/libsycl/unittests/program_manager/CMakeLists.txt
@@ -1,4 +1,5 @@
add_sycl_unittest(ProgramManagerTests
register_and_unregister.cpp
program_cache.cpp
+ image_compatibility.cpp
)
diff --git a/libsycl/unittests/program_manager/image_compatibility.cpp b/libsycl/unittests/program_manager/image_compatibility.cpp
new file mode 100644
index 00000000000000..0b25f9d60dd018
--- /dev/null
+++ b/libsycl/unittests/program_manager/image_compatibility.cpp
@@ -0,0 +1,100 @@
+//===----------------------------------------------------------------------===//
+//
+// Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
+// See https://llvm.org/LICENSE.txt for license information.
+// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
+//
+//===----------------------------------------------------------------------===//
+///
+/// \file
+/// Tests for matching a registered device image against a device: only the
+/// architecture component of the image's target triple takes part in it.
+///
+//===----------------------------------------------------------------------===//
+
+#include <common/scoped_binary_registration.hpp>
+#include <mock/helpers.hpp>
+
+#include <detail/context_impl.hpp>
+#include <detail/device_impl.hpp>
+#include <detail/program_manager.hpp>
+
+#include <sycl/__impl/device.hpp>
+#include <sycl/__impl/exception.hpp>
+#include <sycl/__impl/property_list.hpp>
+
+#include <gmock/gmock.h>
+#include <gtest/gtest.h>
+
+#include <memory>
+#include <string>
+#include <string_view>
+#include <vector>
+
+using namespace sycl;
+using namespace ::testing;
+
+namespace {
+
+/// Creates a context over a single device, bypassing sycl::context, which
+/// exposes a platform's default context only.
+std::shared_ptr<detail::ContextImpl> createContext(const device &Device) {
+ std::vector<detail::DeviceImpl *> Devices = {detail::getSyclObjImpl(Device)};
+ return detail::ContextImpl::create(
+ std::move(Devices), detail::defaultAsyncHandler, property_list{});
+}
+
+/// Allows the liboffload calls that these tests do not assert on, so that
+/// context and program lifetime does not produce uninteresting call warnings.
+void allowContextAndProgramLifetimeCalls(mock::MockLiboffload &Mock) {
+ EXPECT_CALL(Mock, olCreateContext(_, _, _)).Times(AnyNumber());
+ EXPECT_CALL(Mock, olDestroyContext(_)).Times(AnyNumber());
+ EXPECT_CALL(Mock, olCreateProgram(_, _, _, _, _)).Times(AnyNumber());
+ EXPECT_CALL(Mock, olDestroyProgram(_)).Times(AnyNumber());
+ EXPECT_CALL(Mock, olGetSymbol(_, _, _, _)).Times(AnyNumber());
+}
+
+ol_symbol_handle_t
+getKernel(const std::shared_ptr<detail::ContextImpl> &Context,
+ const device &Device, std::string_view KernelName) {
+ detail::ProgramAndKernelManager &Manager =
+ detail::ProgramAndKernelManager::getInstance();
+ return Manager.getOrCreateKernel(Manager.getDeviceKernelInfo(KernelName),
+ Context, *detail::getSyclObjImpl(Device));
+}
+
+} // namespace
+
+// An image whose triple carries an extra environment component is still usable.
+TEST(ImageCompatibility, ArchitectureAloneSelectsImage) {
+ mock::MockWrapper Mock;
+ allowContextAndProgramLifetimeCalls(Mock.get());
+
+ const std::string KernelName = "kernel";
+ sycl::unittests::ScopedKernelRegistration Registration(
+ KernelName, "spirv64-unknown-unknown-elf");
+
+ const device Device;
+ std::shared_ptr<detail::ContextImpl> Context = createContext(Device);
+
+ EXPECT_NE(getKernel(Context, Device, KernelName), nullptr);
+}
+
+// An image built for another architecture must not be handed to the device.
+TEST(ImageCompatibility, ImageForAnotherArchitectureIsRejected) {
+ mock::MockWrapper Mock;
+ allowContextAndProgramLifetimeCalls(Mock.get());
+
+ const std::string KernelName = "kernel";
+ sycl::unittests::ScopedKernelRegistration Registration(
+ KernelName, "spirv32-unknown-unknown");
+
+ const device Device;
+ std::shared_ptr<detail::ContextImpl> Context = createContext(Device);
+
+ EXPECT_THAT([&]() { getKernel(Context, Device, KernelName); },
+ Throws<sycl::exception>(AllOf(
+ Property(&sycl::exception::what,
+ HasSubstr("No compatible image for " + KernelName)),
+ Property(&sycl::exception::code, Eq(sycl::errc::runtime)))));
+}
More information about the cfe-commits
mailing list