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

Method operator()

src/backend/oneapi/kernel/index.hpp:56–110  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

54 }
55
56 void operator()(sycl::nd_item<3> it) const {
57 // retrieve index pointers
58 // these can be 0 where af_array index is not used
59 sycl::group g = it.get_group();
60 const uint* ptr0 = p.ptr[0].get_pointer();
61 const uint* ptr1 = p.ptr[1].get_pointer();
62 const uint* ptr2 = p.ptr[2].get_pointer();
63 const uint* ptr3 = p.ptr[3].get_pointer();
64 // retrive booleans that tell us which index to use
65 const bool s0 = p.isSeq[0];
66 const bool s1 = p.isSeq[1];
67 const bool s2 = p.isSeq[2];
68 const bool s3 = p.isSeq[3];
69
70 const int gz = g.get_group_id(0) / nBBS0;
71 const int gx = g.get_local_range(0) * (g.get_group_id(0) - gz * nBBS0) +
72 it.get_local_id(0);
73
74 const int gw =
75 (g.get_group_id(1) + g.get_group_id(2) * g.get_group_range(1)) /
76 nBBS1;
77 const int gy =
78 g.get_local_range(1) * ((g.get_group_id(1) +
79 g.get_group_id(2) * g.get_group_range(1)) -
80 gw * nBBS1) +
81 it.get_local_id(1);
82
83 size_t odims0 = outp.dims[0];
84 size_t odims1 = outp.dims[1];
85 size_t odims2 = outp.dims[2];
86 size_t odims3 = outp.dims[3];
87
88 if (gx < odims0 && gy < odims1 && gz < odims2 && gw < odims3) {
89 // calculate pointer offsets for input
90 int i = p.strds[0] *
91 trimIndex(s0 ? gx * p.steps[0] + p.offs[0] : ptr0[gx],
92 inp.dims[0]);
93 int j = p.strds[1] *
94 trimIndex(s1 ? gy * p.steps[1] + p.offs[1] : ptr1[gy],
95 inp.dims[1]);
96 int k = p.strds[2] *
97 trimIndex(s2 ? gz * p.steps[2] + p.offs[2] : ptr2[gz],
98 inp.dims[2]);
99 int l = p.strds[3] *
100 trimIndex(s3 ? gw * p.steps[3] + p.offs[3] : ptr3[gw],
101 inp.dims[3]);
102 // offset input and output pointers
103 const T* src = (const T*)in.get_pointer() + (i + j + k + l);
104 T* dst = (T*)out.get_pointer() +
105 (gx * outp.strides[0] + gy * outp.strides[1] +
106 gz * outp.strides[2] + gw * outp.strides[3]);
107 // set the output
108 dst[0] = src[0];
109 }
110 }
111};
112
113template<typename T>

Callers

nothing calls this directly

Calls 1

trimIndexFunction · 0.70

Tested by

no test coverage detected