| 215 | |
| 216 | template<typename Tk, typename To> |
| 217 | void compactDim(sycl::buffer<int> &reduced_block_sizes, Array<Tk> &keys_out, |
| 218 | Array<To> &vals_out, const Array<Tk> &keys, |
| 219 | const Array<To> &vals, const int numBlocks, const int threads_x, |
| 220 | const int dim, std::vector<int> dim_ordering) { |
| 221 | sycl::range<3> local(threads_x, 1, 1); |
| 222 | sycl::range<3> global( |
| 223 | local[0] * numBlocks, vals_out.dims()[dim_ordering[1]], |
| 224 | vals_out.dims()[dim_ordering[2]] * vals_out.dims()[dim_ordering[3]]); |
| 225 | |
| 226 | auto keys_out_get = keys_out.get(); |
| 227 | auto vals_out_get = vals_out.get(); |
| 228 | auto keys_get = keys.get(); |
| 229 | auto vals_get = vals.get(); |
| 230 | getQueue().submit([&](sycl::handler &h) { |
| 231 | read_accessor<int> reduced_block_sizes_acc{reduced_block_sizes, h}; |
| 232 | write_accessor<Tk> keys_out_acc{*keys_out_get, h}; |
| 233 | write_accessor<To> vals_out_acc{*vals_out_get, h}; |
| 234 | read_accessor<Tk> keys_acc{*keys_get, h}; |
| 235 | read_accessor<To> vals_acc{*vals_get, h}; |
| 236 | |
| 237 | h.parallel_for( |
| 238 | sycl::nd_range<3>(global, local), |
| 239 | kernel::compactDimKernel<Tk, To>( |
| 240 | reduced_block_sizes_acc, keys_out_acc, keys_out, vals_out_acc, |
| 241 | vals_out, keys_acc, keys, vals_acc, vals, |
| 242 | static_cast<int>(vals_out.dims()[dim_ordering[2]]), dim)); |
| 243 | }); |
| 244 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 245 | } |
| 246 | |
| 247 | template<typename Tk> |
| 248 | void testNeedsReduction(sycl::buffer<int> needs_reduction, |