[clang] [AMDGPU][Clang] Handle instantiation-dependent fence arguments (PR #214294)
Shilei Tian via cfe-commits
cfe-commits at lists.llvm.org
Wed Aug 5 10:48:12 PDT 2026
https://github.com/shiltian created https://github.com/llvm/llvm-project/pull/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.
>From dd70646e1a66c349543c7608fcd2b422ea1b016c Mon Sep 17 00:00:00 2001
From: Shilei Tian <i at tianshilei.me>
Date: Wed, 5 Aug 2026 13:43:50 -0400
Subject: [PATCH] [AMDGPU][Clang] Handle instantiation-dependent fence
arguments
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.
---
clang/lib/Sema/SemaAMDGPU.cpp | 97 ++++++++++-----------
clang/test/SemaHIP/amdgpu-builtin-fence.hip | 33 +++++++
2 files changed, 80 insertions(+), 50 deletions(-)
create mode 100644 clang/test/SemaHIP/amdgpu-builtin-fence.hip
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
+//
More information about the cfe-commits
mailing list