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

C/C++ OpenACC: acc_pcopyin, acc_pcreate 

libgomp/
	* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
	of preprocessor definitions.
	* libgomp.h (strong_alias): Guard by "#ifdef
	HAVE_ATTRIBUTE_ALIAS".
	* oacc-mem.c: Provide "acc_pcreate" as alias for
	"acc_present_or_create", and "acc_pcopyin" as alias for
	"acc_present_or_copyin".
	* libgomp.map: New version "OACC_2.0.1".
	(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
	* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
	its content into...
	* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
	Extend testing.

From-SVN: r248410
This commit is contained in:
Thomas Schwinge 2017-05-24 15:23:34 +02:00 committed by Thomas Schwinge
parent 3f3fb6c913
commit 9b94fbc7e4
7 changed files with 273 additions and 88 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

14
libgomp/ChangeLog
View File

@ -1,5 +1,19 @@
2017-05-24 Thomas Schwinge <thomas@codesourcery.com>
* openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes instead
of preprocessor definitions.
* libgomp.h (strong_alias): Guard by "#ifdef
HAVE_ATTRIBUTE_ALIAS".
* oacc-mem.c: Provide "acc_pcreate" as alias for
"acc_present_or_create", and "acc_pcopyin" as alias for
"acc_present_or_copyin".
* libgomp.map: New version "OACC_2.0.1".
(OACC_2.0.1): Add "acc_pcopyin", and "acc_pcreate".
* testsuite/libgomp.oacc-c-c++-common/lib-38.c: Remove, merging
its content into...
* testsuite/libgomp.oacc-c-c++-common/lib-32.c: ... this file.
Extend testing.
* plugin/plugin-nvptx.c (nvptx_get_num_devices): Debugging output
when disabling nvptx offloading.

5
libgomp/libgomp.h
View File

@ -1060,8 +1060,6 @@ extern void gomp_set_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern void gomp_unset_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
# define strong_alias(fn, al) \
extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define omp_lock_symver(fn) \
__asm (".symver g" #fn "_30, " #fn "@@OMP_3.0"); \
__asm (".symver g" #fn "_25, " #fn "@OMP_1.0");
@ -1085,6 +1083,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t *) __GOMP_NOTHROW;
#endif
#ifdef HAVE_ATTRIBUTE_ALIAS
# define strong_alias(fn, al) \
extern __typeof (fn) al __attribute__ ((alias (#fn)));
# define ialias_ulp ialias_str1(__USER_LABEL_PREFIX__)
# define ialias_str1(x) ialias_str2(x)
# define ialias_str2(x) #x

6
libgomp/libgomp.map
View File

@ -378,6 +378,12 @@ OACC_2.0 {
acc_set_cuda_stream;
};
OACC_2.0.1 {
global:
acc_pcopyin;
acc_pcreate;
} OACC_2.0;
GOACC_2.0 {
global:
GOACC_data_end;

22
libgomp/oacc-mem.c
View File

@ -514,12 +514,34 @@ acc_present_or_create (void *h, size_t s)
return present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s);
}
/* acc_pcreate is acc_present_or_create by a different name. */
#ifdef HAVE_ATTRIBUTE_ALIAS
strong_alias (acc_present_or_create, acc_pcreate)
#else
void *
acc_pcreate (void *h, size_t s)
{
return acc_present_or_create (h, s);
}
#endif
void *
acc_present_or_copyin (void *h, size_t s)
{
return present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s);
}
/* acc_pcopyin is acc_present_or_copyin by a different name. */
#ifdef HAVE_ATTRIBUTE_ALIAS
strong_alias (acc_present_or_copyin, acc_pcopyin)
#else
void *
acc_pcopyin (void *h, size_t s)
{
return acc_present_or_copyin (h, s);
}
#endif
#define FLAG_COPYOUT (1 << 0)
static void

7
libgomp/openacc.h
View File

@ -91,8 +91,10 @@ void acc_free (void *) __GOACC_NOTHROW;
the standard specifies otherwise. */
void *acc_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_copyin (void *, size_t) __GOACC_NOTHROW;
void *acc_pcopyin (void *, size_t) __GOACC_NOTHROW;
void *acc_create (void *, size_t) __GOACC_NOTHROW;
void *acc_present_or_create (void *, size_t) __GOACC_NOTHROW;
void *acc_pcreate (void *, size_t) __GOACC_NOTHROW;
void acc_copyout (void *, size_t) __GOACC_NOTHROW;
void acc_delete (void *, size_t) __GOACC_NOTHROW;
void acc_update_device (void *, size_t) __GOACC_NOTHROW;
@ -105,11 +107,6 @@ int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
/* Old names. OpenACC does not specify whether these can or must
not be macros, inlines or aliases for the new names. */
#define acc_pcreate acc_present_or_create
#define acc_pcopyin acc_present_or_copyin
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
void *acc_get_current_cuda_context (void) __GOACC_NOTHROW;

243
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c
View File

@ -1,36 +1,245 @@
/* { dg-do run } */
/* acc_present_or_create, acc_present_or_copyin, etc. */
#include <stdbool.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
const int N = 256;
unsigned char *h;
void *d1, *d2;
int *h, *d;
const int N = 10000;
const int S = N * sizeof *h;
bool shared_mem;
h = (unsigned char *) malloc (N);
h = (int *) malloc (S);
if (!h)
abort ();
for (int i = 0; i < N; ++i)
h[i] = i + 0;
d1 = acc_present_or_create (h, N);
if (!d1)
shared_mem = acc_is_present (h, S);
d = (int *) acc_present_or_create (h, S);
if (!d)
abort ();
if (shared_mem)
if (h != d)
abort ();
if (!acc_is_present (h, S))
abort ();
d2 = acc_present_or_create (h, N);
if (!d2)
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
d[i] = i + 1;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 1 : 0))
abort ();
h[i] = i + 2;
}
{
int *d_ = (int *) acc_present_or_create (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 2 : 1))
abort ();
d[i] = i + 3;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 3 : 2))
abort ();
h[i] = i + 4;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 4 : 3))
abort ();
d[i] = i + 5;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 5 : 4))
abort ();
h[i] = i + 6;
}
{
int *d_ = (int *) acc_present_or_copyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 6 : 5))
abort ();
d[i] = i + 7;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 7 : 6))
abort ();
h[i] = i + 8;
}
{
int *d_ = (int *) acc_pcopyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 8 : 7))
abort ();
d[i] = i + 9;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 9 : 8))
abort ();
h[i] = i + 10;
}
acc_copyout (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
abort ();
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 10 : 9))
abort ();
}
d = (int *) acc_pcopyin (h, S);
if (!d)
abort ();
if (shared_mem)
if (h != d)
abort ();
if (!acc_is_present (h, S))
abort ();
if (d1 != d2)
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 10 : 9))
abort ();
d[i] = i + 11;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 11 : 9))
abort ();
h[i] = i + 12;
}
{
int *d_ = (int *) acc_pcopyin (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 12 : 11))
abort ();
d[i] = i + 13;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 13 : 12))
abort ();
h[i] = i + 14;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 14 : 13))
abort ();
d[i] = i + 15;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 15 : 14))
abort ();
h[i] = i + 16;
}
{
int *d_ = (int *) acc_pcreate (h, S);
if (d_ != d)
abort ();
}
#pragma acc parallel loop deviceptr (d)
for (int i = 0; i < N; ++i)
{
if (d[i] != i + (shared_mem ? 16 : 15))
abort ();
d[i] = i + 17;
}
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 17 : 16))
abort ();
h[i] = i + 18;
}
acc_update_self (h, S);
if (!acc_is_present (h, S))
abort ();
d2 = acc_pcreate (h, N);
if (!d2)
abort ();
for (int i = 0; i < N; ++i)
{
if (h[i] != i + (shared_mem ? 18 : 17))
abort ();
}
if (d1 != d2)
abort ();
acc_delete (h, N);
acc_delete (h, S);
d = NULL;
if (!shared_mem)
if (acc_is_present (h, S))
abort();
free (h);

64
libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c
View File

@ -1,64 +0,0 @@
/* { dg-do run } */
/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */
#include <string.h>
#include <stdlib.h>
#include <openacc.h>
int
main (int argc, char **argv)
{
const int N = 256;
int i;
unsigned char *h;
void *d1, *d2;
h = (unsigned char *) malloc (N);
for (i = 0; i < N; i++)
{
h[i] = i;
}
d1 = acc_present_or_copyin (h, N);
if (!d1)
abort ();
for (i = 0; i < N; i++)
{
h[i] = 0xab;
}
d2 = acc_present_or_copyin (h, N);
if (!d2)
abort ();
if (d1 != d2)
abort ();
memset (&h[0], 0, N);
acc_copyout (h, N);
for (i = 0; i < N; i++)
{
if (h[i] != i)
abort ();
}
d2 = acc_pcopyin (h, N);
if (!d2)
abort ();
acc_copyout (h, N);
for (i = 0; i < N; i++)
{
if (h[i] != i)
abort ();
}
free (h);
return 0;
}