FreeChainXenon/gcc
1
0
Fork 0
Code Issues Pull requests Projects Releases Packages Wiki Activity Actions

[OpenMP, Fortran] Add structure/derived-type element mapping

Browse source 
gcc/fortran/ChangeLog:

	* openmp.c (gfc_match_omp_clauses): Match also derived-type
	component refs in OMP_CLAUSE_MAP.
	(resolve_omp_clauses): Resolve those.
	* trans-openmp.c (gfc_trans_omp_array_section, gfc_trans_omp_clauses):
	Handle OpenMP structure-element mapping.
	(gfc_trans_oacc_construct, gfc_trans_oacc_executable_directive,
	(gfc_trans_oacc_combined_directive, gfc_trans_oacc_declare): Update
	add openacc=true in gfc_trans_omp_clauses call.

gcc/testsuite/ChangeLog:

	* gfortran.dg/goacc/finalize-1.f: Update dump scan pattern.
	* gfortran.dg/gomp/map-1.f90: Update dg-error.
	* gfortran.dg/gomp/map-2.f90: New test.

libgomp/ChangeLog:

	* testsuite/libgomp.fortran/struct-elem-map-1.f90: New test.
This commit is contained in:
Tobias Burnus 2020-07-14 13:39:46 +02:00
parent 174e79bf73
commit 102502e32e
6 changed files with 595 additions and 118 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

5
gcc/fortran/openmp.c
View file

@ -1464,7 +1464,7 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask,
head = NULL;
if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP],
false, NULL, &head,
true) == MATCH_YES)
true, true) == MATCH_YES)
{
gfc_omp_namelist *n;
for (n = *head; n; n = n->next)
@ -4553,7 +4553,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
/* Look through component refs to find last array
reference. */
if (openacc && resolved)
if (resolved)
{
/* The "!$acc cache" directive allows rectangular
subarrays to be specified, with some restrictions
@ -4563,6 +4563,7 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses,
arr(-n:n,-n:n) could be contiguous even if it looks
like it may not be. */
if (list != OMP_LIST_CACHE
&& list != OMP_LIST_DEPEND
&& !gfc_is_simply_contiguous (n->expr, false, true)
&& gfc_is_not_contiguous (n->expr))
gfc_error ("Array is not contiguous at %L",

332
gcc/fortran/trans-openmp.c
View file

@ -2092,10 +2092,11 @@ static vec<tree, va_heap, vl_embed> *doacross_steps;
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)
tree &node, tree &node2, tree &node3, tree &node4)
{
gfc_se se;
tree ptr, ptr2;
tree elemsz = NULL_TREE;
gfc_init_se (&se, NULL);
@ -2104,7 +2105,8 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
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));
OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr)));
elemsz = OMP_CLAUSE_SIZE (node);
}
else
{
@ -2114,14 +2116,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
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 = 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);
gcc_assert (se.post.head == NULL_TREE);
ptr = fold_convert (build_pointer_type (char_type_node), ptr);
OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (ptr);
ptr = fold_convert (ptrdiff_type_node, ptr);
if (POINTER_TYPE_P (TREE_TYPE (decl))
&& GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl)))
@ -2134,28 +2137,71 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
OMP_CLAUSE_SIZE (node4) = size_int (0);
decl = build_fold_indirect_ref (decl);
}
ptr = fold_convert (sizetype, ptr);
else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER
&& n->expr->ts.type == BT_CHARACTER
&& n->expr->ts.deferred)
{
gomp_map_kind map_kind;
if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
map_kind = GOMP_MAP_TO;
else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE
|| OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE)
map_kind = OMP_CLAUSE_MAP_KIND (node);
else
map_kind = GOMP_MAP_ALLOC;
gcc_assert (se.string_length);
node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
OMP_CLAUSE_DECL (node4) = se.string_length;
OMP_CLAUSE_SIZE (node4) = TYPE_SIZE_UNIT (gfc_charlen_type_node);
}
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl)))
{
tree desc_node;
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);
desc_node = build_omp_clause (input_location, OMP_CLAUSE_MAP);
OMP_CLAUSE_DECL (desc_node) = decl;
OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
if (ptr_kind == GOMP_MAP_ALWAYS_POINTER)
{
OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO);
node2 = node;
node = desc_node; /* Needs to come first. */
}
else
{
OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET);
node2 = desc_node;
}
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);
/* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra
cast prevents gimplify.c from recognising it as being part of the
struct – and adding an 'alloc: for the 'desc.data' pointer, which
would break as the 'desc' (the descriptor) is also mapped
(see node4 above). */
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);
{
tree offset;
ptr2 = build_fold_addr_expr (decl);
offset = fold_build2 (MINUS_EXPR, ptrdiff_type_node, ptr,
fold_convert (ptrdiff_type_node, ptr2));
offset = build2 (TRUNC_DIV_EXPR, ptrdiff_type_node,
offset, fold_convert (ptrdiff_type_node, elemsz));
offset = build4_loc (input_location, ARRAY_REF,
TREE_TYPE (TREE_TYPE (decl)),
decl, offset, NULL_TREE, NULL_TREE);
OMP_CLAUSE_DECL (node) = offset;
}
else
{
gcc_assert (POINTER_TYPE_P (TREE_TYPE (decl)));
@ -2166,14 +2212,15 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n,
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);
ptr2 = fold_convert (ptrdiff_type_node, ptr2);
OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node,
ptr, ptr2);
}
static tree
gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
locus where, bool declare_simd = false)
locus where, bool declare_simd = false,
bool openacc = false)
{
tree omp_clauses = NULL_TREE, chunk_size, c;
int list, ifc;
@ -2488,6 +2535,67 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
tree node2 = NULL_TREE;
tree node3 = NULL_TREE;
tree node4 = NULL_TREE;
switch (n->u.map_op)
{
case OMP_MAP_ALLOC:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALLOC);
break;
case OMP_MAP_IF_PRESENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
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;
case OMP_MAP_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
break;
case OMP_MAP_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
break;
case OMP_MAP_ALWAYS_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
break;
case OMP_MAP_ALWAYS_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
break;
case OMP_MAP_ALWAYS_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
break;
case OMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
break;
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;
case OMP_MAP_FORCE_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
break;
case OMP_MAP_FORCE_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
break;
case OMP_MAP_FORCE_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
break;
case OMP_MAP_FORCE_PRESENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
break;
case OMP_MAP_FORCE_DEVICEPTR:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
break;
default:
gcc_unreachable ();
}
tree decl = gfc_trans_omp_variable (n->sym, false);
if (DECL_P (decl))
TREE_ADDRESSABLE (decl) = 1;
@ -2496,7 +2604,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
&& n->expr->ref->u.ar.type == AR_FULL))
{
tree present = gfc_omp_check_optional_argument (decl, true);
if (n->sym->ts.type == BT_CLASS)
if (openacc && n->sym->ts.type == BT_CLASS)
{
tree type = TREE_TYPE (decl);
if (n->sym->attr.optional)
@ -2724,8 +2832,42 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
/* Last component is a scalar. */
gfc_conv_expr (&se, n->expr);
gfc_add_block_to_block (block, &se.pre);
OMP_CLAUSE_DECL (node) = se.expr;
/* For BT_CHARACTER a pointer is returned. */
OMP_CLAUSE_DECL (node)
= POINTER_TYPE_P (TREE_TYPE (se.expr))
? build_fold_indirect_ref (se.expr) : se.expr;
gfc_add_block_to_block (block, &se.post);
if (sym_attr.pointer || sym_attr.allocatable)
{
node2 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node2,
openacc
? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_DECL (node2)
= POINTER_TYPE_P (TREE_TYPE (se.expr))
? se.expr : gfc_build_addr_expr (NULL, se.expr);
OMP_CLAUSE_SIZE (node2) = size_int (0);
if (!openacc
&& n->expr->ts.type == BT_CHARACTER
&& n->expr->ts.deferred)
{
gcc_assert (se.string_length);
tree tmp = gfc_get_char_type (n->expr->ts.kind);
OMP_CLAUSE_SIZE (node)
= fold_build2 (MULT_EXPR, size_type_node,
fold_convert (size_type_node,
se.string_length),
TYPE_SIZE_UNIT (tmp));
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO);
OMP_CLAUSE_DECL (node3) = se.string_length;
OMP_CLAUSE_SIZE (node3)
= TYPE_SIZE_UNIT (gfc_charlen_type_node);
}
}
goto finalize_map_clause;
}
@ -2752,7 +2894,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (lastcomp->u.c.component->ts.type == BT_DERIVED
|| lastcomp->u.c.component->ts.type == BT_CLASS)
{
if (sym_attr.allocatable || sym_attr.pointer)
if (sym_attr.pointer || (openacc && sym_attr.allocatable))
{
tree data, size;
@ -2773,7 +2915,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
node2 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node2,
GOMP_MAP_ATTACH_DETACH);
openacc
? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_DECL (node2) = data;
OMP_CLAUSE_SIZE (node2) = size_int (0);
}
@ -2800,32 +2944,82 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner)))
{
gomp_map_kind map_kind;
tree desc_node;
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));
if (GOMP_MAP_COPY_TO_P (OMP_CLAUSE_MAP_KIND (node)))
map_kind = GOMP_MAP_TO;
else if (n->u.map_op == OMP_MAP_RELEASE
|| n->u.map_op == OMP_MAP_DELETE)
map_kind = OMP_CLAUSE_MAP_KIND (node);
else
map_kind = GOMP_MAP_ALLOC;
if (!openacc
&& n->expr->ts.type == BT_CHARACTER
&& n->expr->ts.deferred)
{
gcc_assert (se.string_length);
tree len = fold_convert (size_type_node,
se.string_length);
elemsz = gfc_get_char_type (n->expr->ts.kind);
elemsz = TYPE_SIZE_UNIT (elemsz);
elemsz = fold_build2 (MULT_EXPR, size_type_node,
len, elemsz);
node4 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node4, map_kind);
OMP_CLAUSE_DECL (node4) = se.string_length;
OMP_CLAUSE_SIZE (node4)
= TYPE_SIZE_UNIT (gfc_charlen_type_node);
}
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);
desc_node = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
if (openacc)
OMP_CLAUSE_SET_MAP_KIND (desc_node,
GOMP_MAP_TO_PSET);
else
OMP_CLAUSE_SET_MAP_KIND (desc_node, map_kind);
OMP_CLAUSE_DECL (desc_node) = inner;
OMP_CLAUSE_SIZE (desc_node) = TYPE_SIZE_UNIT (type);
if (openacc)
node2 = desc_node;
else
{
node2 = node;
node = desc_node; /* Put first. */
}
node3 = build_omp_clause (input_location,
OMP_CLAUSE_MAP);
OMP_CLAUSE_SET_MAP_KIND (node3,
openacc
? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
OMP_CLAUSE_DECL (node3)
= gfc_conv_descriptor_data_get (inner);
/* Similar to gfc_trans_omp_array_section (details
there), we add/keep the cast for OpenMP to prevent
that an 'alloc:' gets added for node3 ('desc.data')
as that is part of the whole descriptor (node3).
TODO: Remove once the ME handles this properly. */
if (!openacc)
OMP_CLAUSE_DECL (node3)
= fold_convert (TREE_TYPE (TREE_OPERAND(ptr, 0)),
OMP_CLAUSE_DECL (node3));
else
STRIP_NOPS (OMP_CLAUSE_DECL (node3));
OMP_CLAUSE_SIZE (node3) = size_int (0);
}
else
OMP_CLAUSE_DECL (node) = inner;
@ -2837,9 +3031,11 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
&& lastcomp->next->type == REF_ARRAY
&& lastcomp->next->u.ar.type == AR_ELEMENT);
gomp_map_kind kind = (openacc ? GOMP_MAP_ATTACH_DETACH
: GOMP_MAP_ALWAYS_POINTER);
gfc_trans_omp_array_section (block, n, inner, element,
GOMP_MAP_ATTACH_DETACH,
node, node2, node3, node4);
kind, node, node2, node3,
node4);
}
}
else /* An array element or array section. */
@ -2851,65 +3047,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses,
}
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_IF_PRESENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_IF_PRESENT);
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;
case OMP_MAP_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FROM);
break;
case OMP_MAP_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_TOFROM);
break;
case OMP_MAP_ALWAYS_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TO);
break;
case OMP_MAP_ALWAYS_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_FROM);
break;
case OMP_MAP_ALWAYS_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_ALWAYS_TOFROM);
break;
case OMP_MAP_RELEASE:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE);
break;
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;
case OMP_MAP_FORCE_TO:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TO);
break;
case OMP_MAP_FORCE_FROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_FROM);
break;
case OMP_MAP_FORCE_TOFROM:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_TOFROM);
break;
case OMP_MAP_FORCE_PRESENT:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_PRESENT);
break;
case OMP_MAP_FORCE_DEVICEPTR:
OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_FORCE_DEVICEPTR);
break;
default:
gcc_unreachable ();
}
omp_clauses = gfc_trans_add_clause (node, omp_clauses);
if (node2)
omp_clauses = gfc_trans_add_clause (node2, omp_clauses);
@ -3661,7 +3799,7 @@ gfc_trans_oacc_construct (gfc_code *code)
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
code->loc);
code->loc, false, true);
stmt = gfc_trans_omp_code (code->block->next, true);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);
@ -3697,7 +3835,7 @@ gfc_trans_oacc_executable_directive (gfc_code *code)
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses,
code->loc);
code->loc, false, true);
stmt = build1_loc (input_location, construct_code, void_type_node,
oacc_clauses);
gfc_add_expr_to_block (&block, stmt);
@ -4522,7 +4660,7 @@ gfc_trans_oacc_combined_directive (gfc_code *code)
if (construct_code == OACC_KERNELS)
construct_clauses.lists[OMP_LIST_REDUCTION] = NULL;
oacc_clauses = gfc_trans_omp_clauses (&block, &construct_clauses,
code->loc);
code->loc, false, true);
}
if (!loop_clauses.seq)
pblock = &block;
@ -5703,7 +5841,7 @@ gfc_trans_oacc_declare (gfc_code *code)
gfc_start_block (&block);
oacc_clauses = gfc_trans_omp_clauses (&block, code->ext.oacc_declare->clauses,
code->loc);
code->loc, false, true);
stmt = gfc_trans_omp_code (code->block->next, true);
stmt = build2_loc (input_location, construct_code, void_type_node, stmt,
oacc_clauses);

