Skip to content

[SYCL][Bindless] Update spirv read/fetch from sampled image and sampled image array #14493

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Jul 9, 2024
Merged

[SYCL][Bindless] Update spirv read/fetch from sampled image and sampled image array #14493

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
85a5373
[SYCL][Bindless] Update spirv read/fetch from sampled image and sampl…
wenju-he Jul 9, 2024
62f49ad
add missing CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY, fix lit
wenju-he Jul 9, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
71 changes: 60 additions & 11 deletions sycl/include/sycl/ext/oneapi/bindless_images.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -812,6 +812,11 @@ template <int NDims>
using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type<
NDims, sycl::access::mode::write, sycl::access::target::image_array>::type;

template <int NDims>
using OCLSampledImageArrayTyRead =
typename sycl::detail::sampled_opencl_image_type<
detail::OCLImageArrayTyRead<NDims>>::type;

// Macros are required because it is not legal for a function to return
// a variable of type 'opencl_image_type'.
#if defined(__SPIR__)
Expand All @@ -823,34 +828,64 @@ using OCLImageArrayTyWrite = typename sycl::detail::opencl_image_type<
typename sycl::detail::sampled_opencl_image_type< \
detail::OCLImageTyRead<NDims>>::type>(raw_handle)

#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) \
__spirv_ConvertHandleToSampledImageINTEL< \
typename sycl::detail::sampled_opencl_image_type< \
detail::OCLImageArrayTyRead<NDims>>::type>(raw_handle)

#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \
__invoke__ImageRead<DataT>(raw_handle, coords)

#define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \
__invoke__ImageRead<DataT>(raw_handle, coords)
__invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)

#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \
__invoke__ImageReadLod<DataT>(raw_handle, coords, 0.f)

#define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \
__invoke__ImageRead<DataT>(raw_handle, coordsLayer)

#define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \
__invoke__ImageWrite(raw_handle, coordsLayer, color)

#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
coordsLayer) \
__invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)

#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
coordsLayer) \
__invoke__ImageReadLod<DataT>(raw_handle, coordsLayer, 0.f)

#else
#define CONVERT_HANDLE_TO_IMAGE(raw_handle, ImageType) raw_handle

#define CONVERT_HANDLE_TO_SAMPLED_IMAGE(raw_handle, NDims) raw_handle

#define CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(raw_handle, NDims) raw_handle

#define FETCH_UNSAMPLED_IMAGE(DataT, raw_handle, coords) \
__invoke__ImageFetch<DataT>(raw_handle, coords)

#define FETCH_SAMPLED_IMAGE(DataT, raw_handle, coords) \
__invoke__SampledImageFetch<DataT>(raw_handle, coords)

#define SAMPLE_IMAGE_READ(DataT, raw_handle, coords) \
__invoke__ImageRead<DataT>(raw_handle, coords)

#define FETCH_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, coordsLayer) \
__invoke__ImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)

#define WRITE_IMAGE_ARRAY(raw_handle, coords, arrayLayer, coordsLayer, color) \
__invoke__ImageArrayWrite(raw_handle, coords, arrayLayer, color)

#define FETCH_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
coordsLayer) \
__invoke__SampledImageArrayFetch<DataT>(raw_handle, coords, arrayLayer)

#define READ_SAMPLED_IMAGE_ARRAY(DataT, raw_handle, coords, arrayLayer, \
coordsLayer) \
__invoke__ImageArrayRead<DataT>(raw_handle, coords, arrayLayer)

#endif

