From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: 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 Content-Transfer-Encoding: 8bit Content-Type: text/plain 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 List-Id: 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" } } -- 2.25.1