Skip to content

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

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Open
juanchuletas wants to merge 6 commits into
base: sycl
Choose a base branch
from
Open

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

Show file tree
Hide file tree
Changes from 3 commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
2900c7a
[SYCL][Bindless][UR] Add support for sRGB image formats
juanchuletas Jun 9, 2026
15fc9ae
[SYCL][Bindless][UR] applying clang format
juanchuletas Jun 9, 2026
26ae9ea
[SYCL][Bindless][UR] Address review feedback
juanchuletas Jun 10, 2026
69d0137
[SYCL][Bindless][UR] adding the image_color_space struct
juanchuletas Jun 22, 2026
3d89061
[SYCL][Bindless][UR] clang format
juanchuletas Jun 22, 2026
b284a0d
[SYCL][Bindless][UR] moving srgb check to pupulate_ur_structs
juanchuletas Jun 23, 2026
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
19 changes: 16 additions & 3 deletions sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -13,7 +13,8 @@
#include <sycl/image.hpp> // for image_channel_order, image_channel_type
#include <sycl/range.hpp> // for range

#include <algorithm> // for max
#include <algorithm> // for max

@dyniols dyniols Jun 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

nit: re-align the include comments and add // for std::optional.

#include <optional>
#include <stddef.h> // for size_t
#include <system_error> // for error_code

Expand Down Expand Up @@ -61,7 +62,7 @@ struct image_descriptor {
image_type type{image_type::standard};
unsigned int num_levels{1};
unsigned int array_size{1};

std::optional<image_channel_order> channel_order;

@0x12CC 0x12CC Jun 9, 2026

Copy link
Copy Markdown
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

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

image_descriptor() = default;

image_descriptor(range<1> dims, unsigned int num_channels,
Expand Down Expand Up @@ -127,7 +128,19 @@ struct image_descriptor {
throw sycl::exception(sycl::errc::invalid,
"Images must have 1, 2, 3, or 4 channels.");
}

if (channel_order.has_value() &&
channel_order.value() == image_channel_order::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");
}
}
switch (this->type) {
case image_type::standard:
if (this->array_size > 1) {
Expand Down
7 changes: 4 additions & 3 deletions sycl/source/detail/bindless_images.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -50,9 +50,10 @@ void populate_ur_structs(const image_descriptor &desc, ur_image_desc_t &urDesc,

urFormat = {};
urFormat.channelType = sycl::detail::convertChannelType(desc.channel_type);
urFormat.channelOrder = sycl::detail::convertChannelOrder(
sycl::ext::oneapi::experimental::detail::get_image_default_channel_order(
desc.num_channels));
urFormat.channelOrder =
sycl::detail::convertChannelOrder(desc.channel_order.value_or(
sycl::ext::oneapi::experimental::detail::
get_image_default_channel_order(desc.num_channels)));
}

detail::image_mem_impl::image_mem_impl(const image_descriptor &desc,
Expand Down
27 changes: 24 additions & 3 deletions unified-runtime/source/adapters/level_zero/image_common.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -296,6 +296,11 @@ ur_result_t bindlessImagesCreateImpl(ur_context_handle_t hContext,
ZeStruct<ze_image_desc_t> ZeImageDesc;
UR_CALL(ur2zeImageDesc(pImageFormat, pImageDesc, ZeImageDesc));

ze_srgb_ext_desc_t ZeSrgbDesc{};
ZeSrgbDesc.stype = ZE_STRUCTURE_TYPE_SRGB_EXT_DESC;
ZeSrgbDesc.pNext = nullptr;
ZeSrgbDesc.sRGB = true;

ZeStruct<ze_image_bindless_exp_desc_t> BindlessDesc;
BindlessDesc.flags = ZE_IMAGE_BINDLESS_EXP_FLAG_BINDLESS;
ZeImageDesc.pNext = &BindlessDesc;
Expand All @@ -305,8 +310,15 @@ ur_result_t bindlessImagesCreateImpl(ur_context_handle_t hContext,
if (Sampled) {
ze_api_version_t ZeApiVersion = hContext->getPlatform()->ZeApiVersion;
UR_CALL(ur2zeSamplerDesc(ZeApiVersion, pSamplerDesc, ZeSamplerDesc));
BindlessDesc.pNext = &ZeSamplerDesc;
if (pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_SRGBA) {
ZeSrgbDesc.pNext = reinterpret_cast<ze_base_desc_t *>(&ZeSamplerDesc);
BindlessDesc.pNext = &ZeSrgbDesc;
} else {
BindlessDesc.pNext = &ZeSamplerDesc;
}
BindlessDesc.flags |= ZE_IMAGE_BINDLESS_EXP_FLAG_SAMPLED_IMAGE;
} else if (pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_SRGBA) {
BindlessDesc.pNext = &ZeSrgbDesc;
}

v2::raii::ze_image_handle_t ZeImage;
Expand Down Expand Up @@ -597,7 +609,8 @@ ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat,
case UR_IMAGE_CHANNEL_ORDER_RGBX:
case UR_IMAGE_CHANNEL_ORDER_BGRA:
case UR_IMAGE_CHANNEL_ORDER_ARGB:
case UR_IMAGE_CHANNEL_ORDER_ABGR: {
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;
Expand Down Expand Up @@ -658,6 +671,7 @@ ur_result_t ur2zeImageDesc(const ur_image_format_t *ImageFormat,
ZE_IMAGE_FORMAT_SWIZZLE_X, ZE_IMAGE_FORMAT_SWIZZLE_1};
break;
case UR_IMAGE_CHANNEL_ORDER_RGBA:
case UR_IMAGE_CHANNEL_ORDER_SRGBA:
Comment thread
juanchuletas marked this conversation as resolved.
ZeFormatDesc = {ZeImageFormatLayout, ZeImageFormatType,
ZE_IMAGE_FORMAT_SWIZZLE_R, ZE_IMAGE_FORMAT_SWIZZLE_G,
ZE_IMAGE_FORMAT_SWIZZLE_B, ZE_IMAGE_FORMAT_SWIZZLE_A};
Expand Down Expand Up @@ -1105,10 +1119,17 @@ ur_result_t urBindlessImagesImageAllocateExp(

ZeStruct<ze_image_desc_t> ZeImageDesc;
UR_CALL(ur2zeImageDesc(pImageFormat, pImageDesc, ZeImageDesc));
ze_srgb_ext_desc_t ZeSrgbDesc{};
ZeSrgbDesc.stype = ZE_STRUCTURE_TYPE_SRGB_EXT_DESC;
ZeSrgbDesc.pNext = nullptr;
ZeSrgbDesc.sRGB = true;

ze_image_bindless_exp_desc_t ZeImageBindlessDesc;
ZeImageBindlessDesc.stype = ZE_STRUCTURE_TYPE_BINDLESS_IMAGE_EXP_DESC;
ZeImageBindlessDesc.pNext = nullptr;
ZeImageBindlessDesc.pNext =
(pImageFormat->channelOrder == UR_IMAGE_CHANNEL_ORDER_SRGBA)
? static_cast<void *>(&ZeSrgbDesc)
: nullptr;
ZeImageBindlessDesc.flags = ZE_IMAGE_BINDLESS_EXP_FLAG_BINDLESS;
ZeImageDesc.pNext = &ZeImageBindlessDesc;

Expand Down