Skip to content

[SYCL] Fix explicit copy operation for host device #2627

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 2 commits into from
Oct 14, 2020
Merged

[SYCL] Fix explicit copy operation for host device #2627

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
1707b62
[SYCL] Fix explicit copy operation for host device
sergey-semenov Oct 12, 2020
eaaf1d2
Add const to LinearIndex
sergey-semenov Oct 12, 2020
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
8 changes: 2 additions & 6 deletions sycl/include/CL/sycl/handler.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -641,9 +641,7 @@ class __SYCL_EXPORT handler {
range<Dim> Range = Src.get_range();
parallel_for<class __copyAcc2Ptr<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
(Range, [=](id<Dim> Index) {
size_t LinearIndex = Index[0];
for (int I = 1; I < Dim; ++I)
LinearIndex += Range[I] * Index[I];
const size_t LinearIndex = detail::getLinearIndex(Index, Range);
using TSrcNonConst = typename std::remove_const<TSrc>::type;
(reinterpret_cast<TSrcNonConst *>(Dst))[LinearIndex] = Src[Index];
});
Expand Down Expand Up @@ -678,9 +676,7 @@ class __SYCL_EXPORT handler {
range<Dim> Range = Dst.get_range();
parallel_for<class __copyPtr2Acc<TSrc, TDst, Dim, AccMode, AccTarget, IsPH>>
(Range, [=](id<Dim> Index) {
size_t LinearIndex = Index[0];
for (int I = 1; I < Dim; ++I)
LinearIndex += Range[I] * Index[I];
const size_t LinearIndex = detail::getLinearIndex(Index, Range);
Dst[Index] = (reinterpret_cast<const TDst *>(Src))[LinearIndex];
});
}
Expand Down
62 changes: 62 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 @@ -34,7 +34,9 @@ template <typename T> struct point {

template <typename T> void test_fill(T Val);
template <typename T> void test_copy_ptr_acc();
template <typename T> void test_3D_copy_ptr_acc();
template <typename T> void test_copy_acc_ptr();
template <typename T> void test_3D_copy_acc_ptr();
template <typename T> void test_copy_shared_ptr_acc();
template <typename T> void test_copy_shared_ptr_const_acc();
template <typename T> void test_copy_acc_shared_ptr();
Expand Down Expand Up @@ -72,6 +74,14 @@ int main() {
test_copy_ptr_acc<point<int>>();
test_copy_ptr_acc<point<float>>();
}
// handler.copy(ptr, acc) 3D
{
test_3D_copy_ptr_acc<int>();
test_3D_copy_ptr_acc<int>();
test_3D_copy_ptr_acc<point<int>>();
test_3D_copy_ptr_acc<point<int>>();
test_3D_copy_ptr_acc<point<float>>();
}
// handler.copy(acc, ptr)
{
test_copy_acc_ptr<int>();
Expand All @@ -80,6 +90,14 @@ int main() {
test_copy_acc_ptr<point<int>>();
test_copy_acc_ptr<point<float>>();
}
// handler.copy(acc, ptr) 3D
{
test_3D_copy_acc_ptr<int>();
test_3D_copy_acc_ptr<int>();
test_3D_copy_acc_ptr<point<int>>();
test_3D_copy_acc_ptr<point<int>>();
test_3D_copy_acc_ptr<point<float>>();
}
// handler.copy(shared_ptr, acc)
{
test_copy_shared_ptr_acc<int>();
Expand Down Expand Up @@ -277,6 +295,28 @@ template <typename T> void test_copy_ptr_acc() {
assert(DstValue == 99);
}

template <typename T> void test_3D_copy_ptr_acc() {
const range<3> Range{2, 3, 4};
const size_t Size = 2 * 3 * 4;
T Data[Size] = {0};
T Values[Size] = {0};
for (size_t I = 0; I < Size; ++I)
Values[I] = I;

{
buffer<T, 3> Buffer(Data, Range);
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 3, access::mode::write, access::target::global_buffer>
Accessor(Buffer, Cgh, Range);
Cgh.copy(Values, Accessor);
});
}

for (int I = 0; I < Size; ++I)
assert(Data[I] == Values[I]);
}

template <typename T> void test_copy_acc_ptr() {
const size_t Size = 10;
T Data[Size] = {0};
Expand Down Expand Up @@ -345,6 +385,28 @@ template <typename T> void test_copy_acc_ptr() {
assert(DstValue == 77);
}

template <typename T> void test_3D_copy_acc_ptr() {
const range<3> Range{2, 3, 4};
const size_t Size = 2 * 3 * 4;
T Data[Size] = {0};
T Values[Size] = {0};
for (size_t I = 0; I < Size; ++I)
Data[I] = I;

{
buffer<T, 3> Buffer(Data, Range);
queue Queue;
Queue.submit([&](handler &Cgh) {
accessor<T, 3, access::mode::read, access::target::global_buffer>
Accessor(Buffer, Cgh, Range);
Cgh.copy(Accessor, Values);
});
}

for (size_t I = 0; I < Size; ++I)
assert(Data[I] == Values[I]);
}

template <typename T> void test_copy_shared_ptr_acc() {
const size_t Size = 10;
T Data[Size] = {0};
Expand Down