From b408f1293e29a009ba70a3fda7b800277e1f310a Mon Sep 17 00:00:00 2001 From: Alexander Monakov Date: Tue, 1 Dec 2015 13:37:14 +0300 Subject: 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. --- libgomp/ChangeLog.gomp-nvptx | 7 +++++++ 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,3 +1,10 @@ +2015-12-09 Alexander Monakov + + * 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 * plugin/plugin-nvptx.c (nvptx_open_device): Adjust heap size. 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 @@ -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 -- cgit v1.2.1