diff --git a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc index ba920c485a176..3ad0ca825cc58 100644 --- a/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc +++ b/sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc @@ -191,23 +191,29 @@ struct image_descriptor { image_type type{image_type::standard}; unsigned int num_levels{1}; unsigned int array_size{1}; + unsigned int num_samples{0}; + size_t row_pitch{0}; + size_t slice_pitch{0}; image_descriptor() = default; - 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(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, + unsigned int num_samples = 0); - 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(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, + unsigned int num_samples = 0, size_t row_pitch = 0, size_t slice_pitch = 0); - 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); + 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, + unsigned int num_samples = 0, size_t row_pitch = 0, size_t slice_pitch = 0); image_descriptor get_mip_level_desc(unsigned int level) const; @@ -279,6 +285,36 @@ type of the returned `image_descriptor` will be `image_type::standard`. Only array image types support more than one array layer. +The `num_samples` member specifies the number of samples per pixel for +multisampled images, and defaults to `0` (single-sample). It is only meaningful +on backends that support multisampled images. + +The `row_pitch` and `slice_pitch` members specify, in bytes, the stride between +successive rows and slices of pixel data respectively. They default to `0`, +which is interpreted as a tightly-packed layout (rows are +`width * num_channels * sizeof(channel)` bytes, slices are `row_pitch * height` +bytes). Non-zero values allow callers to describe images whose backing memory +has additional padding to meet device alignment requirements. `row_pitch` is +meaningful only for 2D and 3D images. `slice_pitch` is meaningful only for 3D +images and for 2D image arrays (where each array layer is a "slice"). Both +members are accepted on the `range<2>` constructor to cover the 2D-array case +uniformly. + +The `verify` member function will reject descriptors that set these members on +image shapes where they have no meaning (a non-zero `row_pitch` on a 1D image, +or a non-zero `slice_pitch` on anything other than a 3D image or a 2D image +array), and will also reject a non-zero `slice_pitch` that is smaller than +`row_pitch * height`. + +[NOTE] +==== +When an API also accepts an explicit pitch parameter (for example, +`ext_oneapi_copy` overloads that take `SrcPitch`/`DestPitch` arguments), the +explicit parameter takes precedence when it is non-zero. The `row_pitch` value +carried in the `image_descriptor` is used only when no non-zero explicit pitch +is supplied. +==== + ==== Querying image support Not all devices support all combinations of image channel type, the number of diff --git a/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp index 8d2fdd0a03eb1..ae6662964a20d 100644 --- a/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp +++ b/sycl/include/sycl/ext/oneapi/bindless_images_descriptor.hpp @@ -58,38 +58,52 @@ struct image_descriptor { 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}; + // The following fields map to ur_image_desc_t on the bindless-images path + // only. Core UR image creation requires numMipLevel/numSamples to be 0; the + // bindless-images adapters (e.g. L0's ze_image_desc_t::miplevels) honor them. + unsigned int num_levels{1}; // -- ur_image_desc_t::numMipLevel (bindless) + unsigned int array_size{1}; // -- ur_image_desc_t::arraySize + unsigned int num_samples{0}; // -- ur_image_desc_t::numSamples (bindless) + size_t row_pitch{0}; // -- ur_image_desc_t::rowPitch + size_t slice_pitch{0}; // -- ur_image_desc_t::slicePitch image_descriptor() = default; 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) + unsigned int num_levels = 1, unsigned int array_size = 1, + unsigned int num_samples = 0) : 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) { + array_size(array_size), num_samples(num_samples) { verify(); } 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) + unsigned int num_levels = 1, unsigned int array_size = 1, + unsigned int num_samples = 0, size_t row_pitch = 0, + size_t slice_pitch = 0) : 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) { + array_size(array_size), num_samples(num_samples), row_pitch(row_pitch), + slice_pitch(slice_pitch) { verify(); } 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) + unsigned int num_levels = 1, unsigned int array_size = 1, + unsigned int num_samples = 0, size_t row_pitch = 0, + size_t slice_pitch = 0) : width(dims[0]), height(dims[1]), depth(dims[2]), num_channels(num_channels), channel_type(channel_type), type(type), - num_levels(num_levels), array_size(array_size) { + num_levels(num_levels), array_size(array_size), + num_samples(num_samples), row_pitch(row_pitch), + slice_pitch(slice_pitch) { verify(); }; @@ -129,6 +143,30 @@ struct image_descriptor { "Images must have 1, 2, 3, or 4 channels."); } + // row_pitch is meaningful only for 2D and 3D images. + if (this->row_pitch != 0 && this->height == 0) { + throw sycl::exception( + sycl::errc::invalid, + "row_pitch is meaningful only for 2D and 3D images."); + } + + // slice_pitch is meaningful only for 3D images and 2D image arrays. + const bool is_3d = this->depth != 0; + const bool is_2d_array = + this->type == image_type::array && this->height != 0; + if (this->slice_pitch != 0 && !is_3d && !is_2d_array) { + throw sycl::exception( + sycl::errc::invalid, + "slice_pitch is meaningful only for 3D images and 2D image arrays."); + } + + // A non-zero slice_pitch must describe non-overlapping slices. + if (this->slice_pitch != 0 && this->row_pitch != 0 && this->height != 0 && + this->slice_pitch < this->row_pitch * this->height) { + throw sycl::exception(sycl::errc::invalid, + "slice_pitch must be at least row_pitch * height."); + } + switch (this->type) { case image_type::standard: if (this->array_size > 1) { diff --git a/sycl/source/detail/bindless_images.cpp b/sycl/source/detail/bindless_images.cpp index bdf0007b14801..6768f7f45667c 100644 --- a/sycl/source/detail/bindless_images.cpp +++ b/sycl/source/detail/bindless_images.cpp @@ -21,8 +21,15 @@ namespace sycl { inline namespace _V1 { namespace ext::oneapi::experimental { -void populate_ur_structs(const image_descriptor &desc, ur_image_desc_t &urDesc, - ur_image_format_t &urFormat, size_t pitch = 0) { +// If the caller passes a non-null UserPitchMarker, and the image_descriptor's +// pitch fields were set by the user, chain the marker off urDesc.pNext so +// adapters that can honor a user-supplied pitch (e.g. L0) know the values in +// urDesc.rowPitch/slicePitch are a user request rather than a driver-computed +// pitch bookkept back through the descriptor. +void populate_ur_structs( + const image_descriptor &desc, ur_image_desc_t &urDesc, + ur_image_format_t &urFormat, size_t pitch = 0, + ur_exp_image_user_pitch_desc_t *UserPitchMarker = nullptr) { urDesc = {}; urDesc.stype = UR_STRUCTURE_TYPE_IMAGE_DESC; urDesc.width = desc.width; @@ -30,7 +37,6 @@ void populate_ur_structs(const image_descriptor &desc, ur_image_desc_t &urDesc, urDesc.depth = desc.depth; if (desc.array_size > 1) { - // Image array or cubemap urDesc.type = desc.type == image_type::cubemap ? UR_MEM_TYPE_IMAGE_CUBEMAP_EXP : desc.type == image_type::gather ? UR_MEM_TYPE_IMAGE_GATHER_EXP @@ -42,11 +48,22 @@ void populate_ur_structs(const image_descriptor &desc, ur_image_desc_t &urDesc, : UR_MEM_TYPE_IMAGE1D); } - urDesc.rowPitch = pitch; + // Use 'pitch' arg if provided, otherwise use descriptor row_pitch + urDesc.rowPitch = (pitch != 0) ? pitch : desc.row_pitch; + urDesc.arraySize = desc.array_size; - urDesc.slicePitch = 0; + urDesc.slicePitch = desc.slice_pitch; urDesc.numMipLevel = (desc.type == image_type::mipmap) ? desc.num_levels : 0; - urDesc.numSamples = 0; + urDesc.numSamples = desc.num_samples; + + if (UserPitchMarker != nullptr && + (desc.row_pitch != 0 || desc.slice_pitch != 0)) { + *UserPitchMarker = {}; + UserPitchMarker->stype = UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC; + UserPitchMarker->rowPitch = desc.row_pitch; + UserPitchMarker->slicePitch = desc.slice_pitch; + urDesc.pNext = UserPitchMarker; + } urFormat = {}; urFormat.channelType = sycl::detail::convertChannelType(desc.channel_type); @@ -241,7 +258,8 @@ create_image(image_mem_handle memHandle, const image_descriptor &desc, ur_image_desc_t urDesc; ur_image_format_t urFormat; - populate_ur_structs(desc, urDesc, urFormat); + ur_exp_image_user_pitch_desc_t UserPitchMarker; + populate_ur_structs(desc, urDesc, urFormat, /*pitch=*/0, &UserPitchMarker); // Call impl. ur_exp_image_native_handle_t urImageHandle = 0; @@ -369,7 +387,8 @@ create_image(void *devPtr, size_t pitch, const bindless_image_sampler &sampler, ur_image_desc_t urDesc; ur_image_format_t urFormat; - populate_ur_structs(desc, urDesc, urFormat, pitch); + ur_exp_image_user_pitch_desc_t UserPitchMarker; + populate_ur_structs(desc, urDesc, urFormat, pitch, &UserPitchMarker); // Call impl. ur_exp_image_native_handle_t urImageHandle = 0; diff --git a/sycl/source/detail/image_impl.cpp b/sycl/source/detail/image_impl.cpp index 8ffc6181b127b..46e975c4bb624 100644 --- a/sycl/source/detail/image_impl.cpp +++ b/sycl/source/detail/image_impl.cpp @@ -326,8 +326,7 @@ void *image_impl::allocateMem(context_impl *Context, bool InitFromUserData, BaseT::determineHostPtr(Context, InitFromUserData, HostPtr, HostPtrReadOnly); ur_image_desc_t Desc = getImageDesc(HostPtr != nullptr); - assert(checkImageDesc(Desc, Context, HostPtr) && - "The check an image desc failed."); + assert(checkImageDesc(Desc, Context) && "The check an image desc failed."); ur_image_format_t Format = getImageFormat(); assert(checkImageFormat(Format, Context) && @@ -340,7 +339,7 @@ void *image_impl::allocateMem(context_impl *Context, bool InitFromUserData, } bool image_impl::checkImageDesc(const ur_image_desc_t &Desc, - context_impl *Context, void *UserPtr) { + context_impl *Context) { devices_range Devices = Context ? Context->getDevices() : devices_range{}; if (checkAny(Desc.type, UR_MEM_TYPE_IMAGE1D, UR_MEM_TYPE_IMAGE1D_ARRAY, UR_MEM_TYPE_IMAGE2D_ARRAY, UR_MEM_TYPE_IMAGE2D) && @@ -378,22 +377,6 @@ bool image_impl::checkImageDesc(const ur_image_desc_t &Desc, "For a 3D image, the depth must be a Value >= 1 and <= " "info::device::image2d_max_depth"); - if ((nullptr == UserPtr) && (0 != Desc.rowPitch)) - throw exception(make_error_code(errc::invalid), - "The row_pitch must be 0 if host_ptr is nullptr."); - - if ((nullptr == UserPtr) && (0 != Desc.slicePitch)) - throw exception(make_error_code(errc::invalid), - "The slice_pitch must be 0 if host_ptr is nullptr."); - - if (0 != Desc.numMipLevel) - throw exception(make_error_code(errc::invalid), - "The mip_levels must be 0."); - - if (0 != Desc.numSamples) - throw exception(make_error_code(errc::invalid), - "The num_samples must be 0."); - return true; } diff --git a/sycl/source/detail/image_impl.hpp b/sycl/source/detail/image_impl.hpp index 87405fd8d0720..c62a2e4d6d9ad 100644 --- a/sycl/source/detail/image_impl.hpp +++ b/sycl/source/detail/image_impl.hpp @@ -329,8 +329,7 @@ class image_impl final : public SYCLMemObjT { return Desc; } - bool checkImageDesc(const ur_image_desc_t &Desc, context_impl *Context, - void *UserPtr); + bool checkImageDesc(const ur_image_desc_t &Desc, context_impl *Context); ur_image_format_t getImageFormat() { ur_image_format_t Format = {}; diff --git a/sycl/source/handler.cpp b/sycl/source/handler.cpp index 6c79c6dce743a..26e77a0c9e683 100644 --- a/sycl/source/handler.cpp +++ b/sycl/source/handler.cpp @@ -235,6 +235,12 @@ fill_image_desc(const ext::oneapi::experimental::image_descriptor &ImgDesc) { UrDesc.height = ImgDesc.height; UrDesc.depth = ImgDesc.depth; UrDesc.arraySize = ImgDesc.array_size; + + UrDesc.rowPitch = ImgDesc.row_pitch; + UrDesc.slicePitch = ImgDesc.slice_pitch; + UrDesc.numSamples = ImgDesc.num_samples; + UrDesc.numMipLevel = ImgDesc.num_levels; + return UrDesc; } @@ -304,8 +310,12 @@ static void fill_copy_args( impl->MDstImageDesc.depth = DestExtent[2]; } - impl->MSrcImageDesc.rowPitch = SrcPitch; - impl->MDstImageDesc.rowPitch = DestPitch; + // Explicit pitch arg wins when non-zero; otherwise keep the descriptor's + // row_pitch already set by fill_image_desc(). + if (SrcPitch != 0) + impl->MSrcImageDesc.rowPitch = SrcPitch; + if (DestPitch != 0) + impl->MDstImageDesc.rowPitch = DestPitch; } static void @@ -319,8 +329,35 @@ fill_copy_args(detail::handler_impl *impl, sycl::range<3> DestExtent = {0, 0, 0}, sycl::range<3> CopyExtent = {0, 0, 0}) { - size_t SrcPitch = SrcExtent[0] * Desc.num_channels * get_channel_size(Desc); - size_t DestPitch = DestExtent[0] * Desc.num_channels * get_channel_size(Desc); + // The UR adapters interpret pSrcImageDesc->rowPitch / pDstImageDesc->rowPitch + // as the row stride of whichever side is host/USM memory (see e.g. + // unified-runtime/source/adapters/level_zero/image_common.cpp). Host memory + // is always tightly packed, so for MEM_TO_IMAGE / IMAGE_TO_MEM we must not + // let the image descriptor's row_pitch (a device stride) reach the memory + // side. Derive the host pitch from the caller-supplied extent when present, + // otherwise from the descriptor width. + auto TightHostPitch = [&](sycl::range<3> Extent) { + size_t W = Extent[0] != 0 ? Extent[0] : Desc.width; + return W * Desc.num_channels * get_channel_size(Desc); + }; + + size_t SrcPitch = 0; + size_t DestPitch = 0; + switch (ImageCopyInputTypes) { + case UR_EXP_IMAGE_COPY_INPUT_TYPES_MEM_TO_IMAGE: + SrcPitch = TightHostPitch(SrcExtent); + break; + case UR_EXP_IMAGE_COPY_INPUT_TYPES_IMAGE_TO_MEM: + DestPitch = TightHostPitch(DestExtent); + break; + default: + // MEM_TO_MEM and IMAGE_TO_IMAGE do not reach this overload from the + // handler; preserve the pre-existing extent-driven derivation as a + // safe fallback. + SrcPitch = SrcExtent[0] * Desc.num_channels * get_channel_size(Desc); + DestPitch = DestExtent[0] * Desc.num_channels * get_channel_size(Desc); + break; + } fill_copy_args(impl, Desc, Desc, ImageCopyFlags, ImageCopyInputTypes, SrcPitch, DestPitch, SrcOffset, SrcExtent, DestOffset, diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp index 7fa71442bb168..f28b4e72fcc14 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_setup.hpp @@ -313,6 +313,15 @@ size_t getRowPitch(VulkanContext &ctx, VkImage image) { return layout.rowPitch; } +// vkGetImageSubresourceLayout is only valid for VK_IMAGE_TILING_LINEAR images; +// call only when the VkImage was created with linear tiling. +size_t getSlicePitch(VulkanContext &ctx, VkImage image) { + VkSubresourceLayout layout; + VkImageSubresource subResource{VK_IMAGE_ASPECT_COLOR_BIT, 0, 0}; + vkGetImageSubresourceLayout(ctx.device, image, &subResource, &layout); + return layout.depthPitch; +} + inline uint32_t findMemoryType(VkPhysicalDevice physicalDevice, uint32_t typeFilter, VkMemoryPropertyFlags properties) { diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp index 25ef04c0c4ddc..bdecc738b7467 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_2d_arithmetic.cpp @@ -344,10 +344,10 @@ int runTest( channels, // num_channels syclType, // channel_type syclexp::image_type::standard, // type (default) - 1 //, // num_levels (default) - // 1, // array_size (default) - // 0, // num_samples (default) - // pitchA // pitch + 1, // num_levels (default) + 1, // array_size (default) + 0, // num_samples (default) + pitchA // row_pitch (0 = tightly packed) ); auto imgMemA = syclexp::map_external_image_memory( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp index cbc0a72abeb07..925ba602e0947 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_2d.cpp @@ -289,10 +289,20 @@ int runTest( ? syclOverride.value() : getSyclChannelType(); + // When the Vulkan image was created with LINEAR tiling, its row stride is + // not guaranteed to match SYCL's tightly-packed default. Query Vulkan for + // the actual row pitch and forward it to the image_descriptor so adapters + // that can honor a user-supplied pitch (e.g. L0) use the right stride. + // vkGetImageSubresourceLayout is only defined for LINEAR tiling; leave the + // pitch at 0 (tightly-packed) for OPTIMAL tiling. + size_t rowPitch = useLinear ? getRowPitch(vkCtx, imgRes.image) : 0; + // bindless image use (x,y,z) order, // differening from SYCL 2020 "fastest incrementing" convention. - syclexp::image_descriptor imgDesc(sycl::range<2>(width, height), channels, - syclType); + syclexp::image_descriptor imgDesc( + sycl::range<2>(width, height), channels, syclType, + syclexp::image_type::standard, /*num_levels=*/1, /*array_size=*/1, + /*num_samples=*/0, /*row_pitch=*/rowPitch); // Map external memory syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp index 94889fac53dfd..6f62453eece48 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_read_3d.cpp @@ -188,10 +188,22 @@ int runTest( ? syclOverride.value() : getSyclChannelType(); + // When the Vulkan image was created with LINEAR tiling, its row/slice + // strides are not guaranteed to match SYCL's tightly-packed default. Query + // Vulkan for the actual pitches and forward them to the image_descriptor + // so adapters that can honor a user-supplied pitch (e.g. L0) use the right + // strides. vkGetImageSubresourceLayout is only defined for LINEAR tiling; + // leave the pitches at 0 (tightly-packed) for OPTIMAL tiling. + size_t rowPitch = useLinear ? getRowPitch(vkCtx, imgRes.image) : 0; + size_t slicePitch = useLinear ? getSlicePitch(vkCtx, imgRes.image) : 0; + // bindless image use (x,y,z) order, // differening from SYCL 2020 "fastest incrementing" convention. - syclexp::image_descriptor imgDesc(sycl::range<3>(width, height, depth), - channels, syclType); + syclexp::image_descriptor imgDesc( + sycl::range<3>(width, height, depth), channels, syclType, + syclexp::image_type::standard, /*num_levels=*/1, /*array_size=*/1, + /*num_samples=*/0, /*row_pitch=*/rowPitch, + /*slice_pitch=*/slicePitch); // Map external memory syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp index ed3a66a8d6955..e79f68407a98a 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_2d_unsampled.cpp @@ -248,10 +248,21 @@ int runTest( sycl::image_channel_type syclType = syclOverride.has_value() ? syclOverride.value() : getSyclChannelType(); + + // When the Vulkan image was created with LINEAR tiling, its row stride is + // not guaranteed to match SYCL's tightly-packed default. Query Vulkan for + // the actual row pitch and forward it to the image_descriptor so adapters + // that can honor a user-supplied pitch (e.g. L0) use the right stride. + // vkGetImageSubresourceLayout is only defined for LINEAR tiling; leave the + // pitch at 0 (tightly-packed) for OPTIMAL tiling. + size_t rowPitch = useLinear ? getRowPitch(vkCtx, imgRes.image) : 0; + // bindless image ranges use (x,y,z) order, // differening from SYCL 2020 "fastest incrementing" convention. - syclexp::image_descriptor imgDesc(sycl::range<2>(width, height), channels, - syclType); + syclexp::image_descriptor imgDesc( + sycl::range<2>(width, height), channels, syclType, + syclexp::image_type::standard, /*num_levels=*/1, /*array_size=*/1, + /*num_samples=*/0, /*row_pitch=*/rowPitch); syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory( extMem, imgDesc, q.get_device(), q.get_context()); syclexp::unsampled_image_handle unsampledHandle = syclexp::create_image( diff --git a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp index 3edc11a7142d4..381d0caa1f717 100644 --- a/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp +++ b/sycl/test-e2e/bindless_images/vulkan_interop/vulkan_sycl_image_interop_write_3d_unsampled.cpp @@ -199,10 +199,23 @@ int runTest( sycl::image_channel_type syclType = syclOverride.has_value() ? syclOverride.value() : getSyclChannelType(); + + // When the Vulkan image was created with LINEAR tiling, its row/slice + // strides are not guaranteed to match SYCL's tightly-packed default. Query + // Vulkan for the actual pitches and forward them to the image_descriptor + // so adapters that can honor a user-supplied pitch (e.g. L0) use the right + // strides. vkGetImageSubresourceLayout is only defined for LINEAR tiling; + // leave the pitches at 0 (tightly-packed) for OPTIMAL tiling. + size_t rowPitch = useLinear ? getRowPitch(vkCtx, imgRes.image) : 0; + size_t slicePitch = useLinear ? getSlicePitch(vkCtx, imgRes.image) : 0; + // bindless image ranges use (x,y,z) order, // differening from SYCL 2020 "fastest incrementing" convention. - syclexp::image_descriptor imgDesc(sycl::range<3>(width, height, depth), - channels, syclType); + syclexp::image_descriptor imgDesc( + sycl::range<3>(width, height, depth), channels, syclType, + syclexp::image_type::standard, /*num_levels=*/1, /*array_size=*/1, + /*num_samples=*/0, /*row_pitch=*/rowPitch, + /*slice_pitch=*/slicePitch); syclexp::image_mem_handle devHandle = syclexp::map_external_image_memory( extMem, imgDesc, q.get_device(), q.get_context()); syclexp::unsampled_image_handle unsampledHandle = syclexp::create_image( diff --git a/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt b/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt index cbc028cb65903..8fa0f6133daca 100644 --- a/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt +++ b/sycl/unittests/Extensions/BindlessImages/CMakeLists.txt @@ -1,4 +1,5 @@ add_sycl_unittest(BindlessImagesExtensionTests OBJECT + ImageDescriptors.cpp MipLevels.cpp Semaphores.cpp ) diff --git a/sycl/unittests/Extensions/BindlessImages/ImageDescriptors.cpp b/sycl/unittests/Extensions/BindlessImages/ImageDescriptors.cpp new file mode 100644 index 0000000000000..f8e07e0f1431f --- /dev/null +++ b/sycl/unittests/Extensions/BindlessImages/ImageDescriptors.cpp @@ -0,0 +1,229 @@ +#include +#include + +#include +#include +#include +#include +#include + +namespace syclexp = sycl::ext::oneapi::experimental; + +thread_local size_t ExpectedRowPitch = 0; +thread_local size_t ExpectedSlicePitch = 0; +thread_local uint32_t ExpectedNumSamples = 0; +thread_local int MapExternalArrayCallCounter = 0; + +thread_local size_t ImageCopySrcRowPitch = 0; +thread_local size_t ImageCopyDstRowPitch = 0; +thread_local size_t ImageCopyDstSlicePitch = 0; +thread_local int ImageCopyCallCounter = 0; + +// ----- +// Mocks +// ----- +inline ur_result_t urBindlessImagesMapExternalArrayExp_replace(void *pParams) { + ++MapExternalArrayCallCounter; + auto Params = + *reinterpret_cast( + pParams); + + const ur_image_desc_t *urDesc = *Params.ppImageDesc; + + EXPECT_EQ(urDesc->rowPitch, ExpectedRowPitch); + EXPECT_EQ(urDesc->slicePitch, ExpectedSlicePitch); + EXPECT_EQ(urDesc->numSamples, ExpectedNumSamples); + + if (Params.pphImageMem && *Params.pphImageMem) { + **Params.pphImageMem = + mock::createDummyHandle(); + } + return UR_RESULT_SUCCESS; +} + +inline ur_result_t urBindlessImagesImageCopyExp_replace(void *pParams) { + ++ImageCopyCallCounter; + auto Params = + *reinterpret_cast(pParams); + + const ur_image_desc_t *urSrcDesc = *Params.ppSrcImageDesc; + const ur_image_desc_t *urDstDesc = *Params.ppDstImageDesc; + ImageCopySrcRowPitch = urSrcDesc->rowPitch; + ImageCopyDstRowPitch = urDstDesc->rowPitch; + ImageCopyDstSlicePitch = urDstDesc->slicePitch; + + return UR_RESULT_SUCCESS; +} + +// ----------------------------------------------------------------------------- +// Tests +// ----------------------------------------------------------------------------- +TEST(BindlessImagesExtensionTests, ImageDescriptorPropagatesLayout) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_replace_callback( + "urBindlessImagesMapExternalArrayExp", + &urBindlessImagesMapExternalArrayExp_replace); + + sycl::queue Q; + + MapExternalArrayCallCounter = 0; + ExpectedRowPitch = 1024; + ExpectedSlicePitch = 0; + ExpectedNumSamples = 4; + + syclexp::image_descriptor Desc(sycl::range<2>{32, 32}, 4, + sycl::image_channel_type::fp32, + syclexp::image_type::standard, 1, 1, + ExpectedNumSamples, ExpectedRowPitch); + + syclexp::external_mem_descriptor ExtMemDesc{ + {123}, syclexp::external_mem_handle_type::opaque_fd, 0}; + + auto MemHandle = syclexp::import_external_memory(ExtMemDesc, Q.get_device(), + Q.get_context()); + + try { + auto ImgHandle = syclexp::map_external_image_memory(MemHandle, Desc, Q); + (void)ImgHandle; + } catch (const sycl::exception &e) { + FAIL() << "Caught unexpected SYCL exception: " << e.what(); + } + + EXPECT_EQ(MapExternalArrayCallCounter, 1); +} + +// Regression test: descriptor row_pitch must propagate through the +// ext_oneapi_copy path when the overload has no explicit pitch argument. +// Previously fill_copy_args unconditionally clobbered rowPitch with the +// (derived-from-extent) SrcPitch/DestPitch values, dropping the descriptor's +// pitch. +TEST(BindlessImagesExtensionTests, ImageDescriptorCopyPropagatesPitch) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_replace_callback( + "urBindlessImagesImageCopyExp", &urBindlessImagesImageCopyExp_replace); + + sycl::queue Q; + + ImageCopyCallCounter = 0; + ImageCopySrcRowPitch = 0; + ImageCopyDstRowPitch = 0; + ImageCopyDstSlicePitch = 0; + + constexpr size_t DescRowPitch = 4096; + // 32 * 4 channels * sizeof(float) = 512 + constexpr size_t TightHostPitch = 32 * 4 * sizeof(float); + + syclexp::image_descriptor Desc( + sycl::range<2>{32, 32}, 4, sycl::image_channel_type::fp32, + syclexp::image_type::standard, 1, 1, 0, DescRowPitch); + + syclexp::image_mem_handle DstHandle = + syclexp::alloc_image_mem(Desc, Q.get_device(), Q.get_context()); + + std::vector HostSrc(32 * 32 * 4, 0.0f); + + try { + Q.ext_oneapi_copy(HostSrc.data(), DstHandle, Desc); + Q.wait(); + } catch (const sycl::exception &e) { + syclexp::free_image_mem(DstHandle, syclexp::image_type::standard, Q); + FAIL() << "Caught unexpected SYCL exception: " << e.what(); + } + + EXPECT_EQ(ImageCopyCallCounter, 1); + // Image side (dst) gets the descriptor's row_pitch. + EXPECT_EQ(ImageCopyDstRowPitch, DescRowPitch); + // Memory side (src) must be the tight host stride, NOT the descriptor's + // row_pitch. UR adapters use pSrcImageDesc->rowPitch as the host stride. + EXPECT_EQ(ImageCopySrcRowPitch, TightHostPitch); + + syclexp::free_image_mem(DstHandle, syclexp::image_type::standard, Q); +} + +TEST(BindlessImagesExtensionTests, ImageDescriptorPitchValidation) { + // row_pitch on a 1D image is rejected. The 1D ctor doesn't take row_pitch, + // so we default-construct and set the field directly. + { + syclexp::image_descriptor Desc; + Desc.width = 32; + Desc.num_channels = 4; + Desc.channel_type = sycl::image_channel_type::fp32; + Desc.row_pitch = 256; + EXPECT_THROW(Desc.verify(), sycl::exception); + } + + // slice_pitch on a plain 2D standard image is rejected. + { + syclexp::image_descriptor Desc; + Desc.width = 32; + Desc.height = 32; + Desc.num_channels = 4; + Desc.channel_type = sycl::image_channel_type::fp32; + Desc.slice_pitch = 8192; + EXPECT_THROW(Desc.verify(), sycl::exception); + } + + // slice_pitch on a 2D image array is accepted. + { + syclexp::image_descriptor Desc; + Desc.width = 32; + Desc.height = 32; + Desc.num_channels = 4; + Desc.channel_type = sycl::image_channel_type::fp32; + Desc.type = syclexp::image_type::array; + Desc.array_size = 4; + Desc.row_pitch = 512; + Desc.slice_pitch = 512 * 32; + EXPECT_NO_THROW(Desc.verify()); + } + + // slice_pitch smaller than row_pitch * height is rejected (3D case). + { + syclexp::image_descriptor Desc; + Desc.width = 16; + Desc.height = 16; + Desc.depth = 4; + Desc.num_channels = 1; + Desc.channel_type = sycl::image_channel_type::fp32; + Desc.row_pitch = 512; + Desc.slice_pitch = 2048; // < 512 * 16 = 8192 + EXPECT_THROW(Desc.verify(), sycl::exception); + } +} + +TEST(BindlessImagesExtensionTests, ImageDescriptorPropagatesSlicePitch) { + sycl::unittest::UrMock<> Mock; + mock::getCallbacks().set_replace_callback( + "urBindlessImagesMapExternalArrayExp", + &urBindlessImagesMapExternalArrayExp_replace); + + sycl::queue Q; + + MapExternalArrayCallCounter = 0; + ExpectedRowPitch = 512; + // slice_pitch is the byte stride between successive slices; it must be at + // least row_pitch * height (512 * 16 = 8192) to describe a non-overlapping + // layout. + ExpectedSlicePitch = 8192; + ExpectedNumSamples = 1; + + syclexp::image_descriptor Desc( + sycl::range<3>{16, 16, 4}, 1, sycl::image_channel_type::fp32, + syclexp::image_type::standard, 1, 1, ExpectedNumSamples, ExpectedRowPitch, + ExpectedSlicePitch); + + syclexp::external_mem_descriptor ExtMemDesc{ + {456}, syclexp::external_mem_handle_type::opaque_fd, 0}; + + auto MemHandle = syclexp::import_external_memory(ExtMemDesc, Q.get_device(), + Q.get_context()); + + try { + auto ImgHandle = syclexp::map_external_image_memory(MemHandle, Desc, Q); + (void)ImgHandle; + } catch (const sycl::exception &e) { + FAIL() << "Caught unexpected SYCL exception: " << e.what(); + } + + EXPECT_EQ(MapExternalArrayCallCounter, 1); +} diff --git a/unified-runtime/include/unified-runtime/ur_api.h b/unified-runtime/include/unified-runtime/ur_api.h index 4624a5c9acb0d..00c195853de27 100644 --- a/unified-runtime/include/unified-runtime/ur_api.h +++ b/unified-runtime/include/unified-runtime/ur_api.h @@ -649,6 +649,8 @@ typedef enum ur_structure_type_t { UR_STRUCTURE_TYPE_EXP_IMAGE_COPY_REGION = 0x2007, /// ::ur_exp_win32_name_t UR_STRUCTURE_TYPE_EXP_WIN32_NAME = 0x2008, + /// ::ur_exp_image_user_pitch_desc_t + UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC = 0x2009, /// ::ur_exp_async_usm_alloc_properties_t UR_STRUCTURE_TYPE_EXP_ASYNC_USM_ALLOC_PROPERTIES = 0x2050, /// ::ur_exp_enqueue_native_command_properties_t @@ -10208,6 +10210,33 @@ typedef struct ur_exp_image_copy_region_t { } ur_exp_image_copy_region_t; +/////////////////////////////////////////////////////////////////////////////// +/// @brief User-declared image pitch signal for image creation. +/// +/// @details +/// - Chain this off ::ur_image_desc_t::pNext when the caller has explicitly +/// declared a custom row/slice pitch (e.g. via SYCL's +/// image_descriptor.row_pitch / slice_pitch). The presence of this +/// structure disambiguates a user-supplied pitch from a driver-computed +/// pitch that some SYCL entry points bookkeep back through +/// ::ur_image_desc_t::rowPitch. Adapters that can override the driver's +/// computed pitch (e.g. Level Zero via ze_custom_pitch_exp_desc_t) should +/// key off the presence of this struct rather than +/// ::ur_image_desc_t::rowPitch alone. Adapters that always consume +/// ::ur_image_desc_t::rowPitch as-is may ignore it. +typedef struct ur_exp_image_user_pitch_desc_t { + /// [in] type of this structure, must be + /// ::UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC + ur_structure_type_t stype; + /// [in][optional] pointer to extension-specific structure + const void *pNext; + /// [in] user-declared row pitch, in bytes + size_t rowPitch; + /// [in] user-declared slice pitch, in bytes + size_t slicePitch; + +} ur_exp_image_user_pitch_desc_t; + /////////////////////////////////////////////////////////////////////////////// /// @brief USM allocate pitched memory /// diff --git a/unified-runtime/include/unified-runtime/ur_print.h b/unified-runtime/include/unified-runtime/ur_print.h index b78b93d23f090..01ff032df465a 100644 --- a/unified-runtime/include/unified-runtime/ur_print.h +++ b/unified-runtime/include/unified-runtime/ur_print.h @@ -1392,6 +1392,16 @@ UR_APIEXPORT ur_result_t UR_APICALL urPrintExpImageCopyRegion( const struct ur_exp_image_copy_region_t params, char *buffer, const size_t buff_size, size_t *out_size); +/////////////////////////////////////////////////////////////////////////////// +/// @brief Print ur_exp_image_user_pitch_desc_t struct +/// @returns +/// - ::UR_RESULT_SUCCESS +/// - ::UR_RESULT_ERROR_INVALID_SIZE +/// - `buff_size < out_size` +UR_APIEXPORT ur_result_t UR_APICALL urPrintExpImageUserPitchDesc( + const struct ur_exp_image_user_pitch_desc_t params, char *buffer, + const size_t buff_size, size_t *out_size); + /////////////////////////////////////////////////////////////////////////////// /// @brief Print ur_exp_program_flag_t enum /// @returns diff --git a/unified-runtime/include/unified-runtime/ur_print.hpp b/unified-runtime/include/unified-runtime/ur_print.hpp index 5570ed1d7b0be..b60db495713b1 100644 --- a/unified-runtime/include/unified-runtime/ur_print.hpp +++ b/unified-runtime/include/unified-runtime/ur_print.hpp @@ -609,6 +609,9 @@ inline std::ostream &operator<<( inline std::ostream & operator<<(std::ostream &os, [[maybe_unused]] const struct ur_exp_image_copy_region_t params); +inline std::ostream & +operator<<(std::ostream &os, + [[maybe_unused]] const struct ur_exp_image_user_pitch_desc_t params); inline std::ostream &operator<<(std::ostream &os, enum ur_exp_program_flag_t value); inline std::ostream & @@ -1597,6 +1600,9 @@ inline std::ostream &operator<<(std::ostream &os, case UR_STRUCTURE_TYPE_EXP_WIN32_NAME: os << "UR_STRUCTURE_TYPE_EXP_WIN32_NAME"; break; + case UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC: + os << "UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC"; + break; case UR_STRUCTURE_TYPE_EXP_ASYNC_USM_ALLOC_PROPERTIES: os << "UR_STRUCTURE_TYPE_EXP_ASYNC_USM_ALLOC_PROPERTIES"; break; @@ -1940,6 +1946,12 @@ inline ur_result_t printStruct(std::ostream &os, const void *ptr) { printPtr(os, pstruct); } break; + case UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC: { + const ur_exp_image_user_pitch_desc_t *pstruct = + (const ur_exp_image_user_pitch_desc_t *)ptr; + printPtr(os, pstruct); + } break; + case UR_STRUCTURE_TYPE_EXP_ASYNC_USM_ALLOC_PROPERTIES: { const ur_exp_async_usm_alloc_properties_t *pstruct = (const ur_exp_async_usm_alloc_properties_t *)ptr; @@ -12516,6 +12528,37 @@ operator<<(std::ostream &os, const struct ur_exp_image_copy_region_t params) { return os; } /////////////////////////////////////////////////////////////////////////////// +/// @brief Print operator for the ur_exp_image_user_pitch_desc_t type +/// @returns +/// std::ostream & +inline std::ostream & +operator<<(std::ostream &os, + const struct ur_exp_image_user_pitch_desc_t params) { + os << "(struct ur_exp_image_user_pitch_desc_t){"; + + os << ".stype = "; + + os << (params.stype); + + os << ", "; + os << ".pNext = "; + + ur::details::printStruct(os, (params.pNext)); + + os << ", "; + os << ".rowPitch = "; + + os << (params.rowPitch); + + os << ", "; + os << ".slicePitch = "; + + os << (params.slicePitch); + + os << "}"; + return os; +} +/////////////////////////////////////////////////////////////////////////////// /// @brief Print operator for the ur_exp_program_flag_t type /// @returns /// std::ostream & diff --git a/unified-runtime/scripts/core/exp-bindless-images.yml b/unified-runtime/scripts/core/exp-bindless-images.yml index 387ebc9361932..47e92ce7c1f1c 100644 --- a/unified-runtime/scripts/core/exp-bindless-images.yml +++ b/unified-runtime/scripts/core/exp-bindless-images.yml @@ -148,6 +148,9 @@ etors: - name: EXP_WIN32_NAME desc: $x_exp_win32_name_t value: "0x2008" + - name: EXP_IMAGE_USER_PITCH_DESC + desc: $x_exp_image_user_pitch_desc_t + value: "0x2009" --- #-------------------------------------------------------------------------- type: enum extend: true @@ -374,6 +377,31 @@ members: name: copyExtent desc: "[in] the extent (region) of the image to copy" --- #-------------------------------------------------------------------------- +type: struct +desc: | + User-declared image pitch signal for image creation. +details: + - "Chain this off $x_image_desc_t::pNext when the caller has explicitly + declared a custom row/slice pitch (e.g. via SYCL's + image_descriptor.row_pitch / slice_pitch). The presence of this + structure disambiguates a user-supplied pitch from a driver-computed + pitch that some SYCL entry points bookkeep back through + $x_image_desc_t::rowPitch. Adapters that can override the driver's + computed pitch (e.g. Level Zero via ze_custom_pitch_exp_desc_t) should + key off the presence of this struct rather than $x_image_desc_t::rowPitch + alone. Adapters that always consume $x_image_desc_t::rowPitch as-is may + ignore it." +class: $xBindlessImages +name: $x_exp_image_user_pitch_desc_t +base: $x_base_desc_t +members: + - type: size_t + name: rowPitch + desc: "[in] user-declared row pitch, in bytes" + - type: size_t + name: slicePitch + desc: "[in] user-declared slice pitch, in bytes" +--- #-------------------------------------------------------------------------- type: function desc: "USM allocate pitched memory" class: $xUSM diff --git a/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.cpp b/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.cpp index 8e452447c2cad..8da11c603447d 100644 --- a/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.cpp +++ b/unified-runtime/source/adapters/level_zero/common/helpers/shared_helpers.cpp @@ -235,6 +235,10 @@ template <> ze_structure_type_t getZeStructureType() { return ZE_STRUCTURE_TYPE_PITCHED_IMAGE_EXP_DESC; } +template <> +ze_structure_type_t getZeStructureType() { + return ZE_STRUCTURE_TYPE_CUSTOM_PITCH_EXP_DESC; +} template <> ze_structure_type_t getZeStructureType() { return ZE_STRUCTURE_TYPE_MODULE_DESC; } diff --git a/unified-runtime/source/adapters/level_zero/common/image_common.cpp b/unified-runtime/source/adapters/level_zero/common/image_common.cpp index d153fcc5a43f1..52ea9b8270c49 100644 --- a/unified-runtime/source/adapters/level_zero/common/image_common.cpp +++ b/unified-runtime/source/adapters/level_zero/common/image_common.cpp @@ -342,6 +342,28 @@ ur_result_t bindlessImagesCreateImpl(ur_context_handle_t hContext, } else { BindlessDesc.pNext = &PitchedDesc; } + + // Only chain a ze_custom_pitch_exp_desc_t when SYCL has signalled a + // user-declared pitch via ur_exp_image_user_pitch_desc_t on the pNext + // chain. A non-zero pImageDesc->rowPitch alone is ambiguous: it may be + // a driver-computed pitch (e.g. from urUSMPitchedAllocExp) bookkept back + // through the descriptor, in which case letting L0 compute the pitch + // itself is what we want. + ZeStruct CustomPitchDesc; + for (const auto *Base = + static_cast(pImageDesc->pNext); + Base != nullptr; + Base = static_cast(Base->pNext)) { + if (Base->stype == UR_STRUCTURE_TYPE_EXP_IMAGE_USER_PITCH_DESC) { + const auto *UserPitch = + reinterpret_cast(Base); + CustomPitchDesc.rowPitch = UserPitch->rowPitch; + CustomPitchDesc.slicePitch = UserPitch->slicePitch; + PitchedDesc.pNext = &CustomPitchDesc; + break; + } + } + try { ZE2UR_CALL_THROWS(zeImageCreate, (zeCtx, hDevice->ZeDevice, &ZeImageDesc, ZeImage.ptr())); diff --git a/unified-runtime/source/common/stype_map_helpers.def b/unified-runtime/source/common/stype_map_helpers.def index 35742d7471eaa..f402733ae835c 100644 --- a/unified-runtime/source/common/stype_map_helpers.def +++ b/unified-runtime/source/common/stype_map_helpers.def @@ -170,6 +170,9 @@ template <> struct stype_map : stype_map_impl {}; template <> +struct stype_map + : stype_map_impl {}; +template <> struct stype_map : stype_map_impl {}; template <> diff --git a/unified-runtime/source/loader/loader.def.in b/unified-runtime/source/loader/loader.def.in index 57c8fb82b22ed..0afc088823d3d 100644 --- a/unified-runtime/source/loader/loader.def.in +++ b/unified-runtime/source/loader/loader.def.in @@ -385,6 +385,7 @@ EXPORTS urPrintExpImageCopyInputTypes urPrintExpImageCopyRegion urPrintExpImageMemType + urPrintExpImageUserPitchDesc urPrintExpKernelArgMemObjTuple urPrintExpKernelArgProperties urPrintExpKernelArgType diff --git a/unified-runtime/source/loader/loader.map.in b/unified-runtime/source/loader/loader.map.in index 5296424a405fe..58a90ce0b4605 100644 --- a/unified-runtime/source/loader/loader.map.in +++ b/unified-runtime/source/loader/loader.map.in @@ -385,6 +385,7 @@ urPrintExpImageCopyInputTypes; urPrintExpImageCopyRegion; urPrintExpImageMemType; + urPrintExpImageUserPitchDesc; urPrintExpKernelArgMemObjTuple; urPrintExpKernelArgProperties; urPrintExpKernelArgType; diff --git a/unified-runtime/source/loader/ur_print.cpp b/unified-runtime/source/loader/ur_print.cpp index 8004688ec7f57..705a8e366e2d0 100644 --- a/unified-runtime/source/loader/ur_print.cpp +++ b/unified-runtime/source/loader/ur_print.cpp @@ -1120,6 +1120,15 @@ urPrintExpImageCopyRegion(const struct ur_exp_image_copy_region_t params, return str_copy(&ss, buffer, buff_size, out_size); } +ur_result_t +urPrintExpImageUserPitchDesc(const struct ur_exp_image_user_pitch_desc_t params, + char *buffer, const size_t buff_size, + size_t *out_size) { + std::stringstream ss; + ss << params; + return str_copy(&ss, buffer, buff_size, out_size); +} + ur_result_t urPrintExpProgramFlags(enum ur_exp_program_flag_t value, char *buffer, const size_t buff_size, size_t *out_size) {