libgomp/testsuite/libgomp.oacc-fortran/declare-allocatable-array_descriptor-1-directive.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', do
  14 ! '!$acc enter data create'/'!$acc exit data delete' 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   !$acc enter data create (b)
 106   ! This is now OpenACC "present":
 107   if (.not.acc_is_present (b)) error stop
 108   ! ..., and got the actual array descriptor installed:
 109   !$acc serial
 110   call verify_n1_allocated
 111   !$acc end serial
 112
 113   do i = n1_lb, n1_ub
 114      b(i) = i - 1
 115   end do
 116
 117   ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
 118   ! that host-to-device copy doesn't touch the device-side (still initial)
 119   ! array descriptor (but it does copy the array data").  This is here not
 120   ! applicable anymore, as we've already gotten the actual array descriptor
 121   ! installed.  Thus now verify that it does copy the array data.
 122   call acc_update_device (b)
 123   !$acc serial
 124   call verify_n1_allocated
 125   !$acc end serial
 126
 127   b = 40
 128
 129   !$acc parallel copyout (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
 130   call verify_n1_values (-1)
 131   id1_1 = 0
 132   !$acc end parallel
 133   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(from:id1_1\)$} 1 original } }
 134   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(from:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
 135
 136   !$acc parallel copy (b) copyout (id1_2)
 137   ! As already present, 'copy (b)' doesn't copy; addend is still '-1'.
 138   call verify_n1_values (-1)
 139   id1_2 = 0
 140   !$acc end parallel
 141   ! { 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\(from:id1_2\)$} 1 original } }
 142   !TODO ..., but without an actual use of 'b', the gimplifier removes the
 143   !TODO 'GOMP_MAP_TO_PSET':
 144   ! { 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\(from:id1_2 \[len: [0-9]+\]\)$} 1 gimple } }
 145
 146   ! In 'declare-allocatable-array_descriptor-1-runtime.f90', this does "verify
 147   ! that device-to-host copy doesn't touch the host-side array descriptor,
 148   ! doesn't copy out the device-side (still initial) array descriptor (but it
 149   ! does copy the array data)".  This is here not applicable anymore, as we've
 150   ! already gotten the actual array descriptor installed.  Thus now verify that
 151   ! it does copy the array data.
 152   call acc_update_self (b)
 153   call verify_n1_allocated
 154
 155   do i = n1_lb, n1_ub
 156      if (b(i) /= i - 1) error stop
 157      b(i) = b(i) + 2
 158   end do
 159
 160   ! The same using the OpenACC 'update' directive.
 161
 162   !$acc update device (b) self (id1_1)
 163   ! We do have 'GOMP_MAP_TO_PSET' here:
 164   ! { 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 } }
 165   ! { 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 } }
 166   ! ..., but it's silently skipped in 'GOACC_update'.
 167   !$acc serial
 168   call verify_n1_allocated
 169   !$acc end serial
 170
 171   b = 41
 172
 173   !$acc parallel
 174   call verify_n1_values (1)
 175   !$acc end parallel
 176
 177   !$acc parallel copy (b)
 178   call verify_n1_values (1)
 179   !$acc end parallel
 180
 181   !$acc update self (b) self (id1_2)
 182   ! We do have 'GOMP_MAP_TO_PSET' here:
 183   ! { 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 } }
 184   ! { 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 } }
 185   ! ..., but it's silently skipped in 'GOACC_update'.
 186   call verify_n1_allocated
 187
 188   do i = n1_lb, n1_ub
 189      if (b(i) /= i + 1) error stop
 190      b(i) = b(i) + 2
 191   end do
 192
 193   ! Now test that (potentially re-)installing the actual array descriptor is a
 194   ! no-op, via a data clause for 'b' (explicit or implicit): must get a
 195   ! 'GOMP_MAP_TO_PSET'.
 196   !$acc serial present (b) copyin (id1_2)
 197   call verify_n1_allocated
 198   !TODO Use of 'b':
 199   id1_2 = ubound (b, 1)
 200   !$acc end serial
 201   ! { 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 } }
 202   ! { 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 } }
 203
 204   !$acc parallel copyin (id1_1) ! No data clause for 'b' (explicit or implicit): no 'GOMP_MAP_TO_PSET'.
 205   call verify_n1_values (1)
 206   id1_1 = 0
 207   !$acc end parallel
 208   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma acc parallel map\(to:id1_1\)$} 1 original } }
 209   ! { dg-final { scan-tree-dump-times {(?n)^ *#pragma omp target oacc_parallel map\(to:id1_1 \[len: [0-9]+\]\)$} 1 gimple } }
 210
 211   !$acc parallel copy (b) copyin (id1_2)
 212   ! As already present, 'copy (b)' doesn't copy; addend is still '1'.
 213   call verify_n1_values (1)
 214   id1_2 = 0
 215   !$acc end parallel
 216   ! { 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 } }
 217   !TODO ..., but without an actual use of 'b', the gimplifier removes the
 218   !TODO 'GOMP_MAP_TO_PSET':
 219   ! { 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 } }
 220
 221   call verify_n1_allocated
 222   if (.not.acc_is_present (b)) error stop
 223
 224   !$acc exit data delete (b)
 225   if (.not.allocated (b)) error stop
 226   if (acc_is_present (b)) error stop
 227   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 228   ! as "allocated":
 229   !$acc serial
 230   call verify_n1_allocated
 231   !$acc end serial
 232
 233   deallocate (b)
 234   call verify_n1_deallocated (.false.)
 235   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 236   ! as "allocated":
 237   !$acc serial
 238   call verify_n1_allocated
 239   !$acc end serial
 240
 241   ! Now try to install the actual array descriptor, via a data clause for 'b'
 242   ! (explicit or implicit): must get a 'GOMP_MAP_TO_PSET', which then in
 243   ! 'gomp_map_vars_internal' is handled as 'declare target', but because of
 244   ! '*(void **) hostaddrs[i] == NULL', we've got 'has_always_ptrset == false',
 245   ! 'always_to_cnt == 0', and therefore 'gomp_map_vars_existing' doesn't update
 246   ! the 'GOMP_MAP_TO_PSET'.
 247   ! The device-side array descriptor doesn't get updated, so 'b' still appears
 248   ! as "allocated":
 249   !TODO Why does 'present (b)' still work here?
 250   !$acc serial present (b) copyout (id1_2)
 251   call verify_n1_deallocated (.true.)
 252   !TODO Use of 'b'.
 253   id1_2 = ubound (b, 1)
 254   !$acc end serial
 255   ! { 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 } }
 256   ! { 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 } }
 257
 258
 259   ! Restart the procedure, with different array dimensions.
 260
 261   allocate (b(n2_lb:n2_ub))
 262   call verify_n2_allocated
 263   if (acc_is_present (b)) error stop
 264   !$acc enter data create (b)
 265   if (.not.acc_is_present (b)) error stop
 266   ! ..., and got the actual array descriptor installed:
 267   !$acc serial
 268   call verify_n2_allocated
 269   !$acc end serial
 270
 271   do i = n2_lb, n2_ub
 272      b(i) = i + 20
 273   end do
 274
 275   call acc_update_device (b)
 276   !$acc serial
 277   call verify_n2_allocated
 278   !$acc end serial
 279
 280   b = -40
 281
 282   !$acc parallel
 283   call verify_n2_values (20)
 284   !$acc end parallel
 285
 286   !$acc parallel copy (b)
 287   call verify_n2_values (20)
 288   !$acc end parallel
 289
 290   call acc_update_self (b)
 291   call verify_n2_allocated
 292
 293   do i = n2_lb, n2_ub
 294      if (b(i) /= i + 20) error stop
 295      b(i) = b(i) - 40
 296   end do
 297
 298   !$acc update device (b)
 299   !$acc serial
 300   call verify_n2_allocated
 301   !$acc end serial
 302
 303   b = -41
 304
 305   !$acc parallel
 306   call verify_n2_values (-20)
 307   !$acc end parallel
 308
 309   !$acc parallel copy (b)
 310   call verify_n2_values (-20)
 311   !$acc end parallel
 312
 313   !$acc update self (b)
 314   call verify_n2_allocated
 315
 316   do i = n2_lb, n2_ub
 317      if (b(i) /= i - 20) error stop
 318      b(i) = b(i) + 10
 319   end do
 320
 321   !$acc serial present (b) copy (id1_2)
 322   call verify_n2_allocated
 323   !TODO Use of 'b':
 324   id1_2 = ubound (b, 1)
 325   !$acc end serial
 326
 327   !$acc parallel
 328   call verify_n2_values (-20)
 329   !$acc end parallel
 330
 331   !$acc parallel copy (b)
 332   call verify_n2_values (-20)
 333   !$acc end parallel
 334
 335   call verify_n2_allocated
 336   if (.not.acc_is_present (b)) error stop
 337
 338   !$acc exit data delete (b)
 339   if (.not.allocated (b)) error stop
 340   if (acc_is_present (b)) error stop
 341   !$acc serial
 342   call verify_n2_allocated
 343   !$acc end serial
 344
 345   deallocate (b)
 346   call verify_n2_deallocated (.false.)
 347   !$acc serial
 348   call verify_n2_allocated
 349   !$acc end serial
 350
 351   !$acc serial present (b) copy (id1_2)
 352   call verify_n2_deallocated (.true.)
 353   !TODO Use of 'b':
 354   id1_2 = ubound (b, 1)
 355   !$acc end serial
 356
 357 end program test
 358
 359
 360 subroutine verify_initial
 361   use vars
 362   implicit none
 363   !$acc routine seq
 364
 365   if (allocated (b)) error stop "verify_initial allocated"
 366   if (any (lbound (b) /= [0])) error stop "verify_initial lbound"
 367   if (any (ubound (b) /= [0])) error stop "verify_initial ubound"
 368 end subroutine verify_initial
 369
 370 subroutine verify_n1_allocated
 371   use vars
 372   implicit none
 373   !$acc routine seq
 374
 375   if (.not.allocated (b)) error stop "verify_n1_allocated allocated"
 376   if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_allocated lbound"
 377   if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_allocated ubound"
 378 end subroutine verify_n1_allocated
 379
 380 subroutine verify_n1_values (addend)
 381   use vars
 382   implicit none
 383   !$acc routine gang
 384   integer, value :: addend
 385   integer :: i
 386
 387   !$acc loop
 388   do i = n1_lb, n1_ub
 389      if (b(i) /= i + addend) error stop
 390   end do
 391 end subroutine verify_n1_values
 392
 393 subroutine verify_n1_deallocated (expect_allocated)
 394   use vars
 395   implicit none
 396   !$acc routine seq
 397   logical, value :: expect_allocated
 398
 399   if (allocated(b) .neqv. expect_allocated) error stop "verify_n1_deallocated allocated"
 400   ! Apparently 'deallocate'ing doesn't unset the bounds.
 401   if (any (lbound (b) /= [n1_lb])) error stop "verify_n1_deallocated lbound"
 402   if (any (ubound (b) /= [n1_ub])) error stop "verify_n1_deallocated ubound"
 403 end subroutine verify_n1_deallocated
 404
 405 subroutine verify_n2_allocated
 406   use vars
 407   implicit none
 408   !$acc routine seq
 409
 410   if (.not.allocated(b)) error stop "verify_n2_allocated allocated"
 411   if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_allocated lbound"
 412   if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_allocated ubound"
 413 end subroutine verify_n2_allocated
 414
 415 subroutine verify_n2_values (addend)
 416   use vars
 417   implicit none
 418   !$acc routine gang
 419   integer, value :: addend
 420   integer :: i
 421
 422   !$acc loop
 423   do i = n2_lb, n2_ub
 424      if (b(i) /= i + addend) error stop
 425   end do
 426 end subroutine verify_n2_values
 427
 428 subroutine verify_n2_deallocated (expect_allocated)
 429   use vars
 430   implicit none
 431   !$acc routine seq
 432   logical, value :: expect_allocated
 433
 434   if (allocated(b) .neqv. expect_allocated) error stop "verify_n2_deallocated allocated"
 435   ! Apparently 'deallocate'ing doesn't unset the bounds.
 436   if (any (lbound (b) /= [n2_lb])) error stop "verify_n2_deallocated lbound"
 437   if (any (ubound (b) /= [n2_ub])) error stop "verify_n2_deallocated ubound"
 438 end subroutine verify_n2_deallocated