[clang] [CIR][AMDGPU] Add support for AMDGCN fcmp builtins (PR #198135)
Ayokunle Amodu via cfe-commits
cfe-commits at lists.llvm.org
Thu Jul 30 08:14:04 PDT 2026
https://github.com/ayokunle321 updated https://github.com/llvm/llvm-project/pull/198135
>From fbc52065b05949f0e47f95399d99d71d5c277606 Mon Sep 17 00:00:00 2001
From: Ayokunle Amodu <ayokunle321 at gmail.com>
Date: Thu, 30 Jul 2026 11:13:41 -0400
Subject: [PATCH] correct semantics by adding ballot
---
clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 81 +++++++++++++++----
.../CIR/CodeGenHIP/builtins-amdgcn-wave32.hip | 28 +++++++
.../CIR/CodeGenHIP/builtins-amdgcn-wave64.hip | 28 +++++++
clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip | 60 ++++++++++++++
4 files changed, 181 insertions(+), 16 deletions(-)
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
create mode 100644 clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 63bc9d8b6b144..4b9c84f87f56f 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -345,12 +345,10 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
return mlir::Value{};
}
case AMDGPU::BI__builtin_amdgcn_ballot_w32:
- case AMDGPU::BI__builtin_amdgcn_ballot_w64: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_ballot_w64:
+ return emitBuiltinWithOneOverloadedType<1>(expr, "amdgcn.ballot",
+ convertType(expr->getType()))
+ .getValue();
case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w32:
case AMDGPU::BI__builtin_amdgcn_inverse_ballot_w64: {
cgm.errorNYI(expr->getSourceRange(),
@@ -369,18 +367,69 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
case AMDGPU::BI__builtin_amdgcn_uicmp:
case AMDGPU::BI__builtin_amdgcn_uicmpl:
case AMDGPU::BI__builtin_amdgcn_sicmp:
- case AMDGPU::BI__builtin_amdgcn_sicmpl: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
+ case AMDGPU::BI__builtin_amdgcn_sicmpl:
case AMDGPU::BI__builtin_amdgcn_fcmp:
case AMDGPU::BI__builtin_amdgcn_fcmpf: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
+ mlir::Value lhs = emitScalarExpr(expr->getArg(0));
+ mlir::Value rhs = emitScalarExpr(expr->getArg(1));
+
+ // The predicate is a compile-time constant. Evaluate it from the AST
+ // rather than the emitted value: a non-literal constant expression such as
+ // `39 - 1` lowers to a `cir.binop`, not a `cir.const`, so reading it back
+ // off the IR would fail.
+ uint64_t imm =
+ expr->getArg(2)->EvaluateKnownConstInt(getContext()).getZExtValue();
+
+ cir::CmpOpKind pred;
+ switch (imm) {
+ case 0x1: // FCMP_OEQ
+ case 0x20: // ICMP_EQ
+ pred = cir::CmpOpKind::eq;
+ break;
+ case 0xe: // FCMP_UNE
+ case 0x21: // ICMP_NE
+ pred = cir::CmpOpKind::ne;
+ break;
+ case 0x2: // FCMP_OGT
+ case 0x22: // ICMP_UGT
+ case 0x26: // ICMP_SGT
+ pred = cir::CmpOpKind::gt;
+ break;
+ case 0x3: // FCMP_OGE
+ case 0x23: // ICMP_UGE
+ case 0x27: // ICMP_SGE
+ pred = cir::CmpOpKind::ge;
+ break;
+ case 0x4: // FCMP_OLT
+ case 0x24: // ICMP_ULT
+ case 0x28: // ICMP_SLT
+ pred = cir::CmpOpKind::lt;
+ break;
+ case 0x5: // FCMP_OLE
+ case 0x25: // ICMP_ULE
+ case 0x29: // ICMP_SLE
+ pred = cir::CmpOpKind::le;
+ break;
+ case 0x6: // FCMP_ONE
+ pred = cir::CmpOpKind::one;
+ break;
+ case 0x8: // FCMP_UNO
+ pred = cir::CmpOpKind::uno;
+ break;
+ default:
+ cgm.errorNYI(expr->getSourceRange(),
+ "amdgcn compare with unsupported predicate");
+ return mlir::Value{};
+ }
+
+ mlir::Location loc = getLoc(expr->getExprLoc());
+ mlir::Value cmp = builder.createCompare(loc, pred, lhs, rhs);
+
+ // FIXME-GFX10: The builtin's return type is fixed at uint64_t, so the
+ // ballot mask is always 64 bits wide even on wave32 targets where only the
+ // low 32 bits are meaningful.
+ return builder.emitIntrinsicCallOp(loc, "amdgcn.ballot",
+ convertType(expr->getType()), cmp);
}
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
new file mode 100644
index 0000000000000..3b6a32f9f8715
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip
@@ -0,0 +1,28 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
+// RUN: -target-cpu gfx1100 -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+#define __device__ __attribute__((device))
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU wave32 builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z15test_ballot_w32Pjii
+// CIR: [[CMP:%.+]] = cir.cmp eq %{{.+}}, %{{.+}} : !s32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u32i
+// LLVM: define{{.*}} void @_Z15test_ballot_w32Pjii
+// LLVM: [[CMP:%.+]] = icmp eq i32 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i32 @llvm.amdgcn.ballot.i32(i1 [[CMP]])
+__device__ void test_ballot_w32(unsigned int* out, int a, int b) {
+ *out = __builtin_amdgcn_ballot_w32(a == b);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
new file mode 100644
index 0000000000000..6fd27e9a2140c
--- /dev/null
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip
@@ -0,0 +1,28 @@
+// REQUIRES: amdgpu-registered-target
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu tahiti -fcuda-is-device -emit-cir %s -o %t.cir
+// RUN: FileCheck --check-prefix=CIR --input-file=%t.cir %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 -fclangir \
+// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t-cir.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t-cir.ll %s
+
+// RUN: %clang_cc1 -triple amdgcn-amd-amdhsa -x hip -std=c++11 \
+// RUN: -target-cpu tahiti -fcuda-is-device -emit-llvm %s -o %t.ll
+// RUN: FileCheck --check-prefix=LLVM --input-file=%t.ll %s
+
+#define __device__ __attribute__((device))
+
+//===----------------------------------------------------------------------===//
+// Test AMDGPU wave64 builtins
+//===----------------------------------------------------------------------===//
+
+// CIR-LABEL: @_Z15test_ballot_w64Pmii
+// CIR: [[CMP:%.+]] = cir.cmp eq %{{.+}}, %{{.+}} : !s32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z15test_ballot_w64Pmii
+// LLVM: [[CMP:%.+]] = icmp eq i32 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_ballot_w64(unsigned long* out, int a, int b) {
+ *out = __builtin_amdgcn_ballot_w64(a == b);
+}
diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
index 5dfc95490209b..8acd246ebf3f5 100644
--- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
+++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip
@@ -223,3 +223,63 @@ __device__ void test_frexp_mant_f32(float* out, float a) {
__device__ void test_frexp_mant_f64(double* out, double a) {
*out = __builtin_amdgcn_frexp_mant(a);
}
+
+// CIR-LABEL: @_Z14test_sicmp_i32Pmii
+// CIR: [[CMP:%.+]] = cir.cmp eq %{{.+}}, %{{.+}} : !s32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z14test_sicmp_i32Pmii
+// LLVM: [[CMP:%.+]] = icmp eq i32 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_sicmp_i32(unsigned long* out, int a, int b) {
+ *out = __builtin_amdgcn_sicmp(a, b, 32);
+}
+
+// CIR-LABEL: @_Z14test_uicmp_i32Pmjj
+// CIR: [[CMP:%.+]] = cir.cmp eq %{{.+}}, %{{.+}} : !u32i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z14test_uicmp_i32Pmjj
+// LLVM: [[CMP:%.+]] = icmp eq i32 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_uicmp_i32(unsigned long* out, unsigned a, unsigned b) {
+ *out = __builtin_amdgcn_uicmp(a, b, 32);
+}
+
+// CIR-LABEL: @_Z14test_sicmp_i64Pmll
+// CIR: [[CMP:%.+]] = cir.cmp gt %{{.+}}, %{{.+}} : !s64i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z14test_sicmp_i64Pmll
+// LLVM: [[CMP:%.+]] = icmp sgt i64 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_sicmp_i64(unsigned long* out, long a, long b) {
+ *out = __builtin_amdgcn_sicmpl(a, b, 39 - 1);
+}
+
+// CIR-LABEL: @_Z14test_uicmp_i64Pmmm
+// CIR: [[CMP:%.+]] = cir.cmp ge %{{.+}}, %{{.+}} : !u64i
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z14test_uicmp_i64Pmmm
+// LLVM: [[CMP:%.+]] = icmp uge i64 %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_uicmp_i64(unsigned long* out, unsigned long a, unsigned long b) {
+ *out = __builtin_amdgcn_uicmpl(a, b, 30 + 5);
+}
+
+// CIR-LABEL: @_Z13test_fcmp_f32Pmff
+// CIR: [[CMP:%.+]] = cir.cmp le %{{.+}}, %{{.+}} : !cir.float
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z13test_fcmp_f32Pmff
+// LLVM: [[CMP:%.+]] = fcmp {{.*}}ole float %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_fcmp_f32(unsigned long* out, float a, float b) {
+ *out = __builtin_amdgcn_fcmpf(a, b, 5);
+}
+
+// CIR-LABEL: @_Z13test_fcmp_f64Pmdd
+// CIR: [[CMP:%.+]] = cir.cmp one %{{.+}}, %{{.+}} : !cir.double
+// CIR: cir.call_llvm_intrinsic "amdgcn.ballot" [[CMP]] : (!cir.bool) -> !u64i
+// LLVM: define{{.*}} void @_Z13test_fcmp_f64Pmdd
+// LLVM: [[CMP:%.+]] = fcmp {{.*}}one double %{{.+}}, %{{.+}}
+// LLVM: call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]])
+__device__ void test_fcmp_f64(unsigned long* out, double a, double b) {
+ *out = __builtin_amdgcn_fcmp(a, b, 3 + 3);
+}
More information about the cfe-commits
mailing list