[llvm] AMDGPU/GlobalISel: RegBankLegalize rules for s_bitreplicate (PR #189138)
via llvm-commits
llvm-commits at lists.llvm.org
Tue Jun 30 20:28:28 PDT 2026
https://github.com/vangthao95 updated https://github.com/llvm/llvm-project/pull/189138
>From fa81c2426255a66b90afbfe7bd175e47cfde912d Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Fri, 27 Mar 2026 21:58:22 -0400
Subject: [PATCH 1/7] AMDGPU/GlobalISel: RegBankLegalize rules for
s_bitreplicate
---
llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp | 4 ++++
llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll | 2 +-
2 files changed, 5 insertions(+), 1 deletion(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 7f7225ad6d311..85f36f44d467b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1762,6 +1762,10 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({returnaddress}).Any({{UniP0}, {{SgprP0}, {}}});
+ addRulesForIOpcs({amdgcn_s_bitreplicate}, Standard)
+ .Uni(S64, {{Sgpr64}, {IntrId, Sgpr32}})
+ .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, SgprB32_ReadFirstLane}});
+
// Note: amdgcn.icmp with i1 inputs is legalized to ballot in the legalizer,
// so no S1 rules are needed here.
addRulesForIOpcs({amdgcn_icmp})
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index 041399c66418e..631fdc7406918 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -1,5 +1,5 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
-; xUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=1 < %s | FileCheck -check-prefixes=GFX11 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=1 < %s | FileCheck -check-prefixes=GFX11 %s
; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=0 < %s | FileCheck -check-prefixes=GFX11 %s
declare i64 @llvm.amdgcn.s.bitreplicate(i32)
>From 8787abdafeb4b7d45b408d691adc69ce909f4f86 Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Tue, 31 Mar 2026 14:12:21 -0400
Subject: [PATCH 2/7] Change to WF and fix WF for SALUs with a copy to VGPRs
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 16 +++++++--
.../AMDGPU/AMDGPURegBankLegalizeHelper.h | 3 ++
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 2 +-
.../AMDGPU/llvm.amdgcn.bitreplicate.ll | 36 ++++++++++++++-----
4 files changed, 45 insertions(+), 12 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index b9554d7858d32..da249c5267754 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -66,8 +66,9 @@ bool RegBankLegalizeHelper::findRuleAndApplyMapping(MachineInstr &MI) {
WaterfallInfo WFI;
unsigned OpIdx = 0;
+ OldNextMI = std::next(MI.getIterator());
if (!Mapping->DstOpMapping.empty()) {
- B.setInsertPt(*MI.getParent(), std::next(MI.getIterator()));
+ B.setInsertPt(*MI.getParent(), OldNextMI);
if (!applyMappingDst(MI, OpIdx, Mapping->DstOpMapping))
return false;
}
@@ -2431,8 +2432,19 @@ bool RegBankLegalizeHelper::applyMappingSrc(
if (RB != SgprRB) {
WFI.SgprWaterfallOperandRegs.insert(Reg);
if (!WFI.Start.isValid()) {
+ // Waterfall range [WFI.Start, WFI.End). Use OldNextMI so that
+ // any instructions inserted by applyMappingDst are included.
WFI.Start = MI.getIterator();
- WFI.End = std::next(MI.getIterator());
+ WFI.End = OldNextMI;
+
+ // Mark any COPY as exec-dependent since machine-sink may move
+ // it out of the loop body.
+ MCRegister ExecReg = ST.getRegisterInfo()->getExec();
+ for (auto It = std::next(MI.getIterator()); It != OldNextMI; ++It) {
+ if (It->isCopy())
+ It->addOperand(MachineOperand::CreateReg(ExecReg, /*isDef=*/false,
+ /*isImp=*/true));
+ }
}
}
break;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
index 25c8c3a0f6127..ae96d4de99224 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
@@ -52,6 +52,9 @@ class RegBankLegalizeHelper {
MachineOptimizationRemarkEmitter MORE;
const RegBankLegalizeRules &RBLRules;
const bool IsWave32;
+ // The original next instruction after MI, saved before applyMappingDst may
+ // insert instructions. Used by Sgpr*_WF to set the waterfall range.
+ MachineBasicBlock::iterator OldNextMI;
const RegisterBank *SgprRB;
const RegisterBank *VgprRB;
const RegisterBank *AgprRB;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 85f36f44d467b..fc66d967dbe7b 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1764,7 +1764,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({amdgcn_s_bitreplicate}, Standard)
.Uni(S64, {{Sgpr64}, {IntrId, Sgpr32}})
- .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, SgprB32_ReadFirstLane}});
+ .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, Sgpr32_WF}});
// Note: amdgcn.icmp with i1 inputs is legalized to ballot in the legalizer,
// so no S1 rules are needed here.
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index 631fdc7406918..af3c84801c0d9 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -1,6 +1,6 @@
; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 2
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=1 < %s | FileCheck -check-prefixes=GFX11 %s
-; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=0 < %s | FileCheck -check-prefixes=GFX11 %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=1 < %s | FileCheck -check-prefixes=GFX11,GFX11-GISEL %s
+; RUN: llc -mtriple=amdgcn -mcpu=gfx1100 -amdgpu-enable-delay-alu=0 -global-isel=0 < %s | FileCheck -check-prefixes=GFX11,GFX11-SDAG %s
declare i64 @llvm.amdgcn.s.bitreplicate(i32)
@@ -76,13 +76,31 @@ entry:
}
define i64 @test_s_bitreplicate_vgpr(i32 %mask) {
-; GFX11-LABEL: test_s_bitreplicate_vgpr:
-; GFX11: ; %bb.0: ; %entry
-; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
-; GFX11-NEXT: s_setpc_b64 s[30:31]
+; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr:
+; GFX11-GISEL: ; %bb.0: ; %entry
+; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v2, v0
+; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
+; GFX11-GISEL-NEXT: .LBB6_1: ; =>This Inner Loop Header: Depth=1
+; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v2
+; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
+; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v2
+; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr2
+; GFX11-GISEL-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
+; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
+; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB6_1
+; GFX11-GISEL-NEXT: ; %bb.2:
+; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
+; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr:
+; GFX11-SDAG: ; %bb.0: ; %entry
+; GFX11-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-SDAG-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-SDAG-NEXT: s_setpc_b64 s[30:31]
entry:
%br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
ret i64 %br
>From 198133aa7709cd24dc2f5e98ebf51545787b3d8c Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Wed, 1 Apr 2026 18:22:11 -0400
Subject: [PATCH 3/7] Add PHI tied operand
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 42 +++++++---
.../AMDGPU/GlobalISel/regbankselect-call.ll | 4 +-
.../regbankselect-waterfall-call.mir | 16 +++-
.../AMDGPU/llvm.amdgcn.bitreplicate.ll | 81 ++++++++++++++++++-
4 files changed, 126 insertions(+), 17 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index da249c5267754..a55387a08d7b1 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -201,6 +201,39 @@ bool RegBankLegalizeHelper::executeInWaterfallLoop(MachineIRBuilder &B,
auto NewEnd = BodyBB->end();
assert(std::distance(NewBegin, NewEnd) == OrigRangeSize);
+ // Create loop-carried dependencies for VGPR defs that lack inherent exec
+ // dependency, preventing machine-sink from moving them out of the loop.
+ for (MachineInstr &MI : make_range(NewBegin, NewEnd)) {
+ if (MI.mayLoadOrStore() || MI.hasUnmodeledSideEffects())
+ continue;
+ for (unsigned I = 0, E = MI.getNumDefs(); I < E; ++I) {
+ if (!MI.getOperand(I).isReg())
+ continue;
+ Register DefReg = MI.getOperand(I).getReg();
+ if (!DefReg.isVirtual() || MRI.getRegBank(DefReg) != VgprRB)
+ continue;
+
+ LLT Ty = MRI.getType(DefReg);
+
+ B.setInsertPt(MBB, MBB.end());
+ Register InitReg = MRI.createVirtualRegister({VgprRB, Ty});
+ B.buildInstr(TargetOpcode::IMPLICIT_DEF).addDef(InitReg);
+
+ Register PhiReg = MRI.createVirtualRegister({VgprRB, Ty});
+ B.setInsertPt(*LoopBB, LoopBB->begin());
+ B.buildInstr(TargetOpcode::G_PHI)
+ .addDef(PhiReg)
+ .addReg(InitReg)
+ .addMBB(&MBB)
+ .addReg(DefReg)
+ .addMBB(BodyBB);
+
+ MI.addOperand(MachineOperand::CreateReg(PhiReg, /*isDef=*/false,
+ /*isImp=*/true));
+ MI.tieOperands(I, MI.getNumOperands() - 1);
+ }
+ }
+
B.setMBB(*LoopBB);
Register CondReg;
@@ -2436,15 +2469,6 @@ bool RegBankLegalizeHelper::applyMappingSrc(
// any instructions inserted by applyMappingDst are included.
WFI.Start = MI.getIterator();
WFI.End = OldNextMI;
-
- // Mark any COPY as exec-dependent since machine-sink may move
- // it out of the loop body.
- MCRegister ExecReg = ST.getRegisterInfo()->getExec();
- for (auto It = std::next(MI.getIterator()); It != OldNextMI; ++It) {
- if (It->isCopy())
- It->addOperand(MachineOperand::CreateReg(ExecReg, /*isDef=*/false,
- /*isImp=*/true));
- }
}
}
break;
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
index b35907de70b38..66677fa5e542f 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
@@ -76,11 +76,13 @@ define amdgpu_ps i32 @test_divergent_indirect_call_p0_with_args(ptr %fptr, i32 %
; CHECK-NEXT: [[MV:%[0-9]+]]:vgpr(p0) = G_MERGE_VALUES [[COPY]](s32), [[COPY1]](s32)
; CHECK-NEXT: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr2
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(s32) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(s32) = G_PHI [[DEF1]](s32), %bb.1, %5(s32), %bb.3
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES [[MV]](p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -97,7 +99,7 @@ define amdgpu_ps i32 @test_divergent_indirect_call_p0_with_args(ptr %fptr, i32 %
; CHECK-NEXT: [[COPY3:%[0-9]+]]:sgpr(<4 x s32>) = COPY $private_rsrc_reg
; CHECK-NEXT: $sgpr0_sgpr1_sgpr2_sgpr3 = COPY [[COPY3]](<4 x s32>)
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[MV1]](p0), 0, csr_amdgpu_si_gfx_gfx90ainsts, implicit $vgpr0, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit-def $vgpr0
- ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vgpr(s32) = COPY $vgpr0, implicit [[PHI]](tied-def 0)(s32)
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: $exec = S_XOR_B64_term $exec, [[S_AND_SAVEEXEC_B64_]], implicit-def $scc
; CHECK-NEXT: SI_WATERFALL_LOOP %bb.2, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
index 5207d992ea74d..934f6f72452ad 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
@@ -13,11 +13,13 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p0) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p0) = G_PHI [[DEF1]](p0), %bb.0, %2(p0), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -31,7 +33,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p0) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0), implicit [[PHI]](tied-def 0)(p0)
; CHECK-NEXT: %func_ptr:vgpr(p0) = G_LOAD [[COPY]](p0) :: (load (p0))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p0), 0, csr_amdgpu
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -66,11 +68,13 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p4) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p4) = G_PHI [[DEF1]](p4), %bb.0, %2(p4), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p4)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -84,7 +88,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p4) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4), implicit [[PHI]](tied-def 0)(p4)
; CHECK-NEXT: %func_ptr:vgpr(p4) = G_LOAD [[COPY]](p4) :: (load (p4))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p4), 0, csr_amdgpu
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -119,11 +123,13 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p0) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p0) = G_PHI [[DEF1]](p0), %bb.0, %2(p0), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -137,7 +143,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p0) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0), implicit [[PHI]](tied-def 0)(p0)
; CHECK-NEXT: %func_ptr:vgpr(p0) = G_LOAD [[COPY]](p0) :: (load (p0))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p0), 0, csr_amdgpu, implicit $sgpr4, implicit $sgpr5, implicit-def $vgpr0
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -172,11 +178,13 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
+ ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p4) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
+ ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p4) = G_PHI [[DEF1]](p4), %bb.0, %2(p4), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p4)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -190,7 +198,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p4) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4), implicit [[PHI]](tied-def 0)(p4)
; CHECK-NEXT: %func_ptr:vgpr(p4) = G_LOAD [[COPY]](p4) :: (load (p4))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p4), 0, csr_amdgpu, implicit $sgpr4, implicit $sgpr5, implicit-def $vgpr0
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index af3c84801c0d9..121212a1dc1b1 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -75,21 +75,96 @@ entry:
ret void
}
+define amdgpu_cs void @test_s_bitreplicate_vgpr_store(i32 %mask, ptr addrspace(1) %out) {
+; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_store:
+; GFX11-GISEL: ; %bb.0: ; %entry
+; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr3_vgpr4
+; GFX11-GISEL-NEXT: .LBB6_1: ; =>This Inner Loop Header: Depth=1
+; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
+; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v0
+; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v4, s3
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v3, s2
+; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
+; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB6_1
+; GFX11-GISEL-NEXT: ; %bb.2:
+; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
+; GFX11-GISEL-NEXT: global_store_b64 v[1:2], v[3:4], off
+; GFX11-GISEL-NEXT: s_endpgm
+;
+; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr_store:
+; GFX11-SDAG: ; %bb.0: ; %entry
+; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-SDAG-NEXT: v_dual_mov_b32 v4, s1 :: v_dual_mov_b32 v3, s0
+; GFX11-SDAG-NEXT: global_store_b64 v[1:2], v[3:4], off
+; GFX11-SDAG-NEXT: s_endpgm
+entry:
+ %br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
+ store i64 %br, ptr addrspace(1) %out
+ ret void
+}
+
+define i64 @test_s_bitreplicate_vgpr_multi_use(i32 %mask) {
+; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_multi_use:
+; GFX11-GISEL: ; %bb.0: ; %entry
+; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr1_vgpr2
+; GFX11-GISEL-NEXT: .LBB7_1: ; =>This Inner Loop Header: Depth=1
+; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v0
+; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
+; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v0
+; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, s2
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v2, s3
+; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
+; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB7_1
+; GFX11-GISEL-NEXT: ; %bb.2:
+; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
+; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v1, v2
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr_multi_use:
+; GFX11-SDAG: ; %bb.0: ; %entry
+; GFX11-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-SDAG-NEXT: v_mov_b32_e32 v1, 0
+; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-SDAG-NEXT: v_add_nc_u32_e64 v0, s0, s1
+; GFX11-SDAG-NEXT: s_setpc_b64 s[30:31]
+entry:
+ %br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
+ %lo = trunc i64 %br to i32
+ %hi = lshr i64 %br, 32
+ %hi32 = trunc i64 %hi to i32
+ %sum = add i32 %lo, %hi32
+ %result = zext i32 %sum to i64
+ ret i64 %result
+}
+
define i64 @test_s_bitreplicate_vgpr(i32 %mask) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr:
; GFX11-GISEL: ; %bb.0: ; %entry
; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
; GFX11-GISEL-NEXT: v_mov_b32_e32 v2, v0
; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
-; GFX11-GISEL-NEXT: .LBB6_1: ; =>This Inner Loop Header: Depth=1
+; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0_vgpr1
+; GFX11-GISEL-NEXT: .LBB8_1: ; =>This Inner Loop Header: Depth=1
; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v2
; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v2
; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
; GFX11-GISEL-NEXT: ; implicit-def: $vgpr2
-; GFX11-GISEL-NEXT: v_dual_mov_b32 v0, s2 :: v_dual_mov_b32 v1, s3
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v0, s2
+; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, s3
; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
-; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB6_1
+; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB8_1
; GFX11-GISEL-NEXT: ; %bb.2:
; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
>From 55158bc98d585efa59971669b205734a94264e89 Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Mon, 6 Apr 2026 17:01:08 -0400
Subject: [PATCH 4/7] Remove waterfall, expand divergent bitreplicate
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 116 ++++++++++++------
.../AMDGPU/AMDGPURegBankLegalizeHelper.h | 4 +-
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 2 +-
.../AMDGPU/AMDGPURegBankLegalizeRules.h | 3 +-
.../AMDGPU/GlobalISel/regbankselect-call.ll | 4 +-
.../regbankselect-waterfall-call.mir | 16 +--
.../AMDGPU/llvm.amdgcn.bitreplicate.ll | 93 +++++++-------
7 files changed, 136 insertions(+), 102 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index a55387a08d7b1..f011d59089f68 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -66,9 +66,8 @@ bool RegBankLegalizeHelper::findRuleAndApplyMapping(MachineInstr &MI) {
WaterfallInfo WFI;
unsigned OpIdx = 0;
- OldNextMI = std::next(MI.getIterator());
if (!Mapping->DstOpMapping.empty()) {
- B.setInsertPt(*MI.getParent(), OldNextMI);
+ B.setInsertPt(*MI.getParent(), std::next(MI.getIterator()));
if (!applyMappingDst(MI, OpIdx, Mapping->DstOpMapping))
return false;
}
@@ -201,39 +200,6 @@ bool RegBankLegalizeHelper::executeInWaterfallLoop(MachineIRBuilder &B,
auto NewEnd = BodyBB->end();
assert(std::distance(NewBegin, NewEnd) == OrigRangeSize);
- // Create loop-carried dependencies for VGPR defs that lack inherent exec
- // dependency, preventing machine-sink from moving them out of the loop.
- for (MachineInstr &MI : make_range(NewBegin, NewEnd)) {
- if (MI.mayLoadOrStore() || MI.hasUnmodeledSideEffects())
- continue;
- for (unsigned I = 0, E = MI.getNumDefs(); I < E; ++I) {
- if (!MI.getOperand(I).isReg())
- continue;
- Register DefReg = MI.getOperand(I).getReg();
- if (!DefReg.isVirtual() || MRI.getRegBank(DefReg) != VgprRB)
- continue;
-
- LLT Ty = MRI.getType(DefReg);
-
- B.setInsertPt(MBB, MBB.end());
- Register InitReg = MRI.createVirtualRegister({VgprRB, Ty});
- B.buildInstr(TargetOpcode::IMPLICIT_DEF).addDef(InitReg);
-
- Register PhiReg = MRI.createVirtualRegister({VgprRB, Ty});
- B.setInsertPt(*LoopBB, LoopBB->begin());
- B.buildInstr(TargetOpcode::G_PHI)
- .addDef(PhiReg)
- .addReg(InitReg)
- .addMBB(&MBB)
- .addReg(DefReg)
- .addMBB(BodyBB);
-
- MI.addOperand(MachineOperand::CreateReg(PhiReg, /*isDef=*/false,
- /*isImp=*/true));
- MI.tieOperands(I, MI.getNumOperands() - 1);
- }
- }
-
B.setMBB(*LoopBB);
Register CondReg;
@@ -1445,6 +1411,80 @@ bool RegBankLegalizeHelper::lowerGetRounding(MachineInstr &MI) {
return true;
}
+bool RegBankLegalizeHelper::lowerBitReplicateToVALU(MachineInstr &MI) {
+ // Lower divergent s_bitreplicate (bit i -> bits 2i, 2i+1) to VALU by
+ // splitting into two 16-bit halves and spreading bits via:
+ // lo = v_perm_b32(0, src, 0x0C010C00) // [00][B1][00][B0]
+ // hi = v_perm_b32(0, src, 0x0C030C02) // [00][B3][00][B2]
+ // M4 = (x | (x << 4)) & 0x0F0F0F0F // 4-bit spread
+ // M2 = (M4 | (M4 << 2)) & 0x33333333 // 2-bit spread
+ // M1 = (M2 | (M2 << 1)) & 0x55555555 // 1-bit spread
+ // R = M1 | (M1 << 1) // duplicate each bit
+ Register Dst = MI.getOperand(0).getReg();
+ Register Src = MI.getOperand(2).getReg();
+
+ // v_perm_b32 splits the input into 16-bit halves and spreads each half's
+ // bytes with zero-byte gaps in one instruction. Selector byte values:
+ // 0x00-0x03 = src byte N, 0x0C = zero byte
+ // e.g. 0x134C3A92 -> Hi: [0013][004C], Lo: [003A][0092]
+ auto Zero = B.buildConstant(VgprRB_S32, 0);
+ Register LoByteSpread =
+ B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(B.buildConstant(VgprRB_S32, 0x0C010C00).getReg(0))
+ .getReg(0);
+ Register HiByteSpread =
+ B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(B.buildConstant(VgprRB_S32, 0x0C030C02).getReg(0))
+ .getReg(0);
+
+ // Spread masks: each keeps the low N bits of every 2N-bit group, clearing
+ // the garbage left by the shift+OR.
+ auto Mask4 =
+ B.buildConstant(VgprRB_S32, 0x0F0F0F0F); // 0000_1111: low 4 bits per byte
+ auto Mask2 = B.buildConstant(
+ VgprRB_S32, 0x33333333); // 0011_0011: low 2 bits per 4-bit group
+ auto Mask1 =
+ B.buildConstant(VgprRB_S32, 0x55555555); // 0101_0101: every other bit
+
+ auto Spread = [&](Register In) -> Register {
+ // 4-bit spread: separate each byte into its two 4-bit halves.
+ // e.g. [003A][0092] -> [03][0A][09][02]
+ auto S4 = B.buildShl(VgprRB_S32, In, B.buildConstant(VgprRB_S32, 4));
+ auto Or4 = B.buildOr(VgprRB_S32, In, S4);
+ auto M4 = B.buildAnd(VgprRB_S32, Or4, Mask4);
+
+ // 2-bit spread: separate each 4-bit group into two 2-bit pairs.
+ // e.g. [03][0A][09][02] -> [0000][0011][0010][0010][0010][0001][0000][0010]
+ auto S2 = B.buildShl(VgprRB_S32, M4, B.buildConstant(VgprRB_S32, 2));
+ auto Or2 = B.buildOr(VgprRB_S32, M4, S2);
+ auto M2 = B.buildAnd(VgprRB_S32, Or2, Mask2);
+
+ // 1-bit spread: separate each 2-bit pair into individual bits.
+ // e.g. -> [00][00][01][01][01][00][01][00][01][00][00][01][00][00][01][00]
+ auto S1 = B.buildShl(VgprRB_S32, M2, B.buildConstant(VgprRB_S32, 1));
+ auto Or1 = B.buildOr(VgprRB_S32, M2, S1);
+ auto M1 = B.buildAnd(VgprRB_S32, Or1, Mask1);
+
+ // Duplicate: double each isolated bit.
+ // e.g. -> [00][00][11][11][11][00][11][00][11][00][00][11][00][00][11][00]
+ auto Dup = B.buildShl(VgprRB_S32, M1, B.buildConstant(VgprRB_S32, 1));
+ auto Result = B.buildOr(VgprRB_S32, M1, Dup);
+
+ // e.g. -> return 0x0FCCC30C.
+ return Result.getReg(0);
+ };
+
+ Register LoResult = Spread(LoByteSpread);
+ Register HiResult = Spread(HiByteSpread);
+ B.buildMergeLikeInstr(Dst, {LoResult, HiResult});
+ MI.eraseFromParent();
+ return true;
+}
+
bool RegBankLegalizeHelper::lower(MachineInstr &MI,
const RegBankLLTMapping &Mapping,
WaterfallInfo &WFI) {
@@ -1825,6 +1865,8 @@ bool RegBankLegalizeHelper::lower(MachineInstr &MI,
return lowerSetRounding(MI);
case LowerGetRounding:
return lowerGetRounding(MI);
+ case BitReplicateToVALU:
+ return lowerBitReplicateToVALU(MI);
}
return true;
@@ -2465,10 +2507,8 @@ bool RegBankLegalizeHelper::applyMappingSrc(
if (RB != SgprRB) {
WFI.SgprWaterfallOperandRegs.insert(Reg);
if (!WFI.Start.isValid()) {
- // Waterfall range [WFI.Start, WFI.End). Use OldNextMI so that
- // any instructions inserted by applyMappingDst are included.
WFI.Start = MI.getIterator();
- WFI.End = OldNextMI;
+ WFI.End = std::next(MI.getIterator());
}
}
break;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
index ae96d4de99224..f974bc9bb7794 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
@@ -52,9 +52,6 @@ class RegBankLegalizeHelper {
MachineOptimizationRemarkEmitter MORE;
const RegBankLegalizeRules &RBLRules;
const bool IsWave32;
- // The original next instruction after MI, saved before applyMappingDst may
- // insert instructions. Used by Sgpr*_WF to set the waterfall range.
- MachineBasicBlock::iterator OldNextMI;
const RegisterBank *SgprRB;
const RegisterBank *VgprRB;
const RegisterBank *AgprRB;
@@ -149,6 +146,7 @@ class RegBankLegalizeHelper {
bool lowerSplitTo32Select(MachineInstr &MI);
bool lowerSplitTo32SExtInReg(MachineInstr &MI);
bool lowerSplitBitCount64To32(MachineInstr &MI);
+ bool lowerBitReplicateToVALU(MachineInstr &MI);
bool lowerUnpackMinMax(MachineInstr &MI);
bool lowerUnpackAExt(MachineInstr &MI);
bool lowerSBufToBuf(MachineInstr &MI, WaterfallInfo &WFI);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index fc66d967dbe7b..09896093cb9ef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1764,7 +1764,7 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({amdgcn_s_bitreplicate}, Standard)
.Uni(S64, {{Sgpr64}, {IntrId, Sgpr32}})
- .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, Sgpr32_WF}});
+ .Div(S64, {{Vgpr64}, {IntrId, Vgpr32}, BitReplicateToVALU});
// Note: amdgcn.icmp with i1 inputs is legalized to ballot in the legalizer,
// so no S1 rules are needed here.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
index 16a5634a8eb62..fd6bfd9ae7128 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
@@ -349,7 +349,8 @@ enum LoweringMethodID {
DynStackAlloc,
DeletePrefetch,
LowerSetRounding,
- LowerGetRounding
+ LowerGetRounding,
+ BitReplicateToVALU
};
enum FastRulesTypes {
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
index 66677fa5e542f..b35907de70b38 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-call.ll
@@ -76,13 +76,11 @@ define amdgpu_ps i32 @test_divergent_indirect_call_p0_with_args(ptr %fptr, i32 %
; CHECK-NEXT: [[MV:%[0-9]+]]:vgpr(p0) = G_MERGE_VALUES [[COPY]](s32), [[COPY1]](s32)
; CHECK-NEXT: [[COPY2:%[0-9]+]]:vgpr(s32) = COPY $vgpr2
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(s32) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: bb.2:
; CHECK-NEXT: successors: %bb.3(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(s32) = G_PHI [[DEF1]](s32), %bb.1, %5(s32), %bb.3
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES [[MV]](p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -99,7 +97,7 @@ define amdgpu_ps i32 @test_divergent_indirect_call_p0_with_args(ptr %fptr, i32 %
; CHECK-NEXT: [[COPY3:%[0-9]+]]:sgpr(<4 x s32>) = COPY $private_rsrc_reg
; CHECK-NEXT: $sgpr0_sgpr1_sgpr2_sgpr3 = COPY [[COPY3]](<4 x s32>)
; CHECK-NEXT: $sgpr30_sgpr31 = noconvergent G_SI_CALL [[MV1]](p0), 0, csr_amdgpu_si_gfx_gfx90ainsts, implicit $vgpr0, implicit $sgpr0_sgpr1_sgpr2_sgpr3, implicit-def $vgpr0
- ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vgpr(s32) = COPY $vgpr0, implicit [[PHI]](tied-def 0)(s32)
+ ; CHECK-NEXT: [[COPY4:%[0-9]+]]:vgpr(s32) = COPY $vgpr0
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
; CHECK-NEXT: $exec = S_XOR_B64_term $exec, [[S_AND_SAVEEXEC_B64_]], implicit-def $scc
; CHECK-NEXT: SI_WATERFALL_LOOP %bb.2, implicit $exec
diff --git a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
index 934f6f72452ad..5207d992ea74d 100644
--- a/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
+++ b/llvm/test/CodeGen/AMDGPU/GlobalISel/regbankselect-waterfall-call.mir
@@ -13,13 +13,11 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p0) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p0) = G_PHI [[DEF1]](p0), %bb.0, %2(p0), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -33,7 +31,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p0) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0), implicit [[PHI]](tied-def 0)(p0)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0)
; CHECK-NEXT: %func_ptr:vgpr(p0) = G_LOAD [[COPY]](p0) :: (load (p0))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p0), 0, csr_amdgpu
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -68,13 +66,11 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p4) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p4) = G_PHI [[DEF1]](p4), %bb.0, %2(p4), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p4)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -88,7 +84,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p4) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4), implicit [[PHI]](tied-def 0)(p4)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4)
; CHECK-NEXT: %func_ptr:vgpr(p4) = G_LOAD [[COPY]](p4) :: (load (p4))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p4), 0, csr_amdgpu
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -123,13 +119,11 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p0) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p0) = G_PHI [[DEF1]](p0), %bb.0, %2(p0), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p0)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -143,7 +137,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p0) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0), implicit [[PHI]](tied-def 0)(p0)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p0) = COPY %g_ptr(p0)
; CHECK-NEXT: %func_ptr:vgpr(p0) = G_LOAD [[COPY]](p0) :: (load (p0))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p0), 0, csr_amdgpu, implicit $sgpr4, implicit $sgpr5, implicit-def $vgpr0
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
@@ -178,13 +172,11 @@ body: |
; CHECK-NEXT: liveins: $sgpr0_sgpr1
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: [[DEF:%[0-9]+]]:sreg_64_xexec = IMPLICIT_DEF
- ; CHECK-NEXT: [[DEF1:%[0-9]+]]:vgpr(p4) = IMPLICIT_DEF
; CHECK-NEXT: [[S_MOV_B64_:%[0-9]+]]:sreg_64_xexec = S_MOV_B64 $exec
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: .1:
; CHECK-NEXT: successors: %bb.2(0x80000000)
; CHECK-NEXT: {{ $}}
- ; CHECK-NEXT: [[PHI:%[0-9]+]]:vgpr(p4) = G_PHI [[DEF1]](p4), %bb.0, %2(p4), %bb.2
; CHECK-NEXT: [[UV:%[0-9]+]]:vgpr(s32), [[UV1:%[0-9]+]]:vgpr(s32) = G_UNMERGE_VALUES %func_ptr(p4)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV]](s32)
; CHECK-NEXT: [[INTRINSIC_CONVERGENT1:%[0-9]+]]:sgpr(s32) = G_INTRINSIC_CONVERGENT intrinsic(@llvm.amdgcn.readfirstlane), [[UV1]](s32)
@@ -198,7 +190,7 @@ body: |
; CHECK-NEXT: {{ $}}
; CHECK-NEXT: ADJCALLSTACKUP 0, 0, implicit-def $scc
; CHECK-NEXT: %g_ptr:sgpr(p4) = COPY $sgpr0_sgpr1
- ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4), implicit [[PHI]](tied-def 0)(p4)
+ ; CHECK-NEXT: [[COPY:%[0-9]+]]:vgpr(p4) = COPY %g_ptr(p4)
; CHECK-NEXT: %func_ptr:vgpr(p4) = G_LOAD [[COPY]](p4) :: (load (p4))
; CHECK-NEXT: $sgpr2_sgpr3 = G_SI_CALL [[MV]](p4), 0, csr_amdgpu, implicit $sgpr4, implicit $sgpr5, implicit-def $vgpr0
; CHECK-NEXT: ADJCALLSTACKDOWN 0, 0, implicit-def $scc
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index 121212a1dc1b1..d27204019a065 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -78,20 +78,22 @@ entry:
define amdgpu_cs void @test_s_bitreplicate_vgpr_store(i32 %mask, ptr addrspace(1) %out) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_store:
; GFX11-GISEL: ; %bb.0: ; %entry
-; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr3_vgpr4
-; GFX11-GISEL-NEXT: .LBB6_1: ; =>This Inner Loop Header: Depth=1
-; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v0
-; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
-; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v0
-; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v4, s3
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v3, s2
-; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
-; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB6_1
-; GFX11-GISEL-NEXT: ; %bb.2:
-; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
+; GFX11-GISEL-NEXT: v_perm_b32 v3, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 4, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0xf0f0f0f, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 2, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x33333333, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x55555555, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v4, v0, 1, v0
; GFX11-GISEL-NEXT: global_store_b64 v[1:2], v[3:4], off
; GFX11-GISEL-NEXT: s_endpgm
;
@@ -112,21 +114,23 @@ define i64 @test_s_bitreplicate_vgpr_multi_use(i32 %mask) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_multi_use:
; GFX11-GISEL: ; %bb.0: ; %entry
; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr1_vgpr2
-; GFX11-GISEL-NEXT: .LBB7_1: ; =>This Inner Loop Header: Depth=1
-; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v0
-; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
-; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v0
-; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, s2
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v2, s3
-; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
-; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB7_1
-; GFX11-GISEL-NEXT: ; %bb.2:
-; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
-; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v1, v2
+; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v1, v0
; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, 0
; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
;
@@ -152,21 +156,22 @@ define i64 @test_s_bitreplicate_vgpr(i32 %mask) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr:
; GFX11-GISEL: ; %bb.0: ; %entry
; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v2, v0
-; GFX11-GISEL-NEXT: s_mov_b32 s0, exec_lo
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr0_vgpr1
-; GFX11-GISEL-NEXT: .LBB8_1: ; =>This Inner Loop Header: Depth=1
-; GFX11-GISEL-NEXT: v_readfirstlane_b32 s2, v2
-; GFX11-GISEL-NEXT: s_mov_b32 s1, exec_lo
-; GFX11-GISEL-NEXT: v_cmpx_eq_u32_e64 s2, v2
-; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[2:3], s2
-; GFX11-GISEL-NEXT: ; implicit-def: $vgpr2
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v0, s2
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, s3
-; GFX11-GISEL-NEXT: s_xor_b32 exec_lo, exec_lo, s1
-; GFX11-GISEL-NEXT: s_cbranch_execnz .LBB8_1
-; GFX11-GISEL-NEXT: ; %bb.2:
-; GFX11-GISEL-NEXT: s_mov_b32 exec_lo, s0
+; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v2, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v2, 1, v2
; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr:
>From 1029b10d713150d355dc2921194a3bd653f6b4bf Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Thu, 14 May 2026 19:56:26 -0400
Subject: [PATCH 5/7] Add offload execution test
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 24 +++----
.../offloading/ompx_bare_s_bitreplicate.c | 63 +++++++++++++++++++
2 files changed, 75 insertions(+), 12 deletions(-)
create mode 100644 offload/test/offloading/ompx_bare_s_bitreplicate.c
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index f011d59089f68..4955670a61a96 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -1428,18 +1428,18 @@ bool RegBankLegalizeHelper::lowerBitReplicateToVALU(MachineInstr &MI) {
// 0x00-0x03 = src byte N, 0x0C = zero byte
// e.g. 0x134C3A92 -> Hi: [0013][004C], Lo: [003A][0092]
auto Zero = B.buildConstant(VgprRB_S32, 0);
- Register LoByteSpread =
- B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
- .addUse(Zero.getReg(0))
- .addUse(Src)
- .addUse(B.buildConstant(VgprRB_S32, 0x0C010C00).getReg(0))
- .getReg(0);
- Register HiByteSpread =
- B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
- .addUse(Zero.getReg(0))
- .addUse(Src)
- .addUse(B.buildConstant(VgprRB_S32, 0x0C030C02).getReg(0))
- .getReg(0);
+ auto LoSel = B.buildConstant(VgprRB_S32, 0x0C010C00);
+ auto HiSel = B.buildConstant(VgprRB_S32, 0x0C030C02);
+ Register LoByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(LoSel.getReg(0))
+ .getReg(0);
+ Register HiByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(HiSel.getReg(0))
+ .getReg(0);
// Spread masks: each keeps the low N bits of every 2N-bit group, clearing
// the garbage left by the shift+OR.
diff --git a/offload/test/offloading/ompx_bare_s_bitreplicate.c b/offload/test/offloading/ompx_bare_s_bitreplicate.c
new file mode 100644
index 0000000000000..722dac5389d95
--- /dev/null
+++ b/offload/test/offloading/ompx_bare_s_bitreplicate.c
@@ -0,0 +1,63 @@
+// End-to-end test for the divergent VALU expansion of the AMDGPU
+// s_bitreplicate intrinsic (lowerBitReplicateToVALU in GlobalISel
+// RegBankLegalize).
+//
+// RUN: %libomptarget-compile-generic \
+// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mllvm=-global-isel \
+// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mllvm=-new-reg-bank-select
+// RUN: %libomptarget-run-generic | %fcheck-generic
+//
+// REQUIRES: amdgpu
+
+#include <assert.h>
+#include <ompx.h>
+#include <stdint.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+// Host-side reference: replicate each bit of x into two adjacent output bits.
+static uint64_t bitreplicate_ref(uint32_t x) {
+ uint64_t r = 0;
+ for (int i = 0; i < 32; ++i)
+ if (x & (1u << i))
+ r |= ((uint64_t)0x3) << (2 * i);
+ return r;
+}
+
+// Call llvm.amdgcn.s.bitreplicate directly via __asm__ symbol rename so we
+// don't need a clang builtin. Guard the host side with declare variant.
+#pragma omp begin declare variant match(device = {arch(amdgcn)})
+uint64_t
+__amdgcn_s_bitreplicate(uint32_t) __asm__("llvm.amdgcn.s.bitreplicate");
+static uint64_t s_bitreplicate(uint32_t x) {
+ return __amdgcn_s_bitreplicate(x);
+}
+#pragma omp end declare variant
+#pragma omp begin declare variant match(device = {kind(cpu)})
+static uint64_t s_bitreplicate(uint32_t x) { return bitreplicate_ref(x); }
+#pragma omp end declare variant
+
+int main(int argc, char *argv[]) {
+ const int num_blocks = 1;
+ const int block_size = 256;
+ const int N = num_blocks * block_size;
+ uint64_t *res = (uint64_t *)malloc(N * sizeof(uint64_t));
+
+#pragma omp target teams ompx_bare num_teams(num_blocks) \
+ thread_limit(block_size) map(from : res[0 : N])
+ {
+ int bid = ompx_block_id_x();
+ int bdim = ompx_block_dim_x();
+ int tid = ompx_thread_id_x();
+ int idx = bid * bdim + tid;
+ res[idx] = s_bitreplicate((uint32_t)idx + 1);
+ }
+ for (int i = 0; i < N; ++i)
+ assert(res[i] == bitreplicate_ref((uint32_t)i + 1));
+
+ // CHECK: PASS
+ printf("PASS\n");
+
+ free(res);
+ return 0;
+}
>From 0816af0ebdae3ae461769247729ea783fb358ac0 Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Tue, 26 May 2026 17:28:16 -0400
Subject: [PATCH 6/7] Add TODO and revert valu expansion
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 76 ----------------
.../AMDGPU/AMDGPURegBankLegalizeHelper.h | 1 -
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 4 +-
.../AMDGPU/AMDGPURegBankLegalizeRules.h | 3 +-
.../AMDGPU/llvm.amdgcn.bitreplicate.ll | 91 ++++---------------
.../offloading/ompx_bare_s_bitreplicate.c | 63 -------------
6 files changed, 22 insertions(+), 216 deletions(-)
delete mode 100644 offload/test/offloading/ompx_bare_s_bitreplicate.c
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index 4955670a61a96..b9554d7858d32 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -1411,80 +1411,6 @@ bool RegBankLegalizeHelper::lowerGetRounding(MachineInstr &MI) {
return true;
}
-bool RegBankLegalizeHelper::lowerBitReplicateToVALU(MachineInstr &MI) {
- // Lower divergent s_bitreplicate (bit i -> bits 2i, 2i+1) to VALU by
- // splitting into two 16-bit halves and spreading bits via:
- // lo = v_perm_b32(0, src, 0x0C010C00) // [00][B1][00][B0]
- // hi = v_perm_b32(0, src, 0x0C030C02) // [00][B3][00][B2]
- // M4 = (x | (x << 4)) & 0x0F0F0F0F // 4-bit spread
- // M2 = (M4 | (M4 << 2)) & 0x33333333 // 2-bit spread
- // M1 = (M2 | (M2 << 1)) & 0x55555555 // 1-bit spread
- // R = M1 | (M1 << 1) // duplicate each bit
- Register Dst = MI.getOperand(0).getReg();
- Register Src = MI.getOperand(2).getReg();
-
- // v_perm_b32 splits the input into 16-bit halves and spreads each half's
- // bytes with zero-byte gaps in one instruction. Selector byte values:
- // 0x00-0x03 = src byte N, 0x0C = zero byte
- // e.g. 0x134C3A92 -> Hi: [0013][004C], Lo: [003A][0092]
- auto Zero = B.buildConstant(VgprRB_S32, 0);
- auto LoSel = B.buildConstant(VgprRB_S32, 0x0C010C00);
- auto HiSel = B.buildConstant(VgprRB_S32, 0x0C030C02);
- Register LoByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
- .addUse(Zero.getReg(0))
- .addUse(Src)
- .addUse(LoSel.getReg(0))
- .getReg(0);
- Register HiByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
- .addUse(Zero.getReg(0))
- .addUse(Src)
- .addUse(HiSel.getReg(0))
- .getReg(0);
-
- // Spread masks: each keeps the low N bits of every 2N-bit group, clearing
- // the garbage left by the shift+OR.
- auto Mask4 =
- B.buildConstant(VgprRB_S32, 0x0F0F0F0F); // 0000_1111: low 4 bits per byte
- auto Mask2 = B.buildConstant(
- VgprRB_S32, 0x33333333); // 0011_0011: low 2 bits per 4-bit group
- auto Mask1 =
- B.buildConstant(VgprRB_S32, 0x55555555); // 0101_0101: every other bit
-
- auto Spread = [&](Register In) -> Register {
- // 4-bit spread: separate each byte into its two 4-bit halves.
- // e.g. [003A][0092] -> [03][0A][09][02]
- auto S4 = B.buildShl(VgprRB_S32, In, B.buildConstant(VgprRB_S32, 4));
- auto Or4 = B.buildOr(VgprRB_S32, In, S4);
- auto M4 = B.buildAnd(VgprRB_S32, Or4, Mask4);
-
- // 2-bit spread: separate each 4-bit group into two 2-bit pairs.
- // e.g. [03][0A][09][02] -> [0000][0011][0010][0010][0010][0001][0000][0010]
- auto S2 = B.buildShl(VgprRB_S32, M4, B.buildConstant(VgprRB_S32, 2));
- auto Or2 = B.buildOr(VgprRB_S32, M4, S2);
- auto M2 = B.buildAnd(VgprRB_S32, Or2, Mask2);
-
- // 1-bit spread: separate each 2-bit pair into individual bits.
- // e.g. -> [00][00][01][01][01][00][01][00][01][00][00][01][00][00][01][00]
- auto S1 = B.buildShl(VgprRB_S32, M2, B.buildConstant(VgprRB_S32, 1));
- auto Or1 = B.buildOr(VgprRB_S32, M2, S1);
- auto M1 = B.buildAnd(VgprRB_S32, Or1, Mask1);
-
- // Duplicate: double each isolated bit.
- // e.g. -> [00][00][11][11][11][00][11][00][11][00][00][11][00][00][11][00]
- auto Dup = B.buildShl(VgprRB_S32, M1, B.buildConstant(VgprRB_S32, 1));
- auto Result = B.buildOr(VgprRB_S32, M1, Dup);
-
- // e.g. -> return 0x0FCCC30C.
- return Result.getReg(0);
- };
-
- Register LoResult = Spread(LoByteSpread);
- Register HiResult = Spread(HiByteSpread);
- B.buildMergeLikeInstr(Dst, {LoResult, HiResult});
- MI.eraseFromParent();
- return true;
-}
-
bool RegBankLegalizeHelper::lower(MachineInstr &MI,
const RegBankLLTMapping &Mapping,
WaterfallInfo &WFI) {
@@ -1865,8 +1791,6 @@ bool RegBankLegalizeHelper::lower(MachineInstr &MI,
return lowerSetRounding(MI);
case LowerGetRounding:
return lowerGetRounding(MI);
- case BitReplicateToVALU:
- return lowerBitReplicateToVALU(MI);
}
return true;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
index f974bc9bb7794..25c8c3a0f6127 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
@@ -146,7 +146,6 @@ class RegBankLegalizeHelper {
bool lowerSplitTo32Select(MachineInstr &MI);
bool lowerSplitTo32SExtInReg(MachineInstr &MI);
bool lowerSplitBitCount64To32(MachineInstr &MI);
- bool lowerBitReplicateToVALU(MachineInstr &MI);
bool lowerUnpackMinMax(MachineInstr &MI);
bool lowerUnpackAExt(MachineInstr &MI);
bool lowerSBufToBuf(MachineInstr &MI, WaterfallInfo &WFI);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 09896093cb9ef..7d92e1771a263 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1762,9 +1762,11 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({returnaddress}).Any({{UniP0}, {{SgprP0}, {}}});
+ // TODO: Divergent case can be expanded to VALU operations, however the
+ // results should probably be verified with a GPU execution test.
addRulesForIOpcs({amdgcn_s_bitreplicate}, Standard)
.Uni(S64, {{Sgpr64}, {IntrId, Sgpr32}})
- .Div(S64, {{Vgpr64}, {IntrId, Vgpr32}, BitReplicateToVALU});
+ .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, SgprB32_ReadFirstLane}});
// Note: amdgcn.icmp with i1 inputs is legalized to ballot in the legalizer,
// so no S1 rules are needed here.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
index fd6bfd9ae7128..16a5634a8eb62 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
@@ -349,8 +349,7 @@ enum LoweringMethodID {
DynStackAlloc,
DeletePrefetch,
LowerSetRounding,
- LowerGetRounding,
- BitReplicateToVALU
+ LowerGetRounding
};
enum FastRulesTypes {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index d27204019a065..759d7cf69bd25 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -76,34 +76,13 @@ entry:
}
define amdgpu_cs void @test_s_bitreplicate_vgpr_store(i32 %mask, ptr addrspace(1) %out) {
-; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_store:
-; GFX11-GISEL: ; %bb.0: ; %entry
-; GFX11-GISEL-NEXT: v_perm_b32 v3, 0, v0, 0xc010c00
-; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 4, v3
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0xf0f0f0f, v3
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 2, v3
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x33333333, v3
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x55555555, v3
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v4, v0, 1, v0
-; GFX11-GISEL-NEXT: global_store_b64 v[1:2], v[3:4], off
-; GFX11-GISEL-NEXT: s_endpgm
-;
-; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr_store:
-; GFX11-SDAG: ; %bb.0: ; %entry
-; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-SDAG-NEXT: v_dual_mov_b32 v4, s1 :: v_dual_mov_b32 v3, s0
-; GFX11-SDAG-NEXT: global_store_b64 v[1:2], v[3:4], off
-; GFX11-SDAG-NEXT: s_endpgm
+; GFX11-LABEL: test_s_bitreplicate_vgpr_store:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-NEXT: v_dual_mov_b32 v4, s1 :: v_dual_mov_b32 v3, s0
+; GFX11-NEXT: global_store_b64 v[1:2], v[3:4], off
+; GFX11-NEXT: s_endpgm
entry:
%br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
store i64 %br, ptr addrspace(1) %out
@@ -114,23 +93,10 @@ define i64 @test_s_bitreplicate_vgpr_multi_use(i32 %mask) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_multi_use:
; GFX11-GISEL: ; %bb.0: ; %entry
; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
-; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
-; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v1, v0
+; GFX11-GISEL-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v0, v1
; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, 0
; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
;
@@ -153,34 +119,13 @@ entry:
}
define i64 @test_s_bitreplicate_vgpr(i32 %mask) {
-; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr:
-; GFX11-GISEL: ; %bb.0: ; %entry
-; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
-; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
-; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
-; GFX11-GISEL-NEXT: v_and_b32_e32 v2, 0x55555555, v0
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v1, 1, v1
-; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v2, 1, v2
-; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
-;
-; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr:
-; GFX11-SDAG: ; %bb.0: ; %entry
-; GFX11-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-SDAG-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
-; GFX11-SDAG-NEXT: s_setpc_b64 s[30:31]
+; GFX11-LABEL: test_s_bitreplicate_vgpr:
+; GFX11: ; %bb.0: ; %entry
+; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-NEXT: s_setpc_b64 s[30:31]
entry:
%br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
ret i64 %br
diff --git a/offload/test/offloading/ompx_bare_s_bitreplicate.c b/offload/test/offloading/ompx_bare_s_bitreplicate.c
deleted file mode 100644
index 722dac5389d95..0000000000000
--- a/offload/test/offloading/ompx_bare_s_bitreplicate.c
+++ /dev/null
@@ -1,63 +0,0 @@
-// End-to-end test for the divergent VALU expansion of the AMDGPU
-// s_bitreplicate intrinsic (lowerBitReplicateToVALU in GlobalISel
-// RegBankLegalize).
-//
-// RUN: %libomptarget-compile-generic \
-// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mllvm=-global-isel \
-// RUN: -Xopenmp-target=amdgcn-amd-amdhsa -mllvm=-new-reg-bank-select
-// RUN: %libomptarget-run-generic | %fcheck-generic
-//
-// REQUIRES: amdgpu
-
-#include <assert.h>
-#include <ompx.h>
-#include <stdint.h>
-#include <stdio.h>
-#include <stdlib.h>
-
-// Host-side reference: replicate each bit of x into two adjacent output bits.
-static uint64_t bitreplicate_ref(uint32_t x) {
- uint64_t r = 0;
- for (int i = 0; i < 32; ++i)
- if (x & (1u << i))
- r |= ((uint64_t)0x3) << (2 * i);
- return r;
-}
-
-// Call llvm.amdgcn.s.bitreplicate directly via __asm__ symbol rename so we
-// don't need a clang builtin. Guard the host side with declare variant.
-#pragma omp begin declare variant match(device = {arch(amdgcn)})
-uint64_t
-__amdgcn_s_bitreplicate(uint32_t) __asm__("llvm.amdgcn.s.bitreplicate");
-static uint64_t s_bitreplicate(uint32_t x) {
- return __amdgcn_s_bitreplicate(x);
-}
-#pragma omp end declare variant
-#pragma omp begin declare variant match(device = {kind(cpu)})
-static uint64_t s_bitreplicate(uint32_t x) { return bitreplicate_ref(x); }
-#pragma omp end declare variant
-
-int main(int argc, char *argv[]) {
- const int num_blocks = 1;
- const int block_size = 256;
- const int N = num_blocks * block_size;
- uint64_t *res = (uint64_t *)malloc(N * sizeof(uint64_t));
-
-#pragma omp target teams ompx_bare num_teams(num_blocks) \
- thread_limit(block_size) map(from : res[0 : N])
- {
- int bid = ompx_block_id_x();
- int bdim = ompx_block_dim_x();
- int tid = ompx_thread_id_x();
- int idx = bid * bdim + tid;
- res[idx] = s_bitreplicate((uint32_t)idx + 1);
- }
- for (int i = 0; i < N; ++i)
- assert(res[i] == bitreplicate_ref((uint32_t)i + 1));
-
- // CHECK: PASS
- printf("PASS\n");
-
- free(res);
- return 0;
-}
>From 5f7fa0618399d156fdcdcdf1e85264cc0222bf04 Mon Sep 17 00:00:00 2001
From: Vang Thao <Vang.Thao at amd.com>
Date: Tue, 30 Jun 2026 22:54:21 -0400
Subject: [PATCH 7/7] Re-add valu expansion without offload test
---
.../AMDGPU/AMDGPURegBankLegalizeHelper.cpp | 76 +++++++++++++++
.../AMDGPU/AMDGPURegBankLegalizeHelper.h | 1 +
.../AMDGPU/AMDGPURegBankLegalizeRules.cpp | 4 +-
.../AMDGPU/AMDGPURegBankLegalizeRules.h | 3 +-
.../AMDGPU/llvm.amdgcn.bitreplicate.ll | 92 +++++++++++++++----
5 files changed, 153 insertions(+), 23 deletions(-)
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
index b9554d7858d32..4955670a61a96 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.cpp
@@ -1411,6 +1411,80 @@ bool RegBankLegalizeHelper::lowerGetRounding(MachineInstr &MI) {
return true;
}
+bool RegBankLegalizeHelper::lowerBitReplicateToVALU(MachineInstr &MI) {
+ // Lower divergent s_bitreplicate (bit i -> bits 2i, 2i+1) to VALU by
+ // splitting into two 16-bit halves and spreading bits via:
+ // lo = v_perm_b32(0, src, 0x0C010C00) // [00][B1][00][B0]
+ // hi = v_perm_b32(0, src, 0x0C030C02) // [00][B3][00][B2]
+ // M4 = (x | (x << 4)) & 0x0F0F0F0F // 4-bit spread
+ // M2 = (M4 | (M4 << 2)) & 0x33333333 // 2-bit spread
+ // M1 = (M2 | (M2 << 1)) & 0x55555555 // 1-bit spread
+ // R = M1 | (M1 << 1) // duplicate each bit
+ Register Dst = MI.getOperand(0).getReg();
+ Register Src = MI.getOperand(2).getReg();
+
+ // v_perm_b32 splits the input into 16-bit halves and spreads each half's
+ // bytes with zero-byte gaps in one instruction. Selector byte values:
+ // 0x00-0x03 = src byte N, 0x0C = zero byte
+ // e.g. 0x134C3A92 -> Hi: [0013][004C], Lo: [003A][0092]
+ auto Zero = B.buildConstant(VgprRB_S32, 0);
+ auto LoSel = B.buildConstant(VgprRB_S32, 0x0C010C00);
+ auto HiSel = B.buildConstant(VgprRB_S32, 0x0C030C02);
+ Register LoByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(LoSel.getReg(0))
+ .getReg(0);
+ Register HiByteSpread = B.buildIntrinsic(Intrinsic::amdgcn_perm, {VgprRB_S32})
+ .addUse(Zero.getReg(0))
+ .addUse(Src)
+ .addUse(HiSel.getReg(0))
+ .getReg(0);
+
+ // Spread masks: each keeps the low N bits of every 2N-bit group, clearing
+ // the garbage left by the shift+OR.
+ auto Mask4 =
+ B.buildConstant(VgprRB_S32, 0x0F0F0F0F); // 0000_1111: low 4 bits per byte
+ auto Mask2 = B.buildConstant(
+ VgprRB_S32, 0x33333333); // 0011_0011: low 2 bits per 4-bit group
+ auto Mask1 =
+ B.buildConstant(VgprRB_S32, 0x55555555); // 0101_0101: every other bit
+
+ auto Spread = [&](Register In) -> Register {
+ // 4-bit spread: separate each byte into its two 4-bit halves.
+ // e.g. [003A][0092] -> [03][0A][09][02]
+ auto S4 = B.buildShl(VgprRB_S32, In, B.buildConstant(VgprRB_S32, 4));
+ auto Or4 = B.buildOr(VgprRB_S32, In, S4);
+ auto M4 = B.buildAnd(VgprRB_S32, Or4, Mask4);
+
+ // 2-bit spread: separate each 4-bit group into two 2-bit pairs.
+ // e.g. [03][0A][09][02] -> [0000][0011][0010][0010][0010][0001][0000][0010]
+ auto S2 = B.buildShl(VgprRB_S32, M4, B.buildConstant(VgprRB_S32, 2));
+ auto Or2 = B.buildOr(VgprRB_S32, M4, S2);
+ auto M2 = B.buildAnd(VgprRB_S32, Or2, Mask2);
+
+ // 1-bit spread: separate each 2-bit pair into individual bits.
+ // e.g. -> [00][00][01][01][01][00][01][00][01][00][00][01][00][00][01][00]
+ auto S1 = B.buildShl(VgprRB_S32, M2, B.buildConstant(VgprRB_S32, 1));
+ auto Or1 = B.buildOr(VgprRB_S32, M2, S1);
+ auto M1 = B.buildAnd(VgprRB_S32, Or1, Mask1);
+
+ // Duplicate: double each isolated bit.
+ // e.g. -> [00][00][11][11][11][00][11][00][11][00][00][11][00][00][11][00]
+ auto Dup = B.buildShl(VgprRB_S32, M1, B.buildConstant(VgprRB_S32, 1));
+ auto Result = B.buildOr(VgprRB_S32, M1, Dup);
+
+ // e.g. -> return 0x0FCCC30C.
+ return Result.getReg(0);
+ };
+
+ Register LoResult = Spread(LoByteSpread);
+ Register HiResult = Spread(HiByteSpread);
+ B.buildMergeLikeInstr(Dst, {LoResult, HiResult});
+ MI.eraseFromParent();
+ return true;
+}
+
bool RegBankLegalizeHelper::lower(MachineInstr &MI,
const RegBankLLTMapping &Mapping,
WaterfallInfo &WFI) {
@@ -1791,6 +1865,8 @@ bool RegBankLegalizeHelper::lower(MachineInstr &MI,
return lowerSetRounding(MI);
case LowerGetRounding:
return lowerGetRounding(MI);
+ case BitReplicateToVALU:
+ return lowerBitReplicateToVALU(MI);
}
return true;
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
index 25c8c3a0f6127..f974bc9bb7794 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeHelper.h
@@ -146,6 +146,7 @@ class RegBankLegalizeHelper {
bool lowerSplitTo32Select(MachineInstr &MI);
bool lowerSplitTo32SExtInReg(MachineInstr &MI);
bool lowerSplitBitCount64To32(MachineInstr &MI);
+ bool lowerBitReplicateToVALU(MachineInstr &MI);
bool lowerUnpackMinMax(MachineInstr &MI);
bool lowerUnpackAExt(MachineInstr &MI);
bool lowerSBufToBuf(MachineInstr &MI, WaterfallInfo &WFI);
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
index 7d92e1771a263..09896093cb9ef 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.cpp
@@ -1762,11 +1762,9 @@ RegBankLegalizeRules::RegBankLegalizeRules(const GCNSubtarget &_ST,
addRulesForIOpcs({returnaddress}).Any({{UniP0}, {{SgprP0}, {}}});
- // TODO: Divergent case can be expanded to VALU operations, however the
- // results should probably be verified with a GPU execution test.
addRulesForIOpcs({amdgcn_s_bitreplicate}, Standard)
.Uni(S64, {{Sgpr64}, {IntrId, Sgpr32}})
- .Div(S64, {{Sgpr64ToVgprDst}, {IntrId, SgprB32_ReadFirstLane}});
+ .Div(S64, {{Vgpr64}, {IntrId, Vgpr32}, BitReplicateToVALU});
// Note: amdgcn.icmp with i1 inputs is legalized to ballot in the legalizer,
// so no S1 rules are needed here.
diff --git a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
index 16a5634a8eb62..fd6bfd9ae7128 100644
--- a/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
+++ b/llvm/lib/Target/AMDGPU/AMDGPURegBankLegalizeRules.h
@@ -349,7 +349,8 @@ enum LoweringMethodID {
DynStackAlloc,
DeletePrefetch,
LowerSetRounding,
- LowerGetRounding
+ LowerGetRounding,
+ BitReplicateToVALU
};
enum FastRulesTypes {
diff --git a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
index 759d7cf69bd25..ebd7df565f661 100644
--- a/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
+++ b/llvm/test/CodeGen/AMDGPU/llvm.amdgcn.bitreplicate.ll
@@ -76,13 +76,34 @@ entry:
}
define amdgpu_cs void @test_s_bitreplicate_vgpr_store(i32 %mask, ptr addrspace(1) %out) {
-; GFX11-LABEL: test_s_bitreplicate_vgpr_store:
-; GFX11: ; %bb.0: ; %entry
-; GFX11-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-NEXT: v_dual_mov_b32 v4, s1 :: v_dual_mov_b32 v3, s0
-; GFX11-NEXT: global_store_b64 v[1:2], v[3:4], off
-; GFX11-NEXT: s_endpgm
+; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_store:
+; GFX11-GISEL: ; %bb.0: ; %entry
+; GFX11-GISEL-NEXT: v_perm_b32 v3, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 4, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0xf0f0f0f, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 2, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x33333333, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v3, 0x55555555, v3
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v3, v3, 1, v3
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v4, v0, 1, v0
+; GFX11-GISEL-NEXT: global_store_b64 v[1:2], v[3:4], off
+; GFX11-GISEL-NEXT: s_endpgm
+;
+; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr_store:
+; GFX11-SDAG: ; %bb.0: ; %entry
+; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-SDAG-NEXT: v_dual_mov_b32 v4, s1 :: v_dual_mov_b32 v3, s0
+; GFX11-SDAG-NEXT: global_store_b64 v[1:2], v[3:4], off
+; GFX11-SDAG-NEXT: s_endpgm
entry:
%br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
store i64 %br, ptr addrspace(1) %out
@@ -93,11 +114,23 @@ define i64 @test_s_bitreplicate_vgpr_multi_use(i32 %mask) {
; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr_multi_use:
; GFX11-GISEL: ; %bb.0: ; %entry
; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-GISEL-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-GISEL-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-GISEL-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
-; GFX11-GISEL-NEXT: v_add_nc_u32_e32 v0, v0, v1
-; GFX11-GISEL-NEXT: v_mov_b32_e32 v1, 0
+; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_dual_mov_b32 v1, 0 :: v_dual_add_nc_u32 v0, v1, v0
; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
;
; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr_multi_use:
@@ -119,13 +152,34 @@ entry:
}
define i64 @test_s_bitreplicate_vgpr(i32 %mask) {
-; GFX11-LABEL: test_s_bitreplicate_vgpr:
-; GFX11: ; %bb.0: ; %entry
-; GFX11-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
-; GFX11-NEXT: v_readfirstlane_b32 s0, v0
-; GFX11-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
-; GFX11-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
-; GFX11-NEXT: s_setpc_b64 s[30:31]
+; GFX11-GISEL-LABEL: test_s_bitreplicate_vgpr:
+; GFX11-GISEL: ; %bb.0: ; %entry
+; GFX11-GISEL-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-GISEL-NEXT: v_perm_b32 v1, 0, v0, 0xc010c00
+; GFX11-GISEL-NEXT: v_perm_b32 v0, 0, v0, 0xc030c02
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 4, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 4, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0xf0f0f0f, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0xf0f0f0f, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 2, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 2, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x33333333, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v0, 0x33333333, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v0, 1, v0
+; GFX11-GISEL-NEXT: v_and_b32_e32 v1, 0x55555555, v1
+; GFX11-GISEL-NEXT: v_and_b32_e32 v2, 0x55555555, v0
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v0, v1, 1, v1
+; GFX11-GISEL-NEXT: v_lshl_or_b32 v1, v2, 1, v2
+; GFX11-GISEL-NEXT: s_setpc_b64 s[30:31]
+;
+; GFX11-SDAG-LABEL: test_s_bitreplicate_vgpr:
+; GFX11-SDAG: ; %bb.0: ; %entry
+; GFX11-SDAG-NEXT: s_waitcnt vmcnt(0) expcnt(0) lgkmcnt(0)
+; GFX11-SDAG-NEXT: v_readfirstlane_b32 s0, v0
+; GFX11-SDAG-NEXT: s_bitreplicate_b64_b32 s[0:1], s0
+; GFX11-SDAG-NEXT: v_dual_mov_b32 v0, s0 :: v_dual_mov_b32 v1, s1
+; GFX11-SDAG-NEXT: s_setpc_b64 s[30:31]
entry:
%br = call i64 @llvm.amdgcn.s.bitreplicate(i32 %mask)
ret i64 %br
More information about the llvm-commits
mailing list