Skip to content
New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Jump to bottom

[SYCL] Add reduction overloads accepting span #6019

Merged
merged 22 commits into from
May 2, 2022
Merged

[SYCL] Add reduction overloads accepting span #6019

Changes from 1 commit
Commits
Show all changes
22 commits
Select commit Hold shift + click to select a range
c0de8c6
[SYCL] Add reduction overloads accepting span
Pennycook Apr 18, 2022
1642d2f
[SYCL] Add specialization of reducer for spans
Pennycook Apr 18, 2022
86e9ef5
[SYCL] Add specialization of reduction_impl
Pennycook Apr 18, 2022
c0e0cec
[SYCL] Enable span reductions
Pennycook Apr 18, 2022
1c97a16
[SYCL] Disable span reductions before C++17
Pennycook Apr 18, 2022
807440c
Remove unused argument
Pennycook Apr 21, 2022
111ea2f
[SYCL][NFC] Refactor to avoid code duplication
Pennycook Apr 20, 2022
f672248
[SYCL] Fix bug introduced during refactor
Pennycook Apr 26, 2022
5207769
[SYCL] Add span support to packs of reductions
Pennycook Apr 26, 2022
cab1b2f
[SYCL] Generalize esimd_verify test for reducers
Pennycook Apr 27, 2022
a9d18cc
Remove incorrect comment
Pennycook Apr 29, 2022
0236111
Remove unused argument
Pennycook Apr 29, 2022
7608d97
Remove unused arguments
Pennycook Apr 29, 2022
ca998d7
Clarify comment describing View template parameter
Pennycook Apr 29, 2022
af22456
Fix comment
Pennycook Apr 29, 2022
657ff53
Add comment to first reducer specialization
Pennycook Apr 29, 2022
9553b7f
Remove redundant SFINAE from getIdentity()
Pennycook Apr 29, 2022
12056fc
clang-format
Pennycook Apr 29, 2022
cb030f2
[SYCL] Revert to sycl::detail::tuple
Pennycook Apr 29, 2022
c2663db
Remove old TODO
Pennycook Apr 29, 2022
90deb60
Remove commented out code
Pennycook Apr 29, 2022
6843a67
Fix IsArrayReduction
Pennycook Apr 29, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
[SYCL] Revert to sycl::detail::tuple 
std::tuple is still used for reducers (as in the original code).

Signed-off-by: John Pennycook <john.pennycook@intel.com>
  • Loading branch information
@Pennycook
Pennycook committed Apr 29, 2022
commit cb030f279bbbc962f2d1231c7701fdc448898929
14 changes: 8 additions & 6 deletions sycl/include/sycl/ext/oneapi/reduction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -94,11 +94,13 @@ using IsReduOptForFastReduce =
sycl::detail::IsMaximum<T, BinaryOperation>::value)>;
#endif

// Using std::tuple allows usage of functions like std::tuple_element_t
// Switching tuple can be done by changing the definitions below
template <typename... Ts> using ReduTupleT = std::tuple<Ts...>;
// std::tuple seems to be a) too heavy and b) not copyable to device now
// Thus sycl::detail::tuple is used instead.
// Switching from sycl::device::tuple to std::tuple can be done by re-defining
// the ReduTupleT type and makeReduTupleT() function below.
template <typename... Ts> using ReduTupleT = sycl::detail::tuple<Ts...>;
template <typename... Ts> ReduTupleT<Ts...> makeReduTupleT(Ts... Elements) {
return std::make_tuple(Elements...);
return sycl::detail::make_tuple(Elements...);
}

__SYCL_EXPORT size_t reduGetMaxWGSize(std::shared_ptr<queue_impl> Queue,
Expand Down Expand Up @@ -2010,7 +2012,7 @@ template <bool Pow2WG, bool IsOneWG, typename... Reductions, int Dims,
typename... Ts, typename... BOPsT, size_t... Is>
void reduCGFuncImplScalar(
nd_item<Dims> NDIt, ReduTupleT<LocalAccT...> LocalAccsTuple,
ReduTupleT<OutAccT...> OutAccsTuple, ReduTupleT<ReducerT...> &ReducersTuple,
ReduTupleT<OutAccT...> OutAccsTuple, std::tuple<ReducerT...> &ReducersTuple,
ReduTupleT<Ts...> IdentitiesTuple, ReduTupleT<BOPsT...> BOPsTuple,
std::array<bool, sizeof...(Reductions)> InitToIdentityProps,
std::index_sequence<Is...> ReduIndices) {
Expand Down Expand Up @@ -2112,7 +2114,7 @@ template <bool Pow2WG, bool IsOneWG, typename... Reductions, int Dims,
typename... Ts, typename... BOPsT, size_t... Is>
void reduCGFuncImplArray(
nd_item<Dims> NDIt, ReduTupleT<LocalAccT...> LocalAccsTuple,
ReduTupleT<OutAccT...> OutAccsTuple, ReduTupleT<ReducerT...> &ReducersTuple,
ReduTupleT<OutAccT...> OutAccsTuple, std::tuple<ReducerT...> &ReducersTuple,
ReduTupleT<Ts...> IdentitiesTuple, ReduTupleT<BOPsT...> BOPsTuple,
std::array<bool, sizeof...(Reductions)> InitToIdentityProps,
std::index_sequence<Is...> ReduIndices) {
Expand Down