From patchwork Mon Jan 17 08:04:26 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 50095 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id DF3563858436 for ; Mon, 17 Jan 2022 08:04:56 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 8F2433858402 for ; Mon, 17 Jan 2022 08:04:36 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 8F2433858402 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com IronPort-SDR: 06pvXxfwXTp0reOeQ80RhdpVGDxWxJx6iCpkLrslP10uo0cr6o1UBL2M6M7MXl5cWok7xo0RTK FQMH3/cwHdrfJsoDzSynv4g3dtRdR0Z0FKVB4tOq0IVFlEbx/gKvBAueev4phZJcja2noBDGxp Z9jMxVpfJINtrgcqEJOXTgcY+qpo3OeMtId1a2BhmDsgwjccuutGliJY5tvler109ACgO6idSN oFeCaNIV/JDEjQrQn1IT9xFChYvSuYy7zbIV5grJCofzJKZCH4tiujLV+MhM3XxXT17FeqgcIQ Ups= X-IronPort-AV: E=Sophos;i="5.88,295,1635235200"; d="scan'208,223";a="70698159" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa3.mentor.iphmx.com with ESMTP; 17 Jan 2022 00:04:34 -0800 IronPort-SDR: +lZvbGwgB8J0mZ1XLRC7DlRvRJKO94i4rrtQ+EWdq6/arLq0kFMEvuYdeG+GRweGwAMVdF1AFI keWCdf44zWJh4ATIqfavqrfSvZnfVVyYk9G2/fZWjSPCoz7HUTlrJVKMSCPYbRzDosf5JS4EdI aSuAJEyQzDX/x7SKeRsVDZzaJWZslI7aeCUAVLcgZUs4sm27oC5SMYbbT6sor6V31NJTWsmW6/ U0rJfMUW+ab3eOCFKUIQeaqeWhjZ/cJxD31HOiteEpqAr9HyPIVdz5p0oLSPXk9D/AfuSgn5B2 XQI= From: Thomas Schwinge To: Subject: Extend test cases for references in OpenACC 'private' clauses (was: Test cases for references in OpenACC 'private' clauses) In-Reply-To: <87lezedc3x.fsf@euler.schwinge.homeip.net> References: <20190920211734.28104-1-julian@codesourcery.com> <87lezedc3x.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/27.1 (x86_64-pc-linux-gnu) Date: Mon, 17 Jan 2022 09:04:26 +0100 Message-ID: <87iluidbzp.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-14.mgc.mentorg.com (139.181.222.14) To svr-ies-mbx-01.mgc.mentorg.com (139.181.222.1) X-Spam-Status: No, score=-12.1 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.4 X-Spam-Checker-Version: SpamAssassin 3.4.4 (2020-01-24) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Cc: Julian Brown Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2022-01-17T09:01:54+0100, I wrote: > I've now pushed to master branch in > commit fbb438808e9b53a6e6b179a5787d609443acaad6 > "Test cases for references in OpenACC 'private' clauses", see attached. ... extended by commit b75aab194e3fe40b594d9a70eb7068dc9950bcf0 "Extend test cases for references in OpenACC 'private' clauses", see attached. Grüße Thomas ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 From b75aab194e3fe40b594d9a70eb7068dc9950bcf0 Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Tue, 24 Aug 2021 18:33:04 +0200 Subject: [PATCH] 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. --- .../libgomp.oacc-c++/privatized-ref-2.C | 128 +++++++++++++- .../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 -- 2.34.1