diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index dba5739f26d20..30fe37f05ceda 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 0000000000000..b73449b5428cc --- /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 +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; +}