projects / thirdparty / gcc.git / blame
| Commit | Line | Data |
|---|---|---|
| cbe34bb5 | 1 | /* Copyright (C) 2015-2017 Free Software Foundation, Inc. |
| 6103184e AM |
2 | Contributed by Alexander Monakov <amonakov@ispras.ru> |
| 3 | ||
| 4 | This file is part of the GNU Offloading and Multi Processing Library | |
| 5 | (libgomp). | |
| 6 | ||
| 7 | Libgomp is free software; you can redistribute it and/or modify it | |
| 8 | under the terms of the GNU General Public License as published by | |
| 9 | the Free Software Foundation; either version 3, or (at your option) | |
| 10 | any later version. | |
| 11 | ||
| 12 | Libgomp is distributed in the hope that it will be useful, but WITHOUT ANY | |
| 13 | WARRANTY; without even the implied warranty of MERCHANTABILITY or FITNESS | |
| 14 | FOR A PARTICULAR PURPOSE. See the GNU General Public License for | |
| 15 | more details. | |
| 16 | ||
| 17 | Under Section 7 of GPL version 3, you are granted additional | |
| 18 | permissions described in the GCC Runtime Library Exception, version | |
| 19 | 3.1, as published by the Free Software Foundation. | |
| 20 | ||
| 21 | You should have received a copy of the GNU General Public License and | |
| 22 | a copy of the GCC Runtime Library Exception along with this program; | |
| 23 | see the files COPYING3 and COPYING.RUNTIME respectively. If not, see | |
| 24 | <http://www.gnu.org/licenses/>. */ | |
| 25 | ||
| 26 | /* This is a simplified barrier that is suitable for thread pool | |
| 27 | synchronizaton. Only a subset of full barrier API (bar.h) is exposed. | |
| 28 | Here in the NVPTX-specific implementation, we expect that thread pool | |
| 29 | corresponds to a PTX CTA (thread block). */ | |
| 30 | ||
| 31 | #ifndef GOMP_SIMPLE_BARRIER_H | |
| 32 | #define GOMP_SIMPLE_BARRIER_H 1 | |
| 33 | ||
| 34 | typedef struct | |
| 35 | { | |
| 36 | unsigned count; | |
| 37 | } gomp_simple_barrier_t; | |
| 38 | ||
| 39 | static inline void | |
| 40 | gomp_simple_barrier_init (gomp_simple_barrier_t *bar, unsigned count) | |
| 41 | { | |
| 42 | bar->count = count * 32; | |
| 43 | } | |
| 44 | ||
| 45 | /* Unused on NVPTX. | |
| 46 | static inline void | |
| 47 | gomp_simple_barrier_reinit (gomp_simple_barrier_t *bar, unsigned count) | |
| 48 | { | |
| 49 | bar->count = count * 32; | |
| 50 | } | |
| 51 | */ | |
| 52 | ||
| 53 | static inline void | |
| 54 | gomp_simple_barrier_destroy (gomp_simple_barrier_t *bar) | |
| 55 | { | |
| 56 | } | |
| 57 | ||
| 58 | static inline void | |
| 59 | gomp_simple_barrier_wait (gomp_simple_barrier_t *bar) | |
| 60 | { | |
| 61 | asm volatile ("bar.sync 0, %0;" : : "r" (bar->count) : "memory"); | |
| 62 | } | |
| 63 | ||
| 64 | static inline void | |
| 65 | gomp_simple_barrier_wait_last (gomp_simple_barrier_t *bar) | |
| 66 | { | |
| 67 | asm volatile ("bar.arrive 0, %0;" : : "r" (bar->count) : "memory"); | |
| 68 | } | |
| 69 | ||
| 70 | #endif /* GOMP_SIMPLE_BARRIER_H */ |