[SYCL] Throw an exception for unsupported aspect

The runtime must throw an exception whenever the application attempts to submit
a kernel to a device where the kernel uses a feature that is not compatible
with the device.

Author: KornevNikita

sycl/include/sycl/aspects.hpp

@@ -12,47 +12,13 @@
 namespace sycl {
 __SYCL_INLINE_VER_NAMESPACE(_V1) {
 
+#define __SYCL_ASPECT(A, B, C) A = B,
+#define __SYCL_SPECIAL_ASPECT(A, B, C) A = B,
 enum class __SYCL_TYPE(aspect) aspect {
-  host = 0,
-  cpu = 1,
-  gpu = 2,
-  accelerator = 3,
-  custom = 4,
-  fp16 = 5,
-  fp64 = 6,
-  int64_base_atomics __SYCL2020_DEPRECATED("use atomic64 instead") = 7,
-  int64_extended_atomics __SYCL2020_DEPRECATED("use atomic64 instead") = 8,
-  image = 9,
-  online_compiler = 10,
-  online_linker = 11,
-  queue_profiling = 12,
-  usm_device_allocations = 13,
-  usm_host_allocations = 14,
-  usm_shared_allocations = 15,
-  usm_restricted_shared_allocations = 16,
-  usm_system_allocations = 17,
-  usm_system_allocator __SYCL2020_DEPRECATED(
-      "use usm_system_allocations instead") = usm_system_allocations,
-  ext_intel_pci_address = 18,
-  ext_intel_gpu_eu_count = 19,
-  ext_intel_gpu_eu_simd_width = 20,
-  ext_intel_gpu_slices = 21,
-  ext_intel_gpu_subslices_per_slice = 22,
-  ext_intel_gpu_eu_count_per_subslice = 23,
-  ext_intel_max_mem_bandwidth = 24,
-  ext_intel_mem_channel = 25,
-  usm_atomic_host_allocations = 26,
-  usm_atomic_shared_allocations = 27,
-  atomic64 = 28,
-  ext_intel_device_info_uuid = 29,
-  ext_oneapi_srgb = 30,
-  ext_oneapi_native_assert = 31,
-  host_debuggable = 32,
-  ext_intel_gpu_hw_threads_per_eu = 33,
-  ext_oneapi_cuda_async_barrier = 34,
-  ext_oneapi_bfloat16 = 35,
-  ext_intel_free_memory = 36,
+#include <sycl/info/aspects.def>
 };
+#undef __SYCL_SPECIAL_ASPECT
+#undef __SYCL_ASPECT
 
 } // __SYCL_INLINE_VER_NAMESPACE(_V1)
 } // namespace sycl

sycl/include/sycl/detail/pi.h

@@ -752,6 +752,9 @@ static const uint8_t PI_DEVICE_BINARY_OFFLOAD_KIND_SYCL = 4;
 #define __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS "SYCL/exported symbols"
 /// PropertySetRegistry::SYCL_DEVICE_GLOBALS defined in PropertySetIO.h
 #define __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS "SYCL/device globals"
+/// PropertySetRegistry::SYCL_DEVICE_REQUIREMENTS defined in PropertySetIO.h
+#define __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS                        \
+  "SYCL/device requirements"
 
 /// Program metadata tags recognized by the PI backends. For kernels the tag
 /// must appear after the kernel name.

sycl/include/sycl/info/aspects.def

