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.

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

[SYCL] Support seq_cst memory order in atomic_ref for AMDGPU targets #12938

Closed
wants to merge 14 commits into from
Closed

[SYCL] Support seq_cst memory order in atomic_ref for AMDGPU targets #12938

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
14 commits
Select commit Hold shift + click to select a range
48c0333
[SYCL] Support seq_cst memory order in atomic_ref for AMDGPU targets
GeorgeWeb Mar 7, 2024
d01ae60
Update UR commit tag
GeorgeWeb Mar 7, 2024
06e2ddb
Merge remote-tracking branch 'upstream/sycl' into georgi/sycl-atomic-…
GeorgeWeb Jun 7, 2024
5e47c9c
Update UR hip adapter tag
GeorgeWeb Jun 7, 2024
4c90ed5
Fix formatting issue
GeorgeWeb Jun 7, 2024
f885d44
Merge remote-tracking branch 'upstream/sycl' into georgi/sycl-atomic-…
GeorgeWeb Jun 10, 2024
dac66e6
Merge remote-tracking branch 'upstream/sycl' into georgi/sycl-atomic-…
GeorgeWeb Jun 10, 2024
da66913
Update Hip UR tag
GeorgeWeb Jun 10, 2024
76d59db
Merge remote-tracking branch 'upstream/sycl' into georgi/sycl-atomic-…
GeorgeWeb Jun 13, 2024
6910d2c
Update UR tag again
GeorgeWeb Jun 13, 2024
8f7105a
Apply the correct UR commit tag
GeorgeWeb Jun 13, 2024
4804ba8
Update the acq_rel and seq_cst tests for cuda-arch to require CUDA
GeorgeWeb Jun 13, 2024
fb07094
Merge branch 'sycl' into georgi/sycl-atomic-ref-sc
GeorgeWeb Jun 13, 2024
171fe5a
Update the acq_rel and seq_cst memory order tests to run for both HIP…
GeorgeWeb Jun 13, 2024
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
6 changes: 0 additions & 6 deletions sycl/include/sycl/atomic_ref.hpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -126,12 +126,6 @@ class atomic_ref_base {
detail::IsValidDefaultOrder<DefaultOrder>::value,
"Invalid default memory_order for atomics. Valid defaults are: "
"relaxed, acq_rel, seq_cst");
#ifdef __AMDGPU__
// FIXME should this query device's memory capabilities at runtime?
static_assert(DefaultOrder != sycl::memory_order::seq_cst,
"seq_cst memory order is not supported on AMDGPU");
#endif


public:
using value_type = T;
Expand Down
12 changes: 6 additions & 6 deletions sycl/plugins/unified_runtime/CMakeLists.txt
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -131,13 +131,13 @@ if(SYCL_PI_UR_USE_FETCH_CONTENT)

fetch_adapter_source(hip
${UNIFIED_RUNTIME_REPO}
GeorgeWeb marked this conversation as resolved.
Show resolved Hide resolved
# commit 2c4303c25b026f7edb215accdccb1bc5ae2e237b
# Merge: abe85cc9 3e011c70
# commit 0e32bb59e0846cbe8153c294da6d8f1cfa88075b
# Merge: 2c4303c2 08b19b22
# Author: Kenneth Benzie (Benie) <k.benzie@codeplay.com>
# Date: Thu Jun 13 10:22:34 2024 +0100
# Merge pull request #1414 from GeorgeWeb/georgi/hip-fences
# [HIP] Enable more ordering and scope capabilities for atomic fences
2c4303c25b026f7edb215accdccb1bc5ae2e237b
# Date: Thu Jun 13 14:15:42 2024 +0100
# Merge pull request #1415 from GeorgeWeb/georgi/hip-atomic-memory-caps
# [HIP] Enable more ordering and scope capabilities for atomic memory ops
0e32bb59e0846cbe8153c294da6d8f1cfa88075b
)

fetch_adapter_source(native_cpu
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/AtomicRef/atomic_memory_order_acq_rel.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %{build} -O3 -o %t.out -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70
// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %if any-device-is-cuda %{ -Xsycl-target-backend --cuda-gpu-arch=sm_70 %} %s -o %t.out
// RUN: %{run} %t.out
Copy link
Contributor

@aelovikov-intel aelovikov-intel Jun 13, 2024 •
edited
Loading

Choose a reason for hiding this comment

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

Why can't you use %{build}?


// NOTE: Tests fetch_add for acquire and release memory ordering.
Expand Down
2 changes: 1 addition & 1 deletion sycl/test-e2e/AtomicRef/atomic_memory_order_seq_cst.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
// RUN: %{build} -O3 -o %t.out -Xsycl-target-backend=nvptx64-nvidia-cuda --cuda-gpu-arch=sm_70
// RUN: %clangxx -fsycl -fsycl-targets=%{sycl_triple} %if any-device-is-cuda %{ -Xsycl-target-backend --cuda-gpu-arch=sm_70 %} %s -o %t.out
// RUN: %{run} %t.out

#include "atomic_memory_order.h"
Expand Down
Loading