| 67 | template <class T, class Kernel, |
| 68 | class BinaryReductionOp> |
| 69 | sycl::event wg_model_reduction(sycl::queue &q, |
| 70 | util::allocation_group &scratch_allocations, |
| 71 | T *output, T init, std::size_t target_num_groups, |
| 72 | std::size_t local_size, std::size_t problem_size, |
| 73 | Kernel k, BinaryReductionOp op, |
| 74 | const std::vector<sycl::event>& deps = {}) { |
| 75 | assert(target_num_groups > 0); |
| 76 | |
| 77 | sycl::event last_event; |
| 78 | auto ndrange_launcher = |
| 79 | [&](std::size_t num_groups, std::size_t wg_size, std::size_t global_size, |
| 80 | std::size_t local_mem, auto kernel) { |
| 81 | last_event = q.submit([&](sycl::handler &cgh) { |
| 82 | // This is just there to register the appropriate amount of local |
| 83 | // memory; the reduction engine will access it directly without going |
| 84 | // through the accessor. |
| 85 | cgh.depends_on(last_event); |
| 86 | sycl::local_accessor<char> acc{sycl::range<1>{local_mem}, cgh}; |
| 87 | cgh.parallel_for(sycl::nd_range<1>{wg_size * num_groups, wg_size}, |
| 88 | kernel); |
| 89 | }); |
| 90 | }; |
| 91 | |
| 92 | auto operator_config = get_reduction_operator_configuration<T>(op); |
| 93 | auto reduction_descriptor = reduction::reduction_descriptor{ |
| 94 | operator_config, init, output}; |
| 95 | |
| 96 | using group_reduction_type = |
| 97 | reduction::wg_model::group_reductions::generic_local_memory< |
| 98 | std::decay_t<decltype(reduction_descriptor)>>; |
| 99 | |
| 100 | // The reduction engine will update this value with the |
| 101 | // appropriate amount of local memory for the main kernel. |
| 102 | std::size_t main_kernel_local_mem = 0; |
| 103 | reduction::wg_model::group_horizontal_reducer<group_reduction_type> |
| 104 | horizontal_reducer{ |
| 105 | group_reduction_type{main_kernel_local_mem, local_size}}; |
| 106 | reduction::wg_hierarchical_reduction_engine engine{horizontal_reducer, |
| 107 | &scratch_allocations}; |
| 108 | |
| 109 | util::data_streamer streamer{q.get_device(), problem_size, local_size}; |
| 110 | |
| 111 | const std::size_t dispatched_global_size = |
| 112 | streamer.get_required_global_size(); |
| 113 | auto plan = engine.create_plan(dispatched_global_size, local_size, |
| 114 | reduction_descriptor); |
| 115 | |
| 116 | auto main_kernel = engine.make_main_reducing_kernel( |
| 117 | [=](sycl::nd_item<1> idx, auto &reducer) { |
| 118 | |
| 119 | util::data_streamer::run(problem_size, idx, [&](sycl::id<1> i){ |
| 120 | k(i, reducer); |
| 121 | }); |
| 122 | }, |
| 123 | plan); |
| 124 | |
| 125 | last_event = q.submit([&](sycl::handler &cgh) { |
| 126 | sycl::local_accessor<char> acc{sycl::range<1>{main_kernel_local_mem}, cgh}; |
no test coverage detected