[clang] [CIR][AMDGPU] Implement inverse_ballot and read_exec codegen (PR #221661)
Steffen Larsen via cfe-commits
cfe-commits at lists.llvm.org
Mon Sep 7 00:19:43 PDT 2026
https://github.com/steffenlarsen created https://github.com/llvm/llvm-project/pull/221661
This commit implements the CIR codegen for the following AMDGPU builtins:
- __builtin_amdgcn_inverse_ballot_w32
- __builtin_amdgcn_inverse_ballot_w64
- __builtin_amdgcn_read_exec
- __builtin_amdgcn_read_exec_lo
- __builtin_amdgcn_read_exec_hi
inverse_ballot_w32/w64 map to llvm.amdgcn.inverse.ballot. read_exec, read_exec_lo and read_exec_hi read the exec mask as a ballot over an all-true predicate, at least as wide as the wavefront.
>From ed7ace17940bd619ff9be5ce0cfd5f2622511282 Mon Sep 17 00:00:00 2001
From: Steffen Holst Larsen <sholstla at amd.com>
Date: Tue, 1 Sep 2026 09:00:57 -0500
Subject: [PATCH] [CIR][AMDGPU] Implement inverse_ballot and read_exec codegen
This commit implements the CIR codegen for the following AMDGPU builtins:
- __builtin_amdgcn_inverse_ballot_w32
- __builtin_amdgcn_inverse_ballot_w64
- __builtin_amdgcn_read_exec
- __builtin_amdgcn_read_exec_lo
- __builtin_amdgcn_read_exec_hi
inverse_ballot_w32/w64 map to llvm.amdgcn.inverse.ballot.
read_exec, read_exec_lo and read_exec_hi read the exec mask as a ballot
over an all-true predicate, at least as wide as the wavefront.
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 31 +++++++++-----
.../CIR/CodeGenHIP/builtins-amdgcn-wave32.hip | 40 ++++++++++++++++++
.../CIR/CodeGenHIP/builtins-amdgcn-wave64.hip | 42 +++++++++++++++++++
3 files changed, 103 insertions(+), 10 deletions(-)
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0280d38b93bd3..2b09ac92e482c 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -364,12 +364,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
convertType(expr->getType()))
.getValue();
case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
- case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64:
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.inverse.ballot",
+ convertType(expr->getType()))
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_tanhf:
case AMDGPU::BI__builtin_amdgcn_tanhh:
case AMDGPU::BI__builtin_amdgcn_tanh_bf16: {
@@ -543,10 +541,23 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_read_exec:
case AMDGPU::BI__builtin_amdgcn_read_exec_lo:
case AMDGPU::BI__builtin_amdgcn_read_exec_hi: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ // The exec mask is read as a ballot over an all-true predicate. The
+ // ballot is at least as wide as the wavefront, so that a wave64 target
+ // still reports both halves for the _lo/_hi forms.
+ mlir::Location loc = getLoc(expr->getExprLoc());
+ unsigned registerWidth =
+ builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_lo ? 32 : 64;
+ unsigned ballotWidth =
+ std::max(getTarget().getGridValue().GV_Warp_Size, registerWidth);
+ cir::IntType ballotTy = builder.getUIntNTy(ballotWidth);
+
+ mlir::Value truePred = builder.getBool(true, loc).getResult();
+ mlir::Value result =
+ builder.emitIntrinsicCallOp(loc, "amdgcn.ballot", ballotTy, truePred);
+
+ if (builtinId == AMDGPU::BI__builtin_amdgcn_read_exec_hi)
+ result = builder.createShiftRight(loc, result, 32);
+ return builder.createIntCast(result, convertType(expr->getType()));
}
case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray:
case AMDGPU::BI__builtin_amdgcn_image_bvh_intersect_ray_h:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
index f06a0f8e2658f..34a930ca3f44c 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
@@ -26,3 +26,43 @@
__device__ void test_ballot_w32(unsigned int* out, int a, int b) {
*out = __builtin_amdgcn_ballot_w32(a == b);
}
+
+// CIR-LABEL: @_Z4ib64y
+// CIR: cir.call_llvm_intrinsic "amdgcn.inverse.ballot" %{{.+}} : (!u32i) -> !cir.bool
+// LLVM-LABEL: @_Z4ib64y
+// LLVM: call i1 @llvm.amdgcn.inverse.ballot.i32(i32
+__device__ bool ib64(unsigned long long m) {
+ return __builtin_amdgcn_inverse_ballot_w32(m);
+}
+
+// CIR-LABEL: @_Z2exv
+// CIR: [[TRUE1:%.+]] = cir.const #true
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE1]] : (!cir.bool) -> !u64i
+// LLVM-LABEL: @_Z2exv
+// LLVM: call i64 @llvm.amdgcn.ballot.i64(i1 true)
+__device__ unsigned long long ex() {
+ return __builtin_amdgcn_read_exec();
+}
+
+// CIR-LABEL: @_Z4exlov
+// CIR: [[TRUE2:%.+]] = cir.const #true
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE2]] : (!cir.bool) -> !u32i
+// LLVM-LABEL: @_Z4exlov
+// LLVM: call i32 @llvm.amdgcn.ballot.i32(i1 true)
+__device__ unsigned int exlo() {
+ return __builtin_amdgcn_read_exec_lo();
+}
+
+// CIR-LABEL: @_Z4exhiv
+// CIR: [[TRUE3:%.+]] = cir.const #true
+// CIR: [[BRES:%.+]] = cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE3]] : (!cir.bool) -> !u64i
+// CIR: [[C32:%.+]] = cir.const #cir.int<32> : !u64i
+// CIR: [[SRES:%.+]] = cir.shift(right, [[BRES]] : !u64i, [[C32]] : !u64i) -> !u64i
+// CIR: cir.cast integral [[SRES]] : !u64i -> !u32i
+// LLVM-LABEL: @_Z4exhiv
+// LLVM: [[BRES:%.+]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
+// LLVM: [[SRES:%.+]] = lshr i64 [[BRES]], 32
+// LLVM: trunc i64 [[SRES]] to i32
+__device__ unsigned int exhi() {
+ return __builtin_amdgcn_read_exec_hi();
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
index f3e75af41830e..afe43f0cc6271 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
@@ -26,3 +26,45 @@
__device__ void test_ballot_w64(unsigned long* out, int a, int b) {
*out = __builtin_amdgcn_ballot_w64(a == b);
}
+
+// CIR-LABEL: @_Z4ib64y
+// CIR: cir.call_llvm_intrinsic "amdgcn.inverse.ballot" %{{.+}} : (!u64i) -> !cir.bool
+// LLVM-LABEL: @_Z4ib64y
+// LLVM: call i1 @llvm.amdgcn.inverse.ballot.i64(i64
+__device__ bool ib64(unsigned long long m) {
+ return __builtin_amdgcn_inverse_ballot_w64(m);
+}
+
+// CIR-LABEL: @_Z2exv
+// CIR: [[TRUE1:%.+]] = cir.const #true
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE1]] : (!cir.bool) -> !u64i
+// LLVM-LABEL: @_Z2exv
+// LLVM: call i64 @llvm.amdgcn.ballot.i64(i1 true)
+__device__ unsigned long long ex() {
+ return __builtin_amdgcn_read_exec();
+}
+
+// CIR-LABEL: @_Z4exlov
+// CIR: [[TRUE2:%.+]] = cir.const #true
+// CIR: [[BRES1:%.+]] = cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE2]] : (!cir.bool) -> !u64i
+// CIR: cir.cast integral [[BRES1]] : !u64i -> !u32i
+// LLVM-LABEL: @_Z4exlov
+// LLVM: [[BRES1:%.+]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
+// LLVM: trunc i64 [[BRES1]] to i32
+__device__ unsigned int exlo() {
+ return __builtin_amdgcn_read_exec_lo();
+}
+
+// CIR-LABEL: @_Z4exhiv
+// CIR: [[TRUE3:%.+]] = cir.const #true
+// CIR: [[BRES2:%.+]] = cir.call_llvm_intrinsic "amdgcn.ballot" [[TRUE3]] : (!cir.bool) -> !u64i
+// CIR: [[C32:%.+]] = cir.const #cir.int<32> : !u64i
+// CIR: [[SRES:%.+]] = cir.shift(right, [[BRES2]] : !u64i, [[C32]] : !u64i) -> !u64i
+// CIR: cir.cast integral [[SRES]] : !u64i -> !u32i
+// LLVM-LABEL: @_Z4exhiv
+// LLVM: [[BRES2:%.+]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
+// LLVM: [[SRES:%.+]] = lshr i64 [[BRES2]], 32
+// LLVM: trunc i64 [[SRES]] to i32
+__device__ unsigned int exhi() {
+ return __builtin_amdgcn_read_exec_hi();
+}
More information about the cfe-commits
mailing list