https://gcc.gnu.org/g:bbd9b7bbd704b8a58fc93685fefb839b4708c700
commit bbd9b7bbd704b8a58fc93685fefb839b4708c700 Author: Kwok Cheung Yeung <[email protected]> Date: Sun Jun 7 01:37:16 2026 +0000 openmp, fortran: Add support for map iterators in OpenMP target construct (Fortran) This adds support for iterators in map clauses within OpenMP 'target' constructs in Fortran. Some special handling for struct field maps has been added to libgomp in order to handle arrays of derived types. gcc/fortran/ * dump-parse-tree.cc (show_omp_namelist): Add iterator support for OMP_LIST_MAP. * match.cc (gfc_free_omp_namelist): Free namespace in namelist for OMP_LIST_MAP. * openmp.cc (gfc_match_omp_clauses): Parse 'iterator' modifier for 'map' clause. (resolve_omp_clauses): Resolve iterators for OMP_LIST_MAP. * trans-openmp.cc: Include tree-ssa-loop-niter.h. (gfc_trans_omp_array_section): Add iterator argument. Replace instances of iterator variables with the initial value when computing biases. (gfc_trans_omp_clauses): Handle iterators in OMP_LIST_MAP clauses. Add expressions to iter_block rather than block. Do not apply iterators to firstprivate maps. Pass iterator to gfc_trans_omp_array_section. gcc/ * gimplify.cc (compute_omp_iterator_count): Account for difference in loop boundaries in Fortran. (build_omp_iterator_loop): Change upper boundary condition for Fortran. Insert block statements into innermost loop. (remove_unused_omp_iterator_vars): Copy block subblocks of old iterator to new iterator and remove original. (extract_base_bit_offset): Add iterator argument. Remove iterator variables from base. (omp_accumulate_sibling_list): Add iterator argument to extract_base_bit_offset. * tree-pretty-print.cc (dump_block_node): Ignore BLOCK_SUBBLOCKS containing iterator block statements. gcc/testsuite/ * gfortran.dg/gomp/target-map-iterators-1.f90: New. * gfortran.dg/gomp/target-map-iterators-2.f90: New. * gfortran.dg/gomp/target-map-iterators-3.f90: New. * gfortran.dg/gomp/target-map-iterators-4.f90: New. * gfortran.dg/gomp/target-map-iterators-6.f90: New. libgomp/ * target.c (kind_to_name): Handle GOMP_MAP_STRUCT and GOMP_MAP_STRUCT_UNORD. (gomp_add_map): New. (gomp_merge_iterator_maps): Expand fields of a struct mapping breadth-first. * testsuite/libgomp.fortran/target-map-iterators-1.f90: New. * testsuite/libgomp.fortran/target-map-iterators-2.f90: New. * testsuite/libgomp.fortran/target-map-iterators-3.f90: New. Co-authored-by: Andrew Stubbs <[email protected]> Co-authored-by: Sandra Loosemore <[email protected]> (cherry picked from commit 665394525f551fa36c43b2b1c110bab368c27376) Diff: --- gcc/fortran/dump-parse-tree.cc | 9 +- gcc/fortran/match.cc | 3 +- gcc/fortran/openmp.cc | 33 ++++++- gcc/fortran/trans-openmp.cc | 108 ++++++++++++++++----- gcc/gimplify.cc | 70 ++++++++++--- .../gfortran.dg/gomp/target-map-iterators-1.f90 | 26 +++++ .../gfortran.dg/gomp/target-map-iterators-2.f90 | 30 ++++++ .../gfortran.dg/gomp/target-map-iterators-3.f90 | 24 +++++ .../gfortran.dg/gomp/target-map-iterators-4.f90 | 28 ++++++ .../gfortran.dg/gomp/target-map-iterators-6.f90 | 26 +++++ gcc/tree-pretty-print.cc | 4 +- libgomp/target.c | 82 +++++++++++----- .../libgomp.fortran/target-map-iterators-1.f90 | 45 +++++++++ .../libgomp.fortran/target-map-iterators-2.f90 | 45 +++++++++ .../libgomp.fortran/target-map-iterators-3.f90 | 50 ++++++++++ 15 files changed, 516 insertions(+), 67 deletions(-) diff --git a/gcc/fortran/dump-parse-tree.cc b/gcc/fortran/dump-parse-tree.cc index 64757d16c801..2263f0a8bd01 100644 --- a/gcc/fortran/dump-parse-tree.cc +++ b/gcc/fortran/dump-parse-tree.cc @@ -1483,7 +1483,8 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) for (; n; n = n->next) { gfc_current_ns = ns_curr; - if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND) + if (list_type == OMP_LIST_AFFINITY || list_type == OMP_LIST_DEPEND + || list_type == OMP_LIST_MAP) { gfc_current_ns = n->u2.ns ? n->u2.ns : ns_curr; if (n->u2.ns != ns_iter) @@ -1495,8 +1496,12 @@ show_omp_namelist (int list_type, gfc_omp_namelist *n) fputs ("AFFINITY (", dumpfile); else if (n->u.depend_doacross_op == OMP_DOACROSS_SINK_FIRST) fputs ("DOACROSS (", dumpfile); - else + else if (list_type == OMP_LIST_DEPEND) fputs ("DEPEND (", dumpfile); + else if (list_type == OMP_LIST_MAP) + fputs ("MAP (", dumpfile); + else + gcc_unreachable (); } if (n->u2.ns) { diff --git a/gcc/fortran/match.cc b/gcc/fortran/match.cc index bc09883ca2f4..0c5e21338f93 100644 --- a/gcc/fortran/match.cc +++ b/gcc/fortran/match.cc @@ -6539,7 +6539,8 @@ gfc_free_namelist (gfc_namelist *name) void gfc_free_omp_namelist (gfc_omp_namelist *name, enum gfc_omp_list_type list) { - bool free_ns = (list == OMP_LIST_AFFINITY || list == OMP_LIST_DEPEND); + bool free_ns = (list == OMP_LIST_AFFINITY || list == OMP_LIST_DEPEND + || list == OMP_LIST_MAP); bool free_align_allocator = (list == OMP_LIST_ALLOCATE); bool free_mem_traits_space = (list == OMP_LIST_USES_ALLOCATORS); bool free_init = (list == OMP_LIST_INIT); diff --git a/gcc/fortran/openmp.cc b/gcc/fortran/openmp.cc index db76d29cac11..6fd4fb6ec1d6 100644 --- a/gcc/fortran/openmp.cc +++ b/gcc/fortran/openmp.cc @@ -3682,9 +3682,12 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, int always_modifier = 0; int close_modifier = 0; int present_modifier = 0; + int iterator_modifier = 0; + gfc_namespace *ns_iter = NULL, *ns_curr = gfc_current_ns; locus second_always_locus = old_loc2; locus second_close_locus = old_loc2; locus second_present_locus = old_loc2; + locus second_iterator_locus = old_loc2; for (;;) { @@ -3704,6 +3707,11 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, if (present_modifier++ == 1) second_present_locus = current_locus; } + else if (gfc_match_iterator (&ns_iter, true) == MATCH_YES) + { + if (iterator_modifier++ == 1) + second_iterator_locus = current_locus; + } else break; if (gfc_match (", ") != MATCH_YES) @@ -3765,15 +3773,29 @@ gfc_match_omp_clauses (gfc_omp_clauses **cp, const omp_mask mask, &second_present_locus); break; } + if (iterator_modifier > 1) + { + gfc_error ("too many %<iterator%> modifiers at %L", + &second_iterator_locus); + break; + } head = NULL; - if (gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP], - false, NULL, &head, - true, true) == MATCH_YES) + if (ns_iter) + gfc_current_ns = ns_iter; + m = gfc_match_omp_variable_list ("", &c->lists[OMP_LIST_MAP], + false, NULL, &head, true, true); + gfc_current_ns = ns_curr; + if (m == MATCH_YES) { gfc_omp_namelist *n; for (n = *head; n; n = n->next) - n->u.map.op = map_op; + { + n->u.map.op = map_op; + n->u2.ns = ns_iter; + if (ns_iter) + ns_iter->refs++; + } continue; } gfc_current_locus = old_loc; @@ -9776,7 +9798,8 @@ resolve_omp_clauses (gfc_code *code, gfc_omp_clauses *omp_clauses, case OMP_LIST_CACHE: for (; n != NULL; n = n->next) { - if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY) + if ((list == OMP_LIST_DEPEND || list == OMP_LIST_AFFINITY + || list == OMP_LIST_MAP) && n->u2.ns && !n->u2.ns->resolved) { n->u2.ns->resolved = 1; diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 30b8e604486d..5f2673bb2e19 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -27,6 +27,7 @@ along with GCC; see the file COPYING3. If not see #include "gfortran.h" #include "basic-block.h" #include "tree-ssa.h" +#include "tree-ssa-loop-niter.h" /* for simplify_replace_tree. */ #include "function.h" #include "gimple.h" #include "gimple-expr.h" @@ -3339,7 +3340,8 @@ static void gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, gfc_omp_namelist *n, tree decl, bool element, bool openmp, gomp_map_kind ptr_kind, tree &node, - tree &node2, tree &node3, tree &node4) + tree &node2, tree &node3, tree &node4, + tree iterator) { gfc_se se; tree ptr, ptr2; @@ -3494,6 +3496,13 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, OMP_CLAUSE_DECL (node3) = decl; } ptr2 = fold_convert (ptrdiff_type_node, ptr2); + for (tree it = iterator; it; it = TREE_CHAIN (it)) + { + ptr = simplify_replace_tree (ptr, TREE_VEC_ELT (it, 0), + TREE_VEC_ELT (it, 1)); + ptr2 = simplify_replace_tree (ptr2, TREE_VEC_ELT (it, 0), + TREE_VEC_ELT (it, 1)); + } OMP_CLAUSE_SIZE (node3) = fold_build2 (MINUS_EXPR, ptrdiff_type_node, ptr, ptr2); } @@ -4306,6 +4315,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } break; case OMP_LIST_MAP: + iterator = NULL_TREE; + prev = NULL; + prev_clauses = omp_clauses; for (; n != NULL; n = n->next) { if (!openacc) @@ -4341,6 +4353,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || n->sym->attr.flavor == FL_PARAMETER) continue; + if (iterator && prev->u2.ns != n->u2.ns) + { + /* Finish previous iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + OMP_CLAUSE_ITERATORS (c) = iterator; + prev_clauses = omp_clauses; + iterator = NULL_TREE; + } + if (n->u2.ns && (!prev || prev->u2.ns != n->u2.ns)) + { + /* Start a new iterator group. */ + gfc_init_block (&iter_block); + tree_block = make_node (BLOCK); + TREE_USED (tree_block) = 1; + BLOCK_VARS (tree_block) = NULL_TREE; + prev_clauses = omp_clauses; + iterator = handle_iterator (n->u2.ns, block, tree_block); + } + if (!iterator) + gfc_init_block (&iter_block); + prev = n; + location_t map_loc = gfc_get_location (&n->where); bool always_modifier = false; tree node = build_omp_clause (map_loc, OMP_CLAUSE_MAP); @@ -4551,7 +4590,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, TRUTH_NOT_EXPR, boolean_type_node, present); - gfc_add_expr_to_block (block, + gfc_add_expr_to_block (&iter_block, build3_loc (input_location, COND_EXPR, void_type_node, @@ -4621,7 +4660,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree type = TREE_TYPE (decl); tree ptr = gfc_conv_descriptor_data_get (decl); if (present) - ptr = gfc_build_cond_assign_expr (block, present, ptr, + ptr = gfc_build_cond_assign_expr (&iter_block, + present, ptr, null_pointer_node); gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr))); ptr = build_fold_indirect_ref (ptr); @@ -4648,7 +4688,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ptr = gfc_conv_descriptor_data_get (decl); ptr = gfc_build_addr_expr (NULL, ptr); ptr = gfc_build_cond_assign_expr ( - block, present, ptr, null_pointer_node); + &iter_block, present, ptr, null_pointer_node); ptr = build_fold_indirect_ref (ptr); OMP_CLAUSE_DECL (node3) = ptr; } @@ -4737,7 +4777,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, TRUTH_ANDIF_EXPR, boolean_type_node, present, cond); - gfc_add_expr_to_block (block, + gfc_add_expr_to_block (&iter_block, build3_loc (input_location, COND_EXPR, void_type_node, @@ -4766,12 +4806,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree cond = build3_loc (input_location, COND_EXPR, void_type_node, present, cond_body, NULL_TREE); - gfc_add_expr_to_block (block, cond); + gfc_add_expr_to_block (&iter_block, cond); OMP_CLAUSE_SIZE (node) = var; } else { - gfc_add_block_to_block (block, &cond_block); + gfc_add_block_to_block (&iter_block, &cond_block); OMP_CLAUSE_SIZE (node) = size; } } @@ -4783,8 +4823,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, /* A single indirectref is handled by the middle end. */ gcc_assert (!POINTER_TYPE_P (TREE_TYPE (decl))); tree tmp = TREE_OPERAND (decl, 0); - tmp = gfc_build_cond_assign_expr (block, present, tmp, - null_pointer_node); + tmp = gfc_build_cond_assign_expr (&iter_block, + present, tmp, + null_pointer_node); OMP_CLAUSE_DECL (node) = build_fold_indirect_ref (tmp); } else @@ -4817,7 +4858,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, size_type_node, cond, size, size_zero_node); - size = gfc_evaluate_now (size, block); + size = gfc_evaluate_now (size, &iter_block); OMP_CLAUSE_SIZE (node) = size; } if ((TREE_CODE (decl) != PARM_DECL @@ -4833,7 +4874,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tmp = DECL_P (decl) ? DECL_SIZE_UNIT (decl) : TYPE_SIZE_UNIT (TREE_TYPE (decl)); tree var = gfc_create_var (TREE_TYPE (tmp), NULL); - gfc_add_modify_loc (input_location, block, var, tmp); + gfc_add_modify_loc (input_location, &iter_block, + var, tmp); OMP_CLAUSE_SIZE (node) = var; gfc_allocate_lang_decl (var); if (TREE_CODE (decl) == INDIRECT_REF) @@ -4863,9 +4905,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, && !(POINTER_TYPE_P (type) && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type)))) k = GOMP_MAP_FIRSTPRIVATE_POINTER; - gfc_trans_omp_array_section (block, op, n, decl, element, - !openacc, k, node, node2, - node3, node4); + gfc_trans_omp_array_section (&iter_block, op, n, decl, + element, !openacc, k, + node, node2, node3, node4, + iterator); } else if (n->expr && n->expr->expr_type == EXPR_VARIABLE @@ -4881,12 +4924,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_init_se (&se, NULL); gfc_conv_expr (&se, n->expr); - gfc_add_block_to_block (block, &se.pre); + gfc_add_block_to_block (&iter_block, &se.pre); /* 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); + gfc_add_block_to_block (&iter_block, &se.post); if (pointer || allocatable) { /* If it's a bare attach/detach clause, we just want @@ -4947,7 +4990,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, ? DECL_SIZE_UNIT (se.expr) : TYPE_SIZE_UNIT (TREE_TYPE (se.expr))); tree var = gfc_create_var (TREE_TYPE (tmp), NULL); - gfc_add_modify_loc (input_location, block, var, tmp); + gfc_add_modify_loc (input_location, &iter_block, + var, tmp); OMP_CLAUSE_SIZE (node) = var; gfc_allocate_lang_decl (var); if (TREE_CODE (se.expr) == INDIRECT_REF) @@ -5106,7 +5150,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to ensure that it is not gimplified + is a decl. */ tree tmp = OMP_CLAUSE_SIZE (node); tree var = gfc_create_var (TREE_TYPE (tmp), NULL); - gfc_add_modify_loc (input_location, block, var, tmp); + gfc_add_modify_loc (input_location, &iter_block, + var, tmp); OMP_CLAUSE_SIZE (node) = var; gfc_allocate_lang_decl (var); if (TREE_CODE (inner) == INDIRECT_REF) @@ -5139,7 +5184,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { bool drop_mapping = gfc_map_array_descriptor ( node, node2, node3, node4, inner, openacc, map_loc, - block, op, n, sym_rooted_nl, se, clauses, false); + &iter_block, op, n, sym_rooted_nl, se, clauses, + false); if (drop_mapping) continue; } @@ -5151,9 +5197,10 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, /* An array element or section. */ bool element = lastref->u.ar.type == AR_ELEMENT; gomp_map_kind kind = GOMP_MAP_ATTACH_DETACH; - gfc_trans_omp_array_section (block, op, n, inner, element, - !openacc, kind, node, node2, - node3, node4); + gfc_trans_omp_array_section (&iter_block, op, n, inner, + element, !openacc, kind, + node, node2, node3, node4, + iterator); } else gcc_unreachable (); @@ -5171,7 +5218,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree node4 = NULL_TREE; gfc_map_array_descriptor (node1, node2, node3, node4, mid_descr[i], openacc, - map_loc, block, op, n, + map_loc, &iter_block, + op, n, sym_rooted_nl, se, clauses, true); @@ -5194,6 +5242,9 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, finalize_map_clause: + if (!iterator) + gfc_add_block_to_block (block, &iter_block); + omp_clauses = gfc_trans_add_clause (node, omp_clauses); if (node2) omp_clauses = gfc_trans_add_clause (node2, omp_clauses); @@ -5204,6 +5255,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, if (node5) omp_clauses = gfc_trans_add_clause (node5, omp_clauses); } + if (iterator) + { + /* Finish last iterator group. */ + BLOCK_SUBBLOCKS (tree_block) = gfc_finish_block (&iter_block); + TREE_VEC_ELT (iterator, 5) = tree_block; + for (tree c = omp_clauses; c != prev_clauses; + c = OMP_CLAUSE_CHAIN (c)) + if (OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_POINTER + && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_FIRSTPRIVATE_REFERENCE) + OMP_CLAUSE_ITERATORS (c) = iterator; + } break; case OMP_LIST_TO: case OMP_LIST_FROM: diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index c8f9adf04626..6b21116ce56e 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -75,6 +75,7 @@ along with GCC; see the file COPYING3. If not see #include "recog.h" #include "output.h" #include "gimplify_reg_info.h" +#include "tree-ssa-loop-niter.h" /* For simplify_replace_tree. */ /* Identifier for a basic condition, mapping it to other basic conditions of its Boolean expression. Basic conditions given the same uid (in the same @@ -9885,10 +9886,17 @@ compute_omp_iterator_count (tree it, gimple_seq *pre_p) endmbegin = fold_build2_loc (loc, POINTER_DIFF_EXPR, stype, end, begin); else endmbegin = fold_build2_loc (loc, MINUS_EXPR, type, end, begin); - tree stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step, - build_int_cst (stype, 1)); - tree stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, - build_int_cst (stype, 1)); + /* Account for iteration stopping on the end value in Fortran rather + than before it. */ + tree stepm1 = step; + tree stepp1 = step; + if (!lang_GNU_Fortran ()) + { + stepm1 = fold_build2_loc (loc, MINUS_EXPR, stype, step, + build_int_cst (stype, 1)); + stepp1 = fold_build2_loc (loc, PLUS_EXPR, stype, step, + build_int_cst (stype, 1)); + } tree pos = fold_build2_loc (loc, PLUS_EXPR, stype, unshare_expr (endmbegin), stepm1); pos = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, pos, step); @@ -9900,10 +9908,11 @@ compute_omp_iterator_count (tree it, gimple_seq *pre_p) } neg = fold_build2_loc (loc, TRUNC_DIV_EXPR, stype, neg, step); step = NULL_TREE; - tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, begin, end); + tree_code cmp_op = lang_GNU_Fortran () ? LE_EXPR : LT_EXPR; + tree cond = fold_build2_loc (loc, cmp_op, boolean_type_node, begin, end); pos = fold_build3_loc (loc, COND_EXPR, stype, cond, pos, build_int_cst (stype, 0)); - cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, end, begin); + cond = fold_build2_loc (loc, cmp_op, boolean_type_node, end, begin); neg = fold_build3_loc (loc, COND_EXPR, stype, cond, neg, build_int_cst (stype, 0)); tree osteptype = TREE_TYPE (orig_step); @@ -9932,6 +9941,7 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) if (*last_bind) gimplify_and_add (*last_bind, pre_p); tree block = TREE_VEC_ELT (it, 5); + tree block_stmts = lang_GNU_Fortran () ? BLOCK_SUBBLOCKS (block) : NULL_TREE; *last_bind = build3 (BIND_EXPR, void_type_node, BLOCK_VARS (block), NULL, block); TREE_SIDE_EFFECTS (*last_bind) = 1; @@ -9943,6 +9953,7 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) tree end = TREE_VEC_ELT (it, 2); tree step = TREE_VEC_ELT (it, 3); tree orig_step = TREE_VEC_ELT (it, 4); + block = TREE_VEC_ELT (it, 5); tree type = TREE_TYPE (var); location_t loc = DECL_SOURCE_LOCATION (var); /* Emit: @@ -9953,9 +9964,9 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) var = var + step; cond_label: if (orig_step > 0) { - if (var < end) goto beg_label; + if (var < end) goto beg_label; // <= for Fortran } else { - if (var > end) goto beg_label; + if (var > end) goto beg_label; // >= for Fortran } for each iterator, with inner iterators added to the ... above. */ @@ -9981,10 +9992,12 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) append_to_statement_list_force (tem, p); tem = build1 (LABEL_EXPR, void_type_node, cond_label); append_to_statement_list (tem, p); - tree cond = fold_build2_loc (loc, LT_EXPR, boolean_type_node, var, end); + tree cond = fold_build2_loc (loc, lang_GNU_Fortran () ? LE_EXPR : LT_EXPR, + boolean_type_node, var, end); tree pos = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, build_and_jump (&beg_label), void_node); - cond = fold_build2_loc (loc, GT_EXPR, boolean_type_node, var, end); + cond = fold_build2_loc (loc, lang_GNU_Fortran () ? GE_EXPR : GT_EXPR, + boolean_type_node, var, end); tree neg = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, build_and_jump (&beg_label), void_node); tree osteptype = TREE_TYPE (orig_step); @@ -9993,6 +10006,11 @@ build_omp_iterator_loop (tree it, gimple_seq *pre_p, tree *last_bind) tem = fold_build3_loc (loc, COND_EXPR, void_type_node, cond, pos, neg); append_to_statement_list_force (tem, p); p = &BIND_EXPR_BODY (bind); + /* The Fortran front-end stashes statements into the BLOCK_SUBBLOCKS + of the last element of the first iterator. These should go into the + body of the innermost loop. */ + if (!TREE_CHAIN (it)) + append_to_statement_list_force (block_stmts, p); } return p; @@ -10126,8 +10144,14 @@ remove_unused_omp_iterator_vars (tree *list_p) i++; } } + tree old_block = TREE_VEC_ELT (OMP_CLAUSE_ITERATORS (c), 5); tree new_block = make_node (BLOCK); BLOCK_VARS (new_block) = new_vars; + if (BLOCK_SUBBLOCKS (old_block)) + { + BLOCK_SUBBLOCKS (new_block) = BLOCK_SUBBLOCKS (old_block); + BLOCK_SUBBLOCKS (old_block) = NULL_TREE; + } TREE_VEC_ELT (new_iters, 5) = new_block; new_iterators.safe_push (new_iters); iter_vars.safe_push (vars.copy ()); @@ -10803,7 +10827,8 @@ build_omp_struct_comp_nodes (enum tree_code code, tree grp_start, tree grp_end, static tree extract_base_bit_offset (tree base, poly_int64 *bitposp, poly_offset_int *poffsetp, - bool *variable_offset) + bool *variable_offset, + tree iterator) { tree offset; poly_int64 bitsize, bitpos; @@ -10813,6 +10838,18 @@ extract_base_bit_offset (tree base, poly_int64 *bitposp, STRIP_NOPS (base); + if (iterator) + { + /* Replace any iterator variables with the lower bound of the iterator. + This will give us the nominal offset and bit position of the first + element, which is all we should need to lay out the mappings. The + actual locations of the iterated mappings are elsewhere. */ + tree it; + for (it = iterator; it; it = TREE_CHAIN (it)) + base = simplify_replace_tree (base, TREE_VEC_ELT (it, 0), + TREE_VEC_ELT (it, 1)); + } + base = get_inner_reference (base, &bitsize, &bitpos, &offset, &mode, &unsignedp, &reversep, &volatilep); @@ -12639,8 +12676,11 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, } bool variable_offset; + tree iterators = OMP_CLAUSE_HAS_ITERATORS (grp_end) + ? OMP_CLAUSE_ITERATORS (grp_end) : NULL_TREE; tree base - = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset); + = extract_base_bit_offset (ocd, &cbitpos, &coffset, &variable_offset, + iterators); int base_token; for (base_token = addr_tokens.length () - 1; base_token >= 0; base_token--) @@ -12969,8 +13009,12 @@ omp_accumulate_sibling_list (enum omp_region_type region_type, sc_decl = TREE_OPERAND (sc_decl, 0); bool variable_offset2; + tree iterators2 = OMP_CLAUSE_HAS_ITERATORS (*sc) + ? OMP_CLAUSE_ITERATORS (*sc) : NULL_TREE; + tree base2 = extract_base_bit_offset (sc_decl, &bitpos, &offset, - &variable_offset2); + &variable_offset2, + iterators2); if (!base2 || !operand_equal_p (base2, base, 0)) break; if (scp) diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 new file mode 100644 index 000000000000..25abbaf741ef --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-1.f90 @@ -0,0 +1,26 @@ +! { dg-do compile } +! { dg-options "-fopenmp" } + +program main + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 39 + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1) + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:)) + !$omp end target + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:), y(i)%ptr(:)) + !$omp end target + + !$omp target map (iterator(i=1:DIM1), to: x(i)%ptr(:) + 3) ! { dg-error "Syntax error in OpenMP variable list at .1." } + !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement at .1." } + + !$omp target map(iterator(i=1:DIM1), iterator(j=1:DIM2), to: x(i)%ptr(j)) ! { dg-error "too many 'iterator' modifiers at .1." } + !$omp end target ! { dg-error "Unexpected \\\!\\\$OMP END TARGET statement at .1." } +end program diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90 new file mode 100644 index 000000000000..59c1d739cbf9 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-2.f90 @@ -0,0 +1,30 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-gimple" } + +program main + implicit none + + integer, parameter :: DIM = 40 + type :: array_ptr + integer, pointer :: ptr(:) + end type + + type (array_ptr) :: x(DIM), y(DIM), z(DIM) + + !$omp target map(iterator(i=1:10), to: x) + ! Add a reference to x to ensure that the 'to' clause does not get dropped. + x(1)%ptr(1) = 0 + !$omp end target + + !$omp target map(iterator(i2=1:10, j2=1:20), from: x(i2)) + !$omp end target + + !$omp target map(iterator(i3=1:10, j3=1:20, k3=1:30), to: x(i3+j3), y(j3+k3), z(k3+i3)) + !$omp end target +end program + +! { dg-final { scan-tree-dump-times "map\\\(to:x" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\) i2=1:10:1, loop_label=\[^\\\)\]+\\\):from:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\) j3=1:20:1, integer\\\(kind=4\\\) i3=1:10:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\) k3=1:30:1, integer\\\(kind=4\\\) j3=1:20:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\\(iterator\\\(integer\\\(kind=4\\\) k3=1:30:1, integer\\\(kind=4\\\) i3=1:10:1, loop_label=\[^\\\)\]+\\\):to:" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 new file mode 100644 index 000000000000..7dad2a69c611 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-3.f90 @@ -0,0 +1,24 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-gimple" } + +program main + implicit none + + integer, parameter :: DIM1 = 17 + integer, parameter :: DIM2 = 27 + type :: ptr_t + integer, pointer :: ptr(:) + end type + + type (ptr_t) :: x(DIM1), y(DIM2) + + !$omp target map(iterator(i=1:DIM1), to: x(i)%ptr(:)) map(iterator(i=1:DIM2), from: y(i)%ptr(:)) + !$omp end target +end program + +! { dg-final { scan-tree-dump-times "if \\(i <= 17\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "if \\(i <= 27\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):to:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):from:MEM <\[^>\]+> \\\[\\\(\[^ \]+ \\\*\\\)D\\\.\[0-9\]+\\\]" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):attach:x\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "map\\(iterator\\(integer\\(kind=4\\) i=1:27:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\.\[0-9\]+\\):attach:y\\\[D\\\.\[0-9\]+\\\]\.ptr\.data" 1 "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90 new file mode 100644 index 000000000000..3359eaad0b48 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-4.f90 @@ -0,0 +1,28 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-gimple" } + +module m + !$omp declare target (baz) + interface + subroutine baz (x, p) + integer, intent(in) :: x + integer, pointer :: p(:) + end subroutine + integer function bar (x, i) + integer :: x, i + end function + end interface +contains + subroutine foo (x, p) + integer :: x + integer, pointer :: p(:) + + !$omp target map (iterator (i=1:4), to: p(bar (x, i))) + call baz (x, p) + !$omp end target + end subroutine +end module + +! { dg-final { scan-tree-dump "firstprivate\\\(x\\\)" "gimple" } } +! { dg-final { scan-tree-dump-times "bar \\\(x, &" 2 "gimple" } } +! { dg-final { scan-tree-dump "map\\\(iterator\\\(integer\\\(kind=4\\\) i=1:4:1, loop_label=" "gimple" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90 b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90 new file mode 100644 index 000000000000..b83481920b60 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/target-map-iterators-6.f90 @@ -0,0 +1,26 @@ +! { dg-do compile } +! { dg-options "-fopenmp -fdump-tree-omplower" } + +module m + implicit none + integer, parameter :: DIM = 17 + type :: array_ptr + integer, pointer :: ptr(:) + end type +contains + subroutine f(x) + type (array_ptr) :: x(DIM) + integer :: i + + !$omp target data map(iterator(it = 1:DIM), tofrom: x(it)%ptr(:)) + do i = 1, DIM + x(i)%ptr(1) = i + end do + !$omp end target data + end subroutine +end module + +! { dg-final { scan-tree-dump "if \\(it <= 17\\) goto <D\\\.\[0-9\]+>; else goto <D\\\.\[0-9\]+>;" "omplower" } } +! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\) it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):to:\\(\\\*x\\\.\[0-9\]+\\)\\\[D\\\.\[0-9\]+\\\]\\\.ptr" "omplower" } } +! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\) it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):tofrom:MEM" "omplower" } } +! { dg-final { scan-tree-dump "map\\(iterator\\(integer\\(kind=\[0-9\]+\\) it=1:17:1, loop_label=<D\\\.\[0-9\]+>, elems=omp_iter_data\\\.\[0-9\]+, index=D\\\.\[0-9\]+\\):attach:\\(\\\*x\\\.\[0-9\]+\\)\\\[D\\\.\[0-9\]+\\\]\\\.ptr.data" "omplower" } } diff --git a/gcc/tree-pretty-print.cc b/gcc/tree-pretty-print.cc index 79e0a734dbb7..0674a27501bf 100644 --- a/gcc/tree-pretty-print.cc +++ b/gcc/tree-pretty-print.cc @@ -1848,7 +1848,9 @@ dump_block_node (pretty_printer *pp, tree block, int spc, dump_flags_t flags) newline_and_indent (pp, spc + 2); } - if (BLOCK_SUBBLOCKS (block)) + if (BLOCK_SUBBLOCKS (block) + && (!lang_GNU_Fortran () + || TREE_CODE (BLOCK_SUBBLOCKS (block)) != STATEMENT_LIST)) { pp_string (pp, "SUBBLOCKS: "); for (t = BLOCK_SUBBLOCKS (block); t; t = BLOCK_CHAIN (t)) diff --git a/libgomp/target.c b/libgomp/target.c index cd25a5b8a4ed..fdc488e94f7a 100644 --- a/libgomp/target.c +++ b/libgomp/target.c @@ -1060,6 +1060,41 @@ kind_to_name (unsigned short kind, bool short_mapkind) } } +static void +gomp_add_map (size_t idx, size_t *new_idx, + void ***hostaddrs, size_t **sizes, unsigned short **skinds, + void ***new_hostaddrs, size_t **new_sizes, + unsigned short **new_kinds, size_t *iterator_count) +{ + if ((*sizes)[idx] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[idx]; + size_t count = *iterator_array++; + for (size_t i = 0; i < count; i++) + { + (*new_hostaddrs)[*new_idx] = (void *) *iterator_array++; + (*new_sizes)[*new_idx] = *iterator_array++; + (*new_kinds)[*new_idx] = (*skinds)[idx]; + iterator_count[*new_idx] = i + 1; + gomp_debug (1, + "Expanding map %u <%s>: " + "hostaddrs[%u] = %p, sizes[%u] = %lu\n", + (int) idx, kind_to_name ((*new_kinds)[*new_idx], true), + (int) *new_idx, (*new_hostaddrs)[*new_idx], + (int) *new_idx, (unsigned long) (*new_sizes)[*new_idx]); + (*new_idx)++; + } + } + else + { + (*new_hostaddrs)[*new_idx] = (*hostaddrs)[idx]; + (*new_sizes)[*new_idx] = (*sizes)[idx]; + (*new_kinds)[*new_idx] = (*skinds)[idx]; + iterator_count[*new_idx] = 0; + (*new_idx)++; + } +} + /* When GCC encounters a clause with an iterator, e.g.: #pragma omp target map (iterator(i=0:4), to: x[i]) @@ -1093,6 +1128,7 @@ gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, bool iterator_p = false; size_t map_count = 0; unsigned short **skinds = (unsigned short **) kinds; + const int typemask = 0xff; for (size_t i = 0; i < *mapnum; i++) if ((*sizes)[i] == SIZE_MAX) @@ -1119,33 +1155,35 @@ gomp_merge_iterator_maps (size_t *mapnum, void ***hostaddrs, size_t **sizes, for (size_t i = 0; i < *mapnum; i++) { - if ((*sizes)[i] == SIZE_MAX) + int map_type = get_kind (true, *skinds, i) & typemask; + if (map_type == GOMP_MAP_STRUCT || map_type == GOMP_MAP_STRUCT_UNORD) { - uintptr_t *iterator_array = (*hostaddrs)[i]; - size_t count = *iterator_array++; - for (size_t j = 0; j < count; j++) + size_t field_count = (*sizes)[i]; + size_t idx_i = new_idx; + + gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, + *iterator_count); + + for (size_t j = i + 1; j <= i + field_count; j++) { - new_hostaddrs[new_idx] = (void *) *iterator_array++; - new_sizes[new_idx] = *iterator_array++; - new_kinds[new_idx] = (*skinds)[i]; - (*iterator_count)[new_idx] = j + 1; - gomp_debug (1, - "Expanding map %u <%s>: " - "hostaddrs[%u] = %p, sizes[%u] = %lu\n", - (int) i, kind_to_name (new_kinds[new_idx], true), - (int) new_idx, new_hostaddrs[new_idx], - (int) new_idx, (unsigned long) new_sizes[new_idx]); - new_idx++; + if ((*sizes)[j] == SIZE_MAX) + { + uintptr_t *iterator_array = (*hostaddrs)[j]; + size_t count = iterator_array[0]; + new_sizes[idx_i] += count - 1; + } + gomp_add_map (j, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, + *iterator_count); } + gomp_debug (1, "Map %u: new field count = %lu\n", + (int) i, (unsigned long) new_sizes[idx_i]); + i += field_count; } else - { - new_hostaddrs[new_idx] = (*hostaddrs)[i]; - new_sizes[new_idx] = (*sizes)[i]; - new_kinds[new_idx] = (*skinds)[i]; - (*iterator_count)[new_idx] = 0; - new_idx++; - } + gomp_add_map (i, &new_idx, hostaddrs, sizes, skinds, + &new_hostaddrs, &new_sizes, &new_kinds, *iterator_count); } *mapnum = map_count; diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 new file mode 100644 index 000000000000..80e077e69fd4 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-1.f90 @@ -0,0 +1,45 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays to target using map +! iterators. + +program test + implicit none + + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: expected, sum, i, j + + expected = mkarray () + + !$omp target map(iterator(i=1:DIM1), to: x(i)%arr(:)) map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + !$omp end target + + if (sum .ne. expected) stop 1 +contains + integer function mkarray () + integer :: exp = 0 + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + exp = exp + x(i)%arr(j) + end do + end do + + mkarray = exp + end function +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 new file mode 100644 index 000000000000..cf0e7fbd9b33 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-2.f90 @@ -0,0 +1,45 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays from target using map +! iterators. + +program test + implicit none + + integer, parameter :: DIM1 = 8 + integer, parameter :: DIM2 = 15 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1) + integer :: expected, sum, i, j + + call mkarray + + !$omp target map(iterator(i=1:DIM1), from: x(i)%arr(:)) map(from: expected) + expected = 0 + do i = 1, DIM1 + do j = 1, DIM2 + x(i)%arr(j) = (i+1) * (j+1) + expected = expected + x(i)%arr(j) + end do + end do + !$omp end target + + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) + end do + end do + + if (sum .ne. expected) stop 1 +contains + subroutine mkarray + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + end do + end subroutine +end program diff --git a/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 new file mode 100644 index 000000000000..b7ba912a8868 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-map-iterators-3.f90 @@ -0,0 +1,50 @@ +! { dg-do run } + +! Test transfer of dynamically-allocated arrays to target using map +! iterators, with multiple iterators and function calls in the iterator +! expression. + +program test + implicit none + + integer, parameter :: DIM1 = 16 + integer, parameter :: DIM2 = 4 + + type :: array_ptr + integer, pointer :: arr(:) + end type + + type (array_ptr) :: x(DIM1), y(DIM1) + integer :: expected, sum, i, j, k + + expected = mkarrays () + + !$omp target map(iterator(i=0:DIM1/4-1, j=0:3), to: x(i * 4 + j + 1)%arr(:)) & + !$omp map(iterator(k=1:DIM1), to: y(k)%arr(:)) & + !$omp map(from: sum) + sum = 0 + do i = 1, DIM1 + do j = 1, DIM2 + sum = sum + x(i)%arr(j) * y(i)%arr(j) + end do + end do + !$omp end target + + if (sum .ne. expected) stop 1 +contains + integer function mkarrays () + integer :: exp = 0 + + do i = 1, DIM1 + allocate (x(i)%arr(DIM2)) + allocate (y(i)%arr(DIM2)) + do j = 1, DIM2 + x(i)%arr(j) = i * j + y(i)%arr(j) = i + j + exp = exp + x(i)%arr(j) * y(i)%arr(j) + end do + end do + + mkarrays = exp + end function +end program
