Skip to content

Commit 57b50ef

Browse files
authored
[AMDGPU] Add InstCombine rule for ballot.i64 intrinsic in wave32 mode. (#71556)
Substitute with zero-extended to i64 ballot.i32 intrinsic.
1 parent 7f7bbb9 commit 57b50ef

File tree

4 files changed

+20
-9
lines changed

4 files changed

+20
-9
lines changed

clang/test/CodeGenOpenCL/builtins-amdgcn-wave32.cl

Lines changed: 2 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -24,23 +24,19 @@ void test_ballot_wave32_target_attr(global uint* out, int a, int b)
2424
}
2525

2626
// CHECK-LABEL: @test_read_exec(
27-
// CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
27+
// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)
2828
void test_read_exec(global uint* out) {
2929
*out = __builtin_amdgcn_read_exec();
3030
}
3131

32-
// CHECK: declare i64 @llvm.amdgcn.ballot.i64(i1) #[[$NOUNWIND_READONLY:[0-9]+]]
33-
3432
// CHECK-LABEL: @test_read_exec_lo(
3533
// CHECK: call i32 @llvm.amdgcn.ballot.i32(i1 true)
3634
void test_read_exec_lo(global uint* out) {
3735
*out = __builtin_amdgcn_read_exec_lo();
3836
}
3937

4038
// CHECK-LABEL: @test_read_exec_hi(
41-
// CHECK: call i64 @llvm.amdgcn.ballot.i64(i1 true)
42-
// CHECK: lshr i64 [[A:%.*]], 32
43-
// CHECK: trunc i64 [[B:%.*]] to i32
39+
// CHECK: store i32 0, ptr addrspace(1) %out
4440
void test_read_exec_hi(global uint* out) {
4541
*out = __builtin_amdgcn_read_exec_hi();
4642
}

llvm/lib/Target/AMDGPU/AMDGPUISelDAGToDAG.cpp

Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -2382,7 +2382,7 @@ void AMDGPUDAGToDAGISel::SelectBRCOND(SDNode *N) {
23822382
auto CC = cast<CondCodeSDNode>(Cond->getOperand(2))->get();
23832383
if ((CC == ISD::SETEQ || CC == ISD::SETNE) &&
23842384
isNullConstant(Cond->getOperand(1)) &&
2385-
// TODO: make condition below an assert after fixing ballot bitwidth.
2385+
// We may encounter ballot.i64 in wave32 mode on -O0.
23862386
VCMP.getValueType().getSizeInBits() == ST->getWavefrontSize()) {
23872387
// %VCMP = i(WaveSize) AMDGPUISD::SETCC ...
23882388
// %C = i1 ISD::SETCC %VCMP, 0, setne/seteq

llvm/lib/Target/AMDGPU/AMDGPUInstCombineIntrinsic.cpp

Lines changed: 13 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -990,6 +990,19 @@ GCNTTIImpl::instCombineIntrinsic(InstCombiner &IC, IntrinsicInst &II) const {
990990
return IC.replaceInstUsesWith(II, Constant::getNullValue(II.getType()));
991991
}
992992
}
993+
if (ST->isWave32() && II.getType()->getIntegerBitWidth() == 64) {
994+
// %b64 = call i64 ballot.i64(...)
995+
// =>
996+
// %b32 = call i32 ballot.i32(...)
997+
// %b64 = zext i32 %b32 to i64
998+
Value *Call = IC.Builder.CreateZExt(
999+
IC.Builder.CreateIntrinsic(Intrinsic::amdgcn_ballot,
1000+
{IC.Builder.getInt32Ty()},
1001+
{II.getArgOperand(0)}),
1002+
II.getType());
1003+
Call->takeName(&II);
1004+
return IC.replaceInstUsesWith(II, Call);
1005+
}
9931006
break;
9941007
}
9951008
case Intrinsic::amdgcn_wqm_vote: {

llvm/test/Transforms/InstCombine/AMDGPU/amdgcn-intrinsics.ll

Lines changed: 4 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -2599,7 +2599,8 @@ declare i32 @llvm.amdgcn.ballot.i32(i1) nounwind readnone convergent
25992599

26002600
define i64 @ballot_nocombine_64(i1 %i) {
26012601
; CHECK-LABEL: @ballot_nocombine_64(
2602-
; CHECK-NEXT: [[B:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 [[I:%.*]])
2602+
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.ballot.i32(i1 [[I:%.*]])
2603+
; CHECK-NEXT: [[B:%.*]] = zext i32 [[TMP1]] to i64
26032604
; CHECK-NEXT: ret i64 [[B]]
26042605
;
26052606
%b = call i64 @llvm.amdgcn.ballot.i64(i1 %i)
@@ -2616,7 +2617,8 @@ define i64 @ballot_zero_64() {
26162617

26172618
define i64 @ballot_one_64() {
26182619
; CHECK-LABEL: @ballot_one_64(
2619-
; CHECK-NEXT: [[B:%.*]] = call i64 @llvm.amdgcn.ballot.i64(i1 true)
2620+
; CHECK-NEXT: [[TMP1:%.*]] = call i32 @llvm.amdgcn.ballot.i32(i1 true)
2621+
; CHECK-NEXT: [[B:%.*]] = zext i32 [[TMP1]] to i64
26202622
; CHECK-NEXT: ret i64 [[B]]
26212623
;
26222624
%b = call i64 @llvm.amdgcn.ballot.i64(i1 1)

0 commit comments

Comments
 (0)