@@ -0,0 +1,39 @@
+  __SYCL_ASPECT(host, 0, "host")
+  __SYCL_ASPECT(cpu, 1, "cpu")
+  __SYCL_ASPECT(gpu, 2, "gpu")
+  __SYCL_ASPECT(accelerator, 3, "accelerator")
+  __SYCL_ASPECT(custom, 4, "custom")
+  __SYCL_ASPECT(fp16, 5, "fp16")
+  __SYCL_ASPECT(fp64, 6, "fp64")
+  __SYCL_ASPECT(int64_base_atomics __SYCL2020_DEPRECATED("use atomic64 instead"), 7, "int64_base_atomics")
+  __SYCL_ASPECT(int64_extended_atomics __SYCL2020_DEPRECATED("use atomic64 instead"), 8, "int64_extended_atomics")
+  __SYCL_ASPECT(image, 9, "image")
+  __SYCL_ASPECT(online_compiler, 10, "online_compiler")
+  __SYCL_ASPECT(online_linker, 11, "online_linker")
+  __SYCL_ASPECT(queue_profiling, 12, "queue_profiling")
+  __SYCL_ASPECT(usm_device_allocations, 13, "usm_device_allocations")
+  __SYCL_ASPECT(usm_host_allocations, 14, "usm_host_allocations")
+  __SYCL_ASPECT(usm_shared_allocations, 15, "usm_shared_allocations")
+  __SYCL_ASPECT(usm_restricted_shared_allocations, 16, "usm_restricted_shared_allocations")
+  __SYCL_ASPECT(usm_system_allocations, 17, "usm_system_allocations")
+  __SYCL_SPECIAL_ASPECT(usm_system_allocator __SYCL2020_DEPRECATED("use usm_system_allocations instead"), usm_system_allocations, "usm_system_allocations")
+  __SYCL_ASPECT(ext_intel_pci_address, 18, "ext_intel_pci_address")
+  __SYCL_ASPECT(ext_intel_gpu_eu_count, 19, "ext_intel_gpu_eu_count")
+  __SYCL_ASPECT(ext_intel_gpu_eu_simd_width, 20, "ext_intel_gpu_eu_simd_width")
+  __SYCL_ASPECT(ext_intel_gpu_slices, 21, "ext_intel_gpu_slices")
+  __SYCL_ASPECT(ext_intel_gpu_subslices_per_slice, 22, "ext_intel_gpu_subslices_per_slice")
+  __SYCL_ASPECT(ext_intel_gpu_eu_count_per_subslice, 23, "ext_intel_gpu_eu_count_per_subslice")
+  __SYCL_ASPECT(ext_intel_max_mem_bandwidth, 24, "ext_intel_max_mem_bandwidth")
+  __SYCL_ASPECT(ext_intel_mem_channel, 25, "ext_intel_mem_channel")
+  __SYCL_ASPECT(usm_atomic_host_allocations, 26, "usm_atomic_host_allocations")
+  __SYCL_ASPECT(usm_atomic_shared_allocations, 27, "usm_atomic_shared_allocations")
+  __SYCL_ASPECT(atomic64, 28, "atomic64")
+  __SYCL_ASPECT(ext_intel_device_info_uuid, 29, "ext_intel_device_info_uuid")
+  __SYCL_ASPECT(ext_oneapi_srgb, 30, "ext_oneapi_srgb")
+  __SYCL_ASPECT(ext_oneapi_native_assert, 31, "ext_oneapi_native_assert")
+  __SYCL_ASPECT(host_debuggable, 32, "host_debuggable")
+  __SYCL_ASPECT(ext_intel_gpu_hw_threads_per_eu, 33, "ext_intel_gpu_hw_threads_per_eu")
+  __SYCL_ASPECT(ext_oneapi_cuda_async_barrier, 34, "ext_oneapi_cuda_async_barrier")
+  __SYCL_ASPECT(ext_oneapi_bfloat16, 35, "ext_oneapi_bfloat16")
+  __SYCL_ASPECT(ext_intel_free_memory, 36, "ext_intel_free_memory")
+  

sycl/source/detail/device_binary_image.cpp

@@ -175,6 +175,7 @@ void RTDeviceBinaryImage::init(pi_device_binary Bin) {
   ProgramMetadata.init(Bin, __SYCL_PI_PROPERTY_SET_PROGRAM_METADATA);
   ExportedSymbols.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_EXPORTED_SYMBOLS);
   DeviceGlobals.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_GLOBALS);
