| 180 | |
| 181 | template<typename T> |
| 182 | void resize(Param<T> out, const Param<T> in, const af_interp_type method) { |
| 183 | constexpr int RESIZE_TX = 16; |
| 184 | constexpr int RESIZE_TY = 16; |
| 185 | |
| 186 | auto local = sycl::range(RESIZE_TX, RESIZE_TY); |
| 187 | |
| 188 | int blocksPerMatX = divup(out.info.dims[0], local[0]); |
| 189 | int blocksPerMatY = divup(out.info.dims[1], local[1]); |
| 190 | auto global = sycl::range(local[0] * blocksPerMatX * in.info.dims[2], |
| 191 | local[1] * blocksPerMatY * in.info.dims[3]); |
| 192 | |
| 193 | double xd = (double)in.info.dims[0] / (double)out.info.dims[0]; |
| 194 | double yd = (double)in.info.dims[1] / (double)out.info.dims[1]; |
| 195 | |
| 196 | float xf = (float)xd, yf = (float)yd; |
| 197 | |
| 198 | getQueue().submit([&](auto& h) { |
| 199 | read_accessor<T> d_in{*in.data, h}; |
| 200 | write_accessor<T> d_out{*out.data, h}; |
| 201 | switch (method) { |
| 202 | case AF_INTERP_NEAREST: |
| 203 | h.parallel_for(sycl::nd_range{global, local}, |
| 204 | resizeCreateKernel<T, AF_INTERP_NEAREST>( |
| 205 | d_out, out.info, d_in, in.info, |
| 206 | blocksPerMatX, blocksPerMatY, xf, yf)); |
| 207 | break; |
| 208 | case AF_INTERP_BILINEAR: |
| 209 | h.parallel_for(sycl::nd_range{global, local}, |
| 210 | resizeCreateKernel<T, AF_INTERP_BILINEAR>( |
| 211 | d_out, out.info, d_in, in.info, |
| 212 | blocksPerMatX, blocksPerMatY, xf, yf)); |
| 213 | break; |
| 214 | case AF_INTERP_LOWER: |
| 215 | h.parallel_for(sycl::nd_range{global, local}, |
| 216 | resizeCreateKernel<T, AF_INTERP_LOWER>( |
| 217 | d_out, out.info, d_in, in.info, |
| 218 | blocksPerMatX, blocksPerMatY, xf, yf)); |
| 219 | break; |
| 220 | default: break; |
| 221 | } |
| 222 | }); |
| 223 | |
| 224 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 225 | } |
| 226 | } // namespace kernel |
| 227 | } // namespace oneapi |
| 228 | } // namespace arrayfire |