@@ -7241,6 +7241,11 @@ extern tree finish_omp_structured_block (tree);
extern tree finish_oacc_data (tree, tree);
extern tree finish_oacc_host_data (tree, tree);
extern tree finish_omp_construct (enum tree_code, tree, tree);
+
+extern tree finish_omp_target (location_t, tree, tree, bool);
+extern void set_omp_target_this_expr (tree);
+
+
extern tree begin_omp_parallel (void);
extern tree finish_omp_parallel (tree, tree);
extern tree begin_omp_task (void);
@@ -842,6 +842,9 @@ lambda_expr_this_capture (tree lambda, int add_capture_p)
type cast (_expr.cast_ 5.4) to the type of 'this'. [ The cast
ensures that the transformed expression is an rvalue. ] */
result = rvalue (result);
+
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (result);
}
return result;
@@ -40742,8 +40742,6 @@ static bool
cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
enum pragma_context context, bool *if_p)
{
- tree *pc = NULL, stmt;
-
if (flag_openmp)
omp_requires_mask
= (enum omp_requires) (omp_requires_mask | OMP_REQUIRES_TARGET_USED);
@@ -40796,6 +40794,7 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
keep_next_level (true);
tree sb = begin_omp_structured_block (), ret;
unsigned save = cp_parser_begin_omp_structured_block (parser);
+ set_omp_target_this_expr (NULL_TREE);
switch (ccode)
{
case OMP_TEAMS:
@@ -40847,15 +40846,9 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
cclauses[C_OMP_CLAUSE_SPLIT_TARGET] = tc;
}
}
- tree stmt = make_node (OMP_TARGET);
- TREE_TYPE (stmt) = void_type_node;
- OMP_TARGET_CLAUSES (stmt) = cclauses[C_OMP_CLAUSE_SPLIT_TARGET];
- OMP_TARGET_BODY (stmt) = body;
- OMP_TARGET_COMBINED (stmt) = 1;
- SET_EXPR_LOCATION (stmt, pragma_tok->location);
- add_stmt (stmt);
- pc = &OMP_TARGET_CLAUSES (stmt);
- goto check_clauses;
+ finish_omp_target (pragma_tok->location,
+ cclauses[C_OMP_CLAUSE_SPLIT_TARGET], body, true);
+ return true;
}
else if (!flag_openmp) /* flag_openmp_simd */
{
@@ -40892,46 +40885,13 @@ cp_parser_omp_target (cp_parser *parser, cp_token *pragma_tok,
return false;
}
- stmt = make_node (OMP_TARGET);
- TREE_TYPE (stmt) = void_type_node;
-
- OMP_TARGET_CLAUSES (stmt)
- = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
- "#pragma omp target", pragma_tok);
- pc = &OMP_TARGET_CLAUSES (stmt);
+ tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", pragma_tok);
keep_next_level (true);
- OMP_TARGET_BODY (stmt) = cp_parser_omp_structured_block (parser, if_p);
+ set_omp_target_this_expr (NULL_TREE);
+ tree body = cp_parser_omp_structured_block (parser, if_p);
- SET_EXPR_LOCATION (stmt, pragma_tok->location);
- add_stmt (stmt);
-
-check_clauses:
- while (*pc)
- {
- if (OMP_CLAUSE_CODE (*pc) == OMP_CLAUSE_MAP)
- switch (OMP_CLAUSE_MAP_KIND (*pc))
- {
- case GOMP_MAP_TO:
- case GOMP_MAP_ALWAYS_TO:
- case GOMP_MAP_FROM:
- case GOMP_MAP_ALWAYS_FROM:
- case GOMP_MAP_TOFROM:
- case GOMP_MAP_ALWAYS_TOFROM:
- case GOMP_MAP_ALLOC:
- case GOMP_MAP_FIRSTPRIVATE_POINTER:
- case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
- case GOMP_MAP_ALWAYS_POINTER:
- break;
- default:
- error_at (OMP_CLAUSE_LOCATION (*pc),
- "%<#pragma omp target%> with map-type other "
- "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
- "on %<map%> clause");
- *pc = OMP_CLAUSE_CHAIN (*pc);
- continue;
- }
- pc = &OMP_CLAUSE_CHAIN (*pc);
- }
+ finish_omp_target (pragma_tok->location, clauses, body, false);
return true;
}
@@ -61,6 +61,8 @@ static hash_map<tree, tree> *omp_private_member_map;
static vec<tree> omp_private_member_vec;
static bool omp_private_member_ignore_next;
+/* Used for OpenMP target region 'this' references. */
+static tree omp_target_this_expr = NULL_TREE;
/* Deferred Access Checking Overview
---------------------------------
@@ -1958,6 +1960,7 @@ tree
finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
{
gcc_assert (TREE_CODE (decl) == FIELD_DECL);
+ tree orig_object = object;
bool try_omp_private = !object && omp_private_member_map;
tree ret;
@@ -1996,6 +1999,14 @@ finish_non_static_data_member (tree decl, tree object, tree qualifying_scope)
return error_mark_node;
}
+ if (orig_object == NULL_TREE)
+ {
+ tree this_expr = TREE_OPERAND (object, 0);
+
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (this_expr);
+ }
+
if (current_class_ptr)
TREE_USED (current_class_ptr) = 1;
if (processing_template_decl)
@@ -2783,8 +2794,15 @@ finish_this_expr (void)
}
if (result)
- /* The keyword 'this' is a prvalue expression. */
- return rvalue (result);
+ {
+ /* The keyword 'this' is a prvalue expression. */
+ result = rvalue (result);
+
+ /* Acknowledge to OpenMP target that 'this' was referenced. */
+ set_omp_target_this_expr (result);
+
+ return result;
+ }
tree fn = current_nonlambda_function ();
if (fn && DECL_STATIC_FUNCTION_P (fn))
@@ -8519,6 +8537,109 @@ finish_omp_construct (enum tree_code code, tree body, tree clauses)
return add_stmt (stmt);
}
+void
+set_omp_target_this_expr (tree this_val)
+{
+ omp_target_this_expr = this_val;
+}
+
+tree
+finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
+{
+ if (DECL_LAMBDA_FUNCTION_P (current_function_decl))
+ {
+ tree closure = DECL_ARGUMENTS (current_function_decl);
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TO);
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (closure);
+ OMP_CLAUSE_SIZE (c) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (closure)));
+
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_DECL (c2) = closure;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = clauses;
+ OMP_CLAUSE_CHAIN (c) = c2;
+ clauses = c;
+
+ if (omp_target_this_expr)
+ {
+ STRIP_NOPS (omp_target_this_expr);
+ gcc_assert (DECL_HAS_VALUE_EXPR_P (omp_target_this_expr));
+ omp_target_this_expr = DECL_VALUE_EXPR (omp_target_this_expr);
+ tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_TOFROM);
+ OMP_CLAUSE_DECL (c3) = build_simple_mem_ref (omp_target_this_expr);
+ OMP_CLAUSE_SIZE (c3)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+
+ tree c4 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c4, GOMP_MAP_ALWAYS_POINTER);
+
+ OMP_CLAUSE_DECL (c4) = omp_target_this_expr;
+ OMP_CLAUSE_SIZE (c4) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c3) = c4;
+ OMP_CLAUSE_CHAIN (c4) = OMP_CLAUSE_CHAIN (c2);
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ omp_target_this_expr = NULL_TREE;
+ }
+ }
+ else if (omp_target_this_expr)
+ {
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_TOFROM);
+ OMP_CLAUSE_DECL (c) = build_simple_mem_ref (omp_target_this_expr);
+ OMP_CLAUSE_SIZE (c)
+ = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (omp_target_this_expr)));
+
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
+ STRIP_NOPS (omp_target_this_expr);
+ OMP_CLAUSE_DECL (c2) = omp_target_this_expr;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ OMP_CLAUSE_CHAIN (c2) = clauses;
+ OMP_CLAUSE_CHAIN (c) = c2;
+ clauses = c;
+ omp_target_this_expr = NULL_TREE;
+ }
+
+ tree stmt = make_node (OMP_TARGET);
+ TREE_TYPE (stmt) = void_type_node;
+ OMP_TARGET_CLAUSES (stmt) = clauses;
+ OMP_TARGET_BODY (stmt) = body;
+ OMP_TARGET_COMBINED (stmt) = combined_p;
+ SET_EXPR_LOCATION (stmt, loc);
+
+ tree c = clauses;
+ while (c)
+ {
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (c))
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_ALWAYS_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_ALWAYS_TOFROM:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_FIRSTPRIVATE_POINTER:
+ case GOMP_MAP_FIRSTPRIVATE_REFERENCE:
+ case GOMP_MAP_ALWAYS_POINTER:
+ break;
+ default:
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "%<#pragma omp target%> with map-type other "
+ "than %<to%>, %<from%>, %<tofrom%> or %<alloc%> "
+ "on %<map%> clause");
+ break;
+ }
+ c = OMP_CLAUSE_CHAIN (c);
+ }
+ return add_stmt (stmt);
+}
+
tree
finish_omp_parallel (tree clauses, tree body)
{
new file mode 100644
@@ -0,0 +1,33 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+extern "C" void abort ();
+
+struct S
+{
+ int a, b, c, d;
+
+ int sum (void)
+ {
+ int val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ int sum_offload (void)
+ {
+ int val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
new file mode 100644
@@ -0,0 +1,49 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-do compile }
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+
+extern "C" void abort ();
+
+struct T
+{
+ int x, y;
+
+ auto sum_func (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+ auto sum_func_offload (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ #pragma omp target map(from:v)
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+};
+
+int main (void)
+{
+ T a = { 1, 2 };
+
+ auto s1 = a.sum_func (3);
+ auto s2 = a.sum_func_offload (3);
+
+ if (s1 (1) != s2 (1))
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {map\(to:\*__closure \[len: [0-9]+\]\) map\(firstprivate:__closure \[pointer assign, bias: 0\]\) map\(tofrom:\*_[0-9]+ \[len: [0-9]+\]\) map\(always_pointer:__closure->__this \[pointer assign, bias: 0\]\)} "gimple" } } */
new file mode 100644
@@ -0,0 +1,29 @@
+extern "C" void abort ();
+
+struct S
+{
+ int a, b, c, d;
+
+ int sum (void)
+ {
+ int val = 0;
+ val += a + b + this->c + this->d;
+ return val;
+ }
+
+ int sum_offload (void)
+ {
+ int val = 0;
+ #pragma omp target map(val)
+ val += a + b + this->c + this->d;
+ return val;
+ }
+};
+
+int main (void)
+{
+ S s = { 1, 2, 3, 4 };
+ if (s.sum () != s.sum_offload ())
+ abort ();
+ return 0;
+}
new file mode 100644
@@ -0,0 +1,47 @@
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+
+extern "C" void abort ();
+
+struct T
+{
+ int x, y;
+
+ auto sum_func (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+ auto sum_func_offload (int n)
+ {
+ auto fn = [=](int m) -> int
+ {
+ int v;
+ #pragma omp target map(from:v)
+ v = (x + y) * n + m;
+ return v;
+ };
+ return fn;
+ }
+
+};
+
+int main (void)
+{
+ T a = { 1, 2 };
+
+ auto s1 = a.sum_func (3);
+ auto s2 = a.sum_func_offload (3);
+
+ if (s1 (1) != s2 (1))
+ abort ();
+
+ return 0;
+}
Hi Jakub, this patch implements automatically adding map(tofrom: this[:1]) to omp target regions inside non-static member functions, as specified in OpenMP 5.0. This patch factors away some parts of cp_parser_omp_target, into a new finish_omp_target function, and implements the new clause adding there. For target regions in normal non-static member functions, the case is more simple. For the inside lambda function case, this is implemented by copying the entire __closure as a "to" map first (and yeah, this patch allows target regions inside lambda functions to largely work, but since it's just a copying of __closure, the capture by reference case still shouldn't work yet). __closure->this is then implemented by an always_pointer map clause. I've added two testcases, as both compiler scan testcases and libgomp executable test. Testing of g++ and libgomp both are regression free with nvptx offloading. Is this okay for trunk? Thanks, Chung-Lin 2020-09-16 Chung-Lin Tang <cltang@codesourcery.com> PR middle-end/92120 gcc/cp/ * cp-tree.h (finish_omp_target): New declaration. (set_omp_target_this_expr): Likewise. * lambda.c (lambda_expr_this_capture): Add call to set_omp_target_this_expr. * parser.c (cp_parser_omp_target): Factor out code, change to call finish_omp_target, add re-initing call to set_omp_target_this_expr. * semantics.c (omp_target_this_expr): New static variable. (finish_non_static_data_member): Add call to set_omp_target_this_expr. (finish_this_expr): Likewise. (set_omp_target_this_expr): New function to set omp_target_this_expr. (finish_omp_target): New function with code merged from cp_parser_omp_target, plus code to add this and __closure map clauses for OpenMP. gcc/testsuite/ * g++.dg/gomp/target-this-1.C: New testcase. * g++.dg/gomp/target-this-2.C: New testcase. libgomp/testsuite/ * libgomp.c++/target-this-1.C: New testcase. * libgomp.c++/target-this-2.C: New testcase.