aboutsummaryrefslogtreecommitdiff
path: root/libgomp/config
diff options
context:
space:
mode:
authorAlexander Monakov <amonakov@ispras.ru>2016-11-23 21:36:41 +0300
committerAlexander Monakov <amonakov@gcc.gnu.org>2016-11-23 21:36:41 +0300
commit6103184e81c0b6a8b1f4e072e0c32d9bb86fcc15 (patch)
treeafbc9d83e9693712185a8afd9d6593f23eef2734 /libgomp/config
parent6251fe936f3eab1542c1581646254e4bae0f6688 (diff)
downloadgcc-6103184e81c0b6a8b1f4e072e0c32d9bb86fcc15.zip
gcc-6103184e81c0b6a8b1f4e072e0c32d9bb86fcc15.tar.gz
gcc-6103184e81c0b6a8b1f4e072e0c32d9bb86fcc15.tar.bz2
OpenMP offloading to NVPTX: libgomp changes
* Makefile.am (libgomp_la_SOURCES): Add atomic.c, icv.c, icv-device.c. * Makefile.in. Regenerate. * configure.ac [nvptx*-*-*] (libgomp_use_pthreads): Set and use it... (LIBGOMP_USE_PTHREADS): ...here; new define. * configure: Regenerate. * config.h.in: Likewise. * config/posix/affinity.c: Move to... * affinity.c: ...here (new file). Guard use of Pthreads-specific interface by LIBGOMP_USE_PTHREADS. * critical.c: Split out GOMP_atomic_{start,end} into... * atomic.c: ...here (new file). * env.c: Split out ICV definitions into... * icv.c: ...here (new file) and... * icv-device.c: ...here. New file. * config/linux/lock.c (gomp_init_lock_30): Move to generic lock.c. (gomp_destroy_lock_30): Ditto. (gomp_set_lock_30): Ditto. (gomp_unset_lock_30): Ditto. (gomp_test_lock_30): Ditto. (gomp_init_nest_lock_30): Ditto. (gomp_destroy_nest_lock_30): Ditto. (gomp_set_nest_lock_30): Ditto. (gomp_unset_nest_lock_30): Ditto. (gomp_test_nest_lock_30): Ditto. * lock.c: New. * config/nvptx/lock.c: New. * config/nvptx/bar.c: New. * config/nvptx/bar.h: New. * config/nvptx/doacross.h: New. * config/nvptx/error.c: New. * config/nvptx/icv-device.c: New. * config/nvptx/mutex.h: New. * config/nvptx/pool.h: New. * config/nvptx/proc.c: New. * config/nvptx/ptrlock.h: New. * config/nvptx/sem.h: New. * config/nvptx/simple-bar.h: New. * config/nvptx/target.c: New. * config/nvptx/task.c: New. * config/nvptx/team.c: New. * config/nvptx/time.c: New. * config/posix/simple-bar.h: New. * libgomp.h: Guard pthread.h inclusion. Include simple-bar.h. (gomp_num_teams_var): Declare. (struct gomp_thread_pool): Change threads_dock member to gomp_simple_barrier_t. [__nvptx__] (gomp_thread): New implementation. (gomp_thread_attr): Guard by LIBGOMP_USE_PTHREADS. (gomp_thread_destructor): Ditto. (gomp_init_thread_affinity): Ditto. * team.c: Guard uses of Pthreads-specific interfaces by LIBGOMP_USE_PTHREADS. Adjust all uses of threads_dock. (gomp_free_thread) [__nvptx__]: Do not call 'free'. * config/nvptx/alloc.c: Delete. * config/nvptx/barrier.c: Ditto. * config/nvptx/fortran.c: Ditto. * config/nvptx/iter.c: Ditto. * config/nvptx/iter_ull.c: Ditto. * config/nvptx/loop.c: Ditto. * config/nvptx/loop_ull.c: Ditto. * config/nvptx/ordered.c: Ditto. * config/nvptx/parallel.c: Ditto. * config/nvptx/priority_queue.c: Ditto. * config/nvptx/sections.c: Ditto. * config/nvptx/single.c: Ditto. * config/nvptx/splay-tree.c: Ditto. * config/nvptx/work.c: Ditto. * testsuite/libgomp.fortran/fortran.exp (lang_link_flags): Pass -foffload=-lgfortran in addition to -lgfortran. * testsuite/libgomp.oacc-fortran/fortran.exp (lang_link_flags): Ditto. * plugin/plugin-nvptx.c: Include <limits.h>. (struct targ_fn_descriptor): Add new fields. (struct ptx_device): Ditto. Set them... (nvptx_open_device): ...here. (nvptx_adjust_launch_bounds): New. (nvptx_host2dev): Allow NULL 'nvthd'. (nvptx_dev2host): Ditto. (GOMP_OFFLOAD_get_caps): Add GOMP_OFFLOAD_CAP_OPENMP_400. (link_ptx): Adjust log sizes. (nvptx_host2dev): Allow NULL 'nvthd'. (nvptx_dev2host): Ditto. (nvptx_set_clocktick): New. Use it... (GOMP_OFFLOAD_load_image): ...here. Set new targ_fn_descriptor fields. (GOMP_OFFLOAD_dev2dev): New. (nvptx_adjust_launch_bounds): New. (nvptx_stacks_size): New. (nvptx_stacks_alloc): New. (nvptx_stacks_free): New. (GOMP_OFFLOAD_run): New. (GOMP_OFFLOAD_async_run): New (stub). Co-Authored-By: Dmitry Melnik <dm@ispras.ru> Co-Authored-By: Jakub Jelinek <jakub@redhat.com> From-SVN: r242789
Diffstat (limited to 'libgomp/config')
-rw-r--r--libgomp/config/linux/lock.c94
-rw-r--r--libgomp/config/nvptx/alloc.c0
-rw-r--r--libgomp/config/nvptx/bar.c206
-rw-r--r--libgomp/config/nvptx/bar.h166
-rw-r--r--libgomp/config/nvptx/barrier.c0
-rw-r--r--libgomp/config/nvptx/doacross.h60
-rw-r--r--libgomp/config/nvptx/error.c42
-rw-r--r--libgomp/config/nvptx/icv-device.c74
-rw-r--r--libgomp/config/nvptx/iter.c0
-rw-r--r--libgomp/config/nvptx/iter_ull.c0
-rw-r--r--libgomp/config/nvptx/lock.c41
-rw-r--r--libgomp/config/nvptx/loop.c0
-rw-r--r--libgomp/config/nvptx/loop_ull.c0
-rw-r--r--libgomp/config/nvptx/mutex.h60
-rw-r--r--libgomp/config/nvptx/ordered.c0
-rw-r--r--libgomp/config/nvptx/parallel.c0
-rw-r--r--libgomp/config/nvptx/pool.h (renamed from libgomp/config/nvptx/fortran.c)39
-rw-r--r--libgomp/config/nvptx/priority_queue.c1
-rw-r--r--libgomp/config/nvptx/proc.c41
-rw-r--r--libgomp/config/nvptx/ptrlock.h73
-rw-r--r--libgomp/config/nvptx/sections.c0
-rw-r--r--libgomp/config/nvptx/sem.h65
-rw-r--r--libgomp/config/nvptx/simple-bar.h70
-rw-r--r--libgomp/config/nvptx/single.c0
-rw-r--r--libgomp/config/nvptx/splay-tree.c0
-rw-r--r--libgomp/config/nvptx/target.c49
-rw-r--r--libgomp/config/nvptx/task.c43
-rw-r--r--libgomp/config/nvptx/team.c178
-rw-r--r--libgomp/config/nvptx/time.c49
-rw-r--r--libgomp/config/nvptx/work.c0
-rw-r--r--libgomp/config/posix/affinity.c140
-rw-r--r--libgomp/config/posix/simple-bar.h69
32 files changed, 1312 insertions, 248 deletions
diff --git a/libgomp/config/linux/lock.c b/libgomp/config/linux/lock.c
index a671b7e..07845c9 100644
--- a/libgomp/config/linux/lock.c
+++ b/libgomp/config/linux/lock.c
@@ -32,98 +32,8 @@
#include <sys/syscall.h>
#include "wait.h"
-
-/* The internal gomp_mutex_t and the external non-recursive omp_lock_t
- have the same form. Re-use it. */
-
-void
-gomp_init_lock_30 (omp_lock_t *lock)
-{
- gomp_mutex_init (lock);
-}
-
-void
-gomp_destroy_lock_30 (omp_lock_t *lock)
-{
- gomp_mutex_destroy (lock);
-}
-
-void
-gomp_set_lock_30 (omp_lock_t *lock)
-{
- gomp_mutex_lock (lock);
-}
-
-void
-gomp_unset_lock_30 (omp_lock_t *lock)
-{
- gomp_mutex_unlock (lock);
-}
-
-int
-gomp_test_lock_30 (omp_lock_t *lock)
-{
- int oldval = 0;
-
- return __atomic_compare_exchange_n (lock, &oldval, 1, false,
- MEMMODEL_ACQUIRE, MEMMODEL_RELAXED);
-}
-
-void
-gomp_init_nest_lock_30 (omp_nest_lock_t *lock)
-{
- memset (lock, '\0', sizeof (*lock));
-}
-
-void
-gomp_destroy_nest_lock_30 (omp_nest_lock_t *lock)
-{
-}
-
-void
-gomp_set_nest_lock_30 (omp_nest_lock_t *lock)
-{
- void *me = gomp_icv (true);
-
- if (lock->owner != me)
- {
- gomp_mutex_lock (&lock->lock);
- lock->owner = me;
- }
-
- lock->count++;
-}
-
-void
-gomp_unset_nest_lock_30 (omp_nest_lock_t *lock)
-{
- if (--lock->count == 0)
- {
- lock->owner = NULL;
- gomp_mutex_unlock (&lock->lock);
- }
-}
-
-int
-gomp_test_nest_lock_30 (omp_nest_lock_t *lock)
-{
- void *me = gomp_icv (true);
- int oldval;
-
- if (lock->owner == me)
- return ++lock->count;
-
- oldval = 0;
- if (__atomic_compare_exchange_n (&lock->lock, &oldval, 1, false,
- MEMMODEL_ACQUIRE, MEMMODEL_RELAXED))
- {
- lock->owner = me;
- lock->count = 1;
- return 1;
- }
-
- return 0;
-}
+/* Reuse the generic implementation in terms of gomp_mutex_t. */
+#include "../../lock.c"
#ifdef LIBGOMP_GNU_SYMBOL_VERSIONING
/* gomp_mutex_* can be safely locked in one thread and
diff --git a/libgomp/config/nvptx/alloc.c b/libgomp/config/nvptx/alloc.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/alloc.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/bar.c b/libgomp/config/nvptx/bar.c
index e69de29..820affb 100644
--- a/libgomp/config/nvptx/bar.c
+++ b/libgomp/config/nvptx/bar.c
@@ -0,0 +1,206 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is an NVPTX specific implementation of a barrier synchronization
+ mechanism for libgomp. This type is private to the library. This
+ implementation uses atomic instructions and bar.sync instruction. */
+
+#include <limits.h>
+#include "libgomp.h"
+
+
+void
+gomp_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+ if (__builtin_expect (state & BAR_WAS_LAST, 0))
+ {
+ /* Next time we'll be awaiting TOTAL threads again. */
+ bar->awaited = bar->total;
+ __atomic_store_n (&bar->generation, bar->generation + BAR_INCR,
+ MEMMODEL_RELEASE);
+ }
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+}
+
+void
+gomp_barrier_wait (gomp_barrier_t *bar)
+{
+ gomp_barrier_wait_end (bar, gomp_barrier_wait_start (bar));
+}
+
+/* Like gomp_barrier_wait, except that if the encountering thread
+ is not the last one to hit the barrier, it returns immediately.
+ The intended usage is that a thread which intends to gomp_barrier_destroy
+ this barrier calls gomp_barrier_wait, while all other threads
+ call gomp_barrier_wait_last. When gomp_barrier_wait returns,
+ the barrier can be safely destroyed. */
+
+void
+gomp_barrier_wait_last (gomp_barrier_t *bar)
+{
+ /* Deferring to gomp_barrier_wait does not use the optimization opportunity
+ allowed by the interface contract for all-but-last participants. The
+ original implementation in config/linux/bar.c handles this better. */
+ gomp_barrier_wait (bar);
+}
+
+void
+gomp_team_barrier_wake (gomp_barrier_t *bar, int count)
+{
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+}
+
+void
+gomp_team_barrier_wait_end (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+ unsigned int generation, gen;
+
+ if (__builtin_expect (state & BAR_WAS_LAST, 0))
+ {
+ /* Next time we'll be awaiting TOTAL threads again. */
+ struct gomp_thread *thr = gomp_thread ();
+ struct gomp_team *team = thr->ts.team;
+
+ bar->awaited = bar->total;
+ team->work_share_cancelled = 0;
+ if (__builtin_expect (team->task_count, 0))
+ {
+ gomp_barrier_handle_tasks (state);
+ state &= ~BAR_WAS_LAST;
+ }
+ else
+ {
+ state &= ~BAR_CANCELLED;
+ state += BAR_INCR - BAR_WAS_LAST;
+ __atomic_store_n (&bar->generation, state, MEMMODEL_RELEASE);
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+ return;
+ }
+ }
+
+ generation = state;
+ state &= ~BAR_CANCELLED;
+ do
+ {
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+ gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ if (__builtin_expect (gen & BAR_TASK_PENDING, 0))
+ {
+ gomp_barrier_handle_tasks (state);
+ gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ }
+ generation |= gen & BAR_WAITING_FOR_TASK;
+ }
+ while (gen != state + BAR_INCR);
+}
+
+void
+gomp_team_barrier_wait (gomp_barrier_t *bar)
+{
+ gomp_team_barrier_wait_end (bar, gomp_barrier_wait_start (bar));
+}
+
+void
+gomp_team_barrier_wait_final (gomp_barrier_t *bar)
+{
+ gomp_barrier_state_t state = gomp_barrier_wait_final_start (bar);
+ if (__builtin_expect (state & BAR_WAS_LAST, 0))
+ bar->awaited_final = bar->total;
+ gomp_team_barrier_wait_end (bar, state);
+}
+
+bool
+gomp_team_barrier_wait_cancel_end (gomp_barrier_t *bar,
+ gomp_barrier_state_t state)
+{
+ unsigned int generation, gen;
+
+ if (__builtin_expect (state & BAR_WAS_LAST, 0))
+ {
+ /* Next time we'll be awaiting TOTAL threads again. */
+ /* BAR_CANCELLED should never be set in state here, because
+ cancellation means that at least one of the threads has been
+ cancelled, thus on a cancellable barrier we should never see
+ all threads to arrive. */
+ struct gomp_thread *thr = gomp_thread ();
+ struct gomp_team *team = thr->ts.team;
+
+ bar->awaited = bar->total;
+ team->work_share_cancelled = 0;
+ if (__builtin_expect (team->task_count, 0))
+ {
+ gomp_barrier_handle_tasks (state);
+ state &= ~BAR_WAS_LAST;
+ }
+ else
+ {
+ state += BAR_INCR - BAR_WAS_LAST;
+ __atomic_store_n (&bar->generation, state, MEMMODEL_RELEASE);
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+ return false;
+ }
+ }
+
+ if (__builtin_expect (state & BAR_CANCELLED, 0))
+ return true;
+
+ generation = state;
+ do
+ {
+ asm ("bar.sync 1, %0;" : : "r" (32 * bar->total));
+ gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ if (__builtin_expect (gen & BAR_CANCELLED, 0))
+ return true;
+ if (__builtin_expect (gen & BAR_TASK_PENDING, 0))
+ {
+ gomp_barrier_handle_tasks (state);
+ gen = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ }
+ generation |= gen & BAR_WAITING_FOR_TASK;
+ }
+ while (gen != state + BAR_INCR);
+
+ return false;
+}
+
+bool
+gomp_team_barrier_wait_cancel (gomp_barrier_t *bar)
+{
+ return gomp_team_barrier_wait_cancel_end (bar, gomp_barrier_wait_start (bar));
+}
+
+void
+gomp_team_barrier_cancel (struct gomp_team *team)
+{
+ gomp_mutex_lock (&team->task_lock);
+ if (team->barrier.generation & BAR_CANCELLED)
+ {
+ gomp_mutex_unlock (&team->task_lock);
+ return;
+ }
+ team->barrier.generation |= BAR_CANCELLED;
+ gomp_mutex_unlock (&team->task_lock);
+ gomp_team_barrier_wake (&team->barrier, INT_MAX);
+}
diff --git a/libgomp/config/nvptx/bar.h b/libgomp/config/nvptx/bar.h
new file mode 100644
index 0000000..757edf1
--- /dev/null
+++ b/libgomp/config/nvptx/bar.h
@@ -0,0 +1,166 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is an NVPTX specific implementation of a barrier synchronization
+ mechanism for libgomp. This type is private to the library. This
+ implementation uses atomic instructions and bar.sync instruction. */
+
+#ifndef GOMP_BARRIER_H
+#define GOMP_BARRIER_H 1
+
+#include "mutex.h"
+
+typedef struct
+{
+ unsigned total;
+ unsigned generation;
+ unsigned awaited;
+ unsigned awaited_final;
+} gomp_barrier_t;
+
+typedef unsigned int gomp_barrier_state_t;
+
+/* The generation field contains a counter in the high bits, with a few
+ low bits dedicated to flags. Note that TASK_PENDING and WAS_LAST can
+ share space because WAS_LAST is never stored back to generation. */
+#define BAR_TASK_PENDING 1
+#define BAR_WAS_LAST 1
+#define BAR_WAITING_FOR_TASK 2
+#define BAR_CANCELLED 4
+#define BAR_INCR 8
+
+static inline void gomp_barrier_init (gomp_barrier_t *bar, unsigned count)
+{
+ bar->total = count;
+ bar->awaited = count;
+ bar->awaited_final = count;
+ bar->generation = 0;
+}
+
+static inline void gomp_barrier_reinit (gomp_barrier_t *bar, unsigned count)
+{
+ __atomic_add_fetch (&bar->awaited, count - bar->total, MEMMODEL_ACQ_REL);
+ bar->total = count;
+}
+
+static inline void gomp_barrier_destroy (gomp_barrier_t *bar)
+{
+}
+
+extern void gomp_barrier_wait (gomp_barrier_t *);
+extern void gomp_barrier_wait_last (gomp_barrier_t *);
+extern void gomp_barrier_wait_end (gomp_barrier_t *, gomp_barrier_state_t);
+extern void gomp_team_barrier_wait (gomp_barrier_t *);
+extern void gomp_team_barrier_wait_final (gomp_barrier_t *);
+extern void gomp_team_barrier_wait_end (gomp_barrier_t *,
+ gomp_barrier_state_t);
+extern bool gomp_team_barrier_wait_cancel (gomp_barrier_t *);
+extern bool gomp_team_barrier_wait_cancel_end (gomp_barrier_t *,
+ gomp_barrier_state_t);
+extern void gomp_team_barrier_wake (gomp_barrier_t *, int);
+struct gomp_team;
+extern void gomp_team_barrier_cancel (struct gomp_team *);
+
+static inline gomp_barrier_state_t
+gomp_barrier_wait_start (gomp_barrier_t *bar)
+{
+ unsigned int ret = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ ret &= -BAR_INCR | BAR_CANCELLED;
+ /* A memory barrier is needed before exiting from the various forms
+ of gomp_barrier_wait, to satisfy OpenMP API version 3.1 section
+ 2.8.6 flush Construct, which says there is an implicit flush during
+ a barrier region. This is a convenient place to add the barrier,
+ so we use MEMMODEL_ACQ_REL here rather than MEMMODEL_ACQUIRE. */
+ if (__atomic_add_fetch (&bar->awaited, -1, MEMMODEL_ACQ_REL) == 0)
+ ret |= BAR_WAS_LAST;
+ return ret;
+}
+
+static inline gomp_barrier_state_t
+gomp_barrier_wait_cancel_start (gomp_barrier_t *bar)
+{
+ return gomp_barrier_wait_start (bar);
+}
+
+/* This is like gomp_barrier_wait_start, except it decrements
+ bar->awaited_final rather than bar->awaited and should be used
+ for the gomp_team_end barrier only. */
+static inline gomp_barrier_state_t
+gomp_barrier_wait_final_start (gomp_barrier_t *bar)
+{
+ unsigned int ret = __atomic_load_n (&bar->generation, MEMMODEL_ACQUIRE);
+ ret &= -BAR_INCR | BAR_CANCELLED;
+ /* See above gomp_barrier_wait_start comment. */
+ if (__atomic_add_fetch (&bar->awaited_final, -1, MEMMODEL_ACQ_REL) == 0)
+ ret |= BAR_WAS_LAST;
+ return ret;
+}
+
+static inline bool
+gomp_barrier_last_thread (gomp_barrier_state_t state)
+{
+ return state & BAR_WAS_LAST;
+}
+
+/* All the inlines below must be called with team->task_lock
+ held. */
+
+static inline void
+gomp_team_barrier_set_task_pending (gomp_barrier_t *bar)
+{
+ bar->generation |= BAR_TASK_PENDING;
+}
+
+static inline void
+gomp_team_barrier_clear_task_pending (gomp_barrier_t *bar)
+{
+ bar->generation &= ~BAR_TASK_PENDING;
+}
+
+static inline void
+gomp_team_barrier_set_waiting_for_tasks (gomp_barrier_t *bar)
+{
+ bar->generation |= BAR_WAITING_FOR_TASK;
+}
+
+static inline bool
+gomp_team_barrier_waiting_for_tasks (gomp_barrier_t *bar)
+{
+ return (bar->generation & BAR_WAITING_FOR_TASK) != 0;
+}
+
+static inline bool
+gomp_team_barrier_cancelled (gomp_barrier_t *bar)
+{
+ return __builtin_expect ((bar->generation & BAR_CANCELLED) != 0, 0);
+}
+
+static inline void
+gomp_team_barrier_done (gomp_barrier_t *bar, gomp_barrier_state_t state)
+{
+ bar->generation = (state & -BAR_INCR) + BAR_INCR;
+}
+
+#endif /* GOMP_BARRIER_H */
diff --git a/libgomp/config/nvptx/barrier.c b/libgomp/config/nvptx/barrier.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/barrier.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/doacross.h b/libgomp/config/nvptx/doacross.h
new file mode 100644
index 0000000..fd011d4
--- /dev/null
+++ b/libgomp/config/nvptx/doacross.h
@@ -0,0 +1,60 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is the NVPTX implementation of doacross spinning. */
+
+#ifndef GOMP_DOACROSS_H
+#define GOMP_DOACROSS_H 1
+
+#include "libgomp.h"
+
+static int zero;
+
+static inline int
+cpu_relax (void)
+{
+ int r;
+ /* Here we need a long-latency operation to make the current warp yield.
+ We could use ld.cv, uncached load from system (host) memory, but that
+ would require allocating locked memory in the plugin. Alternatively,
+ we can use ld.cg, which evicts from L1 and caches in L2. */
+ asm volatile ("ld.cg.s32 %0, [%1];" : "=r" (r) : "i" (&zero) : "memory");
+ return r;
+}
+
+static inline void doacross_spin (unsigned long *addr, unsigned long expected,
+ unsigned long cur)
+{
+ /* Prevent compiler from optimizing based on bounds of containing object. */
+ asm ("" : "+r" (addr));
+ do
+ {
+ int i = cpu_relax ();
+ cur = addr[i];
+ }
+ while (cur <= expected);
+}
+
+#endif /* GOMP_DOACROSS_H */
diff --git a/libgomp/config/nvptx/error.c b/libgomp/config/nvptx/error.c
index e69de29..f3e74cd 100644
--- a/libgomp/config/nvptx/error.c
+++ b/libgomp/config/nvptx/error.c
@@ -0,0 +1,42 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file contains routines used to signal errors. On NVPTX, we have
+ one default output stream (stdout), so redirect everything there. */
+
+#include "libgomp.h"
+#include <stdarg.h>
+#include <stdio.h>
+#include <stdlib.h>
+
+#undef vfprintf
+#undef fputs
+#undef fputc
+
+#define vfprintf(stream, fmt, list) vprintf (fmt, list)
+#define fputs(s, stream) printf ("%s", s)
+#define fputc(c, stream) printf ("%c", c)
+
+#include "../../error.c"
diff --git a/libgomp/config/nvptx/icv-device.c b/libgomp/config/nvptx/icv-device.c
new file mode 100644
index 0000000..831ba11
--- /dev/null
+++ b/libgomp/config/nvptx/icv-device.c
@@ -0,0 +1,74 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file defines OpenMP API entry points that accelerator targets are
+ expected to replace. */
+
+#include "libgomp.h"
+
+void
+omp_set_default_device (int device_num __attribute__((unused)))
+{
+}
+
+int
+omp_get_default_device (void)
+{
+ return 0;
+}
+
+int
+omp_get_num_devices (void)
+{
+ return 0;
+}
+
+int
+omp_get_num_teams (void)
+{
+ return gomp_num_teams_var + 1;
+}
+
+int
+omp_get_team_num (void)
+{
+ int ctaid;
+ asm ("mov.u32 %0, %%ctaid.x;" : "=r" (ctaid));
+ return ctaid;
+}
+
+int
+omp_is_initial_device (void)
+{
+ /* NVPTX is an accelerator-only target. */
+ return 0;
+}
+
+ialias (omp_set_default_device)
+ialias (omp_get_default_device)
+ialias (omp_get_num_devices)
+ialias (omp_get_num_teams)
+ialias (omp_get_team_num)
+ialias (omp_is_initial_device)
diff --git a/libgomp/config/nvptx/iter.c b/libgomp/config/nvptx/iter.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/iter.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/iter_ull.c b/libgomp/config/nvptx/iter_ull.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/iter_ull.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/lock.c b/libgomp/config/nvptx/lock.c
index e69de29..7731704 100644
--- a/libgomp/config/nvptx/lock.c
+++ b/libgomp/config/nvptx/lock.c
@@ -0,0 +1,41 @@
+/* Copyright (C) 2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is a NVPTX specific implementation of the public OpenMP locking
+ primitives. */
+
+/* Reuse the generic implementation in terms of gomp_mutex_t. */
+#include "../../lock.c"
+
+ialias (omp_init_lock)
+ialias (omp_init_nest_lock)
+ialias (omp_destroy_lock)
+ialias (omp_destroy_nest_lock)
+ialias (omp_set_lock)
+ialias (omp_set_nest_lock)
+ialias (omp_unset_lock)
+ialias (omp_unset_nest_lock)
+ialias (omp_test_lock)
+ialias (omp_test_nest_lock)
diff --git a/libgomp/config/nvptx/loop.c b/libgomp/config/nvptx/loop.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/loop.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/loop_ull.c b/libgomp/config/nvptx/loop_ull.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/loop_ull.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/mutex.h b/libgomp/config/nvptx/mutex.h
new file mode 100644
index 0000000..e408ca7
--- /dev/null
+++ b/libgomp/config/nvptx/mutex.h
@@ -0,0 +1,60 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is an NVPTX specific implementation of a mutex synchronization
+ mechanism for libgomp. This type is private to the library. This
+ implementation uses atomic instructions and busy waiting. */
+
+#ifndef GOMP_MUTEX_H
+#define GOMP_MUTEX_H 1
+
+typedef int gomp_mutex_t;
+
+#define GOMP_MUTEX_INIT_0 1
+
+static inline void
+gomp_mutex_init (gomp_mutex_t *mutex)
+{
+ *mutex = 0;
+}
+
+static inline void
+gomp_mutex_destroy (gomp_mutex_t *mutex)
+{
+}
+
+static inline void
+gomp_mutex_lock (gomp_mutex_t *mutex)
+{
+ while (__sync_lock_test_and_set (mutex, 1))
+ /* spin */ ;
+}
+
+static inline void
+gomp_mutex_unlock (gomp_mutex_t *mutex)
+{
+ __sync_lock_release (mutex);
+}
+#endif /* GOMP_MUTEX_H */
diff --git a/libgomp/config/nvptx/ordered.c b/libgomp/config/nvptx/ordered.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/ordered.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/parallel.c b/libgomp/config/nvptx/parallel.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/parallel.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/fortran.c b/libgomp/config/nvptx/pool.h
index 71ba6ed..70e233c 100644
--- a/libgomp/config/nvptx/fortran.c
+++ b/libgomp/config/nvptx/pool.h
@@ -1,8 +1,5 @@
-/* OpenACC Runtime Fortran wrapper routines
-
- Copyright (C) 2014-2016 Free Software Foundation, Inc.
-
- Contributed by Mentor Embedded.
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
This file is part of the GNU Offloading and Multi Processing Library
(libgomp).
@@ -26,15 +23,27 @@
see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
<http://www.gnu.org/licenses/>. */
-/* Temporary hack; this will be provided by libgfortran. */
+/* This is the NVPTX implementation of the thread pool management
+ for libgomp. This type is private to the library. */
+
+#ifndef GOMP_POOL_H
+#define GOMP_POOL_H 1
+
+#include "libgomp.h"
+
+/* Get the thread pool. */
+
+static inline struct gomp_thread_pool *
+gomp_get_thread_pool (struct gomp_thread *thr, unsigned nthreads)
+{
+ /* NVPTX is running with a fixed pool of pre-started threads. */
+ return thr->thread_pool;
+}
-extern void _gfortran_abort (void);
+static inline void
+gomp_release_thread_pool (struct gomp_thread_pool *pool)
+{
+ /* Do nothing. */
+}
-__asm__ ("// BEGIN GLOBAL FUNCTION DECL: _gfortran_abort\n"
- ".visible .func _gfortran_abort;\n"
- "// BEGIN GLOBAL FUNCTION DEF: _gfortran_abort\n"
- ".visible .func _gfortran_abort\n"
- "{\n"
- "trap;\n"
- "ret;\n"
- "}\n");
+#endif /* GOMP_POOL_H */
diff --git a/libgomp/config/nvptx/priority_queue.c b/libgomp/config/nvptx/priority_queue.c
deleted file mode 100644
index 63aecd2..0000000
--- a/libgomp/config/nvptx/priority_queue.c
+++ /dev/null
@@ -1 +0,0 @@
-/* Empty stub for omp task priority support. */
diff --git a/libgomp/config/nvptx/proc.c b/libgomp/config/nvptx/proc.c
index e69de29..8c1c366 100644
--- a/libgomp/config/nvptx/proc.c
+++ b/libgomp/config/nvptx/proc.c
@@ -0,0 +1,41 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file contains system specific routines related to counting
+ online processors and dynamic load balancing. */
+
+#include "libgomp.h"
+
+unsigned
+gomp_dynamic_max_threads (void)
+{
+ return gomp_icv (false)->nthreads_var;
+}
+
+int
+omp_get_num_procs (void)
+{
+ return gomp_icv (false)->nthreads_var;
+}
diff --git a/libgomp/config/nvptx/ptrlock.h b/libgomp/config/nvptx/ptrlock.h
new file mode 100644
index 0000000..c2eae75
--- /dev/null
+++ b/libgomp/config/nvptx/ptrlock.h
@@ -0,0 +1,73 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is an NVPTX specific implementation of a mutex synchronization
+ mechanism for libgomp. This type is private to the library. This
+ implementation uses atomic instructions and busy waiting.
+
+ A ptrlock has four states:
+ 0/NULL Initial
+ 1 Owned by me, I get to write a pointer to ptrlock.
+ 2 Some thread is waiting on the ptrlock.
+ >2 Ptrlock contains a valid pointer.
+ It is not valid to gain the ptrlock and then write a NULL to it. */
+
+#ifndef GOMP_PTRLOCK_H
+#define GOMP_PTRLOCK_H 1
+
+typedef void *gomp_ptrlock_t;
+
+static inline void gomp_ptrlock_init (gomp_ptrlock_t *ptrlock, void *ptr)
+{
+ *ptrlock = ptr;
+}
+
+static inline void *gomp_ptrlock_get (gomp_ptrlock_t *ptrlock)
+{
+ uintptr_t v = (uintptr_t) __atomic_load_n (ptrlock, MEMMODEL_ACQUIRE);
+ if (v > 2)
+ return (void *) v;
+
+ if (v == 0
+ && __atomic_compare_exchange_n (ptrlock, &v, 1, false,
+ MEMMODEL_ACQUIRE, MEMMODEL_ACQUIRE))
+ return NULL;
+
+ while (v == 1)
+ v = (uintptr_t) __atomic_load_n (ptrlock, MEMMODEL_ACQUIRE);
+
+ return (void *) v;
+}
+
+static inline void gomp_ptrlock_set (gomp_ptrlock_t *ptrlock, void *ptr)
+{
+ __atomic_store_n (ptrlock, ptr, MEMMODEL_RELEASE);
+}
+
+static inline void gomp_ptrlock_destroy (gomp_ptrlock_t *ptrlock)
+{
+}
+
+#endif /* GOMP_PTRLOCK_H */
diff --git a/libgomp/config/nvptx/sections.c b/libgomp/config/nvptx/sections.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/sections.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/sem.h b/libgomp/config/nvptx/sem.h
new file mode 100644
index 0000000..82c0dfb
--- /dev/null
+++ b/libgomp/config/nvptx/sem.h
@@ -0,0 +1,65 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is an NVPTX specific implementation of a semaphore synchronization
+ mechanism for libgomp. This type is private to the library. This
+ semaphore implementation uses atomic instructions and busy waiting. */
+
+#ifndef GOMP_SEM_H
+#define GOMP_SEM_H 1
+
+typedef int gomp_sem_t;
+
+static inline void
+gomp_sem_init (gomp_sem_t *sem, int value)
+{
+ *sem = value;
+}
+
+static inline void
+gomp_sem_destroy (gomp_sem_t *sem)
+{
+}
+
+static inline void
+gomp_sem_wait (gomp_sem_t *sem)
+{
+ int count = __atomic_load_n (sem, MEMMODEL_ACQUIRE);
+ for (;;)
+ {
+ while (count == 0)
+ count = __atomic_load_n (sem, MEMMODEL_ACQUIRE);
+ if (__atomic_compare_exchange_n (sem, &count, count - 1, false,
+ MEMMODEL_ACQUIRE, MEMMODEL_RELAXED))
+ return;
+ }
+}
+
+static inline void
+gomp_sem_post (gomp_sem_t *sem)
+{
+ (void) __atomic_add_fetch (sem, 1, MEMMODEL_RELEASE);
+}
+#endif /* GOMP_SEM_H */
diff --git a/libgomp/config/nvptx/simple-bar.h b/libgomp/config/nvptx/simple-bar.h
new file mode 100644
index 0000000..e7b56d9
--- /dev/null
+++ b/libgomp/config/nvptx/simple-bar.h
@@ -0,0 +1,70 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is a simplified barrier that is suitable for thread pool
+ synchronizaton. Only a subset of full barrier API (bar.h) is exposed.
+ Here in the NVPTX-specific implementation, we expect that thread pool
+ corresponds to a PTX CTA (thread block). */
+
+#ifndef GOMP_SIMPLE_BARRIER_H
+#define GOMP_SIMPLE_BARRIER_H 1
+
+typedef struct
+{
+ unsigned count;
+} gomp_simple_barrier_t;
+
+static inline void
+gomp_simple_barrier_init (gomp_simple_barrier_t *bar, unsigned count)
+{
+ bar->count = count * 32;
+}
+
+/* Unused on NVPTX.
+static inline void
+gomp_simple_barrier_reinit (gomp_simple_barrier_t *bar, unsigned count)
+{
+ bar->count = count * 32;
+}
+*/
+
+static inline void
+gomp_simple_barrier_destroy (gomp_simple_barrier_t *bar)
+{
+}
+
+static inline void
+gomp_simple_barrier_wait (gomp_simple_barrier_t *bar)
+{
+ asm volatile ("bar.sync 0, %0;" : : "r" (bar->count) : "memory");
+}
+
+static inline void
+gomp_simple_barrier_wait_last (gomp_simple_barrier_t *bar)
+{
+ asm volatile ("bar.arrive 0, %0;" : : "r" (bar->count) : "memory");
+}
+
+#endif /* GOMP_SIMPLE_BARRIER_H */
diff --git a/libgomp/config/nvptx/single.c b/libgomp/config/nvptx/single.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/single.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/splay-tree.c b/libgomp/config/nvptx/splay-tree.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/splay-tree.c
+++ /dev/null
diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c
index e69de29..38ea7f7 100644
--- a/libgomp/config/nvptx/target.c
+++ b/libgomp/config/nvptx/target.c
@@ -0,0 +1,49 @@
+/* Copyright (C) 2013-2016 Free Software Foundation, Inc.
+ Contributed by Jakub Jelinek <jakub@redhat.com>.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+#include "libgomp.h"
+#include <limits.h>
+
+void
+GOMP_teams (unsigned int num_teams, unsigned int thread_limit)
+{
+ if (thread_limit)
+ {
+ struct gomp_task_icv *icv = gomp_icv (true);
+ icv->thread_limit_var
+ = thread_limit > INT_MAX ? UINT_MAX : thread_limit;
+ }
+ unsigned int num_blocks, block_id;
+ asm ("mov.u32 %0, %%nctaid.x;" : "=r" (num_blocks));
+ asm ("mov.u32 %0, %%ctaid.x;" : "=r" (block_id));
+ if (!num_teams || num_teams >= num_blocks)
+ num_teams = num_blocks;
+ else if (block_id >= num_teams)
+ {
+ gomp_free_thread (nvptx_thrs);
+ asm ("exit;");
+ }
+ gomp_num_teams_var = num_teams - 1;
+}
diff --git a/libgomp/config/nvptx/task.c b/libgomp/config/nvptx/task.c
index e69de29..c183716 100644
--- a/libgomp/config/nvptx/task.c
+++ b/libgomp/config/nvptx/task.c
@@ -0,0 +1,43 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>.
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file handles the maintainence of tasks in response to task
+ creation and termination. */
+
+#ifdef __nvptx_softstack__
+
+#include "libgomp.h"
+
+/* NVPTX is an accelerator-only target, so this should never be called. */
+
+bool
+gomp_target_task_fn (void *data)
+{
+ __builtin_unreachable ();
+}
+
+#include "../../task.c"
+
+#endif
diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c
index e69de29..f7b5e3e 100644
--- a/libgomp/config/nvptx/team.c
+++ b/libgomp/config/nvptx/team.c
@@ -0,0 +1,178 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file handles maintainance of threads on NVPTX. */
+
+#if defined __nvptx_softstack__ && defined __nvptx_unisimt__
+
+#include "libgomp.h"
+#include <stdlib.h>
+#include <string.h>
+
+struct gomp_thread *nvptx_thrs __attribute__((shared,nocommon));
+
+static void gomp_thread_start (struct gomp_thread_pool *);
+
+
+/* This externally visible function handles target region entry. It
+ sets up a per-team thread pool and transfers control by calling FN (FN_DATA)
+ in the master thread or gomp_thread_start in other threads.
+
+ The name of this function is part of the interface with the compiler: for
+ each target region, GCC emits a PTX .kernel function that sets up soft-stack
+ and uniform-simt state and calls this function, passing in FN the original
+ function outlined for the target region. */
+
+void
+gomp_nvptx_main (void (*fn) (void *), void *fn_data)
+{
+ int tid, ntids;
+ asm ("mov.u32 %0, %%tid.y;" : "=r" (tid));
+ asm ("mov.u32 %0, %%ntid.y;" : "=r" (ntids));
+ if (tid == 0)
+ {
+ gomp_global_icv.nthreads_var = ntids;
+ /* Starting additional threads is not supported. */
+ gomp_global_icv.dyn_var = true;
+
+ nvptx_thrs = alloca (ntids * sizeof (*nvptx_thrs));
+ memset (nvptx_thrs, 0, ntids * sizeof (*nvptx_thrs));
+
+ struct gomp_thread_pool *pool = alloca (sizeof (*pool));
+ pool->threads = alloca (ntids * sizeof (*pool->threads));
+ for (tid = 0; tid < ntids; tid++)
+ pool->threads[tid] = nvptx_thrs + tid;
+ pool->threads_size = ntids;
+ pool->threads_used = ntids;
+ pool->threads_busy = 1;
+ pool->last_team = NULL;
+ gomp_simple_barrier_init (&pool->threads_dock, ntids);
+
+ nvptx_thrs[0].thread_pool = pool;
+ asm ("bar.sync 0;");
+ fn (fn_data);
+
+ gomp_free_thread (nvptx_thrs);
+ }
+ else
+ {
+ asm ("bar.sync 0;");
+ gomp_thread_start (nvptx_thrs[0].thread_pool);
+ }
+}
+
+/* This function contains the idle loop in which a thread waits
+ to be called up to become part of a team. */
+
+static void
+gomp_thread_start (struct gomp_thread_pool *pool)
+{
+ struct gomp_thread *thr = gomp_thread ();
+
+ gomp_sem_init (&thr->release, 0);
+ thr->thread_pool = pool;
+
+ do
+ {
+ gomp_simple_barrier_wait (&pool->threads_dock);
+ if (!thr->fn)
+ continue;
+ thr->fn (thr->data);
+ thr->fn = NULL;
+
+ struct gomp_task *task = thr->task;
+ gomp_team_barrier_wait_final (&thr->ts.team->barrier);
+ gomp_finish_task (task);
+ }
+ /* Work around an NVIDIA driver bug: when generating sm_50 machine code,
+ it can trash stack pointer R1 in loops lacking exit edges. Add a cheap
+ artificial exit that the driver would not be able to optimize out. */
+ while (nvptx_thrs);
+}
+
+/* Launch a team. */
+
+void
+gomp_team_start (void (*fn) (void *), void *data, unsigned nthreads,
+ unsigned flags, struct gomp_team *team)
+{
+ struct gomp_thread *thr, *nthr;
+ struct gomp_task *task;
+ struct gomp_task_icv *icv;
+ struct gomp_thread_pool *pool;
+ unsigned long nthreads_var;
+
+ thr = gomp_thread ();
+ pool = thr->thread_pool;
+ task = thr->task;
+ icv = task ? &task->icv : &gomp_global_icv;
+
+ /* Always save the previous state, even if this isn't a nested team.
+ In particular, we should save any work share state from an outer
+ orphaned work share construct. */
+ team->prev_ts = thr->ts;
+
+ thr->ts.team = team;
+ thr->ts.team_id = 0;
+ ++thr->ts.level;
+ if (nthreads > 1)
+ ++thr->ts.active_level;
+ thr->ts.work_share = &team->work_shares[0];
+ thr->ts.last_work_share = NULL;
+ thr->ts.single_count = 0;
+ thr->ts.static_trip = 0;
+ thr->task = &team->implicit_task[0];
+ nthreads_var = icv->nthreads_var;
+ gomp_init_task (thr->task, task, icv);
+ team->implicit_task[0].icv.nthreads_var = nthreads_var;
+
+ if (nthreads == 1)
+ return;
+
+ /* Release existing idle threads. */
+ for (unsigned i = 1; i < nthreads; ++i)
+ {
+ nthr = pool->threads[i];
+ nthr->ts.team = team;
+ nthr->ts.work_share = &team->work_shares[0];
+ nthr->ts.last_work_share = NULL;
+ nthr->ts.team_id = i;
+ nthr->ts.level = team->prev_ts.level + 1;
+ nthr->ts.active_level = thr->ts.active_level;
+ nthr->ts.single_count = 0;
+ nthr->ts.static_trip = 0;
+ nthr->task = &team->implicit_task[i];
+ gomp_init_task (nthr->task, task, icv);
+ team->implicit_task[i].icv.nthreads_var = nthreads_var;
+ nthr->fn = fn;
+ nthr->data = data;
+ team->ordered_release[i] = &nthr->release;
+ }
+
+ gomp_simple_barrier_wait (&pool->threads_dock);
+}
+
+#include "../../team.c"
+#endif
diff --git a/libgomp/config/nvptx/time.c b/libgomp/config/nvptx/time.c
index e69de29..88fb130 100644
--- a/libgomp/config/nvptx/time.c
+++ b/libgomp/config/nvptx/time.c
@@ -0,0 +1,49 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Dmitry Melnik <dm@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This file implements timer routines for NVPTX. It uses the %clock64 cycle
+ counter. */
+
+#include "libgomp.h"
+
+/* This is set from host in plugin-nvptx.c. */
+double __nvptx_clocktick = 0;
+
+double
+omp_get_wtime (void)
+{
+ uint64_t clock;
+ asm ("mov.u64 %0, %%clock64;" : "=r" (clock));
+ return clock * __nvptx_clocktick;
+}
+
+double
+omp_get_wtick (void)
+{
+ return __nvptx_clocktick;
+}
+
+ialias (omp_get_wtime)
+ialias (omp_get_wtick)
diff --git a/libgomp/config/nvptx/work.c b/libgomp/config/nvptx/work.c
deleted file mode 100644
index e69de29..0000000
--- a/libgomp/config/nvptx/work.c
+++ /dev/null
diff --git a/libgomp/config/posix/affinity.c b/libgomp/config/posix/affinity.c
deleted file mode 100644
index 32986fa..0000000
--- a/libgomp/config/posix/affinity.c
+++ /dev/null
@@ -1,140 +0,0 @@
-/* Copyright (C) 2006-2016 Free Software Foundation, Inc.
- Contributed by Jakub Jelinek <jakub@redhat.com>.
-
- This file is part of the GNU Offloading and Multi Processing Library
- (libgomp).
-
- Libgomp is free software; you can redistribute it and/or modify it
- under the terms of the GNU General Public License as published by
- the Free Software Foundation; either version 3, or (at your option)
- any later version.
-
- Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
- WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
- FOR A PARTICULAR PURPOSE. See the GNU General Public License for
- more details.
-
- Under Section 7 of GPL version 3, you are granted additional
- permissions described in the GCC Runtime Library Exception, version
- 3.1, as published by the Free Software Foundation.
-
- You should have received a copy of the GNU General Public License and
- a copy of the GCC Runtime Library Exception along with this program;
- see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
- <http://www.gnu.org/licenses/>. */
-
-/* This is a generic stub implementation of a CPU affinity setting. */
-
-#include "libgomp.h"
-
-void
-gomp_init_affinity (void)
-{
-}
-
-void
-gomp_init_thread_affinity (pthread_attr_t *attr, unsigned int place)
-{
- (void) attr;
- (void) place;
-}
-
-void **
-gomp_affinity_alloc (unsigned long count, bool quiet)
-{
- (void) count;
- if (!quiet)
- gomp_error ("Affinity not supported on this configuration");
- return NULL;
-}
-
-void
-gomp_affinity_init_place (void *p)
-{
- (void) p;
-}
-
-bool
-gomp_affinity_add_cpus (void *p, unsigned long num,
- unsigned long len, long stride, bool quiet)
-{
- (void) p;
- (void) num;
- (void) len;
- (void) stride;
- (void) quiet;
- return false;
-}
-
-bool
-gomp_affinity_remove_cpu (void *p, unsigned long num)
-{
- (void) p;
- (void) num;
- return false;
-}
-
-bool
-gomp_affinity_copy_place (void *p, void *q, long stride)
-{
- (void) p;
- (void) q;
- (void) stride;
- return false;
-}
-
-bool
-gomp_affinity_same_place (void *p, void *q)
-{
- (void) p;
- (void) q;
- return false;
-}
-
-bool
-gomp_affinity_finalize_place_list (bool quiet)
-{
- (void) quiet;
- return false;
-}
-
-bool
-gomp_affinity_init_level (int level, unsigned long count, bool quiet)
-{
- (void) level;
- (void) count;
- (void) quiet;
- if (!quiet)
- gomp_error ("Affinity not supported on this configuration");
- return NULL;
-}
-
-void
-gomp_affinity_print_place (void *p)
-{
- (void) p;
-}
-
-int
-omp_get_place_num_procs (int place_num)
-{
- (void) place_num;
- return 0;
-}
-
-void
-omp_get_place_proc_ids (int place_num, int *ids)
-{
- (void) place_num;
- (void) ids;
-}
-
-void
-gomp_get_place_proc_ids_8 (int place_num, int64_t *ids)
-{
- (void) place_num;
- (void) ids;
-}
-
-ialias(omp_get_place_num_procs)
-ialias(omp_get_place_proc_ids)
diff --git a/libgomp/config/posix/simple-bar.h b/libgomp/config/posix/simple-bar.h
new file mode 100644
index 0000000..b77491c
--- /dev/null
+++ b/libgomp/config/posix/simple-bar.h
@@ -0,0 +1,69 @@
+/* Copyright (C) 2015-2016 Free Software Foundation, Inc.
+ Contributed by Alexander Monakov <amonakov@ispras.ru>
+
+ This file is part of the GNU Offloading and Multi Processing Library
+ (libgomp).
+
+ Libgomp is free software; you can redistribute it and/or modify it
+ under the terms of the GNU General Public License as published by
+ the Free Software Foundation; either version 3, or (at your option)
+ any later version.
+
+ Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY
+ WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS
+ FOR A PARTICULAR PURPOSE. See the GNU General Public License for
+ more details.
+
+ Under Section 7 of GPL version 3, you are granted additional
+ permissions described in the GCC Runtime Library Exception, version
+ 3.1, as published by the Free Software Foundation.
+
+ You should have received a copy of the GNU General Public License and
+ a copy of the GCC Runtime Library Exception along with this program;
+ see the files COPYING3 and COPYING.RUNTIME respectively. If not, see
+ <http://www.gnu.org/licenses/>. */
+
+/* This is a simplified barrier that is suitable for thread pool
+ synchronizaton. Only a subset of full barrier API (bar.h) is exposed. */
+
+#ifndef GOMP_SIMPLE_BARRIER_H
+#define GOMP_SIMPLE_BARRIER_H 1
+
+#include "bar.h"
+
+typedef struct
+{
+ gomp_barrier_t bar;
+} gomp_simple_barrier_t;
+
+static inline void
+gomp_simple_barrier_init (gomp_simple_barrier_t *bar, unsigned count)
+{
+ gomp_barrier_init (&bar->bar, count);
+}
+
+static inline void
+gomp_simple_barrier_reinit (gomp_simple_barrier_t *bar, unsigned count)
+{
+ gomp_barrier_reinit (&bar->bar, count);
+}
+
+static inline void
+gomp_simple_barrier_destroy (gomp_simple_barrier_t *bar)
+{
+ gomp_barrier_destroy (&bar->bar);
+}
+
+static inline void
+gomp_simple_barrier_wait (gomp_simple_barrier_t *bar)
+{
+ gomp_barrier_wait (&bar->bar);
+}
+
+static inline void
+gomp_simple_barrier_wait_last (gomp_simple_barrier_t *bar)
+{
+ gomp_barrier_wait_last (&bar->bar);
+}
+
+#endif /* GOMP_SIMPLE_BARRIER_H */