[llvm] d63ab54 - [NVPTX] don't erase CopyToRegs when folding movs into loads (#149393)
via llvm-commits
llvm-commits at lists.llvm.org
Fri Jul 18 14:11:35 PDT 2025
Author: Princeton Ferro
Date: 2025-07-18T14:11:31-07:00
New Revision: d63ab5467dcae0492e2f4def336ddbb73ce10dc5
URL: https://github.com/llvm/llvm-project/commit/d63ab5467dcae0492e2f4def336ddbb73ce10dc5
DIFF: https://github.com/llvm/llvm-project/commit/d63ab5467dcae0492e2f4def336ddbb73ce10dc5.diff
LOG: [NVPTX] don't erase CopyToRegs when folding movs into loads (#149393)
We may still need to keep CopyToReg even after folding uses into vector
loads, since the original register may be used in other blocks.
Partially reverts 1fdbe6984976d9e85ab3b1a93e8de434a85c5646
Added:
llvm/test/CodeGen/NVPTX/pr126337.ll
Modified:
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
Removed:
################################################################################
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7aa06f9079b09..31b236a6126ad 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5060,12 +5060,6 @@ combineUnpackingMovIntoLoad(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
return !U.getUser()->use_empty();
}
- // Handle CopyToReg nodes that will become dead after our replacement
- if (U.getUser()->getOpcode() == ISD::CopyToReg) {
- DeadCopyToRegs.push_back(U.getUser());
- return true;
- }
-
// Otherwise, this use prevents us from splitting a value.
return false;
}))
@@ -5132,10 +5126,6 @@ combineUnpackingMovIntoLoad(SDNode *N, TargetLowering::DAGCombinerInfo &DCI) {
for (unsigned I : seq(NewLoad->getNumValues() - NewNumOutputs))
Results.push_back(NewLoad.getValue(NewNumOutputs + I));
- // Remove dead CopyToReg nodes by folding them into the chain they reference
- for (SDNode *CTR : DeadCopyToRegs)
- DCI.CombineTo(CTR, CTR->getOperand(0));
-
return DCI.DAG.getMergeValues(Results, DL);
}
@@ -6544,4 +6534,4 @@ void NVPTXTargetLowering::computeKnownBitsForTargetNode(
default:
break;
}
-}
\ No newline at end of file
+}
diff --git a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
index d0e2c1817f696..3baefde072be7 100644
--- a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
@@ -45,11 +45,12 @@ define <2 x half> @test_ret_const() #0 {
define half @test_extract_0(<2 x half> %a) #0 {
; CHECK-LABEL: test_extract_0(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_0_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];
+; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-NEXT: ret;
%e = extractelement <2 x half> %a, i32 0
@@ -59,12 +60,13 @@ define half @test_extract_0(<2 x half> %a) #0 {
define half @test_extract_1(<2 x half> %a) #0 {
; CHECK-LABEL: test_extract_1(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<3>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_1_param_0];
-; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
+; CHECK-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];
+; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }
+; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-NEXT: ret;
%e = extractelement <2 x half> %a, i32 1
ret half %e
@@ -80,8 +82,9 @@ define half @test_extract_i(<2 x half> %a, i64 %idx) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_i_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];
; CHECK-NEXT: setp.eq.b64 %p1, %rd1, 0;
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
; CHECK-NEXT: ret;
@@ -107,14 +110,16 @@ define <2 x half> @test_fadd(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<10>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fadd_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fadd_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fadd_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: add.rn.f32 %r5, %r4, %r3;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NOF16-NEXT: add.rn.f32 %r8, %r7, %r6;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -143,7 +148,8 @@ define <2 x half> @test_fadd_imm_0(<2 x half> %a) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fadd_imm_0_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_0_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: add.rn.f32 %r3, %r2, 0f40000000;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -175,7 +181,8 @@ define <2 x half> @test_fadd_imm_1(<2 x half> %a) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fadd_imm_1_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fadd_imm_1_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: add.rn.f32 %r3, %r2, 0f40000000;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -207,14 +214,16 @@ define <2 x half> @test_fsub(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<10>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fsub_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fsub_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fsub_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fsub_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: sub.rn.f32 %r5, %r4, %r3;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NOF16-NEXT: sub.rn.f32 %r8, %r7, %r6;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -242,7 +251,8 @@ define <2 x half> @test_fneg(<2 x half> %a) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<8>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fneg_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fneg_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: mov.b32 %r3, 0f00000000;
; CHECK-NOF16-NEXT: sub.rn.f32 %r4, %r3, %r2;
@@ -275,14 +285,16 @@ define <2 x half> @test_fmul(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<10>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmul_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmul_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fmul_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fmul_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: mul.rn.f32 %r5, %r4, %r3;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NOF16-NEXT: mul.rn.f32 %r8, %r7, %r6;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NOF16-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -299,14 +311,16 @@ define <2 x half> @test_fdiv(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];
-; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];
-; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NEXT: ld.param.b32 %r2, [test_fdiv_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fdiv_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;
; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NEXT: div.rn.f32 %r8, %r7, %r6;
; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -331,10 +345,12 @@ define <2 x half> @test_frem(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NEXT: .reg .b32 %r<18>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_frem_param_0];
-; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_frem_param_1];
-; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NEXT: ld.param.b32 %r2, [test_frem_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [test_frem_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;
; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;
; CHECK-NEXT: neg.f32 %r7, %r6;
@@ -342,8 +358,8 @@ define <2 x half> @test_frem(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NEXT: testp.infinite.f32 %p1, %r3;
; CHECK-NEXT: selp.f32 %r9, %r4, %r8, %p1;
; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r9;
-; CHECK-NEXT: cvt.f32.f16 %r10, %rs3;
-; CHECK-NEXT: cvt.f32.f16 %r11, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r10, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r11, %rs3;
; CHECK-NEXT: div.rn.f32 %r12, %r11, %r10;
; CHECK-NEXT: cvt.rzi.f32.f32 %r13, %r12;
; CHECK-NEXT: neg.f32 %r14, %r13;
@@ -535,11 +551,13 @@ define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, <
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
; CHECK-F16-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];
-; CHECK-F16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];
+; CHECK-F16-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];
+; CHECK-F16-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r3, %r4;
-; CHECK-F16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_1];
-; CHECK-F16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;
-; CHECK-F16-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;
+; CHECK-F16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-F16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-F16-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;
+; CHECK-F16-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;
; CHECK-F16-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; CHECK-F16-NEXT: ret;
;
@@ -550,18 +568,22 @@ define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, <
; CHECK-NOF16-NEXT: .reg .b32 %r<9>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_3];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_select_cc_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs5;
+; CHECK-NOF16-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
+; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r4;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r6, %r5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs6;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs4;
; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r8, %r7;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs7, %rs8}, [test_select_cc_param_1];
-; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs2, %rs8, %p2;
-; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs1, %rs7, %p1;
+; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; CHECK-NOF16-NEXT: selp.b16 %rs9, %rs8, %rs6, %p2;
+; CHECK-NOF16-NEXT: selp.b16 %rs10, %rs7, %rs5, %p1;
; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
; CHECK-NOF16-NEXT: ret;
%cc = fcmp une <2 x half> %c, %d
@@ -579,11 +601,13 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b,
; CHECK-F16-NEXT: // %bb.0:
; CHECK-F16-NEXT: ld.param.b32 %r2, [test_select_cc_f32_f16_param_3];
; CHECK-F16-NEXT: ld.param.b32 %r1, [test_select_cc_f32_f16_param_2];
-; CHECK-F16-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f16_param_0];
+; CHECK-F16-NEXT: ld.param.b64 %rd2, [test_select_cc_f32_f16_param_1];
+; CHECK-F16-NEXT: ld.param.b64 %rd1, [test_select_cc_f32_f16_param_0];
; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r1, %r2;
-; CHECK-F16-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f32_f16_param_1];
-; CHECK-F16-NEXT: selp.f32 %r7, %r4, %r6, %p2;
-; CHECK-F16-NEXT: selp.f32 %r8, %r3, %r5, %p1;
+; CHECK-F16-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-F16-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-F16-NEXT: selp.f32 %r7, %r6, %r4, %p2;
+; CHECK-F16-NEXT: selp.f32 %r8, %r5, %r3, %p1;
; CHECK-F16-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-F16-NEXT: ret;
;
@@ -595,18 +619,22 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b,
; CHECK-NOF16-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f16_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_f32_f16_param_3];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_f32_f16_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r6, %r5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs2;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs4;
-; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r8, %r7;
-; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r9, %r10}, [test_select_cc_f32_f16_param_1];
-; CHECK-NOF16-NEXT: selp.f32 %r11, %r4, %r10, %p2;
-; CHECK-NOF16-NEXT: selp.f32 %r12, %r3, %r9, %p1;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_select_cc_f32_f16_param_3];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_select_cc_f32_f16_param_2];
+; CHECK-NOF16-NEXT: ld.param.b64 %rd2, [test_select_cc_f32_f16_param_1];
+; CHECK-NOF16-NEXT: ld.param.b64 %rd1, [test_select_cc_f32_f16_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs1;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs3;
+; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r4, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs4;
+; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r6, %r5;
+; CHECK-NOF16-NEXT: mov.b64 {%r7, %r8}, %rd2;
+; CHECK-NOF16-NEXT: mov.b64 {%r9, %r10}, %rd1;
+; CHECK-NOF16-NEXT: selp.f32 %r11, %r10, %r8, %p2;
+; CHECK-NOF16-NEXT: selp.f32 %r12, %r9, %r7, %p1;
; CHECK-NOF16-NEXT: st.param.v2.b32 [func_retval0], {%r12, %r11};
; CHECK-NOF16-NEXT: ret;
<2 x half> %c, <2 x half> %d) #0 {
@@ -624,14 +652,18 @@ define <2 x half> @test_select_cc_f16_f32(<2 x half> %a, <2 x half> %b,
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_f16_f32_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f16_f32_param_2];
-; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f16_f32_param_3];
-; CHECK-NEXT: setp.neu.f32 %p1, %r3, %r5;
-; CHECK-NEXT: setp.neu.f32 %p2, %r4, %r6;
-; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_f16_f32_param_1];
-; CHECK-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;
-; CHECK-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_select_cc_f16_f32_param_3];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_select_cc_f16_f32_param_2];
+; CHECK-NEXT: ld.param.b32 %r2, [test_select_cc_f16_f32_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [test_select_cc_f16_f32_param_0];
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NEXT: setp.neu.f32 %p1, %r5, %r3;
+; CHECK-NEXT: setp.neu.f32 %p2, %r6, %r4;
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;
+; CHECK-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; CHECK-NEXT: ret;
<2 x float> %c, <2 x float> %d) #0 {
@@ -664,13 +696,15 @@ define <2 x i1> @test_fcmp_une(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_une_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_une_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_une_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_une_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -705,13 +739,15 @@ define <2 x i1> @test_fcmp_ueq(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ueq_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ueq_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ueq_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ueq_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.equ.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.equ.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -746,13 +782,15 @@ define <2 x i1> @test_fcmp_ugt(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ugt_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ugt_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ugt_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ugt_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.gtu.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.gtu.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -787,13 +825,15 @@ define <2 x i1> @test_fcmp_uge(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_uge_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_uge_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uge_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_uge_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.geu.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.geu.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -828,13 +868,15 @@ define <2 x i1> @test_fcmp_ult(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ult_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ult_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ult_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ult_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.ltu.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.ltu.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -869,13 +911,15 @@ define <2 x i1> @test_fcmp_ule(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ule_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ule_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ule_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ule_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.leu.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.leu.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -911,13 +955,15 @@ define <2 x i1> @test_fcmp_uno(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_uno_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_uno_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_uno_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_uno_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.nan.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -952,13 +998,15 @@ define <2 x i1> @test_fcmp_one(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_one_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_one_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_one_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_one_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.ne.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.ne.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -993,13 +1041,15 @@ define <2 x i1> @test_fcmp_oeq(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_oeq_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_oeq_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oeq_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_oeq_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.eq.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.eq.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1034,13 +1084,15 @@ define <2 x i1> @test_fcmp_ogt(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ogt_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ogt_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ogt_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ogt_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.gt.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1075,13 +1127,15 @@ define <2 x i1> @test_fcmp_oge(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_oge_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_oge_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_oge_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_oge_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.ge.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.ge.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1116,13 +1170,15 @@ define <2 x i1> @test_fcmp_olt(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_olt_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_olt_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_olt_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_olt_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.lt.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1157,13 +1213,15 @@ define <2 x i1> @test_fcmp_ole(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ole_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ole_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ole_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ole_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.le.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.le.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1198,13 +1256,15 @@ define <2 x i1> @test_fcmp_ord(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fcmp_ord_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ord_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fcmp_ord_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fcmp_ord_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: setp.num.f32 %p1, %r4, %r3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs3;
; CHECK-NOF16-NEXT: setp.num.f32 %p2, %r6, %r5;
; CHECK-NOF16-NEXT: selp.b16 %rs5, -1, 0, %p2;
; CHECK-NOF16-NEXT: st.param.b8 [func_retval0], %rs5;
@@ -1222,7 +1282,8 @@ define <2 x i32> @test_fptosi_i32(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_i32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fptosi_i32_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rzi.s32.f16 %r2, %rs2;
; CHECK-NEXT: cvt.rzi.s32.f16 %r3, %rs1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -1239,7 +1300,8 @@ define <2 x i64> @test_fptosi_i64(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_i64_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fptosi_i64_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rzi.s64.f16 %rd1, %rs2;
; CHECK-NEXT: cvt.rzi.s64.f16 %rd2, %rs1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -1255,7 +1317,8 @@ define <2 x i32> @test_fptoui_2xi32(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rzi.u32.f16 %r2, %rs2;
; CHECK-NEXT: cvt.rzi.u32.f16 %r3, %rs1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -1272,7 +1335,8 @@ define <2 x i64> @test_fptoui_2xi64(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rzi.u64.f16 %rd1, %rs2;
; CHECK-NEXT: cvt.rzi.u64.f16 %rd2, %rs1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -1369,16 +1433,17 @@ define <2 x half> @test_uitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_uitofp_2xi32_fadd_param_1];
; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs1, %r1;
; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs2, %r2;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_uitofp_2xi32_fadd_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs2;
-; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r4, %r5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
+; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r5, %r4;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r6;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;
-; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r7, %r8;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r8, %r7;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;
; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};
; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r10;
@@ -1411,16 +1476,17 @@ define <2 x half> @test_sitofp_2xi32_fadd(<2 x i32> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sitofp_2xi32_fadd_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_sitofp_2xi32_fadd_param_1];
; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs1, %r1;
; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs2, %r2;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_sitofp_2xi32_fadd_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs2;
-; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r4, %r5;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
+; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r5, %r4;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs5, %r6;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;
-; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r7, %r8;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r8, %r7;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;
; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};
; CHECK-NOF16-NEXT: st.param.b32 [func_retval0], %r10;
@@ -1468,7 +1534,8 @@ define <2 x float> @test_fpext_2xfloat(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fpext_2xfloat_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xfloat_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -1485,7 +1552,8 @@ define <2 x double> @test_fpext_2xdouble(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f64.f16 %rd1, %rs2;
; CHECK-NEXT: cvt.f64.f16 %rd2, %rs1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -1578,7 +1646,8 @@ define <2 x half> @test_sqrt(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sqrt_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_sqrt_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: sqrt.rn.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1606,7 +1675,8 @@ define <2 x half> @test_sin(<2 x half> %a) #0 #1 {
; CHECK-NEXT: .reg .b32 %r<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sin_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_sin_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: sin.approx.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1627,7 +1697,8 @@ define <2 x half> @test_cos(<2 x half> %a) #0 #1 {
; CHECK-NEXT: .reg .b32 %r<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_cos_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_cos_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: cos.approx.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1703,17 +1774,20 @@ define <2 x half> @test_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<13>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fma_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fma_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_fma_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs6;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fma_param_2];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fma_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fma_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
+; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs6;
; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs7, %r7;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs5;
; CHECK-NOF16-NEXT: fma.rn.f32 %r11, %r10, %r9, %r8;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs8, %r11;
; CHECK-NOF16-NEXT: mov.b32 %r12, {%rs8, %rs7};
@@ -1740,7 +1814,8 @@ define <2 x half> @test_fabs(<2 x half> %a) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<7>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fabs_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fabs_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: abs.f32 %r3, %r2;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1761,14 +1836,16 @@ define <2 x half> @test_minnum(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_minnum_param_0];
-; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_minnum_param_1];
-; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NEXT: ld.param.b32 %r2, [test_minnum_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [test_minnum_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NEXT: min.f32 %r5, %r4, %r3;
; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NEXT: min.f32 %r8, %r7, %r6;
; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -1785,14 +1862,16 @@ define <2 x half> @test_maxnum(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NEXT: .reg .b32 %r<10>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_maxnum_param_0];
-; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_maxnum_param_1];
-; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
-; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NEXT: ld.param.b32 %r2, [test_maxnum_param_1];
+; CHECK-NEXT: ld.param.b32 %r1, [test_maxnum_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NEXT: cvt.f32.f16 %r3, %rs2;
+; CHECK-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; CHECK-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NEXT: max.f32 %r5, %r4, %r3;
; CHECK-NEXT: cvt.rn.f16.f32 %rs5, %r5;
-; CHECK-NEXT: cvt.f32.f16 %r6, %rs3;
-; CHECK-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r6, %rs1;
+; CHECK-NEXT: cvt.f32.f16 %r7, %rs3;
; CHECK-NEXT: max.f32 %r8, %r7, %r6;
; CHECK-NEXT: cvt.rn.f16.f32 %rs6, %r8;
; CHECK-NEXT: mov.b32 %r9, {%rs6, %rs5};
@@ -1822,13 +1901,15 @@ define <2 x half> @test_copysign(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<3>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_copysign_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_copysign_param_1];
-; CHECK-NOF16-NEXT: and.b16 %rs5, %rs4, -32768;
-; CHECK-NOF16-NEXT: and.b16 %rs6, %rs2, 32767;
-; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs5;
-; CHECK-NOF16-NEXT: and.b16 %rs8, %rs3, -32768;
-; CHECK-NOF16-NEXT: and.b16 %rs9, %rs1, 32767;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_copysign_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, -32768;
+; CHECK-NOF16-NEXT: mov.b32 {%rs4, %rs5}, %r1;
+; CHECK-NOF16-NEXT: and.b16 %rs6, %rs5, 32767;
+; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs3;
+; CHECK-NOF16-NEXT: and.b16 %rs8, %rs1, -32768;
+; CHECK-NOF16-NEXT: and.b16 %rs9, %rs4, 32767;
; CHECK-NOF16-NEXT: or.b16 %rs10, %rs9, %rs8;
; CHECK-NOF16-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs7};
; CHECK-NOF16-NEXT: ret;
@@ -1844,8 +1925,9 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 {
; CHECK-F16-NEXT: .reg .b64 %rd<2>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.b64 %rd1, [test_copysign_f32_param_1];
; CHECK-F16-NEXT: ld.param.b32 %r1, [test_copysign_f32_param_0];
-; CHECK-F16-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_copysign_f32_param_1];
+; CHECK-F16-NEXT: mov.b64 {%r2, %r3}, %rd1;
; CHECK-F16-NEXT: cvt.rn.f16.f32 %rs1, %r3;
; CHECK-F16-NEXT: cvt.rn.f16.f32 %rs2, %r2;
; CHECK-F16-NEXT: mov.b32 %r4, {%rs2, %rs1};
@@ -1862,8 +1944,10 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b64 %rd<2>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_copysign_f32_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_copysign_f32_param_1];
+; CHECK-NOF16-NEXT: ld.param.b64 %rd1, [test_copysign_f32_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_f32_param_0];
+; CHECK-NOF16-NEXT: mov.b64 {%r2, %r3}, %rd1;
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, 32767;
; CHECK-NOF16-NEXT: and.b32 %r4, %r3, -2147483648;
; CHECK-NOF16-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs4}, %r4; }
@@ -1906,7 +1990,8 @@ define <2 x half> @test_copysign_f64(<2 x half> %a, <2 x double> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_copysign_f64_param_1];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_copysign_f64_param_0];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_f64_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, 32767;
; CHECK-NOF16-NEXT: and.b64 %rd3, %rd2, -9223372036854775808;
; CHECK-NOF16-NEXT: shr.u64 %rd4, %rd3, 48;
@@ -1948,13 +2033,15 @@ define <2 x float> @test_copysign_extended(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b32 %r<5>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_copysign_extended_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_copysign_extended_param_1];
-; CHECK-NOF16-NEXT: and.b16 %rs5, %rs3, -32768;
-; CHECK-NOF16-NEXT: and.b16 %rs6, %rs1, 32767;
-; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs5;
-; CHECK-NOF16-NEXT: and.b16 %rs8, %rs4, -32768;
-; CHECK-NOF16-NEXT: and.b16 %rs9, %rs2, 32767;
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_copysign_extended_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_copysign_extended_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; CHECK-NOF16-NEXT: and.b16 %rs3, %rs1, -32768;
+; CHECK-NOF16-NEXT: mov.b32 {%rs4, %rs5}, %r1;
+; CHECK-NOF16-NEXT: and.b16 %rs6, %rs4, 32767;
+; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs3;
+; CHECK-NOF16-NEXT: and.b16 %rs8, %rs2, -32768;
+; CHECK-NOF16-NEXT: and.b16 %rs9, %rs5, 32767;
; CHECK-NOF16-NEXT: or.b16 %rs10, %rs9, %rs8;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs10;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs7;
@@ -1972,7 +2059,8 @@ define <2 x half> @test_floor(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_floor_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_floor_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rmi.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rmi.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -1988,7 +2076,8 @@ define <2 x half> @test_ceil(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_ceil_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_ceil_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rpi.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rpi.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2004,7 +2093,8 @@ define <2 x half> @test_trunc(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_trunc_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_trunc_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rzi.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rzi.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2020,7 +2110,8 @@ define <2 x half> @test_rint(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_rint_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_rint_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2036,7 +2127,8 @@ define <2 x half> @test_nearbyint(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_nearbyint_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_nearbyint_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2052,7 +2144,8 @@ define <2 x half> @test_roundeven(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_roundeven_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_roundeven_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rni.f16.f16 %rs3, %rs2;
; CHECK-NEXT: cvt.rni.f16.f16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2070,7 +2163,8 @@ define <2 x half> @test_round(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<21>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_round_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_round_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: or.b32 %r4, %r3, 1056964608;
@@ -2121,17 +2215,20 @@ define <2 x half> @test_fmuladd(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0
; CHECK-NOF16-NEXT: .reg .b32 %r<13>;
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmuladd_param_0];
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmuladd_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
-; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_fmuladd_param_1];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs6;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs2;
+; CHECK-NOF16-NEXT: ld.param.b32 %r3, [test_fmuladd_param_2];
+; CHECK-NOF16-NEXT: ld.param.b32 %r2, [test_fmuladd_param_1];
+; CHECK-NOF16-NEXT: ld.param.b32 %r1, [test_fmuladd_param_0];
+; CHECK-NOF16-NEXT: mov.b32 {%rs1, %rs2}, %r3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: mov.b32 {%rs3, %rs4}, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
+; CHECK-NOF16-NEXT: mov.b32 {%rs5, %rs6}, %r1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs6;
; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs7, %r7;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs5;
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs3;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs5;
; CHECK-NOF16-NEXT: fma.rn.f32 %r11, %r10, %r9, %r8;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs8, %r11;
; CHECK-NOF16-NEXT: mov.b32 %r12, {%rs8, %rs7};
@@ -2148,7 +2245,8 @@ define <2 x half> @test_shufflevector(<2 x half> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_shufflevector_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
; CHECK-NEXT: ret;
%s = shufflevector <2 x half> %a, <2 x half> undef, <2 x i32> <i32 1, i32 0>
@@ -2158,12 +2256,13 @@ define <2 x half> @test_shufflevector(<2 x half> %a) #0 {
define <2 x half> @test_insertelement(<2 x half> %a, half %x) #0 {
; CHECK-LABEL: test_insertelement(
; CHECK: {
-; CHECK-NEXT: .reg .b16 %rs<4>;
+; CHECK-NEXT: .reg .b16 %rs<3>;
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
-; CHECK-NEXT: ld.param.v2.b16 {%rs2, %rs3}, [test_insertelement_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
+; CHECK-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
; CHECK-NEXT: ret;
%i = insertelement <2 x half> %a, half %x, i64 1
@@ -2177,7 +2276,8 @@ define <2 x half> @test_sitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_sitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rn.f16.s16 %rs3, %rs2;
; CHECK-NEXT: cvt.rn.f16.s16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -2193,7 +2293,8 @@ define <2 x half> @test_uitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {
; CHECK-NEXT: .reg .b32 %r<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_uitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: ld.param.b32 %r1, [test_uitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; CHECK-NEXT: cvt.rn.f16.u16 %rs3, %rs2;
; CHECK-NEXT: cvt.rn.f16.u16 %rs4, %rs1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
diff --git a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
index af3cb63082e78..da9e2d8cba139 100644
--- a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
@@ -28,29 +28,53 @@ define <2 x float> @test_ret_const() #0 {
}
define float @test_extract_0(<2 x float> %a) #0 {
-; CHECK-LABEL: test_extract_0(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_0_param_0];
-; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
-; CHECK-NEXT: ret;
+; CHECK-NOF32X2-LABEL: test_extract_0(
+; CHECK-NOF32X2: {
+; CHECK-NOF32X2-NEXT: .reg .b32 %r<2>;
+; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-NOF32X2-EMPTY:
+; CHECK-NOF32X2-NEXT: // %bb.0:
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_extract_0_param_0];
+; CHECK-NOF32X2-NEXT: { .reg .b32 tmp; mov.b64 {%r1, tmp}, %rd1; }
+; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NOF32X2-NEXT: ret;
+;
+; CHECK-F32X2-LABEL: test_extract_0(
+; CHECK-F32X2: {
+; CHECK-F32X2-NEXT: .reg .b32 %r<2>;
+; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-F32X2-EMPTY:
+; CHECK-F32X2-NEXT: // %bb.0:
+; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_extract_0_param_0];
+; CHECK-F32X2-NEXT: mov.b64 {%r1, _}, %rd1;
+; CHECK-F32X2-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-F32X2-NEXT: ret;
%e = extractelement <2 x float> %a, i32 0
ret float %e
}
define float @test_extract_1(<2 x float> %a) #0 {
-; CHECK-LABEL: test_extract_1(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<3>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_1_param_0];
-; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
-; CHECK-NEXT: ret;
+; CHECK-NOF32X2-LABEL: test_extract_1(
+; CHECK-NOF32X2: {
+; CHECK-NOF32X2-NEXT: .reg .b32 %r<2>;
+; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-NOF32X2-EMPTY:
+; CHECK-NOF32X2-NEXT: // %bb.0:
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_extract_1_param_0];
+; CHECK-NOF32X2-NEXT: { .reg .b32 tmp; mov.b64 {tmp, %r1}, %rd1; }
+; CHECK-NOF32X2-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-NOF32X2-NEXT: ret;
+;
+; CHECK-F32X2-LABEL: test_extract_1(
+; CHECK-F32X2: {
+; CHECK-F32X2-NEXT: .reg .b32 %r<2>;
+; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-F32X2-EMPTY:
+; CHECK-F32X2-NEXT: // %bb.0:
+; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_extract_1_param_0];
+; CHECK-F32X2-NEXT: mov.b64 {_, %r1}, %rd1;
+; CHECK-F32X2-NEXT: st.param.b32 [func_retval0], %r1;
+; CHECK-F32X2-NEXT: ret;
%e = extractelement <2 x float> %a, i32 1
ret float %e
}
@@ -70,10 +94,12 @@ define <2 x float> @test_fadd(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_param_1];
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fadd_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -98,7 +124,8 @@ define <2 x float> @test_fadd_imm_0(<2 x float> %a) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_0_param_0];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40000000;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -128,7 +155,8 @@ define <2 x float> @test_fadd_imm_1(<2 x float> %a) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_1_param_0];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40000000;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f3F800000;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -158,13 +186,17 @@ define <4 x float> @test_fadd_v4(<4 x float> %a, <4 x float> %b) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<5>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_v4_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_param_1];
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r9, %r4, %r8;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r10, %r3, %r7;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r11, %r2, %r6;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r12, %r1, %r5;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r10, %r9};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, %r1;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r7, %r8}, %rd3;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r9, %r10}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r11, %r10, %r8;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r12, %r9, %r7;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_v4(
@@ -189,12 +221,14 @@ define <4 x float> @test_fadd_imm_0_v4(<4 x float> %a) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_0_v4_param_0];
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, 0f40800000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, 0f40400000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r2, 0f40000000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40800000;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f40400000;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r6, 0f40000000;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r5, 0f3F800000;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r4, %r3};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_imm_0_v4(
@@ -225,12 +259,14 @@ define <4 x float> @test_fadd_imm_1_v4(<4 x float> %a) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_1_v4_param_0];
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r4, 0f40800000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r3, 0f40400000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r2, 0f40000000;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r1, 0f3F800000;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r3, %r2, 0f40800000;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r4, %r1, 0f40400000;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r6, 0f40000000;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r5, 0f3F800000;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r4, %r3};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_imm_1_v4(
@@ -261,10 +297,12 @@ define <2 x float> @test_fsub(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fsub_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_param_1];
-; CHECK-NOF32X2-NEXT: sub.rn.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: sub.rn.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fsub_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fsub_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: sub.rn.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: sub.rn.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -289,7 +327,8 @@ define <2 x float> @test_fneg(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fneg_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: neg.f32 %r3, %r2;
; CHECK-NEXT: neg.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -305,10 +344,12 @@ define <2 x float> @test_fmul(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmul_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_param_1];
-; CHECK-NOF32X2-NEXT: mul.rn.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: mul.rn.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fmul_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fmul_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: mul.rn.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: mul.rn.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -333,10 +374,12 @@ define <2 x float> @test_fdiv(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_param_1];
-; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
-; CHECK-NEXT: div.rn.f32 %r6, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fdiv_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fdiv_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: div.rn.f32 %r5, %r4, %r2;
+; CHECK-NEXT: div.rn.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NEXT: ret;
%r = fdiv <2 x float> %a, %b
@@ -351,20 +394,22 @@ define <2 x float> @test_frem(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_param_1];
-; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_frem_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_frem_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: div.rn.f32 %r5, %r4, %r2;
; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;
; CHECK-NEXT: neg.f32 %r7, %r6;
-; CHECK-NEXT: fma.rn.f32 %r8, %r7, %r4, %r2;
-; CHECK-NEXT: testp.infinite.f32 %p1, %r4;
-; CHECK-NEXT: selp.f32 %r9, %r2, %r8, %p1;
-; CHECK-NEXT: div.rn.f32 %r10, %r1, %r3;
+; CHECK-NEXT: fma.rn.f32 %r8, %r7, %r2, %r4;
+; CHECK-NEXT: testp.infinite.f32 %p1, %r2;
+; CHECK-NEXT: selp.f32 %r9, %r4, %r8, %p1;
+; CHECK-NEXT: div.rn.f32 %r10, %r3, %r1;
; CHECK-NEXT: cvt.rzi.f32.f32 %r11, %r10;
; CHECK-NEXT: neg.f32 %r12, %r11;
-; CHECK-NEXT: fma.rn.f32 %r13, %r12, %r3, %r1;
-; CHECK-NEXT: testp.infinite.f32 %p2, %r3;
-; CHECK-NEXT: selp.f32 %r14, %r1, %r13, %p2;
+; CHECK-NEXT: fma.rn.f32 %r13, %r12, %r1, %r3;
+; CHECK-NEXT: testp.infinite.f32 %p2, %r1;
+; CHECK-NEXT: selp.f32 %r14, %r3, %r13, %p2;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};
; CHECK-NEXT: ret;
%r = frem <2 x float> %a, %b
@@ -378,10 +423,12 @@ define <2 x float> @test_fadd_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_ftz_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_ftz_param_1];
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fadd_ftz_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -406,7 +453,8 @@ define <2 x float> @test_fadd_imm_0_ftz(<2 x float> %a) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_0_ftz_param_0];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_0_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40000000;
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f3F800000;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -436,7 +484,8 @@ define <2 x float> @test_fadd_imm_1_ftz(<2 x float> %a) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_imm_1_ftz_param_0];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fadd_imm_1_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40000000;
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f3F800000;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -466,13 +515,17 @@ define <4 x float> @test_fadd_v4_ftz(<4 x float> %a, <4 x float> %b) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<5>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_v4_ftz_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_ftz_param_1];
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r9, %r4, %r8;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r10, %r3, %r7;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r11, %r2, %r6;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r12, %r1, %r5;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r10, %r9};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_fadd_v4_ftz_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_v4_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, %r1;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r7, %r8}, %rd3;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r9, %r10}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r11, %r10, %r8;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r12, %r9, %r7;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r12, %r11, %r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_v4_ftz(
@@ -497,12 +550,14 @@ define <4 x float> @test_fadd_imm_0_v4_ftz(<4 x float> %a) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_0_v4_ftz_param_0];
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, 0f40800000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, 0f40400000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r2, 0f40000000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r1, 0f3F800000;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_0_v4_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40800000;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f40400000;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r6, 0f40000000;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r5, 0f3F800000;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r4, %r3};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_imm_0_v4_ftz(
@@ -533,12 +588,14 @@ define <4 x float> @test_fadd_imm_1_v4_ftz(<4 x float> %a) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [test_fadd_imm_1_v4_ftz_param_0];
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r4, 0f40800000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r3, 0f40400000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r2, 0f40000000;
-; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r1, 0f3F800000;
-; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r6, %r5};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_fadd_imm_1_v4_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r3, %r2, 0f40800000;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r4, %r1, 0f40400000;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r7, %r6, 0f40000000;
+; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r8, %r5, 0f3F800000;
+; CHECK-NOF32X2-NEXT: st.param.v4.b32 [func_retval0], {%r8, %r7, %r4, %r3};
; CHECK-NOF32X2-NEXT: ret;
;
; CHECK-F32X2-LABEL: test_fadd_imm_1_v4_ftz(
@@ -569,10 +626,12 @@ define <2 x float> @test_fsub_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fsub_ftz_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_ftz_param_1];
-; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fsub_ftz_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fsub_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -597,7 +656,8 @@ define <2 x float> @test_fneg_ftz(<2 x float> %a) #2 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_ftz_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fneg_ftz_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: neg.ftz.f32 %r3, %r2;
; CHECK-NEXT: neg.ftz.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -613,10 +673,12 @@ define <2 x float> @test_fmul_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<3>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmul_ftz_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_ftz_param_1];
-; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r5, %r2, %r4;
-; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r6, %r1, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fmul_ftz_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fmul_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r5, %r4, %r2;
+; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r6, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -641,11 +703,14 @@ define <2 x float> @test_fma_ftz(<2 x float> %a, <2 x float> %b, <2 x float> %c)
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<4>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fma_ftz_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_ftz_param_1];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_ftz_param_2];
-; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r7, %r2, %r4, %r6;
-; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r8, %r1, %r3, %r5;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd3, [test_fma_ftz_param_2];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fma_ftz_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fma_ftz_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd3;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r7, %r6, %r4, %r2;
+; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r8, %r5, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -671,10 +736,12 @@ define <2 x float> @test_fdiv_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_ftz_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_ftz_param_1];
-; CHECK-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;
-; CHECK-NEXT: div.rn.ftz.f32 %r6, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fdiv_ftz_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fdiv_ftz_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: div.rn.ftz.f32 %r5, %r4, %r2;
+; CHECK-NEXT: div.rn.ftz.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NEXT: ret;
%r = fdiv <2 x float> %a, %b
@@ -689,20 +756,22 @@ define <2 x float> @test_frem_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_ftz_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_ftz_param_1];
-; CHECK-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_frem_ftz_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_frem_ftz_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: div.rn.ftz.f32 %r5, %r4, %r2;
; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r6, %r5;
; CHECK-NEXT: neg.ftz.f32 %r7, %r6;
-; CHECK-NEXT: fma.rn.ftz.f32 %r8, %r7, %r4, %r2;
-; CHECK-NEXT: testp.infinite.f32 %p1, %r4;
-; CHECK-NEXT: selp.f32 %r9, %r2, %r8, %p1;
-; CHECK-NEXT: div.rn.ftz.f32 %r10, %r1, %r3;
+; CHECK-NEXT: fma.rn.ftz.f32 %r8, %r7, %r2, %r4;
+; CHECK-NEXT: testp.infinite.f32 %p1, %r2;
+; CHECK-NEXT: selp.f32 %r9, %r4, %r8, %p1;
+; CHECK-NEXT: div.rn.ftz.f32 %r10, %r3, %r1;
; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r11, %r10;
; CHECK-NEXT: neg.ftz.f32 %r12, %r11;
-; CHECK-NEXT: fma.rn.ftz.f32 %r13, %r12, %r3, %r1;
-; CHECK-NEXT: testp.infinite.f32 %p2, %r3;
-; CHECK-NEXT: selp.f32 %r14, %r1, %r13, %p2;
+; CHECK-NEXT: fma.rn.ftz.f32 %r13, %r12, %r1, %r3;
+; CHECK-NEXT: testp.infinite.f32 %p2, %r1;
+; CHECK-NEXT: selp.f32 %r14, %r3, %r13, %p2;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r14, %r9};
; CHECK-NEXT: ret;
%r = frem <2 x float> %a, %b
@@ -877,14 +946,18 @@ define <2 x float> @test_select_cc(<2 x float> %a, <2 x float> %b, <2 x float> %
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_param_2];
-; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_param_3];
-; CHECK-NEXT: setp.neu.f32 %p1, %r3, %r5;
-; CHECK-NEXT: setp.neu.f32 %p2, %r4, %r6;
-; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [test_select_cc_param_1];
-; CHECK-NEXT: selp.f32 %r9, %r2, %r8, %p2;
-; CHECK-NEXT: selp.f32 %r10, %r1, %r7, %p1;
+; CHECK-NEXT: ld.param.b64 %rd4, [test_select_cc_param_3];
+; CHECK-NEXT: ld.param.b64 %rd3, [test_select_cc_param_2];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_select_cc_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_select_cc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd3;
+; CHECK-NEXT: setp.neu.f32 %p1, %r3, %r1;
+; CHECK-NEXT: setp.neu.f32 %p2, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r5, %r6}, %rd2;
+; CHECK-NEXT: mov.b64 {%r7, %r8}, %rd1;
+; CHECK-NEXT: selp.f32 %r9, %r8, %r6, %p2;
+; CHECK-NEXT: selp.f32 %r10, %r7, %r5, %p1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};
; CHECK-NEXT: ret;
%cc = fcmp une <2 x float> %c, %d
@@ -902,10 +975,12 @@ define <2 x double> @test_select_cc_f64_f32(<2 x double> %a, <2 x double> %b, <2
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f64_f32_param_1];
; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_select_cc_f64_f32_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f64_f32_param_2];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f64_f32_param_3];
-; CHECK-NEXT: setp.neu.f32 %p1, %r1, %r3;
-; CHECK-NEXT: setp.neu.f32 %p2, %r2, %r4;
+; CHECK-NEXT: ld.param.b64 %rd6, [test_select_cc_f64_f32_param_3];
+; CHECK-NEXT: ld.param.b64 %rd5, [test_select_cc_f64_f32_param_2];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd6;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd5;
+; CHECK-NEXT: setp.neu.f32 %p1, %r3, %r1;
+; CHECK-NEXT: setp.neu.f32 %p2, %r4, %r2;
; CHECK-NEXT: selp.f64 %rd7, %rd2, %rd4, %p2;
; CHECK-NEXT: selp.f64 %rd8, %rd1, %rd3, %p1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd8, %rd7};
@@ -925,12 +1000,14 @@ define <2 x float> @test_select_cc_f32_f64(<2 x float> %a, <2 x float> %b, <2 x
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b64 {%rd5, %rd6}, [test_select_cc_f32_f64_param_3];
; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [test_select_cc_f32_f64_param_2];
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_f64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_select_cc_f32_f64_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_select_cc_f32_f64_param_0];
; CHECK-NEXT: setp.neu.f64 %p1, %rd3, %rd5;
; CHECK-NEXT: setp.neu.f64 %p2, %rd4, %rd6;
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f64_param_1];
-; CHECK-NEXT: selp.f32 %r5, %r2, %r4, %p2;
-; CHECK-NEXT: selp.f32 %r6, %r1, %r3, %p1;
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: selp.f32 %r5, %r4, %r2, %p2;
+; CHECK-NEXT: selp.f32 %r6, %r3, %r1, %p1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NEXT: ret;
%cc = fcmp une <2 x double> %c, %d
@@ -947,10 +1024,12 @@ define <2 x i1> @test_fcmp_une(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_une_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_une_param_1];
-; CHECK-NEXT: setp.neu.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.neu.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_une_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_une_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.neu.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.neu.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -969,10 +1048,12 @@ define <2 x i1> @test_fcmp_ueq(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ueq_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ueq_param_1];
-; CHECK-NEXT: setp.equ.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.equ.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ueq_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ueq_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.equ.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.equ.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -991,10 +1072,12 @@ define <2 x i1> @test_fcmp_ugt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ugt_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ugt_param_1];
-; CHECK-NEXT: setp.gtu.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.gtu.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ugt_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ugt_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.gtu.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.gtu.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1013,10 +1096,12 @@ define <2 x i1> @test_fcmp_uge(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uge_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uge_param_1];
-; CHECK-NEXT: setp.geu.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.geu.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_uge_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_uge_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.geu.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.geu.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1035,10 +1120,12 @@ define <2 x i1> @test_fcmp_ult(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ult_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ult_param_1];
-; CHECK-NEXT: setp.ltu.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.ltu.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ult_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ult_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.ltu.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.ltu.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1057,10 +1144,12 @@ define <2 x i1> @test_fcmp_ule(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ule_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ule_param_1];
-; CHECK-NEXT: setp.leu.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.leu.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ule_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ule_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.leu.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.leu.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1079,10 +1168,12 @@ define <2 x i1> @test_fcmp_uno(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uno_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uno_param_1];
-; CHECK-NEXT: setp.nan.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.nan.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_uno_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_uno_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.nan.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.nan.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1101,10 +1192,12 @@ define <2 x i1> @test_fcmp_one(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_one_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_one_param_1];
-; CHECK-NEXT: setp.ne.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.ne.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_one_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_one_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.ne.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.ne.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1123,10 +1216,12 @@ define <2 x i1> @test_fcmp_oeq(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oeq_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oeq_param_1];
-; CHECK-NEXT: setp.eq.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.eq.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_oeq_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_oeq_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.eq.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.eq.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1145,10 +1240,12 @@ define <2 x i1> @test_fcmp_ogt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ogt_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ogt_param_1];
-; CHECK-NEXT: setp.gt.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.gt.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ogt_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ogt_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.gt.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.gt.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1167,10 +1264,12 @@ define <2 x i1> @test_fcmp_oge(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oge_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oge_param_1];
-; CHECK-NEXT: setp.ge.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.ge.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_oge_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_oge_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.ge.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.ge.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1189,10 +1288,12 @@ define <2 x i1> @test_fcmp_olt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_olt_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_olt_param_1];
-; CHECK-NEXT: setp.lt.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.lt.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_olt_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_olt_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.lt.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.lt.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1211,10 +1312,12 @@ define <2 x i1> @test_fcmp_ole(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ole_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ole_param_1];
-; CHECK-NEXT: setp.le.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.le.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ole_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ole_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.le.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.le.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1233,10 +1336,12 @@ define <2 x i1> @test_fcmp_ord(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ord_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ord_param_1];
-; CHECK-NEXT: setp.num.f32 %p1, %r2, %r4;
-; CHECK-NEXT: setp.num.f32 %p2, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_fcmp_ord_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fcmp_ord_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: setp.num.f32 %p1, %r4, %r2;
+; CHECK-NEXT: setp.num.f32 %p2, %r3, %r1;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
; CHECK-NEXT: st.param.b8 [func_retval0], %rs1;
; CHECK-NEXT: selp.b16 %rs2, -1, 0, %p1;
@@ -1253,7 +1358,8 @@ define <2 x i32> @test_fptosi_i32(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i32_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fptosi_i32_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rzi.s32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rzi.s32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1269,7 +1375,8 @@ define <2 x i64> @test_fptosi_i64(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fptosi_i64_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rzi.s64.f32 %rd2, %r2;
; CHECK-NEXT: cvt.rzi.s64.f32 %rd3, %r1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};
@@ -1285,7 +1392,8 @@ define <2 x i32> @test_fptoui_2xi32(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rzi.u32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rzi.u32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1301,7 +1409,8 @@ define <2 x i64> @test_fptoui_2xi64(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rzi.u64.f32 %rd2, %r2;
; CHECK-NEXT: cvt.rzi.u64.f32 %rd3, %r1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};
@@ -1380,9 +1489,10 @@ define <2 x float> @test_uitofp_2xi32_fadd(<2 x i32> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_uitofp_2xi32_fadd_param_0];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_uitofp_2xi32_fadd_param_1];
; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r3, %r1;
; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r4, %r2;
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_uitofp_2xi32_fadd_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r6, %r4;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r5, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
@@ -1431,7 +1541,8 @@ define <2 x double> @test_fpext_2xdouble(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<4>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.f64.f32 %rd2, %r2;
; CHECK-NEXT: cvt.f64.f32 %rd3, %r1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};
@@ -1499,7 +1610,8 @@ define <2 x float> @test_sqrt(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sqrt_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sqrt_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: sqrt.rn.f32 %r3, %r2;
; CHECK-NEXT: sqrt.rn.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1522,7 +1634,8 @@ define <2 x float> @test_sin(<2 x float> %a) #0 #1 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sin_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_sin_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: sin.approx.f32 %r3, %r2;
; CHECK-NEXT: sin.approx.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1538,7 +1651,8 @@ define <2 x float> @test_cos(<2 x float> %a) #0 #1 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_cos_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_cos_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cos.approx.f32 %r3, %r2;
; CHECK-NEXT: cos.approx.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1597,11 +1711,14 @@ define <2 x float> @test_fma(<2 x float> %a, <2 x float> %b, <2 x float> %c) #0
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<4>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fma_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_param_1];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_param_2];
-; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r4, %r6;
-; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r3, %r5;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd3, [test_fma_param_2];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fma_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fma_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd3;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r5, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -1627,7 +1744,8 @@ define <2 x float> @test_fabs(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fabs_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_fabs_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: abs.f32 %r3, %r2;
; CHECK-NEXT: abs.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1643,10 +1761,12 @@ define <2 x float> @test_minnum(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_minnum_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_minnum_param_1];
-; CHECK-NEXT: min.f32 %r5, %r2, %r4;
-; CHECK-NEXT: min.f32 %r6, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_minnum_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_minnum_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: min.f32 %r5, %r4, %r2;
+; CHECK-NEXT: min.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NEXT: ret;
%r = call <2 x float> @llvm.minnum(<2 x float> %a, <2 x float> %b)
@@ -1660,10 +1780,12 @@ define <2 x float> @test_maxnum(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_maxnum_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_maxnum_param_1];
-; CHECK-NEXT: max.f32 %r5, %r2, %r4;
-; CHECK-NEXT: max.f32 %r6, %r1, %r3;
+; CHECK-NEXT: ld.param.b64 %rd2, [test_maxnum_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_maxnum_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd2;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd1;
+; CHECK-NEXT: max.f32 %r5, %r4, %r2;
+; CHECK-NEXT: max.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
; CHECK-NEXT: ret;
%r = call <2 x float> @llvm.maxnum(<2 x float> %a, <2 x float> %b)
@@ -1677,8 +1799,10 @@ define <2 x float> @test_copysign(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<3>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_param_1];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_copysign_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_copysign_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
; CHECK-NEXT: copysign.f32 %r5, %r4, %r2;
; CHECK-NEXT: copysign.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -1696,18 +1820,19 @@ define <2 x float> @test_copysign_f64(<2 x float> %a, <2 x double> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_copysign_f64_param_1];
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_f64_param_0];
-; CHECK-NEXT: abs.f32 %r3, %r2;
-; CHECK-NEXT: neg.f32 %r4, %r3;
+; CHECK-NEXT: ld.param.b64 %rd1, [test_copysign_f64_param_0];
; CHECK-NEXT: shr.u64 %rd4, %rd3, 63;
; CHECK-NEXT: and.b64 %rd5, %rd4, 1;
; CHECK-NEXT: setp.ne.b64 %p1, %rd5, 0;
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
+; CHECK-NEXT: abs.f32 %r3, %r2;
+; CHECK-NEXT: neg.f32 %r4, %r3;
; CHECK-NEXT: selp.f32 %r5, %r4, %r3, %p1;
-; CHECK-NEXT: abs.f32 %r6, %r1;
-; CHECK-NEXT: neg.f32 %r7, %r6;
; CHECK-NEXT: shr.u64 %rd6, %rd2, 63;
; CHECK-NEXT: and.b64 %rd7, %rd6, 1;
; CHECK-NEXT: setp.ne.b64 %p2, %rd7, 0;
+; CHECK-NEXT: abs.f32 %r6, %r1;
+; CHECK-NEXT: neg.f32 %r7, %r6;
; CHECK-NEXT: selp.f32 %r8, %r7, %r6, %p2;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r5};
; CHECK-NEXT: ret;
@@ -1723,8 +1848,10 @@ define <2 x double> @test_copysign_extended(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_extended_param_0];
-; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_extended_param_1];
+; CHECK-NEXT: ld.param.b64 %rd2, [test_copysign_extended_param_1];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_copysign_extended_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
; CHECK-NEXT: copysign.f32 %r5, %r3, %r1;
; CHECK-NEXT: copysign.f32 %r6, %r4, %r2;
; CHECK-NEXT: cvt.f64.f32 %rd3, %r6;
@@ -1743,7 +1870,8 @@ define <2 x float> @test_floor(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_floor_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_floor_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rmi.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rmi.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1759,7 +1887,8 @@ define <2 x float> @test_ceil(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_ceil_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_ceil_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rpi.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rpi.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1775,7 +1904,8 @@ define <2 x float> @test_trunc(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_trunc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rzi.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rzi.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1791,7 +1921,8 @@ define <2 x float> @test_rint(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_rint_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_rint_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rni.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rni.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1807,7 +1938,8 @@ define <2 x float> @test_nearbyint(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_nearbyint_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_nearbyint_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rni.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rni.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1823,7 +1955,8 @@ define <2 x float> @test_roundeven(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_roundeven_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_roundeven_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: cvt.rni.f32.f32 %r3, %r2;
; CHECK-NEXT: cvt.rni.f32.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1841,7 +1974,8 @@ define <2 x float> @test_round(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_round_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_round_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: or.b32 %r4, %r3, 1056964608;
; CHECK-NEXT: add.rn.f32 %r5, %r2, %r4;
@@ -1875,11 +2009,14 @@ define <2 x float> @test_fmuladd(<2 x float> %a, <2 x float> %b, <2 x float> %c)
; CHECK-NOF32X2-NEXT: .reg .b64 %rd<4>;
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmuladd_param_0];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmuladd_param_1];
-; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fmuladd_param_2];
-; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r4, %r6;
-; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r3, %r5;
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd3, [test_fmuladd_param_2];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd2, [test_fmuladd_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_fmuladd_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 {%r1, %r2}, %rd3;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NOF32X2-NEXT: mov.b64 {%r5, %r6}, %rd1;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r6, %r4, %r2;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r5, %r3, %r1;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -1905,7 +2042,8 @@ define <2 x float> @test_shufflevector(<2 x float> %a) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_shufflevector_param_0];
+; CHECK-NEXT: ld.param.b64 %rd1, [test_shufflevector_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
; CHECK-NEXT: ret;
%s = shufflevector <2 x float> %a, <2 x float> poison, <2 x i32> <i32 1, i32 0>
@@ -1913,16 +2051,29 @@ define <2 x float> @test_shufflevector(<2 x float> %a) #0 {
}
define <2 x float> @test_insertelement(<2 x float> %a, float %x) #0 {
-; CHECK-LABEL: test_insertelement(
-; CHECK: {
-; CHECK-NEXT: .reg .b32 %r<4>;
-; CHECK-NEXT: .reg .b64 %rd<2>;
-; CHECK-EMPTY:
-; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];
-; CHECK-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_insertelement_param_0];
-; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
-; CHECK-NEXT: ret;
+; CHECK-NOF32X2-LABEL: test_insertelement(
+; CHECK-NOF32X2: {
+; CHECK-NOF32X2-NEXT: .reg .b32 %r<3>;
+; CHECK-NOF32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-NOF32X2-EMPTY:
+; CHECK-NOF32X2-NEXT: // %bb.0:
+; CHECK-NOF32X2-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];
+; CHECK-NOF32X2-NEXT: ld.param.b64 %rd1, [test_insertelement_param_0];
+; CHECK-NOF32X2-NEXT: { .reg .b32 tmp; mov.b64 {%r2, tmp}, %rd1; }
+; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
+; CHECK-NOF32X2-NEXT: ret;
+;
+; CHECK-F32X2-LABEL: test_insertelement(
+; CHECK-F32X2: {
+; CHECK-F32X2-NEXT: .reg .b32 %r<3>;
+; CHECK-F32X2-NEXT: .reg .b64 %rd<2>;
+; CHECK-F32X2-EMPTY:
+; CHECK-F32X2-NEXT: // %bb.0:
+; CHECK-F32X2-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];
+; CHECK-F32X2-NEXT: ld.param.b64 %rd1, [test_insertelement_param_0];
+; CHECK-F32X2-NEXT: mov.b64 {%r2, _}, %rd1;
+; CHECK-F32X2-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
+; CHECK-F32X2-NEXT: ret;
%i = insertelement <2 x float> %a, float %x, i64 1
ret <2 x float> %i
}
diff --git a/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
index 1a61498b10142..2b7a06c33d948 100644
--- a/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
@@ -32,31 +32,57 @@ define <2 x i16> @test_ret_const() #0 {
}
define i16 @test_extract_0(<2 x i16> %a) #0 {
-; COMMON-LABEL: test_extract_0(
-; COMMON: {
-; COMMON-NEXT: .reg .b16 %rs<3>;
-; COMMON-NEXT: .reg .b32 %r<3>;
-; COMMON-EMPTY:
-; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_0_param_0];
-; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
-; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
-; COMMON-NEXT: ret;
+; I16x2-LABEL: test_extract_0(
+; I16x2: {
+; I16x2-NEXT: .reg .b16 %rs<2>;
+; I16x2-NEXT: .reg .b32 %r<3>;
+; I16x2-EMPTY:
+; I16x2-NEXT: // %bb.0:
+; I16x2-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];
+; I16x2-NEXT: mov.b32 {%rs1, _}, %r1;
+; I16x2-NEXT: cvt.u32.u16 %r2, %rs1;
+; I16x2-NEXT: st.param.b32 [func_retval0], %r2;
+; I16x2-NEXT: ret;
+;
+; NO-I16x2-LABEL: test_extract_0(
+; NO-I16x2: {
+; NO-I16x2-NEXT: .reg .b16 %rs<2>;
+; NO-I16x2-NEXT: .reg .b32 %r<3>;
+; NO-I16x2-EMPTY:
+; NO-I16x2-NEXT: // %bb.0:
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_extract_0_param_0];
+; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {%rs1, tmp}, %r1; }
+; NO-I16x2-NEXT: cvt.u32.u16 %r2, %rs1;
+; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;
+; NO-I16x2-NEXT: ret;
%e = extractelement <2 x i16> %a, i32 0
ret i16 %e
}
define i16 @test_extract_1(<2 x i16> %a) #0 {
-; COMMON-LABEL: test_extract_1(
-; COMMON: {
-; COMMON-NEXT: .reg .b16 %rs<3>;
-; COMMON-NEXT: .reg .b32 %r<3>;
-; COMMON-EMPTY:
-; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_1_param_0];
-; COMMON-NEXT: cvt.u32.u16 %r2, %rs2;
-; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
-; COMMON-NEXT: ret;
+; I16x2-LABEL: test_extract_1(
+; I16x2: {
+; I16x2-NEXT: .reg .b16 %rs<2>;
+; I16x2-NEXT: .reg .b32 %r<3>;
+; I16x2-EMPTY:
+; I16x2-NEXT: // %bb.0:
+; I16x2-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];
+; I16x2-NEXT: mov.b32 {_, %rs1}, %r1;
+; I16x2-NEXT: cvt.u32.u16 %r2, %rs1;
+; I16x2-NEXT: st.param.b32 [func_retval0], %r2;
+; I16x2-NEXT: ret;
+;
+; NO-I16x2-LABEL: test_extract_1(
+; NO-I16x2: {
+; NO-I16x2-NEXT: .reg .b16 %rs<2>;
+; NO-I16x2-NEXT: .reg .b32 %r<3>;
+; NO-I16x2-EMPTY:
+; NO-I16x2-NEXT: // %bb.0:
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_extract_1_param_0];
+; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {tmp, %rs1}, %r1; }
+; NO-I16x2-NEXT: cvt.u32.u16 %r2, %rs1;
+; NO-I16x2-NEXT: st.param.b32 [func_retval0], %r2;
+; NO-I16x2-NEXT: ret;
%e = extractelement <2 x i16> %a, i32 1
ret i16 %e
}
@@ -71,8 +97,9 @@ define i16 @test_extract_i(<2 x i16> %a, i64 %idx) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_i_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_extract_i_param_0];
; COMMON-NEXT: setp.eq.b64 %p1, %rd1, 0;
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
; COMMON-NEXT: cvt.u32.u16 %r2, %rs3;
; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
@@ -99,10 +126,12 @@ define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<3>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_add_param_0];
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_add_param_1];
-; NO-I16x2-NEXT: add.s16 %rs5, %rs2, %rs4;
-; NO-I16x2-NEXT: add.s16 %rs6, %rs1, %rs3;
+; NO-I16x2-NEXT: ld.param.b32 %r2, [test_add_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; NO-I16x2-NEXT: add.s16 %rs5, %rs4, %rs2;
+; NO-I16x2-NEXT: add.s16 %rs6, %rs3, %rs1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; NO-I16x2-NEXT: ret;
%r = add <2 x i16> %a, %b
@@ -128,7 +157,8 @@ define <2 x i16> @test_add_imm_0(<2 x i16> %a) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<2>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_add_imm_0_param_0];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_0_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;
; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -155,7 +185,8 @@ define <2 x i16> @test_add_imm_1(<2 x i16> %a) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<2>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_add_imm_1_param_0];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_add_imm_1_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; NO-I16x2-NEXT: add.s16 %rs3, %rs2, 2;
; NO-I16x2-NEXT: add.s16 %rs4, %rs1, 1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -171,10 +202,12 @@ define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) #0 {
; COMMON-NEXT: .reg .b32 %r<3>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sub_param_0];
-; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_sub_param_1];
-; COMMON-NEXT: sub.s16 %rs5, %rs2, %rs4;
-; COMMON-NEXT: sub.s16 %rs6, %rs1, %rs3;
+; COMMON-NEXT: ld.param.b32 %r2, [test_sub_param_1];
+; COMMON-NEXT: ld.param.b32 %r1, [test_sub_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; COMMON-NEXT: sub.s16 %rs5, %rs4, %rs2;
+; COMMON-NEXT: sub.s16 %rs6, %rs3, %rs1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; COMMON-NEXT: ret;
%r = sub <2 x i16> %a, %b
@@ -199,10 +232,12 @@ define <2 x i16> @test_smax(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<3>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_smax_param_0];
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_smax_param_1];
-; NO-I16x2-NEXT: max.s16 %rs5, %rs2, %rs4;
-; NO-I16x2-NEXT: max.s16 %rs6, %rs1, %rs3;
+; NO-I16x2-NEXT: ld.param.b32 %r2, [test_smax_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_smax_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; NO-I16x2-NEXT: max.s16 %rs5, %rs4, %rs2;
+; NO-I16x2-NEXT: max.s16 %rs6, %rs3, %rs1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; NO-I16x2-NEXT: ret;
%cmp = icmp sgt <2 x i16> %a, %b
@@ -228,10 +263,12 @@ define <2 x i16> @test_umax(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<3>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_umax_param_0];
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_umax_param_1];
-; NO-I16x2-NEXT: max.u16 %rs5, %rs2, %rs4;
-; NO-I16x2-NEXT: max.u16 %rs6, %rs1, %rs3;
+; NO-I16x2-NEXT: ld.param.b32 %r2, [test_umax_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_umax_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; NO-I16x2-NEXT: max.u16 %rs5, %rs4, %rs2;
+; NO-I16x2-NEXT: max.u16 %rs6, %rs3, %rs1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; NO-I16x2-NEXT: ret;
%cmp = icmp ugt <2 x i16> %a, %b
@@ -257,10 +294,12 @@ define <2 x i16> @test_smin(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<3>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_smin_param_0];
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_smin_param_1];
-; NO-I16x2-NEXT: min.s16 %rs5, %rs2, %rs4;
-; NO-I16x2-NEXT: min.s16 %rs6, %rs1, %rs3;
+; NO-I16x2-NEXT: ld.param.b32 %r2, [test_smin_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_smin_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; NO-I16x2-NEXT: min.s16 %rs5, %rs4, %rs2;
+; NO-I16x2-NEXT: min.s16 %rs6, %rs3, %rs1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; NO-I16x2-NEXT: ret;
%cmp = icmp sle <2 x i16> %a, %b
@@ -286,10 +325,12 @@ define <2 x i16> @test_umin(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-NEXT: .reg .b32 %r<3>;
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_umin_param_0];
-; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_umin_param_1];
-; NO-I16x2-NEXT: min.u16 %rs5, %rs2, %rs4;
-; NO-I16x2-NEXT: min.u16 %rs6, %rs1, %rs3;
+; NO-I16x2-NEXT: ld.param.b32 %r2, [test_umin_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_umin_param_0];
+; NO-I16x2-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; NO-I16x2-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; NO-I16x2-NEXT: min.u16 %rs5, %rs4, %rs2;
+; NO-I16x2-NEXT: min.u16 %rs6, %rs3, %rs1;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; NO-I16x2-NEXT: ret;
%cmp = icmp ule <2 x i16> %a, %b
@@ -304,10 +345,12 @@ define <2 x i16> @test_mul(<2 x i16> %a, <2 x i16> %b) #0 {
; COMMON-NEXT: .reg .b32 %r<3>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_mul_param_0];
-; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_mul_param_1];
-; COMMON-NEXT: mul.lo.s16 %rs5, %rs2, %rs4;
-; COMMON-NEXT: mul.lo.s16 %rs6, %rs1, %rs3;
+; COMMON-NEXT: ld.param.b32 %r2, [test_mul_param_1];
+; COMMON-NEXT: ld.param.b32 %r1, [test_mul_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; COMMON-NEXT: mul.lo.s16 %rs5, %rs4, %rs2;
+; COMMON-NEXT: mul.lo.s16 %rs6, %rs3, %rs1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; COMMON-NEXT: ret;
%r = mul <2 x i16> %a, %b
@@ -686,14 +729,18 @@ define <2 x i16> @test_select_cc(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x
; COMMON-NEXT: .reg .b32 %r<5>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];
-; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_2];
-; COMMON-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_select_cc_param_3];
-; COMMON-NEXT: setp.ne.b16 %p1, %rs3, %rs5;
-; COMMON-NEXT: setp.ne.b16 %p2, %rs4, %rs6;
-; COMMON-NEXT: ld.param.v2.b16 {%rs7, %rs8}, [test_select_cc_param_1];
-; COMMON-NEXT: selp.b16 %rs9, %rs2, %rs8, %p2;
-; COMMON-NEXT: selp.b16 %rs10, %rs1, %rs7, %p1;
+; COMMON-NEXT: ld.param.b32 %r4, [test_select_cc_param_3];
+; COMMON-NEXT: ld.param.b32 %r3, [test_select_cc_param_2];
+; COMMON-NEXT: ld.param.b32 %r2, [test_select_cc_param_1];
+; COMMON-NEXT: ld.param.b32 %r1, [test_select_cc_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r4;
+; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r3;
+; COMMON-NEXT: setp.ne.b16 %p1, %rs3, %rs1;
+; COMMON-NEXT: setp.ne.b16 %p2, %rs4, %rs2;
+; COMMON-NEXT: mov.b32 {%rs5, %rs6}, %r2;
+; COMMON-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; COMMON-NEXT: selp.b16 %rs9, %rs8, %rs6, %p2;
+; COMMON-NEXT: selp.b16 %rs10, %rs7, %rs5, %p1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
; COMMON-NEXT: ret;
%cc = icmp ne <2 x i16> %c, %d
@@ -711,10 +758,12 @@ define <2 x i32> @test_select_cc_i32_i16(<2 x i32> %a, <2 x i32> %b,
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_i32_i16_param_1];
; COMMON-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_i32_i16_param_0];
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_i32_i16_param_2];
-; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_i32_i16_param_3];
-; COMMON-NEXT: setp.ne.b16 %p1, %rs1, %rs3;
-; COMMON-NEXT: setp.ne.b16 %p2, %rs2, %rs4;
+; COMMON-NEXT: ld.param.b32 %r6, [test_select_cc_i32_i16_param_3];
+; COMMON-NEXT: ld.param.b32 %r5, [test_select_cc_i32_i16_param_2];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r6;
+; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r5;
+; COMMON-NEXT: setp.ne.b16 %p1, %rs3, %rs1;
+; COMMON-NEXT: setp.ne.b16 %p2, %rs4, %rs2;
; COMMON-NEXT: selp.b32 %r7, %r2, %r4, %p2;
; COMMON-NEXT: selp.b32 %r8, %r1, %r3, %p1;
; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
@@ -735,12 +784,14 @@ define <2 x i16> @test_select_cc_i16_i32(<2 x i16> %a, <2 x i16> %b,
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_i16_i32_param_3];
; COMMON-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_i16_i32_param_2];
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_i16_i32_param_0];
+; COMMON-NEXT: ld.param.b32 %r2, [test_select_cc_i16_i32_param_1];
+; COMMON-NEXT: ld.param.b32 %r1, [test_select_cc_i16_i32_param_0];
; COMMON-NEXT: setp.ne.b32 %p1, %r3, %r5;
; COMMON-NEXT: setp.ne.b32 %p2, %r4, %r6;
-; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_i16_i32_param_1];
-; COMMON-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;
-; COMMON-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; COMMON-NEXT: mov.b32 {%rs3, %rs4}, %r1;
+; COMMON-NEXT: selp.b16 %rs5, %rs4, %rs2, %p2;
+; COMMON-NEXT: selp.b16 %rs6, %rs3, %rs1, %p1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
; COMMON-NEXT: ret;
<2 x i32> %c, <2 x i32> %d) #0 {
@@ -851,7 +902,8 @@ define <2 x i32> @test_zext_2xi32(<2 x i16> %a) #0 {
; COMMON-NEXT: .reg .b32 %r<4>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_zext_2xi32_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_zext_2xi32_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: cvt.u32.u16 %r2, %rs2;
; COMMON-NEXT: cvt.u32.u16 %r3, %rs1;
; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -868,7 +920,8 @@ define <2 x i64> @test_zext_2xi64(<2 x i16> %a) #0 {
; COMMON-NEXT: .reg .b64 %rd<3>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_zext_2xi64_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_zext_2xi64_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: cvt.u64.u16 %rd1, %rs2;
; COMMON-NEXT: cvt.u64.u16 %rd2, %rs1;
; COMMON-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -926,7 +979,8 @@ define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 {
; COMMON-NEXT: .reg .b32 %r<2>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_shufflevector_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_shufflevector_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
; COMMON-NEXT: ret;
%s = shufflevector <2 x i16> %a, <2 x i16> undef, <2 x i32> <i32 1, i32 0>
@@ -934,16 +988,29 @@ define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 {
}
define <2 x i16> @test_insertelement(<2 x i16> %a, i16 %x) #0 {
-; COMMON-LABEL: test_insertelement(
-; COMMON: {
-; COMMON-NEXT: .reg .b16 %rs<4>;
-; COMMON-NEXT: .reg .b32 %r<2>;
-; COMMON-EMPTY:
-; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
-; COMMON-NEXT: ld.param.v2.b16 {%rs2, %rs3}, [test_insertelement_param_0];
-; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
-; COMMON-NEXT: ret;
+; I16x2-LABEL: test_insertelement(
+; I16x2: {
+; I16x2-NEXT: .reg .b16 %rs<3>;
+; I16x2-NEXT: .reg .b32 %r<2>;
+; I16x2-EMPTY:
+; I16x2-NEXT: // %bb.0:
+; I16x2-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
+; I16x2-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
+; I16x2-NEXT: mov.b32 {%rs2, _}, %r1;
+; I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; I16x2-NEXT: ret;
+;
+; NO-I16x2-LABEL: test_insertelement(
+; NO-I16x2: {
+; NO-I16x2-NEXT: .reg .b16 %rs<3>;
+; NO-I16x2-NEXT: .reg .b32 %r<2>;
+; NO-I16x2-EMPTY:
+; NO-I16x2-NEXT: // %bb.0:
+; NO-I16x2-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
+; NO-I16x2-NEXT: ld.param.b32 %r1, [test_insertelement_param_0];
+; NO-I16x2-NEXT: { .reg .b16 tmp; mov.b32 {%rs2, tmp}, %r1; }
+; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
+; NO-I16x2-NEXT: ret;
%i = insertelement <2 x i16> %a, i16 %x, i64 1
ret <2 x i16> %i
}
@@ -955,7 +1022,8 @@ define <2 x i16> @test_fptosi_2xhalf_to_2xi16(<2 x half> %a) #0 {
; COMMON-NEXT: .reg .b32 %r<2>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_fptosi_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;
; COMMON-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
@@ -971,7 +1039,8 @@ define <2 x i16> @test_fptoui_2xhalf_to_2xi16(<2 x half> %a) #0 {
; COMMON-NEXT: .reg .b32 %r<2>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
-; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: ld.param.b32 %r1, [test_fptoui_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: mov.b32 {%rs1, %rs2}, %r1;
; COMMON-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;
; COMMON-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs4, %rs3};
diff --git a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
index aba20e6b0f27f..9891e33151f8a 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -1935,16 +1935,18 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 {
; O0-NEXT: .reg .b32 %r<12>;
; O0-EMPTY:
; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptosi_4xhalf_to_4xi8_param_0];
-; O0-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;
-; O0-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;
-; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
-; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3;
-; O0-NEXT: cvt.u32.u16 %r4, %rs8;
-; O0-NEXT: cvt.u32.u16 %r5, %rs7;
+; O0-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_4xhalf_to_4xi8_param_0];
+; O0-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; O0-NEXT: cvt.rzi.s16.f16 %rs3, %rs2;
+; O0-NEXT: cvt.rzi.s16.f16 %rs4, %rs1;
+; O0-NEXT: mov.b32 %r3, {%rs4, %rs3};
+; O0-NEXT: mov.b32 {%rs5, %rs6}, %r3;
+; O0-NEXT: cvt.u32.u16 %r4, %rs6;
+; O0-NEXT: cvt.u32.u16 %r5, %rs5;
; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
-; O0-NEXT: cvt.rzi.s16.f16 %rs9, %rs2;
-; O0-NEXT: cvt.rzi.s16.f16 %rs10, %rs1;
+; O0-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; O0-NEXT: cvt.rzi.s16.f16 %rs9, %rs8;
+; O0-NEXT: cvt.rzi.s16.f16 %rs10, %rs7;
; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};
; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;
; O0-NEXT: cvt.u32.u16 %r8, %rs12;
@@ -1989,16 +1991,18 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 {
; O0-NEXT: .reg .b32 %r<12>;
; O0-EMPTY:
; O0-NEXT: // %bb.0:
-; O0-NEXT: ld.param.v4.b16 {%rs1, %rs2, %rs3, %rs4}, [test_fptoui_4xhalf_to_4xi8_param_0];
-; O0-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;
-; O0-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;
-; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
-; O0-NEXT: mov.b32 {%rs7, %rs8}, %r3;
-; O0-NEXT: cvt.u32.u16 %r4, %rs8;
-; O0-NEXT: cvt.u32.u16 %r5, %rs7;
+; O0-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_4xhalf_to_4xi8_param_0];
+; O0-NEXT: mov.b32 {%rs1, %rs2}, %r2;
+; O0-NEXT: cvt.rzi.u16.f16 %rs3, %rs2;
+; O0-NEXT: cvt.rzi.u16.f16 %rs4, %rs1;
+; O0-NEXT: mov.b32 %r3, {%rs4, %rs3};
+; O0-NEXT: mov.b32 {%rs5, %rs6}, %r3;
+; O0-NEXT: cvt.u32.u16 %r4, %rs6;
+; O0-NEXT: cvt.u32.u16 %r5, %rs5;
; O0-NEXT: prmt.b32 %r6, %r5, %r4, 0x3340U;
-; O0-NEXT: cvt.rzi.u16.f16 %rs9, %rs2;
-; O0-NEXT: cvt.rzi.u16.f16 %rs10, %rs1;
+; O0-NEXT: mov.b32 {%rs7, %rs8}, %r1;
+; O0-NEXT: cvt.rzi.u16.f16 %rs9, %rs8;
+; O0-NEXT: cvt.rzi.u16.f16 %rs10, %rs7;
; O0-NEXT: mov.b32 %r7, {%rs10, %rs9};
; O0-NEXT: mov.b32 {%rs11, %rs12}, %r7;
; O0-NEXT: cvt.u32.u16 %r8, %rs12;
diff --git a/llvm/test/CodeGen/NVPTX/pr126337.ll b/llvm/test/CodeGen/NVPTX/pr126337.ll
new file mode 100644
index 0000000000000..32e411584b0e5
--- /dev/null
+++ b/llvm/test/CodeGen/NVPTX/pr126337.ll
@@ -0,0 +1,41 @@
+; NOTE: Assertions have been autogenerated by utils/update_llc_test_checks.py UTC_ARGS: --version 5
+; RUN: llc < %s -mtriple=nvptx64 -mcpu=sm_70 | FileCheck %s
+; RUN: %if ptxas %{ llc < %s -mtriple=nvptx64 -mcpu=sm_70 | %ptxas-verify %}
+
+; This IR should compile without triggering assertions in LICM
+; when the CopyToReg from %0 in the first BB gets eliminated
+; but we still use its result in the second BB.
+; Technically the problem happens in MIR, but there are multiple
+; passes involved, so testing with the IR reproducer is more convenient.
+; https://github.com/llvm/llvm-project/pull/126337#issuecomment-3081431594
+
+target datalayout = "e-p6:32:32-i64:64-i128:128-v16:16-v32:32-n16:32:64"
+target triple = "nvptx64-nvidia-cuda"
+
+define ptx_kernel void @Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel(<2 x float> %0) {
+; CHECK-LABEL: Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel(
+; CHECK: {
+; CHECK-NEXT: .reg .pred %p<2>;
+; CHECK-NEXT: .reg .b16 %rs<2>;
+; CHECK-NEXT: .reg .b32 %r<2>;
+; CHECK-NEXT: .reg .b64 %rd<3>;
+; CHECK-EMPTY:
+; CHECK-NEXT: // %bb.0: // %.preheader15
+; CHECK-NEXT: ld.param.b64 %rd1, [Equal_GPU_DT_COMPLEX64_DT_BOOL_kernel_param_0];
+; CHECK-NEXT: { .reg .b32 tmp; mov.b64 {%r1, tmp}, %rd1; }
+; CHECK-NEXT: setp.eq.f32 %p1, %r1, 0f00000000;
+; CHECK-NEXT: selp.b16 %rs1, 1, 0, %p1;
+; CHECK-NEXT: $L__BB0_1: // =>This Inner Loop Header: Depth=1
+; CHECK-NEXT: mov.b64 %rd2, 0;
+; CHECK-NEXT: st.b8 [%rd2], %rs1;
+; CHECK-NEXT: bra.uni $L__BB0_1;
+.preheader15:
+ br label %1
+
+1: ; preds = %1, %.preheader15
+ %2 = fcmp oeq <2 x float> %0, zeroinitializer
+ %3 = extractelement <2 x i1> %2, i64 0
+ store i1 %3, ptr null, align 4
+ br label %1
+}
+
diff --git a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
index 87f965c84b6b6..92cb51b17f0c8 100644
--- a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
@@ -117,16 +117,20 @@ define float @reduce_fadd_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_param_0];
-; CHECK-NEXT: add.rn.f32 %r9, %r5, 0f00000000;
-; CHECK-NEXT: add.rn.f32 %r10, %r9, %r6;
-; CHECK-NEXT: add.rn.f32 %r11, %r10, %r7;
-; CHECK-NEXT: add.rn.f32 %r12, %r11, %r8;
-; CHECK-NEXT: add.rn.f32 %r13, %r12, %r1;
-; CHECK-NEXT: add.rn.f32 %r14, %r13, %r2;
-; CHECK-NEXT: add.rn.f32 %r15, %r14, %r3;
-; CHECK-NEXT: add.rn.f32 %r16, %r15, %r4;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fadd_float_param_0+16];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd3;
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fadd_float_param_0];
+; CHECK-NEXT: mov.b64 {%r5, %r6}, %rd2;
+; CHECK-NEXT: mov.b64 {%r7, %r8}, %rd1;
+; CHECK-NEXT: add.rn.f32 %r9, %r7, 0f00000000;
+; CHECK-NEXT: add.rn.f32 %r10, %r9, %r8;
+; CHECK-NEXT: add.rn.f32 %r11, %r10, %r5;
+; CHECK-NEXT: add.rn.f32 %r12, %r11, %r6;
+; CHECK-NEXT: add.rn.f32 %r13, %r12, %r3;
+; CHECK-NEXT: add.rn.f32 %r14, %r13, %r4;
+; CHECK-NEXT: add.rn.f32 %r15, %r14, %r1;
+; CHECK-NEXT: add.rn.f32 %r16, %r15, %r2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r16;
; CHECK-NEXT: ret;
%res = call float @llvm.vector.reduce.fadd(float 0.0, <8 x float> %in)
@@ -140,14 +144,18 @@ define float @reduce_fadd_float_reassoc(<8 x float> %in) {
; CHECK-SM80-NEXT: .reg .b64 %rd<5>;
; CHECK-SM80-EMPTY:
; CHECK-SM80-NEXT: // %bb.0:
-; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_reassoc_param_0+16];
-; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_reassoc_param_0];
-; CHECK-SM80-NEXT: add.rn.f32 %r9, %r7, %r3;
-; CHECK-SM80-NEXT: add.rn.f32 %r10, %r5, %r1;
-; CHECK-SM80-NEXT: add.rn.f32 %r11, %r8, %r4;
-; CHECK-SM80-NEXT: add.rn.f32 %r12, %r6, %r2;
+; CHECK-SM80-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fadd_float_reassoc_param_0+16];
+; CHECK-SM80-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fadd_float_reassoc_param_0];
+; CHECK-SM80-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-SM80-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-SM80-NEXT: add.rn.f32 %r5, %r3, %r1;
+; CHECK-SM80-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-SM80-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-SM80-NEXT: add.rn.f32 %r10, %r8, %r6;
+; CHECK-SM80-NEXT: add.rn.f32 %r11, %r4, %r2;
+; CHECK-SM80-NEXT: add.rn.f32 %r12, %r9, %r7;
; CHECK-SM80-NEXT: add.rn.f32 %r13, %r12, %r11;
-; CHECK-SM80-NEXT: add.rn.f32 %r14, %r10, %r9;
+; CHECK-SM80-NEXT: add.rn.f32 %r14, %r10, %r5;
; CHECK-SM80-NEXT: add.rn.f32 %r15, %r14, %r13;
; CHECK-SM80-NEXT: add.rn.f32 %r16, %r15, 0f00000000;
; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r16;
@@ -321,15 +329,19 @@ define float @reduce_fmul_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_param_0];
-; CHECK-NEXT: mul.rn.f32 %r9, %r5, %r6;
-; CHECK-NEXT: mul.rn.f32 %r10, %r9, %r7;
-; CHECK-NEXT: mul.rn.f32 %r11, %r10, %r8;
-; CHECK-NEXT: mul.rn.f32 %r12, %r11, %r1;
-; CHECK-NEXT: mul.rn.f32 %r13, %r12, %r2;
-; CHECK-NEXT: mul.rn.f32 %r14, %r13, %r3;
-; CHECK-NEXT: mul.rn.f32 %r15, %r14, %r4;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmul_float_param_0+16];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd3;
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmul_float_param_0];
+; CHECK-NEXT: mov.b64 {%r5, %r6}, %rd2;
+; CHECK-NEXT: mov.b64 {%r7, %r8}, %rd1;
+; CHECK-NEXT: mul.rn.f32 %r9, %r7, %r8;
+; CHECK-NEXT: mul.rn.f32 %r10, %r9, %r5;
+; CHECK-NEXT: mul.rn.f32 %r11, %r10, %r6;
+; CHECK-NEXT: mul.rn.f32 %r12, %r11, %r3;
+; CHECK-NEXT: mul.rn.f32 %r13, %r12, %r4;
+; CHECK-NEXT: mul.rn.f32 %r14, %r13, %r1;
+; CHECK-NEXT: mul.rn.f32 %r15, %r14, %r2;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
; CHECK-NEXT: ret;
%res = call float @llvm.vector.reduce.fmul(float 1.0, <8 x float> %in)
@@ -343,14 +355,18 @@ define float @reduce_fmul_float_reassoc(<8 x float> %in) {
; CHECK-SM80-NEXT: .reg .b64 %rd<5>;
; CHECK-SM80-EMPTY:
; CHECK-SM80-NEXT: // %bb.0:
-; CHECK-SM80-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_reassoc_param_0+16];
-; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_reassoc_param_0];
-; CHECK-SM80-NEXT: mul.rn.f32 %r9, %r7, %r3;
-; CHECK-SM80-NEXT: mul.rn.f32 %r10, %r5, %r1;
-; CHECK-SM80-NEXT: mul.rn.f32 %r11, %r8, %r4;
-; CHECK-SM80-NEXT: mul.rn.f32 %r12, %r6, %r2;
+; CHECK-SM80-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmul_float_reassoc_param_0+16];
+; CHECK-SM80-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmul_float_reassoc_param_0];
+; CHECK-SM80-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-SM80-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-SM80-NEXT: mul.rn.f32 %r5, %r3, %r1;
+; CHECK-SM80-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-SM80-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-SM80-NEXT: mul.rn.f32 %r10, %r8, %r6;
+; CHECK-SM80-NEXT: mul.rn.f32 %r11, %r4, %r2;
+; CHECK-SM80-NEXT: mul.rn.f32 %r12, %r9, %r7;
; CHECK-SM80-NEXT: mul.rn.f32 %r13, %r12, %r11;
-; CHECK-SM80-NEXT: mul.rn.f32 %r14, %r10, %r9;
+; CHECK-SM80-NEXT: mul.rn.f32 %r14, %r10, %r5;
; CHECK-SM80-NEXT: mul.rn.f32 %r15, %r14, %r13;
; CHECK-SM80-NEXT: st.param.b32 [func_retval0], %r15;
; CHECK-SM80-NEXT: ret;
@@ -494,13 +510,17 @@ define float @reduce_fmax_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_param_0];
-; CHECK-NEXT: max.f32 %r9, %r8, %r4;
-; CHECK-NEXT: max.f32 %r10, %r6, %r2;
-; CHECK-NEXT: max.f32 %r11, %r10, %r9;
-; CHECK-NEXT: max.f32 %r12, %r7, %r3;
-; CHECK-NEXT: max.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmax_float_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmax_float_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: max.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: max.f32 %r10, %r9, %r7;
+; CHECK-NEXT: max.f32 %r11, %r10, %r5;
+; CHECK-NEXT: max.f32 %r12, %r3, %r1;
+; CHECK-NEXT: max.f32 %r13, %r8, %r6;
; CHECK-NEXT: max.f32 %r14, %r13, %r12;
; CHECK-NEXT: max.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -517,13 +537,17 @@ define float @reduce_fmax_float_reassoc(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_reassoc_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_reassoc_param_0];
-; CHECK-NEXT: max.f32 %r9, %r8, %r4;
-; CHECK-NEXT: max.f32 %r10, %r6, %r2;
-; CHECK-NEXT: max.f32 %r11, %r10, %r9;
-; CHECK-NEXT: max.f32 %r12, %r7, %r3;
-; CHECK-NEXT: max.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmax_float_reassoc_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmax_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: max.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: max.f32 %r10, %r9, %r7;
+; CHECK-NEXT: max.f32 %r11, %r10, %r5;
+; CHECK-NEXT: max.f32 %r12, %r3, %r1;
+; CHECK-NEXT: max.f32 %r13, %r8, %r6;
; CHECK-NEXT: max.f32 %r14, %r13, %r12;
; CHECK-NEXT: max.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -628,13 +652,17 @@ define float @reduce_fmin_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_param_0];
-; CHECK-NEXT: min.f32 %r9, %r8, %r4;
-; CHECK-NEXT: min.f32 %r10, %r6, %r2;
-; CHECK-NEXT: min.f32 %r11, %r10, %r9;
-; CHECK-NEXT: min.f32 %r12, %r7, %r3;
-; CHECK-NEXT: min.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmin_float_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmin_float_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: min.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: min.f32 %r10, %r9, %r7;
+; CHECK-NEXT: min.f32 %r11, %r10, %r5;
+; CHECK-NEXT: min.f32 %r12, %r3, %r1;
+; CHECK-NEXT: min.f32 %r13, %r8, %r6;
; CHECK-NEXT: min.f32 %r14, %r13, %r12;
; CHECK-NEXT: min.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -651,13 +679,17 @@ define float @reduce_fmin_float_reassoc(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_reassoc_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_reassoc_param_0];
-; CHECK-NEXT: min.f32 %r9, %r8, %r4;
-; CHECK-NEXT: min.f32 %r10, %r6, %r2;
-; CHECK-NEXT: min.f32 %r11, %r10, %r9;
-; CHECK-NEXT: min.f32 %r12, %r7, %r3;
-; CHECK-NEXT: min.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmin_float_reassoc_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmin_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: min.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: min.f32 %r10, %r9, %r7;
+; CHECK-NEXT: min.f32 %r11, %r10, %r5;
+; CHECK-NEXT: min.f32 %r12, %r3, %r1;
+; CHECK-NEXT: min.f32 %r13, %r8, %r6;
; CHECK-NEXT: min.f32 %r14, %r13, %r12;
; CHECK-NEXT: min.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -762,13 +794,17 @@ define float @reduce_fmaximum_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_param_0];
-; CHECK-NEXT: max.NaN.f32 %r9, %r8, %r4;
-; CHECK-NEXT: max.NaN.f32 %r10, %r6, %r2;
-; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r9;
-; CHECK-NEXT: max.NaN.f32 %r12, %r7, %r3;
-; CHECK-NEXT: max.NaN.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmaximum_float_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmaximum_float_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: max.NaN.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: max.NaN.f32 %r10, %r9, %r7;
+; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r5;
+; CHECK-NEXT: max.NaN.f32 %r12, %r3, %r1;
+; CHECK-NEXT: max.NaN.f32 %r13, %r8, %r6;
; CHECK-NEXT: max.NaN.f32 %r14, %r13, %r12;
; CHECK-NEXT: max.NaN.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -785,13 +821,17 @@ define float @reduce_fmaximum_float_reassoc(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_reassoc_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_reassoc_param_0];
-; CHECK-NEXT: max.NaN.f32 %r9, %r8, %r4;
-; CHECK-NEXT: max.NaN.f32 %r10, %r6, %r2;
-; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r9;
-; CHECK-NEXT: max.NaN.f32 %r12, %r7, %r3;
-; CHECK-NEXT: max.NaN.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fmaximum_float_reassoc_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fmaximum_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: max.NaN.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: max.NaN.f32 %r10, %r9, %r7;
+; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r5;
+; CHECK-NEXT: max.NaN.f32 %r12, %r3, %r1;
+; CHECK-NEXT: max.NaN.f32 %r13, %r8, %r6;
; CHECK-NEXT: max.NaN.f32 %r14, %r13, %r12;
; CHECK-NEXT: max.NaN.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -896,13 +936,17 @@ define float @reduce_fminimum_float(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_param_0];
-; CHECK-NEXT: min.NaN.f32 %r9, %r8, %r4;
-; CHECK-NEXT: min.NaN.f32 %r10, %r6, %r2;
-; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r9;
-; CHECK-NEXT: min.NaN.f32 %r12, %r7, %r3;
-; CHECK-NEXT: min.NaN.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fminimum_float_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fminimum_float_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: min.NaN.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: min.NaN.f32 %r10, %r9, %r7;
+; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r5;
+; CHECK-NEXT: min.NaN.f32 %r12, %r3, %r1;
+; CHECK-NEXT: min.NaN.f32 %r13, %r8, %r6;
; CHECK-NEXT: min.NaN.f32 %r14, %r13, %r12;
; CHECK-NEXT: min.NaN.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
@@ -919,13 +963,17 @@ define float @reduce_fminimum_float_reassoc(<8 x float> %in) {
; CHECK-NEXT: .reg .b64 %rd<5>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
-; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_reassoc_param_0+16];
-; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_reassoc_param_0];
-; CHECK-NEXT: min.NaN.f32 %r9, %r8, %r4;
-; CHECK-NEXT: min.NaN.f32 %r10, %r6, %r2;
-; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r9;
-; CHECK-NEXT: min.NaN.f32 %r12, %r7, %r3;
-; CHECK-NEXT: min.NaN.f32 %r13, %r5, %r1;
+; CHECK-NEXT: ld.param.v2.b64 {%rd3, %rd4}, [reduce_fminimum_float_reassoc_param_0+16];
+; CHECK-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [reduce_fminimum_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 {%r1, %r2}, %rd4;
+; CHECK-NEXT: mov.b64 {%r3, %r4}, %rd2;
+; CHECK-NEXT: min.NaN.f32 %r5, %r4, %r2;
+; CHECK-NEXT: mov.b64 {%r6, %r7}, %rd3;
+; CHECK-NEXT: mov.b64 {%r8, %r9}, %rd1;
+; CHECK-NEXT: min.NaN.f32 %r10, %r9, %r7;
+; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r5;
+; CHECK-NEXT: min.NaN.f32 %r12, %r3, %r1;
+; CHECK-NEXT: min.NaN.f32 %r13, %r8, %r6;
; CHECK-NEXT: min.NaN.f32 %r14, %r13, %r12;
; CHECK-NEXT: min.NaN.f32 %r15, %r14, %r11;
; CHECK-NEXT: st.param.b32 [func_retval0], %r15;
More information about the llvm-commits
mailing list