#endif // __SYCL_DEVICE_ONLY__
Expand Down Expand Up @@ -1030,11 +1065,13 @@ DataT sample_image(const sampled_image_handle &imageHandle [[maybe_unused]],

#ifdef __SYCL_DEVICE_ONLY__
if constexpr (detail::is_recognized_standard_type<DataT>()) {
return __invoke__ImageRead<DataT>(
return SAMPLE_IMAGE_READ(
DataT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
coords);
} else {
return sycl::bit_cast<DataT>(__invoke__ImageRead<HintT>(
return sycl::bit_cast<DataT>(SAMPLE_IMAGE_READ(
HintT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE(imageHandle.raw_handle, coordSize),
coords));
}
Expand Down Expand Up @@ -1410,17 +1447,23 @@ DataT fetch_image_array(const sampled_image_handle &imageHandle
"and 2D images respectively.");

#ifdef __SYCL_DEVICE_ONLY__
sycl::vec<int, coordSize + 1> coordsLayer{coords, arrayLayer};
if constexpr (detail::is_recognized_standard_type<DataT>()) {
return __invoke__SampledImageArrayFetch<DataT>(imageHandle.raw_handle,
coords, arrayLayer);
return FETCH_SAMPLED_IMAGE_ARRAY(DataT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
imageHandle.raw_handle, coordSize),
coords, arrayLayer, coordsLayer);
} else {
static_assert(sizeof(HintT) == sizeof(DataT),
"When trying to fetch a user-defined type, HintT must be of "
"the same size as the user-defined DataT.");
static_assert(detail::is_recognized_standard_type<HintT>(),
"HintT must always be a recognized standard type");
return sycl::bit_cast<DataT>(__invoke__SampledImageArrayFetch<HintT>(
imageHandle.raw_handle, coords, arrayLayer));
return sycl::bit_cast<DataT>(
FETCH_SAMPLED_IMAGE_ARRAY(HintT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
imageHandle.raw_handle, coordSize),
coords, arrayLayer, coordsLayer));
}
#else
assert(false); // Bindless images not yet implemented on host.
Expand Down Expand Up @@ -1454,17 +1497,23 @@ DataT sample_image_array(const sampled_image_handle &imageHandle
"and 2D images respectively.");

#ifdef __SYCL_DEVICE_ONLY__
sycl::vec<float, coordSize + 1> coordsLayer{coords, arrayLayer};
if constexpr (detail::is_recognized_standard_type<DataT>()) {
return __invoke__ImageArrayRead<DataT>(imageHandle.raw_handle, coords,
arrayLayer);
return READ_SAMPLED_IMAGE_ARRAY(DataT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
imageHandle.raw_handle, coordSize),
coords, arrayLayer, coordsLayer);
} else {
static_assert(sizeof(HintT) == sizeof(DataT),
"When trying to fetch a user-defined type, HintT must be of "
"the same size as the user-defined DataT.");
static_assert(detail::is_recognized_standard_type<HintT>(),
"HintT must always be a recognized standard type");
return sycl::bit_cast<DataT>(__invoke__ImageArrayRead<HintT>(
imageHandle.raw_handle, coords, arrayLayer));
return sycl::bit_cast<DataT>(
READ_SAMPLED_IMAGE_ARRAY(HintT,
CONVERT_HANDLE_TO_SAMPLED_IMAGE_ARRAY(
imageHandle.raw_handle, coordSize),
coords, arrayLayer, coordsLayer));
}
#else
assert(false); // Bindless images not yet implemented on host.
Expand Down
4 changes: 2 additions & 2 deletions sycl/test/check_device_code/extensions/bindless_images.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -4,7 +4,7 @@
// RUN: %clangxx -S -emit-llvm -fsycl -fsycl-device-only -fsycl-targets=spir64-unknown-unknown %s -o - | FileCheck %s --check-prefix=CHECK-LLVM

// 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 noundef <4 x float> @_Z30__spirv_ImageSampleExplicitLod
// 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
Expand Down Expand Up @@ -50,7 +50,7 @@

// Read sampled image
// Arguments: Result Type, Result, Image, Coords
// CHECK-SPIRV-NEXT: ImageRead [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}}
// CHECK-SPIRV-NEXT: ImageSampleExplicitLod [[PIXELTYPE]] {{[0-9]+}} [[SAMPIMAGEVAR]] {{[0-9]+}}

// Convert handle to SPIR-V image
// Arguments: Result Type, Result, Handle
Expand Down
Loading