Fix synchronization issues in radix sort and histogram #2054

dmitriy-sobolev · 2025-02-07T16:42:21Z

Value (__buf_val) and Count (__buf_count) buffers store data in either local or global memory. Let's use group barriers with proper fences (local or global memory fences) to avoid memory contention issues, which are observed on Xe2 architectures.

The fix changes how __group_barrier is defined, so the PR additionally includes what is done in #1988 to avoid conflicts and fix another issue related to the barriers: we use SYCL 1.2.1 barriers, but oneDPL claims to be SYCL 2020 compatible.

SergeyKopienko · 2025-02-07T16:52:28Z

Taking into account the discussion from #1679
I think that probably the switcher which barrier exactly we should use in all oneDPL code - may be incorrect approach...

adamfidel · 2025-02-07T17:17:19Z

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h

+                            else
+                                __dpl_sycl::__group_barrier(__it, __dpl_sycl::__fence_space_global_and_local{});
+                        };
+


This just a thought, but could we use a single lambda here and pass the boolean constant as a parameter to it? Something like:

auto __mem_adjusted_barrier = [__it](auto __is_slm) { if constexpr (decltype(__is_slm)::value) __dpl_sycl::__group_barrier(__it); else __dpl_sycl::__group_barrier(__it, __dpl_sycl::__fence_space_global_and_local{}); };

And when used:

__mem_adjusted_barrier(_SLM_tag_val{});

__mem_adjusted_barrier(_SLM_counter{});

See https://godbolt.org/z/GMx64dTb5.

Why do we need a compile time tag here?
Whether is it simpler to call with a constant, wrapped into __dpl_sycl::name_xx, like?

__dpl_sycl::__group_barrier(__it, __dpl_sycl::fence_global_and_local);

@MikeDvorskiy, I do not understand how your suggestion is supposed to work. We need to use different fence arguments depending on an SLM tag. Could you be more specific? Perhaps I just do not get your idea.

Meanwhile, I've implemented what Adam suggested.

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h

MikeDvorskiy · 2025-02-07T17:19:35Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

+
+#if ONEDPL_SYCL121_GROUP_BARRIER
+template <sycl::access::fence_space _Space>
+struct __fence_space


I would suggest no writing 3floor C++ code....
Actually, we already have the simple approach for wrapping sycl constants. For example,
see __target_device usage.
https://github.com/uxlfoundation/oneDPL/pull/2054/files#diff-521f7db06a55567c1a1dffa855dde585fb6bc5fbe8633d19b528356e3a501ea0R481

https://github.com/uxlfoundation/oneDPL/blob/main/include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_utils.h#L539C9-L539C109

The problem is that I need to provide some no-op alternative in contrast to the case with __target_device. Defining structures helps to achieve it. Could you provide an example in this context?

Perhaps you can do this a little simpler...

#if ONEDPL_SYCL121_GROUP_BARRIER constexpr sycl::access::fence_space __fence_space_local = sycl::access::fence_space::local_space; constexpr sycl::access::fence_space __fence_space_global = sycl::access::fence_space::global_space; constexpr sycl::access::fence_space __fence_space_global_and_local = sycl::access::fence_space::global_and_local; #else constexpr int __fence_space_local = 0; constexpr int __fence_space_global = 0; constexpr int __fence_space_global_and_local = 0; #endif // ONEDPL_SYCL121_GROUP_BARRIER template <typename _Item, typename _Space = decltype(__fence_space_local)> void __group_barrier(_Item __item, [[maybe_unused]] _Space __space = __fence_space_local) { #if ONEDPL_SYCL121_GROUP_BARRIER __item.barrier(__space); #elif _ONEDPL_SYCL2020_GROUP_BARRIER_PRESENT sycl::group_barrier(__item.get_group(), sycl::memory_scope::work_group); #else # error "sycl::group_barrier is not supported, and no alternative is available" #endif }

Does this work? maybe there is a better "dummy" type / value than int, but otherwise I think it would work.

You would just need to remove the {} from __dpl_sycl::__fence_space_global{} above I think.

I guess if we are following the example of __target_device we can do a similar trick to create a typename __fence_space_t which is either sycl::access::fence_space or int depending on ONEDPL_SYCL121_GROUP_BARRIER.

