https://github.com/jayfoad created https://github.com/llvm/llvm-project/pull/208231
Use `llvm.amdgcn.ballot` instead of the deprecated intrinsics `llvm.amdgcn.icmp` and `llvm.amdgcn.fcmp`. >From 1cff29ab173184ed5aad52abdeafbd453db135cf Mon Sep 17 00:00:00 2001 From: Jay Foad <[email protected]> Date: Wed, 8 Jul 2026 15:18:51 +0100 Subject: [PATCH] [AMDGPU] Reimplement icmp and fcmp builtins using ballot Use `llvm.amdgcn.ballot` instead of the deprecated intrinsics `llvm.amdgcn.icmp` and `llvm.amdgcn.fcmp`. --- clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp | 24 +++++++------------ clang/test/CodeGenCUDA/builtins-amdgcn.cu | 7 +++--- .../test/CodeGenCUDA/builtins-spirv-amdgcn.cu | 14 ++++++----- clang/test/CodeGenOpenCL/builtins-amdgcn.cl | 22 ++++++++++------- 4 files changed, 34 insertions(+), 33 deletions(-) diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp index 6a17f4c55f80d..227a2c4d33bdd 100644 --- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp +++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp @@ -775,26 +775,18 @@ Value *CodeGenFunction::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: { - llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); - llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); - llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); - - // FIXME-GFX10: How should 32 bit mask be handled? - Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_icmp, - { Builder.getInt64Ty(), Src0->getType() }); - return Builder.CreateCall(F, { Src0, Src1, Src2 }); - } + case AMDGPU::BI__builtin_amdgcn_sicmpl: case AMDGPU::BI__builtin_amdgcn_fcmp: case AMDGPU::BI__builtin_amdgcn_fcmpf: { - llvm::Value *Src0 = EmitScalarExpr(E->getArg(0)); - llvm::Value *Src1 = EmitScalarExpr(E->getArg(1)); - llvm::Value *Src2 = EmitScalarExpr(E->getArg(2)); + Value *LHS = EmitScalarExpr(E->getArg(0)); + Value *RHS = EmitScalarExpr(E->getArg(1)); + CmpInst::Predicate Pred = static_cast<CmpInst::Predicate>( + cast<ConstantInt>(EmitScalarExpr(E->getArg(2)))->getZExtValue()); // FIXME-GFX10: How should 32 bit mask be handled? - Function *F = CGM.getIntrinsic(Intrinsic::amdgcn_fcmp, - { Builder.getInt64Ty(), Src0->getType() }); - return Builder.CreateCall(F, { Src0, Src1, Src2 }); + return Builder.CreateIntrinsic(Builder.getInt64Ty(), + Intrinsic::amdgcn_ballot, + Builder.CreateCmp(Pred, LHS, RHS)); } case AMDGPU::BI__builtin_amdgcn_class: case AMDGPU::BI__builtin_amdgcn_classf: diff --git a/clang/test/CodeGenCUDA/builtins-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-amdgcn.cu index 35673773ec80c..47c6ba57ec2f2 100644 --- a/clang/test/CodeGenCUDA/builtins-amdgcn.cu +++ b/clang/test/CodeGenCUDA/builtins-amdgcn.cu @@ -182,9 +182,10 @@ __global__ void endpgm() { // CHECK-NEXT: store i64 [[B:%.*]], ptr [[B_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr [[A_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr [[B_ADDR_ASCAST]], align 8 -// CHECK-NEXT: [[TMP2:%.*]] = call i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP0]], i64 [[TMP1]], i32 35) -// CHECK-NEXT: [[TMP3:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 -// CHECK-NEXT: store i64 [[TMP2]], ptr [[TMP3]], align 8 +// CHECK-NEXT: [[TMP2:%.*]] = icmp uge i64 [[TMP0]], [[TMP1]] +// CHECK-NEXT: [[TMP3:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 [[TMP2]]) +// CHECK-NEXT: [[TMP4:%.*]] = load ptr, ptr [[OUT_ADDR_ASCAST]], align 8 +// CHECK-NEXT: store i64 [[TMP3]], ptr [[TMP4]], align 8 // CHECK-NEXT: ret void // __global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b) diff --git a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu index 94338f9027db1..ccd76037b03cf 100644 --- a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu +++ b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu @@ -275,9 +275,10 @@ __global__ void endpgm() { // CHECK-NEXT: store i64 [[B:%.*]], ptr addrspace(4) [[B_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // CHECK-NEXT: [[TMP1:%.*]] = load i64, ptr addrspace(4) [[B_ADDR_ASCAST]], align 8 -// CHECK-NEXT: [[TMP2:%.*]] = call addrspace(4) i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP0]], i64 [[TMP1]], i32 35) -// CHECK-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8 -// CHECK-NEXT: store i64 [[TMP2]], ptr addrspace(4) [[TMP3]], align 8 +// CHECK-NEXT: [[TMP2:%.*]] = icmp uge i64 [[TMP0]], [[TMP1]] +// CHECK-NEXT: [[TMP3:%.*]] = call addrspace(4) i64 @llvm.amdgcn.ballot.i64(i1 [[TMP2]]) +// CHECK-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8 +// CHECK-NEXT: store i64 [[TMP3]], ptr addrspace(4) [[TMP4]], align 8 // CHECK-NEXT: ret void // // AMDGCNSPIRV-LABEL: @_Z14test_uicmp_i64Pyyy( @@ -296,9 +297,10 @@ __global__ void endpgm() { // AMDGCNSPIRV-NEXT: store i64 [[B:%.*]], ptr addrspace(4) [[B_ADDR_ASCAST]], align 8 // AMDGCNSPIRV-NEXT: [[TMP0:%.*]] = load i64, ptr addrspace(4) [[A_ADDR_ASCAST]], align 8 // AMDGCNSPIRV-NEXT: [[TMP1:%.*]] = load i64, ptr addrspace(4) [[B_ADDR_ASCAST]], align 8 -// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = call addrspace(4) i64 @llvm.amdgcn.icmp.i64.i64(i64 [[TMP0]], i64 [[TMP1]], i32 35) -// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8 -// AMDGCNSPIRV-NEXT: store i64 [[TMP2]], ptr addrspace(4) [[TMP3]], align 8 +// AMDGCNSPIRV-NEXT: [[TMP2:%.*]] = icmp uge i64 [[TMP0]], [[TMP1]] +// AMDGCNSPIRV-NEXT: [[TMP3:%.*]] = call addrspace(4) i64 @llvm.amdgcn.ballot.i64(i1 [[TMP2]]) +// AMDGCNSPIRV-NEXT: [[TMP4:%.*]] = load ptr addrspace(4), ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8 +// AMDGCNSPIRV-NEXT: store i64 [[TMP3]], ptr addrspace(4) [[TMP4]], align 8 // AMDGCNSPIRV-NEXT: ret void // __global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b) diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl index 11fbfdde92fa1..0d00994af1d5d 100644 --- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl +++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl @@ -252,28 +252,34 @@ void test_fract_f64(global int* out, double a) } // CHECK-LABEL: @test_sicmp_i32 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32) +// CHECK: [[CMP:%.+]] = icmp eq i32 %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_sicmp_i32(global ulong* out, int a, int b) { *out = __builtin_amdgcn_sicmp(a, b, 32); } +// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1){{.*}} #[[$NOUNWIND_READONLY_NOPOISON:[0-9]+]] + // CHECK-LABEL: @test_uicmp_i32 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32) +// CHECK: [[CMP:%.+]] = icmp eq i32 %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_uicmp_i32(global ulong* out, uint a, uint b) { *out = __builtin_amdgcn_uicmp(a, b, 32); } // CHECK-LABEL: @test_sicmp_i64 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 38) +// CHECK: [[CMP:%.+]] = icmp sgt i64 %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_sicmp_i64(global ulong* out, long a, long b) { *out = __builtin_amdgcn_sicmpl(a, b, 39-1); } // CHECK-LABEL: @test_uicmp_i64 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i64(i64 %a, i64 %b, i32 35) +// CHECK: [[CMP:%.+]] = icmp uge i64 %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_uicmp_i64(global ulong* out, ulong a, ulong b) { *out = __builtin_amdgcn_uicmpl(a, b, 30+5); @@ -322,14 +328,16 @@ void test_wave_shuffle(global int* out, int a, int b) } // CHECK-LABEL: @test_fcmp_f32 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 5) +// CHECK: [[CMP:%.+]] = fcmp ole float %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_fcmp_f32(global ulong* out, float a, float b) { *out = __builtin_amdgcn_fcmpf(a, b, 5); } // CHECK-LABEL: @test_fcmp_f64 -// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f64(double %a, double %b, i32 6) +// CHECK: [[CMP:%.+]] = fcmp one double %a, %b +// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.ballot.i64(i1 [[CMP]]) void test_fcmp_f64(global ulong* out, double a, double b) { *out = __builtin_amdgcn_fcmp(a, b, 3+3); @@ -1039,8 +1047,6 @@ void test_read_exec(global ulong* out) { *out = __builtin_amdgcn_read_exec(); } -// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1){{.*}} #[[$NOUNWIND_READONLY_NOPOISON:[0-9]+]] - // CHECK-LABEL: @test_read_exec_lo( // CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ballot.i32(i1 true) void test_read_exec_lo(global uint* out) { _______________________________________________ cfe-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
