Hi!
On 2020-05-26T04:08:44-0300, Alexandre Oliva <ol...@adacore.com> wrote:
> Thanks, here's the combined patch I'm checking in.
>
> revamp dump and aux output names
For libgomp offloading testing, I'm seeing a number of failures like:
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
[-PASS:-]{+UNRESOLVED:+}
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
scan-offload-rtl-dump mach "Merging loop .* into "
[...]
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
[-PASS:-]{+FAIL:+} 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
[...]
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
[-PASS:-]{+UNRESOLVED:+}
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\\)"
[...]
UNSUPPORTED: libgomp.oacc-c/vec.c -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O0
PASS: libgomp.oacc-c/vec.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0
-foffload=nvptx-none -O2 (test for excess errors)
PASS: libgomp.oacc-c/vec.c -DACC_DEVICE_TYPE_nvidia=1 -DACC_MEM_SHARED=0
-foffload=nvptx-none -O2 execution test
[-PASS:-]{+UNRESOLVED:+} libgomp.oacc-c/vec.c -DACC_DEVICE_TYPE_nvidia=1
-DACC_MEM_SHARED=0 -foffload=nvptx-none -O2 scan-offload-tree-dump slp1
"vector\\(2\\) long long unsigned int"
[...]
(A number of similar ones skipped.)
A word of warning: some of that testing indeed is fragile/ugly ;-) -- but
we have to restore it anyway. That means, we shall gladly accept
suggestions how to do it better.
Consider 'libgomp.oacc-c-c++-common/nvptx-merged-loop.c':
/* { dg-options "-foffload=-fdump-rtl-mach" } */
/* { dg-final { scan-offload-rtl-dump "Merging loop .* into " "mach" } } */
Previously, for '-foffload=nvptx-none -foffload=-fdump-rtl-mach
-save-temps -o ./nvptx-merged-loop.exe', GCC produced the expected
'nvptx-merged-loop.o.307r.mach'.
Now, I find the file at '/tmp/cc6nlTN9.mkoffload-cc4ebUn4.o.307r.mach'
(The '-save-temps' comes from
'libgomp/testsuite/lib/libgomp-dg.exp:libgomp-dg-test', by the way.)
(Yes, that's all a bit ugly, how the 'mkoffload's etc. handle their
files...)
Similar for 'libgomp.oacc-c-c++-common/vector-length-128-1.c':
/* { dg-additional-options "-foffload=-fdump-tree-oaccdevlow" } */
/* { dg-final { scan-offload-tree-dump "__attribute__\\(\\(oacc function
\\(1, 1, 128\\)" "oaccdevlow" } } */
Similar for 'libgomp.oacc-c/vec.c':
/* { dg-additional-options "-std=c99 -ftree-slp-vectorize
-foffload=-fdump-tree-slp1" } */
/* { dg-final { scan-offload-tree-dump "vector\\(2\\) long long unsigned
int" "slp1" } } */
(Thus not analyed these any further.)
Consider 'libgomp.oacc-c-c++-common/pr85381-2.c':
/* { dg-additional-options "-save-temps" } */
/* { dg-final { scan-assembler-times "bar.sync" 2 } } */
This expects to scan the PTX offloading compilation assembler code (not
host code!), expecting that nvptx offloading code assembly is produced
after the host code, and thus overwrites the latter file. (Yes, that's
certainly ugly/fragile...)
Previously, for '-foffload=nvptx-none -save-temps -o ./pr85381-2.exe',
GCC left the '-foffload=nvptx-none' assembly in the expected
'pr85381-2.s'.
Now, I find the file at '/tmp/cc8JsG4H.mkoffload-pr85381-2.s'.
Are you able to easily create/suggest patches for these? (You're
probably not set up for offloading compilation...) Can you suggest
how/where to adjust: producer-side (GCC driver, 'mkoffload's?), or
consumer-side (testsuite: offload tree scanning machinery etc., or have
to put some '-dumpbase' or similar into all such test case files?)?
Grüße
Thomas
-----------------
Mentor Graphics (Deutschland) GmbH, Arnulfstraße 201, 80634 München / Germany
Registergericht München HRB 106955, Geschäftsführer: Thomas Heurung, Alexander
Walter
