[gomp-nvptx 4/5] libgomp: remove __nvptx_stacks setup code

[Alexander Monakov <amonakov at ispras dot ru>]

This patch implements the NVPTX libgomp part of the transition to
host-allocated soft stacks.  The wrapper around gomp_nvptx_main previously
responsible for that is no longer needed.

This mostly reverts commit b408f1293e29a009ba70a3fda7b800277e1f310a.

	* config/nvptx/team.c: (gomp_nvptx_main_1): Rename back to...
	(gomp_nvptx_main): ...this; delete the wrapper.

diff --git a/libgomp/config/nvptx/team.c b/libgomp/config/nvptx/team.c
index bc8c4e6..b9f9f9f 100644
--- a/libgomp/config/nvptx/team.c
+++ b/libgomp/config/nvptx/team.c
@@ -34,9 +34,12 @@ struct gomp_thread *nvptx_thrs __attribute__((shared));
 
 static void gomp_thread_start (struct gomp_thread_pool *);
 
-static void __attribute__((noinline))
-gomp_nvptx_main_1 (void (*fn) (void *), void *fn_data, int ntids, int tid)
+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;
@@ -69,30 +72,6 @@ gomp_nvptx_main_1 (void (*fn) (void *), void *fn_data, int ntids, int tid)
     }
 }
 
-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.  */