* [PATCH 3/6] [og8] OpenACC 2.6 manual deep copy support (attach/detach)
2018-11-20 21:55 [PATCH 0/6] [og8] OpenACC attach/detach Julian Brown
2018-11-20 21:55 ` [PATCH 2/6] [og8] Factor out duplicate code in gimplify_scan_omp_clauses Julian Brown
@ 2018-11-20 21:55 ` Julian Brown
2018-11-22 15:48 ` Bernhard Reutner-Fischer
2018-11-20 21:55 ` [PATCH 4/6] [og8] Interaction of dynamic/multidimensional arrays with attach/detach Julian Brown
` (3 subsequent siblings)
5 siblings, 1 reply; 8+ messages in thread
From: Julian Brown @ 2018-11-20 21:55 UTC (permalink / raw)
To: gcc-patches; +Cc: cltang, Catherine_Moore, jakub
[-- Attachment #1: Type: text/plain, Size: 9082 bytes --]
Previously posted upstream:
https://gcc.gnu.org/ml/gcc-patches/2018-11/msg00826.html
gcc/c/
* c-parser.c (c_parser_omp_variable_list): Allow deref (->) in
variable lists.
(c_parser_oacc_all_clauses): Re-alphabetize cases.
* c-typeck.c (handle_omp_array_sections_1): Support deref.
gcc/cp/
* parser.c (cp_parser_omp_var_list_no_open): Support deref.
(cp_parser_oacc_all_clauses): Re-alphabetize cases.
* semantics.c (finish_omp_clauses): Allow "this" for OpenACC data
clauses. Support deref.
gcc/fortran/
* gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH.
* openmp.c (omp_mask2): Add OMP_CLAUSE_ATTACH, OMP_CLAUSE_DETACH.
(gfc_match_omp_clauses): Remove allow_derived parameter, infer from
clause mask. Support attach and detach. Slight reformatting.
(OACC_PARALLEL_CLAUSES, OACC_KERNELS_CLAUSES, OACC_DATA_CLAUSES)
(OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH.
(match_acc): Remove derived_types parameter, and don't pass to
gfc_match_omp_clauses.
(gfc_match_oacc_update): Don't pass allow_derived argument.
(gfc_match_oacc_enter_data): Likewise.
(gfc_match_oacc_exit_data): Likewise.
(check_symbol_not_pointer): Don't disallow pointer objects of derived
type.
(resolve_oacc_data_clauses): Don't disallow allocatable derived types.
(resolve_omp_clauses): Perform duplicate checking only for non-derived
type component accesses (plain variables and arrays or array sections).
Support component refs.
* trans-openmp.c (gfc_omp_privatize_by_reference): Support component
refs.
(gfc_trans_omp_clauses_1): Support component refs, attach and detach
clauses.
gcc/
* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
(insert_struct_component_mapping): Support derived-type member mappings
for arrays with descriptors which use GOMP_MAP_TO_PSET.
(gimplify_scan_omp_clauses): Rewrite GOMP_MAP_ALWAYS_POINTER to
GOMP_MAP_ATTACH for OpenACC struct/derived-type component pointers.
Handle pointer mappings that use GOMP_MAP_TO_PSET. Handle attach/detach
clauses.
(gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
attach/detach clauses.
(gimplify_omp_target_update): Handle finalize for detach.
gcc/testsuite/
* c-c++-common/goacc/mdc-1.c: Update scan tests.
* gfortran.dg/goacc/data-clauses.f95: Remove expected errors.
* gfortran.dg/goacc/derived-types.f90: Likewise.
* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
libgomp/
* libgomp.h (struct target_var_desc): Add do_detach flag.
(struct splay_tree_key_s): Add attach_count field.
(struct gomp_coalesce_buf): Add forward declaration.
(gomp_map_val, gomp_attach_pointer, gomp_detach_pointer): Add
prototypes.
(gomp_unmap_vars): Add finalize parameter.
* libgomp.map (OACC_2.6): New section. Add acc_attach, acc_attach_async,
acc_detach, acc_detach_async, acc_detach_finalize,
acc_detach_finalize_async.
* oacc-async.c (goacc_async_copyout_unmap_vars): Add finalize parameter.
Pass to gomp_unmap_vars_async.
* oacc-init.c (acc_shutdown_1): Update call to gomp_unmap_vars.
* oacc-int.h (goacc_async_copyout_unmap_vars): Add finalize parameter.
* oacc-mem.c (acc_unmap_data): Update call to gomp_unmap_vars.
(present_create_copy): Initialise attach_count.
(delete_copyout): Likewise.
(gomp_acc_insert_pointer): Likewise.
(gomp_acc_remove_pointer): Update calls to gomp_unmap_vars,
goacc_async_copyout_unmap_vars.
(acc_attach_async, acc_attach, goacc_detach_internal, acc_detach)
(acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): New
functions.
* oacc-parallel.c (find_pointer): Support attach/detach. Make a little
more strict.
(GOACC_parallel_keyed_internal): Use gomp_map_val to calculate device
addresses. Update calls to gomp_unmap_vars,
goacc_async_copyout_unmap_vars.
(GOACC_data_end): Update call to gomp_unmap_vars.
(GOACC_enter_exit_data): Support attach/detach and GOMP_MAP_STRUCT.
* openacc.h (acc_attach, acc_attach_async, acc_detach)
(acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
prototypes.
* target.c (limits.h): Include.
(gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc.
(gomp_attach_pointer, gomp_detach_pointer): New functions.
(gomp_map_val): Make global.
(gomp_map_vars_async): Support attach and detach.
(gomp_remove_var): Free attach count array if present.
(gomp_unmap_vars): Add finalize parameter. Update call to
gomp_unmap_vars_async.
(gomp_unmap_vars_async): Add finalize parameter. Add pointer detaching
support.
(GOMP_target): Update call to gomp_unmap_vars.
(GOMP_target_ext): Likewise.
(gomp_exit_data): Free attach count array if present.
(gomp_target_task_fn): Update call to gomp_unmap_vars.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-1.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-2.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-3.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-4.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-5.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-6.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-7.c: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-8.c: New test.
* testsuite/libgomp.oacc-fortran/derived-type-1.f90: Update test to use
stop <n>.
* testsuite/libgomp.oacc-fortran/update-2.f90: Likewise.
---
gcc/c/c-parser.c | 15 +-
gcc/c/c-typeck.c | 4 +
gcc/cp/parser.c | 16 +-
gcc/cp/semantics.c | 6 +-
gcc/fortran/gfortran.h | 2 +
gcc/fortran/openmp.c | 126 +++++++++-----
gcc/fortran/trans-openmp.c | 143 ++++++---------
gcc/gimplify.c | 82 +++++++--
gcc/testsuite/c-c++-common/goacc/mdc-1.c | 10 +-
gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 | 38 ++--
gcc/testsuite/gfortran.dg/goacc/derived-types.f90 | 23 +--
.../gfortran.dg/goacc/enter-exit-data.f95 | 24 ++--
libgomp/libgomp.h | 23 ++-
libgomp/libgomp.map | 10 +
libgomp/oacc-async.c | 4 +-
libgomp/oacc-init.c | 2 +-
libgomp/oacc-int.h | 2 +-
libgomp/oacc-mem.c | 86 +++++++++-
libgomp/oacc-parallel.c | 190 +++++++++++++++-----
libgomp/openacc.h | 6 +
libgomp/target.c | 189 ++++++++++++++++++-
.../libgomp.oacc-c-c++-common/deep-copy-1.c | 24 +++
.../libgomp.oacc-c-c++-common/deep-copy-2.c | 29 +++
.../libgomp.oacc-c-c++-common/deep-copy-3.c | 34 ++++
.../libgomp.oacc-c-c++-common/deep-copy-4.c | 87 +++++++++
.../libgomp.oacc-c-c++-common/deep-copy-5.c | 81 +++++++++
.../testsuite/libgomp.oacc-fortran/deep-copy-1.f90 | 35 ++++
.../testsuite/libgomp.oacc-fortran/deep-copy-2.f90 | 33 ++++
.../testsuite/libgomp.oacc-fortran/deep-copy-3.f90 | 34 ++++
.../testsuite/libgomp.oacc-fortran/deep-copy-4.f90 | 49 +++++
.../testsuite/libgomp.oacc-fortran/deep-copy-5.f90 | 57 ++++++
.../testsuite/libgomp.oacc-fortran/deep-copy-6.f90 | 61 +++++++
.../testsuite/libgomp.oacc-fortran/deep-copy-7.f90 | 89 +++++++++
.../testsuite/libgomp.oacc-fortran/deep-copy-8.f90 | 41 +++++
.../libgomp.oacc-fortran/derived-type-1.f90 | 6 +-
.../testsuite/libgomp.oacc-fortran/update-2.f90 | 44 +++---
36 files changed, 1407 insertions(+), 298 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0003-og8-OpenACC-2.6-manual-deep-copy-support-attach-deta.patch --]
[-- Type: text/x-patch; name="0003-og8-OpenACC-2.6-manual-deep-copy-support-attach-deta.patch", Size: 91395 bytes --]
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index ffc5fe9..4b6ab84 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -11553,9 +11553,12 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
- while (c_parser_next_token_is (parser, CPP_DOT))
+ while (c_parser_next_token_is (parser, CPP_DOT)
+ || c_parser_next_token_is (parser, CPP_DEREF))
{
location_t op_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_next_token_is (parser, CPP_DEREF))
+ t = build_simple_mem_ref (t);
c_parser_consume_token (parser);
if (!c_parser_next_token_is (parser, CPP_NAME))
{
@@ -11679,7 +11682,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
}
/* OpenACC 2.5:
- attach (variable-list )
+ attach ( variable-list )
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
@@ -14090,15 +14093,15 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_clause_async (parser, clauses);
c_name = "async";
break;
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "attach";
+ break;
case PRAGMA_OACC_CLAUSE_AUTO:
clauses = c_parser_oacc_simple_clause (parser, here, OMP_CLAUSE_AUTO,
clauses);
c_name = "auto";
break;
- case PRAGMA_OACC_CLAUSE_ATTACH:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "attach";
- break;
case PRAGMA_OACC_CLAUSE_BIND:
clauses = c_parser_oacc_clause_bind (parser, clauses);
c_name = "bind";
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index ab6819c..1a18867 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12446,6 +12446,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
t = TREE_OPERAND (t, 0);
}
+ if (TREE_CODE (t) == MEM_REF)
+ t = TREE_OPERAND (t, 0);
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
@@ -13750,6 +13752,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (remove)
break;
+ if (TREE_CODE (t) == MEM_REF)
+ t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index 8161d63..79c03d2 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -31563,15 +31563,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
- while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+ while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+ || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))
{
+ cpp_ttype ttype
+ = cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+ ? CPP_DOT : CPP_DEREF;
location_t loc
= cp_lexer_peek_token (parser->lexer)->location;
cp_id_kind idk = CP_ID_KIND_NONE;
cp_lexer_consume_token (parser->lexer);
decl = convert_from_reference (decl);
decl
- = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+ = cp_parser_postfix_dot_deref_expression (parser, ttype,
decl, false,
&idk, loc);
}
@@ -33858,15 +33862,15 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_clause_async (parser, clauses);
c_name = "async";
break;
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "attach";
+ break;
case PRAGMA_OACC_CLAUSE_AUTO:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO,
clauses, here);
c_name = "auto";
break;
- case PRAGMA_OACC_CLAUSE_ATTACH:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "attach";
- break;
case PRAGMA_OACC_CLAUSE_BIND:
clauses = cp_parser_oacc_clause_bind (parser, clauses);
c_name = "bind";
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 796ae7f..7cbcb34 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -6724,7 +6724,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
error ("%qE is not a variable in %<depend%> clause", t);
remove = true;
}
- else if (ort != C_ORT_ACC && t == current_class_ptr)
+ else if (t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6810,6 +6810,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
+ if (ort == C_ORT_ACC
+ && TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
if (TREE_CODE (t) == COMPONENT_REF
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index 3a9e45b..14b5def 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1183,10 +1183,12 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_ATTACH,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
OMP_MAP_DELETE,
+ OMP_MAP_DETACH,
OMP_MAP_FORCE_ALLOC,
OMP_MAP_FORCE_TO,
OMP_MAP_FORCE_FROM,
diff --git a/gcc/fortran/openmp.c b/gcc/fortran/openmp.c
index 654ceb6..f120e3d 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -808,7 +808,7 @@ enum omp_mask1
OMP_MASK1_LAST
};
-/* OpenACC 2.0 specific clauses. */
+/* OpenACC 2.0+ specific clauses. */
enum omp_mask2
{
OMP_CLAUSE_ASYNC,
@@ -837,6 +837,8 @@ enum omp_mask2
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_DEVICE_TYPE,
+ OMP_CLAUSE_ATTACH,
+ OMP_CLAUSE_DETACH,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -964,10 +966,18 @@ static match
gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
const omp_mask dtype_mask,
bool first = true, bool needs_space = true,
- bool openacc = false, bool allow_derived = false)
+ bool openacc = false)
{
gfc_omp_clauses *base_clauses, *c = gfc_get_omp_clauses ();
locus old_loc;
+ /* Determine whether we're dealing with an OpenACC directive that permits
+ derived type member accesses. This in particular disallows
+ "!$acc declare" from using such accesses, because it's not clear if/how
+ that should work. */
+ bool allow_derived = (openacc
+ && ((mask & OMP_CLAUSE_ATTACH)
+ || (mask & OMP_CLAUSE_DETACH)
+ || (mask & OMP_CLAUSE_HOST_SELF)));
base_clauses = c;
@@ -1043,6 +1053,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
needs_space = true;
continue;
}
+ if ((mask & OMP_CLAUSE_ATTACH)
+ && gfc_match ("attach ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_ATTACH, false,
+ allow_derived))
+ continue;
break;
case 'b':
if ((mask & OMP_CLAUSE_BIND) && c->routine_bind == NULL
@@ -1098,8 +1114,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, true,
- allow_derived))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
@@ -1109,8 +1124,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, true,
- allow_derived))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
break;
case 'd':
@@ -1190,6 +1204,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
else
gfc_current_locus = old_loc;
}
+ if ((mask & OMP_CLAUSE_DETACH)
+ && gfc_match ("detach ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_DETACH, false,
+ allow_derived))
+ continue;
if ((mask & OMP_CLAUSE_DEVICE)
&& !openacc
&& c->device == NULL
@@ -1784,8 +1804,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if (gfc_match_omp_variable_list (" :",
&c->lists[OMP_LIST_REDUCTION],
- false, NULL, &head,
- openacc) == MATCH_YES)
+ false, NULL, &head, openacc,
+ allow_derived) == MATCH_YES)
{
gfc_omp_namelist *n;
if (rop == OMP_REDUCTION_NONE)
@@ -2053,7 +2073,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
- | OMP_CLAUSE_DEFAULT)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \
@@ -2063,12 +2083,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_DEFAULT)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
- | OMP_CLAUSE_DEVICEPTR)
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_ATTACH)
#define OACC_HOST_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_USE_DEVICE))
#define OACC_LOOP_CLAUSES \
@@ -2098,12 +2118,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
#define OACC_ENTER_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
+ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH)
#define OACC_EXIT_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE \
- | OMP_CLAUSE_FINALIZE)
+ | OMP_CLAUSE_FINALIZE | OMP_CLAUSE_DETACH)
#define OACC_ROUTINE_CLAUSES \
(omp_mask (OMP_CLAUSE_GANG) | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \
| OMP_CLAUSE_SEQ \
@@ -2139,12 +2159,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
static match
-match_acc (gfc_exec_op op, const omp_mask mask, const omp_mask dtype_mask,
- bool derived_types=false)
+match_acc (gfc_exec_op op, const omp_mask mask, const omp_mask dtype_mask)
{
gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, mask, dtype_mask, false, false, true,
- derived_types)
+ if (gfc_match_omp_clauses (&c, mask, dtype_mask, false, false, true)
!= MATCH_YES)
return MATCH_ERROR;
new_st.op = op;
@@ -2309,7 +2327,8 @@ gfc_match_oacc_update (void)
if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES,
OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK, false,
- false, true, true) != MATCH_YES)
+ false, true)
+ != MATCH_YES)
return MATCH_ERROR;
if (!c->lists[OMP_LIST_MAP])
@@ -2329,7 +2348,7 @@ match
gfc_match_oacc_enter_data (void)
{
return match_acc (EXEC_OACC_ENTER_DATA, OACC_ENTER_DATA_CLAUSES,
- OMP_MASK2_LAST, true);
+ OMP_MASK2_LAST);
}
@@ -2337,7 +2356,7 @@ match
gfc_match_oacc_exit_data (void)
{
return match_acc (EXEC_OACC_EXIT_DATA, OACC_EXIT_DATA_CLAUSES,
- OMP_MASK2_LAST, true);
+ OMP_MASK2_LAST);
}
@@ -4017,9 +4036,6 @@ resolve_nonnegative_int_expr (gfc_expr *expr, const char *clause)
static void
check_symbol_not_pointer (gfc_symbol *sym, locus loc, const char *name)
{
- if (sym->ts.type == BT_DERIVED && sym->attr.pointer)
- gfc_error ("POINTER object %qs of derived type in %s clause at %L",
- sym->name, name, &loc);
if (sym->ts.type == BT_DERIVED && sym->attr.cray_pointer)
gfc_error ("Cray pointer object %qs of derived type in %s clause at %L",
sym->name, name, &loc);
@@ -4060,9 +4076,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
static void
resolve_oacc_data_clauses (gfc_symbol *sym, locus loc, const char *name)
{
- if (sym->ts.type == BT_DERIVED && sym->attr.allocatable)
- gfc_error ("ALLOCATABLE object %qs of derived type in %s clause at %L",
- sym->name, name, &loc);
if ((sym->ts.type == BT_ASSUMED && sym->attr.allocatable)
|| (sym->ts.type == BT_CLASS && CLASS_DATA (sym)
&& CLASS_DATA (sym)->attr.allocatable))
@@ -4408,11 +4421,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
&& (list != OMP_LIST_REDUCTION || !openacc))
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, &n->where);
- else
- n->sym->mark = 1;
+ bool array_only_p = true;
+ /* Disallow duplicate bare variable references and multiple
+ subarrays of the same array here, but allow multiple components of
+ the same (e.g. derived-type) variable. For the latter, duplicate
+ components are detected elsewhere. */
+ if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE)
+ for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
+ if (ref->type != REF_ARRAY)
+ array_only_p = false;
+ if (array_only_p)
+ {
+ if (n->sym->mark)
+ gfc_error ("Symbol %qs present on multiple clauses at %L",
+ n->sym->name, &n->where);
+ else
+ n->sym->mark = 1;
+ }
}
gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
@@ -4603,26 +4628,41 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"are allowed on ORDERED directive at %L",
&n->where);
}
+ gfc_ref *array_ref = NULL;
+ bool resolved = false;
if (n->expr)
{
- if (!gfc_resolve_expr (n->expr)
+ array_ref = n->expr->ref;
+ resolved = gfc_resolve_expr (n->expr);
+
+ /* Look through component refs to find last array
+ reference. */
+ while (resolved
+ && array_ref
+ && (array_ref->type == REF_COMPONENT
+ || (array_ref->type == REF_ARRAY
+ && array_ref->next
+ && array_ref->next->type == REF_COMPONENT)))
+ array_ref = array_ref->next;
+ }
+ if (array_ref
+ || (n->expr
+ && (!resolved || n->expr->expr_type != EXPR_VARIABLE)))
+ {
+ if (!resolved
|| n->expr->expr_type != EXPR_VARIABLE
- || n->expr->ref == NULL
- || n->expr->ref->next
- || n->expr->ref->type != REF_ARRAY)
- {
- if (n->sym->ts.type != BT_DERIVED)
- gfc_error ("%qs in %s clause at %L is not a proper "
- "array section", n->sym->name, name,
- &n->where);
- }
- else if (n->expr->ref->u.ar.codimen)
+ || array_ref->next
+ || array_ref->type != REF_ARRAY)
+ gfc_error ("%qs in %s clause at %L is not a proper "
+ "array section", n->sym->name, name,
+ &n->where);
+ else if (array_ref->u.ar.codimen)
gfc_error ("Coarrays not supported in %s clause at %L",
name, &n->where);
else
{
int i;
- gfc_array_ref *ar = &n->expr->ref->u.ar;
+ gfc_array_ref *ar = &array_ref->u.ar;
for (i = 0; i < ar->dimen; i++)
if (ar->stride[i])
{
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index 8840fd2..98f40d1 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -62,6 +62,9 @@ gfc_omp_privatize_by_reference (const_tree decl)
if (TREE_CODE (type) == POINTER_TYPE)
{
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 1);
+
/* Array POINTER/ALLOCATABLE have aggregate types, all user variables
that have POINTER_TYPE type and aren't scalar pointers, scalar
allocatables, Cray pointees or C pointers are supposed to be
@@ -2121,69 +2124,35 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
tree decl = gfc_get_symbol_decl (n->sym);
if (DECL_P (decl))
TREE_ADDRESSABLE (decl) = 1;
- /* Handle derived-typed members for OpenACC Update. */
- if (n->sym->ts.type == BT_DERIVED
- && n->expr != NULL && n->expr->ref != NULL
- && (n->expr->ref->next == NULL
- || (n->expr->ref->next != NULL
- && n->expr->ref->next->type == REF_ARRAY
- && n->expr->ref->next->u.ar.type == AR_FULL))
- && (n->expr->ref->type == REF_ARRAY
- && n->expr->ref->u.ar.type != AR_SECTION))
- {
- gfc_ref *ref = n->expr->ref;
- gfc_component *c = ref->u.c.component;
- tree field;
- tree context;
- tree ptr;
- tree type;
- tree scratch;
- if (c->backend_decl == NULL_TREE
- && ref->u.c.sym != NULL)
- gfc_get_derived_type (ref->u.c.sym);
+ gfc_ref *ref = n->expr ? n->expr->ref : NULL;
+ symbol_attribute *sym_attr = &n->sym->attr;
+ gomp_map_kind ptr_map_kind = GOMP_MAP_POINTER;
- field = c->backend_decl;
- gcc_assert (field && TREE_CODE (field) == FIELD_DECL);
- context = DECL_FIELD_CONTEXT (field);
-
- type = TREE_TYPE (decl);
- if (POINTER_TYPE_P (type))
- type = TREE_TYPE (type);
+ if (ref && n->sym->ts.type == BT_DERIVED)
+ {
+ if (gfc_omp_privatize_by_reference (decl))
+ decl = build_fold_indirect_ref (decl);
- if (context != type)
+ for (; ref && ref->type == REF_COMPONENT; ref = ref->next)
{
- tree f2 = c->norestrict_decl;
- if (!f2 || DECL_FIELD_CONTEXT (f2) != type)
- for (f2 = TYPE_FIELDS (TREE_TYPE (decl)); f2;
- f2 = DECL_CHAIN (f2))
- if (TREE_CODE (f2) == FIELD_DECL
- && DECL_NAME (f2) == DECL_NAME (field))
- break;
- gcc_assert (f2);
- c->norestrict_decl = f2;
- field = f2;
+ tree field = ref->u.c.component->backend_decl;
+ gcc_assert (field && TREE_CODE (field) == FIELD_DECL);
+ decl = fold_build3 (COMPONENT_REF, TREE_TYPE (field),
+ decl, field, NULL_TREE);
+ sym_attr = &ref->u.c.component->attr;
}
- if (POINTER_TYPE_P (TREE_TYPE (decl)))
- decl = build_fold_indirect_ref_loc (input_location,
- decl);
-
- scratch = fold_build3_loc (input_location, COMPONENT_REF,
- TREE_TYPE (field), decl, field,
- NULL_TREE);
- type = TREE_TYPE (scratch);
- ptr = gfc_create_var (pvoid_type_node, NULL);
- scratch = fold_convert (pvoid_type_node,
- build_fold_addr_expr (scratch));
- gfc_add_modify (block, ptr, scratch);
- OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (type);
- OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+ ptr_map_kind = GOMP_MAP_ALWAYS_POINTER;
}
- else if ((n->sym->ts.type == BT_DERIVED && n->expr == NULL)
- || (n->expr == NULL
- || n->expr->ref->u.ar.type == AR_FULL))
+
+ if (ref == NULL || ref->u.ar.type == AR_FULL)
{
+ tree field = decl;
+
+ while (TREE_CODE (field) == COMPONENT_REF)
+ field = TREE_OPERAND (field, 1);
+
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
{
@@ -2192,18 +2161,18 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
}
else if (POINTER_TYPE_P (TREE_TYPE (decl))
&& (gfc_omp_privatize_by_reference (decl)
- || GFC_DECL_GET_SCALAR_POINTER (decl)
- || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
- || GFC_DECL_CRAY_POINTEE (decl)
+ || GFC_DECL_GET_SCALAR_POINTER (field)
+ || GFC_DECL_GET_SCALAR_ALLOCATABLE (field)
+ || GFC_DECL_CRAY_POINTEE (field)
|| GFC_DESCRIPTOR_TYPE_P
- (TREE_TYPE (TREE_TYPE (decl)))))
+ (TREE_TYPE (TREE_TYPE (field)))))
{
tree orig_decl = decl;
enum gomp_map_kind gmk = GOMP_MAP_FIRSTPRIVATE_POINTER;
if (GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
&& (n->sym->attr.oacc_declare_create)
&& clauses->update_allocatable)
- gmk = GOMP_MAP_ALWAYS_POINTER;
+ gmk = ptr_map_kind;
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, gmk);
@@ -2216,7 +2185,7 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
{
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3) = decl;
OMP_CLAUSE_SIZE (node3) = size_int (0);
decl = build_fold_indirect_ref (decl);
@@ -2225,7 +2194,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
}
- if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
+ && n->u.map_op != OMP_MAP_ATTACH
+ && n->u.map_op != OMP_MAP_DETACH)
{
tree type = TREE_TYPE (decl);
tree ptr = gfc_conv_descriptor_data_get (decl);
@@ -2238,14 +2209,16 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
+ if (ptr_map_kind == GOMP_MAP_ALWAYS_POINTER)
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
OMP_CLAUSE_SIZE (node3) = size_int (0);
/* We have to check for n->sym->attr.dimension because
of scalar coarrays. */
- if (n->sym->attr.pointer && n->sym->attr.dimension)
+ if (sym_attr->pointer && sym_attr->dimension)
{
stmtblock_t cond_block;
tree size
@@ -2275,11 +2248,11 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
else_b));
OMP_CLAUSE_SIZE (node) = size;
}
- else if (n->sym->attr.dimension)
+ else if (sym_attr->dimension)
OMP_CLAUSE_SIZE (node)
= gfc_full_array_size (block, decl,
GFC_TYPE_ARRAY_RANK (type));
- if (n->sym->attr.dimension)
+ if (sym_attr->dimension)
{
tree elemsz
= TYPE_SIZE_UNIT (gfc_get_element_type (type));
@@ -2292,31 +2265,17 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
else
OMP_CLAUSE_DECL (node) = decl;
}
- else
+ else if (ref)
{
tree ptr, ptr2;
gfc_init_se (&se, NULL);
- if ((n->sym->ts.type == BT_DERIVED
- && n->expr->rank == 0)
- || (n->sym->ts.type != BT_DERIVED
- && n->expr->ref->u.ar.type == AR_ELEMENT))
+ if (ref->u.ar.type == AR_ELEMENT)
{
gfc_conv_expr_reference (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
ptr = se.expr;
- tree type = TREE_TYPE (ptr);
- if (n->sym->ts.type == BT_DERIVED)
- {
- tree t = gfc_create_var (build_pointer_type
- (void_type_node),
- NULL);
- ptr = fold_convert (pvoid_type_node, ptr);
- gfc_add_modify (block, t, ptr);
- ptr = t;
- type = TREE_TYPE (type);
- }
OMP_CLAUSE_SIZE (node)
- = TYPE_SIZE_UNIT (type);
+ = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
}
else
{
@@ -2337,14 +2296,12 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
gfc_add_block_to_block (block, &se.post);
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
- if (n->sym->ts.type == BT_DERIVED)
- goto finalize_map_clause;
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
{
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node4, ptr_map_kind);
OMP_CLAUSE_DECL (node4) = decl;
OMP_CLAUSE_SIZE (node4) = size_int (0);
decl = build_fold_indirect_ref (decl);
@@ -2361,9 +2318,11 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
+ if (ptr_map_kind == GOMP_MAP_ALWAYS_POINTER)
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
}
else
{
@@ -2376,7 +2335,7 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3) = decl;
}
ptr2 = fold_convert (sizetype, ptr2);
@@ -2384,11 +2343,16 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
finalize_map_clause:;
}
+ else
+ gcc_unreachable ();
switch (n->u.map_op)
{
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_ATTACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
@@ -2413,6 +2377,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_DELETE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
break;
+ case OMP_MAP_DETACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
+ break;
case OMP_MAP_FORCE_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
break;
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 824e020..40bf586 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -111,6 +111,10 @@ enum gimplify_omp_var_data
/* Flag for OpenACC deviceptrs. */
GOVD_DEVICEPTR = (1<<21),
+ /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
+ fields. */
+ GOVD_MAP_HAS_ATTACHMENTS = (1<<22),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -7692,7 +7696,13 @@ insert_struct_component_mapping (enum tree_code code, tree c, tree struct_node,
OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));
OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;
- OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
+ if (OMP_CLAUSE_CHAIN (prev_node) != c
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+ == GOMP_MAP_TO_PSET))
+ OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));
+ else
+ OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
if (struct_node)
OMP_CLAUSE_CHAIN (struct_node) = c2;
@@ -8245,7 +8255,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
- if (DECL_P (decl))
+ if (DECL_P (decl)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+ && code != OACC_UPDATE)
{
if (error_operand_p (decl))
{
@@ -8297,17 +8309,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
= splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
bool ptr = (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_ALWAYS_POINTER);
- if ((n == NULL || (n->value & GOVD_MAP) == 0)
- && code != OACC_UPDATE)
+ bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH;
+ bool has_attachments = false;
+ /* For OpenACC, pointers in structs should trigger an
+ attach action. */
+ if (ptr && (region_type & ORT_ACC) != 0)
+ {
+ /* Turning a GOMP_MAP_ALWAYS_POINTER clause into a
+ GOMP_MAP_ATTACH clause after we have detected a case
+ that needs a GOMP_MAP_STRUCT mapping adding. */
+ OMP_CLAUSE_SET_MAP_KIND (c,
+ (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ has_attachments = true;
+ }
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
{
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+ OMP_CLAUSE_SET_MAP_KIND (l, attach
+ ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT);
if (!base_eq_orig_base)
OMP_CLAUSE_DECL (l) = unshare_expr (orig_base);
else
OMP_CLAUSE_DECL (l) = decl;
- OMP_CLAUSE_SIZE (l) = size_int (1);
+ OMP_CLAUSE_SIZE (l) = attach
+ ? (DECL_P (OMP_CLAUSE_DECL (l))
+ ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+ : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))))
+ : size_int (1);
if (struct_map_to_clause == NULL)
struct_map_to_clause = new hash_map<tree, tree>;
struct_map_to_clause->put (decl, l);
@@ -8339,9 +8370,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags = GOVD_MAP | GOVD_EXPLICIT;
if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
flags |= GOVD_SEEN;
+ if (has_attachments)
+ flags |= GOVD_MAP_HAS_ATTACHMENTS;
goto do_add_decl;
}
- else
+ else if (struct_map_to_clause)
{
tree *osc = struct_map_to_clause->get (decl);
tree *sc = NULL, *scp = NULL;
@@ -8350,8 +8383,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
sc = &OMP_CLAUSE_CHAIN (*osc);
if (*sc != c
&& (OMP_CLAUSE_MAP_KIND (*sc)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
sc = &OMP_CLAUSE_CHAIN (*sc);
+ /* Here "prev_list_p" is the end of the inserted
+ alloc/release nodes after the struct node, OSC. */
for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
if (ptr && sc == prev_list_p)
break;
@@ -8410,9 +8445,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (remove)
break;
- OMP_CLAUSE_SIZE (*osc)
- = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
- size_one_node);
+ if (!attach)
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
if (ptr)
{
tree cl
@@ -8444,11 +8480,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_CHAIN (c)
&& OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ALWAYS_POINTER))
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ALWAYS_POINTER)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_TO_PSET)))
prev_list_p = list_p;
+
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -9020,6 +9060,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
return 0;
if ((flags & GOVD_SEEN) == 0)
return 0;
+ if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
+ return 0;
if (flags & GOVD_DEBUG_PRIVATE)
{
gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
@@ -9509,8 +9551,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && (code == OMP_TARGET_EXIT_DATA
- || code == OACC_EXIT_DATA))
+ && code == OMP_TARGET_EXIT_DATA)
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -11218,10 +11259,15 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
finalize_marked = true;
break;
+ case GOMP_MAP_STRUCT:
+ case GOMP_MAP_FORCE_PRESENT:
+ /* Skip over an initial struct or force_present mapping. */
+ break;
default:
- /* Check consistency: libgomp relies on the very first data
- mapping clause being marked, so make sure we did that before
- any other mapping clauses. */
+ /* Check consistency: libgomp relies on the very first
+ non-struct, non-force-present data mapping clause being
+ marked, so make sure we did that before any other mapping
+ clauses. */
gcc_assert (finalize_marked);
break;
}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
index c20b94d..84a44af 100644
--- a/gcc/testsuite/c-c++-common/goacc/mdc-1.c
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -42,13 +42,13 @@ t1 ()
}
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.always_pointer:s.a .pointer assign, bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .len: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.force_present:s .len: 32.. map.detach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_present:s .len: 32.. map.force_detach:s.a .len: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
index b94214e..1a4a671 100644
--- a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95
@@ -39,9 +39,9 @@ contains
!$acc end data
- !$acc parallel copy (tip) ! { dg-error "POINTER" }
+ !$acc parallel copy (tip)
!$acc end parallel
- !$acc parallel copy (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel copy (tia)
!$acc end parallel
!$acc parallel deviceptr (i) copy (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -54,9 +54,9 @@ contains
!$acc end data
- !$acc parallel copyin (tip) ! { dg-error "POINTER" }
+ !$acc parallel copyin (tip)
!$acc end parallel
- !$acc parallel copyin (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel copyin (tia)
!$acc end parallel
!$acc parallel deviceptr (i) copyin (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -71,9 +71,9 @@ contains
!$acc end data
- !$acc parallel copyout (tip) ! { dg-error "POINTER" }
+ !$acc parallel copyout (tip)
!$acc end parallel
- !$acc parallel copyout (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel copyout (tia)
!$acc end parallel
!$acc parallel deviceptr (i) copyout (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -90,9 +90,9 @@ contains
!$acc end data
- !$acc parallel create (tip) ! { dg-error "POINTER" }
+ !$acc parallel create (tip)
!$acc end parallel
- !$acc parallel create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel create (tia)
!$acc end parallel
!$acc parallel deviceptr (i) create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -111,9 +111,9 @@ contains
!$acc end data
- !$acc parallel present (tip) ! { dg-error "POINTER" }
+ !$acc parallel present (tip)
!$acc end parallel
- !$acc parallel present (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel present (tia)
!$acc end parallel
!$acc parallel deviceptr (i) present (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -144,9 +144,9 @@ contains
!$acc end parallel
- !$acc parallel present_or_copy (tip) ! { dg-error "POINTER" }
+ !$acc parallel present_or_copy (tip)
!$acc end parallel
- !$acc parallel present_or_copy (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel present_or_copy (tia)
!$acc end parallel
!$acc parallel deviceptr (i) present_or_copy (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -169,9 +169,9 @@ contains
!$acc end data
- !$acc parallel present_or_copyin (tip) ! { dg-error "POINTER" }
+ !$acc parallel present_or_copyin (tip)
!$acc end parallel
- !$acc parallel present_or_copyin (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel present_or_copyin (tia)
!$acc end parallel
!$acc parallel deviceptr (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -196,9 +196,9 @@ contains
!$acc end data
- !$acc parallel present_or_copyout (tip) ! { dg-error "POINTER" }
+ !$acc parallel present_or_copyout (tip)
!$acc end parallel
- !$acc parallel present_or_copyout (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel present_or_copyout (tia)
!$acc end parallel
!$acc parallel deviceptr (i) present_or_copyout (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -225,9 +225,9 @@ contains
!$acc end data
- !$acc parallel present_or_create (tip) ! { dg-error "POINTER" }
+ !$acc parallel present_or_create (tip)
!$acc end parallel
- !$acc parallel present_or_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc parallel present_or_create (tia)
!$acc end parallel
!$acc parallel deviceptr (i) present_or_create (i) ! { dg-error "multiple clauses" }
!$acc end parallel
@@ -256,4 +256,4 @@ contains
!$acc end data
end subroutine foo
-end module test
\ No newline at end of file
+end module test
diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
index 11d055a..5fb2981 100644
--- a/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
+++ b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
@@ -33,48 +33,45 @@ program derived_acc
!$acc exit data copyout(var)
!$acc exit data copyout(var%a)
- !$acc data copy(var%a) ! { dg-error "Syntax error in OpenMP" }
- !$acc end data ! { dg-error "Unexpected ..ACC END DATA" }
-
!$acc data copy(var)
!$acc end data
- !$acc data copyout(var%a) ! { dg-error "Syntax error in OpenMP" }
- !$acc end data ! { dg-error "Unexpected ..ACC END" }
+ !$acc data copyout(var%a)
+ !$acc end data
!$acc parallel loop pcopyout(var)
do i = 1, 10
end do
!$acc end parallel loop
- !$acc parallel loop copyout(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc parallel loop copyout(var%a)
do i = 1, 10
end do
- !$acc end parallel loop ! { dg-error "Unexpected ..ACC END" }
+ !$acc end parallel loop
!$acc parallel pcopy(var)
!$acc end parallel
- !$acc parallel pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc parallel pcopy(var%a)
do i = 1, 10
end do
- !$acc end parallel ! { dg-error "Unexpected ..ACC END" }
+ !$acc end parallel
!$acc kernels pcopyin(var)
!$acc end kernels
- !$acc kernels pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc kernels pcopy(var%a)
do i = 1, 10
end do
- !$acc end kernels ! { dg-error "Unexpected ..ACC END" }
+ !$acc end kernels
!$acc kernels loop pcopyin(var)
do i = 1, 10
end do
!$acc end kernels loop
- !$acc kernels loop pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc kernels loop pcopy(var%a)
do i = 1, 10
end do
- !$acc end kernels loop ! { dg-error "Unexpected ..ACC END" }
+ !$acc end kernels loop
end program derived_acc
diff --git a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95 b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
index 805459c..b616b39 100644
--- a/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
+++ b/gcc/testsuite/gfortran.dg/goacc/enter-exit-data.f95
@@ -44,14 +44,14 @@ contains
!$acc enter data wait (i, 1)
!$acc enter data wait (a) ! { dg-error "INTEGER" }
!$acc enter data wait (b(5:6)) ! { dg-error "INTEGER" }
- !$acc enter data copyin (tip) ! { dg-error "POINTER" }
- !$acc enter data copyin (tia) ! { dg-error "ALLOCATABLE" }
- !$acc enter data create (tip) ! { dg-error "POINTER" }
- !$acc enter data create (tia) ! { dg-error "ALLOCATABLE" }
- !$acc enter data present_or_copyin (tip) ! { dg-error "POINTER" }
- !$acc enter data present_or_copyin (tia) ! { dg-error "ALLOCATABLE" }
- !$acc enter data present_or_create (tip) ! { dg-error "POINTER" }
- !$acc enter data present_or_create (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc enter data copyin (tip)
+ !$acc enter data copyin (tia)
+ !$acc enter data create (tip)
+ !$acc enter data create (tia)
+ !$acc enter data present_or_copyin (tip)
+ !$acc enter data present_or_copyin (tia)
+ !$acc enter data present_or_create (tip)
+ !$acc enter data present_or_create (tia)
!$acc enter data copyin (i) create (i) ! { dg-error "multiple clauses" }
!$acc enter data copyin (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
!$acc enter data create (i) present_or_copyin (i) ! { dg-error "multiple clauses" }
@@ -79,10 +79,10 @@ contains
!$acc exit data wait (i, 1)
!$acc exit data wait (a) ! { dg-error "INTEGER" }
!$acc exit data wait (b(5:6)) ! { dg-error "INTEGER" }
- !$acc exit data copyout (tip) ! { dg-error "POINTER" }
- !$acc exit data copyout (tia) ! { dg-error "ALLOCATABLE" }
- !$acc exit data delete (tip) ! { dg-error "POINTER" }
- !$acc exit data delete (tia) ! { dg-error "ALLOCATABLE" }
+ !$acc exit data copyout (tip)
+ !$acc exit data copyout (tia)
+ !$acc exit data delete (tip)
+ !$acc exit data delete (tia)
!$acc exit data copyout (i) delete (i) ! { dg-error "multiple clauses" }
!$acc exit data finalize
!$acc exit data finalize copyout (i)
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index acf7f8f..17fe0d3 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -806,6 +806,8 @@ struct target_var_desc {
bool copy_from;
/* True if data always should be copied from device to host at the end. */
bool always_copy_from;
+ /* True if variable should be detached at end of region. */
+ bool do_detach;
/* Relative offset against key host_start. */
uintptr_t offset;
/* Actual length. */
@@ -860,6 +862,8 @@ struct splay_tree_key_s {
uintptr_t refcount;
/* Dynamic reference count. */
uintptr_t dynamic_refcount;
+ /* For a block with attached pointers, the attachment counters for each. */
+ unsigned short *attach_count;
/* Pointer to the original mapping of "omp declare target link" object. */
splay_tree_key link_key;
};
@@ -1003,6 +1007,8 @@ enum gomp_map_vars_kind
GOMP_MAP_VARS_ENTER_DATA
};
+struct gomp_coalesce_buf;
+
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *, int);
extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
@@ -1013,8 +1019,17 @@ extern void gomp_copy_host2dev (struct gomp_device_descr *,
void *, const void *, size_t,
struct gomp_coalesce_buf *);
extern void gomp_copy_dev2host (struct gomp_device_descr *,
- struct goacc_asyncqueue *,
- void *, const void *, size_t);
+ struct goacc_asyncqueue *, void *, const void *,
+ size_t);
+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 *);
+extern void gomp_detach_pointer (struct gomp_device_descr *,
+ struct goacc_asyncqueue *, splay_tree_key,
+ uintptr_t, bool, struct gomp_coalesce_buf *);
+
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
size_t, void **, void **,
size_t *, void *, bool,
@@ -1025,9 +1040,9 @@ extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *,
size_t *, void *, bool,
enum gomp_map_vars_kind);
extern void gomp_unmap_tgt (struct target_mem_desc *);
-extern void gomp_unmap_vars (struct target_mem_desc *, bool);
+extern void gomp_unmap_vars (struct target_mem_desc *, bool, bool);
extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
- struct goacc_asyncqueue *);
+ struct goacc_asyncqueue *, bool);
extern void gomp_init_device (struct gomp_device_descr *);
extern bool gomp_fini_device (struct gomp_device_descr *);
extern void gomp_unload_device (struct gomp_device_descr *);
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index 595b988..cc1ce2a 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -440,6 +440,16 @@ OACC_2.5 {
acc_update_self_async_array_h_;
} OACC_2.0.1;
+OACC_2.6 {
+ global:
+ acc_attach;
+ acc_attach_async;
+ acc_detach;
+ acc_detach_async;
+ acc_detach_finalize;
+ acc_detach_finalize_async;
+} OACC_2.5;
+
GOACC_2.0 {
global:
GOACC_data_end;
diff --git a/libgomp/oacc-async.c b/libgomp/oacc-async.c
index bb00279..6c12c82 100644
--- a/libgomp/oacc-async.c
+++ b/libgomp/oacc-async.c
@@ -373,14 +373,14 @@ goacc_async_unmap_tgt (void *ptr)
attribute_hidden void
goacc_async_copyout_unmap_vars (struct target_mem_desc *tgt,
- struct goacc_asyncqueue *aq)
+ struct goacc_asyncqueue *aq, bool finalize)
{
struct gomp_device_descr *devicep = tgt->device_descr;
/* Increment reference to delay freeing of device memory until callback
has triggered. */
tgt->refcount++;
- gomp_unmap_vars_async (tgt, true, aq);
+ gomp_unmap_vars_async (tgt, true, aq, finalize);
devicep->openacc.async.queue_callback_func (aq, goacc_async_unmap_tgt,
(void *) tgt);
}
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 48c9646..e1938c5 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -391,7 +391,7 @@ acc_shutdown_1 (acc_device_t d)
{
struct target_mem_desc *tgt = walk->dev->mem_map.root->key.tgt;
- gomp_unmap_vars (tgt, false);
+ gomp_unmap_vars (tgt, false, false);
}
walk->dev = NULL;
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 1f6c62c..878f0f4 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -112,7 +112,7 @@ void goacc_host_init (void);
void goacc_init_asyncqueues (struct gomp_device_descr *);
bool goacc_fini_asyncqueues (struct gomp_device_descr *);
void goacc_async_copyout_unmap_vars (struct target_mem_desc *,
- struct goacc_asyncqueue *);
+ struct goacc_asyncqueue *, bool);
void goacc_async_free (struct gomp_device_descr *,
struct goacc_asyncqueue *, void *);
struct goacc_asyncqueue *get_goacc_asyncqueue (int);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index e5ee956..76ba914 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -518,7 +518,7 @@ acc_unmap_data (void *h)
gomp_mutex_unlock (&acc_dev->lock);
- gomp_unmap_vars (t, true);
+ gomp_unmap_vars (t, true, false);
if (profiling_setup_p)
{
@@ -612,6 +612,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
&kinds, true, GOMP_MAP_VARS_OPENACC);
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
+ tgt->list[0].key->attach_count = NULL;
gomp_mutex_lock (&acc_dev->lock);
@@ -750,6 +751,7 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
{
n->refcount = 0;
n->dynamic_refcount = 0;
+ n->attach_count = NULL;
}
if (n->refcount < n->dynamic_refcount)
{
@@ -997,6 +999,7 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
+ tgt->list[0].key->attach_count = NULL;
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
@@ -1084,11 +1087,11 @@ gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
/* If running synchronously, unmap immediately. */
if (async < acc_async_noval)
- gomp_unmap_vars (t, true);
+ gomp_unmap_vars (t, true, finalize);
else
{
goacc_aq aq = get_goacc_asyncqueue (async);
- goacc_async_copyout_unmap_vars (t, aq);
+ goacc_async_copyout_unmap_vars (t, aq, finalize);
}
}
@@ -1096,3 +1099,80 @@ gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
}
+
+
+void
+acc_attach_async (void **hostaddr, int async)
+{
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+ goacc_aq aq = get_goacc_asyncqueue (async);
+
+ struct splay_tree_key_s cur_node;
+ splay_tree_key n;
+
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
+ cur_node.host_start = (uintptr_t) hostaddr;
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ gomp_fatal ("struct not mapped for acc_attach");
+
+ gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
+ 0, NULL);
+}
+
+void
+acc_attach (void **hostaddr)
+{
+ acc_attach_async (hostaddr, acc_async_sync);
+}
+
+static void
+goacc_detach_internal (void **hostaddr, int async, bool finalize)
+{
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+ struct splay_tree_key_s cur_node;
+ splay_tree_key n;
+ struct goacc_asyncqueue *aq = get_goacc_asyncqueue (async);
+
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
+ cur_node.host_start = (uintptr_t) hostaddr;
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ gomp_fatal ("struct not mapped for acc_detach");
+
+ gomp_detach_pointer (acc_dev, aq, n, (uintptr_t) hostaddr, finalize, NULL);
+}
+
+void
+acc_detach (void **hostaddr)
+{
+ goacc_detach_internal (hostaddr, acc_async_sync, false);
+}
+
+void
+acc_detach_async (void **hostaddr, int async)
+{
+ goacc_detach_internal (hostaddr, async, false);
+}
+
+void
+acc_detach_finalize (void **hostaddr)
+{
+ goacc_detach_internal (hostaddr, acc_async_sync, true);
+}
+
+void
+acc_detach_finalize_async (void **hostaddr, int async)
+{
+ goacc_detach_internal (hostaddr, async, true);
+}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index 15b1462..f6c9114 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -50,12 +50,29 @@ find_pointer (int pos, size_t mapnum, unsigned short *kinds)
if (pos + 1 >= mapnum)
return 0;
- unsigned char kind = kinds[pos+1] & 0xff;
+ unsigned char kind0 = kinds[pos] & 0xff;
- if (kind == GOMP_MAP_TO_PSET)
- return 3;
- else if (kind == GOMP_MAP_POINTER)
- return 2;
+ switch (kind0)
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ {
+ unsigned char kind1 = kinds[pos + 1] & 0xff;
+ if (kind1 == GOMP_MAP_POINTER
+ || kind1 == GOMP_MAP_ALWAYS_POINTER
+ || kind1 == GOMP_MAP_ATTACH
+ || kind1 == GOMP_MAP_DETACH)
+ return 2;
+ else if (kind1 == GOMP_MAP_TO_PSET)
+ return 3;
+ }
+ default:
+ /* empty. */;
+ }
return 0;
}
@@ -355,14 +372,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++)
- {
- if (tgt->list[i].key != NULL)
- devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset
- + tgt->list[i].offset);
- else
- devaddrs[i] = NULL;
- }
+ devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, i);
if (aq == NULL)
{
@@ -382,7 +392,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
&api_info);
}
/* If running synchronously, unmap immediately. */
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
if (profiling_dispatch_p)
{
prof_info.event_type = acc_ev_exit_data_end;
@@ -400,7 +410,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
else
acc_dev->openacc.async.exec_func (tgt_fn, mapnum, hostaddrs,
devaddrs, dims, tgt, aq);
- goacc_async_copyout_unmap_vars (tgt, aq);
+ goacc_async_copyout_unmap_vars (tgt, aq, false);
}
out:
@@ -637,7 +647,7 @@ GOACC_data_end (void)
gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
thr->mapped_data = tgt->prev;
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
if (profiling_dispatch_p)
@@ -668,6 +678,10 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (mapnum > 0)
{
unsigned char kind = kinds[0] & 0xff;
+
+ if (kind == GOMP_MAP_STRUCT || kind == GOMP_MAP_FORCE_PRESENT)
+ kind = kinds[1] & 0xff;
+
if (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM)
finalize = true;
@@ -678,11 +692,14 @@ GOACC_enter_exit_data (int device, size_t mapnum,
{
unsigned char kind = kinds[i] & 0xff;
- if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+ if (kind == GOMP_MAP_POINTER
+ || kind == GOMP_MAP_TO_PSET
+ || kind == GOMP_MAP_STRUCT
+ || kind == GOMP_MAP_FORCE_PRESENT)
continue;
if (kind == GOMP_MAP_FORCE_ALLOC
- || kind == GOMP_MAP_FORCE_PRESENT
+ || kind == GOMP_MAP_ATTACH
|| kind == GOMP_MAP_FORCE_TO
|| kind == GOMP_MAP_TO
|| kind == GOMP_MAP_ALLOC
@@ -694,6 +711,8 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (kind == GOMP_MAP_RELEASE
|| kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_DETACH
+ || kind == GOMP_MAP_FORCE_DETACH
|| kind == GOMP_MAP_FROM
|| kind == GOMP_MAP_FORCE_FROM
|| kind == GOMP_MAP_DECLARE_DEALLOCATE)
@@ -809,6 +828,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
case GOMP_MAP_ALLOC:
acc_present_or_create (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_FORCE_PRESENT:
+ break;
case GOMP_MAP_FORCE_ALLOC:
acc_create (hostaddrs[i], sizes[i]);
break;
@@ -818,6 +840,27 @@ GOACC_enter_exit_data (int device, size_t mapnum,
case GOMP_MAP_FORCE_TO:
acc_copyin (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_STRUCT:
+ {
+ int elems = sizes[i];
+ struct splay_tree_key_s k;
+ splay_tree_key str;
+ k.host_start = (uintptr_t) hostaddrs[i];
+ k.host_end = k.host_start + 1;
+ gomp_mutex_lock (&acc_dev->lock);
+ str = splay_tree_lookup (&acc_dev->mem_map, &k);
+ gomp_mutex_unlock (&acc_dev->lock);
+ /* We increment the dynamic reference count for the struct
+ itself by the number of struct elements that we
+ mapped. */
+ if (str->refcount != REFCOUNT_INFINITY)
+ {
+ str->refcount += elems;
+ str->dynamic_refcount += elems;
+ }
+ i += elems;
+ }
+ break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
@@ -839,16 +882,57 @@ GOACC_enter_exit_data (int device, size_t mapnum,
i += pointer - 1;
}
}
+
+ /* This loop only handles explicit "attach" clauses that are not an
+ implicit part of a copy{,in,out}, etc. mapping. */
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ /* Scan for pointers and PSETs. */
+ int pointer = find_pointer (i, mapnum, kinds);
+
+ if (!pointer)
+ {
+ if (kind == GOMP_MAP_ATTACH)
+ acc_attach (hostaddrs[i]);
+ else if (kind == GOMP_MAP_STRUCT)
+ i += sizes[i];
+ }
+ else
+ i += pointer - 1;
+ }
}
else
- for (i = 0; i < mapnum; ++i)
- {
- unsigned char kind = kinds[i] & 0xff;
+ {
+ /* This loop only handles explicit "detach" clauses that are not an
+ implicit part of a copy{,in,out}, etc. mapping. */
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
- int pointer = find_pointer (i, mapnum, kinds);
+ int pointer = find_pointer (i, mapnum, kinds);
- if (!pointer)
- {
+ if (!pointer)
+ {
+ if (kind == GOMP_MAP_DETACH)
+ acc_detach (hostaddrs[i]);
+ else if (kind == GOMP_MAP_FORCE_DETACH)
+ acc_detach_finalize (hostaddrs[i]);
+ else if (kind == GOMP_MAP_STRUCT)
+ i += sizes[i];
+ }
+ else
+ i += pointer - 1;
+ }
+
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ int pointer = find_pointer (i, mapnum, kinds);
+
+ if (!pointer)
switch (kind)
{
case GOMP_MAP_RELEASE:
@@ -861,6 +945,10 @@ GOACC_enter_exit_data (int device, size_t mapnum,
acc_delete_async (hostaddrs[i], sizes[i], async);
}
break;
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
+ case GOMP_MAP_FORCE_PRESENT:
+ break;
case GOMP_MAP_DECLARE_DEALLOCATE:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
@@ -869,28 +957,48 @@ GOACC_enter_exit_data (int device, size_t mapnum,
else
acc_copyout_async (hostaddrs[i], sizes[i], async);
break;
+ case GOMP_MAP_STRUCT:
+ {
+ int elems = sizes[i];
+ struct splay_tree_key_s k;
+ splay_tree_key str;
+ k.host_start = (uintptr_t) hostaddrs[i];
+ k.host_end = k.host_start + 1;
+ gomp_mutex_lock (&acc_dev->lock);
+ str = splay_tree_lookup (&acc_dev->mem_map, &k);
+ gomp_mutex_unlock (&acc_dev->lock);
+ /* Decrement dynamic reference count for the struct by the
+ number of elements that we are unmapping. */
+ if (str->dynamic_refcount >= elems)
+ {
+ str->dynamic_refcount -= elems;
+ str->refcount -= elems;
+ }
+ i += elems;
+ }
+ break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
break;
}
- }
- else
- {
- if (kind == GOMP_MAP_DECLARE_DEALLOCATE)
- gomp_acc_declare_allocate (false, pointer, &hostaddrs[i],
- &sizes[i], &kinds[i]);
- else
- {
- bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
- || kind == GOMP_MAP_FROM);
- gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
- finalize, pointer);
- /* See the above comment. */
- }
- i += pointer - 1;
- }
- }
+ else
+ {
+ if (kind == GOMP_MAP_DECLARE_DEALLOCATE)
+ gomp_acc_declare_allocate (false, pointer, &hostaddrs[i],
+ &sizes[i], &kinds[i]);
+ else
+ {
+ bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_FROM);
+ gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom,
+ async, finalize, pointer);
+ /* See the above comment. */
+ }
+ i += pointer - 1;
+ }
+ }
+ }
out:
if (profiling_dispatch_p)
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 261636c..41dd514 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -113,6 +113,10 @@ void *acc_hostptr (void *) __GOACC_NOTHROW;
int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_attach (void **) __GOACC_NOTHROW;
+void acc_attach_async (void **, int) __GOACC_NOTHROW;
+void acc_detach (void **) __GOACC_NOTHROW;
+void acc_detach_async (void **, int) __GOACC_NOTHROW;
/* Async functions, specified in OpenACC 2.5. */
void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW;
@@ -129,6 +133,8 @@ void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
+void acc_detach_finalize (void **) __GOACC_NOTHROW;
+void acc_detach_finalize_async (void **, int) __GOACC_NOTHROW;
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
diff --git a/libgomp/target.c b/libgomp/target.c
index 7220ac6..d9d42eb 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -39,6 +39,7 @@
#include <string.h>
#include <assert.h>
#include <errno.h>
+#include <limits.h>
#ifdef PLUGIN_SUPPORT
#include <dlfcn.h>
@@ -373,6 +374,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
tgt_var->key = oldn;
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
+ tgt_var->do_detach = false;
tgt_var->offset = newn->host_start - oldn->host_start;
tgt_var->length = newn->host_end - newn->host_start;
@@ -539,7 +541,128 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
(void *) cur_node.host_end);
}
-static inline uintptr_t
+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 splay_tree_key_s s;
+ size_t size, idx;
+
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("enclosing struct not mapped for attach");
+ }
+
+ size = (n->host_end - n->host_start + sizeof (void *) - 1) / sizeof (void *);
+ /* We might have a pointer in a packed struct: however we cannot have more
+ than one such pointer in each pointer-sized portion of the struct, so
+ this is safe. */
+ idx = (attach_to - n->host_start) / sizeof (void *);
+
+ if (!n->attach_count)
+ n->attach_count = gomp_malloc_cleared (sizeof (*n->attach_count) * size);
+
+ if (n->attach_count[idx] < USHRT_MAX)
+ n->attach_count[idx]++;
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count overflow");
+ }
+
+ if (n->attach_count[idx] == 1)
+ {
+ uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + attach_to
+ - n->host_start;
+ uintptr_t target = (uintptr_t) *(void **) attach_to;
+ splay_tree_key tn;
+ uintptr_t data;
+
+ if ((void *) target == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attempt to attach null pointer");
+ }
+
+ s.host_start = target + bias;
+ s.host_end = s.host_start + 1;
+ tn = splay_tree_lookup (mem_map, &s);
+
+ if (!tn)
+ {
+ 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;
+
+ gomp_debug (1,
+ "%s: attaching host %p, target %p (struct base %p) to %p\n",
+ __FUNCTION__, (void *) attach_to, (void *) devptr,
+ (void *) (n->tgt->tgt_start + n->tgt_offset), (void *) data);
+
+ gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
+ sizeof (void *), cbufp);
+ }
+ else
+ gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+ (void *) attach_to, n->attach_count[idx]);
+}
+
+void
+gomp_detach_pointer (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, splay_tree_key n,
+ uintptr_t detach_from, bool finalize,
+ struct gomp_coalesce_buf *cbufp)
+{
+ size_t idx;
+
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("enclosing struct not mapped for detach");
+ }
+
+ idx = (detach_from - n->host_start) / sizeof (void *);
+
+ if (!n->attach_count)
+ gomp_fatal ("no attachment counters for struct");
+
+ if (finalize)
+ n->attach_count[idx] = 1;
+
+ if (n->attach_count[idx] == 0)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count underflow");
+ }
+ else
+ n->attach_count[idx]--;
+
+ if (n->attach_count[idx] == 0)
+ {
+ uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + detach_from
+ - n->host_start;
+ uintptr_t target = (uintptr_t) *(void **) detach_from;
+
+ gomp_debug (1,
+ "%s: detaching host %p, target %p (struct base %p) to %p\n",
+ __FUNCTION__, (void *) detach_from, (void *) devptr,
+ (void *) (n->tgt->tgt_start + n->tgt_offset),
+ (void *) target);
+
+ gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &target,
+ sizeof (void *), cbufp);
+ }
+ else
+ gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+ (void *) detach_from, n->attach_count[idx]);
+}
+
+uintptr_t
gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
{
if (tgt->list[i].key != NULL)
@@ -883,7 +1006,12 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
da->map_index = i;
continue;
}
-
+ else if ((kind & typemask) == GOMP_MAP_ATTACH)
+ {
+ tgt->list[i].key = NULL;
+ has_firstprivate = true;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -1141,6 +1269,30 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_ATTACH:
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n != NULL)
+ {
+ tgt->list[i].key = n;
+ tgt->list[i].offset = cur_node.host_start - n->host_start;
+ tgt->list[i].length = n->host_end - n->host_start;
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = true;
+ }
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("outer struct not mapped for attach");
+ }
+ gomp_attach_pointer (devicep, aq, mem_map, n,
+ (uintptr_t) hostaddrs[i], sizes[i],
+ cbufp);
+ continue;
+ }
default:
break;
}
@@ -1194,10 +1346,12 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+ tgt->list[i].do_detach = false;
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
k->dynamic_refcount = 0;
+ k->attach_count = NULL;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1482,6 +1636,8 @@ gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
is_tgt_unmapped = true;
gomp_unmap_tgt (k->tgt);
}
+ if (k->attach_count)
+ free (k->attach_count);
return is_tgt_unmapped;
}
@@ -1490,14 +1646,14 @@ gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
has been done already. */
attribute_hidden void
-gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
+gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, bool finalize)
{
- gomp_unmap_vars_async (tgt, do_copyfrom, NULL);
+ gomp_unmap_vars_async (tgt, do_copyfrom, NULL, finalize);
}
attribute_hidden void
gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
- struct goacc_asyncqueue *aq)
+ struct goacc_asyncqueue *aq, bool finalize)
{
struct gomp_device_descr *devicep = tgt->device_descr;
@@ -1517,10 +1673,23 @@ gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
}
size_t i;
+
+ /* We must perform detachments before any copies back to the host. */
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
- if (k == NULL)
+
+ if (k != NULL && tgt->list[i].do_detach)
+ gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
+ + tgt->list[i].offset, finalize,
+ NULL);
+ }
+
+ for (i = 0; i < tgt->list_count; i++)
+ {
+ splay_tree_key k = tgt->list[i].key;
+
+ if (k == NULL || tgt->list[i].do_detach)
continue;
bool do_unmap = false;
@@ -2139,7 +2308,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
GOMP_MAP_VARS_TARGET);
devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start,
NULL);
- gomp_unmap_vars (tgt_vars, true);
+ gomp_unmap_vars (tgt_vars, true, false);
}
/* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
@@ -2283,7 +2452,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
args);
if (tgt_vars)
- gomp_unmap_vars (tgt_vars, true);
+ gomp_unmap_vars (tgt_vars, true, false);
}
/* Host fallback for GOMP_target_data{,_ext} routines. */
@@ -2352,7 +2521,7 @@ GOMP_target_end_data (void)
{
struct target_mem_desc *tgt = icv->target_data;
icv->target_data = tgt->prev;
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
}
}
@@ -2587,7 +2756,7 @@ gomp_target_task_fn (void *data)
if (ttask->state == GOMP_TARGET_TASK_FINISHED)
{
if (ttask->tgt)
- gomp_unmap_vars (ttask->tgt, true);
+ gomp_unmap_vars (ttask->tgt, true, false);
return false;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
new file mode 100644
index 0000000..d8d7067
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-1.c
@@ -0,0 +1,24 @@
+#include <stdlib.h>
+#include <assert.h>
+
+struct dc
+{
+ int a;
+ int *b;
+};
+
+int
+main ()
+{
+ int n = 100, i;
+ struct dc v = { .a = 3, .b = (int *) malloc (sizeof (int) * n) };
+
+#pragma acc parallel loop copy(v.a, v.b[:n])
+ for (i = 0; i < n; i++)
+ v.b[i] = v.a;
+
+ for (i = 0; i < 10; i++)
+ assert (v.b[i] == v.a);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
new file mode 100644
index 0000000..7e26e9a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-2.c
@@ -0,0 +1,29 @@
+#include <assert.h>
+#include <stdlib.h>
+
+int
+main(int argc, char* argv[])
+{
+ struct foo {
+ int *a, *b, c, d, *e;
+ } s;
+
+ s.a = (int *) malloc (16 * sizeof (int));
+ s.b = (int *) malloc (16 * sizeof (int));
+ s.e = (int *) malloc (16 * sizeof (int));
+
+ #pragma acc data copy(s)
+ {
+ #pragma acc data copy(s.a[0:10])
+ {
+ #pragma acc parallel loop attach(s.a)
+ for (int i = 0; i < 10; i++)
+ s.a[i] = i;
+ }
+ }
+
+ for (int i = 0; i < 10; i++)
+ assert (s.a[i] == i);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
new file mode 100644
index 0000000..cec764b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-3.c
@@ -0,0 +1,34 @@
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+int
+main ()
+{
+ int n = 100, i;
+ int *a = (int *) malloc (sizeof (int) * n);
+ int *b;
+
+ for (i = 0; i < n; i++)
+ a[i] = i+1;
+
+#pragma acc enter data copyin(a[:n]) create(b)
+
+ b = a;
+ acc_attach ((void **)&b);
+
+#pragma acc parallel loop present (b[:n])
+ for (i = 0; i < n; i++)
+ b[i] = i+1;
+
+ acc_detach ((void **)&b);
+
+#pragma acc exit data copyout(a[:n], b)
+
+ for (i = 0; i < 10; i++)
+ assert (a[i] == b[i]);
+
+ free (a);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
new file mode 100644
index 0000000..8874ca0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-4.c
@@ -0,0 +1,87 @@
+#include <assert.h>
+#include <stdlib.h>
+
+#define LIST_LENGTH 10
+
+struct node
+{
+ struct node *next;
+ int val;
+};
+
+int
+sum_nodes (struct node *head)
+{
+ int i = 0, sum = 0;
+
+#pragma acc parallel reduction(+:sum) present(head[:1])
+ {
+ for (; head != NULL; head = head->next)
+ sum += head->val;
+ }
+
+ return sum;
+}
+
+void
+insert (struct node *head, int val)
+{
+ struct node *n = (struct node *) malloc (sizeof (struct node));
+
+ if (head->next)
+ {
+#pragma acc exit data detach(head->next)
+ }
+
+ n->val = val;
+ n->next = head->next;
+ head->next = n;
+
+#pragma acc enter data copyin(n[:1])
+#pragma acc enter data attach(head->next)
+ if (n->next)
+ {
+#pragma acc enter data attach(n->next)
+ }
+}
+
+void
+destroy (struct node *head)
+{
+ while (head->next != NULL)
+ {
+#pragma acc exit data detach(head->next)
+ struct node * n = head->next;
+ head->next = n->next;
+ if (n->next)
+ {
+#pragma acc exit data detach(n->next)
+ }
+#pragma acc exit data delete (n[:1])
+ if (head->next)
+ {
+#pragma acc enter data attach(head->next)
+ }
+ free (n);
+ }
+}
+
+int
+main ()
+{
+ struct node list = { .next = NULL, .val = 0 };
+ int i;
+
+#pragma acc enter data copyin(list)
+
+ for (i = 0; i < LIST_LENGTH; i++)
+ insert (&list, i + 1);
+
+ assert (sum_nodes (&list) == (LIST_LENGTH * LIST_LENGTH + LIST_LENGTH) / 2);
+
+ destroy (&list);
+
+#pragma acc exit data delete(list)
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
new file mode 100644
index 0000000..89cafbb
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-5.c
@@ -0,0 +1,81 @@
+#include <assert.h>
+#include <stdlib.h>
+#include <openacc.h>
+
+struct node
+{
+ struct node *next;
+ int val;
+};
+
+int
+sum_nodes (struct node *head)
+{
+ int i = 0, sum = 0;
+
+#pragma acc parallel reduction(+:sum) present(head[:1])
+ {
+ for (; head != NULL; head = head->next)
+ sum += head->val;
+ }
+
+ return sum;
+}
+
+void
+insert (struct node *head, int val)
+{
+ struct node *n = (struct node *) malloc (sizeof (struct node));
+
+ if (head->next)
+ acc_detach ((void **) &head->next);
+
+ n->val = val;
+ n->next = head->next;
+ head->next = n;
+
+ acc_copyin (n, sizeof (struct node));
+ acc_attach((void **) &head->next);
+
+ if (n->next)
+ acc_attach ((void **) &n->next);
+}
+
+void
+destroy (struct node *head)
+{
+ while (head->next != NULL)
+ {
+ acc_detach ((void **) &head->next);
+ struct node * n = head->next;
+ head->next = n->next;
+ if (n->next)
+ acc_detach ((void **) &n->next);
+
+ acc_delete (n, sizeof (struct node));
+ if (head->next)
+ acc_attach((void **) &head->next);
+
+ free (n);
+ }
+}
+
+int
+main ()
+{
+ struct node list = { .next = NULL, .val = 0 };
+ int i;
+
+ acc_copyin (&list, sizeof (struct node));
+
+ for (i = 0; i < 10; i++)
+ insert (&list, 2);
+
+ assert (sum_nodes (&list) == 10 * 2);
+
+ destroy (&list);
+
+ acc_delete (&list, sizeof (struct node));
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
new file mode 100644
index 0000000..c4cea11
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-1.f90
@@ -0,0 +1,35 @@
+! { dg-do run }
+
+! Test of attach/detach with "acc data".
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+
+!$acc data copy(var)
+!$acc data copy(var%a)
+
+!$acc parallel loop
+ do i = 1,n
+ var%a(i) = i
+ end do
+!$acc end parallel loop
+
+!$acc end data
+!$acc end data
+
+ do i = 1,n
+ if (i .ne. var%a(i)) stop 1
+ end do
+
+ deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
new file mode 100644
index 0000000..3593661
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-2.f90
@@ -0,0 +1,33 @@
+! { dg-do run }
+
+! Test of attach/detach with "acc data", two clauses at once.
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+
+!$acc data copy(var) copy(var%a)
+
+!$acc parallel loop
+ do i = 1,n
+ var%a(i) = i
+ end do
+!$acc end parallel loop
+
+!$acc end data
+
+ do i = 1,n
+ if (i .ne. var%a(i)) stop 1
+ end do
+
+ deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
new file mode 100644
index 0000000..667d944
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-3.f90
@@ -0,0 +1,34 @@
+! { dg-do run }
+
+! Test of attach/detach with "acc parallel".
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ integer, allocatable :: b(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+ allocate(var%b(1:n))
+
+!$acc parallel loop copy(var) copy(var%a(1:n)) copy(var%b(1:n))
+ do i = 1,n
+ var%a(i) = i
+ var%b(i) = i
+ end do
+!$acc end parallel loop
+
+ do i = 1,n
+ if (i .ne. var%a(i)) stop 1
+ if (i .ne. var%b(i)) stop 2
+ end do
+
+ deallocate(var%a)
+ deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
new file mode 100644
index 0000000..6949e12
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-4.f90
@@ -0,0 +1,49 @@
+! { dg-do run }
+
+! Test of attach/detach with "acc enter/exit data".
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ integer, allocatable :: b(:)
+ end type mytype
+ integer, allocatable :: r(:)
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+ allocate(var%b(1:n))
+ allocate(r(1:n))
+
+!$acc enter data copyin(var)
+
+!$acc enter data copyin(var%a, var%b, r)
+
+!$acc parallel loop
+ do i = 1,n
+ var%a(i) = i
+ var%b(i) = i * 2
+ r(i) = i * 3
+ end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a)
+!$acc exit data copyout(var%b)
+!$acc exit data copyout(r)
+
+ do i = 1,n
+ if (i .ne. var%a(i)) stop 1
+ if (i * 2 .ne. var%b(i)) stop 2
+ if (i * 3 .ne. r(i)) stop 3
+ end do
+
+!$acc exit data delete(var)
+
+ deallocate(var%a)
+ deallocate(var%b)
+ deallocate(r)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
new file mode 100644
index 0000000..6843cf1
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-5.f90
@@ -0,0 +1,57 @@
+! { dg-do run }
+
+! Test of attach/detach, "enter data" inside "data", and subarray.
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ integer, allocatable :: b(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+ allocate(var%b(1:n))
+
+!$acc data copy(var)
+
+ do i = 1, n
+ var%a(i) = 0
+ var%b(i) = 0
+ end do
+
+!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5))
+
+!$acc parallel loop
+ do i = 5,n - 5
+ var%a(i) = i
+ var%b(i) = i * 2
+ end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5))
+
+!$acc end data
+
+ do i = 1,4
+ if (var%a(i) .ne. 0) stop 1
+ if (var%b(i) .ne. 0) stop 2
+ end do
+
+ do i = 5,n - 5
+ if (i .ne. var%a(i)) stop 3
+ if (i * 2 .ne. var%b(i)) stop 4
+ end do
+
+ do i = n - 4,n
+ if (var%a(i) .ne. 0) stop 5
+ if (var%b(i) .ne. 0) stop 6
+ end do
+
+ deallocate(var%a)
+ deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
new file mode 100644
index 0000000..12910d0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-6.f90
@@ -0,0 +1,61 @@
+! { dg-do run }
+
+! Test of attachment counters and finalize.
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ integer, allocatable :: b(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+ allocate(var%b(1:n))
+
+!$acc data copy(var)
+
+ do i = 1, n
+ var%a(i) = 0
+ var%b(i) = 0
+ end do
+
+!$acc enter data copyin(var%a(5:n - 5), var%b(5:n - 5))
+
+ do i = 1,20
+ !$acc enter data attach(var%a)
+ end do
+
+!$acc parallel loop
+ do i = 5,n - 5
+ var%a(i) = i
+ var%b(i) = i * 2
+ end do
+!$acc end parallel loop
+
+!$acc exit data copyout(var%a(5:n - 5), var%b(5:n - 5)) finalize
+
+!$acc end data
+
+ do i = 1,4
+ if (var%a(i) .ne. 0) stop 1
+ if (var%b(i) .ne. 0) stop 2
+ end do
+
+ do i = 5,n - 5
+ if (i .ne. var%a(i)) stop 3
+ if (i * 2 .ne. var%b(i)) stop 4
+ end do
+
+ do i = n - 4,n
+ if (var%a(i) .ne. 0) stop 5
+ if (var%b(i) .ne. 0) stop 6
+ end do
+
+ deallocate(var%a)
+ deallocate(var%b)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
new file mode 100644
index 0000000..ab44f0a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-7.f90
@@ -0,0 +1,89 @@
+! { dg-do run }
+
+! Test of attach/detach with scalar elements and nested derived types.
+
+program dtype
+ implicit none
+ integer, parameter :: n = 512
+ type subtype
+ integer :: g, h
+ integer, allocatable :: q(:)
+ end type subtype
+ type mytype
+ integer, allocatable :: a(:)
+ integer, allocatable :: c, d
+ integer, allocatable :: b(:)
+ integer :: f
+ type(subtype) :: s
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+ allocate(var%b(1:n))
+ allocate(var%c)
+ allocate(var%d)
+ allocate(var%s%q(1:n))
+
+ var%c = 16
+ var%d = 20
+ var%f = 7
+ var%s%g = 21
+ var%s%h = 38
+
+!$acc enter data copyin(var)
+
+ do i = 1, n
+ var%a(i) = 0
+ var%b(i) = 0
+ var%s%q(i) = 0
+ end do
+
+!$acc data copy(var%a(5:n - 5), var%b(5:n - 5), var%c, var%d) &
+!$acc & copy(var%s%q)
+
+!$acc parallel loop default(none) present(var)
+ do i = 5,n - 5
+ var%a(i) = i
+ var%b(i) = i * 2
+ var%s%q(i) = i * 3
+ var%s%g = 100
+ var%s%h = 101
+ end do
+!$acc end parallel loop
+
+!$acc end data
+
+!$acc exit data copyout(var)
+
+ do i = 1,4
+ if (var%a(i) .ne. 0) stop 1
+ if (var%b(i) .ne. 0) stop 2
+ if (var%s%q(i) .ne. 0) stop 3
+ end do
+
+ do i = 5,n - 5
+ if (i .ne. var%a(i)) stop 4
+ if (i * 2 .ne. var%b(i)) stop 5
+ if (i * 3 .ne. var%s%q(i)) stop 6
+ end do
+
+ do i = n - 4,n
+ if (var%a(i) .ne. 0) stop 7
+ if (var%b(i) .ne. 0) stop 8
+ if (var%s%q(i) .ne. 0) stop 9
+ end do
+
+ if (var%c .ne. 16) stop 10
+ if (var%d .ne. 20) stop 11
+ if (var%s%g .ne. 100 .or. var%s%h .ne. 101) stop 12
+ if (var%f .ne. 7) stop 13
+
+ deallocate(var%a)
+ deallocate(var%b)
+ deallocate(var%c)
+ deallocate(var%d)
+ deallocate(var%s%q)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90 b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
new file mode 100644
index 0000000..d142763
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/deep-copy-8.f90
@@ -0,0 +1,41 @@
+! { dg-do run }
+
+! Test of explicit attach/detach clauses and attachment counters. There are no
+! acc_attach/acc_detach API routines in Fortran.
+
+program dtype
+ use openacc
+ implicit none
+ integer, parameter :: n = 512
+ type mytype
+ integer, allocatable :: a(:)
+ end type mytype
+ integer i
+
+ type(mytype) :: var
+
+ allocate(var%a(1:n))
+
+ call acc_copyin(var)
+ call acc_copyin(var%a)
+
+ !$acc enter data attach(var%a)
+
+!$acc parallel loop attach(var%a)
+ do i = 1,n
+ var%a(i) = i
+ end do
+!$acc end parallel loop
+
+ !$acc exit data detach(var%a)
+
+ call acc_copyout(var%a)
+ call acc_copyout(var)
+
+ do i = 1,n
+ if (i .ne. var%a(i)) stop 1
+ end do
+
+ deallocate(var%a)
+
+end program dtype
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
index 1ec4784..eb7812d 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
@@ -20,9 +20,9 @@
!$acc end data
do i = 1, n
- if (d(i)%a /= i) call abort
- if (d(i)%b /= i-1) call abort
- if (d(i)%c /= i+1) call abort
+ if (d(i)%a /= i) stop 1
+ if (d(i)%b /= i-1) stop 2
+ if (d(i)%c /= i+1) stop 3
end do
end program
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
index a37d526..c3c8a07 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
@@ -37,7 +37,7 @@ program derived_acc
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 1
var%b = 100
@@ -51,7 +51,7 @@ program derived_acc
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 2
!$acc parallel loop present (var)
do i = 1, n
@@ -64,7 +64,7 @@ program derived_acc
var%a = -1
do i = 1, n
- if (var%c(i) /= i) call abort
+ if (var%c(i) /= i) stop 3
var%c(i) = var%a
end do
@@ -78,7 +78,7 @@ program derived_acc
if (var%c(i) /= var%a) res = res + 1
end do
- if (res /= 0) call abort
+ if (res /= 0) stop 4
var%c(:) = 0
@@ -93,8 +93,8 @@ program derived_acc
!$acc update host(var%c(5))
do i = 1, n
- if (i /= 5 .and. var%c(i) /= 0) call abort
- if (i == 5 .and. var%c(i) /= 1) call abort
+ if (i /= 5 .and. var%c(i) /= 0) stop 5
+ if (i == 5 .and. var%c(i) /= 1) stop 6
end do
!$acc parallel loop present(var)
@@ -106,7 +106,7 @@ program derived_acc
!$acc update host(var%in%d)
do i = 1, n
- if (var%in%d(i) /= var%a) call abort
+ if (var%in%d(i) /= var%a) stop 7
end do
var%c(:) = 0
@@ -124,8 +124,8 @@ program derived_acc
!$acc update host(var%c(n/2:n))
do i = 1,n
- if (i < n/2 .and. var%c(i) /= -1) call abort
- if (i >= n/2 .and. var%c(i) /= i) call abort
+ if (i < n/2 .and. var%c(i) /= -1) stop 8
+ if (i >= n/2 .and. var%c(i) /= i) stop 9
end do
var%in%d(:) = 0
@@ -140,8 +140,8 @@ program derived_acc
!$acc update host(var%in%d(5))
do i = 1, n
- if (i /= 5 .and. var%in%d(i) /= 0) call abort
- if (i == 5 .and. var%in%d(i) /= 1) call abort
+ if (i /= 5 .and. var%in%d(i) /= 0) stop 10
+ if (i == 5 .and. var%in%d(i) /= 1) stop 11
end do
!$acc exit data delete(var)
@@ -173,7 +173,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 12
var%b = 100
@@ -187,7 +187,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 13
!$acc parallel loop present (var)
do i = 1, n
@@ -200,7 +200,7 @@ subroutine derived_acc_subroutine(var)
var%a = -1
do i = 1, n
- if (var%c(i) /= i) call abort
+ if (var%c(i) /= i) stop 14
var%c(i) = var%a
end do
@@ -214,7 +214,7 @@ subroutine derived_acc_subroutine(var)
if (var%c(i) /= var%a) res = res + 1
end do
- if (res /= 0) call abort
+ if (res /= 0) stop 15
var%c(:) = 0
@@ -229,8 +229,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%c(5))
do i = 1, n
- if (i /= 5 .and. var%c(i) /= 0) call abort
- if (i == 5 .and. var%c(i) /= 1) call abort
+ if (i /= 5 .and. var%c(i) /= 0) stop 16
+ if (i == 5 .and. var%c(i) /= 1) stop 17
end do
!$acc parallel loop present(var)
@@ -242,7 +242,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%in%d)
do i = 1, n
- if (var%in%d(i) /= var%a) call abort
+ if (var%in%d(i) /= var%a) stop 18
end do
var%c(:) = 0
@@ -260,8 +260,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%c(n/2:n))
do i = 1,n
- if (i < n/2 .and. var%c(i) /= -1) call abort
- if (i >= n/2 .and. var%c(i) /= i) call abort
+ if (i < n/2 .and. var%c(i) /= -1) stop 19
+ if (i >= n/2 .and. var%c(i) /= i) stop 20
end do
var%in%d(:) = 0
@@ -276,8 +276,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%in%d(5))
do i = 1, n
- if (i /= 5 .and. var%in%d(i) /= 0) call abort
- if (i == 5 .and. var%in%d(i) /= 1) call abort
+ if (i /= 5 .and. var%in%d(i) /= 0) stop 21
+ if (i == 5 .and. var%in%d(i) /= 1) stop 22
end do
!$acc exit data delete(var)
^ permalink raw reply [flat|nested] 8+ messages in thread* [PATCH 6/6] [og8] OpenACC refcounting refresh
2018-11-20 21:55 [PATCH 0/6] [og8] OpenACC attach/detach Julian Brown
` (4 preceding siblings ...)
2018-11-20 21:56 ` [PATCH 5/6] [og8] Backport parts of upstream declare-allocate patch Julian Brown
@ 2018-11-20 21:57 ` Julian Brown
5 siblings, 0 replies; 8+ messages in thread
From: Julian Brown @ 2018-11-20 21:57 UTC (permalink / raw)
To: gcc-patches; +Cc: cltang, Catherine_Moore, jakub
[-- Attachment #1: Type: text/plain, Size: 3827 bytes --]
This patch represents a mild overhaul of reference counting for OpenACC
in libgomp. It's been partly automatically checked (using code not yet
quite finished nor submitted upstream), but it's already more precise
than the pre-patch implementation (as demonstrated by adjustments to
previously-erroneous tests, included).
I have a few more changes planned, but those are still tbd.
libgomp/
* libgomp.h (gomp_device_descr): Add GOMP_MAP_VARS_OPENACC_ENTER_DATA.
(gomp_acc_remove_pointer): Update prototype.
(gomp_acc_data_env_remove_tgt): Add prototype.
(gomp_unmap_vars, gomp_map_vars_async): Update prototype.
* oacc-int.h (goacc_async_copyout_unmap_vars): Update prototype.
* oacc-async.c (goacc_async_copyout_unmap_vars): Remove finalize
parameter.
* oacc-init.c (acc_shutdown_1): Remove finalize argument to
gomp_unmap_vars call.
* oacc-mem.c (lookup_dev_1): New helper function.
(lookup_dev): Rewrite in terms of above.
(acc_free): Update calls to lookup_dev.
(acc_map_data): Likewise. Don't add data mapped this way to OpenACC
data environment list.
(gomp_acc_data_env_remove, gomp_acc_data_env_remove_tgt): New functions.
(acc_unmap_data): Rewrite using splay tree functions directly. Don't
call gomp_unmap_vars. Fix refcount handling.
(present_create_copy): Use GOMP_MAP_VARS_OPENACC_ENTER_DATA in
gomp_map_vars_async call. Adjust refcount handling.
(delete_copyout): Remove dubious handling of target_mem_desc refcount.
(gomp_acc_insert_pointer): Use GOMP_MAP_VARS_OPENACC_ENTER_DATA in
gomp_map_vars_async call. Update refcount handling.
(gomp_acc_remove_pointer): Reimplement. Fix detach and refcount
handling.
* oacc-parallel.c (find_pointer): Handle more mapping types. Update
calls to gomp_unmap_vars and goacc_async_copyout_unmap_vars.
(GOACC_enter_exit_data): Update refcount handling.
libgomp/
* target.c (gomp_detach_pointer): Unlock device on error path.
(gomp_map_vars_async): Support GOMP_MAP_VARS_OPENACC_ENTER_DATA and
mapping size fix GOMP_MAP_ATTACH.
(gomp_unmap_tgt): Call gomp_acc_data_env_remove_tgt.
(gomp_unmap_vars): Remove finalize parameter.
(gomp_unmap_vars_async): Likewise. Adjust detach handling.
(GOMP_target, GOMP_target_ext, GOMP_target_end_data)
(gomp_target_task_fn): Update calls to gomp_unmap_vars.
* testsuite/libgomp.oacc-c-c++-common/context-2.c: Use correct API to
unmap data.
* testsuite/libgomp.oacc-c-c++-common/context-4.c: Likewise.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c: New test.
* testsuite/libgomp.oacc-fortran/data-2.f90: Fix for unmap semantics.
---
libgomp/libgomp.h | 10 +-
libgomp/oacc-async.c | 4 +-
libgomp/oacc-init.c | 2 +-
libgomp/oacc-int.h | 2 +-
libgomp/oacc-mem.c | 387 ++++++++++----------
libgomp/oacc-parallel.c | 76 +++--
libgomp/target.c | 35 ++-
.../libgomp.oacc-c-c++-common/context-2.c | 6 +-
.../libgomp.oacc-c-c++-common/context-4.c | 6 +-
.../libgomp.oacc-c-c++-common/deep-copy-6.c | 59 +++
.../libgomp.oacc-c-c++-common/deep-copy-7.c | 42 +++
.../libgomp.oacc-c-c++-common/deep-copy-8.c | 53 +++
libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 | 20 +-
13 files changed, 445 insertions(+), 257 deletions(-)
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
create mode 100644 libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
[-- Warning: decoded text below may be mangled, UTF-8 assumed --]
[-- Attachment #2: 0006-og8-OpenACC-refcounting-refresh.patch --]
[-- Type: text/x-patch; name="0006-og8-OpenACC-refcounting-refresh.patch", Size: 34989 bytes --]
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index 17fe0d3..568e260 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -1002,6 +1002,7 @@ struct gomp_device_descr
enum gomp_map_vars_kind
{
GOMP_MAP_VARS_OPENACC,
+ GOMP_MAP_VARS_OPENACC_ENTER_DATA,
GOMP_MAP_VARS_TARGET,
GOMP_MAP_VARS_DATA,
GOMP_MAP_VARS_ENTER_DATA
@@ -1010,7 +1011,8 @@ enum gomp_map_vars_kind
struct gomp_coalesce_buf;
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *, int);
-extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
+extern void gomp_acc_remove_pointer (void **, size_t *, unsigned short *,
+ int, void *, bool, int);
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
unsigned short *);
struct gomp_coalesce_buf;
@@ -1039,10 +1041,12 @@ extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *,
size_t, void **, void **,
size_t *, void *, bool,
enum gomp_map_vars_kind);
+extern void gomp_acc_data_env_remove_tgt (struct target_mem_desc **,
+ struct target_mem_desc *);
extern void gomp_unmap_tgt (struct target_mem_desc *);
-extern void gomp_unmap_vars (struct target_mem_desc *, bool, bool);
+extern void gomp_unmap_vars (struct target_mem_desc *, bool);
extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
- struct goacc_asyncqueue *, bool);
+ struct goacc_asyncqueue *);
extern void gomp_init_device (struct gomp_device_descr *);
extern bool gomp_fini_device (struct gomp_device_descr *);
extern void gomp_unload_device (struct gomp_device_descr *);
diff --git a/libgomp/oacc-async.c b/libgomp/oacc-async.c
index 6c12c82..bb00279 100644
--- a/libgomp/oacc-async.c
+++ b/libgomp/oacc-async.c
@@ -373,14 +373,14 @@ goacc_async_unmap_tgt (void *ptr)
attribute_hidden void
goacc_async_copyout_unmap_vars (struct target_mem_desc *tgt,
- struct goacc_asyncqueue *aq, bool finalize)
+ struct goacc_asyncqueue *aq)
{
struct gomp_device_descr *devicep = tgt->device_descr;
/* Increment reference to delay freeing of device memory until callback
has triggered. */
tgt->refcount++;
- gomp_unmap_vars_async (tgt, true, aq, finalize);
+ gomp_unmap_vars_async (tgt, true, aq);
devicep->openacc.async.queue_callback_func (aq, goacc_async_unmap_tgt,
(void *) tgt);
}
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index e1938c5..48c9646 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -391,7 +391,7 @@ acc_shutdown_1 (acc_device_t d)
{
struct target_mem_desc *tgt = walk->dev->mem_map.root->key.tgt;
- gomp_unmap_vars (tgt, false, false);
+ gomp_unmap_vars (tgt, false);
}
walk->dev = NULL;
diff --git a/libgomp/oacc-int.h b/libgomp/oacc-int.h
index 878f0f4..1f6c62c 100644
--- a/libgomp/oacc-int.h
+++ b/libgomp/oacc-int.h
@@ -112,7 +112,7 @@ void goacc_host_init (void);
void goacc_init_asyncqueues (struct gomp_device_descr *);
bool goacc_fini_asyncqueues (struct gomp_device_descr *);
void goacc_async_copyout_unmap_vars (struct target_mem_desc *,
- struct goacc_asyncqueue *, bool);
+ struct goacc_asyncqueue *);
void goacc_async_free (struct gomp_device_descr *,
struct goacc_asyncqueue *, void *);
struct goacc_asyncqueue *get_goacc_asyncqueue (int);
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 76ba914..3202f06 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -52,6 +52,25 @@ lookup_host (struct gomp_device_descr *dev, void *h, size_t s)
return key;
}
+/* Helper for lookup_dev. Iterate over splay tree. */
+
+static splay_tree_key
+lookup_dev_1 (splay_tree_node node, uintptr_t d, size_t s)
+{
+ splay_tree_key k = &node->key;
+ struct target_mem_desc *t = k->tgt;
+
+ if (d >= t->tgt_start && d + s <= t->tgt_end)
+ return k;
+
+ if (node->left)
+ return lookup_dev_1 (node->left, d, s);
+ if (node->right)
+ return lookup_dev_1 (node->right, d, s);
+
+ return NULL;
+}
+
/* Return block containing [D->S), or NULL if not contained.
The list isn't ordered by device address, so we have to iterate
over the whole array. This is not expected to be a common
@@ -59,35 +78,12 @@ lookup_host (struct gomp_device_descr *dev, void *h, size_t s)
remains locked on exit. */
static splay_tree_key
-lookup_dev (struct target_mem_desc *tgt, void *d, size_t s)
+lookup_dev (splay_tree mem_map, void *d, size_t s)
{
- int i;
- struct target_mem_desc *t;
-
- if (!tgt)
- return NULL;
-
- for (t = tgt; t != NULL; t = t->prev)
- {
- if (t->tgt_start <= (uintptr_t) d && t->tgt_end >= (uintptr_t) d + s)
- break;
- }
-
- if (!t)
+ if (!mem_map || !mem_map->root)
return NULL;
- for (i = 0; i < t->list_count; i++)
- {
- void * offset;
-
- splay_tree_key k = &t->array[i].key;
- offset = d - t->tgt_start + k->tgt_offset;
-
- if (k->host_start + offset <= (void *) k->host_end)
- return k;
- }
-
- return NULL;
+ return lookup_dev_1 (mem_map->root, (uintptr_t) d, s);
}
/* OpenACC is silent on how memory exhaustion is indicated. We return
@@ -165,7 +161,7 @@ acc_free (void *d)
/* We don't have to call lazy open here, as the ptr value must have
been returned by acc_malloc. It's not permitted to pass NULL in
(unless you got that null from acc_malloc). */
- if ((k = lookup_dev (acc_dev->openacc.data_environ, d, 1)))
+ if ((k = lookup_dev (&acc_dev->mem_map, d, 1)))
{
void *offset;
@@ -325,7 +321,7 @@ acc_hostptr (void *d)
gomp_mutex_lock (&acc_dev->lock);
- n = lookup_dev (acc_dev->openacc.data_environ, d, 1);
+ n = lookup_dev (&acc_dev->mem_map, d, 1);
if (!n)
{
@@ -422,7 +418,7 @@ acc_map_data (void *h, void *d, size_t s)
(int)s);
}
- if (lookup_dev (thr->dev->openacc.data_environ, d, s))
+ if (lookup_dev (&thr->dev->mem_map, d, s))
{
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("device address [%p, +%d] is already mapped", (void *)d,
@@ -436,11 +432,6 @@ acc_map_data (void *h, void *d, size_t s)
tgt->list[0].key->refcount = REFCOUNT_INFINITY;
}
- gomp_mutex_lock (&acc_dev->lock);
- tgt->prev = acc_dev->openacc.data_environ;
- acc_dev->openacc.data_environ = tgt;
- gomp_mutex_unlock (&acc_dev->lock);
-
if (profiling_setup_p)
{
thr->prof_info = NULL;
@@ -448,11 +439,83 @@ acc_map_data (void *h, void *d, size_t s)
}
}
+/* Remove the target_mem_desc holding the mapping for MAPNUM HOSTADDRS from
+ the OpenACC data environment pointed to by DATA_ENV. The device lock
+ should be held before calling, and remains locked on exit. */
+
+static void
+gomp_acc_data_env_remove (struct gomp_device_descr *acc_dev,
+ struct target_mem_desc **data_env, void **hostaddrs,
+ int mapnum)
+{
+ struct target_mem_desc *t, *tp;
+
+ for (tp = NULL, t = *data_env; t != NULL; tp = t, t = t->prev)
+ {
+ bool all_match = true;
+
+ /* We must locate the target descriptor by "value", matching each
+ hostaddr that it describes. */
+ if (t->list_count != mapnum)
+ continue;
+
+ for (int i = 0; i < t->list_count; i++)
+ if (t->list[i].key
+ && (t->list[i].key->host_start + t->list[i].offset
+ != (uintptr_t) hostaddrs[i]))
+ {
+ all_match = false;
+ break;
+ }
+
+ if (all_match)
+ {
+ if (t->refcount > 1)
+ t->refcount--;
+ else
+ {
+ if (tp)
+ tp->prev = t->prev;
+ else
+ *data_env = t->prev;
+ }
+ return;
+ }
+ }
+
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("cannot find data mapping to remove in data environment");
+}
+
+/* Similar, but removes target_mem_desc REMOVE from the DATA_ENV, in case its
+ reference count drops to zero resulting in it being unmapped (in
+ target.c:gomp_unmap_tgt). Unlike the above function it is not an error if
+ REMOVE is not present in the environment. The device lock should be held
+ before calling, and remains locked on exit. */
+
+attribute_hidden void
+gomp_acc_data_env_remove_tgt (struct target_mem_desc **data_env,
+ struct target_mem_desc *remove)
+{
+ struct target_mem_desc *t, *tp;
+
+ for (tp = NULL, t = *data_env; t != NULL; tp = t, t = t->prev)
+ if (t == remove)
+ {
+ if (tp)
+ tp->prev = t->prev;
+ else
+ *data_env = t->prev;
+ return;
+ }
+}
+
void
acc_unmap_data (void *h)
{
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ struct splay_tree_key_s cur_node;
/* No need to call lazy open, as the address must have been mapped. */
@@ -466,12 +529,11 @@ acc_unmap_data (void *h)
= __builtin_expect (goacc_profiling_setup_p (thr, &prof_info, &api_info),
false);
- size_t host_size;
-
gomp_mutex_lock (&acc_dev->lock);
- splay_tree_key n = lookup_host (acc_dev, h, 1);
- struct target_mem_desc *t;
+ cur_node.host_start = (uintptr_t) h;
+ cur_node.host_end = cur_node.host_start + 1;
+ splay_tree_key n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
if (!n)
{
@@ -479,47 +541,28 @@ acc_unmap_data (void *h)
gomp_fatal ("%p is not a mapped block", (void *)h);
}
- host_size = n->host_end - n->host_start;
-
if (n->host_start != (uintptr_t) h)
{
+ size_t host_size = n->host_end - n->host_start;
gomp_mutex_unlock (&acc_dev->lock);
gomp_fatal ("[%p,%d] surrounds %p",
(void *) n->host_start, (int) host_size, (void *) h);
}
- /* Mark for removal. */
- n->refcount = 1;
+ splay_tree_remove (&acc_dev->mem_map, n);
- t = n->tgt;
+ struct target_mem_desc *tgt = n->tgt;
- if (t->refcount == 2)
+ if (tgt->refcount > 0)
+ tgt->refcount--;
+ else
{
- struct target_mem_desc *tp;
-
- /* This is the last reference, so pull the descriptor off the
- chain. This avoids gomp_unmap_vars via gomp_unmap_tgt from
- freeing the device memory. */
- t->tgt_end = 0;
- t->to_free = 0;
-
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
-
- break;
- }
+ free (tgt->array);
+ free (tgt);
}
gomp_mutex_unlock (&acc_dev->lock);
- gomp_unmap_vars (t, true, false);
-
if (profiling_setup_p)
{
thr->prof_info = NULL;
@@ -585,6 +628,24 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
n->refcount++;
n->dynamic_refcount++;
}
+
+ struct target_mem_desc *tgt = gomp_malloc (sizeof (*tgt)
+ + sizeof (tgt->list[0]));
+ tgt->refcount = 1;
+ tgt->tgt_start = 0;
+ tgt->tgt_end = 0;
+ tgt->to_free = NULL;
+ tgt->prev = acc_dev->openacc.data_environ;
+ tgt->list_count = 1;
+ tgt->device_descr = acc_dev;
+ tgt->list[0].key = n;
+ tgt->list[0].copy_from = false;
+ tgt->list[0].always_copy_from = false;
+ tgt->list[0].do_detach = false;
+ tgt->list[0].offset = (uintptr_t) h - n->host_start;
+ tgt->list[0].length = 0;
+ acc_dev->openacc.data_environ = tgt;
+
gomp_mutex_unlock (&acc_dev->lock);
}
else if (!(f & FLAG_CREATE))
@@ -609,18 +670,19 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
goacc_aq aq = get_goacc_asyncqueue (async);
tgt = gomp_map_vars_async (acc_dev, aq, mapnum, &hostaddrs, NULL, &s,
- &kinds, true, GOMP_MAP_VARS_OPENACC);
- /* Initialize dynamic refcount. */
- tgt->list[0].key->dynamic_refcount = 1;
- tgt->list[0].key->attach_count = NULL;
+ &kinds, true,
+ GOMP_MAP_VARS_OPENACC_ENTER_DATA);
- gomp_mutex_lock (&acc_dev->lock);
+ for (int i = 0; i < tgt->list_count; i++)
+ if (tgt->list[i].key)
+ tgt->list[i].key->dynamic_refcount++;
- d = tgt->to_free;
+ gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
acc_dev->openacc.data_environ = tgt;
-
gomp_mutex_unlock (&acc_dev->lock);
+
+ d = tgt->to_free;
}
if (profiling_setup_p)
@@ -753,11 +815,6 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
n->dynamic_refcount = 0;
n->attach_count = NULL;
}
- if (n->refcount < n->dynamic_refcount)
- {
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
- }
if (f & FLAG_FINALIZE)
{
@@ -772,21 +829,6 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
if (n->refcount == 0)
{
- if (n->tgt->refcount == 2)
- {
- struct target_mem_desc *tp, *t;
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
- break;
- }
- }
-
if (f & FLAG_COPYOUT)
{
goacc_aq aq = get_goacc_asyncqueue (async);
@@ -968,38 +1010,16 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
- if (acc_is_present (*hostaddrs, *sizes))
- {
- splay_tree_key n;
- gomp_mutex_lock (&acc_dev->lock);
- n = lookup_host (acc_dev, *hostaddrs, *sizes);
- gomp_mutex_unlock (&acc_dev->lock);
-
- tgt = n->tgt;
- for (size_t i = 0; i < tgt->list_count; i++)
- if (tgt->list[i].key == n)
- {
- for (size_t j = 0; j < mapnum; j++)
- if (i + j < tgt->list_count && tgt->list[i + j].key)
- {
- tgt->list[i + j].key->refcount++;
- tgt->list[i + j].key->dynamic_refcount++;
- }
- return;
- }
- /* Should not reach here. */
- gomp_fatal ("Dynamic refcount incrementing failed for pointer/pset");
- }
-
gomp_debug (0, " %s: prepare mappings\n", __FUNCTION__);
goacc_aq aq = get_goacc_asyncqueue (async);
tgt = gomp_map_vars_async (acc_dev, aq, mapnum, hostaddrs,
- NULL, sizes, kinds, true, GOMP_MAP_VARS_OPENACC);
+ NULL, sizes, kinds, true,
+ GOMP_MAP_VARS_OPENACC_ENTER_DATA);
gomp_debug (0, " %s: mappings prepared\n", __FUNCTION__);
- /* Initialize dynamic refcount. */
- tgt->list[0].key->dynamic_refcount = 1;
- tgt->list[0].key->attach_count = NULL;
+ for (size_t i = 0; i < tgt->list_count; i++)
+ if (tgt->list[i].key)
+ tgt->list[i].key->dynamic_refcount++;
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
@@ -1008,96 +1028,83 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
}
void
-gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
- int finalize, int mapnum)
+gomp_acc_remove_pointer (void **hostaddrs, size_t *sizes, unsigned short *kinds,
+ int async, void *detach_from, bool finalize,
+ int mapnum)
{
struct goacc_thread *thr = goacc_thread ();
struct gomp_device_descr *acc_dev = thr->dev;
+ struct splay_tree_key_s cur_node;
splay_tree_key n;
- struct target_mem_desc *t;
- int minrefs = (mapnum == 1) ? 2 : 3;
-
- if (!acc_is_present (h, s))
- return;
gomp_mutex_lock (&acc_dev->lock);
- n = lookup_host (acc_dev, h, 1);
-
- if (!n)
+ if (detach_from)
{
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("%p is not a mapped block", (void *)h);
+ splay_tree_key n2 = lookup_host (acc_dev, detach_from, 1);
+ goacc_aq aq = get_goacc_asyncqueue (async);
+ gomp_detach_pointer (acc_dev, aq, n2, (uintptr_t) detach_from, finalize,
+ NULL);
}
- gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
-
- t = n->tgt;
+ gomp_acc_data_env_remove (acc_dev, &acc_dev->openacc.data_environ, hostaddrs,
+ mapnum);
- if (n->refcount < n->dynamic_refcount)
+ for (int i = 0; i < mapnum; i++)
{
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_fatal ("Dynamic reference counting assert fail\n");
- }
-
- if (finalize)
- {
- n->refcount -= n->dynamic_refcount;
- n->dynamic_refcount = 0;
- }
- else if (n->dynamic_refcount)
- {
- n->dynamic_refcount--;
- n->refcount--;
- }
+ int kind = kinds[i] & 0xff;
+ bool copyfrom = false;
- gomp_mutex_unlock (&acc_dev->lock);
-
- if (n->refcount == 0)
- {
- if (t->refcount == minrefs)
- {
- /* This is the last reference, so pull the descriptor off the
- chain. This prevents gomp_unmap_vars via gomp_unmap_tgt from
- freeing the device memory. */
- struct target_mem_desc *tp;
- for (tp = NULL, t = acc_dev->openacc.data_environ; t != NULL;
- tp = t, t = t->prev)
+ switch (kind)
+ {
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALWAYS_FROM:
+ copyfrom = true;
+ /* Fallthrough. */
+ case GOMP_MAP_TO_PSET:
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_DELETE:
+ case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start
+ + ((kind == GOMP_MAP_DETACH
+ || kind == GOMP_MAP_FORCE_DETACH
+ || kind == GOMP_MAP_POINTER)
+ ? sizeof (void *) : sizes[i]);
+ n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+ if (n == NULL)
+ continue;
+ if (finalize)
{
- if (n->tgt == t)
- {
- if (tp)
- tp->prev = t->prev;
- else
- acc_dev->openacc.data_environ = t->prev;
- break;
- }
+ n->refcount -= n->dynamic_refcount;
+ n->dynamic_refcount = 0;
}
- }
-
- /* Set refcount to 1 to allow gomp_unmap_vars to unmap it. */
- n->refcount = 1;
- t->refcount = minrefs;
- for (size_t i = 0; i < t->list_count; i++)
- if (t->list[i].key == n)
- {
- t->list[i].copy_from = force_copyfrom ? 1 : 0;
- break;
- }
-
- /* If running synchronously, unmap immediately. */
- if (async < acc_async_noval)
- gomp_unmap_vars (t, true, finalize);
- else
- {
- goacc_aq aq = get_goacc_asyncqueue (async);
- goacc_async_copyout_unmap_vars (t, aq, finalize);
+ else if (n->refcount > 0 && n->refcount != REFCOUNT_INFINITY)
+ {
+ n->refcount--;
+ n->dynamic_refcount--;
+ }
+ if (copyfrom)
+ gomp_copy_dev2host (acc_dev, NULL, (void *) cur_node.host_start,
+ (void *) (n->tgt->tgt_start + n->tgt_offset
+ + cur_node.host_start
+ - n->host_start),
+ cur_node.host_end - cur_node.host_start);
+ if (n->refcount == 0)
+ gomp_remove_var (acc_dev, n);
+ break;
+ default:
+ gomp_mutex_unlock (&acc_dev->lock);
+ gomp_fatal ("gomp_acc_remove_pointer unhandled kind 0x%.2x",
+ kind);
}
}
- gomp_mutex_unlock (&acc_dev->lock);
- gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
+ gomp_mutex_unlock (&acc_dev->lock);
}
diff --git a/libgomp/oacc-parallel.c b/libgomp/oacc-parallel.c
index f6c9114..8a3c65b 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -58,8 +58,12 @@ find_pointer (int pos, size_t mapnum, unsigned short *kinds)
case GOMP_MAP_FORCE_TO:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_TOFROM:
+ case GOMP_MAP_FORCE_TOFROM:
case GOMP_MAP_ALLOC:
case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DECLARE_ALLOCATE:
+ case GOMP_MAP_DECLARE_DEALLOCATE:
{
unsigned char kind1 = kinds[pos + 1] & 0xff;
if (kind1 == GOMP_MAP_POINTER
@@ -392,7 +396,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
&api_info);
}
/* If running synchronously, unmap immediately. */
- gomp_unmap_vars (tgt, true, false);
+ gomp_unmap_vars (tgt, true);
if (profiling_dispatch_p)
{
prof_info.event_type = acc_ev_exit_data_end;
@@ -410,7 +414,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
else
acc_dev->openacc.async.exec_func (tgt_fn, mapnum, hostaddrs,
devaddrs, dims, tgt, aq);
- goacc_async_copyout_unmap_vars (tgt, aq, false);
+ goacc_async_copyout_unmap_vars (tgt, aq);
}
out:
@@ -647,7 +651,7 @@ GOACC_data_end (void)
gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
thr->mapped_data = tgt->prev;
- gomp_unmap_vars (tgt, true, false);
+ gomp_unmap_vars (tgt, true);
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
if (profiling_dispatch_p)
@@ -845,18 +849,39 @@ GOACC_enter_exit_data (int device, size_t mapnum,
int elems = sizes[i];
struct splay_tree_key_s k;
splay_tree_key str;
- k.host_start = (uintptr_t) hostaddrs[i];
- k.host_end = k.host_start + 1;
+ uintptr_t elems_lo = (uintptr_t) hostaddrs[i + 1];
+ uintptr_t elems_hi = (uintptr_t) hostaddrs[i + elems]
+ + sizes[i + elems];
+ k.host_start = elems_lo;
+ k.host_end = elems_hi;
gomp_mutex_lock (&acc_dev->lock);
str = splay_tree_lookup (&acc_dev->mem_map, &k);
gomp_mutex_unlock (&acc_dev->lock);
- /* We increment the dynamic reference count for the struct
- itself by the number of struct elements that we
- mapped. */
- if (str->refcount != REFCOUNT_INFINITY)
+ if (str == NULL)
{
- str->refcount += elems;
- str->dynamic_refcount += elems;
+ size_t mapsize = elems_hi - elems_lo;
+ goacc_aq aq = get_goacc_asyncqueue (async);
+ struct target_mem_desc *tgt;
+ unsigned short thiskind = GOMP_MAP_ALLOC;
+ int j;
+ for (j = 0; j < elems; j++)
+ if ((kinds[i + j] & 0xff) != GOMP_MAP_ALLOC)
+ {
+ thiskind = GOMP_MAP_TO;
+ break;
+ }
+ tgt = gomp_map_vars_async (acc_dev, aq, 1,
+ &hostaddrs[i + 1], NULL, &mapsize, &thiskind,
+ true, GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+
+ for (j = 0; j < tgt->list_count; j++)
+ if (tgt->list[j].key)
+ tgt->list[j].key->dynamic_refcount++;
+
+ gomp_mutex_lock (&acc_dev->lock);
+ tgt->prev = acc_dev->openacc.data_environ;
+ acc_dev->openacc.data_environ = tgt;
+ gomp_mutex_unlock (&acc_dev->lock);
}
i += elems;
}
@@ -962,18 +987,17 @@ GOACC_enter_exit_data (int device, size_t mapnum,
int elems = sizes[i];
struct splay_tree_key_s k;
splay_tree_key str;
- k.host_start = (uintptr_t) hostaddrs[i];
- k.host_end = k.host_start + 1;
+ uintptr_t elems_lo = (uintptr_t) hostaddrs[i + 1];
+ uintptr_t elems_hi = (uintptr_t) hostaddrs[i + elems]
+ + sizes[i + elems];
+ k.host_start = elems_lo;
+ k.host_end = elems_hi;
gomp_mutex_lock (&acc_dev->lock);
str = splay_tree_lookup (&acc_dev->mem_map, &k);
gomp_mutex_unlock (&acc_dev->lock);
- /* Decrement dynamic reference count for the struct by the
- number of elements that we are unmapping. */
- if (str->dynamic_refcount >= elems)
- {
- str->dynamic_refcount -= elems;
- str->refcount -= elems;
- }
+ if (str == NULL)
+ gomp_fatal ("[%p,%ld] is not mapped", (void *) elems_lo,
+ (unsigned long) (elems_hi - elems_lo));
i += elems;
}
break;
@@ -989,10 +1013,14 @@ GOACC_enter_exit_data (int device, size_t mapnum,
&sizes[i], &kinds[i]);
else
{
- bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
- || kind == GOMP_MAP_FROM);
- gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom,
- async, finalize, pointer);
+ unsigned short ptrkind = kinds[i + pointer - 1] & 0xff;
+ bool detach = (ptrkind == GOMP_MAP_DETACH
+ || ptrkind == GOMP_MAP_FORCE_DETACH);
+ void *detach_from = detach ? hostaddrs[i + pointer - 1]
+ : NULL;
+ gomp_acc_remove_pointer (&hostaddrs[i], &sizes[i], &kinds[i],
+ async, detach_from, finalize,
+ pointer);
/* See the above comment. */
}
i += pointer - 1;
diff --git a/libgomp/target.c b/libgomp/target.c
index da51291..bb5e1e9 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -629,7 +629,10 @@ gomp_detach_pointer (struct gomp_device_descr *devicep,
idx = (detach_from - n->host_start) / sizeof (void *);
if (!n->attach_count)
- gomp_fatal ("no attachment counters for struct");
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("no attachment counters for struct");
+ }
if (finalize)
n->attach_count[idx] = 1;
@@ -1013,7 +1016,8 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
continue;
}
cur_node.host_start = (uintptr_t) hostaddrs[i];
- if (!GOMP_MAP_POINTER_P (kind & typemask))
+ if (!GOMP_MAP_POINTER_P (kind & typemask)
+ && (kind & typemask) != GOMP_MAP_ATTACH)
cur_node.host_end = cur_node.host_start + sizes[i];
else
cur_node.host_end = cur_node.host_start + sizeof (void *);
@@ -1281,7 +1285,9 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
tgt->list[i].length = n->host_end - n->host_start;
tgt->list[i].copy_from = false;
tgt->list[i].always_copy_from = false;
- tgt->list[i].do_detach = true;
+ tgt->list[i].do_detach
+ = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ n->refcount++;
}
else
{
@@ -1622,6 +1628,8 @@ gomp_unmap_tgt (struct target_mem_desc *tgt)
if (tgt->tgt_end)
gomp_free_device_memory (tgt->device_descr, tgt->to_free);
+ gomp_acc_data_env_remove_tgt (&tgt->device_descr->openacc.data_environ, tgt);
+
free (tgt->array);
free (tgt);
}
@@ -1650,17 +1658,18 @@ gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
has been done already. */
attribute_hidden void
-gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, bool finalize)
+gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
{
- gomp_unmap_vars_async (tgt, do_copyfrom, NULL, finalize);
+ gomp_unmap_vars_async (tgt, do_copyfrom, NULL);
}
attribute_hidden void
gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
- struct goacc_asyncqueue *aq, bool finalize)
+ struct goacc_asyncqueue *aq)
{
struct gomp_device_descr *devicep = tgt->device_descr;
+
if (tgt->list_count == 0)
{
free (tgt);
@@ -1685,15 +1694,15 @@ gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
if (k != NULL && tgt->list[i].do_detach)
gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
- + tgt->list[i].offset, finalize,
- NULL);
+ + tgt->list[i].offset,
+ k->refcount == 1, NULL);
}
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
- if (k == NULL || tgt->list[i].do_detach)
+ if (k == NULL)
continue;
bool do_unmap = false;
@@ -2314,7 +2323,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
GOMP_MAP_VARS_TARGET);
devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start,
NULL);
- gomp_unmap_vars (tgt_vars, true, false);
+ gomp_unmap_vars (tgt_vars, true);
}
/* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
@@ -2458,7 +2467,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
args);
if (tgt_vars)
- gomp_unmap_vars (tgt_vars, true, false);
+ gomp_unmap_vars (tgt_vars, true);
}
/* Host fallback for GOMP_target_data{,_ext} routines. */
@@ -2527,7 +2536,7 @@ GOMP_target_end_data (void)
{
struct target_mem_desc *tgt = icv->target_data;
icv->target_data = tgt->prev;
- gomp_unmap_vars (tgt, true, false);
+ gomp_unmap_vars (tgt, true);
}
}
@@ -2762,7 +2771,7 @@ gomp_target_task_fn (void *data)
if (ttask->state == GOMP_TARGET_TASK_FINISHED)
{
if (ttask->tgt)
- gomp_unmap_vars (ttask->tgt, true, false);
+ gomp_unmap_vars (ttask->tgt, true);
return false;
}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
index 6a52f74..6bdcfe7 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-2.c
@@ -182,13 +182,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
index 71365e8..b403a5c 100644
--- a/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/context-4.c
@@ -176,13 +176,13 @@ main (int argc, char **argv)
exit (EXIT_FAILURE);
}
+ acc_delete (&h_X[0], N * sizeof (float));
+ acc_delete (&h_Y1[0], N * sizeof (float));
+
free (h_X);
free (h_Y1);
free (h_Y2);
- acc_free (d_X);
- acc_free (d_Y);
-
context_check (pctx);
s = cublasDestroy (h);
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
new file mode 100644
index 0000000..81c1c5e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c
@@ -0,0 +1,59 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+ int a;
+ int **b;
+};
+
+int
+main ()
+{
+ int n = 100, i, j, k;
+ struct dc v = { .a = 3 };
+
+ v.b = (int **) malloc (sizeof (int *) * n);
+ for (i = 0; i < n; i++)
+ v.b[i] = (int *) malloc (sizeof (int) * n);
+
+ for (k = 0; k < 16; k++)
+ {
+#pragma acc data copy(v)
+ {
+#pragma acc data copy(v.b[:n])
+ {
+ for (i = 0; i < n; i++)
+ {
+ acc_copyin (v.b[i], sizeof (int) * n);
+ acc_attach ((void **) &v.b[i]);
+ }
+
+#pragma acc parallel loop
+ for (i = 0; i < n; i++)
+ for (j = 0; j < n; j++)
+ v.b[i][j] = v.a + i + j;
+
+ for (i = 0; i < n; i++)
+ {
+ acc_detach ((void **) &v.b[i]);
+ acc_copyout (v.b[i], sizeof (int) * n);
+ }
+ }
+ }
+
+ for (i = 0; i < n; i++)
+ for (j = 0; j < n; j++)
+ assert (v.b[i][j] == v.a + i + j);
+
+ assert (!acc_is_present (&v, sizeof (v)));
+ assert (!acc_is_present (v.b, sizeof (int *) * n));
+ for (i = 0; i < n; i++)
+ assert (!acc_is_present (v.b[i], sizeof (int) * n));
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
new file mode 100644
index 0000000..3a970a0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
@@ -0,0 +1,42 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+ int a;
+ int *b;
+};
+
+int
+main ()
+{
+ int n = 100, i, j, k;
+ struct dc v = { .a = 3 };
+
+ v.b = (int *) malloc (sizeof (int) * n);
+
+ for (k = 0; k < 16; k++)
+ {
+#pragma acc enter data copyin(v.a, v.b[0:n])
+
+#pragma acc enter data pcopyin(v.b[0:n])
+
+#pragma acc parallel loop attach(v.b)
+ for (i = 0; i < n; i++)
+ v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n])
+#pragma acc exit data delete(v) finalize
+
+ for (i = 0; i < n; i++)
+ assert (v.b[i] == v.a + i);
+
+ assert (!acc_is_present (&v, sizeof (v)));
+ assert (!acc_is_present (v.b, sizeof (int *) * n));
+ }
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
new file mode 100644
index 0000000..54f553b
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
@@ -0,0 +1,53 @@
+/* { dg-do run { target { ! openacc_host_selected } } } */
+
+#include <stdlib.h>
+#include <assert.h>
+#include <openacc.h>
+
+struct dc
+{
+ int a;
+ int *b;
+ int *c;
+ int *d;
+};
+
+int
+main ()
+{
+ int n = 100, i, j, k;
+ struct dc v = { .a = 3 };
+
+ v.b = (int *) malloc (sizeof (int) * n);
+ v.c = (int *) malloc (sizeof (int) * n);
+ v.d = (int *) malloc (sizeof (int) * n);
+
+#pragma acc enter data copyin(v)
+
+ for (k = 0; k < 16; k++)
+ {
+#pragma acc enter data copyin(v.a, v.b[:n], v.c[:n], v.d[:n])
+
+#pragma acc parallel loop
+ for (i = 0; i < n; i++)
+ v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n])
+#pragma acc exit data copyout(v.c[:n])
+#pragma acc exit data copyout(v.d[:n])
+
+ for (i = 0; i < n; i++)
+ assert (v.b[i] == v.a + i);
+
+ assert (acc_is_present (&v, sizeof (v)));
+ assert (!acc_is_present (v.b, sizeof (int *) * n));
+ assert (!acc_is_present (v.c, sizeof (int *) * n));
+ assert (!acc_is_present (v.d, sizeof (int *) * n));
+ }
+
+#pragma acc exit data copyout(v)
+
+ assert (!acc_is_present (&v, sizeof (v)));
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90 b/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90
index db80413..a58e465 100644
--- a/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90
+++ b/libgomp/testsuite/libgomp.oacc-fortran/data-2.f90
@@ -92,10 +92,6 @@ program test
if (acc_is_present (c) .eqv. .TRUE.) call abort
- !$acc exit data delete (c(0:N))
-
- if (acc_is_present (c) .eqv. .TRUE.) call abort
-
do i = 1, N
if (c(i) .ne. 3.0) call abort
end do
@@ -113,11 +109,6 @@ program test
if (acc_is_present (c) .eqv. .TRUE.) call abort
if (acc_is_present (d) .eqv. .TRUE.) call abort
- !$acc exit data delete (c(0:N), d(0:N))
-
- if (acc_is_present (c) .eqv. .TRUE.) call abort
- if (acc_is_present (d) .eqv. .TRUE.) call abort
-
do i = 1, N
if (c(i) .ne. 5.0) call abort
if (d(i) .ne. 9.0) call abort
@@ -177,8 +168,8 @@ program test
!$acc exit data delete (c(0:N), d(0:N))
- !if (acc_is_present (c) .eqv. .TRUE.) call abort
- !if (acc_is_present (d) .eqv. .TRUE.) call abort
+ if (acc_is_present (c) .eqv. .FALSE.) call abort
+ if (acc_is_present (d) .eqv. .FALSE.) call abort
!$acc exit data delete (c(0:N), d(0:N))
@@ -190,12 +181,7 @@ program test
if (acc_is_present (c) .eqv. .FALSE.) call abort
if (acc_is_present (d) .eqv. .TRUE.) call abort
- !$acc exit data delete (c(0:N), d(0:N))
-
- if (acc_is_present (c) .eqv. .TRUE.) call abort
- if (acc_is_present (d) .eqv. .TRUE.) call abort
-
- !$acc exit data delete (c(0:N), d(0:N))
+ !$acc exit data delete (c(0:N))
if (acc_is_present (c) .eqv. .TRUE.) call abort
if (acc_is_present (d) .eqv. .TRUE.) call abort
^ permalink raw reply [flat|nested] 8+ messages in thread