4
gcc/testsuite/gfortran.dg/goacc/finalize-1.f
View file

@ -20,7 +20,7 @@
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:del_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5))
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm\\.0\\.data - \\(integer\\(kind=8\\)\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_r)
@ -32,6 +32,6 @@
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:cpo_f \\\[len: \[0-9\]+\\\]\\) finalize$" 1 "gimple" } }
!$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(integer\\(kind=8\\)\\) parm\\.1\\.data - \\(integer\\(kind=8\\)\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } }
! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } }
END SUBROUTINE f

35
gcc/testsuite/gfortran.dg/gomp/map-1.f90
View file

@ -57,18 +57,20 @@ subroutine test(aas)
!$omp target map(j(:))
!$omp end target
!$omp target map(j(1:9:2)) ! { dg-error "Stride should not be specified for array section in MAP clause" }
!$omp target map(j(1:9:2))
! { dg-error "Array is not contiguous" "" { target *-*-* } 60 }
! { dg-error "Stride should not be specified for array section in MAP clause" "" { target *-*-* } 60 }
!$omp end target
!$omp target map(aas(5:))
!$omp end target
! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 63 }
! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 63 }
! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 65 }
! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 65 }
!$omp target map(aas(:))
!$omp end target
! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 68 }
! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 68 }
! { dg-error "Rightmost upper bound of assumed size array section not specified" "" { target *-*-* } 70 }
! { dg-error "'aas' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 70 }
!$omp target map(aas) ! { dg-error "Assumed size array" }
!$omp end target
@ -81,29 +83,28 @@ subroutine test(aas)
!$omp target map(k(5:))
!$omp end target
! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 82 }
! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 82 }
! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 84 }
! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 84 }
!$omp target map(k(5:,:,3))
!$omp end target
! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 87 }
! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 87 }
! { dg-error "Rank mismatch in array reference" "" { target *-*-* } 89 }
! { dg-error "'k' in MAP clause at \\\(1\\\) is not a proper array section" "" { target *-*-* } 89 }
!$omp target map(tt)
!$omp end target
!$omp target map(tt%i) ! { dg-error "Syntax error in OpenMP variable list" }
!$omp target map(tt%k) ! { dg-error "not a member of" }
!$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
!$omp target map(tt%j) ! { dg-error "Syntax error in OpenMP variable list" }
!$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
!$omp target map(tt%j)
!$omp end target
! broken test
!$omp target map(tt%j(1)) ! { dg-error "Syntax error in OpenMP variable list" }
!$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
!$omp target map(tt%j(1))
!$omp end target
!$omp target map(tt%j(1:)) ! { dg-error "Syntax error in OpenMP variable list" }
!$omp end target ! { dg-error "Unexpected !\\\$OMP END TARGET statement" }
!$omp target map(tt%j(1:))
!$omp end target
!$omp target map(tp) ! { dg-error "THREADPRIVATE object 'tp' in MAP clause" }
!$omp end target

