@@ -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)
new file mode 100644
@@ -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" } } */
new file mode 100644
@@ -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 ();
+}
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