libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-runtime.f90

   1 ! Test OpenACC 'declare create' with allocatable arrays.
   2
   3 ! { dg-do run }
   4
   5 ! Note that we're not testing OpenACC semantics here, but rather documenting
   6 ! current GCC behavior, specifically, behavior concerning updating of
   7 ! host/device array descriptors.
   8 ! { dg-skip-if n/a { *-*-* } { -DACC_MEM_SHARED=1 } }
   9
  10 !TODO-OpenACC-declare-allocate
  11 ! Missing support for OpenACC "Changes from Version 2.0 to 2.5":
  12 ! "The 'declare create' directive with a Fortran 'allocatable' has new behavior".
  13 ! Thus, after 'allocate'/before 'deallocate', call 'acc_create'/'acc_delete'
  14 ! manually.
  15
  16
  17 !TODO { dg-additional-options -fno-inline } for stable results regarding OpenACC 'routine'.
  18
  19
  20 !TODO OpenACC 'serial' vs. GCC/nvptx:
  21 !TODO { dg-prune-output {using 'vector_length \(32\)', ignoring 1} }
  22
  23
  24 ! { dg-additional-options -fdump-tree-original }
  25 ! { dg-additional-options -fdump-tree-gimple }
  26
  27
  28 module vars
  29   implicit none
  30   integer, parameter :: n1_lb = -3
  31   integer, parameter :: n1_ub = 6
  32   integer, parameter :: n2_lb = -9999
  33   integer, parameter :: n2_ub = 22222
  34
  35   integer, allocatable :: b(:)
  36   !$acc declare create (b)
  37
  38 end module vars
  39
  40 program test
  41   use vars
  42   use openacc
  43   implicit none
  44   integer :: i
  45
  46   ! Identifiers for purposes of reliable '-fdump-tree-[...]' scanning.
  47   integer :: id1_1, id1_2
  48
  49   interface
  50
  51      subroutine verify_initial
  52        implicit none
  53        !$acc routine seq
  54      end subroutine verify_initial
  55
  56      subroutine verify_n1_allocated
  57        implicit none
  58        !$acc routine seq
  59      end subroutine verify_n1_allocated
  60
  61      subroutine verify_n1_values (addend)
  62        implicit none
  63        !$acc routine gang
  64        integer, value :: addend
  65      end subroutine verify_n1_values
  66
  67      subroutine verify_n1_deallocated (expect_allocated)
  68        implicit none
  69        !$acc routine seq
  70        logical, value :: expect_allocated
  71      end subroutine verify_n1_deallocated
  72
  73      subroutine verify_n2_allocated
  74        implicit none
  75        !$acc routine seq
  76      end subroutine verify_n2_allocated
  77
  78      subroutine verify_n2_values (addend)
  79        implicit none
  80        !$acc routine gang
  81        integer, value :: addend
  82      end subroutine verify_n2_values
  83
  84      subroutine verify_n2_deallocated (expect_allocated)
  85        implicit none
  86        !$acc routine seq
  87        logical, value :: expect_allocated
  88      end subroutine verify_n2_deallocated
  89
  90   end interface
  91
  92   call acc_create (id1_1)
  93   call acc_create (id1_2)
  94
  95   call verify_initial
  96   ! It is important here (and similarly, following) that there is no data
  97   ! clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
  98   !$acc serial
  99   call verify_initial
 100   !$acc end serial
 101
 102   allocate (b(n1_lb:n1_ub))
 103   call verify_n1_allocated
 104   if (acc_is_present (b)) error stop
 105   call acc_create (b)
 106   ! This is now OpenACC "present":
 107   if (.not.acc_is_present (b)) error stop
 108   ! This still has the initial array descriptor:
 109   !$acc serial
 110   call verify_initial
 111   !$acc end serial
 112
 113   do i = n1_lb, n1_ub
 114      b(i) = i - 1
 115   end do
 116
 117   ! Verify that host-to-device copy doesn't touch the device-side (still
 118   ! initial) array descriptor (but it does copy the array data).
 119   call acc_update_device (b)
 120   !$acc serial
 121   call verify_initial
 122   !$acc end serial
 123
 124   b = 40
 125
 126   ! Verify that device-to-host copy doesn't touch the host-side array
 127   ! descriptor, doesn't copy out the device-side (still initial) array
 128   ! descriptor (but it does copy the array data).
 129   call acc_update_self (b)
 130   call verify_n1_allocated
 131
 132   do i = n1_lb, n1_ub
 133      if (b(i) /= i - 1) error stop
 134      b(i) = b(i) + 2
 135   end do
 136
 137   ! The same using the OpenACC 'update' directive.
 138
 139   !$acc update device (b) self (id1_1)
 140   ! We do have 'GOMP_MAP_TO_PSET' here:
 141   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_to:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1\);$} 1 original } }
 142   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_to:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
 143   ! ..., but it's silently skipped in 'GOACC_update'.
 144   !$acc serial
 145   call verify_initial
 146   !$acc end serial
 147
 148   b = 41
 149
 150   !$acc update self (b) self (id1_2)
 151   ! We do have 'GOMP_MAP_TO_PSET' here:
 152   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc update map\(force_from:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2\);$} 1 original } }
 153   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_update map\(force_from:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(force_from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
 154   ! ..., but it's silently skipped in 'GOACC_update'.
 155   call verify_n1_allocated
 156
 157   do i = n1_lb, n1_ub
 158      if (b(i) /= i + 1) error stop
 159      b(i) = b(i) + 2
 160   end do
 161
 162   ! Now install the actual array descriptor, via a data clause for 'b'
 163   ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
 164   ! 'gomp_map_vars_internal' is handled as 'declare target', and because of
 165   ! '*(void **) hostaddrs[i] != NULL', we've got 'has_always_ptrset == true',
 166   ! 'always_to_cnt == 1', and therefore 'gomp_map_vars_existing' does update
 167   ! the 'GOMP_MAP_TO_PSET'.
 168   !$acc serial present (b) copyin (id1_1)
 169   call verify_initial
 170   id1_1 = 0
 171   !$acc end serial
 172   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1\)$} 1 original } }
 173   !TODO ..., but without an actual use of 'b', the gimplifier removes the
 174   !TODO 'GOMP_MAP_TO_PSET':
 175   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
 176   !$acc serial present (b) copyin (id1_2)
 177   call verify_n1_allocated
 178   !TODO Use of 'b':
 179   id1_2 = ubound (b, 1)
 180   !$acc end serial
 181   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
 182   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
 183
 184   !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
 185   call verify_n1_values (1)
 186   id1_1 = 0
 187   !$acc end parallel
 188   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } }
 189   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
 190
 191   !$acc parallel copy (b) copyin (id1_2)
 192   ! As already present, 'copy (b)' doesn't copy; addend is still '1'.
 193   call verify_n1_values (1)
 194   id1_2 = 0
 195   !$acc end parallel
 196   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(tofrom:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2\)$} 1 original } }
 197   !TODO ..., but without an actual use of 'b', the gimplifier removes the
 198   !TODO 'GOMP_MAP_TO_PSET':
 199   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(tofrom:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(to:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
 200
 201   call verify_n1_allocated
 202   if (.not.acc_is_present (b)) error stop
 203
 204   call acc_delete (b)
 205   if (.not.allocated (b)) error stop
 206   if (acc_is_present (b)) error stop
 207   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 208   ! as "allocated":
 209   !$acc serial
 210   call verify_n1_allocated
 211   !$acc end serial
 212
 213   deallocate (b)
 214   call verify_n1_deallocated (.false.)
 215   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 216   ! as "allocated":
 217   !$acc serial
 218   call verify_n1_allocated
 219   !$acc end serial
 220
 221   ! Now try to install the actual array descriptor, via a data clause for 'b'
 222   ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
 223   ! 'gomp_map_vars_internal' is handled as 'declare target', but because of
 224   ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false',
 225   ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update
 226   ! the 'GOMP_MAP_TO_PSET'.
 227   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 228   ! as "allocated":
 229   !TODO Why does 'present (b)' still work here?
 230   !$acc serial present (b) copyout (id1_2)
 231   call verify_n1_deallocated (.true.)
 232   !TODO Use of 'b'.
 233   id1_2 = ubound (b, 1)
 234   !$acc end serial
 235   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc serial map\(force_present:\*\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:\(integer\(kind=[0-9]+\)\[0:\] \* restrict\) b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2\)$} 1 original } }
 236   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_serial map\(force_present:MEM <integer\(kind=[0-9]+\)\[0:\]> \[\(integer\(kind=[0-9]+\)\[0:\] \*\)[^\]]+\] \[len: [^\]]+\]\) map\(to:b \[pointer set, len: [0-9]+\]\) map\(alloc:b\.data \[pointer assign, bias: 0\]\) map\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
 237
 238
 239   ! Restart the procedure, with different array dimensions.
 240
 241   allocate (b(n2_lb:n2_ub))
 242   call verify_n2_allocated
 243   if (acc_is_present (b)) error stop
 244   call acc_create (b)
 245   if (.not.acc_is_present (b)) error stop
 246   ! This still has the previous (n1) array descriptor:
 247   !$acc serial
 248   call verify_n1_deallocated (.true.)
 249   !$acc end serial
 250
 251   do i = n2_lb, n2_ub
 252      b(i) = i + 20
 253   end do
 254
 255   call acc_update_device (b)
 256   !$acc serial
 257   call verify_n1_deallocated (.true.)
 258   !$acc end serial
 259
 260   b = -40
 261
 262   call acc_update_self (b)
 263   call verify_n2_allocated
 264
 265   do i = n2_lb, n2_ub
 266      if (b(i) /= i + 20) error stop
 267      b(i) = b(i) - 40
 268   end do
 269
 270   !$acc update device (b)
 271   !$acc serial
 272   call verify_n1_deallocated (.true.)
 273   !$acc end serial
 274
 275   b = -41
 276
 277   !$acc update self (b)
 278   call verify_n2_allocated
 279
 280   do i = n2_lb, n2_ub
 281      if (b(i) /= i - 20) error stop
 282      b(i) = b(i) + 10
 283   end do
 284
 285   !$acc serial present (b) copy (id1_2)
 286   call verify_n2_allocated
 287   !TODO Use of 'b':
 288   id1_2 = ubound (b, 1)
 289   !$acc end serial
 290
 291   !$acc parallel
 292   call verify_n2_values (-20)
 293   !$acc end parallel
 294
 295   !$acc parallel copy (b)
 296   call verify_n2_values (-20)
 297   !$acc end parallel
 298
 299   call verify_n2_allocated
 300   if (.not.acc_is_present (b)) error stop
 301
 302   call acc_delete (b)
 303   if (.not.allocated (b)) error stop
 304   if (acc_is_present (b)) error stop
 305   !$acc serial
 306   call verify_n2_allocated
 307   !$acc end serial
 308
 309   deallocate (b)
 310   call verify_n2_deallocated (.false.)
 311   !$acc serial
 312   call verify_n2_allocated
 313   !$acc end serial
 314
 315   !$acc serial present (b) copy (id1_2)
 316   call verify_n2_deallocated (.true.)
 317   !TODO Use of 'b':
 318   id1_2 = ubound (b, 1)
 319   !$acc end serial
 320
 321 end program test
 322
 323
 324 subroutine verify_initial
 325   use vars
 326   implicit none
 327   !$acc routine seq
 328
 329   if (allocated (b)) error stop "verify_initial allocated"
 330   if (any (lbound (b) /= [0])) error stop "verify_initial lbound"
 331   if (any (ubound (b) /= [0])) error stop "verify_initial ubound"
 332 end subroutine verify_initial
 333
 334 subroutine verify_n1_allocated
 335   use vars
 336   implicit none
 337   !$acc routine seq
 338
 339   if (.not.allocated (b)) error stop "verify_n1_allocated allocated"
 340   if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound"
 341   if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound"
 342 end subroutine verify_n1_allocated
 343
 344 subroutine verify_n1_values (addend)
 345   use vars
 346   implicit none
 347   !$acc routine gang
 348   integer, value :: addend
 349   integer :: i
 350
 351   !$acc loop
 352   do i = n1_lb, n1_ub
 353      if (b(i) /= i + addend) error stop
 354   end do
 355 end subroutine verify_n1_values
 356
 357 subroutine verify_n1_deallocated (expect_allocated)
 358   use vars
 359   implicit none
 360   !$acc routine seq
 361   logical, value :: expect_allocated
 362
 363   if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated"
 364   ! Apparently 'deallocate'ing doesn't unset the bounds.
 365   if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound"
 366   if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound"
 367 end subroutine verify_n1_deallocated
 368
 369 subroutine verify_n2_allocated
 370   use vars
 371   implicit none
 372   !$acc routine seq
 373
 374   if (.not.allocated(b)) error stop "verify_n2_allocated allocated"
 375   if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound"
 376   if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound"
 377 end subroutine verify_n2_allocated
 378
 379 subroutine verify_n2_values (addend)
 380   use vars
 381   implicit none
 382   !$acc routine gang
 383   integer, value :: addend
 384   integer :: i
 385
 386   !$acc loop
 387   do i = n2_lb, n2_ub
 388      if (b(i) /= i + addend) error stop
 389   end do
 390 end subroutine verify_n2_values
 391
 392 subroutine verify_n2_deallocated (expect_allocated)
 393   use vars
 394   implicit none
 395   !$acc routine seq
 396   logical, value :: expect_allocated
 397
 398   if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated"
 399   ! Apparently 'deallocate'ing doesn't unset the bounds.
 400   if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound"
 401   if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound"
 402 end subroutine verify_n2_deallocated