Implementation issue of Reduction upon --acpp-targets='generic'

[Jiuxiaoyunhai]

#define sycl_reduction_plus(argus) sycl::reduction(&(argus), sycl::plus<real_t>())
#define sycl_reduction_max(argus) sycl::reduction(&(argus), sycl::maximum<real_t>())
#define sycl_reduction_min(argus) sycl::reduction(&(argus), sycl::minimum<real_t>())

the reduction expression located at: /XFLUIDS/src/solver_GetDt/GlobalDt_block.hpp:33
q.submit([&](sycl::handler &h) {
auto reduction_max_x = sycl_reduction_max(uvw_c_max[0]);
h.parallel_for(sycl::nd_range<1>(global_ndrange, local_ndrange), reduction_max_x, [=](nd_item<1> index, auto &temp_max_x) {
auto id = index.get_global_id();
temp_max_x.combine(sycl::fabs(u[id]) + c[id]);
});
});

get an compile time error:

In file included from /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/sycl.hpp:31:
In file included from /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/sycl.hpp:93:
In file included from /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/queue.hpp:51:
In file included from /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/handler.hpp:67:
In file included from /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/kernel_launcher_factory.hpp:57:
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:207:5: error: no matching function for call to object of type 'const (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86)'
_k(this_item);
^~
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:63:5: note: in instantiation of member function 'hipsycl::glue::__sscp_dispatch::ndrange_parallel_for<(lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>::operator()' requested here
k();
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:430:7: note: in instantiation of function template specialization '__hipsycl_sscp_kernel<hipsycl::glue::__sscp_dispatch::ndrange_parallel_for<(lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>>' requested here
__hipsycl_sscp_kernel(k);
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:413:31: note: in instantiation of function template specialization 'hipsycl::glue::sscp_kernel_launcher::generate_kernel<hipsycl::glue::__sscp_dispatch::ndrange_parallel_for<(lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>>' requested here
std::string kernel_name = generate_kernel(k);
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:379:5: note: in instantiation of function template specialization 'hipsycl::glue::sscp_kernel_launcher::launch_kernel<hipsycl::glue::__sscp_dispatch::ndrange_parallel_for<(lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>>' requested here
launch_kernel(k, op, num_groups, selected_group_size, local_mem_size);
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/llvm-sscp/sscp_kernel_launcher.hpp:295:11: note: in instantiation of function template specialization 'hipsycl::glue::sscp_kernel_launcher::launch_kernel_with_global_range<hipsycl::glue::__sscp_dispatch::ndrange_parallel_for<(lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>, 1>' requested here
launch_kernel_with_global_range(
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/hipSYCL/glue/kernel_launcher_factory.hpp:109:15: note: in instantiation of function template specialization 'hipsycl::glue::sscp_kernel_launcher::bind<hipsycl::glue::kernel_name_traits<__hipsycl_unnamed_kernel, (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86)>, hipsycl::rt::kernel_type::ndrange_parallel_for, 1, (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), hipsycl::sycl::detail::pointer_reduction_descriptor<double, hipsycl::sycl::maximum>>' requested here
launcher->bind<name_traits, Type>(offset, global_range, local_range,
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/handler.hpp:804:15: note: in instantiation of function template specialization 'hipsycl::glue::make_kernel_launchers<__hipsycl_unnamed_kernel, hipsycl::rt::kernel_type::ndrange_parallel_for, 1, (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), hipsycl::sycl::detail::pointer_reduction_descriptor<double, hipsycl::sycl::maximum>>' requested here
glue::make_kernel_launchers<KernelName, KernelType>(
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/handler.hpp:374:13: note: in instantiation of function template specialization 'hipsycl::sycl::handler::submit_kernel<__hipsycl_unnamed_kernel, hipsycl::rt::kernel_type::ndrange_parallel_for, (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1, hipsycl::sycl::detail::pointer_reduction_descriptor<double, hipsycl::sycl::maximum>>' requested here
this->submit_kernel<KernelName, rt::kernel_type::ndrange_parallel_for>(
^
/home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/handler.hpp:380:13: note: in instantiation of function template specialization 'hipsycl::sycl::detail::separate_last_argument_and_apply<(lambda at /home/ljl/Downloads/XFLUIDS-Test/external/install/AdaptiveCpp/bin/../include/AdaptiveCpp/sycl/../hipSYCL/sycl/handler.hpp:373:20) &, const hipsycl::sycl::detail::pointer_reduction_descriptor<double, hipsycl::sycl::maximum> &, const (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86) &>' requested here
detail::separate_last_argument_and_apply(invoker, redu_kernel...);
^
/home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:6: note: in instantiation of function template specialization 'hipsycl::sycl::handler::parallel_for<__hipsycl_unnamed_kernel, hipsycl::sycl::detail::pointer_reduction_descriptor<double, hipsycl::sycl::maximum>, (lambda at /home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86), 1>' requested here
h.parallel_for(sycl::nd_range<1>(global_ndrange, local_ndrange), reduction_max_x, [=](nd_item<1> index, auto &temp_max_x) { //
^
/home/ljl/Downloads/XFLUIDS-Test/src/solver_GetDt/GlobalDt_block.hpp:33:86: note: candidate function template not viable: requires 2 arguments, but 1 was provided
h.parallel_for(sycl::nd_range<1>(global_ndrange, local_ndrange), reduction_max_x, [=](nd_item<1> index, auto &temp_max_x) { //
^

[illuhad]

sycl::reduction is not yet implemented for generic SSCP compiler. It is also only partially and incompletely implemented in the other compilation flows.
You can use algorithms::transform_reduce, which is supported on all compilation flows:

AdaptiveCpp/include/hipSYCL/algorithms/numeric.hpp

Line 220 in 37d1dcd

transform_reduce(sycl::queue &q, util::allocation_group &scratch_allocations,

[svtcli]

Would you know if reductions will be supported at some point?
It's the only error preventing us from compiling DPEcho...
Thanks for any info!

[illuhad]

@svtcli You are in luck, please check out this PR: #1453. We need to do some more testing on this one, but it might already work for you.
Unfortunately, the reduction API in SYCL 2020 is so needlessly complex that supporting this feature with all corner cases is an almost infeasible endeavour.