[gomp4 0/8] NVPTX: initial OpenMP offloading

[Alexander Monakov <amonakov at ispras dot ru>]

Hello,

This patch series implements some minimally required changes to have OpenMP
offloading working for NVPTX target on the gomp4 branch.  '#pragma omp target'
and data updates should work, but all parallel execution functionality remains
stubbed out (uses of '#pragma omp parallel' in target regions yield a link
error).

I'd like to get feedback on the patches, and approval for the gomp-4_0-branch
where possible.

Patches 1-2 unbreak compilation with offloading, patch 4 allows to invoke a
target region on the accelerator, patches 5-8 unbreak libgomp.h and allow
env.c to be compiled for the accelerator.

  nvptx: remove assumption of OpenACC attrs presence
  nvptx mkoffload: do not restrict to OpenACC
  libgomp: provide target-to-host fallback diagnostic
  libgomp: minimal OpenMP support in plugin-nvptx.c
  libgomp: provide sem.h, mutex.h, ptrlock.h on nvptx
  libgomp: provide stub bar.h on nvptx
  libgomp: work around missing pthread_attr_t on nvptx
  libgomp: provide ICVs via env.c on nvptx

 gcc/config/nvptx/mkoffload.c   |   7 +-
 gcc/config/nvptx/nvptx.c       |  19 ++--
 libgomp/config/nvptx/bar.h     |  38 +++++++
 libgomp/config/nvptx/env.c     | 219 +++++++++++++++++++++++++++++++++++++++++
 libgomp/config/nvptx/mutex.h   |  67 +++++++++++++
 libgomp/config/nvptx/ptrlock.h |  73 ++++++++++++++
 libgomp/config/nvptx/sem.h     |  65 ++++++++++++
 libgomp/libgomp.h              |   5 +
 libgomp/plugin/plugin-nvptx.c  |  30 +++++-
 libgomp/target.c               |   1 +
 10 files changed, 508 insertions(+), 16 deletions(-)
 create mode 100644 libgomp/config/nvptx/bar.h
 create mode 100644 libgomp/config/nvptx/mutex.h
 create mode 100644 libgomp/config/nvptx/ptrlock.h
 create mode 100644 libgomp/config/nvptx/sem.h