diff options
| author | Alexander Monakov <amonakov@ispras.ru> | 2015-12-01 13:37:14 +0300 |
|---|---|---|
| committer | Alexander Monakov <amonakov@ispras.ru> | 2015-12-09 19:31:43 +0300 |
| commit | b408f1293e29a009ba70a3fda7b800277e1f310a (patch) | |
| tree | fb687b72a8cea1b9042e70362abbb9d3659753ab | |
| parent | cae79672cf02c532c7ca0343869f6ae264984967 (diff) | |
| download | gcc-b408f1293e29a009ba70a3fda7b800277e1f310a.tar.gz | |
libgomp: update gomp_nvptx_main for -mgomp
Update gomp_nvptx_main to set up shared memory arrays __nvptx_stacks and
__nvptx_uni for -mgomp. Since it makes sense only for -mgomp multilib, I've
wrapped the whole file under #ifdef that checks corresponding built-in macros.
* config/nvptx/team.c (gomp_nvptx_main): Rename to...
(gomp_nvptx_main_1): ... this and mark noinline.
(gomp_nvptx_main): Wrap the above, set up __nvptx_uni and
__nvptx_stacks.
| -rw-r--r-- | libgomp/ChangeLog.gomp-nvptx | 7 | ||||
| -rw-r--r-- | libgomp/config/nvptx/team.c | 37 |
2 files changed, 36 insertions, 8 deletions
diff --git a/libgomp/ChangeLog.gomp-nvptx b/libgomp/ChangeLog.gomp-nvptx index e2c3f493fc0..bc100478b48 100644 --- a/libgomp/ChangeLog.gomp-nvptx +++ b/libgomp/ChangeLog.gomp-nvptx @@ -1,5 +1,12 @@ 2015-12-09 Alexander Monakov <amonakov@ispras.ru> + * config/nvptx/team.c (gomp_nvptx_main): Rename to... + (gomp_nvptx_main_1): ... this and mark noinline. + (gomp_nvptx_main): Wrap the above, set up __nvptx_uni and + __nvptx_stacks. + +2015-12-09 Alexander Monakov <amonakov@ispras.ru> + * plugin/plugin-nvptx.c (nvptx_open_device): Adjust heap size. 2015-12-09 Alexander Monakov <amonakov@ispras.ru> diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c index 88d1d346315..c18517a6cd0 100644 --- a/libgomp/config/nvptx/team.c +++ b/libgomp/config/nvptx/team.c @@ -24,6 +24,8 @@ /* This file handles the maintainence of threads on NVPTX. */ +#if defined __nvptx_softstack__ && defined __nvptx_unisimt__ + #include "libgomp.h" #include <stdlib.h> @@ -31,15 +33,9 @@ struct gomp_thread *nvptx_thrs; static void gomp_thread_start (struct gomp_thread_pool *); -void -gomp_nvptx_main (void (*fn) (void *), void *fn_data) +static void __attribute__((noinline)) +gomp_nvptx_main_1 (void (*fn) (void *), void *fn_data, int ntids, int tid) { - int ntids, tid, laneid; - asm ("mov.u32 %0, %%laneid;" : "=r" (laneid)); - if (laneid) - return; - 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; @@ -72,6 +68,30 @@ gomp_nvptx_main (void (*fn) (void *), void *fn_data) } } +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)); + char *stacks = 0; + int *__nvptx_uni; + asm ("cvta.shared.u64 %0, __nvptx_uni;" : "=r" (__nvptx_uni)); + __nvptx_uni[tid] = 0; + if (tid == 0) + { + size_t stacksize = 131072; + stacks = gomp_malloc (stacksize * ntids); + char **__nvptx_stacks = 0; + asm ("cvta.shared.u64 %0, __nvptx_stacks;" : "=r" (__nvptx_stacks)); + for (int i = 0; i < ntids; i++) + __nvptx_stacks[i] = stacks + stacksize * (i + 1); + } + asm ("bar.sync 0;"); + gomp_nvptx_main_1 (fn, fn_data, ntids, tid); + free (stacks); +} + /* This function is a pthread_create entry point. This contains the idle loop in which a thread waits to be called up to become part of a team. */ @@ -160,3 +180,4 @@ gomp_team_start (void (*fn) (void *), void *data, unsigned nthreads, } #include "../../team.c" +#endif |