[SYCL] Improve error message when no compatible device image is found

sycl/source/detail/device_binary_image.cpp

@@ -90,6 +90,27 @@ const char *DeviceBinaryProperty::asCString() const {
   return ur::cast<const char *>(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");

sycl/source/detail/device_binary_image.hpp

@@ -20,6 +20,7 @@
 #include <cstring>
 #include <memory>
 #include <mutex>
+#include <string_view>
 
 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,

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 <typename StorageKey>
 const RTDeviceBinaryImage *getBinImageFromMultiMap(
     const std::unordered_multimap<StorageKey, const RTDeviceBinaryImage *>
@@ -1308,6 +1324,19 @@ const RTDeviceBinaryImage *getBinImageFromMultiMap(
   return DeviceFilteredImgs[ImgInd];
 }
 
+std::string ProgramManager::getKernelTargetList(const kernel_id &KernelID) {
+  std::lock_guard<std::mutex> 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<kernel_id> FoundKernelID;
   {
     std::lock_guard<std::mutex> 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";
@@ -3474,12 +3514,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<const char *>(&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

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<RTDeviceBinaryImage>;
   using DynRTDeviceBinaryImageUPtr = std::unique_ptr<DynRTDeviceBinaryImage>;

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<NoDeviceKernel>([]() {});
+    queue{}.single_task<NoKernelDevice>([]() {});
   } 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]."));
   }
 }