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>

8 files changed, 424 insertions, 0 deletions

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 0000000..a57d6fd
--- /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 0000000..9b5e5cf
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/metadirective-2.c
@@ -0,0 +1,41 @@
+/* { 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;
+
+  f (a, M_E);
+
+  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 0000000..b6879c9
--- /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 0000000..62eb3e9
--- /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 0000000..e869932
--- /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 0000000..e9dcd3f
--- /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 0000000..af333bf
--- /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 0000000..7c2f7e5
--- /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;
+}