<table border="1" cellspacing="0" cellpadding="8">
    <tr>
        <th>Issue</th>
        <td>
            <a href=https://github.com/llvm/llvm-project/issues/89205>89205</a>
        </td>
    </tr>

    <tr>
        <th>Summary</th>
        <td>
            [llvm][NVPTX] `cvta.local.u64` used with generic address as input
        </td>
    </tr>

    <tr>
      <th>Labels</th>
      <td>
            new issue
      </td>
    </tr>

    <tr>
      <th>Assignees</th>
      <td>
      </td>
    </tr>

    <tr>
      <th>Reporter</th>
      <td>
          JOE1994
      </td>
    </tr>
</table>

<pre>
    PTX instruction [`cvta.local.u64`](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#data-movement-and-conversion-instructions-cvta) converts a 64-bit _**local address**_ to a 64-bit _**generic address**_ .

With the below input IR, `llc` emits PTX code that uses `cvta.local.u64` with an already _**generic address**_ as input.
This produces an invalid memory address, since [`cvta.local.u64` adds the base address of the "local address space window" to its input address](https://docs.nvidia.com/cuda/parallel-thread-execution/index.html#generic-addressing). A program using this invalid address can crash with **`CUDA_ERROR_ILLEGAL_ADDRESS`** runtime error.



### Minimal Reproducible Example
* Tested with llvm 19 trunk at `e90bc9cfd4d22c89dd993f62ede700ae25df49c5`

#### Input LLVM IR
```llvm
; mre.ll

target datalayout = "e-i64:64-i128:128-v16:16-v32:32-n16:32:64"
target triple = "nvptx64-nvidia-cuda"

declare hidden i32 @BAR(ptr %n)

define ptx_kernel void @FOO(ptr %aarg) noinline {
"BB0":
  %aloha = alloca i32, align 4
  store i32 25, ptr %aloha, align 4
  %r = ptrtoint ptr %aloha to i64
  %r2 = inttoptr i64 %r to ptr addrspace(5)
  %r3 = addrspacecast ptr addrspace(5) %r2 to ptr
  %r4 = call i32 @BAR(ptr %r3)
  store i32 %r4, ptr %aarg, align 4
  ret void
}
```

#### Output PTX code (from `llc -mcpu=sm_80 mre.ll`)

* Problematic code line is `cvta.local.u64  %rd5, %rd3;`
```ptx
.version 7.0
.target sm_80
.address_size 64

 // .globl FOO                     // -- Begin function FOO
.extern .func (.param .b32 func_retval0) BAR
(
  .param .b64 BAR_param_0
)
;
                                        // @FOO
.visible .entry FOO(
  .param .u64 FOO_param_0
)
{
 .local .align 4 .b8   __local_depot0[4];
  .reg .b64   %SP;
  .reg .b64 %SPL;
  .reg .b32   %r<4>;
  .reg .b64   %rd<6>;

// %bb.0: // %BB0
  mov.u64   %SPL, __local_depot0;
 cvta.local.u64  %SP, %SPL;  // %SP holds generic address 
  ld.param.u64 %rd1, [FOO_param_0];
  cvta.to.global.u64  %rd2, %rd1;
  add.u64 %rd3, %SP, 0;  // %rd3 holds generic address
  add.u64   %rd4, %SPL, 0;
  mov.b32   %r1, 25;
  st.local.u32  [%rd4], %r1;
 cvta.local.u64  %rd5, %rd3;  // PROBLEMATIC: %rd3 holds a generic address, but used as input to `cvta.local.u64`
  { // callseq 0, 0
  .param .b64 param0;
  st.param.b64  [param0+0], %rd5;
  .param .b32 retval0;
 call.uni (retval0),
  BAR,
  (
  param0
  );
  ld.param.b32 %r2, [retval0+0];
  } // callseq 0
  st.global.u32   [%rd2], %r2;
 ret;
                                        // -- End function
}
```
</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJy0V11z6roV_TXiZY89Rv4AP_CAIXROJ6dkctLb-8bIlgD1yhKVZJL013ck2WAIaTudaSaTgLQ_1l57aVsmxvCDZGyB8grl6wnp7FHpxZ-3T9OyzCa1op-Ll7ffgUtjdddYriQ40yJpzpbEQjVExF2RoSJB-Rrh-dHak0HpEuENwhuqGhPLM6ecxI1qEd40HSUIb05EEyGYiOxRM0Ij9sGazkVHeMMlZR_x0bYC4ZQSS6JWnVnLpI2IpFGj5Jlpw5WMRqhM5AAhXELYtwYIFFlUcws7hJcILz1YIJRqZkxY2oFVX-wOTDLNm3vLGCVrlCzD379xewR7ZFAzod6By1Nn4ccrwitARSJEg4oEWMutAUdfoygDeyQWOsMMPKIP3l1IIoEIx8jnf0BDTEjao3o7cgMnrWjXMOOicHkmglNoWav059V7BYbLhn3XRGdoQmHEsMEN1N6vIYxvSARzIg2Ddy6pekcYOzZdyYGOIef_RRc9LVGfhMsDwmUMS0fCQZMWOrcG1vEycDGgboiERhNzDJwHSlGRrP66Xu6eXl-3r7sfz89Pf1o-75br9evTr19O3t4KdCctbxkwrZW-kcTNX5yGX_jJJW-JgFcWusNrweDpg7QnwQbbJbwxYxkNcIQ4tzAtwepO_gHEOrWwMqmbstnTjGLczEtKyzLdF5hRNksSwnBO91nZ5A7nIxgOyQ_flOfn3346oYb9Igm_Lme_lFbQahYLMQ5kiT4wC-4wCvKpOgsoXTs5sIgXGUqXRRbxKZ6jdDnF8-g8LdynIjqnGKXLFEfSr_hvRYYwvglrNT8JNoSU55P9KLIoyCMKwsBjNJQ1gmgGR04pk8BTDChLquUrwvOT1YBwLhEub132XDI42Y_dH0xLJuCsOHVum-326kaIdjoCqbgUzgHNqoFKXFWJA5L2IcE7CHUkHjkR7mg4MO6QEcEPErLB0lilmQeKc7c9pHPeD8wRzrUPerLaKi7tjYM_ZcWNMfbWXFqrnCUvshDDKu_phO-PKsLz_MJMcE0D-sGiIcY-9OnzhIjjAJkP0BAhHnZCp6OEVx6855gKz_wXJjSzvlN9E2brO-F-q_ZtZ53cL8MX4fleq7afzhC1zalD6dq0u3kyCN4d8vI24hJetKoFa4nlTYjkdcEfDPFAB_UN9p9SlFZXjAPkk_0IK3H_GINZ3NvE_YHwqPqlfmjtDP8ng6HrPT1hmEJ8EKoWsNlu4dFPbxVFULEDl7DvZHiQO-mHJOzDMi0hdluOqtjN4RbiOsXefKeZPRPh-AHX3Z6d-dCmi32Ruf2d_7obKr-QmlaDw3_30yPvT2lPGjd-hsZMWv0J4fje43Dd2Gy33-EYDjWE5kHcqw7ieg4Au51f3lF2UjZBeZW5Z9gVfKzZIdTqW_7r5eGe33n-upXi4KZRuspQ-vR9YE1RuirGJkMhgRac13WcoPQiBIRzN6T6aK06B1kOUPDqvrRL6gdK_vXSCzlUAaMkv17gqAQ1cHc5gSG1oKETPpivZOqD5dW4KzekegRWeS3fnCd8OU_TkTmh9Bo8vSB1H5I7tJqmj-Hex-oTZqO6-3BjSq8N9DXhfLRv7ECiM3J3rBDQXYFWvc-_4_x-elzKeHndVs9PP5dvP1ah36OqyJe68Arqzt816eWi6Cb3w2v7MMpn1ZDMzXLD_gFJKP_BGfcfk9vCQ8O9eFFe9Ra4SkbF0zFX4xkzjJcrN0SIuJPcDaPr7EF4NXj7Z8zq-hy6jIA-8WWjHKW8yLLun0C4V-UlRXWnSjRbf2XlWvMg1iCJvt14VDG-xtLM_q8DMIrgSdLL4P7uYTihi5SWaUkmbDGdTdMsLbJ8NjkuaF3nWYKzfUnnWTPf7xnBbMboLK9JQWd0whc4wVmSTefTJM2mszjDmBRZUU6nGSlLkqMsYS3hInY3xVjpw4Qb07HFvMRJPhGkZsL4d0iMJXsHv-kuS_l6ohfOJ6q7g0FZIrix5hrFciv8y6e_geZrlFd_-e3l7XeUrx-_JXlN-3vy_ewZhD7ptFjcvm4cuD12df-m4TOFf9FJq7-zxrqXC4fYILzxFf0rAAD__-pnRHI">