Skip to content

[SYCL] Fix undefined symbols in async_work_group_copy #1243

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
Mar 17, 2020
Merged

[SYCL] Fix undefined symbols in async_work_group_copy #1243

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
b59e6c2
[SYCL] Fix undefined symbols in async_work_group_copy
dnmokhov Feb 13, 2020
e6e02a3
Added tests of async_work_group_copy without stride.
dnmokhov Mar 6, 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
45 changes: 29 additions & 16 deletions sycl/include/CL/sycl/group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -12,6 +12,7 @@
#include <CL/__spirv/spirv_types.hpp>
#include <CL/__spirv/spirv_vars.hpp>
#include <CL/sycl/detail/common.hpp>
#include <CL/sycl/detail/generic_type_traits.hpp>
#include <CL/sycl/detail/helpers.hpp>
#include <CL/sycl/device_event.hpp>
#include <CL/sycl/h_item.hpp>
Expand Down Expand Up @@ -264,21 +265,27 @@ template <int dimensions = 1> class group {
device_event async_work_group_copy(local_ptr<dataT> dest,
global_ptr<dataT> src,
size_t numElements) const {
__ocl_event_t e =
OpGroupAsyncCopyGlobalToLocal<dataT>(
__spv::Scope::Workgroup,
dest.get(), src.get(), numElements, 1, 0);
using T = detail::ConvertToOpenCLType_t<dataT>;
using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;

__ocl_event_t e = OpGroupAsyncCopyGlobalToLocal<T>(
__spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
numElements, 1, 0);
return device_event(&e);
}

template <typename dataT>
device_event async_work_group_copy(global_ptr<dataT> dest,
local_ptr<dataT> src,
size_t numElements) const {
__ocl_event_t e =
OpGroupAsyncCopyLocalToGlobal<dataT>(
__spv::Scope::Workgroup,
dest.get(), src.get(), numElements, 1, 0);
using T = detail::ConvertToOpenCLType_t<dataT>;
using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;

__ocl_event_t e = OpGroupAsyncCopyLocalToGlobal<T>(
__spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
numElements, 1, 0);
return device_event(&e);
}

Expand All @@ -287,10 +294,13 @@ template <int dimensions = 1> class group {
global_ptr<dataT> src,
size_t numElements,
size_t srcStride) const {
__ocl_event_t e =
OpGroupAsyncCopyGlobalToLocal<dataT>(
__spv::Scope::Workgroup,
dest.get(), src.get(), numElements, srcStride, 0);
using T = detail::ConvertToOpenCLType_t<dataT>;
using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;

__ocl_event_t e = OpGroupAsyncCopyGlobalToLocal<T>(
__spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
numElements, srcStride, 0);
return device_event(&e);
}

Expand All @@ -299,10 +309,13 @@ template <int dimensions = 1> class group {
local_ptr<dataT> src,
size_t numElements,
size_t destStride) const {
__ocl_event_t e =
OpGroupAsyncCopyLocalToGlobal<dataT>(
__spv::Scope::Workgroup,
dest.get(), src.get(), numElements, destStride, 0);
using T = detail::ConvertToOpenCLType_t<dataT>;
using DestT = detail::ConvertToOpenCLType_t<decltype(dest)>;
using SrcT = detail::ConvertToOpenCLType_t<decltype(src)>;

__ocl_event_t e = OpGroupAsyncCopyLocalToGlobal<T>(
__spv::Scope::Workgroup, DestT(dest.get()), SrcT(src.get()),
numElements, destStride, 0);
return device_event(&e);
}

Expand Down
86 changes: 86 additions & 0 deletions sycl/test/regression/group.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -162,10 +162,96 @@ bool group__get_linear_id() {
return Pass;
}

// Tests group::async_work_group_copy()
bool group__async_work_group_copy() {
std::cout << "+++ Running group::async_work_group_copy() test...\n";
constexpr int DIMS = 2;
bool Pass = true;

std::vector<std::pair<range<DIMS>, range<DIMS>>> ranges;
ranges.push_back({{3, 1}, {2, 3}});
ranges.push_back({{1, 3}, {3, 2}});

for (const auto &i : ranges) {
const auto LocalRange = i.first;
const auto GroupRange = i.second;
const range<DIMS> GlobalRange = LocalRange * GroupRange;
using DataType = vec<size_t, DIMS>;
const int DataLen = GlobalRange.size();
std::unique_ptr<DataType[]> Data(new DataType[DataLen]);
std::memset(Data.get(), 0, DataLen * sizeof(DataType));

try {
buffer<DataType, 1> Buf(Data.get(), DataLen);
queue Q(AsyncHandler{});

Q.submit([&](handler &cgh) {
auto AccGlobal = Buf.get_access<access::mode::read_write>(cgh);
accessor<DataType, DIMS, access::mode::read_write,
access::target::local>
AccLocal(LocalRange, cgh);

cgh.parallel_for<class group__async_work_group_copy>(
nd_range<2>{GlobalRange, LocalRange},
[=](nd_item<DIMS> I) {
const auto Group = I.get_group();
const auto NumElem = AccLocal.get_count();
const auto Off = Group[0] * I.get_group_range(1) * NumElem +
Group[1] * I.get_local_range(1);
auto PtrGlobal = AccGlobal.get_pointer() + Off;
auto PtrLocal = AccLocal.get_pointer();
if (I.get_local_range(0) == 1) {
Group.async_work_group_copy(PtrLocal, PtrGlobal, NumElem);
} else {
Group.async_work_group_copy(PtrLocal, PtrGlobal, NumElem,
I.get_global_range(1));
}
AccLocal[I.get_local_id()][0] += I.get_global_id(0);
AccLocal[I.get_local_id()][1] += I.get_global_id(1);
if (I.get_local_range(0) == 1) {
Group.async_work_group_copy(PtrGlobal, PtrLocal, NumElem);
} else {
Group.async_work_group_copy(PtrGlobal, PtrLocal, NumElem,
I.get_global_range(1));
}
});
});
} catch (cl::sycl::exception const &E) {
std::cout << "SYCL exception caught: " << E.what() << '\n';
return 2;
}
const size_t SIZE_Y = GlobalRange.get(0);
const size_t SIZE_X = GlobalRange.get(1);
int ErrCnt = 0;

for (size_t Y = 0; Y < SIZE_Y; Y++) {
for (size_t X = 0; X < SIZE_X; X++) {
const size_t Ind = Y * SIZE_X + X;
const auto Test0 = Data[Ind][0];
const auto Test1 = Data[Ind][1];
const auto Gold0 = Y;
const auto Gold1 = X;
const bool Ok = (Test0 == Gold0 && Test1 == Gold1);
Pass &= Ok;

if (!Ok && ErrCnt++ < 10) {
std::cout << "*** ERROR at [" << Y << "][" << X << "]: ";
std::cout << Test0 << " " << Test1 << " != ";
std::cout << Gold0 << " " << Gold1 << "\n";
}
}
}
}
if (Pass)
std::cout << " pass\n";
return Pass;
}

int main() {
bool Pass = 1;
Pass &= group__get_group_range();
Pass &= group__get_linear_id();
Pass &= group__async_work_group_copy();

if (!Pass) {
std::cout << "FAILED\n";
Expand Down