From: Sandra Loosemore <sandra@codesourcery.com>
To: <gcc-patches@gcc.gnu.org>
Cc: <jakub@rethat.com>, <julian@codesourcery.com>, <kcy@codesourcery.com>
Subject: [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective"
Date: Fri, 18 Aug 2023 11:39:35 -0600 [thread overview]
Message-ID: <20230818173938.430758-2-sandra@codesourcery.com> (raw)
In-Reply-To: <20230818173938.430758-1-sandra@codesourcery.com>
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 <sandra@codesourcery.com>
+
+ * 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 <julian@codesourcery.com>
* 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 <sandra@codesourcery.com>
+
+ * 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 <julian@codesourcery.com>
* 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
next prev parent reply other threads:[~2023-08-18 17:40 UTC|newest]
Thread overview: 4+ messages / expand[flat|nested] mbox.gz Atom feed top
2023-08-18 17:39 [OG13, committed 0/3] C++ attribute syntax fixes/testcases Sandra Loosemore
2023-08-18 17:39 ` Sandra Loosemore [this message]
2023-08-18 17:39 ` [OG13, committed 2/3] OpenMP: C++ attribute syntax fixes/testcases for "declare mapper" Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 3/3] OpenMP: C++ attribute syntax fixes/testcases for loop transformations Sandra Loosemore
Reply instructions:
You may reply publicly to this message via plain-text email
using any one of the following methods:
* Save the following mbox file, import it into your mail client,
and reply-to-all from there: mbox
Avoid top-posting and favor interleaved quoting:
https://en.wikipedia.org/wiki/Posting_style#Interleaved_style
* Reply using the --to, --cc, and --in-reply-to
switches of git-send-email(1):
git send-email \
--in-reply-to=20230818173938.430758-2-sandra@codesourcery.com \
--to=sandra@codesourcery.com \
--cc=gcc-patches@gcc.gnu.org \
--cc=jakub@rethat.com \
--cc=julian@codesourcery.com \
--cc=kcy@codesourcery.com \
/path/to/YOUR_REPLY
https://kernel.org/pub/software/scm/git/docs/git-send-email.html
* If your mail client supports setting the In-Reply-To header
via mailto: links, try the mailto: link
Be sure your reply has a Subject: header at the top and a blank line
before the message body.
This is a public inbox, see mirroring instructions
for how to clone and mirror all data and code used for this inbox;
as well as URLs for read-only IMAP folder(s) and NNTP newsgroup(s).