[libc-commits] [libc] [llvm] [libc] Support AMDGPU device interrupts for the RPC interface (PR #188067)

[Joel E. Denny via libc-commits]

jdenny-ornl wrote:

Starting at this commit (4961700c1004) and persisting through today's bf2a97a0dd27, the following program hangs for me when offloading to an MI100:

```
$ cat test.c
#include <stdio.h>
int main() {
  int x = 0;
  #pragma omp target teams num_teams(2) reduction(+:x)
  x += 1;
  printf("Hello World: %d\n", x);
  return 0;
}
$ clang -fopenmp --offload-arch=native test.c
$ ./a.out # never terminates
```

If I remove the reduction clause or compile with >= -O1, it behaves correctly.  Also, check-all sees many offloading tests hang.  I configured using `offload/cmake/caches/Offload.cmake`.

When targeting an MI60, the above program and check-all appear to additionally crash the driver for everyone, and I have to ask our sys admins to restart.

However, when targeting an MI300A, I do not see these problems.

I understand AMD is phasing out support for some older GPUs.  Should we assume that upstream LLVM's offload support is no longer expected to be stable for MI60 and MI100?

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