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 <san...@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 <jul...@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 <san...@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 <jul...@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
