[NVPTX] Remove few more unneeded fp16 instruction variants

[tra]

Cleanup.

llvm/lib/Target/NVPTX/NVPTXISelDAGToDAG.cpp

for (auto *U : Vector.getNode()->uses()) {
     }
   }
 
   // There's no point scattering f16x2 if we only ever access one
   // element of it.
   if (E0.empty() || E1.empty())
     return false;
 
-  unsigned Op = NVPTX::SplitF16x2;
-  // If the vector has been BITCAST'ed from i32, we can use original
-  // value directly and avoid register-to-register move.
-  SDValue Source = Vector;
-  if (Vector->getOpcode() == ISD::BITCAST) {
-    Op = NVPTX::SplitI32toF16x2;
-    Source = Vector->getOperand(0);
-  }
   // Merge (f16 extractelt(V, 0), f16 extractelt(V,1))
   // into f16,f16 SplitF16x2(V)
-  SDNode *ScatterOp =
-      CurDAG->getMachineNode(Op, SDLoc(N), MVT::f16, MVT::f16, Source);
+  SDNode *ScatterOp = CurDAG->getMachineNode(NVPTX::I32toV2I16, SDLoc(N),
+                                             MVT::f16, MVT::f16, Vector);
   for (auto *Node : E0)
     ReplaceUses(SDValue(Node, 0), SDValue(ScatterOp, 0));
   for (auto *Node : E1)
     ReplaceUses(SDValue(Node, 0), SDValue(ScatterOp, 1));
 
   return true;
 }
 

llvm/lib/Target/NVPTX/NVPTXInstrInfo.td

