7a5e036b61
... so that it matches what we analyze and what we action on.
Fix-up for commit 29a2f51806
"openacc:
Add support for gang local storage allocation in shared memory [PR90115]".
PR middle-end/90115
PR middle-end/102330
PR middle-end/104774
gcc/
* omp-low.cc (oacc_privatization_candidate_p)
(oacc_privatization_scan_clause_chain)
(oacc_privatization_scan_decl_chain, lower_oacc_private_marker):
Analyze 'lookup_decl'-translated DECL.
gcc/testsuite/
* c-c++-common/goacc/kernels-decompose-pr104061-1-3.c: Adjust.
* c-c++-common/goacc/kernels-decompose-pr104061-1-4.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104132-1.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104133-1.c: Likewise.
* c-c++-common/goacc/kernels-decompose-pr104774-1.c: Likewise.
* c-c++-common/goacc/privatization-1-compute-loop.c: Likewise.
* c-c++-common/goacc/privatization-1-compute.c: Likewise.
* c-c++-common/goacc/privatization-1-routine_gang-loop.c:
Likewise.
* c-c++-common/goacc/privatization-1-routine_gang.c: Likewise.
* gfortran.dg/goacc-gomp/pr102330-1.f90: Likewise, and subsume...
* gfortran.dg/goacc-gomp/pr102330-2.f90: ... this file, and...
* gfortran.dg/goacc-gomp/pr102330-3.f90: ... this file.
* gfortran.dg/goacc/privatization-1-compute-loop.f90: Adjust.
* gfortran.dg/goacc/privatization-1-compute.f90: Likewise.
* gfortran.dg/goacc/privatization-1-routine_gang-loop.f90:
Likewise.
* gfortran.dg/goacc/privatization-1-routine_gang.f90: Likewise.
libgomp/
* testsuite/libgomp.oacc-c-c++-common/kernels-decompose-1.c:
Enhance.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-1.c:
Adjust.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-local-worker-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-1.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-vector-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-2.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-3.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-4.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-5.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-6.c:
Likewise.
* testsuite/libgomp.oacc-c-c++-common/kernels-private-vars-loop-worker-7.c:
Likewise.
* testsuite/libgomp.oacc-fortran/optional-private.f90: Likewise.
* testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: Likewise.
* testsuite/libgomp.oacc-fortran/privatized-ref-2.f90: Likewise.
89 lines
4.1 KiB
C
89 lines
4.1 KiB
C
/* Test of worker-private variables declared on a loop directive, broadcasting
|
|
to vector-partitioned mode. Aggregate worker variable. */
|
|
|
|
/* { dg-additional-options "--param=openacc-kernels=decompose" } */
|
|
|
|
/* { dg-additional-options "-fopt-info-omp-all" }
|
|
{ dg-additional-options "-foffload=-fopt-info-omp-all" } */
|
|
|
|
/* { dg-additional-options "--param=openacc-privatization=noisy" }
|
|
{ dg-additional-options "-foffload=--param=openacc-privatization=noisy" }
|
|
Prune a few: uninteresting:
|
|
{ dg-prune-output {note: variable 'D\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} } */
|
|
|
|
/* 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_i 0 c_loop_j 0 c_loop_k 0] }
|
|
{ dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid
|
|
"WARNING: dg-line var l_dummy defined, but not used". */
|
|
|
|
#include <assert.h>
|
|
|
|
typedef struct
|
|
{
|
|
int x, y;
|
|
} vec2;
|
|
|
|
int
|
|
main (int argc, char* argv[])
|
|
{
|
|
int i, arr[32 * 32 * 32];
|
|
vec2 pt;
|
|
|
|
for (i = 0; i < 32 * 32 * 32; i++)
|
|
arr[i] = i;
|
|
|
|
#pragma acc kernels copy(arr) /* { dg-line l_compute[incr c_compute] } */
|
|
/* [PR104784] For some reason, for C++, the OpenACC 'kernels' decomposition
|
|
decides that a data region is needed for 'j', and subsequently requests it
|
|
to be made addressable.
|
|
{ dg-note {OpenACC 'kernels' decomposition: variable 'j' declared in block requested to be made addressable} {} { target c++ } l_compute$c_compute }
|
|
{ dg-note {variable 'j' made addressable} {} { target c++ } l_compute$c_compute }
|
|
{ dg-note {variable 'j' declared in block is candidate for adjusting OpenACC privatization level} {} { target c++ } l_compute$c_compute } */
|
|
{
|
|
int j;
|
|
|
|
/* { dg-note {forwarded loop nest in OpenACC 'kernels' region to 'parloops' for analysis} {} { target *-*-* } .+1 } */
|
|
#pragma acc loop gang(num:32) /* { dg-line l_loop_i[incr c_loop_i] } */
|
|
/* { dg-note {variable 'j' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target c } l_loop_i$c_loop_i } */
|
|
/* { dg-note {variable 'i' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_i$c_loop_i } */
|
|
for (i = 0; i < 32; i++)
|
|
{
|
|
#pragma acc loop worker(num:32) private(pt) /* { dg-line l_loop_j[incr c_loop_j] } */
|
|
/* { dg-note {variable 'pt' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
|
|
/* { dg-note {variable 'j' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
|
|
/* { dg-note {variable 'k' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_j$c_loop_j } */
|
|
for (j = 0; j < 32; j++)
|
|
{
|
|
int k;
|
|
|
|
pt.x = i ^ j * 3;
|
|
pt.y = i | j * 5;
|
|
|
|
#pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
|
|
/* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
|
|
for (k = 0; k < 32; k++)
|
|
arr[i * 1024 + j * 32 + k] += pt.x * k;
|
|
|
|
#pragma acc loop vector(length:32) /* { dg-line l_loop_k[incr c_loop_k] } */
|
|
/* { dg-note {variable 'k' in 'private' clause isn't candidate for adjusting OpenACC privatization level: not addressable} {} { target *-*-* } l_loop_k$c_loop_k } */
|
|
for (k = 0; k < 32; k++)
|
|
arr[i * 1024 + j * 32 + k] += pt.y * k;
|
|
}
|
|
}
|
|
/* { dg-optimized {assigned OpenACC seq loop parallelism} {} { target *-*-* } l_loop_i$c_loop_i } */
|
|
}
|
|
|
|
for (i = 0; i < 32; i++)
|
|
for (int j = 0; j < 32; j++)
|
|
for (int k = 0; k < 32; k++)
|
|
{
|
|
int idx = i * 1024 + j * 32 + k;
|
|
assert (arr[idx] == idx + (i ^ j * 3) * k + (i | j * 5) * k);
|
|
}
|
|
|
|
return 0;
|
|
}
|