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);
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. */
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. */
-! { dg-additional-options "-fdump-tree-original" }
+! { dg-additional-options "-fdump-tree-original -fdump-tree-gimple" }
program att
implicit none
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
! { 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(:)
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
projects / gcc.git / commitdiff