From 5309b942c5c293b73a6cd5ffb1d32c5ab99121d6 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Wed, 28 Aug 2024 11:20:32 +0100 Subject: [PATCH 1/4] [SYCL][Bindless] Allow 3-channel image formats This patch doesn't implement support for 3-channel formats, it just allows it in the spec and implementation so that a backend could implement it later (e.g. Level Zero). * Removed check from `image_descriptor::verify` * Updated spec to allow 3-channel formats * Very simple test, just verifies that CUDA backend still fails * Missing some PTX intrinsics, use `float4` instead of `float4` when compiling for CUDA. The more important part is checking that the CUDA backend throws an exception when trying to create image. --- .../sycl_ext_oneapi_bindless_images.asciidoc | 5 +- .../ext/oneapi/bindless_images_descriptor.hpp | 12 +- .../bindless_images/3_channel_format.cpp | 109 ++++++++++++++++++ 3 files changed, 118 insertions(+), 8 deletions(-) create mode 100644 sycl/test-e2e/bindless_images/3_channel_format.cpp 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 7218217298237..7cb33f92a01f1 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -239,7 +239,9 @@ descriptor against the limitations outlined below. If the given descriptor is deemed invalid, then a `sycl::exception` will be thrown with error code `sycl::errc::invalid`. -For all image types, the value of `num_channels` must be `1`, `2`, or `4`. +The value of `num_channels` supported by all image types and backends +is `1`, `2`, or `4`. +Some backends also support `num_channels` to be `3`. For the `standard` image type, the value of `num_levels` and `array_size` must both be `1`. @@ -2884,4 +2886,5 @@ These features still need to be handled: handles and the imported `interop_xxx_handle`. |5.17|2024-07-30| - Add support for mapping external memory to linear USM using `map_external_linear_memory`. +|5.18|2024-08-27| - Allow 3-channel image formats on some backends. |====================== diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp index e0635c99a776c..ba7015acb52e3 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp @@ -23,13 +23,15 @@ namespace ext::oneapi::experimental { namespace detail { -inline image_channel_order +constexpr image_channel_order get_image_default_channel_order(unsigned int num_channels) { switch (num_channels) { case 1: return image_channel_order::r; case 2: return image_channel_order::rg; + case 3: + return image_channel_order::rgb; case 4: return image_channel_order::rgba; default: @@ -120,13 +122,9 @@ struct image_descriptor { } void verify() const { - - if (this->num_channels != 1 && this->num_channels != 2 && - this->num_channels != 4) { - // Images can only have 1, 2, or 4 channels. + if ((this->num_channels < 1) || (this->num_channels > 4)) { throw sycl::exception(sycl::errc::invalid, - "Images must have only 1, 2, or 4 channels! Use a " - "valid number of channels instead."); + "Images must have 1, 2, 3, or 4 channels."); } switch (this->type) { diff --git a/sycl/test-e2e/bindless_images/3_channel_format.cpp b/sycl/test-e2e/bindless_images/3_channel_format.cpp new file mode 100644 index 0000000000000..f227bdcf1cade --- /dev/null +++ b/sycl/test-e2e/bindless_images/3_channel_format.cpp @@ -0,0 +1,109 @@ +// REQUIRES: cuda + +// RUN: %{build} -o %t.out +// RUN: %t.out + +#include +#include + +#include + +// Uncomment to print additional test information +// #define VERBOSE_PRINT + +class image_kernel; + +namespace syclexp = sycl::ext::oneapi::experimental; + +int main() { + sycl::device dev; + sycl::queue q(dev); + auto ctxt = q.get_context(); + + constexpr size_t width = 512; + std::vector out(width); + std::vector expected(width); + std::vector dataIn(width); + float exp = 512; + for (int i = 0; i < width; i++) { + expected[i] = exp; + dataIn[i] = sycl::float3(exp, width, i); + } + + try { + // Main point of this test is to check creating an image + // with a 3-channel format + syclexp::image_descriptor desc({width}, 3, sycl::image_channel_type::fp32); + + syclexp::image_mem imgMem(desc, dev, ctxt); + + q.ext_oneapi_copy(dataIn.data(), imgMem.get_handle(), desc); + q.wait_and_throw(); + + // Some backends don't support 3-channel formats + // We still try to create the image, + // but we expect it to fail with UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT + syclexp::unsampled_image_handle imgHandle = + sycl::ext::oneapi::experimental::create_image(imgMem, desc, dev, ctxt); + + sycl::buffer buf(out.data(), width); + + q.submit([&](sycl::handler &cgh) { + sycl::accessor outAcc{buf}; + + cgh.parallel_for(width, [=](sycl::id<1> id) { +#if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) + // This shouldn't be hit anyway since CUDA doesn't support + // 3-channel formats, but we need to ensure the kernel can compile + using pixel_t = sycl::float4; +#else + using pixel_t = sycl::float3; +#endif + auto pixel = syclexp::fetch_image(imgHandle, int(id[0])); + outAcc[id] = pixel[0]; + }); + }); + q.wait_and_throw(); + + } catch (const sycl::exception &ex) { + const std::string_view errMsg(ex.what()); + if (ctxt.get_backend() == sycl::backend::ext_oneapi_cuda) { + if (errMsg.find("UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT") != + std::string::npos) { + std::cout << "CUDA doesn't support 3-channel formats, test passed." + << std::endl; + return 0; + } + } + std::cerr << "Unexpected SYCL exception: " << errMsg << "\n"; + return 1; + } catch (...) { + std::cerr << "Unknown exception caught!\n"; + return 2; + } + + bool validated = true; + for (int i = 0; i < width; i++) { + bool mismatch = false; + if (out[i] != expected[i]) { + mismatch = true; + validated = false; + } + + if (mismatch) { +#ifdef VERBOSE_PRINT + std::cout << "Result mismatch! Expected: " << expected[i] + << ", Actual: " << out[i] << std::endl; +#else + break; +#endif + } + } + if (validated) { + std::cout << "Test passed!" << std::endl; + return 0; + } + + std::cout << "Test failed!" << std::endl; + return 3; +} From e0f186ef65dc42e3569c137874181131eaf15c81 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Thu, 29 Aug 2024 08:37:30 +0100 Subject: [PATCH 2/4] Update sycl/test-e2e/bindless_images/3_channel_format.cpp Co-authored-by: Wenju He --- sycl/test-e2e/bindless_images/3_channel_format.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/sycl/test-e2e/bindless_images/3_channel_format.cpp b/sycl/test-e2e/bindless_images/3_channel_format.cpp index f227bdcf1cade..40750ba691081 100644 --- a/sycl/test-e2e/bindless_images/3_channel_format.cpp +++ b/sycl/test-e2e/bindless_images/3_channel_format.cpp @@ -49,7 +49,7 @@ int main() { sycl::buffer buf(out.data(), width); q.submit([&](sycl::handler &cgh) { - sycl::accessor outAcc{buf}; + sycl::accessor outAcc{buf, cgh}; cgh.parallel_for(width, [=](sycl::id<1> id) { #if defined(__SYCL_DEVICE_ONLY__) && defined(__NVPTX__) From f30fbaaae2b5a2ae6b155a9972a038fa030aa181 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Thu, 29 Aug 2024 15:04:29 +0100 Subject: [PATCH 3/4] Issue note on missing query Currently no way to query whether a backends supports 3 channel formats. --- .../experimental/sycl_ext_oneapi_bindless_images.asciidoc | 6 ++++++ 1 file changed, 6 insertions(+) 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 7cb33f92a01f1..b4638971452fe 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2683,6 +2683,12 @@ There are dimension specific limitations: * 3D - No support at the moment. Possible support in non CUDA backends in the future. +=== 3 channel format support + +The ability to create an image with 3 channels depends on the backend. +There is currently no way to query a backend whether it supports this feature. +This query should be added in a later revision of the proposal. + === Not supported yet These features still need to be handled: From 42faaef86fc3784ffa75c794dc427e233146607b Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?Peter=20=C5=BDu=C5=BEek?= Date: Fri, 6 Sep 2024 10:03:09 +0100 Subject: [PATCH 4/4] Revision 6.1 --- .../experimental/sycl_ext_oneapi_bindless_images.asciidoc | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) 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 d2869d9dbc5ba..1a8323dde36da 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -2894,7 +2894,7 @@ These features still need to be handled: handles and the imported `interop_xxx_handle`. |5.17|2024-07-30| - Add support for mapping external memory to linear USM using `map_external_linear_memory`. -|6 |2024-08-29 | - Collated all changes since revision 5. +|6 |2024-08-05 | - Collated all changes since revision 5. - Bumped SYCL_EXT_ONEAPI_BINDLESS_IMAGES to number 6. - - Allow 3-channel image formats on some backends. +|6.1 |2024-09-06| - Allow 3-channel image formats on some backends. |======================