diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index cd86e32bc66e8..c0249d04903dc 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -1036,15 +1036,19 @@ DataT fetch_image(const sampled_image_handle &imageHandle [[maybe_unused]], "HintT must always be a recognized standard type"); #ifdef __SYCL_DEVICE_ONLY__ + // Convert the raw handle to an image and use FETCH_UNSAMPLED_IMAGE since + // fetch_image should not use the sampler if constexpr (detail::is_recognized_standard_type()) { - return FETCH_SAMPLED_IMAGE( + return FETCH_UNSAMPLED_IMAGE( DataT, - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + detail::OCLImageTyRead), coords); } else { - return sycl::bit_cast(FETCH_SAMPLED_IMAGE( + return sycl::bit_cast(FETCH_UNSAMPLED_IMAGE( HintT, - CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize), + CONVERT_HANDLE_TO_IMAGE(imageHandle.raw_handle, + detail::OCLImageTyRead), coords)); } #else diff --git a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp index a3f253c5e04b5..037c18e8d91d6 100644 --- a/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp +++ b/sycl/test-e2e/bindless_images/sampled_fetch/fetch_2D_USM_device.cpp @@ -11,9 +11,10 @@ #include #include -class kernel_sampled_fetch; +namespace { -int main() { +template +static int testSampledImageFetch() { sycl::device dev; sycl::queue q(dev); @@ -23,9 +24,9 @@ int main() { constexpr size_t width = 5; constexpr size_t height = 6; constexpr size_t N = width * height; - std::vector> out(N); - std::vector> expected(N); - std::vector> dataIn(N); + std::vector> out(N); + std::vector> expected(N); + std::vector> dataIn(N); for (int i = 0; i < width; i++) { for (int j = 0; j < height; j++) { auto index = i + (width * j); @@ -43,8 +44,7 @@ int main() { sycl::filtering_mode::linear); // Extension: image descriptor - syclexp::image_descriptor desc({width, height}, 4, - sycl::image_channel_type::unsigned_int16); + syclexp::image_descriptor desc({width, height}, 4, ChanType); size_t pitch = 0; // Extension: returns the device pointer to USM allocated pitched memory @@ -65,21 +65,20 @@ int main() { sycl::buffer buf(out.data(), sycl::range{height, width}); q.submit([&](sycl::handler &cgh) { - auto outAcc = buf.get_access( + auto outAcc = buf.template get_access( cgh, sycl::range<2>{height, width}); - cgh.parallel_for( - sycl::nd_range<2>{{width, height}, {width, height}}, - [=](sycl::nd_item<2> it) { - size_t dim0 = it.get_local_id(0); - size_t dim1 = it.get_local_id(1); + cgh.parallel_for(sycl::nd_range<2>{{width, height}, {width, height}}, + [=](sycl::nd_item<2> it) { + size_t dim0 = it.get_local_id(0); + size_t dim1 = it.get_local_id(1); - // Extension: fetch data from sampled image handle - auto px1 = syclexp::fetch_image>( - imgHandle, sycl::int2(dim0, dim1)); + // Extension: fetch data from sampled image handle + auto px1 = syclexp::fetch_image>( + imgHandle, sycl::int2(dim0, dim1)); - outAcc[sycl::id<2>{dim1, dim0}] = px1; - }); + outAcc[sycl::id<2>{dim1, dim0}] = px1; + }); }); q.wait_and_throw(); @@ -121,3 +120,23 @@ int main() { std::cout << "Test failed!" << std::endl; return 3; } + +} // namespace + +int main() { + if (int err = + testSampledImageFetch()) { + return err; + } + if (int err = + testSampledImageFetch()) { + return err; + } + if (int err = + testSampledImageFetch()) { + return err; + } + return 0; +}