The option with global constants looks good. I've implemented it with adding a dummy type.

SergeyKopienko · 2025-02-07T17:20:14Z

My assumptions that the changes from this PR is a little bit overcomplicated.
Let me propose alternative approach for this fix: #2055

MikeDvorskiy · 2025-02-07T17:33:40Z

include/oneapi/dpl/pstl/hetero/dpcpp/sycl_defs.h

    sycl::group_barrier(__item.get_group(), sycl::memory_scope::work_group);
 #else
-    __item.barrier(sycl::access::fence_space::local_space);
+#    error "sycl::group_barrier is not supported, and no alternative is available"


Whether do we get potential issues here with user's code in case when neither ONEDPL_SYCL121_GROUP_BARRIER neither _ONEDPL_SYCL2020_GROUP_BARRIER_PRESENT is defined?

If neither macro is defined, this error will appear. _ONEDPL_SYCL2020_GROUP_BARRIER_PRESENT is defined for any compiler irrespective to a version, and for icpx starting with some old release (some pre-2022 release). Is that an issue? Did I get you right?

SergeyKopienko · 2025-02-07T17:46:02Z

About
"The fix changes how __group_barrier is defined, so the PR additionally includes what is done in #1988 to avoid conflicts and fix another issue related to the barriers: we use SYCL 1.2.1 barriers, but oneDPL claims to be SYCL 2020 compatible."

I believe it's incorrect moment to change that before release. Let's do this change later in the next release.

dmitriy-sobolev · 2025-02-07T19:01:51Z

Taking into account the discussion from #1679 I think that probably the switcher which barrier exactly we should use in all oneDPL code - may be incorrect approach...

I briefly checked the other places: they all seem to be using local memory. Hence the current approach should be fine (with this radix-sort as an exception). Another state with SYCL 2020 barrier, which has stronger guarantees, also looks grand. But let me delve deeper.

Co-authored-by: Dan Hoeflinger <[email protected]>

SergeyKopienko · 2025-02-07T20:41:25Z

If somebody will have interest for #2055, please reopen it.

danhoeflinger · 2025-02-07T20:59:26Z

I briefly checked the other places: they all seem to be using local memory.
But let me delve deeper.

It looks like one implementation of histogram may have a similar issue (my fault):

__clear_wglocal_histograms when called from __histogram_general_private_global_atomics_submitter is operating on global data, also the barrier in the top level parallel for lambda for that same kernel should be a global fence.

I can come up with a fix if you want, its also something which could probably wait until after the release.

danhoeflinger · 2025-02-07T21:22:18Z

I've added #2056 in case we want to include it here. The rest of the usages I believe should be in the local memory space. If we have any question here with the changes, lets wait because this implementation requires a very large number of bins to be used.

SergeyKopienko · 2025-02-10T09:19:27Z

I think as far as @dmitriy-sobolev introduced three new types here, probably required to declare that they are device-copyable (is_device_copyable):

__fence_space_local
__fence_space_global
__fence_space_global_and_local

SergeyKopienko · 2025-02-10T09:26:56Z

Additional question: in the current state of this PR the value __fence_space_global_and_local is never used.
So should we really declare this value __fence_space_global_and_local ?

Co-authored-by: Dmitriy Sobolev <[email protected]>

dmitriy-sobolev · 2025-02-10T13:33:51Z

Additional question: in the current state of this PR the value __fence_space_global_and_local is never used. So should we really declare this value __fence_space_global_and_local ?

They were indeed unnecessary, I've removed them.

dmitriy-sobolev · 2025-02-10T14:47:02Z

I think as far as @dmitriy-sobolev introduced three new types here, probably required to declare that they are device-copyable (is_device_copyable):

__fence_space_local

__fence_space_global

__fence_space_global_and_local

I see a couple of problems here:

We should not copy these object onto a device.
sycl::access::fence_space (enum) and __fence_space_dummy are a trivially copyable types, so they should also be device copyable.

We discussed it offline, and the issue was work-arounded by declaring get_fence as static (a function this issue occurred), but I still have no understanding why the lack of static triggered that issue...

danhoeflinger

I agree with the first clang formatting complaint, not the others. Otherwise, LGTM.

Probably makes sense to move the TODO above the line.

danhoeflinger

LGTM

