From patchwork Thu Mar 23 09:28:26 2023 Content-Type: text/plain; charset="utf-8" MIME-Version: 1.0 Content-Transfer-Encoding: 8bit X-Patchwork-Submitter: Tobias Burnus X-Patchwork-Id: 66790 Return-Path: X-Original-To: patchwork@sourceware.org Delivered-To: patchwork@sourceware.org Received: from server2.sourceware.org (localhost [IPv6:::1]) by sourceware.org (Postfix) with ESMTP id E6828387604E for ; Thu, 23 Mar 2023 09:29:03 +0000 (GMT) X-Original-To: gcc-patches@gcc.gnu.org Delivered-To: gcc-patches@gcc.gnu.org Received: from esa3.mentor.iphmx.com (esa3.mentor.iphmx.com [68.232.137.180]) by sourceware.org (Postfix) with ESMTPS id C29BF385840F; Thu, 23 Mar 2023 09:28:37 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org C29BF385840F Authentication-Results: sourceware.org; dmarc=none (p=none dis=none) header.from=codesourcery.com Authentication-Results: sourceware.org; spf=pass smtp.mailfrom=mentor.com X-IronPort-AV: E=Sophos;i="5.98,283,1673942400"; d="diff'?scan'208";a="81081" Received: from orw-gwy-02-in.mentorg.com ([192.94.38.167]) by esa3.mentor.iphmx.com with ESMTP; 23 Mar 2023 01:28:33 -0800 IronPort-SDR: gX4RXtRfIr/jIs6P14YlNDmuZ87QzOpNamhla6xlBoNLvhiGuwuXk4ZDhxZSIjbIoRFG/9ZfNj drxyh/eaHPdwBRVlmf+sK2m1s2dU5/3tqV3o5wXzzzjizlTrbjqZF6ZLrSr3jGMsror3Q8R4Ln HrwqcH8T69jRpeV/l1i6z3PAvvRGxhXnsHyvYd7CMzuRUfgEZqJaZRf38t5udynroycLAnsXVg H4NmZHpP2ewZy7ybLIVgdDZMCeTfnpNccMwbXYS+xOFwC42+0H+weqyVwHVUTb3oy30Rg0A3Pq R/U= Message-ID: <0b7ce95f-d86d-b960-3c20-4a62bdc2be9c@codesourcery.com> Date: Thu, 23 Mar 2023 10:28:26 +0100 MIME-Version: 1.0 User-Agent: Mozilla/5.0 (X11; Linux x86_64; rv:102.0) Gecko/20100101 Thunderbird/102.9.0 Subject: [Patch,v4] Fortran/OpenMP: Fix mapping of array descriptors and deferred-length strings Content-Language: en-US From: Tobias Burnus To: gcc-patches , Jakub Jelinek , fortran References: <57880257-d50f-e72b-5423-932af6c64e92@codesourcery.com> In-Reply-To: X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-12.mgc.mentorg.com (139.181.222.12) X-Spam-Status: No, score=-11.3 required=5.0 tests=BAYES_00, GIT_PATCH_0, HEADER_FROM_DIFFERENT_DOMAINS, KAM_DMARC_STATUS, KAM_SHORT, SPF_HELO_PASS, SPF_PASS, TXREP autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org X-BeenThere: gcc-patches@gcc.gnu.org X-Mailman-Version: 2.1.29 Precedence: list List-Id: Gcc-patches mailing list List-Unsubscribe: , List-Archive: List-Post: List-Help: List-Subscribe: , Errors-To: gcc-patches-bounces+patchwork=sourceware.org@gcc.gnu.org Sender: "Gcc-patches" [GCC 13 vs GCC 14] I am unsure whether this should still go to GCC 13 or not. It is somewhat larger albeit well contained (Fortran, only OpenMP, ...) and fixes real-world bugs, but it is not a regression - and we are meanwhile slowly approaching the release. An alternative would be to go to GCC 14 and then be eventually backported to GCC 14 (albeit I am not sure whether early testing would be better). Or just to GCC 14, hmm. Thoughts? * * * Another update - fixing an independent issue which makes sense to be part of this patch. Allocatable/pointer scalars are currently mapped as: #pragma omp target enter data map(to:*var.1_1 [len: 4]) map(alloc:var [pointer assign, bias: 0]) #pragma omp target exit data map(from:*var.2_2 [len: 4]) where 'GOMP_MAP_POINTER' is removed in gimplify.cc. In v3 (and v4) of this patch, this kind of handling moved from gimplify.cc to fortran/trans-openmp.cc; however, v3 has the same problem. For allocatable arrays, we have PSET + POINTER and the PSET part is changed/set to RELEASE/DELETE for 'exit data' But for scalars, the map was still left on the stack. Besides having a stale map, this could lead to fails when the stack was popped, especially when attempting to later map another stack variable with the same stack address, partially overlapping with the stale POINTER. Side remark: I found this for testcase that is part of an upcoming deep-mapping follow-up patch; that test failed with -O1 but worked with -O0/-Og due to changed stack usage. (Deep-mapping of allocatable components is on the OG12 branch; it is scheduled for mainline integration after stage1 opened.) The updated mainline patch is included; map-10.f90 is the new testcase. If anyone wants to see it separately, the patch has been committed to OG12 as https://gcc.gnu.org/g:8ea805840200f7dfd2c11b37abf5fbfe479c2fe2 Comments/thoughts/remarks to this patch? Tobias PS: For the rest of the patch, see a short description below - or with some longer remarks previous in this thread. On 27.02.23 13:15, Tobias Burnus wrote: > And another re-diff for GCC 13/mainline, updating gcc/testsuite/ > > (The last change is related to the "[OG12,committed] Update dg-dump-scan > for ..." discussion + OG12 https://gcc.gnu.org/g:e4de87a2309 / > https://gcc.gnu.org/pipermail/gcc-patches/2023-February/612871.html ) > > On 23.02.23 17:42, Tobias Burnus wrote: >> On 21.02.23 12:57, Tobias Burnus wrote: >>> This patch moves some generic code for Fortran out of gimplify.cc >>> to trans-openmp.cc and fixes several issues related to mapping. >>> >>> Tested with nvptx offloading. >>> OK for mainline? > Tobias >>> Caveats: >>> >>> Besides the issues shown in the comment-out code, there remains also an >>> issue with implicit mapping - at least for deferred-length strings, >>> but I wouldn't be surprised if - at least depending on the used >>> 'defaultmap' value (e.g. 'alloc') - there are also issues with array >>> descriptors. >>> >>> Note: >>> >>> Regarding the declare target check for mapping: Without declare >>> target, my assumption is that the hidden length variable will >>> get implicitly mapped if needed. Independent of deferred-length >>> or not, there is probably an issue with 'defaultmap(none)' and >>> the hidden variable. - In any case, I prefer to defer all those >>> issues to later (by having them captured in one/several PR). >>> >>> >>> Tobias >>> >>> PS: This patch is a follow up to >>> [Patch] Fortran/OpenMP: Fix DT struct-component with 'alloc' and >>> array descr >>> https://gcc.gnu.org/pipermail/gcc-patches/2022-November/604887.html >>> which fixed part of the problems. But as discussed on IRC, it did >>> treat 'alloc' >>> as special and missed some other map types. - In addition, this patch >>> has a >>> much extended test coverage and fixes some more issues found that way. ----------------- Siemens Electronic Design Automation GmbH; Anschrift: Arnulfstraße 201, 80634 München; Gesellschaft mit beschränkter Haftung; Geschäftsführer: Thomas Heurung, Frank Thürauf; Sitz der Gesellschaft: München; Registergericht München, HRB 106955 Fortran/OpenMP: Fix mapping of array descriptors and deferred-length strings Previously, array descriptors might have been mapped as 'alloc' instead of 'to' for 'alloc', not updating the array bounds. The 'alloc' could also appear for 'data exit', failing with a libgomp assert. In some cases, either array descriptors or deferred-length string's length variable was not mapped. And, finally, some offset calculations with array-sections mappings went wrong. Additionally, the patch now unmaps for scalar allocatables/pointers the GOMP_MAP_POINTER, avoiding stale mappings. The testcases contain some comment-out tests which require follow-up work and for which PR exist. Those mostly relate to deferred-length strings which have several issues beyong OpenMP support. gcc/fortran/ChangeLog: * trans-decl.cc (gfc_get_symbol_decl): Add attributes such as 'declare target' also to hidden artificial variable for deferred-length character variables. * trans-openmp.cc (gfc_trans_omp_array_section, gfc_trans_omp_clauses, gfc_trans_omp_target_exit_data): Improve mapping of array descriptors and deferred-length string variables. gcc/ChangeLog: * gimplify.cc (gimplify_scan_omp_clauses): Remove Fortran special case. libgomp/ChangeLog: * testsuite/libgomp.fortran/target-enter-data-3.f90: Uncomment 'target exit data'. * testsuite/libgomp.fortran/target-enter-data-4.f90: New test. * testsuite/libgomp.fortran/target-enter-data-5.f90: New test. * testsuite/libgomp.fortran/target-enter-data-6.f90: New test. * testsuite/libgomp.fortran/target-enter-data-7.f90: New test. gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update dg-tree; shows a fix for 'finalize' as a ptr is now 'delete' instead of 'release'. * gfortran.dg/gomp/pr78260-2.f90: Likewise as elem-size calc moved to if (allocated) block * gfortran.dg/gomp/target-exit-data.f90: Likewise as a var is now a replaced by a MEM< _25 > expression. * gfortran.dg/gomp/map-9.f90: Update dg-scan-tree-dump. * gfortran.dg/gomp/map-10.f90: New test. gcc/fortran/trans-decl.cc | 2 + gcc/fortran/trans-openmp.cc | 336 +++++++++---- gcc/gimplify.cc | 42 +- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 +- gcc/testsuite/gfortran.dg/gomp/map-10.f90 | 69 +++ gcc/testsuite/gfortran.dg/gomp/map-9.f90 | 2 +- gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 | 6 +- .../gfortran.dg/gomp/target-exit-data.f90 | 4 +- .../libgomp.fortran/target-enter-data-3.f90 | 2 +- .../libgomp.fortran/target-enter-data-4.f90 | 540 +++++++++++++++++++++ .../libgomp.fortran/target-enter-data-5.f90 | 540 +++++++++++++++++++++ .../libgomp.fortran/target-enter-data-6.f90 | 392 +++++++++++++++ .../libgomp.fortran/target-enter-data-7.f90 | 78 +++ 13 files changed, 1877 insertions(+), 140 deletions(-) diff --git a/gcc/fortran/trans-decl.cc b/gcc/fortran/trans-decl.cc index 77610df340b..24d3a65a156 100644 --- a/gcc/fortran/trans-decl.cc +++ b/gcc/fortran/trans-decl.cc @@ -1824,6 +1824,8 @@ gfc_get_symbol_decl (gfc_symbol * sym) /* Add attributes to variables. Functions are handled elsewhere. */ attributes = add_attributes_to_decl (sym->attr, NULL_TREE); decl_attributes (&decl, attributes, 0); + if (sym->ts.deferred) + decl_attributes (&length, attributes, 0); /* Symbols from modules should have their assembler names mangled. This is done here rather than in gfc_finish_var_decl because it diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 84c0184f48e..5f83e476ce6 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -2403,33 +2403,50 @@ static vec *doacross_steps; /* Translate an array section or array element. */ static void -gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, - tree decl, bool element, gomp_map_kind ptr_kind, - tree &node, tree &node2, tree &node3, tree &node4) +gfc_trans_omp_array_section (stmtblock_t *block, gfc_exec_op op, + gfc_omp_namelist *n, tree decl, bool element, + gomp_map_kind ptr_kind, tree &node, tree &node2, + tree &node3, tree &node4) { gfc_se se; tree ptr, ptr2; tree elemsz = NULL_TREE; gfc_init_se (&se, NULL); - if (element) { gfc_conv_expr_reference (&se, n->expr); gfc_add_block_to_block (block, &se.pre); ptr = se.expr; - OMP_CLAUSE_SIZE (node) = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr))); - elemsz = OMP_CLAUSE_SIZE (node); } else { gfc_conv_expr_descriptor (&se, n->expr); ptr = gfc_conv_array_data (se.expr); + } + if (n->expr->ts.type == BT_CHARACTER && n->expr->ts.deferred) + { + gcc_assert (se.string_length); + tree len = gfc_evaluate_now (se.string_length, block); + elemsz = gfc_get_char_type (n->expr->ts.kind); + elemsz = TYPE_SIZE_UNIT (elemsz); + elemsz = fold_build2 (MULT_EXPR, size_type_node, + fold_convert (size_type_node, len), elemsz); + } + if (element) + { + if (!elemsz) + elemsz = TYPE_SIZE_UNIT (TREE_TYPE (TREE_TYPE (ptr))); + OMP_CLAUSE_SIZE (node) = elemsz; + } + else + { tree type = TREE_TYPE (se.expr); gfc_add_block_to_block (block, &se.pre); OMP_CLAUSE_SIZE (node) = gfc_full_array_size (block, se.expr, GFC_TYPE_ARRAY_RANK (type)); - elemsz = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + if (!elemsz) + 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); @@ -2441,7 +2458,11 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, if (POINTER_TYPE_P (TREE_TYPE (decl)) && GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (TREE_TYPE (decl))) - && ptr_kind == GOMP_MAP_POINTER) + && ptr_kind == GOMP_MAP_POINTER + && op != EXEC_OMP_TARGET_EXIT_DATA + && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_RELEASE + && OMP_CLAUSE_MAP_KIND (node) != GOMP_MAP_DELETE) + { node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP); @@ -2455,13 +2476,13 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, && 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) + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE) map_kind = OMP_CLAUSE_MAP_KIND (node); + else if (op == EXEC_OMP_TARGET_EXIT_DATA + || OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE) + map_kind = GOMP_MAP_RELEASE; else - map_kind = GOMP_MAP_ALLOC; + map_kind = GOMP_MAP_TO; gcc_assert (se.string_length); node4 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (node4, map_kind); @@ -2476,7 +2497,18 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, 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) + if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_DELETE) + { + OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_DELETE); + node2 = desc_node; + } + else if (OMP_CLAUSE_MAP_KIND (node) == GOMP_MAP_RELEASE + || op == EXEC_OMP_TARGET_EXIT_DATA) + { + OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_RELEASE); + node2 = desc_node; + } + else if (ptr_kind == GOMP_MAP_ALWAYS_POINTER) { OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO); node2 = node; @@ -2487,11 +2519,11 @@ gfc_trans_omp_array_section (stmtblock_t *block, gfc_omp_namelist *n, OMP_CLAUSE_SET_MAP_KIND (desc_node, GOMP_MAP_TO_PSET); node2 = desc_node; } - node3 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); + if (op == EXEC_OMP_TARGET_EXIT_DATA) + return; + 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); + OMP_CLAUSE_DECL (node3) = gfc_conv_descriptor_data_get (decl); /* This purposely does not include GOMP_MAP_ALWAYS_POINTER. The extra cast prevents gimplify.cc from recognising it as being part of the struct - and adding an 'alloc: for the 'desc.data' pointer, which @@ -2595,7 +2627,7 @@ handle_iterator (gfc_namespace *ns, stmtblock_t *iter_block, tree block) static tree gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, locus where, bool declare_simd = false, - bool openacc = false) + bool openacc = false, gfc_exec_op op = EXEC_NOP) { tree omp_clauses = NULL_TREE, prev_clauses, chunk_size, c; tree iterator = NULL_TREE; @@ -3026,6 +3058,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tree node2 = NULL_TREE; tree node3 = NULL_TREE; tree node4 = NULL_TREE; + tree node5 = NULL_TREE; /* OpenMP: automatically map pointer targets with the pointer; hence, always update the descriptor/pointer itself. */ @@ -3130,6 +3163,24 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, || (n->expr->ref->type == REF_ARRAY && n->expr->ref->u.ar.type == AR_FULL)) { + gomp_map_kind map_kind; + tree type = TREE_TYPE (decl); + if (n->sym->ts.type == BT_CHARACTER + && n->sym->ts.deferred + && n->sym->attr.omp_declare_target + && (always_modifier || n->sym->attr.pointer) + && op != EXEC_OMP_TARGET_EXIT_DATA + && n->u.map_op != OMP_MAP_DELETE + && n->u.map_op != OMP_MAP_RELEASE) + { + gcc_assert (n->sym->ts.u.cl->backend_decl); + node5 = build_omp_clause (input_location, OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node5, GOMP_MAP_ALWAYS_TO); + OMP_CLAUSE_DECL (node5) = n->sym->ts.u.cl->backend_decl; + OMP_CLAUSE_SIZE (node5) + = TYPE_SIZE_UNIT (gfc_charlen_type_node); + } + tree present = gfc_omp_check_optional_argument (decl, true); if (openacc && n->sym->ts.type == BT_CLASS) { @@ -3145,13 +3196,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node2) = size_int (0); goto finalize_map_clause; } - else if (POINTER_TYPE_P (TREE_TYPE (decl)) + else if (POINTER_TYPE_P (type) && (gfc_omp_privatize_by_reference (decl) || GFC_DECL_GET_SCALAR_POINTER (decl) || GFC_DECL_GET_SCALAR_ALLOCATABLE (decl) || GFC_DECL_CRAY_POINTEE (decl) - || GFC_DESCRIPTOR_TYPE_P - (TREE_TYPE (TREE_TYPE (decl))) + || GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (type)) || (n->sym->ts.type == BT_DERIVED && (n->sym->ts.u.derived->ts.f90_type != BT_VOID)))) @@ -3164,7 +3214,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, to avoid accessing undefined variables. We cannot use a temporary variable here as otherwise the replacement of the variables in omp-low.cc will not work. */ - if (present && GFC_ARRAY_TYPE_P (TREE_TYPE (decl))) + if (present && GFC_ARRAY_TYPE_P (type)) { tree tmp = fold_build2_loc (input_location, MODIFY_EXPR, @@ -3181,22 +3231,51 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, cond, tmp, NULL_TREE)); } - node4 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node4, GOMP_MAP_POINTER); - OMP_CLAUSE_DECL (node4) = decl; - OMP_CLAUSE_SIZE (node4) = size_int (0); + /* For descriptor types, the unmapping happens below. */ + if (op != EXEC_OMP_TARGET_EXIT_DATA + || !GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (decl))) + { + enum gomp_map_kind gmk = GOMP_MAP_POINTER; + if (op == EXEC_OMP_TARGET_EXIT_DATA + && n->u.map_op == OMP_MAP_DELETE) + gmk = GOMP_MAP_DELETE; + else if (op == EXEC_OMP_TARGET_EXIT_DATA) + gmk = GOMP_MAP_RELEASE; + tree size; + if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE) + size = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + else + size = size_int (0); + node4 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + OMP_CLAUSE_SET_MAP_KIND (node4, gmk); + OMP_CLAUSE_DECL (node4) = decl; + OMP_CLAUSE_SIZE (node4) = size; + } decl = build_fold_indirect_ref (decl); if ((TREE_CODE (TREE_TYPE (orig_decl)) == REFERENCE_TYPE || gfc_omp_is_optional_argument (orig_decl)) && (GFC_DECL_GET_SCALAR_POINTER (orig_decl) || GFC_DECL_GET_SCALAR_ALLOCATABLE (orig_decl))) { + enum gomp_map_kind gmk; + if (op == EXEC_OMP_TARGET_EXIT_DATA + && n->u.map_op == OMP_MAP_DELETE) + gmk = GOMP_MAP_DELETE; + else if (op == EXEC_OMP_TARGET_EXIT_DATA) + gmk = GOMP_MAP_RELEASE; + else + gmk = GOMP_MAP_POINTER; + tree size; + if (gmk == GOMP_MAP_RELEASE || gmk == GOMP_MAP_DELETE) + size = TYPE_SIZE_UNIT (TREE_TYPE (decl)); + else + size = size_int (0); node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); + OMP_CLAUSE_SET_MAP_KIND (node3, gmk); OMP_CLAUSE_DECL (node3) = decl; - OMP_CLAUSE_SIZE (node3) = size_int (0); + OMP_CLAUSE_SIZE (node3) = size; decl = build_fold_indirect_ref (decl); } } @@ -3210,56 +3289,70 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gcc_assert (POINTER_TYPE_P (TREE_TYPE (ptr))); 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); + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_DECL (node2) = decl; OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); - node3 = build_omp_clause (input_location, - OMP_CLAUSE_MAP); - if (present) - { - 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); - ptr = build_fold_indirect_ref (ptr); - OMP_CLAUSE_DECL (node3) = ptr; - } + if (n->u.map_op == OMP_MAP_DELETE) + map_kind = GOMP_MAP_DELETE; + else if (op == EXEC_OMP_TARGET_EXIT_DATA + || n->u.map_op == OMP_MAP_RELEASE) + map_kind = GOMP_MAP_RELEASE; else - OMP_CLAUSE_DECL (node3) - = gfc_conv_descriptor_data_get (decl); - OMP_CLAUSE_SIZE (node3) = size_int (0); - if (n->u.map_op == OMP_MAP_ATTACH) - { - /* Standalone attach clauses used with arrays with - descriptors must copy the descriptor to the target, - else they won't have anything to perform the - attachment onto (see OpenACC 2.6, "2.6.3. Data - Structures with Pointers"). */ - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); - /* We don't want to map PTR at all in this case, so - delete its node and shuffle the others down. */ - node = node2; - node2 = node3; - node3 = NULL; - goto finalize_map_clause; - } - else if (n->u.map_op == OMP_MAP_DETACH) + map_kind = GOMP_MAP_TO_PSET; + OMP_CLAUSE_SET_MAP_KIND (node2, map_kind); + + if (op != EXEC_OMP_TARGET_EXIT_DATA + && n->u.map_op != OMP_MAP_DELETE + && n->u.map_op != OMP_MAP_RELEASE) { - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); - /* Similarly to above, we don't want to unmap PTR - here. */ - node = node2; - node2 = node3; - node3 = NULL; - goto finalize_map_clause; + node3 = build_omp_clause (input_location, + OMP_CLAUSE_MAP); + if (present) + { + 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); + ptr = build_fold_indirect_ref (ptr); + OMP_CLAUSE_DECL (node3) = ptr; + } + else + OMP_CLAUSE_DECL (node3) + = gfc_conv_descriptor_data_get (decl); + OMP_CLAUSE_SIZE (node3) = size_int (0); + + if (n->u.map_op == OMP_MAP_ATTACH) + { + /* Standalone attach clauses used with arrays with + descriptors must copy the descriptor to the + target, else they won't have anything to + perform the attachment onto (see OpenACC 2.6, + "2.6.3. Data Structures with Pointers"). */ + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); + /* We don't want to map PTR at all in this case, + so delete its node and shuffle the others + down. */ + node = node2; + node2 = node3; + node3 = NULL; + goto finalize_map_clause; + } + else if (n->u.map_op == OMP_MAP_DETACH) + { + OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); + /* Similarly to above, we don't want to unmap PTR + here. */ + node = node2; + node2 = node3; + node3 = NULL; + goto finalize_map_clause; + } + else + OMP_CLAUSE_SET_MAP_KIND (node3, + always_modifier + ? GOMP_MAP_ALWAYS_POINTER + : GOMP_MAP_POINTER); } - else - OMP_CLAUSE_SET_MAP_KIND (node3, - always_modifier - ? GOMP_MAP_ALWAYS_POINTER - : GOMP_MAP_POINTER); /* We have to check for n->sym->attr.dimension because of scalar coarrays. */ @@ -3275,6 +3368,23 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, tem = gfc_full_array_size (&cond_block, decl, GFC_TYPE_ARRAY_RANK (type)); + tree elemsz; + if (n->sym->ts.type == BT_CHARACTER + && n->sym->ts.deferred) + { + tree len = n->sym->ts.u.cl->backend_decl; + len = fold_convert (size_type_node, len); + elemsz = gfc_get_char_type (n->sym->ts.kind); + elemsz = TYPE_SIZE_UNIT (elemsz); + elemsz = fold_build2 (MULT_EXPR, size_type_node, + len, elemsz); + } + else + elemsz + = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + elemsz = fold_convert (gfc_array_index_type, elemsz); + tem = fold_build2 (MULT_EXPR, gfc_array_index_type, + tem, elemsz); gfc_add_modify (&cond_block, size, tem); then_b = gfc_finish_block (&cond_block); gfc_init_block (&cond_block); @@ -3305,6 +3415,12 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_init_block (&cond_block); tree size = gfc_full_array_size (&cond_block, decl, GFC_TYPE_ARRAY_RANK (type)); + tree elemsz + = TYPE_SIZE_UNIT (gfc_get_element_type (type)); + elemsz = fold_convert (gfc_array_index_type, elemsz); + size = fold_build2 (MULT_EXPR, gfc_array_index_type, + size, elemsz); + size = gfc_evaluate_now (size, &cond_block); if (present) { tree var = gfc_create_var (gfc_array_index_type, @@ -3323,15 +3439,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node) = size; } } - if (n->sym->attr.dimension) - { - tree elemsz - = TYPE_SIZE_UNIT (gfc_get_element_type (type)); - elemsz = fold_convert (gfc_array_index_type, elemsz); - OMP_CLAUSE_SIZE (node) - = fold_build2 (MULT_EXPR, gfc_array_index_type, - OMP_CLAUSE_SIZE (node), elemsz); - } } else if (present && TREE_CODE (decl) == INDIRECT_REF @@ -3347,6 +3454,37 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, } else OMP_CLAUSE_DECL (node) = decl; + + if (!n->sym->attr.dimension + && n->sym->ts.type == BT_CHARACTER + && n->sym->ts.deferred) + { + if (!DECL_P (decl)) + { + gcc_assert (TREE_CODE (decl) == INDIRECT_REF); + decl = TREE_OPERAND (decl, 0); + } + tree cond = fold_build2_loc (input_location, NE_EXPR, + boolean_type_node, + decl, null_pointer_node); + if (present) + cond = fold_build2_loc (input_location, + TRUTH_ANDIF_EXPR, + boolean_type_node, + present, cond); + tree len = n->sym->ts.u.cl->backend_decl; + len = fold_convert (size_type_node, len); + tree size = gfc_get_char_type (n->sym->ts.kind); + size = TYPE_SIZE_UNIT (size); + size = fold_build2 (MULT_EXPR, size_type_node, len, size); + size = build3_loc (input_location, + COND_EXPR, + size_type_node, + cond, size, + size_zero_node); + size = gfc_evaluate_now (size, block); + OMP_CLAUSE_SIZE (node) = size; + } } else if (n->expr && n->expr->expr_type == EXPR_VARIABLE @@ -3363,7 +3501,7 @@ 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, n, decl, element, k, + gfc_trans_omp_array_section (block, op, n, decl, element, k, node, node2, node3, node4); } else if (n->expr @@ -3411,9 +3549,15 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, fold_convert (size_type_node, se.string_length), TYPE_SIZE_UNIT (tmp)); + if (n->u.map_op == OMP_MAP_DELETE) + kind = GOMP_MAP_DELETE; + else if (op == EXEC_OMP_TARGET_EXIT_DATA) + kind = GOMP_MAP_RELEASE; + else + kind = GOMP_MAP_TO; node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_TO); + OMP_CLAUSE_SET_MAP_KIND (node3, kind); OMP_CLAUSE_DECL (node3) = se.string_length; OMP_CLAUSE_SIZE (node3) = TYPE_SIZE_UNIT (gfc_charlen_type_node); @@ -3517,11 +3661,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, = 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; + map_kind = OMP_CLAUSE_MAP_KIND (node); + if (GOMP_MAP_COPY_TO_P (map_kind) + || map_kind == GOMP_MAP_ALLOC) + map_kind = ((GOMP_MAP_ALWAYS_P (map_kind) + || gfc_expr_attr (n->expr).pointer) + ? GOMP_MAP_ALWAYS_TO : 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 if (op == EXEC_OMP_TARGET_EXIT_DATA) + map_kind = GOMP_MAP_RELEASE; else map_kind = GOMP_MAP_ALLOC; if (!openacc @@ -3562,6 +3712,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, node2 = node; node = desc_node; /* Put first. */ } + if (op == EXEC_OMP_TARGET_EXIT_DATA) + goto finalize_map_clause; node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); OMP_CLAUSE_SET_MAP_KIND (node3, @@ -3592,7 +3744,7 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, bool element = lastref->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, + gfc_trans_omp_array_section (block, op, n, inner, element, kind, node, node2, node3, node4); } @@ -3611,6 +3763,8 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, omp_clauses = gfc_trans_add_clause (node3, omp_clauses); if (node4) omp_clauses = gfc_trans_add_clause (node4, omp_clauses); + if (node5) + omp_clauses = gfc_trans_add_clause (node5, omp_clauses); } break; case OMP_LIST_TO: @@ -7475,7 +7629,7 @@ gfc_trans_omp_target_exit_data (gfc_code *code) gfc_start_block (&block); omp_clauses = gfc_trans_omp_clauses (&block, code->ext.omp_clauses, - code->loc); + code->loc, false, false, code->op); stmt = build1_loc (input_location, OMP_TARGET_EXIT_DATA, void_type_node, omp_clauses); gfc_add_expr_to_block (&block, stmt); diff --git a/gcc/gimplify.cc b/gcc/gimplify.cc index ade6e335da7..94d5dd3515e 100644 --- a/gcc/gimplify.cc +++ b/gcc/gimplify.cc @@ -10722,7 +10722,7 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, { struct gimplify_omp_ctx *ctx, *outer_ctx; tree c; - tree *prev_list_p = NULL, *orig_list_p = list_p; + tree *orig_list_p = list_p; int handled_depend_iterators = -1; int nowait = -1; @@ -11149,31 +11149,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, default: break; } - /* For Fortran, not only the pointer to the data is mapped but also - the address of the pointer, the array descriptor etc.; for - 'exit data' - and in particular for 'delete:' - having an 'alloc:' - does not make sense. Likewise, for 'update' only transferring the - data itself is needed as the rest has been handled in previous - directives. However, for 'exit data', the array descriptor needs - to be delete; hence, we turn the MAP_TO_PSET into a MAP_DELETE. - - NOTE: Generally, it is not safe to perform "enter data" operations - on arrays where the data *or the descriptor* may go out of scope - before a corresponding "exit data" operation -- and such a - descriptor may be synthesized temporarily, e.g. to pass an - explicit-shape array to a function expecting an assumed-shape - argument. Performing "enter data" inside the called function - would thus be problematic. */ - if (code == OMP_TARGET_EXIT_DATA - && OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET) - OMP_CLAUSE_SET_MAP_KIND (c, OMP_CLAUSE_MAP_KIND (*prev_list_p) - == GOMP_MAP_DELETE - ? GOMP_MAP_DELETE : GOMP_MAP_RELEASE); - else if ((code == OMP_TARGET_EXIT_DATA || code == OMP_TARGET_UPDATE) - && (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET)) - remove = true; - if (remove) break; if (DECL_P (decl) && outer_ctx && (region_type & ORT_ACC)) @@ -11433,21 +11408,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, remove = true; break; } - - if (!remove - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ALWAYS_POINTER - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_ATTACH_DETACH - && OMP_CLAUSE_MAP_KIND (c) != GOMP_MAP_TO_PSET - && OMP_CLAUSE_CHAIN (c) - && OMP_CLAUSE_CODE (OMP_CLAUSE_CHAIN (c)) == OMP_CLAUSE_MAP - && ((OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ALWAYS_POINTER) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_ATTACH_DETACH) - || (OMP_CLAUSE_MAP_KIND (OMP_CLAUSE_CHAIN (c)) - == GOMP_MAP_TO_PSET))) - prev_list_p = list_p; - break; } flags = GOVD_MAP | GOVD_EXPLICIT; diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 1e5bf0ba1e6..23f0ffc627e 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -20,8 +20,8 @@ ! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_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:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) 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: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(release:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(.*int.*\\) parm\\.0\\.data - \\(.*int.*\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(delete:MEM <\[^>\]+> \\\[\\(integer\\(kind=.\\)\\\[0:\\\] \\*\\)_\[0-9\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(delete:del_f_p \\\[len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-10.f90 b/gcc/testsuite/gfortran.dg/gomp/map-10.f90 new file mode 100644 index 00000000000..c12bf25ad19 --- /dev/null +++ b/gcc/testsuite/gfortran.dg/gomp/map-10.f90 @@ -0,0 +1,69 @@ +! { dg-additional-options "-fdump-tree-omplower" } + +! If enter data adds a (GOMP_MAP_)POINTER attachment, exit data needs to remove +! it again. If not there can be all kind of issues, in particular when +! stack memory was mapped, reused later and mapped again. + +subroutine test_aa (aa2, aa3) + integer(kind=4), allocatable :: aa1, aa2, aa3 + optional :: aa3 + !$omp target enter data map(aa1) + !$omp target exit data map(aa1) + !$omp target enter data map(aa2) + !$omp target exit data map(aa2) + !$omp target enter data map(aa3) + !$omp target exit data map(aa3) +end + +subroutine test_pp (pp2, pp3) + integer(kind=4), allocatable :: pp1, pp2, pp3 + optional :: pp3 + !$omp target enter data map(pp1) + !$omp target exit data map(pp1) + !$omp target enter data map(pp2) + !$omp target exit data map(pp2) + !$omp target enter data map(pp3) + !$omp target exit data map(pp3) +end + +subroutine test_pprelease (rp2, rp3) + integer(kind=4), allocatable :: rp1, rp2, rp3 + optional :: rp3 + !$omp target enter data map(rp1) + !$omp target exit data map(release:rp1) + !$omp target enter data map(rp2) + !$omp target exit data map(release:rp2) + !$omp target enter data map(rp3) + !$omp target exit data map(release:rp3) +end + +subroutine test_ppdelete (dp2, dp3) + integer(kind=4), allocatable :: dp1, dp2, dp3 + optional :: dp3 + !$omp target enter data map(dp1) + !$omp target exit data map(delete:dp1) + !$omp target enter data map(dp2) + !$omp target exit data map(delete:dp2) + !$omp target enter data map(dp3) + !$omp target exit data map(delete:dp3) +end + + +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:aa1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*aa1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa1 \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa2 \\\[len: .\\\]\\) map\\(release:\\*aa2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:aa3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:aa3 \\\[len: .\\\]\\) map\\(release:\\*aa3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:pp1 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*pp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp1 \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp2 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*_\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp2 \\\[len: .\\\]\\) map\\(release:\\*pp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(to:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(alloc:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:pp3 \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(from:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:pp3 \\\[len: .\\\]\\) map\\(release:\\*pp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp1 \\\[len: .\\\]\\) map\\(release:\\*rp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:rp2 \\\[len: .\\\]\\) map\\(release:\\*rp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(release:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(release:rp3 \\\[len: .\\\]\\) map\\(release:\\*rp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp1 \\\[len: .\\\]\\) map\\(delete:\\*dp1.\[0-9\]+_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:dp2 \\\[len: .\\\]\\) map\\(delete:\\*dp2.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\) map\\(delete:\\*_\[0-9\]+ \\\[len: 4\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:\\*D.\[0-9\]+ \\\[len: 4\\\]\\) map\\(delete:dp3 \\\[len: .\\\]\\) map\\(delete:\\*dp3.\[0-9\]+_\[0-9\]+ \\\[len: .\\\]\\)" "omplower" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 b/gcc/testsuite/gfortran.dg/gomp/map-9.f90 index 9e7b811c8af..b770b931bee 100644 --- a/gcc/testsuite/gfortran.dg/gomp/map-9.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/map-9.f90 @@ -2,7 +2,7 @@ ! PR fortran/108545 -! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target enter data map\\(struct:x \\\[len: 1\\\]\\) map\\(always,to:x.a \\\[len: \[0-9\]+\\\]\\) map\\(to:MEM \\\[\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\)_\[0-9\]+] \\\[len: _\[0-9\]+\\\]\\) map\\(always_pointer:x.a.data \\\[pointer assign, bias: 0\\\]\\)" "omplower" } } program p type t diff --git a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 index f5d888592b9..cd771b33a25 100644 --- a/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/pr78260-2.f90 @@ -48,9 +48,11 @@ contains end subroutine sub end module m -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } +! Check for multiplication: len = arrays_size * 4: +! { dg-final { scan-tree-dump-times "D\\.\[0-9\]+ = D\\.\[0-9\]+ \\* 4;" 2 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:arr \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) arr.data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } } -! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } +! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+\\\]\\) map\\(to:\\*__result \\\[pointer set, len: ..\\\]\\) map\\(alloc:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[pointer assign, bias: 0\\\]\\) map\\(alloc:__result \\\[pointer assign, bias: 0\\\]\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) __result->data \\\[len: D.\[0-9\]+ \\* 4\\\]\\)" 1 "original" } } ! { dg-final { scan-tree-dump-times "#pragma omp target data map\\(tofrom:\\*__result.0\\) map\\(alloc:__result.0 \\\[pointer assign, bias: 0\\\]\\)" 2 "original" } } ! { dg-final { scan-tree-dump-times "#pragma omp target update to\\(\\*__result.0\\)" 2 "original" } } diff --git a/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90 b/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90 index ed57d0072d7..219dc467c46 100644 --- a/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90 +++ b/gcc/testsuite/gfortran.dg/gomp/target-exit-data.f90 @@ -15,6 +15,6 @@ integer, allocatable :: one(:), two(:), three(:) !$omp target exit data map(from:three) end -! { dg-final { scan-tree-dump "omp target exit data map\\(delete:.*\\) map\\(delete:one \\\[len: .*\\\]\\)" "omplower" } } -! { dg-final { scan-tree-dump "omp target exit data map\\(release:.*\\) map\\(release:two \\\[len: .*\\\]\\)" "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(delete:one \\\[len: \[0-9\]+\\\]\\) map\\(delete:MEM " "omplower" } } +! { dg-final { scan-tree-dump "#pragma omp target exit data map\\(release:two \\\[len: \[0-9\]+\\\]\\) map\\(release:MEM " "omplower" } } ! { dg-final { scan-tree-dump "omp target exit data map\\(from:.*\\) map\\(release:three \\\[len: .*\\\]\\)" "omplower" } } diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 index 5d97566c66c..1b3cdf9e76a 100644 --- a/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-3.f90 @@ -17,6 +17,6 @@ var%p2 = [46,679,54] if (any (var%p1 /= [22,53,28,6,4])) stop 3 if (any (var%p2 /= [46,679,54])) stop 4 !$omp end target -!!$omp target exit data map(from:var%p1, var%p2) +!$omp target exit data map(from:var%p1, var%p2) end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 new file mode 100644 index 00000000000..6192bf29d9a --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-4.f90 @@ -0,0 +1,540 @@ +! Check that 'map(alloc:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5), ic2 + character(len=11) :: ccstr(3:4), ccstr2 + character(len=11,kind=4) :: cc4str(3:7), cc4str2 + integer, pointer :: pc(:), pc2 + character(len=:), pointer :: pcstr(:), pcstr2 + character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +end type t + +type(t) :: dt + +integer :: ii(5), ii2 +character(len=11) :: clstr(-1:1), clstr2 +character(len=11,kind=4) :: cl4str(0:3), cl4str2 +integer, pointer :: ip(:), ip2 +integer, allocatable :: ia(:), ia2 +character(len=:), pointer :: pstr(:), pstr2 +character(len=:), allocatable :: astr(:), astr2 +character(len=:,kind=4), pointer :: p4str(:), p4str2 +character(len=:,kind=4), allocatable :: a4str(:), a4str2 + +allocate(dt%pc(5), dt%pc2) +allocate(character(len=2) :: dt%pcstr(2)) +allocate(character(len=4) :: dt%pcstr2) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +allocate(character(len=5,kind=4) :: dt%pc4str2) + +allocate(ip(5), ip2, ia(8), ia2) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=4) :: pstr2) +allocate(character(len=6) :: astr(3:5)) +allocate(character(len=8) :: astr2) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=5,kind=4) :: p4str2) +allocate(character(len=7,kind=4) :: a4str(-2:3)) +allocate(character(len=9,kind=4) :: a4str2) + + +! integer :: ic(2:5), ic2 + +!$omp target enter data map(alloc: dt%ic) +!$omp target map(alloc: dt%ic) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%ic2) +!$omp target map(alloc: dt%ic2) + dt%ic2 = 42 +!$omp end target +!$omp target exit data map(from: dt%ic2) +if (dt%ic2 /= 42) error stop + + +! character(len=11) :: ccstr(3:4), ccstr2 + +!$omp target enter data map(alloc: dt%ccstr) +!$omp target map(alloc: dt%ccstr) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr = ["12345678901", "abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + +!$omp target enter data map(alloc: dt%ccstr2) +!$omp target map(alloc: dt%ccstr2) + if (len(dt%ccstr2) /= 11) error stop + dt%ccstr2 = "ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%ccstr2) +if (len(dt%ccstr2) /= 11) error stop +if (dt%ccstr2 /= "ABCDEFGHIJK") error stop + + +! character(len=11,kind=4) :: cc4str(3:7), cc4str2 + +! Value check fails +!$omp target enter data map(alloc: dt%cc4str) +!$omp target map(alloc: dt%cc4str) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +!$omp target enter data map(alloc: dt%cc4str2) +!$omp target map(alloc: dt%cc4str2) + if (len(dt%cc4str2) /= 11) error stop + dt%cc4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%cc4str2) +if (len(dt%cc4str2) /= 11) error stop +if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop + + +! integer, pointer :: pc(:), pc2 +! allocate(dt%pc(5), dt%pc2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pc) +!$omp target map(alloc: dt%pc) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: dt%pc2) +!$omp target map(alloc: dt%pc2) + if (.not. associated(dt%pc2)) error stop + dt%pc2 = 99 +!$omp end target +!$omp target exit data map(from: dt%pc2) +if (dt%pc2 /= 99) error stop +if (.not. associated(dt%pc2)) error stop + + +! character(len=:), pointer :: pcstr(:), pcstr2 +! allocate(character(len=2) :: dt%pcstr(2)) +! allocate(character(len=4) :: dt%pcstr2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pcstr) +!$omp target map(alloc: dt%pcstr) + if (.not. associated(dt%pcstr)) error stop + if (len(dt%pcstr) /= 2) error stop + if (size(dt%pcstr) /= 2) error stop + if (lbound(dt%pcstr, 1) /= 1) error stop + if (ubound(dt%pcstr, 1) /= 2) error stop + dt%pcstr = ["01", "jk"] +!$omp end target +!$omp target exit data map(from: dt%pcstr) +if (.not. associated(dt%pcstr)) error stop +if (len(dt%pcstr) /= 2) error stop +if (size(dt%pcstr) /= 2) error stop +if (lbound(dt%pcstr, 1) /= 1) error stop +if (ubound(dt%pcstr, 1) /= 2) error stop +if (any (dt%pcstr /= ["01", "jk"])) error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(alloc: dt%pcstr2) +!$omp target map(alloc: dt%pcstr2) + if (.not. associated(dt%pcstr2)) error stop + if (len(dt%pcstr2) /= 4) error stop + dt%pcstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: dt%pcstr2) +if (.not. associated(dt%pcstr2)) error stop +if (len(dt%pcstr2) /= 4) error stop +if (dt%pcstr2 /= "HIJK") error stop + + +! character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +! allocate(character(len=5,kind=4) :: dt%pc4str2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 +! structure element when other mapped elements from the same structure weren't mapped together with it +!$omp target enter data map(alloc: dt%pc4str) +!$omp target map(alloc: dt%pc4str) + if (.not. associated(dt%pc4str)) error stop + if (len(dt%pc4str) /= 3) error stop + if (size(dt%pc4str) /= 2) error stop + if (lbound(dt%pc4str, 1) /= 2) error stop + if (ubound(dt%pc4str, 1) /= 3) error stop + dt%pc4str = [4_"456", 4_"tzu"] +!$omp end target +!$omp target exit data map(from: dt%pc4str) +if (.not. associated(dt%pc4str)) error stop +if (len(dt%pc4str) /= 3) error stop +if (size(dt%pc4str) /= 2) error stop +if (lbound(dt%pc4str, 1) /= 2) error stop +if (ubound(dt%pc4str, 1) /= 3) error stop +if (dt%pc4str(2) /= 4_"456") error stop +if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(alloc: dt%pc4str2) +!$omp target map(alloc: dt%pc4str2) + if (.not. associated(dt%pc4str2)) error stop + if (len(dt%pc4str2) /= 5) error stop + dt%pc4str2 = 4_"98765" +!$omp end target +!$omp target exit data map(from: dt%pc4str2) +if (.not. associated(dt%pc4str2)) error stop +if (len(dt%pc4str2) /= 5) error stop +if (dt%pc4str2 /= 4_"98765") error stop + + +! integer :: ii(5), ii2 + +!$omp target enter data map(alloc: ii) +!$omp target map(alloc: ii) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii = [-1, -2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + +!$omp target enter data map(alloc: ii2) +!$omp target map(alloc: ii2) + ii2 = -410 +!$omp end target +!$omp target exit data map(from: ii2) +if (ii2 /= -410) error stop + + +! character(len=11) :: clstr(-1:1), clstr2 + +!$omp target enter data map(alloc: clstr) +!$omp target map(alloc: clstr) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +!$omp target enter data map(alloc: clstr2) +!$omp target map(alloc: clstr2) + if (len(clstr2) /= 11) error stop + clstr2 = "ABCDEFghijk" +!$omp end target +!$omp target exit data map(from: clstr2) +if (len(clstr2) /= 11) error stop +if (clstr2 /= "ABCDEFghijk") error stop + + +! character(len=11,kind=4) :: cl4str(0:3), cl4str2 + +!$omp target enter data map(alloc: cl4str) +!$omp target map(alloc: cl4str) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + +!$omp target enter data map(alloc: cl4str2) +!$omp target map(alloc: cl4str2) + if (len(cl4str2) /= 11) error stop + cl4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: cl4str2) +if (len(cl4str2) /= 11) error stop +if (cl4str2 /= 4_"ABCDEFGHIJK") error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ip) +!$omp target map(alloc: ip) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(alloc: ip2) +!$omp target map(alloc: ip2) + if (.not. associated(ip2)) error stop + ip2 = 99 +!$omp end target +!$omp target exit data map(from: ip2) +if (ip2 /= 99) error stop +if (.not. associated(ip2)) error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(alloc: ia) +!$omp target map(alloc: ia) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia = [1,2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + +!$omp target enter data map(alloc: ia2) +!$omp target map(alloc: ia2) + if (.not. allocated(ia2)) error stop + ia2 = 102 +!$omp end target +!$omp target exit data map(from: ia2) +if (ia2 /= 102) error stop +if (.not. allocated(ia2)) error stop + + +! character(len=:), pointer :: pstr(:), pstr2 +! allocate(character(len=2) :: pstr(-2:0)) +! allocate(character(len=4) :: pstr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: pstr) +!$omp target map(alloc: pstr) + if (.not. associated(pstr)) error stop + if (len(pstr) /= 2) error stop + if (size(pstr) /= 3) error stop + if (lbound(pstr, 1) /= -2) error stop + if (ubound(pstr, 1) /= 0) error stop + pstr = ["01", "jk", "aq"] +!$omp end target +!$omp target exit data map(from: pstr) +if (.not. associated(pstr)) error stop +if (len(pstr) /= 2) error stop +if (size(pstr) /= 3) error stop +if (lbound(pstr, 1) /= -2) error stop +if (ubound(pstr, 1) /= 0) error stop +if (any (pstr /= ["01", "jk", "aq"])) error stop + +!$omp target enter data map(alloc: pstr2) +!$omp target map(alloc: pstr2) + if (.not. associated(pstr2)) error stop + if (len(pstr2) /= 4) error stop + pstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: pstr2) +if (.not. associated(pstr2)) error stop +if (len(pstr2) /= 4) error stop +if (pstr2 /= "HIJK") error stop + + +! character(len=:), allocatable :: astr(:), astr2 +! allocate(character(len=6) :: astr(3:5)) +! allocate(character(len=8) :: astr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: astr) +!$omp target map(alloc: astr) + if (.not. allocated(astr)) error stop + if (len(astr) /= 6) error stop + if (size(astr) /= 3) error stop + if (lbound(astr, 1) /= 3) error stop + if (ubound(astr, 1) /= 5) error stop + astr = ["01db45", "jk$D%S", "zutg47"] +!$omp end target +!$omp target exit data map(from: astr) +if (.not. allocated(astr)) error stop +if (len(astr) /= 6) error stop +if (size(astr) /= 3) error stop +if (lbound(astr, 1) /= 3) error stop +if (ubound(astr, 1) /= 5) error stop +if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(alloc: astr2) +!$omp target map(alloc: astr2) + if (.not. allocated(astr2)) error stop + if (len(astr2) /= 8) error stop + astr2 = "HIJKhijk" +!$omp end target +!$omp target exit data map(from: astr2) +if (.not. allocated(astr2)) error stop +if (len(astr2) /= 8) error stop +if (astr2 /= "HIJKhijk") error stop + + +! character(len=:,kind=4), pointer :: p4str(:), p4str2 +! allocate(character(len=3,kind=4) :: p4str(2:4)) +! allocate(character(len=5,kind=4) :: p4str2) + +! FAILS with value check + +!$omp target enter data map(alloc: p4str) +!$omp target map(alloc: p4str) + if (.not. associated(p4str)) error stop + if (len(p4str) /= 3) error stop + if (size(p4str) /= 3) error stop + if (lbound(p4str, 1) /= 2) error stop + if (ubound(p4str, 1) /= 4) error stop + p4str(:) = [4_"f85", 4_"8af", 4_"A%F"] +!$omp end target +!$omp target exit data map(from: p4str) +if (.not. associated(p4str)) error stop +if (len(p4str) /= 3) error stop +if (size(p4str) /= 3) error stop +if (lbound(p4str, 1) /= 2) error stop +if (ubound(p4str, 1) /= 4) error stop +if (p4str(2) /= 4_"f85") error stop +if (p4str(3) /= 4_"8af") error stop +if (p4str(4) /= 4_"A%F") error stop + +!$omp target enter data map(alloc: p4str2) +!$omp target map(alloc: p4str2) + if (.not. associated(p4str2)) error stop + if (len(p4str2) /= 5) error stop + p4str2 = 4_"9875a" +!$omp end target +!$omp target exit data map(from: p4str2) +if (.not. associated(p4str2)) error stop +if (len(p4str2) /= 5) error stop +if (p4str2 /= 4_"9875a") error stop + + +! character(len=:,kind=4), allocatable :: a4str(:), a4str2 +! allocate(character(len=7,kind=4) :: a4str(-2:3)) +! allocate(character(len=9,kind=4) :: a4str2) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +!$omp target enter data map(alloc: a4str) +!$omp target map(alloc: a4str) + if (.not. allocated(a4str)) error stop + if (len(a4str) /= 7) error stop + if (size(a4str) /= 6) error stop + if (lbound(a4str, 1) /= -2) error stop + if (ubound(a4str, 1) /= 3) error stop + ! See PR fortran/107508 why '(:)' is required + a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!$omp end target +!$omp target exit data map(from: a4str) +if (.not. allocated(a4str)) error stop +if (len(a4str) /= 7) error stop +if (size(a4str) /= 6) error stop +if (lbound(a4str, 1) /= -2) error stop +if (ubound(a4str, 1) /= 3) error stop +if (a4str(-2) /= 4_"sf456aq") error stop +if (a4str(-1) /= 4_"3dtzu24") error stop +if (a4str(0) /= 4_"_4fh7sm") error stop +if (a4str(1) /= 4_"=ff85s7") error stop +if (a4str(2) /= 4_"j=8af4d") error stop +if (a4str(3) /= 4_".,A%Fsz") error stop + +!$omp target enter data map(alloc: a4str2) +!$omp target map(alloc: a4str2) + if (.not. allocated(a4str2)) error stop + if (len(a4str2) /= 9) error stop + a4str2 = 4_"98765a23d" +!$omp end target +!$omp target exit data map(from: a4str2) +if (.not. allocated(a4str2)) error stop +if (len(a4str2) /= 9) error stop +if (a4str2 /= 4_"98765a23d") error stop + + +deallocate(dt%pc, dt%pc2) +deallocate(dt%pcstr) +deallocate(dt%pcstr2) + +deallocate(dt%pc4str) +deallocate(dt%pc4str2) + +deallocate(ip, ip2, ia, ia2) +deallocate(pstr) +deallocate(pstr2) +deallocate(astr) +deallocate(astr2) + +deallocate(p4str) +deallocate(p4str2) +deallocate(a4str) +deallocate(a4str2) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 new file mode 100644 index 00000000000..cf759346a84 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-5.f90 @@ -0,0 +1,540 @@ +! Check that 'map((to)from:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5), ic2 + character(len=11) :: ccstr(3:4), ccstr2 + character(len=11,kind=4) :: cc4str(3:7), cc4str2 + integer, pointer :: pc(:), pc2 + character(len=:), pointer :: pcstr(:), pcstr2 + character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +end type t + +type(t) :: dt + +integer :: ii(5), ii2 +character(len=11) :: clstr(-1:1), clstr2 +character(len=11,kind=4) :: cl4str(0:3), cl4str2 +integer, pointer :: ip(:), ip2 +integer, allocatable :: ia(:), ia2 +character(len=:), pointer :: pstr(:), pstr2 +character(len=:), allocatable :: astr(:), astr2 +character(len=:,kind=4), pointer :: p4str(:), p4str2 +character(len=:,kind=4), allocatable :: a4str(:), a4str2 + +allocate(dt%pc(5), dt%pc2) +allocate(character(len=2) :: dt%pcstr(2)) +allocate(character(len=4) :: dt%pcstr2) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +allocate(character(len=5,kind=4) :: dt%pc4str2) + +allocate(ip(5), ip2, ia(8), ia2) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=4) :: pstr2) +allocate(character(len=6) :: astr(3:5)) +allocate(character(len=8) :: astr2) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=5,kind=4) :: p4str2) +allocate(character(len=7,kind=4) :: a4str(-2:3)) +allocate(character(len=9,kind=4) :: a4str2) + + +! integer :: ic(2:5), ic2 + +!$omp target enter data map(tofrom: dt%ic) +!$omp target map(from: dt%ic) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: dt%ic2) +!$omp target map(from: dt%ic2) + dt%ic2 = 42 +!$omp end target +!$omp target exit data map(from: dt%ic2) +if (dt%ic2 /= 42) error stop + + +! character(len=11) :: ccstr(3:4), ccstr2 + +!$omp target enter data map(tofrom: dt%ccstr) +!$omp target map(from: dt%ccstr) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr = ["12345678901", "abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + +!$omp target enter data map(tofrom: dt%ccstr2) +!$omp target map(from: dt%ccstr2) + if (len(dt%ccstr2) /= 11) error stop + dt%ccstr2 = "ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%ccstr2) +if (len(dt%ccstr2) /= 11) error stop +if (dt%ccstr2 /= "ABCDEFGHIJK") error stop + + +! character(len=11,kind=4) :: cc4str(3:7), cc4str2 + +! Value check fails +!$omp target enter data map(tofrom: dt%cc4str) +!$omp target map(from: dt%cc4str) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +!$omp target enter data map(tofrom: dt%cc4str2) +!$omp target map(from: dt%cc4str2) + if (len(dt%cc4str2) /= 11) error stop + dt%cc4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: dt%cc4str2) +if (len(dt%cc4str2) /= 11) error stop +if (dt%cc4str2 /= 4_"ABCDEFGHIJK") error stop + + +! integer, pointer :: pc(:), pc2 +! allocate(dt%pc(5), dt%pc2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pc) +!$omp target map(from: dt%pc) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: dt%pc2) +!$omp target map(from: dt%pc2) + if (.not. associated(dt%pc2)) error stop + dt%pc2 = 99 +!$omp end target +!$omp target exit data map(from: dt%pc2) +if (dt%pc2 /= 99) error stop +if (.not. associated(dt%pc2)) error stop + + +! character(len=:), pointer :: pcstr(:), pcstr2 +! allocate(character(len=2) :: dt%pcstr(2)) +! allocate(character(len=4) :: dt%pcstr2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pcstr) +!$omp target map(from: dt%pcstr) + if (.not. associated(dt%pcstr)) error stop + if (len(dt%pcstr) /= 2) error stop + if (size(dt%pcstr) /= 2) error stop + if (lbound(dt%pcstr, 1) /= 1) error stop + if (ubound(dt%pcstr, 1) /= 2) error stop + dt%pcstr = ["01", "jk"] +!$omp end target +!$omp target exit data map(from: dt%pcstr) +if (.not. associated(dt%pcstr)) error stop +if (len(dt%pcstr) /= 2) error stop +if (size(dt%pcstr) /= 2) error stop +if (lbound(dt%pcstr, 1) /= 1) error stop +if (ubound(dt%pcstr, 1) /= 2) error stop +if (any (dt%pcstr /= ["01", "jk"])) error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(tofrom: dt%pcstr2) +!$omp target map(from: dt%pcstr2) + if (.not. associated(dt%pcstr2)) error stop + if (len(dt%pcstr2) /= 4) error stop + dt%pcstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: dt%pcstr2) +if (.not. associated(dt%pcstr2)) error stop +if (len(dt%pcstr2) /= 4) error stop +if (dt%pcstr2 /= "HIJK") error stop + + +! character(len=:,kind=4), pointer :: pc4str(:), pc4str2 +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) +! allocate(character(len=5,kind=4) :: dt%pc4str2) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(tofrom: dt%pc4str) +!$omp target map(from: dt%pc4str) + if (.not. associated(dt%pc4str)) error stop + if (len(dt%pc4str) /= 3) error stop + if (size(dt%pc4str) /= 2) error stop + if (lbound(dt%pc4str, 1) /= 2) error stop + if (ubound(dt%pc4str, 1) /= 3) error stop + dt%pc4str = [4_"456", 4_"tzu"] +!$omp end target +!$omp target exit data map(from: dt%pc4str) +if (.not. associated(dt%pc4str)) error stop +if (len(dt%pc4str) /= 3) error stop +if (size(dt%pc4str) /= 2) error stop +if (lbound(dt%pc4str, 1) /= 2) error stop +if (ubound(dt%pc4str, 1) /= 3) error stop +if (dt%pc4str(2) /= 4_"456") error stop +if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +!$omp target enter data map(tofrom: dt%pc4str2) +!$omp target map(from: dt%pc4str2) + if (.not. associated(dt%pc4str2)) error stop + if (len(dt%pc4str2) /= 5) error stop + dt%pc4str2 = 4_"98765" +!$omp end target +!$omp target exit data map(from: dt%pc4str2) +if (.not. associated(dt%pc4str2)) error stop +if (len(dt%pc4str2) /= 5) error stop +if (dt%pc4str2 /= 4_"98765") error stop + + +! integer :: ii(5), ii2 + +!$omp target enter data map(tofrom: ii) +!$omp target map(from: ii) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii = [-1, -2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + +!$omp target enter data map(tofrom: ii2) +!$omp target map(from: ii2) + ii2 = -410 +!$omp end target +!$omp target exit data map(from: ii2) +if (ii2 /= -410) error stop + + +! character(len=11) :: clstr(-1:1), clstr2 + +!$omp target enter data map(tofrom: clstr) +!$omp target map(from: clstr) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr = ["12345678901", "abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +!$omp target enter data map(tofrom: clstr2) +!$omp target map(from: clstr2) + if (len(clstr2) /= 11) error stop + clstr2 = "ABCDEFghijk" +!$omp end target +!$omp target exit data map(from: clstr2) +if (len(clstr2) /= 11) error stop +if (clstr2 /= "ABCDEFghijk") error stop + + +! character(len=11,kind=4) :: cl4str(0:3), cl4str2 + +!$omp target enter data map(tofrom: cl4str) +!$omp target map(from: cl4str) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str = [4_"12345678901", 4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + +!$omp target enter data map(tofrom: cl4str2) +!$omp target map(from: cl4str2) + if (len(cl4str2) /= 11) error stop + cl4str2 = 4_"ABCDEFGHIJK" +!$omp end target +!$omp target exit data map(from: cl4str2) +if (len(cl4str2) /= 11) error stop +if (cl4str2 /= 4_"ABCDEFGHIJK") error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(tofrom: ip) +!$omp target map(from: ip) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip = [11, 22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +!$omp target enter data map(tofrom: ip2) +!$omp target map(from: ip2) + if (.not. associated(ip2)) error stop + ip2 = 99 +!$omp end target +!$omp target exit data map(from: ip2) +if (ip2 /= 99) error stop +if (.not. associated(ip2)) error stop + + +! allocate(ip(5), ip2, ia(8), ia2) + +!$omp target enter data map(tofrom: ia) +!$omp target map(from: ia) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia = [1,2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + +!$omp target enter data map(tofrom: ia2) +!$omp target map(from: ia2) + if (.not. allocated(ia2)) error stop + ia2 = 102 +!$omp end target +!$omp target exit data map(from: ia2) +if (ia2 /= 102) error stop +if (.not. allocated(ia2)) error stop + + +! character(len=:), pointer :: pstr(:), pstr2 +! allocate(character(len=2) :: pstr(-2:0)) +! allocate(character(len=4) :: pstr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: pstr) +!$omp target map(from: pstr) + if (.not. associated(pstr)) error stop + if (len(pstr) /= 2) error stop + if (size(pstr) /= 3) error stop + if (lbound(pstr, 1) /= -2) error stop + if (ubound(pstr, 1) /= 0) error stop + pstr = ["01", "jk", "aq"] +!$omp end target +!$omp target exit data map(from: pstr) +if (.not. associated(pstr)) error stop +if (len(pstr) /= 2) error stop +if (size(pstr) /= 3) error stop +if (lbound(pstr, 1) /= -2) error stop +if (ubound(pstr, 1) /= 0) error stop +if (any (pstr /= ["01", "jk", "aq"])) error stop + +!$omp target enter data map(tofrom: pstr2) +!$omp target map(from: pstr2) + if (.not. associated(pstr2)) error stop + if (len(pstr2) /= 4) error stop + pstr2 = "HIJK" +!$omp end target +!$omp target exit data map(from: pstr2) +if (.not. associated(pstr2)) error stop +if (len(pstr2) /= 4) error stop +if (pstr2 /= "HIJK") error stop + + +! character(len=:), allocatable :: astr(:), astr2 +! allocate(character(len=6) :: astr(3:5)) +! allocate(character(len=8) :: astr2) + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: astr) +!$omp target map(from: astr) + if (.not. allocated(astr)) error stop + if (len(astr) /= 6) error stop + if (size(astr) /= 3) error stop + if (lbound(astr, 1) /= 3) error stop + if (ubound(astr, 1) /= 5) error stop + astr = ["01db45", "jk$D%S", "zutg47"] +!$omp end target +!$omp target exit data map(from: astr) +if (.not. allocated(astr)) error stop +if (len(astr) /= 6) error stop +if (size(astr) /= 3) error stop +if (lbound(astr, 1) /= 3) error stop +if (ubound(astr, 1) /= 5) error stop +if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop + +! libgomp: nvptx_alloc error: out of memory + +!$omp target enter data map(tofrom: astr2) +!$omp target map(from: astr2) + if (.not. allocated(astr2)) error stop + if (len(astr2) /= 8) error stop + astr2 = "HIJKhijk" +!$omp end target +!$omp target exit data map(from: astr2) +if (.not. allocated(astr2)) error stop +if (len(astr2) /= 8) error stop +if (astr2 /= "HIJKhijk") error stop + + +! character(len=:,kind=4), pointer :: p4str(:), p4str2 +! allocate(character(len=3,kind=4) :: p4str(2:4)) +! allocate(character(len=5,kind=4) :: p4str2) + +! FAILS with value check + +!$omp target enter data map(tofrom: p4str) +!$omp target map(from: p4str) + if (.not. associated(p4str)) error stop + if (len(p4str) /= 3) error stop + if (size(p4str) /= 3) error stop + if (lbound(p4str, 1) /= 2) error stop + if (ubound(p4str, 1) /= 4) error stop + p4str(:) = [4_"f85", 4_"8af", 4_"A%F"] +!$omp end target +!$omp target exit data map(from: p4str) +if (.not. associated(p4str)) error stop +if (len(p4str) /= 3) error stop +if (size(p4str) /= 3) error stop +if (lbound(p4str, 1) /= 2) error stop +if (ubound(p4str, 1) /= 4) error stop +if (p4str(2) /= 4_"f85") error stop +if (p4str(3) /= 4_"8af") error stop +if (p4str(4) /= 4_"A%F") error stop + +!$omp target enter data map(tofrom: p4str2) +!$omp target map(from: p4str2) + if (.not. associated(p4str2)) error stop + if (len(p4str2) /= 5) error stop + p4str2 = 4_"9875a" +!$omp end target +!$omp target exit data map(from: p4str2) +if (.not. associated(p4str2)) error stop +if (len(p4str2) /= 5) error stop +if (p4str2 /= 4_"9875a") error stop + + +! character(len=:,kind=4), allocatable :: a4str(:), a4str2 +! allocate(character(len=7,kind=4) :: a4str(-2:3)) +! allocate(character(len=9,kind=4) :: a4str2) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +!$omp target enter data map(tofrom: a4str) +!$omp target map(from: a4str) + if (.not. allocated(a4str)) error stop + if (len(a4str) /= 7) error stop + if (size(a4str) /= 6) error stop + if (lbound(a4str, 1) /= -2) error stop + if (ubound(a4str, 1) /= 3) error stop + ! See PR fortran/107508 why '(:)' is required + a4str(:) = [4_"sf456aq", 4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!$omp end target +!$omp target exit data map(from: a4str) +if (.not. allocated(a4str)) error stop +if (len(a4str) /= 7) error stop +if (size(a4str) /= 6) error stop +if (lbound(a4str, 1) /= -2) error stop +if (ubound(a4str, 1) /= 3) error stop +if (a4str(-2) /= 4_"sf456aq") error stop +if (a4str(-1) /= 4_"3dtzu24") error stop +if (a4str(0) /= 4_"_4fh7sm") error stop +if (a4str(1) /= 4_"=ff85s7") error stop +if (a4str(2) /= 4_"j=8af4d") error stop +if (a4str(3) /= 4_".,A%Fsz") error stop + +!$omp target enter data map(tofrom: a4str2) +!$omp target map(from: a4str2) + if (.not. allocated(a4str2)) error stop + if (len(a4str2) /= 9) error stop + a4str2 = 4_"98765a23d" +!$omp end target +!$omp target exit data map(from: a4str2) +if (.not. allocated(a4str2)) error stop +if (len(a4str2) /= 9) error stop +if (a4str2 /= 4_"98765a23d") error stop + + +deallocate(dt%pc, dt%pc2) +deallocate(dt%pcstr) +deallocate(dt%pcstr2) + +deallocate(dt%pc4str) +deallocate(dt%pc4str2) + +deallocate(ip, ip2, ia, ia2) +deallocate(pstr) +deallocate(pstr2) +deallocate(astr) +deallocate(astr2) + +deallocate(p4str) +deallocate(p4str2) +deallocate(a4str) +deallocate(a4str2) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 new file mode 100644 index 00000000000..80d30edbfc5 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-6.f90 @@ -0,0 +1,392 @@ +! Check that 'map(alloc:' properly works with +! - deferred-length character strings +! - arrays with array descriptors +! For those, the array descriptor / string length must be mapped with 'to:' + +program main +implicit none + +type t + integer :: ic(2:5) + character(len=11) :: ccstr(3:4) + character(len=11,kind=4) :: cc4str(3:7) + integer, pointer :: pc(:) + character(len=:), pointer :: pcstr(:) + character(len=:,kind=4), pointer :: pc4str(:) +end type t + +type(t) :: dt + +integer :: ii(5) +character(len=11) :: clstr(-1:1) +character(len=11,kind=4) :: cl4str(0:3) +integer, pointer :: ip(:) +integer, allocatable :: ia(:) +character(len=:), pointer :: pstr(:) +character(len=:), allocatable :: astr(:) +character(len=:,kind=4), pointer :: p4str(:) +character(len=:,kind=4), allocatable :: a4str(:) + +allocate(dt%pc(5)) +allocate(character(len=2) :: dt%pcstr(2)) + +allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) + +allocate(ip(5), ia(8)) +allocate(character(len=2) :: pstr(-2:0)) +allocate(character(len=6) :: astr(3:5)) + +allocate(character(len=3,kind=4) :: p4str(2:4)) +allocate(character(len=7,kind=4) :: a4str(-2:3)) + + +! integer :: ic(2:5) + +!$omp target enter data map(alloc: dt%ic(3:5)) +dt%ic(2) = 22 +!$omp target map(alloc: dt%ic(3:5)) + if (size(dt%ic) /= 4) error stop + if (lbound(dt%ic, 1) /= 2) error stop + if (ubound(dt%ic, 1) /= 5) error stop + dt%ic(3:5) = [33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%ic(3:5)) +if (size(dt%ic) /= 4) error stop +if (lbound(dt%ic, 1) /= 2) error stop +if (ubound(dt%ic, 1) /= 5) error stop +if (any (dt%ic /= [22, 33, 44, 55])) error stop + +! character(len=11) :: ccstr(3:4) + +!$omp target enter data map(alloc: dt%ccstr(4:4)) +dt%ccstr(3) = "12345678901" +!$omp target map(alloc: dt%ccstr(4:4)) + if (len(dt%ccstr) /= 11) error stop + if (size(dt%ccstr) /= 2) error stop + if (lbound(dt%ccstr, 1) /= 3) error stop + if (ubound(dt%ccstr, 1) /= 4) error stop + dt%ccstr(4:4) = ["abcdefghijk"] +!$omp end target +!$omp target exit data map(from: dt%ccstr(4:4)) +if (len(dt%ccstr) /= 11) error stop +if (size(dt%ccstr) /= 2) error stop +if (lbound(dt%ccstr, 1) /= 3) error stop +if (ubound(dt%ccstr, 1) /= 4) error stop +if (any (dt%ccstr /= ["12345678901", "abcdefghijk"])) error stop + + +! character(len=11,kind=4) :: cc4str(3:7) + +! Value check fails +!$omp target enter data map(alloc: dt%cc4str(4:7)) +dt%cc4str(3) = 4_"12345678901" +!$omp target map(alloc: dt%cc4str(4:7)) + if (len(dt%cc4str) /= 11) error stop + if (size(dt%cc4str) /= 5) error stop + if (lbound(dt%cc4str, 1) /= 3) error stop + if (ubound(dt%cc4str, 1) /= 7) error stop + dt%cc4str(4:7) = [4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4.", & + 4_"45ngwj56sj2"] +!$omp end target +!$omp target exit data map(from: dt%cc4str(4:7)) +if (len(dt%cc4str) /= 11) error stop +if (size(dt%cc4str) /= 5) error stop +if (lbound(dt%cc4str, 1) /= 3) error stop +if (ubound(dt%cc4str, 1) /= 7) error stop +if (dt%cc4str(3) /= 4_"12345678901") error stop +if (dt%cc4str(4) /= 4_"abcdefghijk") error stop +if (dt%cc4str(5) /= 4_"qerftcea6ds") error stop +if (dt%cc4str(6) /= 4_"a1f9g37ga4.") error stop +if (dt%cc4str(7) /= 4_"45ngwj56sj2") error stop + +! integer, pointer :: pc(:) +! allocate(dt%pc(5)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +!$omp target enter data map(alloc: dt%pc(2:5)) +dt%pc(1) = 11 +!$omp target map(alloc: dt%pc(2:5)) + if (.not. associated(dt%pc)) error stop + if (size(dt%pc) /= 5) error stop + if (lbound(dt%pc, 1) /= 1) error stop + if (ubound(dt%pc, 1) /= 5) error stop + dt%pc(2:5) = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: dt%pc(2:5)) +if (.not. associated(dt%pc)) error stop +if (size(dt%pc) /= 5) error stop +if (lbound(dt%pc, 1) /= 1) error stop +if (ubound(dt%pc, 1) /= 5) error stop +if (any (dt%pc /= [11, 22, 33, 44, 55])) error stop + +! character(len=:), pointer :: pcstr(:) +! allocate(character(len=2) :: dt%pcstr(2)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 + +! FIXME: Disabled befause of PR108837 +! +!!$omp target enter data map(alloc: dt%pcstr(2:2)) +!dt%pcstr(1) = "01" +!!$omp target map(alloc: dt%pcstr(2:2)) +! if (.not. associated(dt%pcstr)) error stop +! if (len(dt%pcstr) /= 2) error stop +! if (size(dt%pcstr) /= 2) error stop +! if (lbound(dt%pcstr, 1) /= 1) error stop +! if (ubound(dt%pcstr, 1) /= 2) error stop +! dt%pcstr(2:2) = ["jk"] +!!$omp end target +!!$omp target exit data map(from: dt%pcstr(2:2)) +!if (.not. associated(dt%pcstr)) error stop +!if (len(dt%pcstr) /= 2) error stop +!if (size(dt%pcstr) /= 2) error stop +!if (lbound(dt%pcstr, 1) /= 1) error stop +!if (ubound(dt%pcstr, 1) /= 2) error stop +!if (any (dt%pcstr /= ["01", "jk"])) error stop + + +! character(len=:,kind=4), pointer :: pc4str(:) +! allocate(character(len=3,kind=4) :: dt%pc4str(2:3)) + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x00 +! structure element when other mapped elements from the same structure weren't mapped together with it + +! FIXME: Disabled befause of PR108837 +! +!!$omp target enter data map(alloc: dt%pc4str(3:3)) +!dt%pc4str(2) = 4_"456" +!!$omp target map(alloc: dt%pc4str(3:3)) +! if (.not. associated(dt%pc4str)) error stop +! if (len(dt%pc4str) /= 3) error stop +! if (size(dt%pc4str) /= 2) error stop +! if (lbound(dt%pc4str, 1) /= 2) error stop +! if (ubound(dt%pc4str, 1) /= 3) error stop +! dt%pc4str(3:3) = [4_"tzu"] +!!$omp end target +!!$omp target exit data map(from: dt%pc4str(3:3)) +!if (.not. associated(dt%pc4str)) error stop +!if (len(dt%pc4str) /= 3) error stop +!if (size(dt%pc4str) /= 2) error stop +!if (lbound(dt%pc4str, 1) /= 2) error stop +!if (ubound(dt%pc4str, 1) /= 3) error stop +!if (dt%pc4str(2) /= 4_"456") error stop +!if (dt%pc4str(3) /= 4_"tzu") error stop + +! libgomp: GOMP_target_enter_exit_data unhandled kind 0x01 + +! integer :: ii(5) + +!$omp target enter data map(alloc: ii(2:5)) +ii(1) = -1 +!$omp target map(alloc: ii(2:5)) + if (size(ii) /= 5) error stop + if (lbound(ii, 1) /= 1) error stop + if (ubound(ii, 1) /= 5) error stop + ii(2:5) = [-2, -3, -4, -5] +!$omp end target +!$omp target exit data map(from: ii(2:5)) +if (size(ii) /= 5) error stop +if (lbound(ii, 1) /= 1) error stop +if (ubound(ii, 1) /= 5) error stop +if (any (ii /= [-1, -2, -3, -4, -5])) error stop + + +! character(len=11) :: clstr(-1:1) + +!$omp target enter data map(alloc: clstr(0:1)) +clstr(-1) = "12345678901" +!$omp target map(alloc: clstr(0:1)) + if (len(clstr) /= 11) error stop + if (size(clstr) /= 3) error stop + if (lbound(clstr, 1) /= -1) error stop + if (ubound(clstr, 1) /= 1) error stop + clstr(0:1) = ["abcdefghijk", "ABCDEFGHIJK"] +!$omp end target +!$omp target exit data map(from: clstr(0:1)) +if (len(clstr) /= 11) error stop +if (size(clstr) /= 3) error stop +if (lbound(clstr, 1) /= -1) error stop +if (ubound(clstr, 1) /= 1) error stop +if (any (clstr /= ["12345678901", "abcdefghijk", "ABCDEFGHIJK"])) error stop + +! character(len=11,kind=4) :: cl4str(0:3) + +!$omp target enter data map(alloc: cl4str(1:3)) +cl4str(0) = 4_"12345678901" +!$omp target map(alloc: cl4str(1:3)) + if (len(cl4str) /= 11) error stop + if (size(cl4str) /= 4) error stop + if (lbound(cl4str, 1) /= 0) error stop + if (ubound(cl4str, 1) /= 3) error stop + cl4str(1:3) = [4_"abcdefghijk", & + 4_"qerftcea6ds", 4_"a1f9g37ga4."] +!$omp end target +!$omp target exit data map(from: cl4str(1:3)) +if (len(cl4str) /= 11) error stop +if (size(cl4str) /= 4) error stop +if (lbound(cl4str, 1) /= 0) error stop +if (ubound(cl4str, 1) /= 3) error stop +if (cl4str(0) /= 4_"12345678901") error stop +if (cl4str(1) /= 4_"abcdefghijk") error stop +if (cl4str(2) /= 4_"qerftcea6ds") error stop +if (cl4str(3) /= 4_"a1f9g37ga4.") error stop + + +! allocate(ip(5), ia(8)) + +!$omp target enter data map(alloc: ip(2:5)) +ip(1) = 11 +!$omp target map(alloc: ip(2:5)) + if (.not. associated(ip)) error stop + if (size(ip) /= 5) error stop + if (lbound(ip, 1) /= 1) error stop + if (ubound(ip, 1) /= 5) error stop + ip(2:5) = [22, 33, 44, 55] +!$omp end target +!$omp target exit data map(from: ip(2:5)) +if (.not. associated(ip)) error stop +if (size(ip) /= 5) error stop +if (lbound(ip, 1) /= 1) error stop +if (ubound(ip, 1) /= 5) error stop +if (any (ip /= [11, 22, 33, 44, 55])) error stop + +! allocate(ip(5), ia(8)) + +!$omp target enter data map(alloc: ia(2:8)) +ia(1) = 1 +!$omp target map(alloc: ia(2:8)) + if (.not. allocated(ia)) error stop + if (size(ia) /= 8) error stop + if (lbound(ia, 1) /= 1) error stop + if (ubound(ia, 1) /= 8) error stop + ia(2:8) = [2,3,4,5,6,7,8] +!$omp end target +!$omp target exit data map(from: ia(2:8)) +if (.not. allocated(ia)) error stop +if (size(ia) /= 8) error stop +if (lbound(ia, 1) /= 1) error stop +if (ubound(ia, 1) /= 8) error stop +if (any (ia /= [1,2,3,4,5,6,7,8])) error stop + + +! character(len=:), pointer :: pstr(:) +! allocate(character(len=2) :: pstr(-2:0)) + +! libgomp: nvptx_alloc error: out of memory + +! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR +! +!!$omp target enter data map(alloc: pstr(-1:0)) +!pstr(-2) = "01" +!!$omp target map(alloc: pstr(-1:0)) +! if (.not. associated(pstr)) error stop +! if (len(pstr) /= 2) error stop +! if (size(pstr) /= 3) error stop +! if (lbound(pstr, 1) /= -2) error stop +! if (ubound(pstr, 1) /= 0) error stop +! pstr(-1:0) = ["jk", "aq"] +!!$omp end target +!!$omp target exit data map(from: pstr(-1:0)) +!if (.not. associated(pstr)) error stop +!if (len(pstr) /= 2) error stop +!if (size(pstr) /= 3) error stop +!if (lbound(pstr, 1) /= -2) error stop +!if (ubound(pstr, 1) /= 0) error stop +!if (any (pstr /= ["01", "jk", "aq"])) error stop + + +! character(len=:), allocatable :: astr(:) +! allocate(character(len=6) :: astr(3:5)) + +! libgomp: nvptx_alloc error: out of memory + +! FIXME +!!$omp target enter data map(alloc: astr(4:5)) +!astr(3) = "01db45" +!!$omp target map(alloc: astr(4:5)) +! if (.not. allocated(astr)) error stop +! if (len(astr) /= 6) error stop +! if (size(astr) /= 3) error stop +! if (lbound(astr, 1) /= 3) error stop +! if (ubound(astr, 1) /= 5) error stop +!!! astr(4:5) = ["jk$D%S", "zutg47"] +!!$omp end target +!!!$omp target exit data map(from: astr(4:5)) +!!if (.not. allocated(astr)) error stop +!!!if (len(astr) /= 6) error stop +!if (size(astr) /= 3) error stop +!if (lbound(astr, 1) /= 3) error stop +!if (ubound(astr, 1) /= 5) error stop +!if (any (astr /= ["01db45", "jk$D%S", "zutg47"])) error stop +! + +! character(len=:,kind=4), pointer :: p4str(:) +! allocate(character(len=3,kind=4) :: p4str(2:4)) + +! FAILS with value check + +! FIXME: array offset wrongly calculated as it uses TYPE_SIZE_UNIT, which is a SAVE_EXPR +! +!!$omp target enter data map(alloc: p4str(3:4)) +!p4str(2) = 4_"f85" +!!$omp target map(alloc: p4str(3:4)) +! if (.not. associated(p4str)) error stop +! if (len(p4str) /= 3) error stop +! if (size(p4str) /= 3) error stop +! if (lbound(p4str, 1) /= 2) error stop +! if (ubound(p4str, 1) /= 4) error stop +! p4str(3:4) = [4_"8af", 4_"A%F"] +!!$omp end target +!!$omp target exit data map(from: p4str(3:4)) +!if (.not. associated(p4str)) error stop +!if (len(p4str) /= 3) error stop +!if (size(p4str) /= 3) error stop +!if (lbound(p4str, 1) /= 2) error stop +!if (ubound(p4str, 1) /= 4) error stop +!if (p4str(2) /= 4_"f85") error stop +!if (p4str(3) /= 4_"8af") error stop +!if (p4str(4) /= 4_"A%F") error stop + +! character(len=:,kind=4), allocatable :: a4str(:) +! allocate(character(len=7,kind=4) :: a4str(-2:3)) + +! libgomp: Trying to map into device [0x1027ba0..0x251050bb9c9ebba0) object when [0x7ffd026e6708..0x7ffd026e6710) is already mapped + +! FIXME: Disabled befause of PR108838 +!!$omp target enter data map(alloc: a4str(-1:3)) +!!a4str(-2) = 4_"sf456aq" +!!$omp target map(alloc: a4str(-1:3)) +! if (.not. allocated(a4str)) error stop +! if (len(a4str) /= 7) error stop +! if (size(a4str) /= 6) error stop +! if (lbound(a4str, 1) /= -2) error stop +! if (ubound(a4str, 1) /= 3) error stop +! a4str(-1:3) = [4_"3dtzu24", 4_"_4fh7sm", 4_"=ff85s7", 4_"j=8af4d", 4_".,A%Fsz"] +!!$omp end target +!!$omp target exit data map(from: a4str(-1:3)) +!if (.not. allocated(a4str)) error stop +!if (len(a4str) /= 7) error stop +!if (size(a4str) /= 6) error stop +!if (lbound(a4str, 1) /= -2) error stop +!if (ubound(a4str, 1) /= 3) error stop +!if (a4str(-2) /= 4_"sf456aq") error stop +!if (a4str(-1) /= 4_"3dtzu24") error stop +!if (a4str(0) /= 4_"_4fh7sm") error stop +!if (a4str(1) /= 4_"=ff85s7") error stop +!if (a4str(2) /= 4_"j=8af4d") error stop +!if (a4str(3) /= 4_".,A%Fsz") error stop + +deallocate(dt%pc) +deallocate(dt%pcstr) + +deallocate(dt%pc4str) + +deallocate(ip, ia) +deallocate(pstr) +deallocate(astr) + +deallocate(p4str) +deallocate(a4str) +end diff --git a/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 new file mode 100644 index 00000000000..f129f559336 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/target-enter-data-7.f90 @@ -0,0 +1,78 @@ +module m + implicit none + character(len=:), allocatable :: strA(:), strA2 + character(len=:), pointer :: strP(:), strP2 + !$omp declare target enter(strA,strA2,strP,strP2) +contains + subroutine opt_map(str1, str2, str3) + character(len=:), allocatable :: str1, str2, str3, str4 + optional :: str2, str3 + + if (.not.present(str2)) error stop + if (present(str3)) error stop + + !$omp target map(str1,str2,str3,str4) + if (allocated(str1)) error stop + if (allocated(str2)) error stop + if (present(str3)) error stop + if (allocated(str4)) error stop + !$omp end target + end + subroutine call_opt() + character(len=:), allocatable :: str1, str2 + call opt_map(str1, str2) + end + subroutine test + !$omp declare target + if (.not. allocated(strA)) error stop + !if (.not. allocated(strA2)) error stop + if (.not. associated(strP)) error stop + !if (.not. associated(strP2)) error stop + + ! ensure length was updated as well + if (len(strA) /= 3) error stop + if (len(strA2) /= 5) error stop + if (len(strP) /= 4) error stop + if (len(strP2) /= 8) error stop +! if (any (strA /= ['Hav', 'e f', 'un!'])) error stop +! if (strA2 /= 'Hello') error stop +! if (any (strP /= ['abcd', 'efgh', 'ijkl'])) error stop +! if (strP2 /= 'TestCase') error stop +! +! strA = ['123', '456', '789'] +! strA2 = 'World' +! strP = ['ABCD', 'EFGH', 'IJKL'] +! strP2 = 'Passed!!' + end +end + +program main + use m + implicit none + call call_opt + + strA = ['Hav', 'e f', 'un!'] + strA2 = 'Hello' + allocate(character(len=4) :: strP(3)) + strP = ['abcd', 'efgh', 'ijkl'] + allocate(character(len=8) :: strP2) + strP2 = 'TestCase' + + !$omp target enter data map(always, to: strA, strA2) + !$omp target enter data map(to: strP, strP2) + !$omp target + call test() + !$omp end target + !$omp target exit data map(always, from: strA, strA2, strP, strP2) + + if (len(strA) /= 3) error stop + if (len(strA2) /= 5) error stop + if (len(strP) /= 4) error stop + if (len(strP2) /= 8) error stop +! if (any (strA /= ['123', '456', '789'])) error stop +! if (strA2 /= 'World') error stop +! if (any(strP /= ['ABCD', 'EFGH', 'IJKL'])) error stop +! if (strP2 /= 'Passed!!') error stop + +! deallocate(strP, strP2, strA, strA2) +end