From patchwork Wed Oct 25 09:29:52 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Thomas Schwinge X-Patchwork-Id: 157956 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a59:ce89:0:b0:403:3b70:6f57 with SMTP id p9csp2473551vqx; Wed, 25 Oct 2023 02:30:38 -0700 (PDT) X-Google-Smtp-Source: AGHT+IH7ztWvvWr7LRQPCDCNHYOiVjmpFNGkS1A5LittlcoMGBZRbGRoZrvXIile0m8zLA30a+GX X-Received: by 2002:a05:6214:27cb:b0:658:3a12:9949 with SMTP id ge11-20020a05621427cb00b006583a129949mr18304134qvb.53.1698226237957; Wed, 25 Oct 2023 02:30:37 -0700 (PDT) ARC-Seal: i=2; a=rsa-sha256; t=1698226237; cv=pass; d=google.com; s=arc-20160816; b=wv1QJBY9e0sjxYxrdMUKeqzFLQS6VgCtOCT3nB3soz251cloRiR0v+3+xQuOfBuZw7 ej3EcGTWEwOLffmxrA6uY4bnai160N98csSx6N4KdXz2No+TfH2L8zLbGIMWVicY1Ttk rWlDLz14cnv5ch34xfm74ovMIBktLb8sOdE/L2Aq3D6obsFBHGQ1RWR4fV56MjhdseUC T9EyeEMKsyTXrkFxTbHIyX/k2ok1H99OcH43AmfKPbnMKrflitSN5dcJVnzQ6AxpoWiN 35vT47dWSXWJ0PInZiMEbT9Z33UkWsqTc/Ln+nSw2pBRqI487XXkR6zgsjAzX0j8dLtT dzMA== ARC-Message-Signature: i=2; a=rsa-sha256; c=relaxed/relaxed; d=google.com; s=arc-20160816; h=errors-to:list-subscribe:list-help:list-post:list-archive :list-unsubscribe:list-id:precedence:mime-version:message-id:date :user-agent:references:in-reply-to:subject:cc:to:from:ironport-sdr :arc-filter:dmarc-filter:delivered-to; bh=Tlnrw1Dqo1aHNlQe/UwmQG5TvlWdOQ/hlhECOFNrR/w=; fh=fUaghwsITalRfY5ECywWs9xQrYlBehvJbVsj5wI2vdM=; b=aWeDIsIrKzZ4PJp4doCyVbNGvpGf4YOEODBDMtqfZIWyq7gJkwDszzr0XPUpXm550h wXAGxtlc1m/TicW7dTiDbwiqMVjDZveTNZ7ghTS5DQoyT0uyBaaacdWeh7E3uBE4HcUf oPJwIWySvpCyE3W5s32rMvxTy6qNBJ1YrG1Y2Mqt8nX6SPzMJyTC8p7qpNHVmCa4aO2h Oz1cN0hnY3+VMMRSgHDLPP5eAkcBrBkm2L6rlw2J9YyBXjoI3oTtxZRgRSkGWZx9Ld4T dRMhD7o/V837zuRX3B0K2NMKrE2ZdOrpmJUbzb93hdnU29OiCwsaOfEp5sR6lCwb6+0z jzWw== ARC-Authentication-Results: i=2; mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (server2.sourceware.org. [8.43.85.97]) by mx.google.com with ESMTPS id oo28-20020a056214451c00b0065611b26a34si8339725qvb.126.2023.10.25.02.30.37 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 25 Oct 2023 02:30:37 -0700 (PDT) Received-SPF: pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) client-ip=8.43.85.97; Authentication-Results: mx.google.com; arc=pass (i=1); spf=pass (google.com: domain of gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org designates 8.43.85.97 as permitted sender) smtp.mailfrom="gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org" Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id B14493857BA4 for ; Wed, 25 Oct 2023 09:30:37 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 6C0103858281; Wed, 25 Oct 2023 09:30:04 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6C0103858281 Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com ARC-Filter: OpenARC Filter v1.0.0 sourceware.org 6C0103858281 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.141.98 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698226210; cv=none; b=suG7gzboO1EYzoNwALwe8jwfkBm2M+Er75H1rsfeuMblakyTBRx5tmN85JlIzswbRfV3h/RazI19Ay0gtLHPFllImY21dlhRitnvspSZJ/cN2wPUZ8y4M3s5/5ghVNW0EPeFYNNTMI9YtlfAY9EwBn2XvVxuQ0Nl8eEua89Xrvk= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1698226210; c=relaxed/simple; bh=jXThVVH3UEVIt44B73pzODJRCSoGXxDJ5uKKmJ/jlBg=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=irOfvcoHkmfmj5kN0wXIEjMVUNQhpSM9zpQBJE20KcWXEpgGG2BgM9TUeVwOrw5VY/qDli/d+rc/qye1Es2i6TGjDtOiU1LzJDnw1v6nutsMvP7fAcsnEUHxvVmmOSCOVj0802pPmVd2m1lD4YtJnxi42jJ3kYaLGqTo2FSBHkg= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: Glf/cX+FSjqBUx1wZ6CfkQ== X-CSE-MsgGUID: BgwEOiVjTlGlI3vvICKC5A== X-IronPort-AV: E=Sophos;i="6.03,250,1694764800"; d="scan'208,223";a="23200093" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 25 Oct 2023 01:30:03 -0800 IronPort-SDR: q8GaMWBk8gpKXCteScL+8FTV9RFRoC8l4nFXHro7ATknhq0wCW15VsYFW56/V1NvmFPjT1WDXz oR0j9nXGnX9UMshttjopL7lDpPb0aNz2gcbgYViphGDGB5y009Vt7CVncjJdfqO7ttuyZNQAZi x7fBuzUMi0GElf9Ge0C6r3lNICkDY7bRYPm372HBY3/GNAxPxtkAOmA7M6b5mrdi8DwizkX8UF 3dQ2cBynFxvpOjviSSyT4o5/Q64xzFvbVirgzFOWBjLgHcdsHvdvhoRfbQLslIbW3farQ3N0wu M+w= From: Thomas Schwinge To: Chung-Lin Tang , , CC: Catherine Moore , Tobias Burnus Subject: Extend test suite coverage for OpenACC 'self' clause for compute constructs (was: [PATCH, OpenACC 2.7] Implement self clause for compute constructs) In-Reply-To: <87pm13w04d.fsf@euler.schwinge.homeip.net> References: <87pm13w04d.fsf@euler.schwinge.homeip.net> User-Agent: Notmuch/0.29.3+94~g74c3f1b (https://notmuchmail.org) Emacs/28.2 (x86_64-pc-linux-gnu) Date: Wed, 25 Oct 2023 11:29:52 +0200 Message-ID: <87edhjvylr.fsf@euler.schwinge.homeip.net> MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) To svr-ies-mbx-10.mgc.mentorg.com (139.181.222.10) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_LOTSOFHASH, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.30 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+ouuuleilei=gmail.com@gcc.gnu.org X-getmail-retrieved-from-mailbox: INBOX X-GMAIL-THRID: 1780719275286227933 X-GMAIL-MSGID: 1780719275286227933 Hi! On 2023-10-25T10:57:06+0200, I wrote: > With minor textual conflicts resolved, I've pushed this to master branch > in commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a > "OpenACC 2.7: Implement self clause for compute constructs", see > attached. > > > I'll then apply/submit a number of follow-on commits. > From 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a Mon Sep 17 00:00:00 2001 > From: Chung-Lin Tang > Date: Tue, 13 Jun 2023 08:44:31 -0700 > Subject: [PATCH] OpenACC 2.7: Implement self clause for compute constructs > .../c-c++-common/goacc/self-clause-1.c | 22 + > .../c-c++-common/goacc/self-clause-2.c | 17 + > gcc/testsuite/gfortran.dg/goacc/self.f95 | 53 + > .../libgomp.oacc-c-c++-common/self-1.c | 962 ++++++++++++++++++ I found that insufficient, and added some more. Pushed to master branch commit 047841a68ebf5f991e842961f9e54f3c10b94f2c "Extend test suite coverage for OpenACC 'self' clause for compute constructs", see attached. This is mostly just adapting and cross-linking some existing 'if' clause test cases. (..., which turned up a problem when the 'self' clause is used with OpenACC 'kernels'.) 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 047841a68ebf5f991e842961f9e54f3c10b94f2c Mon Sep 17 00:00:00 2001 From: Thomas Schwinge Date: Mon, 23 Oct 2023 14:53:29 +0200 Subject: [PATCH] Extend test suite coverage for OpenACC 'self' clause for compute constructs ... on top of what was provided in recent commit 3a3596389c2e539cb8fd5dc5784a4e2afe193a2a "OpenACC 2.7: Implement self clause for compute constructs". gcc/testsuite/ * c-c++-common/goacc/if-clause-2.c: Enhance. * c-c++-common/goacc/self-clause-1.c: Likewise. * c-c++-common/goacc/self-clause-2.c: Likewise. * gfortran.dg/goacc/if.f95: Likewise. * gfortran.dg/goacc/kernels-tree.f95: Likewise. * gfortran.dg/goacc/parallel-tree.f95: Likewise. * gfortran.dg/goacc/self.f95: Likewise. libgomp/ * testsuite/libgomp.oacc-c-c++-common/if-1.c: Enhance. * testsuite/libgomp.oacc-c-c++-common/self-1.c: Likewise. * testsuite/libgomp.oacc-fortran/if-1.f90: Likewise. * testsuite/libgomp.oacc-c-c++-common/if-self-1.c: New. * testsuite/libgomp.oacc-fortran/self-1.f90: Likewise. --- .../c-c++-common/goacc/if-clause-2.c | 2 + .../c-c++-common/goacc/self-clause-1.c | 6 + .../c-c++-common/goacc/self-clause-2.c | 20 + gcc/testsuite/gfortran.dg/goacc/if.f95 | 10 +- .../gfortran.dg/goacc/kernels-tree.f95 | 5 +- .../gfortran.dg/goacc/parallel-tree.f95 | 3 +- gcc/testsuite/gfortran.dg/goacc/self.f95 | 8 + .../libgomp.oacc-c-c++-common/if-1.c | 4 + .../libgomp.oacc-c-c++-common/if-self-1.c | 36 + .../libgomp.oacc-c-c++-common/self-1.c | 5 + .../testsuite/libgomp.oacc-fortran/if-1.f90 | 4 + .../testsuite/libgomp.oacc-fortran/self-1.f90 | 996 ++++++++++++++++++ 12 files changed, 1094 insertions(+), 5 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 diff --git a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c index a48072509e1..71475521758 100644 --- a/gcc/testsuite/c-c++-common/goacc/if-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/if-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'self-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ /* { dg-additional-options "--param=openacc-kernels=decompose" } { dg-additional-options "-fdump-tree-omp_oacc_kernels_decompose" } */ diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c index fe892bea210..28de3dc0584 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-1.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-1.c @@ -5,6 +5,8 @@ f (int b) { struct { int i; } *p; +#pragma acc parallel self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc parallel self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc parallel self(*p) @@ -12,6 +14,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc kernels self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc kernels self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc kernels self(*p) @@ -19,6 +23,8 @@ f (int b) { dg-error {could not convert '\* p' from 'f\(int\)::' to 'bool'} {} { target c++ } .-2 } */ ; +#pragma acc serial self(0) self(b) /* { dg-error "too many 'self' clauses" } */ + ; #pragma acc serial self self(b) /* { dg-error "too many 'self' clauses" } */ ; #pragma acc serial self(*p) diff --git a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c index d932ac9a4a6..769694baec9 100644 --- a/gcc/testsuite/c-c++-common/goacc/self-clause-2.c +++ b/gcc/testsuite/c-c++-common/goacc/self-clause-2.c @@ -1,3 +1,5 @@ +/* See also 'if-clause-2.c'. */ + /* { dg-additional-options "-fdump-tree-gimple" } */ void @@ -15,3 +17,21 @@ f (short c) /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:c \[len: [0-9]+\]\) self\(_[0-9]+\)$} 1 "gimple" } } */ ++c; } + +/* The same, but with implicit 'true' condition-argument. */ + +void +g (short d) +{ +#pragma acc parallel self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc kernels self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_kernels map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; + +#pragma acc serial self copy(d) + /* { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_serial map\(tofrom:d \[len: [0-9]+\]\) self\(1\)$} 1 "gimple" } } */ + ++d; +} diff --git a/gcc/testsuite/gfortran.dg/goacc/if.f95 b/gcc/testsuite/gfortran.dg/goacc/if.f95 index 56f3711f320..753ef8251c2 100644 --- a/gcc/testsuite/gfortran.dg/goacc/if.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/if.f95 @@ -1,3 +1,5 @@ +! See also 'self.f95'. + ! { dg-do compile } program test @@ -12,12 +14,14 @@ program test !$acc end parallel !$acc parallel if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end parallel - !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } - !$acc end kernels + !$acc kernels if ! { dg-error "Expected '\\(' after 'if'" } !$acc kernels if () ! { dg-error "Invalid character" } + !$acc kernels if (i) ! { dg-error "scalar LOGICAL expression" } + !$acc end kernels !$acc kernels if (1) ! { dg-error "scalar LOGICAL expression" } !$acc end kernels + !$acc data if ! { dg-error "Expected '\\(' after 'if'" } !$acc data if () ! { dg-error "Invalid character" } !$acc data if (i) ! { dg-error "scalar LOGICAL expression" } @@ -36,12 +40,14 @@ program test !$acc end parallel !$acc parallel if (i.gt.1) !$acc end parallel + !$acc kernels if (x) !$acc end kernels !$acc kernels if (.true.) !$acc end kernels !$acc kernels if (i.gt.1) !$acc end kernels + !$acc data if (x) !$acc end data !$acc data if (.true.) diff --git a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 index ceb07fbb9e9..1ba04a84e12 100644 --- a/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/kernels-tree.f95 @@ -12,6 +12,7 @@ program test logical :: l = .true. !$acc kernels if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -27,7 +28,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(from:k\\)" 1 "original" } } @@ -42,4 +43,4 @@ end program test ! { dg-final { scan-tree-dump-times "map\\(force_deviceptr:u\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_data_kernels if\((?:D\.|_)[0-9]+\)$} 1 "omp_oacc_kernels_decompose" } } -! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } +! { dg-final { scan-tree-dump-times {(?n)#pragma omp target oacc_parallel_kernels_gang_single num_gangs\(1\) if\((?:D\.|_)[0-9]+\) self\(1\) async\(-1\)$} 1 "omp_oacc_kernels_decompose" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 index 6110d93b91e..0d4ec1133af 100644 --- a/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/parallel-tree.f95 @@ -14,6 +14,7 @@ program test logical :: l = .true. !$acc parallel if(l) async num_gangs(i) num_workers(i) vector_length(i) & + !$acc self & !$acc reduction(max:q), copy(i), copyin(j), copyout(k), create(m) & !$acc no_create(n) & !$acc present(o), pcopy(p), pcopyin(r), pcopyout(s), pcreate(t) & @@ -33,7 +34,7 @@ end program test ! { dg-final { scan-tree-dump-times "num_gangs" 1 "original" } } ! { dg-final { scan-tree-dump-times "num_workers" 1 "original" } } ! { dg-final { scan-tree-dump-times "vector_length" 1 "original" } } - +! { dg-final { scan-tree-dump-times "self\\(1\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "reduction\\(max:q\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(tofrom:i\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "map\\(to:j\\)" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/self.f95 b/gcc/testsuite/gfortran.dg/goacc/self.f95 index 4817f16be56..aa0f6fe88c5 100644 --- a/gcc/testsuite/gfortran.dg/goacc/self.f95 +++ b/gcc/testsuite/gfortran.dg/goacc/self.f95 @@ -1,3 +1,5 @@ +! See also 'if.f95'. + ! { dg-do compile } program test @@ -29,6 +31,8 @@ program test !$acc kernels self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } !$acc serial self (.false.) self (.false.) { dg-error "Duplicated 'self' clause" } + !$acc parallel self + !$acc end parallel !$acc parallel self (x) !$acc end parallel !$acc parallel self (.true.) @@ -36,6 +40,8 @@ program test !$acc parallel self (i.gt.1) !$acc end parallel + !$acc kernels self + !$acc end kernels !$acc kernels self (x) !$acc end kernels !$acc kernels self (.true.) @@ -43,6 +49,8 @@ program test !$acc kernels self (i.gt.1) !$acc end kernels + !$acc serial self + !$acc end serial !$acc serial self (x) !$acc end serial !$acc serial self (.true.) diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c index 5398905f411..f04c02fdb89 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-1.c @@ -1,3 +1,7 @@ +/* OpenACC 'if' clause. */ + +/* See also 'self-1.c'. */ + #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c new file mode 100644 index 00000000000..f77edda83e3 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/if-self-1.c @@ -0,0 +1,36 @@ +/* Test 'if' and 'self' clause appearing together. */ + +#include + +static int test(float i, long double s) +{ + int ret; +#pragma acc serial copyout(ret) if(i) self(s) + /* { dg-bogus {using 'vector_length \(32\)', ignoring 1} {} { xfail openacc_nvidia_accel_selected } .-1 } */ + { + ret = acc_on_device(acc_device_host); + } + return ret; +} + +int main() +{ + if (!test(0, 0)) + __builtin_abort(); + + if (!test(0, 1)) + __builtin_abort(); + +#if ACC_MEM_SHARED + if (!test(1, 0)) + __builtin_abort(); +#else + if (test(1, 0)) + __builtin_abort(); +#endif + + if (!test(1, 1)) + __builtin_abort(); + + return 0; +} diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c index 752e16e8545..d17e27c8af3 100644 --- a/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c +++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/self-1.c @@ -1,3 +1,8 @@ +/* OpenACC 'self' clause. */ + +/* This is 'if-1.c' with 'self(!cond)' instead of 'if(cond)' on compute + constructs. */ + #include #include #include diff --git a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 index e0cfd912d0f..5ba6509749e 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/if-1.f90 @@ -1,3 +1,7 @@ +! OpenACC 'if' clause. + +! See also 'self-1.f90'. + ! { dg-do run } ! { dg-additional-options "-cpp" } diff --git a/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 new file mode 100644 index 00000000000..b9ec9de08d9 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/self-1.f90 @@ -0,0 +1,996 @@ +! OpenACC 'self' clause. + +! This is 'if-1.f90' with 'self(!cond)' instead of 'if(cond)' on compute +! constructs. +! ..., which the exception of certain 'kernels' constructs. + +! { dg-do run } +! { dg-additional-options "-cpp" } + +! { dg-additional-options "--param=openacc-kernels=decompose" } + +! { 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" } +! Prune a few: uninteresting, and potentially varying depending on GCC configuration (data types): +! { 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] } +! { dg-message "dummy" "" { target iN-VAl-Id } l_dummy } to avoid +! "WARNING: dg-line var l_dummy defined, but not used". */ + +program main + use openacc + implicit none + + integer, parameter :: N = 8 + integer, parameter :: one = 1 + integer, parameter :: zero = 0 + integer i, nn + real, allocatable :: a(:), b(:) + real exp, exp2 + + i = 0 + + allocate (a(N)) + allocate (b(N)) + + a(:) = 4.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + !TODO Unhandled 'CONST_DECL' instances for constant argument in 'acc_on_device' call. + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 1 + end do + + a(:) = 16.0 + + !$acc parallel self (0 /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 17.0) STOP 2 + end do + + a(:) = 8.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 3 + end do + + a(:) = 22.0 + + !$acc parallel self (zero /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 23.0) STOP 4 + end do + + a(:) = 16.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 5 + end do + + a(:) = 76.0 + + !$acc parallel self (.TRUE.) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 77.0) STOP 6 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 7 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc parallel self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 19.0) STOP 8 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 9 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 39.0) STOP 10 + end do + + a(:) = 91.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (.NOT. (-2 > 0)) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 92.0) STOP 11 + end do + + a(:) = 43.0 + + !$acc parallel copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 12 + end do + + a(:) = 87.0 + + !$acc parallel self (one /= 0) ! { dg-line l_compute[incr c_compute] } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end parallel + + do i = 1, N + if (b(i) .ne. 88.0) STOP 13 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 14 + if (b(i) .ne. exp2) STOP 15 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 16 + if (b(i) .ne. exp2) STOP 17 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 18 + if (b(i) .ne. 0.0) STOP 19 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 20 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 21 + if (acc_is_present (b) .eqv. .TRUE.) STOP 22 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 23 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 24 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc parallel present (a(1:N)) present (b(1:N)) + do i = 1, N + b(i) = a(i) + end do + !$acc end parallel + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 25 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 26 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 27 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 28 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 29 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 30 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 31 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 32 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 33 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 34 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 35 +#endif + + a(:) = 4.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (1 /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 5.0 +#else + exp = 4.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 36 + end do + + a(:) = 16.0 + + !$acc kernels if (0 == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 17.0) STOP 37 + end do + + a(:) = 8.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 9.0 +#else + exp = 8.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 38 + end do + + a(:) = 22.0 + + !$acc kernels if (zero == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 23.0) STOP 39 + end do + + a(:) = 16.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 17.0; +#else + exp = 16.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 40 + end do + + a(:) = 76.0 + + !$acc kernels if (.FALSE.) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 77.0) STOP 41 + end do + + a(:) = 22.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (nn /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 23.0; +#else + exp = 22.0; +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 42 + end do + + a(:) = 18.0 + + nn = 0 + + !$acc kernels if (nn == 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 19.0) STOP 43 + end do + + a(:) = 49.0 + + nn = 1 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (.NOT. ((nn + nn) > 0)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 50.0 +#else + exp = 49.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 44 + end do + + a(:) = 38.0 + + nn = 0; + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if ((nn + nn) > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 39.0) STOP 45 + end do + + a(:) = 91.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) if (-2 > 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 92.0) STOP 46 + end do + + a(:) = 43.0 + + !$acc kernels copyin (a(1:N)) copyout (b(1:N)) self (one /= 1) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + +#if ACC_MEM_SHARED + exp = 44.0 +#else + exp = 43.0 +#endif + + do i = 1, N + if (b(i) .ne. exp) STOP 47 + end do + + a(:) = 87.0 + + !$acc kernels if (one == 0) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + ! { dg-note {variable 'C\.[0-9]+' declared in block potentially has improper OpenACC privatization level: 'const_decl'} "TODO" { target *-*-* } l_compute$c_compute } + if (acc_on_device (acc_device_host) .eqv. .TRUE.) then + b(i) = a(i) + 1 + else + b(i) = a(i) + end if + end do + !$acc end kernels + + do i = 1, N + if (b(i) .ne. 88.0) STOP 48 + end do + + a(:) = 3.0 + b(:) = 9.0 + +#if ACC_MEM_SHARED + exp = 0.0 + exp2 = 0.0 +#else + call acc_copyin (a, sizeof (a)) + call acc_copyin (b, sizeof (b)) + exp = 3.0; + exp2 = 9.0; +#endif + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 49 + if (b(i) .ne. exp2) STOP 50 + end do + + a(:) = 6.0 + b(:) = 12.0 + + !$acc update device (a(1:N), b(1:N)) if (0 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (1 == 1) + + do i = 1, N + if (a(i) .ne. exp) STOP 51 + if (b(i) .ne. exp2) STOP 52 + end do + + a(:) = 26.0 + b(:) = 21.0 + + !$acc update device (a(1:N), b(1:N)) if (1 == 1) + + a(:) = 0.0 + b(:) = 0.0 + + !$acc update host (a(1:N), b(1:N)) if (0 == 1) + + do i = 1, N + if (a(i) .ne. 0.0) STOP 53 + if (b(i) .ne. 0.0) STOP 54 + end do + +#if !ACC_MEM_SHARED + call acc_copyout (a, sizeof (a)) + call acc_copyout (b, sizeof (b)) +#endif + + a(:) = 4.0 + b(:) = 0.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + !$acc end data + + do i = 1, N + if (b(i) .ne. 4.0) STOP 55 + end do + + a(:) = 8.0 + b(:) = 1.0 + + !$acc data copyin (a(1:N)) copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-1 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .TRUE.) STOP 56 + if (acc_is_present (b) .eqv. .TRUE.) STOP 57 +#endif + + !$acc end data + + a(:) = 18.0 + b(:) = 21.0 + + !$acc data copyin (a(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } + +#if !ACC_MEM_SHARED + if (acc_is_present (a) .eqv. .FALSE.) STOP 58 +#endif + + !$acc data copyout (b(1:N)) if (0 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: artificial} "" { target { ! openacc_host_selected } } .-2 } +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 59 +#endif + !$acc data copyout (b(1:N)) if (1 == 1) + ! { dg-note {variable 'parm\.[0-9]+' declared in block isn't candidate for adjusting OpenACC privatization level: not addressable} "" { target *-*-* } .-1 } + + !$acc kernels present (a(1:N)) present (b(1:N)) ! { dg-line l_compute[incr c_compute] } + ! { dg-note {OpenACC 'kernels' decomposition: variable 'i' in 'copy' clause requested to be made addressable} {} { target *-*-* } l_compute$c_compute } + ! { dg-note {variable 'i' already made addressable} {} { target *-*-* } l_compute$c_compute } */ + ! { dg-note {beginning 'parloops' part in OpenACC 'kernels' region} "" { target *-*-* } .+1 } + do i = 1, N + b(i) = a(i) + end do + !$acc end kernels + + !$acc end data + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 60 +#endif + !$acc end data + !$acc end data + + do i = 1, N + if (b(1) .ne. 18.0) STOP 61 + end do + + !$acc enter data copyin (b(1:N)) if (0 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 62 +#endif + + !$acc exit data delete (b(1:N)) if (0 == 1) + + !$acc enter data copyin (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 63 +#endif + + !$acc exit data delete (b(1:N)) if (1 == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 64 +#endif + + !$acc enter data copyin (b(1:N)) if (zero == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 65 +#endif + + !$acc exit data delete (b(1:N)) if (zero == 1) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 66 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 67 +#endif + + !$acc enter data copyin (b(1:N)) if (one == 0) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 68 +#endif + + !$acc exit data delete (b(1:N)) if (one == 0) + + !$acc enter data copyin (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .FALSE.) STOP 69 +#endif + + !$acc exit data delete (b(1:N)) if (one == 1) + +#if !ACC_MEM_SHARED + if (acc_is_present (b) .eqv. .TRUE.) STOP 70 +#endif + +end program main -- 2.34.1