Skip to content

[SYCL] Support 3-, 16-elements vectors in SG load/store #3617

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 5 commits into from
Apr 27, 2021
Merged

[SYCL] Support 3-, 16-elements vectors in SG load/store #3617

Changes from 3 commits
Commits
Show all changes
5 commits
Select commit Hold shift + click to select a range
0f6e5f2
[SYCL] Support 3- and 16-elememnt vectors for sub-group load/store
vladimirlaz Apr 25, 2021
7be5439
Decrease code ducplication for load
vladimirlaz Apr 25, 2021
db5415c
Reduce code duplication for store function
vladimirlaz Apr 26, 2021
89aeb49
Apply review comments
vladimirlaz Apr 26, 2021
27456b5
Apply review comments
vladimirlaz Apr 26, 2021
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
142 changes: 91 additions & 51 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 @@ -295,29 +295,70 @@ struct sub_group {
PI_INVALID_DEVICE);
#endif
}

#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N != 1,
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
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
}
#else // __NVPTX__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N != 1 && N != 3 && N != 16,
vec<T, N>>
load(const multi_ptr<T, Space> src) const {
return sycl::detail::sub_group::load<N, T>(src);
#endif // __NVPTX__
#else
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 16,
vec<T, 16>>
load(const multi_ptr<T, Space> src) const {
return {sycl::detail::sub_group::load<8, T>(src),
sycl::detail::sub_group::load<8, T>(src +
8 * get_max_local_range()[0])};
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 3,
vec<T, 3>>
load(const multi_ptr<T, Space> src) const {
auto res = sycl::detail::sub_group::load<4, T>(src);
return {res.s2(), res.s1(), res.s0()};
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 1,
vec<T, 1>>
load(const multi_ptr<T, Space> src) const {
return sycl::detail::sub_group::load(src);
}
#endif // ___NVPTX___
#else // __SYCL_DEVICE_ONLY__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value,
vec<T, N>>
load(const multi_ptr<T, Space> src) const {
(void)src;
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}
#endif // __SYCL_DEVICE_ONLY__

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
Expand All @@ -337,25 +378,6 @@ struct sub_group {
#endif
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 1,
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.",
PI_INVALID_DEVICE);
#endif
}

#ifdef __SYCL_DEVICE_ONLY__
// Method for decorated pointer
template <typename T>
Expand Down Expand Up @@ -437,45 +459,63 @@ struct sub_group {
#endif
}

#ifdef __SYCL_DEVICE_ONLY__
#ifdef __NVPTX__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space> dst, const vec<T, N> &x) const {
for (int i = 0; i < N; ++i) {
*(dst.get() + i * get_max_local_range()[0] + get_local_id()[0]) = x[i];
}
}
#else // __NVPTX__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N != 1 && N != 3 && N != 16>
store(multi_ptr<T, Space> dst, const vec<T, N> &x) const {
sycl::detail::sub_group::store(dst, x);
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
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;
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
sycl::detail::sub_group::store(dst, x);
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
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);
N == 3>
store(multi_ptr<T, Space> dst, const vec<T, 3> &x) const {
store<1, T, Space>(dst, x.s0());
store<2, T, Space>(dst + get_max_local_range()[0], {x.s1(), x.s2()});
}

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value &&
N == 16>
store(multi_ptr<T, Space> dst, const vec<T, 16> &x) const {
store<8, T, Space>(dst, x.lo());
store<8, T, Space>(dst + 8 * get_max_local_range()[0], x.hi());
}

#endif // __NVPTX__
#else
#else // __SYCL_DEVICE_ONLY__
template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
sycl::detail::sub_group::AcceptableForGlobalLoadStore<T, Space>::value>
store(multi_ptr<T, Space> dst, const vec<T, N> &x) const {
(void)dst;
(void)x;
throw runtime_error("Sub-groups are not supported on host device.",
PI_INVALID_DEVICE);
#endif
}
#endif // __SYCL_DEVICE_ONLY__

template <int N, typename T, access::address_space Space>
sycl::detail::enable_if_t<
Expand Down