| 160 | |
| 161 | template<typename Tk, typename To, af_op_t op> |
| 162 | void finalBoundaryReduceDim(sycl::buffer<int> &reduced_block_sizes, |
| 163 | Array<Tk> keys, Array<To> vals_out, const int n, |
| 164 | const int numBlocks, const int threads_x, |
| 165 | const int dim, std::vector<int> dim_ordering) { |
| 166 | sycl::range<3> local(threads_x, 1, 1); |
| 167 | sycl::range<3> global( |
| 168 | local[0] * numBlocks, vals_out.dims()[dim_ordering[1]], |
| 169 | vals_out.dims()[dim_ordering[2]] * vals_out.dims()[dim_ordering[3]]); |
| 170 | |
| 171 | auto vals_out_get = vals_out.get(); |
| 172 | auto keys_get = keys.get(); |
| 173 | getQueue().submit([&](sycl::handler &h) { |
| 174 | write_accessor<int> reduced_block_sizes_acc{reduced_block_sizes, h}; |
| 175 | read_accessor<Tk> keys_acc{*keys_get, h}; |
| 176 | sycl::accessor<To> vals_out_acc{*vals_out_get, h}; |
| 177 | |
| 178 | // TODO: fold 3,4 dimensions |
| 179 | h.parallel_for( |
| 180 | sycl::nd_range<3>(global, local), |
| 181 | kernel::finalBoundaryReduceDimKernel<Tk, To, op>( |
| 182 | reduced_block_sizes_acc, keys_acc, keys, vals_out_acc, vals_out, |
| 183 | n, vals_out.dims()[dim_ordering[2]])); |
| 184 | }); |
| 185 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 186 | } |
| 187 | |
| 188 | template<typename Tk, typename To> |
| 189 | void compact(sycl::buffer<int> reduced_block_sizes, Array<Tk> &keys_out, |