[SYCL] Implement user-defined reduction extension #7587
Conversation
|
PR #7436 closed due to unsuccessful merge conflict resolving. This is a copy of that PR. |
|
/verify with intel/llvm-test-suite#1395 |
dm-vodopyanov
requested review from
aelovikov-intel,
a team and
AlexeySachkov
| if (g.leader()) { | ||
| for (int i = 1; i < g.get_local_linear_range(); i++) { | ||
| result = binary_op(result, Memory[i]); | ||
| } | ||
| } |
There was a problem hiding this comment.
I doubt people would be happy with this naive implementation. Might be fine for the initial commit though.
There was a problem hiding this comment.
This is intentional. Making this through SPIR-V extensions is too long (spec, impl for them). The plan is to re-write the current using these extensions without rush if customers would not be happy with the performance.
sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp
Outdated
Show resolved
Hide resolved
sycl/include/sycl/ext/oneapi/experimental/user_defined_reductions.hpp
Outdated
Show resolved
Hide resolved
| auto g = group_helper.get_group(); | ||
| T partial = *(first + g.get_local_linear_id()); | ||
| Ptr second = first + g.get_local_linear_range(); | ||
| sycl::detail::for_each(g, second, last, |
There was a problem hiding this comment.
Would it work with non-pointer iterators?
There was a problem hiding this comment.
Probably not, SYCL2020 sycl::joint_reduce also uses this approach. User passes sycl::accessor::get_pointer() to this func.
There was a problem hiding this comment.
Can we claim the extension as supported with that implementation though?
There was a problem hiding this comment.
At the moment this won't work with non-pointer iterators simply because we have enable_if at the beginning of the function, which requires Ptr to be a pointer type.
I would suggest that we go with the current approach as a first iteration of the extension: we will anyway only claim it as experimental for now.
@dm-vodopyanov, we should raise this question with spec writers, because both extension and core spec do not say whether Ptr can be a generic iterator or is it a mere pointer.
|
/verify with intel/llvm-test-suite#1395 |
Spec: intel/llvm#7202 Implementation: intel/llvm#7587
Spec: #7202
Tests: intel/llvm-test-suite#1395