Skip to content

SYCL: support 0-dim acc in handler::copy(accessor, accessor) #1551

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
Apr 22, 2020
Merged

SYCL: support 0-dim acc in handler::copy(accessor, accessor) #1551

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
134 changes: 108 additions & 26 deletions sycl/include/CL/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -9,6 +9,7 @@
#pragma once

#include <CL/sycl/access/access.hpp>
#include <CL/sycl/atomic.hpp>
#include <CL/sycl/context.hpp>
#include <CL/sycl/detail/cg.hpp>
#include <CL/sycl/detail/export.hpp>
Expand Down Expand Up @@ -366,6 +367,107 @@ class __SYCL_EXPORT handler {
return true;
}

/// Handles some special cases of the copy operation from one accessor
/// to another accessor. Returns true if the copy is handled here.
///
/// \param Src is a source SYCL accessor.
/// \param Dst is a destination SYCL accessor.
// TODO: support atomic accessor in Src or/and Dst.
template <typename TSrc, int DimSrc, access::mode ModeSrc,
access::target TargetSrc, typename TDst, int DimDst,
access::mode ModeDst, access::target TargetDst,
access::placeholder IsPHSrc, access::placeholder IsPHDst>
detail::enable_if_t<(DimSrc > 0) && (DimDst > 0), bool>
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Should we check here if DimSrc == DimDst?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I don't know if it is correct to copy from 1D accessor to 3D. It is probably correct. So, that check is not needed.
In my patch I do not change any constraints.
My patch only separates the existing code to this routine to make it possible to copy to/from 0-dim accessor.

copyAccToAccHelper(accessor<TSrc, DimSrc, ModeSrc, TargetSrc, IsPHSrc> Src,
accessor<TDst, DimDst, ModeDst, TargetDst, IsPHDst> Dst) {
if (!MIsHost &&
IsCopyingRectRegionAvailable(Src.get_range(), Dst.get_range()))
return false;

range<1> LinearizedRange(Src.get_count());
parallel_for<class __copyAcc2Acc<TSrc, DimSrc, ModeSrc, TargetSrc,
TDst, DimDst, ModeDst, TargetDst,
IsPHSrc, IsPHDst>>
(LinearizedRange, [=](id<1> Id) {
size_t Index = Id[0];
id<DimSrc> SrcIndex = getDelinearizedIndex(Src.get_range(), Index);
id<DimDst> DstIndex = getDelinearizedIndex(Dst.get_range(), Index);
Dst[DstIndex] = Src[SrcIndex];
});
return true;
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, T>
readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
atomic<T, access::address_space::global_space> AtomicSrc = Src;
return AtomicSrc.load();
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<(Dim > 0) && Mode == access::mode::atomic, T>
readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
id<Dim> Id = getDelinearizedIndex(Src.get_range(), 0);
return Src[Id].load();
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<Mode != access::mode::atomic, T>
readFromFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Src) const {
return *(Src.get_pointer());
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<Dim == 0 && Mode == access::mode::atomic, void>
writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
atomic<T, access::address_space::global_space> AtomicDst = Dst;
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Shouldn't it be a reference to atomic?

Copy link
Contributor Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

The accessor operator[] and operator() return a value, not a reference.

93 /* Available only when: accessMode == access::mode::atomic && dimensions ==
94 0 /
95 operator atomic<dataT, access::address_space::global_space> () const;
96
97 / Available only when: accessMode == access::mode::atomic && dimensions >
98 0 */
99 atomic<dataT, access::address_space::global_space> operator[](
100 id index) const;

Please clarify if I understood your comment incorrectly.
Did you think about adding & sign right before AtomicDst?
atomic<T, access::address_space::global_space> &AtomicDst = Dst;
If so it would be incorrect code, I think.

AtomicDst.store(V);
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<(Dim > 0) && Mode == access::mode::atomic, void>
writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
id<Dim> Id = getDelinearizedIndex(Dst.get_range(), 0);
Dst[Id].store(V);
}

