diff mbox series

[committed,nvptx,libgomp] Move rtl-dump test-cases to libgomp

Message ID 20181219142839.GA5247@delia
State New
Headers show
Series [committed,nvptx,libgomp] Move rtl-dump test-cases to libgomp | expand

Commit Message

Tom de Vries Dec. 19, 2018, 2:28 p.m. UTC 
Hi,

The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing
during linking due to missing libgomp.spec.

Move them to the libgomp testsuite.

Build and reg-tested on x86_64 with nvptx accelerator.

Committed to trunk.

Thanks,
- Tom

[nvptx, libgomp] Move rtl-dump test-cases to libgomp

2018-12-19  Tom de Vries  <tdevries@suse.de>

	* gcc.dg/goacc/nvptx-merged-loop.c: Move to
	libgomp/testsuite/libgomp.oacc-c-c++-common.
	* gcc.dg/goacc/nvptx-sese-1.c: Same.

	* testsuite/lib/libgomp.exp: Add load_lib of scanoffloadrtl.exp.
	* testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c: Move from
	gcc/testsuite/gcc.dg/goacc.
	* testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Same.

---
 libgomp/testsuite/lib/libgomp.exp                                 | 1 +
 .../testsuite/libgomp.oacc-c-c++-common}/nvptx-merged-loop.c      | 8 ++++----
 .../testsuite/libgomp.oacc-c-c++-common}/nvptx-sese-1.c           | 8 ++++----
 3 files changed, 9 insertions(+), 8 deletions(-)

Comments

Thomas Schwinge Feb. 22, 2019, 11:09 a.m. UTC | #1 
Hi!

On Wed, 19 Dec 2018 15:28:41 +0100, Tom de Vries <tdevries@suse.de> wrote:
> The goacc.exp test-cases nvptx-merged-loop.c and nvptx-sese-1.c are failing
> during linking due to missing libgomp.spec.

Right; see also
<http://mid.mail-archive.com/87bnalrwdl.fsf@kepler.schwinge.homeip.net>,
from three years ago.  ;-)

> Move them to the libgomp testsuite.
> 
> Build and reg-tested on x86_64 with nvptx accelerator.
> 
> Committed to trunk.

ACK, thanks.

> diff --git a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
> similarity index 59%
> rename from gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c
> rename to libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
> index 3ff537c1d97..8a2117e1624 100644
> --- a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c
> +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
> @@ -1,6 +1,6 @@
> -/* { dg-do link } */
> -/* { dg-require-effective-target offload_nvptx } */
> -/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */
> +/* { dg-do run { target openacc_nvidia_accel_selected } } */
> +/* { dg-options "-foffload=-fdump-rtl-mach" } */
> +/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
> [...]
> -/* { dg-final { scan-rtl-dump "Merging loop .* into " "mach" } } */
> +/* { dg-final { scan-offload-rtl-dump "Merging loop .* into " "mach" } } */

There is however still a problem here: if you have more than one offload
compiler configured, all of them will react to
'-foffload=-fdump-rtl-mach', and will overwrite their dump files, or if
different compiler passes configured, that may produce more than one
"mach" dump file, all this resulting in "funny" testing FAILs/ERRORs.
But: we can't just use '-foffload=[nvptx]=-fdump-rtl-mach', because
'[nvptx]' can be 'nvptx', or 'nvptx-none', etc.

Similar problem for a lot of other test cases that you've recently added
(... but thanks for all these, of course!).

And, "dg-require-effective-target offload_nvptx" is now obsolete, and
should be removed.

All this I'm addressing in the commits I've just done, see attached:

  - r269106 "[libgomp] In OpenACC offloading testing, be more explicit in
    what is supported, and what is not, or why not"
  - r269107 "[libgomp] Clarify difference between offload target, offload
    plugin, and OpenACC device type"
  - r269108 "[libgomp] In OpenACC testing, cycle though all offload
    targets"
  - r269109 "[libgomp] In OpenACC testing, by default only build for the
    offload target that we're actually going to test"
  - r269110 "Remove unused check_effective_target_offload_nvptx"

