D38798 does not work correctly and there is no stable Sema-based solution to do this absolutely correctly. We need to adapt the codegen so that it emits metadata for the implicit declare target functions.

In D50845#1204216, @Hahnfeld wrote:

In D50845#1204210, @ABataev wrote:

Right, warning wasn't a good thought. We really want strict checking and would have to error out when we find a function that wasn't implicitly declare target on the host.
I meant to ask how common that would be? If that's only some known functions we could handle them separately.

Again, it does not matter how common is this situation. We cannot rely on the probability here, we need to make the compiler to work correctly in all possible situation, no matter how often they can occur.

Got that, I agree on the conservative approach: If we find a function to be called that wasn't checked (because it wasn't implicitly declare target on the host) the compiler can error out. That should be correct in all cases, shouldn't it?

There's a trade-off here:

• How many TUs pass full analysis and how many don't? (today's situation; we know that some headers don't work)
• How many TUs pass when we only check called functions (and error if we call non-checked ones) and how many regress compared to today's situation? If the number of regressions is zero for all practical situations but we can compile some important cases, that should be a win.

Right, warning wasn't a good thought. We really want strict checking and would have to error out when we find a function that wasn't implicitly declare target on the host.
I meant to ask how common that would be? If that's only some known functions we could handle them separately.

In D50845#1203973, @Hahnfeld wrote:

In D50845#1203967, @ABataev wrote:

In D50845#1203961, @Hahnfeld wrote:

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

In D50845#1203961, @Hahnfeld wrote:

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

If I understand it correctly, the root cause of this exercise is that we want to compile for GPU using plain C. CUDA avoids this issue by separating device and host code via target attributes and clang has few special cases to ignore inline assembly errors in the host code if we're compiling for device. For OpenMP there's no such separation, not in the system headers, at least.

Yes, that's one of the nice properties of CUDA (for the compiler). There used to be the same restriction for OpenMP where all functions used in target regions needed to be put in declare target. However that was relaxed in favor of implicitly marking all called functions in that TU to be declare target.
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.

In D50845#1202550, @Hahnfeld wrote:

In D50845#1202540, @ABataev wrote:

Maybe for device compilation we also should define __NO_MATH_INLINES and __NO_STRING_INLINES macros to disable inline assembly in glibc?

The problem is that __NO_MATH_INLINES doesn't even avoid all inline assembly from bits/mathinline.h :-( incidentally Clang already defines __NO_MATH_INLINES for x86 (due to an old bug which has been fixed long ago) - and on CentOS we still have problems as described in PR38464.

As a second thought: This might be valid for NVPTX, but I don't think it's a good idea for x86-like offloading targets - they might well profit from inline assembly code.

Maybe for device compilation we also should define __NO_MATH_INLINES and __NO_STRING_INLINES macros to disable inline assembly in glibc?

In D50840#1202486, @RKSimon wrote:

It could be the kind of thing we should do in slp @ABataev what do you think?

Yes, looks like the opportunity for the SLP Vectorizer.

In D46021#1200884, @tra wrote:

Is there a way to control this from clang command line?

Added an extra test.

It really requires a test.

LG

Reworked the patch to change the way we get file id to make it more robust.

Reworked according to Eric's comments

In D50158#1185919, @Hahnfeld wrote:

In D50158#1185906, @RaviNarayanaswamy wrote:

We should not be having calls to _-kmpc_omp_taskwait in libomptarget.
The wait should be done in the user code and then libomptarget routine should be called.

There is no way to implement asynchronous offloading without having knowledge of the dependencies. Intel failed to keep the separation strict for liboffload, so I don't think it helps to complain about the situation.

Reworked according to Eric's comments.

Updated and reworked to latest changes

Address David's comments

Reworked to control output by DICompileUnit emission kind.

nvcc always emits debug location before the very first label that represents function start. So, we can consider this as the requirement for NVPTX.

In D45784#1181208, @echristo wrote:

But you have other locations inside of functions so I don't get it.

eric

In D45784#1181119, @echristo wrote:

Can you explain this a bit more with some examples of wrong relocations?

Thanks!

-eric

Eric accepted the patch offline.

Updated to latest version

Added maximal length of the merged directives.

Updated to the latest revision.

Address ERic's comments.

Yes, it was committed

Fixed SPMD mode barrier for initialized runtime.

Added checks for all possible debug options.

In D44186#1168001, @grokos wrote:

@ABataev: Can you put a link to the clang-side patch in the description so that we link the two patches together? Also, please let me know when you commit the clang patch so that I commit this one as well.

1. The new tool definitely requires an RFC at first.
2. The patch is too big and should be split into several small patches