MCPcopy Create free account
hub / github.com/arrayfire/arrayfire / operator()

Method operator()

src/backend/oneapi/kernel/mean.hpp:74–188  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

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;

Callers

nothing calls this directly

Calls 3

transformFunction · 0.70
stable_meanFunction · 0.70
initFunction · 0.50

Tested by

no test coverage detected