[all-commits] [llvm/llvm-project] f5e499: Revert "[HIP] use offload wrapper for non-device-o...

[Joseph Huber via All-commits]

  Branch: refs/heads/main
  Home:   https://github.com/llvm/llvm-project
  Commit: f5e499a3383c1e3b9f60e60151075e8d9c1c3166
      https://github.com/llvm/llvm-project/commit/f5e499a3383c1e3b9f60e60151075e8d9c1c3166
  Author: Joseph Huber <huberjn at outlook.com>
  Date:   2025-06-09 (Mon, 09 Jun 2025)

  Changed paths:
    M clang/lib/CodeGen/CGCUDANV.cpp
    M clang/lib/Driver/Driver.cpp
    M clang/lib/Driver/ToolChains/Clang.cpp
    M clang/test/Driver/hip-binding.hip
    M clang/test/Driver/hip-phases.hip
    M clang/test/Driver/hip-toolchain-no-rdc.hip

  Log Message:
  -----------
  Revert "[HIP] use offload wrapper for non-device-only non-rdc (#132869)" (#143432)

This breaks a lot of new driver HIP compilation. We should probably
revert this for now until we can make a fixed version.

```c++

static __global__ void print() { printf("%s\n", "foo"); }

void b();

int main() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
  b();
}
```
```c++

static __global__ void print() { printf("%s\n", "bar"); }

void b() {
  hipLaunchKernelGGL(print, dim3(1), dim3(1), 0, 0);
  auto y = hipDeviceSynchronize();
}
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver
$ ./a.out
foo
foo
```
```console
$ clang++ a.hip b.hip --offload-arch=gfx1030 --offload-new-driver -flto
<crash>
```

This reverts commit d54c28b9c1396fa92d9347ac1135da7907121cb8.



To unsubscribe from these emails, change your notification settings at https://github.com/llvm/llvm-project/settings/notifications