[SYCL][Bindless][UR] Add support for sRGB image formats

[juanchuletas]

Summary

This PR adds native sRGB decode support for bindless images through the full DPC++ stack: SYCL API, Unified Runtime, and the Level Zero adapter.

Description

The hardware sampler on Intel Arc GPUs supports native sRGB decode via ze_srgb_ext_desc_t, confirmed by direct Level Zero testing against driver 26.18.38308.1. The driver correctly applies the IEC 61966-2-1 piecewise formula on fetch, returning linear values to the kernel. This capability was not exposed anywhere in the DPC++ or UR stack.

The SYCL bindless image API had no way to request sRGB color space decoding. The image_descriptor only accepted num_channels and channel_type, with the channel order always derived internally. This made it impossible to distinguish between a linear RGBA texture and an sRGB-encoded RGBA texture, forcing applications to manually decode sRGB values on the CPU before upload at a significant memory cost.

Proposed Changes

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

Added a new image_color_space enum (linear / srgb) and a color_space field on image_descriptor, gated behind __INTEL_PREVIEW_BREAKING_CHANGES since it changes the descriptor's constructor signatures. Under that macro, all three range-based constructors (1D/2D/3D) gain a color_space parameter defaulting to image_color_space::linear, preserving existing call sites when the flag is off. Added a validation check in verify() (also macro-gated) requiring num_channels == 4 and channel_type == unorm_int8 whenever color_space == srgb, matching the only configuration the Level Zero driver accepts for sRGB images.

sycl/source/detail/bindless_images.cpp

Modified populate_ur_structs so that, under __INTEL_PREVIEW_BREAKING_CHANGES, urFormat.channelOrder is set to UR_IMAGE_CHANNEL_ORDER_SRGBA when desc.color_space == image_color_space::srgb, and falls back to the default channel order derived from num_channels otherwise. Without the preview flag, behavior is unchanged.

Proposed Usage

sycl::ext::oneapi::experimental::image_descriptor desc(
    sycl::range<2>(width, height),
    4,
    sycl::image_channel_type::unorm_int8,
    sycl::ext::oneapi::experimental::image_color_space::srgb
);

sycl::ext::oneapi::experimental::image_mem imgMem(desc, q);

Note: image_color_space and the color_space constructor parameter are only available when building with __INTEL_PREVIEW_BREAKING_CHANGES defined.

[juanchuletas]

Update,

I tested the changes using:

srgb_vs_rgb.cpp

and I got:

Device: Intel(R) Arc(TM) Graphics

Input raw pixel value:            186 (0xba)
Input normalized (unorm_int8):    0.729412
Expected linear :  0.491021

RGBA  result (no decode):  R=0.729412 G=0.729412 B=0.729412 A=1
sRGBA result (hardware decoded): R=0.490745 G=0.490745 B=0.490745 A=1

Difference in R channel:   0.238666

PASS: RGBA returns raw unorm, sRGBA returns decoded linear

[kswiecicki]

L0 adapter side LGTM, 2 nitpicks.

[0x12CC]

Thanks, @juanchuletas. Please add relevant tests.

[0x12CC]

These changes need corresponding changes in sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc to ensure the implementation and specification remain consistent.

[juanchuletas]

@0x12CC @kswiecicki

I was thinking about the current implementation and some scenarios came out. The current implementation exposes channel_order as an std::optional<image_channel_order> on image_descriptor.

Nothing stops the user from doing the following:

sycl::ext::oneapi::experimental::image_descriptor desc(
    sycl::range<2>(width, height),
    4,
    sycl::image_channel_type::unorm_int8
);
desc.channel_order = sycl::image_channel_order::rgx;

We can prohibit that kind of usage from the verify() method

if (channel_order.has_value()) {
  if (channel_order.value() != image_channel_order::ext_oneapi_srgba) {
    throw sycl::exception(
        sycl::errc::invalid,
        "channel_order only supports ext_oneapi_srgba");
  }
  if (num_channels != 4) {
    throw sycl::exception(
        sycl::errc::invalid,
        "ext_oneapi_srgba channel order requires num_channels == 4");
  }
  if (channel_type != image_channel_type::unorm_int8) {
    throw sycl::exception(
        sycl::errc::invalid,
        "ext_oneapi_srgba channel order requires unorm_int8 channel type");
  }
}

but it feels weird to expose desc.channel_order and only allowing srgb. Other approach is to Make channel_order fully general: validate channel count consistency in verify() for any value set.

[dyniols]

This PR is ABI breaking because of changes in image_descriptor struct layout.

I wonder if we could re-work image_descriptor back to support channel_order since it was replaced with num_channels due to lack of support in CUDA.
Here is PR when this change was introduced: #13745
I believe that Level Zero, OpenCL backend should support channel order.

