[openacc,PR85411] Move GOMP_OPENACC_DIM parsing out of nvptx plugin

Message ID 694a714b-d250-ad64-e015-567e253ef301@mentor.com
State New
Headers show
Series
  • [openacc,PR85411] Move GOMP_OPENACC_DIM parsing out of nvptx plugin
Related show

Commit Message

Tom de Vries April 16, 2018, 9:41 a.m.
Hi,

this patch moves the parsing of the GOMP_OPENACC_DIM environment 
variable from the nvptx target plugin to the libgomp library.

The variable is not part of the OpenACC standard, but it is specific for 
the gcc implementation of OpenACC, so it makes sense to share the part 
handling the parsing, rather than having each target plugin duplicate it.

Build on x86_64 with nvptx accelerator and reg-tested libgomp.

OK for stage1?

Thanks,
- Tom

Comments

Jakub Jelinek April 16, 2018, 9:43 a.m. | #1
On Mon, Apr 16, 2018 at 11:41:35AM +0200, Tom de Vries wrote:
> Hi,
> 
> this patch moves the parsing of the GOMP_OPENACC_DIM environment variable
> from the nvptx target plugin to the libgomp library.
> 
> The variable is not part of the OpenACC standard, but it is specific for the
> gcc implementation of OpenACC, so it makes sense to share the part handling
> the parsing, rather than having each target plugin duplicate it.
> 
> Build on x86_64 with nvptx accelerator and reg-tested libgomp.
> 
> OK for stage1?

Ok.

> [openacc] Move GOMP_OPENACC_DIM parsing out of nvptx plugin
> 
> 2018-04-15  Tom de Vries  <tom@codesourcery.com>
> 
> 	PR libgomp/85411
> 	* plugin/plugin-nvptx.c (notify_var): Remove no longer used function.
> 	(nvptx_exec): Move parsing of
> 	GOMP_OPENACC_DIM ...
> 	* env.c (parse_gomp_openacc_dim): ... here.  New function.
> 	(initialize_env): Call parse_gomp_openacc_dim.
> 	(goacc_default_dims): Define.
> 	* libgomp.h (goacc_default_dims): Declare.
> 	* oacc-plugin.c (GOMP_PLUGIN_acc_default_dim): New function.
> 	* oacc-plugin.h (GOMP_PLUGIN_acc_default_dim): Declare.
> 	* libgomp.map: New version "GOMP_PLUGIN_1.2". Add
> 	GOMP_PLUGIN_acc_default_dim.
> 	* testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c: New test.
> 	* testsuite/libgomp.oacc-c-c++-common/loop-default.h: New test.

	Jakub

Patch

[openacc] Move GOMP_OPENACC_DIM parsing out of nvptx plugin

2018-04-15  Tom de Vries  <tom@codesourcery.com>

	PR libgomp/85411
	* plugin/plugin-nvptx.c (notify_var): Remove no longer used function.
	(nvptx_exec): Move parsing of
	GOMP_OPENACC_DIM ...
	* env.c (parse_gomp_openacc_dim): ... here.  New function.
	(initialize_env): Call parse_gomp_openacc_dim.
	(goacc_default_dims): Define.
	* libgomp.h (goacc_default_dims): Declare.
	* oacc-plugin.c (GOMP_PLUGIN_acc_default_dim): New function.
	* oacc-plugin.h (GOMP_PLUGIN_acc_default_dim): Declare.
	* libgomp.map: New version "GOMP_PLUGIN_1.2". Add
	GOMP_PLUGIN_acc_default_dim.
	* testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/loop-default.h: New test.

---
 libgomp/env.c                                      |  32 +++++
 libgomp/libgomp.h                                  |   2 +
 libgomp/libgomp.map                                |   5 +
 libgomp/oacc-plugin.c                              |  11 ++
 libgomp/oacc-plugin.h                              |   1 +
 libgomp/plugin/plugin-nvptx.c                      |  38 +-----
 .../loop-default-runtime.c                         |  16 +++
 .../libgomp.oacc-c-c++-common/loop-default.h       | 144 +++++++++++++++++++++
 8 files changed, 213 insertions(+), 36 deletions(-)

diff --git a/libgomp/env.c b/libgomp/env.c
index 871a3e4..18c90bb 100644
--- a/libgomp/env.c
+++ b/libgomp/env.c
@@ -90,6 +90,7 @@  int gomp_debug_var;
 unsigned int gomp_num_teams_var;
 char *goacc_device_type;
 int goacc_device_num;
+int goacc_default_dims[GOMP_DIM_MAX];
 
 #ifndef LIBGOMP_OFFLOADED_ONLY
 
@@ -1066,6 +1067,36 @@  parse_acc_device_type (void)
 }
 
 static void
