Skip to content

[SYCL][Bindless] Device 'image_mem_handle' to 'image_mem_handle' Sub-Region Copy #15579

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

Merged
merged 6 commits into from
Oct 15, 2024
Merged

[SYCL][Bindless] Device 'image_mem_handle' to 'image_mem_handle' Sub-Region Copy #15579

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
6 commits
Select commit Hold shift + click to select a range
46666c1
[SYCL][Bindless] Device 'image_mem_handle' to 'image_mem_handle' Sub-…
DBDuncan Aug 12, 2024
7700111
Rename 'SrcImageDesc' and 'DestImageDesc' to 'SrcImgDesc' and 'DescIm…
DBDuncan Oct 7, 2024
9befa55
Merge remote-tracking branch 'origin/sycl' into duncan/sub-copy
DBDuncan Oct 7, 2024
0ec04d4
Update windows symbols test
DBDuncan Oct 8, 2024
ab38b3d
Fix order of new functions
DBDuncan Oct 8, 2024
de540bd
Add missing full stops
DBDuncan Oct 11, 2024
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
40 changes: 40 additions & 0 deletions sycl/doc/extensions/experimental/sycl_ext_oneapi_bindless_images.asciidoc
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -813,6 +813,16 @@ public:
const ext::oneapi::experimental::image_mem_handle Src,
ext::oneapi::experimental::image_mem_handle Dest,
const ext::oneapi::experimental::image_descriptor &ImageDesc);

// Device to device copy with offsets and extent
void ext_oneapi_copy(
Copy link
Contributor

@cperkinsintel cperkinsintel Oct 3, 2024

Choose a reason for hiding this comment

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

are we sure about this name? For something so generic like ext_oneapi_copy I would expect something that handles any type of memory (not just image data) and maybe across all address spaces (host, device). Much like queue.copy() is now.

But, the comments on this routine suggest its capability is much more focused. If so, I think it's name should express that. maybe something like ext_oneapi_d2d_img_copy or ext_d2d_img_copy ?

Copy link
Contributor Author

@DBDuncan DBDuncan Oct 4, 2024

Choose a reason for hiding this comment

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

We already have a lot of overloads that perform copies. The reason we chose this originally is we just intended to extend the existing copy functionality to support new bindless images. Also, our overloads are not just images to images but also memory to memory copies.

Copy link
Contributor Author

@DBDuncan DBDuncan Oct 4, 2024 •
edited
Loading

Choose a reason for hiding this comment

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

There are other new overloads planned to be added. Such as host to host and device usm to device image_mem_handle.

Copy link
Contributor Author

@DBDuncan DBDuncan Oct 8, 2024

Choose a reason for hiding this comment

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

I have been thinking that perhaps the amount of ext_oneapi_copy overloads we have is getting a bit unwieldy. Especially considering a number additional ones need to be added to support more device to device copy variants.

I don't think it would be a good idea to start making changes involving other copy functions here. But something to be looked into later.

Copy link
Contributor

@cperkinsintel cperkinsintel Oct 14, 2024

Choose a reason for hiding this comment

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

sorry for not responding to your comments earlier. If there are other overloads that handle other data types, it wouldn't hurt to mention that in the comment. Even just "this overload of ext_copy is for bindless images".

Copy link
Contributor Author

@DBDuncan DBDuncan Oct 15, 2024 •
edited
Loading

Choose a reason for hiding this comment

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

No problem, thanks for the approval! On your suggestion, it would likely be best to apply a comment like you suggest to all the overloads at the same time which is a bit out of scope it feels for this PR. As there are 50~ of them (Have to triple them in queue.hpp...). queue.hpp and handler.hpp need a bit of a refactor which is very tentatively planned early next year. Double checking the order of the functions and their corresponding comments. Comments stating they are overloads for bindless images can be added then.

const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent)
};

class queue {
Expand Down Expand Up @@ -954,6 +964,34 @@ public:
ext::oneapi::experimental::image_mem_handle Dest,
const ext::oneapi::experimental::image_descriptor &ImageDesc,
const std::vector<event> &DepEvents);

