From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: by sourceware.org (Postfix, from userid 1643) id 5BFFF3858006; Mon, 17 Jan 2022 08:00:00 +0000 (GMT) DKIM-Filter: OpenDKIM Filter v2.11.0 sourceware.org 5BFFF3858006 MIME-Version: 1.0 Content-Transfer-Encoding: 7bit Content-Type: text/plain; charset="utf-8" From: Thomas Schwinge To: gcc-cvs@gcc.gnu.org Subject: [gcc r12-6624] Extend test cases for references in OpenACC 'private' clauses X-Act-Checkin: gcc X-Git-Author: Thomas Schwinge X-Git-Refname: refs/heads/master X-Git-Oldrev: fbb438808e9b53a6e6b179a5787d609443acaad6 X-Git-Newrev: b75aab194e3fe40b594d9a70eb7068dc9950bcf0 Message-Id: <20220117080000.5BFFF3858006@sourceware.org> Date: Mon, 17 Jan 2022 08:00:00 +0000 (GMT) X-BeenThere: gcc-cvs@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-cvs mailing list List-Unsubscribe: , List-Archive: List-Help: List-Subscribe: , X-List-Received-Date: Mon, 17 Jan 2022 08:00:00 -0000 https://gcc.gnu.org/g:b75aab194e3fe40b594d9a70eb7068dc9950bcf0 commit r12-6624-gb75aab194e3fe40b594d9a70eb7068dc9950bcf0 Author: Thomas Schwinge Date: Tue Aug 24 18:33:04 2021 +0200 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. Diff: --- .../testsuite/libgomp.oacc-c++/privatized-ref-2.C | 128 +++++++++++++++-- .../testsuite/libgomp.oacc-c++/privatized-ref-3.C | 159 +++++++++++++++++++-- .../libgomp.oacc-fortran/privatized-ref-1.f95 | 99 ++++++++++++- 3 files changed, 364 insertions(+), 22 deletions(-) diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C index 7091091cac2..520016ab59d 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C @@ -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 +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; } diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C index 478876e3596..cb7085a01e8 100644 --- a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C @@ -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 +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; } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 index bb0910b1006..a8230561fc9 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 @@ -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