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" <[email protected]> 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 <[email protected]> --- .../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" <[email protected]> Date: Thu, 15 Oct 2026 03:23:31 -0700 Subject: [PATCH 2/2] redesign Signed-off-by: Tikhomirova, Kseniya <[email protected]> --- .../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))))); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
