| 187 | |
| 188 | template<typename Tk, typename To> |
| 189 | void compact(sycl::buffer<int> reduced_block_sizes, Array<Tk> &keys_out, |
| 190 | Array<To> &vals_out, const Array<Tk> &keys, const Array<To> &vals, |
| 191 | const int numBlocks, const int threads_x) { |
| 192 | sycl::range<3> local(threads_x, 1, 1); |
| 193 | sycl::range<3> global(local[0] * numBlocks, vals_out.dims()[1], |
| 194 | vals_out.dims()[2] * vals_out.dims()[3]); |
| 195 | |
| 196 | auto keys_out_get = keys_out.get(); |
| 197 | auto vals_out_get = vals_out.get(); |
| 198 | auto keys_get = keys.get(); |
| 199 | auto vals_get = vals.get(); |
| 200 | getQueue().submit([&](sycl::handler &h) { |
| 201 | read_accessor<int> reduced_block_sizes_acc{reduced_block_sizes, h}; |
| 202 | write_accessor<Tk> keys_out_acc{*keys_out_get, h}; |
| 203 | write_accessor<To> vals_out_acc{*vals_out_get, h}; |
| 204 | read_accessor<Tk> keys_acc{*keys_get, h}; |
| 205 | read_accessor<To> vals_acc{*vals_get, h}; |
| 206 | |
| 207 | h.parallel_for(sycl::nd_range<3>(global, local), |
| 208 | kernel::compactKernel<Tk, To>( |
| 209 | reduced_block_sizes_acc, keys_out_acc, keys_out, |
| 210 | vals_out_acc, vals_out, keys_acc, keys, vals_acc, |
| 211 | vals, static_cast<int>(vals_out.dims()[2]))); |
| 212 | }); |
| 213 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 214 | } |
| 215 | |
| 216 | template<typename Tk, typename To> |
| 217 | void compactDim(sycl::buffer<int> &reduced_block_sizes, Array<Tk> &keys_out, |