From mboxrd@z Thu Jan 1 00:00:00 1970 Return-Path: Received: from esa2.mentor.iphmx.com (esa2.mentor.iphmx.com [68.232.141.98]) by sourceware.org (Postfix) with ESMTPS id 3D5183858D37; Thu, 27 Apr 2023 18:37:02 +0000 (GMT) DMARC-Filter: OpenDMARC Filter v1.4.2 sourceware.org 3D5183858D37 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.99,232,1677571200"; d="scan'208";a="4157311" Received: from orw-gwy-01-in.mentorg.com ([192.94.38.165]) by esa2.mentor.iphmx.com with ESMTP; 27 Apr 2023 10:37:00 -0800 IronPort-SDR: G9CCjqZ5JkbQHdr/TonNf6+xhvuBln4ew9zLxpnjr4OJXVwJ83Qek2svYyCt3rQDW1sCGyVRiZ UEtFYdZqWtvSxsy/mleFa2Ir29btIKEDbJ9m/IasSwtDz3AG800byN1TrlWZxzskgUgbH2Cht3 +pzDXX/XjtF/0qOjy/Sy+HlWnxduy4+I+BwTgdMKjZ2mnP/UokFlCt2/qOdylXfUgd3BtQjcrf nmUl8z2CUwfx6jP+N3X7TY75iJuWQGKbo/IAjkwbSqOwXGlL+mbGNjmLEqKv4TTl6r9rad8rMu XaY= From: Julian Brown To: CC: , , Subject: [PATCH] OpenACC: Stand-alone attach/detach clause fixes for Fortran [PR109622] Date: Thu, 27 Apr 2023 11:36:47 -0700 Message-ID: <20230427183647.99112-1-julian@codesourcery.com> X-Mailer: git-send-email 2.29.2 MIME-Version: 1.0 Content-Transfer-Encoding: 8bit Content-Type: text/plain X-Originating-IP: [137.202.0.90] X-ClientProxiedBy: svr-ies-mbx-15.mgc.mentorg.com (139.181.222.15) To svr-ies-mbx-11.mgc.mentorg.com (139.181.222.11) X-Spam-Status: No, score=-11.7 required=5.0 tests=BAYES_00,GIT_PATCH_0,HEADER_FROM_DIFFERENT_DOMAINS,KAM_DMARC_STATUS,RCVD_IN_MSPIKE_H2,SPF_HELO_PASS,SPF_PASS,TXREP,T_SCC_BODY_TEXT_LINE autolearn=ham autolearn_force=no version=3.4.6 X-Spam-Checker-Version: SpamAssassin 3.4.6 (2021-04-09) on server2.sourceware.org List-Id: This patch fixes several cases where multiple attach or detach mapping nodes were being created for stand-alone attach or detach clauses in Fortran. After the introduction of stricter checking later during compilation, these extra nodes could cause ICEs, as seen in the PR. The patch also fixes cases that "happened to work" previously where the user attaches/detaches a pointer to array using a descriptor, and (I think!) the "_data" field has offset zero, hence the same address as the descriptor as a whole. Tested with offloading to nvptx. OK? Thanks, Julian 2023-04-27 Julian Brown PR fortran/109622 gcc/fortran/ * trans-openmp.cc (gfc_trans_omp_clauses): Attach/detach clause fixes. gcc/testsuite/ * gfortran.dg/goacc/attach-descriptor.f90: Adjust expected output. libgomp/ * testsuite/libgomp.fortran/pr109622.f90: New test. * testsuite/libgomp.fortran/pr109622-2.f90: New test. * testsuite/libgomp.fortran/pr109622-3.f90: New test. --- gcc/fortran/trans-openmp.cc | 36 +++++++++++++++++-- .../gfortran.dg/goacc/attach-descriptor.f90 | 12 +++---- .../testsuite/libgomp.fortran/pr109622-2.f90 | 32 +++++++++++++++++ .../testsuite/libgomp.fortran/pr109622-3.f90 | 32 +++++++++++++++++ .../testsuite/libgomp.fortran/pr109622.f90 | 32 +++++++++++++++++ 5 files changed, 135 insertions(+), 9 deletions(-) create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622-2.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622-3.f90 create mode 100644 libgomp/testsuite/libgomp.fortran/pr109622.f90 diff --git a/gcc/fortran/trans-openmp.cc b/gcc/fortran/trans-openmp.cc index 4ff9c59df5cb..dbb4a335ab57 100644 --- a/gcc/fortran/trans-openmp.cc +++ b/gcc/fortran/trans-openmp.cc @@ -3388,6 +3388,17 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, gfc_add_block_to_block (block, &se.post); if (pointer || allocatable) { + /* If it's a bare attach/detach clause, we just want + to perform a single attach/detach operation, of the + pointer itself, not of the pointed-to object. */ + if (openacc + && (n->u.map_op == OMP_MAP_ATTACH + || n->u.map_op == OMP_MAP_DETACH)) + { + OMP_CLAUSE_SIZE (node) = size_zero_node; + goto finalize_map_clause; + } + node2 = build_omp_clause (input_location, OMP_CLAUSE_MAP); gomp_map_kind kind @@ -3458,6 +3469,19 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, { if (pointer || (openacc && allocatable)) { + /* If it's a bare attach/detach clause, we just want + to perform a single attach/detach operation, of the + pointer itself, not of the pointed-to object. */ + if (openacc + && (n->u.map_op == OMP_MAP_ATTACH + || n->u.map_op == OMP_MAP_DETACH)) + { + OMP_CLAUSE_DECL (node) + = build_fold_addr_expr (inner); + OMP_CLAUSE_SIZE (node) = size_zero_node; + goto finalize_map_clause; + } + tree data, size; if (lastref->u.c.component->ts.type == BT_CLASS) @@ -3494,12 +3518,18 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, else if (lastref->type == REF_ARRAY && lastref->u.ar.type == AR_FULL) { - /* Just pass the (auto-dereferenced) decl through for - bare attach and detach clauses. */ + /* Bare attach and detach clauses don't want any + additional nodes. */ if (n->u.map_op == OMP_MAP_ATTACH || n->u.map_op == OMP_MAP_DETACH) { - OMP_CLAUSE_DECL (node) = inner; + if (GFC_DESCRIPTOR_TYPE_P (TREE_TYPE (inner))) + { + tree ptr = gfc_conv_descriptor_data_get (inner); + OMP_CLAUSE_DECL (node) = ptr; + } + else + OMP_CLAUSE_DECL (node) = inner; OMP_CLAUSE_SIZE (node) = size_zero_node; goto finalize_map_clause; } diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 8c2ee4a5cca4..734afbe6ca48 100644 --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -11,19 +11,19 @@ program att integer, pointer :: myptr(:) !$acc enter data attach(myvar%arr2, myptr) -! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc enter data map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_data map\\(attach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(attach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } !$acc exit data detach(myvar%arr2, myptr) -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\);$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(detach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:myptr\\.data \\\[bias: 0\\\]\\)$" 1 "gimple" } } ! Test valid usage and processing of the finalize clause. !$acc exit data detach(myvar%arr2, myptr) finalize -! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\* restrict\\) myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(detach:\\(integer\\(kind=4\\)\\\[0:\\\] \\*\\) myptr\\.data \\\[bias: 0\\\]\\) finalize;$" 1 "original" } } ! For array-descriptor detaches, we no longer generate a "release" mapping ! for the pointed-to data for gimplify.c to turn into "delete". Make sure ! the mapping still isn't there. -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2 \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_exit_data map\\(force_detach:myvar\\.arr2\\.data \\\[bias: 0\\\]\\) map\\(to:myptr \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(force_detach:myptr\\.data \\\[bias: 0\\\]\\) finalize$" 1 "gimple" } } end program att diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 new file mode 100644 index 000000000000..8c5f373f39f7 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/pr109622-2.f90 @@ -0,0 +1,32 @@ +! { dg-do run } + +type t +integer :: foo +integer, pointer :: bar +end type t + +type(t) :: var +integer, target :: tgt + +var%bar => tgt + +var%foo = 99 +tgt = 199 + +!$acc enter data copyin(var) + +!$acc enter data attach(var%bar) + +!$acc serial +var%foo = 5 +var%bar = 7 +!$acc end serial + +!$acc exit data detach(var%bar) + +!$acc exit data copyout(var) + +if (var%foo.ne.5) stop 1 +if (tgt.ne.7) stop 2 + +end diff --git a/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 new file mode 100644 index 000000000000..3ee1b43a7464 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/pr109622-3.f90 @@ -0,0 +1,32 @@ +! { dg-do run } + +type t +integer :: foo +integer, pointer :: bar(:) +end type t + +type(t) :: var +integer, target :: tgt(20) + +var%bar => tgt + +var%foo = 99 +tgt = 199 + +!$acc enter data copyin(var, tgt) + +!$acc enter data attach(var%bar) + +!$acc serial +var%foo = 5 +var%bar = 7 +!$acc end serial + +!$acc exit data detach(var%bar) + +!$acc exit data copyout(var, tgt) + +if (var%foo.ne.5) stop 1 +if (any(tgt.ne.7)) stop 2 + +end diff --git a/libgomp/testsuite/libgomp.fortran/pr109622.f90 b/libgomp/testsuite/libgomp.fortran/pr109622.f90 new file mode 100644 index 000000000000..5b8c4102f768 --- /dev/null +++ b/libgomp/testsuite/libgomp.fortran/pr109622.f90 @@ -0,0 +1,32 @@ +! { dg-do run } + +type t +integer :: value +type(t), pointer :: chain +end type t + +type(t), target :: var, var2 + +var%value = 99 +var2%value = 199 + +var%chain => var2 +nullify(var2%chain) + +!$acc enter data copyin(var, var2) + +!$acc enter data attach(var%chain) + +!$acc serial +var%value = 5 +var%chain%value = 7 +!$acc end serial + +!$acc exit data detach(var%chain) + +!$acc exit data copyout(var, var2) + +if (var%value.ne.5) stop 1 +if (var2%value.ne.7) stop 2 + +end -- 2.29.2