+parse_gomp_openacc_dim (void)
+{
+  /* The syntax is the same as for the -fopenacc-dim compilation option.  */
+  const char *var_name = "GOMP_OPENACC_DIM";
+  const char *env_var = getenv (var_name);
+  if (!env_var)
+    return;
+
+  const char *pos = env_var;
+  int i;
+  for (i = 0; *pos && i != GOMP_DIM_MAX; i++)
+    {
+      if (i && *pos++ != ':')
+	break;
+
+      if (*pos == ':')
+	continue;
+
+      const char *eptr;
+      errno = 0;
+      long val = strtol (pos, (char **)&eptr, 10);
+      if (errno || val < 0 || (unsigned)val != val)
+	break;
+
+      goacc_default_dims[i] = (int)val;
+      pos = eptr;
+    }
+}
+
+static void
 handle_omp_display_env (unsigned long stacksize, int wait_policy)
 {
   const char *env;
@@ -1336,6 +1367,7 @@  initialize_env (void)
     goacc_device_num = 0;
 
   parse_acc_device_type ();
+  parse_gomp_openacc_dim ();
 
   goacc_runtime_initialize ();
 }
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index d659cd2..10ea894 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -44,6 +44,7 @@ 
 #include "config.h"
 #include "gstdint.h"
 #include "libgomp-plugin.h"
+#include "gomp-constants.h"
 
 #ifdef HAVE_PTHREAD_H
 #include <pthread.h>
@@ -367,6 +368,7 @@  extern unsigned int gomp_num_teams_var;
 extern int gomp_debug_var;
 extern int goacc_device_num;
 extern char *goacc_device_type;
+extern int goacc_default_dims[GOMP_DIM_MAX];
 
 enum gomp_task_kind
 {
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index f9044ae..8752348 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -420,3 +420,8 @@  GOMP_PLUGIN_1.1 {
   global:
 	GOMP_PLUGIN_target_task_completion;
 } GOMP_PLUGIN_1.0;
+
+GOMP_PLUGIN_1.2 {
+  global:
+	GOMP_PLUGIN_acc_default_dim;
+} GOMP_PLUGIN_1.1;
diff --git a/libgomp/oacc-plugin.c b/libgomp/oacc-plugin.c
index 475f357..c04db90 100644
--- a/libgomp/oacc-plugin.c
+++ b/libgomp/oacc-plugin.c
@@ -49,3 +49,14 @@  GOMP_PLUGIN_acc_thread (void)
   struct goacc_thread *thr = goacc_thread ();
   return thr ? thr->target_tls : NULL;
 }
+
+int
+GOMP_PLUGIN_acc_default_dim (unsigned int i)
+{
+  if (i >= GOMP_DIM_MAX)
+    {
+      gomp_fatal ("invalid dimension argument: %d", i);
+      return -1;
+    }
+  return goacc_default_dims[i];
+}
diff --git a/libgomp/oacc-plugin.h b/libgomp/oacc-plugin.h
index ae152aa..0a183bb 100644
--- a/libgomp/oacc-plugin.h
+++ b/libgomp/oacc-plugin.h
@@ -29,5 +29,6 @@ 
 
 extern void GOMP_PLUGIN_async_unmap_vars (void *, int);
 extern void *GOMP_PLUGIN_acc_thread (void);
+extern int GOMP_PLUGIN_acc_default_dim (unsigned int);
 
 #endif
diff --git a/libgomp/plugin/plugin-nvptx.c b/libgomp/plugin/plugin-nvptx.c
index 9ae6095..365b45e 100644
--- a/libgomp/plugin/plugin-nvptx.c
+++ b/libgomp/plugin/plugin-nvptx.c
@@ -867,15 +867,6 @@  nvptx_get_num_devices (void)
   return n;
 }
 
-static void
-notify_var (const char *var_name, const char *env_var)
-{
-  if (env_var == NULL)
-    GOMP_PLUGIN_debug (0, "%s: <Not defined>\n", var_name);
-  else
-    GOMP_PLUGIN_debug (0, "%s: '%s'\n", var_name, env_var);
-}
-
 static bool
 link_ptx (CUmodule *module, const struct targ_ptx_obj *ptx_objs,
 	  unsigned num_objs)
