| 292 | |
| 293 | template<typename inType, typename outType> |
| 294 | void copy(Param<outType> dst, const Param<inType> src, const int ndims, |
| 295 | const outType default_value, const double factor, |
| 296 | const bool same_dims) { |
| 297 | using std::string; |
| 298 | |
| 299 | sycl::range<2> local(DIM0, DIM1); |
| 300 | size_t local_size[] = {DIM0, DIM1}; |
| 301 | |
| 302 | local_size[0] *= local_size[1]; |
| 303 | if (ndims == 1) { local_size[1] = 1; } |
| 304 | |
| 305 | int blk_x = divup(dst.info.dims[0], local_size[0]); |
| 306 | int blk_y = divup(dst.info.dims[1], local_size[1]); |
| 307 | |
| 308 | sycl::range<2> global(blk_x * dst.info.dims[2] * DIM0, |
| 309 | blk_y * dst.info.dims[3] * DIM1); |
| 310 | |
| 311 | sycl::nd_range<2> ndrange(global, local); |
| 312 | |
| 313 | dims_t trgt_dims; |
| 314 | if (same_dims) { |
| 315 | trgt_dims = {{dst.info.dims[0], dst.info.dims[1], dst.info.dims[2], |
| 316 | dst.info.dims[3]}}; |
| 317 | } else { |
| 318 | dim_t trgt_l = std::min(dst.info.dims[3], src.info.dims[3]); |
| 319 | dim_t trgt_k = std::min(dst.info.dims[2], src.info.dims[2]); |
| 320 | dim_t trgt_j = std::min(dst.info.dims[1], src.info.dims[1]); |
| 321 | dim_t trgt_i = std::min(dst.info.dims[0], src.info.dims[0]); |
| 322 | trgt_dims = {{trgt_i, trgt_j, trgt_k, trgt_l}}; |
| 323 | } |
| 324 | |
| 325 | getQueue().submit([&](sycl::handler &h) { |
| 326 | write_accessor<outType> dst_acc{*dst.data, h}; |
| 327 | read_accessor<inType> src_acc{ |
| 328 | *const_cast<sycl::buffer<inType> *>(src.data), h}; |
| 329 | |
| 330 | if (same_dims) { |
| 331 | h.parallel_for(ndrange, |
| 332 | reshapeCopy<inType, outType, true>( |
| 333 | dst_acc, dst.info, src_acc, src.info, |
| 334 | default_value, factor, trgt_dims, blk_x, blk_y)); |
| 335 | } else { |
| 336 | h.parallel_for(ndrange, |
| 337 | reshapeCopy<inType, outType, false>( |
| 338 | dst_acc, dst.info, src_acc, src.info, |
| 339 | default_value, factor, trgt_dims, blk_x, blk_y)); |
| 340 | } |
| 341 | }); |
| 342 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 343 | } |
| 344 | |
| 345 | } // namespace kernel |
| 346 | } // namespace oneapi |