2aea19bdb1
gcc/ChangeLog: * config/nvptx/nvptx.cc (nvptx_goacc_validate_dims_1): Update warning messages. libgomp/ChangeLog: * testsuite/libgomp.oacc-c++/privatized-ref-2.C: Update scanning patterns. * testsuite/libgomp.oacc-c++/privatized-ref-3.C: Likewise. * testsuite/libgomp.oacc-c-c++-common/acc_prof-kernels-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/kernels-loop-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/parallel-dims.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr85486.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/pr95270-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/routine-nohost-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-1.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/struct-copyout-2.c: Likewise. * testsuite/libgomp.oacc-c-c++-common/vector-length-64-1.c: Likewise. * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/derivedtypes-arrays-1.f90: Likewise. * testsuite/libgomp.oacc-fortran/kernels-loop-2.f95: Likewise. * testsuite/libgomp.oacc-fortran/parallel-dims.f90: Likewise. * testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise. Co-authored-by: Thomas Schwinge <thomas@codesourcery.com>
212 lines
13 KiB
C
212 lines
13 KiB
C
/* { dg-do run } */
|
|
|
|
/* { dg-additional-options "-fopt-info-note-omp" }
|
|
{ dg-additional-options "-foffload=-fopt-info-note-omp" } */
|
|
|
|
/* { dg-additional-options "--param=openacc-privatization=noisy" }
|
|
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" } */
|
|
|
|
/* { dg-additional-options "-Wuninitialized" } */
|
|
|
|
/*TODO
|
|
{ dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */
|
|
|
|
/* It's only with Tcl 8.5 (released in 2007) that "the variable 'varName'
|
|
passed to 'incr' may be unset, and in that case, it will be set to [...]",
|
|
so to maintain compatibility with earlier Tcl releases, we manually
|
|
initialize counter variables:
|
|
{ dg-line l_dummy[variable c_compute 0 c_loop 0] }
|
|
{ dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
|
|
"WARNING: dg-line var l_dummy defined, but not used". */
|
|
|
|
#include <stdlib.h>
|
|
|
|
void gangs (void)
|
|
{
|
|
double res[65536];
|
|
int i;
|
|
|
|
#pragma acc parallel copyout(res) num_gangs(64) /* { dg-line l_compute[incr c_compute] } */
|
|
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
|
|
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
|
|
No longer having address taken: tmpvar
|
|
Now a gimple register: tmpvar
|
|
However, 'tmpvar' remains in the candidate set:
|
|
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
|
|
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
|
|
For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
|
|
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
|
|
*/
|
|
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
{
|
|
int i, j;
|
|
int tmpvar;
|
|
int &tmpref = tmpvar;
|
|
#pragma acc loop collapse(2) gang private(tmpref) /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (i = 0; i < 256; i++)
|
|
{
|
|
for (j = 0; j < 256; j++)
|
|
{
|
|
tmpref = (i * 256 + j) * 97;
|
|
res[i * 256 + j] = tmpref;
|
|
}
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < 65536; i++)
|
|
if (res[i] != i * 97)
|
|
abort ();
|
|
}
|
|
|
|
void workers (void)
|
|
{
|
|
double res[65536];
|
|
int i;
|
|
|
|
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
|
|
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
|
|
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
|
|
No longer having address taken: tmpvar
|
|
Now a gimple register: tmpvar
|
|
However, 'tmpvar' remains in the candidate set:
|
|
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
|
|
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
|
|
For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
|
|
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
|
|
*/
|
|
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
|
{
|
|
int i, j;
|
|
int tmpvar;
|
|
int &tmpref = tmpvar;
|
|
#pragma acc loop gang /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (i = 0; i < 256; i++)
|
|
{
|
|
#pragma acc loop worker private(tmpref) /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (j = 0; j < 256; j++)
|
|
{
|
|
tmpref = (i * 256 + j) * 99;
|
|
res[i * 256 + j] = tmpref;
|
|
}
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < 65536; i++)
|
|
if (res[i] != i * 99)
|
|
abort ();
|
|
}
|
|
|
|
void vectors (void)
|
|
{
|
|
double res[65536];
|
|
int i;
|
|
|
|
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
|
|
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
|
|
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
|
|
No longer having address taken: tmpvar
|
|
Now a gimple register: tmpvar
|
|
However, 'tmpvar' remains in the candidate set:
|
|
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
|
|
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
|
|
For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
|
|
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
|
|
*/
|
|
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
|
{
|
|
int i, j;
|
|
int tmpvar;
|
|
int &tmpref = tmpvar;
|
|
#pragma acc loop gang worker /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (i = 0; i < 256; i++)
|
|
{
|
|
#pragma acc loop vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (j = 0; j < 256; j++)
|
|
{
|
|
tmpref = (i * 256 + j) * 101;
|
|
res[i * 256 + j] = tmpref;
|
|
}
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < 65536; i++)
|
|
if (res[i] != i * 101)
|
|
abort ();
|
|
}
|
|
|
|
void gangs_workers_vectors (void)
|
|
{
|
|
double res[65536];
|
|
int i;
|
|
|
|
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) /* { dg-line l_compute[incr c_compute] } */
|
|
/* { dg-note {variable 'i' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_compute$c_compute }
|
|
But, with optimizations enabled, per the '*.ssa' dump ('gcc/tree-ssa.c:execute_update_addresses_taken'):
|
|
No longer having address taken: tmpvar
|
|
Now a gimple register: tmpvar
|
|
However, 'tmpvar' remains in the candidate set:
|
|
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'gang'} "" { target *-*-* } l_compute$c_compute }
|
|
Now, for GCN offloading, 'adjust_private_decl' does the privatization change right away:
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target openacc_radeon_accel_selected } l_compute$c_compute }
|
|
For nvptx offloading however, we first mark up 'tmpvar', and then later apply the privatization change -- or, with optimizations enabled, don't, because we then don't actually call 'expand_var_decl'.
|
|
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && { ! __OPTIMIZE__ } } } l_compute$c_compute }
|
|
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_compute$c_compute }
|
|
*/
|
|
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_compute$c_compute } */
|
|
/* { dg-warning "using .num_workers \\(32\\)., ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
|
{
|
|
int i, j;
|
|
int tmpvar;
|
|
int &tmpref = tmpvar;
|
|
#pragma acc loop collapse(2) gang worker vector private(tmpref) /* { dg-line l_loop[incr c_loop] } */
|
|
/* { dg-note {variable 'tmpref' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
|
for (i = 0; i < 256; i++)
|
|
{
|
|
for (j = 0; j < 256; j++)
|
|
{
|
|
tmpref = (i * 256 + j) * 103;
|
|
res[i * 256 + j] = tmpref;
|
|
}
|
|
}
|
|
}
|
|
|
|
for (i = 0; i < 65536; i++)
|
|
if (res[i] != i * 103)
|
|
abort ();
|
|
}
|
|
|
|
int main (int argc, char *argv[])
|
|
{
|
|
gangs ();
|
|
workers ();
|
|
vectors ();
|
|
gangs_workers_vectors ();
|
|
return 0;
|
|
}
|