| 49 | |
| 50 | template<typename T> |
| 51 | Array<T> copyArray(const Array<T> &A) { |
| 52 | Array<T> out = createEmptyArray<T>(A.dims()); |
| 53 | if (A.elements() == 0) { return out; } |
| 54 | |
| 55 | dim_t offset = A.getOffset(); |
| 56 | if (A.isReady()) { |
| 57 | if (A.isLinear()) { |
| 58 | // FIXME: Add checks |
| 59 | |
| 60 | sycl::buffer<T> *A_buf = A.get(); |
| 61 | sycl::buffer<T> *out_buf = out.get(); |
| 62 | |
| 63 | size_t aelem = A.elements(); |
| 64 | getQueue().submit([&](sycl::handler &h) { |
| 65 | range rr(aelem); |
| 66 | id offset_id(offset); |
| 67 | accessor offset_acc_A = |
| 68 | A_buf->template get_access<access_mode::read>(h, rr, |
| 69 | offset_id); |
| 70 | accessor acc_out = |
| 71 | out_buf->template get_access<access_mode::write>(h); |
| 72 | |
| 73 | h.copy(offset_acc_A, acc_out); |
| 74 | }); |
| 75 | } else { |
| 76 | kernel::memcopy<T>(out.get(), out.strides().get(), A.get(), |
| 77 | A.dims().get(), A.strides().get(), offset, |
| 78 | (uint)A.ndims()); |
| 79 | } |
| 80 | } else { |
| 81 | Param info = {out.get(), |
| 82 | {{A.dims().dims[0], A.dims().dims[1], A.dims().dims[2], |
| 83 | A.dims().dims[3]}, |
| 84 | {out.strides().dims[0], out.strides().dims[1], |
| 85 | out.strides().dims[2], out.strides().dims[3]}, |
| 86 | 0}}; |
| 87 | evalNodes(info, A.getNode().get()); |
| 88 | } |
| 89 | return out; |
| 90 | } |
| 91 | |
| 92 | template<typename T> |
| 93 | void multiply_inplace(Array<T> &in, double val) { |
no test coverage detected