This is the mail archive of the
gcc-patches@gcc.gnu.org
mailing list for the GCC project.
Re: [PATCH 7/n] OpenMP 4.0 offloading infrastructure: testsuite
- From: Jakub Jelinek <jakub at redhat dot com>
- To: Ilya Verbin <iverbin at gmail dot com>
- Cc: gcc-patches at gcc dot gnu dot org, Bernd Schmidt <bernds at codesourcery dot com>, Thomas Schwinge <thomas at codesourcery dot com>, Kirill Yukhin <kirill dot yukhin at gmail dot com>, Andrey Turetskiy <andrey dot turetskiy at gmail dot com>
- Date: Fri, 17 Oct 2014 16:14:55 +0200
- Subject: Re: [PATCH 7/n] OpenMP 4.0 offloading infrastructure: testsuite
- Authentication-results: sourceware.org; auth=none
- References: <20141015145752 dot GB46277 at msticlxl57 dot ims dot intel dot com> <20141015150543 dot GR10376 at tucnak dot redhat dot com> <20141015152837 dot GC46277 at msticlxl57 dot ims dot intel dot com> <20141015153518 dot GS10376 at tucnak dot redhat dot com> <20141017140211 dot GA28310 at msticlxl57 dot ims dot intel dot com>
- Reply-to: Jakub Jelinek <jakub at redhat dot com>
On Fri, Oct 17, 2014 at 06:02:11PM +0400, Ilya Verbin wrote:
> --- /dev/null
> +++ b/libgomp/testsuite/libgomp.c++/examples-4/e.53.2.C
> @@ -0,0 +1,43 @@
> +// { dg-do run }
> +// { dg-require-effective-target offload_device }
Well, this test actually relies not only offload_device,
but also on non-shared address space (so, if we ever have HSA
backend, it would fail there). So, perhaps not immediately,
but eventually we'll want an effective target whether
address space is shared or not between offloading device and
host.
> --- a/libgomp/testsuite/libgomp.c/target-7.c
> +++ b/libgomp/testsuite/libgomp.c/target-7.c
> @@ -1,7 +1,9 @@
> +// { dg-require-effective-target offload_device }
> +
Why? The test was specially written such that it tests
host fallback (if f is true) too.
> #include <omp.h>
> #include <stdlib.h>
>
> -volatile int v;
> +volatile int v = 0;
Why?
> void
> foo (int f)
> @@ -18,7 +20,7 @@ foo (int f)
> if (omp_get_level () != 0 || !omp_is_initial_device ())
> abort ();
> #pragma omp target if (v <= 1)
> - if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
> + if (omp_get_level () != 0 || omp_is_initial_device ())
> abort ();
> #pragma omp target device (d) if (v <= 1)
> if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
> @@ -30,7 +32,7 @@ foo (int f)
> if (omp_get_level () != 0 || !omp_is_initial_device ())
> abort ();
> #pragma omp target if (1)
> - if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
> + if (omp_get_level () != 0 || omp_is_initial_device ())
> abort ();
> #pragma omp target device (d) if (1)
> if (omp_get_level () != 0 || (f && !omp_is_initial_device ()))
> @@ -59,7 +61,7 @@ foo (int f)
> #pragma omp target data if (v <= 1) map (to: h)
> {
> #pragma omp target if (v <= 1)
> - if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 8)
> + if (omp_get_level () != 0 || omp_is_initial_device () || h++ != 8)
> abort ();
> #pragma omp target update if (v <= 1) from (h)
> }
> @@ -87,7 +89,7 @@ foo (int f)
> #pragma omp target data if (1) map (to: h)
> {
> #pragma omp target if (1)
> - if (omp_get_level () != 0 || (f && !omp_is_initial_device ()) || h++ != 12)
> + if (omp_get_level () != 0 || omp_is_initial_device () || h++ != 12)
> abort ();
> #pragma omp target update if (1) from (h)
> }
I don't understand any of these changes.
Otherwise it LGTM.
Jakub