@@ -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
Here's how I've updated 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. Reaching those shared memory arrays is awkward. I cannot declare them with toplevel asms because the compiler implicitely declares them too, and ptxas does not handle duplicated declaration. Ideally I'd like to be able to say: extern char *__shared __nvptx_stacks[32]; Bernd, is your position on exposing shared memory as first-class address space on NVPTX subject to change? Do you remember what middle-end issues you've encountered when trying that? * 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/config/nvptx/team.c | 37 +++++++++++++++++++++++++++++-------- 1 file changed, 29 insertions(+), 8 deletions(-)