| 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; |