mmichel11

LGTM

Co-authored-by: Dan Hoeflinger <[email protected]> Co-authored-by: Sergey Kopienko <[email protected]>

Co-authored-by: Dmitriy Sobolev <[email protected]> Co-authored-by: Dan Hoeflinger <[email protected]> Co-authored-by: Sergey Kopienko <[email protected]>

dmitriy-sobolev added 6 commits February 7, 2025 14:29

Use global mem barriers

ac04fb3

Generalize for SYCL 1.2.1. and 2020

df03fd7

Abstract barrier

2062f9e

Abstract barrier, lambda

e6e2a92

Abstract barrier, conservative approach

0df22d0

Allow specifying SYCL barrier type

bd882dc

dmitriy-sobolev added the bug label Feb 7, 2025

dmitriy-sobolev requested review from MikeDvorskiy, SergeyKopienko, akukanov, danhoeflinger, mmichel11 and adamfidel February 7, 2025 16:42

dmitriy-sobolev added this to the 2022.8.0 milestone Feb 7, 2025

dmitriy-sobolev and others added 3 commits February 7, 2025 17:01

ONEDPL_USE_SYCL121_GROUP_BARRIER -> ONEDPL_SYCL121_GROUP_BARRIER

aaaa029

Fix a type in the macro description

5c28f9f

clang-format

8a5d554

adamfidel reviewed Feb 7, 2025

View reviewed changes

mmichel11 reviewed Feb 7, 2025

View reviewed changes

include/oneapi/dpl/pstl/hetero/dpcpp/parallel_backend_sycl_radix_sort_one_wg.h Outdated Show resolved Hide resolved

dmitriy-sobolev changed the title ~~Fix memory contention in radix sort~~ Fix synchronization issues in radix sort Feb 7, 2025

MikeDvorskiy reviewed Feb 7, 2025

View reviewed changes

Refactor: __mem_adjusted_barrier

4834d3e

MikeDvorskiy reviewed Feb 7, 2025

View reviewed changes

SergeyKopienko mentioned this pull request Feb 7, 2025

Alternative approach of working with barriers #2055

Closed

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/radix-fix branch 2 times, most recently from 008e9a8 to 4834d3e Compare February 7, 2025 18:37

Fix a typo

e2d9f10

Co-authored-by: Dan Hoeflinger <[email protected]>

SergeyKopienko removed their request for review February 7, 2025 20:38

SergeyKopienko mentioned this pull request Feb 10, 2025

Simplify fence usage in radix sort #2057

Merged

SergeyKopienko self-requested a review February 10, 2025 09:28

Remove unused fence alias

affaabe

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/radix-fix branch from bb31221 to affaabe Compare February 10, 2025 11:01