template <typename T, int Dim, access::mode Mode, access::target Target,
access::placeholder IsPH>
detail::enable_if_t<Mode != access::mode::atomic, void>
writeToFirstAccElement(accessor<T, Dim, Mode, Target, IsPH> Dst, T V) const {
*(Dst.get_pointer()) = V;
}

/// Handles some special cases of the copy operation from one accessor
/// to another accessor. Returns true if the copy is handled here.
///
/// Source must have at least as many bytes as the range accessed by Dst.
///
/// \param Src is a source SYCL accessor.
/// \param Dst is a destination SYCL accessor.
template <typename TSrc, int DimSrc, access::mode ModeSrc,
access::target TargetSrc, typename TDst, int DimDst,
access::mode ModeDst, access::target TargetDst,
access::placeholder IsPHSrc, access::placeholder IsPHDst>
detail::enable_if_t<DimSrc == 0 || DimDst == 0, bool>
copyAccToAccHelper(accessor<TSrc, DimSrc, ModeSrc, TargetSrc, IsPHSrc> Src,
accessor<TDst, DimDst, ModeDst, TargetDst, IsPHDst> Dst) {
if (!MIsHost)
return false;

single_task<class __copyAcc2Acc<TSrc, DimSrc, ModeSrc, TargetSrc,
TDst, DimDst, ModeDst, TargetDst,
IsPHSrc, IsPHDst>> ([=]() {
writeToFirstAccElement(Dst, readFromFirstAccElement(Src));
});
return true;
}

