-
Notifications
You must be signed in to change notification settings - Fork 114
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.
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
Alternative approach of working with barriers #2055
Conversation
@@ -74,6 +74,9 @@ struct __subgroup_radix_sort | |||
uint16_t __buf_size; | |||
|
|||
public: | |||
|
|||
static constexpr sycl::access::fence_space _space = sycl::access::fence_space::local_space; |
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think the problem with this approach is that sycl::access::fence_space
is not a part of SYCL2020 at all, so we cannot expect a compiler to implement it.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I think it's simple not documented now, but not deprecated too.
And we haven't any warnings about that.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
Practically we are using this enum in oneDPL code in main branch too, but only in one place:
template <typename _Item>
constexpr void
__group_barrier(_Item __item)
{
#if 0 // !defined(_ONEDPL_LIBSYCL_VERSION) || _ONEDPL_LIBSYCL_VERSION >= 50300
//TODO: usage of sycl::group_barrier: probably, we have to revise SYCL parallel patterns which use a group_barrier.
// 1) sycl::group_barrier() implementation is not ready
// 2) sycl::group_barrier and sycl::item::group_barrier are not quite equivalent
sycl::group_barrier(__item.get_group(), sycl::memory_scope::work_group);
#else
__item.barrier(sycl::access::fence_space::local_space); // <<<<<<<<<<<<< HERE
#endif
}
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
I see, so I guess the current state is probably not compatible with non-dpcpp SYCL compilers unless they also have this SYCL 1.2.1 support.
There was a problem hiding this comment.
Choose a reason for hiding this comment
The reason will be displayed to describe this comment to others. Learn more.
May be. But in any case, we are not making anything worse at this place.
…_one_wg.h - apply GitHUB clang format
…_one_wg.h - apply GitHUB clang format
As far as I understood this implementation is not preferable, so I close the PR. |
The initial issue has been detected and investigated by @dmitriy-sobolev in the PR #2054
This PR simple make the fit for that issue by a little bit another way.