Hi! In <openacc.h>, we currently describe acc_pcopyin, acc_pcreate as "old names", but they're not "old" but really "alternative names", with the intention to provide them at symbol level, not via "#define"s. OK for trunk?
commit 89c6599cc343a2f3b087caf9cb47c2c42bb02074 Author: Thomas Schwinge <tho...@codesourcery.com> Date: Wed Apr 19 15:37:11 2017 +0200 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 (OACC_2.0): 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. --- libgomp/libgomp.h | 5 +- libgomp/libgomp.map | 2 + libgomp/oacc-mem.c | 22 ++ libgomp/openacc.h | 7 +- .../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++++-- .../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------ 6 files changed, 254 insertions(+), 87 deletions(-) diff --git libgomp/libgomp.h libgomp/libgomp.h index 1769a48..940b5b8 100644 --- libgomp/libgomp.h +++ 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 libgomp/libgomp.map libgomp/libgomp.map index 4d42c42..c7bf245 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -335,6 +335,7 @@ OACC_2.0 { acc_copyin_64_h_; acc_copyin_array_h_; acc_present_or_copyin; + acc_pcopyin; acc_present_or_copyin_32_h_; acc_present_or_copyin_64_h_; acc_present_or_copyin_array_h_; @@ -343,6 +344,7 @@ OACC_2.0 { acc_create_64_h_; acc_create_array_h_; acc_present_or_create; + acc_pcreate; acc_present_or_create_32_h_; acc_present_or_create_64_h_; acc_present_or_create_array_h_; diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c index 2df2202..ff3ed49 100644 --- libgomp/oacc-mem.c +++ 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 libgomp/openacc.h libgomp/openacc.h index 53d0c39..ebccb18 100644 --- libgomp/openacc.h +++ 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 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c index e3f87a8..6a9e995 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/lib-32.c +++ 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 <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); - - 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 libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c deleted file mode 100644 index 05d8498..0000000 --- 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 <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; -} Grüße Thomas