OpenE2K/gcc
2
0
Fork 0
Code Issues Pull Requests Projects Releases Wiki Activity

[openacc] Move GOMP_OPENACC_DIM parsing out of nvptx plugin

Browse Source 
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.

From-SVN: r259852
This commit is contained in:
Tom de Vries 2018-05-02 17:53:56 +00:00 committed by Tom de Vries
parent 92243e7cea
commit ec00d3faf4
9 changed files with 227 additions and 27 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

16
libgomp/ChangeLog
View File

@ -1,3 +1,19 @@
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

32
libgomp/env.c
View File

@ -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 ();
}

2
libgomp/libgomp.h
View File

@ -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
{

5
libgomp/libgomp.map
View File

@ -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;

11
libgomp/oacc-plugin.c
View File

@ -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];
}

1
libgomp/oacc-plugin.h
View File

@ -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

29
libgomp/plugin/plugin-nvptx.c
View File

@ -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;

13
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default-runtime.c Normal file
View File

@ -0,0 +1,13 @@
/* { 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;
}

145
libgomp/testsuite/libgomp.oacc-c-c++-common/loop-default.h Normal file
View File

@ -0,0 +1,145 @@
#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;
}