[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#1203973, @Hahnfeld wrote:

> In https://reviews.llvm.org/D50845#1203967, @ABataev wrote:
>
> > 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.
>
>
> How common is that for functions that are used in `target` regions? In the worst case we can make my fourth point a warning and lose Sema checking for those functions.


It does not matter how common is it or not. If the bad situation can happen, it will happen.
Warning won't work here, because, again, you may end up with the code that may cause compiler crash for the device. For example, if the system function uses throw/catch stmts, we may emit the warning for this function, but will have troubles during the codegen.

> 
> 
>> 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.
> 
> Once we know which functions are used, they can be checked as usual.
> 
>> The best solution is to use only device specific header files. Device compilation phase should use system header files for the host at all.
> 
> You mean "shouldn't use system header files for the host"? I think that may be hard to achieve, especially if we want to Sema check all of the source code during device compilation.

Yes, I mean should not. Yes, this is hard to achieve but that's the only complete and correct solution. Everything else looks like a non-stable hack.


Repository:
  rC Clang

https://reviews.llvm.org/D50845