| 54 | } |
| 55 | |
| 56 | void operator()(sycl::nd_item<3> it) const { |
| 57 | // retrieve index pointers |
| 58 | // these can be 0 where af_array index is not used |
| 59 | sycl::group g = it.get_group(); |
| 60 | const uint* ptr0 = p.ptr[0].get_pointer(); |
| 61 | const uint* ptr1 = p.ptr[1].get_pointer(); |
| 62 | const uint* ptr2 = p.ptr[2].get_pointer(); |
| 63 | const uint* ptr3 = p.ptr[3].get_pointer(); |
| 64 | // retrive booleans that tell us which index to use |
| 65 | const bool s0 = p.isSeq[0]; |
| 66 | const bool s1 = p.isSeq[1]; |
| 67 | const bool s2 = p.isSeq[2]; |
| 68 | const bool s3 = p.isSeq[3]; |
| 69 | |
| 70 | const int gz = g.get_group_id(0) / nBBS0; |
| 71 | const int gx = g.get_local_range(0) * (g.get_group_id(0) - gz * nBBS0) + |
| 72 | it.get_local_id(0); |
| 73 | |
| 74 | const int gw = |
| 75 | (g.get_group_id(1) + g.get_group_id(2) * g.get_group_range(1)) / |
| 76 | nBBS1; |
| 77 | const int gy = |
| 78 | g.get_local_range(1) * ((g.get_group_id(1) + |
| 79 | g.get_group_id(2) * g.get_group_range(1)) - |
| 80 | gw * nBBS1) + |
| 81 | it.get_local_id(1); |
| 82 | |
| 83 | size_t odims0 = outp.dims[0]; |
| 84 | size_t odims1 = outp.dims[1]; |
| 85 | size_t odims2 = outp.dims[2]; |
| 86 | size_t odims3 = outp.dims[3]; |
| 87 | |
| 88 | if (gx < odims0 && gy < odims1 && gz < odims2 && gw < odims3) { |
| 89 | // calculate pointer offsets for input |
| 90 | int i = p.strds[0] * |
| 91 | trimIndex(s0 ? gx * p.steps[0] + p.offs[0] : ptr0[gx], |
| 92 | inp.dims[0]); |
| 93 | int j = p.strds[1] * |
| 94 | trimIndex(s1 ? gy * p.steps[1] + p.offs[1] : ptr1[gy], |
| 95 | inp.dims[1]); |
| 96 | int k = p.strds[2] * |
| 97 | trimIndex(s2 ? gz * p.steps[2] + p.offs[2] : ptr2[gz], |
| 98 | inp.dims[2]); |
| 99 | int l = p.strds[3] * |
| 100 | trimIndex(s3 ? gw * p.steps[3] + p.offs[3] : ptr3[gw], |
| 101 | inp.dims[3]); |
| 102 | // offset input and output pointers |
| 103 | const T* src = (const T*)in.get_pointer() + (i + j + k + l); |
| 104 | T* dst = (T*)out.get_pointer() + |
| 105 | (gx * outp.strides[0] + gy * outp.strides[1] + |
| 106 | gz * outp.strides[2] + gw * outp.strides[3]); |
| 107 | // set the output |
| 108 | dst[0] = src[0]; |
| 109 | } |
| 110 | } |
| 111 | }; |
| 112 | |
| 113 | template<typename T> |
nothing calls this directly
no test coverage detected