From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 65298 invoked by alias); 24 May 2017 13:40:43 -0000 Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk List-Id: List-Archive: List-Post: List-Help: Sender: gcc-patches-owner@gcc.gnu.org Received: (qmail 65281 invoked by uid 89); 24 May 2017 13:40:42 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-24.5 required=5.0 tests=AWL,BAYES_00,GIT_PATCH_0,GIT_PATCH_1,GIT_PATCH_2,GIT_PATCH_3,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy= X-HELO: relay1.mentorg.com Received: from relay1.mentorg.com (HELO relay1.mentorg.com) (192.94.38.131) by sourceware.org (qpsmtpd/0.93/v0.84-503-g423c35a) with ESMTP; Wed, 24 May 2017 13:40:37 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-04.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1dDWWk-0000bJ-D9 from Thomas_Schwinge@mentor.com ; Wed, 24 May 2017 06:40:38 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.87) by SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Wed, 24 May 2017 14:40:35 +0100 From: Thomas Schwinge To: Jakub Jelinek , Subject: Re: C/C++ OpenACC: acc_pcopyin, acc_pcreate In-Reply-To: <20170524112054.GA8499@tucnak> References: <87bmqlynkn.fsf@hertz.schwinge.homeip.net> <20170523110715.GQ8499@tucnak> <8737buy15s.fsf@hertz.schwinge.homeip.net> <20170524112054.GA8499@tucnak> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.5.1 (x86_64-pc-linux-gnu) Date: Wed, 24 May 2017 13:41:00 -0000 Message-ID: <87poeywey8.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-ClientProxiedBy: svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-SW-Source: 2017-05/txt/msg01863.txt.bz2 Hi! On Wed, 24 May 2017 13:20:54 +0200, Jakub Jelinek wrote: > On Wed, May 24, 2017 at 12:55:27PM +0200, Thomas Schwinge wrote: > > On Tue, 23 May 2017 13:07:15 +0200, Jakub Jelinek wr= ote: > > > new symbols should never be added to an existing symbol > > > version after a GCC version with that symbol version has been release= d. > > would you please remind me where this > > requirement is coming from? >=20 > There are many reasons Thanks for the explanation! > one is e.g. that the dynamic linker uses the symbol > version names during library loading to verify the library satisfies the > needs of a binary or some other library; if you add further symbols into > existing symbol version, then it will be satisfied even by library that > doesn't provide the symbol (in the above case say acc_pcopyin), and you c= an > get a failure either much later (during relocation processing, with lazy > binding that can be much later), or not at all (with lazy binding and if > the function isn't called at all during the runtime of the program). Some might consider that a feature. ;-) > Another one is e.g. rpm uses symbol versions in dependencies, but only > those and not at the required/provided symbol granularity. So, if you > have acc_pcopyin@@OACC_2.0.1 and use it in some program, the program > will have a dependency on libgomp.so.1()(OACC_2.0.1) and thus you > automatically get the right library that provides it. If you add to an > existing symbol version, the dependency will be satisfied even for librar= ies > that don't provide it at all. Ah, yes, I see -- a fine reason. > All in all, symbol versions aren't there just for documentation purposes, > but serve various other roles, including the possibility to change ABI > of existing symbols, but not limited to that. Sure, I was aware of that specific use -- but that's not relevant to the case discussed here: a symbol's (non-)existence will always be detected reliably, whether lazily (see above), or at load time. That's where I had gotten the idea from that it's fine to use existing versions for these new symbols. (But you reminded me why it's better to use a new version here, too.) > The policy of not adding > new symbols into symbol versions that were already shipped in released > libraries is something that is strictly followed in glibc, gcc libraries > (I remember in libstdc++ a bug happened some years ago and some symbols > were unintentionally exported at an earlier version, which has been cured > by moving into a new symbol version and providing the old one just as an @ > (vs. @@) symver for compatibility only; e.g. the current libstdc++ check-= abi > stuff verifies this stuff). > For more info see e.g. https://www.akkadia.org/drepper/symbol-versioning Yeah, thanks, I had read that a few years ago. > > But that will only redirect them at the user side. The intention here = is >=20 > The question is if the OpenACC standard allows that or not. Let's put it this way: it doesn't explicitly disallow that. > > to also care for users providing their own declarations instead of using > > , or using "dlsym", and so on -- quite "pathological", I kno= w, > > but... > > OK for trunk using "OACC_2.0.1" symbol version? >=20 > Ok. Thanks. As posted, committed to trunk in r248410: commit e4d15e02dfbaf61fb5655595423baaff4a345a60 Author: tschwinge Date: Wed May 24 13:23:34 2017 +0000 C/C++ OpenACC: acc_pcopyin, acc_pcreate =20=20=20=20 libgomp/ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes inst= ead 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. =20=20=20=20 git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@248410 138bc75d-0d04-04= 10-961f-82ee72b054a4 --- libgomp/ChangeLog | 14 ++ libgomp/libgomp.h | 5 +- libgomp/libgomp.map | 6 + 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 ------ 7 files changed, 272 insertions(+), 87 deletions(-) diff --git libgomp/ChangeLog libgomp/ChangeLog index 14e95ef..4a95755 100644 --- libgomp/ChangeLog +++ libgomp/ChangeLog @@ -1,5 +1,19 @@ 2017-05-24 Thomas Schwinge =20 + * 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. =20 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; =20 -# 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 =20 #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..b43c6de 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -378,6 +378,12 @@ OACC_2.0 { acc_set_cuda_stream; }; =20 +OACC_2.0.1 { + global: + acc_pcopyin; + acc_pcreate; +} OACC_2.0; + GOACC_2.0 { global: GOACC_data_end; 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); } =20 +/* 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= ); } =20 +/* 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) =20 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; =20 -/* 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/tes= tsuite/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. */ =20 +#include #include #include =20 int main (int argc, char **argv) { - const int N =3D 256; - unsigned char *h; - void *d1, *d2; + int *h, *d; + const int N =3D 10000; + const int S =3D N * sizeof *h; + bool shared_mem; =20 - h =3D (unsigned char *) malloc (N); - - d1 =3D acc_present_or_create (h, N); - if (!d1) + h =3D (int *) malloc (S); + if (!h) abort (); + for (int i =3D 0; i < N; ++i) + h[i] =3D i + 0; =20 - d2 =3D acc_present_or_create (h, N); - if (!d2) - abort (); + shared_mem =3D acc_is_present (h, S); =20 - if (d1 !=3D d2) + d =3D (int *) acc_present_or_create (h, S); + if (!d) abort (); + if (shared_mem) + if (h !=3D d) + abort (); + if (!acc_is_present (h, S)) + abort (); + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + d[i] =3D i + 1; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 1 : 0)) + abort (); + h[i] =3D i + 2; + } + + { + int *d_ =3D (int *) acc_present_or_create (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 2 : 1)) + abort (); + d[i] =3D i + 3; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 3 : 2)) + abort (); + h[i] =3D i + 4; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 4 : 3)) + abort (); + d[i] =3D i + 5; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 5 : 4)) + abort (); + h[i] =3D i + 6; + } + + { + int *d_ =3D (int *) acc_present_or_copyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 6 : 5)) + abort (); + d[i] =3D i + 7; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 7 : 6)) + abort (); + h[i] =3D i + 8; + } + + { + int *d_ =3D (int *) acc_pcopyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 8 : 7)) + abort (); + d[i] =3D i + 9; + } =20 - d2 =3D acc_pcreate (h, N); - if (!d2) + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 9 : 8)) + abort (); + h[i] =3D i + 10; + } + + acc_copyout (h, S); + d =3D NULL; + if (!shared_mem) + if (acc_is_present (h, S)) + abort (); + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 10 : 9)) + abort (); + } + + d =3D (int *) acc_pcopyin (h, S); + if (!d) + abort (); + if (shared_mem) + if (h !=3D d) + abort (); + if (!acc_is_present (h, S)) abort (); =20 - if (d1 !=3D d2) +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 10 : 9)) + abort (); + d[i] =3D i + 11; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 11 : 9)) + abort (); + h[i] =3D i + 12; + } + + { + int *d_ =3D (int *) acc_pcopyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 12 : 11)) + abort (); + d[i] =3D i + 13; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 13 : 12)) + abort (); + h[i] =3D i + 14; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 14 : 13)) + abort (); + d[i] =3D i + 15; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 15 : 14)) + abort (); + h[i] =3D i + 16; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 16 : 15)) + abort (); + d[i] =3D i + 17; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 17 : 16)) + abort (); + h[i] =3D i + 18; + } + + acc_update_self (h, S); + if (!acc_is_present (h, S)) abort (); =20 - acc_delete (h, N); + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 18 : 17)) + abort (); + } + + acc_delete (h, S); + d =3D NULL; + if (!shared_mem) + if (acc_is_present (h, S)) + abort(); =20 free (h); =20 diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/tes= tsuite/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=3D0" } } */ - -#include -#include -#include - -int -main (int argc, char **argv) -{ - const int N =3D 256; - int i; - unsigned char *h; - void *d1, *d2; - - h =3D (unsigned char *) malloc (N); - - for (i =3D 0; i < N; i++) - { - h[i] =3D i; - } - - d1 =3D acc_present_or_copyin (h, N); - if (!d1) - abort (); - - for (i =3D 0; i < N; i++) - { - h[i] =3D 0xab; - } - - d2 =3D acc_present_or_copyin (h, N); - if (!d2) - abort (); - - if (d1 !=3D d2) - abort (); - - memset (&h[0], 0, N); - - acc_copyout (h, N); - - for (i =3D 0; i < N; i++) - { - if (h[i] !=3D i) - abort (); - } - - d2 =3D acc_pcopyin (h, N); - if (!d2) - abort (); - - acc_copyout (h, N); - - for (i =3D 0; i < N; i++) - { - if (h[i] !=3D i) - abort (); - } - - free (h); - - return 0; -} Committed to gomp-4_0-branch in r248414: commit 705bc833826559e64a7ccd821fcf419bfaf2b727 Author: tschwinge Date: Wed May 24 13:25:42 2017 +0000 C/C++ OpenACC: acc_pcopyin, acc_pcreate =20=20=20=20 libgomp/ * openacc.h (acc_pcopyin, acc_pcreate): Provide prototypes inst= ead 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. =20=20=20=20 trunk r248410 =20=20=20=20 git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@2484= 14 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog.gomp | 14 ++ libgomp/libgomp.h | 5 +- libgomp/libgomp.map | 8 +- libgomp/oacc-mem.c | 24 +- libgomp/openacc.h | 7 +- .../testsuite/libgomp.oacc-c-c++-common/lib-32.c | 241 +++++++++++++++++= ++-- .../testsuite/libgomp.oacc-c-c++-common/lib-38.c | 64 ------ 7 files changed, 271 insertions(+), 92 deletions(-) diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index 15e0c58..4ab1004 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,5 +1,19 @@ 2017-05-24 Thomas Schwinge =20 + * 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. =20 diff --git libgomp/libgomp.h libgomp/libgomp.h index a07a604..50916f2 100644 --- libgomp/libgomp.h +++ libgomp/libgomp.h @@ -1049,8 +1049,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; =20 -# 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"); @@ -1074,6 +1072,9 @@ extern int gomp_test_nest_lock_25 (omp_nest_lock_25_t= *) __GOMP_NOTHROW; #endif =20 #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 cee0bc9..2461c00 100644 --- libgomp/libgomp.map +++ libgomp/libgomp.map @@ -378,6 +378,12 @@ OACC_2.0 { acc_set_cuda_stream; }; =20 +OACC_2.0.1 { + global: + acc_pcopyin; + acc_pcreate; +} OACC_2.0; + OACC_2.5 { global: acc_copyin_async; @@ -430,7 +436,7 @@ OACC_2.5 { acc_update_self_async_32_h_; acc_update_self_async_64_h_; acc_update_self_async_array_h_; -} OACC_2.0; +} OACC_2.0.1; =20 GOACC_2.0 { global: diff --git libgomp/oacc-mem.c libgomp/oacc-mem.c index 6a1b25d..59a6f93 100644 --- libgomp/oacc-mem.c +++ libgomp/oacc-mem.c @@ -654,15 +654,23 @@ acc_create_async (void *h, size_t s, int async) present_create_copy (FLAG_PRESENT | FLAG_CREATE, h, s, async); } =20 +/* acc_present_or_create used to be what acc_create is now. */ +/* acc_pcreate is acc_present_or_create by a different name. */ #ifdef HAVE_ATTRIBUTE_ALIAS -extern void *acc_present_or_create (void *, size_t) - __attribute__((alias ("acc_create"))); +strong_alias (acc_create, acc_present_or_create) +strong_alias (acc_create, acc_pcreate) #else void * acc_present_or_create (void *h, size_t s) { return acc_create (h, s); } + +void * +acc_pcreate (void *h, size_t s) +{ + return acc_create (h, s); +} #endif =20 void * @@ -678,15 +686,23 @@ acc_copyin_async (void *h, size_t s, int async) present_create_copy (FLAG_PRESENT | FLAG_CREATE | FLAG_COPY, h, s, async= ); } =20 +/* acc_present_or_copyin used to be what acc_copyin is now. */ +/* acc_pcopyin is acc_present_or_copyin by a different name. */ #ifdef HAVE_ATTRIBUTE_ALIAS -extern void *acc_present_or_copyin (void *, size_t) - __attribute__((alias ("acc_copyin"))); +strong_alias (acc_copyin, acc_present_or_copyin) +strong_alias (acc_copyin, acc_pcopyin) #else void * acc_present_or_copyin (void *h, size_t s) { return acc_copyin (h, s); } + +void * +acc_pcopyin (void *h, size_t s) +{ + return acc_copyin (h, s); +} #endif =20 #define FLAG_COPYOUT (1 << 0) diff --git libgomp/openacc.h libgomp/openacc.h index eb7c888..67d8639 100644 --- libgomp/openacc.h +++ libgomp/openacc.h @@ -94,8 +94,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; @@ -124,11 +126,6 @@ void acc_copyout_finalize_async (void *, size_t, int) = __GOACC_NOTHROW; void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW; void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW; =20=20=20 -/* 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/tes= tsuite/libgomp.oacc-c-c++-common/lib-32.c index e3f87a8..666d4b8 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. */ =20 +#include #include #include =20 int main (int argc, char **argv) { - const int N =3D 256; - unsigned char *h; - void *d1, *d2; + int *h, *d; + const int N =3D 10000; + const int S =3D N * sizeof *h; + bool shared_mem; =20 - h =3D (unsigned char *) malloc (N); - - d1 =3D acc_present_or_create (h, N); - if (!d1) + h =3D (int *) malloc (S); + if (!h) abort (); + for (int i =3D 0; i < N; ++i) + h[i] =3D i + 0; =20 - d2 =3D acc_present_or_create (h, N); - if (!d2) - abort (); + shared_mem =3D acc_is_present (h, S); =20 - if (d1 !=3D d2) + d =3D (int *) acc_present_or_create (h, S); + if (!d) abort (); + if (shared_mem) + if (h !=3D d) + abort (); + if (!acc_is_present (h, S)) + abort (); + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + d[i] =3D i + 1; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 1 : 0)) + abort (); + h[i] =3D i + 2; + } + + { + int *d_ =3D (int *) acc_present_or_create (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 2 : 1)) + abort (); + d[i] =3D i + 3; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 3 : 2)) + abort (); + h[i] =3D i + 4; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 4 : 3)) + abort (); + d[i] =3D i + 5; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 5 : 4)) + abort (); + h[i] =3D i + 6; + } + + { + int *d_ =3D (int *) acc_present_or_copyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 6 : 5)) + abort (); + d[i] =3D i + 7; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 7 : 6)) + abort (); + h[i] =3D i + 8; + } + + { + int *d_ =3D (int *) acc_pcopyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 8 : 7)) + abort (); + d[i] =3D i + 9; + } =20 - d2 =3D acc_pcreate (h, N); - if (!d2) + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 9 : 8)) + abort (); + h[i] =3D i + 10; + } + + acc_copyout_finalize (h, S); + d =3D NULL; + if (!shared_mem) + if (acc_is_present (h, S)) + abort (); + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 10 : 9)) + abort (); + } + + d =3D (int *) acc_pcopyin (h, S); + if (!d) + abort (); + if (shared_mem) + if (h !=3D d) + abort (); + if (!acc_is_present (h, S)) abort (); =20 - if (d1 !=3D d2) +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 10 : 9)) + abort (); + d[i] =3D i + 11; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 11 : 9)) + abort (); + h[i] =3D i + 12; + } + + { + int *d_ =3D (int *) acc_pcopyin (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 12 : 11)) + abort (); + d[i] =3D i + 13; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 13 : 12)) + abort (); + h[i] =3D i + 14; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 14 : 13)) + abort (); + d[i] =3D i + 15; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 15 : 14)) + abort (); + h[i] =3D i + 16; + } + + { + int *d_ =3D (int *) acc_pcreate (h, S); + if (d_ !=3D d) + abort (); + } + +#pragma acc parallel loop deviceptr (d) + for (int i =3D 0; i < N; ++i) + { + if (d[i] !=3D i + (shared_mem ? 16 : 15)) + abort (); + d[i] =3D i + 17; + } + + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 17 : 16)) + abort (); + h[i] =3D i + 18; + } + + acc_update_self (h, S); + if (!acc_is_present (h, S)) abort (); =20 - acc_delete (h, N); + for (int i =3D 0; i < N; ++i) + { + if (h[i] !=3D i + (shared_mem ? 18 : 17)) + abort (); + } + + acc_delete_finalize (h, S); + d =3D NULL; + if (!shared_mem) + if (acc_is_present (h, S)) + abort(); =20 free (h); =20 diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/lib-38.c libgomp/tes= tsuite/libgomp.oacc-c-c++-common/lib-38.c deleted file mode 100644 index 1756867..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=3D0" } } */ - -#include -#include -#include - -int -main (int argc, char **argv) -{ - const int N =3D 256; - int i; - unsigned char *h; - void *d1, *d2; - - h =3D (unsigned char *) malloc (N); - - for (i =3D 0; i < N; i++) - { - h[i] =3D i; - } - - d1 =3D acc_present_or_copyin (h, N); - if (!d1) - abort (); - - for (i =3D 0; i < N; i++) - { - h[i] =3D 0xab; - } - - d2 =3D acc_present_or_copyin (h, N); - if (!d2) - abort (); - - if (d1 !=3D d2) - abort (); - - memset (&h[0], 0, N); - - acc_copyout_finalize (h, N); - - for (i =3D 0; i < N; i++) - { - if (h[i] !=3D i) - abort (); - } - - d2 =3D acc_pcopyin (h, N); - if (!d2) - abort (); - - acc_copyout (h, N); - - for (i =3D 0; i < N; i++) - { - if (h[i] !=3D i) - abort (); - } - - free (h); - - return 0; -} Gr=C3=BC=C3=9Fe Thomas