From 85a5373b578006f3e1e9471ab0fbe150bb8a4641 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 9 Jul 2024 13:35:01 +0800 Subject: [PATCH 1/2] [SYCL][Bindless] Update spirv read/fetch from sampled image and sampled image array For spirv path, call __invoke__ImageReadLod to generate image builtin represented as ImageSampleExplicitLod instruction in SPIRV that can be consumed by Intel Graphics Compiler. --- .../sycl/ext/oneapi/bindless_images.hpp | 69 ++++++++++++++++--- 1 file changed, 58 insertions(+), 11 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index 9e1381abab76b..842693c18ffcb 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -812,6 +812,11 @@ template using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type< NDims, sycl::access::mode::write, sycl::access::target::image_array>::type; +template +using OCLSampledImageArrayTyRead = + typename sycl::detail::sampled_opencl_image_type< + detail::OCLImageArrayTyRead>::type; + // Macros are required because it is not legal for a function to return // a variable of type 'opencl_image_type'. #if defined(__SPIR__) @@ -823,11 +828,19 @@ using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type< typename sycl::detail::sampled_opencl_image_type< \ detail::OCLImageTyRead>::type>(raw_handle) +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) \ + __spirv_ConvertHandleToSampledImageINTEL< \ + typename sycl::detail::sampled_opencl_image_type< \ + detail::OCLImageArrayTyRead>::type>(raw_handle) + #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ __invoke__ImageRead(raw_handle, coords) #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \ - __invoke__ImageRead(raw_handle, coords) + __invoke__ImageReadLod(raw_handle, coords, 0.f) + +#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \ + __invoke__ImageReadLod(raw_handle, coords, 0.f) #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \ __invoke__ImageRead(raw_handle, coordsLayer) @@ -835,6 +848,14 @@ using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type< #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \ __invoke__ImageWrite(raw_handle, coordsLayer, color) +#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \ + coordsLayer) \ + __invoke__ImageReadLod(raw_handle, coordsLayer, 0.f) + +#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \ + coordsLayer) \ + __invoke__ImageReadLod(raw_handle, coordsLayer, 0.f) + #else #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle @@ -846,11 +867,23 @@ using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type< #define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \ __invoke__SampledImageFetch(raw_handle, coords) +#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \ + __invoke__ImageRead(raw_handle, coords) + #define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \ __invoke__ImageArrayFetch(raw_handle, coords, arrayLayer) #define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \ __invoke__ImageArrayWrite(raw_handle, coords, arrayLayer, color) + +#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \ + coordsLayer) \ + __invoke__SampledImageArrayFetch(raw_handle, coords, arrayLayer) + +#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \ + coordsLayer) \ + __invoke__ImageArrayRead(raw_handle, coords, arrayLayer) + #endif #endif // __SYCL_DEVICE_ONLY__ @@ -1030,11 +1063,13 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageRead( + return SAMPLE_IMAGE_READ( + DataT, CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords); } else { - return sycl::bit_cast(__invoke__ImageRead( + return sycl::bit_cast(SAMPLE_IMAGE_READ( + HintT, CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords)); } @@ -1410,17 +1445,23 @@ DataT fetch_image_array(const sampled_image_handle &imageHandle "and 2D images respectively."); #ifdef __SYCL_DEVICE_ONLY__ + sycl::vec coordsLayer{coords, arrayLayer}; if constexpr (detail::is_recognized_standard_type()) { - return __invoke__SampledImageArrayFetch(imageHandle.raw_handle, - coords, arrayLayer); + return FETCH_SAMPLED_IMAGE_ARRAY(DataT, + CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY( + imageHandle.raw_handle, coordSize), + coords, arrayLayer, coordsLayer); } 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 must always be a recognized standard type"); - return sycl::bit_cast(__invoke__SampledImageArrayFetch( - imageHandle.raw_handle, coords, arrayLayer)); + return sycl::bit_cast( + FETCH_SAMPLED_IMAGE_ARRAY(HintT, + CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY( + imageHandle.raw_handle, coordSize), + coords, arrayLayer, coordsLayer)); } #else assert(false); // Bindless images not yet implemented on host. @@ -1454,17 +1495,23 @@ DataT sample_image_array(const sampled_image_handle &imageHandle "and 2D images respectively."); #ifdef __SYCL_DEVICE_ONLY__ + sycl::vec coordsLayer{coords, arrayLayer}; if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageArrayRead(imageHandle.raw_handle, coords, - arrayLayer); + return READ_SAMPLED_IMAGE_ARRAY(DataT, + CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY( + imageHandle.raw_handle, coordSize), + coords, arrayLayer, coordsLayer); } 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 must always be a recognized standard type"); - return sycl::bit_cast(__invoke__ImageArrayRead( - imageHandle.raw_handle, coords, arrayLayer)); + return sycl::bit_cast( + READ_SAMPLED_IMAGE_ARRAY(HintT, + CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY( + imageHandle.raw_handle, coordSize), + coords, arrayLayer, coordsLayer)); } #else assert(false); // Bindless images not yet implemented on host. From 62f49ad2bfa0568bc2838f22c6bf2539dc5af925 Mon Sep 17 00:00:00 2001 From: Wenju He Date: Tue, 9 Jul 2024 15:36:21 +0800 Subject: [PATCH 2/2] add missing CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY, fix lit --- sycl/include/sycl/ext/oneapi/bindless_images.hpp | 2 ++ sycl/test/check_device_code/extensions/bindless_images.cpp | 4 ++-- 2 files changed, 4 insertions(+), 2 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index 842693c18ffcb..75e7f06297d4e 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -861,6 +861,8 @@ using OCLSampledImageArrayTyRead = #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) raw_handle + #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ __invoke__ImageFetch(raw_handle, coords) diff --git a/sycl/test/check_device_code/extensions/bindless_images.cpp b/sycl/test/check_device_code/extensions/bindless_images.cpp index 03ad366d08720..b98f3a1434592 100644 --- a/sycl/test/check_device_code/extensions/bindless_images.cpp +++ b/sycl/test/check_device_code/extensions/bindless_images.cpp @@ -4,7 +4,7 @@ // RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-LLVM // CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4 -// CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4 +// CHECK-LLVM: tail call spir_func noundef <4 x float> @_Z30__spirv_ImageSampleExplicitLod // CHECK-LLVM: tail call spir_func void @_Z18__spirv_ImageWriteI14 // RUN: %clangxx -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o %t.out @@ -50,7 +50,7 @@ // Read sampled image // Arguments: Result Type, Result, Image, Coords -// CHECK-SPIRV-NEXT: ImageRead [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}} +// CHECK-SPIRV-NEXT: ImageSampleExplicitLod [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}} // Convert handle to SPIR-V image // Arguments: Result Type, Result, Handle