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

Drew Kersnar via llvm-commits llvm-commits at lists.llvm.org
Thu Dec 12 15:44:14 PST 2024


================
@@ -194,6 +194,156 @@ define void @generic_4xi8(ptr %a) {
   ret void
 }
 
+define void @generic_8xi8(ptr %a) {
+; CHECK-LABEL: generic_8xi8(
+; CHECK:       {
+; CHECK-NEXT:    .reg .b16 %rs<17>;
+; CHECK-NEXT:    .reg .b32 %r<25>;
+; CHECK-NEXT:    .reg .b64 %rd<2>;
+; CHECK-EMPTY:
+; CHECK-NEXT:  // %bb.0:
+; CHECK-NEXT:    ld.param.u64 %rd1, [generic_8xi8_param_0];
+; CHECK-NEXT:    ld.v2.b32 {%r1, %r2}, [%rd1];
+; CHECK-NEXT:    bfe.u32 %r3, %r2, 24, 8;
+; CHECK-NEXT:    cvt.u16.u32 %rs1, %r3;
+; CHECK-NEXT:    add.s16 %rs2, %rs1, 1;
+; CHECK-NEXT:    cvt.u32.u16 %r4, %rs2;
+; CHECK-NEXT:    bfe.u32 %r5, %r2, 16, 8;
+; CHECK-NEXT:    cvt.u16.u32 %rs3, %r5;
+; CHECK-NEXT:    add.s16 %rs4, %rs3, 1;
+; CHECK-NEXT:    cvt.u32.u16 %r6, %rs4;
----------------
dakersnar wrote:

Good idea. And especially since this affects the existing v4i8 lowering, I agree that this would be better in a separate patch. Do you want to open a tracking bug?

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


More information about the llvm-commits mailing list