Skip to content

Commit

Permalink
[Bindless][Exp][NFC] Deprecate read_image for more descriptive naming
Browse files Browse the repository at this point in the history 
- The `read_image` and `read_mipmap` APIs have been deprecated
- They are replaced with the more descriptive `fetch_image`,
  `sample_image`, and `sample_mipmap` (for unsampled reads, sampled
  reads, and mipmap sampled reads, respectively).
- This change is made in preperation for future functionality of
  fetching data from sampled images.
- The reason behind this change is to avoid determining the
  underlying image read operation based on the coordinate type
  passed, and instead making it more transparent for the user which
  operation is performed based on the name of the function.
- The extension document, bindless images headers, and all bindless
  images tests have all been updated.
- The specification revision history has been updated to include a
  missed changelog entry for PR intel#12581
  • Loading branch information
@przemektmalon
przemektmalon committed Feb 19, 2024
1 parent 1a98c4c commit 99a7430
Show file tree
Hide file tree
Showing 29 changed files with 272 additions and 193 deletions.
86 changes: 47 additions & 39 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc
View file
Original file line number Diff line number Diff line change
Expand Up @@ -986,39 +986,41 @@ listed above caused the failure.
namespace sycl::ext::oneapi::experimental {

template <typename DataT, typename HintT = DataT, typename CoordT>
DataT read_image(const unsampled_image_handle &ImageHandle,
const CoordT &Coords);
DataT fetch_image(const unsampled_image_handle &ImageHandle,
const CoordT &Coords);
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT read_image(const sampled_image_handle &ImageHandle,
const CoordT &Coords);
DataT sample_image(const sampled_image_handle &ImageHandle,
const CoordT &Coords);

template <typename DataT, typename CoordT>
void write_image(unsampled_image_handle &ImageHandle,
const CoordT &Coords, const DataT &Color);
}
```

Inside a kernel, it's possible to read an image via `read_image`, passing
the image handle. For the form that takes `unsampled_image_handle`, image data
will be fetched exactly as is in device memory. For the form that takes a
`sampled_image_handle`, the image will be sampled according to the
Inside a kernel, it's possible to retrieve data from an image via `fetch_image`
or `sample_image`, passing the appropirate image handle. The `fetch_image` API
is only applicable to unsampled images, and the data will be fetched exactly as
is in device memory. The `sample_image` API is only applicable to sampled
images, the image data will be sampled according to the
`bindless_image_sampler` that was passed to the image upon construction.

The user is required to pass a `DataT` template parameter, which specifies the
return type of the `read_image` function. If `DataT` is not a recognized
standard type, as defined in <<recognized_standard_types>>, and instead a
user-defined type, the user must provide a `HintT` template parameter to the
`read_image` function, to allow the backend to select the correct device
intrinsic to fetch or sample their data.
return type of the `fetch_image` and `sample_image` functions. If `DataT` is
not a recognized standard type, as defined in <<recognized_standard_types>>,
and instead a user-defined type, the user must provide a `HintT` template
parameter to the `fetch_image` and `sample_image` functions, to allow the
backend to select the correct device intrinsic to fetch or sample their data.

`HintT` must be one of the the <<recognized_standard_types>>, and must be the
same size as `DataT`.
If `DataT` is a recognized standard type, and `HintT` is also passed, `HintT`
will be ignored.

When reading a texture backed by a normalized integer channel type, either
`DataT` must be a 32-bit or 16-bit floating point value, a `sycl::vec` of
32-bit or 16-bit floating point values, or, in the case `DataT` is not one of
the above, then `HintT` must be one of the above, and be of the same size as
When fetching or sampling an image backed by a normalized integer channel type,
either `DataT` must be a 32-bit or 16-bit floating point value, a `sycl::vec`
of 32-bit or 16-bit floating point values, or, in the case `DataT` is not one
of the above, then `HintT` must be one of the above, and be of the same size as
`DataT`.

It's possible to write to an unsampled image via `write_image` passing the
Expand All @@ -1029,8 +1031,8 @@ of the <<recognized_standard_types>>.

Sampled images cannot be written to using `write_image`.

For reading and writing of unsampled images, coordinates are specified by `int`,
`sycl::vec<int, 2>`, and `sycl::vec<int, 3>` for 1D, 2D, and 3D images,
For fetching and writing of unsampled images, coordinates are specified by
`int`, `sycl::vec<int, 2>`, and `sycl::vec<int, 3>` for 1D, 2D, and 3D images,
respectively.

Sampled image reads take `float`, `sycl::vec<float, 2>`, and
Expand All @@ -1046,8 +1048,8 @@ kernel must be submitted for the written data to be accessible.

[NOTE]
====
Attempting to read an image with `read_mipmap` or any other defined read
function will result in undefined behaviour.
Attempting to sample a standard sampled image with `sample_mipmap` or any other
defined sampling function will result in undefined behaviour.
====

=== Recognized standard types [[recognized_standard_types]]
Expand All @@ -1057,7 +1059,8 @@ standard types.

* All POD types (`char`, `short`, `int`, `float`, etc.) excluding `double`
* `sycl::half`
* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`, `2`, or `3`
* Variants of `sycl::vec<T, N>` where `T` is one of the above, and `N` is `1`,
`2`, or `3`

