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

Method operator()

src/backend/oneapi/kernel/bilateral.hpp:55–136  ·  view source on GitHub ↗

Source from the content-addressed store, hash-verified

53 , nBBS0_(nBBS0)
54 , nBBS1_(nBBS1) {}
55 void operator()(sycl::nd_item<2> it) const {
56 sycl::group g = it.get_group();
57 const int radius = sycl::max((int)(sigma_space_ * 1.5f), 1);
58 const int padding = 2 * radius;
59 const int window_size = padding + 1;
60 const int shrdLen = g.get_local_range(0) + padding;
61 const float variance_range = sigma_color_ * sigma_color_;
62 const float variance_space = sigma_space_ * sigma_space_;
63 const float variance_space_neg2 = -2.0 * variance_space;
64 const float inv_variance_range_neg2 = -0.5 / (variance_range);
65
66 // gfor batch offsets
67 unsigned b2 = g.get_group_id(0) / nBBS0_;
68 unsigned b3 = g.get_group_id(1) / nBBS1_;
69
70 const inType* in =
71 d_src_.get_pointer() +
72 (b2 * iInfo_.strides[2] + b3 * iInfo_.strides[3] + iInfo_.offset);
73 outType* out = d_dst_.get_pointer() +
74 (b2 * oInfo_.strides[2] + b3 * oInfo_.strides[3]);
75
76 int lx = it.get_local_id(0);
77 int ly = it.get_local_id(1);
78
79 const int gx =
80 g.get_local_range(0) * (g.get_group_id(0) - b2 * nBBS0_) + lx;
81 const int gy =
82 g.get_local_range(1) * (g.get_group_id(1) - b3 * nBBS1_) + ly;
83
84 // generate gauss2d_ spatial variance values for block
85 if (lx < window_size && ly < window_size) {
86 int x = lx - radius;
87 int y = ly - radius;
88 gauss2d_[ly * window_size + lx] =
89 exp_native_nonnative<outType, USE_NATIVE_EXP>(
90 ((x * x) + (y * y)) / variance_space_neg2);
91 }
92
93 int s0 = iInfo_.strides[0];
94 int s1 = iInfo_.strides[1];
95 int d0 = iInfo_.dims[0];
96 int d1 = iInfo_.dims[1];
97 // pull image to local memory
98 for (int b = ly, gy2 = gy; b < shrdLen;
99 b += g.get_local_range(1), gy2 += g.get_local_range(1)) {
100 // move row_set g.get_local_range(1) along coloumns
101 for (int a = lx, gx2 = gx; a < shrdLen;
102 a += g.get_local_range(0), gx2 += g.get_local_range(0)) {
103 load2LocalMem(localMem_, in, a, b, shrdLen, d0, d1,
104 gx2 - radius, gy2 - radius, s1, s0);
105 }
106 }
107
108 it.barrier();
109
110 if (gx < iInfo_.dims[0] && gy < iInfo_.dims[1]) {
111 lx += radius;
112 ly += radius;

Callers

nothing calls this directly

Calls 1

maxFunction · 0.50

Tested by

no test coverage detected