From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 49966 invoked by alias); 22 May 2017 14:26:58 -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 49946 invoked by uid 89); 22 May 2017 14:26:57 -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=d2, sk:present 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; Mon, 22 May 2017 14:26:55 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-MBX-03.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1dCoIR-00063I-5s from Thomas_Schwinge@mentor.com ; Mon, 22 May 2017 07:26:55 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.87) by SVR-IES-MBX-03.mgc.mentorg.com (139.181.222.3) with Microsoft SMTP Server (TLS) id 15.0.1210.3; Mon, 22 May 2017 15:26:52 +0100 From: Thomas Schwinge To: , Jakub Jelinek Subject: C/C++ OpenACC: acc_pcopyin, acc_pcreate User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.5.1 (x86_64-pc-linux-gnu) Date: Mon, 22 May 2017 14:28:00 -0000 Message-ID: <87bmqlynkn.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-03.mgc.mentorg.com (139.181.222.3) X-SW-Source: 2017-05/txt/msg01677.txt.bz2 Hi! In , 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 Date: Wed Apr 19 15:37:11 2017 +0200 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 (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; =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..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); } =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; -} Gr=C3=BC=C3=9Fe Thomas