From patchwork Mon Jan 17 08:01:54 2022 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 1580618 Return-Path: X-Original-To: incoming@patchwork.ozlabs.org Delivered-To: patchwork-incoming@bilbo.ozlabs.org Authentication-Results: ozlabs.org; spf=pass (sender SPF authorized) smtp.mailfrom=gcc.gnu.org (client-ip=2620:52:3:1:0:246e:9693:128c; helo=sourceware.org; envelope-from=gcc-patches-bounces+incoming=patchwork.ozlabs.org@gcc.gnu.org; receiver=) Received: from sourceware.org (server2.sourceware.org [IPv6:2620:52:3:1:0:246e:9693:128c]) (using TLSv1.3 with cipher TLS_AES_256_GCM_SHA384 (256/256 bits) key-exchange X25519 server-signature RSA-PSS (4096 bits) server-digest SHA256) (No client certificate requested) by bilbo.ozlabs.org (Postfix) with ESMTPS id 4JcksR5zZHz9s9c for ; Mon, 17 Jan 2022 19:02:34 +1100 (AEDT) Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id CF7F83858023 for ; Mon, 17 Jan 2022 08:02:31 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa4.mentor.iphmx.com (esa4.mentor.iphmx.com [68.232.137.252]) by sourceware.org (Postfix) with ESMTPS id 8D4EF3858402 for ; Mon, 17 Jan 2022 08:02:06 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 8D4EF3858402 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: zQ0BDbyuwDs/BxCBjbrl/LbFnlkGgcdIwlBZRhO71YCma18FDWkrfM06bWJsJ5/y2apRavJf9C WhLn3iBPNFq44RuyU8DQqRfgSRBjOAwB8kfMTyVpycAwf+xRQ4jZjNEdB50BcvgYYf72V+e/s4 9SvkS9LBVh0z98X55ofFYbXNWf+bat3ihM2SQ8IOuO7wTA9ZUx1/68JCsrQ6TTZEhUCxhjYH/C KMJ+3MfqQrMR2lICRTeNmPkNuJw3B99xBhu/XMD89bdtDBU9yJ6HHVxjYmAg8HyoeRqgpHJHtu 7F6U9lWOJfwU5SpdM870Wr7F X-IronPort-AV: E=Sophos;i="5.88,295,1635235200"; d="scan'208,223";a="70822357" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa4.mentor.iphmx.com with ESMTP; 17 Jan 2022 00:02:04 -0800 IronPort-SDR: Hvo3+htd/TAl6kmvK0huMroz+SqgV+zGmi3ow1tTiIFm23uYlsR8aV8dOBeJVlyKpSBFkI9Hbv awo2jQ9/tFVlTqNGO3G9m9H4zbl1+0GjnI5XAJFjjupFaxaNWLGqUUlznl3Bj5HZ/f7Ul9gR70 J3p01WYir/QACUBFI4o95vTwiEvgfuGDsaGk3R86b2usJGdNJiTWb15xFBfQi5y1Yrn0rZHg+X VSmHXsg3YWJB/zuksevKmsNHmHyDfSNwxwtGd3nY9s4ordwYZyhdwQNGu1t1YOhfBSlaGp3ylh r7I= From: Thomas Schwinge To: Subject: Test cases for references in OpenACC 'private' clauses In-Reply-To: <20190920211734.28104-1-julian@codesourcery.com> References: <20190920211734.28104-1-julian@codesourcery.com> 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:01:54 +0100 Message-ID: <87lezedc3x.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) 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+incoming=patchwork.ozlabs.org@gcc.gnu.org Sender: "Gcc-patches" Hi! On 2019-09-20T14:17:33-0700, Julian Brown wrote: > This patch ['Handle references in OpenACC "private" clauses'] [...] ..., and its prerequisite changes, in particular... > [...] is intended as a somewhat temporary solution: it works for the > newly-included tests, but is not very elegant. ..., and breaks other things, as discussed internally a while ago. This will have to be done differently/analyzed in more detail. The test cases however, amend with: #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 } */ ... etc., and 'libgomp.oacc-c++/privatized-ref-3.C', 'libgomp.oacc-fortran/privatized-ref-1.f95' with: +/*TODO + { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */ ... I've now pushed to master branch in commit fbb438808e9b53a6e6b179a5787d609443acaad6 "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 fbb438808e9b53a6e6b179a5787d609443acaad6 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Fri, 20 Sep 2019 13:53:10 -0700 Subject: [PATCH] Test cases for references in OpenACC 'private' clauses libgomp/ * testsuite/libgomp.oacc-fortran/privatized-ref-1.f95: New test. * testsuite/libgomp.oacc-c++/privatized-ref-2.C: New test. * testsuite/libgomp.oacc-c++/privatized-ref-3.C: New test. Co-authored-by: Thomas Schwinge --- .../libgomp.oacc-c++/privatized-ref-2.C | 66 ++++++++++++++++ .../libgomp.oacc-c++/privatized-ref-3.C | 69 +++++++++++++++++ .../libgomp.oacc-fortran/privatized-ref-1.f95 | 76 +++++++++++++++++++ 3 files changed, 211 insertions(+) create mode 100644 libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C create mode 100644 libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 diff --git a/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C new file mode 100644 index 00000000000..7091091cac2 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-2.C @@ -0,0 +1,66 @@ +/* { dg-do run } */ + +#include + +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 } */ + { + int i, j; +#pragma acc loop gang + for (i = 0; i < 256; i++) + { +#pragma acc loop worker + for (j = 0; j < 256; j++) + { + int tmpvar; + int &tmpref = tmpvar; + tmpref = (i * 256 + j) * 99; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 99) + abort (); +} + +void vectors (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; +#pragma acc loop gang worker + for (i = 0; i < 256; i++) + { +#pragma acc loop vector + for (j = 0; j < 256; j++) + { + int tmpvar; + int &tmpref = tmpvar; + tmpref = (i * 256 + j) * 101; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 101) + abort (); +} + +int main (int argc, char *argv[]) +{ + 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 new file mode 100644 index 00000000000..478876e3596 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c++/privatized-ref-3.C @@ -0,0 +1,69 @@ +/* { dg-do run } */ + +/*TODO + { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } */ + +#include + +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 } */ + { + int i, j; + int tmpvar; + int &tmpref = tmpvar; +#pragma acc loop gang + for (i = 0; i < 256; i++) + { +#pragma acc loop worker private(tmpref) + for (j = 0; j < 256; j++) + { + tmpref = (i * 256 + j) * 99; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 99) + abort (); +} + +void vectors (void) +{ + double res[65536]; + int i; + +#pragma acc parallel copyout(res) num_gangs(64) num_workers(64) + /* { dg-warning "using num_workers \\(32\\), ignoring 64" "" { target openacc_nvidia_accel_selected } .-1 } */ + { + int i, j; + int tmpvar; + int &tmpref = tmpvar; +#pragma acc loop gang worker + for (i = 0; i < 256; i++) + { +#pragma acc loop vector private(tmpref) + for (j = 0; j < 256; j++) + { + tmpref = (i * 256 + j) * 101; + res[i * 256 + j] = tmpref; + } + } + } + + for (i = 0; i < 65536; i++) + if (res[i] != i * 101) + abort (); +} + +int main (int argc, char *argv[]) +{ + 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 new file mode 100644 index 00000000000..bb0910b1006 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/privatized-ref-1.f95 @@ -0,0 +1,76 @@ +! { dg-do run } + +!TODO +! { dg-xfail-run-if TODO { openacc_radeon_accel_selected && { ! __OPTIMIZE__ } } } + +program main + implicit none + integer :: myint + integer :: i + real :: res(65536), tmp + + res(:) = 0.0 + + myint = 5 + call workers(myint, res) + + do i=1,65536 + tmp = i * 99 + if (res(i) .ne. tmp) stop 1 + end do + + res(:) = 0.0 + + myint = 7 + call vectors(myint, res) + + do i=1,65536 + tmp = i * 101 + if (res(i) .ne. tmp) stop 2 + end do + +contains + + subroutine workers(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 gang + do i=0,255 + !$acc loop worker private(t1) + do j=1,256 + t1 = (i * 256 + j) * 99 + res(i * 256 + j) = t1 + end do + end do + + !$acc end parallel + end subroutine workers + + subroutine 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 gang worker + do i=0,255 + !$acc loop vector private(t1) + do j=1,256 + t1 = (i * 256 + j) * 101 + res(i * 256 + j) = t1 + end do + end do + + !$acc end parallel + end subroutine vectors + +end program main -- 2.34.1