[SYCL][Bindless][UR] Add support for sRGB image formats - #22257
[SYCL][Bindless][UR] Add support for sRGB image formats#22257juanchuletas wants to merge 18 commits into
Conversation
8e6a729 to
15fc9ae
Compare
Update,I tested the changes using: 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
left a comment
There was a problem hiding this comment.
L0 adapter side LGTM, 2 nitpicks.
0x12CC
left a comment
There was a problem hiding this comment.
Thanks, @juanchuletas. Please add relevant tests.
There was a problem hiding this comment.
These changes need corresponding changes in sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc to ensure the implementation and specification remain consistent.
7d07faf to
26ae9ea
Compare
|
I was thinking about the current implementation and some scenarios came out. The current implementation exposes 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
left a comment
There was a problem hiding this comment.
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?
I am open to work on any of the approaches |
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
I lean towards replacement of |
|
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 |
Sorry for the late reply. I think this direction makes sense, and using a new |
Great! Working on it! |
d16f9d1 to
b284a0d
Compare
|
Hi! @dyniols, @kswiecicki, @0x12CC On my last commit I moved the check to pupulate_ur_structs because adding the sRGB check inside // 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 What should we do here: -Should we make an effort and add color_space to the constructor? |
|
I think adding Regarding |
…er __INTEL_PREVIEW_BREAKING_CHANGES
|
Hi! I made the changes, 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 |
There was a problem hiding this comment.
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_spaceand 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
SRGBAlayouts/swizzles and appendze_srgb_ext_desc_tinto thepNextchain 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_SRGBAis currently treated like generic 4-channel formats and will map to 16/32-bit layouts depending onchannelType. 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 (returnUR_RESULT_ERROR_UNSUPPORTED_IMAGE_FORMATwhenSRGBAis 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;
iclsrc
left a comment
There was a problem hiding this comment.
This PR adds image_color_space::srgb support to bindless images, gating the new color_space field and constructors behind __INTEL_PREVIEW_BREAKING_CHANGES and wiring UR_IMAGE_CHANNEL_ORDER_SRGBA through to the Level Zero ze_srgb_ext_desc_t extension. The core UR and SYCL runtime plumbing looks correct. Three issues are noted below: one correctness bug in get_mip_level_desc, one test-coverage gap (only R channel validated), and a suggestion for a device-independent unit test.
… sRGB constraints
|
@iclsrc I saw several tests are failing, is there any else I can do ? |
|
@juanchuletas I will be gone on vacations and I will be back after 26th July FYI. You might need toping other bindless image reviewers. |
Can you merge the latest intel/llvm into your PR branch? I will review but I would prefer to re-run CI with newest intel/llvm changes. |
Merge done! :) |
|
Hi @iclsrc , is anything else I can do here? |
Oh, this is a Claude bot, not a real person. I will take a look. I see Pre Commit is still failing. |
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 8 out of 8 changed files in this pull request and generated 3 comments.
Comments suppressed due to low confidence (4)
sycl/test-e2e/bindless_images/srgb/fetch_unsampled_srgb.cpp:20
- Missing standard headers: this test uses
uint8_tandstd::vectorbut does not include<cstdint>/<vector>, which can break compilation depending on transitive includes.
#include <cmath>
#include <iostream>
#include <sycl/detail/core.hpp>
#include <sycl/ext/oneapi/bindless_images.hpp>
sycl/test-e2e/bindless_images/srgb/fetch_unsampled_srgb.cpp:96
- Bindless image coordinates are
(x, y)for a 2D descriptor created as{width, height}; the kernel currently passes(y, x)tofetch_image, which becomes out-of-bounds ifwidth != height.
acc[sycl::id<2>(dim0, dim1)] = syclexp::fetch_image<sycl::float4>(
linearImg, sycl::int2(dim0, dim1));
sycl/test-e2e/bindless_images/srgb/fetch_sampled_srgb.cpp:111
- For a 2D descriptor created as
{width, height}, normalized sample coordinates are(x/width, y/height). The test currently passes(y/height, x/width), which is inconsistent with the bindless image coordinate conventions (seecoordinate_order_2D.cpp).
float fdim0 = (float(dim0) + 0.5f) / float(height);
float fdim1 = (float(dim1) + 0.5f) / float(width);
acc[sycl::id<2>(dim0, dim1)] =
syclexp::sample_image<sycl::float4>(
linearImg, sycl::float2(fdim0, fdim1));
sycl/test-e2e/bindless_images/srgb/fetch_pitched_srgb.cpp:121
- For a 2D descriptor created as
{width, height}, normalized sample coordinates are(x/width, y/height). The test currently passes(y/height, x/width), which is inconsistent with the bindless image coordinate conventions (seecoordinate_order_2D.cpp).
float fdim0 = (float(dim0) + 0.5f) / float(height);
float fdim1 = (float(dim1) + 0.5f) / float(width);
acc[sycl::id<2>(dim0, dim1)] =
syclexp::sample_image<sycl::float4>(
linearImg, sycl::float2(fdim0, fdim1));
|
Changes suggested by copilot: Done |
|
@juanchuletas Can you resolve conflicts? Thanks and please ping me I will restart CI. |
There was a problem hiding this comment.
Pull request overview
Copilot reviewed 8 out of 8 changed files in this pull request and generated 1 comment.
Suppressed comments (3)
unified-runtime/source/adapters/level_zero/image_common.cpp:623
- Adding SRGBA to this shared translator also makes non-bindless
urMemImageCreateaccept the format. Bothmemory.cpp:1514-1523andv2/memory.cpp:549-557callur2zeImageDescand thenzeImageCreatewithout chainingze_srgb_ext_desc_t, so those APIs now return a linear RGBA image for an SRGBA request instead of decoding sRGB (or rejecting it). Keep SRGBA acceptance bindless-only, or attach the sRGB descriptor at every image-creation caller.
case UR_IMAGE_CHANNEL_ORDER_SRGBA: {
sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp:56
- This adds a public enum, descriptor member, constructor parameters, validation rules, and Level Zero-only behavior, but
sycl_ext_oneapi_bindless_images.asciidocstill specifies the oldimage_descriptorAPI, whilesycl_ext_oneapi_srgb.asciidocstill states that sRGB is OpenCL-only. Update the extension specifications so users can determine the API contract, preview gating, supported formats, and backend limitations.
/// image color space enum
enum class image_color_space : uint32_t {
linear = 0,
srgb = 1,
};
sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp:56
- The PR description says
image_color_spaceitself is available only with__INTEL_PREVIEW_BREAKING_CHANGES, but the enum is currently exposed in non-preview builds. Gate this declaration with the same macro as the field and constructors so the documented preview boundary matches the API.
This issue also appears on line 52 of the same file.
/// image color space enum
enum class image_color_space : uint32_t {
linear = 0,
srgb = 1,
};
|
@dyniols Hi Daniel, I fixed the issue that copilot found. |
|
@juanchuletas Can you merge the latest intel/llvm into your PR branch and resolve conflicts in |
Done! :) |
dyniols
left a comment
There was a problem hiding this comment.
Okay, changes look good for me except few comments.
Let's wait for CI and see if all tests are passing.
| syclexp::image_color_space::linear)); | ||
| } | ||
|
|
||
| #endif // __INTEL_PREVIEW_BREAKING_CHANGES No newline at end of file |
|
|
||
| std::cout << "Test passed!" << std::endl; | ||
| return 0; | ||
| } No newline at end of file |
|
|
||
| std::cout << "Test passed!" << std::endl; | ||
| return 0; | ||
| } No newline at end of file |
|
|
||
| std::cout << "Test passed!" << std::endl; | ||
| return 0; | ||
| } No newline at end of file |
| // REQUIRES: aspect-ext_oneapi_bindless_images | ||
| // REQUIRES: preview-breaking-changes-supported | ||
|
|
||
| // UNSUPPORTED: cuda, hip |
There was a problem hiding this comment.
It's supported only in Level Zero. I would use REQUIRE: level_zero instead.
On second thought I think we should add new aspect that will help to check for support of SRGB channel order in bindless image extension. However it can be done as follow up in other PR.
| }); | ||
| }); | ||
|
|
||
| q.wait_and_throw(); |
There was a problem hiding this comment.
This wait can be omitted since SYCL buffer's are emplaced within braces and destructor should handle synchronization and copying data back from device to host.
| }); | ||
| }); | ||
|
|
||
| q.wait_and_throw(); |
| }); | ||
| }); | ||
|
|
||
| q.wait_and_throw(); |
|
I saw several tests failing on CI, so I checked and found that the test files hadn't been updated after I added image_color_space as the 4th argument to the image_descriptor constructor. They were still passing image_type in that slot. So, I edited the tests to pass image_color_space::linear before image_type, matching the constructor. I also fixed the comments you left: replaced UNSUPPORTED: cuda, hip with REQUIRES: level_zero in the sRGB tests, since sRGB decode is only supported on Level Zero for now and I added the trailing newline. |
…ructor Update image_descriptor calls to pass image_color_space before image_type, and require level_zero for sRGB tests instead of excluding cuda and hip.
|
Some tests are still failing :( but I think the reason is teh failing ones have no the "preview-mode", other failing test is something about AMDGPU_TARGETS was not set, and system GPU detection was unsuccsesful |
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_descriptoronly 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.hppAdded a new
image_color_spaceenum (linear/srgb) and acolor_spacefield onimage_descriptor, gated behind__INTEL_PREVIEW_BREAKING_CHANGESsince it changes the descriptor's constructor signatures. Under that macro, all three range-based constructors (1D/2D/3D) gain acolor_spaceparameter defaulting toimage_color_space::linear, preserving existing call sites when the flag is off. Added a validation check inverify()(also macro-gated) requiringnum_channels == 4andchannel_type == unorm_int8whenevercolor_space == srgb, matching the only configuration the Level Zero driver accepts for sRGB images.sycl/source/detail/bindless_images.cppModified
populate_ur_structsso that, under__INTEL_PREVIEW_BREAKING_CHANGES,urFormat.channelOrderis set toUR_IMAGE_CHANNEL_ORDER_SRGBAwhendesc.color_space == image_color_space::srgb, and falls back to the default channel order derived fromnum_channelsotherwise. Without the preview flag, behavior is unchanged.Proposed Usage