[C++,OpenMP,5.0,OG11] Fixes for lambda in offload regions

Message ID 0a267fdd-2348-3aba-d5ce-6920daabebdd@codesourcery.com
State New
Headers show
Series
  • [C++,OpenMP,5.0,OG11] Fixes for lambda in offload regions
Related show

Commit Message

Chung-Lin Tang June 17, 2021, 1:55 p.m.
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.

Added another testcase to test if lambda works in these cases.
Tested without regressions on devel/omp/gcc-11, pushed there.

Jakub, this technically is a further bug fix for the PR92120 v3 patch.
I'll submit a v4 for mainline trunk later, or this patch independently
in case the v3 patch is already reviewed by then.

Thanks,
Chung-Lin

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.
From dbf5d72f4c077215330e5b06fbb9b3311b807c2a Mon Sep 17 00:00:00 2001
From: Chung-Lin Tang <cltang@codesourcery.com>
Date: Thu, 17 Jun 2021 21:53:10 +0800
Subject: [PATCH] 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.
---
 gcc/cp/semantics.c                            | 22 ++++++++++--
 gcc/testsuite/g++.dg/gomp/target-lambda-2.C   | 35 +++++++++++++++++++
 .../testsuite/libgomp.c++/target-lambda-2.C   | 30 ++++++++++++++++
 3 files changed, 85 insertions(+), 2 deletions(-)
 create mode 100644 gcc/testsuite/g++.dg/gomp/target-lambda-2.C
 create mode 100644 libgomp/testsuite/libgomp.c++/target-lambda-2.C

Patch

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 ();
+}