[SYCL][UR][Bindless][Doc] Fix copy docs and implementation. (#19093)

This patch fixes the implementation of bindless image copies.
Previously, source and destination pitch values were not being set
correctly.

This patch also updates the wording around the requirements for
`ext_oneapi_copy`.

A missing requirement was added to the specification. Namely that the
`CopyExtent` parameter in the `ext_oneapi_copy` functions that take it,
must not have `0` values in any of the three dimensions, they must be
greater than or equal to `1`.

The requirements for `ext_oneapi_copy` have also been re-written to
prescribe what the functions expect, instead of providing a list of
cases in which the function may fail. This should hopefully make it
clearer and more prescriptive, rather than saying the copy function may
fail if some condition is not met, we now say that the functions require
that certain conditions be met.

The coverage for sub-region copy testing has also been extended to
prevent future regressions.

Author: przemektmalon

sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc

@@ -1335,51 +1335,63 @@ For the forms that take a USM pointer, the image memory must also have been
 allocated within the same context and device of the `queue`. The USM memory 
 must be accessible on the queue's device.
 
-The `ext_oneapi_copy` function variants that don't take offsets and extents may 
-fail in the following scenarios:
+The `ext_oneapi_copy` function variants that do not take offsets and extents
+must ensure that the following conditions are met:
 
-1. The `Src` and `Dest` memory was not allocated on the same device and 
-context of the queue.
+1. The `Src` and `Dest` memory was allocated on the same device and context.
 
-2. The `Src` and `Dest` memory regions, where `Src` or `Dest` can be either 
-on the host or device, do not have the same memory capacity, where the capacity 
-is calculate from the `width`, `height`, `depth`, `channel_order`, and 
+2. The `Src` and `Dest` memory regions, where `Src` or `Dest` can be either
+on the host or device, have the same memory capacity, where the capacity
+is calculated from the `width`, `height`, `depth`, `channel_order`, and
 `channel_type` members of the `image_descriptor` parameter.
 
-The `ext_oneapi_copy` function variants that do take offsets and extents may 
-fail in the following scenarios:
+The `ext_oneapi_copy` function variants that do take offsets and extents must
+ensure that the following conditions are met. If a condition names a specific
+parameter, it is only applicable to the function variants that take that
+parameter.
 
-1. The `Src` and `Dest` memory was not allocated on the same device and 
-context of the queue.
+1. The `Src` and `Dest` memory was allocated on the same device and context.
 
-2. The image descriptor passed does not match the image descriptor used to 
-allocate the image on the device.
+2. The image descriptors passed match the image descriptors used to allocate
+the image's memory on the device.
 
-3. the `CopyExtent` describes a memory region larger than that which was 
-allocated on either the host or the device.
+3. The `CopyExtent` describes a memory region that is not larger than that which
+was allocated on either the host or the device.
 
-4. The `HostExtent` describes a memory region larger than that which was 
-allocated on the host.
+4. The `HostExtent` describes a memory region that is not larger than that which
+was allocated on the host.
 
-5. The `SrcExtent` describes a memory region larger than that which was 
-allocated, where `Src` can be either the host or device.
+5. The `SrcExtent` describes a memory region that is not larger than that which
+was allocated, where `Src` can be either on the host or on the device.
 
-6. The `DestExtent` describes a memory region larger than that which was 
-allocated, where `Dest` can be either the host or device.
+6. The `DestExtent` describes a memory region that is not larger than that which
+was allocated, where `Dest` can be either on the host or on the device.
 
-7. If `SrcOffset + CopyExtent` moves the memory sub-region outside the bounds 
-of the memory described by `Src`, irrespective of whether `Src` is on the host 
-or the device.
+7. The `DeviceRowPitch` adheres to the alignment requirements outlined in the
+"Pitch alignment restrictions and queries" section.
 
-8. If `DestOffset + CopyExtent` moves the memory sub-region outside the bounds 
-of the memory described by `Dest`, irrespective of whether `Dest` is on the 
-host or the device.
+8. The `DeviceRowPitch` is greater than or equal to the width of the image on
+the device.
 
-9. The `DeviceRowPitch` does not adhere to the alignment requirements 
-outlined in section "Pitch alignment restrictions and queries"
+9. For the relevant dimensions, `SrcOffset + CopyExtent` does not move the
+memory sub-region outside the bounds of the memory described by `Src`, 
+irrespective of whether `Src` is on the host or the device. The relevant
+dimensions are `x` for 1D images; `x` and `y` for 2D images; and `x`, `y`, and
+`z` for 3D images. `x`, `y`, and `z` correspond to indices `0`, `1`, and `2` of
+the `SrcOffset` and `CopyExtent` parameters, respectively.
+
+10. For the relevant dimensions, `DestOffset + CopyExtent` does not move the
+memory sub-region outside the bounds of the memory described by `Dest`,
+irrespective of whether `Dest` is on the host or the device. The relevant
+dimensions are `x` for 1D images, `x` and `y` for 2D images, and `x`, `y`, and
+`z` for 3D images. `x`, `y`, and `z` correspond to indices `0`, `1`, and `2` of
+the `SrcOffset` and `CopyExtent` parameters, respectively.
+
+11. The `CopyExtent`'s' `x`, `y`, and `z` dimensions must not be `0`. They must
+be greater than or equal to `1`. Even if the image is 1D or 2D, the remaining
+non-relevant dimension's values must be set to `1` in the `CopyExtent`
+parameter.
 
-10. The value of `DeviceRowPitch` is smaller than the width of the image on 
-the device.
 
 If copying of an image fails, `ext_oneapi_copy` will throw a `sycl::exception` 
 with error code `sycl::errc::invalid`, and relay an error message back to the 

sycl/source/handler.cpp

@@ -98,6 +98,33 @@ void *getValueFromDynamicParameter(
 
 // Bindless image helpers
 
+constexpr size_t get_channel_size(
+    const sycl::ext::oneapi::experimental::image_descriptor &Desc) {
+  switch (Desc.channel_type) {
+  case sycl::image_channel_type::fp16:
+    return sizeof(sycl::half);
+  case sycl::image_channel_type::fp32:
+    return sizeof(float);
+  case sycl::image_channel_type::snorm_int8:
+  case sycl::image_channel_type::unorm_int8:
+  case sycl::image_channel_type::signed_int8:
+  case sycl::image_channel_type::unsigned_int8:
+    return sizeof(uint8_t);
+  case sycl::image_channel_type::snorm_int16:
+  case sycl::image_channel_type::unorm_int16:
+  case sycl::image_channel_type::signed_int16:
+  case sycl::image_channel_type::unsigned_int16:
+    return sizeof(uint16_t);
+  case sycl::image_channel_type::signed_int32:
+  case sycl::image_channel_type::unsigned_int32:
+    return sizeof(uint32_t);
+  default:
+    throw sycl::exception(make_error_code(errc::invalid),
+                          "Unsupported channel type");
+    return 0;
+  }
+}
+
 // Fill image type and return depth or array_size
 static unsigned int
 fill_image_type(const ext::oneapi::experimental::image_descriptor &Desc,
@@ -257,16 +284,8 @@ fill_copy_args(detail::handler_impl *impl,
     impl->MDstImageDesc.depth = DestExtent[2];
   }
 
-  if (impl->MImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) {
-    impl->MSrcImageDesc.rowPitch = 0;
-    impl->MDstImageDesc.rowPitch = DestPitch;
-  } else if (impl->MImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
-    impl->MSrcImageDesc.rowPitch = SrcPitch;
-    impl->MDstImageDesc.rowPitch = 0;
-  } else {
-    impl->MSrcImageDesc.rowPitch = SrcPitch;
-    impl->MDstImageDesc.rowPitch = DestPitch;
-  }
+  impl->MSrcImageDesc.rowPitch = SrcPitch;
+  impl->MDstImageDesc.rowPitch = DestPitch;
 }
 
 static void
@@ -279,9 +298,11 @@ fill_copy_args(detail::handler_impl *impl,
                sycl::range<3> DestExtent = {0, 0, 0},
                sycl::range<3> CopyExtent = {0, 0, 0}) {
 
-  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, 0 /*SrcPitch*/,
-                 0 /*DestPitch*/, SrcOffset, SrcExtent, DestOffset, DestExtent,
-                 CopyExtent);
+  size_t SrcPitch = SrcExtent[0] * Desc.num_channels * get_channel_size(Desc);
+  size_t DestPitch = DestExtent[0] * Desc.num_channels * get_channel_size(Desc);
+
+  fill_copy_args(impl, Desc, Desc, ImageCopyFlags, SrcPitch, DestPitch,
+                 SrcOffset, SrcExtent, DestOffset, DestExtent, CopyExtent);
 }
 
 static void
@@ -309,8 +330,13 @@ fill_copy_args(detail::handler_impl *impl,
                sycl::range<3> DestExtent = {0, 0, 0},
                sycl::range<3> CopyExtent = {0, 0, 0}) {
 
-  fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags, 0 /*SrcPitch*/,
-                 0 /*DestPitch*/, SrcOffset, SrcExtent, DestOffset, DestExtent,
+  size_t SrcPitch =
+      SrcExtent[0] * SrcImgDesc.num_channels * get_channel_size(SrcImgDesc);
+  size_t DestPitch =
+      DestExtent[0] * DestImgDesc.num_channels * get_channel_size(DestImgDesc);
+
+  fill_copy_args(impl, SrcImgDesc, DestImgDesc, ImageCopyFlags, SrcPitch,
+                 DestPitch, SrcOffset, SrcExtent, DestOffset, DestExtent,
                  CopyExtent);
 }
 
@@ -1618,10 +1644,17 @@ void handler::ext_oneapi_copy(
       get_pointer_type(Dest,
                        createSyclObjFromImpl<context>(impl->get_context())));
 
-  if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE ||
-      ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
-    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, DeviceRowPitch,
+  // Calculate host pitch, where host memory is always assumed to be tightly
+  // packed.
+  size_t HostRowPitch =
+      Desc.width * Desc.num_channels * detail::get_channel_size(Desc);
+
+  if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) {
+    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, HostRowPitch,
                            DeviceRowPitch);
