From patchwork Wed Dec 20 14:48:29 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Abid Qadeer X-Patchwork-Id: 181665 Return-Path: Delivered-To: ouuuleilei@gmail.com Received: by 2002:a05:7300:24d3:b0:fb:cd0c:d3e with SMTP id r19csp2688313dyi; Wed, 20 Dec 2023 06:49:21 -0800 (PST) X-Google-Smtp-Source: AGHT+IGK1UMMdmmI6QyOAHed39IAZnjZZnpXl0ZLgl/ZG0jIrgQBBi1h+CIAEfeS1uq9ln2xYwAn X-Received: by 2002:a05:6808:6541:b0:3ba:1080:f8e8 with SMTP id fn1-20020a056808654100b003ba1080f8e8mr18245332oib.69.1703083761364; Wed, 20 Dec 2023 06:49:21 -0800 (PST) ARC-Seal: i=2; a=rsa-sha256; t=1703083761; cv=pass; d=google.com; s=arc-20160816; b=YdMgf0u79/cTBrzSq1obUMQ74zD/T0igR9tmnU8HSgwDEfczBiZUD4PgJOpKL1SSAj 3B3ozkRhf/bQqz1s1mqH/87XVXCj0CksTSeyZwVZUw2m8y4JiIbjFgEVpsIRtQaJVFrn cw6xJRKDldF6u/oBzRygL6kqYCgK8J73kwrkiHs0vndFFXdQKEHYmSrTD1g2ONsDR8Ft vM3q8J+QA5CgX7uY94wT/NSlbm6Kp/+T/OlE+pyw5OxTmp1pcFJPuJmBRWoRiaEXUCYA 2KSePsQiJQ57/n/rhis8qoyhPld9BChEFNbMnO7k5yoY2zkGl1JNR0BfU30kHkDk2RG3 VGMQ== 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:content-transfer-encoding :mime-version:message-id:date:subject:cc:to:from:ironport-sdr :arc-filter:dmarc-filter:delivered-to; bh=mr7jDVkwQWUElcEZi2g+f/uRkgbQGY1TRH/pwNcKAz0=; fh=lGV3fxH6Z71xVqz7LPO5z7kRU1jIXyhlGCWs7Mp/Ju8=; b=gBbUXINPcbOyNV+t2swa57W+HOI8ieQvk0uVsbPifw1Rf1z8OMswJWq7uoZuJxpXnq SsR+aC4qZQtPjJrq+7yMw+7Z2lhyc7QPZ1hfvvnSurnMSyA+b2LBL5bkaJ3K9hKANQlt IkuUd9VZNtgqrmQK62K/gOPtm2LnN0AVmjku8wBzO+GKQo2h8cPnjJGLshO+EqOsm9pa e2TrAnEz0r2BVkCai6uAMryHZo/DWXVZh8zio8FoYO39IOx8zkCS7dytERVnyEA3mEwM WJp95KoRmRuLMUldB8BAhKIpv/o9Rje7U8zauukw6ivP4WVLGCpbBCKm1iD1C/MJ+xfR VGjQ== 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 j23-20020a05620a147700b0078116d4a4easi717855qkl.28.2023.12.20.06.49.21 for (version=TLS1_3 cipher=TLS_AES_256_GCM_SHA384 bits=256/256); Wed, 20 Dec 2023 06:49:21 -0800 (PST) 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 09F26385E006 for ; Wed, 20 Dec 2023 14:49:21 +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 6E5853858D20; Wed, 20 Dec 2023 14:48:44 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 6E5853858D20 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 6E5853858D20 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=1703083735; cv=none; b=A9B0VFXJpQiPSLtzkDdAKC7LAdQQ+VuICkn9tIdfF7fTcHM2BO/EDfjjipoT1/xp0uimqtv34qSUBbWVapVQrnsSMaq+YpLVXDI5NN3L/HbPIpRZF5nKm0M0QcUpMQwPiTb3JCgfXpJsdPSrbzi8gx3ZSk71tsPynb8dWudBe50= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1703083735; c=relaxed/simple; bh=GvpQd5ycLnomdK+K0CxaeXmpbQzmqZHtxcrrS0RK09w=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=UjM6pJ1xyN94s9eNW5i1qgD8H5dfibhZeYW0VjWPyW87IgAPkxGdLWSBGAiSIgSnio9geNukmSOSxK/8EpivH8BQAzNZuNGr1oP/mBuwq2kWPJerBf2SIqRkYt8VZBjjbxTAuriLlFZfnQa7mP3XR34Yl+79GOreZrycwbFNZ1I= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: jdtCYfZGQqqvg7WFbRFB7w== X-CSE-MsgGUID: cakPgwDfSAyEr8cU3yya8Q== X-IronPort-AV: E=Sophos;i="6.04,291,1695715200"; d="scan'208";a="28878890" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 20 Dec 2023 06:48:42 -0800 IronPort-SDR: AiM896yLiTtJW6Hk9aUXhqcoFFES+hJqOcx1TlJBZPoLAKkaugTViBGpI8O+0sPsWxEJx05emD NJyx9ONI3cNfRtqy5Otx5d0lonj0Ki0nKbl7qto4kzzRxqnllJXjfQchanJygdrJpvC1yjH3sx jOs2QV02k3PFujkLAVkX6x/1G+O6SZ0pro4mqa6qP6h5T/6ig0n0XqrKFLxXk/CzjrthbZyTT3 JnaigELf1c//97Y/eGHM41mMNdk9rAUnokhLhRnTQGtQt4/+eeWqvh2jJW39r6nrtlChdzr+IL t3o= From: Abid Qadeer To: , CC: , Subject: [PATCH] [OpenACC] Add tests for implied copy of variables in reduction clause. Date: Wed, 20 Dec 2023 14:48:29 +0000 Message-ID: <20231220144829.765056-1-abidh@codesourcery.com> X-Mailer: git-send-email 2.25.1 MIME-Version: 1.0 X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-13.mgc.mentorg.com (139.181.222.13) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-12.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, SPF_HELO_PASS, SPF_PASS, TXREP, T_SCC_BODY_TEXT_LINE 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: 1785812758019056716 X-GMAIL-MSGID: 1785812758019056716 From: Hafiz Abid Qadeer The OpenACC reduction clause on compute construct implies a copy clause for each reduction variable [1]. This patch adds tests to check if the implied copy is being generated. The check covers various types and operators as described in the specification. gcc/testsuite/ChangeLog: * c-c++-common/goacc/implied-copy-1.c: New test. * c-c++-common/goacc/implied-copy-2.c: New test. * g++.dg/goacc/implied-copy.C: New test. * gcc.dg/goacc/implied-copy.c: New test. * gfortran.dg/goacc/implied-copy-1.f90: New test. * gfortran.dg/goacc/implied-copy-2.f90: New test. [1] OpenACC 2.7 Specification section 2.5.13 --- .../c-c++-common/goacc/implied-copy-1.c | 33 ++++ .../c-c++-common/goacc/implied-copy-2.c | 121 +++++++++++++ gcc/testsuite/g++.dg/goacc/implied-copy.C | 24 +++ gcc/testsuite/gcc.dg/goacc/implied-copy.c | 29 ++++ .../gfortran.dg/goacc/implied-copy-1.f90 | 35 ++++ .../gfortran.dg/goacc/implied-copy-2.f90 | 160 ++++++++++++++++++ 6 files changed, 402 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/goacc/implied-copy-1.c create mode 100644 gcc/testsuite/c-c++-common/goacc/implied-copy-2.c create mode 100644 gcc/testsuite/g++.dg/goacc/implied-copy.C create mode 100644 gcc/testsuite/gcc.dg/goacc/implied-copy.c create mode 100644 gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/goacc/implied-copy-2.f90 diff --git a/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c b/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c new file mode 100644 index 00000000000..ae06339dc2d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/implied-copy-1.c @@ -0,0 +1,33 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +/* Test for implied copy of reduction variable on combined construct. */ +void test1 (void) +{ + int i, sum = 0, prod = 1, a[100]; + + #pragma acc kernels loop reduction(+:sum) reduction(*:prod) + for (int i = 0; i < 10; ++i) + { + sum += a[i]; + prod *= a[i]; + } + + #pragma acc parallel loop reduction(+:sum) reduction(*:prod) + for (int i = 0; i < 10; ++i) + { + sum += a[i]; + prod *= a[i]; + } + + #pragma acc serial loop reduction(+:sum) reduction(*:prod) + for (int i = 0; i < 10; ++i) + { + sum += a[i]; + prod *= a[i]; + } +} + +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(force_tofrom:prod \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:prod \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/goacc/implied-copy-2.c b/gcc/testsuite/c-c++-common/goacc/implied-copy-2.c new file mode 100644 index 00000000000..124f128964d --- /dev/null +++ b/gcc/testsuite/c-c++-common/goacc/implied-copy-2.c @@ -0,0 +1,121 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +/* Test that reduction on compute construct implies a copy of the reduction + variable . */ + +#define n 1000 + +#if __cplusplus + typedef bool BOOL; +#else + typedef _Bool BOOL; +#endif + +int +main(void) +{ + int i; + int sum = 0; + int prod = 1; + int result = 0; + int tmp = 1; + int array[n]; + + double sumd = 0.0; + double arrayd[n]; + + float sumf = 0.0; + float arrayf[n]; + + char sumc; + char arrayc[n]; + + BOOL lres; + +#pragma acc parallel reduction(+:sum, sumf, sumd, sumc) reduction(*:prod) + for (i = 0; i < n; i++) + { + sum += array[i]; + sumf += arrayf[i]; + sumd += arrayd[i]; + sumc += arrayc[i]; + prod *= array[i]; + } + +#pragma acc parallel reduction (max:result) + for (i = 0; i < n; i++) + result = result > array[i] ? result : array[i]; + +#pragma acc parallel reduction (min:result) + for (i = 0; i < n; i++) + result = result < array[i] ? result : array[i]; + +#pragma acc parallel reduction (&:result) + for (i = 0; i < n; i++) + result &= array[i]; + +#pragma acc parallel reduction (|:result) + for (i = 0; i < n; i++) + result |= array[i]; + +#pragma acc parallel reduction (^:result) + for (i = 0; i < n; i++) + result ^= array[i]; + +#pragma acc parallel reduction (&&:lres) copy(tmp) + for (i = 0; i < n; i++) + lres = lres && (tmp > array[i]); + +#pragma acc parallel reduction (||:lres) copy(tmp) + for (i = 0; i < n; i++) + lres = lres || (tmp > array[i]); + + /* Same checks on serial construct. */ +#pragma acc serial reduction(+:sum, sumf, sumd, sumc) reduction(*:prod) + for (i = 0; i < n; i++) + { + sum += array[i]; + sumf += arrayf[i]; + sumd += arrayd[i]; + sumc += arrayc[i]; + prod *= array[i]; + } + +#pragma acc serial reduction (max:result) + for (i = 0; i < n; i++) + result = result > array[i] ? result : array[i]; + +#pragma acc serial reduction (min:result) + for (i = 0; i < n; i++) + result = result < array[i] ? result : array[i]; + +#pragma acc serial reduction (&:result) + for (i = 0; i < n; i++) + result &= array[i]; + +#pragma acc serial reduction (|:result) + for (i = 0; i < n; i++) + result |= array[i]; + +#pragma acc serial reduction (^:result) + for (i = 0; i < n; i++) + result ^= array[i]; + +#pragma acc serial reduction (&&:lres) copy(tmp) + for (i = 0; i < n; i++) + lres = lres && (tmp > array[i]); + +#pragma acc serial reduction (||:lres) copy(tmp) + for (i = 0; i < n; i++) + lres = lres || (tmp > array[i]); + + return 0; +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom:sum \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:sumf \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:sumd \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:sumc \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:lres \\\[len: \[0-9\]+\\\]\\)" 4 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:result \\\[len: \[0-9\]+\\\]\\)" 10 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:prod \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/g++.dg/goacc/implied-copy.C b/gcc/testsuite/g++.dg/goacc/implied-copy.C new file mode 100644 index 00000000000..cfc4e133e48 --- /dev/null +++ b/gcc/testsuite/g++.dg/goacc/implied-copy.C @@ -0,0 +1,24 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +/* Test for wchar_t type. */ +int +main(void) +{ + int i; + wchar_t a[100], s; + +#pragma acc parallel reduction (+:s) + for (i = 0; i < 10; i++) + { + s += a[i]; + } + +#pragma acc serial reduction (+:s) + for (i = 0; i < 10; i++) + { + s += a[i]; + } + return 0; +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom:s \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/gcc.dg/goacc/implied-copy.c b/gcc/testsuite/gcc.dg/goacc/implied-copy.c new file mode 100644 index 00000000000..98adbcb1987 --- /dev/null +++ b/gcc/testsuite/gcc.dg/goacc/implied-copy.c @@ -0,0 +1,29 @@ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +/* Test for float _Complex and double _Complex types. */ +int +main(void) +{ + int i; + float _Complex fc[100]; + float _Complex s1; + double _Complex dc[100]; + double _Complex s2; + +#pragma acc parallel reduction (+:s1, s2) + for (i = 0; i < 10; i++) + { + s1 += fc[i]; + s2 += dc[i]; + } +#pragma acc serial reduction (+:s1, s2) + for (i = 0; i < 10; i++) + { + s1 += fc[i]; + s2 += dc[i]; + } + return 0; +} + +/* { dg-final { scan-tree-dump-times "map\\(tofrom:s1 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ +/* { dg-final { scan-tree-dump-times "map\\(tofrom:s2 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } */ diff --git a/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 b/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 new file mode 100644 index 00000000000..7f07c8afbec --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/implied-copy-1.f90 @@ -0,0 +1,35 @@ +! { dg-additional-options "-fdump-tree-gimple" } + +! Test for implied copy of reduction variable on combined construct. + +subroutine test + implicit none + integer a(100), i, s, p + p = 1 + + !$acc parallel loop reduction(+:s) reduction(*:p) + do i = 1, 100 + s = s + a(i) + p = p * a(i) + end do + !$acc end parallel loop + + !$acc serial loop reduction(+:s) reduction(*:p) + do i = 1, 100 + s = s + a(i) + p = p * a(i) + end do + !$acc end serial loop + + !$acc kernels loop reduction(+:s) reduction(*:p) + do i = 1, 100 + s = s + a(i) + p = p * a(i) + end do + !$acc end kernels loop +end subroutine test + +! { dg-final { scan-tree-dump-times "map\\(force_tofrom:s \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(force_tofrom:p \\\[len: \[0-9\]+\\\]\\)" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:s \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:p \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/goacc/implied-copy-2.f90 b/gcc/testsuite/gfortran.dg/goacc/implied-copy-2.f90 new file mode 100644 index 00000000000..24e5610723b --- /dev/null +++ b/gcc/testsuite/gfortran.dg/goacc/implied-copy-2.f90 @@ -0,0 +1,160 @@ +! { dg-additional-options "-fdump-tree-gimple" } + +! Test that reduction on compute construct implies a copy of the reduction variable + +subroutine test + implicit none + integer i + integer a(100), s1, p1 + integer r1 + real b(100), s2 + logical c(100), r2 + double precision d(100), s3 + complex e(100), s4 + p1 = 1 + + !$acc parallel reduction(+:s1, s2, s3, s4) reduction(*:p1) + do i = 1, 100 + s1 = s1 + a(i) + p1 = p1 * a(i) + s2 = s2 + b(i) + s3 = s3 + d(i) + s4 = s4 + e(i) + end do + !$acc end parallel + + !$acc parallel reduction (max:r1) + do i = 1,10 + if (r1 <= a(i)) then + r1 = a(i) + end if + end do + !$acc end parallel + + !$acc parallel reduction (min:r1) + do i = 1,10 + if (r1 >= a(i)) then + r1 = a(i) + end if + end do + !$acc end parallel + + !$acc parallel reduction (iand:r1) + do i = 1,10 + r1 = iand (r1, a(i)) + end do + !$acc end parallel + + !$acc parallel reduction (ior:r1) + do i = 1,10 + r1 = ior (r1, a(i)) + end do + !$acc end parallel + + !$acc parallel reduction (ieor:r1) + do i = 1,10 + r1 = ieor (r1, a(i)) + end do + !$acc end parallel + + !$acc parallel reduction (.and.:r2) + do i = 1,10 + r2 = r2 .and. c(i) + end do + !$acc end parallel + + !$acc parallel reduction (.or.:r2) + do i = 1,10 + r2 = r2 .or. c(i) + end do + !$acc end parallel + + !$acc parallel reduction (.eqv.:r2) + do i = 1,10 + r2 = r2 .eqv. c(i) + end do + !$acc end parallel + + !$acc parallel reduction (.neqv.:r2) + do i = 1,10 + r2 = r2 .neqv. c(i) + end do + !$acc end parallel + + !$acc serial reduction(+:s1, s2, s3, s4) reduction(*:p1) + do i = 1, 100 + s1 = s1 + a(i) + p1 = p1 * a(i) + s2 = s2 + b(i) + s3 = s3 + d(i) + s4 = s4 + e(i) + end do + !$acc end serial + + !$acc serial reduction (max:r1) + do i = 1,10 + if (r1 <= a(i)) then + r1 = a(i) + end if + end do + !$acc end serial + + !$acc serial reduction (min:r1) + do i = 1,10 + if (r1 >= a(i)) then + r1 = a(i) + end if + end do + !$acc end serial + + !$acc serial reduction (iand:r1) + do i = 1,10 + r1 = iand (r1, a(i)) + end do + !$acc end serial + + !$acc serial reduction (ior:r1) + do i = 1,10 + r1 = ior (r1, a(i)) + end do + !$acc end serial + + !$acc serial reduction (ieor:r1) + do i = 1,10 + r1 = ieor (r1, a(i)) + end do + !$acc end serial + + !$acc serial reduction (.and.:r2) + do i = 1,10 + r2 = r2 .and. c(i) + end do + !$acc end serial + + !$acc serial reduction (.or.:r2) + do i = 1,10 + r2 = r2 .or. c(i) + end do + !$acc end serial + + !$acc serial reduction (.eqv.:r2) + do i = 1,10 + r2 = r2 .eqv. c(i) + end do + !$acc end serial + + !$acc serial reduction (.neqv.:r2) + do i = 1,10 + r2 = r2 .neqv. c(i) + end do + !$acc end serial + +end subroutine test + +! { dg-final { scan-tree-dump-times "map\\(tofrom:s1 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:r1 \\\[len: \[0-9\]+\\\]\\)" 10 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:r2 \\\[len: \[0-9\]+\\\]\\)" 8 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:s2 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:s3 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:s4 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(tofrom:p1 \\\[len: \[0-9\]+\\\]\\)" 2 "gimple" } }