Skip to content

Commit

Permalink
[Sema] atomic_compare_exchange: check failure memory order (#74959)
Browse files Browse the repository at this point in the history 
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, 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.
  • Loading branch information
@MaskRay
MaskRay authored Dec 14, 2023
1 parent 419c45a commit fed5644
Show file tree
Hide file tree
Showing 4 changed files with 69 additions and 22 deletions.
2 changes: 1 addition & 1 deletion clang/include/clang/Basic/DiagnosticSemaKinds.td
View file
Original file line number Diff line number Diff line change
Expand Up @@ -8729,7 +8729,7 @@ 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">,
"%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">;
Expand Down
34 changes: 26 additions & 8 deletions clang/lib/Sema/SemaChecking.cpp
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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:
Expand Down Expand Up @@ -8177,13 +8177,31 @@ 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)
<< /*success=*/(Form == C11CmpXchg || Form == GNUCmpXchg)
<< 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_memory_order)
<< /*failure=*/2 << SubExprs[3]->getSourceRange();
}
}
}
}

if (auto ScopeModel = AtomicExpr::getScopeModel(Op)) {
Expand Down
53 changes: 41 additions & 12 deletions clang/test/Sema/atomic-ops.c
View file
Original file line number Diff line number Diff line change
Expand Up @@ -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);
Expand Down Expand Up @@ -427,19 +427,35 @@ 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);
(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 {{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 {{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);
(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 {{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 {{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);
Expand Down Expand Up @@ -595,19 +611,32 @@ 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);
(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, -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);
(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) {
Expand Down
2 changes: 1 addition & 1 deletion clang/test/SemaCUDA/atomic-ops.cu
View file
Original file line number Diff line number Diff line change
Expand Up @@ -76,7 +76,7 @@ __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_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_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);
Expand Down

0 comments on commit fed5644

Please sign in to comment.