// Device to device copy with offsets and extent
void ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent)
void ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent
event DepEvent)
void ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent
const std::vector<event> &DepEvents)
};
}
```
Expand Down Expand Up @@ -2903,4 +2941,6 @@ These features still need to be handled:
|6.1|2024-09-09| - Update for image-array sub-region copy support.
|6.2|2024-09-26| - Added addressing mode `ext_oneapi_clamp_to_border` value,
equivalent to `clamp`, to match with external APIs.
|6.3|2024-10-02| - Add support for `image_mem_handle` to `image_mem_handle`
sub-region copies.
|======================
52 changes: 52 additions & 0 deletions sycl/include/sycl/ext/oneapi/bindless_images.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1609,6 +1609,58 @@ inline event queue::ext_oneapi_copy(
CodeLoc);
}

inline event queue::ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent, const detail::code_location &CodeLoc) {
detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
return submit(
[&](handler &CGH) {
CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
DestImgDesc, CopyExtent);
},
CodeLoc);
}

inline event queue::ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent, event DepEvent,
const detail::code_location &CodeLoc) {
detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
return submit(
[&](handler &CGH) {
CGH.depends_on(DepEvent);
CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
DestImgDesc, CopyExtent);
},
CodeLoc);
}

inline event queue::ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
const detail::code_location &CodeLoc) {
detail::tls_code_loc_t TlsCodeLocCapture(CodeLoc);
return submit(
[&](handler &CGH) {
CGH.depends_on(DepEvents);
CGH.ext_oneapi_copy(Src, SrcOffset, SrcImgDesc, Dest, DestOffset,
DestImgDesc, CopyExtent);
},
CodeLoc);
}

inline event queue::ext_oneapi_copy(
const void *Src, sycl::range<3> SrcOffset, void *Dest,
sycl::range<3> DestOffset,
Expand Down
28 changes: 27 additions & 1 deletion sycl/include/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -3143,7 +3143,7 @@ class __SYCL_EXPORT handler {
/// incomplete.
///
/// \param Src is an opaque image memory handle to the source memory.
/// \param SrcOffset is an offset from the origin of source measured in pixels
/// \param SrcOffset is an offset from the source origin measured in pixels
/// (pixel size determined by \p SrcImgDesc )
/// \param SrcImgDesc is the source image descriptor
/// \param Dest is a USM pointer to the destination memory.
Expand Down Expand Up @@ -3189,6 +3189,32 @@ class __SYCL_EXPORT handler {
ext::oneapi::experimental::image_mem_handle Dest,
const ext::oneapi::experimental::image_descriptor &ImageDesc);

/// Copies data from device to device memory, where \p Src and \p Dest
/// are opaque image memory handles. Allows for a sub-region copy, where
/// \p SrcOffset, \p DestOffset and \p CopyExtent are used to determine the
/// sub-region. Pixel size is determined by \p SrcImgDesc
/// An exception is thrown if either \p Src or \p Dest is incomplete.
///
/// \param Src is an opaque image memory handle to the source memory.
/// \param SrcOffset is an offset from the source origin measured in pixels
/// (pixel size determined by \p SrcImgDesc )
/// \param SrcImgDesc is the source image descriptor
/// \param Dest is an opaque image memory handle to the destination memory.
/// \param DestOffset is an offset from the destination origin measured in
/// pixels (pixel size determined by \p DestImgDesc )
/// \param DestImgDesc is the destination image descriptor
/// \param CopyExtent is the width, height, and depth of the region to copy
/// measured in pixels (pixel size determined by
/// \p SrcImgDesc )
void ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent);

/// Copies data from one memory region to another, where \p Src and \p Dest
/// are USM pointers. Allows for a sub-region copy, where \p SrcOffset ,
/// \p DestOffset , and \p Extent are used to determine the sub-region.
Expand Down
87 changes: 87 additions & 0 deletions sycl/include/sycl/queue.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1757,6 +1757,93 @@ class __SYCL_EXPORT queue : public detail::OwnerLessBase<queue> {
const ext::oneapi::experimental::image_descriptor &ImageDesc,
const detail::code_location &CodeLoc = detail::code_location::current());

/// Copies data from device to device memory, where \p Src and \p Dest
/// are opaque image memory handles. Allows for a sub-region copy, where
/// \p SrcOffset, \p DestOffset and \p CopyExtent are used to determine the
/// sub-region. Pixel size is determined by \p SrcImgDesc
/// An exception is thrown if either \p Src or \p Dest is incomplete.
///
/// \param Src is an opaque image memory handle to the source memory.
/// \param SrcOffset is an offset from the origin of source measured in pixels
/// (pixel size determined by \p SrcImgDesc )
/// \param SrcImgDesc is the source image descriptor
/// \param Dest is an opaque image memory handle to the destination memory.
/// \param DestOffset is an offset from the origin of destination measured in
/// pixels (pixel size determined by \p DestImgDesc )
/// \param DestImgDesc is the destination image descriptor
/// \param CopyExtent is the width, height, and depth of the region to copy
/// measured in pixels (pixel size determined by
/// \p SrcImgDesc )
/// \return an event representing the copy operation.
event ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent,
const detail::code_location &CodeLoc = detail::code_location::current());

/// Copies data from device to device memory, where \p Src and \p Dest
/// are opaque image memory handles. Allows for a sub-region copy, where
/// \p SrcOffset, \p DestOffset and \p CopyExtent are used to determine the
/// sub-region. Pixel size is determined by \p SrcImgDesc
/// An exception is thrown if either \p Src or \p Dest is incomplete.
///
/// \param Src is an opaque image memory handle to the source memory.
/// \param SrcOffset is an offset from the origin of source measured in pixels
/// (pixel size determined by \p SrcImgDesc )
/// \param SrcImgDesc is the source image descriptor
/// \param Dest is an opaque image memory handle to the destination memory.
/// \param DestOffset is an offset from the origin of destination measured in
/// pixels (pixel size determined by \p DestImgDesc )
/// \param DestImgDesc is the destination image descriptor
/// \param CopyExtent is the width, height, and depth of the region to copy
/// measured in pixels (pixel size determined by
/// \p SrcImgDesc )
/// \param DepEvent is an event that specifies the kernel dependencies.
/// \return an event representing the copy operation.
event ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent, event DepEvent,
const detail::code_location &CodeLoc = detail::code_location::current());

/// Copies data from device to device memory, where \p Src and \p Dest
/// are opaque image memory handles. Allows for a sub-region copy, where
/// \p SrcOffset, \p DestOffset and \p CopyExtent are used to determine the
/// sub-region. Pixel size is determined by \p SrcImgDesc
/// An exception is thrown if either \p Src or \p Dest is incomplete.
///
/// \param Src is an opaque image memory handle to the source memory.
/// \param SrcOffset is an offset from the origin of source measured in pixels
/// (pixel size determined by \p SrcImgDesc )
/// \param srcImgDesc is the source image descriptor
/// \param Dest is an opaque image memory handle to the destination memory.
/// \param DestOffset is an offset from the origin of destination measured in
/// pixels (pixel size determined by \p DestImgDesc )
/// \param DestImgDesc is the destination image descriptor
/// \param CopyExtent is the width, height, and depth of the region to copy
/// measured in pixels (pixel size determined by
/// \p SrcImgDesc )
/// \param DepEvents is a vector of events that specifies the kernel
/// dependencies.
/// \return an event representing the copy operation.
event ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest,
sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent, const std::vector<event> &DepEvents,
const detail::code_location &CodeLoc = detail::code_location::current());

/// Copies data from one memory region to another, where \p Src and \p Dest
/// are USM pointers. Allows for a sub-region copy, where \p SrcOffset ,
/// \p DestOffset , and \p Extent are used to determine the sub-region.
Expand Down
99 changes: 99 additions & 0 deletions sycl/source/handler.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1248,6 +1248,105 @@ void handler::ext_oneapi_copy(
setType(detail::CGType::CopyImage);
}

void handler::ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
const ext::oneapi::experimental::image_descriptor &SrcImgDesc,
ext::oneapi::experimental::image_mem_handle Dest, sycl::range<3> DestOffset,
const ext::oneapi::experimental::image_descriptor &DestImgDesc,
sycl::range<3> CopyExtent) {
throwIfGraphAssociated<
ext::oneapi::experimental::detail::UnsupportedGraphFeatures::
sycl_ext_oneapi_bindless_images>();
SrcImgDesc.verify();
DestImgDesc.verify();

auto isOutOfRange = [](const sycl::range<3> &range,
const sycl::range<3> &offset,
const sycl::range<3> &copyExtent) {
sycl::range<3> result = (range > 0UL && ((offset + copyExtent) > range));

return (static_cast<bool>(result[0]) || static_cast<bool>(result[1]) ||
static_cast<bool>(result[2]));
};

sycl::range<3> SrcImageSize = {SrcImgDesc.width, SrcImgDesc.height,
SrcImgDesc.depth};
sycl::range<3> DestImageSize = {DestImgDesc.width, DestImgDesc.height,
DestImgDesc.depth};

if (isOutOfRange(SrcImageSize, SrcOffset, CopyExtent) ||
isOutOfRange(DestImageSize, DestOffset, CopyExtent)) {
throw sycl::exception(
make_error_code(errc::invalid),
"Image copy attempted to access out of bounds memory!");
}

MSrcPtr = reinterpret_cast<void *>(Src.raw_handle);
MDstPtr = reinterpret_cast<void *>(Dest.raw_handle);

ur_image_desc_t UrSrcDesc = {};
UrSrcDesc.width = SrcImgDesc.width;
UrSrcDesc.height = SrcImgDesc.height;
UrSrcDesc.depth = SrcImgDesc.depth;
UrSrcDesc.arraySize = SrcImgDesc.array_size;

ur_image_desc_t UrDestDesc = {};
UrDestDesc.width = DestImgDesc.width;
UrDestDesc.height = DestImgDesc.height;
UrDestDesc.depth = DestImgDesc.depth;
UrDestDesc.arraySize = DestImgDesc.array_size;

auto fill_image_type =
[](const ext::oneapi::experimental::image_descriptor &Desc,
ur_image_desc_t &UrDesc) {
if (Desc.array_size > 1) {
// Image Array.
UrDesc.type = Desc.height > 0 ? UR_MEM_TYPE_IMAGE2D_ARRAY
: UR_MEM_TYPE_IMAGE1D_ARRAY;

// Cubemap.
UrDesc.type =
Desc.type == sycl::ext::oneapi::experimental::image_type::cubemap
? UR_MEM_TYPE_IMAGE_CUBEMAP_EXP
: UrDesc.type;
} else {
UrDesc.type = Desc.depth > 0
? UR_MEM_TYPE_IMAGE3D
: (Desc.height > 0 ? UR_MEM_TYPE_IMAGE2D
: UR_MEM_TYPE_IMAGE1D);
}
};

fill_image_type(SrcImgDesc, UrSrcDesc);
fill_image_type(DestImgDesc, UrDestDesc);

auto fill_format = [](const ext::oneapi::experimental::image_descriptor &Desc,
ur_image_format_t &UrFormat) {
UrFormat.channelType =
sycl::_V1::detail::convertChannelType(Desc.channel_type);
UrFormat.channelOrder = sycl::detail::convertChannelOrder(
sycl::_V1::ext::oneapi::experimental::detail::
get_image_default_channel_order(Desc.num_channels));
};

ur_image_format_t UrSrcFormat;
ur_image_format_t UrDestFormat;

fill_format(SrcImgDesc, UrSrcFormat);
fill_format(DestImgDesc, UrDestFormat);

impl->MSrcOffset = {SrcOffset[0], SrcOffset[1], SrcOffset[2]};
impl->MDestOffset = {DestOffset[0], DestOffset[1], DestOffset[2]};
impl->MCopyExtent = {CopyExtent[0], CopyExtent[1], CopyExtent[2]};
impl->MSrcImageDesc = UrSrcDesc;
impl->MDstImageDesc = UrDestDesc;
impl->MSrcImageFormat = UrSrcFormat;
impl->MDstImageFormat = UrDestFormat;
impl->MImageCopyFlags = UR_EXP_IMAGE_COPY_FLAG_DEVICE_TO_DEVICE;
setType(detail::CGType::CopyImage);
}

void handler::ext_oneapi_copy(
const ext::oneapi::experimental::image_mem_handle Src,
sycl::range<3> SrcOffset,
Expand Down
Loading
Loading