Skip to content

Commit 4d1b7c4

Browse files
TejaX-AlaghariTejaX-Alaghari
committed
Fixed bindless_images bugs with copy_extent
1 parent ff80033 commit 4d1b7c4

File tree

1 file changed

+34
-18
lines changed

1 file changed

+34
-18
lines changed

Diff for: ‎clang/runtime/dpct-rt/include/dpct/bindless_images.hpp

+34-18
Original file line numberDiff line numberDiff line change
@@ -963,9 +963,9 @@ dpct_memcpy(sycl::ext::oneapi::experimental::image_mem_handle src,
963963
sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
964964
const auto dest_offset = sycl::range<3>(0, 0, 0);
965965
const auto dest_extend = sycl::range<3>(p / ele_size, 0, 0);
966-
const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
966+
const auto copy_extend = sycl::range<3>(w / ele_size, h, 1);
967967
return q.ext_oneapi_copy(src, src_offset, desc_src, dest, dest_offset,
968-
dest_extend, copy_extend);
968+
desc_src, p, copy_extend);
969969
}
970970

971971
static inline std::vector<sycl::event> dpct_memcpy_to_host(
@@ -983,7 +983,7 @@ static inline std::vector<sycl::event> dpct_memcpy_to_host(
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);
985985
const auto copy_extend =
986-
sycl::range<3>((w - w_offset_src) / ele_size, 1, 0);
986+
sycl::range<3>((w - w_offset_src) / ele_size, 1, 1);
987987
event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
988988
dest_host_ptr, dest_offset,
989989
dest_extend, copy_extend));
@@ -996,7 +996,7 @@ static inline std::vector<sycl::event> dpct_memcpy_to_host(
996996
sycl::range<3>(w_offset_src / ele_size, h_offset_src, 0);
997997
const auto dest_offset = sycl::range<3>(offset_dest / ele_size, 0, 0);
998998
const auto dest_extend = sycl::range<3>(0, 0, 0);
999-
const auto copy_extend = sycl::range<3>((s - offset_dest) / ele_size, 1, 0);
999+
const auto copy_extend = sycl::range<3>((s - offset_dest) / ele_size, 1, 1);
10001000
event_list.push_back(q.ext_oneapi_copy(src, src_offset, desc_src,
10011001
dest_host_ptr, dest_offset,
10021002
dest_extend, copy_extend));
@@ -1035,7 +1035,7 @@ dpct_memcpy(const void *src,
10351035
const auto src_extend = sycl::range<3>(p / ele_size, 0, 0);
10361036
const auto dest_offset =
10371037
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
1038-
const auto copy_extend = sycl::range<3>(w / ele_size, h, 0);
1038+
const auto copy_extend = sycl::range<3>(w / ele_size, h, 1);
10391039
// TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
10401040
return q.ext_oneapi_copy(const_cast<void *>(src), src_offset, src_extend,
10411041
dest, dest_offset, desc_dest, copy_extend);
@@ -1056,7 +1056,7 @@ static inline std::vector<sycl::event> dpct_memcpy_from_host(
10561056
const auto dest_offset =
10571057
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
10581058
const auto copy_extend =
1059-
sycl::range<3>((w - w_offset_dest) / ele_size, 1, 0);
1059+
sycl::range<3>((w - w_offset_dest) / 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,
@@ -1070,7 +1070,7 @@ static inline std::vector<sycl::event> dpct_memcpy_from_host(
10701070
const auto src_extend = sycl::range<3>(0, 0, 0);
10711071
const auto dest_offset =
10721072
sycl::range<3>(w_offset_dest / ele_size, h_offset_dest, 0);
1073-
const auto copy_extend = sycl::range<3>((s - offset_src) / ele_size, 1, 0);
1073+
const auto copy_extend = sycl::range<3>((s - offset_src) / ele_size, 1, 1);
10741074
// TODO: Remove const_cast after refining the signature of ext_oneapi_copy.
10751075
event_list.push_back(q.ext_oneapi_copy(
10761076
const_cast<void *>(src_host_ptr), src_offset, src_extend, dest,
@@ -1115,7 +1115,7 @@ dpct_memcpy(const image_mem_wrapper *src, const sycl::id<3> &src_id,
11151115
sycl::range<3>(dest.get_pitch() / ele_size, dest.get_y(), 1);
11161116
const auto copy_extend = sycl::range<3>(
11171117
copy_x_size_byte != 0 ? copy_x_size_byte / ele_size : size[0], size[1],
1118-
size[2]);
1118+
size[2] != 0 ? size[2] : 1);
11191119
return q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
11201120
dest.get_data_ptr(), dest_offset, dest_extend,
11211121
copy_extend);
@@ -1946,11 +1946,19 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
19461946
size_t w_offset_src, size_t h_offset_src,
19471947
size_t w, size_t h,
19481948
sycl::queue q = get_default_queue()) {
1949-
auto temp = (void *)sycl::malloc_device(w * h, q);
1950-
// TODO: Need change logic when sycl support image_mem to image_mem copy.
1951-
dpct_memcpy(temp, w, src, w_offset_src, h_offset_src, w, h, q);
1952-
dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, w, w, h, q);
1953-
sycl::free(temp, q);
1949+
const auto from_ele_size = detail::get_ele_size(src->get_desc());
1950+
const auto src_offset = sycl::range<3>(
1951+
w_offset_src != 0 ? w_offset_src / from_ele_size : w_offset_src,
1952+
h_offset_src, 0);
1953+
const auto to_ele_size = detail::get_ele_size(dest->get_desc());
1954+
const auto dest_offset = sycl::range<3>(
1955+
w_offset_dest != 0 ? w_offset_dest / to_ele_size : w_offset_dest,
1956+
h_offset_dest, 0);
1957+
const auto copy_extent = sycl::range<3>(w / from_ele_size, h, 1);
1958+
q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
1959+
dest->get_handle(), dest_offset, dest->get_desc(),
1960+
copy_extent)
1961+
.wait();
19541962
}
19551963

19561964
/// Synchronously copies from image memory to the image memory, The function
@@ -1968,11 +1976,19 @@ static inline void dpct_memcpy(image_mem_wrapper *dest, size_t w_offset_dest,
19681976
const image_mem_wrapper *src,
19691977
size_t w_offset_src, size_t h_offset_src,
19701978
size_t s, sycl::queue q = get_default_queue()) {
1971-
auto temp = (void *)sycl::malloc_device(s, q);
1972-
// TODO: Need change logic when sycl support image_mem to image_mem copy.
1973-
dpct_memcpy(temp, src, w_offset_src, h_offset_src, s, q);
1974-
dpct_memcpy(dest, w_offset_dest, h_offset_dest, temp, s, q);
1975-
sycl::free(temp, q);
1979+
const auto from_ele_size = detail::get_ele_size(src->get_desc());
1980+
const auto src_offset = sycl::range<3>(
1981+
w_offset_src != 0 ? w_offset_src / from_ele_size : w_offset_src,
1982+
h_offset_src, 0);
1983+
const auto to_ele_size = detail::get_ele_size(dest->get_desc());
1984+
const auto dest_offset = sycl::range<3>(
1985+
w_offset_dest != 0 ? w_offset_dest / to_ele_size : w_offset_dest,
1986+
h_offset_dest, 0);
1987+
const auto copy_extent = sycl::range<3>(s / from_ele_size, 1, 1);
1988+
q.ext_oneapi_copy(src->get_handle(), src_offset, src->get_desc(),
1989+
dest->get_handle(), dest_offset, dest->get_desc(),
1990+
copy_extent)
1991+
.wait();
19761992
}
19771993
// A wrapper for sycl fetch_image function for the byte addressing image.
19781994
template <typename DataT, typename HintT = DataT, typename CoordT>

0 commit comments

Comments
 (0)