[llvm] [NVPTX][InferAS] assume alloca instructions are in local AS (PR #121710)

Alex MacLean via llvm-commits llvm-commits at lists.llvm.org
Tue Jan 7 09:33:34 PST 2025


================
@@ -9,38 +9,43 @@ define dso_local void @caller_St8x4(ptr nocapture noundef readonly byval(%struct
 ; CHECK-NEXT:    .local .align 8 .b8 __local_depot0[32];
 ; CHECK-NEXT:    .reg .b32 %SP;
 ; CHECK-NEXT:    .reg .b32 %SPL;
-; CHECK-NEXT:    .reg .b32 %r<2>;
-; CHECK-NEXT:    .reg .b64 %rd<13>;
+; CHECK-NEXT:    .reg .b32 %r<4>;
+; CHECK-NEXT:    .reg .b64 %rd<17>;
 ; CHECK-EMPTY:
 ; CHECK-NEXT:  // %bb.0:
 ; CHECK-NEXT:    mov.u32 %SPL, __local_depot0;
 ; CHECK-NEXT:    cvta.local.u32 %SP, %SPL;
 ; CHECK-NEXT:    ld.param.u32 %r1, [caller_St8x4_param_1];
+; CHECK-NEXT:    add.u32 %r3, %SPL, 0;
----------------
AlexMaclean wrote:

We're now directly accessing the %SPL in the `st.local` and this is a quirk of how that is represented. Would be nice to fix for readability but likely not an issue for ptxas, and out of score for this MR. 

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


More information about the llvm-commits mailing list