Hahnfeld added a comment.

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.

> 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.


Repository:
  rC Clang

https://reviews.llvm.org/D50845



_______________________________________________
cfe-commits mailing list
cfe-commits@lists.llvm.org
http://lists.llvm.org/cgi-bin/mailman/listinfo/cfe-commits

Reply via email to