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

Author: Prince781

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

@@ -2861,9 +2861,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;",
@@ -2904,6 +2913,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 {%f1, tmp}, %rd1; }
+; CHECK-NEXT:    mov.b64 {%f1, _}, %rd1;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %f1;
 ; 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, %f1}, %rd1; }
+; CHECK-NEXT:    mov.b64 {_, %f1}, %rd1;
 ; CHECK-NEXT:    st.param.b32 [func_retval0], %f1;
 ; CHECK-NEXT:    ret;
   %e = extractelement <2 x float> %a, i32 1