[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)

Message ID	871t6g7b5d.fsf@hertz.schwinge.homeip.net
State	New
Headers	show Return-Path: <gcc-patches-return-424538-incoming=patchwork.ozlabs.org@gcc.gnu.org> DomainKey-Signature: a=rsa-sha1; c=nofws; d=gcc.gnu.org; h=list-id :list-unsubscribe:list-archive:list-post:list-help:sender:from :to:subject:in-reply-to:references:date:message-id:mime-version :content-type:content-transfer-encoding; q=dns; s=default; b=ihn 3h8vm+RJEJzfl+CVp99Ldjdi7tJhbw0O42iGqObFbxeTuyYVH31VkympNwKpG4hc jLqCmrGX0eSYaI/4PZuLVg1KZU04vP1m6XcJkeaAyC6q1OvRZQ2n1JSa9p0FucOO rU8jXGkxB/MLTdXcQItmXHOonOXsfG81XjFvUoMQ= Mailing-List: contact gcc-patches-help@gcc.gnu.org; run by ezmlm Precedence: bulk Sender: gcc-patches-owner@gcc.gnu.org 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) 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, 8 Apr 2016 10:33:18 +0200 Message-ID: <871t6g7b5d.fsf@hertz.schwinge.homeip.net> MIME-Version: 1.0 Content-Type: text/plain; charset="utf-8" Content-Transfer-Encoding: quoted-printable

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; }

[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)

Commit Message

Patch