| 72 | , output_weight_(output_weight) {} |
| 73 | |
| 74 | void operator()(sycl::nd_item<2> it) const { |
| 75 | sycl::group g = it.get_group(); |
| 76 | const uint lidx = it.get_local_id(0); |
| 77 | const uint lidy = it.get_local_id(1); |
| 78 | const uint lid = lidy * g.get_local_range(0) + lidx; |
| 79 | |
| 80 | const uint zid = g.get_group_id(0) / groups_x_; |
| 81 | const uint wid = g.get_group_id(1) / groups_y_; |
| 82 | const uint groupIdx_x = g.get_group_id(0) - (groups_x_)*zid; |
| 83 | const uint groupIdx_y = g.get_group_id(1) - (groups_y_)*wid; |
| 84 | const uint xid = groupIdx_x * g.get_local_range(0) + lidx; |
| 85 | const uint yid = |
| 86 | groupIdx_y; // yid of output. updated for input later. |
| 87 | |
| 88 | uint ids[4] = {xid, yid, zid, wid}; |
| 89 | |
| 90 | const Ti *iptr = in_.get_pointer(); |
| 91 | To *optr = out_.get_pointer(); |
| 92 | |
| 93 | uint ooffset = ids[3] * oInfo_.strides[3] + ids[2] * oInfo_.strides[2] + |
| 94 | ids[1] * oInfo_.strides[1] + ids[0] + oInfo_.offset; |
| 95 | // There is only one element per block for out |
| 96 | // There are blockDim.y elements per block for in |
| 97 | // Hence increment ids[dim] just after offseting out and before |
| 98 | // offsetting in |
| 99 | optr += ooffset; |
| 100 | |
| 101 | const uint blockIdx_dim = ids[dim]; |
| 102 | ids[dim] = ids[dim] * g.get_local_range(1) + lidy; |
| 103 | |
| 104 | uint ioffset = ids[3] * iInfo_.strides[3] + ids[2] * iInfo_.strides[2] + |
| 105 | ids[1] * iInfo_.strides[1] + ids[0] + iInfo_.offset; |
| 106 | iptr += ioffset; |
| 107 | |
| 108 | const Tw *iwptr = nullptr; |
| 109 | Tw *owptr = nullptr; |
| 110 | |
| 111 | if (output_weight_) owptr = owt_.get_pointer() + ooffset; |
| 112 | if (input_weight_) iwptr = iwt_.get_pointer() + ioffset; |
| 113 | |
| 114 | const uint id_dim_in = ids[dim]; |
| 115 | const uint istride_dim = iInfo_.strides[dim]; |
| 116 | |
| 117 | bool is_valid = (ids[0] < iInfo_.dims[0]) && |
| 118 | (ids[1] < iInfo_.dims[1]) && |
| 119 | (ids[2] < iInfo_.dims[2]) && (ids[3] < iInfo_.dims[3]); |
| 120 | |
| 121 | common::Transform<Ti, compute_t<To>, af_add_t> transform; |
| 122 | |
| 123 | compute_t<To> val = common::Binary<compute_t<To>, af_add_t>::init(); |
| 124 | compute_t<Tw> weight = common::Binary<compute_t<Tw>, af_add_t>::init(); |
| 125 | |
| 126 | if (is_valid && id_dim_in < iInfo_.dims[dim]) { |
| 127 | val = transform(*iptr); |
| 128 | if (iwptr) { |
| 129 | weight = *iwptr; |
| 130 | } else { |
| 131 | weight = (Tw)1; |
nothing calls this directly
no test coverage detected