SYCL custom reductions required for ArborX #3544
Conversation
masterleinad
force-pushed
the
sycl_custom_reductions_clean
branch
2 times, most recently
from
2ed9f3f
to
3a6f355
Compare
|
It seems that the detection mechanism for |
| if constexpr (ReduceFunctorHasJoin<Functor>::value) { | ||
| return cl::sycl::ONEAPI::reduction( | ||
| result_ptr, identity, | ||
| [=](value_type& old_value, const value_type& new_value) { |
There was a problem hiding this comment.
Please comment why you are capturing by value here.
There was a problem hiding this comment.
/var/jenkins/workspace/Kokkos/core/src/SYCL/Kokkos_SYCL_Parallel_Reduce.hpp:170:19: error: 'const Kokkos::Impl::ParallelReduce<Test::TestViewOperator_LeftAndRight<int [2][3][4][2][3][4], Kokkos::Experimental::SYCL, 6>, Kokkos::RangePolicy<Kokkos::Experimental::SYCL>, Kokkos::InvalidType, Kokkos::Experimental::SYCL>::ExtendedReferenceWrapper<Test::TestViewOperator_LeftAndRight<int [2][3][4][2][3][4], Kokkos::Experimental::SYCL, 6>> &' cannot be used as the type of a kernel parameter
functor.join(old_value, new_value);
^
9f67b70
to
e21ba68
Compare
| @@ -83,7 +83,7 @@ SYCL::SYCL() : m_space_instance(&Impl::SYCLInternal::singleton()) { | |||
|
|
|||
| int SYCL::concurrency() { | |||
| // FIXME_SYCL | |||
| return 1; | |||
| return 2; | |||
There was a problem hiding this comment.
This is still a fixme but one of the test was failing because asserting that the concurrency is greater than one
There was a problem hiding this comment.
I'd like to see a more detailed comment than "FIXME" here, something like "this is to work around some tests that assert concurrency must be greater than 1"
|
I find this change in AnalyzeFunctor highly suspect. On the other hand that whole thing is highly suspect and Daisy wants to rewrite it from scratch so ... |
|
|
||
| value_type host_result = identity; | ||
| if constexpr (ReduceFunctorHasInit<Functor>::value) | ||
| ValueInit::init(functor, &host_result); | ||
| q.memcpy(result_ptr, &host_result, sizeof(host_result)); |
There was a problem hiding this comment.
I think we should just use assignment here *result_ptr = host_result;, mainly because the rest of the code also takes advantage of it being in USM shared memory (*m_result_ptr = *result_ptr;, for instance).
There was a problem hiding this comment.
Improved in e297544 (#3544)
|
Plz cleanup history and I will merge |
masterleinad
force-pushed
the
sycl_custom_reductions_clean
branch
from
32b9967
to
0aa76bc
Compare
|
@dalg24 Here you go! |
Based on #3490 and slightly on #3542.