Any other types are classified as user-defined types.

Expand Down Expand Up @@ -1168,26 +1171,26 @@ level of a given top-level descriptor.

=== Reading a mipmap

Inside the kernel, it's possible to read a mipmap via `read_mipmap`, passing the
`sampled_image_handle`, the coordinates, and either the level or anisotropic
gradient values.
Inside the kernel, it's possible to sample a mipmap via `sample_mipmap`,
passing the `sampled_image_handle`, the coordinates, and either the level or
anisotropic gradient values.

The method of sampling a mipmap is different based on which `read_mipmap`
The method of sampling a mipmap is different based on which `sample_mipmap`
function is used, and the sampler attributes passed upon creation of the
mipmap.

```c++
// Nearest/linear filtering between mip levels
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT read_mipmap(const sampled_image_handle &ImageHandle,
const CoordT &Coords,
const float Level);
DataT sample_mipmap(const sampled_image_handle &ImageHandle,
const CoordT &Coords,
const float Level);

// Anisotropic filtering
template <typename DataT, typename HintT = DataT, typename CoordT>
DataT read_mipmap(const sampled_image_handle &ImageHandle,
const CoordT &Coords,
const CoordT &Dx, const CoordT &Dy);
DataT sample_mipmap(const sampled_image_handle &ImageHandle,
const CoordT &Coords,
const CoordT &Dx, const CoordT &Dy);
```

Reading a mipmap follows the same restrictions on what coordinate types may be
Expand All @@ -1199,8 +1202,8 @@ the restrictions as laid out in <<reading_writing_inside_kernel>>.

[NOTE]
====
Attempting to read a mipmap with `read_image` or any other defined read function
will result in undefined behaviour.
Attempting to sample a mipmap with `sample_image` or any other defined sample
function will result in undefined behaviour.
====

== Interoperability
Expand Down Expand Up @@ -1544,7 +1547,7 @@ try {

cgh.parallel_for(width, [=](sycl::id<1> id) {
// Extension: read image data from handle
float pixel = sycl::ext::oneapi::experimental::read_image<float>(
float pixel = sycl::ext::oneapi::experimental::fetch_image<float>(
imgIn, int(id[0]));

// Extension: write to image data using handle
Expand Down Expand Up @@ -1646,7 +1649,7 @@ try {
float sum = 0;
for (int i = 0; i < numImages; i++) {
// Extension: read image data from handle
sum += (sycl::ext::oneapi::experimental::read_image<float>(
sum += (sycl::ext::oneapi::experimental::fetch_image<float>(
imgHandleAcc[i], sycl::vec<int, 2>(dim0, dim1)));
}
outAcc[sycl::id{dim1, dim0}] = sum;
Expand Down Expand Up @@ -1736,9 +1739,9 @@ try {
float x = (static_cast<float>(id[0]) + 0.5f) / static_cast<float>(width);
// Read mipmap level 0 with anisotropic filtering
// and level 1 with level filtering
float px1 = sycl::ext::oneapi::experimental::read_mipmap<float>(
float px1 = sycl::ext::oneapi::experimental::sample_mipmap<float>(
mipHandle, x, 0.0f, 0.0f);
float px2 = sycl::ext::oneapi::experimental::read_mipmap<float>(
float px2 = sycl::ext::oneapi::experimental::sample_mipmap<float>(
mipHandle, x, 1.0f);

sum = px1 + px2;
Expand Down Expand Up @@ -1874,7 +1877,7 @@ try {

// Extension: read image data from handle to imported image
uint32_t pixel =
sycl::ext::oneapi::experimental::read_image<uint32_t>(
sycl::ext::oneapi::experimental::fetch_image<uint32_t>(
img_input, sycl::vec<int, 2>(dim0, dim1));

// Modify the data before writing back
Expand Down Expand Up @@ -2076,4 +2079,9 @@ These features still need to be handled:
user-defined type.
|5.1|2023-12-06| - Added unique addressing modes per dimension to the
`bindless_image_sampler`
|5.2|2024-02-14| - Image read and write functions now accept 3-component
coordinates for 3D reads, instead of 4-component coordinates.
|5.3|2024-02-16| - Replace `read_image` and `read_mipmap` APIs in favor of more
descriptive naming, with `fetch_image`, `sample_image`, and
`sample_mipmap`.
|======================
Loading

0 comments on commit 99a7430

Please sign in to comment.