[SYCL][Bindless] Refactor existing bindless aspects for missing L0 bindless functionality

llvm/include/llvm/SYCLLowerIR/DeviceConfigFile.td

@@ -54,20 +54,21 @@ def AspectExt_oneapi_bindless_images_shared_usm : Aspect<"ext_oneapi_bindless_im
 def AspectExt_oneapi_bindless_images_1d_usm : Aspect<"ext_oneapi_bindless_images_1d_usm">;
 def AspectExt_oneapi_bindless_images_2d_usm : Aspect<"ext_oneapi_bindless_images_2d_usm">;
 def AspectExt_oneapi_interop_memory_import : Aspect<"ext_oneapi_interop_memory_import">;
-def AspectExt_oneapi_interop_memory_export : Aspect<"ext_oneapi_interop_memory_export">;
 def AspectExt_oneapi_interop_semaphore_import : Aspect<"ext_oneapi_interop_semaphore_import">;
-def AspectExt_oneapi_interop_semaphore_export : Aspect<"ext_oneapi_interop_semaphore_export">;
 def AspectExt_oneapi_mipmap : Aspect<"ext_oneapi_mipmap">;
 def AspectExt_oneapi_mipmap_anisotropy : Aspect<"ext_oneapi_mipmap_anisotropy">;
 def AspectExt_oneapi_mipmap_level_reference : Aspect<"ext_oneapi_mipmap_level_reference">;
 def AspectExt_oneapi_bindless_sampled_image_fetch_1d_usm : Aspect<"ext_oneapi_bindless_sampled_image_fetch_1d_usm">;
 def AspectExt_oneapi_bindless_sampled_image_fetch_1d : Aspect<"ext_oneapi_bindless_sampled_image_fetch_1d">;
 def AspectExt_oneapi_bindless_sampled_image_fetch_2d_usm : Aspect<"ext_oneapi_bindless_sampled_image_fetch_2d_usm">;
 def AspectExt_oneapi_bindless_sampled_image_fetch_2d : Aspect<"ext_oneapi_bindless_sampled_image_fetch_2d">;
-def AspectExt_oneapi_bindless_sampled_image_fetch_3d_usm : Aspect<"ext_oneapi_bindless_sampled_image_fetch_3d_usm">;
 def AspectExt_oneapi_bindless_sampled_image_fetch_3d : Aspect<"ext_oneapi_bindless_sampled_image_fetch_3d">;
 def AspectExt_oneapi_cubemap : Aspect<"ext_oneapi_cubemap">;
 def AspectExt_oneapi_cubemap_seamless_filtering : Aspect<"ext_oneapi_cubemap_seamless_filtering">;
+def AspectExt_oneapi_image_array : Aspect<"ext_oneapi_image_array">;
+def AspectExt_oneapi_unique_addressing_per_dim : Aspect<"ext_oneapi_unique_addressing_per_dim">;
+def AspectExt_oneapi_bindless_images_sample_1d_usm : Aspect<"ext_oneapi_bindless_images_sample_1d_usm">;
+def AspectExt_oneapi_bindless_images_sample_2d_usm : Aspect<"ext_oneapi_bindless_images_sample_2d_usm">;
 def AspectExt_intel_esimd : Aspect<"ext_intel_esimd">;
 def AspectExt_oneapi_ballot_group : Aspect<"ext_oneapi_ballot_group">;
 def AspectExt_oneapi_fixed_size_group : Aspect<"ext_oneapi_fixed_size_group">;
@@ -128,13 +129,16 @@ def : TargetInfo<"__TestAspectList",
     AspectExt_intel_device_id, AspectExt_intel_memory_clock_rate, AspectExt_intel_memory_bus_width, AspectEmulated,
     AspectExt_intel_legacy_image, AspectExt_oneapi_bindless_images,
     AspectExt_oneapi_bindless_images_shared_usm, AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm,
-    AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_memory_export,
-    AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_interop_semaphore_export,
-    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference, AspectExt_oneapi_cubemap, 
+    AspectExt_oneapi_interop_memory_import, AspectExt_oneapi_interop_semaphore_import,
+    AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy, AspectExt_oneapi_mipmap_level_reference,
+    AspectExt_oneapi_bindless_sampled_image_fetch_3d, AspectExt_oneapi_cubemap, 
     AspectExt_oneapi_cubemap_seamless_filtering, 
+    AspectExt_oneapi_image_array,
+    AspectExt_oneapi_unique_addressing_per_dim,
+    AspectExt_oneapi_bindless_images_sample_1d_usm,
+    AspectExt_oneapi_bindless_images_sample_2d_usm,
     AspectExt_oneapi_bindless_sampled_image_fetch_1d_usm, AspectExt_oneapi_bindless_sampled_image_fetch_1d,
     AspectExt_oneapi_bindless_sampled_image_fetch_2d_usm, AspectExt_oneapi_bindless_sampled_image_fetch_2d,
