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

    <tr>
        <th>Summary</th>
        <td>
            [AMDGPU] Casting to addrspace 5 casing mem fault
        </td>
    </tr>

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

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

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

<pre>
    ```
__device__ void foo(int* xxx) {
 auto *p = (__attribute__((address_space(5)))int*)xxx;
    *p = 1;
}
__device__ void bar(int* xxx) {
    *xxx = 1;
}

__managed__ int y;
extern "C" __global__ void kernel() {
  int x;
  foo(&x);  // memfault
  //bar(&x); // works
  y = x;
}

int main() {
  kernel<<<1,1>>>();
 hipDeviceSynchronize();
  printf("%d\n", y);
}

```
the above code cause memfault. xxx is actually pointing to a private memory. it is OK to be accessed by a flat pointer but memfault when accessed by casting the flat pointer to private pointer.

compare ISA of foo and bar (https://godbolt.org/z/YqT4dhGcE ) shows casting xxx to a private pointer then accessing it generates buffer store which depends on s[0-3]. However, the kernel does not know that and does not set up s[0-3], which makes buffer store not working. flat store does not depend on s[0-3], therefore it still works.

This issue is revealed by blender crash due to enabling AAInstanceInfo by 7a68449

@arsenm @shiltian 


</pre>
<img width="1px" height="1px" alt="" src="http://email.email.llvm.org/o/eJx8VU-P47YP_TTKhZjAlu3EPviQSTb7G_xQtMBuDz0ZtEXH6siSK8n5s5--kJ1kkmm3gIAEzOPjI_mkoHPyoIlKlr2ybLfA0XfGlpezw17JcVEbcSnZKrqeaMeiTVUJOsqGqgqORgpojWE8l9ozvoHz-cx4AWz9OoMBR2-A8c0ALNkB43lVofdW1qOnqmI8ZzxHISw5V7kBG2I8zxgv5jOzMl4E3uTGCfDBGN_DbL37d4E12v8QOJOdz-ef0d1Ie9R4IFFVILWHyx1IZ09WA-N8yziHqjooU6O6VX8nq0lNjT6VDSSPPc1jZHwV9LHkNcjaM76HnvoWR-VvwDk8N_UAv6JPxr67G_Qy9XT-WU9BQo9S_1PcVXSynU_M-DZmyZfrmeAf0js57KaBf7voprNGyx_0GQODldq3U5gzngmWbfX0dQuXR-QniZ-85zsCrM2RoDGCoMHR0X1Ay7BckA6w8SMqdYHBSO2lPoA3gEHCEf2EN_ayBOkD-Nf_h19rAmwaco4E1BdAaBX6OZ8s1KO_V4FTR_oJ3KCbi3T0nObNveY1tHzsrTH9gJbg7dsGTBsMAKgnu4Z70nk_OJZs5r0ejKiN8ktjD4zvfzC-_-Ov76novjZfIOzOdebk7krCHJ5aviv60B5w0sOBNFn05KAe25YsOG8swamTTQeCBtLCgdHgWPYavSQs2y3hf-ZER7Jhd6Hn2SwgDDnQxsO7NifwHfqpnXvYkYdxeCAK-XOdHt8_CwgZwctSH5bzVOf4nW7W9iztqshSG6Ay5Eil5jvxNPrvnXQgnRspWMDSkVDNy6wVaUEWGouuAzFSGCRprFUY2Gbzpp1H3dCbbk3Ar3GVp2nx5Nk0QutI98DSyHVSeYkaHhELUSaiSApcUBmvebJK8zjii66Mk6igrIhXeV5HNWbJepWseNGKtKasSJuFLHnE0yiP0ziLch4vKc9x3azTlmO8zleCpRH1KNVSqWMf7LKY2izjKI2SbKGwJuWm555zTad5COEiZruFLUPSSz0eHEsjJZ13HzReejX9T2x-2X397XeW7WB7M76B8IRPzzdkwYUh2lMP041ZjFaVn-wsfTfWy8b0jO9DievHy2DNn9R4xveTMMf4_qr8WPK_AwAA__8NBQm_">