diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index f7d1b76b19f2a..072b8bfc191a6 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -1506,13 +1506,13 @@ sampling depends on the sampler attributes passed upon creation of the cubemap. ```c++ // Unsampled cubemap read -template +template DataT fetch_cubemap(const unsampled_image_handle &ImageHandle, const int2 &Coords, const int Face); // Sampled cubemap read -template +template DataT sample_cubemap(const sampled_image_handle &ImageHandle, const float3 &Vec); @@ -2684,4 +2684,6 @@ These features still need to be handled: parameter. |5.7|2024-04-09| - Allow fetching of sampled image data through the `fetch_image` API. +|5.8|2024-05-09| - Add missing cubemap `HintT` template parameter to + `fetch_cubemap` and `sample_cubemap`. |====================== diff --git a/sycl/include/sycl/ext/oneapi/bindless_images.hpp b/sycl/include/sycl/ext/oneapi/bindless_images.hpp index e11903a085ab5..51d3dbb88b80c 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images.hpp @@ -1244,22 +1244,30 @@ DataT fetch_image_array(const unsampled_image_handle &imageHandle * @brief Fetch data from an unsampled cubemap image using its handle * * @tparam DataT The return type + * @tparam HintT A hint type that can be used to select for a specialized + * backend intrinsic when a user-defined type is passed as `DataT`. + * HintT should be a `sycl::vec` type, `sycl::half` type, or POD type. + * HintT must also have the same size as DataT. * * @param imageHandle The image handle * @param coords The coordinates at which to fetch image data (int2 only) * @param face The cubemap face at which to fetch * @return Image data */ -template +template DataT fetch_cubemap(const unsampled_image_handle &imageHandle, const int2 &coords, const unsigned int face) { - return fetch_image_array(imageHandle, coords, face); + return fetch_image_array(imageHandle, coords, face); } /** * @brief Sample a cubemap image using its handle * * @tparam DataT The return type + * @tparam HintT A hint type that can be used to select for a specialized + * backend intrinsic when a user-defined type is passed as `DataT`. + * HintT should be a `sycl::vec` type, `sycl::half` type, or POD type. + * HintT must also have the same size as DataT. * * @param imageHandle The image handle * @param dirVec The direction vector at which to sample image data (float3 @@ -1280,7 +1288,7 @@ DataT sample_cubemap(const sampled_image_handle &imageHandle [[maybe_unused]], "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__ImageReadCubemap( + return sycl::bit_cast(__invoke__ImageReadCubemap( imageHandle.raw_handle, dirVec)); } #else diff --git a/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp b/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp index d2ba5da27d5f3..d47df77000112 100644 --- a/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp +++ b/sycl/test-e2e/bindless_images/cubemap/cubemap_sampled.cpp @@ -4,6 +4,7 @@ // RUN: %{build} -o %t.out // RUN: %{run} %t.out +#include "../user_types/user_types_common.hpp" #include #include @@ -27,12 +28,12 @@ int main() { size_t N = width * height; std::vector out(N); std::vector expected(N); - std::vector dataIn1(N * 6); + std::vector dataIn(N * 6); for (int i = 0; i < width; i++) { for (int j = 0; j < height; j++) { for (int k = 0; k < 6; k++) { - dataIn1[i + width * (j + height * k)] = {i + width * (j + height * k), - 0, 0, 0}; + dataIn[i + width * (j + height * k)] = {i + width * (j + height * k), 0, + 0, 0}; } } } @@ -63,7 +64,7 @@ int main() { // Extension: copy over data to device (handler variant). q.submit([&](sycl::handler &cgh) { - cgh.ext_oneapi_copy(dataIn1.data(), imgMem.get_handle(), desc); + cgh.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc); }); q.wait_and_throw(); @@ -95,11 +96,13 @@ int main() { float fdim2 = (((float(dim1) / (float)height) * 1.98) - 0.99) + (1.f / (float)height); + my_float4 myPixel{}; + // Extension: sample cubemap data from handle. - sycl::float4 px = syclexp::sample_cubemap( + myPixel = syclexp::sample_cubemap( imgHandle, sycl::float3(fdim0, fdim1, fdim2)); - outAcc[sycl::id<2>{dim0, dim1}] = px[0]; + outAcc[sycl::id<2>{dim0, dim1}] = myPixel.x; }); }); q.wait_and_throw();