From 761dd9cd9d4af4b26085a28be9d0c66c251957a9 Mon Sep 17 00:00:00 2001 From: Jinsong Ji Date: Tue, 14 Nov 2023 14:19:02 -0800 Subject: [PATCH 1/2] [SYCL][Bindless] Enable SPIRV path for bindless_images This is 1st patch to enable SPIRV path for bindless_images so that we can start generating SPIRV for testing purpose. --- .../sycl/ext/oneapi/bindless_images.hpp | 26 +--------- sycl/test/extensions/bindless_images.cpp | 51 +++++++++++++++++++ 2 files changed, 52 insertions(+), 25 deletions(-) create mode 100644 sycl/test/extensions/bindless_images.cpp diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index dba5739f26d2..30fe37f05ced 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -761,11 +761,7 @@ DataT read_image(const unsampled_image_handle &imageHandle [[maybe_unused]], "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageRead(imageHandle.raw_handle, coords); -#else - // TODO: add SPIRV part for unsampled image read -#endif #else assert(false); // Bindless images not yet implemented on host #endif @@ -797,11 +793,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageRead(imageHandle.raw_handle, coords); -#else - // TODO: add SPIRV part for sampled image read -#endif #else assert(false); // Bindless images not yet implemented on host. #endif @@ -829,11 +821,7 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageReadLod(imageHandle.raw_handle, coords, level); -#else - // TODO: add SPIRV for mipmap level read -#endif #else assert(false); // Bindless images not yet implemented on host #endif @@ -863,11 +851,7 @@ DataT read_mipmap(const sampled_image_handle &imageHandle [[maybe_unused]], "components for 1D, 2D, and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageReadGrad(imageHandle.raw_handle, coords, dX, dY); -#else - // TODO: add SPIRV part for mipmap grad read -#endif #else assert(false); // Bindless images not yet implemented on host #endif @@ -898,11 +882,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], "for 1D, 2D and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageReadLod(imageHandle.raw_handle, coords, level); -#else - // TODO: add SPIRV for mipmap level read -#endif #else assert(false); // Bindless images not yet implemented on host #endif @@ -935,11 +915,7 @@ DataT read_image(const sampled_image_handle &imageHandle [[maybe_unused]], "components for 1D, 2D, and 3D images, respectively."); #ifdef __SYCL_DEVICE_ONLY__ -#if defined(__NVPTX__) return __invoke__ImageReadGrad(imageHandle.raw_handle, coords, dX, dY); -#else - // TODO: add SPIRV part for mipmap grad read -#endif #else assert(false); // Bindless images not yet implemented on host #endif @@ -969,7 +945,7 @@ void write_image(const unsampled_image_handle &imageHandle [[maybe_unused]], __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, detail::convert_color_nvptx(color)); #else - // TODO: add SPIRV part for unsampled image write + __invoke__ImageWrite((uint64_t)imageHandle.raw_handle, coords, color); #endif #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 new file mode 100644 index 000000000000..924f7ab80218 --- /dev/null +++ b/sycl/test/extensions/bindless_images.cpp @@ -0,0 +1,51 @@ +// RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s + +#include +#include + +// CHECK: spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI10image_readEE +// CHECK: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4_fmiET_T0_T1_ +using namespace sycl::ext::oneapi::experimental; +class image_read; +int main() { + + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + constexpr size_t width = 512; + std::vector out(width); + std::vector dataIn1(width); + for (int i = 0; i < width; i++) { + dataIn1[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); + + q.ext_oneapi_copy(dataIn1.data(), imgMem0.get_handle(), desc); + q.wait_and_throw(); + + sycl::buffer buf((float *)out.data(), width); + q.submit([&](sycl::handler &cgh) { + auto outAcc = buf.get_access(cgh, width); + + cgh.parallel_for(width, [=](sycl::id<1> id) { + sycl::float4 px1 = read_image(imgHandle1, int(id[0])); + outAcc[id] = px1[0]; + }); + }); + + q.wait_and_throw(); + destroy_image_handle(imgHandle1, dev, ctxt); + } + + for (int i = 0; i < width; i++) { + std::cout << "Actual: " << out[i] << std::endl; + } + return 0; +} From 328748fde1e328d6bb9d3759b4a7bb8d95386e2b Mon Sep 17 00:00:00 2001 From: Jinsong Ji Date: Wed, 15 Nov 2023 13:00:26 -0800 Subject: [PATCH 2/2] Allow both ul and ull in function call On windows we are getting _Z17__spirv_ImageReadIDv4_fyiET_T0_T1_ instead. --- sycl/test/extensions/bindless_images.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test/extensions/bindless_images.cpp b/sycl/test/extensions/bindless_images.cpp index 924f7ab80218..b73449b5428c 100644 --- a/sycl/test/extensions/bindless_images.cpp +++ b/sycl/test/extensions/bindless_images.cpp @@ -4,7 +4,7 @@ #include // CHECK: spir_kernel void @_ZTSN4sycl3_V16detail19__pf_kernel_wrapperI10image_readEE -// CHECK: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4_fmiET_T0_T1_ +// CHECK: tail call spir_func noundef <4 x float> @_Z17__spirv_ImageReadIDv4 using namespace sycl::ext::oneapi::experimental; class image_read; int main() {