[prev in list] [next in list] [prev in thread] [next in thread]
List: gcc-fortran
Subject: Re: [PATCH] OpenACC 2.6 manual deep copy support (attach/detach)
From: Julian Brown <julian () codesourcery ! com>
Date: 2019-11-26 2:45:02
Message-ID: 20191126024502.10808ed5 () squid ! athome
[Download RAW message or body]
On Mon, 25 Nov 2019 11:27:35 +0100
Tobias Burnus <tobias@codesourcery.com> wrote:
> Hi Julian,
>
> On 11/23/19 12:42 AM, Julian Brown wrote:
> >
> > gcc/fortran/
> > * trans-expr.c
> > (gfc_auto_dereference_var): New function, broken out
> > of... […]
> > +gfc_auto_dereference_var (location_t loc, gfc_symbol *sym, tree
> > var,
> > + bool descriptor_only_p, bool
> > is_classarray) […]
> > + var = build_fold_indirect_ref_loc (input_location, var);
> As your patch takes a location as argument – it also makes sense to
> use that location. (Alternatively, one could remove the argument as
> both callers explicitly pass 'input_location' as argument.)
>
> One could do either way, but the current variant does not make sense
> – and, additionally, the current variant causes a compile-time
> warning.
>
> See also my OG9 commit 500483e6ced44e2e0fea6a37e4f8c267ebaf826a where
> do s/input_location/loc/g in that function.
Thanks -- I'd missed that. This version of the patch uses your
alternative suggestion -- i.e., removing the location_t parameter from
gfc_auto_dereference_var. That gets rid of the warning.
The rest of the patch is as the last-posted version.
Cheers,
Julian
["attach-detach-fsf-20191126.diff" (text/x-patch)]
commit e647a5c0a5be59fb40f9ed1f9a1085fc62fabee2
Author: Julian Brown <julian@codesourcery.com>
Date: Mon Sep 30 14:14:11 2019 -0700
OpenACC 2.6 manual deep copy support (attach/detach)
gcc/c-family/
* c-common.h (c_omp_map_clause_name): Add prototype.
* c-omp.c (c_omp_map_clause_name): New function.
* c-pragma.h (pragma_omp_clause): Add PRAGMA_OACC_CLAUSE_ATTACH and
PRAGMA_OACC_CLAUSE_DETACH.
gcc/c/
* c-parser.c (c_parser_omp_clause_name): Add parsing of attach and
detach clauses.
(c_parser_omp_variable_list): Add ALLOW_DEREF optional parameter.
Allow deref (->) in variable lists if true.
(c_parser_omp_var_list_parens): Add ALLOW_DEREF optional parameter.
Pass to c_parser_omp_variable_list.
(c_parser_oacc_data_clause): Support attach and detach clauses. Update
call to c_parser_omp_variable_list.
(c_parser_oacc_all_clauses): Support attach and detach clauses.
(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
OACC_SERIAL_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
* c-typeck.c (handle_omp_array_sections_1): Reject subarrays for attach
and detach. Support deref.
(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
GOMP_MAP_ALWAYS_POINTER for OpenACC.
(c_oacc_check_attachments): New function.
(c_finish_omp_clauses): Check attach/detach arguments for being
pointers using above. Support deref.
gcc/cp/
* parser.c (cp_parser_omp_clause_name): Support attach and detach
clauses.
(cp_parser_omp_var_list_no_open): Add ALLOW_DEREF optional parameter.
Parse deref if true.
(cp_parser_omp_var_list): Add ALLOW_DEREF optional parameter. Pass to
cp_parser_omp_var_list_no_open.
(cp_parser_oacc_data_clause): Support attach and detach clauses.
Update call to cp_parser_omp_var_list_no_open.
(cp_parser_oacc_all_clauses): Support attach and detach.
(OACC_DATA_CLAUSE_MASK, OACC_ENTER_DATA_CLAUSE_MASK,
OACC_KERNELS_CLAUSE_MASK, OACC_PARALLEL_CLAUSE_MASK,
OACC_SERIAL_CLAUSE_MASK): Add
PRAGMA_OACC_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSE_MASK): Add PRAGMA_OACC_CLAUSE_DETACH.
* semantics.c (handle_omp_array_sections_1): Reject subarrays for
attach and detach.
(handle_omp_array_sections): Use GOMP_MAP_ATTACH_DETACH instead of
GOMP_MAP_ALWAYS_POINTER for OpenACC.
(cp_oacc_check_attachments): New function.
(finish_omp_clauses): Use above function. Allow structure fields and
class members to appear in OpenACC data clauses. Support
GOMP_MAP_ATTACH_DETACH. Support deref.
gcc/fortran/
* gfortran.h (gfc_omp_map_op): Add OMP_MAP_ATTACH, OMP_MAP_DETACH.
* openmp.c (gfc_match_omp_variable_list): Add allow_derived parameter.
Parse derived-type member accesses if true.
(omp_mask2): Add OMP_CLAUSE_ATTACH and OMP_CLAUSE_DETACH.
(gfc_match_omp_map_clause): Add allow_derived parameter. Pass to
gfc_match_omp_variable_list.
(gfc_match_omp_clauses): Support attach and detach. Support derived
types for appropriate OpenACC directives.
(OACC_PARALLEL_CLAUSES, OACC_SERIAL_CLAUSES, OACC_KERNELS_CLAUSES,
OACC_DATA_CLAUSES, OACC_ENTER_DATA_CLAUSES): Add OMP_CLAUSE_ATTACH.
(OACC_EXIT_DATA_CLAUSES): Add OMP_CLAUSE_DETACH.
(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-expr.c (gfc_conv_component_ref,
conv_parent_component_references): Make global.
(gfc_auto_dereference_var): New function, broken out of...
(gfc_conv_variable): ...here. Call above function.
* trans-openmp.c (gfc_omp_privatize_by_reference): Support component
refs.
(gfc_trans_omp_array_section): New function, broken out of...
(gfc_trans_omp_clauses): ...here. Support component refs/derived
types, attach and detach clauses.
* trans.h (gfc_conv_component_ref, conv_parent_component_references,
gfc_auto_dereference_var): Add prototypes.
gcc/
* gimplify.c (gimplify_omp_var_data): Add GOVD_MAP_HAS_ATTACHMENTS.
(insert_struct_comp_map): Support derived-type member mappings
for arrays with descriptors which use GOMP_MAP_TO_PSET. Support
GOMP_MAP_ATTACH_DETACH.
(gimplify_scan_omp_clauses): Tidy up OACC_ENTER_DATA/OACC_EXIT_DATA
mappings. Handle attach/detach clauses and component references.
(gimplify_adjust_omp_clauses_1): Skip adjustments for explicit
attach/detach clauses.
(gimplify_omp_target_update): Handle finalize for detach.
* omp-low.c (lower_omp_target): Support GOMP_MAP_ATTACH,
GOMP_MAP_DETACH, GOMP_MAP_FORCE_DETACH.
* tree-pretty-print.c (dump_omp_clause): Likewise, plus
GOMP_MAP_ATTACH_DETACH.
include/
* gomp-constants.h (GOMP_MAP_FLAG_SPECIAL_4, GOMP_MAP_DEEP_COPY):
Define.
(gomp_map_kind): Add GOMP_MAP_ATTACH, GOMP_MAP_DETACH,
GOMP_MAP_FORCE_DETACH.
gcc/testsuite/
* c-c++-common/goacc/deep-copy-arrayofstruct.c: New test.
* c-c++-common/goacc/mdc-1.c: New test.
* c-c++-common/goacc/mdc-2.c: New test.
* gcc.dg/goacc/mdc.C: New test.
* gfortran.dg/goacc/derived-types.f90: New test.
* gfortran.dg/goacc/derived-types-2.f90: New test.
* gfortran.dg/goacc/data-clauses.f95: Adjust for expected errors.
* gfortran.dg/goacc/enter-exit-data.f95: Likewise.
libgomp/
* libgomp.h (struct target_var_desc): Add do_detach flag.
(struct splay_tree_aux): New.
(struct splay_tree_key_s): Replace link_key field with aux pointer.
(gomp_attach_pointer, gomp_detach_pointer): Add prototypes.
* 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-init.c (acc_shutdown_1): Free aux block if present.
* oacc-mem.c (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_group_last): Add SIZES parameter. Support
struct components. Tidy up and add some new checks.
(goacc_enter_data_internal): Update call to find_group_last.
(goacc_exit_data_internal): Support detach operations and
GOMP_MAP_STRUCT.
(GOACC_enter_exit_data): Handle initial GOMP_MAP_STRUCT or
GOMP_MAP_FORCE_PRESENT in finalization detection code. Handle
attach/detach in enter/exit data detection code.
* openacc.h (acc_attach, acc_attach_async, acc_detach,
(acc_detach_async, acc_detach_finalize, acc_detach_finalize_async): Add
prototypes.
* target.c (dump_tgt): Support aux field.
(gomp_map_vars_existing): Initialise do_detach field of tgt_var_desc.
(gomp_attach_pointer, gomp_detach_pointer): New functions.
(gomp_map_vars_internal): Support attach and detach.
(gomp_remove_var_internal): Free aux block and attachment counts if
present.
(gomp_unmap_vars_internal): Support detach.
(gomp_load_image_to_device): Zero-initialise aux field instead of
link_key field.
(gomp_exit_data): Handle link key in aux field. Free aux field when
appropriate.
(omp_target_associate_ptr): Zero-initialize aux field instead of
link_key.
* 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-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-c-c++-common/deep-copy-9.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c: New test.
* testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c: New test.
* testsuite/libgomp.oacc-c++/deep-copy-12.C: New test.
* testsuite/libgomp.oacc-c++/deep-copy-13.C: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-1.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-2.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-3.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-4.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-5.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-6.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-7.f90: New test.
* testsuite/libgomp.oacc-fortran/deep-copy-8.f90: New test.
* testsuite/libgomp.oacc-fortran/derived-type-1.f90: New test.
* testsuite/libgomp.oacc-fortran/derivedtype-1.f95: New test.
* testsuite/libgomp.oacc-fortran/derivedtype-2.f95: New test.
* testsuite/libgomp.oacc-fortran/multidim-slice.f95: New test.
* testsuite/libgomp.oacc-fortran/update-2.f90: New test.
diff --git a/gcc/c-family/c-common.h b/gcc/c-family/c-common.h
index f3478d39beb..673e07a6177 100644
--- a/gcc/c-family/c-common.h
+++ b/gcc/c-family/c-common.h
@@ -1204,6 +1204,7 @@ extern bool c_omp_predefined_variable (tree);
extern enum omp_clause_default_kind c_omp_predetermined_sharing (tree);
extern tree c_omp_check_context_selector (location_t, tree);
extern void c_omp_mark_declare_variant (location_t, tree, tree);
+extern const char *c_omp_map_clause_name (tree, bool);
/* Return next tree in the chain for chain_next walking of tree nodes. */
static inline tree
diff --git a/gcc/c-family/c-omp.c b/gcc/c-family/c-omp.c
index a4be2d68b9a..04f2c0b0682 100644
--- a/gcc/c-family/c-omp.c
+++ b/gcc/c-family/c-omp.c
@@ -2259,3 +2259,36 @@ c_omp_mark_declare_variant (location_t loc, tree variant, tree \
construct)
error_at (loc, "%qD used as a variant with incompatible %<construct%> "
"selector sets", variant);
}
+
+/* For OpenACC, the OMP_CLAUSE_MAP_KIND of an OMP_CLAUSE_MAP is used internally
+ to distinguish clauses as seen by the user. Return the "friendly" clause
+ name for error messages etc., where possible. See also
+ c/c-parser.c:c_parser_oacc_data_clause and
+ cp/parser.c:cp_parser_oacc_data_clause. */
+
+const char *
+c_omp_map_clause_name (tree clause, bool oacc)
+{
+ if (oacc && OMP_CLAUSE_CODE (clause) == OMP_CLAUSE_MAP)
+ switch (OMP_CLAUSE_MAP_KIND (clause))
+ {
+ case GOMP_MAP_FORCE_ALLOC:
+ case GOMP_MAP_ALLOC: return "create";
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_TO: return "copyin";
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_FROM: return "copyout";
+ case GOMP_MAP_FORCE_TOFROM:
+ case GOMP_MAP_TOFROM: return "copy";
+ case GOMP_MAP_RELEASE: return "delete";
+ case GOMP_MAP_FORCE_PRESENT: return "present";
+ case GOMP_MAP_ATTACH: return "attach";
+ case GOMP_MAP_FORCE_DETACH:
+ case GOMP_MAP_DETACH: return "detach";
+ case GOMP_MAP_DEVICE_RESIDENT: return "device_resident";
+ case GOMP_MAP_LINK: return "link";
+ case GOMP_MAP_FORCE_DEVICEPTR: return "deviceptr";
+ default: break;
+ }
+ return omp_clause_code_name[OMP_CLAUSE_CODE (clause)];
+}
diff --git a/gcc/c-family/c-pragma.h b/gcc/c-family/c-pragma.h
index bfe681bb430..8a04e611bc7 100644
--- a/gcc/c-family/c-pragma.h
+++ b/gcc/c-family/c-pragma.h
@@ -143,11 +143,13 @@ enum pragma_omp_clause {
/* Clauses for OpenACC. */
PRAGMA_OACC_CLAUSE_ASYNC,
+ PRAGMA_OACC_CLAUSE_ATTACH,
PRAGMA_OACC_CLAUSE_AUTO,
PRAGMA_OACC_CLAUSE_COPY,
PRAGMA_OACC_CLAUSE_COPYOUT,
PRAGMA_OACC_CLAUSE_CREATE,
PRAGMA_OACC_CLAUSE_DELETE,
+ PRAGMA_OACC_CLAUSE_DETACH,
PRAGMA_OACC_CLAUSE_DEVICEPTR,
PRAGMA_OACC_CLAUSE_DEVICE_RESIDENT,
PRAGMA_OACC_CLAUSE_FINALIZE,
diff --git a/gcc/c/c-parser.c b/gcc/c/c-parser.c
index 03194b438f2..0ef9e640e9c 100644
--- a/gcc/c/c-parser.c
+++ b/gcc/c/c-parser.c
@@ -12435,6 +12435,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OMP_CLAUSE_ALIGNED;
else if (!strcmp ("async", p))
result = PRAGMA_OACC_CLAUSE_ASYNC;
+ else if (!strcmp ("attach", p))
+ result = PRAGMA_OACC_CLAUSE_ATTACH;
break;
case 'b':
if (!strcmp ("bind", p))
@@ -12461,6 +12463,8 @@ c_parser_omp_clause_name (c_parser *parser)
result = PRAGMA_OACC_CLAUSE_DELETE;
else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
+ else if (!strcmp ("detach", p))
+ result = PRAGMA_OACC_CLAUSE_DETACH;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
@@ -12704,12 +12708,16 @@ c_parser_oacc_wait_list (c_parser *parser, location_t \
clause_loc, tree list) If KIND is nonzero, CLAUSE_LOC is the location of the clause.
If KIND is zero, create a TREE_LIST with the decl in TREE_PURPOSE;
- return the list created. */
+ return the list created.
+
+ The optional ALLOW_DEREF argument is true if list items can use the deref
+ (->) operator. */
static tree
c_parser_omp_variable_list (c_parser *parser,
location_t clause_loc,
- enum omp_clause_code kind, tree list)
+ enum omp_clause_code kind, tree list,
+ bool allow_deref = false)
{
auto_vec<c_token> tokens;
unsigned int tokens_avail = 0;
@@ -12836,9 +12844,13 @@ 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)
+ || (allow_deref
+ && 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))
{
@@ -12960,11 +12972,12 @@ c_parser_omp_variable_list (c_parser *parser,
}
/* Similarly, but expect leading and trailing parenthesis. This is a very
- common case for OpenACC and OpenMP clauses. */
+ common case for OpenACC and OpenMP clauses. The optional ALLOW_DEREF
+ argument is true if list items can use the deref (->) operator. */
static tree
c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
- tree list)
+ tree list, bool allow_deref = false)
{
/* The clauses location. */
location_t loc = c_parser_peek_token (parser)->location;
@@ -12972,18 +12985,20 @@ c_parser_omp_var_list_parens (c_parser *parser, enum \
omp_clause_code kind, matching_parens parens;
if (parens.require_open (parser))
{
- list = c_parser_omp_variable_list (parser, loc, kind, list);
+ list = c_parser_omp_variable_list (parser, loc, kind, list, allow_deref);
parens.skip_until_found_close (parser);
}
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.0+:
+ attach ( variable-list )
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
+ detach ( variable-list )
present ( variable-list ) */
static tree
@@ -12993,6 +13008,9 @@ c_parser_oacc_data_clause (c_parser *parser, \
pragma_omp_clause c_kind, enum gomp_map_kind kind;
switch (c_kind)
{
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ kind = GOMP_MAP_ATTACH;
+ break;
case PRAGMA_OACC_CLAUSE_COPY:
kind = GOMP_MAP_TOFROM;
break;
@@ -13008,6 +13026,9 @@ c_parser_oacc_data_clause (c_parser *parser, \
pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_RELEASE;
break;
+ case PRAGMA_OACC_CLAUSE_DETACH:
+ kind = GOMP_MAP_DETACH;
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
@@ -13027,7 +13048,7 @@ c_parser_oacc_data_clause (c_parser *parser, \
pragma_omp_clause c_kind, gcc_unreachable ();
}
tree nl, c;
- nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list);
+ nl = c_parser_omp_var_list_parens (parser, OMP_CLAUSE_MAP, list, true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -15742,6 +15763,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask \
mask, 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_COLLAPSE:
clauses = c_parser_omp_clause_collapse (parser, clauses);
c_name = "collapse";
@@ -15770,6 +15795,10 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask \
mask, clauses = c_parser_omp_clause_default (parser, clauses, true);
c_name = "default";
break;
+ case PRAGMA_OACC_CLAUSE_DETACH:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "detach";
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@@ -16280,7 +16309,8 @@ c_parser_oacc_cache (location_t loc, c_parser *parser)
*/
#define OACC_DATA_CLAUSE_MASK \
- ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
@@ -16463,6 +16493,7 @@ c_parser_oacc_declare (c_parser *parser)
#define OACC_ENTER_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -16472,6 +16503,7 @@ c_parser_oacc_declare (c_parser *parser)
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -16611,6 +16643,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char \
*p_name,
#define OACC_KERNELS_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@@ -16626,6 +16659,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char \
*p_name,
#define OACC_PARALLEL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@@ -16644,6 +16678,7 @@ c_parser_oacc_loop (location_t loc, c_parser *parser, char \
*p_name,
#define OACC_SERIAL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
diff --git a/gcc/c/c-typeck.c b/gcc/c/c-typeck.c
index 5f74a3b28d9..5bbe3f8923b 100644
--- a/gcc/c/c-typeck.c
+++ b/gcc/c/c-typeck.c
@@ -12897,7 +12897,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> \
&types, return error_mark_node;
}
if (TREE_CODE (t) == COMPONENT_REF
- && ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM))
@@ -12918,6 +12917,15 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> \
&types, return error_mark_node;
}
t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+ {
+ if (maybe_ne (mem_ref_offset (t), 0))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "cannot dereference %qE in %qs clause", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else
+ t = TREE_OPERAND (t, 0);
+ }
}
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
@@ -13003,7 +13011,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> \
&types, length = fold_convert (sizetype, length);
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
-
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ {
+ if (length != integer_one_node)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "expected single pointer in %qs clause",
+ c_omp_map_clause_name (c, ort == C_ORT_ACC));
+ return error_mark_node;
+ }
+ }
if (length != NULL_TREE)
{
if (!integer_nonzerop (length))
@@ -13443,7 +13462,11 @@ handle_omp_array_sections (tree c, enum c_omp_region_type \
ort) if (ort != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else if (TREE_CODE (t) == COMPONENT_REF)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ {
+ gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, k);
+ }
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
if (OMP_CLAUSE_MAP_KIND (c2) != GOMP_MAP_FIRSTPRIVATE_POINTER
@@ -13680,6 +13703,35 @@ c_omp_finish_iterators (tree iter)
return ret;
}
+/* Ensure that pointers are used in OpenACC attach and detach clauses.
+ Return true if an error has been detected. */
+
+static bool
+c_oacc_check_attachments (tree c)
+{
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return false;
+
+ /* OpenACC attach / detach clauses must be pointers. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ {
+ tree t = OMP_CLAUSE_DECL (c);
+
+ while (TREE_CODE (t) == TREE_LIST)
+ t = TREE_CHAIN (t);
+
+ if (TREE_CODE (TREE_TYPE (t)) != POINTER_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
+ c_omp_map_clause_name (c, true));
+ return true;
+ }
+ }
+
+ return false;
+}
+
/* For all elements of CLAUSES, validate them against their constraints.
Remove any elements from the list that are invalid. */
@@ -14433,6 +14485,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type \
ort) }
}
}
+ if (c_oacc_check_attachments (c))
+ remove = true;
break;
}
if (t == error_mark_node)
@@ -14440,8 +14494,13 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type \
ort) remove = true;
break;
}
+ /* OpenACC attach / detach clauses must be pointers. */
+ if (c_oacc_check_attachments (c))
+ {
+ remove = true;
+ break;
+ }
if (TREE_CODE (t) == COMPONENT_REF
- && (ort & C_ORT_OMP)
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
if (DECL_BIT_FIELD (TREE_OPERAND (t, 1)))
@@ -14476,6 +14535,15 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type \
ort) break;
}
t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && TREE_CODE (t) == MEM_REF)
+ {
+ if (maybe_ne (mem_ref_offset (t), 0))
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "cannot dereference %qE in %qs clause", t,
+ omp_clause_code_name[OMP_CLAUSE_CODE (c)]);
+ else
+ t = TREE_OPERAND (t, 0);
+ }
}
if (remove)
break;
diff --git a/gcc/cp/parser.c b/gcc/cp/parser.c
index c473e7fd92f..ffe0d00c2a6 100644
--- a/gcc/cp/parser.c
+++ b/gcc/cp/parser.c
@@ -33016,6 +33016,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_ALIGNED;
else if (!strcmp ("async", p))
result = PRAGMA_OACC_CLAUSE_ASYNC;
+ else if (!strcmp ("attach", p))
+ result = PRAGMA_OACC_CLAUSE_ATTACH;
break;
case 'b':
if (!strcmp ("bind", p))
@@ -33040,6 +33042,8 @@ cp_parser_omp_clause_name (cp_parser *parser)
result = PRAGMA_OMP_CLAUSE_DEFAULTMAP;
else if (!strcmp ("depend", p))
result = PRAGMA_OMP_CLAUSE_DEPEND;
+ else if (!strcmp ("detach", p))
+ result = PRAGMA_OACC_CLAUSE_DETACH;
else if (!strcmp ("device", p))
result = PRAGMA_OMP_CLAUSE_DEVICE;
else if (!strcmp ("deviceptr", p))
@@ -33242,11 +33246,15 @@ check_no_duplicate_clause (tree clauses, enum \
omp_clause_code code,
COLON can be NULL if only closing parenthesis should end the list,
or pointer to bool which will receive false if the list is terminated
- by closing parenthesis or true if the list is terminated by colon. */
+ by closing parenthesis or true if the list is terminated by colon.
+
+ The optional ALLOW_DEREF argument is true if list items can use the deref
+ (->) operator. */
static tree
cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
- tree list, bool *colon)
+ tree list, bool *colon,
+ bool allow_deref = false)
{
cp_token *token;
bool saved_colon_corrects_to_scope_p = parser->colon_corrects_to_scope_p;
@@ -33327,15 +33335,20 @@ 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)
+ || (allow_deref
+ && 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);
}
@@ -33453,19 +33466,23 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum \
omp_clause_code kind, common case for omp clauses. */
static tree
-cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list)
+cp_parser_omp_var_list (cp_parser *parser, enum omp_clause_code kind, tree list,
+ bool allow_deref = false)
{
if (cp_parser_require (parser, CPP_OPEN_PAREN, RT_OPEN_PAREN))
- return cp_parser_omp_var_list_no_open (parser, kind, list, NULL);
+ return cp_parser_omp_var_list_no_open (parser, kind, list, NULL,
+ allow_deref);
return list;
}
-/* OpenACC 2.0:
+/* OpenACC 2.0+:
+ attach ( variable-list )
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
create ( variable-list )
delete ( variable-list )
+ detach ( variable-list )
present ( variable-list ) */
static tree
@@ -33475,6 +33492,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, \
pragma_omp_clause c_kind, enum gomp_map_kind kind;
switch (c_kind)
{
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ kind = GOMP_MAP_ATTACH;
+ break;
case PRAGMA_OACC_CLAUSE_COPY:
kind = GOMP_MAP_TOFROM;
break;
@@ -33490,6 +33510,9 @@ cp_parser_oacc_data_clause (cp_parser *parser, \
pragma_omp_clause c_kind, case PRAGMA_OACC_CLAUSE_DELETE:
kind = GOMP_MAP_RELEASE;
break;
+ case PRAGMA_OACC_CLAUSE_DETACH:
+ kind = GOMP_MAP_DETACH;
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
kind = GOMP_MAP_FORCE_TO;
break;
@@ -33509,7 +33532,7 @@ cp_parser_oacc_data_clause (cp_parser *parser, \
pragma_omp_clause c_kind, gcc_unreachable ();
}
tree nl, c;
- nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list);
+ nl = cp_parser_omp_var_list (parser, OMP_CLAUSE_MAP, list, true);
for (c = nl; c != list; c = OMP_CLAUSE_CHAIN (c))
OMP_CLAUSE_SET_MAP_KIND (c, kind);
@@ -35987,6 +36010,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, \
omp_clause_mask mask, clauses);
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_COLLAPSE:
clauses = cp_parser_omp_clause_collapse (parser, clauses, here);
c_name = "collapse";
@@ -36015,6 +36042,10 @@ cp_parser_oacc_all_clauses (cp_parser *parser, \
omp_clause_mask mask, clauses = cp_parser_omp_clause_default (parser, clauses, here, \
true); c_name = "default";
break;
+ case PRAGMA_OACC_CLAUSE_DETACH:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "detach";
+ break;
case PRAGMA_OACC_CLAUSE_DEVICE:
clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
c_name = "device";
@@ -39863,10 +39894,12 @@ cp_parser_oacc_cache (cp_parser *parser, cp_token \
*pragma_tok) structured-block */
#define OACC_DATA_CLAUSE_MASK \
- ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
+ ( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DEVICEPTR) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_PRESENT) )
@@ -40066,6 +40099,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token \
*pragma_tok)
#define OACC_ENTER_DATA_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_IF) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_CREATE) \
@@ -40076,6 +40110,7 @@ cp_parser_oacc_declare (cp_parser *parser, cp_token \
*pragma_tok) | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DELETE) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_DETACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_FINALIZE) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_WAIT) )
@@ -40183,6 +40218,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, \
char *p_name,
#define OACC_KERNELS_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@@ -40198,6 +40234,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, \
char *p_name,
#define OACC_PARALLEL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
@@ -40216,6 +40253,7 @@ cp_parser_oacc_loop (cp_parser *parser, cp_token *pragma_tok, \
char *p_name,
#define OACC_SERIAL_CLAUSE_MASK \
( (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ASYNC) \
+ | (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_ATTACH) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPY) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYIN) \
| (OMP_CLAUSE_MASK_1 << PRAGMA_OACC_CLAUSE_COPYOUT) \
diff --git a/gcc/cp/semantics.c b/gcc/cp/semantics.c
index 2adc9ef792b..5ac2e7266c6 100644
--- a/gcc/cp/semantics.c
+++ b/gcc/cp/semantics.c
@@ -4752,7 +4752,6 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
t = TREE_OPERAND (t, 0);
ret = t;
if (TREE_CODE (t) == COMPONENT_REF
- && ort == C_ORT_OMP
&& (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_TO
|| OMP_CLAUSE_CODE (c) == OMP_CLAUSE_FROM)
@@ -4776,6 +4775,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
return error_mark_node;
}
t = TREE_OPERAND (t, 0);
+ if (ort == C_ORT_ACC && TREE_CODE (t) == INDIRECT_REF)
+ t = TREE_OPERAND (t, 0);
}
if (REFERENCE_REF_P (t))
t = TREE_OPERAND (t, 0);
@@ -4875,6 +4876,18 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
if (low_bound == NULL_TREE)
low_bound = integer_zero_node;
+ if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH))
+ {
+ if (length != integer_one_node)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c),
+ "expected single pointer in %qs clause",
+ c_omp_map_clause_name (c, ort == C_ORT_ACC));
+ return error_mark_node;
+ }
+ }
if (length != NULL_TREE)
{
if (!integer_nonzerop (length))
@@ -5322,12 +5335,18 @@ handle_omp_array_sections (tree c, enum c_omp_region_type \
ort) if ((ort & C_ORT_OMP_DECLARE_SIMD) != C_ORT_OMP && ort != C_ORT_ACC)
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_POINTER);
else if (TREE_CODE (t) == COMPONENT_REF)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ {
+ gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, k);
+ }
else if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
t = TREE_OPERAND (t, 0);
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ gomp_map_kind k = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, k);
}
else
OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_FIRSTPRIVATE_POINTER);
@@ -6243,6 +6262,41 @@ cp_omp_finish_iterators (tree iter)
return ret;
}
+/* Ensure that pointers are used in OpenACC attach and detach clauses.
+ Return true if an error has been detected. */
+
+static bool
+cp_oacc_check_attachments (tree c)
+{
+ if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
+ return false;
+
+ /* OpenACC attach / detach clauses must be pointers. */
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH)
+ {
+ tree t = OMP_CLAUSE_DECL (c);
+ tree type;
+
+ while (TREE_CODE (t) == TREE_LIST)
+ t = TREE_CHAIN (t);
+
+ type = TREE_TYPE (t);
+
+ if (TREE_CODE (type) == REFERENCE_TYPE)
+ type = TREE_TYPE (type);
+
+ if (TREE_CODE (type) != POINTER_TYPE)
+ {
+ error_at (OMP_CLAUSE_LOCATION (c), "expected pointer in %qs clause",
+ c_omp_map_clause_name (c, true));
+ return true;
+ }
+ }
+
+ return false;
+}
+
/* For all elements of CLAUSES, validate them vs OpenMP constraints.
Remove any elements from the list that are invalid. */
@@ -6507,7 +6561,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = OMP_CLAUSE_DECL (c);
check_dup_generic_t:
if (t == current_class_ptr
- && (ort != C_ORT_OMP_DECLARE_SIMD
+ && ((ort != C_ORT_OMP_DECLARE_SIMD && ort != C_ORT_ACC)
|| (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_LINEAR
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE_UNIFORM)))
{
@@ -6577,8 +6631,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
handle_field_decl:
if (!remove
&& TREE_CODE (t) == FIELD_DECL
- && t == OMP_CLAUSE_DECL (c)
- && ort != C_ORT_ACC)
+ && t == OMP_CLAUSE_DECL (c))
{
OMP_CLAUSE_DECL (c)
= omp_privatize_field (t, (OMP_CLAUSE_CODE (c)
@@ -6645,7 +6698,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
omp_note_field_privatization (t, OMP_CLAUSE_DECL (c));
else
t = OMP_CLAUSE_DECL (c);
- if (t == current_class_ptr)
+ if (ort != C_ORT_ACC && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7134,7 +7187,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (t == error_mark_node)
remove = true;
- else if (t == current_class_ptr)
+ else if (ort != C_ORT_ACC && t == current_class_ptr)
{
error_at (OMP_CLAUSE_LOCATION (c),
"%<this%> allowed in OpenMP only in %<declare simd%>"
@@ -7265,6 +7318,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
}
}
+ if (cp_oacc_check_attachments (c))
+ remove = true;
break;
}
if (t == error_mark_node)
@@ -7272,14 +7327,25 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
remove = true;
break;
}
+ /* OpenACC attach / detach clauses must be pointers. */
+ if (cp_oacc_check_attachments (c))
+ {
+ remove = true;
+ break;
+ }
if (REFERENCE_REF_P (t)
&& TREE_CODE (TREE_OPERAND (t, 0)) == COMPONENT_REF)
{
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_OMP_DECLARE_SIMD) == C_ORT_OMP
+ || ort == C_ORT_ACC)
&& OMP_CLAUSE_CODE (c) != OMP_CLAUSE__CACHE_)
{
if (type_dependent_expression_p (t))
@@ -7329,7 +7395,8 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
break;
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP
&& (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
- || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER))
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH))
break;
if (DECL_P (t))
error_at (OMP_CLAUSE_LOCATION (c),
@@ -7411,7 +7478,9 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
else
bitmap_set_bit (&generic_head, DECL_UID (t));
}
- else if (bitmap_bit_p (&map_head, DECL_UID (t)))
+ else if (bitmap_bit_p (&map_head, DECL_UID (t))
+ && (ort != C_ORT_ACC
+ || !bitmap_bit_p (&map_field_head, DECL_UID (t))))
{
if (OMP_CLAUSE_CODE (c) != OMP_CLAUSE_MAP)
error_at (OMP_CLAUSE_LOCATION (c),
@@ -7466,7 +7535,12 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
if (TREE_CODE (t) == COMPONENT_REF)
- OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALWAYS_POINTER);
+ {
+ gomp_map_kind k
+ = (ort == C_ORT_ACC) ? GOMP_MAP_ATTACH_DETACH
+ : GOMP_MAP_ALWAYS_POINTER;
+ OMP_CLAUSE_SET_MAP_KIND (c2, k);
+ }
else
OMP_CLAUSE_SET_MAP_KIND (c2,
GOMP_MAP_FIRSTPRIVATE_REFERENCE);
diff --git a/gcc/fortran/gfortran.h b/gcc/fortran/gfortran.h
index e962db59bc5..6469c6b26cf 100644
--- a/gcc/fortran/gfortran.h
+++ b/gcc/fortran/gfortran.h
@@ -1192,10 +1192,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 dc0521b40f0..d79f4a90271 100644
--- a/gcc/fortran/openmp.c
+++ b/gcc/fortran/openmp.c
@@ -233,7 +233,8 @@ static match
gfc_match_omp_variable_list (const char *str, gfc_omp_namelist **list,
bool allow_common, bool *end_colon = NULL,
gfc_omp_namelist ***headp = NULL,
- bool allow_sections = false)
+ bool allow_sections = false,
+ bool allow_derived = false)
{
gfc_omp_namelist *head, *tail, *p;
locus old_loc, cur_loc;
@@ -259,7 +260,8 @@ gfc_match_omp_variable_list (const char *str, gfc_omp_namelist \
**list, case MATCH_YES:
gfc_expr *expr;
expr = NULL;
- if (allow_sections && gfc_peek_ascii_char () == '(')
+ if ((allow_sections && gfc_peek_ascii_char () == '(')
+ || (allow_derived && gfc_peek_ascii_char () == '%'))
{
gfc_current_locus = cur_loc;
m = gfc_match_variable (&expr, 0);
@@ -797,7 +799,7 @@ enum omp_mask1
OMP_MASK1_LAST
};
-/* OpenACC 2.0 specific clauses. */
+/* OpenACC 2.0+ specific clauses. */
enum omp_mask2
{
OMP_CLAUSE_ASYNC,
@@ -823,6 +825,8 @@ enum omp_mask2
OMP_CLAUSE_TILE,
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
+ OMP_CLAUSE_ATTACH,
+ OMP_CLAUSE_DETACH,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -927,10 +931,11 @@ omp_inv_mask::omp_inv_mask (const omp_mask &m) : omp_mask (m)
static bool
gfc_match_omp_map_clause (gfc_omp_namelist **list, gfc_omp_map_op map_op,
- bool allow_common)
+ bool allow_common, bool allow_derived)
{
gfc_omp_namelist **head = NULL;
- if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true)
+ if (gfc_match_omp_variable_list ("", list, allow_common, NULL, &head, true,
+ allow_derived)
== MATCH_YES)
{
gfc_omp_namelist *n;
@@ -952,6 +957,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, {
gfc_omp_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)));
gcc_checking_assert (OMP_MASK1_LAST <= 64 && OMP_MASK2_LAST <= 64);
*cp = NULL;
@@ -1025,6 +1038,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const 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 'c':
if ((mask & OMP_CLAUSE_COLLAPSE)
@@ -1052,7 +1071,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, true))
+ OMP_MAP_TOFROM, true,
+ allow_derived))
continue;
if (mask & OMP_CLAUSE_COPYIN)
{
@@ -1060,7 +1080,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, {
if (gfc_match ("copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, true))
+ OMP_MAP_TO, true,
+ allow_derived))
continue;
}
else if (gfc_match_omp_variable_list ("copyin (",
@@ -1071,7 +1092,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const 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))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
@@ -1081,7 +1102,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const 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))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
break;
case 'd':
@@ -1117,7 +1138,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, if ((mask & OMP_CLAUSE_DELETE)
&& gfc_match ("delete ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_RELEASE, true))
+ OMP_MAP_RELEASE, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEPEND)
&& gfc_match ("depend ( ") == MATCH_YES)
@@ -1160,6 +1182,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const 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
@@ -1169,12 +1197,14 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, && openacc
&& gfc_match ("device ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_TO, true))
+ OMP_MAP_FORCE_TO, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICEPTR)
&& gfc_match ("deviceptr ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_DEVICEPTR, false))
+ OMP_MAP_FORCE_DEVICEPTR, false,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_DEVICE_RESIDENT)
&& gfc_match_omp_variable_list
@@ -1252,7 +1282,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("host ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM, true))
+ OMP_MAP_FORCE_FROM, true,
+ allow_derived))
continue;
break;
case 'i':
@@ -1524,47 +1555,49 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("pcopy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, true))
+ OMP_MAP_TOFROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("pcopyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, true))
+ OMP_MAP_TO, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("pcopyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, true))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("pcreate ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, true))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_PRESENT)
&& gfc_match ("present ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_PRESENT, false))
+ OMP_MAP_FORCE_PRESENT, false,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPY)
&& gfc_match ("present_or_copy ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TOFROM, true))
+ OMP_MAP_TOFROM, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYIN)
&& gfc_match ("present_or_copyin ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_TO, true))
+ OMP_MAP_TO, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("present_or_copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, true))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("present_or_create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, true))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_PRIORITY)
&& c->priority == NULL
@@ -1682,8 +1715,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const 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)
@@ -1782,7 +1815,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, if ((mask & OMP_CLAUSE_HOST_SELF)
&& gfc_match ("self ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FORCE_FROM, true))
+ OMP_MAP_FORCE_FROM, true,
+ allow_derived))
continue;
if ((mask & OMP_CLAUSE_SEQ)
&& !c->seq
@@ -1957,23 +1991,23 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, | OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_NUM_GANGS \
| OMP_CLAUSE_NUM_WORKERS | OMP_CLAUSE_VECTOR_LENGTH | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
#define OACC_SERIAL_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_ASYNC | OMP_CLAUSE_REDUCTION \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT | OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE | OMP_CLAUSE_DEFAULT \
- | OMP_CLAUSE_WAIT)
+ | OMP_CLAUSE_WAIT | OMP_CLAUSE_ATTACH)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_COPY \
| OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT | OMP_CLAUSE_CREATE \
- | OMP_CLAUSE_PRESENT)
+ | OMP_CLAUSE_PRESENT | OMP_CLAUSE_ATTACH)
#define OACC_LOOP_CLAUSES \
(omp_mask (OMP_CLAUSE_COLLAPSE) | OMP_CLAUSE_GANG | OMP_CLAUSE_WORKER \
| OMP_CLAUSE_VECTOR | OMP_CLAUSE_SEQ | OMP_CLAUSE_INDEPENDENT \
@@ -1996,10 +2030,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask \
mask, | OMP_CLAUSE_DEVICE | OMP_CLAUSE_WAIT | OMP_CLAUSE_IF_PRESENT)
#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_COPYOUT | OMP_CLAUSE_DELETE | OMP_CLAUSE_FINALIZE \
+ | OMP_CLAUSE_DETACH)
#define OACC_WAIT_CLAUSES \
omp_mask (OMP_CLAUSE_ASYNC)
#define OACC_ROUTINE_CLAUSES \
@@ -3847,9 +3882,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);
@@ -3890,9 +3922,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))
@@ -4275,11 +4304,26 @@ 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;
+ break;
+ }
+ 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);
@@ -4470,23 +4514,43 @@ 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. */
+ if (openacc)
+ 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)
+ || 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 (n->expr->ref->u.ar.codimen)
+ 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-expr.c b/gcc/fortran/trans-expr.c
index fe89c7b02ed..90fbe0ab74f 100644
--- a/gcc/fortran/trans-expr.c
+++ b/gcc/fortran/trans-expr.c
@@ -2423,7 +2423,7 @@ gfc_conv_substring (gfc_se * se, gfc_ref * ref, int kind,
/* Convert a derived type component reference. */
-static void
+void
gfc_conv_component_ref (gfc_se * se, gfc_ref * ref)
{
gfc_component *c;
@@ -2513,7 +2513,7 @@ gfc_conv_component_ref (gfc_se * se, gfc_ref * ref)
/* This function deals with component references to components of the
parent type for derived type extensions. */
-static void
+void
conv_parent_component_references (gfc_se * se, gfc_ref * ref)
{
gfc_component *c;
@@ -2579,6 +2579,95 @@ conv_inquiry (gfc_se * se, gfc_ref * ref, gfc_expr *expr, \
gfc_typespec *ts) se->expr = res;
}
+/* Transparently dereference VAR if it is a pointer, reference, etc.
+ according to Fortran semantics. */
+
+tree
+gfc_auto_dereference_var (gfc_symbol *sym, tree var, bool descriptor_only_p,
+ bool is_classarray)
+{
+ /* Characters are entirely different from other types, they are treated
+ separately. */
+ if (sym->ts.type == BT_CHARACTER)
+ {
+ /* Dereference character pointer dummy arguments
+ or results. */
+ if ((sym->attr.pointer || sym->attr.allocatable)
+ && (sym->attr.dummy
+ || sym->attr.function
+ || sym->attr.result))
+ var = build_fold_indirect_ref_loc (input_location, var);
+ }
+ else if (!sym->attr.value)
+ {
+ /* Dereference temporaries for class array dummy arguments. */
+ if (sym->attr.dummy && is_classarray
+ && GFC_ARRAY_TYPE_P (TREE_TYPE (var)))
+ {
+ if (!descriptor_only_p)
+ var = GFC_DECL_SAVED_DESCRIPTOR (var);
+
+ var = build_fold_indirect_ref_loc (input_location, var);
+ }
+
+ /* Dereference non-character scalar dummy arguments. */
+ if (sym->attr.dummy && !sym->attr.dimension
+ && !(sym->attr.codimension && sym->attr.allocatable)
+ && (sym->ts.type != BT_CLASS
+ || (!CLASS_DATA (sym)->attr.dimension
+ && !(CLASS_DATA (sym)->attr.codimension
+ && CLASS_DATA (sym)->attr.allocatable))))
+ var = build_fold_indirect_ref_loc (input_location, var);
+
+ /* Dereference scalar hidden result. */
+ if (flag_f2c && sym->ts.type == BT_COMPLEX
+ && (sym->attr.function || sym->attr.result)
+ && !sym->attr.dimension && !sym->attr.pointer
+ && !sym->attr.always_explicit)
+ var = build_fold_indirect_ref_loc (input_location, var);
+
+ /* Dereference non-character, non-class pointer variables.
+ These must be dummies, results, or scalars. */
+ if (!is_classarray
+ && (sym->attr.pointer || sym->attr.allocatable
+ || gfc_is_associate_pointer (sym)
+ || (sym->as && sym->as->type == AS_ASSUMED_RANK))
+ && (sym->attr.dummy
+ || sym->attr.function
+ || sym->attr.result
+ || (!sym->attr.dimension
+ && (!sym->attr.codimension || !sym->attr.allocatable))))
+ var = build_fold_indirect_ref_loc (input_location, var);
+ /* Now treat the class array pointer variables accordingly. */
+ else if (sym->ts.type == BT_CLASS
+ && sym->attr.dummy
+ && (CLASS_DATA (sym)->attr.dimension
+ || CLASS_DATA (sym)->attr.codimension)
+ && ((CLASS_DATA (sym)->as
+ && CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK)
+ || CLASS_DATA (sym)->attr.allocatable
+ || CLASS_DATA (sym)->attr.class_pointer))
+ var = build_fold_indirect_ref_loc (input_location, var);
+ /* And the case where a non-dummy, non-result, non-function,
+ non-allotable and non-pointer classarray is present. This case was
+ previously covered by the first if, but with introducing the
+ condition !is_classarray there, that case has to be covered
+ explicitly. */
+ else if (sym->ts.type == BT_CLASS
+ && !sym->attr.dummy
+ && !sym->attr.function
+ && !sym->attr.result
+ && (CLASS_DATA (sym)->attr.dimension
+ || CLASS_DATA (sym)->attr.codimension)
+ && (sym->assoc
+ || !CLASS_DATA (sym)->attr.allocatable)
+ && !CLASS_DATA (sym)->attr.class_pointer)
+ var = build_fold_indirect_ref_loc (input_location, var);
+ }
+
+ return var;
+}
+
/* Return the contents of a variable. Also handles reference/pointer
variables (all Fortran pointer references are implicit). */
@@ -2685,94 +2774,9 @@ gfc_conv_variable (gfc_se * se, gfc_expr * expr)
return;
}
-
- /* Dereference the expression, where needed. Since characters
- are entirely different from other types, they are treated
- separately. */
- if (sym->ts.type == BT_CHARACTER)
- {
- /* Dereference character pointer dummy arguments
- or results. */
- if ((sym->attr.pointer || sym->attr.allocatable)
- && (sym->attr.dummy
- || sym->attr.function
- || sym->attr.result))
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
-
- }
- else if (!sym->attr.value)
- {
- /* Dereference temporaries for class array dummy arguments. */
- if (sym->attr.dummy && is_classarray
- && GFC_ARRAY_TYPE_P (TREE_TYPE (se->expr)))
- {
- if (!se->descriptor_only)
- se->expr = GFC_DECL_SAVED_DESCRIPTOR (se->expr);
-
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
- }
-
- /* Dereference non-character scalar dummy arguments. */
- if (sym->attr.dummy && !sym->attr.dimension
- && !(sym->attr.codimension && sym->attr.allocatable)
- && (sym->ts.type != BT_CLASS
- || (!CLASS_DATA (sym)->attr.dimension
- && !(CLASS_DATA (sym)->attr.codimension
- && CLASS_DATA (sym)->attr.allocatable))))
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
-
- /* Dereference scalar hidden result. */
- if (flag_f2c && sym->ts.type == BT_COMPLEX
- && (sym->attr.function || sym->attr.result)
- && !sym->attr.dimension && !sym->attr.pointer
- && !sym->attr.always_explicit)
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
-
- /* Dereference non-character, non-class pointer variables.
- These must be dummies, results, or scalars. */
- if (!is_classarray
- && (sym->attr.pointer || sym->attr.allocatable
- || gfc_is_associate_pointer (sym)
- || (sym->as && sym->as->type == AS_ASSUMED_RANK))
- && (sym->attr.dummy
- || sym->attr.function
- || sym->attr.result
- || (!sym->attr.dimension
- && (!sym->attr.codimension || !sym->attr.allocatable))))
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
- /* Now treat the class array pointer variables accordingly. */
- else if (sym->ts.type == BT_CLASS
- && sym->attr.dummy
- && (CLASS_DATA (sym)->attr.dimension
- || CLASS_DATA (sym)->attr.codimension)
- && ((CLASS_DATA (sym)->as
- && CLASS_DATA (sym)->as->type == AS_ASSUMED_RANK)
- || CLASS_DATA (sym)->attr.allocatable
- || CLASS_DATA (sym)->attr.class_pointer))
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
- /* And the case where a non-dummy, non-result, non-function,
- non-allotable and non-pointer classarray is present. This case was
- previously covered by the first if, but with introducing the
- condition !is_classarray there, that case has to be covered
- explicitly. */
- else if (sym->ts.type == BT_CLASS
- && !sym->attr.dummy
- && !sym->attr.function
- && !sym->attr.result
- && (CLASS_DATA (sym)->attr.dimension
- || CLASS_DATA (sym)->attr.codimension)
- && (sym->assoc
- || !CLASS_DATA (sym)->attr.allocatable)
- && !CLASS_DATA (sym)->attr.class_pointer)
- se->expr = build_fold_indirect_ref_loc (input_location,
- se->expr);
- }
+ /* Dereference the expression, where needed. */
+ se->expr = gfc_auto_dereference_var (sym, se->expr, se->descriptor_only,
+ is_classarray);
ref = expr->ref;
}
diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c
index d9dfcabc65e..2b31d44983a 100644
--- a/gcc/fortran/trans-openmp.c
+++ b/gcc/fortran/trans-openmp.c
@@ -169,6 +169,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
@@ -1930,6 +1933,92 @@ gfc_convert_expr_to_tree (stmtblock_t *block, gfc_expr *expr)
static vec<tree, va_heap, vl_embed> *doacross_steps;
+
+/* Translate an array section or array element. */
+
+static void
+gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
+ tree decl, bool element, gomp_map_kind ptr_kind,
+ tree node, tree &node2, tree &node3, tree &node4)
+{
+ gfc_se se;
+ tree ptr, ptr2;
+
+ gfc_init_se (&se, NULL);
+
+ if (element)
+ {
+ gfc_conv_expr_reference (&se, n->expr);
+ gfc_add_block_to_block (block, &se.pre);
+ ptr = se.expr;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+ }
+ else
+ {
+ gfc_conv_expr_descriptor (&se, n->expr);
+ ptr = gfc_conv_array_data (se.expr);
+ tree type = TREE_TYPE (se.expr);
+ gfc_add_block_to_block (block, &se.pre);
+ OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr,
+ GFC_TYPE_ARRAY_RANK (type));
+ tree elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node) = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ gfc_add_block_to_block (block, &se.post);
+ ptr = fold_convert (build_pointer_type (char_type_node), ptr);
+ OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+
+ if (POINTER_TYPE_P (TREE_TYPE (decl))
+ && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
+ && ptr_kind == GOMP_MAP_POINTER)
+ {
+ node4 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
+ OMP_CLAUSE_DECL (node4) = decl;
+ OMP_CLAUSE_SIZE (node4) = size_int (0);
+ decl = build_fold_indirect_ref (decl);
+ }
+ ptr = fold_convert (sizetype, ptr);
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ {
+ tree type = TREE_TYPE (decl);
+ ptr2 = gfc_conv_descriptor_data_get (decl);
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = decl;
+ OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (decl);
+ if (ptr_kind == GOMP_MAP_ATTACH_DETACH)
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+ }
+ else
+ {
+ if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
+ ptr2 = build_fold_addr_expr (decl);
+ else
+ {
+ gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
+ ptr2 = decl;
+ }
+ node3 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_kind);
+ OMP_CLAUSE_DECL (node3) = decl;
+ }
+ ptr2 = fold_convert (sizetype, ptr2);
+ OMP_CLAUSE_SIZE (node3)
+ = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+}
+
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false)
@@ -2258,7 +2347,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses \
*clauses, || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
|| GFC_DECL_CRAY_POINTEE (decl)
|| GFC_DESCRIPTOR_TYPE_P
- (TREE_TYPE (TREE_TYPE (decl)))))
+ (TREE_TYPE (TREE_TYPE (decl)))
+ || n->sym->ts.type == BT_DERIVED))
{
tree orig_decl = decl;
node4 = build_omp_clause (input_location,
@@ -2280,10 +2370,13 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses \
*clauses, decl = build_fold_indirect_ref (decl);
}
}
- 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);
+
ptr = fold_convert (build_pointer_type (char_type_node),
ptr);
ptr = build_fold_indirect_ref (ptr);
@@ -2349,88 +2442,152 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses \
*clauses, else
OMP_CLAUSE_DECL (node) = decl;
}
- else
+ else if (n->expr
+ && n->expr->expr_type == EXPR_VARIABLE
+ && n->expr->ref->type == REF_COMPONENT)
{
- tree ptr, ptr2;
+ gfc_ref *lastcomp;
+
+ for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
+ if (ref->type == REF_COMPONENT)
+ lastcomp = ref;
+
+ symbol_attribute sym_attr;
+
+ sym_attr = lastcomp->u.c.component->attr;
+
gfc_init_se (&se, NULL);
- if (n->expr->ref->u.ar.type == AR_ELEMENT)
+
+ if (!sym_attr.dimension
+ && lastcomp->u.c.component->ts.type != BT_DERIVED)
{
- gfc_conv_expr_reference (&se, n->expr);
+ /* Last component is a scalar. */
+ gfc_conv_expr (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
- ptr = se.expr;
- OMP_CLAUSE_SIZE (node)
- = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
+ OMP_CLAUSE_DECL (node) = se.expr;
+ gfc_add_block_to_block (block, &se.post);
+ goto finalize_map_clause;
}
- else
- {
- gfc_conv_expr_descriptor (&se, n->expr);
- ptr = gfc_conv_array_data (se.expr);
- tree type = TREE_TYPE (se.expr);
- gfc_add_block_to_block (block, &se.pre);
- OMP_CLAUSE_SIZE (node)
- = gfc_full_array_size (block, se.expr,
- GFC_TYPE_ARRAY_RANK (type));
- tree elemsz
- = TYPE_SIZE_UNIT (gfc_get_element_type (type));
- elemsz = fold_convert (gfc_array_index_type, elemsz);
- OMP_CLAUSE_SIZE (node)
- = fold_build2 (MULT_EXPR, gfc_array_index_type,
- OMP_CLAUSE_SIZE (node), elemsz);
- }
- gfc_add_block_to_block (block, &se.post);
- ptr = fold_convert (build_pointer_type (char_type_node),
- ptr);
- OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
- if (POINTER_TYPE_P (TREE_TYPE (decl))
- && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
+ se.expr = gfc_auto_dereference_var (n->sym, decl);
+
+ for (gfc_ref *ref = n->expr->ref;
+ ref && ref != lastcomp->next;
+ ref = ref->next)
{
- node4 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
- OMP_CLAUSE_DECL (node4) = decl;
- OMP_CLAUSE_SIZE (node4) = size_int (0);
- decl = build_fold_indirect_ref (decl);
+ if (ref->type == REF_COMPONENT)
+ {
+ if (ref->u.c.sym->attr.extension)
+ conv_parent_component_references (&se, ref);
+
+ gfc_conv_component_ref (&se, ref);
+ }
+ else
+ sorry ("unhandled derived-type component");
}
- ptr = fold_convert (sizetype, ptr);
- if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+
+ tree inner = se.expr;
+
+ /* Last component is a derived type. */
+ if (lastcomp->u.c.component->ts.type == BT_DERIVED)
{
- tree type = TREE_TYPE (decl);
- ptr2 = gfc_conv_descriptor_data_get (decl);
- node2 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
- OMP_CLAUSE_DECL (node2) = decl;
- 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_DECL (node3)
- = gfc_conv_descriptor_data_get (decl);
+ if (sym_attr.allocatable || sym_attr.pointer)
+ {
+ tree data = inner;
+ tree size = TYPE_SIZE_UNIT (TREE_TYPE (inner));
+
+ OMP_CLAUSE_DECL (node)
+ = build_fold_indirect_ref (data);
+ OMP_CLAUSE_SIZE (node) = size;
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2,
+ GOMP_MAP_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (node2) = data;
+ OMP_CLAUSE_SIZE (node2) = size_int (0);
+ }
+ else
+ {
+ OMP_CLAUSE_DECL (node) = decl;
+ OMP_CLAUSE_SIZE (node)
+ = TYPE_SIZE_UNIT (TREE_TYPE (decl));
+ }
}
- else
+ else if (lastcomp->next
+ && lastcomp->next->type == REF_ARRAY
+ && lastcomp->next->u.ar.type == AR_FULL)
{
- if (TREE_CODE (TREE_TYPE (decl)) == ARRAY_TYPE)
- ptr2 = build_fold_addr_expr (decl);
- else
+ /* Just pass the (auto-dereferenced) decl through for
+ bare attach and detach clauses. */
+ if (n->u.map_op == OMP_MAP_ATTACH
+ || n->u.map_op == OMP_MAP_DETACH)
{
- gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
- ptr2 = decl;
+ OMP_CLAUSE_DECL (node) = inner;
+ OMP_CLAUSE_SIZE (node) = size_zero_node;
+ goto finalize_map_clause;
}
- node3 = build_omp_clause (input_location,
- OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
- OMP_CLAUSE_DECL (node3) = decl;
+
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
+ {
+ tree type = TREE_TYPE (inner);
+ tree ptr = gfc_conv_descriptor_data_get (inner);
+ ptr = build_fold_indirect_ref (ptr);
+ OMP_CLAUSE_DECL (node) = ptr;
+ node2 = build_omp_clause (input_location,
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (node2, GOMP_MAP_TO_PSET);
+ OMP_CLAUSE_DECL (node2) = inner;
+ 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_ATTACH_DETACH);
+ OMP_CLAUSE_DECL (node3)
+ = gfc_conv_descriptor_data_get (inner);
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
+ OMP_CLAUSE_SIZE (node3) = size_int (0);
+ int rank = GFC_TYPE_ARRAY_RANK (type);
+ OMP_CLAUSE_SIZE (node)
+ = gfc_full_array_size (block, inner, rank);
+ tree elemsz
+ = TYPE_SIZE_UNIT (gfc_get_element_type (type));
+ elemsz = fold_convert (gfc_array_index_type, elemsz);
+ OMP_CLAUSE_SIZE (node)
+ = fold_build2 (MULT_EXPR, gfc_array_index_type,
+ OMP_CLAUSE_SIZE (node), elemsz);
+ }
+ else
+ OMP_CLAUSE_DECL (node) = inner;
}
- ptr2 = fold_convert (sizetype, ptr2);
- OMP_CLAUSE_SIZE (node3)
- = fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
+ else /* An array element or section. */
+ {
+ bool element
+ = (lastcomp->next
+ && lastcomp->next->type == REF_ARRAY
+ && lastcomp->next->u.ar.type == AR_ELEMENT);
+
+ gfc_trans_omp_array_section (block, n, inner, element,
+ GOMP_MAP_ATTACH_DETACH,
+ node, node2, node3, node4);
+ }
+ }
+ else /* An array element or array section. */
+ {
+ bool element = n->expr->ref->u.ar.type == AR_ELEMENT;
+ gfc_trans_omp_array_section (block, n, decl, element,
+ GOMP_MAP_POINTER, node, node2,
+ node3, node4);
}
+
+ finalize_map_clause:
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;
@@ -2455,6 +2612,9 @@ gfc_trans_omp_clauses (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/fortran/trans.h b/gcc/fortran/trans.h
index 359c7a2561a..24aaf3d4f96 100644
--- a/gcc/fortran/trans.h
+++ b/gcc/fortran/trans.h
@@ -565,6 +565,14 @@ tree gfc_conv_expr_present (gfc_symbol *);
/* Convert a missing, dummy argument into a null or zero. */
void gfc_conv_missing_dummy (gfc_se *, gfc_expr *, gfc_typespec, int);
+/* Lowering of component references. */
+void gfc_conv_component_ref (gfc_se * se, gfc_ref * ref);
+void conv_parent_component_references (gfc_se * se, gfc_ref * ref);
+
+/* Automatically dereference var. */
+tree gfc_auto_dereference_var (gfc_symbol *, tree, bool desc_only = false,
+ bool is_classarray = false);
+
/* Generate code to allocate a string temporary. */
tree gfc_conv_string_tmp (gfc_se *, tree, tree);
/* Get the string length variable belonging to an expression. */
diff --git a/gcc/gimplify.c b/gcc/gimplify.c
index 8f8fb2b54a1..9c78afedf2e 100644
--- a/gcc/gimplify.c
+++ b/gcc/gimplify.c
@@ -123,6 +123,10 @@ enum gimplify_omp_var_data
/* Flag for GOVD_REDUCTION: inscan seen in {in,ex}clusive clause. */
GOVD_REDUCTION_INSCAN = 0x2000000,
+ /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
+ fields. */
+ GOVD_MAP_HAS_ATTACHMENTS = 8388608,
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -8206,20 +8210,33 @@ insert_struct_comp_map (enum tree_code code, tree c, tree \
struct_node, tree prev_node, tree *scp)
{
enum gomp_map_kind mkind
- = code == OMP_TARGET_EXIT_DATA ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
+ = (code == OMP_TARGET_EXIT_DATA || code == OACC_EXIT_DATA)
+ ? GOMP_MAP_RELEASE : GOMP_MAP_ALLOC;
tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
tree cl = scp ? prev_node : c2;
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;
/* We might need to create an additional mapping if we have a reference to a
- pointer (in C++). */
- if (OMP_CLAUSE_CHAIN (prev_node) != c)
+ pointer (in C++). Don't do this if we have something other than a
+ GOMP_MAP_ALWAYS_POINTER though, i.e. a GOMP_MAP_TO_PSET. */
+ 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_ALWAYS_POINTER)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+ == GOMP_MAP_ATTACH_DETACH)))
{
tree c4 = OMP_CLAUSE_CHAIN (prev_node);
tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c), OMP_CLAUSE_MAP);
@@ -8326,6 +8343,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
struct gimplify_omp_ctx *ctx, *outer_ctx;
tree c;
hash_map<tree, tree> *struct_map_to_clause = NULL;
+ hash_set<tree> *struct_deref_set = NULL;
tree *prev_list_p = NULL, *orig_list_p = list_p;
int handled_depend_iterators = -1;
int nowait = -1;
@@ -8728,8 +8746,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
case OMP_TARGET_DATA:
case OMP_TARGET_ENTER_DATA:
case OMP_TARGET_EXIT_DATA:
- case OACC_ENTER_DATA:
- case OACC_EXIT_DATA:
case OACC_HOST_DATA:
if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
|| (OMP_CLAUSE_MAP_KIND (c)
@@ -8738,6 +8754,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
mapped, but not the pointer to it. */
remove = true;
break;
+ case OACC_ENTER_DATA:
+ case OACC_EXIT_DATA:
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER
+ || (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ remove = true;
+ break;
default:
break;
}
@@ -8811,7 +8836,35 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
pd = &TREE_OPERAND (decl, 0);
decl = TREE_OPERAND (decl, 0);
}
- if (TREE_CODE (decl) == COMPONENT_REF)
+ bool indir_p = false;
+ tree orig_decl = decl;
+ tree decl_ref = NULL_TREE;
+ if ((region_type & ORT_ACC) != 0
+ && TREE_CODE (*pd) == COMPONENT_REF
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH
+ && code != OACC_UPDATE)
+ {
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ {
+ decl = TREE_OPERAND (decl, 0);
+ if ((TREE_CODE (decl) == MEM_REF
+ && integer_zerop (TREE_OPERAND (decl, 1)))
+ || INDIRECT_REF_P (decl))
+ {
+ indir_p = true;
+ decl = TREE_OPERAND (decl, 0);
+ }
+ if (TREE_CODE (decl) == INDIRECT_REF
+ && DECL_P (TREE_OPERAND (decl, 0))
+ && (TREE_CODE (TREE_TYPE (TREE_OPERAND (decl, 0)))
+ == REFERENCE_TYPE))
+ {
+ decl_ref = decl;
+ decl = TREE_OPERAND (decl, 0);
+ }
+ }
+ }
+ else if (TREE_CODE (decl) == COMPONENT_REF)
{
while (TREE_CODE (decl) == COMPONENT_REF)
decl = TREE_OPERAND (decl, 0);
@@ -8821,13 +8874,76 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
== REFERENCE_TYPE))
decl = TREE_OPERAND (decl, 0);
}
+ if (decl != orig_decl && DECL_P (decl) && indir_p)
+ {
+ gomp_map_kind k = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH;
+ /* We have a dereference of a struct member. Make this an
+ attach/detach operation, and ensure the base pointer is
+ mapped as a FIRSTPRIVATE_POINTER. */
+ OMP_CLAUSE_SET_MAP_KIND (c, k);
+ flags = GOVD_MAP | GOVD_SEEN | GOVD_EXPLICIT;
+ tree next_clause = OMP_CLAUSE_CHAIN (c);
+ if (k == GOMP_MAP_ATTACH
+ && code != OACC_ENTER_DATA
+ && (!next_clause
+ || (OMP_CLAUSE_CODE (next_clause) != OMP_CLAUSE_MAP)
+ || (OMP_CLAUSE_MAP_KIND (next_clause)
+ != GOMP_MAP_POINTER)
+ || OMP_CLAUSE_DECL (next_clause) != decl)
+ && (!struct_deref_set
+ || !struct_deref_set->contains (decl)))
+ {
+ if (!struct_deref_set)
+ struct_deref_set = new hash_set<tree> ();
+ /* As well as the attach, we also need a
+ FIRSTPRIVATE_POINTER clause to properly map the
+ pointer to the struct base. */
+ tree c2 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c2, GOMP_MAP_ALLOC);
+ OMP_CLAUSE_MAP_MAYBE_ZERO_LENGTH_ARRAY_SECTION (c2)
+ = 1;
+ tree charptr_zero
+ = build_int_cst (build_pointer_type (char_type_node),
+ 0);
+ OMP_CLAUSE_DECL (c2)
+ = build2 (MEM_REF, char_type_node,
+ decl_ref ? decl_ref : decl, charptr_zero);
+ OMP_CLAUSE_SIZE (c2) = size_zero_node;
+ tree c3 = build_omp_clause (OMP_CLAUSE_LOCATION (c),
+ OMP_CLAUSE_MAP);
+ OMP_CLAUSE_SET_MAP_KIND (c3,
+ GOMP_MAP_FIRSTPRIVATE_POINTER);
+ OMP_CLAUSE_DECL (c3) = decl;
+ OMP_CLAUSE_SIZE (c3) = size_zero_node;
+ tree mapgrp = *prev_list_p;
+ *prev_list_p = c2;
+ OMP_CLAUSE_CHAIN (c3) = mapgrp;
+ OMP_CLAUSE_CHAIN (c2) = c3;
+
+ struct_deref_set->add (decl);
+ }
+ goto do_add_decl;
+ }
+ /* An "attach/detach" operation on an update directive should
+ behave as a GOMP_MAP_ALWAYS_POINTER. Beware that
+ unlike attach or detach map kinds, GOMP_MAP_ALWAYS_POINTER
+ depends on the previous mapping. */
+ if (code == OACC_UPDATE
+ && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
+ OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_ALWAYS_POINTER);
if (gimplify_expr (pd, pre_p, NULL, is_gimple_lvalue, fb_lvalue)
== GS_ERROR)
{
remove = true;
break;
}
- if (DECL_P (decl))
+ if (DECL_P (decl)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_DETACH
+ && code != OACC_UPDATE)
{
if (error_operand_p (decl))
{
@@ -8848,7 +8964,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
break;
}
- if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER)
+ if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ALWAYS_POINTER
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH_DETACH)
{
/* Error recovery. */
if (prev_list_p == NULL)
@@ -8881,20 +8998,47 @@ 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);
+ bool attach_detach = (OMP_CLAUSE_MAP_KIND (c)
+ == GOMP_MAP_ATTACH_DETACH);
+ 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 (attach_detach && (region_type & ORT_ACC) != 0)
+ {
+ /* Turn a GOMP_MAP_ATTACH_DETACH clause into a
+ GOMP_MAP_ATTACH or GOMP_MAP_DETACH clause after we
+ have detected a case that needs a GOMP_MAP_STRUCT
+ mapping added. */
+ gomp_map_kind k
+ = (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH;
+ OMP_CLAUSE_SET_MAP_KIND (c, k);
+ 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);
+ gomp_map_kind k = attach ? GOMP_MAP_FORCE_PRESENT
+ : GOMP_MAP_STRUCT;
+
+ OMP_CLAUSE_SET_MAP_KIND (l, k);
if (base_ref)
OMP_CLAUSE_DECL (l) = unshare_expr (base_ref);
else
OMP_CLAUSE_DECL (l) = decl;
- OMP_CLAUSE_SIZE (l) = size_int (1);
+ OMP_CLAUSE_SIZE (l)
+ = (!attach
+ ? size_int (1)
+ : DECL_P (OMP_CLAUSE_DECL (l))
+ ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+ : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))));
if (struct_map_to_clause == NULL)
struct_map_to_clause = new hash_map<tree, tree>;
struct_map_to_clause->put (decl, l);
- if (ptr)
+ if (ptr || attach_detach)
{
insert_struct_comp_map (code, c, l, *prev_list_p,
NULL);
@@ -8920,23 +9064,31 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
OMP_CLAUSE_CHAIN (l) = c2;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+ || ptr
+ || attach_detach)
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;
- if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
+ if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c))
+ || ptr
+ || attach_detach)
n->value |= GOVD_SEEN;
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)
+ if ((ptr || attach_detach) && sc == prev_list_p)
break;
else if (TREE_CODE (OMP_CLAUSE_DECL (*sc))
!= COMPONENT_REF
@@ -8989,7 +9141,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
|| (known_eq (offset1, offsetn)
&& maybe_lt (bitpos1, bitposn)))
{
- if (ptr)
+ if (ptr || attach_detach)
scp = sc;
else
break;
@@ -8997,10 +9149,11 @@ 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 (ptr)
+ if (!attach)
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
+ if (ptr || attach_detach)
{
tree cl = insert_struct_comp_map (code, c, NULL,
*prev_list_p, scp);
@@ -9030,11 +9183,18 @@ 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_ATTACH_DETACH
+ && 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_ATTACH_DETACH)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_TO_PSET)))
prev_list_p = list_p;
+
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -9558,6 +9718,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
gimplify_omp_ctxp = ctx;
if (struct_map_to_clause)
delete struct_map_to_clause;
+ if (struct_deref_set)
+ delete struct_deref_set;
}
/* Return true if DECL is a candidate for shared to firstprivate
@@ -9705,6 +9867,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);
@@ -12759,8 +12923,9 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
&& omp_find_clause (OMP_STANDALONE_CLAUSES (expr),
OMP_CLAUSE_FINALIZE))
{
- /* Use GOMP_MAP_DELETE/GOMP_MAP_FORCE_FROM to denote that "finalize"
- semantics apply to all mappings of this OpenACC directive. */
+ /* Use GOMP_MAP_DELETE, GOMP_MAP_FORCE_DETACH, and
+ GOMP_MAP_FORCE_FROM to denote that "finalize" semantics apply
+ to all mappings of this OpenACC directive. */
bool finalize_marked = false;
for (tree c = OMP_STANDALONE_CLAUSES (expr); c; c = OMP_CLAUSE_CHAIN (c))
if (OMP_CLAUSE_CODE (c) == OMP_CLAUSE_MAP)
@@ -12774,10 +12939,19 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq \
*pre_p) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE);
finalize_marked = true;
break;
+ case GOMP_MAP_DETACH:
+ 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/omp-low.c b/gcc/omp-low.c
index 19132f76da2..309b608332f 100644
--- a/gcc/omp-low.c
+++ b/gcc/omp-low.c
@@ -11439,6 +11439,9 @@ lower_omp_target (gimple_stmt_iterator *gsi_p, omp_context \
*ctx) case GOMP_MAP_FORCE_DEVICEPTR:
case GOMP_MAP_DEVICE_RESIDENT:
case GOMP_MAP_LINK:
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
gcc_assert (is_gimple_omp_oacc (stmt));
break;
default:
diff --git a/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c \
b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c new file mode 100644
index 00000000000..d411bcfa8e7
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/deep-copy-arrayofstruct.c
@@ -0,0 +1,84 @@
+/* { dg-do compile } */
+
+#include <stdlib.h>
+#include <stdio.h>
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ const int S = 32;
+ mystruct *m = (mystruct *) calloc (S, sizeof (*m));
+ int i, j;
+
+ for (i = 0; i < S; i++)
+ {
+ m[i].a = (int *) malloc (N * sizeof (int));
+ m[i].b = (int *) malloc (N * sizeof (int));
+ m[i].c = (int *) malloc (N * sizeof (int));
+ }
+
+ for (j = 0; j < S; j++)
+ for (i = 0; i < N; i++)
+ {
+ m[j].a[i] = 0;
+ m[j].b[i] = 0;
+ m[j].c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j, k;
+ for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].a[0:N]) /* { dg-error "expected .\\\). before \
.\\\.. token" } */ + for (j = 0; j < N; j++)
+ m[k].a[j]++;
+
+ for (k = 0; k < S; k++)
+#pragma acc parallel loop copy(m[k].b[0:N], m[k].c[5:N-10]) /* { dg-error "expected \
.\\\). before .\\\.. token" } */ + /* { dg-error ".m. appears more than once in data \
clauses" "" { target c++ } .-1 } */ + for (j = 0; j < N; j++)
+ {
+ m[k].b[j]++;
+ if (j > 5 && j < N - 5)
+ m[k].c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (j = 0; j < S; j++)
+ {
+ for (i = 0; i < N; i++)
+ {
+ if (m[j].a[i] != 99)
+ abort ();
+ if (m[j].b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m[j].c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m[j].c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m[j].a);
+ free (m[j].b);
+ free (m[j].c);
+ }
+ free (m);
+
+ return 0;
+}
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-1.c \
b/gcc/testsuite/c-c++-common/goacc/mdc-1.c new file mode 100644
index 00000000000..6c6a81ea73a
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-1.c
@@ -0,0 +1,55 @@
+/* Test OpenACC's support for manual deep copy, including the attach
+ and detach clauses. */
+
+/* { dg-do compile { target int32 } } */
+/* { dg-additional-options "-fdump-tree-omplower" } */
+
+void
+t1 ()
+{
+ struct foo {
+ int *a, *b, c, d, *e;
+ } s;
+
+ int *a, *z;
+
+#pragma acc enter data copyin(s)
+ {
+#pragma acc data copy(s.a[0:10]) copy(z[0:10])
+ {
+ s.e = z;
+#pragma acc parallel loop attach(s.e)
+ for (int i = 0; i < 10; i++)
+ s.a[i] = s.e[i];
+
+
+ a = s.e;
+#pragma acc enter data attach(a)
+#pragma acc exit data detach(a)
+ }
+
+#pragma acc enter data copyin(a)
+#pragma acc acc enter data attach(s.e)
+#pragma acc exit data detach(s.e)
+
+#pragma acc data attach(s.e)
+ {
+ }
+#pragma acc exit data delete(a)
+
+#pragma acc exit data detach(a) finalize
+#pragma acc exit data detach(s.a) finalize
+ }
+}
+
+/* { 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.attach:s.a .bias: 0.." 1 \
"omplower" } } */ +/* { dg-final { scan-tree-dump-times "pragma omp target \
oacc_parallel map.attach:s.e .bias: 8.. map.tofrom:s .len: 32" 1 "omplower" } } */ \
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data \
map.attach:a .bias: 8.." 1 "omplower" } } */ +/* { dg-final { scan-tree-dump-times \
"pragma omp target oacc_enter_exit_data map.detach:a .bias: 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 .bias: 8.." 1 "omplower" } } */ +/* { \
dg-final { scan-tree-dump-times "pragma omp target oacc_data map.attach:s.e .bias: \
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 .bias: 8.." 1 "omplower" } } */ +/* { dg-final { \
scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize \
map.force_detach:s.a .bias: 8.." 1 "omplower" } } */
diff --git a/gcc/testsuite/c-c++-common/goacc/mdc-2.c \
b/gcc/testsuite/c-c++-common/goacc/mdc-2.c new file mode 100644
index 00000000000..fae86671fc9
--- /dev/null
+++ b/gcc/testsuite/c-c++-common/goacc/mdc-2.c
@@ -0,0 +1,62 @@
+/* Test OpenACC's support for manual deep copy, including the attach
+ and detach clauses. */
+
+void
+t1 ()
+{
+ struct foo {
+ int *a, *b, c, d, *e;
+ } s;
+
+ int *a, *z, scalar, **y;
+
+#pragma acc enter data copyin(s) detach(z) /* { dg-error ".detach. is not valid for" \
} */ + {
+#pragma acc data copy(s.a[0:10]) copy(z[0:10])
+ {
+ s.e = z;
+#pragma acc parallel loop attach(s.e) detach(s.b) /* { dg-error ".detach. is not \
valid for" } */ + for (int i = 0; i < 10; i++)
+ s.a[i] = s.e[i];
+
+ a = s.e;
+#pragma acc enter data attach(a) detach(s.c) /* { dg-error ".detach. is not valid \
for" } */ +#pragma acc exit data detach(a)
+ }
+
+#pragma acc enter data attach(z[:5]) /* { dg-error "expected single pointer in \
.attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* \
} .-1 } */ +#pragma acc exit data detach(z[:5]) /* { dg-error "expected single \
pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { \
target *-*-* } .-1 } */ +#pragma acc enter data attach(z[1:]) /* { dg-error "expected \
single pointer in .attach. clause" } */ +/* { dg-error "has no data movement clause" \
"" { target *-*-* } .-1 } */ +#pragma acc exit data detach(z[1:]) /* { dg-error \
"expected single pointer in .detach. clause" } */ +/* { dg-error "has no data \
movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data attach(z[:]) /* \
{ dg-error "expected single pointer in .attach. clause" } */ +/* { dg-error "has no \
data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data \
detach(z[:]) /* { dg-error "expected single pointer in .detach. clause" } */ +/* { \
dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc \
enter data attach(z[3]) /* { dg-error "expected pointer in .attach. clause" } */ \
+#pragma acc exit data detach(z[3]) /* { dg-error "expected pointer in .detach. \
clause" } */ +
+#pragma acc acc enter data attach(s.e)
+#pragma acc exit data detach(s.e) attach(z) /* { dg-error ".attach. is not valid \
for" } */ +
+#pragma acc data attach(s.e)
+ {
+ }
+#pragma acc exit data delete(a) attach(s.a) /* { dg-error ".attach. is not valid \
for" } */ +
+#pragma acc enter data attach(scalar) /* { dg-error "expected pointer in .attach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc exit data detach(scalar) /* { dg-error "expected pointer in .detach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc enter data attach(s) /* { dg-error "expected pointer in .attach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc exit data detach(s) /* { dg-error "expected pointer in .detach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ + }
+
+#pragma acc enter data attach(y[10])
+#pragma acc exit data detach(y[10])
+}
diff --git a/gcc/testsuite/g++.dg/goacc/mdc.C b/gcc/testsuite/g++.dg/goacc/mdc.C
new file mode 100644
index 00000000000..b3abab30423
--- /dev/null
+++ b/gcc/testsuite/g++.dg/goacc/mdc.C
@@ -0,0 +1,68 @@
+/* Test OpenACC's support for manual deep copy, including the attach
+ and detach clauses. */
+
+void
+t1 ()
+{
+ struct foo {
+ int *a, *b, c, d, *e;
+ } s;
+
+ struct foo& rs = s;
+
+ int *a, *z, scalar, **y;
+ int* const &ra = a;
+ int* const &rz = z;
+ int& rscalar = scalar;
+ int** const &ry = y;
+
+#pragma acc enter data copyin(rs) detach(rz) /* { dg-error ".detach. is not valid \
for" } */ + {
+#pragma acc data copy(rs.a[0:10]) copy(rz[0:10])
+ {
+ s.e = z;
+#pragma acc parallel loop attach(rs.e) detach(rs.b) /* { dg-error ".detach. is not \
valid for" } */ + for (int i = 0; i < 10; i++)
+ s.a[i] = s.e[i];
+
+ a = s.e;
+#pragma acc enter data attach(ra) detach(rs.c) /* { dg-error ".detach. is not valid \
for" } */ +#pragma acc exit data detach(ra)
+ }
+
+#pragma acc enter data attach(rz[:5]) /* { dg-error "expected single pointer in \
.attach. clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* \
} .-1 } */ +#pragma acc exit data detach(rz[:5]) /* { dg-error "expected single \
pointer in .detach. clause" } */ +/* { dg-error "has no data movement clause" "" { \
target *-*-* } .-1 } */ +#pragma acc enter data attach(rz[1:]) /* { dg-error \
"expected single pointer in .attach. clause" } */ +/* { dg-error "has no data \
movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit data detach(rz[1:]) \
/* { dg-error "expected single pointer in .detach. clause" } */ +/* { dg-error "has \
no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc enter data \
attach(rz[:]) /* { dg-error "expected single pointer in .attach. clause" } */ +/* { \
dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma acc exit \
data detach(rz[:]) /* { dg-error "expected single pointer in .detach. clause" } */ \
+/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } */ +#pragma \
acc enter data attach(rz[3]) /* { dg-error "expected pointer in .attach. clause" } */ \
+#pragma acc exit data detach(rz[3]) /* { dg-error "expected pointer in .detach. \
clause" } */ +
+#pragma acc acc enter data attach(rs.e)
+#pragma acc exit data detach(rs.e) attach(rz) /* { dg-error ".attach. is not valid \
for" } */ +
+#pragma acc data attach(rs.e)
+ {
+ }
+#pragma acc exit data delete(ra) attach(rs.a) /* { dg-error ".attach. is not valid \
for" } */ +
+#pragma acc enter data attach(rscalar) /* { dg-error "expected pointer in .attach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc exit data detach(rscalar) /* { dg-error "expected pointer in .detach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc enter data attach(rs) /* { dg-error "expected pointer in .attach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ +#pragma acc exit data detach(rs) /* { dg-error "expected pointer in .detach. \
clause" } */ +/* { dg-error "has no data movement clause" "" { target *-*-* } .-1 } \
*/ + }
+
+#pragma acc enter data attach(ry[10])
+#pragma acc exit data detach(ry[10])
+}
diff --git a/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 \
b/gcc/testsuite/gfortran.dg/goacc/data-clauses.f95 index b94214e8b63..1a4a6719987 \
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-2.f90 \
b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90 new file mode 100644
index 00000000000..d01583fac89
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/derived-types-2.f90
@@ -0,0 +1,14 @@
+module bar
+ type :: type1
+ real(8), pointer, public :: p(:) => null()
+ end type
+ type :: type2
+ class(type1), pointer :: p => null()
+ end type
+end module
+
+subroutine foo (var)
+ use bar
+ type(type2), intent(inout) :: var
+ !$acc enter data create(var%p%p)
+end subroutine
diff --git a/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 \
b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90 new file mode 100644
index 00000000000..5fb29816c42
--- /dev/null
+++ b/gcc/testsuite/gfortran.dg/goacc/derived-types.f90
@@ -0,0 +1,77 @@
+! Test ACC UPDATE with derived types.
+
+module dt
+ integer, parameter :: n = 10
+ type inner
+ integer :: d(n)
+ end type inner
+ type dtype
+ integer(8) :: a, b, c(n)
+ type(inner) :: in
+ end type dtype
+end module dt
+
+program derived_acc
+ use dt
+
+ implicit none
+ type(dtype):: var
+ integer i
+ !$acc declare create(var)
+ !$acc declare pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+
+ !$acc update host(var)
+ !$acc update host(var%a)
+ !$acc update device(var)
+ !$acc update device(var%a)
+ !$acc update self(var)
+ !$acc update self(var%a)
+
+ !$acc enter data copyin(var)
+ !$acc enter data copyin(var%a)
+
+ !$acc exit data copyout(var)
+ !$acc exit data copyout(var%a)
+
+ !$acc data copy(var)
+ !$acc end data
+
+ !$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)
+ do i = 1, 10
+ end do
+ !$acc end parallel loop
+
+ !$acc parallel pcopy(var)
+ !$acc end parallel
+
+ !$acc parallel pcopy(var%a)
+ do i = 1, 10
+ end do
+ !$acc end parallel
+
+ !$acc kernels pcopyin(var)
+ !$acc end kernels
+
+ !$acc kernels pcopy(var%a)
+ do i = 1, 10
+ end do
+ !$acc end kernels
+
+ !$acc kernels loop pcopyin(var)
+ do i = 1, 10
+ end do
+ !$acc end kernels loop
+
+ !$acc kernels loop pcopy(var%a)
+ do i = 1, 10
+ end do
+ !$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 a414df8d439..c2a49796318 \
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/gcc/tree-pretty-print.c b/gcc/tree-pretty-print.c
index 1cf7a912133..379858d0f1f 100644
--- a/gcc/tree-pretty-print.c
+++ b/gcc/tree-pretty-print.c
@@ -849,6 +849,18 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, \
dump_flags_t flags) case GOMP_MAP_LINK:
pp_string (pp, "link");
break;
+ case GOMP_MAP_ATTACH:
+ pp_string (pp, "attach");
+ break;
+ case GOMP_MAP_DETACH:
+ pp_string (pp, "detach");
+ break;
+ case GOMP_MAP_FORCE_DETACH:
+ pp_string (pp, "force_detach");
+ break;
+ case GOMP_MAP_ATTACH_DETACH:
+ pp_string (pp, "attach_detach");
+ break;
default:
gcc_unreachable ();
}
@@ -870,6 +882,12 @@ dump_omp_clause (pretty_printer *pp, tree clause, int spc, \
dump_flags_t flags) case GOMP_MAP_TO_PSET:
pp_string (pp, " [pointer set, len: ");
break;
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
+ case GOMP_MAP_ATTACH_DETACH:
+ pp_string (pp, " [bias: ");
+ break;
default:
pp_string (pp, " [len: ");
break;
diff --git a/include/gomp-constants.h b/include/gomp-constants.h
index 9e356cdfeec..f40d6069582 100644
--- a/include/gomp-constants.h
+++ b/include/gomp-constants.h
@@ -40,8 +40,11 @@
#define GOMP_MAP_FLAG_SPECIAL_0 (1 << 2)
#define GOMP_MAP_FLAG_SPECIAL_1 (1 << 3)
#define GOMP_MAP_FLAG_SPECIAL_2 (1 << 4)
+#define GOMP_MAP_FLAG_SPECIAL_4 (1 << 6)
#define GOMP_MAP_FLAG_SPECIAL (GOMP_MAP_FLAG_SPECIAL_1 \
| GOMP_MAP_FLAG_SPECIAL_0)
+#define GOMP_MAP_DEEP_COPY (GOMP_MAP_FLAG_SPECIAL_4 \
+ | GOMP_MAP_FLAG_SPECIAL_2)
/* Flag to force a specific behavior (or else, trigger a run-time error). */
#define GOMP_MAP_FLAG_FORCE (1 << 7)
@@ -127,12 +130,23 @@ enum gomp_map_kind
/* Decrement usage count and deallocate if zero. */
GOMP_MAP_RELEASE = (GOMP_MAP_FLAG_SPECIAL_2
| GOMP_MAP_DELETE),
+ /* In OpenACC, attach a pointer to a mapped struct field. */
+ GOMP_MAP_ATTACH = (GOMP_MAP_DEEP_COPY | 0),
+ /* In OpenACC, detach a pointer to a mapped struct field. */
+ GOMP_MAP_DETACH = (GOMP_MAP_DEEP_COPY | 1),
+ /* In OpenACC, detach a pointer to a mapped struct field. */
+ GOMP_MAP_FORCE_DETACH = (GOMP_MAP_DEEP_COPY
+ | GOMP_MAP_FLAG_FORCE | 1),
/* Internal to GCC, not used in libgomp. */
/* Do not map, but pointer assign a pointer instead. */
GOMP_MAP_FIRSTPRIVATE_POINTER = (GOMP_MAP_LAST | 1),
/* Do not map, but pointer assign a reference instead. */
- GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2)
+ GOMP_MAP_FIRSTPRIVATE_REFERENCE = (GOMP_MAP_LAST | 2),
+ /* An attach or detach operation. Rewritten to the appropriate type during
+ gimplification, depending on directive (i.e. "enter data" or
+ parallel/kernels region vs. "exit data"). */
+ GOMP_MAP_ATTACH_DETACH = (GOMP_MAP_LAST | 3)
};
#define GOMP_MAP_COPY_TO_P(X) \
diff --git a/libgomp/libgomp.h b/libgomp/libgomp.h
index bcbc903e32f..0631e5e26f0 100644
--- a/libgomp/libgomp.h
+++ b/libgomp/libgomp.h
@@ -948,6 +948,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. */
@@ -997,6 +999,16 @@ struct target_mem_desc {
#define OFFSET_POINTER (~(uintptr_t) 1)
#define OFFSET_STRUCT (~(uintptr_t) 2)
+/* Auxiliary structure for infrequently-used or API-specific data. */
+
+struct splay_tree_aux {
+ /* Pointer to the original mapping of "omp declare target link" object. */
+ splay_tree_key link_key;
+ /* For a block with attached pointers, the attachment counters for each.
+ Only used for OpenACC. */
+ uintptr_t *attach_count;
+};
+
struct splay_tree_key_s {
/* Address of the host object. */
uintptr_t host_start;
@@ -1017,8 +1029,7 @@ struct splay_tree_key_s {
/* The recalculated reference count, for verification. */
uintptr_t refcount_chk;
#endif
- /* Pointer to the original mapping of "omp declare target link" object. */
- splay_tree_key link_key;
+ struct splay_tree_aux *aux;
};
/* The comparison function. */
@@ -1164,6 +1175,13 @@ extern void gomp_copy_dev2host (struct gomp_device_descr *,
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 *);
#ifdef RC_CHECKING
extern void dump_tgt (const char *, struct target_mem_desc *);
diff --git a/libgomp/libgomp.map b/libgomp/libgomp.map
index c79430f8d8d..63276f7d29b 100644
--- a/libgomp/libgomp.map
+++ b/libgomp/libgomp.map
@@ -484,6 +484,16 @@ OACC_2.5.1 {
acc_register_library;
} OACC_2.5;
+OACC_2.6 {
+ global:
+ acc_attach;
+ acc_attach_async;
+ acc_detach;
+ acc_detach_async;
+ acc_detach_finalize;
+ acc_detach_finalize_async;
+} OACC_2.5.1;
+
GOACC_2.0 {
global:
GOACC_data_end;
diff --git a/libgomp/oacc-init.c b/libgomp/oacc-init.c
index 495e5391b2c..b5096fd813f 100644
--- a/libgomp/oacc-init.c
+++ b/libgomp/oacc-init.c
@@ -344,7 +344,8 @@ acc_shutdown_1 (acc_device_t d)
while (walk->dev->mem_map.root)
{
splay_tree_key k = &walk->dev->mem_map.root->key;
- k->link_key = NULL;
+ if (k->aux)
+ k->aux->link_key = NULL;
gomp_remove_var (walk->dev, k);
}
diff --git a/libgomp/oacc-mem.c b/libgomp/oacc-mem.c
index 34795437b21..4c6d7cd2cac 100644
--- a/libgomp/oacc-mem.c
+++ b/libgomp/oacc-mem.c
@@ -845,3 +845,87 @@ acc_update_self_async (void *h, size_t s, int async)
{
update_dev_host (0, h, s, async);
}
+
+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;
+
+ gomp_mutex_lock (&acc_dev->lock);
+
+ 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);
+
+ gomp_mutex_unlock (&acc_dev->lock);
+}
+
+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;
+
+ gomp_mutex_lock (&acc_dev->lock);
+
+ 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);
+
+ gomp_mutex_unlock (&acc_dev->lock);
+}
+
+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 6332c0a5b03..3188f9d9163 100644
--- a/libgomp/oacc-parallel.c
+++ b/libgomp/oacc-parallel.c
@@ -53,33 +53,48 @@ _Static_assert (GOACC_FLAGS_UNMARSHAL (GOMP_DEVICE_HOST_FALLBACK)
mappings. */
static int
-find_group_last (int pos, size_t mapnum, unsigned short *kinds)
+find_group_last (int pos, size_t mapnum, size_t *sizes, unsigned short *kinds)
{
unsigned char kind0 = kinds[pos] & 0xff;
- int first_pos = pos, last_pos = pos;
+ int first_pos = pos;
- if (kind0 == GOMP_MAP_TO_PSET)
+ switch (kind0)
{
+ case GOMP_MAP_TO_PSET:
while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
- last_pos = ++pos;
+ pos++;
/* We expect at least one GOMP_MAP_POINTER after a GOMP_MAP_TO_PSET. */
- assert (last_pos > first_pos);
- }
- else
- {
+ assert (pos > first_pos);
+ break;
+
+ case GOMP_MAP_STRUCT:
+ pos += sizes[pos];
+ break;
+
+ case GOMP_MAP_POINTER:
+ case GOMP_MAP_ALWAYS_POINTER:
+ /* These mappings are only expected after some other mapping. If we
+ see one by itself, something has gone wrong. */
+ gomp_fatal ("unexpected mapping");
+ break;
+
+ default:
/* GOMP_MAP_ALWAYS_POINTER can only appear directly after some other
mapping. */
- if (pos + 1 < mapnum
- && (kinds[pos + 1] & 0xff) == GOMP_MAP_ALWAYS_POINTER)
- return pos + 1;
+ if (pos + 1 < mapnum)
+ {
+ unsigned char kind1 = kinds[pos + 1] & 0xff;
+ if (kind1 == GOMP_MAP_ALWAYS_POINTER)
+ return pos + 1;
+ }
- /* We can have one or several GOMP_MAP_POINTER mappings after a to/from
+ /* We can have zero or more GOMP_MAP_POINTER mappings after a to/from
(etc.) mapping. */
while (pos + 1 < mapnum && (kinds[pos + 1] & 0xff) == GOMP_MAP_POINTER)
- last_pos = ++pos;
+ pos++;
}
- return last_pos;
+ return pos;
}
/* Handle the mapping pair that are presented when a
@@ -630,7 +645,7 @@ goacc_enter_data_internal (struct gomp_device_descr *acc_dev, \
size_t mapnum, {
for (size_t i = 0; i < mapnum; i++)
{
- int group_last = find_group_last (i, mapnum, kinds);
+ int group_last = find_group_last (i, mapnum, sizes, kinds);
gomp_map_vars_async (acc_dev, aq,
(group_last - i) + 1,
@@ -652,6 +667,33 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, \
size_t mapnum, {
gomp_mutex_lock (&acc_dev->lock);
+ /* Handle "detach" before copyback/deletion of mapped data. */
+ for (size_t i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+ switch (kind)
+ {
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
+ {
+ struct splay_tree_key_s cur_node;
+ uintptr_t hostaddr = (uintptr_t) hostaddrs[i];
+ cur_node.host_start = hostaddr;
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ splay_tree_key n
+ = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ gomp_fatal ("struct not mapped for detach operation");
+
+ gomp_detach_pointer (acc_dev, aq, n, hostaddr, finalize, NULL);
+ }
+ break;
+ default:
+ ;
+ }
+ }
+
for (size_t i = 0; i < mapnum; ++i)
{
unsigned char kind = kinds[i] & 0xff;
@@ -669,6 +711,8 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, \
size_t mapnum, case GOMP_MAP_POINTER:
case GOMP_MAP_DELETE:
case GOMP_MAP_RELEASE:
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
{
struct splay_tree_key_s cur_node;
cur_node.host_start = (uintptr_t) hostaddrs[i];
@@ -712,9 +756,42 @@ goacc_exit_data_internal (struct gomp_device_descr *acc_dev, \
size_t mapnum, gomp_remove_var_async (acc_dev, n, aq);
}
break;
+
+ case GOMP_MAP_STRUCT:
+ {
+ int elems = sizes[i];
+ for (int j = 1; j <= elems; j++)
+ {
+ struct splay_tree_key_s k;
+ k.host_start = (uintptr_t) hostaddrs[i + j];
+ k.host_end = k.host_start + sizes[i + j];
+ splay_tree_key str;
+ str = splay_tree_lookup (&acc_dev->mem_map, &k);
+ if (str)
+ {
+ if (finalize)
+ {
+ str->refcount -= str->virtual_refcount;
+ str->virtual_refcount = 0;
+ }
+ if (str->virtual_refcount > 0)
+ {
+ str->refcount--;
+ str->virtual_refcount--;
+ }
+ else if (str->refcount > 0)
+ str->refcount--;
+ if (str->refcount == 0)
+ gomp_remove_var_async (acc_dev, str, aq);
+ }
+ }
+ i += elems;
+ }
+ break;
+
default:
gomp_fatal (">>>> goacc_exit_data_internal UNHANDLED kind 0x%.2x",
- kind);
+ kind);
}
}
@@ -744,8 +821,13 @@ GOACC_enter_exit_data (int flags_m, 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)
+ || kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_FORCE_DETACH)
finalize = true;
}
@@ -754,11 +836,14 @@ GOACC_enter_exit_data (int flags_m, 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)
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)
@@ -769,6 +854,8 @@ GOACC_enter_exit_data (int flags_m, 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)
break;
diff --git a/libgomp/openacc.h b/libgomp/openacc.h
index 42c861caabf..d2e5c101f7f 100644
--- a/libgomp/openacc.h
+++ b/libgomp/openacc.h
@@ -109,12 +109,18 @@ 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;
/* Finalize versions of copyout/delete functions, specified in OpenACC 2.5. */
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;
/* Async functions, specified in OpenACC 2.5. */
void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW;
diff --git a/libgomp/target.c b/libgomp/target.c
index c86a7404531..b1d7f02fa75 100644
--- a/libgomp/target.c
+++ b/libgomp/target.c
@@ -379,8 +379,13 @@ dump_tgt (const char *where, struct target_mem_desc *tgt)
(int) tgt->list[i].key->refcount);
fprintf (stderr, " key.virtual_refcount=%d\n",
(int) tgt->list[i].key->virtual_refcount);
- fprintf (stderr, " key.link_key=%p\n",
- (void*) tgt->list[i].key->link_key);
+ if (tgt->list[i].key->aux)
+ {
+ fprintf (stderr, " key.aux->link_key=%p\n",
+ (void*) tgt->list[i].key->aux->link_key);
+ fprintf (stderr, " key.aux->attach_count=%p\n",
+ (void*) tgt->list[i].key->aux->attach_count);
+ }
}
}
fprintf (stderr, "\n");
@@ -535,6 +540,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 = kind == GOMP_MAP_ATTACH;
tgt_var->offset = newn->host_start - oldn->host_start;
tgt_var->length = newn->host_end - newn->host_start;
@@ -668,6 +674,134 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
(void *) cur_node.host_end);
}
+attribute_hidden void
+gomp_attach_pointer (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, splay_tree mem_map,
+ splay_tree_key n, uintptr_t attach_to, size_t bias,
+ struct gomp_coalesce_buf *cbufp)
+{
+ struct 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->aux)
+ n->aux = gomp_malloc_cleared (sizeof (struct splay_tree_aux));
+
+ if (!n->aux->attach_count)
+ n->aux->attach_count
+ = gomp_malloc_cleared (sizeof (*n->aux->attach_count) * size);
+
+ if (n->aux->attach_count[idx] < UINTPTR_MAX)
+ n->aux->attach_count[idx]++;
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count overflow");
+ }
+
+ if (n->aux->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, (int) n->aux->attach_count[idx]);
+}
+
+attribute_hidden 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->aux || !n->aux->attach_count)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("no attachment counters for struct");
+ }
+
+ if (finalize)
+ n->aux->attach_count[idx] = 1;
+
+ if (n->aux->attach_count[idx] == 0)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count underflow");
+ }
+ else
+ n->aux->attach_count[idx]--;
+
+ if (n->aux->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, (int) n->aux->attach_count[idx]);
+}
+
attribute_hidden uintptr_t
gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
{
@@ -845,8 +979,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
has_firstprivate = true;
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))
+ 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 *);
@@ -1070,6 +1211,32 @@ gomp_map_vars_internal (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
+ = (pragma_kind != GOMP_MAP_VARS_OPENACC_ENTER_DATA);
+ n->refcount++;
+ }
+ 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;
}
@@ -1085,13 +1252,15 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
kind & typemask, cbufp);
else
{
- k->link_key = NULL;
+ k->aux = NULL;
if (n && n->refcount == REFCOUNT_LINK)
{
/* Replace target address of the pointer with target address
of mapped object in the splay tree. */
splay_tree_remove (mem_map, n);
- k->link_key = n;
+ k->aux
+ = gomp_malloc_cleared (sizeof (struct splay_tree_aux));
+ k->aux->link_key = n;
}
size_t align = (size_t) 1 << (kind >> rshift);
tgt->list[i].key = k;
@@ -1112,10 +1281,12 @@ gomp_map_vars_internal (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->virtual_refcount = 0;
+ k->aux = NULL;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1166,6 +1337,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
tgt->list[j].key = k;
tgt->list[j].copy_from = false;
tgt->list[j].always_copy_from = false;
+ tgt->list[j].do_detach = false;
if (k->refcount != REFCOUNT_INFINITY)
k->refcount++;
gomp_map_pointer (tgt, aq,
@@ -1209,7 +1381,7 @@ gomp_map_vars_internal (struct gomp_device_descr *devicep,
kind);
}
- if (k->link_key)
+ if (k->aux && k->aux->link_key)
{
/* Set link pointer on target to the device address of the
mapped object. */
@@ -1338,8 +1510,15 @@ gomp_remove_var_internal (struct gomp_device_descr *devicep, \
splay_tree_key k, {
bool is_tgt_unmapped = false;
splay_tree_remove (&devicep->mem_map, k);
- if (k->link_key)
- splay_tree_insert (&devicep->mem_map, (splay_tree_node) k->link_key);
+ if (k->aux)
+ {
+ if (k->aux->link_key)
+ splay_tree_insert (&devicep->mem_map,
+ (splay_tree_node) k->aux->link_key);
+ if (k->aux->attach_count)
+ free (k->aux->attach_count);
+ free (k->aux);
+ }
if (aq)
devicep->openacc.async.queue_callback_func (aq, gomp_unref_tgt_void,
(void *) k->tgt);
@@ -1393,6 +1572,18 @@ gomp_unmap_vars_internal (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 && tgt->list[i].do_detach)
+ gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
+ + tgt->list[i].offset,
+ k->refcount == 1, NULL);
+ }
+
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
@@ -1565,7 +1756,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, \
unsigned version, k->tgt_offset = target_table[i].start;
k->refcount = REFCOUNT_INFINITY;
k->virtual_refcount = 0;
- k->link_key = NULL;
+ k->aux = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
@@ -1598,7 +1789,7 @@ gomp_load_image_to_device (struct gomp_device_descr *devicep, \
unsigned version, k->tgt_offset = target_var->start;
k->refcount = target_size & link_bit ? REFCOUNT_LINK : REFCOUNT_INFINITY;
k->virtual_refcount = 0;
- k->link_key = NULL;
+ k->aux = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
@@ -2313,9 +2504,13 @@ gomp_exit_data (struct gomp_device_descr *devicep, size_t \
mapnum, if (k->refcount == 0)
{
splay_tree_remove (&devicep->mem_map, k);
- if (k->link_key)
- splay_tree_insert (&devicep->mem_map,
- (splay_tree_node) k->link_key);
+ if (k->aux)
+ {
+ if (k->aux->link_key)
+ splay_tree_insert (&devicep->mem_map,
+ (splay_tree_node) k->aux->link_key);
+ free (k->aux);
+ }
assert (k->tgt->refcount != REFCOUNT_INFINITY);
if (k->tgt->refcount > 1)
k->tgt->refcount--;
@@ -2854,7 +3049,7 @@ omp_target_associate_ptr (const void *host_ptr, const void \
*device_ptr, k->tgt_offset = (uintptr_t) device_ptr + device_offset;
k->refcount = REFCOUNT_INFINITY;
k->virtual_refcount = 0;
- k->link_key = NULL;
+ k->aux = NULL;
array->left = NULL;
array->right = NULL;
splay_tree_insert (&devicep->mem_map, array);
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C \
b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C new file mode 100644
index 00000000000..a512008685d
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-12.C
@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test attach/detach with dereferences of reference to pointer to struct. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ mystruct *&mref = m;
+ int i;
+
+ mref->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ mref->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(mref->a[0:N])
+ for (j = 0; j < N; j++)
+ mref->a[j]++;
+#pragma acc parallel loop copy(mref->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ mref->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C \
b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C new file mode 100644
index 00000000000..a5194568603
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c++/deep-copy-13.C
@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test array slice with reference to pointer. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ int *&ptr = m->a;
+#pragma acc parallel loop copy(ptr[0:N])
+ for (j = 0; j < N; j++)
+ ptr[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ m->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
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 00000000000..d8d7067e452
--- /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-10.c \
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c new file mode 100644
index 00000000000..573a8214bf0
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-10.c
@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+/* Test asyncronous attach and detach operation. */
+
+typedef struct {
+ int *a;
+ int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct m;
+ int i;
+
+ m.a = (int *) malloc (N * sizeof (int));
+ m.b = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m.a[i] = 0;
+ m.b[i] = 0;
+ }
+
+#pragma acc enter data copyin(m)
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(m.a[0:N]) async(i % 2)
+ for (j = 0; j < N; j++)
+ m.a[j]++;
+#pragma acc parallel loop copy(m.b[0:N]) async((i + 1) % 2)
+ for (j = 0; j < N; j++)
+ m.b[j]++;
+ }
+
+#pragma acc exit data copyout(m) wait(0, 1)
+
+ for (i = 0; i < N; i++)
+ {
+ if (m.a[i] != 99)
+ abort ();
+ if (m.b[i] != 99)
+ abort ();
+ }
+
+ free (m.a);
+ free (m.b);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c \
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c new file mode 100644
index 00000000000..db6012fb352
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-11.c
@@ -0,0 +1,72 @@
+#include <stdlib.h>
+
+/* Test multiple struct dereferences on one directive, and slices starting at
+ non-zero. */
+
+typedef struct {
+ int *a;
+ int *b;
+ int *c;
+} mystruct;
+
+int main(int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+ m->c = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ m->c[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+#pragma acc parallel loop copy(m->a[0:N])
+ for (j = 0; j < N; j++)
+ m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N], m->c[5:N-10])
+ for (j = 0; j < N; j++)
+ {
+ m->b[j]++;
+ if (j > 5 && j < N - 5)
+ m->c[j]++;
+ }
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ if (i > 5 && i < N-5)
+ {
+ if (m->c[i] != 99)
+ abort ();
+ }
+ else
+ {
+ if (m->c[i] != 0)
+ abort ();
+ }
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m->c);
+ free (m);
+
+ return 0;
+}
diff --git a/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c \
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c new file mode 100644
index 00000000000..275fa9ae256
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-14.c
@@ -0,0 +1,63 @@
+#include <openacc.h>
+#include <stdlib.h>
+
+/* Test attach/detach operation with chained dereferences. */
+
+typedef struct mystruct {
+ int *a;
+ struct mystruct *next;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->next = (mystruct *) malloc (sizeof (*m));
+ m->next->a = (int *) malloc (N * sizeof (int));
+ m->next->next = NULL;
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->next->a[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+ acc_copyin (m->next, sizeof (*m));
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ acc_copyin (m->next->a, N * sizeof (int));
+ acc_attach ((void **) &m->next);
+ /* This will attach only the innermost pointer, i.e. "a[0:N]". That's
+ why we have to attach the "m->next" pointer manually above. */
+#pragma acc parallel loop copy(m->next->a[0:N])
+ for (j = 0; j < N; j++)
+ m->next->a[j]++;
+ acc_detach ((void **) &m->next);
+ acc_copyout (m->next->a, N * sizeof (int));
+ }
+
+ acc_copyout (m->next, sizeof (*m));
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 0)
+ abort ();
+ if (m->next->a[i] != 99)
+ abort ();
+ }
+
+ free (m->next->a);
+ free (m->next);
+ free (m->a);
+ free (m);
+
+ 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 00000000000..7e26e9aa8b9
--- /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 00000000000..cec764bd3e7
--- /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 00000000000..8874ca0a504
--- /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 00000000000..89cafbb62ab
--- /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-c-c++-common/deep-copy-6.c \
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-6.c new file mode 100644
index 00000000000..391149459c9
--- /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 00000000000..a59047af520
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-7.c
@@ -0,0 +1,45 @@
+/* { 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++)
+ {
+ /* Here, we do not explicitly copy the enclosing structure, but work
+ with fields directly. Make sure attachment counters and reference
+ counters work properly in that case. */
+#pragma acc enter data copyin(v.a, v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+#pragma acc enter data pcopyin(v.b[0:n])
+
+#pragma acc parallel loop present(v.a, v.b)
+ for (i = 0; i < n; i++)
+ v.b[i] = v.a + i;
+
+#pragma acc exit data copyout(v.b[:n]) finalize
+#pragma acc exit data delete(v.a)
+
+ 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 00000000000..0ca5990b377
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-8.c
@@ -0,0 +1,54 @@
+/* { 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])
+#pragma acc exit data copyout(v.a)
+
+ 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-c-c++-common/deep-copy-9.c \
b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c new file mode 100644
index 00000000000..e86a46bd84a
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-c-c++-common/deep-copy-9.c
@@ -0,0 +1,53 @@
+#include <stdlib.h>
+
+typedef struct {
+ int *a;
+ int *b;
+} mystruct;
+
+int
+main (int argc, char* argv[])
+{
+ const int N = 1024;
+ mystruct *m = (mystruct *) malloc (sizeof (*m));
+ int i;
+
+ m->a = (int *) malloc (N * sizeof (int));
+ m->b = (int *) malloc (N * sizeof (int));
+
+ for (i = 0; i < N; i++)
+ {
+ m->a[i] = 0;
+ m->b[i] = 0;
+ }
+
+#pragma acc enter data copyin(m[0:1])
+
+ for (int i = 0; i < 99; i++)
+ {
+ int j;
+ int *ptr = m->a;
+#pragma acc parallel loop copy(m->a[0:N])
+ for (j = 0; j < N; j++)
+ m->a[j]++;
+#pragma acc parallel loop copy(m->b[0:N])
+ for (j = 0; j < N; j++)
+ m->b[j]++;
+ }
+
+#pragma acc exit data copyout(m[0:1])
+
+ for (i = 0; i < N; i++)
+ {
+ if (m->a[i] != 99)
+ abort ();
+ if (m->b[i] != 99)
+ abort ();
+ }
+
+ free (m->a);
+ free (m->b);
+ free (m);
+
+ 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 00000000000..c4cea11b571
--- /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 00000000000..35936617b87
--- /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 00000000000..667d944fecb
--- /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 00000000000..6949e120c9f
--- /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 00000000000..6843cf1d0fa
--- /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 00000000000..12910d0d655
--- /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 00000000000..ab44f0a73b9
--- /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 00000000000..d142763ae59
--- /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 new file mode 100644
index 00000000000..eb7812d541e
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derived-type-1.f90
@@ -0,0 +1,28 @@
+! Test derived types with subarrays
+
+! { dg-do run }
+
+ implicit none
+ type dtype
+ integer :: a, b, c
+ end type dtype
+ integer, parameter :: n = 100
+ integer i
+ type (dtype), dimension(n) :: d
+
+ !$acc data copy(d(1:n))
+ !$acc parallel loop
+ do i = 1, n
+ d(i)%a = i
+ d(i)%b = i-1
+ d(i)%c = i+1
+ end do
+ !$acc end data
+
+ do i = 1, n
+ 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/derivedtype-1.f95 \
b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95 new file mode 100644
index 00000000000..75ce48ddca2
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-1.f95
@@ -0,0 +1,30 @@
+! { dg-do run }
+
+program main
+ implicit none
+
+ type mytype
+ integer :: a, b, c
+ end type mytype
+
+ type(mytype) :: myvar
+ integer :: i
+
+ myvar%a = 0
+ myvar%b = 0
+ myvar%c = 0
+
+!$acc enter data copyin(myvar)
+
+!$acc parallel present(myvar)
+ myvar%a = 1
+ myvar%b = 2
+ myvar%c = 3
+!$acc end parallel
+
+!$acc exit data copyout(myvar)
+
+ if (myvar%a .ne. 1) stop 1
+ if (myvar%b .ne. 2) stop 2
+ if (myvar%c .ne. 3) stop 3
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 \
b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95 new file mode 100644
index 00000000000..3088b832957
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/derivedtype-2.f95
@@ -0,0 +1,41 @@
+! { dg-do run }
+
+program main
+ implicit none
+
+ type tnest
+ integer :: ia, ib, ic
+ end type tnest
+
+ type mytype
+ type(tnest) :: nest
+ integer :: a, b, c
+ end type mytype
+
+ type(mytype) :: myvar
+ integer :: i
+
+ myvar%a = 0
+ myvar%b = 0
+ myvar%c = 0
+ myvar%nest%ia = 0
+ myvar%nest%ib = 0
+ myvar%nest%ic = 0
+
+!$acc enter data copyin(myvar%nest)
+
+!$acc parallel present(myvar%nest)
+ myvar%nest%ia = 4
+ myvar%nest%ib = 5
+ myvar%nest%ic = 6
+!$acc end parallel
+
+!$acc exit data copyout(myvar%nest)
+
+ if (myvar%a .ne. 0) stop 1
+ if (myvar%b .ne. 0) stop 2
+ if (myvar%c .ne. 0) stop 3
+ if (myvar%nest%ia .ne. 4) stop 4
+ if (myvar%nest%ib .ne. 5) stop 5
+ if (myvar%nest%ic .ne. 6) stop 6
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 \
b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95 new file mode 100644
index 00000000000..a9b40eeab4c
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/multidim-slice.f95
@@ -0,0 +1,50 @@
+! { dg-do run }
+
+program main
+ implicit none
+ real, allocatable :: myarr(:,:,:,:,:)
+ integer i, j, k, l, m
+
+ allocate(myarr(1:10,1:10,1:10,1:10,1:10))
+
+ do i=1,10
+ do j=1,10
+ do k=1,10
+ do l=1,10
+ do m=1,10
+ myarr(m,l,k,j,i) = i+j+k+l+m
+ end do
+ end do
+ end do
+ end do
+ end do
+
+ do i=1,10
+ !$acc data copy(myarr(:,:,:,:,i))
+ !$acc parallel loop collapse(4) present(myarr(:,:,:,:,i))
+ do j=1,10
+ do k=1,10
+ do l=1,10
+ do m=1,10
+ myarr(m,l,k,j,i) = myarr(m,l,k,j,i) + 1
+ end do
+ end do
+ end do
+ end do
+ !$acc end parallel loop
+ !$acc end data
+ end do
+
+ do i=1,10
+ do j=1,10
+ do k=1,10
+ do l=1,10
+ do m=1,10
+ if (myarr(m,l,k,j,i) .ne. i+j+k+l+m+1) stop 1
+ end do
+ end do
+ end do
+ end do
+ end do
+
+end program main
diff --git a/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 \
b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90 new file mode 100644
index 00000000000..c3c8a07868f
--- /dev/null
+++ b/libgomp/testsuite/libgomp.oacc-fortran/update-2.f90
@@ -0,0 +1,284 @@
+! Test ACC UPDATE with derived types.
+
+! { dg-do run }
+
+module dt
+ integer, parameter :: n = 10
+ type inner
+ integer :: d(n)
+ end type inner
+ type mytype
+ integer(8) :: a, b, c(n)
+ type(inner) :: in
+ end type mytype
+end module dt
+
+program derived_acc
+ use dt
+
+ implicit none
+ integer i, res
+ type(mytype) :: var
+
+ var%a = 0
+ var%b = 1
+ var%c(:) = 10
+ var%in%d(:) = 100
+
+ var%c(:) = 10
+
+ !$acc enter data copyin(var)
+
+ !$acc parallel loop present(var)
+ do i = 1, 1
+ var%a = var%b
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%a)
+
+ if (var%a /= var%b) stop 1
+
+ var%b = 100
+
+ !$acc update device(var%b)
+
+ !$acc parallel loop present(var)
+ do i = 1, 1
+ var%a = var%b
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%a)
+
+ if (var%a /= var%b) stop 2
+
+ !$acc parallel loop present (var)
+ do i = 1, n
+ var%c(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c)
+
+ var%a = -1
+
+ do i = 1, n
+ if (var%c(i) /= i) stop 3
+ var%c(i) = var%a
+ end do
+
+ !$acc update device(var%a)
+ !$acc update device(var%c)
+
+ res = 0
+
+ !$acc parallel loop present(var) reduction(+:res)
+ do i = 1, n
+ if (var%c(i) /= var%a) res = res + 1
+ end do
+
+ if (res /= 0) stop 4
+
+ var%c(:) = 0
+
+ !$acc update device(var%c)
+
+ !$acc parallel loop present(var)
+ do i = 5, 5
+ var%c(i) = 1
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c(5))
+
+ do i = 1, n
+ 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)
+ do i = 1, n
+ var%in%d = var%a
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%in%d)
+
+ do i = 1, n
+ if (var%in%d(i) /= var%a) stop 7
+ end do
+
+ var%c(:) = 0
+
+ !$acc update device(var%c)
+
+ var%c(:) = -1
+
+ !$acc parallel loop present(var)
+ do i = n/2, n
+ var%c(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c(n/2:n))
+
+ do i = 1,n
+ 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
+ !$acc update device(var%in%d)
+
+ !$acc parallel loop present(var)
+ do i = 5, 5
+ var%in%d(i) = 1
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%in%d(5))
+
+ do i = 1, n
+ 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)
+
+ call derived_acc_subroutine(var)
+end program derived_acc
+
+subroutine derived_acc_subroutine(var)
+ use dt
+
+ implicit none
+ integer i, res
+ type(mytype) :: var
+
+ var%a = 0
+ var%b = 1
+ var%c(:) = 10
+ var%in%d(:) = 100
+
+ var%c(:) = 10
+
+ !$acc enter data copyin(var)
+
+ !$acc parallel loop present(var)
+ do i = 1, 1
+ var%a = var%b
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%a)
+
+ if (var%a /= var%b) stop 12
+
+ var%b = 100
+
+ !$acc update device(var%b)
+
+ !$acc parallel loop present(var)
+ do i = 1, 1
+ var%a = var%b
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%a)
+
+ if (var%a /= var%b) stop 13
+
+ !$acc parallel loop present (var)
+ do i = 1, n
+ var%c(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c)
+
+ var%a = -1
+
+ do i = 1, n
+ if (var%c(i) /= i) stop 14
+ var%c(i) = var%a
+ end do
+
+ !$acc update device(var%a)
+ !$acc update device(var%c)
+
+ res = 0
+
+ !$acc parallel loop present(var) reduction(+:res)
+ do i = 1, n
+ if (var%c(i) /= var%a) res = res + 1
+ end do
+
+ if (res /= 0) stop 15
+
+ var%c(:) = 0
+
+ !$acc update device(var%c)
+
+ !$acc parallel loop present(var)
+ do i = 5, 5
+ var%c(i) = 1
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c(5))
+
+ do i = 1, n
+ 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)
+ do i = 1, n
+ var%in%d = var%a
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%in%d)
+
+ do i = 1, n
+ if (var%in%d(i) /= var%a) stop 18
+ end do
+
+ var%c(:) = 0
+
+ !$acc update device(var%c)
+
+ var%c(:) = -1
+
+ !$acc parallel loop present(var)
+ do i = n/2, n
+ var%c(i) = i
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%c(n/2:n))
+
+ do i = 1,n
+ 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
+ !$acc update device(var%in%d)
+
+ !$acc parallel loop present(var)
+ do i = 5, 5
+ var%in%d(i) = 1
+ end do
+ !$acc end parallel loop
+
+ !$acc update host(var%in%d(5))
+
+ do i = 1, n
+ 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)
+end subroutine derived_acc_subroutine
[prev in list] [next in list] [prev in thread] [next in thread]
Configure |
About |
News |
Add a list |
Sponsored by KoreLogic