[clang] [Sema] atomic_compare_exchange: check failure memory order (PR #74959)
Fangrui Song via cfe-commits
cfe-commits at lists.llvm.org
Mon Dec 11 15:21:16 PST 2023
https://github.com/MaskRay updated https://github.com/llvm/llvm-project/pull/74959
>From 56a63cd9f910dc78163adda15252e7818cc7a419 Mon Sep 17 00:00:00 2001
From: Fangrui Song <i at maskray.me>
Date: Sat, 9 Dec 2023 16:37:51 -0800
Subject: [PATCH 1/2] [Sema] atomic_compare_exchange: check failure memory
order
For
`__atomic_compare_exchange{,_n}/__c11_atomic_compare_exchange_{strong,weak}`,
GCC checks both the success memory order and the failure memory order
under the default -Winvalid-memory-model ("memory model" is confusing
here and "memory order" is much more common in the atomic context).
* The failure memory order cannot be stronger than the success memory
order.
* The failure memory order, if a constant, must be one of
relaxed/consume/acquire/seq_cst.
Clang checks just the success memory order under the default
-Watomic-memory-ordering. This patch checks the failure memory order.
---
clang/include/clang/Basic/DiagnosticGroups.td | 1 +
.../clang/Basic/DiagnosticSemaKinds.td | 8 ++++-
clang/lib/Sema/SemaChecking.cpp | 32 +++++++++++++++----
clang/test/Sema/atomic-ops.c | 31 ++++++++++++++++--
clang/test/SemaCUDA/atomic-ops.cu | 8 ++---
5 files changed, 66 insertions(+), 14 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index caee2dc6daadb..a282f37292a03 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -822,6 +822,7 @@ def UndeclaredSelector : DiagGroup<"undeclared-selector">;
def ImplicitAtomic : DiagGroup<"implicit-atomic-properties">;
def AtomicAlignment : DiagGroup<"atomic-alignment">;
def CustomAtomic : DiagGroup<"custom-atomic-properties">;
+def AtomicMemoryOrdering : DiagGroup<"atomic-memory-ordering">;
def AtomicProperties : DiagGroup<"atomic-properties",
[ImplicitAtomic, CustomAtomic]>;
def SyncAlignment : DiagGroup<"sync-alignment">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 28d95ca9b1389..03ed23473c8d3 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8728,7 +8728,13 @@ def err_atomic_op_needs_atomic_int : Error<
"%select{|atomic }0integer (%1 invalid)">;
def warn_atomic_op_has_invalid_memory_order : Warning<
"memory order argument to atomic operation is invalid">,
- InGroup<DiagGroup<"atomic-memory-ordering">>;
+ InGroup<AtomicMemoryOrdering>;
+def warn_atomic_op_has_invalid_failure_memory_order : Warning<
+ "failure memory order argument to atomic operation is invalid">,
+ InGroup<AtomicMemoryOrdering>;
+def warn_atomic_op_failure_memory_order_stronger_than_success : Warning<
+ "failure memory order cannot be stronger than success memory order">,
+ InGroup<AtomicMemoryOrdering>;
def err_atomic_op_has_invalid_synch_scope : Error<
"synchronization scope argument to atomic operation is invalid">;
def warn_atomic_implicit_seq_cst : Warning<
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 5c97346184470..4b1f2ccd27c01 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -8181,13 +8181,33 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
break;
}
+ // If the memory orders are constants, check they are valid.
if (SubExprs.size() >= 2 && Form != Init) {
- if (std::optional<llvm::APSInt> Result =
- SubExprs[1]->getIntegerConstantExpr(Context))
- if (!isValidOrderingForOp(Result->getSExtValue(), Op))
- Diag(SubExprs[1]->getBeginLoc(),
- diag::warn_atomic_op_has_invalid_memory_order)
- << SubExprs[1]->getSourceRange();
+ std::optional<llvm::APSInt> Success =
+ SubExprs[1]->getIntegerConstantExpr(Context);
+ if (Success && !isValidOrderingForOp(Success->getSExtValue(), Op))
+ Diag(SubExprs[1]->getBeginLoc(),
+ diag::warn_atomic_op_has_invalid_memory_order)
+ << SubExprs[1]->getSourceRange();
+ if (SubExprs.size() >= 5) {
+ if (std::optional<llvm::APSInt> Failure =
+ SubExprs[3]->getIntegerConstantExpr(Context)) {
+ if (!llvm::is_contained(
+ {llvm::AtomicOrderingCABI::relaxed,
+ llvm::AtomicOrderingCABI::consume,
+ llvm::AtomicOrderingCABI::acquire,
+ llvm::AtomicOrderingCABI::seq_cst},
+ (llvm::AtomicOrderingCABI)Failure->getSExtValue())) {
+ Diag(SubExprs[3]->getBeginLoc(),
+ diag::warn_atomic_op_has_invalid_failure_memory_order)
+ << SubExprs[3]->getSourceRange();
+ } else if (Success && *Success < *Failure) {
+ Diag(SubExprs[3]->getBeginLoc(),
+ diag::warn_atomic_op_failure_memory_order_stronger_than_success)
+ << SubExprs[3]->getSourceRange();
+ }
+ }
+ }
}
if (auto ScopeModel = AtomicExpr::getScopeModel(Op)) {
diff --git a/clang/test/Sema/atomic-ops.c b/clang/test/Sema/atomic-ops.c
index 4fa1223b3038f..d8bf366c44a35 100644
--- a/clang/test/Sema/atomic-ops.c
+++ b/clang/test/Sema/atomic-ops.c
@@ -432,14 +432,28 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_consume, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_release, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acq_rel, memory_order_relaxed);
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_relaxed);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_acquire);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_consume);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_seq_cst);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_relaxed, memory_order_acquire); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acquire, memory_order_seq_cst); // expected-warning {{failure memory order cannot be stronger than success memory order}}
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_consume, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_release, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acq_rel, memory_order_relaxed);
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_relaxed);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_acquire);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_consume);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_seq_cst);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_acquire); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_seq_cst); // expected-warning {{failure memory order cannot be stronger than success memory order}}
(void)__atomic_load_n(p, memory_order_relaxed);
(void)__atomic_load_n(p, memory_order_acquire);
@@ -600,7 +614,12 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_consume, memory_order_relaxed);
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_release, memory_order_relaxed);
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_acq_rel, memory_order_relaxed);
- (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_relaxed);
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_acquire);
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_consume);
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_release); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_seq_cst);
+ (void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_relaxed, memory_order_relaxed);
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_acquire, memory_order_relaxed);
@@ -608,6 +627,12 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_release, memory_order_relaxed);
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_acq_rel, memory_order_relaxed);
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_relaxed);
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_acquire);
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_consume);
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_release); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, memory_order_seq_cst);
+ (void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
}
void nullPointerWarning(void) {
diff --git a/clang/test/SemaCUDA/atomic-ops.cu b/clang/test/SemaCUDA/atomic-ops.cu
index af93b7e1e7944..0b22e81ec9ea3 100644
--- a/clang/test/SemaCUDA/atomic-ops.cu
+++ b/clang/test/SemaCUDA/atomic-ops.cu
@@ -73,10 +73,10 @@ __device__ bool test_hip_atomic_cmpxchg_weak(int *ptr, int val, int desired) {
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SINGLETHREAD);
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_CONSUME, __HIP_MEMORY_SCOPE_SINGLETHREAD);
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SINGLETHREAD);
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQ_REL, __HIP_MEMORY_SCOPE_SINGLETHREAD);
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_CONSUME, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQ_REL, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order argument to atomic operation is invalid}}
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_SEQ_CST, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_CONSUME, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
>From 97c198c1f0c4b0caa4bacc2e6ecddeca8604cc58 Mon Sep 17 00:00:00 2001
From: Fangrui Song <i at maskray.me>
Date: Mon, 11 Dec 2023 14:52:14 -0800
Subject: [PATCH 2/2] Drop "success >= failure" check even if GCC has it. C++
compare_exchange_strong precondition does not require it.
---
clang/include/clang/Basic/DiagnosticGroups.td | 1 -
.../clang/Basic/DiagnosticSemaKinds.td | 10 +----
clang/lib/Sema/SemaChecking.cpp | 16 ++++----
clang/test/Sema/atomic-ops.c | 38 ++++++++++---------
clang/test/SemaCUDA/atomic-ops.cu | 6 +--
5 files changed, 33 insertions(+), 38 deletions(-)
diff --git a/clang/include/clang/Basic/DiagnosticGroups.td b/clang/include/clang/Basic/DiagnosticGroups.td
index a282f37292a03..caee2dc6daadb 100644
--- a/clang/include/clang/Basic/DiagnosticGroups.td
+++ b/clang/include/clang/Basic/DiagnosticGroups.td
@@ -822,7 +822,6 @@ def UndeclaredSelector : DiagGroup<"undeclared-selector">;
def ImplicitAtomic : DiagGroup<"implicit-atomic-properties">;
def AtomicAlignment : DiagGroup<"atomic-alignment">;
def CustomAtomic : DiagGroup<"custom-atomic-properties">;
-def AtomicMemoryOrdering : DiagGroup<"atomic-memory-ordering">;
def AtomicProperties : DiagGroup<"atomic-properties",
[ImplicitAtomic, CustomAtomic]>;
def SyncAlignment : DiagGroup<"sync-alignment">;
diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td
index 03ed23473c8d3..9197a9560efbd 100644
--- a/clang/include/clang/Basic/DiagnosticSemaKinds.td
+++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td
@@ -8727,14 +8727,8 @@ def err_atomic_op_needs_atomic_int : Error<
"address argument to atomic operation must be a pointer to "
"%select{|atomic }0integer (%1 invalid)">;
def warn_atomic_op_has_invalid_memory_order : Warning<
- "memory order argument to atomic operation is invalid">,
- InGroup<AtomicMemoryOrdering>;
-def warn_atomic_op_has_invalid_failure_memory_order : Warning<
- "failure memory order argument to atomic operation is invalid">,
- InGroup<AtomicMemoryOrdering>;
-def warn_atomic_op_failure_memory_order_stronger_than_success : Warning<
- "failure memory order cannot be stronger than success memory order">,
- InGroup<AtomicMemoryOrdering>;
+ "%select{|success |failure }0memory order argument to atomic operation is invalid">,
+ InGroup<DiagGroup<"atomic-memory-ordering">>;
def err_atomic_op_has_invalid_synch_scope : Error<
"synchronization scope argument to atomic operation is invalid">;
def warn_atomic_implicit_seq_cst : Warning<
diff --git a/clang/lib/Sema/SemaChecking.cpp b/clang/lib/Sema/SemaChecking.cpp
index 4b1f2ccd27c01..d36ae43e91add 100644
--- a/clang/lib/Sema/SemaChecking.cpp
+++ b/clang/lib/Sema/SemaChecking.cpp
@@ -5030,14 +5030,14 @@ bool Sema::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID,
if (!llvm::isValidAtomicOrderingCABI(Ord))
return Diag(ArgExpr->getBeginLoc(),
diag::warn_atomic_op_has_invalid_memory_order)
- << ArgExpr->getSourceRange();
+ << 0 << ArgExpr->getSourceRange();
switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) {
case llvm::AtomicOrderingCABI::relaxed:
case llvm::AtomicOrderingCABI::consume:
if (BuiltinID == AMDGPU::BI__builtin_amdgcn_fence)
return Diag(ArgExpr->getBeginLoc(),
diag::warn_atomic_op_has_invalid_memory_order)
- << ArgExpr->getSourceRange();
+ << 0 << ArgExpr->getSourceRange();
break;
case llvm::AtomicOrderingCABI::acquire:
case llvm::AtomicOrderingCABI::release:
@@ -8185,10 +8185,12 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
if (SubExprs.size() >= 2 && Form != Init) {
std::optional<llvm::APSInt> Success =
SubExprs[1]->getIntegerConstantExpr(Context);
- if (Success && !isValidOrderingForOp(Success->getSExtValue(), Op))
+ if (Success && !isValidOrderingForOp(Success->getSExtValue(), Op)) {
Diag(SubExprs[1]->getBeginLoc(),
diag::warn_atomic_op_has_invalid_memory_order)
+ << /*success=*/(Form == C11CmpXchg || Form == GNUCmpXchg)
<< SubExprs[1]->getSourceRange();
+ }
if (SubExprs.size() >= 5) {
if (std::optional<llvm::APSInt> Failure =
SubExprs[3]->getIntegerConstantExpr(Context)) {
@@ -8199,12 +8201,8 @@ ExprResult Sema::BuildAtomicExpr(SourceRange CallRange, SourceRange ExprRange,
llvm::AtomicOrderingCABI::seq_cst},
(llvm::AtomicOrderingCABI)Failure->getSExtValue())) {
Diag(SubExprs[3]->getBeginLoc(),
- diag::warn_atomic_op_has_invalid_failure_memory_order)
- << SubExprs[3]->getSourceRange();
- } else if (Success && *Success < *Failure) {
- Diag(SubExprs[3]->getBeginLoc(),
- diag::warn_atomic_op_failure_memory_order_stronger_than_success)
- << SubExprs[3]->getSourceRange();
+ diag::warn_atomic_op_has_invalid_memory_order)
+ << /*failure=*/2 << SubExprs[3]->getSourceRange();
}
}
}
diff --git a/clang/test/Sema/atomic-ops.c b/clang/test/Sema/atomic-ops.c
index d8bf366c44a35..1d36667d6cf40 100644
--- a/clang/test/Sema/atomic-ops.c
+++ b/clang/test/Sema/atomic-ops.c
@@ -339,18 +339,18 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__c11_atomic_load(Ap, memory_order_relaxed);
(void)__c11_atomic_load(Ap, memory_order_acquire);
(void)__c11_atomic_load(Ap, memory_order_consume);
- (void)__c11_atomic_load(Ap, memory_order_release); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_load(Ap, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_load(Ap, memory_order_release); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_load(Ap, memory_order_acq_rel); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
(void)__c11_atomic_load(Ap, memory_order_seq_cst);
(void)__c11_atomic_load(Ap, val);
- (void)__c11_atomic_load(Ap, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_load(Ap, 42); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_load(Ap, -1); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_load(Ap, 42); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
(void)__c11_atomic_store(Ap, val, memory_order_relaxed);
- (void)__c11_atomic_store(Ap, val, memory_order_acquire); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_store(Ap, val, memory_order_consume); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_store(Ap, val, memory_order_acquire); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_store(Ap, val, memory_order_consume); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
(void)__c11_atomic_store(Ap, val, memory_order_release);
- (void)__c11_atomic_store(Ap, val, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_store(Ap, val, memory_order_acq_rel); // expected-warning-re {{{{^}}memory order argument to atomic operation is invalid}}
(void)__c11_atomic_store(Ap, val, memory_order_seq_cst);
(void)__c11_atomic_fetch_add(Ap, 1, memory_order_relaxed);
@@ -427,6 +427,7 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__c11_atomic_exchange(Ap, val, memory_order_acq_rel);
(void)__c11_atomic_exchange(Ap, val, memory_order_seq_cst);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, -1, memory_order_relaxed); // expected-warning {{success memory order argument to atomic operation is invalid}}
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_relaxed, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acquire, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_consume, memory_order_relaxed);
@@ -434,13 +435,14 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acq_rel, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_acquire);
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_consume);
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{failure memory order argument to atomic operation is invalid}}
(void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, memory_order_seq_cst);
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_relaxed, memory_order_acquire); // expected-warning {{failure memory order cannot be stronger than success memory order}}
- (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acquire, memory_order_seq_cst); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_relaxed, memory_order_acquire);
+ (void)__c11_atomic_compare_exchange_strong(Ap, p, val, memory_order_acquire, memory_order_seq_cst);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, -1, memory_order_relaxed); // expected-warning {{success memory order argument to atomic operation is invalid}}
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_consume, memory_order_relaxed);
@@ -448,12 +450,12 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acq_rel, memory_order_relaxed);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_acquire);
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_consume);
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_release); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_acq_rel); // expected-warning {{failure memory order argument to atomic operation is invalid}}
(void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, memory_order_seq_cst);
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_acquire); // expected-warning {{failure memory order cannot be stronger than success memory order}}
- (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_seq_cst); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_seq_cst, -1); // expected-warning {{failure memory order argument to atomic operation is invalid}}
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_relaxed, memory_order_acquire);
+ (void)__c11_atomic_compare_exchange_weak(Ap, p, val, memory_order_acquire, memory_order_seq_cst);
(void)__atomic_load_n(p, memory_order_relaxed);
(void)__atomic_load_n(p, memory_order_acquire);
@@ -609,6 +611,7 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__atomic_exchange(p, p, p, memory_order_acq_rel);
(void)__atomic_exchange(p, p, p, memory_order_seq_cst);
+ (void)__atomic_compare_exchange(p, p, p, 0, -1, memory_order_relaxed); // expected-warning {{success memory order argument to atomic operation is invalid}}
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_relaxed, memory_order_relaxed);
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_acquire, memory_order_relaxed);
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_consume, memory_order_relaxed);
@@ -621,6 +624,7 @@ void memory_checks(_Atomic(int) *Ap, int *p, int val) {
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, memory_order_seq_cst);
(void)__atomic_compare_exchange(p, p, p, 0, memory_order_seq_cst, -1); // expected-warning {{memory order argument to atomic operation is invalid}}
+ (void)__atomic_compare_exchange_n(p, p, val, 0, -1, memory_order_relaxed); // expected-warning {{success memory order argument to atomic operation is invalid}}
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_relaxed, memory_order_relaxed);
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_acquire, memory_order_relaxed);
(void)__atomic_compare_exchange_n(p, p, val, 0, memory_order_consume, memory_order_relaxed);
diff --git a/clang/test/SemaCUDA/atomic-ops.cu b/clang/test/SemaCUDA/atomic-ops.cu
index 0b22e81ec9ea3..233ed1c10fc11 100644
--- a/clang/test/SemaCUDA/atomic-ops.cu
+++ b/clang/test/SemaCUDA/atomic-ops.cu
@@ -73,9 +73,9 @@ __device__ bool test_hip_atomic_cmpxchg_weak(int *ptr, int val, int desired) {
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_WORKGROUP);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_AGENT);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_CONSUME, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
- flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order cannot be stronger than success memory order}}
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_SEQ_CST, __HIP_MEMORY_SCOPE_SINGLETHREAD);
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_CONSUME, __HIP_MEMORY_SCOPE_SINGLETHREAD);
+ flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQUIRE, __HIP_MEMORY_SCOPE_SINGLETHREAD);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_ACQ_REL, __HIP_MEMORY_SCOPE_SINGLETHREAD); // expected-warning {{failure memory order argument to atomic operation is invalid}}
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_RELAXED, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
flag = __hip_atomic_compare_exchange_weak(ptr, &val, desired, __ATOMIC_SEQ_CST, __ATOMIC_RELAXED, __HIP_MEMORY_SCOPE_SINGLETHREAD);
More information about the cfe-commits
mailing list