Please add tests exercising the new sRGB path with sampled, unsampled and using pitched memory allocation, through pitched_alloc_device.

[juanchuletas]

This PR is ABI breaking because of changes in image_descriptor struct layout.

I wonder if we could re-work image_descriptor back to support channel_order since it was replaced with num_channels due to lack of support in CUDA. Here is PR when this change was introduced: #13745 I believe that Level Zero, OpenCL backend should support channel order.

Please add tests exercising the new sRGB path with sampled, unsampled and using pitched memory allocation, through pitched_alloc_device.

This PR is ABI breaking because of changes in image_descriptor struct layout.

I wonder if we could re-work image_descriptor back to support channel_order since it was replaced with num_channels due to lack of support in CUDA. Here is PR when this change was introduced: #13745 I believe that Level Zero, OpenCL backend should support channel order.

Please add tests exercising the new sRGB path with sampled, unsampled and using pitched memory allocation, through pitched_alloc_device.

Given that restoring channel_order as a primary field would also be an ABI break, and that CUDA does not support channel order natively, what is the preferred approach?

• replacement of num_channels with channel_order

• stay focused on sRGB hardware decode only

I am open to work on any of the approaches

[dyniols]

Given that restoring channel_order as a primary field would also be an ABI break.

I don't think the ABI break is avoidable here. Naturally, a change like this should land with the next major release. In the meantime, we can leverage __INTEL_PREVIEW_BREAKING_CHANGES macro to merge ABI-breaking change without delay and then remove it when possible.

what is the preferred approach?

• replacement of num_channels with channel_order
• stay focused on sRGB hardware decode only

I am open to work on any of the approaches

I lean towards replacement of num_channels with channel_order and using with image_channel_order enum class.
However I would ask @0x12CC on his perspective on such a change.

[juanchuletas]

I've been thinking about another approach. What do you think of this @dyniols , @0x12CC , @kswiecicki ?

Instead of reaching for image_channel_order at all. Why don't treat the sRGB as what it actually is: a property of how the sampler interprets the data:

enum class image_color_space : uint32_t {
    linear = 0,
    srgb   = 1,
};

struct image_descriptor {
    image_type        type          = image_type::standard;
    uint32_t          num_channels  = 0;
    image_channel_type channel_type;
    image_color_space color_space   = image_color_space::linear;  // the new field
   //Other fields
};

and usage like:

image_descriptor desc;
desc.num_channels = 4;
desc.channel_type  = image_channel_type::unorm_int8;
desc.color_space   = image_color_space::srgb;

Actually cuda does something similar, to enable the srgb

[dyniols]

Instead of reaching for image_channel_order at all. Why don't treat the sRGB as what it actually is: a property of how the sampler interprets the data:

enum class image_color_space : uint32_t {
    linear = 0,
    srgb   = 1,
};

struct image_descriptor {
    image_type        type          = image_type::standard;
    uint32_t          num_channels  = 0;
    image_channel_type channel_type;
    image_color_space color_space   = image_color_space::linear;  // the new field
   //Other fields
};

and usage like:

image_descriptor desc;
desc.num_channels = 4;
desc.channel_type  = image_channel_type::unorm_int8;
desc.color_space   = image_color_space::srgb;

Actually cuda does something similar, to enable the srgb

Sorry for the late reply. I think this direction makes sense, and using a new image_color_space field is probably a cleaner way to describe how the color data stored in an image should be interpreted.

[juanchuletas]

Instead of reaching for image_channel_order at all. Why don't treat the sRGB as what it actually is: a property of how the sampler interprets the data:

enum class image_color_space : uint32_t {
    linear = 0,
    srgb   = 1,
};

struct image_descriptor {
    image_type        type          = image_type::standard;
    uint32_t          num_channels  = 0;
    image_channel_type channel_type;
    image_color_space color_space   = image_color_space::linear;  // the new field
   //Other fields
};

and usage like:

image_descriptor desc;
desc.num_channels = 4;
desc.channel_type  = image_channel_type::unorm_int8;
desc.color_space   = image_color_space::srgb;

Actually cuda does something similar, to enable the srgb

Sorry for the late reply. I think this direction makes sense, and using a new image_color_space field is probably a cleaner way to describe how the color data stored in an image should be interpreted.

Great! Working on it!

[juanchuletas]

Hi! @dyniols, @kswiecicki, @0x12CC

On my last commit I moved the check to pupulate_ur_structs because adding the sRGB check inside verify() makes no sense as verify() is called inside the constructor and the current usage is like:

 // sRGBA image: hardware applies IEC 61966-2-1 decode on fetch
    syclexp::image_descriptor srgbaDesc(
        sycl::range<2>(1, 1),
        4,
        sycl::image_channel_type::unorm_int8
    );
    srgbaDesc.color_space = sycl::ext::oneapi::experimental::image_color_space::srgb;

