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