From ec00d3faf4e3d20906e8e6038299343f960dc49e Mon Sep 17 00:00:00 2001 From: Tom de Vries Date: Wed, 2 May 2018 17:53:56 +0000 Subject: [PATCH] [openacc] Move GOMP_OPENACC_DIM parsing out of nvptx plugin 2018-05-02 Tom de Vries 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. From-SVN: r259852 --- libgomp/ChangeLog | 16 ++ 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 | 29 +--- .../loop-default-runtime.c | 13 ++ .../libgomp.oacc-c-c++-common/loop-default.h | 145 ++++++++++++++++++ 9 files changed, 227 insertions(+), 27 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index fd81fa3089f..1d55d8bf361 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,3 +1,19 @@ +2018-05-02 Tom de Vries + + 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 PR testsuite/83791 diff --git a/libgomp/env.c b/libgomp/env.c index 871a3e4cb40..18c90bb09d0 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 @@ -1065,6 +1066,36 @@ parse_acc_device_type (void) 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) { @@ -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 d659cd20379..10ea8940c96 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 @@ -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 f9044ae273b..8752348fbf2 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 475f3571f2f..c04db90691a 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 ae152aaa7b9..0a183bb8834 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 2b875ae2b53..89326e57741 100644 --- a/libgomp/plugin/plugin-nvptx.c +++ b/libgomp/plugin/plugin-nvptx.c @@ -1147,33 +1147,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 00000000000..c6110a150f4 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c @@ -0,0 +1,13 @@ +/* { dg-set-target-env-var GOMP_OPENACC_DIM "8::" } */ + +#include "loop-default.h" +#include + +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 00000000000..a9e26939cf2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h @@ -0,0 +1,145 @@ +#include +#include +#include +#include +#include + +#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; +} -- 2.30.2