[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(-)
@@ -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 ();
}
@@ -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
{
@@ -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;
@@ -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];
+}
@@ -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
@@ -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;
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}