[llvm] [NVPTX] don't erase CopyToRegs when folding movs into loads (PR #149393)
Princeton Ferro via llvm-commits
llvm-commits at lists.llvm.org
Fri Jul 18 00:06:02 PDT 2025
https://github.com/Prince781 updated https://github.com/llvm/llvm-project/pull/149393
>From 20cc9c3d3cc5ea99c9b7d1130a534416fb22998d Mon Sep 17 00:00:00 2001
From: Princeton Ferro <pferro at nvidia.com>
Date: Wed, 16 Jul 2025 18:10:56 -0700
Subject: [PATCH] [NVPTX] don't erase CopyToRegs when folding movs into loads
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
---
llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp | 13 +-
llvm/test/CodeGen/NVPTX/f16x2-instructions.ll | 177 ++++++++++++++----
llvm/test/CodeGen/NVPTX/f32x2-instructions.ll | 175 ++++++++++++++---
llvm/test/CodeGen/NVPTX/i16x2-instructions.ll | 57 ++++--
llvm/test/CodeGen/NVPTX/i8x4-instructions.ll | 4 +
.../CodeGen/NVPTX/reduction-intrinsics.ll | 48 +++++
llvm/test/CodeGen/pr126337.ll | 41 ++++
7 files changed, 427 insertions(+), 88 deletions(-)
create mode 100644 llvm/test/CodeGen/pr126337.ll
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 7aa06f9079b09..bb01357d36a5a 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5060,11 +5060,10 @@ 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());
+ // Peek through CopyToReg nodes. TODO: SelectionDAG needs to be improved
+ // to eliminate these nodes when they're unused in -O0. See PR126337.
+ if (U.getUser()->getOpcode() == ISD::CopyToReg)
return true;
- }
// Otherwise, this use prevents us from splitting a value.
return false;
@@ -5132,10 +5131,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 +6539,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..66166756aecf9 100644
--- a/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f16x2-instructions.ll
@@ -50,6 +50,7 @@ define half @test_extract_0(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_0_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: st.param.b16 [func_retval0], %rs1;
; CHECK-NEXT: ret;
%e = extractelement <2 x half> %a, i32 0
@@ -64,6 +65,7 @@ define half @test_extract_1(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_1_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: st.param.b16 [func_retval0], %rs2;
; CHECK-NEXT: ret;
%e = extractelement <2 x half> %a, i32 1
@@ -79,8 +81,9 @@ define half @test_extract_i(<2 x half> %a, i64 %idx) #0 {
; CHECK-NEXT: .reg .b64 %rd<2>;
; 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: mov.b32 %r1, {%rs1, %rs2};
+; CHECK-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
; CHECK-NEXT: setp.eq.b64 %p1, %rd1, 0;
; CHECK-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
; CHECK-NEXT: st.param.b16 [func_retval0], %rs3;
@@ -108,7 +111,9 @@ define <2 x half> @test_fadd(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fadd_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fadd_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: add.rn.f32 %r5, %r4, %r3;
@@ -144,6 +149,7 @@ define <2 x half> @test_fadd_imm_0(<2 x half> %a) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; 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;
@@ -176,6 +182,7 @@ define <2 x half> @test_fadd_imm_1(<2 x half> %a) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; 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;
@@ -208,7 +215,9 @@ define <2 x half> @test_fsub(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fsub_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fsub_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: sub.rn.f32 %r5, %r4, %r3;
@@ -243,6 +252,7 @@ define <2 x half> @test_fneg(<2 x half> %a) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fneg_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: mov.b32 %r3, 0f00000000;
; CHECK-NOF16-NEXT: sub.rn.f32 %r4, %r3, %r2;
@@ -276,7 +286,9 @@ define <2 x half> @test_fmul(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmul_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmul_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: mul.rn.f32 %r5, %r4, %r3;
@@ -300,7 +312,9 @@ define <2 x half> @test_fdiv(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fdiv_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fdiv_param_1];
+; CHECK-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;
@@ -332,7 +346,9 @@ define <2 x half> @test_frem(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_frem_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_frem_param_1];
+; CHECK-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NEXT: div.rn.f32 %r5, %r4, %r3;
@@ -533,11 +549,13 @@ define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, <
; CHECK-F16-NEXT: .reg .b32 %r<5>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_param_0];
+; CHECK-F16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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: setp.neu.f16x2 %p1|%p2, %r3, %r4;
; CHECK-F16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_1];
+; CHECK-F16-NEXT: mov.b32 %r2, {%rs3, %rs4};
+; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r3, %r4;
; CHECK-F16-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;
; CHECK-F16-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;
; CHECK-F16-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -551,15 +569,19 @@ define <2 x half> @test_select_cc(<2 x half> %a, <2 x half> %b, <2 x half> %c, <
; 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: mov.b32 %r1, {%rs1, %rs2};
; 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: mov.b32 %r4, {%rs3, %rs4};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_select_cc_param_2];
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs5, %rs6};
+; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs7, %rs8}, [test_select_cc_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs7, %rs8};
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs3;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs5;
; 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: 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: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
@@ -577,11 +599,13 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b,
; CHECK-F16-NEXT: .reg .b64 %rd<3>;
; CHECK-F16-EMPTY:
; CHECK-F16-NEXT: // %bb.0:
+; CHECK-F16-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f32_f16_param_0];
+; CHECK-F16-NEXT: mov.b64 %rd1, {%r3, %r4};
; 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: 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: mov.b64 %rd2, {%r5, %r6};
+; CHECK-F16-NEXT: setp.neu.f16x2 %p1|%p2, %r1, %r2;
; CHECK-F16-NEXT: selp.f32 %r7, %r4, %r6, %p2;
; CHECK-F16-NEXT: selp.f32 %r8, %r3, %r5, %p1;
; CHECK-F16-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
@@ -596,17 +620,21 @@ define <2 x float> @test_select_cc_f32_f16(<2 x float> %a, <2 x float> %b,
; 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: mov.b64 %rd1, {%r3, %r4};
; 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: mov.b32 %r2, {%rs1, %rs2};
; 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: mov.b32 %r1, {%rs3, %rs4};
+; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f32_f16_param_1];
+; CHECK-NOF16-NEXT: mov.b64 %rd2, {%r5, %r6};
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
+; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r8, %r7;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r9, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r10, %rs4;
+; CHECK-NOF16-NEXT: setp.neu.f32 %p2, %r10, %r9;
+; CHECK-NOF16-NEXT: selp.f32 %r11, %r4, %r6, %p2;
+; CHECK-NOF16-NEXT: selp.f32 %r12, %r3, %r5, %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 {
@@ -625,11 +653,15 @@ define <2 x half> @test_select_cc_f16_f32(<2 x half> %a, <2 x half> %b,
; 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: mov.b32 %r1, {%rs1, %rs2};
+; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f16_f32_param_3];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_f16_f32_param_2];
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_f16_f32_param_1];
+; CHECK-NEXT: mov.b32 %r2, {%rs3, %rs4};
+; CHECK-NEXT: setp.neu.f32 %p1, %r5, %r3;
+; CHECK-NEXT: setp.neu.f32 %p2, %r6, %r4;
; CHECK-NEXT: selp.b16 %rs5, %rs2, %rs4, %p2;
; CHECK-NEXT: selp.b16 %rs6, %rs1, %rs3, %p1;
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -665,7 +697,9 @@ define <2 x i1> @test_fcmp_une(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_une_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.neu.f32 %p1, %r4, %r3;
@@ -706,7 +740,9 @@ define <2 x i1> @test_fcmp_ueq(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ueq_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.equ.f32 %p1, %r4, %r3;
@@ -747,7 +783,9 @@ define <2 x i1> @test_fcmp_ugt(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ugt_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.gtu.f32 %p1, %r4, %r3;
@@ -788,7 +826,9 @@ define <2 x i1> @test_fcmp_uge(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_uge_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.geu.f32 %p1, %r4, %r3;
@@ -829,7 +869,9 @@ define <2 x i1> @test_fcmp_ult(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ult_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.ltu.f32 %p1, %r4, %r3;
@@ -870,7 +912,9 @@ define <2 x i1> @test_fcmp_ule(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ule_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.leu.f32 %p1, %r4, %r3;
@@ -912,7 +956,9 @@ define <2 x i1> @test_fcmp_uno(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_uno_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.nan.f32 %p1, %r4, %r3;
@@ -953,7 +999,9 @@ define <2 x i1> @test_fcmp_one(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_one_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.ne.f32 %p1, %r4, %r3;
@@ -994,7 +1042,9 @@ define <2 x i1> @test_fcmp_oeq(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_oeq_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.eq.f32 %p1, %r4, %r3;
@@ -1035,7 +1085,9 @@ define <2 x i1> @test_fcmp_ogt(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ogt_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.gt.f32 %p1, %r4, %r3;
@@ -1076,7 +1128,9 @@ define <2 x i1> @test_fcmp_oge(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_oge_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.ge.f32 %p1, %r4, %r3;
@@ -1117,7 +1171,9 @@ define <2 x i1> @test_fcmp_olt(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_olt_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.lt.f32 %p1, %r4, %r3;
@@ -1158,7 +1214,9 @@ define <2 x i1> @test_fcmp_ole(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ole_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.le.f32 %p1, %r4, %r3;
@@ -1199,7 +1257,9 @@ define <2 x i1> @test_fcmp_ord(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fcmp_ord_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NOF16-NEXT: setp.num.f32 %p1, %r4, %r3;
@@ -1223,6 +1283,7 @@ define <2 x i32> @test_fptosi_i32(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_i32_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -1240,6 +1301,7 @@ define <2 x i64> @test_fptosi_i64(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_i64_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -1256,6 +1318,7 @@ define <2 x i32> @test_fptoui_2xi32(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -1273,6 +1336,7 @@ define <2 x i64> @test_fptoui_2xi64(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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,15 +1433,16 @@ 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: 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: ld.param.v2.b16 {%rs1, %rs2}, [test_uitofp_2xi32_fadd_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs1, %rs2};
+; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs3, %r1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.u32 %rs4, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r4, %r5;
; 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: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r7, %r8;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;
; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};
@@ -1411,15 +1476,16 @@ 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: 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: ld.param.v2.b16 {%rs1, %rs2}, [test_sitofp_2xi32_fadd_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs1, %rs2};
+; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs3, %r1;
+; CHECK-NOF16-NEXT: cvt.rn.f16.s32 %rs4, %r2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs2;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs4;
; CHECK-NOF16-NEXT: add.rn.f32 %r6, %r4, %r5;
; 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: cvt.f32.f16 %r7, %rs1;
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r8, %rs3;
; CHECK-NOF16-NEXT: add.rn.f32 %r9, %r7, %r8;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs6, %r9;
; CHECK-NOF16-NEXT: mov.b32 %r10, {%rs6, %rs5};
@@ -1469,6 +1535,7 @@ define <2 x float> @test_fpext_2xfloat(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fpext_2xfloat_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: cvt.f32.f16 %r3, %rs1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -1486,6 +1553,7 @@ define <2 x double> @test_fpext_2xdouble(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f64.f16 %rd1, %rs2;
; CHECK-NEXT: cvt.f64.f16 %rd2, %rs1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -1579,6 +1647,7 @@ define <2 x half> @test_sqrt(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sqrt_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: sqrt.rn.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1607,6 +1676,7 @@ define <2 x half> @test_sin(<2 x half> %a) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sin_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: sin.approx.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1628,6 +1698,7 @@ define <2 x half> @test_cos(<2 x half> %a) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_cos_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: cos.approx.f32 %r3, %r2;
; CHECK-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1704,9 +1775,12 @@ define <2 x half> @test_fma(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fma_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fma_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs3, %rs4};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_fma_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs5, %rs6};
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs6;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs2;
; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;
@@ -1741,6 +1815,7 @@ define <2 x half> @test_fabs(<2 x half> %a) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fabs_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NOF16-NEXT: abs.f32 %r3, %r2;
; CHECK-NOF16-NEXT: cvt.rn.f16.f32 %rs3, %r3;
@@ -1762,7 +1837,9 @@ define <2 x half> @test_minnum(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_minnum_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_minnum_param_1];
+; CHECK-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NEXT: min.f32 %r5, %r4, %r3;
@@ -1786,7 +1863,9 @@ define <2 x half> @test_maxnum(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_maxnum_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_maxnum_param_1];
+; CHECK-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NEXT: cvt.f32.f16 %r3, %rs4;
; CHECK-NEXT: cvt.f32.f16 %r4, %rs2;
; CHECK-NEXT: max.f32 %r5, %r4, %r3;
@@ -1823,7 +1902,9 @@ define <2 x half> @test_copysign(<2 x half> %a, <2 x half> %b) #0 {
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_copysign_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_copysign_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: and.b16 %rs5, %rs4, -32768;
; CHECK-NOF16-NEXT: and.b16 %rs6, %rs2, 32767;
; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs5;
@@ -1846,6 +1927,7 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 {
; CHECK-F16-NEXT: // %bb.0:
; 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 %rd1, {%r2, %r3};
; 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};
@@ -1863,7 +1945,9 @@ define <2 x half> @test_copysign_f32(<2 x half> %a, <2 x float> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b32 {%r2, %r3}, [test_copysign_f32_param_1];
+; CHECK-NOF16-NEXT: mov.b64 %rd1, {%r2, %r3};
; 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; }
@@ -1905,8 +1989,9 @@ define <2 x half> @test_copysign_f64(<2 x half> %a, <2 x double> %b) #0 {
; CHECK-NOF16-NEXT: .reg .b64 %rd<7>;
; 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: mov.b32 %r1, {%rs1, %rs2};
+; CHECK-NOF16-NEXT: ld.param.v2.b64 {%rd1, %rd2}, [test_copysign_f64_param_1];
; CHECK-NOF16-NEXT: and.b16 %rs3, %rs2, 32767;
; CHECK-NOF16-NEXT: and.b64 %rd3, %rd2, -9223372036854775808;
; CHECK-NOF16-NEXT: shr.u64 %rd4, %rd3, 48;
@@ -1949,7 +2034,9 @@ define <2 x float> @test_copysign_extended(<2 x half> %a, <2 x half> %b) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_copysign_extended_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs3, %rs4};
; CHECK-NOF16-NEXT: and.b16 %rs5, %rs3, -32768;
; CHECK-NOF16-NEXT: and.b16 %rs6, %rs1, 32767;
; CHECK-NOF16-NEXT: or.b16 %rs7, %rs6, %rs5;
@@ -1973,6 +2060,7 @@ define <2 x half> @test_floor(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_floor_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -1989,6 +2077,7 @@ define <2 x half> @test_ceil(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_ceil_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2005,6 +2094,7 @@ define <2 x half> @test_trunc(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_trunc_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2021,6 +2111,7 @@ define <2 x half> @test_rint(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_rint_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2037,6 +2128,7 @@ define <2 x half> @test_nearbyint(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_nearbyint_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2053,6 +2145,7 @@ define <2 x half> @test_roundeven(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_roundeven_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2071,6 +2164,7 @@ define <2 x half> @test_round(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_round_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NEXT: cvt.f32.f16 %r2, %rs2;
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: or.b32 %r4, %r3, 1056964608;
@@ -2122,9 +2216,12 @@ define <2 x half> @test_fmuladd(<2 x half> %a, <2 x half> %b, <2 x half> %c) #0
; CHECK-NOF16-EMPTY:
; CHECK-NOF16-NEXT: // %bb.0:
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fmuladd_param_0];
+; CHECK-NOF16-NEXT: mov.b32 %r1, {%rs1, %rs2};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_fmuladd_param_2];
-; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
+; CHECK-NOF16-NEXT: mov.b32 %r3, {%rs3, %rs4};
; CHECK-NOF16-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_fmuladd_param_1];
+; CHECK-NOF16-NEXT: mov.b32 %r2, {%rs5, %rs6};
+; CHECK-NOF16-NEXT: cvt.f32.f16 %r4, %rs4;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r5, %rs6;
; CHECK-NOF16-NEXT: cvt.f32.f16 %r6, %rs2;
; CHECK-NOF16-NEXT: fma.rn.f32 %r7, %r6, %r5, %r4;
@@ -2149,6 +2246,7 @@ define <2 x half> @test_shufflevector(<2 x half> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_shufflevector_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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>
@@ -2162,8 +2260,9 @@ define <2 x half> @test_insertelement(<2 x half> %a, half %x) #0 {
; 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: mov.b32 %r1, {%rs2, %rs3};
+; CHECK-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
; CHECK-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
; CHECK-NEXT: ret;
%i = insertelement <2 x half> %a, half %x, i64 1
@@ -2178,6 +2277,7 @@ define <2 x half> @test_sitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -2194,6 +2294,7 @@ define <2 x half> @test_uitofp_2xi16_to_2xhalf(<2 x i16> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_uitofp_2xi16_to_2xhalf_param_0];
+; CHECK-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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..ce238fbef0685 100644
--- a/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/f32x2-instructions.ll
@@ -35,6 +35,7 @@ define float @test_extract_0(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_0_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: st.param.b32 [func_retval0], %r1;
; CHECK-NEXT: ret;
%e = extractelement <2 x float> %a, i32 0
@@ -49,6 +50,7 @@ define float @test_extract_1(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_extract_1_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: st.param.b32 [func_retval0], %r2;
; CHECK-NEXT: ret;
%e = extractelement <2 x float> %a, i32 1
@@ -71,7 +73,9 @@ define <2 x float> @test_fadd(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fadd_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: add.rn.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: add.rn.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -99,6 +103,7 @@ define <2 x float> @test_fadd_imm_0(<2 x float> %a) #0 {
; 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: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -129,6 +134,7 @@ define <2 x float> @test_fadd_imm_1(<2 x float> %a) #0 {
; 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: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -159,7 +165,11 @@ define <4 x float> @test_fadd_v4(<4 x float> %a, <4 x float> %b) #0 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd4, {%r7, %r8};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd3, {%r5, %r6};
; 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;
@@ -190,6 +200,8 @@ define <4 x float> @test_fadd_imm_0_v4(<4 x float> %a) #0 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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;
@@ -226,6 +238,8 @@ define <4 x float> @test_fadd_imm_1_v4(<4 x float> %a) #0 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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;
@@ -262,7 +276,9 @@ define <2 x float> @test_fsub(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fsub_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: sub.rn.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: sub.rn.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -290,6 +306,7 @@ define <2 x float> @test_fneg(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: neg.f32 %r3, %r2;
; CHECK-NEXT: neg.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -306,7 +323,9 @@ define <2 x float> @test_fmul(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-NOF32X2-EMPTY:
; CHECK-NOF32X2-NEXT: // %bb.0:
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fmul_param_0];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: mul.rn.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: mul.rn.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -334,7 +353,9 @@ define <2 x float> @test_fdiv(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
; CHECK-NEXT: div.rn.f32 %r6, %r1, %r3;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -352,7 +373,9 @@ define <2 x float> @test_frem(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: div.rn.f32 %r5, %r2, %r4;
; CHECK-NEXT: cvt.rzi.f32.f32 %r6, %r5;
; CHECK-NEXT: neg.f32 %r7, %r6;
@@ -379,7 +402,9 @@ define <2 x float> @test_fadd_ftz(<2 x float> %a, <2 x float> %b) #2 {
; 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: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fadd_ftz_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: add.rn.ftz.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -407,6 +432,7 @@ define <2 x float> @test_fadd_imm_0_ftz(<2 x float> %a) #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: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -437,6 +463,7 @@ define <2 x float> @test_fadd_imm_1_ftz(<2 x float> %a) #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: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -467,7 +494,11 @@ define <4 x float> @test_fadd_v4_ftz(<4 x float> %a, <4 x float> %b) #2 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [test_fadd_v4_ftz_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd4, {%r7, %r8};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd3, {%r5, %r6};
; 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;
@@ -498,6 +529,8 @@ define <4 x float> @test_fadd_imm_0_v4_ftz(<4 x float> %a) #2 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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;
@@ -534,6 +567,8 @@ define <4 x float> @test_fadd_imm_1_v4_ftz(<4 x float> %a) #2 {
; 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: mov.b64 %rd2, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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;
@@ -570,7 +605,9 @@ define <2 x float> @test_fsub_ftz(<2 x float> %a, <2 x float> %b) #2 {
; 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: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fsub_ftz_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: sub.rn.ftz.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -598,6 +635,7 @@ define <2 x float> @test_fneg_ftz(<2 x float> %a) #2 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fneg_ftz_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: neg.ftz.f32 %r3, %r2;
; CHECK-NEXT: neg.ftz.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -614,7 +652,9 @@ define <2 x float> @test_fmul_ftz(<2 x float> %a, <2 x float> %b) #2 {
; 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: mov.b64 %rd1, {%r1, %r2};
; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmul_ftz_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r5, %r2, %r4;
; CHECK-NOF32X2-NEXT: mul.rn.ftz.f32 %r6, %r1, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -642,10 +682,13 @@ define <2 x float> @test_fma_ftz(<2 x float> %a, <2 x float> %b, <2 x float> %c)
; 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: mov.b64 %rd1, {%r1, %r2};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_ftz_param_2];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd3, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_ftz_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r5, %r6};
+; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r7, %r2, %r6, %r4;
+; CHECK-NOF32X2-NEXT: fma.rn.ftz.f32 %r8, %r1, %r5, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -672,7 +715,9 @@ define <2 x float> @test_fdiv_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fdiv_ftz_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fdiv_ftz_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;
; CHECK-NEXT: div.rn.ftz.f32 %r6, %r1, %r3;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -690,7 +735,9 @@ define <2 x float> @test_frem_ftz(<2 x float> %a, <2 x float> %b) #2 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_frem_ftz_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_frem_ftz_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: div.rn.ftz.f32 %r5, %r2, %r4;
; CHECK-NEXT: cvt.rzi.ftz.f32.f32 %r6, %r5;
; CHECK-NEXT: neg.ftz.f32 %r7, %r6;
@@ -878,11 +925,15 @@ define <2 x float> @test_select_cc(<2 x float> %a, <2 x float> %b, <2 x float> %
; 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: mov.b64 %rd1, {%r1, %r2};
+; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_param_3];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_select_cc_param_2];
+; CHECK-NEXT: mov.b64 %rd3, {%r5, %r6};
; CHECK-NEXT: ld.param.v2.b32 {%r7, %r8}, [test_select_cc_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: setp.neu.f32 %p1, %r5, %r3;
+; CHECK-NEXT: setp.neu.f32 %p2, %r6, %r4;
; CHECK-NEXT: selp.f32 %r9, %r2, %r8, %p2;
; CHECK-NEXT: selp.f32 %r10, %r1, %r7, %p1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r10, %r9};
@@ -902,10 +953,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.v2.b32 {%r1, %r2}, [test_select_cc_f64_f32_param_3];
+; CHECK-NEXT: mov.b64 %rd6, {%r1, %r2};
+; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_select_cc_f64_f32_param_2];
+; CHECK-NEXT: mov.b64 %rd5, {%r3, %r4};
+; 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};
@@ -923,12 +976,14 @@ define <2 x float> @test_select_cc_f32_f64(<2 x float> %a, <2 x float> %b, <2 x
; CHECK-NEXT: .reg .b64 %rd<7>;
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
+; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_select_cc_f32_f64_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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.v2.b32 {%r3, %r4}, [test_select_cc_f32_f64_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; 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: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -948,7 +1003,9 @@ define <2 x i1> @test_fcmp_une(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_une_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_une_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.neu.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.neu.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -970,7 +1027,9 @@ define <2 x i1> @test_fcmp_ueq(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ueq_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ueq_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.equ.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.equ.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -992,7 +1051,9 @@ define <2 x i1> @test_fcmp_ugt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ugt_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ugt_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.gtu.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.gtu.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1014,7 +1075,9 @@ define <2 x i1> @test_fcmp_uge(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uge_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uge_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.geu.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.geu.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1036,7 +1099,9 @@ define <2 x i1> @test_fcmp_ult(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ult_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ult_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.ltu.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.ltu.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1058,7 +1123,9 @@ define <2 x i1> @test_fcmp_ule(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ule_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ule_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.leu.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.leu.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1080,7 +1147,9 @@ define <2 x i1> @test_fcmp_uno(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_uno_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_uno_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.nan.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.nan.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1102,7 +1171,9 @@ define <2 x i1> @test_fcmp_one(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_one_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_one_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.ne.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.ne.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1124,7 +1195,9 @@ define <2 x i1> @test_fcmp_oeq(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oeq_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oeq_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.eq.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.eq.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1146,7 +1219,9 @@ define <2 x i1> @test_fcmp_ogt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ogt_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ogt_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.gt.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.gt.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1168,7 +1243,9 @@ define <2 x i1> @test_fcmp_oge(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_oge_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_oge_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.ge.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.ge.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1190,7 +1267,9 @@ define <2 x i1> @test_fcmp_olt(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_olt_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_olt_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.lt.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.lt.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1212,7 +1291,9 @@ define <2 x i1> @test_fcmp_ole(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ole_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ole_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.le.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.le.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1234,7 +1315,9 @@ define <2 x i1> @test_fcmp_ord(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fcmp_ord_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fcmp_ord_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: setp.num.f32 %p1, %r2, %r4;
; CHECK-NEXT: setp.num.f32 %p2, %r1, %r3;
; CHECK-NEXT: selp.b16 %rs1, -1, 0, %p2;
@@ -1254,6 +1337,7 @@ define <2 x i32> @test_fptosi_i32(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i32_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1270,6 +1354,7 @@ define <2 x i64> @test_fptosi_i64(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptosi_i64_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1286,6 +1371,7 @@ define <2 x i32> @test_fptoui_2xi32(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi32_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1302,6 +1388,7 @@ define <2 x i64> @test_fptoui_2xi64(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fptoui_2xi64_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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,11 +1467,12 @@ 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: 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: add.rn.f32 %r7, %r6, %r4;
-; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r5, %r3;
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_uitofp_2xi32_fadd_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd1, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r5, %r1;
+; CHECK-NOF32X2-NEXT: cvt.rn.f32.u32 %r6, %r2;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r7, %r4, %r6;
+; CHECK-NOF32X2-NEXT: add.rn.f32 %r8, %r3, %r5;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -1432,6 +1520,7 @@ define <2 x double> @test_fpext_2xdouble(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fpext_2xdouble_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: cvt.f64.f32 %rd2, %r2;
; CHECK-NEXT: cvt.f64.f32 %rd3, %r1;
; CHECK-NEXT: st.param.v2.b64 [func_retval0], {%rd3, %rd2};
@@ -1500,6 +1589,7 @@ define <2 x float> @test_sqrt(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sqrt_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: sqrt.rn.f32 %r3, %r2;
; CHECK-NEXT: sqrt.rn.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1523,6 +1613,7 @@ define <2 x float> @test_sin(<2 x float> %a) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_sin_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: sin.approx.f32 %r3, %r2;
; CHECK-NEXT: sin.approx.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1539,6 +1630,7 @@ define <2 x float> @test_cos(<2 x float> %a) #0 #1 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_cos_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: cos.approx.f32 %r3, %r2;
; CHECK-NEXT: cos.approx.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1598,10 +1690,13 @@ define <2 x float> @test_fma(<2 x float> %a, <2 x float> %b, <2 x float> %c) #0
; 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: mov.b64 %rd1, {%r1, %r2};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fma_param_2];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd3, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fma_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r5, %r6};
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r6, %r4;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r5, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -1628,6 +1723,7 @@ define <2 x float> @test_fabs(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_fabs_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: abs.f32 %r3, %r2;
; CHECK-NEXT: abs.f32 %r4, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r4, %r3};
@@ -1644,7 +1740,9 @@ define <2 x float> @test_minnum(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_minnum_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_minnum_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: min.f32 %r5, %r2, %r4;
; CHECK-NEXT: min.f32 %r6, %r1, %r3;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -1661,7 +1759,9 @@ define <2 x float> @test_maxnum(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_maxnum_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_maxnum_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: max.f32 %r5, %r2, %r4;
; CHECK-NEXT: max.f32 %r6, %r1, %r3;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -1678,7 +1778,9 @@ define <2 x float> @test_copysign(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: copysign.f32 %r5, %r4, %r2;
; CHECK-NEXT: copysign.f32 %r6, %r3, %r1;
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r6, %r5};
@@ -1695,8 +1797,9 @@ define <2 x float> @test_copysign_f64(<2 x float> %a, <2 x double> %b) #0 {
; CHECK-NEXT: .reg .b64 %rd<8>;
; 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: mov.b64 %rd1, {%r1, %r2};
+; CHECK-NEXT: ld.param.v2.b64 {%rd2, %rd3}, [test_copysign_f64_param_1];
; CHECK-NEXT: abs.f32 %r3, %r2;
; CHECK-NEXT: neg.f32 %r4, %r3;
; CHECK-NEXT: shr.u64 %rd4, %rd3, 63;
@@ -1724,7 +1827,9 @@ define <2 x double> @test_copysign_extended(<2 x float> %a, <2 x float> %b) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_copysign_extended_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_copysign_extended_param_1];
+; CHECK-NEXT: mov.b64 %rd2, {%r3, %r4};
; CHECK-NEXT: copysign.f32 %r5, %r3, %r1;
; CHECK-NEXT: copysign.f32 %r6, %r4, %r2;
; CHECK-NEXT: cvt.f64.f32 %rd3, %r6;
@@ -1744,6 +1849,7 @@ define <2 x float> @test_floor(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_floor_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1760,6 +1866,7 @@ define <2 x float> @test_ceil(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_ceil_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1776,6 +1883,7 @@ define <2 x float> @test_trunc(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_trunc_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1792,6 +1900,7 @@ define <2 x float> @test_rint(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_rint_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1808,6 +1917,7 @@ define <2 x float> @test_nearbyint(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_nearbyint_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1824,6 +1934,7 @@ define <2 x float> @test_roundeven(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_roundeven_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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};
@@ -1842,6 +1953,7 @@ define <2 x float> @test_round(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_round_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; CHECK-NEXT: and.b32 %r3, %r2, -2147483648;
; CHECK-NEXT: or.b32 %r4, %r3, 1056964608;
; CHECK-NEXT: add.rn.f32 %r5, %r2, %r4;
@@ -1876,10 +1988,13 @@ define <2 x float> @test_fmuladd(<2 x float> %a, <2 x float> %b, <2 x float> %c)
; 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: mov.b64 %rd1, {%r1, %r2};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r3, %r4}, [test_fmuladd_param_2];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd3, {%r3, %r4};
+; CHECK-NOF32X2-NEXT: ld.param.v2.b32 {%r5, %r6}, [test_fmuladd_param_1];
+; CHECK-NOF32X2-NEXT: mov.b64 %rd2, {%r5, %r6};
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r7, %r2, %r6, %r4;
+; CHECK-NOF32X2-NEXT: fma.rn.f32 %r8, %r1, %r5, %r3;
; CHECK-NOF32X2-NEXT: st.param.v2.b32 [func_retval0], {%r8, %r7};
; CHECK-NOF32X2-NEXT: ret;
;
@@ -1906,6 +2021,7 @@ define <2 x float> @test_shufflevector(<2 x float> %a) #0 {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v2.b32 {%r1, %r2}, [test_shufflevector_param_0];
+; CHECK-NEXT: mov.b64 %rd1, {%r1, %r2};
; 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>
@@ -1919,8 +2035,9 @@ define <2 x float> @test_insertelement(<2 x float> %a, float %x) #0 {
; 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: mov.b64 %rd1, {%r2, %r3};
+; CHECK-NEXT: ld.param.b32 %r1, [test_insertelement_param_1];
; CHECK-NEXT: st.param.v2.b32 [func_retval0], {%r2, %r1};
; CHECK-NEXT: ret;
%i = insertelement <2 x float> %a, float %x, i64 1
diff --git a/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
index 1a61498b10142..937d5427940eb 100644
--- a/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i16x2-instructions.ll
@@ -39,6 +39,7 @@ define i16 @test_extract_0(<2 x i16> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_0_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: cvt.u32.u16 %r2, %rs1;
; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
; COMMON-NEXT: ret;
@@ -54,6 +55,7 @@ define i16 @test_extract_1(<2 x i16> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_extract_1_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: cvt.u32.u16 %r2, %rs2;
; COMMON-NEXT: st.param.b32 [func_retval0], %r2;
; COMMON-NEXT: ret;
@@ -70,8 +72,9 @@ define i16 @test_extract_i(<2 x i16> %a, i64 %idx) #0 {
; COMMON-NEXT: .reg .b64 %rd<2>;
; 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: mov.b32 %r1, {%rs1, %rs2};
+; COMMON-NEXT: ld.param.b64 %rd1, [test_extract_i_param_1];
; COMMON-NEXT: setp.eq.b64 %p1, %rd1, 0;
; COMMON-NEXT: selp.b16 %rs3, %rs1, %rs2, %p1;
; COMMON-NEXT: cvt.u32.u16 %r2, %rs3;
@@ -100,7 +103,9 @@ define <2 x i16> @test_add(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_add_param_0];
+; NO-I16x2-NEXT: mov.b32 %r1, {%rs1, %rs2};
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_add_param_1];
+; NO-I16x2-NEXT: mov.b32 %r2, {%rs3, %rs4};
; NO-I16x2-NEXT: add.s16 %rs5, %rs2, %rs4;
; NO-I16x2-NEXT: add.s16 %rs6, %rs1, %rs3;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -129,6 +134,7 @@ define <2 x i16> @test_add_imm_0(<2 x i16> %a) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -156,6 +162,7 @@ define <2 x i16> @test_add_imm_1(<2 x i16> %a) #0 {
; 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: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -172,7 +179,9 @@ define <2 x i16> @test_sub(<2 x i16> %a, <2 x i16> %b) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_sub_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_sub_param_1];
+; COMMON-NEXT: mov.b32 %r2, {%rs3, %rs4};
; COMMON-NEXT: sub.s16 %rs5, %rs2, %rs4;
; COMMON-NEXT: sub.s16 %rs6, %rs1, %rs3;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -200,7 +209,9 @@ define <2 x i16> @test_smax(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_smax_param_0];
+; NO-I16x2-NEXT: mov.b32 %r1, {%rs1, %rs2};
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_smax_param_1];
+; NO-I16x2-NEXT: mov.b32 %r2, {%rs3, %rs4};
; NO-I16x2-NEXT: max.s16 %rs5, %rs2, %rs4;
; NO-I16x2-NEXT: max.s16 %rs6, %rs1, %rs3;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -229,7 +240,9 @@ define <2 x i16> @test_umax(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_umax_param_0];
+; NO-I16x2-NEXT: mov.b32 %r1, {%rs1, %rs2};
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_umax_param_1];
+; NO-I16x2-NEXT: mov.b32 %r2, {%rs3, %rs4};
; NO-I16x2-NEXT: max.u16 %rs5, %rs2, %rs4;
; NO-I16x2-NEXT: max.u16 %rs6, %rs1, %rs3;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -258,7 +271,9 @@ define <2 x i16> @test_smin(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_smin_param_0];
+; NO-I16x2-NEXT: mov.b32 %r1, {%rs1, %rs2};
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_smin_param_1];
+; NO-I16x2-NEXT: mov.b32 %r2, {%rs3, %rs4};
; NO-I16x2-NEXT: min.s16 %rs5, %rs2, %rs4;
; NO-I16x2-NEXT: min.s16 %rs6, %rs1, %rs3;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -287,7 +302,9 @@ define <2 x i16> @test_umin(<2 x i16> %a, <2 x i16> %b) #0 {
; NO-I16x2-EMPTY:
; NO-I16x2-NEXT: // %bb.0:
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_umin_param_0];
+; NO-I16x2-NEXT: mov.b32 %r1, {%rs1, %rs2};
; NO-I16x2-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_umin_param_1];
+; NO-I16x2-NEXT: mov.b32 %r2, {%rs3, %rs4};
; NO-I16x2-NEXT: min.u16 %rs5, %rs2, %rs4;
; NO-I16x2-NEXT: min.u16 %rs6, %rs1, %rs3;
; NO-I16x2-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -305,7 +322,9 @@ define <2 x i16> @test_mul(<2 x i16> %a, <2 x i16> %b) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_mul_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_mul_param_1];
+; COMMON-NEXT: mov.b32 %r2, {%rs3, %rs4};
; COMMON-NEXT: mul.lo.s16 %rs5, %rs2, %rs4;
; COMMON-NEXT: mul.lo.s16 %rs6, %rs1, %rs3;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -687,11 +706,15 @@ define <2 x i16> @test_select_cc(<2 x i16> %a, <2 x i16> %b, <2 x i16> %c, <2 x
; 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: mov.b32 %r1, {%rs1, %rs2};
+; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_param_3];
+; COMMON-NEXT: mov.b32 %r4, {%rs3, %rs4};
+; COMMON-NEXT: ld.param.v2.b16 {%rs5, %rs6}, [test_select_cc_param_2];
+; COMMON-NEXT: mov.b32 %r3, {%rs5, %rs6};
; COMMON-NEXT: ld.param.v2.b16 {%rs7, %rs8}, [test_select_cc_param_1];
+; COMMON-NEXT: mov.b32 %r2, {%rs7, %rs8};
+; COMMON-NEXT: setp.ne.b16 %p1, %rs5, %rs3;
+; COMMON-NEXT: setp.ne.b16 %p2, %rs6, %rs4;
; COMMON-NEXT: selp.b16 %rs9, %rs2, %rs8, %p2;
; COMMON-NEXT: selp.b16 %rs10, %rs1, %rs7, %p1;
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs10, %rs9};
@@ -711,10 +734,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.v2.b16 {%rs1, %rs2}, [test_select_cc_i32_i16_param_3];
+; COMMON-NEXT: mov.b32 %r6, {%rs1, %rs2};
+; COMMON-NEXT: ld.param.v2.b16 {%rs3, %rs4}, [test_select_cc_i32_i16_param_2];
+; COMMON-NEXT: mov.b32 %r5, {%rs3, %rs4};
+; 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};
@@ -733,12 +758,14 @@ define <2 x i16> @test_select_cc_i16_i32(<2 x i16> %a, <2 x i16> %b,
; COMMON-NEXT: .reg .b32 %r<7>;
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
+; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_select_cc_i16_i32_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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.v2.b16 {%rs3, %rs4}, [test_select_cc_i16_i32_param_1];
+; COMMON-NEXT: mov.b32 %r2, {%rs3, %rs4};
; 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: st.param.v2.b16 [func_retval0], {%rs6, %rs5};
@@ -852,6 +879,7 @@ define <2 x i32> @test_zext_2xi32(<2 x i16> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_zext_2xi32_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: cvt.u32.u16 %r2, %rs2;
; COMMON-NEXT: cvt.u32.u16 %r3, %rs1;
; COMMON-NEXT: st.param.v2.b32 [func_retval0], {%r3, %r2};
@@ -869,6 +897,7 @@ define <2 x i64> @test_zext_2xi64(<2 x i16> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_zext_2xi64_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; COMMON-NEXT: cvt.u64.u16 %rd1, %rs2;
; COMMON-NEXT: cvt.u64.u16 %rd2, %rs1;
; COMMON-NEXT: st.param.v2.b64 [func_retval0], {%rd2, %rd1};
@@ -927,6 +956,7 @@ define <2 x i16> @test_shufflevector(<2 x i16> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_shufflevector_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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>
@@ -940,8 +970,9 @@ define <2 x i16> @test_insertelement(<2 x i16> %a, i16 %x) #0 {
; 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: mov.b32 %r1, {%rs2, %rs3};
+; COMMON-NEXT: ld.param.b16 %rs1, [test_insertelement_param_1];
; COMMON-NEXT: st.param.v2.b16 [func_retval0], {%rs2, %rs1};
; COMMON-NEXT: ret;
%i = insertelement <2 x i16> %a, i16 %x, i64 1
@@ -956,6 +987,7 @@ define <2 x i16> @test_fptosi_2xhalf_to_2xi16(<2 x half> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptosi_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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};
@@ -972,6 +1004,7 @@ define <2 x i16> @test_fptoui_2xhalf_to_2xi16(<2 x half> %a) #0 {
; COMMON-EMPTY:
; COMMON-NEXT: // %bb.0:
; COMMON-NEXT: ld.param.v2.b16 {%rs1, %rs2}, [test_fptoui_2xhalf_to_2xi16_param_0];
+; COMMON-NEXT: mov.b32 %r1, {%rs1, %rs2};
; 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 cbc9f700b1f01..c9eb2362f9e5f 100644
--- a/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
+++ b/llvm/test/CodeGen/NVPTX/i8x4-instructions.ll
@@ -1936,6 +1936,8 @@ define <4 x i8> @test_fptosi_4xhalf_to_4xi8(<4 x half> %a) #0 {
; 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: mov.b32 %r2, {%rs3, %rs4};
+; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};
; O0-NEXT: cvt.rzi.s16.f16 %rs5, %rs4;
; O0-NEXT: cvt.rzi.s16.f16 %rs6, %rs3;
; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
@@ -1990,6 +1992,8 @@ define <4 x i8> @test_fptoui_4xhalf_to_4xi8(<4 x half> %a) #0 {
; 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: mov.b32 %r2, {%rs3, %rs4};
+; O0-NEXT: mov.b32 %r1, {%rs1, %rs2};
; O0-NEXT: cvt.rzi.u16.f16 %rs5, %rs4;
; O0-NEXT: cvt.rzi.u16.f16 %rs6, %rs3;
; O0-NEXT: mov.b32 %r3, {%rs6, %rs5};
diff --git a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
index 87f965c84b6b6..0938ec10e20e3 100644
--- a/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
+++ b/llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll
@@ -118,7 +118,11 @@ define float @reduce_fadd_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fadd_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: add.rn.f32 %r9, %r5, 0f00000000;
; CHECK-NEXT: add.rn.f32 %r10, %r9, %r6;
; CHECK-NEXT: add.rn.f32 %r11, %r10, %r7;
@@ -141,7 +145,11 @@ define float @reduce_fadd_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-SM80-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fadd_float_reassoc_param_0];
+; CHECK-SM80-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-SM80-NEXT: mov.b64 %rd1, {%r5, %r6};
; 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;
@@ -322,7 +330,11 @@ define float @reduce_fmul_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmul_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: mul.rn.f32 %r9, %r5, %r6;
; CHECK-NEXT: mul.rn.f32 %r10, %r9, %r7;
; CHECK-NEXT: mul.rn.f32 %r11, %r10, %r8;
@@ -344,7 +356,11 @@ define float @reduce_fmul_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-SM80-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-SM80-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmul_float_reassoc_param_0];
+; CHECK-SM80-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-SM80-NEXT: mov.b64 %rd1, {%r5, %r6};
; 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;
@@ -495,7 +511,11 @@ define float @reduce_fmax_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmax_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: max.f32 %r9, %r8, %r4;
; CHECK-NEXT: max.f32 %r10, %r6, %r2;
; CHECK-NEXT: max.f32 %r11, %r10, %r9;
@@ -518,7 +538,11 @@ define float @reduce_fmax_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmax_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: max.f32 %r9, %r8, %r4;
; CHECK-NEXT: max.f32 %r10, %r6, %r2;
; CHECK-NEXT: max.f32 %r11, %r10, %r9;
@@ -629,7 +653,11 @@ define float @reduce_fmin_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmin_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: min.f32 %r9, %r8, %r4;
; CHECK-NEXT: min.f32 %r10, %r6, %r2;
; CHECK-NEXT: min.f32 %r11, %r10, %r9;
@@ -652,7 +680,11 @@ define float @reduce_fmin_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmin_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: min.f32 %r9, %r8, %r4;
; CHECK-NEXT: min.f32 %r10, %r6, %r2;
; CHECK-NEXT: min.f32 %r11, %r10, %r9;
@@ -763,7 +795,11 @@ define float @reduce_fmaximum_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fmaximum_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: max.NaN.f32 %r9, %r8, %r4;
; CHECK-NEXT: max.NaN.f32 %r10, %r6, %r2;
; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r9;
@@ -786,7 +822,11 @@ define float @reduce_fmaximum_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fmaximum_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: max.NaN.f32 %r9, %r8, %r4;
; CHECK-NEXT: max.NaN.f32 %r10, %r6, %r2;
; CHECK-NEXT: max.NaN.f32 %r11, %r10, %r9;
@@ -897,7 +937,11 @@ define float @reduce_fminimum_float(<8 x float> %in) {
; CHECK-EMPTY:
; CHECK-NEXT: // %bb.0:
; CHECK-NEXT: ld.param.v4.b32 {%r1, %r2, %r3, %r4}, [reduce_fminimum_float_param_0+16];
+; CHECK-NEXT: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: min.NaN.f32 %r9, %r8, %r4;
; CHECK-NEXT: min.NaN.f32 %r10, %r6, %r2;
; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r9;
@@ -920,7 +964,11 @@ define float @reduce_fminimum_float_reassoc(<8 x float> %in) {
; 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: mov.b64 %rd4, {%r3, %r4};
+; CHECK-NEXT: mov.b64 %rd3, {%r1, %r2};
; CHECK-NEXT: ld.param.v4.b32 {%r5, %r6, %r7, %r8}, [reduce_fminimum_float_reassoc_param_0];
+; CHECK-NEXT: mov.b64 %rd2, {%r7, %r8};
+; CHECK-NEXT: mov.b64 %rd1, {%r5, %r6};
; CHECK-NEXT: min.NaN.f32 %r9, %r8, %r4;
; CHECK-NEXT: min.NaN.f32 %r10, %r6, %r2;
; CHECK-NEXT: min.NaN.f32 %r11, %r10, %r9;
diff --git a/llvm/test/CodeGen/pr126337.ll b/llvm/test/CodeGen/pr126337.ll
new file mode 100644
index 0000000000000..32e411584b0e5
--- /dev/null
+++ b/llvm/test/CodeGen/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
+}
+
More information about the llvm-commits
mailing list