[Bindless Images][Bug] Fixed bindless_images bugs with copy_extent #2689
Conversation
Mark OpenCL CPU known failing device more generally
TejaX-Alaghari
force-pushed
the
bug
branch
2 times, most recently
from
00033c1 to
8b196e1
Compare
There was a problem hiding this comment.
I'm not sure if one of the changes here is correct, as most of the functions are not documented properly as to the copy direction they intend to perform.
There are some bugs in the Bindless Images copy code. See intel/llvm#19093 for the fix to the way we pass pitch values to the SYCL backends. The above linked PR also includes changes to the Bindless Images specification, which was underspecified, in that the CopyExtent parameters require that each dimension of the CopyExtent is greater than or equal to 1.
| const auto copy_extend = sycl::range<3>(w / ele_size, h, 1); | ||
| return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset, | ||
| dest_extend, copy_extend); | ||
| desc_src, p, copy_extend); | ||
| } |
There was a problem hiding this comment.
I'm not sure if this is correct, as there are no comments that describe what this function does (i.e. what the copy direction is).
If this is a (device) image_mem_handle src to (host) void *dest copy, then this should be:
return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset,
dest_extend, copy_extend);But, if this is a (device) image_mem_handle src to (device) void *dest copy, then it should be as it is in this PR (which changes the direction of this copy function).
| @@ -949,9 +949,9 @@ dpct_memcpy(sycl::ext::oneapi::experimental::image_mem_handle src, | |||
| sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0); | |||
| const auto dest_offset = sycl::range<3>(0, 0, 0); | |||
| const auto dest_extend = sycl::range<3>(p / ele_size, 0, 0); | |||
| const auto copy_extend = sycl::range<3>(w / ele_size, h, 0); | |||
| const auto copy_extend = sycl::range<3>(w / ele_size, h, 1); | |||
| return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset, | |||
There was a problem hiding this comment.
If the API that coping to the device is different from copying to the host, we should judge the pointer and then choose the right way. Or add a parameter to show the difference.
This PR aims to update the bindless_image memory copies that are changed in the compiler on Jan31