| 169 | |
| 170 | template<typename inType, typename outType> |
| 171 | void bilateral(Param<outType> out, const Param<inType> in, const float s_sigma, |
| 172 | const float c_sigma) { |
| 173 | constexpr int THREADS_X = 16; |
| 174 | constexpr int THREADS_Y = 16; |
| 175 | constexpr bool UseNativeExp = !std::is_same<inType, double>::value || |
| 176 | std::is_same<inType, cdouble>::value; |
| 177 | |
| 178 | auto local = sycl::range{THREADS_X, THREADS_Y}; |
| 179 | |
| 180 | const int blk_x = divup(in.info.dims[0], THREADS_X); |
| 181 | const int blk_y = divup(in.info.dims[1], THREADS_Y); |
| 182 | |
| 183 | auto global = sycl::range{(size_t)(blk_x * in.info.dims[2] * THREADS_X), |
| 184 | (size_t)(blk_y * in.info.dims[3] * THREADS_Y)}; |
| 185 | |
| 186 | // calculate local memory size |
| 187 | int radius = (int)std::max(s_sigma * 1.5f, 1.f); |
| 188 | int num_shrd_elems = (THREADS_X + 2 * radius) * (THREADS_Y + 2 * radius); |
| 189 | int num_gauss_elems = (2 * radius + 1) * (2 * radius + 1); |
| 190 | size_t localMemSize = (num_shrd_elems + num_gauss_elems) * sizeof(outType); |
| 191 | size_t MaxLocalSize = |
| 192 | getQueue().get_device().get_info<sycl::info::device::local_mem_size>(); |
| 193 | if (localMemSize > MaxLocalSize) { |
| 194 | char errMessage[256]; |
| 195 | snprintf(errMessage, sizeof(errMessage), |
| 196 | "\nOneAPI Bilateral filter doesn't support %f spatial sigma\n", |
| 197 | s_sigma); |
| 198 | ONEAPI_NOT_SUPPORTED(errMessage); |
| 199 | } |
| 200 | |
| 201 | getQueue().submit([&](sycl::handler& h) { |
| 202 | read_accessor<inType> inAcc{*in.data, h}; |
| 203 | write_accessor<outType> outAcc{*out.data, h}; |
| 204 | |
| 205 | auto localMem = sycl::local_accessor<outType, 1>(num_shrd_elems, h); |
| 206 | auto gauss2d = sycl::local_accessor<outType, 1>(num_shrd_elems, h); |
| 207 | |
| 208 | h.parallel_for(sycl::nd_range{global, local}, |
| 209 | bilateralKernel<outType, inType, UseNativeExp>( |
| 210 | outAcc, out.info, inAcc, in.info, localMem, gauss2d, |
| 211 | s_sigma, c_sigma, num_shrd_elems, blk_x, blk_y)); |
| 212 | }); |
| 213 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 214 | } |
| 215 | |
| 216 | } // namespace kernel |
| 217 | } // namespace oneapi |