From: Julian Brown Date: Mon, 27 Jul 2020 13:29:02 +0000 (-0700) Subject: openacc: No attach/detach present/release mappings for array descriptors X-Git-Url: https://git.libre-soc.org/?a=commitdiff_plain;h=f2f4212e20c5dc4422aea355abe9a3bf0e05db71;p=gcc.git openacc: No attach/detach present/release mappings for array descriptors Standalone attach and detach clauses should not create present/release mappings for Fortran array descriptors (e.g. used when we have a pointer to an array), both because it is unnecessary and because those mappings will be incorrectly subject to reference counting. Simply omitting the mappings means we just use GOMP_MAP_TO_PSET and GOMP_MAP_{ATTACH,DETACH} mappings for array descriptors. That requires a tweak in gimplify.c, since we may now see GOMP_MAP_TO_PSET without a preceding data-movement mapping. 2020-08-03 Julian Brown Thomas Schwinge gcc/fortran/ * trans-openmp.c (gfc_trans_omp_clauses): Don't create present/release mappings for array descriptors. gcc/ * gimplify.c (gimplify_omp_target_update): Allow GOMP_MAP_TO_PSET without a preceding data-movement mapping. gcc/testsuite/ * gfortran.dg/goacc/attach-descriptor.f90: Update pattern output. Add scanning of gimplify dump. libgomp/ * testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90: Don't run for shared-memory devices. Extend with further checking. Co-Authored-By: Thomas Schwinge --- diff --git a/gcc/fortran/trans-openmp.c b/gcc/fortran/trans-openmp.c index 076efb03831..98702b1aa22 100644 --- a/gcc/fortran/trans-openmp.c +++ b/gcc/fortran/trans-openmp.c @@ -2718,23 +2718,6 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, OMP_CLAUSE_SIZE (node2) = TYPE_SIZE_UNIT (type); node3 = build_omp_clause (input_location, OMP_CLAUSE_MAP); - 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 (node, GOMP_MAP_ALLOC); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_ATTACH); - } - else if (n->u.map_op == OMP_MAP_DETACH) - { - OMP_CLAUSE_SET_MAP_KIND (node, GOMP_MAP_RELEASE); - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_DETACH); - } - else - OMP_CLAUSE_SET_MAP_KIND (node3, GOMP_MAP_POINTER); if (present) { ptr = gfc_conv_descriptor_data_get (decl); @@ -2748,6 +2731,33 @@ gfc_trans_omp_clauses (stmtblock_t *block, gfc_omp_clauses *clauses, 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, GOMP_MAP_POINTER); /* We have to check for n->sym->attr.dimension because of scalar coarrays. */ diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 15dfee903ab..6a5349cbfdd 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -13013,8 +13013,13 @@ gimplify_omp_target_update (tree *expr_p, gimple_seq *pre_p) OMP_CLAUSE_SET_MAP_KIND (c, GOMP_MAP_DELETE); have_clause = true; break; - case GOMP_MAP_POINTER: case GOMP_MAP_TO_PSET: + /* Fortran arrays with descriptors must map that descriptor when + doing standalone "attach" operations (in OpenACC). In that + case GOMP_MAP_TO_PSET appears by itself with no preceding + clause (see trans-openmp.c:gfc_trans_omp_clauses). */ + break; + case GOMP_MAP_POINTER: /* TODO PR92929: we may see these here, but they'll always follow one of the clauses above, and will be handled by libgomp as one group, so no handling required here. */ diff --git a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 index 9ca36f770c7..373bdcb2114 100644 --- a/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 +++ b/gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90 @@ -1,4 +1,4 @@ -! { dg-additional-options "-fdump-tree-original" } +! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" } program att implicit none @@ -11,8 +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\\(alloc:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) 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 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_exit_data map\\(attach:myvar\\.arr2 \\\[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\\(release:\\*\\(c_char \\*\\) myptr\\.data \\\[len: \[^\\\]\]+\\\]\\) 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 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_enter_exit_data map\\(detach:myvar\\.arr2 \\\[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" } } + ! 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_enter_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" } } + end program att diff --git a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 index 5d79cbc14fc..960b9f94507 100644 --- a/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 +++ b/libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90 @@ -1,8 +1,10 @@ ! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } -program att +subroutine test(variant) use openacc implicit none + integer :: variant type t integer :: arr1(10) integer, allocatable :: arr2(:) @@ -26,28 +28,97 @@ program att myptr => tarr - !$acc enter data attach(myvar%arr2, myptr) + if (variant == 0 & + .or. variant == 3 & + .or. variant == 5) then + !$acc enter data attach(myvar%arr2, myptr) + else if (variant == 1 & + .or. variant == 2 & + .or. variant == 4) then + !$acc enter data attach(myvar%arr2, myptr) + !$acc enter data attach(myvar%arr2, myptr) + else + ! Internal error. + stop 1 + end if ! FIXME: This warning is emitted on the wrong line number. - ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 38 } + ! { dg-warning "using vector_length \\(32\\), ignoring 1" "" { target openacc_nvidia_accel_selected } 52 } !$acc serial present(myvar%arr2) do i=1,10 - myvar%arr1(i) = i - myvar%arr2(i) = i + myvar%arr1(i) = i + variant + myvar%arr2(i) = i - variant end do - myptr(3) = 99 + myptr(3) = 99 - variant !$acc end serial - !$acc exit data detach(myvar%arr2, myptr) + if (variant == 0) then + !$acc exit data detach(myvar%arr2, myptr) + else if (variant == 1) then + !$acc exit data detach(myvar%arr2, myptr) + !$acc exit data detach(myvar%arr2, myptr) + else if (variant == 2) then + !$acc exit data detach(myvar%arr2, myptr) + !$acc exit data detach(myvar%arr2, myptr) finalize + else if (variant == 3 & + .or. variant == 4) then + !$acc exit data detach(myvar%arr2, myptr) finalize + else if (variant == 5) then + ! Do not detach. + else + ! Internal error. + stop 2 + end if + + if (.not. acc_is_present(myvar%arr2)) stop 10 + if (.not. acc_is_present(myvar)) stop 11 + if (.not. acc_is_present(tarr)) stop 12 call acc_copyout(myvar%arr2) + if (acc_is_present(myvar%arr2)) stop 20 + if (.not. acc_is_present(myvar)) stop 21 + if (.not. acc_is_present(tarr)) stop 22 call acc_copyout(myvar) + if (acc_is_present(myvar%arr2)) stop 30 + if (acc_is_present(myvar)) stop 31 + if (.not. acc_is_present(tarr)) stop 32 call acc_copyout(tarr) + if (acc_is_present(myvar%arr2)) stop 40 + if (acc_is_present(myvar)) stop 41 + if (acc_is_present(tarr)) stop 42 do i=1,10 - if (myvar%arr1(i) .ne. i) stop 1 - if (myvar%arr2(i) .ne. i) stop 2 + if (myvar%arr1(i) .ne. i + variant) stop 50 + if (variant == 5) then + ! We have not detached, so have copyied out a device pointer, so cannot + ! access 'myvar%arr2' on the host. + else + if (myvar%arr2(i) .ne. i - variant) stop 51 + end if end do - if (tarr(3) .ne. 99) stop 3 + if (tarr(3) .ne. 99 - variant) stop 52 + + if (variant == 5) then + ! If not explicitly stopping here, we'd in the following try to deallocate + ! the device pointer on the host, SIGSEGV. + stop + end if +end subroutine test + +program att + implicit none + + call test(0) + + call test(1) + + call test(2) + + call test(3) + + call test(4) + call test(5) + ! Make sure that 'test(5)' has stopped the program. + stop 60 end program att