[llvm] [NVPTX] Generalize and extend upsizing when lowering 8/16-bit-element vector loads/stores (PR #119622)

Artem Belevich via llvm-commits llvm-commits at lists.llvm.org
Mon Dec 16 11:58:25 PST 2024


================
@@ -6,19 +6,17 @@ target triple = "nvptx64-unknown-unknown"
 define void @kernel_func(ptr %in.vec, ptr %out.vec0) nounwind {
 ; CHECK-LABEL: kernel_func(
 ; CHECK:       {
-; CHECK-NEXT:    .reg .b32 %r<10>;
+; CHECK-NEXT:    .reg .b32 %r<14>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    ld.param.u32 %r1, [kernel_func_param_0];
-; CHECK-NEXT:    ld.u32 %r2, [%r1+8];
-; CHECK-NEXT:    ld.u32 %r3, [%r1];
-; CHECK-NEXT:    ld.u32 %r4, [%r1+24];
-; CHECK-NEXT:    ld.u32 %r5, [%r1+16];
-; CHECK-NEXT:    ld.param.u32 %r6, [kernel_func_param_1];
-; CHECK-NEXT:    prmt.b32 %r7, %r5, %r4, 0x4000U;
-; CHECK-NEXT:    prmt.b32 %r8, %r3, %r2, 0x40U;
-; CHECK-NEXT:    prmt.b32 %r9, %r8, %r7, 0x7610U;
-; CHECK-NEXT:    st.u32 [%r6], %r9;
+; CHECK-NEXT:    ld.v4.b32 {%r2, %r3, %r4, %r5}, [%r1];
+; CHECK-NEXT:    ld.v4.b32 {%r6, %r7, %r8, %r9}, [%r1+16];
----------------
Artem-B wrote:

SGTM. We'll cross that bridge if/when we get there.

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


More information about the llvm-commits mailing list