From patchwork Fri Dec 10 17:40:36 2021 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Kwok Cheung Yeung X-Patchwork-Id: 48799 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id 217D03857C7A for ; Fri, 10 Dec 2021 17:41:35 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id 7A04C3857C4C for ; Fri, 10 Dec 2021 17:41:08 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.1 sourceware.org 7A04C3857C4C 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: OkB1Fr5h6ORB/P1Su+w/3vVXuLIbiVzYYupZ4PSSSKS2sZaEieHqGmps2PVeu81BhGm7tvGeya M5k32YNkIq048JYMpraD6EULHA7/52LQGfaZiGWfxolCP5Vq76o1UBpnTGq9IgIGPE3Ic24MMJ iYBVDeWkkJicJutSGoOZDAD6QjI095fd05vxX7/4tfRrc4siYbPerOGrUIUE3b+RC5B1y2GT5n aUnCxzmgMXhpwLWMx9rruFY8Aqc0Y7ReitHur4mSzqH4cRsoXKz+Wg5duTQXa4XF/YwcGMixci bpjvge3YTkMxXnieBXiql4Y9 X-IronPort-AV: E=Sophos;i="5.88,196,1635235200"; d="scan'208,223";a="69403716" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 10 Dec 2021 09:41:07 -0800 IronPort-SDR: LkxSkm1dyvwr9idECvBhWfgypiMtxNhrOq3b4lRfZul7dYbVajthmUsjaCQ5L45DqDwDUZQLMT SdcrdOdW5CddwRH5u7EitNDGPAsk6EjPf8g8ZW1uFNk9GGOb0UYVx/zHzQltKJ6WbG9rwSNJf7 d1TN1ywvxMLucWvv8xNEl4k3f03YQHvh23aPdiR3akBf8hgSopFqNZP6f/qjRtUfOEiL48NKkH ytyWCTeGTU6faq9NWwWZw8C3w0OcvrZM2y4fVZemHKZqdRDA4RqvjhC3kBKjlOeLnQfWmwF3Ww NjI= Message-ID: Date: Fri, 10 Dec 2021 17:40:36 +0000 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (Windows NT 10.0; Win64; x64; rv:91.0) Gecko/20100101 Thunderbird/91.4.0 To: gcc-patches , Jakub Jelinek References: Subject: [PATCH 7/7] openmp: Add testcases for metadirectives From: Kwok Cheung Yeung In-Reply-To: X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) To SVR-IES-MBX-04.mgc.mentorg.com (139.181.222.4) X-Spam-Status: No, score=-12.0 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: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" This adds testcases for metadirectives. Kwok From d3f80b603298fb2f3501a28b888acfdbc02a64e7 Mon Sep 17 00:00:00 2001 From: Kwok Cheung Yeung Date: Tue, 7 Dec 2021 11:25:33 +0000 Subject: [PATCH 7/7] openmp: Add testcases for metadirectives 2021-12-10 Kwok Cheung Yeung gcc/testsuite/ * c-c++-common/gomp/metadirective-1.c: New. * c-c++-common/gomp/metadirective-2.c: New. * c-c++-common/gomp/metadirective-3.c: New. * c-c++-common/gomp/metadirective-4.c: New. * c-c++-common/gomp/metadirective-5.c: New. * c-c++-common/gomp/metadirective-6.c: New. * gcc.dg/gomp/metadirective-1.c: New. * gfortran.dg/gomp/metadirective-1.f90: New. * gfortran.dg/gomp/metadirective-2.f90: New. * gfortran.dg/gomp/metadirective-3.f90: New. * gfortran.dg/gomp/metadirective-4.f90: New. * gfortran.dg/gomp/metadirective-5.f90: New. * gfortran.dg/gomp/metadirective-6.f90: New. libgomp/ * testsuite/libgomp.c-c++-common/metadirective-1.c: New. * testsuite/libgomp.c-c++-common/metadirective-2.c: New. * testsuite/libgomp.c-c++-common/metadirective-3.c: New. * testsuite/libgomp.c-c++-common/metadirective-4.c: New. * testsuite/libgomp.fortran/metadirective-1.f90: New. * testsuite/libgomp.fortran/metadirective-2.f90: New. * testsuite/libgomp.fortran/metadirective-3.f90: New. * testsuite/libgomp.fortran/metadirective-4.f90: New. --- .../c-c++-common/gomp/metadirective-1.c | 29 ++++++++ .../c-c++-common/gomp/metadirective-2.c | 74 +++++++++++++++++++ .../c-c++-common/gomp/metadirective-3.c | 31 ++++++++ .../c-c++-common/gomp/metadirective-4.c | 40 ++++++++++ .../c-c++-common/gomp/metadirective-5.c | 24 ++++++ .../c-c++-common/gomp/metadirective-6.c | 31 ++++++++ gcc/testsuite/gcc.dg/gomp/metadirective-1.c | 15 ++++ .../gfortran.dg/gomp/metadirective-1.f90 | 41 ++++++++++ .../gfortran.dg/gomp/metadirective-2.f90 | 59 +++++++++++++++ .../gfortran.dg/gomp/metadirective-3.f90 | 34 +++++++++ .../gfortran.dg/gomp/metadirective-4.f90 | 39 ++++++++++ .../gfortran.dg/gomp/metadirective-5.f90 | 30 ++++++++ .../gfortran.dg/gomp/metadirective-6.f90 | 31 ++++++++ .../libgomp.c-c++-common/metadirective-1.c | 35 +++++++++ .../libgomp.c-c++-common/metadirective-2.c | 41 ++++++++++ .../libgomp.c-c++-common/metadirective-3.c | 34 +++++++++ .../libgomp.c-c++-common/metadirective-4.c | 52 +++++++++++++ .../libgomp.fortran/metadirective-1.f90 | 33 +++++++++ .../libgomp.fortran/metadirective-2.f90 | 40 ++++++++++ .../libgomp.fortran/metadirective-3.f90 | 29 ++++++++ .../libgomp.fortran/metadirective-4.f90 | 46 ++++++++++++ 21 files changed, 788 insertions(+) create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-1.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-2.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-3.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-4.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-5.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-6.c create mode 100644 gcc/testsuite/gcc.dg/gomp/metadirective-1.c create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 create mode 100644 gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-1.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-3.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/metadirective-4.f90 diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-1.c b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c new file mode 100644 index 00000000000..72cf0abbbd7 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c @@ -0,0 +1,29 @@ +/* { dg-do compile } */ + +#define N 100 + +void f (int a[], int b[], int c[]) +{ + #pragma omp metadirective \ + default (teams loop) \ + default (parallel loop) /* { dg-error "there can only be one default clause in a metadirective before '\\(' token" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + default (bad_directive) /* { dg-error "unknown directive name before '\\)' token" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + default (teams loop) \ + where (device={arch("nvptx")}: parallel loop) /* { dg-error "expected 'when' or 'default' before '\\(' token" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + default (teams loop) \ + when (device={arch("nvptx")} parallel loop) /* { dg-error "expected colon before 'parallel'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + default (metadirective default (flush)) /* { dg-error "metadirectives cannot be used as directive variants before 'default'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; +} diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-2.c b/gcc/testsuite/c-c++-common/gomp/metadirective-2.c new file mode 100644 index 00000000000..ea6904c9c12 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-2.c @@ -0,0 +1,74 @@ +/* { dg-do compile } */ + +#define N 100 + +int main (void) +{ + int x = 0; + int y = 0; + + /* Test implicit default (nothing). */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: barrier) + x = 1; + + /* Test with multiple standalone directives. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: barrier) \ + default (flush) + x = 1; + + /* Test combining a standalone directive with one that takes a statement + body. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: parallel) \ + default (barrier) + x = 1; + + /* Test combining a standalone directive with one that takes a for loop. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: parallel for) \ + default (barrier) + for (int i = 0; i < N; i++) + x += i; + + /* Test combining a directive that takes a for loop with one that takes + a regular statement body. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: parallel for) \ + default (parallel) + for (int i = 0; i < N; i++) + x += i; + + /* Test labels inside statement body. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: teams num_teams(512)) \ + when (device={arch("gcn")}: teams num_teams(256)) \ + default (teams num_teams(4)) + { + if (x) + goto l1; + else + goto l2; + l1: ; + l2: ; + } + + /* Test local labels inside statement body. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: teams num_teams(512)) \ + when (device={arch("gcn")}: teams num_teams(256)) \ + default (teams num_teams(4)) + { + //__label__ l1, l2; + + if (x) + goto l1; + else + goto l2; + l1: ; + l2: ; + } + + return 0; +} diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-3.c b/gcc/testsuite/c-c++-common/gomp/metadirective-3.c new file mode 100644 index 00000000000..80c93b1521d --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-3.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ +/* { dg-additional-options "-fdump-tree-optimized" } */ + +#define N 100 + +void f (int x[], int y[], int z[]) +{ + int i; + + #pragma omp target map(to: x, y) map(from: z) + #pragma omp metadirective \ + when (device={arch("nvptx")}: teams loop) \ + default (parallel loop) + for (i = 0; i < N; i++) + z[i] = x[i] * y[i]; +} + +/* The metadirective should be resolved after Gimplification. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(device arch .nvptx.\\):" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c new file mode 100644 index 00000000000..c4b109295db --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c @@ -0,0 +1,40 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 + +#pragma omp declare target +void f(double a[], double x) { + int i; + + #pragma omp metadirective \ + when (construct={target}: distribute parallel for) \ + default (parallel for simd) + for (i = 0; i < N; i++) + a[i] = x * i; +} +#pragma omp end declare target + + int main() +{ + double a[N]; + + #pragma omp target teams map(from: a[0:N]) + f (a, 3.14159); + + /* TODO: This does not execute a version of f with the default clause + active as might be expected. */ + f (a, 2.71828); + + return 0; + } + + /* The metadirective should be resolved during Gimplification. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(construct target.*\\):" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-5.c b/gcc/testsuite/c-c++-common/gomp/metadirective-5.c new file mode 100644 index 00000000000..4a9f1aa85a6 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-5.c @@ -0,0 +1,24 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ + +#define N 100 + +void f (int a[], int flag) +{ + int i; + #pragma omp metadirective \ + when (user={condition(flag)}: \ + target teams distribute parallel for map(from: a[0:N])) \ + default (parallel for) + for (i = 0; i < N; i++) + a[i] = i; +} + +/* The metadirective should be resolved at parse time. */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp target" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp for" 2 "original" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-6.c b/gcc/testsuite/c-c++-common/gomp/metadirective-6.c new file mode 100644 index 00000000000..c77c0065e17 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-6.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-original" } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 100 + +void bar (int a[], int run_parallel, int run_guided) +{ + #pragma omp metadirective \ + when (user={condition(run_parallel)}: parallel) + { + int i; + #pragma omp metadirective \ + when (construct={parallel}, user={condition(run_guided)}: \ + for schedule(guided)) \ + when (construct={parallel}: for schedule(static)) + for (i = 0; i < N; i++) + a[i] = i; + } + } + +/* The outer metadirective should be resolved at parse time. */ +/* The inner metadirective should be resolved during Gimplificiation. */ + +/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 2 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */ +/* { dg-final { scan-tree-dump-times "#pragma omp for" 4 "original" } } */ +/* { dg-final { scan-tree-dump-times "when \\(construct parallel" 4 "original" } } */ +/* { dg-final { scan-tree-dump-times "default:" 2 "original" } } */ + +/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */ diff --git a/gcc/testsuite/gcc.dg/gomp/metadirective-1.c b/gcc/testsuite/gcc.dg/gomp/metadirective-1.c new file mode 100644 index 00000000000..2ac81bfde75 --- /dev/null +++ b/gcc/testsuite/gcc.dg/gomp/metadirective-1.c @@ -0,0 +1,15 @@ +int main (void) +{ + int x, y; + + /* Test nested functions inside statement body. */ + #pragma omp metadirective \ + when (device={arch("nvptx")}: teams num_teams(512)) \ + when (device={arch("gcn")}: teams num_teams(256)) \ + default (teams num_teams(4)) + { + int f (int x) { return x * 3; } + + y = f (x); + } +} diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 new file mode 100644 index 00000000000..aa439fc855e --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-1.f90 @@ -0,0 +1,41 @@ +! { dg-do compile } + +program main + integer, parameter :: N = 10 + integer, dimension(N) :: a + integer, dimension(N) :: b + integer, dimension(N) :: c + integer :: i + + do i = 1, N + a(i) = i * 2 + b(i) = i * 3 + end do + + !$omp metadirective & + !$omp& default (teams loop) & + !$omp& default (parallel loop) ! { dg-error "there can only be one default clause in a metadirective at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective default (xyz) ! { dg-error "Unclassifiable OpenMP directive at .1." } + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp metadirective & + !$omp& default (teams loop) & ! { dg-error "expected 'default' or 'when' at .1." } + !$omp& where (device={arch("nvptx")}: parallel loop) + do i = 1, N + c(i) = a(i) * b(i) + end do + + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (barrier) ! { dg-error "variant directive used in OMP BEGIN METADIRECTIVE at .1. must have a corresponding end directive" } + do i = 1, N + c(i) = a(i) * b(i) + end do + !$omp end metadirective ! { dg-error "Unexpected !OMP END METADIRECTIVE statement at .1." } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 new file mode 100644 index 00000000000..06c324589d0 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-2.f90 @@ -0,0 +1,59 @@ +! { dg-do compile } + +program main + integer, parameter :: N = 100 + integer :: x = 0 + integer :: y = 0 + integer :: i + + ! Test implicit default directive + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: barrier) + x = 1 + + ! Test implicit default directive combined with a directive that takes a + ! do loop. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) + do i = 1, N + x = x + i + end do + + ! Test with multiple standalone directives. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: barrier) & + !$omp& default (flush) + x = 1 + + ! Test combining a standalone directive with one that takes a do loop. + !$omp metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (barrier) + do i = 1, N + x = x + i + end do + + ! Test combining a directive that takes a do loop with one that takes + ! a statement body. + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& default (parallel) + do i = 1, N + x = x + i + end do + !$omp end metadirective + + ! Test labels in the body + !$omp begin metadirective & + !$omp& when (device={arch("nvptx")}: parallel do) & + !$omp& when (device={arch("gcn")}: parallel) + do i = 1, N + x = x + i + if (x .gt. N/2) goto 10 +10 x = x + 1 + goto 20 + x = x + 2 +20 continue + end do + !$omp end metadirective +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 new file mode 100644 index 00000000000..c36a462bf51 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-3.f90 @@ -0,0 +1,34 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-gimple" } +! { dg-additional-options "-fdump-tree-optimized" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: v1, v2) map(from: v3) + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + !$omp end target + end subroutine +end module + +! The metadirective should be resolved after Gimplification. + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } +! { dg-final { scan-tree-dump-times "when \\(device arch .nvptx.\\):" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } +! { dg-final { scan-tree-dump-times "default:" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp loop" 2 "original" } } + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "gimple" } } + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "optimized" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 new file mode 100644 index 00000000000..b82c9ea96d9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-4.f90 @@ -0,0 +1,39 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-gimple" } + +program test + implicit none + integer, parameter :: N = 100 + real :: a(N) + + !$omp target map(from: a) + call f (a, 3.14159) + !$omp end target + + ! TODO: This does not execute a version of f with the default clause + ! active as might be expected. + call f (a, 2.71828) +contains + subroutine f (a, x) + integer :: i + real :: a(N), x + !$omp declare target + + !$omp metadirective & + !$omp& when (construct={target}: distribute parallel do ) & + !$omp& default(parallel do simd) + do i = 1, N + a(i) = x * i + end do + end subroutine +end program + +! The metadirective should be resolved during Gimplification. + +! { dg-final { scan-tree-dump-times "#pragma omp metadirective" 1 "original" } } +! { dg-final { scan-tree-dump-times "when \\(construct target.*\\):" 1 "original" } } +! { dg-final { scan-tree-dump-times "default:" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 new file mode 100644 index 00000000000..03970393eb4 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-5.f90 @@ -0,0 +1,30 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (a, flag) + integer :: a(N) + logical :: flag + integer :: i + + !$omp metadirective & + !$omp& when (user={condition(flag)}: & + !$omp& target teams distribute parallel do map(from: a(1:N))) & + !$omp& default(parallel do) + do i = 1, N + a(i) = i + end do + end subroutine +end module + +! The metadirective should be resolved at parse time, but is currently +! resolved during Gimplification + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp target" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 new file mode 100644 index 00000000000..9b6c371296f --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/metadirective-6.f90 @@ -0,0 +1,31 @@ +! { dg-do compile } +! { dg-additional-options "-fdump-tree-gimple" } + +module test + integer, parameter :: N = 100 +contains + subroutine f (a, run_parallel, run_guided) + integer :: a(N) + logical :: run_parallel, run_guided + integer :: i + + !$omp begin metadirective when(user={condition(run_parallel)}: parallel) + !$omp metadirective & + !$omp& when(construct={parallel}, user={condition(run_guided)}: & + !$omp& do schedule(guided)) & + !$omp& when(construct={parallel}: do schedule(static)) + do i = 1, N + a(i) = i + end do + !$omp end metadirective + end subroutine +end module + +! The outer metadirective should be resolved at parse time, but is +! currently resolved during Gimplification. + +! The inner metadirective should be resolved during Gimplificiation. + +! { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "#pragma omp for" 2 "gimple" } } diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c new file mode 100644 index 00000000000..0de59cbe3d3 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c @@ -0,0 +1,35 @@ +/* { dg-do run } */ + +#define N 100 + +void f (int x[], int y[], int z[]) +{ + int i; + + #pragma omp target map(to: x[0:N], y[0:N]) map(from: z[0:N]) + #pragma omp metadirective \ + when (device={arch("nvptx")}: teams loop) \ + default (parallel loop) + for (i = 0; i < N; i++) + z[i] = x[i] * y[i]; +} + +int main (void) +{ + int x[N], y[N], z[N]; + int i; + + for (i = 0; i < N; i++) + { + x[i] = i; + y[i] = -i; + } + + f (x, y, z); + + for (i = 0; i < N; i++) + if (z[i] != x[i] * y[i]) + return 1; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c new file mode 100644 index 00000000000..cd5c6c5e21a --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c @@ -0,0 +1,41 @@ +/* { dg-do run } */ + +#include + +#define N 100 +#define EPSILON 0.001 + +#pragma omp declare target +void f(double a[], double x) { + int i; + + #pragma omp metadirective \ + when (construct={target}: distribute parallel for) \ + default (parallel for simd) + for (i = 0; i < N; i++) + a[i] = x * i; +} +#pragma omp end declare target + + int main() +{ + double a[N]; + int i; + + #pragma omp target teams map(from: a[0:N]) + f (a, M_PI); + + for (i = 0; i < N; i++) + if (fabs (a[i] - (M_PI * i)) > EPSILON) + return 1; + + /* TODO: This does not execute a version of f with the default clause + active as might be expected. */ + f (a, M_E); + + for (i = 0; i < N; i++) + if (fabs (a[i] - (M_E * i)) > EPSILON) + return 1; + + return 0; + } diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c new file mode 100644 index 00000000000..e31daf2cb64 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c @@ -0,0 +1,34 @@ +/* { dg-do run } */ + +#define N 100 + +int f (int a[], int flag) +{ + int i; + int res = 0; + + #pragma omp metadirective \ + when (user={condition(!flag)}: \ + target teams distribute parallel for \ + map(from: a[0:N]) private(res)) \ + default (parallel for) + for (i = 0; i < N; i++) + { + a[i] = i; + res = 1; + } + + return res; +} + +int main (void) +{ + int a[N]; + + if (f (a, 0)) + return 1; + if (!f (a, 1)) + return 1; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c new file mode 100644 index 00000000000..7fc601eaba6 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c @@ -0,0 +1,52 @@ +/* { dg-do run } */ + +#include + +#define N 100 + +int f (int a[], int run_parallel, int run_static) +{ + int is_parallel = 0; + int is_static = 0; + + #pragma omp metadirective \ + when (user={condition(run_parallel)}: parallel) + { + int i; + + if (omp_in_parallel ()) + is_parallel = 1; + + #pragma omp metadirective \ + when (construct={parallel}, user={condition(!run_static)}: \ + for schedule(guided) private(is_static)) \ + when (construct={parallel}: for schedule(static)) + for (i = 0; i < N; i++) + { + a[i] = i; + is_static = 1; + } + } + + return (is_parallel << 1) | is_static; +} + +int main (void) +{ + int a[N]; + + /* is_static is always set if run_parallel is false. */ + if (f (a, 0, 0) != 1) + return 1; + + if (f (a, 0, 1) != 1) + return 1; + + if (f (a, 1, 0) != 2) + return 1; + + if (f (a, 1, 1) != 3) + return 1; + + return 0; +} diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 new file mode 100644 index 00000000000..9f6a07459e0 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-1.f90 @@ -0,0 +1,33 @@ +! { dg-do run } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: x(N), y(N), z(N) + integer :: i + + do i = 1, N + x(i) = i; + y(i) = -i; + end do + + call f (x, y, z) + + do i = 1, N + if (z(i) .ne. x(i) * y(i)) stop 1 + end do +contains + subroutine f (x, y, z) + integer :: x(N), y(N), z(N) + + !$omp target map (to: x, y) map(from: z) + !$omp metadirective & + !$omp& when(device={arch("nvptx")}: teams loop) & + !$omp& default(parallel loop) + do i = 1, N + z(i) = x(i) * y(i) + enddo + !$omp end target + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 new file mode 100644 index 00000000000..32017a00077 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-2.f90 @@ -0,0 +1,40 @@ +! { dg-do run } + +program test + implicit none + integer, parameter :: N = 100 + real, parameter :: PI_CONST = 3.14159 + real, parameter :: E_CONST = 2.71828 + real, parameter :: EPSILON = 0.001 + integer :: i + real :: a(N) + + !$omp target map(from: a) + call f (a, PI_CONST) + !$omp end target + + do i = 1, N + if (abs (a(i) - (PI_CONST * i)) .gt. EPSILON) stop 1 + end do + + ! TODO: This does not execute a version of f with the default clause + ! active as might be expected. + call f (a, E_CONST) + + do i = 1, N + if (abs (a(i) - (E_CONST * i)) .gt. EPSILON) stop 2 + end do +contains + subroutine f (a, x) + integer :: i + real :: a(N), x + !$omp declare target + + !$omp metadirective & + !$omp& when (construct={target}: distribute parallel do ) & + !$omp& default(parallel do simd) + do i = 1, N + a(i) = x * i + end do + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 new file mode 100644 index 00000000000..693c40bca5a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-3.f90 @@ -0,0 +1,29 @@ +! { dg-do run } + +program test + implicit none + + integer, parameter :: N = 100 + integer :: a(N) + integer :: res + + if (f (a, .false.)) stop 1 + if (.not. f (a, .true.)) stop 2 +contains + logical function f (a, flag) + integer :: a(N) + logical :: flag + logical :: res = .false. + integer :: i + f = .false. + !$omp metadirective & + !$omp& when (user={condition(.not. flag)}: & + !$omp& target teams distribute parallel do & + !$omp& map(from: a(1:N)) private(res)) & + !$omp& default(parallel do) + do i = 1, N + a(i) = i + f = .true. + end do + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 new file mode 100644 index 00000000000..04fdf61489c --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/metadirective-4.f90 @@ -0,0 +1,46 @@ +! { dg-do run } + +program test + use omp_lib + + implicit none + integer, parameter :: N = 100 + integer :: a(N) + logical :: is_parallel, is_static + + ! is_static is always set if run_parallel is false. + call f (a, .false., .false., is_parallel, is_static) + if (is_parallel .or. .not. is_static) stop 1 + + call f (a, .false., .true., is_parallel, is_static) + if (is_parallel .or. .not. is_static) stop 2 + + call f (a, .true., .false., is_parallel, is_static) + if (.not. is_parallel .or. is_static) stop 3 + + call f (a, .true., .true., is_parallel, is_static) + if (.not. is_parallel .or. .not. is_static) stop 4 +contains + subroutine f (a, run_parallel, run_static, is_parallel, is_static) + integer :: a(N) + logical, intent(in) :: run_parallel, run_static + logical, intent(out) :: is_parallel, is_static + integer :: i + + is_parallel = .false. + is_static = .false. + + !$omp begin metadirective when(user={condition(run_parallel)}: parallel) + if (omp_in_parallel ()) is_parallel = .true. + + !$omp metadirective & + !$omp& when(construct={parallel}, user={condition(.not. run_static)}: & + !$omp& do schedule(guided) private(is_static)) & + !$omp& when(construct={parallel}: do schedule(static)) + do i = 1, N + a(i) = i + is_static = .true. + end do + !$omp end metadirective + end subroutine +end program