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 cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits