Author: Ayokunle Amodu Date: 2026-09-04T07:26:17-04:00 New Revision: f02d128c60a938fe85690430f3a9ce009669b662
URL: https://github.com/llvm/llvm-project/commit/f02d128c60a938fe85690430f3a9ce009669b662 DIFF: https://github.com/llvm/llvm-project/commit/f02d128c60a938fe85690430f3a9ce009669b662.diff LOG: [CIR][AMDGPU] Add support for AMDGCN cmp and ballot builtins (#198135) Adds CIRGen for the AMDGCN comparison and ballot builtins: - __builtin_amdgcn_fcmp - __builtin_amdgcn_fcmpf - __builtin_amdgcn_sicmp - __builtin_amdgcn_sicmpl - __builtin_amdgcn_uicmp - __builtin_amdgcn_uicmpl - __builtin_amdgcn_ballot_w32 - __builtin_amdgcn_ballot_w64 The cmp builtins are deprecated in favor of ballot_w32 and ballot_w64, but classic CodeGen still supports them, so CIRGen matches it for parity. The deprecation is a Sema-level warning and does not affect the emitted IR. @ranapratap55 Let me know if you would like a comment noting the deprecation, either in the implementation or the tests. Added: clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave32.hip clang/test/CIR/CodeGenHIP/builtins-amdgcn-wave64.hip Modified: clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip Removed: ################################################################################ diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp index e3a12d825434a..0280d38b93bd3 100644 --- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp +++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp @@ -359,12 +359,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(), @@ -380,18 +378,61 @@ 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)); + + 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); + 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..f06a0f8e2658f --- /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 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..f3e75af41830e --- /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 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 8ba5a354cfefa..7dcd343aabb30 100644 --- a/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip +++ b/clang/test/CIR/CodeGenHIP/builtins-amdgcn.hip @@ -224,6 +224,66 @@ __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); +} + // CIR-LABEL: @_Z19test_trig_preop_f32Pffi // CIR: cir.call_llvm_intrinsic "amdgcn.trig.preop" {{.*}} : (!cir.float, !s32i) -> !cir.float // LLVM: define{{.*}} void @_Z19test_trig_preop_f32Pffi _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
