| 51 | , s_tmp_(s_tmp) {} |
| 52 | |
| 53 | void operator()(sycl::nd_item<2> it) const { |
| 54 | sycl::group g = it.get_group(); |
| 55 | const uint lidx = it.get_local_id(0); |
| 56 | const uint lidy = it.get_local_id(1); |
| 57 | const uint lid = lidy * g.get_local_range(0) + lidx; |
| 58 | |
| 59 | const uint zid = g.get_group_id(0) / groups_x_; |
| 60 | const uint wid = g.get_group_id(1) / groups_y_; |
| 61 | const uint groupId_x = g.get_group_id(0) - (groups_x_)*zid; |
| 62 | const uint groupId_y = g.get_group_id(1) - (groups_y_)*wid; |
| 63 | const uint xid = groupId_x * g.get_local_range(0) + lidx; |
| 64 | const uint yid = groupId_y; |
| 65 | |
| 66 | uint ids[4] = {xid, yid, zid, wid}; |
| 67 | |
| 68 | const Ti *iptr = in_acc_.get_pointer(); |
| 69 | To *optr = out_acc_.get_pointer(); |
| 70 | To *tptr = tmp_acc_.get_pointer(); |
| 71 | |
| 72 | // There is only one element per block for out |
| 73 | // There are blockDim.y elements per block for in |
| 74 | // Hence increment ids[dim] just after offseting out and before |
| 75 | // offsetting in |
| 76 | tptr += ids[3] * tInfo_.strides[3] + ids[2] * tInfo_.strides[2] + |
| 77 | ids[1] * tInfo_.strides[1] + ids[0]; |
| 78 | |
| 79 | const int groupIdx_dim = ids[dim]; |
| 80 | ids[dim] = ids[dim] * g.get_local_range(1) * lim_ + lidy; |
| 81 | |
| 82 | optr += ids[3] * oInfo_.strides[3] + ids[2] * oInfo_.strides[2] + |
| 83 | ids[1] * oInfo_.strides[1] + ids[0]; |
| 84 | iptr += ids[3] * iInfo_.strides[3] + ids[2] * iInfo_.strides[2] + |
| 85 | ids[1] * iInfo_.strides[1] + ids[0] + iInfo_.offset; |
| 86 | int id_dim = ids[dim]; |
| 87 | const int out_dim = oInfo_.dims[dim]; |
| 88 | |
| 89 | bool is_valid = (ids[0] < oInfo_.dims[0]) && |
| 90 | (ids[1] < oInfo_.dims[1]) && |
| 91 | (ids[2] < oInfo_.dims[2]) && (ids[3] < oInfo_.dims[3]); |
| 92 | |
| 93 | const int ostride_dim = oInfo_.strides[dim]; |
| 94 | const int istride_dim = iInfo_.strides[dim]; |
| 95 | |
| 96 | To *sptr = s_val_.get_pointer() + lid; |
| 97 | |
| 98 | common::Transform<Ti, To, op> transform; |
| 99 | common::Binary<To, op> binop; |
| 100 | |
| 101 | const To init = common::Binary<To, op>::init(); |
| 102 | To val = init; |
| 103 | |
| 104 | const bool isLast = (lidy == (DIMY_ - 1)); |
| 105 | |
| 106 | for (int k = 0; k < lim_; k++) { |
| 107 | if (isLast) s_tmp_[lidx] = val; |
| 108 | |
| 109 | bool cond = (is_valid) && (id_dim < out_dim); |
| 110 | val = cond ? transform(*iptr) : init; |