[AMDGPU] Casting to addrspace 5 causing mem fault

[yxsamliu]

__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 when compiled with -fno-inline-functions. 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. AAInstanceInfo will deduce that xxx is private pointer by IPA and cast flat pointer to addr space 5.

@arsenm @shiltian

[llvmbot]

@llvm/issue-subscribers-backend-amdgpu

Author: Yaxun (Sam) Liu (yxsamliu)

``` __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 


</details>

[arsenm]

    buffer_store_dword v1, v0, s[0:3], 0 offen

Somehow this isn't using the correct SP base register, and is instead acting like it's in a kernel with a known 0 base

[yxsamliu]

using private pointer as function parameter has the same issue https://godbolt.org/z/a693E3fca

I am wondering how OpenCL manages that. Or maybe private pointer as function parameter never worked, and we just did not know that since most cases functions are inlined.

[arsenm]

I think I found the issue. foldFrameIndex wants to use 0 as the soffset in the case where the incoming pointer value is a FrameIndex to directly fold into the memory operation. That is wrong in the case where the pointer value is not a frame index, and isn't folding into the instruction

Nevermind, I forgot how this was supposed to work. 0 is the correct base when interpreting an absolute stack address

[arsenm]

This is the nocapture promote alloca bug again. The alloca in the kernel is getting promoted to LDS, and thinks the callee not capturing means it's safe to do so

[arsenm]

I'm wondering if we should just drop addrspacecast handling from PromoteAlloca. In the cases where we can legally, and usefully, promote it, it should have been optimized out already

[shiltian]

This is the nocapture promote alloca bug again.

Probably not. AMDGPUAttributorPass runs after PromoteAlloca.

[arsenm]

This is the nocapture promote alloca bug again.

Probably not. AMDGPUAttributorPass runs after PromoteAlloca.

It is, AMDGPUAttributor isn't involved