[clang] fb559f5 - [AMDGPU][Clang] Handle instantiation-dependent fence arguments (#214294)
via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 5 11:47:33 PDT 2026
Author: Shilei Tian
Date: 2026-08-05T14:47:29-04:00
New Revision: fb559f5f016a35c07ba179b447bb10019db254c1
URL: https://github.com/llvm/llvm-project/commit/fb559f5f016a35c07ba179b447bb10019db254c1
DIFF: https://github.com/llvm/llvm-project/commit/fb559f5f016a35c07ba179b447bb10019db254c1.diff
LOG: [AMDGPU][Clang] Handle instantiation-dependent fence arguments (#214294)
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.
Added:
clang/test/SemaHIP/amdgpu-builtin-fence.hip
Modified:
clang/lib/Sema/SemaAMDGPU.cpp
Removed:
################################################################################
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..3c1d347c89e43
--- /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 amdgpu-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
+//
More information about the cfe-commits
mailing list