6
gcc/testsuite/gfortran.dg/gomp/map-2.f90 Normal file
View file

@ -0,0 +1,6 @@
type t
integer :: i
end type t
type(t) v
!$omp target enter data map(to:v%i, v%i) ! { dg-error "appears more than once in map clauses" }
end

331
libgomp/testsuite/libgomp.fortran/struct-elem-map-1.f90 Normal file
View file

@ -0,0 +1,331 @@
! { dg-do run }
!
! Test OpenMP 4.5 structure-element mapping
! TODO: character(kind=4,...) needs to be tested, but depends on
! PR fortran/95837
! TODO: ...%str4 should be tested but that currently fails due to
! PR fortran/95868 (see commented lined)
! TODO: Test also array-valued var, nested derived types,
! type-extended types.
program main
implicit none
type t2
integer :: a, b
! For complex, assume small integers are exactly representable
complex(kind=8) :: c
integer :: d(10)
integer, pointer :: e => null(), f(:) => null()
character(len=5) :: str1
character(len=5) :: str2(4)
character(len=:), pointer :: str3 => null()
character(len=:), pointer :: str4(:) => null()
end type t2
integer :: i
call one ()
call two ()
call three ()
call four ()
call five ()
call six ()
call seven ()
call eight ()
contains
! Implicitly mapped – but no pointers are mapped
subroutine one()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "one" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%e, source=99)
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str3, source="HelloWorld")
allocate (var%str4, source=["Let's", "Go!!!"])
!$omp target map(tofrom:var)
if (var%a /= 1) stop 1
if (var%b /= 2) stop 2
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
if (var%str1 /= "abcde") stop 5
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
!$omp end target
deallocate(var%e, var%f, var%str3, var%str4)
end subroutine one
! Explicitly mapped – all and full arrays
subroutine two()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "two" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%e, source=99)
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str3, source="HelloWorld")
allocate (var%str4, source=["Let's", "Go!!!"])
!$omp target map(tofrom: var%a, var%b, var%c, var%d, var%e, var%f, &
!$omp& var%str1, var%str2, var%str3, var%str4)
if (var%a /= 1) stop 1
if (var%b /= 2) stop 2
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
if (var%str1 /= "abcde") stop 5
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
if (.not. associated (var%e)) stop 7
if (var%e /= 99) stop 8
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f /= [22, 33, 44, 55])) stop 11
if (.not. associated (var%str3)) stop 12
if (len (var%str3) /= len ("HelloWorld")) stop 13
if (var%str3 /= "HelloWorld") stop 14
if (.not. associated (var%str4)) stop 15
if (len (var%str4) /= 5) stop 16
if (size (var%str4) /= 2) stop 17
if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
!$omp end target
deallocate(var%e, var%f, var%str3, var%str4)
end subroutine two
! Explicitly mapped – one by one but full arrays
subroutine three()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "three" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%e, source=99)
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str3, source="HelloWorld")
allocate (var%str4, source=["Let's", "Go!!!"])
!$omp target map(tofrom: var%a)
if (var%a /= 1) stop 1
!$omp end target
!$omp target map(tofrom: var%b)
if (var%b /= 2) stop 2
!$omp end target
!$omp target map(tofrom: var%c)
if (var%c%re /= -1.0_8 .or. var%c%im /= 2.0_8) stop 3
!$omp end target
!$omp target map(tofrom: var%d)
if (any (var%d /= [(-3*i, i = 1, 10)])) stop 4
!$omp end target
!$omp target map(tofrom: var%str1)
if (var%str1 /= "abcde") stop 5
!$omp end target
!$omp target map(tofrom: var%str2)
if (any (var%str2 /= ["12345", "67890", "ABCDE", "FGHIJ"])) stop 6
!$omp end target
!$omp target map(tofrom: var%e)
if (.not. associated (var%e)) stop 7
if (var%e /= 99) stop 8
!$omp end target
!$omp target map(tofrom: var%f)
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f /= [22, 33, 44, 55])) stop 11
!$omp end target
!$omp target map(tofrom: var%str3)
if (.not. associated (var%str3)) stop 12
if (len (var%str3) /= len ("HelloWorld")) stop 13
if (var%str3 /= "HelloWorld") stop 14
!$omp end target
!$omp target map(tofrom: var%str4)
if (.not. associated (var%str4)) stop 15
if (len (var%str4) /= 5) stop 16
if (size (var%str4) /= 2) stop 17
if (any (var%str4 /= ["Let's", "Go!!!"])) stop 18
!$omp end target
deallocate(var%e, var%f, var%str3, var%str4)
end subroutine three
! Explicitly mapped – all but only subarrays
subroutine four()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "four" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str4, source=["Let's", "Go!!!"])
! !$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3), var%str4(2:2))
!$omp target map(tofrom: var%d(4:7), var%f(2:3), var%str2(2:3))
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f(2:3) /= [33, 44])) stop 11
! if (.not. associated (var%str4)) stop 15
! if (len (var%str4) /= 5) stop 16
! if (size (var%str4) /= 2) stop 17
! if (var%str4(2) /= "Go!!!") stop 18
!$omp end target
deallocate(var%f, var%str4)
end subroutine four
! Explicitly mapped – all but only subarrays and one by one
subroutine five()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "five" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str4, source=["Let's", "Go!!!"])
!$omp target map(tofrom: var%d(4:7))
if (any (var%d(4:7) /= [(-3*i, i = 4, 7)])) stop 4
!$omp end target
!$omp target map(tofrom: var%str2(2:3))
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
!$omp end target
!$omp target map(tofrom: var%f(2:3))
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f(2:3) /= [33, 44])) stop 11
!$omp end target
! !$omp target map(tofrom: var%str4(2:2))
! if (.not. associated (var%str4)) stop 15
! if (len (var%str4) /= 5) stop 16
! if (size (var%str4) /= 2) stop 17
! if (var%str4(2) /= "Go!!!") stop 18
! !$omp end target
deallocate(var%f, var%str4)
end subroutine five
! Explicitly mapped – all but only array elements
subroutine six()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "six" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str4, source=["Let's", "Go!!!"])
! !$omp target map(tofrom: var%d(5), var%f(3), var%str2(3), var%str4(2))
!$omp target map(tofrom: var%d(5), var%f(3), var%str2(3))
if (var%d(5) /= -3*5) stop 4
if (var%str2(3) /= "ABCDE") stop 6
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (var%f(3) /= 44) stop 11
! if (.not. associated (var%str4)) stop 15
! if (len (var%str4) /= 5) stop 16
! if (size (var%str4) /= 2) stop 17
! if (var%str4(2) /= "Go!!!") stop 18
!$omp end target
deallocate(var%f, var%str4)
end subroutine six
! Explicitly mapped – all but only array elements and one by one
subroutine seven()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "seven" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
allocate (var%f, source=[22, 33, 44, 55])
allocate (var%str4, source=["Let's", "Go!!!"])
!$omp target map(tofrom: var%d(5))
if (var%d(5) /= (-3*5)) stop 4
!$omp end target
!$omp target map(tofrom: var%str2(2:3))
if (any (var%str2(2:3) /= ["67890", "ABCDE"])) stop 6
!$omp end target
!$omp target map(tofrom: var%f(2:3))
if (.not. associated (var%f)) stop 9
if (size (var%f) /= 4) stop 10
if (any (var%f(2:3) /= [33, 44])) stop 11
!$omp end target
! !$omp target map(tofrom: var%str4(2:2))
! if (.not. associated (var%str4)) stop 15
! if (len (var%str4) /= 5) stop 16
! if (size (var%str4) /= 2) stop 17
! if (var%str4(2) /= "Go!!!") stop 18
! !$omp end target
deallocate(var%f, var%str4)
end subroutine seven
! Check mapping of NULL pointers
subroutine eight()
type(t2) :: var, var2(4)
type(t2), pointer :: var3, var4(:)
print '(g0)', '==== TESTCASE "eight" ===='
var = t2(a = 1, &
b = 2, c = cmplx(-1.0_8, 2.0_8,kind=8), &
d = [(-3*i, i = 1, 10)], &
str1 = "abcde", &
str2 = ["12345", "67890", "ABCDE", "FGHIJ"])
! !$omp target map(tofrom: var%e, var%f, var%str3, var%str4)
!$omp target map(tofrom: var%e, var%str3)
if (associated (var%e)) stop 1
! if (associated (var%f)) stop 2
if (associated (var%str3)) stop 3
! if (associated (var%str4)) stop 4
!$omp end target
end subroutine eight
end program main