| 134 | |
| 135 | template<typename T> |
| 136 | void unwrap(Param<T> out, const Param<T> in, const dim_t wx, const dim_t wy, |
| 137 | const dim_t sx, const dim_t sy, const dim_t px, const dim_t py, |
| 138 | const dim_t dx, const dim_t dy, const dim_t nx, |
| 139 | const bool IS_COLUMN) { |
| 140 | dim_t TX = 1, TY = 1; |
| 141 | dim_t BX = 1; |
| 142 | const dim_t BY = out.info.dims[2] * out.info.dims[3]; |
| 143 | int reps = 1; |
| 144 | |
| 145 | if (IS_COLUMN) { |
| 146 | TX = std::min(THREADS_PER_BLOCK, nextpow2(out.info.dims[0])); |
| 147 | TY = THREADS_PER_BLOCK / TX; |
| 148 | BX = divup(out.info.dims[1], TY); |
| 149 | reps = divup((wx * wy), TX); |
| 150 | } else { |
| 151 | TX = THREADS_X; |
| 152 | TY = THREADS_Y; |
| 153 | BX = divup(out.info.dims[0], TX); |
| 154 | reps = divup((wx * wy), TY); |
| 155 | } |
| 156 | |
| 157 | auto local = sycl::range(TX, TY); |
| 158 | auto global = sycl::range(local[0] * BX, local[1] * BY); |
| 159 | |
| 160 | getQueue().submit([&](auto &h) { |
| 161 | sycl::accessor d_out{*out.data, h, sycl::write_only, sycl::no_init}; |
| 162 | sycl::accessor d_in{*in.data, h, sycl::read_only}; |
| 163 | h.parallel_for( |
| 164 | sycl::nd_range{global, local}, |
| 165 | unwrapCreateKernel<T>(d_out, out.info, d_in, in.info, wx, wy, sx, |
| 166 | sy, px, py, dx, dy, nx, reps, IS_COLUMN)); |
| 167 | }); |
| 168 | |
| 169 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 170 | } |
| 171 | |
| 172 | } // namespace kernel |
| 173 | } // namespace oneapi |