Skip to content

Commit 2f6f9f0

Browse files
author
git apple-llvm automerger
committed
Merge commit '76c9bfefa416' from llvm.org/main into next
2 parents 9274e0a + 76c9bfe commit 2f6f9f0

File tree

96 files changed

+9137
-7314
lines changed

Some content is hidden

Large Commits have some content hidden by default. Use the searchbox below for content that may be hidden.

96 files changed

+9137
-7314
lines changed

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

Lines changed: 0 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -223,10 +223,6 @@ unsigned NVPTXAsmPrinter::encodeVirtualRegister(unsigned Reg) {
223223
Ret = (3 << 28);
224224
} else if (RC == &NVPTX::Int64RegsRegClass) {
225225
Ret = (4 << 28);
226-
} else if (RC == &NVPTX::Float32RegsRegClass) {
227-
Ret = (5 << 28);
228-
} else if (RC == &NVPTX::Float64RegsRegClass) {
229-
Ret = (6 << 28);
230226
} else if (RC == &NVPTX::Int128RegsRegClass) {
231227
Ret = (7 << 28);
232228
} else {

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

Lines changed: 4 additions & 7 deletions
Original file line numberDiff line numberDiff line change
@@ -606,8 +606,8 @@ NVPTXTargetLowering::NVPTXTargetLowering(const NVPTXTargetMachine &TM,
606606
addRegisterClass(MVT::v4i8, &NVPTX::Int32RegsRegClass);
607607
addRegisterClass(MVT::i32, &NVPTX::Int32RegsRegClass);
608608
addRegisterClass(MVT::i64, &NVPTX::Int64RegsRegClass);
609-
addRegisterClass(MVT::f32, &NVPTX::Float32RegsRegClass);
610-
addRegisterClass(MVT::f64, &NVPTX::Float64RegsRegClass);
609+
addRegisterClass(MVT::f32, &NVPTX::Int32RegsRegClass);
610+
addRegisterClass(MVT::f64, &NVPTX::Int64RegsRegClass);
611611
addRegisterClass(MVT::f16, &NVPTX::Int16RegsRegClass);
612612
addRegisterClass(MVT::v2f16, &NVPTX::Int32RegsRegClass);
613613
addRegisterClass(MVT::bf16, &NVPTX::Int16RegsRegClass);
@@ -4992,24 +4992,21 @@ NVPTXTargetLowering::getRegForInlineAsmConstraint(const TargetRegisterInfo *TRI,
49924992
case 'b':
49934993
return std::make_pair(0U, &NVPTX::Int1RegsRegClass);
49944994
case 'c':
4995-
return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
49964995
case 'h':
49974996
return std::make_pair(0U, &NVPTX::Int16RegsRegClass);
49984997
case 'r':
4998+
case 'f':
49994999
return std::make_pair(0U, &NVPTX::Int32RegsRegClass);
50005000
case 'l':
50015001
case 'N':
5002+
case 'd':
50025003
return std::make_pair(0U, &NVPTX::Int64RegsRegClass);
50035004
case 'q': {
50045005
if (STI.getSmVersion() < 70)
50055006
report_fatal_error("Inline asm with 128 bit operands is only "
50065007
"supported for sm_70 and higher!");
50075008
return std::make_pair(0U, &NVPTX::Int128RegsRegClass);
50085009
}
5009-
case 'f':
5010-
return std::make_pair(0U, &NVPTX::Float32RegsRegClass);
5011-
case 'd':
5012-
return std::make_pair(0U, &NVPTX::Float64RegsRegClass);
50135010
}
50145011
}
50155012
return TargetLowering::getRegForInlineAsmConstraint(TRI, Constraint, VT);

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

Lines changed: 2 additions & 10 deletions
Original file line numberDiff line numberDiff line change
@@ -44,19 +44,11 @@ void NVPTXInstrInfo::copyPhysReg(MachineBasicBlock &MBB,
4444
} else if (DestRC == &NVPTX::Int16RegsRegClass) {
4545
Op = NVPTX::MOV16r;
4646
} else if (DestRC == &NVPTX::Int32RegsRegClass) {
47-
Op = (SrcRC == &NVPTX::Int32RegsRegClass ? NVPTX::IMOV32r
48-
: NVPTX::BITCONVERT_32_F2I);
47+
Op = NVPTX::IMOV32r;
4948
} else if (DestRC == &NVPTX::Int64RegsRegClass) {
50-
Op = (SrcRC == &NVPTX::Int64RegsRegClass ? NVPTX::IMOV64r
51-
: NVPTX::BITCONVERT_64_F2I);
49+
Op = NVPTX::IMOV64r;
5250
} else if (DestRC == &NVPTX::Int128RegsRegClass) {
5351
Op = NVPTX::IMOV128r;
54-
} else if (DestRC == &NVPTX::Float32RegsRegClass) {
55-
Op = (SrcRC == &NVPTX::Float32RegsRegClass ? NVPTX::FMOV32r
56-
: NVPTX::BITCONVERT_32_I2F);
57-
} else if (DestRC == &NVPTX::Float64RegsRegClass) {
58-
Op = (SrcRC == &NVPTX::Float64RegsRegClass ? NVPTX::FMOV64r
59-
: NVPTX::BITCONVERT_64_I2F);
6052
} else {
6153
llvm_unreachable("Bad register copy");
6254
}

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

Lines changed: 0 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -25,10 +25,6 @@ using namespace llvm;
2525

2626
namespace llvm {
2727
StringRef getNVPTXRegClassName(TargetRegisterClass const *RC) {
28-
if (RC == &NVPTX::Float32RegsRegClass)
29-
return ".b32";
30-
if (RC == &NVPTX::Float64RegsRegClass)
31-
return ".b64";
3228
if (RC == &NVPTX::Int128RegsRegClass)
3329
return ".b128";
3430
if (RC == &NVPTX::Int64RegsRegClass)
@@ -63,10 +59,6 @@ StringRef getNVPTXRegClassName(TargetRegisterClass const *RC) {
6359
}
6460

6561
StringRef getNVPTXRegClassStr(TargetRegisterClass const *RC) {
66-
if (RC == &NVPTX::Float32RegsRegClass)
67-
return "%f";
68-
if (RC == &NVPTX::Float64RegsRegClass)
69-
return "%fd";
7062
if (RC == &NVPTX::Int128RegsRegClass)
7163
return "%rq";
7264
if (RC == &NVPTX::Int64RegsRegClass)

‎llvm/lib/Target/NVPTX/NVPTXRegisterInfo.td

Lines changed: 6 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -40,8 +40,6 @@ foreach i = 0...4 in {
4040
def RQ#i : NVPTXReg<"%rq"#i>; // 128-bit
4141
def H#i : NVPTXReg<"%h"#i>; // 16-bit float
4242
def HH#i : NVPTXReg<"%hh"#i>; // 2x16-bit float
43-
def F#i : NVPTXReg<"%f"#i>; // 32-bit float
44-
def FL#i : NVPTXReg<"%fd"#i>; // 64-bit float
4543

4644
// Arguments
4745
def ia#i : NVPTXReg<"%ia"#i>;
@@ -59,14 +57,13 @@ foreach i = 0...31 in {
5957
//===----------------------------------------------------------------------===//
6058
def Int1Regs : NVPTXRegClass<[i1], 8, (add (sequence "P%u", 0, 4))>;
6159
def Int16Regs : NVPTXRegClass<[i16, f16, bf16], 16, (add (sequence "RS%u", 0, 4))>;
62-
def Int32Regs : NVPTXRegClass<[i32, v2f16, v2bf16, v2i16, v4i8], 32,
60+
def Int32Regs : NVPTXRegClass<[i32, v2f16, v2bf16, v2i16, v4i8, f32], 32,
6361
(add (sequence "R%u", 0, 4),
6462
VRFrame32, VRFrameLocal32)>;
65-
def Int64Regs : NVPTXRegClass<[i64], 64, (add (sequence "RL%u", 0, 4), VRFrame64, VRFrameLocal64)>;
63+
def Int64Regs : NVPTXRegClass<[i64, f64], 64, (add (sequence "RL%u", 0, 4), VRFrame64, VRFrameLocal64)>;
6664
// 128-bit regs are not defined as general regs in NVPTX. They are used for inlineASM only.
6765
def Int128Regs : NVPTXRegClass<[i128], 128, (add (sequence "RQ%u", 0, 4))>;
68-
def Float32Regs : NVPTXRegClass<[f32], 32, (add (sequence "F%u", 0, 4))>;
69-
def Float64Regs : NVPTXRegClass<[f64], 64, (add (sequence "FL%u", 0, 4))>;
66+
7067
def Int32ArgRegs : NVPTXRegClass<[i32], 32, (add (sequence "ia%u", 0, 4))>;
7168
def Int64ArgRegs : NVPTXRegClass<[i64], 64, (add (sequence "la%u", 0, 4))>;
7269
def Float32ArgRegs : NVPTXRegClass<[f32], 32, (add (sequence "fa%u", 0, 4))>;
@@ -75,3 +72,6 @@ def Float64ArgRegs : NVPTXRegClass<[f64], 64, (add (sequence "da%u", 0, 4))>;
7572
// Read NVPTXRegisterInfo.cpp to see how VRFrame and VRDepot are used.
7673
def SpecialRegs : NVPTXRegClass<[i32], 32, (add VRFrame32, VRFrameLocal32, VRDepot,
7774
(sequence "ENVREG%u", 0, 31))>;
75+
76+
defvar Float32Regs = Int32Regs;
77+
defvar Float64Regs = Int64Regs;

‎llvm/test/CodeGen/MIR/NVPTX/expected-floating-point-literal.mir

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,8 @@
1212
---
1313
name: test
1414
registers:
15-
- { id: 0, class: float32regs }
16-
- { id: 1, class: float32regs }
15+
- { id: 0, class: int32regs }
16+
- { id: 1, class: int32regs }
1717
body: |
1818
bb.0.entry:
1919
%0 = LD_f32 0, 4, 1, 2, 32, &test_param_0, 0

‎llvm/test/CodeGen/MIR/NVPTX/floating-point-immediate-operands.mir

Lines changed: 18 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -30,24 +30,24 @@
3030
---
3131
name: test
3232
registers:
33-
- { id: 0, class: float32regs }
34-
- { id: 1, class: float64regs }
33+
- { id: 0, class: int32regs }
34+
- { id: 1, class: int64regs }
3535
- { id: 2, class: int32regs }
36-
- { id: 3, class: float64regs }
37-
- { id: 4, class: float32regs }
38-
- { id: 5, class: float32regs }
39-
- { id: 6, class: float32regs }
40-
- { id: 7, class: float32regs }
36+
- { id: 3, class: int64regs }
37+
- { id: 4, class: int32regs }
38+
- { id: 5, class: int32regs }
39+
- { id: 6, class: int32regs }
40+
- { id: 7, class: int32regs }
4141
body: |
4242
bb.0.entry:
4343
%0 = LD_f32 0, 0, 4, 2, 32, &test_param_0, 0
4444
%1 = CVT_f64_f32 %0, 0
4545
%2 = LD_i32 0, 0, 4, 0, 32, &test_param_1, 0
46-
; CHECK: %3:float64regs = FADD_rnf64ri %1, double 3.250000e+00
46+
; CHECK: %3:int64regs = FADD_rnf64ri %1, double 3.250000e+00
4747
%3 = FADD_rnf64ri %1, double 3.250000e+00
4848
%4 = CVT_f32_f64 %3, 5
4949
%5 = CVT_f32_s32 %2, 5
50-
; CHECK: %6:float32regs = FADD_rnf32ri %5, float 6.250000e+00
50+
; CHECK: %6:int32regs = FADD_rnf32ri %5, float 6.250000e+00
5151
%6 = FADD_rnf32ri %5, float 6.250000e+00
5252
%7 = FMUL_rnf32rr %6, %4
5353
StoreRetvalF32 %7, 0
@@ -56,24 +56,24 @@ body: |
5656
---
5757
name: test2
5858
registers:
59-
- { id: 0, class: float32regs }
60-
- { id: 1, class: float64regs }
59+
- { id: 0, class: int32regs }
60+
- { id: 1, class: int64regs }
6161
- { id: 2, class: int32regs }
62-
- { id: 3, class: float64regs }
63-
- { id: 4, class: float32regs }
64-
- { id: 5, class: float32regs }
65-
- { id: 6, class: float32regs }
66-
- { id: 7, class: float32regs }
62+
- { id: 3, class: int64regs }
63+
- { id: 4, class: int32regs }
64+
- { id: 5, class: int32regs }
65+
- { id: 6, class: int32regs }
66+
- { id: 7, class: int32regs }
6767
body: |
6868
bb.0.entry:
6969
%0 = LD_f32 0, 0, 4, 2, 32, &test2_param_0, 0
7070
%1 = CVT_f64_f32 %0, 0
7171
%2 = LD_i32 0, 0, 4, 0, 32, &test2_param_1, 0
72-
; CHECK: %3:float64regs = FADD_rnf64ri %1, double 0x7FF8000000000000
72+
; CHECK: %3:int64regs = FADD_rnf64ri %1, double 0x7FF8000000000000
7373
%3 = FADD_rnf64ri %1, double 0x7FF8000000000000
7474
%4 = CVT_f32_f64 %3, 5
7575
%5 = CVT_f32_s32 %2, 5
76-
; CHECK: %6:float32regs = FADD_rnf32ri %5, float 0x7FF8000000000000
76+
; CHECK: %6:int32regs = FADD_rnf32ri %5, float 0x7FF8000000000000
7777
%6 = FADD_rnf32ri %5, float 0x7FF8000000000000
7878
%7 = FMUL_rnf32rr %6, %4
7979
StoreRetvalF32 %7, 0

‎llvm/test/CodeGen/MIR/NVPTX/floating-point-invalid-type-error.mir

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,8 +12,8 @@
1212
---
1313
name: test
1414
registers:
15-
- { id: 0, class: float32regs }
16-
- { id: 1, class: float32regs }
15+
- { id: 0, class: int32regs }
16+
- { id: 1, class: int32regs }
1717
body: |
1818
bb.0.entry:
1919
%0 = LD_f32 0, 4, 1, 2, 32, &test_param_0, 0

‎llvm/test/CodeGen/NVPTX/LoadStoreVectorizer.ll

Lines changed: 53 additions & 53 deletions
Original file line numberDiff line numberDiff line change
@@ -45,36 +45,36 @@ define half @fh(ptr %p) {
4545
; ENABLED-LABEL: fh(
4646
; ENABLED: {
4747
; ENABLED-NEXT: .reg .b16 %rs<10>;
48-
; ENABLED-NEXT: .reg .b32 %f<13>;
48+
; ENABLED-NEXT: .reg .b32 %r<13>;
4949
; ENABLED-NEXT: .reg .b64 %rd<2>;
5050
; ENABLED-EMPTY:
5151
; ENABLED-NEXT: // %bb.0:
5252
; ENABLED-NEXT: ld.param.b64 %rd1, [fh_param_0];
5353
; ENABLED-NEXT: ld.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [%rd1];
5454
; ENABLED-NEXT: ld.b16 %rs5, [%rd1+8];
55-
; ENABLED-NEXT: cvt.f32.f16 %f1, %rs2;
56-
; ENABLED-NEXT: cvt.f32.f16 %f2, %rs1;
57-
; ENABLED-NEXT: add.rn.f32 %f3, %f2, %f1;
58-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs6, %f3;
59-
; ENABLED-NEXT: cvt.f32.f16 %f4, %rs4;
60-
; ENABLED-NEXT: cvt.f32.f16 %f5, %rs3;
61-
; ENABLED-NEXT: add.rn.f32 %f6, %f5, %f4;
62-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs7, %f6;
63-
; ENABLED-NEXT: cvt.f32.f16 %f7, %rs7;
64-
; ENABLED-NEXT: cvt.f32.f16 %f8, %rs6;
65-
; ENABLED-NEXT: add.rn.f32 %f9, %f8, %f7;
66-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs8, %f9;
67-
; ENABLED-NEXT: cvt.f32.f16 %f10, %rs8;
68-
; ENABLED-NEXT: cvt.f32.f16 %f11, %rs5;
69-
; ENABLED-NEXT: add.rn.f32 %f12, %f10, %f11;
70-
; ENABLED-NEXT: cvt.rn.f16.f32 %rs9, %f12;
55+
; ENABLED-NEXT: cvt.f32.f16 %r1, %rs2;
56+
; ENABLED-NEXT: cvt.f32.f16 %r2, %rs1;
57+
; ENABLED-NEXT: add.rn.f32 %r3, %r2, %r1;
58+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;
59+
; ENABLED-NEXT: cvt.f32.f16 %r4, %rs4;
60+
; ENABLED-NEXT: cvt.f32.f16 %r5, %rs3;
61+
; ENABLED-NEXT: add.rn.f32 %r6, %r5, %r4;
62+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;
63+
; ENABLED-NEXT: cvt.f32.f16 %r7, %rs7;
64+
; ENABLED-NEXT: cvt.f32.f16 %r8, %rs6;
65+
; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r7;
66+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;
67+
; ENABLED-NEXT: cvt.f32.f16 %r10, %rs8;
68+
; ENABLED-NEXT: cvt.f32.f16 %r11, %rs5;
69+
; ENABLED-NEXT: add.rn.f32 %r12, %r10, %r11;
70+
; ENABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;
7171
; ENABLED-NEXT: st.param.b16 [func_retval0], %rs9;
7272
; ENABLED-NEXT: ret;
7373
;
7474
; DISABLED-LABEL: fh(
7575
; DISABLED: {
7676
; DISABLED-NEXT: .reg .b16 %rs<10>;
77-
; DISABLED-NEXT: .reg .b32 %f<13>;
77+
; DISABLED-NEXT: .reg .b32 %r<13>;
7878
; DISABLED-NEXT: .reg .b64 %rd<2>;
7979
; DISABLED-EMPTY:
8080
; DISABLED-NEXT: // %bb.0:
@@ -84,22 +84,22 @@ define half @fh(ptr %p) {
8484
; DISABLED-NEXT: ld.b16 %rs3, [%rd1+4];
8585
; DISABLED-NEXT: ld.b16 %rs4, [%rd1+6];
8686
; DISABLED-NEXT: ld.b16 %rs5, [%rd1+8];
87-
; DISABLED-NEXT: cvt.f32.f16 %f1, %rs2;
88-
; DISABLED-NEXT: cvt.f32.f16 %f2, %rs1;
89-
; DISABLED-NEXT: add.rn.f32 %f3, %f2, %f1;
90-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs6, %f3;
91-
; DISABLED-NEXT: cvt.f32.f16 %f4, %rs4;
92-
; DISABLED-NEXT: cvt.f32.f16 %f5, %rs3;
93-
; DISABLED-NEXT: add.rn.f32 %f6, %f5, %f4;
94-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs7, %f6;
95-
; DISABLED-NEXT: cvt.f32.f16 %f7, %rs7;
96-
; DISABLED-NEXT: cvt.f32.f16 %f8, %rs6;
97-
; DISABLED-NEXT: add.rn.f32 %f9, %f8, %f7;
98-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs8, %f9;
99-
; DISABLED-NEXT: cvt.f32.f16 %f10, %rs8;
100-
; DISABLED-NEXT: cvt.f32.f16 %f11, %rs5;
101-
; DISABLED-NEXT: add.rn.f32 %f12, %f10, %f11;
102-
; DISABLED-NEXT: cvt.rn.f16.f32 %rs9, %f12;
87+
; DISABLED-NEXT: cvt.f32.f16 %r1, %rs2;
88+
; DISABLED-NEXT: cvt.f32.f16 %r2, %rs1;
89+
; DISABLED-NEXT: add.rn.f32 %r3, %r2, %r1;
90+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs6, %r3;
91+
; DISABLED-NEXT: cvt.f32.f16 %r4, %rs4;
92+
; DISABLED-NEXT: cvt.f32.f16 %r5, %rs3;
93+
; DISABLED-NEXT: add.rn.f32 %r6, %r5, %r4;
94+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs7, %r6;
95+
; DISABLED-NEXT: cvt.f32.f16 %r7, %rs7;
96+
; DISABLED-NEXT: cvt.f32.f16 %r8, %rs6;
97+
; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r7;
98+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs8, %r9;
99+
; DISABLED-NEXT: cvt.f32.f16 %r10, %rs8;
100+
; DISABLED-NEXT: cvt.f32.f16 %r11, %rs5;
101+
; DISABLED-NEXT: add.rn.f32 %r12, %r10, %r11;
102+
; DISABLED-NEXT: cvt.rn.f16.f32 %rs9, %r12;
103103
; DISABLED-NEXT: st.param.b16 [func_retval0], %rs9;
104104
; DISABLED-NEXT: ret;
105105
%p.1 = getelementptr half, ptr %p, i32 1
@@ -121,37 +121,37 @@ define half @fh(ptr %p) {
121121
define float @ff(ptr %p) {
122122
; ENABLED-LABEL: ff(
123123
; ENABLED: {
124-
; ENABLED-NEXT: .reg .b32 %f<10>;
124+
; ENABLED-NEXT: .reg .b32 %r<10>;
125125
; ENABLED-NEXT: .reg .b64 %rd<2>;
126126
; ENABLED-EMPTY:
127127
; ENABLED-NEXT: // %bb.0:
128128
; ENABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];
129-
; ENABLED-NEXT: ld.v4.b32 {%f1, %f2, %f3, %f4}, [%rd1];
130-
; ENABLED-NEXT: ld.b32 %f5, [%rd1+16];
131-
; ENABLED-NEXT: add.rn.f32 %f6, %f1, %f2;
132-
; ENABLED-NEXT: add.rn.f32 %f7, %f3, %f4;
133-
; ENABLED-NEXT: add.rn.f32 %f8, %f6, %f7;
134-
; ENABLED-NEXT: add.rn.f32 %f9, %f8, %f5;
135-
; ENABLED-NEXT: st.param.b32 [func_retval0], %f9;
129+
; ENABLED-NEXT: ld.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
130+
; ENABLED-NEXT: ld.b32 %r5, [%rd1+16];
131+
; ENABLED-NEXT: add.rn.f32 %r6, %r1, %r2;
132+
; ENABLED-NEXT: add.rn.f32 %r7, %r3, %r4;
133+
; ENABLED-NEXT: add.rn.f32 %r8, %r6, %r7;
134+
; ENABLED-NEXT: add.rn.f32 %r9, %r8, %r5;
135+
; ENABLED-NEXT: st.param.b32 [func_retval0], %r9;
136136
; ENABLED-NEXT: ret;
137137
;
138138
; DISABLED-LABEL: ff(
139139
; DISABLED: {
140-
; DISABLED-NEXT: .reg .b32 %f<10>;
140+
; DISABLED-NEXT: .reg .b32 %r<10>;
141141
; DISABLED-NEXT: .reg .b64 %rd<2>;
142142
; DISABLED-EMPTY:
143143
; DISABLED-NEXT: // %bb.0:
144144
; DISABLED-NEXT: ld.param.b64 %rd1, [ff_param_0];
145-
; DISABLED-NEXT: ld.b32 %f1, [%rd1];
146-
; DISABLED-NEXT: ld.b32 %f2, [%rd1+4];
147-
; DISABLED-NEXT: ld.b32 %f3, [%rd1+8];
148-
; DISABLED-NEXT: ld.b32 %f4, [%rd1+12];
149-
; DISABLED-NEXT: ld.b32 %f5, [%rd1+16];
150-
; DISABLED-NEXT: add.rn.f32 %f6, %f1, %f2;
151-
; DISABLED-NEXT: add.rn.f32 %f7, %f3, %f4;
152-
; DISABLED-NEXT: add.rn.f32 %f8, %f6, %f7;
153-
; DISABLED-NEXT: add.rn.f32 %f9, %f8, %f5;
154-
; DISABLED-NEXT: st.param.b32 [func_retval0], %f9;
145+
; DISABLED-NEXT: ld.b32 %r1, [%rd1];
146+
; DISABLED-NEXT: ld.b32 %r2, [%rd1+4];
147+
; DISABLED-NEXT: ld.b32 %r3, [%rd1+8];
148+
; DISABLED-NEXT: ld.b32 %r4, [%rd1+12];
149+
; DISABLED-NEXT: ld.b32 %r5, [%rd1+16];
150+
; DISABLED-NEXT: add.rn.f32 %r6, %r1, %r2;
151+
; DISABLED-NEXT: add.rn.f32 %r7, %r3, %r4;
152+
; DISABLED-NEXT: add.rn.f32 %r8, %r6, %r7;
153+
; DISABLED-NEXT: add.rn.f32 %r9, %r8, %r5;
154+
; DISABLED-NEXT: st.param.b32 [func_retval0], %r9;
155155
; DISABLED-NEXT: ret;
156156
%p.1 = getelementptr float, ptr %p, i32 1
157157
%p.2 = getelementptr float, ptr %p, i32 2

0 commit comments

Comments
 (0)