| 64 | , s_idx_(s_idx) {} |
| 65 | |
| 66 | void operator()(sycl::nd_item<2> it) const { |
| 67 | sycl::group g = it.get_group(); |
| 68 | const uint lidx = it.get_local_id(0); |
| 69 | const uint lidy = it.get_local_id(1); |
| 70 | const uint lid = lidy * g.get_local_range(0) + lidx; |
| 71 | |
| 72 | const uint zid = g.get_group_id(0) / groups_x_; |
| 73 | const uint wid = g.get_group_id(1) / groups_y_; |
| 74 | const uint groupId_x = g.get_group_id(0) - (groups_x_)*zid; |
| 75 | const uint groupId_y = g.get_group_id(1) - (groups_y_)*wid; |
| 76 | const uint xid = groupId_x * g.get_local_range(0) + lidx; |
| 77 | const uint yid = groupId_y; |
| 78 | |
| 79 | uint ids[4] = {xid, yid, zid, wid}; |
| 80 | T *optr = out_.get_pointer() + ids[3] * oInfo_.strides[3] + |
| 81 | ids[2] * oInfo_.strides[2] + ids[1] * oInfo_.strides[1] + |
| 82 | ids[0] + oInfo_.offset; |
| 83 | |
| 84 | uint *olptr = oloc_.get_pointer() + ids[3] * oInfo_.strides[3] + |
| 85 | ids[2] * oInfo_.strides[2] + ids[1] * oInfo_.strides[1] + |
| 86 | ids[0] + oInfo_.offset; |
| 87 | |
| 88 | // There is only one element per block for out |
| 89 | // There are blockDim.y elements per block for in |
| 90 | // Hence increment ids[dim] just after offseting out and before |
| 91 | // offsetting in |
| 92 | const bool rlen_valid = |
| 93 | (ids[0] < rlenInfo_.dims[0]) && (ids[1] < rlenInfo_.dims[1]) && |
| 94 | (ids[2] < rlenInfo_.dims[2]) && (ids[3] < rlenInfo_.dims[3]); |
| 95 | const bool rlen_nonnull = rlenValid_; |
| 96 | const uint *rlenptr = |
| 97 | (rlen_nonnull && rlen_valid) |
| 98 | ? rlen_.get_pointer() + ids[3] * rlenInfo_.strides[3] + |
| 99 | ids[2] * rlenInfo_.strides[2] + |
| 100 | ids[1] * rlenInfo_.strides[1] + ids[0] + rlenInfo_.offset |
| 101 | : nullptr; |
| 102 | |
| 103 | const uint groupIdx_dim = ids[dim]; |
| 104 | |
| 105 | // add thread offset for reduced dim for inputs |
| 106 | ids[dim] = ids[dim] * g.get_local_range(1) + lidy; |
| 107 | |
| 108 | const T *iptr = in_.get_pointer() + ids[3] * iInfo_.strides[3] + |
| 109 | ids[2] * iInfo_.strides[2] + |
| 110 | ids[1] * iInfo_.strides[1] + ids[0] + iInfo_.offset; |
| 111 | const uint *ilptr; |
| 112 | if (!is_first) { |
| 113 | ilptr = iloc_.get_pointer() + ids[3] * iInfo_.strides[3] + |
| 114 | ids[2] * iInfo_.strides[2] + ids[1] * iInfo_.strides[1] + |
| 115 | ids[0] + iInfo_.offset; |
| 116 | } |
| 117 | |
| 118 | const uint id_dim_in = ids[dim]; |
| 119 | const uint istride_dim = iInfo_.strides[dim]; |
| 120 | |
| 121 | size_t xlim = iInfo_.dims[0]; |
| 122 | size_t ylim = iInfo_.dims[1]; |
| 123 | size_t zlim = iInfo_.dims[2]; |