From: Thomas Schwinge <thomas@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
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)
Date: Fri, 08 Apr 2016 08:33:00 -0000 [thread overview]
Message-ID: <871t6g7b5d.fsf@hertz.schwinge.homeip.net> (raw)
In-Reply-To: <87oa9l756u.fsf@kepler.schwinge.homeip.net>
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 <bschmidt@redhat.com> wrote:
> > > > Ok with [...].
> > >
> > > Thanks for the review; committed in r234471:
> >
> > > Also test -O0 for OpenACC C, C++ offloading test cases
> >
> > Merged into gomp-4_0-branch in r234664:
>
> > --- 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" } { "" } } */
> >
> > #include <stdlib.h>
> > #include <openacc.h>
>
> Filed <https://gcc.gnu.org/PR70579>.
>
> > --- 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 <openacc.h>
> > #include <stdlib.h>
> > #include <stdbool.h>
>
> Filed <https://gcc.gnu.org/PR70580>.
>
> The both PASS on trunk.
>
> Currently unclear if it's the same underlying problem or not.
Committed to gomp-4_0-branch in r234823:
commit 8a56a354bea62a280ec46db00a8956184403ff6e
Author: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
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
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".
git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/branches/gomp-4_0-branch@234823 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 <thomas@codesourcery.com>
+
+ 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 <thomas@codesourcery.com>
* 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 offloading 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.
A complete description of all OpenACC directives accepted may be found in
the @uref{http://www.openacc.org/, OpenACC} Application Programming
diff --git libgomp/testsuite/libgomp.oacc-c-c++-common/host_data-1.c libgomp/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" } { "" } } */
#include <stdlib.h>
#include <openacc.h>
@@ -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/testsuite/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 <openacc.h>
#include <stdlib.h>
#include <stdbool.h>
@@ -23,6 +21,11 @@ main(int argc, char **argv)
for (i = 0; i < N; i++)
a[i] = 4.0;
+ /* Are we actually offloading? */
+ bool offloading;
+#pragma acc parallel copyout(offloading)
+ offloading = 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)
}
}
-#if ACC_MEM_SHARED
- exp = 5.0;
-#else
- exp = 4.0;
-#endif
+ if (offloading)
+ exp = 4.0;
+ else
+ exp = 5.0;
for (i = 0; i < N; i++)
{
@@ -86,11 +88,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 9.0;
-#else
- exp = 8.0;
-#endif
+ if (offloading)
+ exp = 8.0;
+ else
+ exp = 9.0;
for (i = 0; i < N; i++)
{
@@ -136,11 +137,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 17.0;
-#else
- exp = 16.0;
-#endif
+ if (offloading)
+ exp = 16.0;
+ else
+ exp = 17.0;
for (i = 0; i < N; i++)
{
@@ -188,11 +188,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 23.0;
-#else
- exp = 22.0;
-#endif
+ if (offloading)
+ exp = 22.0;
+ else
+ exp = 23.0;
for (i = 0; i < N; i++)
{
@@ -242,11 +241,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 50.0;
-#else
- exp = 49.0;
-#endif
+ if (offloading)
+ exp = 49.0;
+ else
+ exp = 50.0;
for (i = 0; i < N; i++)
{
@@ -294,11 +292,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 92.0;
-#else
- exp = 91.0;
-#endif
+ if (offloading)
+ exp = 91.0;
+ else
+ exp = 92.0;
for (i = 0; i < N; i++)
{
@@ -322,11 +319,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 44.0;
-#else
- exp = 43.0;
-#endif
+ if (offloading)
+ exp = 43.0;
+ else
+ exp = 44.0;
for (i = 0; i < N; i++)
{
@@ -362,15 +358,18 @@ main(int argc, char **argv)
b[i] = 9.0;
}
-#if ACC_MEM_SHARED
- exp = 0.0;
- exp2 = 0.0;
-#else
- acc_map_data (a, d_a, N * sizeof (float));
- acc_map_data (b, d_b, N * sizeof (float));
- exp = 3.0;
- exp2 = 9.0;
-#endif
+ if (offloading)
+ {
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ exp = 3.0;
+ exp2 = 9.0;
+ }
+ else
+ {
+ exp = 0.0;
+ exp2 = 0.0;
+ }
#pragma acc update device(a[0:N], b[0:N]) if(1)
@@ -441,10 +440,11 @@ main(int argc, char **argv)
abort();
}
-#if !ACC_MEM_SHARED
- acc_unmap_data (a);
- acc_unmap_data (b);
-#endif
+ if (offloading)
+ {
+ acc_unmap_data (a);
+ acc_unmap_data (b);
+ }
acc_free (d_a);
acc_free (d_b);
@@ -481,17 +481,15 @@ main(int argc, char **argv)
}
#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 ();
+ }
+ }
for (i = 0; i < N; i++)
{
@@ -500,18 +498,20 @@ main(int argc, char **argv)
}
#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 ();
+ }
#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 ();
+ }
#pragma acc data copyout(b[0:N]) if(1)
{
@@ -526,10 +526,11 @@ main(int argc, char **argv)
}
}
-#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 ();
+ }
}
}
@@ -541,72 +542,81 @@ main(int argc, char **argv)
#pragma acc enter data copyin (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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (0)
#pragma acc enter data copyin (b[0:N]) if (1)
-#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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (1)
-#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 ();
+ }
#pragma acc enter data copyin (b[0:N]) if (zero)
-#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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (zero)
#pragma acc enter data copyin (b[0:N]) if (one)
-#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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (one)
-#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 ();
+ }
#pragma acc enter data copyin (b[0:N]) if (one == 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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (one == 0)
#pragma acc enter data copyin (b[0:N]) if (one == 1)
-#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 ();
+ }
#pragma acc exit data delete (b[0:N]) if (one == 1)
-#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 ();
+ }
for (i = 0; i < N; i++)
a[i] = 4.0;
@@ -624,11 +634,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 5.0;
-#else
- exp = 4.0;
-#endif
+ if (offloading)
+ exp = 4.0;
+ else
+ exp = 5.0;
for (i = 0; i < N; i++)
{
@@ -674,11 +683,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 9.0;
-#else
- exp = 8.0;
-#endif
+ if (offloading)
+ exp = 8.0;
+ else
+ exp = 9.0;
for (i = 0; i < N; i++)
{
@@ -724,11 +732,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 17.0;
-#else
- exp = 16.0;
-#endif
+ if (offloading)
+ exp = 16.0;
+ else
+ exp = 17.0;
for (i = 0; i < N; i++)
{
@@ -776,11 +783,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 23.0;
-#else
- exp = 22.0;
-#endif
+ if (offloading)
+ exp = 22.0;
+ else
+ exp = 23.0;
for (i = 0; i < N; i++)
{
@@ -830,11 +836,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 50.0;
-#else
- exp = 49.0;
-#endif
+ if (offloading)
+ exp = 49.0;
+ else
+ exp = 50.0;
for (i = 0; i < N; i++)
{
@@ -882,11 +887,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 92.0;
-#else
- exp = 91.0;
-#endif
+ if (offloading)
+ exp = 91.0;
+ else
+ exp = 92.0;
for (i = 0; i < N; i++)
{
@@ -910,11 +914,10 @@ main(int argc, char **argv)
}
}
-#if ACC_MEM_SHARED
- exp = 44.0;
-#else
- exp = 43.0;
-#endif
+ if (offloading)
+ exp = 43.0;
+ else
+ exp = 44.0;
for (i = 0; i < N; i++)
{
@@ -950,15 +953,18 @@ main(int argc, char **argv)
b[i] = 9.0;
}
-#if ACC_MEM_SHARED
- exp = 0.0;
- exp2 = 0.0;
-#else
- acc_map_data (a, d_a, N * sizeof (float));
- acc_map_data (b, d_b, N * sizeof (float));
- exp = 3.0;
- exp2 = 9.0;
-#endif
+ if (offloading)
+ {
+ acc_map_data (a, d_a, N * sizeof (float));
+ acc_map_data (b, d_b, N * sizeof (float));
+ exp = 3.0;
+ exp2 = 9.0;
+ }
+ else
+ {
+ exp = 0.0;
+ exp2 = 0.0;
+ }
return 0;
}
Grüße
Thomas
next prev parent reply other threads:[~2016-04-08 8:33 UTC|newest]
Thread overview: 22+ messages / expand[flat|nested] mbox.gz Atom feed top
2016-03-22 12:14 Also test -O0 for OpenACC C, C++ offloading test cases Thomas Schwinge
2016-03-22 13:04 ` Thomas Schwinge
2016-03-23 0:03 ` Bernd Schmidt
2016-03-23 7:35 ` Thomas Schwinge
2016-03-23 19:19 ` Bernd Schmidt
2016-03-24 23:33 ` Thomas Schwinge
2016-03-29 13:35 ` Thomas Schwinge
2016-04-01 8:56 ` [gomp4] " Thomas Schwinge
2016-04-07 16:30 ` Thomas Schwinge
2016-04-08 8:33 ` Thomas Schwinge [this message]
2016-03-23 19:47 ` Jakub Jelinek
2016-03-25 1:05 ` check-target-libgomp wall time, without vs. with offloading (was: Also test -O0 for OpenACC C, C++ offloading test cases) Thomas Schwinge
2016-04-28 10:44 ` check-target-libgomp wall time, without vs. with offloading Thomas Schwinge
2016-04-29 7:43 ` Jakub Jelinek
2016-05-02 9:56 ` libgomp: In OpenACC testing, cycle though $offload_targets, and by default only build for the offload target that we're actually going to test (was: check-target-libgomp wall time, without vs. with offloading) Thomas Schwinge
2016-05-11 13:45 ` libgomp: In OpenACC testing, cycle though $offload_targets, and by default only build for the offload target that we're actually going to test Thomas Schwinge
2016-05-18 11:42 ` Thomas Schwinge
2016-05-25 8:50 ` Thomas Schwinge
2016-06-29 9:19 ` [PING, two months] " Thomas Schwinge
2016-06-08 13:28 libgomp: Unconfuse offload plugins vs. offload targets Thomas Schwinge
2016-06-08 14:08 ` Jakub Jelinek
2016-06-08 15:46 ` Thomas Schwinge
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=871t6g7b5d.fsf@hertz.schwinge.homeip.net \
--to=thomas@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).