https://github.com/ssahasra updated https://github.com/llvm/llvm-project/pull/199175
>From ad5404e9eb3bace085a8c51838de2e1553e7cf90 Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Thu, 21 May 2026 16:55:40 +0530 Subject: [PATCH 1/2] [AMDGPU][Clang] refactor addrspace and scope checks [NFC] Assisted-By: Claude Opus 4.6 --- .../clang/Basic/DiagnosticSemaKinds.td | 2 +- clang/lib/Sema/SemaAMDGPU.cpp | 54 ++++++++++--------- ...mdgcn-error-gfx1250-cooperative-atomics.cl | 2 +- 3 files changed, 31 insertions(+), 27 deletions(-) diff --git a/clang/include/clang/Basic/DiagnosticSemaKinds.td b/clang/include/clang/Basic/DiagnosticSemaKinds.td index dbe6cb2c3a41c..76bacb7d49c8b 100644 --- a/clang/include/clang/Basic/DiagnosticSemaKinds.td +++ b/clang/include/clang/Basic/DiagnosticSemaKinds.td @@ -14230,7 +14230,7 @@ def note_amdgcn_unguarded_builtin_silence : Note<"enclose %0 in a __builtin_amdgcn_is_invocable check to silence " "this warning">; -def err_amdgcn_coop_atomic_invalid_as : Error<"cooperative atomic requires a global or generic pointer">; +def err_amdgcn_global_or_flat_pointer_required : Error<"builtin requires a global or generic pointer">; def err_amdgcn_dmask_has_too_many_bits_set : Error<"dmask argument cannot have more bits set than there are elements " diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 1d2b3898c92d6..385736c7e1eac 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -452,19 +452,35 @@ bool SemaAMDGPU::checkAtomicOrderingCABIArg(Expr *E, bool MayLoad, return false; } -bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { - bool Fail = false; - - // First argument is a global or generic pointer. +// Check that the first argument to TheCall is a global or generic pointer. +static bool checkGlobalOrFlatPointerArg(SemaAMDGPU &S, CallExpr *TheCall) { Expr *PtrArg = TheCall->getArg(0); QualType PtrTy = PtrArg->getType()->getPointeeType(); - unsigned AS = getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace()); + unsigned AS = + S.getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace()); if (AS != llvm::AMDGPUAS::FLAT_ADDRESS && - AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) { - Fail = true; - Diag(TheCall->getBeginLoc(), diag::err_amdgcn_coop_atomic_invalid_as) - << PtrArg->getSourceRange(); - } + AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) + return S.Diag(TheCall->getBeginLoc(), + diag::err_amdgcn_global_or_flat_pointer_required) + << PtrArg->getSourceRange(); + return false; +} + +static bool checkScopeAsInt(SemaAMDGPU &S, Expr *Scope) { + if (Scope->isValueDependent()) + return false; + auto ScopeModel = AtomicScopeModel::create(AtomicScopeModelKind::Generic); + if (std::optional<llvm::APSInt> Result = + Scope->getIntegerConstantExpr(S.SemaRef.Context)) + if (!ScopeModel->isValid(Result->getZExtValue())) + return S.Diag(Scope->getBeginLoc(), + diag::err_atomic_op_has_invalid_sync_scope) + << Scope->getSourceRange(); + return false; +} + +bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { + bool Fail = checkGlobalOrFlatPointerArg(*this, TheCall); Expr *AO = TheCall->getArg(IsStore ? 2 : 1); Expr *Scope = TheCall->getArg(TheCall->getNumArgs() - 1); @@ -488,27 +504,15 @@ bool SemaAMDGPU::checkCoopAtomicFunctionCall(CallExpr *TheCall, bool IsStore) { } bool SemaAMDGPU::checkAtomicMonitorLoad(CallExpr *TheCall) { - bool Fail = false; - Expr *AO = TheCall->getArg(1); Expr *Scope = TheCall->getArg(TheCall->getNumArgs() - 1); if (AO->isValueDependent() || Scope->isValueDependent()) return false; - Fail |= checkAtomicOrderingCABIArg(TheCall->getArg(1), /*MayLoad=*/true, - /*MayStore=*/false); - - auto ScopeModel = AtomicScopeModel::create(AtomicScopeModelKind::Generic); - if (std::optional<llvm::APSInt> Result = - Scope->getIntegerConstantExpr(SemaRef.Context)) { - if (!ScopeModel->isValid(Result->getZExtValue())) { - Diag(Scope->getBeginLoc(), diag::err_atomic_op_has_invalid_sync_scope) - << Scope->getSourceRange(); - Fail = true; - } - } - + bool Fail = checkAtomicOrderingCABIArg(AO, /*MayLoad=*/true, + /*MayStore=*/false); + Fail |= checkScopeAsInt(*this, Scope); return Fail; } diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-cooperative-atomics.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-cooperative-atomics.cl index 8f02e6775d37a..a440a1c040270 100644 --- a/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-cooperative-atomics.cl +++ b/clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-cooperative-atomics.cl @@ -48,7 +48,7 @@ v4i test_amdgcn_cooperative_atomic_load_8x16B_acq_rel(global v4i* gaddr) void test_amdgcn_cooperative_atomic_store_32x4B__sharedptr(local int* addr, int val) { - __builtin_amdgcn_cooperative_atomic_store_32x4B(addr, val, __ATOMIC_RELAXED, ""); // expected-error {{cooperative atomic requires a global or generic pointer}} + __builtin_amdgcn_cooperative_atomic_store_32x4B(addr, val, __ATOMIC_RELAXED, ""); // expected-error {{builtin requires a global or generic pointer}} } void test_amdgcn_cooperative_atomic_store_32x4B__ordering_not_imm(local int* addr, int ord, int val) >From 2539d071f6bfb5821a4a711501961cee1f7a25af Mon Sep 17 00:00:00 2001 From: Sameer Sahasrabuddhe <[email protected]> Date: Thu, 28 May 2026 12:08:33 +0530 Subject: [PATCH 2/2] add braces to long statements --- clang/lib/Sema/SemaAMDGPU.cpp | 9 ++++++--- 1 file changed, 6 insertions(+), 3 deletions(-) diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 385736c7e1eac..ef6aa97578070 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -459,10 +459,11 @@ static bool checkGlobalOrFlatPointerArg(SemaAMDGPU &S, CallExpr *TheCall) { unsigned AS = S.getASTContext().getTargetAddressSpace(PtrTy.getAddressSpace()); if (AS != llvm::AMDGPUAS::FLAT_ADDRESS && - AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) + AS != llvm::AMDGPUAS::GLOBAL_ADDRESS) { return S.Diag(TheCall->getBeginLoc(), diag::err_amdgcn_global_or_flat_pointer_required) << PtrArg->getSourceRange(); + } return false; } @@ -471,11 +472,13 @@ static bool checkScopeAsInt(SemaAMDGPU &S, Expr *Scope) { return false; auto ScopeModel = AtomicScopeModel::create(AtomicScopeModelKind::Generic); if (std::optional<llvm::APSInt> Result = - Scope->getIntegerConstantExpr(S.SemaRef.Context)) - if (!ScopeModel->isValid(Result->getZExtValue())) + Scope->getIntegerConstantExpr(S.SemaRef.Context)) { + if (!ScopeModel->isValid(Result->getZExtValue())) { return S.Diag(Scope->getBeginLoc(), diag::err_atomic_op_has_invalid_sync_scope) << Scope->getSourceRange(); + } + } return false; } _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
