[NVPTX] Promote v2i8 to v2i16 (#111189)

Promote v2i8 to v2i16, fixes a crash. 
Re-enable a test in NVPTX/vector-returns.ll

#104864

Author: manasij7479

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

@@ -236,6 +236,10 @@ static void ComputePTXValueVTs(const TargetLowering &TLI, const DataLayout &DL,
         // v*i8 are formally lowered as v4i8
         EltVT = MVT::v4i8;
         NumElts = (NumElts + 3) / 4;
+      } else if (EltVT.getSimpleVT() == MVT::i8 && NumElts == 2) {
+        // v2i8 is promoted to v2i16
+        NumElts = 1;
+        EltVT = MVT::v2i16;
       }
       for (unsigned j = 0; j != NumElts; ++j) {
         ValueVTs.push_back(EltVT);

llvm/test/CodeGen/NVPTX/vector-returns.ll

@@ -325,10 +325,17 @@ define <3 x i8> @byte3() {
   ret <3 x i8> zeroinitializer
 }
 
-; FIXME: This test causes a crash. 
-; define <2 x i8> @byte2() {
-;   ret <2 x i8> zeroinitializer
-; }
+define <2 x i8> @byte2() {
+; CHECK-LABEL: byte2(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b32 %r<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    mov.b32 %r1, 0;
+; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r1;
+; CHECK-NEXT:    ret;
+  ret <2 x i8> zeroinitializer
+}
 
 define <1 x i8> @byte1() {
 ; CHECK-LABEL: byte1(