Simplify fence usage in radix sort (#2057)

720ba1d

Co-authored-by: Dmitriy Sobolev <[email protected]>

Make ONEDPL_SYCL121_GROUP_BARRIER internal

8430e0c

dmitriy-sobolev force-pushed the dev/dmitriy-sobolev/radix-fix branch from 7c0a99c to 8430e0c Compare February 10, 2025 16:41

dmitriy-sobolev and others added 2 commits February 10, 2025 17:17

Add more comments

36934cd

Global barrier fence histogram fix (#2056)

35da181

danhoeflinger previously approved these changes Feb 10, 2025

View reviewed changes

clang-format

e925187

dmitriy-sobolev dismissed danhoeflinger’s stale review via e925187 February 10, 2025 17:42

Remove __fence_space_dummy in favor of __fence_space_t

4164c26

dmitriy-sobolev changed the title ~~Fix synchronization issues in radix sort~~ Fix synchronization issues in radix sort and histogram Feb 10, 2025

danhoeflinger approved these changes Feb 10, 2025

View reviewed changes

mmichel11 approved these changes Feb 10, 2025

View reviewed changes

dmitriy-sobolev merged commit 0bc7341 into main Feb 10, 2025
21 of 22 checks passed

dmitriy-sobolev deleted the dev/dmitriy-sobolev/radix-fix branch February 10, 2025 20:18

dmitriy-sobolev mentioned this pull request Feb 11, 2025

Use SYCL 2020 barrier instead of SYCL 1.2.1 one by default #1988

Closed

timmiesmith pushed a commit that referenced this pull request Feb 11, 2025

Fix synchronization issues in radix sort and histogram (#2054)

47c5daf

Co-authored-by: Dan Hoeflinger <[email protected]> Co-authored-by: Sergey Kopienko <[email protected]>

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Fix synchronization issues in radix sort and histogram #2054

Fix synchronization issues in radix sort and histogram #2054

dmitriy-sobolev commented Feb 7, 2025 •

edited

Loading

SergeyKopienko commented Feb 7, 2025 •

edited

Loading

adamfidel Feb 7, 2025

MikeDvorskiy Feb 7, 2025

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

dmitriy-sobolev Feb 7, 2025

MikeDvorskiy Feb 7, 2025 •

edited

Loading

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

danhoeflinger Feb 7, 2025

danhoeflinger Feb 7, 2025

dmitriy-sobolev Feb 7, 2025

SergeyKopienko commented Feb 7, 2025

MikeDvorskiy Feb 7, 2025

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

SergeyKopienko commented Feb 7, 2025

dmitriy-sobolev commented Feb 7, 2025 •

edited

Loading

SergeyKopienko commented Feb 7, 2025

danhoeflinger commented Feb 7, 2025 •

edited

Loading

danhoeflinger commented Feb 7, 2025

SergeyKopienko commented Feb 10, 2025 •

edited

Loading

SergeyKopienko commented Feb 10, 2025

dmitriy-sobolev commented Feb 10, 2025

dmitriy-sobolev commented Feb 10, 2025 •

edited

Loading

danhoeflinger left a comment

danhoeflinger left a comment

mmichel11 left a comment

Fix synchronization issues in radix sort and histogram #2054

Fix synchronization issues in radix sort and histogram #2054

Conversation

dmitriy-sobolev commented Feb 7, 2025 • edited Loading

SergeyKopienko commented Feb 7, 2025 • edited Loading

adamfidel Feb 7, 2025

Choose a reason for hiding this comment

MikeDvorskiy Feb 7, 2025

Choose a reason for hiding this comment

dmitriy-sobolev Feb 7, 2025 • edited Loading

Choose a reason for hiding this comment

dmitriy-sobolev Feb 7, 2025

Choose a reason for hiding this comment

MikeDvorskiy Feb 7, 2025 • edited Loading

Choose a reason for hiding this comment

dmitriy-sobolev Feb 7, 2025 • edited Loading

Choose a reason for hiding this comment

danhoeflinger Feb 7, 2025

Choose a reason for hiding this comment

danhoeflinger Feb 7, 2025

Choose a reason for hiding this comment

dmitriy-sobolev Feb 7, 2025

Choose a reason for hiding this comment

SergeyKopienko commented Feb 7, 2025

MikeDvorskiy Feb 7, 2025

Choose a reason for hiding this comment

dmitriy-sobolev Feb 7, 2025 • edited Loading

Choose a reason for hiding this comment

SergeyKopienko commented Feb 7, 2025

dmitriy-sobolev commented Feb 7, 2025 • edited Loading

SergeyKopienko commented Feb 7, 2025

danhoeflinger commented Feb 7, 2025 • edited Loading

danhoeflinger commented Feb 7, 2025

SergeyKopienko commented Feb 10, 2025 • edited Loading

SergeyKopienko commented Feb 10, 2025

dmitriy-sobolev commented Feb 10, 2025

dmitriy-sobolev commented Feb 10, 2025 • edited Loading

danhoeflinger left a comment

Choose a reason for hiding this comment

danhoeflinger left a comment

Choose a reason for hiding this comment

mmichel11 left a comment

Choose a reason for hiding this comment

dmitriy-sobolev commented Feb 7, 2025 •

edited

Loading

SergeyKopienko commented Feb 7, 2025 •

edited

Loading

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

MikeDvorskiy Feb 7, 2025 •

edited

Loading

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

dmitriy-sobolev Feb 7, 2025 •

edited

Loading

dmitriy-sobolev commented Feb 7, 2025 •

edited

Loading

danhoeflinger commented Feb 7, 2025 •

edited

Loading

SergeyKopienko commented Feb 10, 2025 •

edited

Loading

dmitriy-sobolev commented Feb 10, 2025 •

edited

Loading