From b20097c65d2e74b1901fba1c55c77f0407e542d2 Mon Sep 17 00:00:00 2001 From: Julian Brown Date: Fri, 5 Jun 2020 14:46:41 -0700 Subject: [PATCH] openacc: Don't strip TO_PSET/POINTER for enter/exit data OpenACC 2.6 specifies that the array descriptor (when present) must be copied to the target before attaching pointers in Fortran. This patch reverses the stripping of GOMP_MAP_TO_PSET and GOMP_MAP_POINTER that was introduced by the "OpenACC reference count overhaul" patch. 2020-07-10 Julian Brown Thomas Schwinge gcc/ * gimplify.c (gimplify_scan_omp_clauses): Do not strip GOMP_MAP_TO_PSET/GOMP_MAP_POINTER for OpenACC enter/exit data directives (see also PR92929). gcc/testsuite/ * gfortran.dg/goacc/finalize-1.f: Update expected dump output. libgomp/ * testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90: New test. Co-Authored-By: Thomas Schwinge --- gcc/gimplify.c | 21 ++-- gcc/testsuite/gfortran.dg/goacc/finalize-1.f | 4 +- .../dynamic-pointer-1.f90 | 97 +++++++++++++++++++ 3 files changed, 110 insertions(+), 12 deletions(-) create mode 100644 libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90 diff --git a/gcc/gimplify.c b/gcc/gimplify.c index 09a30cf69a5..15dfee903ab 100644 --- a/gcc/gimplify.c +++ b/gcc/gimplify.c @@ -8768,6 +8768,8 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, case OMP_TARGET_DATA: case OMP_TARGET_ENTER_DATA: case OMP_TARGET_EXIT_DATA: + case OACC_ENTER_DATA: + case OACC_EXIT_DATA: case OACC_HOST_DATA: if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER || (OMP_CLAUSE_MAP_KIND (c) @@ -8776,15 +8778,6 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, mapped, but not the pointer to it. */ remove = true; break; - case OACC_ENTER_DATA: - case OACC_EXIT_DATA: - if (OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_POINTER - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_TO_PSET - || OMP_CLAUSE_MAP_KIND (c) == GOMP_MAP_FIRSTPRIVATE_POINTER - || (OMP_CLAUSE_MAP_KIND (c) - == GOMP_MAP_FIRSTPRIVATE_REFERENCE)) - remove = true; - break; default: break; } @@ -8794,7 +8787,15 @@ gimplify_scan_omp_clauses (tree *list_p, gimple_seq *pre_p, 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. */ + 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) diff --git a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f index 1e2e3e94b8a..ca642156e9f 100644 --- a/gcc/testsuite/gfortran.dg/goacc/finalize-1.f +++ b/gcc/testsuite/gfortran.dg/goacc/finalize-1.f @@ -21,7 +21,7 @@ !$ACC EXIT DATA FINALIZE DELETE (del_f_p(2:5)) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(release:\\*\\(c_char \\*\\) parm\\.0\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) del_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.0\\.data - \\(sizetype\\) del_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(delete:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:del_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:del_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } !$ACC EXIT DATA COPYOUT (cpo_r) ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:cpo_r\\);$" 1 "original" } } @@ -33,5 +33,5 @@ !$ACC EXIT DATA COPYOUT (cpo_f_p(4:10)) FINALIZE ! { dg-final { scan-tree-dump-times "(?n)#pragma acc exit data map\\(from:\\*\\(c_char \\*\\) parm\\.1\\.data \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:\\(integer\\(kind=1\\)\\\[0:\\\] \\* restrict\\) cpo_f_p\\.data \\\[pointer assign, bias: \\(sizetype\\) parm\\.1\\.data - \\(sizetype\\) cpo_f_p\\.data\\\]\\) finalize;$" 1 "original" } } -! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } +! { dg-final { scan-tree-dump-times "(?n)#pragma omp target oacc_enter_exit_data map\\(force_from:MEM\\\[\\(c_char \\*\\)\[^\\\]\]+\\\] \\\[len: \[^\\\]\]+\\\]\\) map\\(to:cpo_f_p \\\[pointer set, len: \[0-9\]+\\\]\\) map\\(alloc:cpo_f_p\\.data \\\[pointer assign, bias: \[^\\\]\]+\\\]\\) finalize$" 1 "gimple" } } END SUBROUTINE f diff --git a/libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90 b/libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90 new file mode 100644 index 00000000000..4f38902ebc0 --- /dev/null +++ b/libgomp/testsuite/libgomp.oacc-fortran/dynamic-pointer-1.f90 @@ -0,0 +1,97 @@ +! Verify that a 'enter data'ed 'pointer' object creates a persistent, visible device copy + +! { dg-do run } +! { dg-skip-if "" { *-*-* } { "*" } { "-DACC_MEM_SHARED=0" } } + +module m + implicit none +contains + + subroutine verify_a (a_ref, a) + implicit none + integer, dimension (:, :, :), allocatable :: a_ref + integer, dimension (:, :, :), pointer :: a + + !$acc routine seq + + if (any (lbound (a) /= lbound (a_ref))) stop 101 + if (any (ubound (a) /= ubound (a_ref))) stop 102 + if (size (a) /= size (a_ref)) stop 103 + end subroutine verify_a + +end module m + +program main + use m + use openacc + implicit none + integer, parameter :: n = 30 + integer, dimension (:, :, :), allocatable, target :: a1, a2 + integer, dimension (:, :, :), pointer :: p + + allocate (a1(1:n, 0:n-1, 10:n/2)) + !$acc enter data create(a1) + allocate (a2(3:n/3, 10:n, n-10:n+10)) + !$acc enter data create(a2) + + p => a1 + call verify_a(a1, p) + + ! 'p' object isn't present on the device. + !$acc parallel ! Implicit 'copy(p)'; creates 'p' object... + call verify_a(a1, p) + !$acc end parallel ! ..., and deletes it again. + + p => a2 + call verify_a(a2, p) + + ! 'p' object isn't present on the device. + !$acc parallel ! Implicit 'copy(p)'; creates 'p' object... + call verify_a(a2, p) + !$acc end parallel ! ..., and deletes it again. + + p => a1 + + !$acc enter data create(p) + ! 'p' object is now present on the device (visible device copy). + !TODO PR96080 if (.not. acc_is_present (p)) stop 1 + + !$acc parallel + ! On the device, got created as 'p => a1'. + call verify_a(a1, p) + !$acc end parallel + call verify_a(a1, p) + + !$acc parallel + p => a2 + ! On the device, 'p => a2' is now set. + call verify_a(a2, p) + !$acc end parallel + ! On the host, 'p => a1' persists. + call verify_a(a1, p) + + !$acc parallel + ! On the device, 'p => a2' persists. + call verify_a(a2, p) + !$acc end parallel + ! On the host, 'p => a1' still persists. + call verify_a(a1, p) + + p => a2 + + !$acc parallel + p => a1 + ! On the device, 'p => a1' is now set. + call verify_a(a1, p) + !$acc end parallel + ! On the host, 'p => a2' persists. + call verify_a(a2, p) + + !$acc parallel + ! On the device, 'p => a1' persists. + call verify_a(a1, p) + !$acc end parallel + ! On the host, 'p => a2' still persists. + call verify_a(a2, p) + +end program main -- 2.30.2