Skip to content

[SYCL][NFC] Use regular accessors for USM reductions, fix incorrect c… #1962

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

Merged
merged 1 commit into from
Jun 24, 2020
Merged

[SYCL][NFC] Use regular accessors for USM reductions, fix incorrect c… #1962

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
2 changes: 1 addition & 1 deletion sycl/include/CL/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -983,7 +983,7 @@ class __SYCL_EXPORT handler {
RWAcc);
this->finalize();

// Copy from RWAcc to some temp memory.
// Copy from RWAcc to user's reduction accessor.
handler CopyHandler(QueueCopy, MIsHost);
CopyHandler.saveCodeLoc(MCodeLoc);
#ifndef __SYCL_DEVICE_ONLY__
Expand Down
39 changes: 14 additions & 25 deletions sycl/include/CL/sycl/intel/reduction.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1001,8 +1001,8 @@ reduAuxCGFunc(handler &CGH, const nd_range<Dims> &Range, size_t NWorkItems,

/// Creates and returns an object implementing the reduction functionality.
/// Accepts 3 arguments: the accessor to buffer to where the computed reduction
/// must be stored \param Acc, identity value \param Identity, and the
/// binary operation that must be used in the reduction \param Combiner.
/// must be stored \param Acc, identity value \param Identity, and the binary
/// operation used in the reduction.
template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
access::placeholder IsPH>
detail::reduction_impl<T, BinaryOperation, Dims, false, AccMode, IsPH>
Expand All @@ -1015,8 +1015,7 @@ reduction(accessor<T, Dims, AccMode, access::target::global_buffer, IsPH> &Acc,

/// Creates and returns an object implementing the reduction functionality.
/// Accepts 2 arguments: the accessor to buffer to where the computed reduction
/// must be stored \param Acc and the binary operation that must be used
/// in the reduction \param Combiner.
/// must be stored \param Acc and the binary operation used in the reduction.
/// The identity value is not passed to this version as it is statically known.
template <typename T, class BinaryOperation, int Dims, access::mode AccMode,
access::placeholder IsPH>
Expand All @@ -1032,37 +1031,27 @@ reduction(accessor<T, Dims, AccMode, access::target::global_buffer, IsPH> &Acc,

/// Creates and returns an object implementing the reduction functionality.
/// Accepts 3 arguments: the reference to the reduction variable to where
/// the computed reduction must be stored \param VarRef, identity value
/// \param Identity, and the binary operation that must be used
/// in the reduction \param Combiner.
/// the computed reduction must be stored \param VarPtr, identity value
/// \param Identity, and the binary operation used in the reduction.
template <typename T, class BinaryOperation>
detail::reduction_impl<T, BinaryOperation, 0, true, access::mode::read_write,
access::placeholder::true_t>
reduction(T *VarPtr, const T &Identity, BinaryOperation Combiner) {
// The Combiner argument was needed only to define the BinaryOperation param.
(void)Combiner;
detail::reduction_impl<T, BinaryOperation, 0, true, access::mode::read_write>
reduction(T *VarPtr, const T &Identity, BinaryOperation) {
return detail::reduction_impl<T, BinaryOperation, 0, true,
access::mode::read_write,
access::placeholder::true_t>(VarPtr, Identity);
access::mode::read_write>(VarPtr, Identity);
}

/// Creates and returns an object implementing the reduction functionality.
/// Accepts 3 arguments: the reference to the reduction variable to where
/// the computed reduction must be stored \param VarRef, identity value
/// \param Identity, and the binary operation that must be used
/// in the reduction \param Combiner.
/// Accepts 2 arguments: the reference to the reduction variable, to where
/// the computed reduction must be stored \param VarPtr, and the binary
/// operation used in the reduction.
/// The identity value is not passed to this version as it is statically known.
template <typename T, class BinaryOperation>
detail::enable_if_t<detail::IsKnownIdentityOp<T, BinaryOperation>::value,
detail::reduction_impl<T, BinaryOperation, 0, true,
access::mode::read_write,
access::placeholder::true_t>>
reduction(T *VarPtr, BinaryOperation Combiner) {
// The Combiner argument was needed only to define the BinaryOperation param.
(void)Combiner;
access::mode::read_write>>
reduction(T *VarPtr, BinaryOperation) {
return detail::reduction_impl<T, BinaryOperation, 0, true,
access::mode::read_write,
access::placeholder::true_t>(VarPtr);
access::mode::read_write>(VarPtr);
}

} // namespace intel
Expand Down