gcc/testsuite/ChangeLog
* c-c++-common/gomp/attrs-metadirective-1.c: New.
* c-c++-common/gomp/attrs-metadirective-2.c: New.
* c-c++-common/gomp/attrs-metadirective-3.c: New.
* c-c++-common/gomp/attrs-metadirective-4.c: New.
* c-c++-common/gomp/attrs-metadirective-5.c: New.
* c-c++-common/gomp/attrs-metadirective-6.c: New.
* c-c++-common/gomp/attrs-metadirective-7.c: New.
* c-c++-common/gomp/attrs-metadirective-8.c: New.
* c-c++-common/gomp/declare-variant-arg-exprs.c: New.
* c-c++-common/gomp/declare-variant-dynamic-1.c: New.
* c-c++-common/gomp/declare-variant-dynamic-2.c: New.
* c-c++-common/gomp/metadirective-1.c: New.
* c-c++-common/gomp/metadirective-2.c: New.
* c-c++-common/gomp/metadirective-3.c: New.
* c-c++-common/gomp/metadirective-4.c: New.
* c-c++-common/gomp/metadirective-5.c: New.
* c-c++-common/gomp/metadirective-6.c: New.
* c-c++-common/gomp/metadirective-7.c: New.
* c-c++-common/gomp/metadirective-8.c: New.
* c-c++-common/gomp/metadirective-construct.c: New.
* c-c++-common/gomp/metadirective-device.c: New.
* c-c++-common/gomp/metadirective-no-score.c: New.
* c-c++-common/gomp/metadirective-target-device.c: New.
libgomp/ChangeLog
* testsuite/libgomp.c-c++-common/metadirective-1.c: New.
* testsuite/libgomp.c-c++-common/metadirective-2.c: New.
* testsuite/libgomp.c-c++-common/metadirective-3.c: New.
* testsuite/libgomp.c-c++-common/metadirective-4.c: New.
* testsuite/libgomp.c-c++-common/metadirective-5.c: New.
* testsuite/libgomp.c-c++-common/metadirective-late-1.c: New.
* testsuite/libgomp.c-c++-common/metadirective-late-2.c: New.
* testsuite/libgomp.c-c++-common/metadirective-target-device.c: New.
Co-Authored-By: Kwok Cheung Yeung <k...@codesourcery.com>
Co-Authored-By: Sandra Loosemore <san...@codesourcery.com>
---
.../c-c++-common/gomp/attrs-metadirective-1.c | 47 +++++
.../c-c++-common/gomp/attrs-metadirective-2.c | 76 ++++++++
.../c-c++-common/gomp/attrs-metadirective-3.c | 24 +++
.../c-c++-common/gomp/attrs-metadirective-4.c | 43 +++++
.../c-c++-common/gomp/attrs-metadirective-5.c | 26 +++
.../c-c++-common/gomp/attrs-metadirective-6.c | 33 ++++
.../c-c++-common/gomp/attrs-metadirective-7.c | 37 ++++
.../c-c++-common/gomp/attrs-metadirective-8.c | 18 ++
.../gomp/declare-variant-arg-exprs.c | 29 +++
.../gomp/declare-variant-dynamic-1.c | 26 +++
.../gomp/declare-variant-dynamic-2.c | 30 +++
.../c-c++-common/gomp/metadirective-1.c | 58 ++++++
.../c-c++-common/gomp/metadirective-2.c | 75 ++++++++
.../c-c++-common/gomp/metadirective-3.c | 23 +++
.../c-c++-common/gomp/metadirective-4.c | 42 +++++
.../c-c++-common/gomp/metadirective-5.c | 25 +++
.../c-c++-common/gomp/metadirective-6.c | 32 ++++
.../c-c++-common/gomp/metadirective-7.c | 36 ++++
.../c-c++-common/gomp/metadirective-8.c | 17 ++
.../gomp/metadirective-construct.c | 178 ++++++++++++++++++
.../c-c++-common/gomp/metadirective-device.c | 149 +++++++++++++++
.../gomp/metadirective-no-score.c | 95 ++++++++++
.../gomp/metadirective-target-device.c | 149 +++++++++++++++
.../libgomp.c-c++-common/metadirective-1.c | 37 ++++
.../libgomp.c-c++-common/metadirective-2.c | 43 +++++
.../libgomp.c-c++-common/metadirective-3.c | 36 ++++
.../libgomp.c-c++-common/metadirective-4.c | 54 ++++++
.../libgomp.c-c++-common/metadirective-5.c | 48 +++++
.../metadirective-late-1.c | 66 +++++++
.../metadirective-late-2.c | 66 +++++++
.../metadirective-target-device.c | 76 ++++++++
31 files changed, 1694 insertions(+)
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-5.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-6.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/attrs-metadirective-8.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-arg-exprs.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-1.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-2.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-3.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-4.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-5.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-6.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-7.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-8.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-construct.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-device.c
create mode 100644 gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c
create mode 100644
gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c
create mode 100644 libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/metadirective-late-1.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/metadirective-late-2.c
create mode 100644
libgomp/testsuite/libgomp.c-c++-common/metadirective-target-device.c
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-1.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-1.c
new file mode 100644
index 00000000000..275eab165ad
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-1.c
@@ -0,0 +1,47 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+
+#define N 100
+
+void
+f (int a[], int b[], int c[])
+{
+ int i;
+
+ [[omp::directive (metadirective
+ default (teams loop)
+ default (parallel loop))]] /* { dg-error "too many 'otherwise' or
'default' clauses in 'metadirective'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ [[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
+ where (device={arch("nvptx")}: parallel loop) /* { dg-error "'where' is
not valid for 'metadirective'" } */
+ default (teams loop))]]
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ [[omp::directive (metadirective
+ otherwise (teams loop)
+ when (device={arch("nvptx")}: parallel loop))]] /* { dg-error
"'otherwise' or 'default' clause must appear last" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ [[omp::directive (metadirective
+ when (device={arch("nvptx")} parallel loop) /* { dg-error "expected ':'
before 'parallel'" } */
+ default (teams loop))]]
+ 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 variants of a 'metadirective' before
'default'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ /* Test improperly nested metadirectives - even though the second
+ metadirective resolves to 'omp nothing', that is not the same as there
+ being literally nothing there. */
+ [[omp::directive (metadirective
+ when (implementation={vendor("gnu")}: parallel for))]]
+ [[omp::directive (metadirective /* { dg-error "loop nest expected" } */
+ when (implementation={vendor("cray")}: parallel for))]]
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-2.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-2.c
new file mode 100644
index 00000000000..ff401c86b09
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-2.c
@@ -0,0 +1,76 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+
+#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/c-c++-common/gomp/attrs-metadirective-3.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-3.c
new file mode 100644
index 00000000000..31dd054922f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-3.c
@@ -0,0 +1,24 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#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];
+}
+
+/* If offload device "nvptx" isn't supported, the front end can eliminate
+ that alternative and not produce a metadirective at all. Otherwise this
+ won't be resolved until late. */
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" {
target { ! offload_nvptx } } } } */
+/* { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target {
offload_nvptx } } } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c
new file mode 100644
index 00000000000..4cae553704a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c
@@ -0,0 +1,43 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+/* { 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 (void)
+{
+ 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 "otherwise:" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-5.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-5.c
new file mode 100644
index 00000000000..b24818a0fc7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-5.c
@@ -0,0 +1,26 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+/* { 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/c-c++-common/gomp/attrs-metadirective-6.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-6.c
new file mode 100644
index 00000000000..b1338a57001
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-6.c
@@ -0,0 +1,33 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+/* { 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 "otherwise:" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c
new file mode 100644
index 00000000000..8e3bf32c538
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c
@@ -0,0 +1,37 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" } */
+
+#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;
+}
+
+/* We expect one "pragma omp target" with "device(num)" for each target_device
+ selector that specifies "device_num(num)". */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num"
3 "gimple" } } */
+
+/* One OMP_TARGET_DEVICE_MATCHES for each kind/arch/isa selector. These
+ are supposed to go away after ompdevlow. */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" }
} */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow"
} } */
diff --git a/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-8.c
b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-8.c
new file mode 100644
index 00000000000..43746b669ff
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-8.c
@@ -0,0 +1,18 @@
+/* { dg-do compile { target { c || c++11 } } } */
+/* { dg-options "-fopenmp -std=c23" { target { c } } } */
+
+#define N 256
+
+void
+f (void)
+{
+ 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;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-arg-exprs.c
b/gcc/testsuite/c-c++-common/gomp/declare-variant-arg-exprs.c
new file mode 100644
index 00000000000..38bfe928c74
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-arg-exprs.c
@@ -0,0 +1,29 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable" } */
+/* { dg-additional-options "-mavx512bw -mavx512vl" { target { i?86-*-*
x86_64-*-* } } } */
+
+/* References to function parameters in dynamic selector expressions for
+ "declare variant" isn't supported yet; see PR 113904. Check to see that
+ a proper error is diagnosed meanwhile and GCC doesn't just wander off
+ into the weeds and ICE. */
+
+extern int frob (int);
+
+void f01 (int, int);
+void f02 (int, int);
+void f03 (int, int);
+#pragma omp declare variant (f01) match (target_device={device_num (devnum),
isa("avx512f","avx512vl")}) /* { dg-message "sorry, unimplemented: reference to
function parameter" } */
+#pragma omp declare variant (f02) match
(implementation={vendor(score(15):gnu)})
+#pragma omp declare variant (f03) match (user={condition(score(11):frob (ok +
42))}) /* { dg-message "sorry, unimplemented: reference to function parameter"
} */
+void f04 (int devnum, int ok);
+
+void
+test1 (void)
+{
+ int i;
+ #pragma omp parallel for
+ for (i = 0; i < 1; i++)
+ f04 (17, 1);
+}
+
+
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-1.c
b/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-1.c
new file mode 100644
index 00000000000..b406a31eb30
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-1.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+extern int foo_p (int);
+extern int bar;
+
+int f01 (int);
+int f02 (int);
+int f03 (int);
+int f04 (int);
+#pragma omp declare variant (f01) match (device={isa("avx512f")}) /* 4 */
+#pragma omp declare variant (f02) match
(implementation={vendor(score(3):gnu)},device={kind(cpu)}) /* 1 + 3 */
+#pragma omp declare variant (f03) match (user={condition(score(9):foo_p
(bar))})
+#pragma omp declare variant (f04) match
(implementation={vendor(score(6):gnu)},device={kind(host)}) /* 1 + 6 */
+int f05 (int);
+
+
+int
+test1 (int x)
+{
+ return f05 (x);
+}
+
+/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f04 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f05 \\\(x" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-2.c
b/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-2.c
new file mode 100644
index 00000000000..c078123e2e6
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/declare-variant-dynamic-2.c
@@ -0,0 +1,30 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+extern int foo_p (int);
+extern int bar;
+extern int omp_get_default_device (void);
+
+int f01 (int);
+int f02 (int);
+int f03 (int);
+int f04 (int);
+#pragma omp declare variant (f01) match
(target_device={device_num(omp_get_default_device()), isa("avx512f")}) /* 4 */
+#pragma omp declare variant (f02) match (user={condition(score(6):0)})
+#pragma omp declare variant (f03) match (user={condition(score(5):foo_p
(bar))})
+#pragma omp declare variant (f04) match (user={condition(score(3):0)})
+int f05 (int);
+
+int
+test1 (int x)
+{
+ return f05 (x);
+}
+
+/* f01 and f03 are the dynamic selectors, the fall-through is f05.
+ f02 and f04 are static selectors and do not match. */
+/* { dg-final { scan-tree-dump "f01 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f03 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump "f05 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f02 \\\(x" "gimple" } } */
+/* { dg-final { scan-tree-dump-not "f04 \\\(x" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-1.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c
new file mode 100644
index 00000000000..b2bfe950818
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-1.c
@@ -0,0 +1,58 @@
+/* { dg-do compile } */
+
+#define N 100
+
+void
+f (int a[], int b[], int c[])
+{
+ int i;
+
+ #pragma omp metadirective \
+ default (teams loop) \
+ default (parallel loop) /* { dg-error "too many 'otherwise' or 'default'
clauses in 'metadirective'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ otherwise (teams loop) \
+ default (parallel loop) /* { dg-error "too many 'otherwise' or 'default'
clauses in 'metadirective'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ otherwise (teams loop) \
+ otherwise (parallel loop) /* { dg-error "too many 'otherwise' or
'default' clauses in 'metadirective'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (bad_directive) /* { dg-error "unknown directive name before
'\\)' token" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ where (device={arch("nvptx")}: parallel loop) /* { dg-error "'where' is
not valid for 'metadirective'" } */ \
+ default (teams loop)
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ otherwise (teams loop) \
+ when (device={arch("nvptx")}: parallel loop) /* { dg-error "'otherwise'
or 'default' clause must appear last" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ when (device={arch("nvptx")} parallel loop) /* { dg-error "expected ':'
before 'parallel'" } */ \
+ default (teams loop)
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ #pragma omp metadirective \
+ default (metadirective default (flush)) /* { dg-error "metadirectives
cannot be used as variants of a 'metadirective' before 'default'" } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+
+ /* Test improperly nested metadirectives - even though the second
+ metadirective resolves to 'omp nothing', that is not the same as there
+ being literally nothing there. */
+ #pragma omp metadirective \
+ when (implementation={vendor("gnu")}: parallel for)
+ #pragma omp metadirective \
+ when (implementation={vendor("cray")}: parallel for)
+ /* { dg-error "loop nest expected before '#pragma'" "" { target c } .-2
} */
+ /* { dg-error "loop nest expected" "" { target c++ } .-3 } */
+ for (i = 0; i < N; i++) c[i] = a[i] * b[i];
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-2.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-2.c
new file mode 100644
index 00000000000..4b05cb96ab2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-2.c
@@ -0,0 +1,75 @@
+/* { dg-do compile } */
+
+#define N 100
+
+int
+main (void)
+{
+ int x = 0;
+ int y = 0;
+
+ /* Test implicit default (nothing). */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: barrier)
+ x = 1;
+
+ /* Test with multiple standalone directives. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: barrier) \
+ default (flush)
+ x = 1;
+
+ /* Test combining a standalone directive with one that takes a statement
+ body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel) \
+ default (barrier)
+ x = 1;
+
+ /* Test combining a standalone directive with one that takes a for loop. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel for) \
+ default (barrier)
+ for (int i = 0; i < N; i++)
+ x += i;
+
+ /* Test combining a directive that takes a for loop with one that takes
+ a regular statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: parallel for) \
+ default (parallel)
+ for (int i = 0; i < N; i++)
+ x += i;
+
+ /* Test labels inside statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams num_teams(512)) \
+ when (device={arch("gcn")}: teams num_teams(256)) \
+ default (teams num_teams(4))
+ {
+ if (x)
+ goto l1;
+ else
+ goto l2;
+ l1: ;
+ l2: ;
+ }
+
+ /* Test local labels inside statement body. */
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams num_teams(512)) \
+ when (device={arch("gcn")}: teams num_teams(256)) \
+ default (teams num_teams(4))
+ {
+ //__label__ l1, l2;
+
+ if (x)
+ goto l1;
+ else
+ goto l2;
+ l1: ;
+ l2: ;
+ }
+
+ return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-3.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-3.c
new file mode 100644
index 00000000000..0ac0d1d329d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-3.c
@@ -0,0 +1,23 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+void
+f (int x[], int y[], int z[])
+{
+ int i;
+
+ #pragma omp target map(to: x, y) map(from: z)
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams loop) \
+ default (parallel loop)
+ for (i = 0; i < N; i++)
+ z[i] = x[i] * y[i];
+}
+
+/* If offload device "nvptx" isn't supported, the front end can eliminate
+ that alternative and not produce a metadirective at all. Otherwise this
+ won't be resolved until late. */
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" {
target { ! offload_nvptx } } } } */
+/* { dg-final { scan-tree-dump "#pragma omp metadirective" "gimple" { target {
offload_nvptx } } } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-4.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c
new file mode 100644
index 00000000000..a257f5894ff
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c
@@ -0,0 +1,42 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+#pragma omp declare target
+void
+f (double a[], double x) {
+ int i;
+
+ #pragma omp metadirective \
+ when (construct={target}: distribute parallel for) \
+ default (parallel for simd)
+ for (i = 0; i < N; i++)
+ a[i] = x * i;
+}
+#pragma omp end declare target
+
+int
+main (void)
+{
+ 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 "otherwise:" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-5.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-5.c
new file mode 100644
index 00000000000..63ce4498cd1
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-5.c
@@ -0,0 +1,25 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+
+#define N 100
+
+void
+f (int a[], int flag)
+{
+ int i;
+ #pragma omp metadirective \
+ when (user={condition(flag)}: \
+ target teams distribute parallel for map(from: a[0:N])) \
+ default (parallel for)
+ for (i = 0; i < N; i++)
+ a[i] = i;
+}
+
+/* The metadirective should be resolved at parse time. */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "original" } }
*/
+/* { dg-final { scan-tree-dump-times "#pragma omp target" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp teams" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp distribute" 1 "original" }
} */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 2 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 2 "original" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-6.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-6.c
new file mode 100644
index 00000000000..5115d8ecbab
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-6.c
@@ -0,0 +1,32 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-original" } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+#define N 100
+
+void
+bar (int a[], int run_parallel, int run_guided)
+{
+ #pragma omp metadirective \
+ when (user={condition(run_parallel)}: parallel)
+ {
+ int i;
+ #pragma omp metadirective \
+ when (construct={parallel}, user={condition(run_guided)}: \
+ for schedule(guided)) \
+ when (construct={parallel}: for schedule(static))
+ for (i = 0; i < N; i++)
+ a[i] = i;
+ }
+ }
+
+/* The outer metadirective should be resolved at parse time. */
+/* The inner metadirective should be resolved during Gimplificiation. */
+
+/* { dg-final { scan-tree-dump-times "#pragma omp metadirective" 2 "original"
} } */
+/* { dg-final { scan-tree-dump-times "#pragma omp parallel" 1 "original" } } */
+/* { dg-final { scan-tree-dump-times "#pragma omp for" 4 "original" } } */
+/* { dg-final { scan-tree-dump-times "when \\(construct = .parallel" 4
"original" } } */
+/* { dg-final { scan-tree-dump-times "otherwise:" 2 "original" } } */
+
+/* { dg-final { scan-tree-dump-not "#pragma omp metadirective" "gimple" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-7.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-7.c
new file mode 100644
index 00000000000..09bf21cb278
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-7.c
@@ -0,0 +1,36 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple -fdump-tree-ompdevlow" } */
+
+#define N 256
+
+void
+f (int a[], int num)
+{
+ int i;
+
+ #pragma omp metadirective \
+ when (target_device={device_num(num), kind("gpu"), arch("nvptx")}: \
+ target parallel for map(tofrom: a[0:N])) \
+ when (target_device={device_num(num), kind("gpu"), \
+ arch("amdgcn"), isa("gfx906")}: \
+ target parallel for) \
+ when (target_device={device_num(num), kind("cpu"), arch("x86_64")}: \
+ parallel for)
+ for (i = 0; i < N; i++)
+ a[i] += i;
+
+ #pragma omp metadirective \
+ when (target_device={kind("gpu"), arch("nvptx")}: \
+ target parallel for map(tofrom: a[0:N]))
+ for (i = 0; i < N; i++)
+ a[i] += i;
+}
+
+/* We expect one "pragma omp target" with "device(num)" for each target_device
+ selector that specifies "device_num(num)". */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num"
3 "gimple" } } */
+
+/* One OMP_TARGET_DEVICE_MATCHES for each kind/arch/isa selector. These
+ are supposed to go away after ompdevlow. */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 9 "gimple" }
} */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow"
} } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-8.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-8.c
new file mode 100644
index 00000000000..0cab9620cd9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-8.c
@@ -0,0 +1,17 @@
+/* { dg-do compile } */
+
+#define N 256
+
+void
+f (void)
+{
+ int i;
+ int a[N];
+
+ #pragma omp metadirective \
+ when( device={kind(nohost)}: nothing ) \
+ when( device={arch("nvptx")}: nothing) \
+ default( parallel for)
+ for (i = 0; i < N; i++)
+ a[i] = i;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c
new file mode 100644
index 00000000000..3103656d71f
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-construct.c
@@ -0,0 +1,178 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */
+
+#include <stdlib.h>
+
+static void
+init (int n, double *a)
+{
+ for (int i = 0; i < n; i++)
+ a[i] = (double) i;
+}
+
+static void
+check (int n, double *a, double s)
+{
+ for (int i = 0; i < n; i++)
+ if (a[i] != (double) i * s)
+ abort ();
+}
+
+typedef void (transform_fn) (int, double *, double);
+
+static void doit (transform_fn *f, int n, double *a, double s)
+{
+ init (n, a);
+ (*f) (n, a, s);
+ check (n, a, s);
+}
+
+/* Check various combinations for enforcing correct ordering of
+ construct matches. */
+static void
+f1 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={target} \
+ : for) \
+ default (error at(execution) message("f1 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f2 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={teams, parallel} \
+ : for) \
+ default (error at(execution) message("f2 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f3 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={target, teams, parallel} \
+ : for) \
+ default (error at(execution) message("f3 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f4 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={target, parallel} \
+ : for) \
+ default (error at(execution) message("f4 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f5 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={target, teams} \
+ : for) \
+ default (error at(execution) message("f5 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Next batch is for things where the construct doesn't match the context. */
+static void
+f6 (int n, double* a, double s)
+{
+#pragma omp target
+#pragma omp teams
+#pragma omp metadirective \
+ when (construct={parallel} \
+ : error at(execution) message("f6 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f7 (int n, double* a, double s)
+{
+#pragma omp target
+#pragma omp teams
+#pragma omp metadirective \
+ when (construct={target, parallel} \
+ : error at(execution) message("f7 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+static void
+f8 (int n, double* a, double s)
+{
+#pragma omp target
+#pragma omp teams
+#pragma omp metadirective \
+ when (construct={parallel, target} \
+ : error at(execution) message("f8 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Next test choosing the best alternative when there are multiple
+ matches. */
+static void
+f9 (int n, double* a, double s)
+{
+#pragma omp target teams
+#pragma omp parallel
+#pragma omp metadirective \
+ when (construct={teams, parallel} \
+ : error at(execution) message("f9 match incorrect 1")) \
+ when (construct={target, teams, parallel} \
+ : for) \
+ when (construct={target, teams} \
+ : error at(execution) message("f9 match incorrect 2")) \
+ default (error at(execution) message("f9 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Note there are no tests for the matching the extended simd clause
+ syntax, which is only useful for "declare variant". */
+
+#define N 10
+#define S 2.0
+
+int
+main (void)
+{
+ double a[N];
+ doit (f1, N, a, S);
+ doit (f2, N, a, S);
+ doit (f3, N, a, S);
+ doit (f4, N, a, S);
+ doit (f5, N, a, S);
+ doit (f6, N, a, S);
+ doit (f7, N, a, S);
+ doit (f8, N, a, S);
+ doit (f9, N, a, S);
+}
+
+/* All the error calls should be optimized away. */
+/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-device.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-device.c
new file mode 100644
index 00000000000..09b795eeabe
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-device.c
@@ -0,0 +1,149 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */
+/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=sse -msse" {
target x86_64-*-* } } */
+
+#include <stdlib.h>
+
+static void
+init (int n, double *a)
+{
+ for (int i = 0; i < n; i++)
+ a[i] = (double) i;
+}
+
+static void
+check (int n, double *a, double s)
+{
+ for (int i = 0; i < n; i++)
+ if (a[i] != (double) i * s)
+ abort ();
+}
+
+typedef void (transform_fn) (int, double *, double);
+
+static void
+doit (transform_fn *f, int n, double *a, double s)
+{
+ init (n, a);
+ (*f) (n, a, s);
+ check (n, a, s);
+}
+
+/* Check kind=host matches (with offloading disabled). */
+static void
+f1 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={kind(host)} \
+ : parallel for) \
+ default (error at(execution) message("f1 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Check kind=nohost does not match (with offloading disabled). */
+static void
+f2 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={kind(nohost)} \
+ : error at(execution) message("f2 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Check arch. Either DEVICE_ARCH is defined by command-line option,
+ or we know it is not x86_64. */
+#ifdef DEVICE_ARCH
+static void
+f3 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch(DEVICE_ARCH)} \
+ : parallel for) \
+ default (error at(execution) message("f3 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f3 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch("x86_64")} \
+ : error at(execution) message("f3 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+/* Check both kind and arch together. */
+#ifdef DEVICE_ARCH
+static void
+f4 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch(DEVICE_ARCH), kind(host)} \
+ : parallel for) \
+ default (error at(execution) message("f4 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f4 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch("x86_64"), kind(host)} \
+ : error at(execution) message("f4 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+/* Check kind, arch, and ISA together. */
+#if defined(DEVICE_ARCH) && defined(DEVICE_ISA)
+static void
+f5 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)}
\
+ : parallel for) \
+ default (error at(execution) message("f5 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f5 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (device={arch("x86_64"), kind(host), isa("sse")} \
+ : error at(execution) message("f5 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+#define N 10
+#define S 2.0
+
+int
+main (void)
+{
+ double a[N];
+ doit (f1, N, a, S);
+ doit (f2, N, a, S);
+ doit (f3, N, a, S);
+ doit (f4, N, a, S);
+ doit (f5, N, a, S);
+}
+
+/* All the metadirectives involving the device selector should be
+ fully resolved and the error calls optimized away. */
+
+/* { dg-final { scan-tree-dump-not "__builtin_GOMP_error" "optimized" } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c
new file mode 100644
index 00000000000..1f1053eaffa
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-no-score.c
@@ -0,0 +1,95 @@
+/* { dg-do compile { target x86_64-*-* } } */
+/* { dg-additional-options "-foffload=disable" } */
+
+/* This test is expected to fail with compile-time errors:
+ "A trait-score cannot be specified in traits from the construct,
+ device or target_device trait-selector-sets." */
+
+/* Define this to avoid dependence on libgomp header files. */
+
+#define omp_initial_device -1
+
+void
+f1 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (device={kind (score(5) : host)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .device.
trait-selector-set" "" { target *-*-*} .-2 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f2 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (device={kind (host), arch (score(6) : x86_64), isa (avx512f)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .device.
trait-selector-set" "" { target *-*-*} .-2 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f3 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (device={kind (host), arch (score(6) : x86_64), \
+ isa (score(7): avx512f)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .device.
trait-selector-set" "" { target *-*-*} .-3 } */
+ /* { dg-error ".score. cannot be specified in traits in the .device.
trait-selector-set" "" { target *-*-*} .-3 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f4 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={device_num (score(42) : omp_initial_device), \
+ kind (host)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .target_device.
trait-selector-set" "" { target *-*-*} .-3 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f5 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={device_num(omp_initial_device), \
+ kind (score(5) : host)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .target_device.
trait-selector-set" "" { target *-*-*} .-2 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f6 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={device_num(omp_initial_device), kind (host), \
+ arch (score(6) : x86_64), isa (avx512f)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .target_device.
trait-selector-set" "" { target *-*-*} .-2 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+void
+f7 (int n, double *a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={device_num(omp_initial_device), kind (host), \
+ arch (score(6) : x86_64), \
+ isa (score(7): avx512f)} \
+ : parallel for)
+ /* { dg-error ".score. cannot be specified in traits in the .target_device.
trait-selector-set" "" { target *-*-*} .-3 } */
+ /* { dg-error ".score. cannot be specified in traits in the .target_device.
trait-selector-set" "" { target *-*-*} .-3 } */
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c
b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c
new file mode 100644
index 00000000000..a95b5fc59b2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device.c
@@ -0,0 +1,149 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-foffload=disable -fdump-tree-optimized" } */
+/* { dg-additional-options "-DDEVICE_ARCH=x86_64 -DDEVICE_ISA=mmx -mmmx" {
target x86_64-*-* } } */
+
+#include <stdlib.h>
+
+static void
+init (int n, double *a)
+{
+ for (int i = 0; i < n; i++)
+ a[i] = (double) i;
+}
+
+static void
+check (int n, double *a, double s)
+{
+ for (int i = 0; i < n; i++)
+ if (a[i] != (double) i * s)
+ abort ();
+}
+
+typedef void (transform_fn) (int, double *, double);
+
+static void
+doit (transform_fn *f, int n, double *a, double s)
+{
+ init (n, a);
+ (*f) (n, a, s);
+ check (n, a, s);
+}
+
+/* Check kind=host matches (with offloading disabled). */
+static void
+f1 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={kind(host)} \
+ : parallel for) \
+ default (error at(execution) message("f1 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Check kind=nohost does not match (with offloading disabled). */
+static void
+f2 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={kind(nohost)} \
+ : error at(execution) message("f2 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+
+/* Check arch. Either DEVICE_ARCH is defined by command-line option,
+ or we know it is not x86_64. */
+#ifdef DEVICE_ARCH
+static void
+f3 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch(DEVICE_ARCH)} \
+ : parallel for) \
+ default (error at(execution) message("f3 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f3 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch("x86_64")} \
+ : error at(execution) message("f3 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+/* Check both kind and arch together. */
+#ifdef DEVICE_ARCH
+static void
+f4 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch(DEVICE_ARCH), kind(host)} \
+ : parallel for) \
+ default (error at(execution) message("f4 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f4 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch("x86_64"), kind(host)} \
+ : error at(execution) message("f4 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+/* Check kind, arch, and ISA together. */
+#if defined(DEVICE_ARCH) && defined(DEVICE_ISA)
+static void
+f5 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch(DEVICE_ARCH), kind(host), isa(DEVICE_ISA)} \
+ : parallel for) \
+ default (error at(execution) message("f5 match failed"))
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#else
+static void
+f5 (int n, double* a, double s)
+{
+#pragma omp metadirective \
+ when (target_device={arch("x86_64"), kind(host), isa("mmx")} \
+ : error at(execution) message("f5 match failed")) \
+ default (parallel for)
+ for (int i = 0; i < n; i++)
+ a[i] = a[i] * s;
+}
+#endif
+
+#define N 10
+#define S 2.0
+
+int
+main (void)
+{
+ double a[N];
+ doit (f1, N, a, S);
+ doit (f2, N, a, S);
+ doit (f3, N, a, S);
+ doit (f4, N, a, S);
+ doit (f5, N, a, S);
+}
+
+/* Since the target_device selector is dynamic, none of the
+ error checks can be optimized away. */
+
+/* { dg-final { scan-tree-dump-times "GOMP_error" 5 "optimized" } } */
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c
new file mode 100644
index 00000000000..a57d6fd5671
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-1.c
@@ -0,0 +1,37 @@
+/* { dg-do run } */
+
+#define N 100
+
+void
+f (int x[], int y[], int z[])
+{
+ int i;
+
+ #pragma omp target map(to: x[0:N], y[0:N]) map(from: z[0:N])
+ #pragma omp metadirective \
+ when (device={arch("nvptx")}: teams loop) \
+ default (parallel loop)
+ for (i = 0; i < N; i++)
+ z[i] = x[i] * y[i];
+}
+
+int
+main (void)
+{
+ int x[N], y[N], z[N];
+ int i;
+
+ for (i = 0; i < N; i++)
+ {
+ x[i] = i;
+ y[i] = -i;
+ }
+
+ f (x, y, z);
+
+ for (i = 0; i < N; i++)
+ if (z[i] != x[i] * y[i])
+ return 1;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c
new file mode 100644
index 00000000000..9f714d61787
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c
@@ -0,0 +1,43 @@
+/* { dg-do run } */
+
+#include <math.h>
+
+#define N 100
+#define EPSILON 0.001
+
+#pragma omp declare target
+void
+f (double a[], double x) {
+ int i;
+
+ #pragma omp metadirective \
+ when (construct={target}: distribute parallel for) \
+ default (parallel for simd)
+ for (i = 0; i < N; i++)
+ a[i] = x * i;
+}
+#pragma omp end declare target
+
+int
+main (void)
+{
+ double a[N];
+ int i;
+
+ #pragma omp target teams map(from: a[0:N])
+ f (a, M_PI);
+
+ for (i = 0; i < N; i++)
+ if (fabs (a[i] - (M_PI * i)) > EPSILON)
+ return 1;
+
+ /* TODO: This does not execute a version of f with the default clause
+ active as might be expected. */
+ f (a, M_E); /* { dg-warning "direct calls to an offloadable function
containing metadirectives with a 'construct={target}' selector may produce
unexpected results" } */
+
+ for (i = 0; i < N; i++)
+ if (fabs (a[i] - (M_E * i)) > EPSILON)
+ return 1;
+
+ return 0;
+ }
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c
new file mode 100644
index 00000000000..b6879c92440
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-3.c
@@ -0,0 +1,36 @@
+/* { dg-do run } */
+
+#define N 100
+
+int
+f (int a[], int flag)
+{
+ int i;
+ int res = 0;
+
+ #pragma omp metadirective \
+ when (user={condition(!flag)}: \
+ target teams distribute parallel for \
+ map(from: a[0:N]) private(res)) \
+ default (parallel for)
+ for (i = 0; i < N; i++)
+ {
+ a[i] = i;
+ res = 1;
+ }
+
+ return res;
+}
+
+int
+main (void)
+{
+ int a[N];
+
+ if (f (a, 0))
+ return 1;
+ if (!f (a, 1))
+ return 1;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c
new file mode 100644
index 00000000000..62eb3e9c0ef
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-4.c
@@ -0,0 +1,54 @@
+/* { dg-do run } */
+
+#include <omp.h>
+
+#define N 100
+
+int
+f (int a[], int run_parallel, int run_static)
+{
+ int is_parallel = 0;
+ int is_static = 0;
+
+ #pragma omp metadirective \
+ when (user={condition(run_parallel)}: parallel)
+ {
+ int i;
+
+ if (omp_in_parallel ())
+ is_parallel = 1;
+
+ #pragma omp metadirective \
+ when (construct={parallel}, user={condition(!run_static)}: \
+ for schedule(guided) private(is_static)) \
+ when (construct={parallel}: for schedule(static))
+ for (i = 0; i < N; i++)
+ {
+ a[i] = i;
+ is_static = 1;
+ }
+ }
+
+ return (is_parallel << 1) | is_static;
+}
+
+int
+main (void)
+{
+ int a[N];
+
+ /* is_static is always set if run_parallel is false. */
+ if (f (a, 0, 0) != 1)
+ return 1;
+
+ if (f (a, 0, 1) != 1)
+ return 1;
+
+ if (f (a, 1, 0) != 2)
+ return 1;
+
+ if (f (a, 1, 1) != 3)
+ return 1;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c
new file mode 100644
index 00000000000..e869932de6d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-5.c
@@ -0,0 +1,48 @@
+/* { dg-do run } */
+
+#define N 100
+
+#include <stdio.h>
+#include <omp.h>
+
+int
+f (int a[], int num)
+{
+ int on_device = 0;
+ int i;
+
+ #pragma omp metadirective \
+ when (target_device={device_num(num), kind("gpu")}: \
+ target parallel for map(to: a[0:N]), map(from: on_device)) \
+ default (parallel for private (on_device))
+ for (i = 0; i < N; i++)
+ {
+ a[i] += i;
+ on_device = 1;
+ }
+
+ return on_device;
+}
+
+int
+main (void)
+{
+ int a[N];
+ int on_device_count = 0;
+ int i;
+
+ for (i = 0; i < N; i++)
+ a[i] = i;
+
+ for (i = 0; i <= omp_get_num_devices (); i++)
+ on_device_count += f (a, i);
+
+ if (on_device_count != omp_get_num_devices ())
+ return 1;
+
+ for (i = 0; i < N; i++)
+ if (a[i] != 2 * i)
+ return 2;
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-1.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-1.c
new file mode 100644
index 00000000000..e9dcd3f23d6
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-1.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+/* { dg-additional-options -O0 } */
+
+/* Test late resolution of metadirectives with dynamic selectors
+ in "declare simd" functions. All the variants do the same thing;
+ the purpose of this test is to ensure that the "condition" predicates
+ are all called, and in the correct order. */
+
+static int pcount = 0;
+
+static int __attribute__ ((noinline))
+ptrue (int n)
+{
+ pcount++;
+ if (pcount != n)
+ __builtin_abort ();
+ return 1;
+}
+
+static int __attribute__ ((noinline))
+pfalse (int n)
+{
+ pcount++;
+ if (pcount != n)
+ __builtin_abort ();
+ return 0;
+}
+
+#define N 256
+
+#pragma omp declare simd
+void
+f (int a[])
+{
+ int i;
+
+#pragma omp metadirective \
+ when (construct={simd}: \
+ nothing) \
+ when (user={condition(score (100): pfalse (1))}: \
+ nothing) \
+ when (user={condition(score (90): pfalse (2))}: \
+ nothing) \
+ when (user={condition(score (70): (ptrue (5) && pfalse (6)))}: \
+ nothing) \
+ when (user={condition(score (80): (pfalse (3) || pfalse (4)))}: \
+ nothing) \
+ when (user={condition(score (60): \
+ (ptrue (7) ? pfalse (8) : ptrue (8)))}: \
+ nothing) \
+ otherwise (nothing)
+ for (i = 0; i < N; i++)
+ a[i] += i;
+}
+
+int a[N];
+
+int
+main (void)
+{
+ f (a);
+ for (int i = 0; i < N; i++)
+ if (a[i] != i)
+ return 1;
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-2.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-2.c
new file mode 100644
index 00000000000..af333bf79c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-late-2.c
@@ -0,0 +1,66 @@
+/* { dg-do run } */
+/* { dg-additional-options -O3 } */
+
+/* Test late resolution of metadirectives with dynamic selectors
+ in "declare simd" functions. All the variants do the same thing;
+ the purpose of this test is to ensure that the "condition" predicates
+ are all called, and in the correct order. */
+
+static int pcount = 0;
+
+static int __attribute__ ((noinline))
+ptrue (int n)
+{
+ pcount++;
+ if (pcount != n)
+ __builtin_abort ();
+ return 1;
+}
+
+static int __attribute__ ((noinline))
+pfalse (int n)
+{
+ pcount++;
+ if (pcount != n)
+ __builtin_abort ();
+ return 0;
+}
+
+#define N 256
+
+#pragma omp declare simd
+void
+f (int a[])
+{
+ int i;
+
+#pragma omp metadirective \
+ when (construct={simd}: \
+ nothing) \
+ when (user={condition(score (100): pfalse (1))}: \
+ nothing) \
+ when (user={condition(score (90): pfalse (2))}: \
+ nothing) \
+ when (user={condition(score (70): (ptrue (5) && pfalse (6)))}: \
+ nothing) \
+ when (user={condition(score (80): (pfalse (3) || pfalse (4)))}: \
+ nothing) \
+ when (user={condition(score (60): \
+ (ptrue (7) ? pfalse (8) : ptrue (8)))}: \
+ nothing) \
+ otherwise (nothing)
+ for (i = 0; i < N; i++)
+ a[i] += i;
+}
+
+int a[N];
+
+int
+main (void)
+{
+ f (a);
+ for (int i = 0; i < N; i++)
+ if (a[i] != i)
+ return 1;
+ return 0;
+}
diff --git
a/libgomp/testsuite/libgomp.c-c++-common/metadirective-target-device.c
b/libgomp/testsuite/libgomp.c-c++-common/metadirective-target-device.c
new file mode 100644
index 00000000000..7c2f7e5ad52
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-target-device.c
@@ -0,0 +1,76 @@
+#include <omp.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+/* PR112779 item (B) */
+
+/* Check that the target_device selector correctly matches device numbers
+ and handles kind=host|nohost|any. */
+
+static int
+check_explicit_device (int d, int expect_host)
+{
+ int ok = 0;
+ if (expect_host)
+ {
+ #pragma omp metadirective \
+ when (target_device={device_num(d), kind("host")} : nothing) \
+ otherwise (error at(execution) message("check_explicit_device host"))
+ ok = 1;
+ }
+ else
+ {
+ #pragma omp metadirective \
+ when (target_device={device_num(d), kind("nohost")} : nothing) \
+ otherwise (error at(execution) message("check_explicit_device nohost"))
+ ok = 1;
+ }
+
+ return ok;
+}
+
+static int
+check_implicit_device (int d, int expect_host)
+{
+ int ok = 0;
+ omp_set_default_device (d);
+
+ if (expect_host)
+ {
+ #pragma omp metadirective \
+ when (target_device={kind("host")} : nothing) \
+ otherwise (error at(execution) message("check_implicit_device host"))
+ ok = 1;
+ }
+ else
+ {
+ #pragma omp metadirective \
+ when (target_device={kind("nohost")} : nothing) \
+ otherwise (error at(execution) message("check_implicit_device nohost"))
+ ok = 1;
+ }
+ #pragma omp metadirective \
+ when (target_device={kind("any")} : nothing) \
+ otherwise (error at(execution) message("check_implicit_device any"))
+ ok = 1;
+ omp_set_default_device (omp_initial_device);
+
+ return ok;
+}
+
+int
+main (void)
+{
+ printf ("Checking omp_initial_device\n");
+ check_explicit_device (omp_initial_device, 1);
+ check_implicit_device (omp_initial_device, 1);
+ int n = omp_get_num_devices ();
+ printf ("There are %d devices\n", n);
+ for (int i = 0; i < n; i++)
+ {
+ printf ("Checking device %d\n", i);
+ check_explicit_device (i, 0);
+ check_implicit_device (i, 0);
+ }
+ return 0;
+}
--
2.25.1
