| 25 | |
| 26 | template<typename T> |
| 27 | void assign(Array<T>& out, const af_index_t idxrs[], const Array<T>& rhs) { |
| 28 | AssignKernelParam p; |
| 29 | std::vector<af_seq> seqs(4, af_span); |
| 30 | // create seq vector to retrieve output |
| 31 | // dimensions, offsets & offsets |
| 32 | for (dim_t x = 0; x < 4; ++x) { |
| 33 | if (idxrs[x].isSeq) { seqs[x] = idxrs[x].idx.seq; } |
| 34 | } |
| 35 | |
| 36 | // retrieve dimensions, strides and offsets |
| 37 | const dim4& dDims = out.dims(); |
| 38 | // retrieve dimensions & strides for array |
| 39 | // to which rhs is being copied to |
| 40 | dim4 dstOffs = toOffset(seqs, dDims); |
| 41 | dim4 dstStrds = toStride(seqs, dDims); |
| 42 | |
| 43 | for (dim_t i = 0; i < 4; ++i) { |
| 44 | p.isSeq[i] = idxrs[i].isSeq; |
| 45 | p.offs[i] = dstOffs[i]; |
| 46 | p.strds[i] = dstStrds[i]; |
| 47 | } |
| 48 | |
| 49 | sycl::buffer<uint>* bPtrs[4]; |
| 50 | |
| 51 | std::vector<Array<uint>> idxArrs(4, createEmptyArray<uint>(dim4())); |
| 52 | // look through indexs to read af_array indexs |
| 53 | for (dim_t x = 0; x < 4; ++x) { |
| 54 | // set index pointers were applicable |
| 55 | if (!p.isSeq[x]) { |
| 56 | idxArrs[x] = castArray<uint>(idxrs[x].idx.arr); |
| 57 | bPtrs[x] = idxArrs[x].get(); |
| 58 | } else { |
| 59 | // alloc an 1-element buffer to avoid OpenCL from failing using |
| 60 | // direct buffer allocation as opposed to mem manager to avoid |
| 61 | // reference count desprepancies between different backends |
| 62 | static auto* empty = new sycl::buffer<uint>(sycl::range{1}); |
| 63 | bPtrs[x] = empty; |
| 64 | } |
| 65 | } |
| 66 | |
| 67 | kernel::assign<T>(out, rhs, p, bPtrs); |
| 68 | return; |
| 69 | } |
| 70 | |
| 71 | #define INSTANTIATE(T) \ |
| 72 | template void assign<T>(Array<T> & out, const af_index_t idxrs[], \ |