[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