Extend test cases for references in OpenACC 'private' clauses
libgomp/ * testsuite/libgomp.oacc-c++/privatized-ref-2.C: Extend. * testsuite/libgomp.oacc-c++/privatized-ref-3.C: Likewise. * testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise.
This commit is contained in:
Thomas Schwinge
parent
fbb438808e
commit
b75aab194e
3 changed files with 364 additions and 22 deletions
|
|
@ -1,20 +1,87 @@
|
|||
/* { 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" } */
|
||||
|
||||
/* 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 } */
|
||||
{
|
||||
int i, j;
|
||||
#pragma acc loop collapse(2) gang /* { dg-line l_loop[incr 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 } */
|
||||
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
|
||||
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_loop$c_loop }
|
||||
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_loop$c_loop }
|
||||
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_loop$c_loop }
|
||||
{ dg-bogus {note: variable 'tmpvar' adjusted for OpenACC privatization level: 'gang'} "" { target { openacc_nvidia_accel_selected && __OPTIMIZE__ } } l_loop$c_loop }
|
||||
*/
|
||||
/* { dg-note {variable 'tmpref' declared in block 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++)
|
||||
{
|
||||
int tmpvar;
|
||||
int &tmpref = tmpvar;
|
||||
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-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
|
||||
#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-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
||||
{
|
||||
int i, j;
|
||||
#pragma acc loop gang
|
||||
#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
|
||||
#pragma acc loop worker /* { dg-line l_loop[incr 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 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'worker'} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'worker'} "TODO" { target { ! openacc_host_selected } xfail *-*-* } l_loop$c_loop } */
|
||||
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
||||
for (j = 0; j < 256; j++)
|
||||
{
|
||||
int tmpvar;
|
||||
|
|
@ -35,14 +102,22 @@ void vectors (void)
|
|||
double res[65536];
|
||||
int i;
|
||||
|
||||
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
|
||||
/* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
|
||||
#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-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
||||
{
|
||||
int i, j;
|
||||
#pragma acc loop gang worker
|
||||
#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
|
||||
#pragma acc loop vector /* { dg-line l_loop[incr 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 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'vector'} "TODO" { target { ! openacc_host_selected } } l_loop$c_loop } */
|
||||
/* { dg-note {variable 'tmpref' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } l_loop$c_loop } */
|
||||
for (j = 0; j < 256; j++)
|
||||
{
|
||||
int tmpvar;
|
||||
|
|
@ -58,9 +133,46 @@ void vectors (void)
|
|||
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-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } l_compute$c_compute } */
|
||||
{
|
||||
int i, j;
|
||||
#pragma acc loop collapse(2) gang worker vector /* { dg-line l_loop[incr 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 } */
|
||||
/* { dg-note {variable 'tmpvar' declared in block is candidate for adjusting OpenACC privatization level} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' ought to be adjusted for OpenACC privatization level: 'vector'} "" { target *-*-* } l_loop$c_loop }
|
||||
{ dg-note {variable 'tmpvar' adjusted for OpenACC privatization level: 'vector'} "TODO" { target { ! openacc_host_selected } } l_loop$c_loop } */
|
||||
/* { dg-note {variable 'tmpref' declared in block 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++)
|
||||
{
|
||||
int tmpvar;
|
||||
int &tmpref = tmpvar;
|
||||
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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,25 +1,103 @@
|
|||
/* { 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-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
|
||||
#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
|
||||
#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)
|
||||
#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;
|
||||
|
|
@ -38,16 +116,34 @@ void vectors (void)
|
|||
double res[65536];
|
||||
int i;
|
||||
|
||||
#pragma acc parallel copyout(res) num_gangs(64) num_workers(64)
|
||||
/* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */
|
||||
#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
|
||||
#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)
|
||||
#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;
|
||||
|
|
@ -61,9 +157,56 @@ void vectors (void)
|
|||
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;
|
||||
}
|
||||
|
|
|
|||
|
|
@ -1,8 +1,24 @@
|
|||
! { 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_loop 0] }
|
||||
! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
|
||||
! "WARNING: dg-line var l_dummy defined, but not used". */
|
||||
|
||||
program main
|
||||
implicit none
|
||||
integer :: myint
|
||||
|
|
@ -11,12 +27,22 @@ program main
|
|||
|
||||
res(:) = 0.0
|
||||
|
||||
myint = 3
|
||||
call gangs(myint, res)
|
||||
|
||||
do i=1,65536
|
||||
tmp = i * 97
|
||||
if (res(i) .ne. tmp) stop 1
|
||||
end do
|
||||
|
||||
res(:) = 0.0
|
||||
|
||||
myint = 5
|
||||
call workers(myint, res)
|
||||
|
||||
do i=1,65536
|
||||
tmp = i * 99
|
||||
if (res(i) .ne. tmp) stop 1
|
||||
if (res(i) .ne. tmp) stop 2
|
||||
end do
|
||||
|
||||
res(:) = 0.0
|
||||
|
|
@ -26,11 +52,43 @@ program main
|
|||
|
||||
do i=1,65536
|
||||
tmp = i * 101
|
||||
if (res(i) .ne. tmp) stop 2
|
||||
if (res(i) .ne. tmp) stop 3
|
||||
end do
|
||||
|
||||
res(:) = 0.0
|
||||
|
||||
myint = 9
|
||||
call gangs_workers_vectors(myint, res)
|
||||
|
||||
do i=1,65536
|
||||
tmp = i * 103
|
||||
if (res(i) .ne. tmp) stop 4
|
||||
end do
|
||||
|
||||
contains
|
||||
|
||||
subroutine gangs(t1, res)
|
||||
implicit none
|
||||
integer :: t1
|
||||
integer :: i, j
|
||||
real, intent(out) :: res(:)
|
||||
|
||||
!$acc parallel copyout(res) num_gangs(64)
|
||||
|
||||
!$acc loop collapse(2) gang private(t1) ! { 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 }
|
||||
! { 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 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
|
||||
do i=0,255
|
||||
do j=1,256
|
||||
t1 = (i * 256 + j) * 97
|
||||
res(i * 256 + j) = t1
|
||||
end do
|
||||
end do
|
||||
|
||||
!$acc end parallel
|
||||
end subroutine gangs
|
||||
|
||||
subroutine workers(t1, res)
|
||||
implicit none
|
||||
integer :: t1
|
||||
|
|
@ -40,9 +98,12 @@ contains
|
|||
!$acc parallel copyout(res) num_gangs(64) num_workers(64)
|
||||
! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 }
|
||||
|
||||
!$acc loop gang
|
||||
!$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 }
|
||||
do i=0,255
|
||||
!$acc loop worker private(t1)
|
||||
!$acc loop worker private(t1) ! { dg-line l_loop[incr 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 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
|
||||
do j=1,256
|
||||
t1 = (i * 256 + j) * 99
|
||||
res(i * 256 + j) = t1
|
||||
|
|
@ -61,9 +122,12 @@ contains
|
|||
!$acc parallel copyout(res) num_gangs(64) num_workers(64)
|
||||
! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 }
|
||||
|
||||
!$acc loop gang worker
|
||||
!$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 }
|
||||
do i=0,255
|
||||
!$acc loop vector private(t1)
|
||||
!$acc loop vector private(t1) ! { dg-line l_loop[incr 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 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
|
||||
do j=1,256
|
||||
t1 = (i * 256 + j) * 101
|
||||
res(i * 256 + j) = t1
|
||||
|
|
@ -73,4 +137,27 @@ contains
|
|||
!$acc end parallel
|
||||
end subroutine vectors
|
||||
|
||||
subroutine gangs_workers_vectors(t1, res)
|
||||
implicit none
|
||||
integer :: t1
|
||||
integer :: i, j
|
||||
real, intent(out) :: res(:)
|
||||
|
||||
!$acc parallel copyout(res) num_gangs(64) num_workers(64)
|
||||
! { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 }
|
||||
|
||||
!$acc loop collapse(2) gang worker vector private(t1) ! { 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 }
|
||||
! { 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 't1' in 'private' clause potentially has improper OpenACC privatization level: 'parm_decl'} "" { target *-*-* } l_loop$c_loop }
|
||||
do i=0,255
|
||||
do j=1,256
|
||||
t1 = (i * 256 + j) * 103
|
||||
res(i * 256 + j) = t1
|
||||
end do
|
||||
end do
|
||||
|
||||
!$acc end parallel
|
||||
end subroutine gangs_workers_vectors
|
||||
|
||||
end program main
|
||||
|
|
|
|||
Loading…
Add table
Reference in a new issue