@@ -1097,33 +1088,8 @@  nvptx_exec (void (*fn), size_t mapnum, void **hostaddrs, void **devaddrs,
       pthread_mutex_lock (&ptx_dev_lock);
       if (!default_dims[0])
 	{
-	  const char *var_name = "GOMP_OPENACC_DIM";
-	  /* We only read the environment variable once.  You can't
-	     change it in the middle of execution.  The syntax  is
-	     the same as for the -fopenacc-dim compilation option.  */
-	  const char *env_var = getenv (var_name);
-	  notify_var (var_name, env_var);
-	  if (env_var)
-	    {
-	      const char *pos = env_var;
-
-	      for (i = 0; *pos && i != GOMP_DIM_MAX; i++)
-		{
-		  if (i && *pos++ != ':')
-		    break;
-		  if (*pos != ':')
-		    {
-		      const char *eptr;
-
-		      errno = 0;
-		      long val = strtol (pos, (char **)&eptr, 10);
-		      if (errno || val < 0 || (unsigned)val != val)
-			break;
-		      default_dims[i] = (int)val;
-		      pos = eptr;
-		    }
-		}
-	    }
+	  for (int i = 0; i < GOMP_DIM_MAX; ++i)
+	    default_dims[i] = GOMP_PLUGIN_acc_default_dim (i);
 
 	  int warp_size, block_size, dev_size, cpu_size;
 	  CUdevice dev = nvptx_thread()->ptx_dev->dev;
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c
new file mode 100644
index 0000000..e47c29f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c
@@ -0,0 +1,16 @@ 
+/* This code uses nvptx inline assembly guarded with acc_on_device, which is
+   not optimized away at -O0, and then confuses the target assembler.
+   { dg-skip-if "" { *-*-* } { "-O0" } { "" } } */
+/* { dg-set-target-env-var GOMP_OPENACC_DIM "8::" } */
+
+#include "loop-default.h"
+#include <stdlib.h>
+
+int
+main ()
+{
+  if (check_gang (8) != 0)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h
new file mode 100644
index 0000000..5a10cfe
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h
@@ -0,0 +1,144 @@ 
+#include <openacc.h>
+#include <alloca.h>
+#include <string.h>
+#include <stdio.h>
+
+#pragma acc routine seq
+static int __attribute__ ((noinline))
+coord (void)
+{
+  int res = 0;
+
+  if (acc_on_device (acc_device_nvidia))
+    {
+      int g = 0, w = 0, v = 0;
+
+      __asm__ volatile ("mov.u32 %0,%%ctaid.x;" : "=r" (g));
+      __asm__ volatile ("mov.u32 %0,%%tid.y;" : "=r" (w));
+      __asm__ volatile ("mov.u32 %0,%%tid.x;" : "=r" (v));
+      res = (1 << 24) | (g << 16) | (w << 8) | v;
+    }
+
+  return res;
+}
+
+static int
+check (const int *ary, int size, int gp, int wp, int vp)
+{
+  int exit = 0;
+  int ix;
+  int *gangs = (int *)alloca (gp * sizeof (int));
+  int *workers = (int *)alloca (wp * sizeof (int));
+  int *vectors = (int *)alloca (vp * sizeof (int));
+  int offloaded = 0;
+
+  memset (gangs, 0, gp * sizeof (int));
+  memset (workers, 0, wp * sizeof (int));
+  memset (vectors, 0, vp * sizeof (int));
+
+  for (ix = 0; ix < size; ix++)
+    {
+      int g = (ary[ix] >> 16) & 0xff;
+      int w = (ary[ix] >> 8) & 0xff;
+      int v = (ary[ix] >> 0) & 0xff;
+
+      if (g >= gp || w >= wp || v >= vp)
+	{
+	  printf ("unexpected cpu %#x used\n", ary[ix]);
+	  exit = 1;
+	}
+      else
+	{
+	  vectors[v]++;
+	  workers[w]++;
+	  gangs[g]++;
+	}
+      offloaded += ary[ix] >> 24;
+    }
+
+  if (!offloaded)
+    return 0;
+
+  if (offloaded != size)
+    {
+      printf ("offloaded %d times,  expected %d\n", offloaded, size);
+      return 1;
+    }
+
+  for (ix = 0; ix < gp; ix++)
+    if (gangs[ix] != gangs[0])
+      {
+	printf ("gang %d not used %d times\n", ix, gangs[0]);
+	exit = 1;
+      }
+
+  for (ix = 0; ix < wp; ix++)
+    if (workers[ix] != workers[0])
+      {
+	printf ("worker %d not used %d times\n", ix, workers[0]);
+	exit = 1;
+      }
+
+  for (ix = 0; ix < vp; ix++)
+    if (vectors[ix] != vectors[0])
+      {
+	printf ("vector %d not used %d times\n", ix, vectors[0]);
+	exit = 1;
+      }
+
+  return exit;
+}
+
+#define N (32 * 32 * 32)
+int ary[N];
+
+static int
+check_gang (int gp)
+{
+#pragma acc parallel copyout (ary)
+  {
+#pragma acc loop gang (static:1)
+    for (int ix = 0; ix < N; ix++)
+      ary[ix] = coord ();
+  }
+
+  return check (ary, N, gp, 1, 1);
+}
+
+static int
+check_worker (int wp)
+{
+#pragma  acc parallel copyout (ary)
+  {
+#pragma acc loop worker
+    for (int ix = 0; ix < N; ix++)
+      ary[ix] = coord ();
+  }
+
+  return check (ary, N, 1, wp, 1);
+}
+
+static int
+check_vector (int vp)
+{
+#pragma  acc parallel copyout (ary)
+  {
+#pragma acc loop vector
+    for (int ix = 0; ix < N; ix++)
+      ary[ix] = coord ();
+  }
+
+  return check (ary, N, 1, 1, vp);
+}
+
+static int
+test_1 (int gp, int wp, int vp)
+{
+  int exit = 0;
+
+  exit |= check_gang (gp);
+  exit |= check_worker (wp);
+  exit |= check_vector (vp);
+
+  return exit;
+}