From patchwork Sat Jan 6 18:52:54 2024 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 7bit X-Patchwork-Submitter: Sandra Loosemore X-Patchwork-Id: 83464 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 5917C3857B8B for ; Sat, 6 Jan 2024 18:56:29 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa1.mentor.iphmx.com (esa1.mentor.iphmx.com [68.232.129.153]) by sourceware.org (Postfix) with ESMTPS id 294143858D32 for ; Sat, 6 Jan 2024 18:54:38 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 294143858D32 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 294143858D32 Authentication-Results: server2.sourceware.org; arc=none smtp.remote-ip=68.232.129.153 ARC-Seal: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1704567287; cv=none; b=qDIJXJkj1o3vz3atTStCMqTKaz0B7JnljidNxQCbEV3Q33pIUV4NaOVdwQ6pB6YJuP6HFMqaEpW8m0y4orbhpNyosGSRQWFTyX9tlk3GpdJoADZqag8YU1qJ3Qsb2wfx5hjOEH8ROcRM4MnjvlVoMid26B4FAzvDWEoXqwr1K3g= ARC-Message-Signature: i=1; a=rsa-sha256; d=sourceware.org; s=key; t=1704567287; c=relaxed/simple; bh=njV8FnwyTPM0Bl1sB1v/d85yQdoq62KgAHKBA4IK+2k=; h=From:To:Subject:Date:Message-ID:MIME-Version; b=ZInKtIY5Nu3nVhtgOTz+KA25qaOrqSR1TiiGEFXA712JMJPouXNCd3y3NDldeNWxUJLAFMoSwuR/l6oYpFxQDpKqGX41K5wJXdMjRfBqNh0p3CgMWpSEh4fUA369wrg6d1CsofjVQFVbeEH2n2fxfoBMB0+P+Vl9sqzITTwgu7U= ARC-Authentication-Results: i=1; server2.sourceware.org X-CSE-ConnectionGUID: A0ufJUGjQOqi70e/GHZ7eg== X-CSE-MsgGUID: gEIBq2qXTZ+qc6So51YHTA== X-IronPort-AV: E=Sophos;i="6.04,337,1695715200"; d="scan'208";a="30827348" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa1.mentor.iphmx.com with ESMTP; 06 Jan 2024 10:54:37 -0800 IronPort-SDR: LgZLriPQ2ZnVOJzvTVJYePjBloGeK/o+ijIXk9FKasHN7Xhll4OO5zVBmvf1ik6dVDE5/kYcBd gfZ1lIRZSBOCapsrQM83c/7ksRmXMxS/SeWKQL5zH55KeX5+UkS7rFzDXAuLBDBjfaDCX01X54 V7Mxm2E/nO8KXa+qeGJfDlQ7AiGFNGd+0rT05Jt18YDbjQp57Ajk3eKQd/Wp3LvCmaCzazQiqa Nz3PjtZnjsV0yDZoDuerjXidBABIwe4VsrzRxGPSp4fvm7FnDDjk1463Mi3a4w9wJ8OYpYmX5Z pZw= From: Sandra Loosemore To: CC: , , Subject: [PATCH 6/8] OpenMP: common c/c++ testcases for metadirectives Date: Sat, 6 Jan 2024 11:52:54 -0700 Message-ID: <20240106185257.126445-7-sandra@codesourcery.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20240106185257.126445-1-sandra@codesourcery.com> References: <20240106185257.126445-1-sandra@codesourcery.com> MIME-Version: 1.0 X-ClientProxiedBy: svr-orw-mbx-10.mgc.mentorg.com (147.34.90.210) To svr-orw-mbx-13.mgc.mentorg.com (147.34.90.213) X-Spam-Status: No, score=-10.1 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+patchwork=sourceware.org@gcc.gnu.org From: Kwok Cheung Yeung gcc/testsuite/ChangeLog * 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. * c-c++-common/gomp/metadirective-7.c: New. * c-c++-common/gomp/metadirective-8.c: New. * c-c++-common/gomp/metadirective-construct.c: New. * c-c++-common/gomp/metadirective-device.c: New. * c-c++-common/gomp/metadirective-no-score.c: New. * c-c++-common/gomp/metadirective-target-device.c: New. libgomp/ChangeLog * 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.c-c++-common/metadirective-5.c: New. Co-Authored-By: Sandra Loosemore --- .../c-c++-common/gomp/metadirective-1.c | 52 +++++ .../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 +++ .../c-c++-common/gomp/metadirective-7.c | 31 +++ .../c-c++-common/gomp/metadirective-8.c | 16 ++ .../gomp/metadirective-construct.c | 177 ++++++++++++++++++ .../c-c++-common/gomp/metadirective-device.c | 147 +++++++++++++++ .../gomp/metadirective-no-score.c | 95 ++++++++++ .../gomp/metadirective-target-device.c | 147 +++++++++++++++ .../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.c-c++-common/metadirective-5.c | 46 +++++ 17 files changed, 1073 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/c-c++-common/gomp/metadirective-7.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-8.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-construct.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-device.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c 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.c-c++-common/metadirective-5.c 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..37b56237531 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c @@ -0,0 +1,52 @@ +/* { dg-do compile } */ + +#define N 100 + +void f (int a[], int b[], int c[]) +{ + int i; + + #pragma omp metadirective \ + default (teams loop) \ + default (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + otherwise (teams loop) \ + default (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + #pragma omp metadirective \ + otherwise (teams loop) \ + otherwise (parallel loop) /* { dg-error "too many 'otherwise' or 'default' clauses in 'metadirective'" } */ + 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 "'where' is not valid for 'metadirective'" } */ + 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 ':' 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 variants of a 'metadirective' before 'default'" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + /* Test improperly nested metadirectives - even though the second + metadirective resolves to 'omp nothing', that is not the same as there + being literally nothing there. */ + #pragma omp metadirective \ + when (implementation={vendor("gnu")}: parallel for) + #pragma omp metadirective \ + when (implementation={vendor("cray")}: parallel for) + /* { dg-error "for statement expected before '#pragma'" "" { target c } .-2 } */ + /* { dg-error "'#pragma' is not allowed here" "" { target c++ } .-3 } */ + 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..7a2818dd710 --- /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..5fa86561a23 --- /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); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ + + 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..26c7a008e95 --- /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/c-c++-common/gomp/metadirective-7.c b/gcc/testsuite/c-c++-common/gomp/metadirective-7.c new file mode 100644 index 00000000000..cf695aa24cb --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-7.c @@ -0,0 +1,31 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 256 + +void f (int a[], int num) +{ + int i; + + #pragma omp metadirective \ + when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: \ + target parallel for map(tofrom: a[0:N])) \ + when (target_device={device_num(num), kind("gpu"), \ + arch("amdgcn"), isa("gfx906")}: \ + target parallel for) \ + when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: \ + parallel for) + for (i = 0; i < N; i++) + a[i] += i; + + #pragma omp metadirective \ + when (target_device={kind("gpu"), arch("nvptx")}: \ + target parallel for map(tofrom: a[0:N])) + for (i = 0; i < N; i++) + a[i] += i; +} + +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu\"\\\[0\\\], &\"amdgcn\"\\\[0\\\], &\"gfx906\"\\\[0\\\]\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(num, &\"cpu\"\\\[0\\\], &\"x86_64\"\\\[0\\\], 0B\\)" "gimple" } } */ +/* { dg-final { scan-tree-dump "__builtin_GOMP_evaluate_target_device \\(-1, &\"gpu\"\\\[0\\\], &\"nvptx\"\\\[0\\\], 0B\\)" "gimple" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-8.c b/gcc/testsuite/c-c++-common/gomp/metadirective-8.c new file mode 100644 index 00000000000..c7d9c31ed73 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-8.c @@ -0,0 +1,16 @@ +/* { dg-do compile } */ + +#define N 256 + +void f () +{ + int i; + int a[N]; + + #pragma omp metadirective \ + when( device={kind(nohost)}: nothing ) \ + when( device={arch("nvptx")}: nothing) \ + default( parallel for) + for (i = 0; i < N; i++) + a[i] = i; +} diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c b/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c new file mode 100644 index 00000000000..a61966aa899 --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c @@ -0,0 +1,177 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */ + +#include + +static void +init (int n, double *a) +{ + for (int i = 0; i < n; i++) + a[i] = (double) i; +} + +static void +check (int n, double *a, double s) +{ + for (int i = 0; i < n; i++) + if (a[i] != (double) i * s) + abort (); +} + +typedef void (transform_fn) (int, double *, double); + +static void doit (transform_fn *f, int n, double *a, double s) +{ + init (n, a); + (*f) (n, a, s); + check (n, a, s); +} + +/* Check various combinations for enforcing correct ordering of + construct matches. */ +static void +f1 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={target} \ + : for) \ + default (error at(execution) message("f1 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f2 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={teams, parallel} \ + : for) \ + default (error at(execution) message("f2 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f3 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={target, teams, parallel} \ + : for) \ + default (error at(execution) message("f3 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f4 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={target, parallel} \ + : for) \ + default (error at(execution) message("f4 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f5 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={target, teams} \ + : for) \ + default (error at(execution) message("f5 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Next batch is for things where the construct doesn't match the context. */ +static void +f6 (int n, double* a, double s) +{ +#pragma omp target +#pragma omp teams +#pragma omp metadirective \ + when (construct={parallel} \ + : error at(execution) message("f6 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f7 (int n, double* a, double s) +{ +#pragma omp target +#pragma omp teams +#pragma omp metadirective \ + when (construct={target, parallel} \ + : error at(execution) message("f7 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +static void +f8 (int n, double* a, double s) +{ +#pragma omp target +#pragma omp teams +#pragma omp metadirective \ + when (construct={parallel, target} \ + : error at(execution) message("f8 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Next test choosing the best alternative when there are multiple + matches. */ +static void +f9 (int n, double* a, double s) +{ +#pragma omp target teams +#pragma omp parallel +#pragma omp metadirective \ + when (construct={teams, parallel} \ + : error at(execution) message("f9 match incorrect 1")) \ + when (construct={target, teams, parallel} \ + : for) \ + when (construct={target, teams} \ + : error at(execution) message("f9 match incorrect 2")) \ + default (error at(execution) message("f9 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Note there are no tests for the matching the extended simd clause + syntax, which is only useful for "declare variant". */ + +#define N 10 +#define S 2.0 + +int main (void) +{ + double a[N]; + doit (f1, N, a, S); + doit (f2, N, a, S); + doit (f3, N, a, S); + doit (f4, N, a, S); + doit (f5, N, a, S); + doit (f6, N, a, S); + doit (f7, N, a, S); + doit (f8, N, a, S); + doit (f9, N, a, S); +} + +/* All the error calls should be optimized away. */ +/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-device.c b/gcc/testsuite/c-c++-common/gomp/metadirective-device.c new file mode 100644 index 00000000000..9261331260c --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-device.c @@ -0,0 +1,147 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */ +/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=sse -msse" { target x86_64-*-* } } */ + +#include + +static void +init (int n, double *a) +{ + for (int i = 0; i < n; i++) + a[i] = (double) i; +} + +static void +check (int n, double *a, double s) +{ + for (int i = 0; i < n; i++) + if (a[i] != (double) i * s) + abort (); +} + +typedef void (transform_fn) (int, double *, double); + +static void doit (transform_fn *f, int n, double *a, double s) +{ + init (n, a); + (*f) (n, a, s); + check (n, a, s); +} + +/* Check kind=host matches (with offloading disabled). */ +static void +f1 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={kind(host)} \ + : parallel for) \ + default (error at(execution) message("f1 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Check kind=nohost does not match (with offloading disabled). */ +static void +f2 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={kind(nohost)} \ + : error at(execution) message("f2 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Check arch. Either DEVICE_ARCH is defined by command-line option, + or we know it is not x86_64. */ +#ifdef DEVICE_ARCH +static void +f3 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch(DEVICE_ARCH)} \ + : parallel for) \ + default (error at(execution) message("f3 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f3 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch("x86_64")} \ + : error at(execution) message("f3 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +/* Check both kind and arch together. */ +#ifdef DEVICE_ARCH +static void +f4 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch(DEVICE_ARCH), kind(host)} \ + : parallel for) \ + default (error at(execution) message("f4 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f4 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch("x86_64"), kind(host)} \ + : error at(execution) message("f4 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +/* Check kind, arch, and ISA together. */ +#if defined(DEVICE_ARCH) && defined(DEVICE_ISA) +static void +f5 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)} \ + : parallel for) \ + default (error at(execution) message("f5 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f5 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (device={arch("x86_64"), kind(host), isa("sse")} \ + : error at(execution) message("f5 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +#define N 10 +#define S 2.0 + +int main (void) +{ + double a[N]; + doit (f1, N, a, S); + doit (f2, N, a, S); + doit (f3, N, a, S); + doit (f4, N, a, S); + doit (f5, N, a, S); +} + +/* All the metadirectives involving the device selector should be + fully resolved and the error calls optimized away. */ + +/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */ diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c b/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c new file mode 100644 index 00000000000..1f1053eaffa --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c @@ -0,0 +1,95 @@ +/* { dg-do compile { target x86_64-*-* } } */ +/* { dg-additional-options "-foffload=disable" } */ + +/* This test is expected to fail with compile-time errors: + "A trait-score cannot be specified in traits from the construct, + device or target_device trait-selector-sets." */ + +/* Define this to avoid dependence on libgomp header files. */ + +#define omp_initial_device -1 + +void +f1 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (device={kind (score(5) : host)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f2 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-2 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f3 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (device={kind (host), arch (score(6) : x86_64), \ + isa (score(7): avx512f)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 } */ + /* { dg-error ".score. cannot be specified in traits in the .device. trait-selector-set" "" { target *-*-*} .-3 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f4 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (target_device={device_num (score(42) : omp_initial_device), \ + kind (host)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f5 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (target_device={device_num(omp_initial_device), \ + kind (score(5) : host)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f6 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (target_device={device_num(omp_initial_device), kind (host), \ + arch (score(6) : x86_64), isa (avx512f)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-2 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +void +f7 (int n, double *a, double s) +{ +#pragma omp metadirective \ + when (target_device={device_num(omp_initial_device), kind (host), \ + arch (score(6) : x86_64), \ + isa (score(7): avx512f)} \ + : parallel for) + /* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */ + /* { dg-error ".score. cannot be specified in traits in the .target_device. trait-selector-set" "" { target *-*-*} .-3 } */ + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c new file mode 100644 index 00000000000..335db416a4b --- /dev/null +++ b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c @@ -0,0 +1,147 @@ +/* { dg-do compile } */ +/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */ +/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=mmx -mmmx" { target x86_64-*-* } } */ + +#include + +static void +init (int n, double *a) +{ + for (int i = 0; i < n; i++) + a[i] = (double) i; +} + +static void +check (int n, double *a, double s) +{ + for (int i = 0; i < n; i++) + if (a[i] != (double) i * s) + abort (); +} + +typedef void (transform_fn) (int, double *, double); + +static void doit (transform_fn *f, int n, double *a, double s) +{ + init (n, a); + (*f) (n, a, s); + check (n, a, s); +} + +/* Check kind=host matches (with offloading disabled). */ +static void +f1 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={kind(host)} \ + : parallel for) \ + default (error at(execution) message("f1 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Check kind=nohost does not match (with offloading disabled). */ +static void +f2 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={kind(nohost)} \ + : error at(execution) message("f2 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} + +/* Check arch. Either DEVICE_ARCH is defined by command-line option, + or we know it is not x86_64. */ +#ifdef DEVICE_ARCH +static void +f3 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch(DEVICE_ARCH)} \ + : parallel for) \ + default (error at(execution) message("f3 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f3 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch("x86_64")} \ + : error at(execution) message("f3 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +/* Check both kind and arch together. */ +#ifdef DEVICE_ARCH +static void +f4 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch(DEVICE_ARCH), kind(host)} \ + : parallel for) \ + default (error at(execution) message("f4 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f4 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch("x86_64"), kind(host)} \ + : error at(execution) message("f4 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +/* Check kind, arch, and ISA together. */ +#if defined(DEVICE_ARCH) && defined(DEVICE_ISA) +static void +f5 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)} \ + : parallel for) \ + default (error at(execution) message("f5 match failed")) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#else +static void +f5 (int n, double* a, double s) +{ +#pragma omp metadirective \ + when (target_device={arch("x86_64"), kind(host), isa("mmx")} \ + : error at(execution) message("f5 match failed")) \ + default (parallel for) + for (int i = 0; i < n; i++) + a[i] = a[i] * s; +} +#endif + +#define N 10 +#define S 2.0 + +int main (void) +{ + double a[N]; + doit (f1, N, a, S); + doit (f2, N, a, S); + doit (f3, N, a, S); + doit (f4, N, a, S); + doit (f5, N, a, S); +} + +/* Since the target_device selector is dynamic, none of + error checks tests can be optimized away. */ + +/* { dg-final { scan-tree-dump-times "__builtin_GOMP_error" 5 "optimized" } } */ 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..55a6098e525 --- /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); /* { dg-warning "direct calls to an offloadable function containing metadirectives with a 'construct={target}' selector may produce unexpected results" } */ + + 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.c-c++-common/metadirective-5.c b/libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c new file mode 100644 index 00000000000..e8ab7ccb166 --- /dev/null +++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c @@ -0,0 +1,46 @@ +/* { dg-do run } */ + +#define N 100 + +#include +#include + +int f(int a[], int num) +{ + int on_device = 0; + int i; + + #pragma omp metadirective \ + when (target_device={device_num(num), kind("gpu")}: \ + target parallel for map(to: a[0:N]), map(from: on_device)) \ + default (parallel for private (on_device)) + for (i = 0; i < N; i++) + { + a[i] += i; + on_device = 1; + } + + return on_device; +} + +int main (void) +{ + int a[N]; + int on_device_count = 0; + int i; + + for (i = 0; i < N; i++) + a[i] = i; + + for (i = 0; i <= omp_get_num_devices (); i++) + on_device_count += f (a, i); + + if (on_device_count != omp_get_num_devices ()) + return 1; + + for (i = 0; i < N; i++) + if (a[i] != 2 * i) + return 2; + + return 0; +}