[AMDGPU] Remove icmp and fcmp intrinsics and builtins#208183
Open
jayfoad wants to merge 1 commit into
Open
Conversation
Remove these intrinsics in favour of llvm.amdgcn.ballot: llvm.amdgcn.icmp llvm.amdgcn.fcmp Remove these builtins in favour of __builtin_amdgcn_ballot_w32 or __builtin_amdgcn_ballot_w64: __builtin_amdgcn_uicmp __builtin_amdgcn_uicmpl __builtin_amdgcn_sicmp __builtin_amdgcn_sicmpl __builtin_amdgcn_fcmp __builtin_amdgcn_fcmpf
jayfoad
requested review from
andykaylor,
bcardosolopes,
vangthao95 and
xlauko
as code owners
llvmorg-github-actions
Bot
added
backend:AMDGPU
clang:frontend
|
@llvm/pr-subscribers-clangir @llvm/pr-subscribers-clang-codegen Author: Jay Foad (jayfoad) ChangesRemove these intrinsics in favour of llvm.amdgcn.ballot: Remove these builtins in favour of __builtin_amdgcn_ballot_w32 or Patch is 513.80 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/208183.diff 38 Files Affected:
diff --git a/clang/include/clang/Basic/BuiltinsAMDGPU.td b/clang/include/clang/Basic/BuiltinsAMDGPU.td
index 306af993fd869..6e364490dbe4c 100644
--- a/clang/include/clang/Basic/BuiltinsAMDGPU.td
+++ b/clang/include/clang/Basic/BuiltinsAMDGPU.td
@@ -313,14 +313,6 @@ def __builtin_amdgcn_ballot_w64 : AMDGPUBuiltin<"uint64_t(bool)", [Const]>;
def __builtin_amdgcn_inverse_ballot_w32 : AMDGPUBuiltin<"bool(uint32_t)", [Const], "wavefrontsize32">;
def __builtin_amdgcn_inverse_ballot_w64 : AMDGPUBuiltin<"bool(uint64_t)", [Const], "wavefrontsize64">;
-// Deprecated intrinsics in favor of __builtin_amdgn_ballot_{w32|w64}
-def __builtin_amdgcn_uicmp : AMDGPUBuiltin<"uint64_t(unsigned int, unsigned int, _Constant int)", [Const]>;
-def __builtin_amdgcn_uicmpl : AMDGPUBuiltin<"uint64_t(uint64_t, uint64_t, _Constant int)", [Const]>;
-def __builtin_amdgcn_sicmp : AMDGPUBuiltin<"uint64_t(int, int, _Constant int)", [Const]>;
-def __builtin_amdgcn_sicmpl : AMDGPUBuiltin<"uint64_t(int64_t, int64_t, _Constant int)", [Const]>;
-def __builtin_amdgcn_fcmp : AMDGPUBuiltin<"uint64_t(double, double, _Constant int)", [Const]>;
-def __builtin_amdgcn_fcmpf : AMDGPUBuiltin<"uint64_t(float, float, _Constant int)", [Const]>;
-
//===----------------------------------------------------------------------===//
// Flat addressing builtins.
//===----------------------------------------------------------------------===//
diff --git a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
index 0a7ba0c194400..38a2797c6a1a4 100644
--- a/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
+++ b/clang/lib/CIR/CodeGen/CIRGenBuiltinAMDGPU.cpp
@@ -391,22 +391,6 @@ CIRGenFunction::emitAMDGPUBuiltinExpr(unsigned builtinId,
getContext().BuiltinInfo.getName(builtinId));
return mlir::Value{};
}
- 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_fcmp:
- case AMDGPU::BI__builtin_amdgcn_fcmpf: {
- cgm.errorNYI(expr->getSourceRange(),
- std::string("unimplemented AMDGPU builtin call: ") +
- getContext().BuiltinInfo.getName(builtinId));
- return mlir::Value{};
- }
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
case AMDGPU::BI__builtin_amdgcn_classh: {
diff --git a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
index 6a17f4c55f80d..19f36ec55c616 100644
--- a/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
+++ b/clang/lib/CodeGen/TargetBuiltins/AMDGPU.cpp
@@ -772,30 +772,6 @@ Value *CodeGenFunction::EmitAMDGPUBuiltinExpr(unsigned BuiltinID,
case AMDGPU::BI__builtin_amdgcn_tanh_bf16:
return emitBuiltinWithOneOverloadedType<1>(*this, E,
Intrinsic::amdgcn_tanh);
- 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_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));
-
- // 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 });
- }
case AMDGPU::BI__builtin_amdgcn_class:
case AMDGPU::BI__builtin_amdgcn_classf:
case AMDGPU::BI__builtin_amdgcn_classh:
diff --git a/clang/test/CodeGenCUDA/builtins-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-amdgcn.cu
index 35673773ec80c..9ad115d999b1c 100644
--- a/clang/test/CodeGenCUDA/builtins-amdgcn.cu
+++ b/clang/test/CodeGenCUDA/builtins-amdgcn.cu
@@ -164,34 +164,6 @@ __global__ void endpgm() {
__builtin_amdgcn_endpgm();
}
-// Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.
-
-// CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[OUT:%.*]] = alloca ptr, align 8, addrspace(5)
-// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr, align 8, addrspace(5)
-// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
-// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8, addrspace(5)
-// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[OUT_ADDR]] to ptr
-// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[A_ADDR]] to ptr
-// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr addrspace(5) [[B_ADDR]] to ptr
-// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr addrspace(5) [[OUT]], align 8
-// CHECK-NEXT: [[OUT1:%.*]] = load ptr, ptr addrspace(5) [[OUT]], align 8
-// CHECK-NEXT: store ptr [[OUT1]], ptr [[OUT_ADDR_ASCAST]], align 8
-// CHECK-NEXT: store i64 [[A:%.*]], ptr [[A_ADDR_ASCAST]], align 8
-// 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: ret void
-//
-__global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)
-{
- *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
-}
-
// Check the 64 bit return value is correctly returned without truncation or assertion.
// CHECK-LABEL: @_Z14test_s_memtimePy(
diff --git a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
index 94338f9027db1..1f6a8cda42c9d 100644
--- a/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
+++ b/clang/test/CodeGenCUDA/builtins-spirv-amdgcn.cu
@@ -257,55 +257,6 @@ __global__ void endpgm() {
__builtin_amdgcn_endpgm();
}
-// Check the 64 bit argument is correctly passed to the intrinsic without truncation or assertion.
-
-// CHECK-LABEL: @_Z14test_uicmp_i64Pyyy(
-// CHECK-NEXT: entry:
-// CHECK-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 8
-// CHECK-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 8
-// CHECK-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8
-// CHECK-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8
-// CHECK-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)
-// CHECK-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr [[A_ADDR]] to ptr addrspace(4)
-// CHECK-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr [[B_ADDR]] to ptr addrspace(4)
-// CHECK-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT]], align 8
-// CHECK-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr [[OUT]], align 8
-// CHECK-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8
-// CHECK-NEXT: store i64 [[A:%.*]], ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
-// 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: ret void
-//
-// AMDGCNSPIRV-LABEL: @_Z14test_uicmp_i64Pyyy(
-// AMDGCNSPIRV-NEXT: entry:
-// AMDGCNSPIRV-NEXT: [[OUT:%.*]] = alloca ptr addrspace(4), align 8
-// AMDGCNSPIRV-NEXT: [[OUT_ADDR:%.*]] = alloca ptr addrspace(4), align 8
-// AMDGCNSPIRV-NEXT: [[A_ADDR:%.*]] = alloca i64, align 8
-// AMDGCNSPIRV-NEXT: [[B_ADDR:%.*]] = alloca i64, align 8
-// AMDGCNSPIRV-NEXT: [[OUT_ADDR_ASCAST:%.*]] = addrspacecast ptr [[OUT_ADDR]] to ptr addrspace(4)
-// AMDGCNSPIRV-NEXT: [[A_ADDR_ASCAST:%.*]] = addrspacecast ptr [[A_ADDR]] to ptr addrspace(4)
-// AMDGCNSPIRV-NEXT: [[B_ADDR_ASCAST:%.*]] = addrspacecast ptr [[B_ADDR]] to ptr addrspace(4)
-// AMDGCNSPIRV-NEXT: store ptr addrspace(1) [[OUT_COERCE:%.*]], ptr [[OUT]], align 8
-// AMDGCNSPIRV-NEXT: [[OUT1:%.*]] = load ptr addrspace(4), ptr [[OUT]], align 8
-// AMDGCNSPIRV-NEXT: store ptr addrspace(4) [[OUT1]], ptr addrspace(4) [[OUT_ADDR_ASCAST]], align 8
-// AMDGCNSPIRV-NEXT: store i64 [[A:%.*]], ptr addrspace(4) [[A_ADDR_ASCAST]], align 8
-// 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: ret void
-//
-__global__ void test_uicmp_i64(unsigned long long *out, unsigned long long a, unsigned long long b)
-{
- *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
-}
-
// Check the 64 bit return value is correctly returned without truncation or assertion.
// CHECK-LABEL: @_Z14test_s_memtimePy(
diff --git a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
index 11fbfdde92fa1..84ebc14f42766 100644
--- a/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
+++ b/clang/test/CodeGenOpenCL/builtins-amdgcn.cl
@@ -251,34 +251,6 @@ void test_fract_f64(global int* out, double a)
*out = __builtin_amdgcn_fract(a);
}
-// CHECK-LABEL: @test_sicmp_i32
-// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32)
-void test_sicmp_i32(global ulong* out, int a, int b)
-{
- *out = __builtin_amdgcn_sicmp(a, b, 32);
-}
-
-// CHECK-LABEL: @test_uicmp_i32
-// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.icmp.i64.i32(i32 %a, i32 %b, i32 32)
-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)
-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)
-void test_uicmp_i64(global ulong* out, ulong a, ulong b)
-{
- *out = __builtin_amdgcn_uicmpl(a, b, 30+5);
-}
-
// CHECK-LABEL: @test_ds_swizzle
// CHECK: {{.*}}call{{.*}} i32 @llvm.amdgcn.ds.swizzle(i32 %a, i32 32)
void test_ds_swizzle(global int* out, int a)
@@ -321,20 +293,6 @@ void test_wave_shuffle(global int* out, int a, int b)
*out = __builtin_amdgcn_wave_shuffle(a, b);
}
-// CHECK-LABEL: @test_fcmp_f32
-// CHECK: {{.*}}call{{.*}} i64 @llvm.amdgcn.fcmp.i64.f32(float %a, float %b, i32 5)
-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)
-void test_fcmp_f64(global ulong* out, double a, double b)
-{
- *out = __builtin_amdgcn_fcmp(a, b, 3+3);
-}
-
// CHECK-LABEL: @test_class_f32
// CHECK: {{.*}}call{{.*}} i1 @llvm.amdgcn.class.f32
void test_class_f32(global float* out, float a, int b)
diff --git a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
index eb1a86bdcdeb0..a19d7ec85b1ee 100644
--- a/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
+++ b/clang/test/SemaOpenCL/builtins-amdgcn-error.cl
@@ -77,36 +77,6 @@ void test_iglp_opt(int x)
__builtin_amdgcn_iglp_opt(x); // expected-error {{argument to '__builtin_amdgcn_iglp_opt' must be a constant integer}}
}
-void test_sicmp_i32(global ulong* out, int a, int b, uint c)
-{
- *out = __builtin_amdgcn_sicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmp' must be a constant integer}}
-}
-
-void test_uicmp_i32(global ulong* out, uint a, uint b, uint c)
-{
- *out = __builtin_amdgcn_uicmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmp' must be a constant integer}}
-}
-
-void test_sicmp_i64(global ulong* out, long a, long b, uint c)
-{
- *out = __builtin_amdgcn_sicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_sicmpl' must be a constant integer}}
-}
-
-void test_uicmp_i64(global ulong* out, ulong a, ulong b, uint c)
-{
- *out = __builtin_amdgcn_uicmpl(a, b, c); // expected-error {{argument to '__builtin_amdgcn_uicmpl' must be a constant integer}}
-}
-
-void test_fcmp_f32(global ulong* out, float a, float b, uint c)
-{
- *out = __builtin_amdgcn_fcmpf(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmpf' must be a constant integer}}
-}
-
-void test_fcmp_f64(global ulong* out, double a, double b, uint c)
-{
- *out = __builtin_amdgcn_fcmp(a, b, c); // expected-error {{argument to '__builtin_amdgcn_fcmp' must be a constant integer}}
-}
-
void test_ds_swizzle(global int* out, int a, int b)
{
*out = __builtin_amdgcn_ds_swizzle(a, b); // expected-error {{argument to '__builtin_amdgcn_ds_swizzle' must be a constant integer}}
diff --git a/llvm/docs/ReleaseNotes.md b/llvm/docs/ReleaseNotes.md
index 3ca9f5166ad95..58f7062f89130 100644
--- a/llvm/docs/ReleaseNotes.md
+++ b/llvm/docs/ReleaseNotes.md
@@ -218,6 +218,16 @@ Makes programs 10x faster by doing Special New Thing.
previously relied on the subtarget feature to enable misaligned buffer merging
must now set the corresponding module flag to `1` (relaxed). An absent flag is
treated as strict by the backend.
+* These intrinsics have been removed in favour of `llvm.amdgcn.ballot`:
+ * `llvm.amdgcn.icmp`
+ * `llvm.amdgcn.fcmp`
+* The corresponding builtins have been removed in favour of
+ `__builtin_amdgcn_ballot_w32` or `__builtin_amdgcn_ballot_w64`:
+ * `__builtin_amdgcn_uicmp`
+ * `__builtin_amdgcn_uicmpl`
+ * `__builtin_amdgcn_sicmpl`
+ * `__builtin_amdgcn_fcmp`
+ * `__builtin_amdgcn_fcmpf`
### Changes to the ARM Backend
diff --git a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
index 21882e247c027..c02d485cc9273 100644
--- a/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
+++ b/llvm/include/llvm/IR/IntrinsicsAMDGPU.td
@@ -2449,16 +2449,6 @@ def int_amdgcn_cvt_pk_u8_f32 :
ClangBuiltin<"__builtin_amdgcn_cvt_pk_u8_f32">,
PureIntrinsic<[llvm_i32_ty], [llvm_float_ty, llvm_i32_ty, llvm_i32_ty]>;
-def int_amdgcn_icmp :
- Intrinsic<[llvm_anyint_ty], [llvm_anyint_ty, LLVMMatchType<1>, llvm_i32_ty],
- [IntrNoMem, IntrConvergent,
- ImmArg<ArgIndex<2>>, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
-
-def int_amdgcn_fcmp :
- Intrinsic<[llvm_anyint_ty], [llvm_anyfloat_ty, LLVMMatchType<1>, llvm_i32_ty],
- [IntrNoMem, IntrConvergent,
- ImmArg<ArgIndex<2>>, IntrWillReturn, IntrNoCallback, IntrNoFree]>;
-
// Returns a bitfield(i32 or i64) containing the result of its i1 argument
// in all active lanes, and zero in all inactive lanes.
def int_amdgcn_ballot :
diff --git a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
index f5d8e43dce9db..eb4efac76fd11 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp
@@ -1661,141 +1661,6 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
break;
}
- case Intrinsic::amdgcn_icmp:
- case Intrinsic::amdgcn_fcmp: {
- const ConstantInt *CC = cast<ConstantInt>(II.getArgOperand(2));
- // Guard against invalid arguments.
- int64_t CCVal = CC->getZExtValue();
- bool IsInteger = IID == Intrinsic::amdgcn_icmp;
- if ((IsInteger && (CCVal < CmpInst::FIRST_ICMP_PREDICATE ||
- CCVal > CmpInst::LAST_ICMP_PREDICATE)) ||
- (!IsInteger && (CCVal < CmpInst::FIRST_FCMP_PREDICATE ||
- CCVal > CmpInst::LAST_FCMP_PREDICATE)))
- break;
-
- Value *Src0 = II.getArgOperand(0);
- Value *Src1 = II.getArgOperand(1);
-
- if (auto *CSrc0 = dyn_cast<Constant>(Src0)) {
- if (auto *CSrc1 = dyn_cast<Constant>(Src1)) {
- Constant *CCmp = ConstantFoldCompareInstOperands(
- (ICmpInst::Predicate)CCVal, CSrc0, CSrc1, DL);
- if (CCmp && CCmp->isNullValue()) {
- return IC.replaceInstUsesWith(
- II, IC.Builder.CreateSExt(CCmp, II.getType()));
- }
-
- // The result of V_ICMP/V_FCMP assembly instructions (which this
- // intrinsic exposes) is one bit per thread, masked with the EXEC
- // register (which contains the bitmask of live threads). So a
- // comparison that always returns true is the same as a read of the
- // EXEC register. ballot(true) reads EXEC at the wave-size width, so
- // zext/trunc the result to the intrinsic's return type.
- Type *WaveTy = IC.Builder.getIntNTy(ST->getWavefrontSize());
- Value *Ballot = IC.Builder.CreateIntrinsic(
- Intrinsic::amdgcn_ballot, WaveTy, IC.Builder.getTrue());
- Value *Result = IC.Builder.CreateZExtOrTrunc(Ballot, II.getType());
- return IC.replaceInstUsesWith(II, Result);
- }
-
- // Canonicalize constants to RHS.
- CmpInst::Predicate SwapPred =
- CmpInst::getSwappedPredicate(static_cast<CmpInst::Predicate>(CCVal));
- II.setArgOperand(0, Src1);
- II.setArgOperand(1, Src0);
- II.setArgOperand(
- 2, ConstantInt::get(CC->getType(), static_cast<int>(SwapPred)));
- return &II;
- }
-
- if (CCVal != CmpInst::ICMP_EQ && CCVal != CmpInst::ICMP_NE)
- break;
-
- // Canonicalize compare eq with true value to compare != 0
- // llvm.amdgcn.icmp(zext (i1 x), 1, eq)
- // -> llvm.amdgcn.icmp(zext (i1 x), 0, ne)
- // llvm.amdgcn.icmp(sext (i1 x), -1, eq)
- // -> llvm.amdgcn.icmp(sext (i1 x), 0, ne)
- Value *ExtSrc;
- if (CCVal == CmpInst::ICMP_EQ &&
- ((match(Src1, PatternMatch::m_One()) &&
- match(Src0, m_ZExt(PatternMatch::m_Value(ExtSrc)))) ||
- (match(Src1, PatternMatch::m_AllOnes()) &&
- match(Src0, m_SExt(PatternMatch::m_Value(ExtSrc))))) &&
- ExtSrc->getType()->isIntegerTy(1)) {
- IC.replaceOperand(II, 1, ConstantInt::getNullValue(Src1->getType()));
- IC.replaceOperand(II, 2,
- ConstantInt::get(CC->getType(), CmpInst::ICMP_NE));
- return &II;
- }
-
- CmpPredicate SrcPred;
- Value *SrcLHS;
- Value *SrcRHS;
-
- // Fold compare eq/ne with 0 from a compare result as the predicate to the
- // intrinsic. The typical use is a wave vote function in the library, which
- // will be fed from a user code condition compared with 0. Fold in the
- // redundant compare.
-
- // llvm.amdgcn.icmp([sz]ext ([if]cmp pred a, b), 0, ne)
- // -> llvm.amdgcn.[if]cmp(a, b, pred)
- //
- // llvm.amdgcn.icmp([sz]ext ([if]cmp pred a, b), 0, eq)
- // -> llvm.amdgcn.[if]cmp(a, b, inv pred)
- if (match(Src1, PatternMatch::m_Zero()) &&
- match(Src0, PatternMatch::m_ZExtOrSExt(
- m_Cmp(SrcPred, PatternMatch::m_Value(SrcLHS),
- PatternMat...
[truncated]
|
Contributor
Author
|
I think these are still used by Mesa but I don't know how to get those uses fixed, other than landing this so that they are forced to react. |
Contributor
Author
@marekolsak @peppsac FYI |
arsenm
reviewed
Jul 8, 2026
arsenm
left a comment
arsenm
left a comment
Contributor
There was a problem hiding this comment.
Shouldn't just rip out the builtins, should reimplement with the ballot intrinsic. Should also have bitcode auto upgrade
|
Are ballots implicitly prevented from being moved into a predecessor or successor basic blocks with a different EXEC mask, and are CFG restructuring passes tamed around ballot to make sure they don't affect the EXEC mask for it? |
Contributor
Author
Yes. The ballot intrinsic is marked as convergent, which should prevent that kind of thing. Of course there could still be bugs, but hopefully they have been flushed out by now because LLPC already uses ballot. |
Contributor
Author
So then they should not be deprecated?
|
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
3 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Remove these intrinsics in favour of llvm.amdgcn.ballot:
llvm.amdgcn.icmp
llvm.amdgcn.fcmp
Remove these builtins in favour of __builtin_amdgcn_ballot_w32 or
__builtin_amdgcn_ballot_w64:
__builtin_amdgcn_uicmp
__builtin_amdgcn_uicmpl
__builtin_amdgcn_sicmp
__builtin_amdgcn_sicmpl
__builtin_amdgcn_fcmp
__builtin_amdgcn_fcmpf