From 9b94fbc7e485117a931dbf438b2dab31bb3f8f13 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Wed, 24 May 2017 15:23:34 +0200 Subject: [PATCH] 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 --- libgomp/ChangeLog | 14 + libgomp/libgomp.h | 5 +- libgomp/libgomp.map | 6 + libgomp/oacc-mem.c | 22 ++ libgomp/openacc.h | 7 +- .../libgomp.oacc-c-c++-common/lib-32.c | 241 ++++++++++++++++-- .../libgomp.oacc-c-c++-common/lib-38.c | 64 ----- 7 files changed, 272 insertions(+), 87 deletions(-) delete mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog index 14e95ef9ba9..4a9575567b3 100644 --- a/libgomp/ChangeLog +++ b/libgomp/ChangeLog @@ -1,5 +1,19 @@ 2017-05-24 Thomas Schwinge + * 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. diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h index 1769a489fe8..940b5b83f70 100644 --- a/libgomp/libgomp.h +++ b/libgomp/libgomp.h @@ -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 diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map index 4d42c42f441..b43c6deb851 100644 --- a/libgomp/libgomp.map +++ b/libgomp/libgomp.map @@ -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; diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c index 2df220201ba..ff3ed49a586 100644 --- a/libgomp/oacc-mem.c +++ b/libgomp/oacc-mem.c @@ -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 diff --git a/libgomp/openacc.h b/libgomp/openacc.h index 53d0c3955c1..ebccb1856d7 100644 --- a/libgomp/openacc.h +++ b/libgomp/openacc.h @@ -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; diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c index e3f87a85672..6a9e995530f 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c @@ -1,36 +1,245 @@ -/* { dg-do run } */ +/* acc_present_or_create, acc_present_or_copyin, etc. */ +#include #include #include 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); - - d1 = acc_present_or_create (h, N); - if (!d1) + h = (int *) malloc (S); + if (!h) abort (); + for (int i = 0; i < N; ++i) + h[i] = i + 0; - d2 = acc_present_or_create (h, N); - if (!d2) - abort (); + shared_mem = acc_is_present (h, S); - if (d1 != d2) + 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 (); + +#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; + } - d2 = acc_pcreate (h, N); - if (!d2) + 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 (); - acc_delete (h, N); + for (int i = 0; i < N; ++i) + { + if (h[i] != i + (shared_mem ? 18 : 17)) + abort (); + } + + acc_delete (h, S); + d = NULL; + if (!shared_mem) + if (acc_is_present (h, S)) + abort(); free (h); diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c deleted file mode 100644 index 05d8498c1f9..00000000000 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c +++ /dev/null @@ -1,64 +0,0 @@ -/* { dg-do run } */ -/* { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } */ - -#include -#include -#include - -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; -} -- 2.30.2