+2018-05-02 Tom de Vries <tom@codesourcery.com>
+
+ PR libgomp/85411
+ * plugin/plugin-nvptx.c (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.
+
2018-05-02 Tom de Vries <tom@codesourcery.com>
PR testsuite/83791
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
goacc_device_type = NULL;
}
+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)
{
goacc_device_num = 0;
parse_acc_device_type ();
+ parse_gomp_openacc_dim ();
goacc_runtime_initialize ();
}
#include "config.h"
#include "gstdint.h"
#include "libgomp-plugin.h"
+#include "gomp-constants.h"
#ifdef HAVE_PTHREAD_H
#include <pthread.h>
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
{
global:
GOMP_PLUGIN_target_task_completion;
} GOMP_PLUGIN_1.0;
+
+GOMP_PLUGIN_1.2 {
+ global:
+ GOMP_PLUGIN_acc_default_dim;
+} GOMP_PLUGIN_1.1;
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];
+}
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
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;
--- /dev/null
+/* { 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;
+}
--- /dev/null
+#include <openacc.h>
+#include <alloca.h>
+#include <string.h>
+#include <stdio.h>
+#include <gomp-constants.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;
+ g = __builtin_goacc_parlevel_id (GOMP_DIM_GANG);
+ w = __builtin_goacc_parlevel_id (GOMP_DIM_WORKER);
+ v = __builtin_goacc_parlevel_id (GOMP_DIM_VECTOR);
+
+ 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;
+}
projects / gcc.git / commitdiff