[llvm-branch-commits] [llvm] [NVPTX] don't erase CopyToRegs when folding movs into loads (#149393) (PR #155685)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Wed Aug 27 13:10:19 PDT 2025
llvmbot wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-backend-nvptx
Author: Artem Belevich (Artem-B)
<details>
<summary>Changes</summary>
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
---
Patch is 150.86 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/155685.diff
7 Files Affected:
- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+1-11)
- (modified) llvm/test/CodeGen/NVPTX/f16x2-instructions.ll (+346-245)
- (modified) llvm/test/CodeGen/NVPTX/f32x2-instructions.ll (+397-246)
- (modified) llvm/test/CodeGen/NVPTX/i16x2-instructions.ll (+151-82)
- (modified) llvm/test/CodeGen/NVPTX/i8x4-instructions.ll (+22-18)
- (added) llvm/test/CodeGen/NVPTX/pr126337.ll (+41)
- (modified) llvm/test/CodeGen/NVPTX/reduction-intrinsics.ll (+137-89)
``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 14f05250ad6b8..5d496f69172c3 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
@@ -5008,12 +5008,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;
}))
@@ -5080,10 +5074,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);
}
@@ -6418,4 +6408,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, %rs...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/155685
More information about the llvm-branch-commits
mailing list