[PATCH] D50845: [CUDA/OpenMP] Define only some host macros during device compilation

[Alexey Bataev via Phabricator via cfe-commits]

ABataev added a comment.

In https://reviews.llvm.org/D50845#1203961, @Hahnfeld wrote:

> In https://reviews.llvm.org/D50845#1202973, @ABataev wrote:
>
> > > So ideally I think Clang should determine which functions are really `declare target` (either explicit or implicit) and only run semantical analysis on them. If a function is then found to be "broken" it's perfectly desirable to error back to the user.
> >
> > It is not possible for OpenMP because we support implicit declare target functions. Clang cannot identify whether the function is going to be used on the device or not during sema analysis.
>
>
> You are right, we can't do this during device compilation because we don't have an AST before Sema.
>
> However I'm currently thinking about the following:
>
> 1. Identify explicit and implicit `declare target` functions during host Sema and CodeGen.
> 2. Attach meta-data for all of them to LLVM IR `.bc` which is passed via  `-fopenmp-host-ir-file-path`. I think we already do something similar for outlined `target` regions?
> 3. During device Sema query that meta-data so Clang knows when a function will be called from within a `target` region. Skip analysis of functions that are not needed for the device, just as CUDA does.
> 4. Check that we don't need functions that weren't marked in 2. That's to catch users doing something like: ```lang=c #pragma omp target { #ifdef __NVPTX__ target_func() #endif } ```
>
>   For now that's just an idea, I didn't start implementing any of this yet. Do you think that could work?


I thought about this approach already. But it won't work in general. The main problem here is that host and device compilation phases may end up with the different set of implicit declare target functions. The main problem here not the user code, but the system libraries, which may use the different set of functions.
Another one problem here is that the user may use the function that has some host assembler inside. In this case we still need to emit error message, otherwise, we may end up with the compiler crash.

The best solution is to use only device specific header files. Device compilation phase should use system header files for the host at all.


Repository:
  rC Clang

https://reviews.llvm.org/D50845