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
_______________________________________________
cfe-commits mailing list
[email protected]
https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
