[SYCL][Bindless] Replace 'image_channel_order' field in 'image_descriptor' with number of channels (#13745)

'image_channel_order' field in 'image_descriptor' is and can only be
used for the number of channels due to CUDA having no notion of image
channel order. Replaced with the number of channels instead.

Also deprecate 'get_channel_order' function from 'image_mem' as it is
redundant because images no longer have or need a notion of channel
order. Only channel size.

Author: DBDuncan

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -159,23 +159,6 @@ of 3D USM images.
 ```cpp
 namespace sycl::ext::oneapi::experimental {
 
-enum class image_channel_order : /* unspecified */ {
-  a,
-  r,
-  rx,
-  rg,
-  rgx,
-  ra,
-  rgb,
-  rgbx,
-  rgba,
-  argb,
-  bgra,
-  intensity,
-  luminance,
-  abgr,
-};
-
 enum class image_channel_type : /* unspecified */ {
   snorm_int8,
   snorm_int16,
@@ -202,25 +185,25 @@ struct image_descriptor {
   size_t width{0};
   size_t height{0};
   size_t depth{0};
-  image_channel_order channel_order{image_channel_order::rgba};
+  unsigned int num_channels{4};
   image_channel_type channel_type{image_channel_type::fp32};
   image_type type{image_type::standard};
   unsigned int num_levels{1};
   unsigned int array_size{1};
 
   image_descriptor() = default;
 
-  image_descriptor(sycl::range<1> dims, image_channel_order channel_order,
+  image_descriptor(sycl::range<1> dims, unsigned int num_channels,
                    image_channel_type channel_type, 
                    image_type type = image_type::standard, 
                    unsigned int num_levels = 1, unsigned int array_size = 1);
 
-  image_descriptor(sycl::range<2> dims, image_channel_order channel_order,
+  image_descriptor(sycl::range<2> dims, unsigned int num_channels,
                    image_channel_type channel_type, 
                    image_type type = image_type::standard, 
                    unsigned int num_levels = 1, unsigned int array_size = 1);
 
-  image_descriptor(sycl::range<3> dims, image_channel_order channel_order,
+  image_descriptor(sycl::range<3> dims, unsigned int num_channels,
                    image_channel_type channel_type, 
                    image_type type = image_type::standard, 
                    unsigned int num_levels = 1, unsigned int array_size = 1);
@@ -233,9 +216,9 @@ struct image_descriptor {
 }
 ```
 
-The image descriptor represents the image dimensions, channel type, and channel 
-order. An `image_type` member is also present to allow for implementation of 
-mipmapped, image array, and cubemapped images.
+The image descriptor represents the image dimensions, number of channels, and
+channel type. An `image_type` member is also present to allow for implementation
+of mipmapped, image array, and cubemapped images.
 
 The `image_descriptor` shall be default constructible and follow by-value 
 semantics.
@@ -256,6 +239,8 @@ 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`.
+
 For the `standard` image type, the value of `num_levels` and `array_size` must
 both be `1`.
 
@@ -311,9 +296,8 @@ public:
   sycl::context get_context() const;
 
   sycl::range<3> get_range() const;
-  sycl::image_channel_type get_image_channel_type() const;
-  sycl::image_channel_type get_image_channel_order() const;
-  unsigned int get_image_num_channels() const;
+  sycl::image_channel_type get_channel_type() const;
+  unsigned int get_num_channels() const;
   image_type get_type() const;
 
   image_mem_handle get_mip_level_mem_handle(unsigned int level) const;
@@ -394,7 +378,7 @@ using the `image_mem_alloc` function. These are similar to the member functions
 provided by `image_mem`. However, since the `image_mem_handle` is a minimal 
 struct representing just the opaque handle the underlying memory object, there 
 is some information that we cannot retrieve from it, namely the `image_type`,
-`image_channel_order`, the `sycl::context` or `sycl::device` the memory was 
+`num_channels`, the `sycl::context` or `sycl::device` the memory was
 allocated in, and the `image_descriptor` used to allocate the memory.
 
 ```cpp
@@ -1072,13 +1056,13 @@ void write_image(unsampled_image_handle ImageHandle,
 ```
 
 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 
+or `sample_image`, passing the appropriate image handle. The `fetch_image` API 
 is applicable to sampled and 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.
 
-When fetching from a sampled image handle, data exatly as is in memory, no 
+When fetching from a sampled image handle, data exactly as is in memory, no 
 sampling operations will be performed, and the `bindless_image_sampler` passed 
 to the image upon creation has no effect on the returned image data. Note that 
 not all devices may support fetching of sampled image data depending on the 
@@ -1669,7 +1653,7 @@ When calling `create_image` with an `image_mem_handle` mapped from an external
 memory object, the user must ensure that the image descriptor they pass to 
 `create_image` has members that match or map to those of the external API. 
 A mismatch between any of the `width`, `height`, `depth`, `image_channel_type`,
-or `image_channel_order` members will result in undefined behavior.
+or `num_channels` members will result in undefined behavior.
 
 Additionally, the `image_type` describing the image must match to the image of 
 the external API. The current supported importable image types are `standard` 
@@ -1851,7 +1835,7 @@ for (int i = 0; i < width; i++) {
 
 // Image descriptor - can use the same for both images
 sycl::ext::oneapi::experimental::image_descriptor desc(
-    sycl::range{width}, sycl::ext::oneapi::experimental::image_channel_order::r,
+    sycl::range{width}, 1,
     sycl::ext::oneapi::experimental::image_channel_type::fp32);
 
 try {
@@ -1930,7 +1914,7 @@ for (int i = 0; i < width; i++) {
 
 // Image descriptor - can use the same for all images
 sycl::ext::oneapi::experimental::image_descriptor desc(
-    {width, height}, sycl::ext::oneapi::experimental::image_channel_order::r,
+    {width, height}, 1,
     sycl::ext::oneapi::experimental::image_channel_type::fp32);
 
 try {
@@ -2030,7 +2014,7 @@ try {
 
   // Image descriptor -- number of levels
   sycl::ext::oneapi::experimental::image_descriptor desc(
-      {width}, sycl::ext::oneapi::experimental::image_channel_order::r, 
+      {width}, 1,
       sycl::ext::oneapi::experimental::image_channel_type::fp32,
       sycl::ext::oneapi::experimental::image_type::mipmap, num_levels);
 
@@ -2130,7 +2114,7 @@ for (int i = 0; i < width; i++) {
 try {
   // Extension: image descriptor -- number of layers
   sycl::ext::oneapi::experimental::image_descriptor desc(
-      {width}, sycl::image_channel_order::rgba, sycl::image_channel_type::fp32,
+      {width}, 4, sycl::image_channel_type::fp32,
       sycl::ext::oneapi::experimental::image_type::array, 1, array_size);
 
   // Extension: allocate image array memory on device
@@ -2259,7 +2243,7 @@ int main() {
 
   // Extension: image descriptor - Cubemap
   syclexp::image_descriptor desc(
-      {width, height}, sycl::image_channel_order::rgba,
+      {width, height}, 4,
       sycl::image_channel_type::fp32, syclexp::image_type::cubemap, 1, 6);
 
   syclexp::bindless_image_sampler samp(
@@ -2362,17 +2346,17 @@ sycl::context context = queue.get_context();
 size_t width = /* passed from external API */;
 size_t height = /* passed from external API */;
 
-sycl::ext::oneapi::experimental::image_channel_order channel_order = 
+unsigned int num_channels = 1;
     /* mapped from external API */
-    /* we assume sycl::image_channel_order::r */;
+    /* we assume there is one channel */;
 
 sycl::ext::oneapi::experimental::image_channel_type channel_type = 
     /* mapped from external API */
     /* we assume sycl::image_channel_type::unsigned_int32 */;
 
 // Image descriptor - mapped to external API image layout
 sycl::ext::oneapi::experimental::image_descriptor desc(
-    {width, height}, channel_order, channel_type);
+    {width, height}, num_channels, channel_type);
 
 size_t img_size_in_bytes = width * height * sizeof(uint32_t);
 
@@ -2689,4 +2673,13 @@ These features still need to be handled:
 |5.8|2024-05-09| - Add missing cubemap `HintT` template parameter to 
                    `fetch_cubemap` and `sample_cubemap`.
 |5.9|2024-05-14| - Default constructor for `image_descriptor`.
+|5.10|2024-05-20| - Replaced `channel_order` field in `image_descriptor` with
+                   `num_channels`.
+                  - Renamed `image_mem` functions `get_image_channel_type()`
+                    to `get_channel_type()` and `get_image_num_channels()` to
+                    `get_num_channels()`.
+                  - Removed `get_channel_order()` function from `image_mem`.
+                    This function is redundant since images don't have a notion
+                    of channel order, only the channel size. Use
+                    `get_num_channels()` instead.
 |======================

sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp

@@ -21,6 +21,25 @@ namespace sycl {
 inline namespace _V1 {
 namespace ext::oneapi::experimental {
 
+namespace detail {
+
+inline 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 4:
+    return image_channel_order::rgba;
+  default:
+    assert(false && "Invalid channel number");
+    return static_cast<image_channel_order>(0);
+  }
+}
+
+} // namespace detail
+
 /// image type enum
 enum class image_type : unsigned int {
   standard = 0,
@@ -34,40 +53,40 @@ struct image_descriptor {
   size_t width{0};
   size_t height{0};
   size_t depth{0};
-  image_channel_order channel_order{image_channel_order::rgba};
+  unsigned int num_channels{4};
   image_channel_type channel_type{image_channel_type::fp32};
   image_type type{image_type::standard};
   unsigned int num_levels{1};
   unsigned int array_size{1};
 
   image_descriptor() = default;
 
-  image_descriptor(range<1> dims, image_channel_order channel_order,
+  image_descriptor(range<1> dims, unsigned int num_channels,
                    image_channel_type channel_type,
                    image_type type = image_type::standard,
                    unsigned int num_levels = 1, unsigned int array_size = 1)
-      : width(dims[0]), height(0), depth(0), channel_order(channel_order),
+      : width(dims[0]), height(0), depth(0), num_channels(num_channels),
         channel_type(channel_type), type(type), num_levels(num_levels),
         array_size(array_size) {
     verify();
   }
 
-  image_descriptor(range<2> dims, image_channel_order channel_order,
+  image_descriptor(range<2> dims, unsigned int num_channels,
                    image_channel_type channel_type,
                    image_type type = image_type::standard,
                    unsigned int num_levels = 1, unsigned int array_size = 1)
-      : width(dims[0]), height(dims[1]), depth(0), channel_order(channel_order),
+      : width(dims[0]), height(dims[1]), depth(0), num_channels(num_channels),
         channel_type(channel_type), type(type), num_levels(num_levels),
         array_size(array_size) {
     verify();
   }
 
-  image_descriptor(range<3> dims, image_channel_order channel_order,
+  image_descriptor(range<3> dims, unsigned int num_channels,
                    image_channel_type channel_type,
                    image_type type = image_type::standard,
                    unsigned int num_levels = 1, unsigned int array_size = 1)
       : width(dims[0]), height(dims[1]), depth(dims[2]),
-        channel_order(channel_order), channel_type(channel_type), type(type),
+        num_channels(num_channels), channel_type(channel_type), type(type),
         num_levels(num_levels), array_size(array_size) {
     verify();
   };
@@ -93,14 +112,23 @@ struct image_descriptor {
 
     // This will generate the new descriptor with image_type standard
     // since individual mip levels are standard images
-    sycl::ext::oneapi::experimental::image_descriptor levelDesc(
-        {width, height, depth}, this->channel_order, this->channel_type);
+    image_descriptor levelDesc({width, height, depth}, this->num_channels,
+                               this->channel_type);
 
     levelDesc.verify();
     return levelDesc;
   }
 
   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.
+      throw sycl::exception(sycl::errc::invalid,
+                            "Images must have only 1, 2, or 4 channels! Use a "
+                            "valid number of channels instead.");
+    }
+
     switch (this->type) {
     case image_type::standard:
       if (this->array_size > 1) {

sycl/include/sycl/info/ext_oneapi_device_traits.def

@@ -66,7 +66,7 @@ __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
                          max_image_linear_height, uint32_t,
                          PI_EXT_ONEAPI_DEVICE_INFO_MAX_IMAGE_LINEAR_HEIGHT)
 
-// Bindles images mipmaps
+// Bindless images mipmaps
 __SYCL_PARAM_TRAITS_SPEC(ext::oneapi::experimental, device,
                          mipmap_max_anisotropy, float,
                          PI_EXT_ONEAPI_DEVICE_INFO_MIPMAP_MAX_ANISOTROPY)

sycl/plugins/unified_runtime/pi2ur.hpp

@@ -2947,7 +2947,7 @@ static void pi2urImageDesc(const pi_image_format *ImageFormat,
                                      UR_IMAGE_CHANNEL_ORDER_SRGBA)
 #undef PI_TO_UR_MAP_IMAGE_CHANNEL_ORDER
   default: {
-    die("piMemImageCreate: unsuppported image_channel_data_type.");
+    die("piMemImageCreate: unsuppported image_channel_order.");
   }
   }
 

sycl/source/detail/bindless_images.cpp

@@ -52,8 +52,9 @@ void populate_pi_structs(const image_descriptor &desc, pi_image_desc &piDesc,
   piFormat = {};
   piFormat.image_channel_data_type =
       sycl::detail::convertChannelType(desc.channel_type);
-  piFormat.image_channel_order =
-      sycl::detail::convertChannelOrder(desc.channel_order);
+  piFormat.image_channel_order = sycl::detail::convertChannelOrder(
+      sycl::ext::oneapi::experimental::detail::get_image_default_channel_order(
+          desc.num_channels));
 }
 
 detail::image_mem_impl::image_mem_impl(const image_descriptor &desc,
@@ -89,13 +90,15 @@ __SYCL_EXPORT sycl::image_channel_type image_mem::get_channel_type() const {
   return impl->get_descriptor().channel_type;
 }
 
-__SYCL_EXPORT sycl::image_channel_order image_mem::get_channel_order() const {
-  return impl->get_descriptor().channel_order;
+__SYCL_EXPORT_DEPRECATED("get_channel_order() is deprecated. "
+                         "Instead use get_channel_num().")
+sycl::image_channel_order image_mem::get_channel_order() const {
+  return sycl::ext::oneapi::experimental::detail::
+      get_image_default_channel_order(impl->get_descriptor().num_channels);
 }
 
 __SYCL_EXPORT unsigned int image_mem::get_num_channels() const {
-  return sycl::detail::getImageNumberChannels(
-      impl->get_descriptor().channel_order);
+  return impl->get_descriptor().num_channels;
 }
 
 __SYCL_EXPORT image_type image_mem::get_type() const {
@@ -785,10 +788,8 @@ __SYCL_EXPORT void *pitched_alloc_device(size_t *resultPitch,
                                          const image_descriptor &desc,
                                          const sycl::device &syclDevice,
                                          const sycl::context &syclContext) {
-  uint8_t numChannels =
-      sycl::detail::getImageNumberChannels(desc.channel_order);
   unsigned int elementSizeBytes =
-      sycl::detail::getImageElementSize(numChannels, desc.channel_type);
+      sycl::detail::getImageElementSize(desc.num_channels, desc.channel_type);
 
   size_t widthInBytes = desc.width * elementSizeBytes;
   size_t height = desc.height;

sycl/source/detail/device_info.hpp

@@ -2203,7 +2203,7 @@ inline uint32_t get_device_info_host<
 template <>
 inline float get_device_info_host<
     ext::oneapi::experimental::info::device::mipmap_max_anisotropy>() {
-  throw runtime_error("Bindless image mipaps are not supported on HOST device",
+  throw runtime_error("Bindless image mipmaps are not supported on HOST device",
                       PI_ERROR_INVALID_DEVICE);
 }