-    AspectExt_oneapi_bindless_sampled_image_fetch_3d_usm, AspectExt_oneapi_bindless_sampled_image_fetch_3d,
     AspectExt_intel_esimd,
     AspectExt_oneapi_ballot_group, AspectExt_oneapi_fixed_size_group, AspectExt_oneapi_opportunistic_group,
     AspectExt_oneapi_tangle_group, AspectExt_intel_matrix, AspectExt_oneapi_is_composite, AspectExt_oneapi_is_component, 
@@ -216,7 +220,9 @@ defvar CudaMinAspects = !listconcat(AllUSMAspects, [AspectGpu, AspectFp64, Aspec
 defvar CudaBindlessImagesAspects = [AspectExt_oneapi_bindless_images, AspectExt_oneapi_bindless_images_shared_usm,
     AspectExt_oneapi_bindless_images_1d_usm, AspectExt_oneapi_bindless_images_2d_usm, AspectExt_oneapi_interop_memory_import,
     AspectExt_oneapi_interop_semaphore_import, AspectExt_oneapi_mipmap, AspectExt_oneapi_mipmap_anisotropy,
-    AspectExt_oneapi_mipmap_level_reference, AspectExt_oneapi_cubemap, AspectExt_oneapi_cubemap_seamless_filtering];
+    AspectExt_oneapi_mipmap_level_reference, AspectExt_oneapi_cubemap, AspectExt_oneapi_cubemap_seamless_filtering,
+    AspectExt_oneapi_image_array, AspectExt_oneapi_unique_addressing_per_dim, AspectExt_oneapi_bindless_images_sample_2d_usm,
+    AspectExt_oneapi_bindless_images_sample_2d_usm];
 
 def : CudaTargetInfo<"nvidia_gpu_sm_50", !listconcat(CudaMinAspects, CudaBindlessImagesAspects)>;
 def : CudaTargetInfo<"nvidia_gpu_sm_52", !listconcat(CudaMinAspects, CudaBindlessImagesAspects)>;

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -135,8 +135,8 @@ The device aspects for these queries are:
 |======================
 |Device descriptor |Description
 |`aspect::ext_oneapi_bindless_images` | Indicates if the device supports 
-creation of bindless images backed by the `image_mem` and `image_mem_handle` 
-APIs.
+bindless images. This includes creating bindless images backed by the
+`image_mem` and `image_mem_handle` APIs.
 |`aspect::ext_oneapi_bindless_images_shared_usm` | Indicates if the device 
 supports the creation of bindless images backed by shared USM memory.
 |`aspect::ext_oneapi_bindless_images_1d_usm` | Indicates if the device supports 
@@ -714,6 +714,8 @@ represents the backend's default addressing mode. On CUDA this is `Wrap`, i.e.
 `addressing[3]` defines the addressing mode per texture dimension. A 
 `bindless_image_sampler` can be constructed with a singular 
 `sycl::addressing_mode`, where this parameter will define all dimensions.
+Not all devices may support unique addressing per dimension.  We provide device
+aspect queries for this in <<querying_unique_addressing_support>>
 
 `mipmap_filtering` dictates the method in which sampling between mipmap 
 levels is performed.
@@ -1069,6 +1071,9 @@ not all devices may support fetching of sampled image data depending on the
 dimension or backing memory type. We provide device aspect queries for this in 
 <<querying_sampled_image_fetch_support>>.
 
+Additionally, not all devices may support sampling of USM images. We provide
+device aspect queries for this in <<querying_usm_sample_support>>
+
 The user is required to pass a `DataT` template parameter, which specifies the
 return type of the `fetch_image` and `sample_image` functions. If `DataT` is 
 not a recognized standard type, as defined in <<recognized_standard_types>>, 
@@ -1171,14 +1176,41 @@ The device aspect descriptors for these queries are:
 |`aspect::ext_oneapi_bindless_sampled_image_fetch_2d` | 
  Indicates if the device is capable of fetching non-USM backed 2D 
  sampled image data.
-|`aspect::ext_oneapi_bindless_sampled_image_fetch_3d_usm` | 
- Indicates if the device is capable of fetching USM backed 3D 
- sampled image data.
 |`aspect::ext_oneapi_bindless_sampled_image_fetch_3d` | 
  Indicates if the device is capable of fetching non-USM backed 3D 
  sampled image data.
 |======================
 
+=== Querying USM sampling support [[querying_usm_sample_support]]
+
+We provide the following device queries to query support for sampling USM
+images.
+
+The device aspect descriptors for these queries are:
+
+[frame="none",options="header"]
+|======================
+|Device descriptor | Description
+|`aspect::ext_oneapi_bindless_images_sample_1d_usm` | Indicates if the device
+supports the sampling of 1D bindless images backed by USM.
+|`aspect::ext_oneapi_bindless_images_sample_2d_usm` | Indicates if the device
+supports the sampling of 2D bindless images backed by USM.
+|======================
+
+=== Querying unique addressing support [[querying_unique_addressing_support]]
+
+We provide the following device queries to query support information for
+unique addressing for each image dimension.
+
+The device aspect descriptor for this query is:
+
+[frame="none",options="header"]
+|======================
+|Device descriptor |Description
+|`aspect::ext_oneapi_unique_addressing_per_dim` | Indicates if the device 
+supports unique addressing per dimension when sampling.
+|======================
+
 == Mipmapped images
 
 So far, we have described how to create and operate on standard bindless images.
@@ -1312,6 +1344,20 @@ every index has the same dimensionality, size, and data type.
 Image arrays may also be referred to as layered images, and the array indices 
 may be referred to layers.
 
+=== Querying image array support
+
+We provide the following device aspect to retrieve support information for
+image arrays.
+
+The device aspect descriptor for this query is:
+
+[frame="none",options="header"]
+|======================
+|Device descriptor |Description
+|`aspect::ext_oneapi_image_array` | Indicates if the device supports
+image arrays.
+|======================
+
 === Allocation of image arrays
 
 Image arrays are allocated in a similar manner to standard images.
@@ -1547,12 +1593,8 @@ The device aspect descriptors for these queries are:
 |Device descriptor |Description
 |`aspect::ext_oneapi_interop_memory_import` | Indicates if the device supports 
 importing external memory resources.
-|`aspect::ext_oneapi_interop_memory_export` | Indicates if the device supports 
-exporting internal memory resources.
 |`aspect::ext_oneapi_interop_semaphore_import`` | Indicates if the device 
 supports importing external semaphore resources.
-|`aspect::ext_oneapi_interop_semaphore_export` | Indicates if the device 
-supports exporting internal event resources.
 |======================
 
 
@@ -2811,4 +2853,10 @@ These features still need to be handled:
                     funcs.
 |5.14|2024-07-17| - Rename `destroy_external_semaphore` to 
                     `release_external_semaphore`.
+|5.15|2024-07-19 - Add missing device queries for image arrays, sampling USM 
+                   images and unique addressing per dimension.
+                 - Remove aspects for semaphore export, memory export and fetch 
+                   3D USM images as they are not supported on any platform.
+                 - Refine the description of `ext_oneapi_bindless_images` aspect
+                   to indicate support for bindless image APIs.
 |======================

sycl/include/sycl/detail/pi.h

@@ -490,9 +490,7 @@ typedef enum {
   PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY = 0x2010A,
   PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_LEVEL_REFERENCE_SUPPORT = 0x2010B,
   PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_MEMORY_IMPORT_SUPPORT = 0x2010C,
-  PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_MEMORY_EXPORT_SUPPORT = 0x2010D,
   PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_SEMAPHORE_IMPORT_SUPPORT = 0x2010E,
-  PI_EXT_ONEAPI_DEVICE_INFO_INTEROP_SEMAPHORE_EXPORT_SUPPORT = 0x2010F,
 
   PI_EXT_ONEAPI_DEVICE_INFO_MATRIX_COMBINATIONS = 0x20110,
 
@@ -513,7 +511,6 @@ typedef enum {
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_1D = 0x20118,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D_USM = 0x20119,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_2D = 0x2011A,
-  PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_3D_USM = 0x2011B,
   PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLED_IMAGE_FETCH_3D = 0x2011C,
 
   // Timestamp enqueue
@@ -527,6 +524,12 @@ typedef enum {
 
   // Return whether cluster launch is supported by device
   PI_EXT_ONEAPI_DEVICE_INFO_CLUSTER_LAUNCH = 0x2021,
+
+  // Bindless image arrays, unique addressing and USM sampling
+  PI_EXT_ONEAPI_DEVICE_INFO_IMAGE_ARRAY_SUPPORT = 0x20122,
+  PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_UNIQUE_ADDRESSING_PER_DIM = 0x20123,
+  PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLE_1D_USM = 0x20124,
+  PI_EXT_ONEAPI_DEVICE_INFO_BINDLESS_SAMPLE_2D_USM = 0x20125,
 } _pi_device_info;
 
 typedef enum {

sycl/include/sycl/device_aspect_macros.hpp

@@ -238,21 +238,11 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_import__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__
-//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47)
-#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_memory_export__ 0
-#endif
-
 #ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__
 //__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48)
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_import__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__
-//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49)
-#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_interop_semaphore_export__ 0
-#endif
-
 #ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__
 //__SYCL_ASPECT(ext_oneapi_mipmap, 50)
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_mipmap__ 0
@@ -355,12 +345,6 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_2d__ 0
 #endif
 
-#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_3d_usm__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d_usm, 70)
-#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_3d_usm__ \
-  0
-#endif
-
 #ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_3d__
 //__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 71)
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_sampled_image_fetch_3d__ 0
@@ -381,6 +365,26 @@
 #define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_cuda_cluster_group__ 0
 #endif
 
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_image_array__
+//__SYCL_ASPECT(ext_oneapi_image_array, 75)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_image_array__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_unique_addressing_per_dim__
+//__SYCL_ASPECT(ext_oneapi_unique_addressing_per_dim, 76)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_unique_addressing_per_dim__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_sample_1d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 77)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_sample_1d_usm__ 0
+#endif
+
+#ifndef __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_sample_2d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 78)
+#define __SYCL_ALL_DEVICES_HAVE_ext_oneapi_bindless_images_sample_2d_usm__ 0
+#endif
+
 #ifndef __SYCL_ANY_DEVICE_HAS_host__
 // __SYCL_ASPECT(host, 0)
 #define __SYCL_ANY_DEVICE_HAS_host__ 0
@@ -611,21 +615,11 @@
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_import__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__
-//__SYCL_ASPECT(ext_oneapi_interop_memory_export, 47)
-#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_memory_export__ 0
-#endif
-
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__
 //__SYCL_ASPECT(ext_oneapi_interop_semaphore_import, 48)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_import__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__
-//__SYCL_ASPECT(ext_oneapi_interop_semaphore_export, 49)
-#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_interop_semaphore_export__ 0
-#endif
-
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__
 //__SYCL_ASPECT(ext_oneapi_mipmap, 50)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_mipmap__ 0
@@ -687,66 +681,81 @@
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_limited_graph__
-// __SYCL_ASPECT(ext_oneapi_limited_graph, 63)
+// __SYCL_ASPECT(ext_oneapi_limited_graph, 62)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_limited_graph__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_private_alloca__
-// __SYCL_ASPECT(ext_oneapi_private_alloca, 64)
+// __SYCL_ASPECT(ext_oneapi_private_alloca, 63)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_private_alloca__ 1
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_cubemap__
-// __SYCL_ASPECT(ext_oneapi_cubemap, 65)
+// __SYCL_ASPECT(ext_oneapi_cubemap, 64)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_cubemap__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_cubemap_seamless_filtering__
-// __SYCL_ASPECT(ext_oneapi_cubemap_seamless_filtering, 66)
+// __SYCL_ASPECT(ext_oneapi_cubemap_seamless_filtering, 65)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_cubemap_seamless_filtering__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_1d_usm__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d_usm, 67)
+//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d_usm, 66)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_1d_usm__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_1d__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d, 68)
+//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_1d, 67)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_1d__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_2d_usm__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d_usm, 69)
+//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d_usm, 68)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_2d_usm__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_2d__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d, 70)
+//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_2d, 69)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_2d__ 0
 #endif
 
-#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_3d_usm__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d_usm, 71)
-#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_3d_usm__ 0
-#endif
-
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_3d__
-//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 72)
+//__SYCL_ASPECT(ext_oneapi_bindless_sampled_image_fetch_3d, 71)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_sampled_image_fetch_3d__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_queue_profiling_tag__
-// __SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 73)
+// __SYCL_ASPECT(ext_oneapi_queue_profiling_tag, 72)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_queue_profiling_tag__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_virtual_mem__
-// __SYCL_ASPECT(ext_oneapi_virtual_mem, 74)
+// __SYCL_ASPECT(ext_oneapi_virtual_mem, 73)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_virtual_mem__ 0
 #endif
 
 #ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_cuda_cluster_group__
-// __SYCL_ASPECT(ext_oneapi_cuda_cluster_group, 75)
+// __SYCL_ASPECT(ext_oneapi_cuda_cluster_group, 74)
 #define __SYCL_ANY_DEVICE_HAS_ext_oneapi_cuda_cluster_group__ 0
 #endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_image_array__
+//__SYCL_ASPECT(ext_oneapi_image_array, 75)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_image_array__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_unique_addressing_per_dim__
+//__SYCL_ASPECT(ext_oneapi_unique_addressing_per_dim, 76)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_unique_addressing_per_dim__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_sample_1d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_sample_1d_usm, 77)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_sample_1d_usm__ 0
+#endif
+
+#ifndef __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_sample_2d_usm__
+//__SYCL_ASPECT(ext_oneapi_bindless_images_sample_2d_usm, 78)
+#define __SYCL_ANY_DEVICE_HAS_ext_oneapi_bindless_images_sample_2d_usm__ 0
+#endif