Fixed calculation of constant v4i8 values.

Author: Artem-B

llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp

@@ -2229,6 +2229,10 @@ SDValue NVPTXTargetLowering::LowerBUILD_VECTOR(SDValue Op,
       Value = cast<ConstantSDNode>(Operand)->getAPIntValue();
     else
       llvm_unreachable("Unsupported type");
+    // i8 values are carried around as i16, so we need to zero out upper bits,
+    // so they do not get in the way of combining individual byte values
+    if (VT == MVT::v4i8)
+      Value = Value.trunc(8);
     return Value.zext(32);
   };
   APInt Value;

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

@@ -17,10 +17,10 @@ define <4 x i8> @test_ret_const() #0 {
 ; CHECK-NEXT:    .reg .b32 %r<3>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
-; CHECK-NEXT:    mov.u32 %r1, 67305985;
+; CHECK-NEXT:    mov.u32 %r1, -66911489;
 ; CHECK-NEXT:    st.param.b32 [func_retval0+0], %r1;
 ; CHECK-NEXT:    ret;
-  ret <4 x i8> <i8 1, i8 2, i8 3, i8 4>
+  ret <4 x i8> <i8 -1, i8 2, i8 3, i8 -4>
 }
 
 define i8 @test_extract_0(<4 x i8> %a) #0 {