llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-clang Author: Sameer Sahasrabuddhe (ssahasra) <details> <summary>Changes</summary> Assisted-By: Claude Opus 4.6 --- Full diff: https://github.com/llvm/llvm-project/pull/199175.diff 3 Files Affected: - (modified) clang/include/clang/Basic/DiagnosticSemaKinds.td (+1-1) - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+29-25) - (modified) clang/test/SemaOpenCL/builtins-amdgcn-error-gfx1250-cooperative-atomics.cl (+1-1) ``````````diff 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) `````````` </details> https://github.com/llvm/llvm-project/pull/199175 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
