Skip to content

Commit

Permalink
Properly set properties of group_load/group_store to striped
Browse files Browse the repository at this point in the history
Doing so exactly recovers the behavior of sub_group::load<vec_sz>,
sub_group::store<vec_sz> and eliminates warnings with 2025.1 and SYCLOS.

With this change, enable use of group_load, group_store for DPC++
compiler with `__SYCL_MAJOR_VERSION >= 8u` which includes
oneAPI DPC++ 2025.0.x compiler and SYCLOS bundle.
oleksandr-pavlyk committed Dec 12, 2024
1 parent 4549117 commit 03910f3
Showing 1 changed file with 9 additions and 5 deletions.
14 changes: 9 additions & 5 deletions dpctl/tensor/libtensor/include/utils/sycl_utils.hpp
Original file line number Diff line number Diff line change
@@ -482,7 +482,7 @@ struct Identity<Op, T, std::enable_if_t<UseBuiltInIdentity<Op, T>::value>>
SYCL_EXT_ONEAPI_GROUP_LOAD_STORE
#define USE_GROUP_LOAD_STORE 1
#else
#if defined(__INTEL_LLVM_COMPILER) && (__INTEL_LLVM_COMPILER > 20250100u)
#if defined(__LIBSYCL_MAJOR_VERSION) && (__LIBSYCL_MAJOR_VERSION >= 8u)
#define USE_GROUP_LOAD_STORE 1
#else
#define USE_GROUP_LOAD_STORE 0
@@ -504,7 +504,8 @@ auto sub_group_load(const sycl::sub_group &sg,
#if (USE_GROUP_LOAD_STORE)
using ValueT = typename std::remove_cv_t<ElementType>;
sycl::vec<ValueT, vec_sz> x{};
ls_ns::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked);
constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped};
ls_ns::group_load(sg, m_ptr, x, striped);
return x;
#else
return sg.load<vec_sz>(m_ptr);
@@ -520,7 +521,8 @@ auto sub_group_load(const sycl::sub_group &sg,
#if (USE_GROUP_LOAD_STORE)
using ValueT = typename std::remove_cv_t<ElementType>;
ValueT x{};
ls_ns::group_load(sg, m_ptr, x, ls_ns::data_placement_blocked);
constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped};
ls_ns::group_load(sg, m_ptr, x, striped);
return x;
#else
return sg.load(m_ptr);
@@ -541,7 +543,8 @@ sub_group_store(const sycl::sub_group &sg,
{
#if (USE_GROUP_LOAD_STORE)
static_assert(std::is_same_v<VecT, ElementType>);
ls_ns::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked);
constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped};
ls_ns::group_store(sg, val, m_ptr, striped);
return;
#else
sg.store<vec_sz>(m_ptr, val);
@@ -561,7 +564,8 @@ sub_group_store(const sycl::sub_group &sg,
sycl::multi_ptr<ElementType, Space, DecorateAddress> m_ptr)
{
#if (USE_GROUP_LOAD_STORE)
ls_ns::group_store(sg, val, m_ptr, ls_ns::data_placement_blocked);
constexpr auto striped = ls_ns::properties{ls_ns::data_placement_striped};
ls_ns::group_store(sg, val, m_ptr, striped);
return;
#else
sg.store(m_ptr, val);

0 comments on commit 03910f3

Please sign in to comment.