sRGB color space is set after the object construction. I thought in adding the image_color_space field to the constructor, but that will lead into a very big change.

Totally apart, working on that I noticed that the user can set an image descriptor like this:

image_descriptor desc(range<2>{w, h}, 4, image_channel_type::unorm_int8);
desc.num_levels = 5;  // bypasses verify()
desc.array_size = 10; // bypasses verify()

And verify() never runs again.

What should we do here:

-Should we make an effort and add color_space to the constructor?
-Should we mark those fields as private ?
-Should we verify in other place?

[dyniols]

Hi @juanchuletas

I think adding image_color_space field to constructor seems the most reasonable but of course it will be API and ABI breaking change. So it will need to be marked as such.
Your thoughts @0x12CC?

Regarding verify() it runs in more places than just the constructor. It's also called at the start of the functions that are using the image descriptor (alloc_image_mem, create_image, map_external_image_memory, and the copy functions in handler.cpp) all call verify() right before passing it to unified runtime.

[juanchuletas]

@dyniols

Hi! I made the changes, if that looks correct for you I can start adding the proper tests.

[dyniols]

if that looks correct for you I can start adding the proper tests.

Sure, I’ll review the current implementation changes.

There will be another requirement for this PR to update sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc so the specification and implementation stay consistent. However, you can wait to do that until the implementation is accepted.

[Copilot]

Pull request overview

Note

Copilot couldn't run its full agentic review because no GitHub Actions runner was available. Make sure your repository has a runner available to run Copilot's review, or add a copilot-setup-steps.yml file specifying one with the runs-on attribute. See the docs for more details.

Adds end-to-end support for native sRGB decode for bindless images, plumbing a new sRGB selection from SYCL down through Unified Runtime to Level Zero via ze_srgb_ext_desc_t.

Changes:

• Extends SYCL bindless image descriptors with a preview-only image_color_space and validation for sRGB constraints.
• Propagates sRGB selection into UR image format (UR_IMAGE_CHANNEL_ORDER_SRGBA) under the preview macro.
• Updates Level Zero adapter to map SRGBA layouts/swizzles and append ze_srgb_ext_desc_t into the pNext chain for bindless creation/allocation.

Reviewed changes

Copilot reviewed 3 out of 3 changed files in this pull request and generated 6 comments.

File	Description
unified-runtime/source/adapters/level_zero/image_common.cpp	Adds UR_IMAGE_CHANNEL_ORDER_SRGBA mapping and chains ze_srgb_ext_desc_t into Level Zero image creation/allocation.
sycl/source/detail/bindless_images.cpp	Propagates SYCL descriptor sRGB intent into UR image format under preview macro.
sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp	Introduces preview-only image_color_space and adds validation for sRGB-required format constraints.

Comments suppressed due to low confidence (1)

unified-runtime/source/adapters/level_zero/image_common.cpp:623

• UR_IMAGE_CHANNEL_ORDER_SRGBA is currently treated like generic 4-channel formats and will map to 16/32-bit layouts depending on channelType. However, the PR description indicates Level Zero sRGB decode support is limited (e.g., unorm_int8 / 8-bit). The UR adapter should enforce this at the mapping layer (return UR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMAT when SRGBA is used with non-8-bit and/or non-UNORM formats) so non-SYCL callers can’t produce a format that will be rejected or misinterpreted by the driver.

  case UR_IMAGE_CHANNEL_ORDER_ABGR:
  case UR_IMAGE_CHANNEL_ORDER_SRGBA: {
    switch (ZeImageFormatTypeSize) {
    case 8:
      ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_8_8_8_8;
      break;
    case 16:
      ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_16_16_16_16;
      break;
    case 32:
      ZeImageFormatLayout = ZE_IMAGE_FORMAT_LAYOUT_32_32_32_32;
      break;

[dyniols]

The bindless image part looks good overall, aside from a few minor issues, but it still needs tests to validate the new functionality.

[dyniols]

@juanchuletas can you sync your PR with current intel/llvm head to resolve conflicts ?

[juanchuletas]

@juanchuletas can you sync your PR with current intel/llvm head to resolve conflicts ?

Sure. Do you have a preference between rebasing onto the latest intel/sycl or merging the latest intel/sycl into my branch?

[dyniols]

Sure. Do you have a preference between rebasing onto the latest intel/sycl or merging the latest intel/sycl into my branch?

I personally prefer the latter.

[dyniols]

I am okay with changes, now we need tests to cover it :)

Could you take a look as well @kswiecicki @0x12CC?

[Copilot]

Pull request overview

Copilot reviewed 3 out of 3 changed files in this pull request and generated 3 comments.

[juanchuletas]

@dyniols

I added some tests. Please tell me if more are needed. Also should I update the sycl/ReleaseNotes.md and sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc ?

[dyniols]

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

Thanks, I will take a look. For now just sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc