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 <[email protected]> Date: Thu, 30 Jul 2026 11:13:41 -0400 Subject: [PATCH 1/4] 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); +} >From 778119a3e100b0b000e3e228867e29f02ccae68a Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 30 Jul 2026 11:14:19 -0400 Subject: [PATCH 2/4] fix style --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 4b9c84f87f56f..952e4c7b37746 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -421,7 +421,7 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, "amdgcn compare with unsupported predicate"); return mlir::Value{}; } - + mlir::Location loc = getLoc(expr->getExprLoc()); mlir::Value cmp = builder.createCompare(loc, pred, lhs, rhs); >From e4313f018ad5ee7c97796dde486bf7a425fbeb63 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 30 Jul 2026 11:21:05 -0400 Subject: [PATCH 3/4] remove uneccessary comments --- clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp | 8 -------- 1 file changed, 8 deletions(-) diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index 952e4c7b37746..bfa61e953593b 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -373,10 +373,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, 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(); @@ -424,10 +420,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId, 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); } >From ef2a16d7ca8e0b79586d5735b6ecdfcea04dd223 Mon Sep 17 00:00:00 2001 From: Ayokunle Amodu <[email protected]> Date: Thu, 30 Jul 2026 11:22:15 -0400 Subject: [PATCH 4/4] make test comments consistent --- clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip | 2 +- clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip | 2 +- 2 files changed, 2 insertions(+), 2 deletions(-) diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip index 3b6a32f9f8715..f06a0f8e2658f 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip @@ -14,7 +14,7 @@ #define __device__ __attribute__((device)) //===----------------------------------------------------------------------===// -// Test AMDGPU wave32 builtins +// Test AMDGPU builtins //===----------------------------------------------------------------------===// // CIR-LABEL: @_Z15test_ballot_w32Pjii diff --git a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip index 6fd27e9a2140c..f3e75af41830e 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip @@ -14,7 +14,7 @@ #define __device__ __attribute__((device)) //===----------------------------------------------------------------------===// -// Test AMDGPU wave64 builtins +// Test AMDGPU builtins //===----------------------------------------------------------------------===// // CIR-LABEL: @_Z15test_ballot_w64Pmii _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
