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 ba3eac3579490..4c3ade5cf1d29 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -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. @@ -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. |====================== diff --git a/sycl/include/sycl/sampler.hpp b/sycl/include/sycl/sampler.hpp index 547fc6a4cfeb7..ed66446f65e62 100644 --- a/sycl/include/sycl/sampler.hpp +++ b/sycl/include/sycl/sampler.hpp @@ -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 { diff --git a/sycl/test-e2e/bindless_images/helpers/sampling.hpp b/sycl/test-e2e/bindless_images/helpers/sampling.hpp index 51b1a58088094..dd562a0282483 100644 --- a/sycl/test-e2e/bindless_images/helpers/sampling.hpp +++ b/sycl/test-e2e/bindless_images/helpers/sampling.hpp @@ -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"; @@ -581,7 +581,7 @@ read(sycl::range<2> globalSize, sycl::vec 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(coords, globalSize, inputImage); } @@ -623,7 +623,7 @@ read(sycl::range<2> globalSize, sycl::vec 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(coords, globalSize, inputImage); } if (SampAddrMode == sycl::addressing_mode::clamp_to_edge) { diff --git a/sycl/test-e2e/bindless_images/read_sampled.cpp b/sycl/test-e2e/bindless_images/read_sampled.cpp index b4eed686605bd..995fc7a1934d5 100644 --- a/sycl/test-e2e/bindless_images/read_sampled.cpp +++ b/sycl/test-e2e/bindless_images/read_sampled.cpp @@ -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}; @@ -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};