openmp: Also implicitly mark as declare target to functions mentioned in target regions

OpenMP 5.0 also specifies that functions referenced from target regions
(except for target regions with device(ancestor:)) are also implicitly declare target to.

This patch implements that.

2020-05-14  Jakub Jelinek  <jakub@redhat.com>

	* function.h (struct function): Add has_omp_target bit.
	* omp-offload.c (omp_discover_declare_target_fn_r): New function,
	old renamed to ...
	(omp_discover_declare_target_tgt_fn_r): ... this.
	(omp_discover_declare_target_var_r): Call
	omp_discover_declare_target_tgt_fn_r instead of
	omp_discover_declare_target_fn_r.
	(omp_discover_implicit_declare_target): Also queue functions with
	has_omp_target bit set, for those walk with
	omp_discover_declare_target_fn_r, for declare target to functions
	walk with omp_discover_declare_target_tgt_fn_r.
gcc/c/
	* c-parser.c (c_parser_omp_target): Set cfun->has_omp_target.
gcc/cp/
	* cp-gimplify.c (cp_genericize_r): Set cfun->has_omp_target.
gcc/fortran/
	* trans-openmp.c: Include function.h.
	(gfc_trans_omp_target): Set cfun->has_omp_target.
libgomp/
	* testsuite/libgomp.c-c++-common/target-40.c: New test.

2 files changed, 55 insertions, 0 deletions

diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 104c527..4bfce69 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,7 @@
+2020-05-14  Jakub Jelinek  <jakub@redhat.com>
+
+	* testsuite/libgomp.c-c++-common/target-40.c: New test.
+
 2020-05-13  Tobias Burnus  <tobias@codesourcery.com>
 
 	PR fortran/94690
diff --git a/libgomp/testsuite/libgomp.c-c++-common/target-40.c b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
new file mode 100644
index 0000000..22bbdd9
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/target-40.c
@@ -0,0 +1,51 @@
+/* { dg-do run } */
+/* { dg-options "-O0" } */
+
+extern
+#ifdef __cplusplus
+"C"
+#endif
+void abort (void);
+volatile int v;
+#pragma omp declare target to (v)
+typedef void (*fnp1) (void);
+typedef fnp1 (*fnp2) (void);
+void f1 (void) { v++; }
+void f2 (void) { v += 4; }
+void f3 (void) { v += 16; f1 (); }
+fnp1 f4 (void) { v += 64; return f2; }
+int a = 1;
+int *b = &a;
+int **c = &b;
+fnp2 f5 (void) { f3 (); return f4; }
+#pragma omp declare target to (c)
+
+int
+main ()
+{
+  int err = 0;
+  #pragma omp target map(from:err)
+  {
+    volatile int xa;
+    int *volatile xb;
+    int **volatile xc;
+    fnp2 xd;
+    fnp1 xe;
+    err = 0;
+    xa = a;
+    err |= xa != 1;
+    xb = b;
+    err |= xb != &a;
+    xc = c;
+    err |= xc != &b;
+    xd = f5 ();
+    err |= v != 17;
+    xe = xd ();
+    err |= v != 81;
+    xe ();
+    err |= v != 85;
+  }
+  if (err)
+    abort ();
+  return 0;
+}