[llvm] [NVPTX] Promote v2i8 to v2i16 (PR #111189)

[via llvm-commits]

llvmbot wrote:


<!--LLVM PR SUMMARY COMMENT-->

@llvm/pr-subscribers-backend-nvptx

Author: Manasij Mukherjee (manasij7479)

<details>
<summary>Changes</summary>

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

https://github.com/llvm/llvm-project/issues/104864

---
Full diff: https://github.com/llvm/llvm-project/pull/111189.diff


2 Files Affected:

- (modified) llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp (+4) 
- (modified) llvm/test/CodeGen/NVPTX/vector-returns.ll (+11-4) 


``````````diff
diff --git a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp b/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
index 8718b7890bf58a..57bc5fe0ac361c 100644
--- a/llvm/lib/Target/NVPTX/NVPTXISelLowering.cpp
+++ b/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);
diff --git a/llvm/test/CodeGen/NVPTX/vector-returns.ll b/llvm/test/CodeGen/NVPTX/vector-returns.ll
index 0d2ad2c9bee750..956f74392ae130 100644
--- a/llvm/test/CodeGen/NVPTX/vector-returns.ll
+++ b/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(

``````````

</details>


https://github.com/llvm/llvm-project/pull/111189