| 55 | |
| 56 | template<typename T> |
| 57 | class memCopy { |
| 58 | public: |
| 59 | memCopy(write_accessor<T> out, dims_t ostrides, int ooffset, |
| 60 | read_accessor<T> in, dims_t idims, dims_t istrides, int ioffset, |
| 61 | int groups_0, int groups_1) |
| 62 | : out_(out) |
| 63 | , ostrides_(ostrides) |
| 64 | , ooffset_(ooffset) |
| 65 | , in_(in) |
| 66 | , idims_(idims) |
| 67 | , istrides_(istrides) |
| 68 | , ioffset_(ioffset) |
| 69 | , groups_0_(groups_0) |
| 70 | , groups_1_(groups_1) {} |
| 71 | |
| 72 | void operator()(sycl::nd_item<2> it) const { |
| 73 | const int lid0 = it.get_local_id(0); |
| 74 | const int lid1 = it.get_local_id(1); |
| 75 | |
| 76 | sycl::group gg = it.get_group(); |
| 77 | const int id2 = gg.get_group_id(0) / groups_0_; |
| 78 | const int id3 = gg.get_group_id(1) / groups_1_; |
| 79 | const int group_id_0 = gg.get_group_id(0) - groups_0_ * id2; |
| 80 | const int group_id_1 = gg.get_group_id(1) - groups_1_ * id3; |
| 81 | const int id0 = group_id_0 * gg.get_local_range(0) + lid0; |
| 82 | const int id1 = group_id_1 * gg.get_local_range(1) + lid1; |
| 83 | |
| 84 | const T *iptr = in_.get_pointer(); |
| 85 | // FIXME: Do more work per work group |
| 86 | |
| 87 | T *optr = out_.get_pointer(); |
| 88 | optr += id3 * ostrides_.dim[3] + id2 * ostrides_.dim[2] + |
| 89 | id1 * ostrides_.dim[1] + ooffset_; |
| 90 | iptr += id3 * istrides_.dim[3] + id2 * istrides_.dim[2] + |
| 91 | id1 * istrides_.dim[1] + ioffset_; |
| 92 | |
| 93 | int istride0 = istrides_.dim[0]; |
| 94 | size_t idd0 = idims_.dim[0]; |
| 95 | size_t idd1 = idims_.dim[1]; |
| 96 | size_t idd2 = idims_.dim[2]; |
| 97 | size_t idd3 = idims_.dim[3]; |
| 98 | |
| 99 | if (id0 < idd0 && id1 < idd1 && id2 < idd2 && id3 < idd3) { |
| 100 | optr[id0] = iptr[id0 * istride0]; |
| 101 | } |
| 102 | } |
| 103 | |
| 104 | protected: |
| 105 | write_accessor<T> out_; |
| 106 | dims_t ostrides_; |
| 107 | int ooffset_; |
| 108 | read_accessor<T> in_; |
| 109 | dims_t idims_, istrides_; |
| 110 | int ioffset_, groups_0_, groups_1_; |
| 111 | }; |
| 112 | |
| 113 | constexpr uint DIM0 = 32; |
| 114 | constexpr uint DIM1 = 8; |