jhuber6 added a comment. In D120353#3340533 <https://reviews.llvm.org/D120353#3340533>, @ABataev wrote:
> In D120353#3338770 <https://reviews.llvm.org/D120353#3338770>, @jhuber6 wrote: > >> In D120353#3338718 <https://reviews.llvm.org/D120353#3338718>, @ABataev >> wrote: >> >>> In D120353#3338647 <https://reviews.llvm.org/D120353#3338647>, @jhuber6 >>> wrote: >>> >>>> But the main reason I made this patch is for interoperability. Without >>>> this if you want to call a CUDA function from the OpenMP device you'd need >>>> a variant and a dummy implementation. If you don't write a dummy >>>> implementation you'll get a linker error, if you don't use a variant >>>> you'll override the CUDA function. >>> >>> Ah, ok, I see. How is supposed to be used? In Cuda code or in plain C/C++ >>> code? >> >> I haven't finalized the implementation, but the basic support I've tested >> was calling a `__device__` function compiled from another file with OpenMP, >> with this patch the source files would look like this for example. I think >> the inverse would also be possible given some code on the CUDA side. Calling >> CUDA kernels would take some extra work. >> >> __device__ int cuda() { return 0; } >> >> >> >> int cuda(void); >> #pragma omp declare target device_type(nohost) to(cuda) >> >> int main() { >> int x = 1; >> #pragma omp target map(from : x) >> x = cuda(); >> >> return x; >> } > > What if we have `#pragma omp target if (...)` or `#pragma omp target > device(...)`? If we have `#pragma omp target if (...)` then that requires a host fallback and violates the assertion the user passed in, it will hit the unreachable and fail. If the user passed in `#pragma omp target device(...)` we will assume that a host implementation exists as well. Repository: rG LLVM Github Monorepo CHANGES SINCE LAST ACTION https://reviews.llvm.org/D120353/new/ https://reviews.llvm.org/D120353 _______________________________________________ cfe-commits mailing list cfe-commits@lists.llvm.org https://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits
