Hi Jakub,
there was a first version of this patch here:
https://gcc.gnu.org/pipermail/gcc-patches/2020-September/554087.html
The attached patch here is a v2 version that adds implementation of
this part in the this[:1] functionality description in the OpenMP 5.0 spec:
"if the [member] variable [accessed in a target region] is of a type pointer
or reference to pointer, it is also treated as if it has appeared in a map
clause as a zero-length array section."
Basically, referencing a pointer member 'ptr' automatically maps it with the
equivalent of 'map(this->ptr[:0])'
To achieve this, two new map kinds GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION,
and GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION were added, which are
basically split from GOMP_MAP_ATTACH and GOMP_MAP_POINTER, except now allowing
the pointer target to be NULL.
This patch has been tested for gcc, g++, gfortran (C and Fortran are not really
affected, but since omp-low.c was slightly touched, tested along for
completeness)
and libgomp on x86_64-linux with nvptx offloading, all without regressions.
Is this okay for trunk?
Thanks,
Chung-Lin
2020-11-13 Chung-Lin Tang <clt...@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.
(omp_target_ptr_members_accessed): New static hash_map for tracking
accessed non-static pointer-type members.
(finish_non_static_data_member): Add call to set_omp_target_this_expr.
Add recording of non-static pointer-type members access.
(finish_this_expr): Add call to set_omp_target_this_expr.
(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 implement this[:1] and __closure map
clauses for OpenMP.
gcc/
* omp-low.c (lower_omp_target):
Handle GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
* tree-pretty-print.c (dump_omp_clause): Likewise.
include/
* gomp-constants.h (enum gomp_map_kind):
Add GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION map kinds.
(GOMP_MAP_POINTER_P):
Include GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION.
libgomp/
* libgomp.h (gomp_attach_pointer): Add bool parameter.
* oacc-mem.c (acc_attach_async): Update call to gomp_attach_pointer.
(goacc_enter_data_internal): Likewise.
* target.c (gomp_map_vars_existing): Update assert condition to
include GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION.
(gomp_map_pointer): Add 'bool allow_zero_length_array_sections'
parameter, add support for mapping a pointer with NULL target.
(gomp_attach_pointer): Add 'bool allow_zero_length_array_sections'
parameter, add support for attaching a pointer with NULL target.
(gomp_map_vars_internal): Update calls to gomp_map_pointer and
gomp_attach_pointer, add handling for
GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION, and
GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION cases.
gcc/testsuite/
* g++.dg/gomp/target-this-1.C: New testcase.
* g++.dg/gomp/target-this-2.C: New testcase.
* g++.dg/gomp/target-this-3.C: New testcase.
* g++.dg/gomp/target-this-4.C: New testcase.
libgomp/
* testsuite/libgomp.c++/target-this-1.C: New testcase.
* testsuite/libgomp.c++/target-this-2.C: New testcase.
* testsuite/libgomp.c++/target-this-3.C: New testcase.
* testsuite/libgomp.c++/target-this-4.C: New testcase.
diff --git a/gcc/cp/cp-tree.h b/gcc/cp/cp-tree.h
index 63724c0e84f..e45540e08f1 100644
--- a/gcc/cp/cp-tree.h
+++ b/gcc/cp/cp-tree.h
@@ -7277,6 +7277,8 @@ extern void record_null_lambda_scope (tree);
extern void finish_lambda_scope (void);
extern tree start_lambda_function (tree fn, tree lambda_expr);
extern void finish_lambda_function (tree body);
+extern tree finish_omp_target (location_t, tree, tree, bool);
+extern void set_omp_target_this_expr (tree);
/* in tree.c */
extern int cp_tree_operand_length (const_tree);
diff --git a/gcc/cp/lambda.c b/gcc/cp/lambda.c
index 1a1647f465e..eb09971f288 100644
--- a/gcc/cp/lambda.c
+++ b/gcc/cp/lambda.c
@@ -841,6 +841,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;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index efcdce04777..0277de212ce 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -41132,8 +41132,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);
@@ -41186,6 +41184,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:
@@ -41237,15 +41236,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 */
{
@@ -41282,49 +41275,14 @@ 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);
- c_omp_adjust_map_clauses (OMP_TARGET_CLAUSES (stmt), true);
-
- pc = &OMP_TARGET_CLAUSES (stmt);
+ tree clauses = cp_parser_omp_all_clauses (parser, OMP_TARGET_CLAUSE_MASK,
+ "#pragma omp target", pragma_tok);
+ c_omp_adjust_map_clauses (clauses, true);
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:
- case GOMP_MAP_ATTACH_DETACH:
- 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;
}
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 0389198dab8..e31d4f19f43 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -61,6 +61,10 @@ 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;
+
+static hash_map<tree, tree> omp_target_ptr_members_accessed;
/* Deferred Access Checking Overview
---------------------------------
@@ -1958,6 +1962,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 +2001,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)
@@ -2056,6 +2069,14 @@ finish_non_static_data_member (tree decl, tree object,
tree qualifying_scope)
if (v)
ret = convert_from_reference (*v);
}
+ else if (omp_target_this_expr
+ && TREE_TYPE (ret)
+ && POINTER_TYPE_P (TREE_TYPE (ret)))
+ {
+ if (omp_target_ptr_members_accessed.get (decl) == NULL)
+ omp_target_ptr_members_accessed.put (decl, ret);
+ }
+
return ret;
}
@@ -2783,8 +2804,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))
@@ -8652,6 +8680,193 @@ 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;
+
+ if (omp_target_this_expr == NULL_TREE)
+ omp_target_ptr_members_accessed.empty ();
+}
+
+tree
+finish_omp_target (location_t loc, tree clauses, tree body, bool combined_p)
+{
+ tree last_inserted_clause = NULL_TREE;
+ 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;
+ last_inserted_clause = 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;
+ last_inserted_clause = c4;
+ 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;
+ last_inserted_clause = c2;
+ omp_target_this_expr = NULL_TREE;
+ }
+
+ if (last_inserted_clause && !omp_target_ptr_members_accessed.is_empty ())
+ for (hash_map<tree, tree>::iterator i
+ = omp_target_ptr_members_accessed.begin ();
+ i != omp_target_ptr_members_accessed.end (); ++i)
+ {
+ /* For each referenced member that is of pointer or reference-to-pointer
+ type, create the equivalent of map(alloc:this->ptr[:0]). */
+ tree field_decl = (*i).first;
+ tree ptr_member = (*i).second;
+ if (!cxx_mark_addressable (ptr_member))
+ gcc_unreachable ();
+
+ if (TREE_CODE (TREE_TYPE (field_decl)) == REFERENCE_TYPE)
+ {
+ /* For reference to pointers, we need to map the referenced pointer
+ first for things to be correct. */
+ tree ptr_member_type = TREE_TYPE (ptr_member);
+
+ /* Map pointer target as zero-length array section. */
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (c)
+ = build1 (INDIRECT_REF, TREE_TYPE (ptr_member_type), ptr_member);
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+ /* Map pointer to zero-length array section. */
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND
+ (c2, GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION);
+ OMP_CLAUSE_DECL (c2) = ptr_member;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+ /* Attach reference-to-pointer field to pointer. */
+ tree c3 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3, GOMP_MAP_ATTACH);
+ OMP_CLAUSE_DECL (c3) = TREE_OPERAND (ptr_member, 0);
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+ OMP_CLAUSE_CHAIN (c3) = OMP_CLAUSE_CHAIN (last_inserted_clause);
+
+ OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
+ last_inserted_clause = c3;
+ }
+ else if (TREE_CODE (TREE_TYPE (field_decl)) == POINTER_TYPE)
+ {
+ /* Map pointer target as zero-length array section. */
+ tree c = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_DECL (c)
+ = build2 (MEM_REF, char_type_node, ptr_member,
+ build_int_cst (build_pointer_type (char_type_node), 0));
+ OMP_CLAUSE_SIZE (c) = size_zero_node;
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c) = 1;
+
+ /* Attach zero-length array section to pointer. */
+ tree c2 = build_omp_clause (loc, OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND
+ (c2, GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
+ OMP_CLAUSE_DECL (c2) = ptr_member;
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+
+ OMP_CLAUSE_CHAIN (c) = c2;
+ OMP_CLAUSE_CHAIN (c2) = OMP_CLAUSE_CHAIN (last_inserted_clause);
+ OMP_CLAUSE_CHAIN (last_inserted_clause) = c;
+ last_inserted_clause = c2;
+ }
+ else
+ gcc_unreachable ();
+ }
+
+ 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:
+ case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ 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)
{
diff --git a/gcc/omp-low.c b/gcc/omp-low.c
index ed805e2e6d2..f63a14d1106 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11652,6 +11652,8 @@ lower_omp_target (gimple_stmt_iterator *gsi_p,
omp_context *ctx)
case GOMP_MAP_ALWAYS_POINTER:
case GOMP_MAP_ATTACH:
case GOMP_MAP_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
break;
case GOMP_MAP_IF_PRESENT:
case GOMP_MAP_FORCE_ALLOC:
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-1.C
b/gcc/testsuite/g++.dg/gomp/target-this-1.C
new file mode 100644
index 00000000000..de93a3e5e57
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-1.C
@@ -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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-2.C
b/gcc/testsuite/g++.dg/gomp/target-this-2.C
new file mode 100644
index 00000000000..a5e832130fb
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-2.C
@@ -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" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-3.C
b/gcc/testsuite/g++.dg/gomp/target-this-3.C
new file mode 100644
index 00000000000..208ea079b95
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-3.C
@@ -0,0 +1,105 @@
+// { dg-do compile }
+// { dg-additional-options "-fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+extern "C" void abort ();
+
+struct S
+{
+ int * ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ bool set_ptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr != NULL)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ }
+
+ bool set_refptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr != NULL)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ S s = { ptr1, N, ptr2, N };
+
+ bool mapped;
+ int val = 123;
+
+ mapped = s.set_ptr (val);
+ if (mapped)
+ abort ();
+ if (s.ptr != ptr1)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ mapped = s.set_refptr (val);
+ if (mapped)
+ abort ();
+ if (s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N])
+ mapped = s.set_ptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_refptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != val)
+ abort ();
+
+ #pragma omp target data map(ptr2[:N])
+ mapped = s.set_refptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_ptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != val)
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign,
bias: 0\]\) map\(alloc:\*this->refptr \[pointer assign, zero-length array
section, bias: 0\]\) map\(attach:this->refptr \[bias: 0\]\) map\(from:mapped
\[len: [0-9]+\]\) map\(alloc:\*_[0-9+] \[len: 0\]\) firstprivate\(n\)} "gimple"
} } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
map\(tofrom:\*this \[len: [0-9]+\]\) map\(firstprivate:this \[pointer assign,
bias: 0\]\) map\(attach_zero_length_array_section:this->ptr \[bias: 0\]\)
map\(from:mapped \[len: [0-9]+\]\) map\(alloc:MEM.* \[len: 0\]\)
firstprivate\(n\)} "gimple" } } */
diff --git a/gcc/testsuite/g++.dg/gomp/target-this-4.C
b/gcc/testsuite/g++.dg/gomp/target-this-4.C
new file mode 100644
index 00000000000..f42cf384541
--- /dev/null
+++ b/gcc/testsuite/g++.dg/gomp/target-this-4.C
@@ -0,0 +1,107 @@
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14 -fdump-tree-gimple" }
+#include <cstdlib>
+#include <cstring>
+
+struct T
+{
+ int *ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ auto set_ptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+
+ auto set_refptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ T a = { ptr1, N, ptr2, N };
+
+ auto p1 = a.set_ptr_func (1);
+ auto r2 = a.set_refptr_func (2);
+
+ if (p1 ())
+ abort ();
+ if (r2 ())
+ abort ();
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N], ptr2[:N])
+ {
+ if (!p1 ())
+ abort ();
+ if (!r2 ())
+ abort ();
+ }
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 1)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 2)
+ abort ();
+
+ return 0;
+}
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
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\]\)
map\(attach_zero_length_array_section:_[0-9]+->ptr \[bias: 0\]\)
map\(from:mapped \[len: 1\]\) map\(alloc:MEM.* \[len: 0\]\) firstprivate\(n\)
map\(alloc:MEM.* \[len: 0\]\) map\(firstprivate:this \[pointer assign, bias:
0\]\)} "gimple" } } */
+
+/* { dg-final { scan-tree-dump {#pragma omp target num_teams.*
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\]\)
map\(alloc:\*_[0-9]+->refptr \[pointer assign, zero-length array section, bias:
0\]\) map\(attach:_3->refptr \[bias: 0\]\) map\(from:mapped \[len: [0-9]+\]\)
map\(alloc:\*_[0-9]+ \[len: 0\]\) firstprivate\(n\) map\(alloc:MEM.* \[len:
0\]\) map\(firstprivate:this \[pointer assign, bias: 0\]\)} "gimple" } } */
diff --git a/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 318f048bcff..f6063923b55 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -798,6 +798,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc,
dump_flags_t flags)
{
case GOMP_MAP_ALLOC:
case GOMP_MAP_POINTER:
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, "alloc");
break;
case GOMP_MAP_IF_PRESENT:
@@ -876,6 +877,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc,
dump_flags_t flags)
case GOMP_MAP_ATTACH_DETACH:
pp_string (pp, "attach_detach");
break;
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
+ pp_string (pp, "attach_zero_length_array_section");
+ break;
default:
gcc_unreachable ();
}
@@ -894,6 +898,9 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc,
dump_flags_t flags)
case GOMP_MAP_ALWAYS_POINTER:
pp_string (pp, " [pointer assign, bias: ");
break;
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ pp_string (pp, " [pointer assign, zero-length array section,
bias: ");
+ break;
case GOMP_MAP_TO_PSET:
pp_string (pp, " [pointer set, len: ");
break;
@@ -901,6 +908,7 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc,
dump_flags_t flags)
case GOMP_MAP_DETACH:
case GOMP_MAP_FORCE_DETACH:
case GOMP_MAP_ATTACH_DETACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
pp_string (pp, " [bias: ");
break;
default:
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 309cbcadd83..f541ac0cba4 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -132,6 +132,11 @@ enum gomp_map_kind
No refcount is bumped by this, and the store is done unconditionally.
*/
GOMP_MAP_ALWAYS_POINTER = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_FLAG_SPECIAL | 1),
+ /* Like GOMP_MAP_POINTER, but allow zero-length array section, i.e. set to
+ NULL if target is not mapped. */
+ GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION
+ = (GOMP_MAP_FLAG_SPECIAL_2
+ | GOMP_MAP_FLAG_SPECIAL | 2),
/* Forced deallocation of zero length array section. */
GOMP_MAP_DELETE_ZERO_LEN_ARRAY_SECTION
= (GOMP_MAP_FLAG_SPECIAL_2
@@ -152,6 +157,12 @@ enum gomp_map_kind
GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY
| GOMP_MAP_FLAG_FORCE | 1),
+ /* Like GOMP_MAP_ATTACH, but allow attaching to zero-length array sections
+ (i.e. set to NULL when array section is not mapped) Currently only used
+ by OpenMP. */
+ GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION
+ = (GOMP_MAP_DEEP_COPY | 2),
+
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
@@ -175,7 +186,8 @@ enum gomp_map_kind
((X) == GOMP_MAP_ALWAYS_POINTER)
#define GOMP_MAP_POINTER_P(X) \
- ((X) == GOMP_MAP_POINTER)
+ ((X) == GOMP_MAP_POINTER \
+ || (X) == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION)
#define GOMP_MAP_ALWAYS_TO_P(X) \
(((X) == GOMP_MAP_ALWAYS_TO) || ((X) == GOMP_MAP_ALWAYS_TOFROM))
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 0cc3f4d406b..28c507fcb46 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1181,7 +1181,7 @@ extern uintptr_t gomp_map_val (struct target_mem_desc *,
void **, size_t);
extern void gomp_attach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree,
splay_tree_key, uintptr_t, size_t,
- struct gomp_coalesce_buf *);
+ struct gomp_coalesce_buf *, bool);
extern void gomp_detach_pointer (struct gomp_device_descr *,
struct goacc_asyncqueue *, splay_tree_key,
uintptr_t, bool, struct gomp_coalesce_buf *);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 4c8f0e0828e..3c5db96632d 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -939,7 +939,7 @@ acc_attach_async (void **hostaddr, int async)
}
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
- 0, NULL);
+ 0, NULL, false);
gomp_mutex_unlock (&acc_dev->lock);
}
@@ -1143,7 +1143,7 @@ goacc_enter_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
if ((kinds[i] & 0xff) == GOMP_MAP_ATTACH)
{
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n,
- (uintptr_t) h, s, NULL);
+ (uintptr_t) h, s, NULL, false);
/* OpenACC 'attach'/'detach' doesn't affect structured/dynamic
reference counts ('n->refcount', 'n->dynamic_refcount'). */
}
@@ -1161,7 +1161,8 @@ goacc_enter_data_internal (struct gomp_device_descr
*acc_dev, size_t mapnum,
splay_tree_key m
= lookup_host (acc_dev, hostaddrs[j], sizeof (void *));
gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, m,
- (uintptr_t) hostaddrs[j], sizes[j], NULL);
+ (uintptr_t) hostaddrs[j], sizes[j], NULL,
+ false);
}
bool processed = false;
diff --git a/libgomp/target.c b/libgomp/target.c
index 6152f58e13d..50f1d5f9b37 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -371,7 +371,8 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
unsigned char kind, bool always_to_flag,
struct gomp_coalesce_buf *cbuf)
{
- assert (kind != GOMP_MAP_ATTACH);
+ assert (kind != GOMP_MAP_ATTACH
+ || kind != GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
tgt_var->key = oldn;
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
@@ -412,7 +413,8 @@ get_kind (bool short_mapkind, void *kinds, int idx)
static void
gomp_map_pointer (struct target_mem_desc *tgt, struct goacc_asyncqueue *aq,
uintptr_t host_ptr, uintptr_t target_offset, uintptr_t bias,
- struct gomp_coalesce_buf *cbuf)
+ struct gomp_coalesce_buf *cbuf,
+ bool allow_zero_length_array_sections)
{
struct gomp_device_descr *devicep = tgt->device_descr;
struct splay_tree_s *mem_map = &devicep->mem_map;
@@ -434,16 +436,24 @@ gomp_map_pointer (struct target_mem_desc *tgt, struct
goacc_asyncqueue *aq,
splay_tree_key n = gomp_map_lookup (mem_map, &cur_node);
if (n == NULL)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("Pointer target of array section wasn't mapped");
- }
- cur_node.host_start -= n->host_start;
- cur_node.tgt_offset
- = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
- /* At this point tgt_offset is target address of the
- array section. Now subtract bias to get what we want
- to initialize the pointer with. */
- cur_node.tgt_offset -= bias;
+ if (allow_zero_length_array_sections)
+ cur_node.tgt_offset = 0;
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("Pointer target of array section wasn't mapped");
+ }
+ }
+ else
+ {
+ cur_node.host_start -= n->host_start;
+ cur_node.tgt_offset
+ = n->tgt->tgt_start + n->tgt_offset + cur_node.host_start;
+ /* At this point tgt_offset is target address of the
+ array section. Now subtract bias to get what we want
+ to initialize the pointer with. */
+ cur_node.tgt_offset -= bias;
+ }
gomp_copy_host2dev (devicep, aq, (void *) (tgt->tgt_start + target_offset),
(void *) &cur_node.tgt_offset, sizeof (void *), cbuf);
}
@@ -514,7 +524,8 @@ attribute_hidden void
gomp_attach_pointer (struct gomp_device_descr *devicep,
struct goacc_asyncqueue *aq, splay_tree mem_map,
splay_tree_key n, uintptr_t attach_to, size_t bias,
- struct gomp_coalesce_buf *cbufp)
+ struct gomp_coalesce_buf *cbufp,
+ bool allow_zero_length_array_sections)
{
struct splay_tree_key_s s;
size_t size, idx;
@@ -566,11 +577,21 @@ gomp_attach_pointer (struct gomp_device_descr *devicep,
if (!tn)
{
- gomp_mutex_unlock (&devicep->lock);
- gomp_fatal ("pointer target not mapped for attach");
+ if (allow_zero_length_array_sections)
+ {
+ /* When allowing attachment to zero-length array sections, we
+ allow attaching to NULL pointers when the target region is not
+ mapped. */
+ data = 0;
+ }
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("pointer target not mapped for attach");
+ }
}
-
- data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+ else
+ data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
gomp_debug (1,
"%s: attaching host %p, target %p (struct base %p) to %p\n",
@@ -824,7 +845,9 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
has_firstprivate = true;
continue;
}
- else if ((kind & typemask) == GOMP_MAP_ATTACH)
+ else if ((kind & typemask) == GOMP_MAP_ATTACH
+ || ((kind & typemask)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION))
{
tgt->list[i].key = NULL;
has_firstprivate = true;
@@ -1070,7 +1093,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
(uintptr_t) *(void **) hostaddrs[j],
k->tgt_offset + ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp);
+ sizes[j], cbufp, false);
}
}
i = j - 1;
@@ -1197,6 +1220,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
++i;
continue;
case GOMP_MAP_ATTACH:
+ case GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION:
{
cur_node.host_start = (uintptr_t) hostaddrs[i];
cur_node.host_end = cur_node.host_start + sizeof (void *);
@@ -1213,9 +1237,12 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
structured/dynamic reference counts ('n->refcount',
'n->dynamic_refcount'). */
+ bool zlas
+ = ((kind & typemask)
+ == GOMP_MAP_ATTACH_ZERO_LENGTH_ARRAY_SECTION);
gomp_attach_pointer (devicep, aq, mem_map, n,
(uintptr_t) hostaddrs[i], sizes[i],
- cbufp);
+ cbufp, zlas);
}
else if ((pragma_kind & GOMP_MAP_VARS_OPENACC) != 0)
{
@@ -1298,9 +1325,12 @@ gomp_map_vars_internal (struct gomp_device_descr
*devicep,
k->host_end - k->host_start, cbufp);
break;
case GOMP_MAP_POINTER:
- gomp_map_pointer (tgt, aq,
- (uintptr_t) *(void **) k->host_start,
- k->tgt_offset, sizes[i], cbufp);
+ case GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION:
+ gomp_map_pointer
+ (tgt, aq, (uintptr_t) *(void **) k->host_start,
+ k->tgt_offset, sizes[i], cbufp,
+ ((kind & typemask)
+ == GOMP_MAP_POINTER_TO_ZERO_LENGTH_ARRAY_SECTION));
break;
case GOMP_MAP_TO_PSET:
gomp_copy_host2dev (devicep, aq,
@@ -1335,7 +1365,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
k->tgt_offset
+ ((uintptr_t) hostaddrs[j]
- k->host_start),
- sizes[j], cbufp);
+ sizes[j], cbufp, false);
}
}
i = j - 1;
diff --git a/libgomp/testsuite/libgomp.c++/target-this-1.C
b/libgomp/testsuite/libgomp.c++/target-this-1.C
new file mode 100644
index 00000000000..a591ea4c564
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-1.C
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-2.C
b/libgomp/testsuite/libgomp.c++/target-this-2.C
new file mode 100644
index 00000000000..8119be8c2c5
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-2.C
@@ -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;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-3.C
b/libgomp/testsuite/libgomp.c++/target-this-3.C
new file mode 100644
index 00000000000..e15f69a1623
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-3.C
@@ -0,0 +1,99 @@
+#include <stdio.h>
+#include <string.h>
+extern "C" void abort ();
+
+struct S
+{
+ int * ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ bool set_ptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr != NULL)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ }
+
+ bool set_refptr (int n)
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr != NULL)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ S s = { ptr1, N, ptr2, N };
+
+ bool mapped;
+ int val = 123;
+
+ mapped = s.set_ptr (val);
+ if (mapped)
+ abort ();
+ if (s.ptr != ptr1)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ mapped = s.set_refptr (val);
+ if (mapped)
+ abort ();
+ if (s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N])
+ mapped = s.set_ptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_refptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != val)
+ abort ();
+
+ #pragma omp target data map(ptr2[:N])
+ mapped = s.set_refptr (val);
+
+ if (!mapped)
+ abort ();
+ if (s.set_ptr (0))
+ abort ();
+ if (s.ptr != ptr1 || s.refptr != ptr2)
+ abort ();
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != val)
+ abort ();
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.c++/target-this-4.C
b/libgomp/testsuite/libgomp.c++/target-this-4.C
new file mode 100644
index 00000000000..9f53677a240
--- /dev/null
+++ b/libgomp/testsuite/libgomp.c++/target-this-4.C
@@ -0,0 +1,104 @@
+
+// We use 'auto' without a function return type, so specify dialect here
+// { dg-additional-options "-std=c++14" }
+#include <cstdlib>
+#include <cstring>
+
+struct T
+{
+ int *ptr;
+ int ptr_len;
+
+ int *&refptr;
+ int refptr_len;
+
+ auto set_ptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (ptr)
+ for (int i = 0; i < ptr_len; i++)
+ ptr[i] = n;
+ mapped = (ptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+
+ auto set_refptr_func (int n)
+ {
+ auto fn = [=](void) -> bool
+ {
+ bool mapped;
+ #pragma omp target map(from:mapped)
+ {
+ if (refptr)
+ for (int i = 0; i < refptr_len; i++)
+ refptr[i] = n;
+ mapped = (refptr != NULL);
+ }
+ return mapped;
+ };
+ return fn;
+ }
+};
+
+int main (void)
+{
+ #define N 10
+ int *ptr1 = new int[N];
+ int *ptr2 = new int[N];
+
+ memset (ptr1, 0, sizeof (int) * N);
+ memset (ptr2, 0, sizeof (int) * N);
+
+ T a = { ptr1, N, ptr2, N };
+
+ auto p1 = a.set_ptr_func (1);
+ auto r2 = a.set_refptr_func (2);
+
+ if (p1 ())
+ abort ();
+ if (r2 ())
+ abort ();
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 0)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 0)
+ abort ();
+
+ #pragma omp target data map(ptr1[:N], ptr2[:N])
+ {
+ if (!p1 ())
+ abort ();
+ if (!r2 ())
+ abort ();
+ }
+
+ if (a.ptr != ptr1)
+ abort ();
+ if (a.refptr != ptr2)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr1[i] != 1)
+ abort ();
+
+ for (int i = 0; i < N; i++)
+ if (ptr2[i] != 2)
+ abort ();
+
+ return 0;
+}
