[llvm] [NVPTX] Lower allocas to the local address space (PR #204346)

[Tim Besard via llvm-commits]

================
@@ -21,18 +21,17 @@ define internal i32 @foo() {
 ; CHECK-NEXT:  prototype_0 : .callprototype (.param .b32 _) _ (.param .align 1 .b8 _[1], .param .b64 _);
 ; CHECK-NEXT:  // %bb.0: // %entry
 ; CHECK-NEXT:    mov.b64 %SPL, __local_depot0;
-; CHECK-NEXT:    cvta.local.u64 %SP, %SPL;
-; CHECK-NEXT:    ld.global.b64 %rd1, [ptr];
+; CHECK-NEXT:    add.u64 %rd1, %SPL, 0;
+; CHECK-NEXT:    cvta.local.u64 %rd2, %rd1;
+; CHECK-NEXT:    ld.global.b64 %rd3, [ptr];
 ; CHECK-NEXT:    { // callseq 0, 0
 ; CHECK-NEXT:    .param .align 1 .b8 param0[1];
 ; CHECK-NEXT:    .param .b64 param1;
 ; CHECK-NEXT:    .param .b32 retval0;
-; CHECK-NEXT:    add.u64 %rd2, %SP, 0;
 ; CHECK-NEXT:    st.param.b64 [param1], %rd2;
-; CHECK-NEXT:    add.u64 %rd3, %SPL, 1;
-; CHECK-NEXT:    ld.local.b8 %rs1, [%rd3];
+; CHECK-NEXT:    ld.b8 %rs1, [%SPL+1];
----------------
maleadt wrote:

Thanks; this actually looks like a miscompile (a generic `ld` of a local `%SPL`), and I could even get this to crash locally. I'm not too familiar with this code, so this is AI assisted, but it looks like inference itself is fine (since the stores became `st.local`), and this is rather `refinePtrAS` emitting a value-less `MachinePointerInfo(ADDRESS_SPACE_LOCAL)` during `byval` lowering, which then gets discarded by `SelectionDAG::getLoad`. I can fix this by having `refinePtrAS` return a pointer info that actually points to the original `alloca`.

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