llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT--> @llvm/pr-subscribers-backend-amdgpu Author: Shilei Tian (shiltian) <details> <summary>Changes</summary> Refactor atomic builtin checks into their switch case and defer constant evaluation of dependent arguments until instantiation, avoiding a potential crash during template definition. Fixes ROCM-29058. --- This change was assisted by AI but I reviewed all changes. --- Full diff: https://github.com/llvm/llvm-project/pull/214294.diff 2 Files Affected: - (modified) clang/lib/Sema/SemaAMDGPU.cpp (+47-50) - (added) clang/test/SemaHIP/amdgpu-builtin-fence.hip (+33) ``````````diff diff --git a/clang/lib/Sema/SemaAMDGPU.cpp b/clang/lib/Sema/SemaAMDGPU.cpp index 11274458dd962..2208865e3237b 100644 --- a/clang/lib/Sema/SemaAMDGPU.cpp +++ b/clang/lib/Sema/SemaAMDGPU.cpp @@ -36,9 +36,6 @@ SemaAMDGPU::SemaAMDGPU(Sema &S) : SemaBase(S) {} bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, CallExpr *TheCall) { - // position of memory order and scope arguments in the builtin - unsigned OrderIndex, ScopeIndex; - const auto *FD = SemaRef.getCurFunctionDecl(/*AllowLambda=*/true); assert(FD && "AMDGPU builtins should not be used outside of a function"); llvm::StringMap<bool> CallerFeatureMap; @@ -110,13 +107,53 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, case AMDGPU::BI__builtin_amdgcn_atomic_inc64: case AMDGPU::BI__builtin_amdgcn_atomic_dec32: case AMDGPU::BI__builtin_amdgcn_atomic_dec64: - OrderIndex = 2; - ScopeIndex = 3; - break; - case AMDGPU::BI__builtin_amdgcn_fence: - OrderIndex = 0; - ScopeIndex = 1; - break; + case AMDGPU::BI__builtin_amdgcn_fence: { + bool IsFence = BuiltinID == AMDGPU::BI__builtin_amdgcn_fence; + unsigned OrderIndex = IsFence ? 0 : 2; + unsigned ScopeIndex = IsFence ? 1 : 3; + Expr *OrderExpr = TheCall->getArg(OrderIndex); + Expr *ScopeExpr = TheCall->getArg(ScopeIndex); + + // Checks requiring constant evaluation are deferred until instantiation. + if (OrderExpr->isInstantiationDependent() || + ScopeExpr->isInstantiationDependent()) + return false; + + Expr::EvalResult OrderResult; + if (!OrderExpr->EvaluateAsInt(OrderResult, getASTContext())) + return Diag(OrderExpr->getExprLoc(), diag::err_typecheck_expect_int) + << OrderExpr->getType(); + uint64_t Ord = OrderResult.Val.getInt().getZExtValue(); + + // Check validity of memory ordering as per C11 / C++11's memory model. + // Only fence needs check. Atomic dec/inc allow all memory orders. + if (!llvm::isValidAtomicOrderingCABI(Ord)) + return Diag(OrderExpr->getBeginLoc(), + diag::warn_atomic_op_has_invalid_memory_order) + << 0 << OrderExpr->getSourceRange(); + switch (static_cast<llvm::AtomicOrderingCABI>(Ord)) { + case llvm::AtomicOrderingCABI::relaxed: + case llvm::AtomicOrderingCABI::consume: + if (IsFence) + return Diag(OrderExpr->getBeginLoc(), + diag::warn_atomic_op_has_invalid_memory_order) + << 0 << OrderExpr->getSourceRange(); + break; + case llvm::AtomicOrderingCABI::acquire: + case llvm::AtomicOrderingCABI::release: + case llvm::AtomicOrderingCABI::acq_rel: + case llvm::AtomicOrderingCABI::seq_cst: + break; + } + + Expr::EvalResult ScopeResult; + // Check that sync scope is a constant literal + if (!ScopeExpr->EvaluateAsConstantExpr(ScopeResult, getASTContext())) + return Diag(ScopeExpr->getExprLoc(), diag::err_expr_not_string_literal) + << ScopeExpr->getType(); + + return false; + } case AMDGPU::BI__builtin_amdgcn_s_setreg: return SemaRef.BuiltinConstantArgRange(TheCall, /*ArgNum=*/0, /*Low=*/0, /*High=*/UINT16_MAX); @@ -410,46 +447,6 @@ bool SemaAMDGPU::CheckAMDGCNBuiltinFunctionCall(unsigned BuiltinID, default: return false; } - - ExprResult Arg = TheCall->getArg(OrderIndex); - auto ArgExpr = Arg.get(); - Expr::EvalResult ArgResult; - - if (!ArgExpr->EvaluateAsInt(ArgResult, getASTContext())) - return Diag(ArgExpr->getExprLoc(), diag::err_typecheck_expect_int) - << ArgExpr->getType(); - auto Ord = ArgResult.Val.getInt().getZExtValue(); - - // Check validity of memory ordering as per C11 / C++11's memory model. - // Only fence needs check. Atomic dec/inc allow all memory orders. - if (!llvm::isValidAtomicOrderingCABI(Ord)) - return Diag(ArgExpr->getBeginLoc(), - diag::warn_atomic_op_has_invalid_memory_order) - << 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) - << 0 << ArgExpr->getSourceRange(); - break; - case llvm::AtomicOrderingCABI::acquire: - case llvm::AtomicOrderingCABI::release: - case llvm::AtomicOrderingCABI::acq_rel: - case llvm::AtomicOrderingCABI::seq_cst: - break; - } - - Arg = TheCall->getArg(ScopeIndex); - ArgExpr = Arg.get(); - Expr::EvalResult ArgResult1; - // Check that sync scope is a constant literal - if (!ArgExpr->EvaluateAsConstantExpr(ArgResult1, getASTContext())) - return Diag(ArgExpr->getExprLoc(), diag::err_expr_not_string_literal) - << ArgExpr->getType(); - - return false; } bool SemaAMDGPU::checkAtomicOrderingCABIArg(Expr *E, bool MayLoad, diff --git a/clang/test/SemaHIP/amdgpu-builtin-fence.hip b/clang/test/SemaHIP/amdgpu-builtin-fence.hip new file mode 100644 index 0000000000000..e032c172e491d --- /dev/null +++ b/clang/test/SemaHIP/amdgpu-builtin-fence.hip @@ -0,0 +1,33 @@ +// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --include-generated-funcs --version 6 +// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -emit-llvm -fcuda-is-device %s -o - | FileCheck %s +// REQUIRES: amdgpu-registered-target + +#define __device__ __attribute__((device)) + +template <int MemoryOrder> +__device__ void test_dependent_order() { + __builtin_amdgcn_fence(MemoryOrder, "agent"); +} + +extern constexpr char AgentScope[] = "agent"; + +template <const char *Scope> +__device__ void test_dependent_scope() { + __builtin_amdgcn_fence(__ATOMIC_SEQ_CST, Scope); +} + +template __device__ void test_dependent_order<__ATOMIC_SEQ_CST>(); +template __device__ void test_dependent_scope<AgentScope>(); +// CHECK-LABEL: define internal void @_Z20test_dependent_orderILi5EEvv( +// CHECK-SAME: ) #[[ATTR0:[0-9]+]] comdat { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: fence syncscope("agent") seq_cst +// CHECK-NEXT: ret void +// +// +// CHECK-LABEL: define internal void @_Z20test_dependent_scopeIXadsoKcL_Z10AgentScopeEEEEvv( +// CHECK-SAME: ) #[[ATTR0]] comdat { +// CHECK-NEXT: [[ENTRY:.*:]] +// CHECK-NEXT: fence syncscope("agent") seq_cst +// CHECK-NEXT: ret void +// `````````` </details> https://github.com/llvm/llvm-project/pull/214294 _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
