Skip to content

[SYCL][E2E] Avoid illegal failure order in compare_exchange_strong test #19513

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

Open
wants to merge 1 commit into
base: sycl
Choose a base branch
from
Open

[SYCL][E2E] Avoid illegal failure order in compare_exchange_strong test #19513

Changes from all commits
Commits
Show all changes
1 commit
Select commit
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
27 changes: 24 additions & 3 deletions sycl/test-e2e/AtomicRef/compare_exchange.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -49,7 +49,14 @@ void compare_exchange_local_test(queue q, size_t N) {
: order,
scope, space > (loc[0]);
T result = T(N); // Avoid copying pointer
bool success = atm.compare_exchange_strong(result, (T)gid, order);
// From SYCL AtomicRef spec: The failure memory order of this atomic
// operation must be relaxed, acquire or seq_cst.
auto failure_order =
(order == memory_order::acq_rel || order == memory_order::release)
? memory_order::relaxed
Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Why don't we just assert that failure_order is allowed and then modify the callers? Is that because we use the same parameter to test multiple functions and other functions can accept that?

If so, I think we should use the strictest memory order and not the weakest instead of the unsupported, because theoretically the caller might be checking for side effects that weaker one can't guarantee.

Copy link
Author

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

I've chosen memory_order::relaxed as the fallback for failure_order, because that is what is used for the default memory order fallback for AtomicRef in this test. I can change this to acquire -- I don't think anything stronger is required since there are no side effects to release on failure.

Another possibility, and my preferred route, is changing the sycl::AtomicRef::compare_exchange_strong() function to change illegal failure memory orders to the next best legal ones: use acquire instead of acq_rel and relaxed instead of release. This would match the behaviour of the C++ std library: https://en.cppreference.com/w/cpp/atomic/atomic_ref/compare_exchange

Re modifying the test callers. The compare_exchange_test_orders_scopes template function generates all the memory orders to test. This order template parameter is ultimately used for both success_order and failure_order in compare_exchange_strong. If we prevent calls with acq_rel and release being used for order in compare_exchange_test_orders_scopes, then success_order will also not be tested with these variants. If we care about all combinations of success_order and failure_order being tested, then we could add another template function that generates all allowed failure orders, in addition to all possible success orders and memory scopes. However, I don't think the E2E are expected to test every possible parameter combination.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

+ @gmlueck for

Another possibility, and my preferred route, is changing the sycl::AtomicRef::compare_exchange_strong() function to change illegal failure memory orders to the next best legal ones: use acquire instead of acq_rel and relaxed instead of release. This would match the behaviour of the C++ std library: https://en.cppreference.com/w/cpp/atomic/atomic_ref/compare_exchange

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Delegating to @Pennycook

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Good catch. This is definitely a specification bug: we say it's equivalent to compare_exchange_strong(expected, desired, order, order, scope), but that would be UB in several cases. I'll open a pull request against the SYCL specification to fix this.

FWIW, here's the equivalent wording in the ISO C++ specification, which is what I'll try to align with:

When only one memory_order argument is supplied, the value of success is order, and the value of failure is order except that a value of memory_order​::​acq_rel shall be replaced by the value memory_order​::​acquire and a value of memory_order​::​release shall be replaced by the value memory_order​::​relaxed.

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

New wording proposed in KhronosGroup/SYCL-Docs#891.

Copy link
Author

@robertszafa robertszafa Jul 22, 2025 •
edited
Loading

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Great, it'd be nice to have this be clear in the SYCL spec.

Do you want me to open a new PR that implements this in atomic_ref_base::compare_exchange_strong(expected, desired, order, order, scope)? This would mean that the change from this PR to the compare_exchange_strong E2E test is not needed.

All the compare_exchange variants call into atomic_ref_base::compare_exchange_strong(expected, desired, order, order, scope), so adding something like the below there should be enough:

    failure = (failure == memory_order::acq_rel)   ? memory_order::acquire
              : (failure == memory_order::release) ? memory_order::relaxed
                                                   : failure;

Copy link
Contributor

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more.

Do you want me to open a new PR that implements this in atomic_ref_base::compare_exchange_strong(expected, desired, order, order, scope)? This would mean that the change from this PR to the compare_exchange_strong E2E test is not needed.

Yes, please do.

We might want to hold off on merging it until KhronosGroup/SYCL-Docs#891 is merged, just to be on the safe side. But given that this is already the behavior of ISO C++, the chances of it being accepted as a bug fix are very high.

: order;
bool success =
atm.compare_exchange_strong(result, (T)gid, order, failure_order);
if (success) {
out[gid] = result;
} else {
Expand Down Expand Up @@ -99,7 +106,14 @@ void compare_exchange_global_test(queue q, size_t N) {
: order,
scope, space > (exc[0]);
T result = T(N); // Avoid copying pointer
bool success = atm.compare_exchange_strong(result, (T)gid, order);
// From SYCL AtomicRef spec: The failure memory order of this atomic
// operation must be relaxed, acquire or seq_cst.
auto failure_order =
(order == memory_order::acq_rel || order == memory_order::release)
? memory_order::relaxed
: order;
bool success =
atm.compare_exchange_strong(result, (T)gid, order, failure_order);
if (success) {
out[gid] = result;
} else {
Expand Down Expand Up @@ -140,7 +154,14 @@ void compare_exchange_global_test_usm_shared(queue q, size_t N) {
: order,
scope, space > (exc[0]);
T result = initial; // Avoid copying pointer
bool success = atm.compare_exchange_strong(result, (T)gid, order);
// From SYCL AtomicRef spec: The failure memory order of this atomic
// operation must be relaxed, acquire or seq_cst.
auto failure_order =
(order == memory_order::acq_rel || order == memory_order::release)
? memory_order::relaxed
: order;
bool success =
atm.compare_exchange_strong(result, (T)gid, order, failure_order);
if (success) {
output[gid] = result;
} else {
Expand Down