openmp: Fix handling of non-addressable shared scalars in parallel nested inside of target [PR93515]

As the following testcase shows, we need to consider even target to be a construct
that forces not to use copy in/out for shared on parallel inside of the target.
E.g. for parallel nested inside another parallel or host teams, we already avoid
copy in/out and we need to treat target the same.

2020-02-06  Jakub Jelinek  <jakub@redhat.com>

	PR libgomp/93515
	* omp-low.c (use_pointer_for_field): For nested constructs, also
	look for map clauses on target construct.
	(scan_omp_1_stmt) <case GIMPLE_OMP_TARGET>: Bump temporarily
	taskreg_nesting_level.

	* testsuite/libgomp.c-c++-common/pr93515.c: New test.

4 files changed, 73 insertions, 7 deletions

diff --git a/gcc/ChangeLog b/gcc/ChangeLog
index df6bbe6..b5b465a 100644
--- a/gcc/ChangeLog
+++ b/gcc/ChangeLog
@@ -1,6 +1,12 @@
 2020-02-06  Jakub Jelinek  <jakub@redhat.com>
 
 	PR libgomp/93515
+	* omp-low.c (use_pointer_for_field): For nested constructs, also
+	look for map clauses on target construct.
+	(scan_omp_1_stmt) <case GIMPLE_OMP_TARGET>: Bump temporarily
+	taskreg_nesting_level.
+
+	PR libgomp/93515
 	* gimplify.c (gimplify_scan_omp_clauses) <do_notice>: If adding
 	shared clause, call omp_notice_variable on outer context if any.
 
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index eb3fe96..67565d6 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -477,18 +477,30 @@ use_pointer_for_field (tree decl, omp_context *shared_ctx)
 	  omp_context *up;
 
 	  for (up = shared_ctx->outer; up; up = up->outer)
-	    if (is_taskreg_ctx (up) && maybe_lookup_decl (decl, up))
+	    if ((is_taskreg_ctx (up)
+		 || (gimple_code (up->stmt) == GIMPLE_OMP_TARGET
+		     && is_gimple_omp_offloaded (up->stmt)))
+		&& maybe_lookup_decl (decl, up))
 	      break;
 
 	  if (up)
 	    {
 	      tree c;
 
-	      for (c = gimple_omp_taskreg_clauses (up->stmt);
-		   c; c = OMP_CLAUSE_CHAIN (c))
-		if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
-		    && OMP_CLAUSE_DECL (c) == decl)
-		  break;
+	      if (gimple_code (up->stmt) == GIMPLE_OMP_TARGET)
+		{
+		  for (c = gimple_omp_target_clauses (up->stmt);
+		       c; c = OMP_CLAUSE_CHAIN (c))
+		    if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+			&& OMP_CLAUSE_DECL (c) == decl)
+		      break;
+		}
+	      else
+		for (c = gimple_omp_taskreg_clauses (up->stmt);
+		     c; c = OMP_CLAUSE_CHAIN (c))
+		  if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_SHARED
+		      && OMP_CLAUSE_DECL (c) == decl)
+		    break;
 
 	      if (c)
 		goto maybe_mark_addressable_and_ret;
@@ -3781,7 +3793,14 @@ scan_omp_1_stmt (gimple_stmt_iterator *gsi, bool *handled_ops_p,
       break;
 
     case GIMPLE_OMP_TARGET:
-      scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+      if (is_gimple_omp_offloaded (stmt))
+	{
+	  taskreg_nesting_level++;
+	  scan_omp_target (as_a <gomp_target *> (stmt), ctx);
+	  taskreg_nesting_level--;
+	}
+      else
+	scan_omp_target (as_a <gomp_target *> (stmt), ctx);
       break;
 
     case GIMPLE_OMP_TEAMS:
diff --git a/libgomp/ChangeLog b/libgomp/ChangeLog
index 56e9da7..ba14005 100644
--- a/libgomp/ChangeLog
+++ b/libgomp/ChangeLog
@@ -1,3 +1,8 @@
+2020-02-06  Jakub Jelinek  <jakub@redhat.com>
+
+	PR libgomp/93515
+	* testsuite/libgomp.c-c++-common/pr93515.c: New test.
+
 2020-02-05  Tobias Burnus  <tobias@codesourcery.com>
 
 	* testsuite/lib/libgomp.exp
diff --git a/libgomp/testsuite/libgomp.c-c++-common/pr93515.c b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c
new file mode 100644
index 0000000..8a69088
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c-c++-common/pr93515.c
@@ -0,0 +1,36 @@
+/* PR libgomp/93515 */
+
+#include <omp.h>
+#include <stdlib.h>
+
+int
+main ()
+{
+  int i;
+  int a = 42;
+#pragma omp target teams distribute parallel for defaultmap(tofrom: scalar)
+  for (i = 0; i < 64; ++i)
+    if (omp_get_team_num () == 0)
+      if (omp_get_thread_num () == 0)
+	a = 142;
+  if (a != 142)
+    __builtin_abort ();
+  a = 42;
+#pragma omp target parallel for defaultmap(tofrom: scalar)
+  for (i = 0; i < 64; ++i)
+    if (omp_get_thread_num () == 0)
+      a = 143;
+  if (a != 143)
+    __builtin_abort ();
+  a = 42;
+#pragma omp target firstprivate(a)
+  {
+    #pragma omp parallel for
+    for (i = 0; i < 64; ++i)
+      if (omp_get_thread_num () == 0)
+	a = 144;
+    if (a != 144)
+      abort ();
+  }
+  return 0;
+}