Skip to content

Commit

Permalink
[SYCL][Bindless][Doc] Add 'addressing_mode::ext_oneapi_clamp_to_borde…
Browse files Browse the repository at this point in the history
…r' to replace 'addressing_mode::clamp' (#15524)

'addressing_mode::clamp' is not descriptive enough and neither does it
align with what is commonly used in other graphics APIs. It is proposed
to use 'addressing_mode::ext_oneapi_clamp_to_border' instead.
  • Loading branch information
DBDuncan authored Sep 30, 2024
1 parent 314fcb4 commit c4d20a7
Show file tree
Hide file tree
Showing 4 changed files with 19 additions and 7 deletions.
Original file line number Diff line number Diff line change
Expand Up @@ -713,6 +713,13 @@ semantics. The value for the addressing mode, `addressing_mode::none`,
represents the backend's default addressing mode. On CUDA this is `Wrap`, i.e.
`addressing_mode::repeat`.

We propose renaming `addressing_mode::clamp` to
`addressing_mode::clamp_to_border`. This name aligns better with terms used in
other APIs and is more descriptive as to what the addressing mode does. Note
that in this extension, the addressing mode will be named
`addressing_mode::ext_oneapi_clamp_to_border` as to comply with extension naming
guidelines.

`addressing[3]` defines the addressing mode per texture dimension. A
`bindless_image_sampler` can be constructed with a singular
`sycl::addressing_mode`, where this parameter will define all dimensions.
Expand Down Expand Up @@ -2894,4 +2901,6 @@ These features still need to be handled:
|6 |2024-08-05 | - Collated all changes since revision 5.
- Bumped SYCL_EXT_ONEAPI_BINDLESS_IMAGES to number 6.
|6.1|2024-09-09| - Update for image-array sub-region copy support.
|6.2|2024-09-26| - Added addressing mode `ext_oneapi_clamp_to_border` value,
equivalent to `clamp`, to match with external APIs.
|======================
3 changes: 2 additions & 1 deletion sycl/include/sycl/sampler.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -25,7 +25,8 @@ enum class addressing_mode : unsigned int {
repeat = 0x1133, // Value of CL_ADDRESS_REPEAT
clamp_to_edge = 0x1131, // Value of CL_ADDRESS_CLAMP_TO_EDGE
clamp = 0x1132, // Value of CL_ADDRESS_CLAMP
none = 0x1130 // Value of CL_ADDRESS_NONE
ext_oneapi_clamp_to_border = clamp,
none = 0x1130 // Value of CL_ADDRESS_NONE
};

enum class filtering_mode : unsigned int {
Expand Down
8 changes: 4 additions & 4 deletions sycl/test-e2e/bindless_images/helpers/sampling.hpp
Original file line number Diff line number Diff line change
Expand Up @@ -136,8 +136,8 @@ printTestInfo(sycl::ext::oneapi::experimental::bindless_image_sampler &samp,
case sycl::addressing_mode::clamp_to_edge:
std::cout << "clamp_to_edge\n";
break;
case sycl::addressing_mode::clamp:
std::cout << "clamp\n";
case sycl::addressing_mode::ext_oneapi_clamp_to_border:
std::cout << "ext_oneapi_clamp_to_border\n";
break;
case sycl::addressing_mode::none:
std::cout << "none\n";
Expand Down Expand Up @@ -581,7 +581,7 @@ read(sycl::range<2> globalSize, sycl::vec<float, 2> coords, float offset,
if (SampFiltMode == sycl::filtering_mode::nearest) {

sycl::addressing_mode SampAddrMode = samp.addressing[0];
if (SampAddrMode == sycl::addressing_mode::clamp) {
if (SampAddrMode == sycl::addressing_mode::ext_oneapi_clamp_to_border) {
return clampNearest<VecType>(coords, globalSize, inputImage);
}

Expand Down Expand Up @@ -623,7 +623,7 @@ read(sycl::range<2> globalSize, sycl::vec<float, 2> coords, float offset,

} else { // linear
sycl::addressing_mode SampAddrMode = samp.addressing[0];
if (SampAddrMode == sycl::addressing_mode::clamp) {
if (SampAddrMode == sycl::addressing_mode::ext_oneapi_clamp_to_border) {
return clampLinear<DType, NChannels>(coords, globalSize, inputImage);
}
if (SampAddrMode == sycl::addressing_mode::clamp_to_edge) {
Expand Down
6 changes: 4 additions & 2 deletions sycl/test-e2e/bindless_images/read_sampled.cpp
Original file line number Diff line number Diff line change
Expand Up @@ -284,7 +284,8 @@ bool runTests(sycl::range<1> dims, sycl::range<1> localSize, float offset,
// normalized and unnormalized coords.
sycl::addressing_mode addrModes[4] = {
sycl::addressing_mode::repeat, sycl::addressing_mode::mirrored_repeat,
sycl::addressing_mode::clamp_to_edge, sycl::addressing_mode::clamp};
sycl::addressing_mode::clamp_to_edge,
sycl::addressing_mode::ext_oneapi_clamp_to_border};

sycl::filtering_mode filtModes[2] = {sycl::filtering_mode::nearest,
sycl::filtering_mode::linear};
Expand Down Expand Up @@ -440,7 +441,8 @@ bool runTests(sycl::range<2> dims, sycl::range<2> localSize, float offset,
// normalized and unnormalized coords.
sycl::addressing_mode addrModes[4] = {
sycl::addressing_mode::repeat, sycl::addressing_mode::mirrored_repeat,
sycl::addressing_mode::clamp_to_edge, sycl::addressing_mode::clamp};
sycl::addressing_mode::clamp_to_edge,
sycl::addressing_mode::ext_oneapi_clamp_to_border};

sycl::filtering_mode filtModes[2] = {sycl::filtering_mode::nearest,
sycl::filtering_mode::linear};
Expand Down

0 comments on commit c4d20a7

Please sign in to comment.