Skip to content
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

Closed
wants to merge 8 commits into from

Conversation

SergeyKopienko
Copy link
Contributor

@SergeyKopienko SergeyKopienko commented Feb 7, 2025

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.

@@ -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;
Copy link
Contributor

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.

Copy link
Contributor Author

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.

Copy link
Contributor Author

@SergeyKopienko SergeyKopienko Feb 7, 2025

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
}

Copy link
Contributor

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.

Copy link
Contributor Author

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.

@SergeyKopienko SergeyKopienko marked this pull request as ready for review February 7, 2025 17:37
@SergeyKopienko
Copy link
Contributor Author

As far as I understood this implementation is not preferable, so I close the PR.

@SergeyKopienko SergeyKopienko deleted the dev/skopienko/tmp branch February 10, 2025 14:36
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment
Labels
Projects
None yet
Development

Successfully merging this pull request may close these issues.

3 participants