@@ -2024,7 +2024,10 @@ tri(sycl::queue &exec_q,
2024
2024
}
2025
2025
2026
2026
// check that arrays do not overlap, and concurrent copying is safe.
2027
+ char *src_data = src.get_data ();
2028
+ char *dst_data = dst.get_data ();
2027
2029
auto src_offsets = src.get_minmax_offsets ();
2030
+ auto dst_offsets = dst.get_minmax_offsets ();
2028
2031
int src_elem_size = src.get_elemsize ();
2029
2032
int dst_elem_size = dst.get_elemsize ();
2030
2033
@@ -2112,9 +2115,8 @@ tri(sycl::queue &exec_q,
2112
2115
int nd = src_nd - 2 ;
2113
2116
const py::ssize_t *shape = src_shape;
2114
2117
const py::ssize_t *p_src_strides = src_strides.data ();
2115
- ;
2116
2118
const py::ssize_t *p_dst_strides = dst_strides.data ();
2117
- ;
2119
+
2118
2120
simplify_iteration_space (nd, shape, p_src_strides, src_itemsize,
2119
2121
is_src_c_contig, is_src_f_contig, p_dst_strides,
2120
2122
dst_itemsize, is_dst_c_contig, is_dst_f_contig,
@@ -2160,14 +2162,14 @@ tri(sycl::queue &exec_q,
2160
2162
if (part == ' l' ) {
2161
2163
auto fn = tril_generic_dispatch_vector[src_typeid];
2162
2164
tri_ev =
2163
- fn (exec_q, inner_range, outer_range, src. get_data (), dst. get_data () ,
2164
- nd, dev_shape_and_strides, k, depends, {copy_shape_and_strides});
2165
+ fn (exec_q, inner_range, outer_range, src_data, dst_data, nd ,
2166
+ dev_shape_and_strides, k, depends, {copy_shape_and_strides});
2165
2167
}
2166
2168
else {
2167
2169
auto fn = triu_generic_dispatch_vector[src_typeid];
2168
2170
tri_ev =
2169
- fn (exec_q, inner_range, outer_range, src. get_data (), dst. get_data () ,
2170
- nd, dev_shape_and_strides, k, depends, {copy_shape_and_strides});
2171
+ fn (exec_q, inner_range, outer_range, src_data, dst_data, nd ,
2172
+ dev_shape_and_strides, k, depends, {copy_shape_and_strides});
2171
2173
}
2172
2174
2173
2175
exec_q.submit ([&](sycl::handler &cgh) {
0 commit comments