def _v4_asi : NVPTXInst<
     "ld${isVol:volatile}${addsp:addsp}${Vec:vec}.${Sign:sign}$fromWidth "
     "\t{{$dst1, $dst2, $dst3, $dst4}}, [$addr+$offset];", []>;
 }
 let mayLoad=1, hasSideEffects=0 in {
   defm LDV_i8  : LD_VEC<Int16Regs>;
   defm LDV_i16 : LD_VEC<Int16Regs>;
   defm LDV_i32 : LD_VEC<Int32Regs>;
   defm LDV_i64 : LD_VEC<Int64Regs>;
-  defm LDV_f16 : LD_VEC<Int16Regs>;
-  defm LDV_f16x2 : LD_VEC<Int32Regs>;
   defm LDV_f32 : LD_VEC<Float32Regs>;
   defm LDV_f64 : LD_VEC<Float64Regs>;
 }
 
 multiclass ST_VEC<NVPTXRegClass regclass> {
   def _v2_avar : NVPTXInst<
     (outs),
     (ins regclass:$src1, regclass:$src2, LdStCode:$isVol, LdStCode:$addsp,
def _v4_asi : NVPTXInst<
     "$fromWidth \t[$addr+$offset], {{$src1, $src2, $src3, $src4}};", []>;
 }
 
 let mayStore=1, hasSideEffects=0 in {
   defm STV_i8  : ST_VEC<Int16Regs>;
   defm STV_i16 : ST_VEC<Int16Regs>;
   defm STV_i32 : ST_VEC<Int32Regs>;
   defm STV_i64 : ST_VEC<Int64Regs>;
-  defm STV_f16 : ST_VEC<Int16Regs>;
-  defm STV_f16x2 : ST_VEC<Int32Regs>;
   defm STV_f32 : ST_VEC<Float32Regs>;
   defm STV_f64 : ST_VEC<Float64Regs>;
 }
 
 //---- Conversion ----
 
 class F_BITCONVERT<string SzStr, ValueType TIn, ValueType TOut,
   NVPTXRegClass regclassIn = ValueToRegClass<TIn>.ret,
let hasSideEffects = false in {
   def F64toV2F32 : NVPTXInst<(outs Float32Regs:$d1, Float32Regs:$d2),
                              (ins Float64Regs:$s),
                              "mov.b64 \t{{$d1, $d2}}, $s;", []>;
 
   def I32toI16H  : NVPTXInst<(outs Int16Regs:$high),
                              (ins Int32Regs:$s),
                              "{{ .reg .b16 tmp; mov.b32 {tmp, $high}, $s; }}",
                              []>;
+  def I32toI16L  : NVPTXInst<(outs Int16Regs:$low),
+                             (ins Int32Regs:$s),
+                             "{{ .reg .b16 tmp; mov.b32 {$low, tmp}, $s; }}",
+                             []>;
   def I64toI32H  : NVPTXInst<(outs Int32Regs:$high),
                              (ins Int64Regs:$s),
                              "{{ .reg .b32 tmp; mov.b64 {tmp, $high}, $s; }}",
                              []>;
 }
 
 // Using partial vectorized move produces better SASS code for extraction of
 // upper/lower parts of an integer.
 def : Pat<(i16 (trunc (srl Int32Regs:$s, (i32 16)))),
           (I32toI16H Int32Regs:$s)>;
 def : Pat<(i16 (trunc (sra Int32Regs:$s, (i32 16)))),
           (I32toI16H Int32Regs:$s)>;
 def : Pat<(i32 (trunc (srl Int64Regs:$s, (i32 32)))),
           (I64toI32H Int64Regs:$s)>;
 def : Pat<(i32 (trunc (sra Int64Regs:$s, (i32 32)))),
           (I64toI32H Int64Regs:$s)>;
 
-let hasSideEffects = false in {
-  // Extract element of f16x2 register. PTX does not provide any way
-  // to access elements of f16x2 vector directly, so we need to
-  // extract it using a temporary register.
-  def F16x2toF16_0 : NVPTXInst<(outs Int16Regs:$dst),
-                               (ins Int32Regs:$src),
-                               "{{ .reg .b16 \t%tmp_hi;\n\t"
-                               "  mov.b32 \t{$dst, %tmp_hi}, $src; }}",
-                               [(set Int16Regs:$dst,
-                                 (extractelt (v2f16 Int32Regs:$src), 0))]>;
-  def F16x2toF16_1 : NVPTXInst<(outs Int16Regs:$dst),
-                               (ins Int32Regs:$src),
-                               "{{ .reg .b16 \t%tmp_lo;\n\t"
-                               "  mov.b32 \t{%tmp_lo, $dst}, $src; }}",
-                               [(set Int16Regs:$dst,
-                                 (extractelt (v2f16 Int32Regs:$src), 1))]>;
-
-  // Coalesce two f16 registers into f16x2
-  def BuildF16x2 : NVPTXInst<(outs Int32Regs:$dst),
-                             (ins Int16Regs:$a, Int16Regs:$b),
-                             "mov.b32 \t$dst, {{$a, $b}};",
-                             [(set (v2f16 Int32Regs:$dst),
-                               (build_vector (f16 Int16Regs:$a), (f16 Int16Regs:$b)))]>;
-
-  // Directly initializing underlying the b32 register is one less SASS
-  // instruction than than vector-packing move.
-  def BuildF16x2i : NVPTXInst<(outs Int32Regs:$dst), (ins i32imm:$src),
-                              "mov.b32 \t$dst, $src;",
-                              []>;
-
-  // Split f16x2 into two f16 registers.
-  def SplitF16x2  : NVPTXInst<(outs Int16Regs:$lo, Int16Regs:$hi),
-                              (ins Int32Regs:$src),
-                              "mov.b32 \t{{$lo, $hi}}, $src;",
-                              []>;
-  // Split an i32 into two f16
-  def SplitI32toF16x2  : NVPTXInst<(outs Int16Regs:$lo, Int16Regs:$hi),
-                                   (ins Int32Regs:$src),
-                                   "mov.b32 \t{{$lo, $hi}}, $src;",
-                                   []>;
-}
+def : Pat<(f16 (extractelt (v2f16 Int32Regs:$src), 0)),
+          (I32toI16L Int32Regs:$src)>;
+def : Pat<(f16 (extractelt (v2f16 Int32Regs:$src), 1)),
+          (I32toI16H Int32Regs:$src)>;
+def : Pat<(v2f16 (build_vector (f16 Int16Regs:$a), (f16 Int16Regs:$b))),
+          (V2I16toI32 Int16Regs:$a, Int16Regs:$b)>;
 
 // Count leading zeros
 let hasSideEffects = false in {
   def CLZr32 : NVPTXInst<(outs Int32Regs:$d), (ins Int32Regs:$a),
                          "clz.b32 \t$d, $a;", []>;
   def CLZr64 : NVPTXInst<(outs Int32Regs:$d), (ins Int64Regs:$a),
                          "clz.b64 \t$d, $a;", []>;
 }

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

 ; CHECK:     st.param.b32    [func_retval0+0], [[R]];
 ; CHECK-NEXT: ret;
 define <2 x half> @test_ret_const() #0 {
   ret <2 x half> <half 1.0, half 2.0>
 }
 
 ; CHECK-LABEL: test_extract_0(
 ; CHECK:      ld.param.b32    [[A:%r[0-9]+]], [test_extract_0_param_0];
-; CHECK:      mov.b32         {[[R:%rs[0-9]+]], %tmp_hi}, [[A]];
+; CHECK:      mov.b32         {[[R:%rs[0-9]+]], tmp}, [[A]];
 ; CHECK:      st.param.b16    [func_retval0+0], [[R]];
 ; CHECK:      ret;
 define half @test_extract_0(<2 x half> %a) #0 {
   %e = extractelement <2 x half> %a, i32 0
   ret half %e
 }
 
 ; CHECK-LABEL: test_extract_1(
 ; CHECK:      ld.param.b32    [[A:%r[0-9]+]], [test_extract_1_param_0];
-; CHECK:      mov.b32         {%tmp_lo, [[R:%rs[0-9]+]]}, [[A]];
+; CHECK:      mov.b32         {tmp, [[R:%rs[0-9]+]]}, [[A]];
 ; CHECK:      st.param.b16    [func_retval0+0], [[R]];
 ; CHECK:      ret;
 define half @test_extract_1(<2 x half> %a) #0 {
   %e = extractelement <2 x half> %a, i32 1
   ret half %e
 }
 
 ; CHECK-LABEL: test_extract_i(

 ; CHECK: mov.b32 {%rs1, %rs2}, %r1;
 ; CHECK: mov.b32 %r2, {%rs2, %rs1};
 define <2 x half> @test_shufflevector(<2 x half> %a) #0 {
   %s = shufflevector <2 x half> %a, <2 x half> undef, <2 x i32> <i32 1, i32 0>
   ret <2 x half> %s
 }
 
 ; CHECK-LABEL: test_insertelement(
-; CHECK: mov.b32 {%rs2, %tmp_hi}, %r1;
+; CHECK: mov.b32 {%rs2, tmp}, %r1;
 ; CHECK: mov.b32 %r2, {%rs2, %rs1};
 define <2 x half> @test_insertelement(<2 x half> %a, half %x) #0 {
   %i = insertelement <2 x half> %a, half %x, i64 1
   ret <2 x half> %i
 }
 
 attributes #0 = { nounwind }
 attributes #1 = { "unsafe-fp-math" = "true" }