'new header file target_impl.h' should be 'new header file target_interface.h' in the commit message

I really like about half of this change.

Still LGTM. A little surprised these are being run on windows, but good to be robust.

• address review comments

The information I've got on the possible race is:
When this patch is applied (by git's automerge, I think) to the rocm stack, a test located at:
https://github.com/ROCm-Developer-Tools/aomp/blob/master/test/smoke/devices/devices.c
fails in unpredictable fashion.

Most (possibly all except the amdgpu cmake) of this change is incorporated in D94745 or already committed to trunk. I'm going to leave this up as a reference for now but don't intend to land it.

I'd like to call the clang/llvm intrinsics instead, plus replace unsigned long long with uint64_t. Easier to check the correctness of that after the call sites are limited to one file.

• rebase

LGTM, thanks!

• fix copy/paste error, support higher arity functions
• rebase, tidy

• fix copy/paste error, support higher arity functions

In D95155#2513810, @jdoerfert wrote:

Interesting. Should we do this instead of the pure macro stuff? Is this going to break if the argument type doesn't match the declaration parameter type, e.g.,

void f(int);

f(0.0);

This breaks the build for me. Compiling openmp via ENABLE_RUNTIMES without any extra options, the include path isn't set up so kmp_runtime.cpp fails to find the header.

Thanks. If you land this I'll rebase D95093 manually

Yeah... That's not gone well. Unless @jdoerfert can point us at a better way to spell these in openmp, let's add 'fix target atomics' to the todo pile and use clang intrinsics for the time being.

Most of the intrinsic functions are defined on a single type, e.g. if we wanted to use the nvvm encoded one for atomic_inc.

#pragma omp declare target

Possibly obsolete?

In D95062#2510535, @jdoerfert wrote:

LG, can you replace it with clang or OpenMP atomics next?

• clang-format

• insert whitespace requested at review

In D94745#2508792, @jdoerfert wrote:

Let's split this. Declare target in all the source files can go first, doesn't hurt anyone.

D95048

Nice. Particularly like getting rid of these because they use field access syntax to make a function call, which clang does with a Microsoft extension.

@jdoerfert this may be a candidate for deriving names of extracted target regions

I don't think introducing everything from the cuda namespace into openmp nvptx offloading is a feature. Inevitably people will call threadIdx.x instead of the openmp or clang equivalent, and this will mask missing functionality in openmp.

Nice! Thank you

• address some review comments
• Run tests under new pass manager

This patch was written, roughly, by:

• copying the known-working openmp driver from rocm into the trunk source tree
• deleting lots of stuff that didn't look necessary
• deleting some stuff that is broadly necessary, but the specifics are up for debate

Fixed the easy parts. Adding a call to the new pass manager by RUN: opt -S -mtriple=amdgcn-- -passes=amdgpu-lower-function-lds < %s | FileCheck %s does't work - it looks like the pass manager runs the module pass on individual functions. Trying to find some documentation about the new pass manager to see where the plumbing error lies.

• address some review comments

@ronlieb tells me an out of tree offloading test (aomp/test/smoke/devices) started crashing (hangs/segv/fp exception) with this patch applied. That doesn't make sense to me since this doesn't appear to change the target offloading logic, but it might be a smoking gun for a lifetime management error somewhere in the above. Does anyone know if the host openmp runtime is expected to be clean under things like valgrind or thread sanitizer?

Great review guys, thanks! With the exception of avoiding inline asm (which I'd like to, but am not sure what to do about) I'll fix the rest when I'm back in the office.

Remaining calls could be replaced with builtins by duplicating parts of
the cuda wrapper infra. This would mean choosing which to call based
on architecture number (>= sm_70) instead of CUDA_VERSION.

• reduce scope

That's interesting. Move some of the current target_impl.h into clang shipped headers.

Note that this is incremental (on the basis that it's already hard enough to review). Cuda.h is still used for CUDA_VERSION here, but also for the atomic functions and a few libc prototypes. I have the library compiling without it locally.