From dd48b8ea4fa8697bb97a7bc6d161077854997bbf Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Fri, 10 May 2024 17:57:38 +0100 Subject: [PATCH 1/7] [SYCL][Bindless] Update and add support for SPV_INTEL_bindless_images extension new revision Add support to emit instructions that convert handles to images, samplers and sampled images --- clang/lib/Driver/ToolChains/Clang.cpp | 3 +- .../SPV_INTEL_bindless_images.asciidoc | 33 +++- sycl/include/CL/__spirv/spirv_ops.hpp | 12 ++ .../sycl/ext/oneapi/bindless_images.hpp | 154 ++++++++++++++---- sycl/test/extensions/bindless_images.cpp | 90 +++++++++- .../extensions/bindless_images_SPIRV_inst.cpp | 126 ++++++++++++++ 6 files changed, 370 insertions(+), 48 deletions(-) create mode 100644 sycl/test/extensions/bindless_images_SPIRV_inst.cpp diff --git a/clang/lib/Driver/ToolChains/Clang.cpp b/clang/lib/Driver/ToolChains/Clang.cpp index acd2fd611b41f..48241d61c2a7f 100644 --- a/clang/lib/Driver/ToolChains/Clang.cpp +++ b/clang/lib/Driver/ToolChains/Clang.cpp @@ -10394,7 +10394,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/sycl/doc/design/spirv-extensions/SPV_INTEL_bindless_images.asciidoc b/sycl/doc/design/spirv-extensions/SPV_INTEL_bindless_images.asciidoc index 45aa3f4593fc0..af399aeedd373 100644 --- 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`. ' 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} | ' Result Type' | 'Result ' | +' 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 index b9f7b28d8ce46..501e4c3059fd4 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -230,6 +230,18 @@ template extern __DPCPP_SYCL_EXTERNAL TempRetT __spirv_ImageSampleCubemap(SampledType, TempArgT); +template +extern __DPCPP_SYCL_EXTERNAL RetT +__spirv_ConvertHandleToImageINTEL(unsigned long); + +template +extern __DPCPP_SYCL_EXTERNAL RetT +__spirv_ConvertHandleToSamplerINTEL(unsigned long); + +template +extern __DPCPP_SYCL_EXTERNAL RetT +__spirv_ConvertHandleToSampledImageINTEL(unsigned long); + #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 index e11903a085ab5..917693f13d5e4 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -50,12 +50,24 @@ 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; }; +// Image types used for generating SPIR-V +#ifdef __SYCL_DEVICE_ONLY__ +template +using OCLImageTyRead = + typename sycl::detail::opencl_image_type::type; + +template +using OCLImageTyWrite = + typename sycl::detail::opencl_image_type::type; +#endif + /** * @brief Allocate image memory based on image_descriptor * @@ -792,6 +804,43 @@ template constexpr bool is_recognized_standard_type() { std::is_floating_point_v || std::is_same_v); } +#ifdef __SYCL_DEVICE_ONLY__ + +// Macros are required because it is not legal for a function to return +// a variable of type 'opencl_image_type'. +#if defined(__NVPTX__) +#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle +#elif defined(__SPIR__) +#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) \ + __spirv_ConvertHandleToImageINTEL(raw_handle) +#else +#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle +#endif + +#if defined(__NVPTX__) +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) raw_handle +#elif defined(__SPIR__) +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) \ + __spirv_ConvertHandleToSampledImageINTEL< \ + typename sycl::detail::sampled_opencl_image_type::type>( \ + raw_handle) +#else +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) raw_handle +#endif + +#if defined(__NVPTX__) +#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ + __invoke__ImageFetch(raw_handle, coords) +#elif defined(__SPIR__) +#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ + __invoke__ImageRead(raw_handle, coords) +#else +#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ + __invoke__ImageFetch(raw_handle, coords) +#endif + +#endif + } // namespace detail /** @@ -826,15 +875,23 @@ DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageFetch(imageHandle.raw_handle, coords); + return FETCH_UNSAMPLED_IMAGE( + DataT, + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + 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 must always be a recognized standard type"); - return sycl::bit_cast( - __invoke__ImageFetch(imageHandle.raw_handle, coords)); + return sycl::bit_cast(FETCH_UNSAMPLED_IMAGE( + HintT, + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords)); } #else assert(false); // Bindless images not yet implemented on host @@ -907,10 +964,15 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__SampledImageFetch(imageHandle.raw_handle, coords); + return __invoke__SampledImageFetch( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords); } else { - return sycl::bit_cast( - __invoke__SampledImageFetch(imageHandle.raw_handle, coords)); + return sycl::bit_cast(__invoke__SampledImageFetch( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords)); } #else assert(false); // Bindless images not yet implemented on host. @@ -954,10 +1016,15 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageRead(imageHandle.raw_handle, coords); + return __invoke__ImageRead( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords); } else { - return sycl::bit_cast( - __invoke__ImageRead(imageHandle.raw_handle, coords)); + return sycl::bit_cast(__invoke__ImageRead( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords)); } #else assert(false); // Bindless images not yet implemented on host. @@ -1026,15 +1093,20 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageReadLod(imageHandle.raw_handle, coords, level); + return __invoke__ImageReadLod( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + 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 must always be a recognized standard type"); - return sycl::bit_cast( - __invoke__ImageReadLod(imageHandle.raw_handle, coords, level)); + return sycl::bit_cast(__invoke__ImageReadLod( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords, level)); } #else assert(false); // Bindless images not yet implemented on host @@ -1070,16 +1142,20 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageReadGrad(imageHandle.raw_handle, coords, dX, - dY); + return __invoke__ImageReadGrad( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + 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 must always be a recognized standard type"); - return sycl::bit_cast( - __invoke__ImageReadGrad(imageHandle.raw_handle, coords, dX, dY)); + return sycl::bit_cast(__invoke__ImageReadGrad( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords, dX, dY)); } #else assert(false); // Bindless images not yet implemented on host @@ -1224,8 +1300,10 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - return __invoke__ImageArrayFetch(imageHandle.raw_handle, coords, - arrayLayer); + return __invoke__ImageArrayFetch( + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords, arrayLayer); } else { static_assert(sizeof(HintT) == sizeof(DataT), "When trying to fetch a user-defined type, HintT must be of " @@ -1233,7 +1311,9 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle static_assert(detail::is_recognized_standard_type(), "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageArrayFetch( - imageHandle.raw_handle, coords, arrayLayer)); + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords, arrayLayer)); } #else assert(false); // Bindless images not yet implemented on host. @@ -1269,11 +1349,14 @@ DataT fetch_cubemap(const unsampled_image_handle &imageHandle, template 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()) { - return __invoke__ImageReadCubemap(imageHandle.raw_handle, - dirVec); + return __invoke__ImageReadCubemap( + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + dirVec); } else { static_assert(sizeof(HintT) == sizeof(DataT), "When trying to read a user-defined type, HintT must be of " @@ -1281,7 +1364,9 @@ DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]], static_assert(detail::is_recognized_standard_type(), "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageReadCubemap( - imageHandle.raw_handle, dirVec)); + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + dirVec)); } #else assert(false); // Bindless images not yet implemented on host @@ -1310,12 +1395,15 @@ void write_image(unsampled_image_handle imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color); + __invoke__ImageWrite(CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyWrite), + 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, + OCLImageTyWrite), + coords, detail::convert_color(color)); } #else assert(false); // Bindless images not yet implemented on host @@ -1346,13 +1434,17 @@ void write_image_array(unsampled_image_handle imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - __invoke__ImageArrayWrite(static_cast(imageHandle.raw_handle), - coords, arrayLayer, color); + __invoke__ImageArrayWrite( + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + coords, arrayLayer, color); } else { // Convert DataT to a supported backend write type when user-defined type is // passed - __invoke__ImageArrayWrite(static_cast(imageHandle.raw_handle), - coords, arrayLayer, detail::convert_color(color)); + __invoke__ImageArrayWrite( + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + OCLImageTyRead), + 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 index eece72c070930..04442da0054c0 100644 --- a/sycl/test/extensions/bindless_images.cpp +++ b/sycl/test/extensions/bindless_images.cpp @@ -1,10 +1,66 @@ -// RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s +// Check SPV_INTEL_bindless_images instructions are emitted correctly in a +// realistic scenario. + +// RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-LLVM + +// Verify the mangled names of the kernel wrapper and image accesses contain +// the expected types. +// CHECK-LLVM: spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI10image_readEE +// 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 void @_Z18__spirv_ImageWriteI14 + +// RUN: %clangxx -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o %t.out +// RUN: llvm-spirv -spirv-ext=+SPV_INTEL_bindless_images %t.out -spirv-text -o %t.out.spv +// RUN: FileCheck %s --input-file %t.out.spv --check-prefix=CHECK-SPIRV #include #include -// CHECK: spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI10image_readEE -// CHECK: tail call spir_func noundef <4 x float> @_Z18__spirv_ImageFetchIDv4 +// Data type of image pixel components +// Arguments: Result, width +// CHECK-SPIRV: TypeFloat [[PIXELCOMPTYPE:[0-9]+]] 32 + +// Generate the appropriate `Result Type` used by `ConvertHandleToImageINTEL` +// and `ConvertHandleToSampledImageINTEL` Operand `7` here represents +// 'TypeVoid`. Must be `TypeVoid` as the type of the image is not known at +// compile time. The last operand is the access qualifier. With 0 read only and +// 1 write only. Arguments: TypeImage, Result, Sampled Type, Dim, Depth, +// Arrayed, MS, Sampled, Image Format. +// CHECK-SPIRV: TypeImage [[IMAGETYPE:[0-9]+]] 7 0 0 0 0 0 0 0 + +// Image pixel data type +// Arguments: Result, Component Type, Component Count +// CHECK-SPIRV: TypeVector [[PIXELTYPE:[0-9]+]] [[PIXELCOMPTYPE]] 4 + +// CHECK-SPIRV: TypeSampledImage [[SAMPIMAGETYPE:[0-9]+]] [[IMAGETYPE]] + +// CHECK-SPIRV: TypeImage [[IMAGETYPEREAD:[0-9]+]] 7 0 0 0 0 0 0 1 + +// Convert handle to SPIR-V image +// Arguments: Result Type, Result, Handle +// CHECK-SPIRV: ConvertHandleToImageINTEL [[IMAGETYPE]] [[IMAGEVARONE:[0-9]+]] {{[0-9]+}} + +// Read image +// Arguments: Result Type, Result, Image, Coords +// CHECK-SPIRV-NEXT: ImageRead [[PIXELTYPE]] {{[0-9]+}} [[IMAGEVARONE]] {{[0-9]+}} + +// Convert handle to SPIR-V sampled image +// Arguments: Result Type, Result, Handle +// CHECK-SPIRV: ConvertHandleToSampledImageINTEL [[SAMPIMAGETYPE]] [[SAMPIMAGEVAR:[0-9]+]] {{[0-9]+}} + +// Read sampled image +// Arguments: Result Type, Result, Image, Coords +// CHECK-SPIRV-NEXT: ImageRead [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}} + +// Convert handle to SPIR-V image +// Arguments: Result Type, Result, Handle +// CHECK-SPIRV: ConvertHandleToImageINTEL [[IMAGETYPEREAD]] [[IMAGEVARTWO:[0-9]+]] {{[0-9]+}} + +// Write unsampled image +// Arguments: Image, Coords, Data +// CHECK-SPIRV: ImageWrite [[IMAGEVARTWO]] {{[0-9]+}} {{[0-9]+}} + using namespace sycl::ext::oneapi::experimental; class image_read; int main() { @@ -15,19 +71,30 @@ int main() { constexpr size_t width = 512; std::vector out(width); - std::vector dataIn1(width); + std::vector dataIn(width); for (int i = 0; i < width; i++) { - dataIn1[i] = sycl::float4(i, i, i, i); + dataIn[i] = sycl::float4(i, i, i, i); } { image_descriptor desc({width}, sycl::image_channel_order::rgba, sycl::image_channel_type::fp32); - image_mem imgMem0(desc, dev, ctxt); - unsampled_image_handle imgHandle1 = create_image(imgMem0, desc, dev, ctxt); + sycl::ext::oneapi::experimental::bindless_image_sampler samp( + sycl::addressing_mode::clamp, + sycl::coordinate_normalization_mode::unnormalized, + sycl::filtering_mode::nearest); + + image_mem imgMem1(desc, dev, ctxt); + image_mem imgMem2(desc, dev, ctxt); - q.ext_oneapi_copy(dataIn1.data(), imgMem0.get_handle(), desc); + unsampled_image_handle imgHandle1 = create_image(imgMem1, desc, dev, ctxt); + + sampled_image_handle imgHandle2 = + create_image(imgMem2, samp, desc, dev, ctxt); + + q.ext_oneapi_copy(dataIn.data(), imgMem1.get_handle(), desc); + q.ext_oneapi_copy(dataIn.data(), imgMem2.get_handle(), desc); q.wait_and_throw(); sycl::buffer buf((float *)out.data(), width); @@ -36,7 +103,12 @@ int main() { cgh.parallel_for(width, [=](sycl::id<1> id) { sycl::float4 px1 = fetch_image(imgHandle1, int(id[0])); - outAcc[id] = px1[0]; + + sycl::float4 px2 = sample_image(imgHandle2, float(id[0])); + + write_image(imgHandle1, int(id[0]), px1 + px2); + + outAcc[id] = px1[0] + px2[0]; }); }); diff --git a/sycl/test/extensions/bindless_images_SPIRV_inst.cpp b/sycl/test/extensions/bindless_images_SPIRV_inst.cpp new file mode 100644 index 0000000000000..3ade6497549e5 --- /dev/null +++ b/sycl/test/extensions/bindless_images_SPIRV_inst.cpp @@ -0,0 +1,126 @@ +// Check SPV_INTEL_bindless_images instructions are emitted correctly. + +// RUN: %clangxx -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o %t.out +// RUN: llvm-spirv -spirv-ext=+SPV_INTEL_bindless_images %t.out -spirv-text -o %t.out.spv +// RUN: FileCheck %s --input-file %t.out.spv + +#include +#include + +// Data type of image, sampler and sampled image handles +// Arguments: Result, width, signedness +// CHECK: TypeInt 2 64 0 + +// Create dummy image and sampled image handles +// 2 here, represents a 64-bit int type +// Arguments: Result Type, Result, Literal Value +// CHECK: Constant 2 [[IMAGEID:[0-9]+]] 123 +// CHECK: Constant 2 [[IMAGETWOID:[0-9]+]] 1234 +// CHECK: Constant 2 [[SAMPLERID:[0-9]+]] 12345 +// CHECK: Constant 2 [[SAMPLEDIMAGEID:[0-9]+]] 123456 + +// Generate the appropriate `Result Type`s used by `ConvertHandleToImageINTEL`, +// `ConvertHandleToSamplerINTEL` and `ConvertHandleToSampledImageINTEL`. Operand +// `7` here represents 'TypeVoid`. Must be `TypeVoid` as the type of the image +// is not known at compile time. The last operand is the access qualifier. With +// 0 read only and 1 write only. +// Arguments: TypeImage, Result, Sampled Type, Dim, Depth, Arrayed, MS, Sampled +// and Image Format. +// CHECK: TypeImage [[IMAGETYPEREAD:[0-9]+]] 7 0 0 0 0 0 0 0 +// CHECK: TypeImage [[IMAGETYPEWRITE:[0-9]+]] 7 0 0 0 0 0 0 1 +// Generate `Result Type` for samplers +// Arguments: Result +// CHECK: TypeSampler [[SAMPLERTYPE:[0-9]+]] +// Generate `Result Type` for sampled images +// Arguments: Result, Image Type +// CHECK: TypeSampledImage [[SAMPLEDIMAGETYPE:[0-9]+]] [[IMAGETYPEREAD]] + +// Convert handles to SPIR-V images, samplers and sampled images +// Arguments: Result Type, Result, Handle +// CHECK: ConvertHandleToImageINTEL [[IMAGETYPEREAD]] {{[0-9]+}} [[IMAGEID]] +// CHECK: ConvertHandleToImageINTEL [[IMAGETYPEWRITE]] {{[0-9]+}} [[IMAGETWOID]] +// CHECK: ConvertHandleToSamplerINTEL [[SAMPLERTYPE]] {{[0-9]+}} [[SAMPLERID]] +// CHECK: ConvertHandleToSampledImageINTEL [[SAMPLEDIMAGETYPE]] {{[0-9]+}} [[SAMPLEDIMAGEID]] + +#include + +#ifdef __SYCL_DEVICE_ONLY__ +template +using OCLImageTyRead = + typename sycl::detail::opencl_image_type::type; + +template +using OCLImageTyWrite = + typename sycl::detail::opencl_image_type::type; + +template +using OCLSampledImageTy = typename sycl::detail::sampled_opencl_image_type< + OCLImageTyRead>::type; +#endif + +template +ReturnT handleToImage(const unsigned long &imageHandle) { +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__SPIR__) + return __spirv_ConvertHandleToImageINTEL(imageHandle); +#endif +#else + assert(false); // Bindless images not yet implemented on host. +#endif +} + +template +ReturnT handleToSampler(const unsigned long &samplerHandle) { +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__SPIR__) + return __spirv_ConvertHandleToSamplerINTEL(samplerHandle); +#endif +#else + assert(false); // Bindless images not yet implemented on host. +#endif +} + +template +ReturnT handleToSampledImage(const unsigned long &imageHandle) { +#ifdef __SYCL_DEVICE_ONLY__ +#if defined(__SPIR__) + return __spirv_ConvertHandleToSampledImageINTEL(imageHandle); +#endif +#else + assert(false); // Bindless images not yet implemented on host. +#endif +} + +class image_addition; + +int main() { + + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + try { + q.submit([&](sycl::handler &cgh) { + cgh.parallel_for(1, [=](sycl::id<1> id) { +#ifdef __SYCL_DEVICE_ONLY__ + OCLImageTyRead<1> imageRead = handleToImage>(123); + + OCLImageTyWrite<1> imageWrite = handleToImage>(1234); + + __ocl_sampler_t sampler = handleToSampler<__ocl_sampler_t>(12345); + + OCLSampledImageTy<1> sampImage = + handleToSampledImage>(123456); +#endif + }); + }); + + } catch (sycl::exception e) { + std::cerr << "SYCL exception caught! : " << e.what() << "\n"; + return 1; + } catch (...) { + std::cerr << "Unknown exception caught!\n"; + return 2; + } +} From 7fd6ce01e66260a0a8a2186d1f7b311f9f698fa2 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Tue, 14 May 2024 15:17:48 +0100 Subject: [PATCH 2/7] Move opencl types to detail namespace and move 'OCLImageTyRead' to inside 'CONVERT_HANDLE_TO_SAMPLED_IMAGE' as sampled images only ever use read image type --- .../sycl/ext/oneapi/bindless_images.hpp | 91 ++++++++----------- 1 file changed, 40 insertions(+), 51 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index 4e222f0b71a3a..c7c9d880a8536 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -55,19 +55,6 @@ struct sampled_image_handle { raw_image_handle_type raw_handle; }; -// Image types used for generating SPIR-V -#ifdef __SYCL_DEVICE_ONLY__ -template -using OCLImageTyRead = - typename sycl::detail::opencl_image_type::type; - -template -using OCLImageTyWrite = - typename sycl::detail::opencl_image_type::type; -#endif - /** * @brief Allocate image memory based on image_descriptor * @@ -806,6 +793,17 @@ template constexpr bool is_recognized_standard_type() { #ifdef __SYCL_DEVICE_ONLY__ +// Image types used for generating SPIR-V +template +using OCLImageTyRead = + typename sycl::detail::opencl_image_type::type; + +template +using OCLImageTyWrite = + typename sycl::detail::opencl_image_type::type; + // Macros are required because it is not legal for a function to return // a variable of type 'opencl_image_type'. #if defined(__NVPTX__) @@ -818,14 +816,14 @@ template constexpr bool is_recognized_standard_type() { #endif #if defined(__NVPTX__) -#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) raw_handle +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle #elif defined(__SPIR__) -#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) \ +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) \ __spirv_ConvertHandleToSampledImageINTEL< \ - typename sycl::detail::sampled_opencl_image_type::type>( \ - raw_handle) + typename sycl::detail::sampled_opencl_image_type< \ + detail::OCLImageTyRead>::type>(raw_handle) #else -#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, ImageType) raw_handle +#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle #endif #if defined(__NVPTX__) @@ -878,7 +876,7 @@ DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]], return FETCH_UNSAMPLED_IMAGE( DataT, CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords); } else { @@ -890,7 +888,7 @@ DataT fetch_image(const unsampled_image_handle &imageHandle [[maybe_unused]], return sycl::bit_cast(FETCH_UNSAMPLED_IMAGE( HintT, CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords)); } #else @@ -965,13 +963,11 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { return __invoke__SampledImageFetch( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords); } else { return sycl::bit_cast(__invoke__SampledImageFetch( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords)); } #else @@ -1017,13 +1013,11 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { return __invoke__ImageRead( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords); } else { return sycl::bit_cast(__invoke__ImageRead( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords)); } #else @@ -1094,8 +1088,7 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { return __invoke__ImageReadLod( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords, level); } else { static_assert(sizeof(HintT) == sizeof(DataT), @@ -1104,8 +1097,7 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], static_assert(detail::is_recognized_standard_type(), "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageReadLod( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords, level)); } #else @@ -1143,8 +1135,7 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { return __invoke__ImageReadGrad( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords, dX, dY); } else { static_assert(sizeof(HintT) == sizeof(DataT), @@ -1153,8 +1144,7 @@ DataT sample_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], static_assert(detail::is_recognized_standard_type(), "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageReadGrad( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), coords, dX, dY)); } #else @@ -1302,7 +1292,7 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle if constexpr (detail::is_recognized_standard_type()) { return __invoke__ImageArrayFetch( CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords, arrayLayer); } else { static_assert(sizeof(HintT) == sizeof(DataT), @@ -1312,7 +1302,7 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageArrayFetch( CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords, arrayLayer)); } #else @@ -1362,9 +1352,7 @@ DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { return __invoke__ImageReadCubemap( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), - dirVec); + 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 " @@ -1372,8 +1360,7 @@ DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]], static_assert(detail::is_recognized_standard_type(), "HintT must always be a recognized standard type"); return sycl::bit_cast(__invoke__ImageReadCubemap( - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, NDims), dirVec)); } #else @@ -1403,15 +1390,17 @@ void write_image(unsampled_image_handle imageHandle [[maybe_unused]], #ifdef __SYCL_DEVICE_ONLY__ if constexpr (detail::is_recognized_standard_type()) { - __invoke__ImageWrite(CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyWrite), - coords, color); + __invoke__ImageWrite( + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + detail::OCLImageTyWrite), + coords, color); } else { // Convert DataT to a supported backend write type when user-defined type is // passed - __invoke__ImageWrite(CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyWrite), - coords, detail::convert_color(color)); + __invoke__ImageWrite( + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + detail::OCLImageTyWrite), + coords, detail::convert_color(color)); } #else assert(false); // Bindless images not yet implemented on host @@ -1444,14 +1433,14 @@ void write_image_array(unsampled_image_handle imageHandle [[maybe_unused]], if constexpr (detail::is_recognized_standard_type()) { __invoke__ImageArrayWrite( CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords, arrayLayer, color); } else { // Convert DataT to a supported backend write type when user-defined type is // passed __invoke__ImageArrayWrite( CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, - OCLImageTyRead), + detail::OCLImageTyRead), coords, arrayLayer, detail::convert_color(color)); } #else From 9fb5c7dab1af6cb5321f3318a31c8b3acc937f25 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Tue, 14 May 2024 17:46:04 +0100 Subject: [PATCH 3/7] merge macro branches --- sycl/include/sycl/ext/oneapi/bindless_images.hpp | 13 +++---------- 1 file changed, 3 insertions(+), 10 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index c7c9d880a8536..b826a4d25a9ab 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -806,18 +806,14 @@ using OCLImageTyWrite = // Macros are required because it is not legal for a function to return // a variable of type 'opencl_image_type'. -#if defined(__NVPTX__) -#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle -#elif defined(__SPIR__) +#if defined(__SPIR__) #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) \ __spirv_ConvertHandleToImageINTEL(raw_handle) #else #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle #endif -#if defined(__NVPTX__) -#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle -#elif defined(__SPIR__) +#if defined(__SPIR__) #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) \ __spirv_ConvertHandleToSampledImageINTEL< \ typename sycl::detail::sampled_opencl_image_type< \ @@ -826,10 +822,7 @@ using OCLImageTyWrite = #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle #endif -#if defined(__NVPTX__) -#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ - __invoke__ImageFetch(raw_handle, coords) -#elif defined(__SPIR__) +#if defined(__SPIR__) #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ __invoke__ImageRead(raw_handle, coords) #else From f674ba9e5eb2769e465d77ee6d3b7883c2343b96 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Tue, 14 May 2024 17:47:29 +0100 Subject: [PATCH 4/7] Template '__spirv_ConvertHandleTo*' builtin handle type --- sycl/include/CL/__spirv/spirv_ops.hpp | 16 +++++++--------- 1 file changed, 7 insertions(+), 9 deletions(-) diff --git a/sycl/include/CL/__spirv/spirv_ops.hpp b/sycl/include/CL/__spirv/spirv_ops.hpp index 501e4c3059fd4..4978a347f2f85 100644 --- a/sycl/include/CL/__spirv/spirv_ops.hpp +++ b/sycl/include/CL/__spirv/spirv_ops.hpp @@ -230,17 +230,15 @@ template extern __DPCPP_SYCL_EXTERNAL TempRetT __spirv_ImageSampleCubemap(SampledType, TempArgT); -template -extern __DPCPP_SYCL_EXTERNAL RetT -__spirv_ConvertHandleToImageINTEL(unsigned long); +template +extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ConvertHandleToImageINTEL(HandleT); -template -extern __DPCPP_SYCL_EXTERNAL RetT -__spirv_ConvertHandleToSamplerINTEL(unsigned long); +template +extern __DPCPP_SYCL_EXTERNAL RetT __spirv_ConvertHandleToSamplerINTEL(HandleT); -template -extern __DPCPP_SYCL_EXTERNAL RetT -__spirv_ConvertHandleToSampledImageINTEL(unsigned long); +template +extern __DPCPP_SYCL_EXTERNAL + RetT __spirv_ConvertHandleToSampledImageINTEL(HandleT); #define __SYCL_OpGroupAsyncCopyGlobalToLocal __spirv_GroupAsyncCopy #define __SYCL_OpGroupAsyncCopyLocalToGlobal __spirv_GroupAsyncCopy From b5646717a910f96030a8ab7ae7e7b43a407e73df Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Wed, 15 May 2024 13:39:30 +0100 Subject: [PATCH 5/7] Add 'SPV_INTEL_bindless_images' to 'sycl-spirv-ext.c' --- clang/test/Driver/sycl-spirv-ext.c | 1 + 1 file changed, 1 insertion(+) diff --git a/clang/test/Driver/sycl-spirv-ext.c b/clang/test/Driver/sycl-spirv-ext.c index 121cdf3861121..4a4544b633991 100644 --- 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 From b7f82d71e39a4615d5521237a15526c022e33bd8 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Tue, 21 May 2024 12:50:18 +0100 Subject: [PATCH 6/7] Fix 'bindless_images_SPIRV_inst.cpp' not working on windows due to single long being 32-bit in length instead of 64-bit --- sycl/test/extensions/bindless_images_SPIRV_inst.cpp | 9 ++++----- 1 file changed, 4 insertions(+), 5 deletions(-) diff --git a/sycl/test/extensions/bindless_images_SPIRV_inst.cpp b/sycl/test/extensions/bindless_images_SPIRV_inst.cpp index 3ade6497549e5..dd91c773ea94b 100644 --- a/sycl/test/extensions/bindless_images_SPIRV_inst.cpp +++ b/sycl/test/extensions/bindless_images_SPIRV_inst.cpp @@ -60,8 +60,7 @@ using OCLSampledImageTy = typename sycl::detail::sampled_opencl_image_type< OCLImageTyRead>::type; #endif -template -ReturnT handleToImage(const unsigned long &imageHandle) { +template ReturnT handleToImage(const uint64_t &imageHandle) { #ifdef __SYCL_DEVICE_ONLY__ #if defined(__SPIR__) return __spirv_ConvertHandleToImageINTEL(imageHandle); @@ -72,7 +71,7 @@ ReturnT handleToImage(const unsigned long &imageHandle) { } template -ReturnT handleToSampler(const unsigned long &samplerHandle) { +ReturnT handleToSampler(const uint64_t &samplerHandle) { #ifdef __SYCL_DEVICE_ONLY__ #if defined(__SPIR__) return __spirv_ConvertHandleToSamplerINTEL(samplerHandle); @@ -83,10 +82,10 @@ ReturnT handleToSampler(const unsigned long &samplerHandle) { } template -ReturnT handleToSampledImage(const unsigned long &imageHandle) { +ReturnT handleToSampledImage(const uint64_t &sampledImageHandle) { #ifdef __SYCL_DEVICE_ONLY__ #if defined(__SPIR__) - return __spirv_ConvertHandleToSampledImageINTEL(imageHandle); + return __spirv_ConvertHandleToSampledImageINTEL(sampledImageHandle); #endif #else assert(false); // Bindless images not yet implemented on host. From f8f391fdc42adebe9972f3aee677236ed87fdd09 Mon Sep 17 00:00:00 2001 From: Duncan Brawley Date: Thu, 23 May 2024 13:54:26 +0100 Subject: [PATCH 7/7] Combine '#ifs' into single '#if' --- sycl/include/sycl/ext/oneapi/bindless_images.hpp | 12 ++++-------- 1 file changed, 4 insertions(+), 8 deletions(-) diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index b826a4d25a9ab..78fb912fc04de 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -809,23 +809,19 @@ using OCLImageTyWrite = #if defined(__SPIR__) #define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) \ __spirv_ConvertHandleToImageINTEL(raw_handle) -#else -#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle -#endif -#if defined(__SPIR__) #define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) \ __spirv_ConvertHandleToSampledImageINTEL< \ typename sycl::detail::sampled_opencl_image_type< \ detail::OCLImageTyRead>::type>(raw_handle) -#else -#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle -#endif -#if defined(__SPIR__) #define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \ __invoke__ImageRead(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(raw_handle, coords) #endif