[NVPTX] use sink symbol for single-element unpacking of v2f32s

Author: Prince781

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

@@ -2933,9 +2933,18 @@ let hasSideEffects = false in {
                              (ins Int64Regs:$s),
                              "{{ .reg .b32 tmp; mov.b64 {$low, tmp}, $s; }}",
                              []>;
-
   // PTX 7.1 lets you avoid a temp register and just use _ as a "sink" for the
   // unused high/low part.
+  def I64toF32H_Sink : NVPTXInst<(outs Float32Regs:$high),
+                             (ins Int64Regs:$s),
+                             "mov.b64 {{_, $high}}, $s;",
+                             []>,
+                             Requires<[hasPTX<71>]>;
+  def I64toF32L_Sink : NVPTXInst<(outs Float32Regs:$low),
+                             (ins Int64Regs:$s),
+                             "mov.b64 {{$low, _}}, $s;",
+                             []>,
+                             Requires<[hasPTX<71>]>;
   def I32toI16H_Sink  : NVPTXInst<(outs Int16Regs:$high),
                              (ins Int32Regs:$s),
                              "mov.b32 \t{{_, $high}}, $s;",
@@ -2976,6 +2985,11 @@ foreach vt = [v2f16, v2bf16, v2i16] in {
   def : Pat<(extractelt vt:$src, 1), (I32toI16H $src)>;
 }
 
+def : Pat<(extractelt v2f32:$src, 0),
+          (I64toF32L_Sink $src)>, Requires<[hasPTX<71>]>;
+def : Pat<(extractelt v2f32:$src, 1),
+          (I64toF32H_Sink $src)>, Requires<[hasPTX<71>]>;
+
 def : Pat<(extractelt v2f32:$src, 0),
           (I64toF32L $src)>;
 def : Pat<(extractelt v2f32:$src, 1),

llvm/test/CodeGen/NVPTX/f32x2-instructions.ll

@@ -33,7 +33,7 @@ define float @test_extract_0(<2 x float> %a) #0 {
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [test_extract_0_param_0];
-; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {%r1, tmp}, %rd1; }
+; CHECK-NEXT:    mov.b64 {%r1, _}, %rd1;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
 ; CHECK-NEXT:    ret;
   %e = extractelement <2 x float> %a, i32 0
@@ -48,7 +48,7 @@ define float @test_extract_1(<2 x float> %a) #0 {
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [test_extract_1_param_0];
-; CHECK-NEXT:    { .reg .b32 tmp; mov.b64 {tmp, %r1}, %rd1; }
+; CHECK-NEXT:    mov.b64 {_, %r1}, %rd1;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %r1;
 ; CHECK-NEXT:    ret;
   %e = extractelement <2 x float> %a, i32 1

llvm/test/CodeGen/NVPTX/ldg-invariant.ll

@@ -128,14 +128,14 @@ define half @ld_global_v8f16(ptr addrspace(1) %ptr) {
 define float @ld_global_v2f32(ptr addrspace(1) %ptr) {
 ; CHECK-LABEL: ld_global_v2f32(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %f<4>;
+; CHECK-NEXT:    .reg .b32 %r<4>;
 ; CHECK-NEXT:    .reg .b64 %rd<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [ld_global_v2f32_param_0];
-; CHECK-NEXT:    ld.global.nc.v2.b32 {%f1, %f2}, [%rd1];
-; CHECK-NEXT:    add.rn.f32 %f3, %f1, %f2;
-; CHECK-NEXT:    st.param.b32 [func_retval0], %f3;
+; CHECK-NEXT:    ld.global.nc.v2.b32 {%r1, %r2}, [%rd1];
+; CHECK-NEXT:    add.rn.f32 %r3, %r1, %r2;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r3;
 ; CHECK-NEXT:    ret;
   %a = load <2 x float>, ptr addrspace(1) %ptr, !invariant.load !0
   %v1 = extractelement <2 x float> %a, i32 0
@@ -147,16 +147,16 @@ define float @ld_global_v2f32(ptr addrspace(1) %ptr) {
 define float @ld_global_v4f32(ptr addrspace(1) %ptr) {
 ; CHECK-LABEL: ld_global_v4f32(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %f<8>;
+; CHECK-NEXT:    .reg .b32 %r<8>;
 ; CHECK-NEXT:    .reg .b64 %rd<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [ld_global_v4f32_param_0];
-; CHECK-NEXT:    ld.global.nc.v4.b32 {%f1, %f2, %f3, %f4}, [%rd1];
-; CHECK-NEXT:    add.rn.f32 %f5, %f1, %f2;
-; CHECK-NEXT:    add.rn.f32 %f6, %f3, %f4;
-; CHECK-NEXT:    add.rn.f32 %f7, %f5, %f6;
-; CHECK-NEXT:    st.param.b32 [func_retval0], %f7;
+; CHECK-NEXT:    ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1];
+; CHECK-NEXT:    add.rn.f32 %r5, %r1, %r2;
+; CHECK-NEXT:    add.rn.f32 %r6, %r3, %r4;
+; CHECK-NEXT:    add.rn.f32 %r7, %r5, %r6;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r7;
 ; CHECK-NEXT:    ret;
   %a = load <4 x float>, ptr addrspace(1) %ptr, !invariant.load !0
   %v1 = extractelement <4 x float> %a, i32 0
@@ -172,17 +172,17 @@ define float @ld_global_v4f32(ptr addrspace(1) %ptr) {
 define float @ld_global_v8f32(ptr addrspace(1) %ptr) {
 ; CHECK-LABEL: ld_global_v8f32(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %f<12>;
+; CHECK-NEXT:    .reg .b32 %r<12>;
 ; CHECK-NEXT:    .reg .b64 %rd<2>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.b64 %rd1, [ld_global_v8f32_param_0];
-; CHECK-NEXT:    ld.global.nc.v4.b32 {%f1, %f2, %f3, %f4}, [%rd1+16];
-; CHECK-NEXT:    ld.global.nc.v4.b32 {%f5, %f6, %f7, %f8}, [%rd1];
-; CHECK-NEXT:    add.rn.f32 %f9, %f5, %f7;
-; CHECK-NEXT:    add.rn.f32 %f10, %f1, %f3;
-; CHECK-NEXT:    add.rn.f32 %f11, %f9, %f10;
-; CHECK-NEXT:    st.param.b32 [func_retval0], %f11;
+; CHECK-NEXT:    ld.global.nc.v4.b32 {%r1, %r2, %r3, %r4}, [%rd1+16];
+; CHECK-NEXT:    ld.global.nc.v4.b32 {%r5, %r6, %r7, %r8}, [%rd1];
+; CHECK-NEXT:    add.rn.f32 %r9, %r5, %r7;
+; CHECK-NEXT:    add.rn.f32 %r10, %r1, %r3;
+; CHECK-NEXT:    add.rn.f32 %r11, %r9, %r10;
+; CHECK-NEXT:    st.param.b32 [func_retval0], %r11;
 ; CHECK-NEXT:    ret;
   %a = load <8 x float>, ptr addrspace(1) %ptr, !invariant.load !0
   %v1 = extractelement <8 x float> %a, i32 0