constexpr static bool isConstOrGlobal(access::target AccessTarget) {
return AccessTarget == access::target::global_buffer ||
AccessTarget == access::target::constant_buffer;
Expand Down Expand Up @@ -985,6 +1087,7 @@ class __SYCL_EXPORT handler {
///
/// \param Src is a source SYCL accessor.
/// \param Dst is a pointer to destination memory.
// TODO: support 0-dimensional and atomic accessors.
template <typename T_Src, typename T_Dst, int Dims, access::mode AccessMode,
access::target AccessTarget,
access::placeholder IsPlaceholder = access::placeholder::false_t>
Expand Down Expand Up @@ -1030,6 +1133,7 @@ class __SYCL_EXPORT handler {
///
/// \param Src is a pointer to source memory.
/// \param Dst is a destination SYCL accessor.
// TODO: support 0-dimensional and atomic accessors.
template <typename T_Src, typename T_Dst, int Dims, access::mode AccessMode,
access::target AccessTarget,
access::placeholder IsPlaceholder = access::placeholder::false_t>
Expand Down Expand Up @@ -1072,7 +1176,7 @@ class __SYCL_EXPORT handler {
/// Copies the contents of memory object accessed by Src to the memory
/// object accessed by Dst.
///
/// Source must have at least as many bytes as the range accessed by Dst.
/// Dst must have at least as many bytes as the range accessed by Src.
///
/// \param Src is a source SYCL accessor.
/// \param Dst is a destination SYCL accessor.
Expand All @@ -1093,32 +1197,10 @@ class __SYCL_EXPORT handler {
"Invalid source accessor target for the copy method.");
static_assert(isValidTargetForExplicitOp(AccessTarget_Dst),
"Invalid destination accessor target for the copy method.");
// TODO replace to get_size() when it will provide correct values.
assert(
(Dst.get_range().size() * sizeof(T_Dst) >=
Src.get_range().size() * sizeof(T_Src)) &&
"dest must have at least as many bytes as the range accessed by src.");
if (MIsHost ||
!IsCopyingRectRegionAvailable(Src.get_range(), Dst.get_range())) {
range<Dims_Src> CopyRange = Src.get_range();
size_t Range = 1;
for (size_t I = 0; I < Dims_Src; ++I)
Range *= CopyRange[I];
range<1> LinearizedRange(Range);
parallel_for< class __copyAcc2Acc< T_Src, Dims_Src, AccessMode_Src,
AccessTarget_Src, T_Dst, Dims_Dst,
AccessMode_Dst, AccessTarget_Dst,
IsPlaceholder_Src,
IsPlaceholder_Dst>>
(LinearizedRange, [=](id<1> Id) {
size_t Index = Id[0];
id<Dims_Src> SrcIndex = getDelinearizedIndex(Src.get_range(), Index);
id<Dims_Dst> DstIndex = getDelinearizedIndex(Dst.get_range(), Index);
Dst[DstIndex] = Src[SrcIndex];
});

assert(Dst.get_size() >= Src.get_size() &&
"The destination accessor does not fit the copied memory.");
if (copyAccToAccHelper(Src, Dst))
return;
}
MCGType = detail::CG::COPY_ACC_TO_ACC;

detail::AccessorBaseHost *AccBaseSrc = (detail::AccessorBaseHost *)&Src;
Expand Down
124 changes: 124 additions & 0 deletions sycl/test/basic_tests/handler/handler_mem_op.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -46,6 +46,10 @@ template <typename T> void test_copy_acc_acc();
template <typename T> void test_update_host();
template <typename T> void test_2D_copy_acc_acc();
template <typename T> void test_3D_copy_acc_acc();
template <typename T>
void test_0D1D_copy_acc_acc();
template <typename T>
void test_0D1D_copy_acc_acc_atomic();
template <typename T> void test_1D2D_copy_acc_acc();
template <typename T> void test_1D3D_copy_acc_acc();
template <typename T> void test_2D1D_copy_acc_acc();
Expand Down Expand Up @@ -140,6 +144,19 @@ int main() {
test_3D_copy_acc_acc<point<float>>();
}

// handler.copy(acc, acc) 0D to/from 1D
{
test_0D1D_copy_acc_acc<int>();
test_0D1D_copy_acc_acc<point<int>>();
test_0D1D_copy_acc_acc<point<float>>();
}

// handler.copy(acc, acc) 0D to/from 1D where one/both acc are atomic
{
test_0D1D_copy_acc_acc_atomic<int>();
test_0D1D_copy_acc_acc_atomic<float>();
}

// handler.copy(acc, acc) 1D to 2D
{
test_1D2D_copy_acc_acc<int>();
Expand Down Expand Up @@ -433,6 +450,113 @@ template <typename T> void test_3D_copy_acc_acc() {
}
}

template <typename T>
void test_0D1D_copy_acc_acc() {
// Copy 1 element from 0-dim accessor to 1-dim accessor
T Src(1), Dst(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 0, access::mode::read, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 1, access::mode::write, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 1);

// Copy 1 element from 1-dim accessor to 0-dim accessor
Src = T(3);
Dst = T(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 1, access::mode::read, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 0, access::mode::write, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 3);

// Copy 1 element from 0-dim accessor to 0-dim accessor
Src = T(7);
Dst = T(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 0, access::mode::read, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 0, access::mode::write, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 7);
}

template <typename T>
void test_0D1D_copy_acc_acc_atomic() {
// Copy 1 element from 0-dim ATOMIC accessor to 1-dim accessor
T Src = T(1);
T Dst = T(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 0, access::mode::atomic, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 1, access::mode::write, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 1);

// Copy 1 element from 1-dim ATOMIC accessor to 0-dim accessor
Src = T(3);
Dst = T(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 1, access::mode::atomic, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 0, access::mode::write, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 3);

// Copy 1 element from 0-dim ATOMIC accessor to 0-dim ATOMIC accessor
Src = T(7);
Dst = T(0);
{
buffer<T, 1> BufferFrom(&Src, range<1>(1));
buffer<T, 1> BufferTo(&Dst, range<1>(1));
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 0, access::mode::atomic, access::target::global_buffer>
AccessorFrom(BufferFrom, Cgh);
accessor<T, 0, access::mode::atomic, access::target::global_buffer>
AccessorTo(BufferTo, Cgh);
Cgh.copy(AccessorFrom, AccessorTo);
});
}
assert(Dst == 7);
}

template <typename T> void test_1D2D_copy_acc_acc() {
const size_t Size = 20;
std::vector<T> Data(Size);
Expand Down