Skip to content

Commit 5af52a1

Browse files
Prince781Prince781
committed
[NVPTX] update combiner rule for more types of loads
Handle more loads, including ones with multiple proxy registers: - i64 = LOAD - i64 = LoadParam - v2f32,v2f32 = LoadParamV2 Also update the test cases. Because this is an optimization, it is not triggered for some of these tests that compile with no optimizations.
1 parent 0060ea0 commit 5af52a1

File tree

7 files changed

+208
-173
lines changed

7 files changed

+208
-173
lines changed

‎llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

Lines changed: 96 additions & 73 deletions
Original file line numberDiff line numberDiff line change
@@ -5089,11 +5089,13 @@ PerformFADDCombineWithOperands(SDNode *N, SDValue N0, SDValue N1,
50895089
return SDValue();
50905090
}
50915091

5092+
/// OverrideVT - allows overriding result and memory type
50925093
static std::optional<std::pair<SDValue, SDValue>>
5093-
convertVectorLoad(SDNode *N, SelectionDAG &DAG, bool BuildVector) {
5094+
convertVectorLoad(SDNode *N, SelectionDAG &DAG, bool BuildVector,
5095+
std::optional<EVT> OverrideVT = std::nullopt) {
50945096
LoadSDNode *LD = cast<LoadSDNode>(N);
5095-
const EVT ResVT = LD->getValueType(0);
5096-
const EVT MemVT = LD->getMemoryVT();
5097+
const EVT ResVT = OverrideVT.value_or(LD->getValueType(0));
5098+
const EVT MemVT = OverrideVT.value_or(LD->getMemoryVT());
50975099

50985100
// If we're doing sign/zero extension as part of the load, avoid lowering to
50995101
// a LoadV node. TODO: consider relaxing this restriction.
@@ -5147,33 +5149,31 @@ convertVectorLoad(SDNode *N, SelectionDAG &DAG, bool BuildVector) {
51475149
// pass along the extension information
51485150
OtherOps.push_back(DAG.getIntPtrConstant(LD->getExtensionType(), DL));
51495151

5150-
SDValue NewLD = DAG.getMemIntrinsicNode(Opcode, DL, LdResVTs, OtherOps,
5151-
LD->getMemoryVT(),
5152+
SDValue NewLD = DAG.getMemIntrinsicNode(Opcode, DL, LdResVTs, OtherOps, MemVT,
51525153
LD->getMemOperand());
5153-
5154-
SmallVector<SDValue> ScalarRes;
5155-
if (EltVT.isVector()) {
5156-
assert(EVT(EltVT.getVectorElementType()) == ResVT.getVectorElementType());
5157-
assert(NumElts * EltVT.getVectorNumElements() ==
5158-
ResVT.getVectorNumElements());
5159-
// Generate EXTRACT_VECTOR_ELTs to split v2[i,f,bf]16/v4i8 subvectors back
5160-
// into individual elements.
5161-
for (const unsigned I : llvm::seq(NumElts)) {
5162-
SDValue SubVector = NewLD.getValue(I);
5163-
DAG.ExtractVectorElements(SubVector, ScalarRes);
5164-
}
5165-
} else {
5166-
for (const unsigned I : llvm::seq(NumElts)) {
5167-
SDValue Res = NewLD.getValue(I);
5168-
if (LoadEltVT != EltVT)
5169-
Res = DAG.getNode(ISD::TRUNCATE, DL, EltVT, Res);
5170-
ScalarRes.push_back(Res);
5171-
}
5172-
}
5173-
51745154
SDValue LoadChain = NewLD.getValue(NumElts);
51755155

51765156
if (BuildVector) {
5157+
SmallVector<SDValue> ScalarRes;
5158+
if (EltVT.isVector()) {
5159+
assert(EVT(EltVT.getVectorElementType()) == ResVT.getVectorElementType());
5160+
assert(NumElts * EltVT.getVectorNumElements() ==
5161+
ResVT.getVectorNumElements());
5162+
// Generate EXTRACT_VECTOR_ELTs to split v2[i,f,bf]16/v4i8 subvectors back
5163+
// into individual elements.
5164+
for (const unsigned I : llvm::seq(NumElts)) {
5165+
SDValue SubVector = NewLD.getValue(I);
5166+
DAG.ExtractVectorElements(SubVector, ScalarRes);
5167+
}
5168+
} else {
5169+
for (const unsigned I : llvm::seq(NumElts)) {
5170+
SDValue Res = NewLD.getValue(I);
5171+
if (LoadEltVT != EltVT)
5172+
Res = DAG.getNode(ISD::TRUNCATE, DL, EltVT, Res);
5173+
ScalarRes.push_back(Res);
5174+
}
5175+
}
5176+
51775177
const MVT BuildVecVT =
51785178
MVT::getVectorVT(EltVT.getScalarType(), ScalarRes.size());
51795179
SDValue BuildVec = DAG.getBuildVector(BuildVecVT, DL, ScalarRes);
@@ -5188,23 +5188,20 @@ convertVectorLoad(SDNode *N, SelectionDAG &DAG, bool BuildVector) {
51885188
static SDValue PerformLoadCombine(SDNode *N,
51895189
TargetLowering::DAGCombinerInfo &DCI) {
51905190
auto *MemN = cast<MemSDNode>(N);
5191-
EVT MemVT = MemN->getMemoryVT();
5192-
5193-
// ignore volatile loads
5194-
if (MemN->isVolatile())
5195-
return SDValue();
5196-
51975191
// only operate on vectors of f32s / i64s
5198-
if (!MemVT.isVector())
5192+
if (EVT MemVT = MemN->getMemoryVT();
5193+
!(MemVT == MVT::i64 ||
5194+
(MemVT.isVector() && (MemVT.getVectorElementType() == MVT::f32 ||
5195+
MemVT.getVectorElementType() == MVT::i64))))
51995196
return SDValue();
52005197

5201-
EVT ElementVT = MemVT.getVectorElementType();
5202-
if (!(ElementVT == MVT::f32 ||
5203-
(ElementVT == MVT::i64 && N->getOpcode() != ISD::LOAD)))
5204-
return SDValue();
5198+
const unsigned OrigNumResults =
5199+
llvm::count_if(N->values(), [](const auto &VT) {
5200+
return VT == MVT::i64 || VT == MVT::f32 || VT.isVector();
5201+
});
52055202

52065203
SmallDenseMap<SDNode *, unsigned> ExtractElts;
5207-
SDNode *ProxyReg = nullptr;
5204+
SmallVector<SDNode *> ProxyRegs(OrigNumResults, nullptr);
52085205
SmallVector<std::pair<SDNode *, unsigned /*offset*/>> WorkList{{N, 0}};
52095206
while (!WorkList.empty()) {
52105207
auto [V, Offset] = WorkList.pop_back_val();
@@ -5217,8 +5214,14 @@ static SDValue PerformLoadCombine(SDNode *N,
52175214

52185215
SDNode *User = U.getUser();
52195216
if (User->getOpcode() == NVPTXISD::ProxyReg) {
5217+
Offset = U.getResNo() * 2;
5218+
SDNode *&ProxyReg = ProxyRegs[Offset / 2];
5219+
5220+
// We shouldn't have multiple proxy regs for the same value from the
5221+
// load, but bail out anyway since we don't handle this.
52205222
if (ProxyReg)
5221-
return SDValue(); // bail out if we've seen a proxy reg?
5223+
return SDValue();
5224+
52225225
ProxyReg = User;
52235226
} else if (User->getOpcode() == ISD::BITCAST &&
52245227
User->getValueType(0) == MVT::v2f32 &&
@@ -5308,9 +5311,18 @@ static SDValue PerformLoadCombine(SDNode *N,
53085311
if (NewGlueIdx)
53095312
NewGlue = NewLoad.getValue(*NewGlueIdx);
53105313
} else if (N->getOpcode() == ISD::LOAD) { // rewrite a load
5311-
if (auto Result = convertVectorLoad(N, DCI.DAG, /*BuildVector=*/false)) {
5314+
std::optional<EVT> CastToType;
5315+
EVT ResVT = N->getValueType(0);
5316+
if (ResVT == MVT::i64) {
5317+
// ld.b64 is treated as a vector by subsequent code
5318+
CastToType = MVT::v2f32;
5319+
}
5320+
if (auto Result =
5321+
convertVectorLoad(N, DCI.DAG, /*BuildVector=*/false, CastToType)) {
53125322
std::tie(NewLoad, NewChain) = *Result;
5313-
NumElts = MemVT.getVectorNumElements();
5323+
NumElts =
5324+
CastToType.value_or(cast<MemSDNode>(NewLoad.getNode())->getMemoryVT())
5325+
.getVectorNumElements();
53145326
if (NewLoad->getValueType(NewLoad->getNumValues() - 1) == MVT::Glue)
53155327
NewGlue = NewLoad.getValue(NewLoad->getNumValues() - 1);
53165328
}
@@ -5322,54 +5334,65 @@ static SDValue PerformLoadCombine(SDNode *N,
53225334
// (3) begin rewriting uses
53235335
SmallVector<SDValue> NewOutputsF32;
53245336

5325-
if (ProxyReg) {
5326-
// scalarize proxyreg, but first rewrite all uses of chain and glue from the
5327-
// old load to the new load
5337+
if (llvm::any_of(ProxyRegs, [](const SDNode *PR) { return PR != nullptr; })) {
5338+
// scalarize proxy regs, but first rewrite all uses of chain and glue from
5339+
// the old load to the new load
53285340
DCI.DAG.ReplaceAllUsesOfValueWith(OldChain, NewChain);
53295341
DCI.DAG.ReplaceAllUsesOfValueWith(OldGlue, NewGlue);
53305342

5331-
// Update the new chain and glue to be old inputs to the proxyreg, if they
5332-
// came from an intervening instruction between this proxyreg and the
5333-
// original load (ex: callseq_end). Other than bitcasts and extractelts, we
5334-
// followed all other nodes by chain and glue accesses.
5335-
if (SDValue OldInChain = ProxyReg->getOperand(0); OldInChain.getNode() != N)
5343+
for (unsigned ProxyI = 0, ProxyE = ProxyRegs.size(); ProxyI != ProxyE;
5344+
++ProxyI) {
5345+
SDNode *ProxyReg = ProxyRegs[ProxyI];
5346+
5347+
// no proxy reg might mean this result is unused
5348+
if (!ProxyReg)
5349+
continue;
5350+
5351+
// Update the new chain and glue to be old inputs to the proxyreg, if they
5352+
// came from an intervening instruction between this proxyreg and the
5353+
// original load (ex: callseq_end). Other than bitcasts and extractelts,
5354+
// we followed all other nodes by chain and glue accesses.
5355+
if (SDValue OldInChain = ProxyReg->getOperand(0);
5356+
OldInChain.getNode() != N)
53365357
NewChain = OldInChain;
5337-
if (SDValue OldInGlue = ProxyReg->getOperand(2); OldInGlue.getNode() != N)
5358+
if (SDValue OldInGlue = ProxyReg->getOperand(2); OldInGlue.getNode() != N)
53385359
NewGlue = OldInGlue;
53395360

5340-
// update OldChain, OldGlue to the outputs of ProxyReg, which we will
5341-
// replace later
5342-
OldChain = SDValue(ProxyReg, 1);
5343-
OldGlue = SDValue(ProxyReg, 2);
5344-
5345-
// generate the scalar proxy regs
5346-
for (unsigned I = 0, E = NumElts; I != E; ++I) {
5347-
SDValue ProxyRegElem =
5348-
DCI.DAG.getNode(NVPTXISD::ProxyReg, SDLoc(ProxyReg),
5349-
DCI.DAG.getVTList(MVT::f32, MVT::Other, MVT::Glue),
5350-
{NewChain, NewLoad.getValue(I), NewGlue});
5351-
NewChain = ProxyRegElem.getValue(1);
5352-
NewGlue = ProxyRegElem.getValue(2);
5353-
NewOutputsF32.push_back(ProxyRegElem);
5361+
// update OldChain, OldGlue to the outputs of ProxyReg, which we will
5362+
// replace later
5363+
OldChain = SDValue(ProxyReg, 1);
5364+
OldGlue = SDValue(ProxyReg, 2);
5365+
5366+
// generate the scalar proxy regs
5367+
for (unsigned I = 0, E = 2; I != E; ++I) {
5368+
SDValue ProxyRegElem = DCI.DAG.getNode(
5369+
NVPTXISD::ProxyReg, SDLoc(ProxyReg),
5370+
DCI.DAG.getVTList(MVT::f32, MVT::Other, MVT::Glue),
5371+
{NewChain, NewLoad.getValue(ProxyI * 2 + I), NewGlue});
5372+
NewChain = ProxyRegElem.getValue(1);
5373+
NewGlue = ProxyRegElem.getValue(2);
5374+
NewOutputsF32.push_back(ProxyRegElem);
5375+
}
5376+
5377+
// replace all uses of the glue and chain from the old proxy reg
5378+
DCI.DAG.ReplaceAllUsesOfValueWith(OldChain, NewChain);
5379+
DCI.DAG.ReplaceAllUsesOfValueWith(OldGlue, NewGlue);
53545380
}
53555381
} else {
53565382
for (unsigned I = 0, E = NumElts; I != E; ++I)
53575383
if (NewLoad->getValueType(I) == MVT::f32)
53585384
NewOutputsF32.push_back(NewLoad.getValue(I));
5385+
5386+
// replace all glue and chain nodes
5387+
DCI.DAG.ReplaceAllUsesOfValueWith(OldChain, NewChain);
5388+
if (OldGlue)
5389+
DCI.DAG.ReplaceAllUsesOfValueWith(OldGlue, NewGlue);
53595390
}
53605391

5361-
// now, for all extractelts, replace them with one of the new outputs
5392+
// replace all extractelts with the new outputs
53625393
for (auto &[Extract, Index] : ExtractElts)
53635394
DCI.CombineTo(Extract, NewOutputsF32[Index], false);
53645395

5365-
// now replace all glue and chain nodes
5366-
DCI.DAG.ReplaceAllUsesOfValueWith(OldChain, NewChain);
5367-
if (OldGlue)
5368-
DCI.DAG.ReplaceAllUsesOfValueWith(OldGlue, NewGlue);
5369-
5370-
// cleanup
5371-
if (ProxyReg)
5372-
DCI.recursivelyDeleteUnusedNodes(ProxyReg);
53735396
return SDValue();
53745397
}
53755398

‎llvm/test/CodeGen/NVPTX/aggregate-return.ll

Lines changed: 1 addition & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -27,9 +27,7 @@ define void @test_v3f32(<3 x float> %input, ptr %output) {
2727
; CHECK-NOT: ld.param.f32 [[E3:%f[0-9]+]], [retval0+12];
2828
store <3 x float> %call, ptr %output, align 8
2929
; CHECK-DAG: st.f32 [{{%rd[0-9]}}+8],
30-
; -- This is suboptimal. We should do st.v2.f32 instead
31-
; of combining 2xf32 info i64.
32-
; CHECK-DAG: st.u64 [{{%rd[0-9]}}],
30+
; CHECK-DAG: st.v2.f32 [{{%rd[0-9]}}], {[[E0]], [[E1]]}
3331
; CHECK: ret;
3432
ret void
3533
}

‎llvm/test/CodeGen/NVPTX/bf16-instructions.ll

Lines changed: 48 additions & 48 deletions
Original file line numberDiff line numberDiff line change
@@ -762,32 +762,32 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {
762762
; SM70-NEXT: // %bb.0:
763763
; SM70-NEXT: ld.param.u64 %rd1, [test_extload_bf16x8_param_0];
764764
; SM70-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
765-
; SM70-NEXT: mov.b32 {%rs1, %rs2}, %r1;
766-
; SM70-NEXT: mov.b32 {%rs3, %rs4}, %r2;
767-
; SM70-NEXT: mov.b32 {%rs5, %rs6}, %r3;
768-
; SM70-NEXT: mov.b32 {%rs7, %rs8}, %r4;
769-
; SM70-NEXT: cvt.u32.u16 %r5, %rs8;
765+
; SM70-NEXT: mov.b32 {%rs1, %rs2}, %r4;
766+
; SM70-NEXT: cvt.u32.u16 %r5, %rs2;
770767
; SM70-NEXT: shl.b32 %r6, %r5, 16;
771768
; SM70-NEXT: mov.b32 %f1, %r6;
772-
; SM70-NEXT: cvt.u32.u16 %r7, %rs7;
769+
; SM70-NEXT: cvt.u32.u16 %r7, %rs1;
773770
; SM70-NEXT: shl.b32 %r8, %r7, 16;
774771
; SM70-NEXT: mov.b32 %f2, %r8;
775-
; SM70-NEXT: cvt.u32.u16 %r9, %rs6;
772+
; SM70-NEXT: mov.b32 {%rs3, %rs4}, %r3;
773+
; SM70-NEXT: cvt.u32.u16 %r9, %rs4;
776774
; SM70-NEXT: shl.b32 %r10, %r9, 16;
777775
; SM70-NEXT: mov.b32 %f3, %r10;
778-
; SM70-NEXT: cvt.u32.u16 %r11, %rs5;
776+
; SM70-NEXT: cvt.u32.u16 %r11, %rs3;
779777
; SM70-NEXT: shl.b32 %r12, %r11, 16;
780778
; SM70-NEXT: mov.b32 %f4, %r12;
781-
; SM70-NEXT: cvt.u32.u16 %r13, %rs4;
779+
; SM70-NEXT: mov.b32 {%rs5, %rs6}, %r2;
780+
; SM70-NEXT: cvt.u32.u16 %r13, %rs6;
782781
; SM70-NEXT: shl.b32 %r14, %r13, 16;
783782
; SM70-NEXT: mov.b32 %f5, %r14;
784-
; SM70-NEXT: cvt.u32.u16 %r15, %rs3;
783+
; SM70-NEXT: cvt.u32.u16 %r15, %rs5;
785784
; SM70-NEXT: shl.b32 %r16, %r15, 16;
786785
; SM70-NEXT: mov.b32 %f6, %r16;
787-
; SM70-NEXT: cvt.u32.u16 %r17, %rs2;
786+
; SM70-NEXT: mov.b32 {%rs7, %rs8}, %r1;
787+
; SM70-NEXT: cvt.u32.u16 %r17, %rs8;
788788
; SM70-NEXT: shl.b32 %r18, %r17, 16;
789789
; SM70-NEXT: mov.b32 %f7, %r18;
790-
; SM70-NEXT: cvt.u32.u16 %r19, %rs1;
790+
; SM70-NEXT: cvt.u32.u16 %r19, %rs7;
791791
; SM70-NEXT: shl.b32 %r20, %r19, 16;
792792
; SM70-NEXT: mov.b32 %f8, %r20;
793793
; SM70-NEXT: st.param.v4.f32 [func_retval0], {%f8, %f7, %f6, %f5};
@@ -804,18 +804,18 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {
804804
; SM80-NEXT: // %bb.0:
805805
; SM80-NEXT: ld.param.u64 %rd1, [test_extload_bf16x8_param_0];
806806
; SM80-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
807-
; SM80-NEXT: mov.b32 {%rs1, %rs2}, %r1;
808-
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r2;
809-
; SM80-NEXT: mov.b32 {%rs5, %rs6}, %r3;
810-
; SM80-NEXT: mov.b32 {%rs7, %rs8}, %r4;
811-
; SM80-NEXT: cvt.f32.bf16 %f1, %rs8;
812-
; SM80-NEXT: cvt.f32.bf16 %f2, %rs7;
813-
; SM80-NEXT: cvt.f32.bf16 %f3, %rs6;
814-
; SM80-NEXT: cvt.f32.bf16 %f4, %rs5;
815-
; SM80-NEXT: cvt.f32.bf16 %f5, %rs4;
816-
; SM80-NEXT: cvt.f32.bf16 %f6, %rs3;
817-
; SM80-NEXT: cvt.f32.bf16 %f7, %rs2;
818-
; SM80-NEXT: cvt.f32.bf16 %f8, %rs1;
807+
; SM80-NEXT: mov.b32 {%rs1, %rs2}, %r4;
808+
; SM80-NEXT: cvt.f32.bf16 %f1, %rs2;
809+
; SM80-NEXT: cvt.f32.bf16 %f2, %rs1;
810+
; SM80-NEXT: mov.b32 {%rs3, %rs4}, %r3;
811+
; SM80-NEXT: cvt.f32.bf16 %f3, %rs4;
812+
; SM80-NEXT: cvt.f32.bf16 %f4, %rs3;
813+
; SM80-NEXT: mov.b32 {%rs5, %rs6}, %r2;
814+
; SM80-NEXT: cvt.f32.bf16 %f5, %rs6;
815+
; SM80-NEXT: cvt.f32.bf16 %f6, %rs5;
816+
; SM80-NEXT: mov.b32 {%rs7, %rs8}, %r1;
817+
; SM80-NEXT: cvt.f32.bf16 %f7, %rs8;
818+
; SM80-NEXT: cvt.f32.bf16 %f8, %rs7;
819819
; SM80-NEXT: st.param.v4.f32 [func_retval0], {%f8, %f7, %f6, %f5};
820820
; SM80-NEXT: st.param.v4.f32 [func_retval0+16], {%f4, %f3, %f2, %f1};
821821
; SM80-NEXT: ret;
@@ -830,18 +830,18 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {
830830
; SM80-FTZ-NEXT: // %bb.0:
831831
; SM80-FTZ-NEXT: ld.param.u64 %rd1, [test_extload_bf16x8_param_0];
832832
; SM80-FTZ-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
833-
; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r1;
834-
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r2;
835-
; SM80-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r3;
836-
; SM80-FTZ-NEXT: mov.b32 {%rs7, %rs8}, %r4;
837-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs8;
838-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs7;
839-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f3, %rs6;
840-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs5;
841-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs4;
842-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f6, %rs3;
843-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f7, %rs2;
844-
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f8, %rs1;
833+
; SM80-FTZ-NEXT: mov.b32 {%rs1, %rs2}, %r4;
834+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f1, %rs2;
835+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f2, %rs1;
836+
; SM80-FTZ-NEXT: mov.b32 {%rs3, %rs4}, %r3;
837+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f3, %rs4;
838+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f4, %rs3;
839+
; SM80-FTZ-NEXT: mov.b32 {%rs5, %rs6}, %r2;
840+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f5, %rs6;
841+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f6, %rs5;
842+
; SM80-FTZ-NEXT: mov.b32 {%rs7, %rs8}, %r1;
843+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f7, %rs8;
844+
; SM80-FTZ-NEXT: cvt.ftz.f32.bf16 %f8, %rs7;
845845
; SM80-FTZ-NEXT: st.param.v4.f32 [func_retval0], {%f8, %f7, %f6, %f5};
846846
; SM80-FTZ-NEXT: st.param.v4.f32 [func_retval0+16], {%f4, %f3, %f2, %f1};
847847
; SM80-FTZ-NEXT: ret;
@@ -856,18 +856,18 @@ define <8 x float> @test_extload_bf16x8(ptr addrspace(3) noundef %arg) #0 {
856856
; SM90-NEXT: // %bb.0:
857857
; SM90-NEXT: ld.param.u64 %rd1, [test_extload_bf16x8_param_0];
858858
; SM90-NEXT: ld.shared.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
859-
; SM90-NEXT: mov.b32 {%rs1, %rs2}, %r1;
860-
; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r2;
861-
; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r3;
862-
; SM90-NEXT: mov.b32 {%rs7, %rs8}, %r4;
863-
; SM90-NEXT: cvt.f32.bf16 %f1, %rs8;
864-
; SM90-NEXT: cvt.f32.bf16 %f2, %rs7;
865-
; SM90-NEXT: cvt.f32.bf16 %f3, %rs6;
866-
; SM90-NEXT: cvt.f32.bf16 %f4, %rs5;
867-
; SM90-NEXT: cvt.f32.bf16 %f5, %rs4;
868-
; SM90-NEXT: cvt.f32.bf16 %f6, %rs3;
869-
; SM90-NEXT: cvt.f32.bf16 %f7, %rs2;
870-
; SM90-NEXT: cvt.f32.bf16 %f8, %rs1;
859+
; SM90-NEXT: mov.b32 {%rs1, %rs2}, %r4;
860+
; SM90-NEXT: cvt.f32.bf16 %f1, %rs2;
861+
; SM90-NEXT: cvt.f32.bf16 %f2, %rs1;
862+
; SM90-NEXT: mov.b32 {%rs3, %rs4}, %r3;
863+
; SM90-NEXT: cvt.f32.bf16 %f3, %rs4;
864+
; SM90-NEXT: cvt.f32.bf16 %f4, %rs3;
865+
; SM90-NEXT: mov.b32 {%rs5, %rs6}, %r2;
866+
; SM90-NEXT: cvt.f32.bf16 %f5, %rs6;
867+
; SM90-NEXT: cvt.f32.bf16 %f6, %rs5;
868+
; SM90-NEXT: mov.b32 {%rs7, %rs8}, %r1;
869+
; SM90-NEXT: cvt.f32.bf16 %f7, %rs8;
870+
; SM90-NEXT: cvt.f32.bf16 %f8, %rs7;
871871
; SM90-NEXT: st.param.v4.f32 [func_retval0], {%f8, %f7, %f6, %f5};
872872
; SM90-NEXT: st.param.v4.f32 [func_retval0+16], {%f4, %f3, %f2, %f1};
873873
; SM90-NEXT: ret;

0 commit comments

Comments
 (0)