From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id EE0B0385084C for ; Fri, 18 Aug 2023 17:40:05 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org EE0B0385084C Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="6.01,183,1684828800"; d="scan'208";a="14812570" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 18 Aug 2023 09:40:02 -0800 IronPort-SDR: hx0DrMv3ebYg1frpp0THhcuKKEg7Nu3aOkQGS/6cBc5ug9dflQuAP1TuJPIeQcFDsHuORJhcQk kzX8W4K9XctAyJr4EGDE9AvLVDvX6VnffIOS915SnYEVJrgNVh0S6CFR76clBkQsf5K5WuaIh+ dx0KWxA89ex2VhxQ50OVxZs4XmC+c7X4/WVXin335Ro7mjzfSiSVZNcnesOygHDBi9are1bjTd P5wWdj/Lq41VUXGF6YnmFM36q/yMv2E1lyQisErVA4krYIt4V1d725w+afydT6rfYGsT0ecoio Li0= From: Sandra Loosemore To: CC: , , Subject: [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective" Date: Fri, 18 Aug 2023 11:39:35 -0600 Message-ID: <20230818173938.430758-2-sandra@codesourcery.com> X-Mailer: git-send-email 2.34.1 In-Reply-To: <20230818173938.430758-1-sandra@codesourcery.com> References: <20230818173938.430758-1-sandra@codesourcery.com> MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain 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=-9.9 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.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: gcc/cp/ChangeLog: * parser.cc (analyze_metadirective_body): Handle CPP_PRAGMA and CPP_PRAGMA_EOL. (cp_parser_omp_metadirective): Allow comma between clauses. gcc/testsuite/ChangeLog: * g++.dg/gomp/attrs-metadirective-1.C: New file. * g++.dg/gomp/attrs-metadirective-2.C: New file. * g++.dg/gomp/attrs-metadirective-3.C: New file. * g++.dg/gomp/attrs-metadirective-4.C: New file. * g++.dg/gomp/attrs-metadirective-5.C: New file. * g++.dg/gomp/attrs-metadirective-6.C: New file. * g++.dg/gomp/attrs-metadirective-7.C: New file. * g++.dg/gomp/attrs-metadirective-8.C: New file. --- gcc/cp/ChangeLog.omp | 6 ++ gcc/cp/parser.cc | 16 ++++ gcc/testsuite/ChangeLog.omp | 11 +++ .../g++.dg/gomp/attrs-metadirective-1.C | 40 ++++++++++ .../g++.dg/gomp/attrs-metadirective-2.C | 74 +++++++++++++++++++ .../g++.dg/gomp/attrs-metadirective-3.C | 31 ++++++++ .../g++.dg/gomp/attrs-metadirective-4.C | 41 ++++++++++ .../g++.dg/gomp/attrs-metadirective-5.C | 24 ++++++ .../g++.dg/gomp/attrs-metadirective-6.C | 31 ++++++++ .../g++.dg/gomp/attrs-metadirective-7.C | 31 ++++++++ .../g++.dg/gomp/attrs-metadirective-8.C | 16 ++++ 11 files changed, 321 insertions(+) create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp index e146f57d57d..6e154ea3426 100644 --- a/gcc/cp/ChangeLog.omp +++ b/gcc/cp/ChangeLog.omp @@ -1,3 +1,9 @@ +2023-08-18 Sandra Loosemore + + * parser.cc (analyze_metadirective_body): Handle CPP_PRAGMA and + CPP_PRAGMA_EOL. + (cp_parser_omp_metadirective): Allow comma between clauses. + 2023-08-10 Julian Brown * parser.cc (cp_parser_omp_var_list_no_open): Support array-shaping diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc index cbbef6470d5..84ec0de6c69 100644 --- a/gcc/cp/parser.cc +++ b/gcc/cp/parser.cc @@ -49458,6 +49458,7 @@ analyze_metadirective_body (cp_parser *parser, int bracket_depth = 0; bool in_case = false; bool in_label_decl = false; + cp_token *pragma_tok = NULL; while (1) { @@ -49501,6 +49502,19 @@ analyze_metadirective_body (cp_parser *parser, /* Local label declarations are terminated by a semicolon. */ in_label_decl = false; goto add; + case CPP_PRAGMA: + parser->lexer->in_pragma = true; + pragma_tok = token; + goto add; + case CPP_PRAGMA_EOL: + /* C++ attribute syntax for OMP directives lexes as a pragma, + but we must reset the associated lexer state when we reach + the end in order to get the tokens for the statement that + come after it. */ + tokens.safe_push (*token); + cp_parser_skip_to_pragma_eol (parser, pragma_tok); + pragma_tok = NULL; + continue; default: add: tokens.safe_push (*token); @@ -49541,6 +49555,8 @@ cp_parser_omp_metadirective (cp_parser *parser, cp_token *pragma_tok, while (cp_lexer_next_token_is_not (parser->lexer, CPP_PRAGMA_EOL)) { + if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA)) + cp_lexer_consume_token (parser->lexer); if (cp_lexer_next_token_is_not (parser->lexer, CPP_NAME) && cp_lexer_next_token_is_not (parser->lexer, CPP_KEYWORD)) { diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp index b8780d17841..a9b4ac3d0a7 100644 --- a/gcc/testsuite/ChangeLog.omp +++ b/gcc/testsuite/ChangeLog.omp @@ -1,3 +1,14 @@ +2023-08-18 Sandra Loosemore + + * g++.dg/gomp/attrs-metadirective-1.C: New file. + * g++.dg/gomp/attrs-metadirective-2.C: New file. + * g++.dg/gomp/attrs-metadirective-3.C: New file. + * g++.dg/gomp/attrs-metadirective-4.C: New file. + * g++.dg/gomp/attrs-metadirective-5.C: New file. + * g++.dg/gomp/attrs-metadirective-6.C: New file. + * g++.dg/gomp/attrs-metadirective-7.C: New file. + * g++.dg/gomp/attrs-metadirective-8.C: New file. + 2023-08-10 Julian Brown * c-c++-common/gomp/declare-mapper-17.c: New test. diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C new file mode 100644 index 00000000000..22edd257084 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-1.C @@ -0,0 +1,40 @@ +// { dg-do compile { target c++11 } } + +#define N 100 + +void f (int a[], int b[], int c[]) +{ + int i; + + [[omp::directive (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]; + + [[omp::directive (metadirective + default (bad_directive))]] /* { dg-error "unknown directive name before '\\)' token" } */ + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; + + [[omp::directive (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]; + + [[omp::directive (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]; + + [[omp::directive (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]; + + /* Test improperly nested metadirectives - even though the second + metadirective resolves to 'omp nothing', that is not the same as there + being literally nothing there. */ + [[omp::directive (metadirective + when (implementation={vendor("gnu")}: parallel for))]] + [[omp::directive (metadirective /* { dg-error "'#pragma' is not allowed here" } */ + when (implementation={vendor("cray")}: parallel for))]] + for (i = 0; i < N; i++) c[i] = a[i] * b[i]; +} diff --git a/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C new file mode 100644 index 00000000000..44c87df1776 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-2.C @@ -0,0 +1,74 @@ +// { dg-do compile { target c++11 } } + +#define N 100 + +int main (void) +{ + int x = 0; + int y = 0; + + /* Test implicit default (nothing). */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: barrier))]] + x = 1; + + /* Test with multiple standalone directives. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: barrier), + default (flush))]] + x = 1; + + /* Test combining a standalone directive with one that takes a statement + body. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: parallel), + default (barrier))]] + x = 1; + + /* Test combining a standalone directive with one that takes a for loop. */ + [[omp::directive (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. */ + [[omp::directive (metadirective, + when (device={arch("nvptx")}: parallel for), + default (parallel))]] + for (int i = 0; i < N; i++) + x += i; + + /* Test labels inside statement body. */ + [[omp::directive (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. */ + [[omp::directive (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/g++.dg/gomp/attrs-metadirective-3.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C new file mode 100644 index 00000000000..05b50f547c9 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-3.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { 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; + + [[omp::sequence (directive (target map(to: x, y) map(from: z)), + directive (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/g++.dg/gomp/attrs-metadirective-4.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C new file mode 100644 index 00000000000..fccc5d4c36d --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-4.C @@ -0,0 +1,41 @@ +// { dg-do compile { target c++11 } } + +/* { 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; + + [[omp::directive (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/g++.dg/gomp/attrs-metadirective-5.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C new file mode 100644 index 00000000000..1a9cee15be3 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-5.C @@ -0,0 +1,24 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-original" } */ + +#define N 100 + +void f (int a[], int flag) +{ + int i; + [[omp::directive (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/g++.dg/gomp/attrs-metadirective-6.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C new file mode 100644 index 00000000000..fbfa334fcc7 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-6.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { 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) +{ + [[omp::directive (metadirective + when (user={condition(run_parallel)}: parallel))]] + { + int i; + [[omp::directive (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/g++.dg/gomp/attrs-metadirective-7.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C new file mode 100644 index 00000000000..161bbf86bd5 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-7.C @@ -0,0 +1,31 @@ +// { dg-do compile { target c++11 } } +/* { dg-additional-options "-fdump-tree-gimple" } */ + +#define N 256 + +void f (int a[], int num) +{ + int i; + + [[omp::directive (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; + + [[omp::directive (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/g++.dg/gomp/attrs-metadirective-8.C b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C new file mode 100644 index 00000000000..dcb3c365b80 --- /dev/null +++ b/gcc/testsuite/g++.dg/gomp/attrs-metadirective-8.C @@ -0,0 +1,16 @@ +// { dg-do compile { target c++11 } } + +#define N 256 + +void f () +{ + int i; + int a[N]; + + [[omp::directive (metadirective + when( device={kind(nohost)}: nothing ) + when( device={arch("nvptx")}: nothing) + default( parallel for))]] + for (i = 0; i < N; i++) + a[i] = i; +} -- 2.31.1