| 250 | // temporary reduced counts/weights to it. |
| 251 | template<typename Ti, typename Tw, typename To, uint DIMX> |
| 252 | __global__ static void mean_first_kernel(Param<To> out, Param<Tw> owt, |
| 253 | CParam<Ti> in, CParam<Tw> iwt, |
| 254 | uint blocks_x, uint blocks_y, |
| 255 | uint repeat) { |
| 256 | const uint tidx = threadIdx.x; |
| 257 | const uint tidy = threadIdx.y; |
| 258 | const uint tid = tidy * blockDim.x + tidx; |
| 259 | |
| 260 | const uint zid = blockIdx.x / blocks_x; |
| 261 | const uint wid = blockIdx.y / blocks_y; |
| 262 | const uint blockIdx_x = blockIdx.x - (blocks_x)*zid; |
| 263 | const uint blockIdx_y = blockIdx.y - (blocks_y)*wid; |
| 264 | const uint xid = blockIdx_x * blockDim.x * repeat + tidx; |
| 265 | const uint yid = blockIdx_y * blockDim.y + tidy; |
| 266 | |
| 267 | const Ti *iptr = in.ptr; |
| 268 | const Tw *iwptr = iwt.ptr; |
| 269 | To *optr = out.ptr; |
| 270 | Tw *owptr = owt.ptr; |
| 271 | |
| 272 | iptr += wid * in.strides[3] + zid * in.strides[2] + yid * in.strides[1]; |
| 273 | if (iwptr != NULL) |
| 274 | iwptr += |
| 275 | wid * iwt.strides[3] + zid * iwt.strides[2] + yid * iwt.strides[1]; |
| 276 | optr += wid * out.strides[3] + zid * out.strides[2] + yid * out.strides[1]; |
| 277 | if (owptr != NULL) |
| 278 | owptr += |
| 279 | wid * out.strides[3] + zid * out.strides[2] + yid * out.strides[1]; |
| 280 | |
| 281 | if (yid >= in.dims[1] || zid >= in.dims[2] || wid >= in.dims[3]) return; |
| 282 | |
| 283 | int lim = min((int)(xid + repeat * DIMX), in.dims[0]); |
| 284 | |
| 285 | common::Transform<Ti, compute_t<To>, af_add_t> transform; |
| 286 | |
| 287 | compute_t<To> val = common::Binary<compute_t<To>, af_add_t>::init(); |
| 288 | compute_t<Tw> weight = common::Binary<compute_t<Tw>, af_add_t>::init(); |
| 289 | |
| 290 | if (xid < lim) { |
| 291 | val = transform(iptr[xid]); |
| 292 | if (iwptr != NULL) { |
| 293 | weight = iwptr[xid]; |
| 294 | } else { |
| 295 | weight = (Tw)1; |
| 296 | } |
| 297 | } |
| 298 | |
| 299 | __shared__ compute_t<To> s_val[THREADS_PER_BLOCK]; |
| 300 | __shared__ compute_t<Tw> s_idx[THREADS_PER_BLOCK]; |
| 301 | |
| 302 | if (iwptr != NULL) { |
| 303 | for (int id = xid + DIMX; id < lim; id += DIMX) { |
| 304 | stable_mean(&val, &weight, transform(iptr[id]), |
| 305 | compute_t<Tw>(iwptr[id])); |
| 306 | } |
| 307 | } else { |
| 308 | for (int id = xid + DIMX; id < lim; id += DIMX) { |
| 309 | // Faster version of stable_mean when iwptr is NULL |
nothing calls this directly
no test coverage detected