+  } else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
+    detail::fill_copy_args(get_impl(), Desc, ImageCopyFlags, DeviceRowPitch,
+                           HostRowPitch);
   } else {
     throw sycl::exception(make_error_code(errc::invalid),
                           "Copy Error: This copy function only performs host "
@@ -1650,14 +1683,19 @@ void handler::ext_oneapi_copy(
       get_pointer_type(Dest,
                        createSyclObjFromImpl<context>(impl->get_context())));
 
+  // Calculate host pitch, where host memory is always assumed to be tightly
+  // packed.
+  size_t HostRowPitch = HostExtent[0] * DeviceImgDesc.num_channels *
+                        detail::get_channel_size(DeviceImgDesc);
+
   // Fill the host extent based on the type of copy.
   if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_HOST_TO_DEVICE) {
     detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags,
-                           DeviceRowPitch, DeviceRowPitch, SrcOffset,
-                           HostExtent, DestOffset, {0, 0, 0}, CopyExtent);
+                           HostRowPitch, DeviceRowPitch, SrcOffset, HostExtent,
+                           DestOffset, {0, 0, 0}, CopyExtent);
   } else if (ImageCopyFlags == UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_HOST) {
     detail::fill_copy_args(get_impl(), DeviceImgDesc, ImageCopyFlags,
-                           DeviceRowPitch, DeviceRowPitch, SrcOffset, {0, 0, 0},
+                           DeviceRowPitch, HostRowPitch, SrcOffset, {0, 0, 0},
                            DestOffset, HostExtent, CopyExtent);
   } else {
     throw sycl::exception(make_error_code(errc::invalid),