| 138 | |
| 139 | template<typename Tk, typename To, af_op_t op> |
| 140 | void finalBoundaryReduce(sycl::buffer<int> &reduced_block_sizes, Array<Tk> keys, |
| 141 | Array<To> vals_out, const int n, const int numBlocks, |
| 142 | const int threads_x) { |
| 143 | sycl::range<1> local(threads_x); |
| 144 | sycl::range<1> global(local[0] * numBlocks); |
| 145 | |
| 146 | auto vals_out_get = vals_out.get(); |
| 147 | auto keys_get = keys.get(); |
| 148 | getQueue().submit([&](sycl::handler &h) { |
| 149 | write_accessor<int> reduced_block_sizes_acc{reduced_block_sizes, h}; |
| 150 | read_accessor<Tk> keys_acc{*keys_get, h}; |
| 151 | sycl::accessor<To> vals_out_acc{*vals_out_get, h}; |
| 152 | |
| 153 | h.parallel_for(sycl::nd_range<1>(global, local), |
| 154 | kernel::finalBoundaryReduceKernel<Tk, To, op>( |
| 155 | reduced_block_sizes_acc, keys_acc, keys, |
| 156 | vals_out_acc, vals_out, n)); |
| 157 | }); |
| 158 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 159 | } |
| 160 | |
| 161 | template<typename Tk, typename To, af_op_t op> |
| 162 | void finalBoundaryReduceDim(sycl::buffer<int> &reduced_block_sizes, |