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 <clt...@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 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 (); +} -- 2.17.1