| 237 | |
| 238 | template<typename T> |
| 239 | void transform(Param<T> out, const Param<T> in, const Param<float> tf, |
| 240 | bool isInverse, bool isPerspective, af_interp_type method, |
| 241 | int order) { |
| 242 | using std::string; |
| 243 | |
| 244 | using BT = typename dtype_traits<T>::base_type; |
| 245 | |
| 246 | constexpr int TX = 16; |
| 247 | constexpr int TY = 16; |
| 248 | // Used for batching images |
| 249 | constexpr int TI = 4; |
| 250 | |
| 251 | const int nImg2 = in.info.dims[2]; |
| 252 | const int nImg3 = in.info.dims[3]; |
| 253 | const int nTfs2 = tf.info.dims[2]; |
| 254 | const int nTfs3 = tf.info.dims[3]; |
| 255 | |
| 256 | auto local = sycl::range(TX, TY, 1); |
| 257 | |
| 258 | int batchImg2 = 1; |
| 259 | if (nImg2 != nTfs2) batchImg2 = fmin(nImg2, TI); |
| 260 | |
| 261 | const int blocksXPerImage = divup(out.info.dims[0], local[0]); |
| 262 | const int blocksYPerImage = divup(out.info.dims[1], local[1]); |
| 263 | |
| 264 | int global_x = local[0] * blocksXPerImage * (nImg2 / batchImg2); |
| 265 | int global_y = local[1] * blocksYPerImage * nImg3; |
| 266 | int global_z = |
| 267 | local[2] * fmax((nTfs2 / nImg2), 1) * fmax((nTfs3 / nImg3), 1); |
| 268 | |
| 269 | auto global = sycl::range(global_x, global_y, global_z); |
| 270 | |
| 271 | #define INVOKE(PERSPECTIVE, INTERP_ORDER) \ |
| 272 | h.parallel_for( \ |
| 273 | sycl::nd_range{global, local}, \ |
| 274 | transformCreateKernel<T, wtype_t<BT>, PERSPECTIVE, INTERP_ORDER>( \ |
| 275 | d_out, out.info, d_in, in.info, d_tf, tf.info, nImg2, nImg3, \ |
| 276 | nTfs2, nTfs3, batchImg2, blocksXPerImage, blocksYPerImage, method, \ |
| 277 | isInverse)); |
| 278 | |
| 279 | getQueue().submit([&](auto &h) { |
| 280 | read_accessor<T> d_in{*in.data, h}; |
| 281 | read_accessor<float> d_tf{*tf.data, h}; |
| 282 | write_accessor<T> d_out{*out.data, h}; |
| 283 | |
| 284 | if (isPerspective == true && order == 1) INVOKE(true, 1); |
| 285 | if (isPerspective == true && order == 2) INVOKE(true, 2); |
| 286 | if (isPerspective == true && order == 3) INVOKE(true, 3); |
| 287 | |
| 288 | if (isPerspective == false && order == 1) INVOKE(false, 1); |
| 289 | if (isPerspective == false && order == 2) INVOKE(false, 2); |
| 290 | if (isPerspective == false && order == 3) INVOKE(false, 3); |
| 291 | }); |
| 292 | |
| 293 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 294 | } |
| 295 | |
| 296 | } // namespace kernel |
no test coverage detected