+  DeviceRequirements.init(Bin, __SYCL_PI_PROPERTY_SET_SYCL_DEVICE_REQUIREMENTS);
 }
 
 DynRTDeviceBinaryImage::DynRTDeviceBinaryImage(

sycl/source/detail/device_binary_image.hpp

@@ -190,6 +190,9 @@ class RTDeviceBinaryImage {
   const PropertyRange &getProgramMetadata() const { return ProgramMetadata; }
   const PropertyRange &getExportedSymbols() const { return ExportedSymbols; }
   const PropertyRange &getDeviceGlobals() const { return DeviceGlobals; }
+  const PropertyRange &getDeviceRequirements() const {
+    return DeviceRequirements;
+  }
 
 protected:
   void init(pi_device_binary Bin);
@@ -207,6 +210,7 @@ class RTDeviceBinaryImage {
   RTDeviceBinaryImage::PropertyRange ProgramMetadata;
   RTDeviceBinaryImage::PropertyRange ExportedSymbols;
   RTDeviceBinaryImage::PropertyRange DeviceGlobals;
+  RTDeviceBinaryImage::PropertyRange DeviceRequirements;
 };
 
 // Dynamically allocated device binary image, which de-allocates its binary

sycl/source/detail/program_manager/program_manager.cpp

@@ -15,6 +15,7 @@
 #include <detail/program_impl.hpp>
 #include <detail/program_manager/program_manager.hpp>
 #include <detail/spec_constant_impl.hpp>
+#include <sycl/aspects.hpp>
 #include <sycl/backend_types.hpp>
 #include <sycl/context.hpp>
 #include <sycl/detail/common.hpp>
@@ -543,14 +544,52 @@ RT::PiProgram ProgramManager::getBuiltPIProgram(
 
   DeviceImplPtr Dev =
       (MustBuildOnSubdevice == PI_TRUE) ? DeviceImpl : RootDevImpl;
-  auto BuildF = [this, &M, &KSId, &ContextImpl, &Dev, Prg, &CompileOpts,
-                 &LinkOpts, &JITCompilationIsRequired, SpecConsts] {
-    auto Context = createSyclObjFromImpl<context>(ContextImpl);
-    auto Device = createSyclObjFromImpl<device>(Dev);
+  auto Context = createSyclObjFromImpl<context>(ContextImpl);
+  auto Device = createSyclObjFromImpl<device>(Dev);
+  const RTDeviceBinaryImage &Img =
+      getDeviceImage(M, KSId, Context, Device, JITCompilationIsRequired);
+
+  // Check that device supports all aspects used by the kernel
+  const RTDeviceBinaryImage::PropertyRange &ARange =
+      Img.getDeviceRequirements();
+
+#define __SYCL_ASPECT(A, B, C)                                                 \
+  case aspect::A:                                                              \
+    return C;
+#define __SYCL_SPECIAL_ASPECT(A, B, C)
+  auto getAspectNameStr = [](aspect AspectNum) -> std::string {
+    switch (AspectNum) {
+#include <sycl/info/aspects.def>
+    default:
+      throw sycl::exception(
+          errc::kernel_not_supported,
+          "Unknown aspect " + std::to_string(static_cast<unsigned>(AspectNum)));
+    }
+  };
+#undef __SYCL_SPECIAL_ASPECT
+#undef __SYCL_ASPECT
 
-    const RTDeviceBinaryImage &Img =
-        getDeviceImage(M, KSId, Context, Device, JITCompilationIsRequired);
+  for (RTDeviceBinaryImage::PropertyRange::ConstIterator It : ARange) {
+    auto KName = std::string((*It)->Name);
+    if (KName != "aspects")
+      continue;
+    ByteArray Aspects = DeviceBinaryProperty(*It).asByteArray();
+    // 8 because we need to skip 64-bits of size of the byte array
+    auto *AIt = reinterpret_cast<const std::uint32_t *>(&Aspects[8]);
+    auto *AEnd =
+        reinterpret_cast<const std::uint32_t *>(&Aspects[0] + Aspects.size());
+    while (AIt != AEnd) {
+      auto Aspect = static_cast<aspect>(*AIt);
+      if (!Dev->has(Aspect))
+        throw sycl::exception(errc::kernel_not_supported,
+                              "Required aspect " + getAspectNameStr(Aspect) +
+                                  " is not supported on the device");
+      ++AIt;
+    }
+  }
 
+  auto BuildF = [this, &Img, &Context, &ContextImpl, &Device, Prg, &CompileOpts,
+                 &LinkOpts, SpecConsts, &KernelName] {
     applyOptionsFromImage(CompileOpts, LinkOpts, Img);
 
     const detail::plugin &Plugin = ContextImpl->getPlugin();