[llvm] [NVPTX] Switch to imm offset variants for LDG and LDU (PR #128270)
Artem Belevich via llvm-commits
llvm-commits at lists.llvm.org
Mon Feb 24 10:24:46 PST 2025
================
@@ -214,34 +214,33 @@ define dso_local i32 @bar() {
; CHECK-PTX-NEXT: .reg .b64 %SPL;
; CHECK-PTX-NEXT: .reg .b16 %rs<8>;
; CHECK-PTX-NEXT: .reg .b32 %r<4>;
-; CHECK-PTX-NEXT: .reg .b64 %rd<6>;
+; CHECK-PTX-NEXT: .reg .b64 %rd<5>;
; CHECK-PTX-EMPTY:
; CHECK-PTX-NEXT: // %bb.0: // %entry
; CHECK-PTX-NEXT: mov.u64 %SPL, __local_depot3;
; CHECK-PTX-NEXT: cvta.local.u64 %SP, %SPL;
; CHECK-PTX-NEXT: add.u64 %rd2, %SPL, 0;
-; CHECK-PTX-NEXT: mov.u64 %rd3, __const_$_bar_$_s1;
-; CHECK-PTX-NEXT: ld.global.nc.u8 %rs1, [%rd3+7];
+; CHECK-PTX-NEXT: ld.global.nc.u8 %rs1, [__const_$_bar_$_s1+7];
; CHECK-PTX-NEXT: cvt.u16.u8 %rs2, %rs1;
; CHECK-PTX-NEXT: st.local.u8 [%rd2+2], %rs2;
-; CHECK-PTX-NEXT: ld.global.nc.u8 %rs3, [%rd3+6];
+; CHECK-PTX-NEXT: ld.global.nc.u8 %rs3, [__const_$_bar_$_s1+6];
----------------
Artem-B wrote:
Given that the changes in this test file are largely incidental, we may want to add a few test cases to load/store tests that will exercise accessing data with ldu/ldg via symbol+offset.
https://github.com/llvm/llvm-project/pull/128270
More information about the llvm-commits
mailing list