From fd8e6e418500e2092d97fa44440ea511c2f07129 Mon Sep 17 00:00:00 2001 From: Krzysztof Weronski Date: Mon, 13 Apr 2026 15:48:42 +0200 Subject: [PATCH] [SYCL] Improve error message when no compatible device image is found ProgramManager::getDeviceImage previously threw the same "No kernel named X was found" error for two distinct situations: - the kernel name is genuinely unknown (not registered at all), and - the kernel is known but no binary image in the fat binary targets the device being used at runtime. The second case is typically caused by compiling for a different device architecture than the one executing the binary, which is a common user mistake. Conflating it with an unknown kernel name made the error misleading and hard to diagnose. Introduce a KernelFound flag to distinguish the two paths and emit a dedicated message for the no-compatible-image case. --- sycl/source/detail/device_binary_image.cpp | 21 +++++++ sycl/source/detail/device_binary_image.hpp | 2 + .../program_manager/program_manager.cpp | 59 +++++++++++++++---- .../program_manager/program_manager.hpp | 4 ++ .../program_manager/CompileTarget.cpp | 14 +++-- 5 files changed, 82 insertions(+), 18 deletions(-) diff --git a/sycl/source/detail/device_binary_image.cpp b/sycl/source/detail/device_binary_image.cpp index a449931967fae..292e1bacfa600 100644 --- a/sycl/source/detail/device_binary_image.cpp +++ b/sycl/source/detail/device_binary_image.cpp @@ -90,6 +90,27 @@ const char *DeviceBinaryProperty::asCString() const { return ur::cast(Prop->ValAddr) + Shift; } +std::string_view DeviceBinaryProperty::asStringView() const { + const char *Str = asCString(); + // ValSize covers the entire blob stored at ValAddr. The two property types + // that can carry string data have different layouts: + // - BYTE_ARRAY: used by PropertyValue (property_set_io.hpp) when serialising + // any byte sequence, including strings. The blob starts with a mandatory + // 8-byte little-endian uint64_t encoding the payload bit-count, followed + // by the actual bytes. asCString() already skips that 8-byte header, so + // we subtract 8 from ValSize to get the true payload length. + // - STRING: a plain null-terminated C string written directly to ValAddr, + // with ValSize counting the bytes including the terminator. asCString() + // returns the start of the string directly, so we subtract 1 to exclude + // the terminator from the view's length. + assert((Prop->Type == SYCL_PROPERTY_TYPE_STRING || + Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY) && + "property type mismatch"); + size_t Len = Prop->Type == SYCL_PROPERTY_TYPE_BYTE_ARRAY ? Prop->ValSize - 8 + : Prop->ValSize - 1; + return {Str, Len}; +} + void RTDeviceBinaryImage::PropertyRange::init(sycl_device_binary Bin, const char *PropSetName) { assert(!this->Begin && !this->End && "already initialized"); diff --git a/sycl/source/detail/device_binary_image.hpp b/sycl/source/detail/device_binary_image.hpp index ed4baeefd6a65..6aa84e6323f53 100644 --- a/sycl/source/detail/device_binary_image.hpp +++ b/sycl/source/detail/device_binary_image.hpp @@ -20,6 +20,7 @@ #include #include #include +#include namespace sycl { inline namespace _V1 { @@ -75,6 +76,7 @@ class DeviceBinaryProperty { uint32_t asUint32() const; ByteArray asByteArray() const; const char *asCString() const; + std::string_view asStringView() const; protected: friend std::ostream &operator<<(std::ostream &Out, diff --git a/sycl/source/detail/program_manager/program_manager.cpp b/sycl/source/detail/program_manager/program_manager.cpp index e700905812343..275fcb2e1082c 100644 --- a/sycl/source/detail/program_manager/program_manager.cpp +++ b/sycl/source/detail/program_manager/program_manager.cpp @@ -1254,6 +1254,22 @@ const char *getArchName(const device_impl &DeviceImpl) { return "unknown"; } +// Get a human-readable target label for a device binary image, usable in +// error messages. +// - For AOT-compiled images with a compile_target property returns the +// architecture name (e.g. "intel_gpu_bdw") +// - For JIT / generic images returns the raw DeviceTargetSpec string +static std::string_view getImageTargetLabel(const RTDeviceBinaryImage &Img) { + auto PropRange = Img.getDeviceRequirements(); + auto PropIt = + std::find_if(PropRange.begin(), PropRange.end(), [](const auto &Prop) { + return Prop->Name == std::string_view("compile_target"); + }); + if (PropIt != PropRange.end()) + return DeviceBinaryProperty(*PropIt).asStringView(); + return Img.getRawData().DeviceTargetSpec; +} + template const RTDeviceBinaryImage *getBinImageFromMultiMap( const std::unordered_multimap @@ -1308,6 +1324,19 @@ const RTDeviceBinaryImage *getBinImageFromMultiMap( return DeviceFilteredImgs[ImgInd]; } +std::string ProgramManager::getKernelTargetList(const kernel_id &KernelID) { + std::lock_guard Guard(m_ImgMapsMutex); + auto [ItBegin, ItEnd] = m_KernelIDs2BinImage.equal_range(KernelID); + assert(ItBegin != ItEnd && "Expected at least one image"); + + std::string TargetList{getImageTargetLabel(*ItBegin->second)}; + for (auto It = std::next(ItBegin); It != ItEnd; ++It) { + TargetList += ", "; + TargetList += getImageTargetLabel(*It->second); + } + return TargetList; +} + const RTDeviceBinaryImage & ProgramManager::getDeviceImage(std::string_view KernelName, context_impl &ContextImpl, @@ -1328,23 +1357,34 @@ ProgramManager::getDeviceImage(std::string_view KernelName, } const RTDeviceBinaryImage *Img = nullptr; + std::optional FoundKernelID; { std::lock_guard Guard(m_DeviceKernelInfoMapMutex); if (auto It = m_DeviceKernelInfoMap.find(KernelName); It != m_DeviceKernelInfoMap.end()) { - Img = getBinImageFromMultiMap(m_KernelIDs2BinImage, - It->second.getKernelID(), ContextImpl, - DeviceImpl); + FoundKernelID = It->second.getKernelID(); + Img = getBinImageFromMultiMap(m_KernelIDs2BinImage, *FoundKernelID, + ContextImpl, DeviceImpl); } } // Decompress the image if it is compressed. CheckAndDecompressImage(Img); - if (!Img) + if (!Img) { + if (!FoundKernelID) + throw exception(make_error_code(errc::runtime), + "No kernel named " + std::string(KernelName) + + " was found"); + // The kernel is registered but none of its images target the selected + // device. Enumerate the available targets so the user can see what the + // binary supports. throw exception(make_error_code(errc::runtime), - "No kernel named " + std::string(KernelName) + - " was found"); + "Kernel " + std::string(KernelName) + + " has no image for the selected device. " + "Its available images target: [" + + getKernelTargetList(*FoundKernelID) + "]."); + } if constexpr (DbgProgMgr > 0) { std::cerr << "selected device image: " << &Img->getRawData() << "\n"; @@ -3476,12 +3516,7 @@ bool doesImageTargetMatchDevice(const RTDeviceBinaryImage &Img, // Device image has the compile_target property, so it is AOT compiled for // some device, check if that architecture is Device's architecture. - auto CompileTargetByteArray = DeviceBinaryProperty(*PropIt).asByteArray(); - // Drop 8 bytes describing the size of the byte array. - CompileTargetByteArray.dropBytes(8); - std::string_view CompileTarget( - reinterpret_cast(&CompileTargetByteArray[0]), - CompileTargetByteArray.size()); + std::string_view CompileTarget = DeviceBinaryProperty(*PropIt).asStringView(); std::string_view ArchName = getArchName(DevImpl); // Note: there are no explicit targets for CPUs, so on x86_64, // intel_cpu_spr, and intel_cpu_gnr, we use a spir64_x86_64 diff --git a/sycl/source/detail/program_manager/program_manager.hpp b/sycl/source/detail/program_manager/program_manager.hpp index 3264a1e1b341b..1401e7d360115 100644 --- a/sycl/source/detail/program_manager/program_manager.hpp +++ b/sycl/source/detail/program_manager/program_manager.hpp @@ -409,6 +409,10 @@ class ProgramManager { bool shouldBF16DeviceImageBeUsed(const RTDeviceBinaryImage *BinImage, const device_impl &DeviceImpl); + /// Returns a comma-separated list of available image target names for the + /// given kernel ID, for use in error messages. + std::string getKernelTargetList(const kernel_id &KernelID); + protected: using RTDeviceBinaryImageUPtr = std::unique_ptr; using DynRTDeviceBinaryImageUPtr = std::unique_ptr; diff --git a/sycl/unittests/program_manager/CompileTarget.cpp b/sycl/unittests/program_manager/CompileTarget.cpp index 9e5e1f8f8e369..98989600fd44f 100644 --- a/sycl/unittests/program_manager/CompileTarget.cpp +++ b/sycl/unittests/program_manager/CompileTarget.cpp @@ -54,14 +54,14 @@ generateImageWithCompileTarget(std::string KernelName, class SingleTaskKernel; class NDRangeKernel; class RangeKernel; -class NoDeviceKernel; +class NoKernelDevice; class JITFallbackKernel; class SKLOnlyKernel; MOCK_INTEGRATION_HEADER(SingleTaskKernel) MOCK_INTEGRATION_HEADER(NDRangeKernel) MOCK_INTEGRATION_HEADER(RangeKernel) -MOCK_INTEGRATION_HEADER(NoDeviceKernel) +MOCK_INTEGRATION_HEADER(NoKernelDevice) MOCK_INTEGRATION_HEADER(JITFallbackKernel) MOCK_INTEGRATION_HEADER(SKLOnlyKernel) @@ -87,7 +87,7 @@ static sycl::unittest::MockDeviceImage Img[] = { "intel_gpu_pvc"), sycl::unittest::generateImageWithCompileTarget("RangeKernel", "intel_gpu_skl"), - sycl::unittest::generateImageWithCompileTarget("NoDeviceKernel", + sycl::unittest::generateImageWithCompileTarget("NoKernelDevice", "intel_gpu_bdw"), sycl::unittest::generateDefaultImage({"JITFallbackKernel"}), sycl::unittest::generateImageWithCompileTarget("JITFallbackKernel", @@ -319,12 +319,14 @@ TEST_F(CompileTargetTest, RangeKernel) { }); } -TEST_F(CompileTargetTest, NoDeviceKernel) { +TEST_F(CompileTargetTest, NoKernelDevice) { try { - queue{}.single_task([]() {}); + queue{}.single_task([]() {}); } catch (sycl::exception &e) { ASSERT_EQ(e.what(), - std::string("No kernel named NoDeviceKernel was found")); + std::string("Kernel NoKernelDevice has no image for the " + "selected device. Its available images target: " + "[intel_gpu_bdw].")); } }