[SYCL][Bindless] Update and add support for SPV_INTEL_bindless_image extension new revision (#13753)

DBDuncan · web-flow · commit 83db85f19643 · 2024-05-24T14:36:30.000+02:00
Add support to emit instructions that convert handles to images,
samplers and sampled images
diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp
@@ -10411,7 +10411,8 @@ static void getOtherSPIRVTransOpts(Compilation &C,
       ",+SPV_INTEL_fpga_argument_interfaces"
       ",+SPV_INTEL_fpga_invocation_pipelining_attributes"
       ",+SPV_INTEL_fpga_latency_control"
-      ",+SPV_INTEL_task_sequence";
+      ",+SPV_INTEL_task_sequence"
+      ",+SPV_INTEL_bindless_images";
   ExtArg = ExtArg + DefaultExtArg + INTELExtArg;
   if (C.getDriver().IsFPGAHWMode())
     // Enable several extensions on FPGA H/W exclusively
diff --git a/clang/test/Driver/sycl-spirv-ext.c b/clang/test/Driver/sycl-spirv-ext.c
@@ -48,6 +48,7 @@
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_invocation_pipelining_attributes
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_fpga_latency_control
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_task_sequence
+// CHECK-DEFAULT-SAME:,+SPV_INTEL_bindless_images
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_token_type
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_bfloat16_conversion
 // CHECK-DEFAULT-SAME:,+SPV_INTEL_joint_matrix
diff --git a/sycl/doc/design/spirv-extensions/SPV_INTEL_bindless_images.asciidoc b/sycl/doc/design/spirv-extensions/SPV_INTEL_bindless_images.asciidoc
@@ -1,6 +1,7 @@
 :capability_token: 6528
 :handle_to_image_token: 6529
 :handle_to_sampler_token: 6530
+:handle_to_sampled_image_token: 6531
 
 SPV_INTEL_bindless_images
 =========================
@@ -37,8 +38,8 @@ In Development
 
 [width="40%",cols="25,25"]
 |========================================
-| Last Modified Date | 2024-03-25
-| Revision           | 6
+| Last Modified Date | 2024-05-01
+| Revision           | 7
 |========================================
 
 == Dependencies
@@ -52,7 +53,7 @@ This extension requires SPIR-V 1.0.
 
 This extension adds support for bindless images.
 This is done by adding support for SPIR-V to convert unsigned integer handles to
-images/samplers.
+images, samplers and sampled images.
 
 Bindless images are a feature that provides flexibility on how images are
 accessed and used, such as removing limitations on how many images can be
@@ -84,6 +85,7 @@ Instructions added under *BindlessImagesINTEL* capability.
 ----
 OpConvertHandleToImageINTEL
 OpConvertHandleToSamplerINTEL
+OpConvertHandleToSampledImageINTEL
 ----
 
 == Token Number Assignments
@@ -93,9 +95,10 @@ OpConvertHandleToSamplerINTEL
 [cols="70%,30%"]
 [grid="rows"]
 |====
-|BindlessImagesINTEL                |{capability_token}
-|OpConvertHandleToImageINTEL        |{handle_to_image_token}
-|OpConvertHandleToSamplerINTEL      |{handle_to_sampler_token}
+|BindlessImagesINTEL                    |{capability_token}
+|OpConvertHandleToImageINTEL            |{handle_to_image_token}
+|OpConvertHandleToSamplerINTEL          |{handle_to_sampler_token}
+|OpConvertHandleToSampledImageINTEL     |{handle_to_sampled_image_token}
 |====
 --
 
@@ -134,6 +137,21 @@ _Result type_ must be an `OpTypeSampler`.
 '<id> Operand'
 |======
 
+[cols="2*1,3*2"]
+|======
+5+|[[OpConvertHandleToSampledImageINTEL]]*OpConvertHandleToSampledImageINTEL* +
+ +
+Converts an unsigned integer pointed by _Operand_ to sampled image type.
+
+Unsigned integer is either a 32 or 64 bit unsigned integer.
+Depending on if the addressing model is set to *Physical32* or *Physical64*.
+
+_Result type_ must be an `OpTypeSampledImage`.
+
+| 4 | {handle_to_sampled_image_token} | '<id> Result Type' | 'Result <id>' |
+'<id> Operand'
+|======
+
 Modify Section 3.31, Capability, adding row to the capability table:
 
 [width="40%"]
@@ -164,6 +182,7 @@ None Yet.
                               instruction and clarify return types
 |6|2024-03-25|Duncan Brawley| Wording/formatting improvements, clarify sections
                               edited, make capability addition explicit and
-                              substitute instruction numbers            
+                              substitute instruction numbers
+|7|2024-05-01|Duncan Brawley| Add OpConvertHandleToSampledImageINTEL instruction
 |========================================
 
diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp
@@ -230,6 +230,16 @@ template <typename SampledType, typename TempRetT, typename TempArgT>
 extern __DPCPP_SYCL_EXTERNAL TempRetT __spirv_ImageSampleCubemap(SampledType,
                                                                  TempArgT);
 
+template <typename RetT, class HandleT>
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ConvertHandleToImageINTEL(HandleT);
+
+template <typename RetT, class HandleT>
+extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ConvertHandleToSamplerINTEL(HandleT);
+
+template <typename RetT, class HandleT>
+extern __DPCPP_SYCL_EXTERNAL
+    RetT __spirv_ConvertHandleToSampledImageINTEL(HandleT);
+
 #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy
 #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy
 
diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp
@@ -50,8 +50,7 @@ struct sampled_image_handle {
 
   sampled_image_handle() : raw_handle(~0) {}
 
-  sampled_image_handle(raw_image_handle_type raw_image_handle)
-      : raw_handle(raw_image_handle) {}
+  sampled_image_handle(raw_image_handle_type handle) : raw_handle(handle) {}
 
   raw_image_handle_type raw_handle;
 };
@@ -792,6 +791,43 @@ template <typename DataT> constexpr bool is_recognized_standard_type() {
           std::is_floating_point_v<DataT> || std::is_same_v<DataT, sycl::half>);
 }
 
+#ifdef __SYCL_DEVICE_ONLY__
+
+// Image types used for generating SPIR-V
+template <int NDims>
+using OCLImageTyRead =
+    typename sycl::detail::opencl_image_type<NDims, sycl::access::mode::read,
+                                             sycl::access::target::image>::type;
+
+template <int NDims>
+using OCLImageTyWrite =
+    typename sycl::detail::opencl_image_type<NDims, sycl::access::mode::write,
+                                             sycl::access::target::image>::type;
+
+// Macros are required because it is not legal for a function to return
+// a variable of type 'opencl_image_type'.
+#if defined(__SPIR__)
+#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType)                         \
+  __spirv_ConvertHandleToImageINTEL<ImageType>(raw_handle)
+
+#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims)                     \
+  __spirv_ConvertHandleToSampledImageINTEL<                                    \
+      typename sycl::detail::sampled_opencl_image_type<                        \
+          detail::OCLImageTyRead<NDims>>::type>(raw_handle)
+
+#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
+  __invoke__ImageRead<DataT>(raw_handle, coords)
+#else
+#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle
+
+#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle
+
+#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords)                       \
+  __invoke__ImageFetch<DataT>(raw_handle, coords)
+#endif
+
+#endif
+
 } // namespace detail
 
 /**
@@ -826,15 +862,23 @@ DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageFetch<DataT>(imageHandle.raw_handle, coords);
+    return FETCH_UNSAMPLED_IMAGE(
+        DataT,
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords);
+
   } else {
     static_assert(sizeof(HintT) == sizeof(DataT),
                   "When trying to read a user-defined type, HintT must be of "
                   "the same size as the user-defined DataT.");
     static_assert(detail::is_recognized_standard_type<HintT>(),
                   "HintT must always be a recognized standard type");
-    return sycl::bit_cast<DataT>(
-        __invoke__ImageFetch<HintT>(imageHandle.raw_handle, coords));
+    return sycl::bit_cast<DataT>(FETCH_UNSAMPLED_IMAGE(
+        HintT,
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords));
   }
 #else
   assert(false); // Bindless images not yet implemented on host
@@ -907,10 +951,13 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__SampledImageFetch<DataT>(imageHandle.raw_handle, coords);
+    return __invoke__SampledImageFetch<DataT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords);
   } else {
-    return sycl::bit_cast<DataT>(
-        __invoke__SampledImageFetch<HintT>(imageHandle.raw_handle, coords));
+    return sycl::bit_cast<DataT>(__invoke__SampledImageFetch<HintT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.
@@ -954,10 +1001,13 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageRead<DataT>(imageHandle.raw_handle, coords);
+    return __invoke__ImageRead<DataT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords);
   } else {
-    return sycl::bit_cast<DataT>(
-        __invoke__ImageRead<HintT>(imageHandle.raw_handle, coords));
+    return sycl::bit_cast<DataT>(__invoke__ImageRead<HintT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.
@@ -1026,15 +1076,18 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageReadLod<DataT>(imageHandle.raw_handle, coords, level);
+    return __invoke__ImageReadLod<DataT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords, level);
   } else {
     static_assert(sizeof(HintT) == sizeof(DataT),
                   "When trying to read a user-defined type, HintT must be of "
                   "the same size as the user-defined DataT.");
     static_assert(detail::is_recognized_standard_type<HintT>(),
                   "HintT must always be a recognized standard type");
-    return sycl::bit_cast<DataT>(
-        __invoke__ImageReadLod<HintT>(imageHandle.raw_handle, coords, level));
+    return sycl::bit_cast<DataT>(__invoke__ImageReadLod<HintT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords, level));
   }
 #else
   assert(false); // Bindless images not yet implemented on host
@@ -1070,16 +1123,18 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageReadGrad<DataT>(imageHandle.raw_handle, coords, dX,
-                                          dY);
+    return __invoke__ImageReadGrad<DataT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords, dX, dY);
   } else {
     static_assert(sizeof(HintT) == sizeof(DataT),
                   "When trying to read a user-defined type, HintT must be of "
                   "the same size as the user-defined DataT.");
     static_assert(detail::is_recognized_standard_type<HintT>(),
                   "HintT must always be a recognized standard type");
-    return sycl::bit_cast<DataT>(
-        __invoke__ImageReadGrad<HintT>(imageHandle.raw_handle, coords, dX, dY));
+    return sycl::bit_cast<DataT>(__invoke__ImageReadGrad<HintT>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
+        coords, dX, dY));
   }
 #else
   assert(false); // Bindless images not yet implemented on host
@@ -1224,16 +1279,20 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageArrayFetch<DataT>(imageHandle.raw_handle, coords,
-                                            arrayLayer);
+    return __invoke__ImageArrayFetch<DataT>(
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords, arrayLayer);
   } else {
     static_assert(sizeof(HintT) == sizeof(DataT),
                   "When trying to fetch a user-defined type, HintT must be of "
                   "the same size as the user-defined DataT.");
     static_assert(detail::is_recognized_standard_type<HintT>(),
                   "HintT must always be a recognized standard type");
     return sycl::bit_cast<DataT>(__invoke__ImageArrayFetch<HintT>(
-        imageHandle.raw_handle, coords, arrayLayer));
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords, arrayLayer));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.
@@ -1277,19 +1336,21 @@ DataT fetch_cubemap(const unsampled_image_handle &imageHandle,
 template <typename DataT, typename HintT = DataT>
 DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]],
                      const sycl::float3 &dirVec [[maybe_unused]]) {
+  [[maybe_unused]] constexpr size_t NDims = 2;
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    return __invoke__ImageReadCubemap<DataT, uint64_t>(imageHandle.raw_handle,
-                                                       dirVec);
+    return __invoke__ImageReadCubemap<DataT, uint64_t>(
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, NDims), dirVec);
   } else {
     static_assert(sizeof(HintT) == sizeof(DataT),
                   "When trying to read a user-defined type, HintT must be of "
                   "the same size as the user-defined DataT.");
     static_assert(detail::is_recognized_standard_type<HintT>(),
                   "HintT must always be a recognized standard type");
     return sycl::bit_cast<DataT>(__invoke__ImageReadCubemap<HintT, uint64_t>(
-        imageHandle.raw_handle, dirVec));
+        CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, NDims),
+        dirVec));
   }
 #else
   assert(false); // Bindless images not yet implemented on host
@@ -1318,12 +1379,17 @@ void write_image(unsampled_image_handle imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color);
+    __invoke__ImageWrite(
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyWrite<coordSize>),
+        coords, color);
   } else {
     // Convert DataT to a supported backend write type when user-defined type is
     // passed
-    __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords,
-                         detail::convert_color(color));
+    __invoke__ImageWrite(
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyWrite<coordSize>),
+        coords, detail::convert_color(color));
   }
 #else
   assert(false); // Bindless images not yet implemented on host
@@ -1354,13 +1420,17 @@ void write_image_array(unsampled_image_handle imageHandle [[maybe_unused]],
 
 #ifdef __SYCL_DEVICE_ONLY__
   if constexpr (detail::is_recognized_standard_type<DataT>()) {
-    __invoke__ImageArrayWrite(static_cast<uint64_t>(imageHandle.raw_handle),
-                              coords, arrayLayer, color);
+    __invoke__ImageArrayWrite(
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords, arrayLayer, color);
   } else {
     // Convert DataT to a supported backend write type when user-defined type is
     // passed
-    __invoke__ImageArrayWrite(static_cast<uint64_t>(imageHandle.raw_handle),
-                              coords, arrayLayer, detail::convert_color(color));
+    __invoke__ImageArrayWrite(
+        CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle,
+                                detail::OCLImageTyRead<coordSize>),
+        coords, arrayLayer, detail::convert_color(color));
   }
 #else
   assert(false); // Bindless images not yet implemented on host.
diff --git a/sycl/test/extensions/bindless_images.cpp b/sycl/test/extensions/bindless_images.cpp
diff --git a/sycl/test/extensions/bindless_images_SPIRV_inst.cpp b/sycl/test/extensions/bindless_images_SPIRV_inst.cpp
