mirror/gcc
2
0
Fork 0
mirror of git://gcc.gnu.org/git/gcc.git synced 2025-04-05 00:31:30 +08:00
Code Issues Packages Projects Releases Wiki Activity

openacc: No attach/detach present/release mappings for array descriptors

Browse Source 
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  <julian@codesourcery.com>
	    Thomas Schwinge  <thomas@codesourcery.com>

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 <thomas@codesourcery.com>
This commit is contained in:
Julian Brown 2020-07-27 06:29:02 -07:00
parent 105fe3e0b8
commit f2f4212e20
4 changed files with 128 additions and 31 deletions
Show Stats Download Patch File Download Diff File Expand all files Collapse all files

44
gcc/fortran/trans-openmp.c
View File

@ -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. */

7
gcc/gimplify.c
View File

@ -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. */

17
gcc/testsuite/gfortran.dg/goacc/attach-descriptor.f90
View File

@ -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

91
libgomp/testsuite/libgomp.oacc-fortran/attach-descriptor-1.f90
View File

@ -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