| 332 | , output_weight_(output_weight) {} |
| 333 | |
| 334 | void operator()(sycl::nd_item<2> it) const { |
| 335 | sycl::group g = it.get_group(); |
| 336 | const uint lidx = it.get_local_id(0); |
| 337 | const uint lidy = it.get_local_id(1); |
| 338 | const uint lid = lidy * DIMX_ + lidx; |
| 339 | |
| 340 | const uint zid = g.get_group_id(0) / groups_x_; |
| 341 | const uint wid = g.get_group_id(1) / groups_y_; |
| 342 | const uint groupIdx_x = g.get_group_id(0) - (groups_x_)*zid; |
| 343 | const uint groupIdx_y = g.get_group_id(1) - (groups_y_)*wid; |
| 344 | const uint xid = groupIdx_x * g.get_local_range(0) * repeat_ + lidx; |
| 345 | const uint yid = groupIdx_y * g.get_local_range(1) + lidy; |
| 346 | |
| 347 | const Ti *iptr = in_.get_pointer(); |
| 348 | To *optr = out_.get_pointer(); |
| 349 | |
| 350 | iptr += wid * iInfo_.strides[3] + zid * iInfo_.strides[2] + |
| 351 | yid * iInfo_.strides[1] + iInfo_.offset; |
| 352 | optr += wid * oInfo_.strides[3] + zid * oInfo_.strides[2] + |
| 353 | yid * oInfo_.strides[1] + oInfo_.offset; |
| 354 | |
| 355 | const Tw *iwptr = nullptr; |
| 356 | Tw *owptr = nullptr; |
| 357 | if (input_weight_) |
| 358 | iwptr = iwt_.get_pointer() + wid * iwInfo_.strides[3] + |
| 359 | zid * iwInfo_.strides[2] + yid * iwInfo_.strides[1] + |
| 360 | iwInfo_.offset; |
| 361 | |
| 362 | if (output_weight_) |
| 363 | owptr = owt_.get_pointer() + wid * owInfo_.strides[3] + |
| 364 | zid * owInfo_.strides[2] + yid * owInfo_.strides[1] + |
| 365 | owInfo_.offset; |
| 366 | |
| 367 | bool cond = (yid < iInfo_.dims[1] && zid < iInfo_.dims[2] && |
| 368 | wid < iInfo_.dims[3]); |
| 369 | |
| 370 | int lim = min((dim_t)(xid + repeat_ * DIMX_), iInfo_.dims[0]); |
| 371 | |
| 372 | common::Transform<Ti, compute_t<To>, af_add_t> transform; |
| 373 | |
| 374 | compute_t<To> val = common::Binary<compute_t<To>, af_add_t>::init(); |
| 375 | compute_t<Tw> weight = common::Binary<compute_t<Tw>, af_add_t>::init(); |
| 376 | |
| 377 | if (cond && xid < lim) { |
| 378 | val = transform(iptr[xid]); |
| 379 | if (input_weight_) { |
| 380 | weight = iwptr[xid]; |
| 381 | } else { |
| 382 | weight = (Tw)1; |
| 383 | } |
| 384 | } |
| 385 | |
| 386 | if (input_weight_) { |
| 387 | for (int id = xid + DIMX_; cond && id < lim; id += DIMX_) { |
| 388 | stable_mean(&val, &weight, transform(iptr[id]), |
| 389 | compute_t<Tw>(iwptr[id])); |
| 390 | } |
| 391 | } else { |
nothing calls this directly
no test coverage detected