Author: Steffen Larsen Date: 2026-09-09T11:11:35+02:00 New Revision: 68a77b6106d375eabb8a7ec2585a0f3e0382c4a2
URL: https://github.com/llvm/llvm-project/commit/68a77b6106d375eabb8a7ec2585a0f3e0382c4a2 DIFF: https://github.com/llvm/llvm-project/commit/68a77b6106d375eabb8a7ec2585a0f3e0382c4a2.diff LOG: [CIR][AMDGPU] Implement inverse_ballot and read_exec codegen (#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. Added: Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 1c23ea142bf8d..94d5a835d60fa 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: { @@ -541,10 +539,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(); +} _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
