8a8efad098
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.
37 lines
1.0 KiB
C
37 lines
1.0 KiB
C
/* { dg-do run { target openacc_nvidia_accel_selected } }
|
|
{ dg-skip-if "" { *-*-* } { "*" } { "-O2" } } */
|
|
/* { dg-additional-options "-foffload=-fdump-rtl-mach" } */
|
|
|
|
int
|
|
main (void)
|
|
{
|
|
int v1;
|
|
|
|
#pragma acc parallel
|
|
#pragma acc loop worker
|
|
for (v1 = 0; v1 < 20; v1 += 2)
|
|
;
|
|
|
|
return 0;
|
|
}
|
|
|
|
/* Todo: Boths bar.syncs can be removed.
|
|
Atm we generate this dead code inbetween forked and joining:
|
|
|
|
mov.u32 %r28, %ntid.y;
|
|
mov.u32 %r29, %tid.y;
|
|
add.u32 %r30, %r29, %r29;
|
|
setp.gt.s32 %r31, %r30, 19;
|
|
@%r31 bra $L2;
|
|
add.u32 %r25, %r28, %r28;
|
|
mov.u32 %r24, %r30;
|
|
$L3:
|
|
add.u32 %r24, %r24, %r25;
|
|
setp.le.s32 %r33, %r24, 19;
|
|
@%r33 bra $L3;
|
|
$L2:
|
|
|
|
so the loop is not recognized as empty loop (which we detect by seeing if
|
|
joining immediately follows forked). */
|
|
/* { dg-final { scan-offload-rtl-dump-times "nvptx_barsync" 2 "mach" } } */
|