These libgomp changes improves testing time, as discussed before, and
clean up the test log as follows:

    UNSUPPORTED: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-merged-loop.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0 
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-merged-loop.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-merged-loop.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    [-ERROR:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-merged-loop.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   [-: error executing dg-final: couldn't open "nvptx-merged-loop.o.304r.mach nvptx-merged-loop.o.307r.mach": no such file or directory-]{+scan-offload-rtl-dump mach "Merging loop .* into "+}

    UNSUPPORTED: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0 
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    [-ERROR:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/nvptx-sese-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   [-: error executing dg-final: couldn't open "nvptx-sese-1.o.307r.mach nvptx-sese-1.o.304r.mach": no such file or directory-]{+scan-offload-rtl-dump mach "SESE regions:.* [0-9]+{[0-9]+->[0-9]+(\\.[0-9]+)+}"+}

    UNSUPPORTED: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0 
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-assembler-times bar.sync 2

    UNSUPPORTED: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0 
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/pr85381-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-assembler-times bar.sync 4

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-1.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 128\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-2.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 128\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 32\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-3.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 1, 32\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-4.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2, 128\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-5.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 2, 128\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-6.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0, 128\\)"

    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O0   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0, 128\\)"
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  (test for excess errors)
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  execution test
    PASS: libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2  output pattern test
    [-FAIL:-]{+PASS:+} libgomp.oacc-c/../libgomp.oacc-c-c++-common/vector-length-128-7.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0 {+-foffload=nvptx-none+}  -O2   scan-offload-tree-dump oaccdevlow "__attribute__\\(\\(oacc function \\(1, 0, 128\\)"

Same for 'libgomp.oacc-c++'.


Grüße
 Thomas
From e03c3144d02850737551a071b6525f13efcde7c0 Mon Sep 17 00:00:00 2001
From: tschwinge <tschwinge@138bc75d-0d04-0410-961f-82ee72b054a4>
Date: Fri, 22 Feb 2019 10:50:49 +0000
Subject: [PATCH 5/9] [libgomp] In OpenACC offloading testing, be more explicit
 in what is supported, and what is not, or why not

	libgomp/
	* testsuite/lib/libgomp.exp: Error out for unknown offload target.
	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.  Report if
	"offloading: supported, but hardware not accessible".
	* testsuite/libgomp.oacc-c/c.exp: Likewise.
	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.

git-svn-id: svn+ssh://gcc.gnu.org/svn/gcc/trunk@269106 138bc75d-0d04-0410-961f-82ee72b054a4
---
 libgomp/ChangeLog                                  | 8 ++++++++
 libgomp/testsuite/lib/libgomp.exp                  | 7 +++++--
 libgomp/testsuite/libgomp.oacc-c++/c++.exp         | 4 ++--
 libgomp/testsuite/libgomp.oacc-c/c.exp             | 4 ++--
 libgomp/testsuite/libgomp.oacc-fortran/fortran.exp | 4 ++--
 5 files changed, 19 insertions(+), 8 deletions(-)

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index f310c2f6c24..32bc21f50fd 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,11 @@
+2019-02-22  Thomas Schwinge  <thomas@codesourcery.com>
+
+	* testsuite/lib/libgomp.exp: Error out for unknown offload target.
+	* testsuite/libgomp.oacc-c++/c++.exp: Likewise.  Report if
+	"offloading: supported, but hardware not accessible".
+	* testsuite/libgomp.oacc-c/c.exp: Likewise.
+	* testsuite/libgomp.oacc-fortran/fortran.exp: Likewise.
+
 2019-02-19  Chung-Lin Tang <cltang@codesourcery.com>
 
 	PR c/87924
diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index 04738a9ce82..2392642c7d6 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -58,6 +58,9 @@ foreach offload_target_openacc $offload_targets_s {
 	hsa {
 	    continue
 	}
+	default {
+	    error "Unknown offload target: $offload_target_openacc"
+	}
     }
     lappend offload_targets_s_openacc "$offload_target_openacc"
 }
@@ -375,7 +378,7 @@ proc check_effective_target_openacc_nvidia_accel_configured { } {
     return [is-effective-target lp64]
 }
 
-# Return 1 if at least one nvidia board is present.
+# Return 1 if at least one Nvidia GPU is accessible.
 
 proc check_effective_target_openacc_nvidia_accel_present { } {
     return [check_runtime openacc_nvidia_accel_present {
@@ -386,7 +389,7 @@ proc check_effective_target_openacc_nvidia_accel_present { } {
     } "" ]
 }
 
-# Return 1 if at least one nvidia board is present, and the nvidia device type
+# Return 1 if at least one Nvidia GPU is accessible, and the nvidia device type
 # is selected by default by means of setting the environment variable
 # ACC_DEVICE_TYPE.
 
diff --git a/libgomp/testsuite/libgomp.oacc-c++/c++.exp b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
index 9beadd6fc1e..d0e5f42b88b 100644
--- a/libgomp/testsuite/libgomp.oacc-c++/c++.exp
+++ b/libgomp/testsuite/libgomp.oacc-c++/c++.exp
@@ -88,7 +88,7 @@ if { $lang_test_file_found } {
 	    nvidia {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
-		    untested "$subdir $offload_target_openacc offloading"
+		    untested "$subdir $offload_target_openacc offloading: supported, but hardware not accessible"
 		    continue
 		}
 
@@ -101,7 +101,7 @@ if { $lang_test_file_found } {
 		set acc_mem_shared 0
 	    }
 	    default {
-		set acc_mem_shared 0
+		error "Unknown OpenACC device type: $offload_target_openacc"
 	    }
 	}
 	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
diff --git a/libgomp/testsuite/libgomp.oacc-c/c.exp b/libgomp/testsuite/libgomp.oacc-c/c.exp
index 4475bf5341f..77e852efaef 100644
--- a/libgomp/testsuite/libgomp.oacc-c/c.exp
+++ b/libgomp/testsuite/libgomp.oacc-c/c.exp
@@ -51,7 +51,7 @@ foreach offload_target_openacc $offload_targets_s_openacc {
 	nvidia {
 	    if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		# Don't bother; execution testing is going to FAIL.
-		untested "$subdir $offload_target_openacc offloading"
+		untested "$subdir $offload_target_openacc offloading: supported, but hardware not accessible"
 		continue
 	    }
 
@@ -64,7 +64,7 @@ foreach offload_target_openacc $offload_targets_s_openacc {
 	    set acc_mem_shared 0
 	}
 	default {
-	    set acc_mem_shared 0
+	    error "Unknown OpenACC device type: $offload_target_openacc"
 	}
     }
     set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
index 663c9323b72..a262a8a2c97 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
+++ b/libgomp/testsuite/libgomp.oacc-fortran/fortran.exp
@@ -76,14 +76,14 @@ if { $lang_test_file_found } {
 	    nvidia {
 		if { ![check_effective_target_openacc_nvidia_accel_present] } {
 		    # Don't bother; execution testing is going to FAIL.
-		    untested "$subdir $offload_target_openacc offloading"
+		    untested "$subdir $offload_target_openacc offloading: supported, but hardware not accessible"
 		    continue
 		}
 
 		set acc_mem_shared 0
 	    }
 	    default {
-		set acc_mem_shared 0
+		error "Unknown OpenACC device type: $offload_target_openacc"
 	    }
 	}
 	set tagopt "$tagopt -DACC_MEM_SHARED=$acc_mem_shared"
Tom de Vries Feb. 22, 2019, 12:45 p.m. UTC | #2 
On 22-02-19 12:09, Thomas Schwinge wrote:
> There is however still a problem here: if you have more than one offload
> compiler configured, 

Ah, right, that's a configuration I'm not testing, evidently.

> all of them will react to
> '-foffload=-fdump-rtl-mach', and will overwrite their dump files, or if
> different compiler passes configured, that may produce more than one
> "mach" dump file, all this resulting in "funny" testing FAILs/ERRORs.
> But: we can't just use '-foffload=[nvptx]=-fdump-rtl-mach', because
> '[nvptx]' can be 'nvptx', or 'nvptx-none', etc.
> 
> Similar problem for a lot of other test cases that you've recently added
> (... but thanks for all these, of course!).
> 
> And, "dg-require-effective-target offload_nvptx" is now obsolete, and
> should be removed.
> 
> All this I'm addressing in the commits I've just done, see attached:

Thanks for fixing that.

- Tom
diff mbox series

Patch

diff --git a/libgomp/testsuite/lib/libgomp.exp b/libgomp/testsuite/lib/libgomp.exp
index c41b3e6dc18..04738a9ce82 100644
--- a/libgomp/testsuite/lib/libgomp.exp
+++ b/libgomp/testsuite/lib/libgomp.exp
@@ -31,6 +31,7 @@  load_gcc_lib scanrtl.exp
 load_gcc_lib scantree.exp
 load_gcc_lib scanltranstree.exp
 load_gcc_lib scanoffloadtree.exp
+load_gcc_lib scanoffloadrtl.exp
 load_gcc_lib scanipa.exp
 load_gcc_lib scanwpaipa.exp
 load_gcc_lib timeout-dg.exp
diff --git a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
similarity index 59%
rename from gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
index 3ff537c1d97..8a2117e1624 100644
--- a/gcc/testsuite/gcc.dg/goacc/nvptx-merged-loop.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-merged-loop.c
@@ -1,6 +1,6 @@ 
-/* { dg-do link } */
-/* { dg-require-effective-target offload_nvptx } */
-/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-merged-loop.c\\ -Wa,--no-verify" } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
 
 #define N (32*32*32+17)
 void __attribute__ ((noinline)) Foo (int *ary)
@@ -27,4 +27,4 @@  int main ()
   return 0;
 }   
 
-/* { dg-final { scan-rtl-dump "Merging loop .* into " "mach" } } */
+/* { dg-final { scan-offload-rtl-dump "Merging loop .* into " "mach" } } */
diff --git a/gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c
similarity index 63%
rename from gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c
rename to libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c
index 7e67fe78f06..9583265c775 100644
--- a/gcc/testsuite/gcc.dg/goacc/nvptx-sese-1.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c
@@ -1,6 +1,6 @@ 
-/* { dg-do link } */
-/* { dg-require-effective-target offload_nvptx } */
-/* { dg-options "-fopenacc -O2 -foffload=-fdump-rtl-mach\\ -dumpbase\\ nvptx-sese-1.c\\ -Wa,--no-verify" } */
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+/* { dg-options "-foffload=-fdump-rtl-mach" } */
+/* { dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
 
 #pragma acc routine  seq
 int __attribute__((noinline)) foo (int x)
@@ -32,4 +32,4 @@  int main ()
 }
 
 /* Match {N->N(.N)+} */
-/* { dg-final { scan-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */
+/* { dg-final { scan-offload-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */