a9de836c2b
That is, adjust for optimization introduced with recent
commit r13-3217-gc4d15dddf6b9eacb36f535807ad2ee364af46e04
"[PR107195] Set range to zero when nonzero mask is 0", where GCC now
understands that after 'r *= 2;', 'r & 1' will never hold here, and thus
transforms/optimizes/"disturbs" the original code such that GCC/nvptx's later
"Neuter whole SESE regions" optimization no longer is applicable to it:
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
[-PASS:-]{+FAIL:+} 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 scan-nvptx-none-offload-rtl-dump mach "SESE regions:.* [0-9]+{[0-9]+->[0-9]+(\\.[0-9]+)+}"
Same for C++.
It's unclear to me if this is an actual "problem", which optimization is "more
important", so I've filed PR107344 "GCC/nvptx SESE region optimization" to
capture this question, and here restore what we intend to be testing (to my
understanding) in 'libgomp.oacc-c-c++-common/nvptx-sese-1.c'.
PR tree-optimization/107195
PR target/107344
libgomp/
* testsuite/libgomp.oacc-c-c++-common/nvptx-sese-1.c: Restore SESE
regions checking.
35 lines
775 B
C
35 lines
775 B
C
/* { 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)
|
|
{
|
|
return x & 2;
|
|
}
|
|
|
|
int main ()
|
|
{
|
|
int r = 0;
|
|
|
|
#pragma acc parallel copy(r) vector_length(32)
|
|
{
|
|
#pragma acc loop vector reduction (+:r)
|
|
for (int i = 00; i < 40; i++)
|
|
r += i;
|
|
|
|
/* This piece is a multi-block SESE region */
|
|
if (foo (r))
|
|
r *= 2;
|
|
|
|
if (r & 8) /* to here. */
|
|
#pragma acc loop vector reduction (+:r)
|
|
for (int i = 00; i < 40; i++)
|
|
r += i;
|
|
}
|
|
|
|
return 0;
|
|
}
|
|
|
|
/* Match {N->N(.N)+} */
|
|
/* { dg-final { scan-offload-rtl-dump "SESE regions:.* \[0-9\]+{\[0-9\]+->\[0-9\]+(\\.\[0-9\]+)+}" "mach" } } */
|