From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: (qmail 74612 invoked by alias); 8 Apr 2016 08:33:56 -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 74588 invoked by uid 89); 8 Apr 2016 08:33:55 -0000 Authentication-Results: sourceware.org; auth=none X-Virus-Found: No X-Spam-SWARE-Status: No, score=-1.9 required=5.0 tests=AWL,BAYES_00,RCVD_IN_DNSWL_NONE,SPF_PASS,URIBL_RED autolearn=ham version=3.3.2 spammy=136,11, 13611, 8.0, 526,11 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 (AES256-GCM-SHA384 encrypted) ESMTPS; Fri, 08 Apr 2016 08:33:44 +0000 Received: from nat-ies.mentorg.com ([192.94.31.2] helo=SVR-IES-FEM-01.mgc.mentorg.com) by relay1.mentorg.com with esmtp id 1aoRrH-0005Bh-Va from Thomas_Schwinge@mentor.com for gcc-patches@gcc.gnu.org; Fri, 08 Apr 2016 01:33:40 -0700 Received: from hertz.schwinge.homeip.net (137.202.0.76) by SVR-IES-FEM-01.mgc.mentorg.com (137.202.0.104) with Microsoft SMTP Server id 14.3.224.2; Fri, 8 Apr 2016 09:33:38 +0100 From: Thomas Schwinge To: Subject: [gomp4] [PR testsuite/70579, PR testsuite/70580] Fix test cases failing because of the compiler's "avoid offloading" decision (was: Also test -O0 for OpenACC C, C++ offloading test cases) In-Reply-To: <87oa9l756u.fsf@kepler.schwinge.homeip.net> References: <87mvpqrets.fsf@kepler.schwinge.homeip.net> <56F1CC9B.1060003@redhat.com> <87mvppbsgw.fsf@hertz.schwinge.homeip.net> <56F2E72E.9090104@redhat.com> <87wpora7gu.fsf@hertz.schwinge.homeip.net> <8737r5aesq.fsf@hertz.schwinge.homeip.net> <87oa9l756u.fsf@kepler.schwinge.homeip.net> User-Agent: Notmuch/0.9-101-g81dad07 (http://notmuchmail.org) Emacs/24.4.1 (x86_64-pc-linux-gnu) Date: Fri, 08 Apr 2016 08:33:00 -0000 Message-ID: <871t6g7b5d.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable X-SW-Source: 2016-04/txt/msg00379.txt.bz2 Hi! On Thu, 7 Apr 2016 18:29:45 +0200, I wrote: > On Fri, 1 Apr 2016 10:55:49 +0200, I wrote: > > On Thu, 24 Mar 2016 22:31:29 +0100, I wrote: > > > On Wed, 23 Mar 2016 19:57:50 +0100, Bernd Schmidt wrote: > > > > Ok with [...]. > > >=20 > > > Thanks for the review; committed in r234471: > >=20 > > > Also test -O0 for OpenACC C, C++ offloading test cases > >=20 > > Merged into gomp-4_0-branch in r234664: >=20 > > --- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c > > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c > > @@ -1,5 +1,6 @@ > > /* { dg-do run { target openacc_nvidia_accel_selected } } */ > > /* { dg-additional-options "-lcuda -lcublas -lcudart" } */ > > +/* { dg-xfail-run-if "TODO" { *-*-* } { "-O0" } { "" } } */ > >=20=20 > > #include > > #include >=20 > Filed . >=20 > > --- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c > > +++ libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c > > @@ -1,3 +1,5 @@ > > +/* { dg-xfail-run-if "TODO" { *-*-* } { "-O0" } { "" } } */ > > + > > #include > > #include > > #include >=20 > Filed . >=20 > The both PASS on trunk. >=20 > Currently unclear if it's the same underlying problem or not. Committed to gomp-4_0-branch in r234823: commit 8a56a354bea62a280ec46db00a8956184403ff6e Author: tschwinge Date: Fri Apr 8 08:31:40 2016 +0000 [PR testsuite/70579, PR testsuite/70580] Fix test cases failing because= of the compiler's "avoid offloading" decision =20=20=20=20 libgomp/ PR testsuite/70579 PR testsuite/70580 * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Initialize the runtime for acc_device_nvidia. * testsuite/libgomp.oacc-c-c++-common/if-1.c: Cope with "avoid offloading" situation. * libgomp.texi (Enabling OpenACC): Elaborate on interactions of "avoid offloading". =20=20=20=20 git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@2348= 23 138bc75d-0d04-0410-961f-82ee72b054a4 --- libgomp/ChangeLog.gomp | 11 + libgomp/libgomp.texi | 9 + .../libgomp.oacc-c-c++-common/host_data-1.c | 4 +- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c | 316 +++++++++++------= ---- 4 files changed, 184 insertions(+), 156 deletions(-) diff --git libgomp/ChangeLog.gomp libgomp/ChangeLog.gomp index f54de33..1c99026 100644 --- libgomp/ChangeLog.gomp +++ libgomp/ChangeLog.gomp @@ -1,3 +1,14 @@ +2016-04-08 Thomas Schwinge + + PR testsuite/70579 + PR testsuite/70580 + * testsuite/libgomp.oacc-c-c++-common/host_data-1.c: Initialize + the runtime for acc_device_nvidia. + * testsuite/libgomp.oacc-c-c++-common/if-1.c: Cope with + "avoid offloading" situation. + * libgomp.texi (Enabling OpenACC): Elaborate on interactions of + "avoid offloading". + 2016-04-04 Thomas Schwinge =20 * testsuite/libgomp.oacc-c-c++-common/asyncwait-1.c: Don't XFAIL. diff --git libgomp/libgomp.texi libgomp/libgomp.texi index 350fb82..7cb677c 100644 --- libgomp/libgomp.texi +++ libgomp/libgomp.texi @@ -1824,6 +1824,15 @@ the option @option{-foffload-force} to force offload= ing in such cases. Alternatively, offloading is also enabled if a specific device type is requested, in a call to @code{acc_init} or by setting the @env{ACC_DEVICE_TYPE} environment variable, for example. +If you expect code to be offloaded no matter the overhead, for +example, when preparing data on a non-shared memory offloading device +for later use in an OpenACC @code{host_data} construct, be sure to +force offloading, to get the device memory initialized. +For example, if calling into the cuBLAS library (@ref{OpenACC Library +Interoperability}) from an OpenACC @code{host_data} construct, this +program is portable only to @code{acc_device_nvidia} devices, so it's +good practice then to initialize the runtime with an explicit +@code{acc_init (acc_device_nvidia)} call. =20 A complete description of all OpenACC directives accepted may be found in= =20 the @uref{http://www.openacc.org/, OpenACC} Application Programming diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c libgom= p/testsuite/libgomp.oacc-c-c++-common/host_data-1.c index b97d877..d19aa20 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c @@ -1,6 +1,5 @@ /* { dg-do run { target openacc_nvidia_accel_selected } } */ /* { dg-additional-options "-lcuda -lcublas -lcudart" } */ -/* { dg-xfail-run-if "TODO" { *-*-* } { "-O0" } { "" } } */ =20 #include #include @@ -30,6 +29,9 @@ saxpy_target (int n, float a, float *x, float *y) int main(int argc, char **argv) { + /* We're using cuBLAS, so excpect to be using a Nvidia GPU. */ + acc_init (acc_device_nvidia); + #define N 8 int i; float x_ref[N], y_ref[N]; diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c libgomp/tests= uite/libgomp.oacc-c-c++-common/if-1.c index ade064d..cc250f8 100644 --- libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c +++ libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c @@ -1,5 +1,3 @@ -/* { dg-xfail-run-if "TODO" { *-*-* } { "-O0" } { "" } } */ - #include #include #include @@ -23,6 +21,11 @@ main(int argc, char **argv) for (i =3D 0; i < N; i++) a[i] =3D 4.0; =20 + /* Are we actually offloading? */ + bool offloading; +#pragma acc parallel copyout(offloading) + offloading =3D acc_on_device (acc_device_not_host); + #pragma acc parallel copyin(a[0:N]) copyout(b[0:N]) if(1) { int ii; @@ -36,11 +39,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 5.0; -#else - exp =3D 4.0; -#endif + if (offloading) + exp =3D 4.0; + else + exp =3D 5.0; =20 for (i =3D 0; i < N; i++) { @@ -86,11 +88,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 9.0; -#else - exp =3D 8.0; -#endif + if (offloading) + exp =3D 8.0; + else + exp =3D 9.0; =20 for (i =3D 0; i < N; i++) { @@ -136,11 +137,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 17.0; -#else - exp =3D 16.0; -#endif + if (offloading) + exp =3D 16.0; + else + exp =3D 17.0; =20 for (i =3D 0; i < N; i++) { @@ -188,11 +188,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 23.0; -#else - exp =3D 22.0; -#endif + if (offloading) + exp =3D 22.0; + else + exp =3D 23.0; =20 for (i =3D 0; i < N; i++) { @@ -242,11 +241,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 50.0; -#else - exp =3D 49.0; -#endif + if (offloading) + exp =3D 49.0; + else + exp =3D 50.0; =20 for (i =3D 0; i < N; i++) { @@ -294,11 +292,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 92.0; -#else - exp =3D 91.0; -#endif + if (offloading) + exp =3D 91.0; + else + exp =3D 92.0; =20 for (i =3D 0; i < N; i++) { @@ -322,11 +319,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 44.0; -#else - exp =3D 43.0; -#endif + if (offloading) + exp =3D 43.0; + else + exp =3D 44.0; =20 for (i =3D 0; i < N; i++) { @@ -362,15 +358,18 @@ main(int argc, char **argv) b[i] =3D 9.0; } =20 -#if ACC_MEM_SHARED - exp =3D 0.0; - exp2 =3D 0.0; -#else - acc_map_data (a, d_a, N * sizeof (float)); - acc_map_data (b, d_b, N * sizeof (float)); - exp =3D 3.0; - exp2 =3D 9.0; -#endif + if (offloading) + { + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + exp =3D 3.0; + exp2 =3D 9.0; + } + else + { + exp =3D 0.0; + exp2 =3D 0.0; + } =20 #pragma acc update device(a[0:N], b[0:N]) if(1) =20 @@ -441,10 +440,11 @@ main(int argc, char **argv) abort(); } =20 -#if !ACC_MEM_SHARED - acc_unmap_data (a); - acc_unmap_data (b); -#endif + if (offloading) + { + acc_unmap_data (a); + acc_unmap_data (b); + } =20 acc_free (d_a); acc_free (d_b); @@ -481,17 +481,15 @@ main(int argc, char **argv) } =20 #pragma acc data copyin(a[0:N]) copyout(b[0:N]) if(0) -{ -#if !ACC_MEM_SHARED - if (acc_is_present (a, N * sizeof (float))) - abort (); -#endif - -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif -} + { + if (offloading) + { + if (acc_is_present (a, N * sizeof (float))) + abort (); + if (acc_is_present (b, N * sizeof (float))) + abort (); + } + } =20 for (i =3D 0; i < N; i++) { @@ -500,18 +498,20 @@ main(int argc, char **argv) } =20 #pragma acc data copyin(a[0:N]) if(1) -{ -#if !ACC_MEM_SHARED - if (!acc_is_present (a, N * sizeof (float))) - abort (); -#endif + { + if (offloading) + { + if (!acc_is_present (a, N * sizeof (float))) + abort (); + } =20 #pragma acc data copyout(b[0:N]) if(0) - { -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + { + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc data copyout(b[0:N]) if(1) { @@ -526,10 +526,11 @@ main(int argc, char **argv) } } =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } } } =20 @@ -541,72 +542,81 @@ main(int argc, char **argv) =20 #pragma acc enter data copyin (b[0:N]) if (0) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (0) =20 #pragma acc enter data copyin (b[0:N]) if (1) =20 -#if !ACC_MEM_SHARED - if (!acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (!acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (1) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc enter data copyin (b[0:N]) if (zero) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (zero) =20 #pragma acc enter data copyin (b[0:N]) if (one) =20 -#if !ACC_MEM_SHARED - if (!acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (!acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (one) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc enter data copyin (b[0:N]) if (one =3D=3D 0) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (one =3D=3D 0) =20 #pragma acc enter data copyin (b[0:N]) if (one =3D=3D 1) =20 -#if !ACC_MEM_SHARED - if (!acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (!acc_is_present (b, N * sizeof (float))) + abort (); + } =20 #pragma acc exit data delete (b[0:N]) if (one =3D=3D 1) =20 -#if !ACC_MEM_SHARED - if (acc_is_present (b, N * sizeof (float))) - abort (); -#endif + if (offloading) + { + if (acc_is_present (b, N * sizeof (float))) + abort (); + } =20 for (i =3D 0; i < N; i++) a[i] =3D 4.0; @@ -624,11 +634,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 5.0; -#else - exp =3D 4.0; -#endif + if (offloading) + exp =3D 4.0; + else + exp =3D 5.0; =20 for (i =3D 0; i < N; i++) { @@ -674,11 +683,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 9.0; -#else - exp =3D 8.0; -#endif + if (offloading) + exp =3D 8.0; + else + exp =3D 9.0; =20 for (i =3D 0; i < N; i++) { @@ -724,11 +732,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 17.0; -#else - exp =3D 16.0; -#endif + if (offloading) + exp =3D 16.0; + else + exp =3D 17.0; =20 for (i =3D 0; i < N; i++) { @@ -776,11 +783,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 23.0; -#else - exp =3D 22.0; -#endif + if (offloading) + exp =3D 22.0; + else + exp =3D 23.0; =20 for (i =3D 0; i < N; i++) { @@ -830,11 +836,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 50.0; -#else - exp =3D 49.0; -#endif + if (offloading) + exp =3D 49.0; + else + exp =3D 50.0; =20 for (i =3D 0; i < N; i++) { @@ -882,11 +887,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 92.0; -#else - exp =3D 91.0; -#endif + if (offloading) + exp =3D 91.0; + else + exp =3D 92.0; =20 for (i =3D 0; i < N; i++) { @@ -910,11 +914,10 @@ main(int argc, char **argv) } } =20 -#if ACC_MEM_SHARED - exp =3D 44.0; -#else - exp =3D 43.0; -#endif + if (offloading) + exp =3D 43.0; + else + exp =3D 44.0; =20 for (i =3D 0; i < N; i++) { @@ -950,15 +953,18 @@ main(int argc, char **argv) b[i] =3D 9.0; } =20 -#if ACC_MEM_SHARED - exp =3D 0.0; - exp2 =3D 0.0; -#else - acc_map_data (a, d_a, N * sizeof (float)); - acc_map_data (b, d_b, N * sizeof (float)); - exp =3D 3.0; - exp2 =3D 9.0; -#endif + if (offloading) + { + acc_map_data (a, d_a, N * sizeof (float)); + acc_map_data (b, d_b, N * sizeof (float)); + exp =3D 3.0; + exp2 =3D 9.0; + } + else + { + exp =3D 0.0; + exp2 =3D 0.0; + } =20 return 0; } Gr=C3=BC=C3=9Fe Thomas