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

Jonas Hahnfeld via Phabricator via cfe-commits cfe-commits at lists.llvm.org
Fri Aug 17 05:59:24 PDT 2018


Hahnfeld added a comment.

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:

  #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?


Repository:
  rC Clang

https://reviews.llvm.org/D50845





More information about the cfe-commits mailing list