OpenMP: Shared metadirective/dynamic selector tests for C and C++

gcc/testsuite/ChangeLog
	* c-c++-common/gomp/adjust-args-6.c: New.
	* 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-1.c: New.
	* c-c++-common/gomp/metadirective-target-device-2.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 <kcy@codesourcery.com>
Co-Authored-By: Sandra Loosemore <sandra@codesourcery.com>

25 files changed, 1434 insertions, 0 deletions

diff --git a/gcc/testsuite/c-c++-common/gomp/adjust-args-6.c b/gcc/testsuite/c-c++-common/gomp/adjust-args-6.c
new file mode 100644
index 0000000..1d5fc4d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/adjust-args-6.c
@@ -0,0 +1,26 @@
+/* { dg-do compile } */
+/* { dg-additional-options "-fdump-tree-gimple" } */
+
+/* Ensure that adjust_args is only applied when variant substitution happens. */
+
+extern int flag;
+
+void h(int *);
+void f(int *);
+#pragma omp declare variant(f) match(construct={dispatch}, user={condition(flag)}) adjust_args(need_device_ptr : x)
+void g(int *x);
+
+void foo(int *y)
+{
+  #pragma omp dispatch
+    h(y);
+  #pragma omp dispatch
+    f(y);
+  #pragma omp dispatch
+    g(y);  /* { dg-bogus "late or dynamic variant resolution" "" { xfail *-*-* } } */
+}
+
+/* { dg-final { scan-tree-dump-times "h \\(y\\);" 1 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "f \\(y\\);" 2 "gimple" } } */
+/* { dg-final { scan-tree-dump-times "D\.\[0-9]+ = __builtin_omp_get_mapped_ptr \\(y, D\.\[0-9]+\\);" 1 "gimple" { xfail *-*-* } } } */
+/* { dg-final { scan-tree-dump-times "f \\(D\.\[0-9]+\\);" 1 "gimple" { xfail *-*-* } } } */
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 0000000..275eab1
--- /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 0000000..ff401c8
--- /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 0000000..31dd054
--- /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 0000000..5804eb5
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-4.c
@@ -0,0 +1,41 @@
+/* { 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);
+
+  f (a, 2.71828);
+
+  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 0000000..b24818a
--- /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 0000000..b1338a5
--- /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 0000000..f12955d
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/attrs-metadirective-7.c
@@ -0,0 +1,42 @@
+/* { 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;
+}
+
+/* For configurations with offloading, we expect one "pragma omp target"
+   with "device(num)" for each target_device selector that specifies
+   "device_num(num)".  Without offloading, there should be zero as the
+   resolution happens during gimplification.  */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 3 "gimple" { target offloading_enabled } } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 0 "gimple" { target { ! offloading_enabled } } } } */
+
+/* For configurations with offloading, expect 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" { target offloading_enabled } } } */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } } */
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 0000000..43746b6
--- /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 0000000..38bfe92
--- /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 0000000..b406a31
--- /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 0000000..c078123
--- /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 0000000..b2bfe95
--- /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 0000000..4b05cb9
--- /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 0000000..0ac0d1d
--- /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 0000000..5cf3dbd
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-4.c
@@ -0,0 +1,40 @@
+/* { 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);
+
+  f (a, 2.71828);
+
+  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 0000000..63ce449
--- /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 0000000..5115d8e
--- /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 0000000..4ad103b
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-7.c
@@ -0,0 +1,41 @@
+/* { 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;
+}
+
+/* For configurations with offloading, we expect one "pragma omp target"
+   with "device(num)" for each target_device selector that specifies
+   "device_num(num)".  Without offloading, there should be zero as the
+   resolution happens during gimplification.  */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 3 "gimple" { target offloading_enabled } } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target\[^\\n\]* device\\(num" 0 "gimple" { target { ! offloading_enabled } } } } */
+
+/* For configurations with offloading, expect 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" { target offloading_enabled } } } */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "gimple" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump-times "OMP_TARGET_DEVICE_MATCHES" 0 "ompdevlow" { target offloading_enabled } } } */
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 0000000..0cab962
--- /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 0000000..3103656
--- /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 0000000..09b795e
--- /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 0000000..1f1053e
--- /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-1.c b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device-1.c
new file mode 100644
index 0000000..6373349
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device-1.c
@@ -0,0 +1,151 @@
+/* { dg-do compile }  */
+/* { dg-additional-options "-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);
+}
+
+/* The gimplifier emits all 5 alternatives.  In configurations without
+   offloading, the dynamic selector can be completely resolved, otherwise
+   the calls to GOMP_error must remain in the generated code pending
+   runtime evaluation of the selectors.  */
+/* { dg-final { scan-tree-dump-times "GOMP_error" 5 "optimized" { target offloading_enabled } } } */
+/* { dg-final { scan-tree-dump-times "GOMP_error" 0 "optimized" { target { ! offloading_enabled } } } } */
diff --git a/gcc/testsuite/c-c++-common/gomp/metadirective-target-device-2.c b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device-2.c
new file mode 100644
index 0000000..24584f2
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/gomp/metadirective-target-device-2.c
@@ -0,0 +1,132 @@
+/* { dg-do compile */
+/* { dg-additional-options "-fdump-tree-optimized" } */
+
+/* In configurations without offloading configured, we can resolve many
+   instances of the target_device context selector at gimplification time
+   instead of waiting until late resolution.  */
+
+/* Device 0 may be either the host or an offloading device, in configurations
+   that support them.  */
+void
+f1 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(0), kind(host)}		\
+	: parallel for)						\
+  default (error at(execution) message("f1 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f1 match failed" "optimized" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump "f1 match failed" "optimized" { target offloading_enabled } } } */
+
+/* Device -1 is always the host, even in offloading configurations.  */
+void
+f2 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(-1), kind(host)}		\
+	: parallel for)						\
+  default (error at(execution) message("f2 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f2 match failed" "optimized" } } */
+
+/* Constant device numbers < -1 cause the match to fail, even in offloading
+   configurations.  */
+void
+f3 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(-42), kind(host)}		\
+	: error at(execution) message("f3 match failed"))	\
+  default (parallel for)
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f3 match failed" "optimized" } } */
+
+/* In a non-offloading configuration, devices > 0 are always invalid and
+   will cause the match to fail.  In an offload configuration, we don't
+   know at compile-time if the device number is valid or if it refers to
+   a host or offload device.  */
+void
+f4 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(42), kind(host)}		\
+	: error at(execution) message("f4 match failed"))	\
+  default (parallel for)
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f4 match failed" "optimized" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump "f4 match failed" "optimized" { target offloading_enabled } } } */
+
+extern int omp_get_initial_device (void);
+extern int omp_get_default_device (void);
+extern int omp_get_num_devices (void);
+extern int omp_get_device_num (void);
+
+/* On a non-offloading configuration, omp_get_default_device() always refers
+   to the host.  With offloading, we know it's a valid device number but
+   not whether it's the host.  */
+void
+f5 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(omp_get_default_device ()), kind(host)} \
+	: parallel for)						\
+  default (error at(execution) message("f5 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f5 match failed" "optimized" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump "f5 match failed" "optimized" { target offloading_enabled } } } */
+
+/* omp_get_initial_device() always refers to the host, whether offloading is
+   configured or not.  */
+void
+f6 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(omp_get_initial_device ()), kind(host)} \
+	: parallel for)						\
+  default (error at(execution) message("f6 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f6 match failed" "optimized" } } */
+
+/* omp_get_num_devices() returns the number of non-host offload devices,
+   therefore using it as a device number always refers to the host, whether
+   offloading is enabled or not.  */
+void
+f7 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(omp_get_num_devices ()), kind(host)} \
+	: parallel for)						\
+  default (error at(execution) message("f7 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f7 match failed" "optimized" } } */
+
+/* omp_get_device_num() always returns a valid device number but we don't
+   know whether it's the host or offload device, in configurations that
+   support offloading.  */
+void
+f8 (int n, double* a, double s)
+{
+#pragma omp metadirective					\
+  when (target_device={device_num(omp_get_device_num ()), kind(host)} \
+	: parallel for)						\
+  default (error at(execution) message("f8 match failed"))
+  for (int i = 0; i < n; i++)
+    a[i] = a[i] * s;
+}
+/* { dg-final { scan-tree-dump-not "f8 match failed" "optimized" { target { ! offloading_enabled } } } } */
+/* { dg-final { scan-tree-dump "f8 match failed" "optimized" { target offloading_enabled } } } */
+