* [OG13, committed 0/3] C++ attribute syntax fixes/testcases
@ 2023-08-18 17:39 Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective" Sandra Loosemore
` (2 more replies)
0 siblings, 3 replies; 4+ messages in thread
From: Sandra Loosemore @ 2023-08-18 17:39 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub, julian, kcy
I've had a task item to ensure that g++ accepts the standard C++
attribute syntax form for all (currently-implemented) OpenMP 5.1
directives, and that there are tests to verify this. I used some
scripting to scan for existing testcases given a list of the
directives, which I extracted from the reference card on the OpenMP
web site. It looked to me that on mainline all the supported
directives had tests already, but on OG13 I found that
"metadirective", "declare mapper", and the loop transforms "tile" and
"unroll" had no tests, and on further investigation all of them had
bugs, too.
I didn't manually examine all the existing tests for other directives,
BTW, but the ones I spot-checked seem to have good coverage. The new
tests are mostly just adapted from a subset of existing pragma-syntax
tests.
-Sandra
Sandra Loosemore (3):
OpenMP: C++ attribute syntax fixes/testcases for "metadirective"
OpenMP: C++ attribute syntax fixes/testcases for "declare mapper"
OpenMP: C++ attribute syntax fixes/testcases for loop transformations
gcc/c-family/ChangeLog.omp | 4 +
gcc/c-family/c-omp.cc | 4 +-
gcc/cp/ChangeLog.omp | 26 +++
gcc/cp/parser.cc | 87 ++++++---
gcc/testsuite/ChangeLog.omp | 30 +++
.../g++.dg/gomp/attrs-declare-mapper-3.C | 31 ++++
.../g++.dg/gomp/attrs-declare-mapper-4.C | 74 ++++++++
.../g++.dg/gomp/attrs-declare-mapper-5.C | 26 +++
.../g++.dg/gomp/attrs-declare-mapper-6.C | 22 +++
.../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 ++
.../gomp/loop-transforms/attrs-tile-1.C | 164 +++++++++++++++++
.../gomp/loop-transforms/attrs-tile-2.C | 174 ++++++++++++++++++
.../gomp/loop-transforms/attrs-tile-3.C | 111 +++++++++++
.../gomp/loop-transforms/attrs-unroll-1.C | 135 ++++++++++++++
.../gomp/loop-transforms/attrs-unroll-2.C | 81 ++++++++
.../gomp/loop-transforms/attrs-unroll-3.C | 20 ++
.../loop-transforms/attrs-unroll-inner-1.C | 15 ++
.../loop-transforms/attrs-unroll-inner-2.C | 29 +++
.../loop-transforms/attrs-unroll-inner-3.C | 71 +++++++
26 files changed, 1366 insertions(+), 26 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-4.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-5.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-6.C
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
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C
--
2.31.1
^ permalink raw reply [flat|nested] 4+ messages in thread
* [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective"
2023-08-18 17:39 [OG13, committed 0/3] C++ attribute syntax fixes/testcases Sandra Loosemore
@ 2023-08-18 17:39 ` Sandra Loosemore
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
2 siblings, 0 replies; 4+ messages in thread
From: Sandra Loosemore @ 2023-08-18 17:39 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub, julian, kcy
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
^ permalink raw reply [flat|nested] 4+ messages in thread
* [OG13, committed 2/3] OpenMP: C++ attribute syntax fixes/testcases for "declare mapper"
2023-08-18 17:39 [OG13, committed 0/3] C++ attribute syntax fixes/testcases Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective" Sandra Loosemore
@ 2023-08-18 17:39 ` Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 3/3] OpenMP: C++ attribute syntax fixes/testcases for loop transformations Sandra Loosemore
2 siblings, 0 replies; 4+ messages in thread
From: Sandra Loosemore @ 2023-08-18 17:39 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub, julian, kcy
gcc/c-family/ChangeLog
* c-omp.cc (c_omp_directives): Uncomment "declare mapper" entry.
gcc/cp/ChangeLog
* parser.cc (cp_parser_omp_declare_mapper): Allow commas between
clauses.
gcc/testsuite/ChangeLog
* g++.dg/gomp/attrs-declare-mapper-3.C: New file.
* g++.dg/gomp/attrs-declare-mapper-4.C: New file.
* g++.dg/gomp/attrs-declare-mapper-5.C: New file.
* g++.dg/gomp/attrs-declare-mapper-6.C: New file.
---
gcc/c-family/ChangeLog.omp | 4 +
gcc/c-family/c-omp.cc | 4 +-
gcc/cp/ChangeLog.omp | 5 ++
gcc/cp/parser.cc | 2 +
gcc/testsuite/ChangeLog.omp | 7 ++
.../g++.dg/gomp/attrs-declare-mapper-3.C | 31 ++++++++
.../g++.dg/gomp/attrs-declare-mapper-4.C | 74 +++++++++++++++++++
.../g++.dg/gomp/attrs-declare-mapper-5.C | 26 +++++++
.../g++.dg/gomp/attrs-declare-mapper-6.C | 22 ++++++
9 files changed, 173 insertions(+), 2 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-4.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-5.C
create mode 100644 gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-6.C
diff --git a/gcc/c-family/ChangeLog.omp b/gcc/c-family/ChangeLog.omp
index 60f5c29276f..40cb8e811e5 100644
--- a/gcc/c-family/ChangeLog.omp
+++ b/gcc/c-family/ChangeLog.omp
@@ -1,3 +1,7 @@
+2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
+
+ * c-omp.cc (c_omp_directives): Uncomment "declare mapper" entry.
+
2023-08-10 Julian Brown <julian@codesourcery.com>
* c-common.h (c_omp_region_type): Add C_ORT_UPDATE and C_ORT_OMP_UPDATE
diff --git a/gcc/c-family/c-omp.cc b/gcc/c-family/c-omp.cc
index 9ff09b59bc6..7bc69e9e2da 100644
--- a/gcc/c-family/c-omp.cc
+++ b/gcc/c-family/c-omp.cc
@@ -5500,8 +5500,8 @@ const struct c_omp_directive c_omp_directives[] = {
C_OMP_DIR_STANDALONE, false },
{ "critical", nullptr, nullptr, PRAGMA_OMP_CRITICAL,
C_OMP_DIR_CONSTRUCT, false },
- /* { "declare", "mapper", nullptr, PRAGMA_OMP_DECLARE,
- C_OMP_DIR_DECLARATIVE, false }, */
+ { "declare", "mapper", nullptr, PRAGMA_OMP_DECLARE,
+ C_OMP_DIR_DECLARATIVE, false },
{ "declare", "reduction", nullptr, PRAGMA_OMP_DECLARE,
C_OMP_DIR_DECLARATIVE, true },
{ "declare", "simd", nullptr, PRAGMA_OMP_DECLARE,
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 6e154ea3426..1b2d71422d8 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,8 @@
+2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
+
+ * parser.cc (cp_parser_omp_declare_mapper): Allow commas between
+ clauses.
+
2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
* parser.cc (analyze_metadirective_body): Handle CPP_PRAGMA and
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 84ec0de6c69..0ce2a7be608 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -50358,6 +50358,8 @@ cp_parser_omp_declare_mapper (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);
pragma_omp_clause c_kind = cp_parser_omp_clause_name (parser);
if (c_kind != PRAGMA_OMP_CLAUSE_MAP)
{
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index a9b4ac3d0a7..a7d11777988 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,10 @@
+2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
+
+ * g++.dg/gomp/attrs-declare-mapper-3.C: New file.
+ * g++.dg/gomp/attrs-declare-mapper-4.C: New file.
+ * g++.dg/gomp/attrs-declare-mapper-5.C: New file.
+ * g++.dg/gomp/attrs-declare-mapper-6.C: New file.
+
2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
* g++.dg/gomp/attrs-metadirective-1.C: New file.
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-3.C b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-3.C
new file mode 100644
index 00000000000..36345fe0dc2
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-3.C
@@ -0,0 +1,31 @@
+// { dg-do compile { target c++11 } }
+// { dg-additional-options "-fdump-tree-gimple" }
+
+#include <stdlib.h>
+
+// Test named mapper invocation.
+
+struct S {
+ int *ptr;
+ int size;
+};
+
+int main (int argc, char *argv[])
+{
+ int N = 1024;
+ [[omp::directive (declare mapper (mapN:struct S s)
+ map(to:s.ptr, s.size)
+ map(s.ptr[:N]))]];
+
+ struct S s;
+ s.ptr = (int *) malloc (sizeof (int) * N);
+
+ [[omp::directive (target map(mapper(mapN), tofrom: s))]]
+// { dg-final { scan-tree-dump {map\(struct:s \[len: 2\]\) map\(alloc:s\.ptr \[len: [0-9]+\]\) map\(to:s\.size \[len: [0-9]+\]\) map\(tofrom:\*_[0-9]+ \[len: _[0-9]+\]\) map\(attach:s\.ptr \[bias: 0\]\)} "gimple" } }
+ {
+ for (int i = 0; i < N; i++)
+ s.ptr[i]++;
+ }
+
+ return 0;
+}
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-4.C b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-4.C
new file mode 100644
index 00000000000..75cb48e8e44
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-4.C
@@ -0,0 +1,74 @@
+/* { dg-do compile { target c++11 } } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+/* Check mapper binding clauses. */
+
+struct Y {
+ int z;
+};
+
+struct Z {
+ int z;
+};
+
+[[omp::directive (declare mapper (struct Y y) map(tofrom: y))]];
+[[omp::directive (declare mapper (struct Z z) map(tofrom: z))]];
+
+int foo (void)
+{
+ struct Y yy;
+ struct Z zz;
+ int dummy;
+
+ [[omp::directive (target data map(dummy))]]
+ {
+ [[omp::directive (target)]]
+ {
+ yy.z++;
+ zz.z++;
+ }
+ yy.z++;
+ }
+ return yy.z;
+}
+
+struct P
+{
+ struct Z *zp;
+};
+
+int bar (void)
+{
+ struct Y yy;
+ struct Z zz;
+ struct P pp;
+ struct Z t;
+ int dummy;
+
+ pp.zp = &t;
+
+ [[omp::directive (declare mapper (struct Y y) map(tofrom: y.z))]];
+ [[omp::directive (declare mapper (struct Z z) map(tofrom: z.z))]];
+
+ [[omp::directive (target data map(dummy))]]
+ {
+ [[omp::directive (target)]]
+ {
+ yy.z++;
+ zz.z++;
+ }
+ yy.z++;
+ }
+
+ [[omp::directive (declare mapper(struct P x) map(to:x.zp), map(tofrom:*x.zp))]];
+
+ [[omp::directive (target)]]
+ {
+ zz = *pp.zp;
+ }
+
+ return zz.z;
+}
+
+/* { dg-final { scan-tree-dump-times {mapper_binding\(struct Y,omp declare mapper ~1Y\) mapper_binding\(struct Z,omp declare mapper ~1Z\)} 2 "original" { target c++ } } } */
+/* { dg-final { scan-tree-dump {mapper_binding\(struct Z,omp declare mapper ~1Z\) mapper_binding\(struct P,omp declare mapper ~1P\)} "original" { target c++ } } } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-5.C b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-5.C
new file mode 100644
index 00000000000..711f058dc2b
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-5.C
@@ -0,0 +1,26 @@
+/* { dg-do compile { target c++11 } } */
+
+typedef struct S_ {
+ int *myarr;
+ int size;
+} S;
+
+[[omp::directive (declare mapper (named: struct S_ v)
+ map(to:v.size, v.myarr)
+ map(tofrom: v.myarr[0:v.size]))]];
+/* { dg-note "'#pragma omp declare mapper \\(named: S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+[[omp::directive (declare mapper (named: S v)
+ map(to:v.size, v.myarr)
+ map(tofrom: v.myarr[0:v.size]))]];
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(named: S\\)'" "" { target c++ } .-3 } */
+
+[[omp::directive (declare mapper (struct S_ v)
+ map(to:v.size, v.myarr)
+ map(tofrom: v.myarr[0:v.size]))]];
+/* { dg-note "'#pragma omp declare mapper \\(S_\\)' previously defined here" "" { target c++ } .-3 } */
+
+[[omp::directive (declare mapper (S v)
+ map(to:v.size, v.myarr)
+ map(tofrom: v.myarr[0:v.size]))]];
+/* { dg-error "redefinition of '#pragma omp declare mapper \\(S\\)'" "" { target c++ } .-3 } */
diff --git a/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-6.C b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-6.C
new file mode 100644
index 00000000000..fb3a0deb5ef
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/attrs-declare-mapper-6.C
@@ -0,0 +1,22 @@
+/* { dg-do compile { target c++11 } } */
+
+int x = 5;
+
+struct Q {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+[[omp::directive (declare mapper (struct Q myq) map(myq.arr2[0:x]))]];
+
+struct R {
+ int *arr1;
+ int *arr2;
+ int *arr3;
+};
+
+[[omp::directive (declare mapper (struct R myr) map(myr.arr3[0:y]))]];
+/* { dg-error "'y' was not declared in this scope" "" { target c++ } .-1 } */
+
+int y = 7;
--
2.31.1
^ permalink raw reply [flat|nested] 4+ messages in thread
* [OG13, committed 3/3] OpenMP: C++ attribute syntax fixes/testcases for loop transformations
2023-08-18 17:39 [OG13, committed 0/3] C++ attribute syntax fixes/testcases Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective" Sandra Loosemore
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 ` Sandra Loosemore
2 siblings, 0 replies; 4+ messages in thread
From: Sandra Loosemore @ 2023-08-18 17:39 UTC (permalink / raw)
To: gcc-patches; +Cc: jakub, julian, kcy
gcc/cp/ChangeLog
* parser.cc (cp_parser_omp_all_clauses): Allow comma before first
clause.
(cp_parser_see_omp_loop_nest): Accept C++ standard attributes
before RID_FOR.
(cp_parser_omp_loop_nest): Process C++ standard attributes like
pragmas. Improve error handling for bad pragmas/attributes.
Use cp_parser_see_omp_loop_nest instead of duplicating what it
does.
(cp_parser_omp_tile_sizes): Permit comma before the clause.
(cp_parser_omp_tile): Assert that this isn't called for inner
directive.
(cp_parser_omp_unroll): Likewise.
gcc/testsuite/ChangeLog
* g++.dg/gomp/loop-transforms/attrs-tile-1.C: New file.
* g++.dg/gomp/loop-transforms/attrs-tile-2.C: New file.
* g++.dg/gomp/loop-transforms/attrs-tile-3.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-1.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-2.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-3.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C: New file.
* g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C: New file.
---
gcc/cp/ChangeLog.omp | 15 ++
gcc/cp/parser.cc | 69 ++++---
gcc/testsuite/ChangeLog.omp | 12 ++
.../gomp/loop-transforms/attrs-tile-1.C | 164 +++++++++++++++++
.../gomp/loop-transforms/attrs-tile-2.C | 174 ++++++++++++++++++
.../gomp/loop-transforms/attrs-tile-3.C | 111 +++++++++++
.../gomp/loop-transforms/attrs-unroll-1.C | 135 ++++++++++++++
.../gomp/loop-transforms/attrs-unroll-2.C | 81 ++++++++
.../gomp/loop-transforms/attrs-unroll-3.C | 20 ++
.../loop-transforms/attrs-unroll-inner-1.C | 15 ++
.../loop-transforms/attrs-unroll-inner-2.C | 29 +++
.../loop-transforms/attrs-unroll-inner-3.C | 71 +++++++
12 files changed, 872 insertions(+), 24 deletions(-)
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-3.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C
create mode 100644 gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C
diff --git a/gcc/cp/ChangeLog.omp b/gcc/cp/ChangeLog.omp
index 1b2d71422d8..fe5ef67a7ad 100644
--- a/gcc/cp/ChangeLog.omp
+++ b/gcc/cp/ChangeLog.omp
@@ -1,3 +1,18 @@
+2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
+
+ * parser.cc (cp_parser_omp_all_clauses): Allow comma before first
+ clause.
+ (cp_parser_see_omp_loop_nest): Accept C++ standard attributes
+ before RID_FOR.
+ (cp_parser_omp_loop_nest): Process C++ standard attributes like
+ pragmas. Improve error handling for bad pragmas/attributes.
+ Use cp_parser_see_omp_loop_nest instead of duplicating what it
+ does.
+ (cp_parser_omp_tile_sizes): Permit comma before the clause.
+ (cp_parser_omp_tile): Assert that this isn't called for inner
+ directive.
+ (cp_parser_omp_unroll): Likewise.
+
2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
* parser.cc (cp_parser_omp_declare_mapper): Allow commas between
diff --git a/gcc/cp/parser.cc b/gcc/cp/parser.cc
index 0ce2a7be608..4871f4511a9 100644
--- a/gcc/cp/parser.cc
+++ b/gcc/cp/parser.cc
@@ -42240,15 +42240,12 @@ cp_parser_omp_all_clauses (cp_parser *parser, omp_clause_mask mask,
if (nested && cp_lexer_next_token_is (parser->lexer, CPP_CLOSE_PAREN))
break;
- if (!first || nested != 2)
- {
- if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
- cp_lexer_consume_token (parser->lexer);
- else if (nested == 2)
- error_at (cp_lexer_peek_token (parser->lexer)->location,
- "clauses in %<simd%> trait should be separated "
- "by %<,%>");
- }
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+ else if (!first && nested == 2)
+ error_at (cp_lexer_peek_token (parser->lexer)->location,
+ "clauses in %<simd%> trait should be separated "
+ "by %<,%>");
token = cp_lexer_peek_token (parser->lexer);
c_kind = cp_parser_omp_clause_name (parser);
@@ -44803,6 +44800,11 @@ cp_parser_see_omp_loop_nest (cp_parser *parser, enum tree_code code,
|| (cp_parser_pragma_kind (cp_lexer_peek_token (parser->lexer))
== PRAGMA_OMP_TILE))
return true;
+ if (cp_lexer_nth_token_is_keyword
+ (parser->lexer,
+ cp_parser_skip_std_attribute_spec_seq (parser, 1),
+ RID_FOR))
+ return true;
if (error_p)
cp_parser_error (parser, "loop nest expected");
}
@@ -44868,6 +44870,11 @@ cp_parser_omp_loop_nest (cp_parser *parser, bool *if_p)
the depth of the *next* loop, not the level of the loop body the
transformation directive appears in. */
+ /* Arrange for C++ standard attribute syntax to be parsed as regular
+ pragmas. */
+ tree std_attrs = cp_parser_std_attribute_spec_seq (parser);
+ std_attrs = cp_parser_handle_statement_omp_attributes (parser, std_attrs);
+
if ((cp_parser_pragma_kind (cp_lexer_peek_token (parser->lexer))
== PRAGMA_OMP_UNROLL)
|| (cp_parser_pragma_kind (cp_lexer_peek_token (parser->lexer))
@@ -44893,19 +44900,29 @@ cp_parser_omp_loop_nest (cp_parser *parser, bool *if_p)
omp_for_parse_state->orig_declv
= grow_tree_vec (omp_for_parse_state->orig_declv, count);
}
- if (cp_parser_see_omp_loop_nest (parser, omp_for_parse_state->code,
- true))
- return cp_parser_omp_loop_nest (parser, if_p);
- else
+ }
+
+ /* Diagnose errors if we don't have a "for" loop following the
+ optional loop transforms. Otherwise, consume the token. */
+ if (!cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR))
+ {
+ omp_for_parse_state->fail = true;
+ cp_token *token = cp_lexer_peek_token (parser->lexer);
+ /* Don't call cp_parser_error here since it overrides the
+ provided message with a more confusing one if there was
+ a bad pragma or attribute directive. */
+ error_at (token->location, "loop nest expected");
+ /* See if we can recover by skipping over bad pragma(s). */
+ while (token->type == CPP_PRAGMA)
{
- /* FIXME: Better error recovery here? */
- omp_for_parse_state->fail = true;
- return NULL_TREE;
+ cp_parser_skip_to_pragma_eol (parser, token);
+ if (cp_parser_see_omp_loop_nest (parser, omp_for_parse_state->code,
+ false))
+ return cp_parser_omp_loop_nest (parser, if_p);
+ token = cp_lexer_peek_token (parser->lexer);
}
+ return NULL_TREE;
}
-
- /* We have already matched the FOR token but not consumed it yet. */
- gcc_assert (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR));
loc = cp_lexer_consume_token (parser->lexer)->location;
/* Forbid break/continue in the loop initializer, condition, and
@@ -45161,11 +45178,8 @@ cp_parser_omp_loop_nest (cp_parser *parser, bool *if_p)
moreloops = depth < omp_for_parse_state->count - 1;
omp_for_parse_state->want_nested_loop = moreloops;
if (moreloops
- && (cp_lexer_next_token_is_keyword (parser->lexer, RID_FOR)
- || (cp_parser_pragma_kind (cp_lexer_peek_token (parser->lexer))
- == PRAGMA_OMP_UNROLL)
- || (cp_parser_pragma_kind (cp_lexer_peek_token (parser->lexer))
- == PRAGMA_OMP_TILE)))
+ && cp_parser_see_omp_loop_nest (parser, omp_for_parse_state->code,
+ false))
{
omp_for_parse_state->depth++;
add_stmt (cp_parser_omp_loop_nest (parser, if_p));
@@ -47323,6 +47337,9 @@ cp_parser_omp_tile_sizes (cp_parser *parser, location_t loc)
tree sizes = NULL_TREE;
cp_lexer *lexer = parser->lexer;
+ if (cp_lexer_next_token_is (parser->lexer, CPP_COMMA))
+ cp_lexer_consume_token (parser->lexer);
+
cp_token *tok = cp_lexer_peek_token (lexer);
if (tok->type != CPP_NAME
|| strcmp ("sizes", IDENTIFIER_POINTER (tok->u.value)))
@@ -47375,6 +47392,8 @@ cp_parser_omp_tile (cp_parser *parser, cp_token *tok, bool *if_p)
tree block;
tree ret = error_mark_node;
+ gcc_assert (!parser->omp_for_parse_state);
+
tree clauses = cp_parser_omp_tile_sizes (parser, tok->location);
cp_parser_require_pragma_eol (parser, tok);
@@ -47553,6 +47572,8 @@ cp_parser_omp_unroll (cp_parser *parser, cp_token *tok, bool *if_p)
static const char *p_name = "#pragma omp unroll";
omp_clause_mask mask = OMP_UNROLL_CLAUSE_MASK;
+ gcc_assert (!parser->omp_for_parse_state);
+
tree clauses = cp_parser_omp_all_clauses (parser, mask, p_name, tok, true);
if (!clauses)
diff --git a/gcc/testsuite/ChangeLog.omp b/gcc/testsuite/ChangeLog.omp
index a7d11777988..3008a8e52eb 100644
--- a/gcc/testsuite/ChangeLog.omp
+++ b/gcc/testsuite/ChangeLog.omp
@@ -1,3 +1,15 @@
+2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
+
+ * g++.dg/gomp/loop-transforms/attrs-tile-1.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-tile-2.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-tile-3.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-1.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-2.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-3.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C: New file.
+ * g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C: New file.
+
2023-08-18 Sandra Loosemore <sandra@codesourcery.com>
* g++.dg/gomp/attrs-declare-mapper-3.C: New file.
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-1.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-1.C
new file mode 100644
index 00000000000..0906ff3bbe8
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-1.C
@@ -0,0 +1,164 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+ [[omp::directive (tile sizes(1))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes(0))]] /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes(-1))]] /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes())]] /* { dg-error {expected primary-expression before} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes)]] /* { dg-error {expected '\(' before end of line} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1) sizes(1))]] /* { dg-error {expected end of line before 'sizes'} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile, sizes(1), sizes(1))]] /* { dg-error {expected end of line before ','} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1)),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1, 2)),
+ directive (tile sizes(1)))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1, 2)),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(5, 6)),
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ for (int k = 0; k < 100; ++k)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1)),
+ directive (unroll partia), /* { dg-error {expected an OpenMP clause before 'partia'} } */
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1)),
+ directive (unroll))]] /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1)),
+ directive (unroll full))]] /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(1)),
+ directive (unroll partial),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(8,8)),
+ directive (unroll partial), /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence
+ (directive (tile sizes(8,8)),
+ directive (unroll partial))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1, 2))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1, 2))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = i; j < 100; ++j)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1, 2))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 2; j < i; ++j)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1, 2, 3))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::directive (tile sizes(1))]]
+ for (int i = 0; i < 100; ++i)
+ {
+ dummy (i);
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+ }
+
+ [[omp::directive (tile sizes(1))]]
+ for (int i = 0; i < 100; ++i)
+ {
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ dummy (i);
+ }
+
+ [[omp::directive (tile sizes(1, 2))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {inner loops must be perfectly nested} } */
+ {
+ dummy (i);
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ }
+
+ [[omp::directive (tile sizes(1, 2))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {inner loops must be perfectly nested} } */
+ {
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ dummy (i);
+ }
+
+ int s;
+ [[omp::directive (tile sizes(s))]] /* { dg-error {'tile sizes' argument needs positive integral constant} "" { target { ! c++98_only } } } */
+ /* { dg-error {the value of 's' is not usable in a constant expression} "" { target { c++ && { ! c++98_only } } } .-1 } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::directive (tile sizes(42.0))]] /* { dg-error {'tile sizes' argument needs integral type} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+}
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-2.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-2.C
new file mode 100644
index 00000000000..ab02924defa
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-2.C
@@ -0,0 +1,174 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(0)))]] /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(-1)))]] /* { dg-error {'tile sizes' argument needs positive integral constant} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes()))]] /* { dg-error {expected primary-expression before} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(,)))]] /* { dg-error {expected primary-expression before} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes))]] /* { dg-error {expected '\(' before end of line} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1) sizes(1)))]] /* { dg-error {expected end of line before 'sizes'} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1)))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(5, 6)),
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ for (int k = 0; k < 100; ++k)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)),
+ directive (unroll partia), /* { dg-error {expected an OpenMP clause before 'partia'} } */
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)),
+ directive (unroll))]] /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)),
+ directive (unroll full))]] /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)),
+ directive (unroll partial),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(8,8)),
+ directive (unroll partial), /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(8,8)),
+ directive (unroll partial))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = i; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 2; j < i; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ {
+ dummy (i);
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+ }
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ {
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ dummy (i);
+ }
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {inner loops must be perfectly nested} } */
+ {
+ dummy (i);
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ }
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {inner loops must be perfectly nested} } */
+ {
+ for (int j = 0; j < 100; ++j)
+ dummy (j);
+ dummy (i);
+ }
+
+ [[omp::sequence (directive (parallel for),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+}
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-3.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-3.C
new file mode 100644
index 00000000000..95a0115b014
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-tile-3.C
@@ -0,0 +1,111 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+ [[omp::sequence (directive (for),
+ directive (tile sizes(1, 2)))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = i; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for),
+ directive (tile sizes(1, 2)))]] /* { dg-error {'tile' loop transformation may not appear on non-rectangular for} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < i; ++j)
+ dummy (i);
+
+
+ [[omp::sequence (directive (for collapse(1)),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1)))]] /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(3)),
+ directive (tile sizes(1, 2)))]] /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(1)),
+ directive (tile sizes(1)),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1)))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(5, 6)),
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(1)),
+ directive (tile sizes(1)),
+ directive (tile sizes(1)))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1)))]] /* { dg-error {nesting depth left after this transformation too low for outer transformation} } */
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(1, 2)),
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(2)),
+ directive (tile sizes(5, 6)),
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ for (int k = 0; k < 100; ++k)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(3)),
+ directive (tile sizes(1, 2)), /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+ directive (tile sizes(1, 2)))]]
+ for (int i = 0; i < 100; ++i) /* { dg-error {not enough nested loops} } */
+ for (int j = 0; j < 100; ++j)
+ dummy (i);
+
+ [[omp::sequence (directive (for collapse(3)),
+ directive (tile sizes(5, 6)), /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+ directive (tile sizes(1, 2, 3)))]]
+ for (int i = 0; i < 100; ++i)
+ for (int j = 0; j < 100; ++j)
+ for (int k = 0; k < 100; ++k)
+ dummy (i);
+}
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-1.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-1.C
new file mode 100644
index 00000000000..5b93b9fa59e
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-1.C
@@ -0,0 +1,135 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test1 ()
+{
+[[omp::directive (unroll partial)]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+}
+
+void
+test2 ()
+{
+[[omp::directive (unroll partial(10))]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+}
+
+void
+test3 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = 0; i < 100; ++i)
+ dummy (i);
+}
+
+void
+test4 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = 0; i > 100; ++i)
+ dummy (i);
+}
+
+void
+test5 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = 1; i <= 100; ++i)
+ dummy (i);
+}
+
+void
+test6 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = 200; i >= 100; i--)
+ dummy (i);
+}
+
+void
+test7 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = -100; i > 100; ++i)
+ dummy (i);
+}
+
+void
+test8 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = 100; i > -200; --i)
+ dummy (i);
+}
+
+void
+test9 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+}
+
+void
+test10 ()
+{
+[[omp::directive (unroll full)]]
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+}
+
+void
+test12 ()
+{
+[[omp::sequence (directive (unroll full),
+ directive (unroll partial),
+ directive (unroll partial))]]
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+}
+
+void
+test13 ()
+{
+ for (int i = 0; i < 100; ++i)
+[[omp::sequence (directive (unroll full),
+ directive (unroll partial),
+ directive (unroll partial))]]
+ for (int j = -300; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test14 ()
+{
+ [[omp::directive (for)]]
+ for (int i = 0; i < 100; ++i)
+ [[omp::sequence (directive (unroll full),
+ directive (unroll partial),
+ directive (unroll partial))]]
+ for (int j = -300; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test15 ()
+{
+ [[omp::directive (for)]]
+ for (int i = 0; i < 100; ++i)
+ {
+
+ dummy (i);
+
+ [[omp::sequence (directive (unroll full),
+ directive (unroll partial),
+ directive (unroll partial))]]
+ for (int j = -300; j != 100; ++j)
+ dummy (j);
+
+ dummy (i);
+ }
+ }
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-2.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-2.C
new file mode 100644
index 00000000000..1a45eadec64
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-2.C
@@ -0,0 +1,81 @@
+/* { dg-prune-output "error: invalid controlling predicate" } */
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+[[omp::sequence (directive (unroll partial),
+ directive (unroll full))]] /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll full), /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+ directive (unroll partial))]]
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll full), /* { dg-error {'full' clause is invalid here; turns loop into non-loop} } */
+ directive (unroll full))]]
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll partial partial))]] /* { dg-error {too many 'partial' clauses} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::directive (unroll full full)]] /* { dg-error {too many 'full' clauses} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (unroll partial),
+ directive (unroll))]] /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll))]] /* { dg-error {'#pragma omp unroll' without 'partial' clause is invalid here; turns loop into non-loop} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+ int i;
+
+[[omp::sequence (directive (for),
+ directive (unroll foo))]] /* { dg-error {expected an OpenMP clause before 'foo'} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::directive (unroll partial(i))]]
+ /* { dg-error {the value of 'i' is not usable in a constant expression} "" { target c++ } .-1 } */
+ /* { dg-error {partial argument needs positive constant integer expression} "" { target *-*-* } .-2 } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::directive (unroll parti)]] /* { dg-error {expected an OpenMP clause before 'parti'} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll partial(1)),
+ directive (unroll parti))]] /* { dg-error {expected an OpenMP clause before 'parti'} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+[[omp::sequence (directive (for),
+ directive (unroll partial(1)),
+ directive (unroll parti))]] /* { dg-error {expected an OpenMP clause before 'parti'} } */
+ for (int i = -300; i != 100; ++i)
+ dummy (i);
+
+int sum = 0;
+[[omp::sequence (directive (parallel for reduction(+ : sum) collapse(2)),
+ directive (unroll partial(1)))]] /* { dg-error {nesting depth left after this transformation too low for loop collapse} } */
+ for (int i = 3; i < 10; ++i)
+ for (int j = -2; j < 7; ++j)
+ sum++;
+}
+
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-3.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-3.C
new file mode 100644
index 00000000000..20c11c0f314
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-3.C
@@ -0,0 +1,20 @@
+// { dg-do compile { target c++11 } }
+
+/* { dg-additional-options "-fdump-tree-omp_transform_loops" }
+ * { dg-additional-options "-fdump-tree-original" } */
+
+extern void dummy (int);
+
+void
+test1 ()
+{
+ int i;
+ [[omp::directive (unroll full)]]
+ for (int i = 0; i < 10; i++)
+ dummy (i);
+}
+
+ /* Loop should be removed with 10 copies of the body remaining
+ * { dg-final { scan-tree-dump-times "dummy" 10 "omp_transform_loops" } }
+ * { dg-final { scan-tree-dump "#pragma omp loop_transform" "original" } }
+ * { dg-final { scan-tree-dump-not "#pragma omp" "omp_transform_loops" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C
new file mode 100644
index 00000000000..234753ad017
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-1.C
@@ -0,0 +1,15 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (unroll, partial)]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C
new file mode 100644
index 00000000000..26cc665007d
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-2.C
@@ -0,0 +1,29 @@
+// { dg-do compile { target c++11 } }
+
+extern void dummy (int);
+
+void
+test ()
+{
+
+#pragma omp target parallel for collapse(2)
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (tile sizes (2))]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i) /* { dg-error {not enough nested loops} } */
+ [[omp::directive (tile sizes(2, 3))]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+
+ [[omp::directive (target parallel for, collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (tile, sizes(2, 3))]]
+ for (int j = 0; j != 100; ++j)
+ for (int k = 0; k != 100; ++k)
+ dummy (i);
+}
+
+
diff --git a/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C
new file mode 100644
index 00000000000..46970b84a24
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/loop-transforms/attrs-unroll-inner-3.C
@@ -0,0 +1,71 @@
+// { dg-do compile { target c++11 } }
+
+// Test that omp::sequence is handled properly in a loop nest, but that
+// invalid attribute specifiers are rejected.
+
+extern void dummy (int);
+
+void
+test1 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::sequence (directive (unroll, partial))]] // OK
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test2 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (masked)]] // { dg-error "loop nest expected" }
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test3 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (unroll, partial)]] // { dg-error "attributes on the same statement" }
+ [[omp::directive (masked)]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test4 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::sequence (directive (unroll, partial),
+ directive (masked))]] // { dg-error "loop nest expected" }
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test5 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::sequence (directive (masked), // { dg-error "loop nest expected" }
+ directive (unroll, partial))]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
+void
+test6 ()
+{
+ [[omp::directive (target parallel for collapse(2))]]
+ for (int i = -300; i != 100; ++i)
+ [[omp::directive (unroll, partial), // { dg-error "attributes on the same statement" }
+ omp::directive (masked)]]
+ for (int j = 0; j != 100; ++j)
+ dummy (i);
+}
+
--
2.31.1
^ permalink raw reply [flat|nested] 4+ messages in thread
end of thread, other threads:[~2023-08-18 17:40 UTC | newest]
Thread overview: 4+ messages (download: mbox.gz / follow: Atom feed)
-- links below jump to the message on this page --
2023-08-18 17:39 [OG13, committed 0/3] C++ attribute syntax fixes/testcases Sandra Loosemore
2023-08-18 17:39 ` [OG13, committed 1/3] OpenMP: C++ attribute syntax fixes/testcases for "metadirective" Sandra Loosemore
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
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).