These test cases use directives similar to:
/* { 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...)
..., and this broke with recent commit
1dedc12d186a110854537e1279b4e6c29f2df35a
"revamp dump and aux output names" plus fix-up commit commit
efc16503ca10bc0e934e0bace5777500e4dc757a "handle dumpbase in offloading, adjust
testsuite" (short summary: file names changed), so let's finally make that
robust.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/pr85381-2.c: Replace fragile
'scan-assembler' with 'scan-offload-rtl'.
* testsuite/libgomp.oacc-c-c++-common/pr85381-3.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381-5.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/pr85381.c: Likewise.
-/* { dg-additional-options "-save-temps" } */
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
int
main (void)
so the loop is not recognized as empty loop (which we detect by seeing if
joining immediately follows forked). */
-/* { dg-final { scan-assembler-times "bar.sync" 2 } } */
+/* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 2 "mach" } } */
-/* { dg-additional-options "-save-temps -w" } */
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
int a;
#pragma acc declare create(a)
return 0;
}
-/* { dg-final { scan-assembler-not "bar.sync" } } */
+/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */
-/* { dg-additional-options "-save-temps -w" } */
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
#define n 1024
/* Atm, %ntid.y is broadcast from one loop to the next, so there are 2 bar.syncs
for that (the other two are there for the same reason as in pr85381-2.c).
Todo: Recompute %ntid.y instead of broadcasting it. */
-/* { dg-final { scan-assembler-times "bar.sync" 4 } } */
+/* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 4 "mach" } } */
-/* { dg-additional-options "-save-temps" } */
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
#define n 1024
return 0;
}
-/* { dg-final { scan-assembler-not "bar.sync" } } */
+/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */
-/* { dg-additional-options "-save-temps" } */
/* { dg-do run { target openacc_nvidia_accel_selected } }
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
+/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
int
main (void)
return 0;
}
-/* { dg-final { scan-assembler-not "bar.sync" } } */
+/* { dg-final { scan-offload-rtl-dump-not "nvptx_barsync" "mach" } } */