Skip to content

[SYCL][CUDA] Enable sub-group loads and stores #2763

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
Jan 13, 2021
Merged

[SYCL][CUDA] Enable sub-group loads and stores #2763

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
30 changes: 30 additions & 0 deletions sycl/include/CL/sycl/ONEAPI/sub_group.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -230,7 +230,11 @@ struct sub_group {
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value, T>
load(const multi_ptr<T, Space> src) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
return src.get()[get_local_id()[0]];
#else
return sycl::detail::sub_group::load(src);
#endif // __NVPTX__
#else
(void)src;
throw runtime_error("Sub-groups are not supported on host device.",
Expand Down Expand Up @@ -258,7 +262,15 @@ struct sub_group {
vec<T, N>>
load(const multi_ptr<T, Space> src) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
vec<T, N> res;
for (int i = 0; i < N; ++i) {
res[i] = *(src.get() + i * get_max_local_range()[0] + get_local_id()[0]);
}
return res;
#else
return sycl::detail::sub_group::load<N, T>(src);
#endif // __NVPTX__
#else
(void)src;
throw runtime_error("Sub-groups are not supported on host device.",
Expand Down Expand Up @@ -291,7 +303,11 @@ struct sub_group {
vec<T, 1>>
load(const multi_ptr<T, Space> src) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
return src.get()[get_local_id()[0]];
#else
return sycl::detail::sub_group::load(src);
#endif // __NVPTX__
#else
(void)src;
throw runtime_error("Sub-groups are not supported on host device.",
Expand All @@ -304,7 +320,11 @@ struct sub_group {
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space> dst, const T &x) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
dst.get()[get_local_id()[0]] = x;
#else
sycl::detail::sub_group::store(dst, x);
#endif // __NVPTX__
#else
(void)dst;
(void)x;
Expand Down Expand Up @@ -333,7 +353,11 @@ struct sub_group {
N == 1>
store(multi_ptr<T, Space> dst, const vec<T, 1> &x) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
dst.get()[get_local_id()[0]] = x[0];
#else
store<T, Space>(dst, x);
#endif // __NVPTX__
#else
(void)dst;
(void)x;
Expand All @@ -348,7 +372,13 @@ struct sub_group {
N != 1>
store(multi_ptr<T, Space> dst, const vec<T, N> &x) const {
#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
for (int i = 0; i < N; ++i) {
*(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
}
#else
sycl::detail::sub_group::store(dst, x);
#endif // __NVPTX__
#else
(void)dst;
(void)x;
Expand Down