jplehr wrote:
I did a bit of an investigation with the help of an AI tool.
The culprit appears to be reproducible with this piece of code
```
#include <hip/hip_runtime.h>
#include <initializer_list>
#include <cstdio>
__device__ int f(std::initializer_list<int> il) { return *il.begin(); }
__global__ void kernel(int* out) {
out[0] = f({42}); // BUG HERE
}
int main() {
int *d_out, h_out = 0;
if (hipMalloc(&d_out, sizeof(int)) != hipSuccess) return 1;
hipLaunchKernelGGL(kernel, 1, 1, 0, 0, d_out);
hipError_t err = hipDeviceSynchronize();
if (err != hipSuccess) {
printf("FAILED: %s\n", hipGetErrorString(err));
hipFree(d_out);
return 1;
}
hipMemcpy(&h_out, d_out, sizeof(int), hipMemcpyDeviceToHost);
hipFree(d_out);
printf("SUCCESS: result = %d (expected 42)\n", h_out);
return h_out == 42 ? 0 : 1;
}
```
when compiled as HIP and in my test case with offload target gfx90a.
```
/opt/botworker/llvm/llvm/build/bin/clang++ \
-x hip test-bug.cpp -o test-bad \
--offload-arch=gfx90a \
--rocm-path=/opt/rocm-7.1.1 \
-I/opt/rocm-7.1.1/include \
-L/opt/rocm-7.1.1/lib \
-lamdhip64
```
It produces a memory access fault as well and the IR shows the same changes
w.r.t. addrspacecast.
Let me know if I can help more.
https://github.com/llvm/llvm-project/pull/197745
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
