Bindless image desc match ur

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 

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) {

sycl/source/detail/bindless_images.cpp

@@ -21,16 +21,22 @@ 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;
   urDesc.height = desc.height;
   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;

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) &&
@@ -386,22 +385,6 @@ bool image_impl::checkImageDesc(const ur_image_desc_t &Desc,
                     "For a 1D and 2D image array, the array_size must be a "
                     "Value >= 1 and <= info::device::image_max_array_size.");
 
-  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;
 }
 

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 = {};

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,

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) {

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(