diff options
Diffstat (limited to 'libgomp/config/nvptx')
31 files changed, 1244 insertions, 16 deletions
diff --git a/libgomp/config/nvptx/affinity.c b/libgomp/config/nvptx/affinity.c deleted file mode 100644 index e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/affinity.c +++ /dev/null diff --git a/libgomp/config/nvptx/alloc.c b/libgomp/config/nvptx/alloc.c deleted file mode 100644 index e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/alloc.c +++ /dev/null diff --git a/libgomp/config/nvptx/bar.c b/libgomp/config/nvptx/bar.c index e69de29bb2d..a0d8a44dacf 100644 --- a/libgomp/config/nvptx/bar.c +++ b/libgomp/config/nvptx/bar.c @@ -0,0 +1,209 @@ +/* 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) +{ +#if 0 + gomp_barrier_state_t state = gomp_barrier_wait_start (bar); + if (state & BAR_WAS_LAST) + gomp_barrier_wait_end (bar, state); +#else + gomp_barrier_wait (bar); +#endif +} + +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 00000000000..757edf1d7eb --- /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 e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/barrier.c +++ /dev/null diff --git a/libgomp/config/nvptx/critical.c b/libgomp/config/nvptx/critical.c deleted file mode 100644 index e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/critical.c +++ /dev/null diff --git a/libgomp/config/nvptx/doacross.h b/libgomp/config/nvptx/doacross.h new file mode 100644 index 00000000000..fd011d492bd --- /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 e69de29bb2d..f3e74cd9449 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 00000000000..831ba110e95 --- /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 e69de29bb2d..00000000000 --- 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 e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/iter_ull.c +++ /dev/null diff --git a/libgomp/config/nvptx/lock.c b/libgomp/config/nvptx/lock.c index e69de29bb2d..7731704c6fb 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 e69de29bb2d..00000000000 --- 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 e69de29bb2d..00000000000 --- 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 00000000000..e408ca72933 --- /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 e69de29bb2d..00000000000 --- 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 e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/parallel.c +++ /dev/null diff --git a/libgomp/config/nvptx/fortran.c b/libgomp/config/nvptx/pool.h index 71ba6ed73f5..70e233c7caa 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 63aecd249fb..00000000000 --- 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 e69de29bb2d..8c1c3664660 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 00000000000..c2eae75720d --- /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 e69de29bb2d..00000000000 --- 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 00000000000..82c0dfbf7e0 --- /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 00000000000..e7b56d98826 --- /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 e69de29bb2d..00000000000 --- 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 e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/splay-tree.c +++ /dev/null diff --git a/libgomp/config/nvptx/target.c b/libgomp/config/nvptx/target.c index e69de29bb2d..38ea7f7aa68 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 e69de29bb2d..c183716c94e 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 e69de29bb2d..f7b5e3e81b5 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 e69de29bb2d..88fb13050c0 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 e69de29bb2d..00000000000 --- a/libgomp/config/nvptx/work.c +++ /dev/null |