[llvm] [NVPTX][InferAS] assume alloca instructions are in local AS (PR #121710)
Artem Belevich via llvm-commits
llvm-commits at lists.llvm.org
Tue Jan 14 14:55:21 PST 2025
================
@@ -29,31 +29,32 @@ define dso_local noundef i32 @non_kernel_function(ptr nocapture noundef readonly
; PTX-NEXT: .reg .pred %p<2>;
; PTX-NEXT: .reg .b16 %rs<3>;
; PTX-NEXT: .reg .b32 %r<11>;
-; PTX-NEXT: .reg .b64 %rd<9>;
+; PTX-NEXT: .reg .b64 %rd<10>;
; PTX-EMPTY:
; PTX-NEXT: // %bb.0: // %entry
; PTX-NEXT: mov.u64 %SPL, __local_depot0;
; PTX-NEXT: cvta.local.u64 %SP, %SPL;
; PTX-NEXT: ld.param.u8 %rs1, [non_kernel_function_param_1];
; PTX-NEXT: and.b16 %rs2, %rs1, 1;
; PTX-NEXT: setp.eq.b16 %p1, %rs2, 1;
-; PTX-NEXT: ld.param.s32 %rd1, [non_kernel_function_param_2];
-; PTX-NEXT: ld.param.u64 %rd2, [non_kernel_function_param_0+8];
-; PTX-NEXT: st.u64 [%SP+8], %rd2;
-; PTX-NEXT: ld.param.u64 %rd3, [non_kernel_function_param_0];
-; PTX-NEXT: st.u64 [%SP], %rd3;
-; PTX-NEXT: mov.u64 %rd4, gi;
-; PTX-NEXT: cvta.global.u64 %rd5, %rd4;
-; PTX-NEXT: add.u64 %rd6, %SP, 0;
-; PTX-NEXT: selp.b64 %rd7, %rd6, %rd5, %p1;
-; PTX-NEXT: add.s64 %rd8, %rd7, %rd1;
-; PTX-NEXT: ld.u8 %r1, [%rd8];
-; PTX-NEXT: ld.u8 %r2, [%rd8+1];
+; PTX-NEXT: add.u64 %rd1, %SP, 0;
+; PTX-NEXT: add.u64 %rd2, %SPL, 0;
----------------
Artem-B wrote:
Ditto.
https://github.com/llvm/llvm-project/pull/121710
More information about the llvm-commits
mailing list