[gcc/devel/omp/gcc-11] Fixes for lambda in offload regions

[gcc/devel/omp/gcc-11] Fixes for lambda in offload regions

[Chung-Lin Tang]

https://gcc.gnu.org/g:dbf5d72f4c077215330e5b06fbb9b3311b807c2a

commit dbf5d72f4c077215330e5b06fbb9b3311b807c2a
Author: Chung-Lin Tang <cltang@codesourcery.com>
Date:   Thu Jun 17 21:53:10 2021 +0800

    Fixes for lambda in offload regions
    
    This patch contains:
    
    (1) Some fixes for lambda capture by-reference to work inside
    offload regions.
    
    (2) Cases where lambda objects declared inside an offload region
    were mistakenly target-mapped on the enclosing target construct,
    causing a gimplify ICE (because it isn't binded at that position),
    added checks to avoid this.
    
    gcc/cp/ChangeLog:
    
            * semantics.c (struct omp_target_walk_data):
            Add 'hash_set<tree> local_decls' member.
            (finish_omp_target_clauses_r): Handle BIND_EXPR case, fill in
            local_decls there. Adjust case to not add locally declared lambda
            objects to data->lambda_objects_accessed.
            (finish_omp_target_clauses): Peel away TARGET_EXPR for lambda objects.
            Adjust map kind to _TOFROM for reference fields in closures.
    
    gcc/testsuite/ChangeLog:
    
            * g++.dg/gomp/target-lambda-2.C: New test.
    
    libgomp/ChangeLog:
    
            * testsuite/libgomp.c++/target-lambda-2.C: New test.

Diff:
---
 gcc/cp/semantics.c                              | 22 ++++++++++++++--
 gcc/testsuite/g++.dg/gomp/target-lambda-2.C     | 35 +++++++++++++++++++++++++
 libgomp/testsuite/libgomp.c++/target-lambda-2.C | 30 +++++++++++++++++++++
 3 files changed, 85 insertions(+), 2 deletions(-)

diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 25fa6cb5305..1f7eacfe701 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -9145,6 +9145,8 @@ struct omp_target_walk_data
 
   tree current_closure;
   hash_set<tree> closure_vars_accessed;
+
+  hash_set<tree> local_decls;
 };
 
 static tree
@@ -9203,12 +9205,25 @@ finish_omp_target_clauses_r (tree *tp, int *walk_subtrees, void *ptr)
       return NULL_TREE;
     }
 
+  if (TREE_CODE (t) == BIND_EXPR)
+    {
+      tree block = BIND_EXPR_BLOCK (t);
+      for (tree var = BLOCK_VARS (block); var; var = DECL_CHAIN (var))
+	if (!data->local_decls.contains (var))
+	  data->local_decls.add (var);
+      return NULL_TREE;
+    }
+
   if (TREE_TYPE(t) && LAMBDA_TYPE_P (TREE_TYPE (t)))
     {
       tree lt = TREE_TYPE (t);
       gcc_assert (CLASS_TYPE_P (lt));
 
-      if (!data->lambda_objects_accessed.contains (t))
+      if (!data->lambda_objects_accessed.contains (t)
+	  /* Do not prepare to create target maps for locally declared
+	     lambdas or anonymous ones.  */
+	  && !data->local_decls.contains (t)
+	  && TREE_CODE (t) != TARGET_EXPR)
 	data->lambda_objects_accessed.add (t);
       *walk_subtrees = 0;
       return NULL_TREE;
@@ -9494,6 +9509,9 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
 	   i != data.lambda_objects_accessed.end (); ++i)
 	{
 	  tree lobj = *i;
+	  if (TREE_CODE (lobj) == TARGET_EXPR)
+	    lobj = TREE_OPERAND (lobj, 0);
+
 	  tree lt = TREE_TYPE (lobj);
 	  gcc_assert (LAMBDA_TYPE_P (lt) && CLASS_TYPE_P (lt));
 
@@ -9530,7 +9548,7 @@ finish_omp_target_clauses (location_t loc, tree body, tree *clauses_ptr)
 		  tree exp = build3 (COMPONENT_REF, TREE_TYPE (fld),
 				     lobj, fld, NULL_TREE);
 		  tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
-		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+		  OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
 		  OMP_CLAUSE_DECL (c)
 		    = build1 (INDIRECT_REF, TREE_TYPE (TREE_TYPE (exp)), exp);
 		  OMP_CLAUSE_SIZE (c)
diff --git a/gcc/testsuite/g++.dg/gomp/target-lambda-2.C b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
new file mode 100644
index 00000000000..bdf2564cd04
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-lambda-2.C
@@ -0,0 +1,35 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}
+
+/* Make sure lambda objects do NOT appear in target maps.  */
+/* { dg-final { scan-tree-dump {(?n)#pragma omp target num_teams.* map\(from:Y \[len: [0-9]+\]\) map\(from:X \[len: [0-9]+\]\)$} "gimple" } } */
diff --git a/libgomp/testsuite/libgomp.c++/target-lambda-2.C b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
new file mode 100644
index 00000000000..1d3561ffbd7
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-lambda-2.C
@@ -0,0 +1,30 @@
+#include <cstdlib>
+
+#define N 10
+int main (void)
+{
+  int X, Y;
+  #pragma omp target map(from: X, Y)
+  {
+    int x = 0, y = 0;
+
+    for (int i = 0; i < N; i++)
+      [&] (int v) { x += v; } (i);
+
+    auto yinc = [&y] { y++; };
+    for (int i = 0; i < N; i++)
+      yinc ();
+
+    X = x;
+    Y = y;
+  }
+
+  int Xs = 0;
+  for (int i = 0; i < N; i++)
+    Xs += i;
+  if (X != Xs)
+    abort ();
+
+  if (Y != N)
+    abort ();
+}