c-pragma.c (oacc_pragmas): Add entry for declare directive.

2015-11-12  James Norris  <jnorris@codesourcery.com>
	    Joseph Myers  <joseph@codesourcery.com>

	gcc/c-family/
	* c-pragma.c (oacc_pragmas): Add entry for declare directive. 
	* c-pragma.h (enum pragma_kind): Add PRAGMA_OACC_DECLARE.
	(enum pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT and
	PRAGMA_OACC_CLAUSE_LINK.

	gcc/c/
	* c-parser.c (c_parser_pragma): Handle PRAGMA_OACC_DECLARE.
	(c_parser_omp_clause_name): Handle 'device_resident' clause.
	(c_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(c_parser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OACC_CLAUSE_LINK.
	(OACC_DECLARE_CLAUSE_MASK): New definition.
	(c_parser_oacc_declare): New function.

	gcc/cp/
	* parser.c (cp_parser_omp_clause_name): Handle 'device_resident'
	clause.
	(cp_parser_oacc_data_clause): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(cp_paser_oacc_all_clauses): Handle PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT
	and PRAGMA_OMP_CLAUSE_LINK.
	(OACC_DECLARE_CLAUSE_MASK): New definition.
	(cp_parser_oacc_declare): New function.
	(cp_parser_pragma): Handle PRAGMA_OACC_DECLARE.
	* pt.c (tsubst_expr): Handle OACC_DECLARE.

	gcc/
	* gimple-pretty-print.c (dump_gimple_omp_target): Handle
	GF_OMP_TARGET_KIND_OACC_DECLARE. 
	* gimple.h (enum gf_mask): Add GF_OMP_TARGET_KIND_OACC_DECLARE.
	(is_gomple_omp_oacc): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
	* gimplify.c (oacc_declare_returns): New.
	(gimplify_bind_expr): Prepend 'exit' stmt to cleanup.
	(device_resident_p): New function.
	(oacc_default_clause): Handle device_resident clause.
	(gimplify_oacc_declare_1, gimplify_oacc_declare): New functions.
	(gimplify_expr): Handle OACC_DECLARE.
	* omp-builtins.def (BUILT_IN_GOACC_DECLARE): New builtin.
	* omp-low.c (expand_omp_target): Handle
	GF_OMP_TARGET_KIND_OACC_DECLARE and BUILTIN_GOACC_DECLARE.
	(build_omp_regions_1): Handlde GF_OMP_TARGET_KIND_OACC_DECLARE.
	(lower_omp_target): Handle GF_OMP_TARGET_KIND_OACC_DECLARE,
	GOMP_MAP_DEVICE_RESIDENT and GOMP_MAP_LINK.
	(make_gimple_omp_edges): Handle GF_OMP_TARGET_KIND_OACC_DECLARE.
	* tree-pretty-print.c (dump_omp_clause): Handle GOMP_MAP_LINK and
	GOMP_MAP_DEVICE_RESIDENT.

	gcc/testsuite
	* c-c++-common/goacc/declare-1.c: New test.
	* c-c++-common/goacc/declare-2.c: Likewise.

	include/
	* gomp-constants.h (enum gomp_map_kind): Add GOMP_MAP_DEVICE_RESIDENT
	and GOMP_MAP_LINK.

	libgomp/
	* libgomp.map (GOACC_2.0.1): Export GOACC_declare.
	* oacc-parallel.c (GOACC_declare): New function.
	* testsuite/libgomp.oacc-c-c++-common/declare-1.c: New test.
	* testsuite/libgomp.oacc-c-c++-common/declare-2.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/declare-4.c: Likewise.
	* testsuite/libgomp.oacc-c-c++-common/declare-5.c: Likewise.
	* testsuite/libgomp.oacc-c++/declare-1.C: Likewise.

Co-Authored-By: Joseph Myers <joseph@codesourcery.com>

From-SVN: r230275

5 files changed, 273 insertions, 0 deletions

diff --git a/libgomp/testsuite/libgomp.oacc-c++/declare-1.C b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
new file mode 100644
index 00000000000..0286955d0c7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/declare-1.C
@@ -0,0 +1,31 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+
+template<class T>
+T foo()
+{
+  T a, b;
+  #pragma acc declare create (a)
+
+  #pragma acc parallel copyout (b)
+  {
+    a = 5;
+    b = a;
+  }
+
+  return b;
+}
+
+int
+main (void)
+{
+  int rc;
+
+  rc = foo<int>();
+
+  if (rc != 5)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
new file mode 100644
index 00000000000..c63a68dbab7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-1.c
@@ -0,0 +1,122 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <openacc.h>
+#include <stdlib.h>
+#include <stdio.h>
+
+#define N 8
+
+void
+subr2 (int *a)
+{
+  int i;
+  int f[N];
+#pragma acc declare copyout (f)
+
+#pragma acc parallel copy (a[0:N])
+  {
+    for (i = 0; i < N; i++)
+      {
+	f[i] = a[i];
+	a[i] = f[i] + f[i] + f[i];
+      }
+  }
+}
+
+void
+subr1 (int *a)
+{
+  int f[N];
+#pragma acc declare copy (f)
+
+#pragma acc parallel copy (a[0:N])
+  {
+    int i;
+
+    for (i = 0; i < N; i++)
+      {
+	f[i] = a[i];
+	a[i] = f[i] + f[i];
+      }
+  }
+}
+
+int b[8];
+#pragma acc declare create (b)
+
+int d[8] = { 1, 2, 3, 4, 5, 6, 7, 8 };
+#pragma acc declare copyin (d)
+
+int
+main (int argc, char **argv)
+{
+  int a[N];
+  int e[N];
+#pragma acc declare create (e)
+  int i;
+
+  for (i = 0; i < N; i++)
+    a[i] = i + 1;
+
+  if (!acc_is_present (&b, sizeof (b)))
+    abort ();
+
+  if (!acc_is_present (&d, sizeof (d)))
+    abort ();
+
+  if (!acc_is_present (&e, sizeof (e)))
+    abort ();
+
+#pragma acc parallel copyin (a[0:N])
+  {
+    for (i = 0; i < N; i++)
+      {
+	b[i] = a[i];
+	a[i] = b[i];
+      }
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != i + 1)
+	abort ();
+    }
+
+#pragma acc parallel copy (a[0:N])
+  {
+    for (i = 0; i < N; i++)
+      {
+	e[i] = a[i] + d[i];
+	a[i] = e[i];
+      }
+  }
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != (i + 1) * 2)
+	abort ();
+    }
+
+  for (i = 0; i < N; i++)
+    {
+      a[i] = 1234;
+    }
+
+  subr1 (&a[0]);
+
+  for (i = 0; i < N; i++)
+    {
+      if (a[i] != 1234 * 2)
+	abort ();
+    }
+
+  subr2 (&a[0]);
+
+  for (i = 0; i < 1; i++)
+    {
+      if (a[i] != 1234 * 6)
+	abort ();
+    }
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
new file mode 100644
index 00000000000..2078a33afa9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-2.c
@@ -0,0 +1,64 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+
+#define N 16
+
+float c[N];
+#pragma acc declare device_resident (c)
+
+#pragma acc routine
+float
+subr2 (float a)
+{
+  int i;
+
+  for (i = 0; i < N; i++)
+    c[i] = 2.0;
+
+  for (i = 0; i < N; i++)
+    a += c[i];
+
+  return a;
+}
+
+float b[N];
+#pragma acc declare copyin (b)
+
+#pragma acc routine
+float
+subr1 (float a)
+{
+  int i;
+
+  for (i = 0; i < N; i++)
+    a += b[i];
+
+  return a;
+}
+
+int
+main (int argc, char **argv)
+{
+  float a;
+  int i;
+
+  for (i = 0; i < 16; i++)
+    b[i] = 1.0;
+
+  a = 0.0;
+
+  a = subr1 (a);
+
+  if (a != 16.0)
+    abort ();
+
+  a = 0.0;
+
+  a = subr2 (a);
+
+  if (a != 32.0)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
new file mode 100644
index 00000000000..013310ecb7d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-4.c
@@ -0,0 +1,41 @@
+/* { dg-do run  { target openacc_nvidia_accel_selected } } */
+
+#include <stdlib.h>
+#include <openacc.h>
+
+float b;
+#pragma acc declare link (b)
+
+#pragma acc routine
+int
+func (int a)
+{
+  b = a + 1;
+
+  return b;
+}
+
+int
+main (int argc, char **argv)
+{
+  float a;
+
+  a = 2.0;
+
+#pragma acc parallel copy (a)
+  {
+    b = a;
+    a = 1.0;
+    a = a + b;
+  }
+
+  if (a != 3.0)
+    abort ();
+
+  a = func (a);
+
+  if (a != 4.0)
+    abort ();
+
+  return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
new file mode 100644
index 00000000000..38c5de063d9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/declare-5.c
@@ -0,0 +1,15 @@
+/* { dg-do run { target openacc_nvidia_accel_selected } } */
+
+#include <stdio.h>
+
+int
+main (int argc, char **argv)
+{
+  int a[8] __attribute__((unused));
+
+  fprintf (stderr, "CheCKpOInT\n");
+#pragma acc declare present (a)
+}
+
+/* { dg-output "CheCKpOInT" } */
+/* { dg-shouldfail "" } */