Skip to content

Commit 9b5b55f

Browse files
Fixed bindless_images bugs with copy_extent
1 parent c725bae commit 9b5b55f

File tree

1 file changed

+34
-18
lines changed

1 file changed

+34
-18
lines changed

clang/runtime/dpct-rt/include/dpct/bindless_images.hpp

Lines changed: 34 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -949,9 +949,9 @@ dpct_memcpy(sycl::ext::oneapi::experimental::image_mem_handle src,
949949
sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
950950
const auto dest_offset = sycl::range<3>(0, 0, 0);
951951
const auto dest_extend = sycl::range<3>(p / ele_size, 0, 0);
952-
const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
952+
const auto copy_extend = sycl::range<3>(w / ele_size, h, 1);
953953
return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset,
954-
dest_extend, copy_extend);
954+
desc_src, p, copy_extend);
955955
}
956956

957957
static inline std::vector<sycl::event> dpct_memcpy_to_host(
@@ -969,7 +969,7 @@ static inline std::vector<sycl::event> dpct_memcpy_to_host(
969969
const auto dest_offset = sycl::range<3>(offset_dest / ele_size, 0, 0);
970970
const auto dest_extend = sycl::range<3>(0, 0, 0);
971971
const auto copy_extend =
972-
sycl::range<3>((w - w_offset_src) / ele_size, 1, 0);
972+
sycl::range<3>((w - w_offset_src) / ele_size, 1, 1);
973973
event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
974974
dest_host_ptr, dest_offset,
975975
dest_extend, copy_extend));
@@ -982,7 +982,7 @@ static inline std::vector<sycl::event> dpct_memcpy_to_host(
982982
sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
983983
const auto dest_offset = sycl::range<3>(offset_dest / ele_size, 0, 0);
984984
const auto dest_extend = sycl::range<3>(0, 0, 0);
985-
const auto copy_extend = sycl::range<3>((s - offset_dest) / ele_size, 1, 0);
985+
const auto copy_extend = sycl::range<3>((s - offset_dest) / ele_size, 1, 1);
986986
event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
987987
dest_host_ptr, dest_offset,
988988
dest_extend, copy_extend));
@@ -1021,7 +1021,7 @@ dpct_memcpy(const void *src,
10211021
const auto src_extend = sycl::range<3>(p / ele_size, 0, 0);
10221022
const auto dest_offset =
10231023
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
1024-
const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
1024+
const auto copy_extend = sycl::range<3>(w / ele_size, h, 1);
10251025
// TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
10261026
return q.ext_oneapi_copy(const_cast<void *>(src), src_offset, src_extend,
10271027
dest, dest_offset, desc_dest, copy_extend);
@@ -1042,7 +1042,7 @@ static inline std::vector<sycl::event> dpct_memcpy_from_host(
10421042
const auto dest_offset =
10431043
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
10441044
const auto copy_extend =
1045-
sycl::range<3>((w - w_offset_dest) / ele_size, 1, 0);
1045+
sycl::range<3>((w - w_offset_dest) / ele_size, 1, 1);
10461046
// TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
10471047
event_list.push_back(q.ext_oneapi_copy(
10481048
const_cast<void *>(src_host_ptr), src_offset, src_extend, dest,
@@ -1056,7 +1056,7 @@ static inline std::vector<sycl::event> dpct_memcpy_from_host(
10561056
const auto src_extend = sycl::range<3>(0, 0, 0);
10571057
const auto dest_offset =
10581058
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
1059-
const auto copy_extend = sycl::range<3>((s - offset_src) / ele_size, 1, 0);
1059+
const auto copy_extend = sycl::range<3>((s - offset_src) / ele_size, 1, 1);
10601060
// TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
10611061
event_list.push_back(q.ext_oneapi_copy(
10621062
const_cast<void *>(src_host_ptr), src_offset, src_extend, dest,
@@ -1101,7 +1101,7 @@ dpct_memcpy(const image_mem_wrapper *src, const sycl::id<3> &src_id,
11011101
sycl::range<3>(dest.get_pitch() / ele_size, dest.get_y(), 1);
11021102
const auto copy_extend = sycl::range<3>(
11031103
copy_x_size_byte != 0 ? copy_x_size_byte / ele_size : size[0], size[1],
1104-
size[2]);
1104+
size[2] != 0 ? size[2] : 1);
11051105
return q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
11061106
dest.get_data_ptr(), dest_offset, dest_extend,
11071107
copy_extend);
@@ -2001,11 +2001,19 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
20012001
size_t w_offset_src, size_t h_offset_src,
20022002
size_t w, size_t h,
20032003
sycl::queue q = get_default_queue()) {
2004-
auto temp = (void *)sycl::malloc_device(w * h, q);
2005-
// TODO: Need change logic when sycl support image_mem to image_mem copy.
2006-
dpct_memcpy(temp, w, src, w_offset_src, h_offset_src, w, h, q);
2007-
dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, w, w, h, q);
2008-
sycl::free(temp, q);
2004+
const auto from_ele_size = detail::get_ele_size(src->get_desc());
2005+
const auto src_offset = sycl::range<3>(
2006+
w_offset_src != 0 ? w_offset_src / from_ele_size : w_offset_src,
2007+
h_offset_src, 0);
2008+
const auto to_ele_size = detail::get_ele_size(dest->get_desc());
2009+
const auto dest_offset = sycl::range<3>(
2010+
w_offset_dest != 0 ? w_offset_dest / to_ele_size : w_offset_dest,
2011+
h_offset_dest, 0);
2012+
const auto copy_extent = sycl::range<3>(w / from_ele_size, h, 1);
2013+
q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
2014+
dest->get_handle(), dest_offset, dest->get_desc(),
2015+
copy_extent)
2016+
.wait();
20092017
}
20102018

20112019
/// Synchronously copies from image memory to the image memory, The function
@@ -2023,11 +2031,19 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
20232031
const image_mem_wrapper *src,
20242032
size_t w_offset_src, size_t h_offset_src,
20252033
size_t s, sycl::queue q = get_default_queue()) {
2026-
auto temp = (void *)sycl::malloc_device(s, q);
2027-
// TODO: Need change logic when sycl support image_mem to image_mem copy.
2028-
dpct_memcpy(temp, src, w_offset_src, h_offset_src, s, q);
2029-
dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, s, q);
2030-
sycl::free(temp, q);
2034+
const auto from_ele_size = detail::get_ele_size(src->get_desc());
2035+
const auto src_offset = sycl::range<3>(
2036+
w_offset_src != 0 ? w_offset_src / from_ele_size : w_offset_src,
2037+
h_offset_src, 0);
2038+
const auto to_ele_size = detail::get_ele_size(dest->get_desc());
2039+
const auto dest_offset = sycl::range<3>(
2040+
w_offset_dest != 0 ? w_offset_dest / to_ele_size : w_offset_dest,
2041+
h_offset_dest, 0);
2042+
const auto copy_extent = sycl::range<3>(s / from_ele_size, 1, 1);
2043+
q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
2044+
dest->get_handle(), dest_offset, dest->get_desc(),
2045+
copy_extent)
2046+
.wait();
20312047
}
20322048
// A wrapper for sycl fetch_image function for the byte addressing image.
20332049
template <typename DataT, typename HintT = DataT, typename CoordT>

0 commit comments

Comments
 (0)