| 246 | |
| 247 | template<typename Tk> |
| 248 | void testNeedsReduction(sycl::buffer<int> needs_reduction, |
| 249 | sycl::buffer<int> needs_boundary, const Array<Tk> &keys, |
| 250 | const int n, const int numBlocks, const int threads_x) { |
| 251 | sycl::range<1> local(threads_x); |
| 252 | sycl::range<1> global(local[0] * numBlocks); |
| 253 | |
| 254 | auto keys_get = keys.get(); |
| 255 | getQueue().submit([&](sycl::handler &h) { |
| 256 | sycl::accessor<int> needs_reduction_acc{needs_reduction, h}; |
| 257 | sycl::accessor<int> needs_boundary_acc{needs_boundary, h}; |
| 258 | read_accessor<Tk> keys_acc{*keys_get, h}; |
| 259 | auto l_keys = sycl::local_accessor<Tk>(threads_x, h); |
| 260 | |
| 261 | h.parallel_for(sycl::nd_range<1>(global, local), |
| 262 | kernel::testNeedsReductionKernel<Tk>( |
| 263 | needs_reduction_acc, needs_boundary_acc, keys_acc, |
| 264 | keys, n, threads_x, l_keys)); |
| 265 | }); |
| 266 | ONEAPI_DEBUG_FINISH(getQueue()); |
| 267 | } |
| 268 | |
| 269 | template<af_op_t op, typename Ti, typename Tk, typename To> |
| 270 | int reduce_by_key_first(Array<Tk> &keys_out, Array<To> &vals_out, |