[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.

Author: ProGTX

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.
 |======================

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) {

sycl/test-e2e/bindless_images/3_channel_format.cpp

@@ -0,0 +1,109 @@
+// REQUIRES: cuda
+
+// RUN: %{build} -o %t.out
+// RUN: %t.out
+
+#include <iostream>
+#include <sycl/detail/core.hpp>
+
+#include <sycl/ext/oneapi/bindless_images.hpp>
+
+// 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<float> out(width);
+  std::vector<float> expected(width);
+  std::vector<sycl::float3> 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<float> buf(out.data(), width);
+
+    q.submit([&](sycl::handler &cgh) {
+      sycl::accessor outAcc{buf};
+
+      cgh.parallel_for<image_kernel>(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<pixel_t>(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;
+}