@@ -11553,9 +11553,12 @@ c_parser_omp_variable_list (c_parser *parser,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
- while (c_parser_next_token_is (parser, CPP_DOT))
+ while (c_parser_next_token_is (parser, CPP_DOT)
+ || c_parser_next_token_is (parser, CPP_DEREF))
{
location_t op_loc = c_parser_peek_token (parser)->location;
+ if (c_parser_next_token_is (parser, CPP_DEREF))
+ t = build_simple_mem_ref (t);
c_parser_consume_token (parser);
if (!c_parser_next_token_is (parser, CPP_NAME))
{
@@ -11679,7 +11682,7 @@ c_parser_omp_var_list_parens (c_parser *parser, enum omp_clause_code kind,
}
/* OpenACC 2.5:
- attach (variable-list )
+ attach ( variable-list )
copy ( variable-list )
copyin ( variable-list )
copyout ( variable-list )
@@ -14090,15 +14093,15 @@ c_parser_oacc_all_clauses (c_parser *parser, omp_clause_mask mask,
clauses = c_parser_oacc_clause_async (parser, clauses);
c_name = "async";
break;
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "attach";
+ break;
case PRAGMA_OACC_CLAUSE_AUTO:
clauses = c_parser_oacc_simple_clause (parser, here, OMP_CLAUSE_AUTO,
clauses);
c_name = "auto";
break;
- case PRAGMA_OACC_CLAUSE_ATTACH:
- clauses = c_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "attach";
- break;
case PRAGMA_OACC_CLAUSE_BIND:
clauses = c_parser_oacc_clause_bind (parser, clauses);
c_name = "bind";
@@ -12446,6 +12446,8 @@ handle_omp_array_sections_1 (tree c, tree t, vec<tree> &types,
}
t = TREE_OPERAND (t, 0);
}
+ if (TREE_CODE (t) == MEM_REF)
+ t = TREE_OPERAND (t, 0);
}
if (!VAR_P (t) && TREE_CODE (t) != PARM_DECL)
{
@@ -13750,6 +13752,8 @@ c_finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
}
if (remove)
break;
+ if (TREE_CODE (t) == MEM_REF)
+ t = TREE_OPERAND (t, 0);
if (VAR_P (t) || TREE_CODE (t) == PARM_DECL)
{
if (bitmap_bit_p (&map_field_head, DECL_UID (t)))
@@ -31563,15 +31563,19 @@ cp_parser_omp_var_list_no_open (cp_parser *parser, enum omp_clause_code kind,
case OMP_CLAUSE_MAP:
case OMP_CLAUSE_FROM:
case OMP_CLAUSE_TO:
- while (cp_lexer_next_token_is (parser->lexer, CPP_DOT))
+ while (cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+ || cp_lexer_next_token_is (parser->lexer, CPP_DEREF))
{
+ cpp_ttype ttype
+ = cp_lexer_next_token_is (parser->lexer, CPP_DOT)
+ ? CPP_DOT : CPP_DEREF;
location_t loc
= cp_lexer_peek_token (parser->lexer)->location;
cp_id_kind idk = CP_ID_KIND_NONE;
cp_lexer_consume_token (parser->lexer);
decl = convert_from_reference (decl);
decl
- = cp_parser_postfix_dot_deref_expression (parser, CPP_DOT,
+ = cp_parser_postfix_dot_deref_expression (parser, ttype,
decl, false,
&idk, loc);
}
@@ -33858,15 +33862,15 @@ cp_parser_oacc_all_clauses (cp_parser *parser, omp_clause_mask mask,
clauses = cp_parser_oacc_clause_async (parser, clauses);
c_name = "async";
break;
+ case PRAGMA_OACC_CLAUSE_ATTACH:
+ clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
+ c_name = "attach";
+ break;
case PRAGMA_OACC_CLAUSE_AUTO:
clauses = cp_parser_oacc_simple_clause (parser, OMP_CLAUSE_AUTO,
clauses, here);
c_name = "auto";
break;
- case PRAGMA_OACC_CLAUSE_ATTACH:
- clauses = cp_parser_oacc_data_clause (parser, c_kind, clauses);
- c_name = "attach";
- break;
case PRAGMA_OACC_CLAUSE_BIND:
clauses = cp_parser_oacc_clause_bind (parser, clauses);
c_name = "bind";
@@ -6724,7 +6724,7 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
error ("%qE is not a variable in %<depend%> clause", t);
remove = true;
}
- else if (ort != C_ORT_ACC && t == current_class_ptr)
+ else if (t == current_class_ptr)
{
error ("%<this%> allowed in OpenMP only in %<declare simd%>"
" clauses");
@@ -6810,6 +6810,10 @@ finish_omp_clauses (tree clauses, enum c_omp_region_type ort)
t = TREE_OPERAND (t, 0);
OMP_CLAUSE_DECL (c) = t;
}
+ if (ort == C_ORT_ACC
+ && TREE_CODE (t) == COMPONENT_REF
+ && TREE_CODE (TREE_OPERAND (t, 0)) == INDIRECT_REF)
+ t = TREE_OPERAND (TREE_OPERAND (t, 0), 0);
if (TREE_CODE (t) == COMPONENT_REF
&& ((ort & C_ORT_OMP_DECLARE_SIMD) == C_ORT_OMP
|| ort == C_ORT_ACC)
@@ -1183,10 +1183,12 @@ enum gfc_omp_depend_op
enum gfc_omp_map_op
{
OMP_MAP_ALLOC,
+ OMP_MAP_ATTACH,
OMP_MAP_TO,
OMP_MAP_FROM,
OMP_MAP_TOFROM,
OMP_MAP_DELETE,
+ OMP_MAP_DETACH,
OMP_MAP_FORCE_ALLOC,
OMP_MAP_FORCE_TO,
OMP_MAP_FORCE_FROM,
@@ -808,7 +808,7 @@ enum omp_mask1
OMP_MASK1_LAST
};
-/* OpenACC 2.0 specific clauses. */
+/* OpenACC 2.0+ specific clauses. */
enum omp_mask2
{
OMP_CLAUSE_ASYNC,
@@ -837,6 +837,8 @@ enum omp_mask2
OMP_CLAUSE_IF_PRESENT,
OMP_CLAUSE_FINALIZE,
OMP_CLAUSE_DEVICE_TYPE,
+ OMP_CLAUSE_ATTACH,
+ OMP_CLAUSE_DETACH,
/* This must come last. */
OMP_MASK2_LAST
};
@@ -964,10 +966,18 @@ static match
gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
const omp_mask dtype_mask,
bool first = true, bool needs_space = true,
- bool openacc = false, bool allow_derived = false)
+ bool openacc = false)
{
gfc_omp_clauses *base_clauses, *c = gfc_get_omp_clauses ();
locus old_loc;
+ /* Determine whether we're dealing with an OpenACC directive that permits
+ derived type member accesses. This in particular disallows
+ "!$acc declare" from using such accesses, because it's not clear if/how
+ that should work. */
+ bool allow_derived = (openacc
+ && ((mask & OMP_CLAUSE_ATTACH)
+ || (mask & OMP_CLAUSE_DETACH)
+ || (mask & OMP_CLAUSE_HOST_SELF)));
base_clauses = c;
@@ -1043,6 +1053,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
needs_space = true;
continue;
}
+ if ((mask & OMP_CLAUSE_ATTACH)
+ && gfc_match ("attach ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_ATTACH, false,
+ allow_derived))
+ continue;
break;
case 'b':
if ((mask & OMP_CLAUSE_BIND) && c->routine_bind == NULL
@@ -1098,8 +1114,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if ((mask & OMP_CLAUSE_COPYOUT)
&& gfc_match ("copyout ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_FROM, true,
- allow_derived))
+ OMP_MAP_FROM, true, allow_derived))
continue;
if ((mask & OMP_CLAUSE_COPYPRIVATE)
&& gfc_match_omp_variable_list ("copyprivate (",
@@ -1109,8 +1124,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if ((mask & OMP_CLAUSE_CREATE)
&& gfc_match ("create ( ") == MATCH_YES
&& gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
- OMP_MAP_ALLOC, true,
- allow_derived))
+ OMP_MAP_ALLOC, true, allow_derived))
continue;
break;
case 'd':
@@ -1190,6 +1204,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
else
gfc_current_locus = old_loc;
}
+ if ((mask & OMP_CLAUSE_DETACH)
+ && gfc_match ("detach ( ") == MATCH_YES
+ && gfc_match_omp_map_clause (&c->lists[OMP_LIST_MAP],
+ OMP_MAP_DETACH, false,
+ allow_derived))
+ continue;
if ((mask & OMP_CLAUSE_DEVICE)
&& !openacc
&& c->device == NULL
@@ -1784,8 +1804,8 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
if (gfc_match_omp_variable_list (" :",
&c->lists[OMP_LIST_REDUCTION],
- false, NULL, &head,
- openacc) == MATCH_YES)
+ false, NULL, &head, openacc,
+ allow_derived) == MATCH_YES)
{
gfc_omp_namelist *n;
if (rop == OMP_REDUCTION_NONE)
@@ -2053,7 +2073,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
| OMP_CLAUSE_PRIVATE | OMP_CLAUSE_FIRSTPRIVATE \
- | OMP_CLAUSE_DEFAULT)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_KERNELS_CLAUSES \
(omp_mask (OMP_CLAUSE_ASYNC) | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_NUM_GANGS | OMP_CLAUSE_NUM_WORKERS \
@@ -2063,12 +2083,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
| OMP_CLAUSE_DEVICEPTR \
- | OMP_CLAUSE_DEFAULT)
+ | OMP_CLAUSE_DEFAULT | OMP_CLAUSE_ATTACH)
#define OACC_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_COPY | OMP_CLAUSE_COPYIN | OMP_CLAUSE_COPYOUT \
| OMP_CLAUSE_CREATE | OMP_CLAUSE_PRESENT \
- | OMP_CLAUSE_DEVICEPTR)
+ | OMP_CLAUSE_DEVICEPTR | OMP_CLAUSE_ATTACH)
#define OACC_HOST_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_USE_DEVICE))
#define OACC_LOOP_CLAUSES \
@@ -2098,12 +2118,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
#define OACC_ENTER_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
- | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE)
+ | OMP_CLAUSE_COPYIN | OMP_CLAUSE_CREATE | OMP_CLAUSE_ATTACH)
#define OACC_EXIT_DATA_CLAUSES \
(omp_mask (OMP_CLAUSE_IF) \
| OMP_CLAUSE_ASYNC | OMP_CLAUSE_WAIT \
| OMP_CLAUSE_COPYOUT | OMP_CLAUSE_DELETE \
- | OMP_CLAUSE_FINALIZE)
+ | OMP_CLAUSE_FINALIZE | OMP_CLAUSE_DETACH)
#define OACC_ROUTINE_CLAUSES \
(omp_mask (OMP_CLAUSE_GANG) | OMP_CLAUSE_WORKER | OMP_CLAUSE_VECTOR \
| OMP_CLAUSE_SEQ \
@@ -2139,12 +2159,10 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, omp_mask mask,
static match
-match_acc (gfc_exec_op op, const omp_mask mask, const omp_mask dtype_mask,
- bool derived_types=false)
+match_acc (gfc_exec_op op, const omp_mask mask, const omp_mask dtype_mask)
{
gfc_omp_clauses *c;
- if (gfc_match_omp_clauses (&c, mask, dtype_mask, false, false, true,
- derived_types)
+ if (gfc_match_omp_clauses (&c, mask, dtype_mask, false, false, true)
!= MATCH_YES)
return MATCH_ERROR;
new_st.op = op;
@@ -2309,7 +2327,8 @@ gfc_match_oacc_update (void)
if (gfc_match_omp_clauses (&c, OACC_UPDATE_CLAUSES,
OACC_UPDATE_CLAUSE_DEVICE_TYPE_MASK, false,
- false, true, true) != MATCH_YES)
+ false, true)
+ != MATCH_YES)
return MATCH_ERROR;
if (!c->lists[OMP_LIST_MAP])
@@ -2329,7 +2348,7 @@ match
gfc_match_oacc_enter_data (void)
{
return match_acc (EXEC_OACC_ENTER_DATA, OACC_ENTER_DATA_CLAUSES,
- OMP_MASK2_LAST, true);
+ OMP_MASK2_LAST);
}
@@ -2337,7 +2356,7 @@ match
gfc_match_oacc_exit_data (void)
{
return match_acc (EXEC_OACC_EXIT_DATA, OACC_EXIT_DATA_CLAUSES,
- OMP_MASK2_LAST, true);
+ OMP_MASK2_LAST);
}
@@ -4017,9 +4036,6 @@ resolve_nonnegative_int_expr (gfc_expr *expr, const char *clause)
static void
check_symbol_not_pointer (gfc_symbol *sym, locus loc, const char *name)
{
- if (sym->ts.type == BT_DERIVED && sym->attr.pointer)
- gfc_error ("POINTER object %qs of derived type in %s clause at %L",
- sym->name, name, &loc);
if (sym->ts.type == BT_DERIVED && sym->attr.cray_pointer)
gfc_error ("Cray pointer object %qs of derived type in %s clause at %L",
sym->name, name, &loc);
@@ -4060,9 +4076,6 @@ check_array_not_assumed (gfc_symbol *sym, locus loc, const char *name)
static void
resolve_oacc_data_clauses (gfc_symbol *sym, locus loc, const char *name)
{
- if (sym->ts.type == BT_DERIVED && sym->attr.allocatable)
- gfc_error ("ALLOCATABLE object %qs of derived type in %s clause at %L",
- sym->name, name, &loc);
if ((sym->ts.type == BT_ASSUMED && sym->attr.allocatable)
|| (sym->ts.type == BT_CLASS && CLASS_DATA (sym)
&& CLASS_DATA (sym)->attr.allocatable))
@@ -4408,11 +4421,23 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
&& (list != OMP_LIST_REDUCTION || !openacc))
for (n = omp_clauses->lists[list]; n; n = n->next)
{
- if (n->sym->mark)
- gfc_error ("Symbol %qs present on multiple clauses at %L",
- n->sym->name, &n->where);
- else
- n->sym->mark = 1;
+ bool array_only_p = true;
+ /* Disallow duplicate bare variable references and multiple
+ subarrays of the same array here, but allow multiple components of
+ the same (e.g. derived-type) variable. For the latter, duplicate
+ components are detected elsewhere. */
+ if (openacc && n->expr && n->expr->expr_type == EXPR_VARIABLE)
+ for (gfc_ref *ref = n->expr->ref; ref; ref = ref->next)
+ if (ref->type != REF_ARRAY)
+ array_only_p = false;
+ if (array_only_p)
+ {
+ if (n->sym->mark)
+ gfc_error ("Symbol %qs present on multiple clauses at %L",
+ n->sym->name, &n->where);
+ else
+ n->sym->mark = 1;
+ }
}
gcc_assert (OMP_LIST_LASTPRIVATE == OMP_LIST_FIRSTPRIVATE + 1);
@@ -4603,26 +4628,41 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
"are allowed on ORDERED directive at %L",
&n->where);
}
+ gfc_ref *array_ref = NULL;
+ bool resolved = false;
if (n->expr)
{
- if (!gfc_resolve_expr (n->expr)
+ array_ref = n->expr->ref;
+ resolved = gfc_resolve_expr (n->expr);
+
+ /* Look through component refs to find last array
+ reference. */
+ while (resolved
+ && array_ref
+ && (array_ref->type == REF_COMPONENT
+ || (array_ref->type == REF_ARRAY
+ && array_ref->next
+ && array_ref->next->type == REF_COMPONENT)))
+ array_ref = array_ref->next;
+ }
+ if (array_ref
+ || (n->expr
+ && (!resolved || n->expr->expr_type != EXPR_VARIABLE)))
+ {
+ if (!resolved
|| n->expr->expr_type != EXPR_VARIABLE
- || n->expr->ref == NULL
- || n->expr->ref->next
- || n->expr->ref->type != REF_ARRAY)
- {
- if (n->sym->ts.type != BT_DERIVED)
- gfc_error ("%qs in %s clause at %L is not a proper "
- "array section", n->sym->name, name,
- &n->where);
- }
- else if (n->expr->ref->u.ar.codimen)
+ || array_ref->next
+ || array_ref->type != REF_ARRAY)
+ gfc_error ("%qs in %s clause at %L is not a proper "
+ "array section", n->sym->name, name,
+ &n->where);
+ else if (array_ref->u.ar.codimen)
gfc_error ("Coarrays not supported in %s clause at %L",
name, &n->where);
else
{
int i;
- gfc_array_ref *ar = &n->expr->ref->u.ar;
+ gfc_array_ref *ar = &array_ref->u.ar;
for (i = 0; i < ar->dimen; i++)
if (ar->stride[i])
{
@@ -62,6 +62,9 @@ gfc_omp_privatize_by_reference (const_tree decl)
if (TREE_CODE (type) == POINTER_TYPE)
{
+ while (TREE_CODE (decl) == COMPONENT_REF)
+ decl = TREE_OPERAND (decl, 1);
+
/* Array POINTER/ALLOCATABLE have aggregate types, all user variables
that have POINTER_TYPE type and aren't scalar pointers, scalar
allocatables, Cray pointees or C pointers are supposed to be
@@ -2121,69 +2124,35 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
tree decl = gfc_get_symbol_decl (n->sym);
if (DECL_P (decl))
TREE_ADDRESSABLE (decl) = 1;
- /* Handle derived-typed members for OpenACC Update. */
- if (n->sym->ts.type == BT_DERIVED
- && n->expr != NULL && n->expr->ref != NULL
- && (n->expr->ref->next == NULL
- || (n->expr->ref->next != NULL
- && n->expr->ref->next->type == REF_ARRAY
- && n->expr->ref->next->u.ar.type == AR_FULL))
- && (n->expr->ref->type == REF_ARRAY
- && n->expr->ref->u.ar.type != AR_SECTION))
- {
- gfc_ref *ref = n->expr->ref;
- gfc_component *c = ref->u.c.component;
- tree field;
- tree context;
- tree ptr;
- tree type;
- tree scratch;
- if (c->backend_decl == NULL_TREE
- && ref->u.c.sym != NULL)
- gfc_get_derived_type (ref->u.c.sym);
+ gfc_ref *ref = n->expr ? n->expr->ref : NULL;
+ symbol_attribute *sym_attr = &n->sym->attr;
+ gomp_map_kind ptr_map_kind = GOMP_MAP_POINTER;
- field = c->backend_decl;
- gcc_assert (field && TREE_CODE (field) == FIELD_DECL);
- context = DECL_FIELD_CONTEXT (field);
-
- type = TREE_TYPE (decl);
- if (POINTER_TYPE_P (type))
- type = TREE_TYPE (type);
+ if (ref && n->sym->ts.type == BT_DERIVED)
+ {
+ if (gfc_omp_privatize_by_reference (decl))
+ decl = build_fold_indirect_ref (decl);
- if (context != type)
+ for (; ref && ref->type == REF_COMPONENT; ref = ref->next)
{
- tree f2 = c->norestrict_decl;
- if (!f2 || DECL_FIELD_CONTEXT (f2) != type)
- for (f2 = TYPE_FIELDS (TREE_TYPE (decl)); f2;
- f2 = DECL_CHAIN (f2))
- if (TREE_CODE (f2) == FIELD_DECL
- && DECL_NAME (f2) == DECL_NAME (field))
- break;
- gcc_assert (f2);
- c->norestrict_decl = f2;
- field = f2;
+ tree field = ref->u.c.component->backend_decl;
+ gcc_assert (field && TREE_CODE (field) == FIELD_DECL);
+ decl = fold_build3 (COMPONENT_REF, TREE_TYPE (field),
+ decl, field, NULL_TREE);
+ sym_attr = &ref->u.c.component->attr;
}
- if (POINTER_TYPE_P (TREE_TYPE (decl)))
- decl = build_fold_indirect_ref_loc (input_location,
- decl);
-
- scratch = fold_build3_loc (input_location, COMPONENT_REF,
- TREE_TYPE (field), decl, field,
- NULL_TREE);
- type = TREE_TYPE (scratch);
- ptr = gfc_create_var (pvoid_type_node, NULL);
- scratch = fold_convert (pvoid_type_node,
- build_fold_addr_expr (scratch));
- gfc_add_modify (block, ptr, scratch);
- OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (type);
- OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
+ ptr_map_kind = GOMP_MAP_ALWAYS_POINTER;
}
- else if ((n->sym->ts.type == BT_DERIVED && n->expr == NULL)
- || (n->expr == NULL
- || n->expr->ref->u.ar.type == AR_FULL))
+
+ if (ref == NULL || ref->u.ar.type == AR_FULL)
{
+ tree field = decl;
+
+ while (TREE_CODE (field) == COMPONENT_REF)
+ field = TREE_OPERAND (field, 1);
+
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& n->u.map_op == OMP_MAP_FORCE_DEVICEPTR)
{
@@ -2192,18 +2161,18 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
}
else if (POINTER_TYPE_P (TREE_TYPE (decl))
&& (gfc_omp_privatize_by_reference (decl)
- || GFC_DECL_GET_SCALAR_POINTER (decl)
- || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
- || GFC_DECL_CRAY_POINTEE (decl)
+ || GFC_DECL_GET_SCALAR_POINTER (field)
+ || GFC_DECL_GET_SCALAR_ALLOCATABLE (field)
+ || GFC_DECL_CRAY_POINTEE (field)
|| GFC_DESCRIPTOR_TYPE_P
- (TREE_TYPE (TREE_TYPE (decl)))))
+ (TREE_TYPE (TREE_TYPE (field)))))
{
tree orig_decl = decl;
enum gomp_map_kind gmk = GOMP_MAP_FIRSTPRIVATE_POINTER;
if (GFC_DECL_GET_SCALAR_ALLOCATABLE (decl)
&& (n->sym->attr.oacc_declare_create)
&& clauses->update_allocatable)
- gmk = GOMP_MAP_ALWAYS_POINTER;
+ gmk = ptr_map_kind;
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, gmk);
@@ -2216,7 +2185,7 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
{
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3) = decl;
OMP_CLAUSE_SIZE (node3) = size_int (0);
decl = build_fold_indirect_ref (decl);
@@ -2225,7 +2194,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
}
- if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
+ if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))
+ && n->u.map_op != OMP_MAP_ATTACH
+ && n->u.map_op != OMP_MAP_DETACH)
{
tree type = TREE_TYPE (decl);
tree ptr = gfc_conv_descriptor_data_get (decl);
@@ -2238,14 +2209,16 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
+ if (ptr_map_kind == GOMP_MAP_ALWAYS_POINTER)
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
OMP_CLAUSE_SIZE (node3) = size_int (0);
/* We have to check for n->sym->attr.dimension because
of scalar coarrays. */
- if (n->sym->attr.pointer && n->sym->attr.dimension)
+ if (sym_attr->pointer && sym_attr->dimension)
{
stmtblock_t cond_block;
tree size
@@ -2275,11 +2248,11 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
else_b));
OMP_CLAUSE_SIZE (node) = size;
}
- else if (n->sym->attr.dimension)
+ else if (sym_attr->dimension)
OMP_CLAUSE_SIZE (node)
= gfc_full_array_size (block, decl,
GFC_TYPE_ARRAY_RANK (type));
- if (n->sym->attr.dimension)
+ if (sym_attr->dimension)
{
tree elemsz
= TYPE_SIZE_UNIT (gfc_get_element_type (type));
@@ -2292,31 +2265,17 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
else
OMP_CLAUSE_DECL (node) = decl;
}
- else
+ else if (ref)
{
tree ptr, ptr2;
gfc_init_se (&se, NULL);
- if ((n->sym->ts.type == BT_DERIVED
- && n->expr->rank == 0)
- || (n->sym->ts.type != BT_DERIVED
- && n->expr->ref->u.ar.type == AR_ELEMENT))
+ if (ref->u.ar.type == AR_ELEMENT)
{
gfc_conv_expr_reference (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
ptr = se.expr;
- tree type = TREE_TYPE (ptr);
- if (n->sym->ts.type == BT_DERIVED)
- {
- tree t = gfc_create_var (build_pointer_type
- (void_type_node),
- NULL);
- ptr = fold_convert (pvoid_type_node, ptr);
- gfc_add_modify (block, t, ptr);
- ptr = t;
- type = TREE_TYPE (type);
- }
OMP_CLAUSE_SIZE (node)
- = TYPE_SIZE_UNIT (type);
+ = TYPE_SIZE_UNIT (TREE_TYPE (ptr));
}
else
{
@@ -2337,14 +2296,12 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
gfc_add_block_to_block (block, &se.post);
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
- if (n->sym->ts.type == BT_DERIVED)
- goto finalize_map_clause;
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))))
{
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node4, ptr_map_kind);
OMP_CLAUSE_DECL (node4) = decl;
OMP_CLAUSE_SIZE (node4) = size_int (0);
decl = build_fold_indirect_ref (decl);
@@ -2361,9 +2318,11 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type);
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (decl);
+ if (ptr_map_kind == GOMP_MAP_ALWAYS_POINTER)
+ STRIP_NOPS (OMP_CLAUSE_DECL (node3));
}
else
{
@@ -2376,7 +2335,7 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER);
+ OMP_CLAUSE_SET_MAP_KIND (node3, ptr_map_kind);
OMP_CLAUSE_DECL (node3) = decl;
}
ptr2 = fold_convert (sizetype, ptr2);
@@ -2384,11 +2343,16 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
= fold_build2 (MINUS_EXPR, sizetype, ptr, ptr2);
finalize_map_clause:;
}
+ else
+ gcc_unreachable ();
switch (n->u.map_op)
{
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
+ case OMP_MAP_ATTACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ATTACH);
+ break;
case OMP_MAP_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TO);
break;
@@ -2413,6 +2377,9 @@ gfc_trans_omp_clauses_1 (stmtblock_t *block, gfc_omp_clauses *clauses,
case OMP_MAP_DELETE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DELETE);
break;
+ case OMP_MAP_DETACH:
+ OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_DETACH);
+ break;
case OMP_MAP_FORCE_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_ALLOC);
break;
@@ -111,6 +111,10 @@ enum gimplify_omp_var_data
/* Flag for OpenACC deviceptrs. */
GOVD_DEVICEPTR = (1<<21),
+ /* Flag for GOVD_MAP: (struct) vars that have pointer attachments for
+ fields. */
+ GOVD_MAP_HAS_ATTACHMENTS = (1<<22),
+
GOVD_DATA_SHARE_CLASS = (GOVD_SHARED | GOVD_PRIVATE | GOVD_FIRSTPRIVATE
| GOVD_LASTPRIVATE | GOVD_REDUCTION | GOVD_LINEAR
| GOVD_LOCAL)
@@ -7692,7 +7696,13 @@ insert_struct_component_mapping (enum tree_code code, tree c, tree struct_node,
OMP_CLAUSE_SET_MAP_KIND (c2, mkind);
OMP_CLAUSE_DECL (c2) = unshare_expr (OMP_CLAUSE_DECL (c));
OMP_CLAUSE_CHAIN (c2) = scp ? *scp : prev_node;
- OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
+ if (OMP_CLAUSE_CHAIN (prev_node) != c
+ && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (prev_node)) == OMP_CLAUSE_MAP
+ && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (prev_node))
+ == GOMP_MAP_TO_PSET))
+ OMP_CLAUSE_SIZE (c2) = OMP_CLAUSE_SIZE (OMP_CLAUSE_CHAIN (prev_node));
+ else
+ OMP_CLAUSE_SIZE (c2) = TYPE_SIZE_UNIT (ptr_type_node);
if (struct_node)
OMP_CLAUSE_CHAIN (struct_node) = c2;
@@ -8245,7 +8255,9 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
remove = true;
break;
}
- if (DECL_P (decl))
+ if (DECL_P (decl)
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
+ && code != OACC_UPDATE)
{
if (error_operand_p (decl))
{
@@ -8297,17 +8309,36 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
= splay_tree_lookup (ctx->variables, (splay_tree_key)decl);
bool ptr = (OMP_CLAUSE_MAP_KIND (c)
== GOMP_MAP_ALWAYS_POINTER);
- if ((n == NULL || (n->value & GOVD_MAP) == 0)
- && code != OACC_UPDATE)
+ bool attach = OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_ATTACH
+ || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_DETACH;
+ bool has_attachments = false;
+ /* For OpenACC, pointers in structs should trigger an
+ attach action. */
+ if (ptr && (region_type & ORT_ACC) != 0)
+ {
+ /* Turning a GOMP_MAP_ALWAYS_POINTER clause into a
+ GOMP_MAP_ATTACH clause after we have detected a case
+ that needs a GOMP_MAP_STRUCT mapping adding. */
+ OMP_CLAUSE_SET_MAP_KIND (c,
+ (code == OACC_EXIT_DATA) ? GOMP_MAP_DETACH
+ : GOMP_MAP_ATTACH);
+ has_attachments = true;
+ }
+ if (n == NULL || (n->value & GOVD_MAP) == 0)
{
tree l = build_omp_clause (OMP_CLAUSE_LOCATION (c),
OMP_CLAUSE_MAP);
- OMP_CLAUSE_SET_MAP_KIND (l, GOMP_MAP_STRUCT);
+ OMP_CLAUSE_SET_MAP_KIND (l, attach
+ ? GOMP_MAP_FORCE_PRESENT : GOMP_MAP_STRUCT);
if (!base_eq_orig_base)
OMP_CLAUSE_DECL (l) = unshare_expr (orig_base);
else
OMP_CLAUSE_DECL (l) = decl;
- OMP_CLAUSE_SIZE (l) = size_int (1);
+ OMP_CLAUSE_SIZE (l) = attach
+ ? (DECL_P (OMP_CLAUSE_DECL (l))
+ ? DECL_SIZE_UNIT (OMP_CLAUSE_DECL (l))
+ : TYPE_SIZE_UNIT (TREE_TYPE (OMP_CLAUSE_DECL (l))))
+ : size_int (1);
if (struct_map_to_clause == NULL)
struct_map_to_clause = new hash_map<tree, tree>;
struct_map_to_clause->put (decl, l);
@@ -8339,9 +8370,11 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
flags = GOVD_MAP | GOVD_EXPLICIT;
if (GOMP_MAP_ALWAYS_P (OMP_CLAUSE_MAP_KIND (c)) || ptr)
flags |= GOVD_SEEN;
+ if (has_attachments)
+ flags |= GOVD_MAP_HAS_ATTACHMENTS;
goto do_add_decl;
}
- else
+ else if (struct_map_to_clause)
{
tree *osc = struct_map_to_clause->get (decl);
tree *sc = NULL, *scp = NULL;
@@ -8350,8 +8383,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
sc = &OMP_CLAUSE_CHAIN (*osc);
if (*sc != c
&& (OMP_CLAUSE_MAP_KIND (*sc)
- == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
+ == GOMP_MAP_FIRSTPRIVATE_REFERENCE))
sc = &OMP_CLAUSE_CHAIN (*sc);
+ /* Here "prev_list_p" is the end of the inserted
+ alloc/release nodes after the struct node, OSC. */
for (; *sc != c; sc = &OMP_CLAUSE_CHAIN (*sc))
if (ptr && sc == prev_list_p)
break;
@@ -8410,9 +8445,10 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (remove)
break;
- OMP_CLAUSE_SIZE (*osc)
- = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
- size_one_node);
+ if (!attach)
+ OMP_CLAUSE_SIZE (*osc)
+ = size_binop (PLUS_EXPR, OMP_CLAUSE_SIZE (*osc),
+ size_one_node);
if (ptr)
{
tree cl
@@ -8444,11 +8480,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p,
}
if (!remove
&& OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER
+ && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET
&& OMP_CLAUSE_CHAIN (c)
&& OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP
- && (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
- == GOMP_MAP_ALWAYS_POINTER))
+ && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_ALWAYS_POINTER)
+ || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c))
+ == GOMP_MAP_TO_PSET)))
prev_list_p = list_p;
+
break;
}
flags = GOVD_MAP | GOVD_EXPLICIT;
@@ -9020,6 +9060,8 @@ gimplify_adjust_omp_clauses_1 (splay_tree_node n, void *data)
return 0;
if ((flags & GOVD_SEEN) == 0)
return 0;
+ if ((flags & GOVD_MAP_HAS_ATTACHMENTS) != 0)
+ return 0;
if (flags & GOVD_DEBUG_PRIVATE)
{
gcc_assert ((flags & GOVD_DATA_SHARE_CLASS) == GOVD_SHARED);
@@ -9509,8 +9551,7 @@ gimplify_adjust_omp_clauses (gimple_seq *pre_p, gimple_seq body, tree *list_p,
}
}
else if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_STRUCT
- && (code == OMP_TARGET_EXIT_DATA
- || code == OACC_EXIT_DATA))
+ && code == OMP_TARGET_EXIT_DATA)
remove = true;
else if (DECL_SIZE (decl)
&& TREE_CODE (DECL_SIZE (decl)) != INTEGER_CST
@@ -11218,10 +11259,15 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p)
OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_FORCE_DETACH);
finalize_marked = true;
break;
+ case GOMP_MAP_STRUCT:
+ case GOMP_MAP_FORCE_PRESENT:
+ /* Skip over an initial struct or force_present mapping. */
+ break;
default:
- /* Check consistency: libgomp relies on the very first data
- mapping clause being marked, so make sure we did that before
- any other mapping clauses. */
+ /* Check consistency: libgomp relies on the very first
+ non-struct, non-force-present data mapping clause being
+ marked, so make sure we did that before any other mapping
+ clauses. */
gcc_assert (finalize_marked);
break;
}
@@ -42,13 +42,13 @@ t1 ()
}
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:s .len: 32.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.always_pointer:s.a .pointer assign, bias: 0.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.tofrom:.z .len: 40.. map.struct:s .len: 1.. map.alloc:s.a .len: 8.. map.tofrom:._1 .len: 40.. map.attach:s.a .len: 0.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_parallel map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.attach:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.to:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.detach:s.e .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.struct:s .len: 1.. map.attach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.force_present:s .len: 32.. map.detach:s.e .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_data map.force_present:s .len: 32.. map.attach:s.e .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data map.release:a .len: 8.." 1 "omplower" } } */
/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:a .len: 8.." 1 "omplower" } } */
-/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_detach:s.a .len: 8.." 1 "omplower" } } */
+/* { dg-final { scan-tree-dump-times "pragma omp target oacc_enter_exit_data finalize map.force_present:s .len: 32.. map.force_detach:s.a .len: 8.." 1 "omplower" } } */
@@ -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
@@ -33,48 +33,45 @@ program derived_acc
!$acc exit data copyout(var)
!$acc exit data copyout(var%a)
- !$acc data copy(var%a) ! { dg-error "Syntax error in OpenMP" }
- !$acc end data ! { dg-error "Unexpected ..ACC END DATA" }
-
!$acc data copy(var)
!$acc end data
- !$acc data copyout(var%a) ! { dg-error "Syntax error in OpenMP" }
- !$acc end data ! { dg-error "Unexpected ..ACC END" }
+ !$acc data copyout(var%a)
+ !$acc end data
!$acc parallel loop pcopyout(var)
do i = 1, 10
end do
!$acc end parallel loop
- !$acc parallel loop copyout(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc parallel loop copyout(var%a)
do i = 1, 10
end do
- !$acc end parallel loop ! { dg-error "Unexpected ..ACC END" }
+ !$acc end parallel loop
!$acc parallel pcopy(var)
!$acc end parallel
- !$acc parallel pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc parallel pcopy(var%a)
do i = 1, 10
end do
- !$acc end parallel ! { dg-error "Unexpected ..ACC END" }
+ !$acc end parallel
!$acc kernels pcopyin(var)
!$acc end kernels
- !$acc kernels pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc kernels pcopy(var%a)
do i = 1, 10
end do
- !$acc end kernels ! { dg-error "Unexpected ..ACC END" }
+ !$acc end kernels
!$acc kernels loop pcopyin(var)
do i = 1, 10
end do
!$acc end kernels loop
- !$acc kernels loop pcopy(var%a) ! { dg-error "Syntax error in OpenMP" }
+ !$acc kernels loop pcopy(var%a)
do i = 1, 10
end do
- !$acc end kernels loop ! { dg-error "Unexpected ..ACC END" }
+ !$acc end kernels loop
end program derived_acc
@@ -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)
@@ -806,6 +806,8 @@ struct target_var_desc {
bool copy_from;
/* True if data always should be copied from device to host at the end. */
bool always_copy_from;
+ /* True if variable should be detached at end of region. */
+ bool do_detach;
/* Relative offset against key host_start. */
uintptr_t offset;
/* Actual length. */
@@ -860,6 +862,8 @@ struct splay_tree_key_s {
uintptr_t refcount;
/* Dynamic reference count. */
uintptr_t dynamic_refcount;
+ /* For a block with attached pointers, the attachment counters for each. */
+ unsigned short *attach_count;
/* Pointer to the original mapping of "omp declare target link" object. */
splay_tree_key link_key;
};
@@ -1003,6 +1007,8 @@ enum gomp_map_vars_kind
GOMP_MAP_VARS_ENTER_DATA
};
+struct gomp_coalesce_buf;
+
extern void gomp_acc_insert_pointer (size_t, void **, size_t *, void *, int);
extern void gomp_acc_remove_pointer (void *, size_t, bool, int, int, int);
extern void gomp_acc_declare_allocate (bool, size_t, void **, size_t *,
@@ -1013,8 +1019,17 @@ extern void gomp_copy_host2dev (struct gomp_device_descr *,
void *, const void *, size_t,
struct gomp_coalesce_buf *);
extern void gomp_copy_dev2host (struct gomp_device_descr *,
- struct goacc_asyncqueue *,
- void *, const void *, size_t);
+ struct goacc_asyncqueue *, void *, const void *,
+ size_t);
+extern uintptr_t gomp_map_val (struct target_mem_desc *, void **, size_t);
+extern void gomp_attach_pointer (struct gomp_device_descr *,
+ struct goacc_asyncqueue *, splay_tree,
+ splay_tree_key, uintptr_t, size_t,
+ struct gomp_coalesce_buf *);
+extern void gomp_detach_pointer (struct gomp_device_descr *,
+ struct goacc_asyncqueue *, splay_tree_key,
+ uintptr_t, bool, struct gomp_coalesce_buf *);
+
extern struct target_mem_desc *gomp_map_vars (struct gomp_device_descr *,
size_t, void **, void **,
size_t *, void *, bool,
@@ -1025,9 +1040,9 @@ extern struct target_mem_desc *gomp_map_vars_async (struct gomp_device_descr *,
size_t *, void *, bool,
enum gomp_map_vars_kind);
extern void gomp_unmap_tgt (struct target_mem_desc *);
-extern void gomp_unmap_vars (struct target_mem_desc *, bool);
+extern void gomp_unmap_vars (struct target_mem_desc *, bool, bool);
extern void gomp_unmap_vars_async (struct target_mem_desc *, bool,
- struct goacc_asyncqueue *);
+ struct goacc_asyncqueue *, bool);
extern void gomp_init_device (struct gomp_device_descr *);
extern bool gomp_fini_device (struct gomp_device_descr *);
extern void gomp_unload_device (struct gomp_device_descr *);
@@ -440,6 +440,16 @@ OACC_2.5 {
acc_update_self_async_array_h_;
} OACC_2.0.1;
+OACC_2.6 {
+ global:
+ acc_attach;
+ acc_attach_async;
+ acc_detach;
+ acc_detach_async;
+ acc_detach_finalize;
+ acc_detach_finalize_async;
+} OACC_2.5;
+
GOACC_2.0 {
global:
GOACC_data_end;
@@ -373,14 +373,14 @@ goacc_async_unmap_tgt (void *ptr)
attribute_hidden void
goacc_async_copyout_unmap_vars (struct target_mem_desc *tgt,
- struct goacc_asyncqueue *aq)
+ struct goacc_asyncqueue *aq, bool finalize)
{
struct gomp_device_descr *devicep = tgt->device_descr;
/* Increment reference to delay freeing of device memory until callback
has triggered. */
tgt->refcount++;
- gomp_unmap_vars_async (tgt, true, aq);
+ gomp_unmap_vars_async (tgt, true, aq, finalize);
devicep->openacc.async.queue_callback_func (aq, goacc_async_unmap_tgt,
(void *) tgt);
}
@@ -391,7 +391,7 @@ acc_shutdown_1 (acc_device_t d)
{
struct target_mem_desc *tgt = walk->dev->mem_map.root->key.tgt;
- gomp_unmap_vars (tgt, false);
+ gomp_unmap_vars (tgt, false, false);
}
walk->dev = NULL;
@@ -112,7 +112,7 @@ void goacc_host_init (void);
void goacc_init_asyncqueues (struct gomp_device_descr *);
bool goacc_fini_asyncqueues (struct gomp_device_descr *);
void goacc_async_copyout_unmap_vars (struct target_mem_desc *,
- struct goacc_asyncqueue *);
+ struct goacc_asyncqueue *, bool);
void goacc_async_free (struct gomp_device_descr *,
struct goacc_asyncqueue *, void *);
struct goacc_asyncqueue *get_goacc_asyncqueue (int);
@@ -518,7 +518,7 @@ acc_unmap_data (void *h)
gomp_mutex_unlock (&acc_dev->lock);
- gomp_unmap_vars (t, true);
+ gomp_unmap_vars (t, true, false);
if (profiling_setup_p)
{
@@ -612,6 +612,7 @@ present_create_copy (unsigned f, void *h, size_t s, int async)
&kinds, true, GOMP_MAP_VARS_OPENACC);
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
+ tgt->list[0].key->attach_count = NULL;
gomp_mutex_lock (&acc_dev->lock);
@@ -750,6 +751,7 @@ delete_copyout (unsigned f, void *h, size_t s, int async, const char *libfnname)
{
n->refcount = 0;
n->dynamic_refcount = 0;
+ n->attach_count = NULL;
}
if (n->refcount < n->dynamic_refcount)
{
@@ -997,6 +999,7 @@ gomp_acc_insert_pointer (size_t mapnum, void **hostaddrs, size_t *sizes,
/* Initialize dynamic refcount. */
tgt->list[0].key->dynamic_refcount = 1;
+ tgt->list[0].key->attach_count = NULL;
gomp_mutex_lock (&acc_dev->lock);
tgt->prev = acc_dev->openacc.data_environ;
@@ -1084,11 +1087,11 @@ gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
/* If running synchronously, unmap immediately. */
if (async < acc_async_noval)
- gomp_unmap_vars (t, true);
+ gomp_unmap_vars (t, true, finalize);
else
{
goacc_aq aq = get_goacc_asyncqueue (async);
- goacc_async_copyout_unmap_vars (t, aq);
+ goacc_async_copyout_unmap_vars (t, aq, finalize);
}
}
@@ -1096,3 +1099,80 @@ gomp_acc_remove_pointer (void *h, size_t s, bool force_copyfrom, int async,
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
}
+
+
+void
+acc_attach_async (void **hostaddr, int async)
+{
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+ goacc_aq aq = get_goacc_asyncqueue (async);
+
+ struct splay_tree_key_s cur_node;
+ splay_tree_key n;
+
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
+ cur_node.host_start = (uintptr_t) hostaddr;
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ gomp_fatal ("struct not mapped for acc_attach");
+
+ gomp_attach_pointer (acc_dev, aq, &acc_dev->mem_map, n, (uintptr_t) hostaddr,
+ 0, NULL);
+}
+
+void
+acc_attach (void **hostaddr)
+{
+ acc_attach_async (hostaddr, acc_async_sync);
+}
+
+static void
+goacc_detach_internal (void **hostaddr, int async, bool finalize)
+{
+ struct goacc_thread *thr = goacc_thread ();
+ struct gomp_device_descr *acc_dev = thr->dev;
+ struct splay_tree_key_s cur_node;
+ splay_tree_key n;
+ struct goacc_asyncqueue *aq = get_goacc_asyncqueue (async);
+
+ if (thr->dev->capabilities & GOMP_OFFLOAD_CAP_SHARED_MEM)
+ return;
+
+ cur_node.host_start = (uintptr_t) hostaddr;
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ n = splay_tree_lookup (&acc_dev->mem_map, &cur_node);
+
+ if (n == NULL)
+ gomp_fatal ("struct not mapped for acc_detach");
+
+ gomp_detach_pointer (acc_dev, aq, n, (uintptr_t) hostaddr, finalize, NULL);
+}
+
+void
+acc_detach (void **hostaddr)
+{
+ goacc_detach_internal (hostaddr, acc_async_sync, false);
+}
+
+void
+acc_detach_async (void **hostaddr, int async)
+{
+ goacc_detach_internal (hostaddr, async, false);
+}
+
+void
+acc_detach_finalize (void **hostaddr)
+{
+ goacc_detach_internal (hostaddr, acc_async_sync, true);
+}
+
+void
+acc_detach_finalize_async (void **hostaddr, int async)
+{
+ goacc_detach_internal (hostaddr, async, true);
+}
@@ -50,12 +50,29 @@ find_pointer (int pos, size_t mapnum, unsigned short *kinds)
if (pos + 1 >= mapnum)
return 0;
- unsigned char kind = kinds[pos+1] & 0xff;
+ unsigned char kind0 = kinds[pos] & 0xff;
- if (kind == GOMP_MAP_TO_PSET)
- return 3;
- else if (kind == GOMP_MAP_POINTER)
- return 2;
+ switch (kind0)
+ {
+ case GOMP_MAP_TO:
+ case GOMP_MAP_FORCE_TO:
+ case GOMP_MAP_FROM:
+ case GOMP_MAP_FORCE_FROM:
+ case GOMP_MAP_ALLOC:
+ case GOMP_MAP_RELEASE:
+ {
+ unsigned char kind1 = kinds[pos + 1] & 0xff;
+ if (kind1 == GOMP_MAP_POINTER
+ || kind1 == GOMP_MAP_ALWAYS_POINTER
+ || kind1 == GOMP_MAP_ATTACH
+ || kind1 == GOMP_MAP_DETACH)
+ return 2;
+ else if (kind1 == GOMP_MAP_TO_PSET)
+ return 3;
+ }
+ default:
+ /* empty. */;
+ }
return 0;
}
@@ -355,14 +372,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
devaddrs = gomp_alloca (sizeof (void *) * mapnum);
for (i = 0; i < mapnum; i++)
- {
- if (tgt->list[i].key != NULL)
- devaddrs[i] = (void *) (tgt->list[i].key->tgt->tgt_start
- + tgt->list[i].key->tgt_offset
- + tgt->list[i].offset);
- else
- devaddrs[i] = NULL;
- }
+ devaddrs[i] = (void *) gomp_map_val (tgt, hostaddrs, i);
if (aq == NULL)
{
@@ -382,7 +392,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
&api_info);
}
/* If running synchronously, unmap immediately. */
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
if (profiling_dispatch_p)
{
prof_info.event_type = acc_ev_exit_data_end;
@@ -400,7 +410,7 @@ GOACC_parallel_keyed_internal (int device, int params, void (*fn) (void *),
else
acc_dev->openacc.async.exec_func (tgt_fn, mapnum, hostaddrs,
devaddrs, dims, tgt, aq);
- goacc_async_copyout_unmap_vars (tgt, aq);
+ goacc_async_copyout_unmap_vars (tgt, aq, false);
}
out:
@@ -637,7 +647,7 @@ GOACC_data_end (void)
gomp_debug (0, " %s: restore mappings\n", __FUNCTION__);
thr->mapped_data = tgt->prev;
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
gomp_debug (0, " %s: mappings restored\n", __FUNCTION__);
if (profiling_dispatch_p)
@@ -668,6 +678,10 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (mapnum > 0)
{
unsigned char kind = kinds[0] & 0xff;
+
+ if (kind == GOMP_MAP_STRUCT || kind == GOMP_MAP_FORCE_PRESENT)
+ kind = kinds[1] & 0xff;
+
if (kind == GOMP_MAP_DELETE
|| kind == GOMP_MAP_FORCE_FROM)
finalize = true;
@@ -678,11 +692,14 @@ GOACC_enter_exit_data (int device, size_t mapnum,
{
unsigned char kind = kinds[i] & 0xff;
- if (kind == GOMP_MAP_POINTER || kind == GOMP_MAP_TO_PSET)
+ if (kind == GOMP_MAP_POINTER
+ || kind == GOMP_MAP_TO_PSET
+ || kind == GOMP_MAP_STRUCT
+ || kind == GOMP_MAP_FORCE_PRESENT)
continue;
if (kind == GOMP_MAP_FORCE_ALLOC
- || kind == GOMP_MAP_FORCE_PRESENT
+ || kind == GOMP_MAP_ATTACH
|| kind == GOMP_MAP_FORCE_TO
|| kind == GOMP_MAP_TO
|| kind == GOMP_MAP_ALLOC
@@ -694,6 +711,8 @@ GOACC_enter_exit_data (int device, size_t mapnum,
if (kind == GOMP_MAP_RELEASE
|| kind == GOMP_MAP_DELETE
+ || kind == GOMP_MAP_DETACH
+ || kind == GOMP_MAP_FORCE_DETACH
|| kind == GOMP_MAP_FROM
|| kind == GOMP_MAP_FORCE_FROM
|| kind == GOMP_MAP_DECLARE_DEALLOCATE)
@@ -809,6 +828,9 @@ GOACC_enter_exit_data (int device, size_t mapnum,
case GOMP_MAP_ALLOC:
acc_present_or_create (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_ATTACH:
+ case GOMP_MAP_FORCE_PRESENT:
+ break;
case GOMP_MAP_FORCE_ALLOC:
acc_create (hostaddrs[i], sizes[i]);
break;
@@ -818,6 +840,27 @@ GOACC_enter_exit_data (int device, size_t mapnum,
case GOMP_MAP_FORCE_TO:
acc_copyin (hostaddrs[i], sizes[i]);
break;
+ case GOMP_MAP_STRUCT:
+ {
+ int elems = sizes[i];
+ struct splay_tree_key_s k;
+ splay_tree_key str;
+ k.host_start = (uintptr_t) hostaddrs[i];
+ k.host_end = k.host_start + 1;
+ gomp_mutex_lock (&acc_dev->lock);
+ str = splay_tree_lookup (&acc_dev->mem_map, &k);
+ gomp_mutex_unlock (&acc_dev->lock);
+ /* We increment the dynamic reference count for the struct
+ itself by the number of struct elements that we
+ mapped. */
+ if (str->refcount != REFCOUNT_INFINITY)
+ {
+ str->refcount += elems;
+ str->dynamic_refcount += elems;
+ }
+ i += elems;
+ }
+ break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
@@ -839,16 +882,57 @@ GOACC_enter_exit_data (int device, size_t mapnum,
i += pointer - 1;
}
}
+
+ /* This loop only handles explicit "attach" clauses that are not an
+ implicit part of a copy{,in,out}, etc. mapping. */
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ /* Scan for pointers and PSETs. */
+ int pointer = find_pointer (i, mapnum, kinds);
+
+ if (!pointer)
+ {
+ if (kind == GOMP_MAP_ATTACH)
+ acc_attach (hostaddrs[i]);
+ else if (kind == GOMP_MAP_STRUCT)
+ i += sizes[i];
+ }
+ else
+ i += pointer - 1;
+ }
}
else
- for (i = 0; i < mapnum; ++i)
- {
- unsigned char kind = kinds[i] & 0xff;
+ {
+ /* This loop only handles explicit "detach" clauses that are not an
+ implicit part of a copy{,in,out}, etc. mapping. */
+ for (i = 0; i < mapnum; i++)
+ {
+ unsigned char kind = kinds[i] & 0xff;
- int pointer = find_pointer (i, mapnum, kinds);
+ int pointer = find_pointer (i, mapnum, kinds);
- if (!pointer)
- {
+ if (!pointer)
+ {
+ if (kind == GOMP_MAP_DETACH)
+ acc_detach (hostaddrs[i]);
+ else if (kind == GOMP_MAP_FORCE_DETACH)
+ acc_detach_finalize (hostaddrs[i]);
+ else if (kind == GOMP_MAP_STRUCT)
+ i += sizes[i];
+ }
+ else
+ i += pointer - 1;
+ }
+
+ for (i = 0; i < mapnum; ++i)
+ {
+ unsigned char kind = kinds[i] & 0xff;
+
+ int pointer = find_pointer (i, mapnum, kinds);
+
+ if (!pointer)
switch (kind)
{
case GOMP_MAP_RELEASE:
@@ -861,6 +945,10 @@ GOACC_enter_exit_data (int device, size_t mapnum,
acc_delete_async (hostaddrs[i], sizes[i], async);
}
break;
+ case GOMP_MAP_DETACH:
+ case GOMP_MAP_FORCE_DETACH:
+ case GOMP_MAP_FORCE_PRESENT:
+ break;
case GOMP_MAP_DECLARE_DEALLOCATE:
case GOMP_MAP_FROM:
case GOMP_MAP_FORCE_FROM:
@@ -869,28 +957,48 @@ GOACC_enter_exit_data (int device, size_t mapnum,
else
acc_copyout_async (hostaddrs[i], sizes[i], async);
break;
+ case GOMP_MAP_STRUCT:
+ {
+ int elems = sizes[i];
+ struct splay_tree_key_s k;
+ splay_tree_key str;
+ k.host_start = (uintptr_t) hostaddrs[i];
+ k.host_end = k.host_start + 1;
+ gomp_mutex_lock (&acc_dev->lock);
+ str = splay_tree_lookup (&acc_dev->mem_map, &k);
+ gomp_mutex_unlock (&acc_dev->lock);
+ /* Decrement dynamic reference count for the struct by the
+ number of elements that we are unmapping. */
+ if (str->dynamic_refcount >= elems)
+ {
+ str->dynamic_refcount -= elems;
+ str->refcount -= elems;
+ }
+ i += elems;
+ }
+ break;
default:
gomp_fatal (">>>> GOACC_enter_exit_data UNHANDLED kind 0x%.2x",
kind);
break;
}
- }
- else
- {
- if (kind == GOMP_MAP_DECLARE_DEALLOCATE)
- gomp_acc_declare_allocate (false, pointer, &hostaddrs[i],
- &sizes[i], &kinds[i]);
- else
- {
- bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
- || kind == GOMP_MAP_FROM);
- gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom, async,
- finalize, pointer);
- /* See the above comment. */
- }
- i += pointer - 1;
- }
- }
+ else
+ {
+ if (kind == GOMP_MAP_DECLARE_DEALLOCATE)
+ gomp_acc_declare_allocate (false, pointer, &hostaddrs[i],
+ &sizes[i], &kinds[i]);
+ else
+ {
+ bool copyfrom = (kind == GOMP_MAP_FORCE_FROM
+ || kind == GOMP_MAP_FROM);
+ gomp_acc_remove_pointer (hostaddrs[i], sizes[i], copyfrom,
+ async, finalize, pointer);
+ /* See the above comment. */
+ }
+ i += pointer - 1;
+ }
+ }
+ }
out:
if (profiling_dispatch_p)
@@ -113,6 +113,10 @@ void *acc_hostptr (void *) __GOACC_NOTHROW;
int acc_is_present (void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_to_device (void *, void *, size_t) __GOACC_NOTHROW;
void acc_memcpy_from_device (void *, void *, size_t) __GOACC_NOTHROW;
+void acc_attach (void **) __GOACC_NOTHROW;
+void acc_attach_async (void **, int) __GOACC_NOTHROW;
+void acc_detach (void **) __GOACC_NOTHROW;
+void acc_detach_async (void **, int) __GOACC_NOTHROW;
/* Async functions, specified in OpenACC 2.5. */
void acc_copyin_async (void *, size_t, int) __GOACC_NOTHROW;
@@ -129,6 +133,8 @@ void acc_copyout_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_copyout_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
void acc_delete_finalize (void *, size_t) __GOACC_NOTHROW;
void acc_delete_finalize_async (void *, size_t, int) __GOACC_NOTHROW;
+void acc_detach_finalize (void **) __GOACC_NOTHROW;
+void acc_detach_finalize_async (void **, int) __GOACC_NOTHROW;
/* CUDA-specific routines. */
void *acc_get_current_cuda_device (void) __GOACC_NOTHROW;
@@ -39,6 +39,7 @@
#include <string.h>
#include <assert.h>
#include <errno.h>
+#include <limits.h>
#ifdef PLUGIN_SUPPORT
#include <dlfcn.h>
@@ -373,6 +374,7 @@ gomp_map_vars_existing (struct gomp_device_descr *devicep,
tgt_var->key = oldn;
tgt_var->copy_from = GOMP_MAP_COPY_FROM_P (kind);
tgt_var->always_copy_from = GOMP_MAP_ALWAYS_FROM_P (kind);
+ tgt_var->do_detach = false;
tgt_var->offset = newn->host_start - oldn->host_start;
tgt_var->length = newn->host_end - newn->host_start;
@@ -539,7 +541,128 @@ gomp_map_fields_existing (struct target_mem_desc *tgt,
(void *) cur_node.host_end);
}
-static inline uintptr_t
+void
+gomp_attach_pointer (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, splay_tree mem_map,
+ splay_tree_key n, uintptr_t attach_to, size_t bias,
+ struct gomp_coalesce_buf *cbufp)
+{
+ struct splay_tree_key_s s;
+ size_t size, idx;
+
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("enclosing struct not mapped for attach");
+ }
+
+ size = (n->host_end - n->host_start + sizeof (void *) - 1) / sizeof (void *);
+ /* We might have a pointer in a packed struct: however we cannot have more
+ than one such pointer in each pointer-sized portion of the struct, so
+ this is safe. */
+ idx = (attach_to - n->host_start) / sizeof (void *);
+
+ if (!n->attach_count)
+ n->attach_count = gomp_malloc_cleared (sizeof (*n->attach_count) * size);
+
+ if (n->attach_count[idx] < USHRT_MAX)
+ n->attach_count[idx]++;
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count overflow");
+ }
+
+ if (n->attach_count[idx] == 1)
+ {
+ uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + attach_to
+ - n->host_start;
+ uintptr_t target = (uintptr_t) *(void **) attach_to;
+ splay_tree_key tn;
+ uintptr_t data;
+
+ if ((void *) target == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attempt to attach null pointer");
+ }
+
+ s.host_start = target + bias;
+ s.host_end = s.host_start + 1;
+ tn = splay_tree_lookup (mem_map, &s);
+
+ if (!tn)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("pointer target not mapped for attach");
+ }
+
+ data = tn->tgt->tgt_start + tn->tgt_offset + target - tn->host_start;
+
+ gomp_debug (1,
+ "%s: attaching host %p, target %p (struct base %p) to %p\n",
+ __FUNCTION__, (void *) attach_to, (void *) devptr,
+ (void *) (n->tgt->tgt_start + n->tgt_offset), (void *) data);
+
+ gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &data,
+ sizeof (void *), cbufp);
+ }
+ else
+ gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+ (void *) attach_to, n->attach_count[idx]);
+}
+
+void
+gomp_detach_pointer (struct gomp_device_descr *devicep,
+ struct goacc_asyncqueue *aq, splay_tree_key n,
+ uintptr_t detach_from, bool finalize,
+ struct gomp_coalesce_buf *cbufp)
+{
+ size_t idx;
+
+ if (n == NULL)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("enclosing struct not mapped for detach");
+ }
+
+ idx = (detach_from - n->host_start) / sizeof (void *);
+
+ if (!n->attach_count)
+ gomp_fatal ("no attachment counters for struct");
+
+ if (finalize)
+ n->attach_count[idx] = 1;
+
+ if (n->attach_count[idx] == 0)
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("attach count underflow");
+ }
+ else
+ n->attach_count[idx]--;
+
+ if (n->attach_count[idx] == 0)
+ {
+ uintptr_t devptr = n->tgt->tgt_start + n->tgt_offset + detach_from
+ - n->host_start;
+ uintptr_t target = (uintptr_t) *(void **) detach_from;
+
+ gomp_debug (1,
+ "%s: detaching host %p, target %p (struct base %p) to %p\n",
+ __FUNCTION__, (void *) detach_from, (void *) devptr,
+ (void *) (n->tgt->tgt_start + n->tgt_offset),
+ (void *) target);
+
+ gomp_copy_host2dev (devicep, aq, (void *) devptr, (void *) &target,
+ sizeof (void *), cbufp);
+ }
+ else
+ gomp_debug (1, "%s: attach count for %p -> %u\n", __FUNCTION__,
+ (void *) detach_from, n->attach_count[idx]);
+}
+
+uintptr_t
gomp_map_val (struct target_mem_desc *tgt, void **hostaddrs, size_t i)
{
if (tgt->list[i].key != NULL)
@@ -883,7 +1006,12 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
da->map_index = i;
continue;
}
-
+ else if ((kind & typemask) == GOMP_MAP_ATTACH)
+ {
+ tgt->list[i].key = NULL;
+ has_firstprivate = true;
+ continue;
+ }
cur_node.host_start = (uintptr_t) hostaddrs[i];
if (!GOMP_MAP_POINTER_P (kind & typemask))
cur_node.host_end = cur_node.host_start + sizes[i];
@@ -1141,6 +1269,30 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
cur_node.tgt_offset = n->tgt->tgt_start + n->tgt_offset
+ cur_node.host_start - n->host_start;
continue;
+ case GOMP_MAP_ATTACH:
+ {
+ cur_node.host_start = (uintptr_t) hostaddrs[i];
+ cur_node.host_end = cur_node.host_start + sizeof (void *);
+ splay_tree_key n = splay_tree_lookup (mem_map, &cur_node);
+ if (n != NULL)
+ {
+ tgt->list[i].key = n;
+ tgt->list[i].offset = cur_node.host_start - n->host_start;
+ tgt->list[i].length = n->host_end - n->host_start;
+ tgt->list[i].copy_from = false;
+ tgt->list[i].always_copy_from = false;
+ tgt->list[i].do_detach = true;
+ }
+ else
+ {
+ gomp_mutex_unlock (&devicep->lock);
+ gomp_fatal ("outer struct not mapped for attach");
+ }
+ gomp_attach_pointer (devicep, aq, mem_map, n,
+ (uintptr_t) hostaddrs[i], sizes[i],
+ cbufp);
+ continue;
+ }
default:
break;
}
@@ -1194,10 +1346,12 @@ gomp_map_vars_async (struct gomp_device_descr *devicep,
tgt->list[i].copy_from = GOMP_MAP_COPY_FROM_P (kind & typemask);
tgt->list[i].always_copy_from
= GOMP_MAP_ALWAYS_FROM_P (kind & typemask);
+ tgt->list[i].do_detach = false;
tgt->list[i].offset = 0;
tgt->list[i].length = k->host_end - k->host_start;
k->refcount = 1;
k->dynamic_refcount = 0;
+ k->attach_count = NULL;
tgt->refcount++;
array->left = NULL;
array->right = NULL;
@@ -1482,6 +1636,8 @@ gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
is_tgt_unmapped = true;
gomp_unmap_tgt (k->tgt);
}
+ if (k->attach_count)
+ free (k->attach_count);
return is_tgt_unmapped;
}
@@ -1490,14 +1646,14 @@ gomp_remove_var (struct gomp_device_descr *devicep, splay_tree_key k)
has been done already. */
attribute_hidden void
-gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom)
+gomp_unmap_vars (struct target_mem_desc *tgt, bool do_copyfrom, bool finalize)
{
- gomp_unmap_vars_async (tgt, do_copyfrom, NULL);
+ gomp_unmap_vars_async (tgt, do_copyfrom, NULL, finalize);
}
attribute_hidden void
gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
- struct goacc_asyncqueue *aq)
+ struct goacc_asyncqueue *aq, bool finalize)
{
struct gomp_device_descr *devicep = tgt->device_descr;
@@ -1517,10 +1673,23 @@ gomp_unmap_vars_async (struct target_mem_desc *tgt, bool do_copyfrom,
}
size_t i;
+
+ /* We must perform detachments before any copies back to the host. */
for (i = 0; i < tgt->list_count; i++)
{
splay_tree_key k = tgt->list[i].key;
- if (k == NULL)
+
+ if (k != NULL && tgt->list[i].do_detach)
+ gomp_detach_pointer (devicep, aq, k, tgt->list[i].key->host_start
+ + tgt->list[i].offset, finalize,
+ NULL);
+ }
+
+ for (i = 0; i < tgt->list_count; i++)
+ {
+ splay_tree_key k = tgt->list[i].key;
+
+ if (k == NULL || tgt->list[i].do_detach)
continue;
bool do_unmap = false;
@@ -2139,7 +2308,7 @@ GOMP_target (int device, void (*fn) (void *), const void *unused,
GOMP_MAP_VARS_TARGET);
devicep->run_func (devicep->target_id, fn_addr, (void *) tgt_vars->tgt_start,
NULL);
- gomp_unmap_vars (tgt_vars, true);
+ gomp_unmap_vars (tgt_vars, true, false);
}
/* Like GOMP_target, but KINDS is 16-bit, UNUSED is no longer present,
@@ -2283,7 +2452,7 @@ GOMP_target_ext (int device, void (*fn) (void *), size_t mapnum,
tgt_vars ? (void *) tgt_vars->tgt_start : hostaddrs,
args);
if (tgt_vars)
- gomp_unmap_vars (tgt_vars, true);
+ gomp_unmap_vars (tgt_vars, true, false);
}
/* Host fallback for GOMP_target_data{,_ext} routines. */
@@ -2352,7 +2521,7 @@ GOMP_target_end_data (void)
{
struct target_mem_desc *tgt = icv->target_data;
icv->target_data = tgt->prev;
- gomp_unmap_vars (tgt, true);
+ gomp_unmap_vars (tgt, true, false);
}
}
@@ -2587,7 +2756,7 @@ gomp_target_task_fn (void *data)
if (ttask->state == GOMP_TARGET_TASK_FINISHED)
{
if (ttask->tgt)
- gomp_unmap_vars (ttask->tgt, true);
+ gomp_unmap_vars (ttask->tgt, true, false);
return false;
}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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;
+}
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
new file mode 100644
@@ -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
@@ -20,9 +20,9 @@
!$acc end data
do i = 1, n
- if (d(i)%a /= i) call abort
- if (d(i)%b /= i-1) call abort
- if (d(i)%c /= i+1) call abort
+ if (d(i)%a /= i) stop 1
+ if (d(i)%b /= i-1) stop 2
+ if (d(i)%c /= i+1) stop 3
end do
end program
@@ -37,7 +37,7 @@ program derived_acc
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 1
var%b = 100
@@ -51,7 +51,7 @@ program derived_acc
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 2
!$acc parallel loop present (var)
do i = 1, n
@@ -64,7 +64,7 @@ program derived_acc
var%a = -1
do i = 1, n
- if (var%c(i) /= i) call abort
+ if (var%c(i) /= i) stop 3
var%c(i) = var%a
end do
@@ -78,7 +78,7 @@ program derived_acc
if (var%c(i) /= var%a) res = res + 1
end do
- if (res /= 0) call abort
+ if (res /= 0) stop 4
var%c(:) = 0
@@ -93,8 +93,8 @@ program derived_acc
!$acc update host(var%c(5))
do i = 1, n
- if (i /= 5 .and. var%c(i) /= 0) call abort
- if (i == 5 .and. var%c(i) /= 1) call abort
+ if (i /= 5 .and. var%c(i) /= 0) stop 5
+ if (i == 5 .and. var%c(i) /= 1) stop 6
end do
!$acc parallel loop present(var)
@@ -106,7 +106,7 @@ program derived_acc
!$acc update host(var%in%d)
do i = 1, n
- if (var%in%d(i) /= var%a) call abort
+ if (var%in%d(i) /= var%a) stop 7
end do
var%c(:) = 0
@@ -124,8 +124,8 @@ program derived_acc
!$acc update host(var%c(n/2:n))
do i = 1,n
- if (i < n/2 .and. var%c(i) /= -1) call abort
- if (i >= n/2 .and. var%c(i) /= i) call abort
+ if (i < n/2 .and. var%c(i) /= -1) stop 8
+ if (i >= n/2 .and. var%c(i) /= i) stop 9
end do
var%in%d(:) = 0
@@ -140,8 +140,8 @@ program derived_acc
!$acc update host(var%in%d(5))
do i = 1, n
- if (i /= 5 .and. var%in%d(i) /= 0) call abort
- if (i == 5 .and. var%in%d(i) /= 1) call abort
+ if (i /= 5 .and. var%in%d(i) /= 0) stop 10
+ if (i == 5 .and. var%in%d(i) /= 1) stop 11
end do
!$acc exit data delete(var)
@@ -173,7 +173,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 12
var%b = 100
@@ -187,7 +187,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%a)
- if (var%a /= var%b) call abort
+ if (var%a /= var%b) stop 13
!$acc parallel loop present (var)
do i = 1, n
@@ -200,7 +200,7 @@ subroutine derived_acc_subroutine(var)
var%a = -1
do i = 1, n
- if (var%c(i) /= i) call abort
+ if (var%c(i) /= i) stop 14
var%c(i) = var%a
end do
@@ -214,7 +214,7 @@ subroutine derived_acc_subroutine(var)
if (var%c(i) /= var%a) res = res + 1
end do
- if (res /= 0) call abort
+ if (res /= 0) stop 15
var%c(:) = 0
@@ -229,8 +229,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%c(5))
do i = 1, n
- if (i /= 5 .and. var%c(i) /= 0) call abort
- if (i == 5 .and. var%c(i) /= 1) call abort
+ if (i /= 5 .and. var%c(i) /= 0) stop 16
+ if (i == 5 .and. var%c(i) /= 1) stop 17
end do
!$acc parallel loop present(var)
@@ -242,7 +242,7 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%in%d)
do i = 1, n
- if (var%in%d(i) /= var%a) call abort
+ if (var%in%d(i) /= var%a) stop 18
end do
var%c(:) = 0
@@ -260,8 +260,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%c(n/2:n))
do i = 1,n
- if (i < n/2 .and. var%c(i) /= -1) call abort
- if (i >= n/2 .and. var%c(i) /= i) call abort
+ if (i < n/2 .and. var%c(i) /= -1) stop 19
+ if (i >= n/2 .and. var%c(i) /= i) stop 20
end do
var%in%d(:) = 0
@@ -276,8 +276,8 @@ subroutine derived_acc_subroutine(var)
!$acc update host(var%in%d(5))
do i = 1, n
- if (i /= 5 .and. var%in%d(i) /= 0) call abort
- if (i == 5 .and. var%in%d(i) /= 1) call abort
+ if (i /= 5 .and. var%in%d(i) /= 0) stop 21
+ if (i == 5 .and. var%in%d(i) /= 1) stop 22
end do
!$acc exit data delete(var)