[clang] [llvm] [NVPTX][Draft] Make `__nvvm_nanosleep` a no-op if unsupported (PR #81033)

Artem Belevich via cfe-commits cfe-commits at lists.llvm.org
Wed Feb 7 14:12:55 PST 2024


Artem-B wrote:

> Okay, `__nvvm_reflect` doesn't work fully here because the `nanosleep` builtin I added requires `sm_70` at the clang level. Either means I'd need to go back to inline assembly or remove that requirement at least from clang so it's a backend failure.

The question is -- who's going to provide a fallback implementation for the nanosleepbuiltin for the older GPUs. I do not think it's LLVM's job, so constraining the builtin is appropriate. However, nothing stops you from providing your own implementation in libc using inline asm. Something along these lines:
```
__device__ void my_nanosleep(int N) {
  if (__nvvm_reflect(SM_70)) {
    asm volatile("nanosleep")
  } else {
     while(N--) {
        volatile asm("something unoptimizable")
     }
  }
}
```

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


More information about the cfe-commits mailing list