OpenACC slicing window

Hello,

I try to implement window slicing in a Fortran application, in order to hide PCIe cost.

I have a four dimensional array, Q, and I compute an other four dimensional array F on a GPU using openACC.
I cut the array Q in eight parts. I send a slice of Q to the device, I compute a kernel on the device and I get back on the host the newly computed slice of F. I try to use the ACC_MAP_DATA but I may miss something.
Here is my example


      PROGRAM PIPELINE
!
      IMPLICIT NONE
!
      INTEGER, PARAMETER  ::  rp = kind(1.0D0)
!
! Sizes of the problem
      INTEGER, PARAMETER  ::  NX = 100
      INTEGER, PARAMETER  ::  NY = 150
      INTEGER, PARAMETER  ::  NZ = 128
!
      REAL(rp), PARAMETER ::  one  = 1.0_rp
      REAL(rp), PARAMETER ::  zero = 0.0_rp
      REAL(rp), PARAMETER ::  half = 0.5_rp
!
! Number of slices and temporary indexes
      INTEGER, PARAMETER  :: nbslice_z = 8
      INTEGER  ::  kMinz, kMaxz
!
! Global dimensions for computing loops
      INTEGER  ::  Imin, Imax, Jmin, Jmax, Kmin, Kmax
!
! Needed for ghost cells of Q
      INTEGER, PARAMETER  :: i1add  = 3
      INTEGER, PARAMETER  :: j1add  = 3
      INTEGER, PARAMETER  :: k1add  = 3
      INTEGER, PARAMETER  :: i2add  = 3
      INTEGER, PARAMETER  :: j2add  = 3
      INTEGER, PARAMETER  :: k2add  = 3
!
! temporary scalar variables like lopp indexes
      INTEGER  ::  i, j, k, L, ik, length
      INTEGER  ::  icode
!
! Arrays of data
      REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE, TARGET     ::  Q, F
      REAL(rp), POINTER,     CONTIGUOUS ::  PQ(:,:,:,:)
      REAL(rp), POINTER,     CONTIGUOUS ::  PF(:,:,:,:)
      INTEGER,  DIMENSION(:),       ALLOCATABLE  ::  kBeg, kEnd
!
!!!
!
! Debut du programme
!
!!!
!
! Initialization for loop indexes
      kMin = 1 ; kMax = Nz
      jMin = 1 ; jMax = Ny
      iMin = 1 ; iMax = Nx
      WRITE (6,'(A,6I8)') 'iMin, jMin, kMin, iMax, jMax, kMax : ', iMin, jMin, kMin, iMax, jMax, kMax
!
! Dynamic Allocation (the final program will use ghost cells)
      ALLOCATE (Q (5, Imin-i1add:Imax+i2add, Jmin-j1add:Jmax+j2add, Kmin-k1add:Kmax+k2add),
     &          F (5, Imin-i1add:Imax+i2add, Jmin-j1add:Jmax+j2add, Kmin-K1add:Kmax+k2add) )
      WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
      WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
!
! Initialization of the data
      DO k = Kmin-k1add, Kmax+k2add
         DO j = Jmin - j1add, Jmax + j2add
            DO i = Imin - i1add, Imax + i2add
               Q(1,i,j,k) = one
               Q(2,i,j,k) = one
               Q(3,i,j,k) = one
               Q(4,i,j,k) = one
               Q(5,i,j,k) = one
!
               F(1,i,j,k) = zero
               F(2,i,j,k) = zero
               F(3,i,j,k) = zero
               F(4,i,j,k) = zero
               F(5,i,j,k) = zero
            END DO
         END DO
      END DO
!
! Output of Q
      WRITE (6,'(A)') 'Part of array Q :'
      DO L = 1, 5
         WRITE (6,'(4I4,ES22.15)') L, L, L, L, Q(L,L,L,L)
      END DO
      CALL Flush (6)
!
! Filling of the slice cutting array
! example : 8 slices
      ALLOCATE (kBeg(nbslice_z), kEnd(nbslice_z) )
      DO ik = 0, nbslice_z-1
         kBeg(1+ik) = 1 +  (ik   *Nz) / nbslice_z
         kEnd(1+ik) =    ( (ik+1)*Nz) / nbslice_z
      END DO
!
! Output of the slicing
      WRITE (6,'(A,I4)') 'nbslice_z = ', nbslice_z
      DO k = 1, nbslice_z
         WRITE (6,'(3(A,I4) )') 'slice ', k, ' : ', kBeg(k), ' to ', kEnd(k)
         CALL Flush (6)
      END DO
!
!!!
! Loop on the slices
!!!
!
      DO ik = 1, nbslice_z
!
! Current slice
         kminz = kBeg(ik)
         kmaxz = kEnd(ik)
         WRITE (6,'(A,3I4)') 'Beginning slice ', ik, kminz, kmaxz
         CALL Flush (6)
!
         PQ => Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add)
         PF => F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz)
         WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
         WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
         WRITE (6,'(A,8I4)') 'PQ : ', LBOUND(PQ), UBOUND(PQ)
         WRITE (6,'(A,8I4)') 'PF : ', LBOUND(PF), UBOUND(PF)
         CALL Flush (6)

!$ACC ENTER DATA CREATE (PF) COPYIN(PQ)
!
         WRITE (6,'(A)') 'Association on the device'
         CALL FLUSH(6)
!
! Association on the device
         length = 8 * SIZE(Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add) )
         WRITE (6,'(A,I9)') 'length(Q) : ', length
         CALL FLUSH(6)
         CALL ACC_MAP_DATA (Q(1,1-i1add, 1-j1add, kminz-k1add), PQ(1,1-i1add, 1-j1add, kminz-k1add), length)
!
         length = 8 * SIZE(F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz) )
         WRITE (6,'(A,I9)') 'length(F) : ', length
         CALL FLUSH(6)
         CALL ACC_MAP_DATA (F(1,1-i1add, 1-j1add, kminz), PF(1,1-i1add, 1-j1add, kminz), length)
!
         WRITE (6,'(A)') 'Computing on the device'
         CALL FLUSH(6)

! Computing of F's slice
!$ACC PARALLEL LOOP COLLAPSE(4) PRIVATE (i,j,k,L) PRESENT(PF,PQ)
         DO k = kminz, kmaxz
            DO j = 1, Ny+j2add+j1add
               DO i = 1, Nx+i2add+i1add
                  DO L = 1, 5
                     PF(L,i,j,k) = PQ(L,i,j,k) * half + one
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!
! Device To Host of PF
!$ACC UPDATE SELF (PF)
!
! End of association
         CALL ACC_UNMAP_DATA (Q)
         CALL ACC_UNMAP_DATA (F)
!
!$ACC EXIT DATA
!
         WRITE (6,'(A,3I4)') 'End of slice ', ik, kminz, kmaxz
         CALL Flush (6)
      END DO
!
!!!
!
! Output of F
      WRITE (6,'(A)') 'Output of F :'
      DO L = 1, 5
         WRITE (6,'(4I4,2ES22.15)') L, L, L, L, F(L,L,L,L)
      END DO
      CALL Flush (6)
!
! Deallocation
      DEALLOCATE (Q, F)
!
!!!
!
!
!!!
!
      STOP
      END PROGRAM PIPELINE

I compile this example with

make
pgf90 -c -Minfo=all -O0 -g -traceback -Mbounds -Mfixed -Mextend ./src/main.F90 -o ./obj_O0/main.o
pipeline:
    144, Conflict or overlap between pf and pq
pgf90 -O0 ./obj_O0/main.o  -o ./run/a_acc.out

When I run the code the get the following output :

./run/a_acc.out 
iMin, jMin, kMin, iMax, jMax, kMax :        1       1       1     100     150     128
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
Part of array Q :
   1   1   1   1 1.000000000000000E+00
   2   2   2   2 1.000000000000000E+00
   3   3   3   3 1.000000000000000E+00
   4   4   4   4 1.000000000000000E+00
   5   5   5   5 1.000000000000000E+00
nbslice_z =    8
slice    1 :    1 to   16
slice    2 :   17 to   32
slice    3 :   33 to   48
slice    4 :   49 to   64
slice    5 :   65 to   80
slice    6 :   81 to   96
slice    7 :   97 to  112
slice    8 :  113 to  128
Beginning slice    1   1  16
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
0: Subscript out of range for array pq (./src/main.F90: 128)
    subscript=-2, lower bound=1, upper bound=106, dimension=2

If I look at the values of LBOUND and UBOUND for PF and PQ, I do not get those of F and Q. How can I preserve them ? I’d like to have the same indexes on the host and and the device.

Is the use ACC_MAP_DATA correct ?

The next step will be the use of multiple queues and asynchronous execution.

Thank you for your comments.

Regards,
Guy.

Hi Guy,

Is the use ACC_MAP_DATA correct?

No, sorry.

“acc_map_data” is a C/C++ routine used to map a host pointer to a device pointer. I use it to create a pool of device memory then map and unmap multiple host pointers.

For example, “devMem” is a device only memory block which each of column of “A” gets map to when that column is used on the device.

    A= (double**) malloc(BLOCKS*sizeof(double*));
    for (i=0; i < BLOCKS; ++i) {
       A[i] = (double*) malloc(nbytes);
    }
    devMem= (char*) acc_malloc(nbytes);

    for (i=0; i < BLOCKS; ++i) {
        tmpA = A[i];
        acc_map_data(tmpA,devMem,nbytes);
        #pragma acc parallel loop present(tmpA)
        for (j=0; j < N; ++j) {
             tmpA[j] = 2.5 + (double) ((j*N)+i);
        }
        #pragma acc update host(tmpA[0:N])
        acc_unmap_data(A[i]);
    }

I don’t have a Fortran example off-hand but conceptionally you could use a CUDA Fortran device array in place of “devMem” and then map a host sub-array to the device memory.

The next step will be the use of multiple queues and asynchronous execution.

You wouldn’t want to use acc_map_data for this since it creates a one-to-one mapping from host to device data. If you tried to do this asynchronously then you’d be overwriting this map and end up with multiple host threads all using the same device data.

Instead, you want to copy all of “Q” and “F” over to the device before the slices loop, associate “PQ” and “PF” to the slice (as you do now), add the “async” clause to the parallel compute regions, then copy back “F” after the slices loop.

Note that your code has a bug in it in that you should be looping over the “z” dimension from 1 to “kmaxz-kminz”, not “kminz” to “kmaxz”. Otherwise on the second slice you’ll be indexing into the 17th element when PQ and PF only have 16 elements.

Double check my corrections, but I think this is what you intended:

% cat test.F

       PROGRAM PIPELINE
!
       IMPLICIT NONE
!
       INTEGER, PARAMETER  ::  rp = kind(1.0D0)
!
! Sizes of the problem
       INTEGER, PARAMETER  ::  NX = 100
       INTEGER, PARAMETER  ::  NY = 150
       INTEGER, PARAMETER  ::  NZ = 128
!
       REAL(rp), PARAMETER ::  one  = 1.0_rp
       REAL(rp), PARAMETER ::  zero = 0.0_rp
       REAL(rp), PARAMETER ::  half = 0.5_rp
!
! Number of slices and temporary indexes
       INTEGER, PARAMETER  :: nbslice_z = 8
       INTEGER  ::  kMinz, kMaxz
!
! Global dimensions for computing loops
       INTEGER  ::  Imin, Imax, Jmin, Jmax, Kmin, Kmax
!
! Needed for ghost cells of Q
       INTEGER, PARAMETER  :: i1add  = 3
       INTEGER, PARAMETER  :: j1add  = 3
       INTEGER, PARAMETER  :: k1add  = 3
       INTEGER, PARAMETER  :: i2add  = 3
       INTEGER, PARAMETER  :: j2add  = 3
       INTEGER, PARAMETER  :: k2add  = 3
!
! temporary scalar variables like lopp indexes
       INTEGER  ::  i, j, k, L, ik, length
       INTEGER  ::  icode
!
! Arrays of data
       REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE, TARGET     ::  Q, F
       REAL(rp), POINTER,     CONTIGUOUS ::  PQ(:,:,:,:)
       REAL(rp), POINTER,     CONTIGUOUS ::  PF(:,:,:,:)
       INTEGER,  DIMENSION(:),       ALLOCATABLE  ::  kBeg, kEnd
!
!!!
!
! Debut du programme
!
!!!
!
! Initialization for loop indexes
       kMin = 1 ; kMax = Nz
       jMin = 1 ; jMax = Ny
       iMin = 1 ; iMax = Nx
       WRITE (6,'(A,6I8)') 'iMin, jMin, kMin, iMax, jMax, kMax : ', iMin, jMin, kMin, iMax, jMax, kMax
!
! Dynamic Allocation (the final program will use ghost cells)
       ALLOCATE (Q (5, Imin-i1add:Imax+i2add,Jmin-j1add:Jmax+j2add,Kmin-k1add:Kmax+k2add),
     &       F (5, Imin-i1add:Imax+i2add, Jmin-j1add:Jmax+j2add, Kmin-K1add:Kmax+k2add) )
       WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
       WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
!
! Initialization of the data
       DO k = Kmin-k1add, Kmax+k2add
          DO j = Jmin - j1add, Jmax + j2add
             DO i = Imin - i1add, Imax + i2add
                Q(1,i,j,k) = one
                Q(2,i,j,k) = one
                Q(3,i,j,k) = one
                Q(4,i,j,k) = one
                Q(5,i,j,k) = one
!
                F(1,i,j,k) = zero
                F(2,i,j,k) = zero
                F(3,i,j,k) = zero
                F(4,i,j,k) = zero
                F(5,i,j,k) = zero
             END DO
          END DO
       END DO
!
! Output of Q
       WRITE (6,'(A)') 'Part of array Q :'
       DO L = 1, 5
          WRITE (6,'(4I4,ES22.15)') L, L, L, L, Q(L,L,L,L)
       END DO
       CALL Flush (6)
!
! Filling of the slice cutting array
! example : 8 slices
       ALLOCATE (kBeg(nbslice_z), kEnd(nbslice_z) )
       DO ik = 0, nbslice_z-1
          kBeg(1+ik) = 1 +  (ik   *Nz) / nbslice_z
          kEnd(1+ik) =    ( (ik+1)*Nz) / nbslice_z
       END DO
!
! Output of the slicing
       WRITE (6,'(A,I4)') 'nbslice_z = ', nbslice_z
       DO k = 1, nbslice_z
          WRITE (6,'(3(A,I4) )') 'slice ', k, ' : ', kBeg(k), ' to ', kEnd(k)
          CALL Flush (6)
       END DO

!$ACC DATA COPYIN(Q) COPYOUT(F)
!!!
! Loop on the slices
!!!
!
       DO ik = 1, nbslice_z
!
! Current slice
          kminz = kBeg(ik)
          kmaxz = kEnd(ik)
          WRITE (6,'(A,3I4)') 'Beginning slice ', ik, kminz, kmaxz
          CALL Flush (6)
!
          PQ => Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add)
          PF => F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz)
          WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
          WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
          WRITE (6,'(A,8I4)') 'PQ : ', LBOUND(PQ), UBOUND(PQ)
          WRITE (6,'(A,8I4)') 'PF : ', LBOUND(PF), UBOUND(PF)
          CALL Flush (6)

          WRITE (6,'(A)') 'Association on the device'
          CALL FLUSH(6)
!
! Association on the device
          length = 8 * SIZE(Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add) )
          WRITE (6,'(A,I9)') 'length(Q) : ', length
          CALL FLUSH(6)
!
          length = 8 * SIZE(F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz) )
          WRITE (6,'(A,I9)') 'length(F) : ', length
!
          WRITE (6,'(A)') 'Computing on the device'
          CALL FLUSH(6)

! Computing of F's slice
!$ACC PARALLEL LOOP COLLAPSE(4) PRESENT(PF,PQ) ASYNC(ik)
          !DO k = kminz, kmaxz
          DO k = 1,kmaxz-kminz
             DO j = 1, Ny+j2add+j1add
                DO i = 1, Nx+i2add+i1add
                   DO L = 1, 5
                      PF(L,i,j,k) = PQ(L,i,j,k) * half + one
                   END DO
                END DO
             END DO
          END DO
!$ACC END PARALLEL LOOP
!
! End of association
          WRITE (6,'(A,3I4)') 'End of slice ', ik, kminz, kmaxz
          CALL Flush (6)
       END DO
!
!!!
!$ACC WAIT
!$ACC END DATA
! Output of F
       WRITE (6,'(A)') 'Output of F :'
       DO L = 1, 5
          WRITE (6,'(4I4,2ES22.15)') L, L, L, L, F(L,L,L,L)
       END DO
       CALL Flush (6)
!
! Deallocation
       DEALLOCATE (Q, F)
!
!!
!
!
!!!
!
       STOP
       END PROGRAM PIPELINE

% pgf90 test.F -Mextend -acc -Minfo=accel
pipeline:
    101, Generating copyout(f(:,:,:,:))
         Generating copyin(q(:,:,:,:))
    137, Generating present(pf(:,:,:,:),pq(:,:,:,:))
         Accelerator kernel generated
         Generating Tesla code
        139, !$acc loop gang, vector(128) collapse(4) ! blockidx%x threadidx%x
        140,   ! blockidx%x threadidx%x collapsed
        141,   ! blockidx%x threadidx%x collapsed
        142,   ! blockidx%x threadidx%x collapsed

% a.out
iMin, jMin, kMin, iMax, jMax, kMax :        1       1       1     100     150     128
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
Part of array Q :
   1   1   1   1 1.000000000000000E+00
   2   2   2   2 1.000000000000000E+00
   3   3   3   3 1.000000000000000E+00
   4   4   4   4 1.000000000000000E+00
   5   5   5   5 1.000000000000000E+00
nbslice_z =    8
slice    1 :    1 to   16
slice    2 :   17 to   32
slice    3 :   33 to   48
slice    4 :   49 to   64
slice    5 :   65 to   80
slice    6 :   81 to   96
slice    7 :   97 to  112
slice    8 :  113 to  128
Beginning slice    1   1  16
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    1   1  16
Beginning slice    2  17  32
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    2  17  32
Beginning slice    3  33  48
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    3  33  48
Beginning slice    4  49  64
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    4  49  64
Beginning slice    5  65  80
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    5  65  80
Beginning slice    6  81  96
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    6  81  96
Beginning slice    7  97 112
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    7  97 112
Beginning slice    8 113 128
Q  :    1  -2  -2  -2   5 103 153 131
F  :    1  -2  -2  -2   5 103 153 131
PQ :    1   1   1   1   5 106 156  22
PF :    1   1   1   1   5 106 156  16
Association on the device
length(Q) :  14551680
length(F) :  10583040
Computing on the device
End of slice    8 113 128
Output of F :
   1   1   1   1 1.500000000000000E+00
   2   2   2   2 1.500000000000000E+00
   3   3   3   3 1.500000000000000E+00
   4   4   4   4 1.500000000000000E+00
   5   5   5   5 1.500000000000000E+00
Warning: ieee_inexact is signaling
FORTRAN STOP

A couple of other small things. No need to privatize your loop index variables. That’s done by default. Also, if you do use unstructured data regions, i.e. “enter data” and “exit data” (Note that I switch the code to use a structured region), be sure to “delete” your data in the exit directive. A structured data region will do this for you since it’s defined to delete all the data associated with the data region at the end. However, an unstructured data region is more like allocate and deallocate where it’s up to the programmer when to create and delete the device data. Without the “delete”, the data would remain on the device.

Hope this helps,
Mat

Hello Mat,

thank you for your long post.
I try to follow your advices. First I have to make somme comments.

Instead, you want to copy all of “Q” and “F” over to the device before the slices loop, associate “PQ” and “PF” to the slice (as you do now), add the “async” clause to the parallel compute regions, then copy back “F” after the slices loop.

I do not want to copy all Q and F arrays before the loop begins. I’d like to work with slices of Q and F to get some overlap. So take a look at this,

1.    H2D   F(Q)   D2H
2.          H2D    F(Q)   D2H
3.                 H2D    F(Q)   D2H
4.                        H2D    F(Q)   D2H
5.                               H2D    F(Q)   D2H
6.                                      H2D    F(Q)   D2H
7.                                             H2D    F(Q)   D2H
8.                                                    H2D    F(Q)   D2H

Each line is the processing of a slice, several steps :

  • transfer a slice of Q (H2D : host to device)
  • compute a slice of F with a slice of Q, so F(Q)
  • transfer a slice of F (D2H : device to host)

Each column represents what can be done at the same time, because the three previous steps are independent if applied on different data.

When the transfer of slice 1 of Q is finished and the computing of slice 1 of F begins, one can proceed to the transfert of slice 2 of Q, so at the same time as the computing of slice 1 of F.
When the computing of slice 1 of F is finished, one can transfer it back to the host. At the same time the computing of slice 2 of F can happen and the transfer from host to device of slice 3 of Q can happen too.
And so on …

Of course, one needs several working buffers on the device for slices of F and for slices of Q, and one has to swap between them, in order to avoid overwriting. And this is why I talked about asynchronous execution : three different steps, for three different slices, may occur at the same time.

Following your advice, I try to use a CUDA Fortran device array for a slice but I get a core dumped when I try to map it (line61).

     CALL ACC_MAP_DATA (Q(1,1,1,1), PQ, length)



use a CUDA Fortran device array in place of “devMem” and then map a host sub-array to the device memory.

Is the following what you think about ?

      PROGRAM SMALL
!
#ifdef _CUDA
      USE cudafor
#endif
!
      IMPLICIT NONE
!
      INTEGER, PARAMETER  ::  rp = kind(1.0D0)
!
! temporary scalar variables like lopp indexes
      INTEGER  ::  i, j, k, L, length
      INTEGER  ::  icode
!
! Arrays of data
      REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE  ::  Q , F
      REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE  ::  PQ, PF
#ifdef _CUDA
      ATTRIBUTES(DEVICE)  ::  PQ, PF
#endif
!
! Dynamic Allocation
      ALLOCATE (Q (5, 100, 100, 100), F (5, 100, 100, 100), STAT=icode)
!
! Initialization of the data
      DO k = 1, 100
         DO j = 1, 100
            DO i = 1, 100
               DO L  1, 5
                  Q(L,i,j,k) = one
               END DO
            END DO
         END DO
      END DO
!
! Allocation on device
      ALLOCATE (PQ(5,100,100,10), PF(5,100,100,10), STAT=icode)
!
      WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
      WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
      WRITE (6,'(A,8I4)') 'PQ : ', LBOUND(PQ), UBOUND(PQ)
      WRITE (6,'(A,8I4)') 'PF : ', LBOUND(PF), UBOUND(PF)
!
! Data region on the device
!$ACC ENTER DATA CREATE (PF, PQ)
!
      WRITE (6,'(A)') 'Association of a slice on the device'
!
! Association on the device
      length = 8 * SIZE(Q(1:5,1:100,1:100,1:10) )
      WRITE (6,'(A,I9)') 'length(Q) : ', length
      CALL FLUSH(6)
      CALL ACC_MAP_DATA (Q(1,1,1,1), PQ, length)
!
      length = 8 * SIZE(F(1:5,1:100,1:100,1:10) )
      WRITE (6,'(A,I9)') 'length(F) : ', length
      CALL FLUSH(6)
      CALL ACC_MAP_DATA (F(1,1,1,1), PF, length)
!
      WRITE (6,'(A)') 'Computing on the device'
      CALL FLUSH(6)

! Computing of F
!$ACC PARALLEL LOOP COLLAPSE(4) DEVICEPTR(PF,PQ)
      DO k = 1, 10
         DO j = 1, 100
            DO i = 1, 100
               DO L = 1, 5
                  PF(L,i,j,k) = PQ(L,i,j,k) * half + one
               END DO
            END DO
         END DO
      END DO
!$ACC END PARALLEL LOOP
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (PF)
!
! End of association
      CALL ACC_UNMAP_DATA (Q)
      CALL ACC_UNMAP_DATA (F)
!
! Free ressources on device
!$ACC EXIT DATA DELETE (PQ, PF)
!
! Output of F
      WRITE (6,'(A)') 'Output of F :'
      DO L = 1, 5
         WRITE (6,'(4I4,2ES22.15)') L, L, L, L, F(L,L,L,L)
      END DO
      CALL Flush (6)
!
! Free ressources on host
      DEALLOCATE (Q, F)
!
      STOP
      END PROGRAM SMALL

Regards,

Hi Guy,

ACC_MAP_DATA is not the correct approach for what you want for pipelining. “acc_map_data” is used when you want to reuse the same device memory, mapping it to different host data during the execution of the program. This will cause issues with mapping multiple slices of Q and F to the same PQ and PF.

Instead, create F on the device, then use “update device(PF) async(ik)” to have each async queue copy the array section.

Note that due to implementation issues, copying back to the host will cause the queues to block waiting for the copy to complete. This is way I’m having “Q” copyout at the end.

For example:

% cat test.F

       PROGRAM PIPELINE
!
       IMPLICIT NONE
!
       INTEGER, PARAMETER  ::  rp = kind(1.0D0)
!
! Sizes of the problem
       INTEGER, PARAMETER  ::  NX = 100
       INTEGER, PARAMETER  ::  NY = 150
       INTEGER, PARAMETER  ::  NZ = 128
!
       REAL(rp), PARAMETER ::  one  = 1.0_rp
       REAL(rp), PARAMETER ::  zero = 0.0_rp
       REAL(rp), PARAMETER ::  half = 0.5_rp
!
! Number of slices and temporary indexes
       INTEGER, PARAMETER  :: nbslice_z = 8
       INTEGER  ::  kMinz, kMaxz
!
! Global dimensions for computing loops
       INTEGER  ::  Imin, Imax, Jmin, Jmax, Kmin, Kmax
!
! Needed for ghost cells of Q
       INTEGER, PARAMETER  :: i1add  = 3
       INTEGER, PARAMETER  :: j1add  = 3
       INTEGER, PARAMETER  :: k1add  = 3
       INTEGER, PARAMETER  :: i2add  = 3
       INTEGER, PARAMETER  :: j2add  = 3
       INTEGER, PARAMETER  :: k2add  = 3
!
! temporary scalar variables like lopp indexes
       INTEGER  ::  i, j, k, L, ik, length
       INTEGER  ::  icode
!
! Arrays of data
       REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE, TARGET     ::  Q, F
       REAL(rp), POINTER,     CONTIGUOUS ::  PQ(:,:,:,:)
       REAL(rp), POINTER,     CONTIGUOUS ::  PF(:,:,:,:)
       INTEGER,  DIMENSION(:),       ALLOCATABLE  ::  kBeg, kEnd
!
!!!
!
! Debut du programme
!
!!!
!
! Initialization for loop indexes
       kMin = 1 ; kMax = Nz
       jMin = 1 ; jMax = Ny
       iMin = 1 ; iMax = Nx
       WRITE (6,'(A,6I8)') 'iMin, jMin, kMin, iMax, jMax, kMax : ', iMin, jMin, kMin, iMax, jMax, kMax
!
! Dynamic Allocation (the final program will use ghost cells)
       ALLOCATE (Q (5, Imin-i1add:Imax+i2add,Jmin-j1add:Jmax+j2add,Kmin-k1add:Kmax+k2add),
     &       F (5, Imin-i1add:Imax+i2add, Jmin-j1add:Jmax+j2add, Kmin-K1add:Kmax+k2add) )
       WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
       WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
!
! Initialization of the data
       DO k = Kmin-k1add, Kmax+k2add
          DO j = Jmin - j1add, Jmax + j2add
             DO i = Imin - i1add, Imax + i2add
                Q(1,i,j,k) = one
                Q(2,i,j,k) = one
                Q(3,i,j,k) = one
                Q(4,i,j,k) = one
                Q(5,i,j,k) = one
!
                F(1,i,j,k) = zero
                F(2,i,j,k) = zero
                F(3,i,j,k) = zero
                F(4,i,j,k) = zero
                F(5,i,j,k) = zero
             END DO
          END DO
       END DO
!
! Output of Q
       WRITE (6,'(A)') 'Part of array Q :'
       DO L = 1, 5
          WRITE (6,'(4I4,ES22.15)') L, L, L, L, Q(L,L,L,L)
       END DO
       CALL Flush (6)
!
! Filling of the slice cutting array
! example : 8 slices
       ALLOCATE (kBeg(nbslice_z), kEnd(nbslice_z) )
       DO ik = 0, nbslice_z-1
          kBeg(1+ik) = 1 +  (ik   *Nz) / nbslice_z
          kEnd(1+ik) =    ( (ik+1)*Nz) / nbslice_z
       END DO
!
! Output of the slicing
       WRITE (6,'(A,I4)') 'nbslice_z = ', nbslice_z
       DO k = 1, nbslice_z
          WRITE (6,'(3(A,I4) )') 'slice ', k, ' : ', kBeg(k), ' to ', kEnd(k)
          CALL Flush (6)
       END DO

!$ACC DATA COPYOUT(F) CREATE(Q)
!!!
! Loop on the slices
!!!
!
       DO ik = 1, nbslice_z
!
! Current slice
          kminz = kBeg(ik)
          kmaxz = kEnd(ik)
          WRITE (6,'(A,3I4)') 'Beginning slice ', ik, kminz, kmaxz
          CALL Flush (6)
!
          PQ => Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add)
          PF => F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz)
          WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
          WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
          WRITE (6,'(A,8I4)') 'PQ : ', LBOUND(PQ), UBOUND(PQ)
          WRITE (6,'(A,8I4)') 'PF : ', LBOUND(PF), UBOUND(PF)
          CALL Flush (6)

          WRITE (6,'(A)') 'Association on the device'
          CALL FLUSH(6)
!
! Association on the device
          length = 8 * SIZE(Q(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz-k1add:kmaxz+k2add) )
          WRITE (6,'(A,I9)') 'length(Q) : ', length
          CALL FLUSH(6)
!
          length = 8 * SIZE(F(1:5,1-i1add:Nx+i2add, 1-j1add:Ny+j2add, kminz:kmaxz) )
          WRITE (6,'(A,I9)') 'length(F) : ', length
!
          WRITE (6,'(A)') 'Computing on the device'
          CALL FLUSH(6)

!$ACC UPDATE DEVICE(PQ) ASYNC(ik)
! Computing of F's slice
!$ACC PARALLEL LOOP COLLAPSE(4) PRESENT(PF,PQ) ASYNC(ik)
          !DO k = kminz, kmaxz
          DO k = 1,kmaxz-kminz
             DO j = 1, Ny+j2add+j1add
                DO i = 1, Nx+i2add+i1add
                   DO L = 1, 5
                      PF(L,i,j,k) = PQ(L,i,j,k) * half + one
                   END DO
                END DO
             END DO
          END DO
!$ACC END PARALLEL LOOP
!
! End of association
          WRITE (6,'(A,3I4)') 'End of slice ', ik, kminz, kmaxz
          CALL Flush (6)
       END DO
!
!!!
!$ACC WAIT

!$ACC END DATA
! Output of F
       WRITE (6,'(A)') 'Output of F :'
       DO L = 1, 5
          WRITE (6,'(4I4,2ES22.15)') L, L, L, L, F(L,L,L,L)
       END DO
       CALL Flush (6)
!
! Deallocation
       DEALLOCATE (Q, F)
!
!!
!
!
!!!
!
       STOP
       END PROGRAM PIPELINE

Also, the reason why you’re getting the seg fault with acc_map_data is because you need to use the OpenACC module in order to get the interface. Otherwise, you’ll link directly to the C version. Also, you need to update “Q” and “F” since the address association is host to device and no need to create “PQ” and “PF” since that’s implied when allocating a CUDA Fortran device array. Here’s the updated code:

 % cat small.F
      PROGRAM SMALL
!
        use openacc
#ifdef _CUDA
       USE cudafor
#endif
!
       IMPLICIT NONE
!
       INTEGER, PARAMETER  ::  rp = kind(1.0D0)
!
! temporary scalar variables like lopp indexes
       INTEGER  ::  i, j, k, L, length
       INTEGER  ::  icode
       real(rp), parameter :: one = 1.0
       real(rp), parameter :: half = 0.5
!
! Arrays of data
       REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE  ::  Q , F
       REAL(rp), DIMENSION(:,:,:,:), ALLOCATABLE  ::  PQ, PF
#ifdef _CUDA
       ATTRIBUTES(DEVICE)  ::  PQ, PF
#endif
!
! Dynamic Allocation
       ALLOCATE (Q (5, 100, 100, 100), F (5, 100, 100, 100), STAT=icode)
!
! Initialization of the data
       DO k = 1, 100
          DO j = 1, 100
             DO i = 1, 100
                DO L =  1, 5
                   Q(L,i,j,k) = one
                END DO
             END DO
          END DO
       END DO
!
! Allocation on device
       ALLOCATE (PQ(5,100,100,10), PF(5,100,100,10), STAT=icode)
!
       WRITE (6,'(A,8I4)') 'Q  : ', LBOUND(Q ), UBOUND(Q )
       WRITE (6,'(A,8I4)') 'F  : ', LBOUND(F ), UBOUND(F )
       WRITE (6,'(A,8I4)') 'PQ : ', LBOUND(PQ), UBOUND(PQ)
       WRITE (6,'(A,8I4)') 'PF : ', LBOUND(PF), UBOUND(PF)
!
! Data region on the device
!
       WRITE (6,'(A)') 'Association of a slice on the device'
!
! Association on the device
       length = 8 * SIZE(Q(1:5,1:100,1:100,1:10) )
       WRITE (6,'(A,I9)') 'length(Q) : ', length
       CALL FLUSH(6)
       CALL ACC_MAP_DATA (Q(1,1,1,1), PQ, length)
!
       length = 8 * SIZE(F(1:5,1:100,1:100,1:10) )
       WRITE (6,'(A,I9)') 'length(F) : ', length
       CALL FLUSH(6)
       CALL ACC_MAP_DATA (F(1,1,1,1), PF, length)
!
       WRITE (6,'(A)') 'Computing on the device'
       CALL FLUSH(6)

!$ACC UPDATE DEVICE (Q(1:5,1:100,1:100,1:10))
! Computing of F
!$ACC PARALLEL LOOP COLLAPSE(4)
       DO k = 1, 10
          DO j = 1, 100
             DO i = 1, 100
                DO L = 1, 5
                   PF(L,i,j,k) = PQ(L,i,j,k) * half + one
                END DO
             END DO
          END DO
       END DO
!$ACC END PARALLEL LOOP
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:100,1:100,1:10))
!
! End of association
       CALL ACC_UNMAP_DATA (Q)
       CALL ACC_UNMAP_DATA (F)
!
! Free ressources on device
!
! Output of F
       WRITE (6,'(A)') 'Output of F :'
       DO L = 1, 5
          WRITE (6,'(4I4,2ES22.15)') L, L, L, L, F(L,L,L,L)
       END DO
       CALL Flush (6)
!
! Free ressources on host
       DEALLOCATE (Q, F)
!
       STOP
       END PROGRAM SMALL

Hello Mat,

thanks to your help, with the right interface for ACC_MAP_DATA, this comes to its end.
So I try to create a pipelining for slicing the computation of F.

1.    H2D   F(Q)   D2H
2.          H2D    F(Q)   D2H
3.                 H2D    F(Q)   D2H
4.                        H2D    F(Q)   D2H
5.                               H2D    F(Q)   D2H
6.                                      H2D    F(Q)   D2H
7.                                             H2D    F(Q)   D2H
8.                                                    H2D    F(Q)   D2H

I remark that I only need two temporary arrays for the slices of Q and two more for the slices of F. I add a dimension to PF and PQ and I use two integer indexes (iold, inew) to access the right part of PF and PQ. It seems to work.
As it is a small example with few computations, nvvp tells me that there is no so much Memcpy/Compute overlap.
As I have to handle the first and the last stages of the pipeline, the source code is a bit longer, I apologize for this.

      PROGRAM SMALL
!
      USE openacc
#ifdef _CUDA
      USE cudafor
#endif
!
      IMPLICIT NONE
!
      INTEGER, PARAMETER  ::  rp = kind(1.0D0)
      INTEGER, PARAMETER  ::  L_SliceZ = 10
      INTEGER, PARAMETER  ::  Nx = 100
      INTEGER, PARAMETER  ::  Ny = 100
      INTEGER, PARAMETER  ::  Nz = 200
!
! temporary scalar variables like lopp indexes
      INTEGER  ::  i, j, k, L, length, ik
      REAL(rp)  ::  checksum
!
      INTEGER, DIMENSION(0:1)  ::  kBeg, kEnd
      INTEGER  ::  inew, iold, itmp, NbSliceZ
!
! Arrays of data
      REAL(rp), DIMENSION(:,:,:,:),   ALLOCATABLE  ::  Q , F , FF
      REAL(rp), DIMENSION(:,:,:,:,:), POINTER      ::  PQ, PF
#ifdef _CUDA
      ATTRIBUTES(DEVICE)  ::  PQ, PF
#endif
!
! Dynamic Allocation
      ALLOCATE (Q (5, Nx, Ny, Nz), F (5, Nx, Ny, Nz), FF (5, Nx, Ny, Nz) )
!
! Initialization of the data
      DO k = 1, Nz
         DO j = 1, Ny
            DO i = 1, Nx
               DO L =  1, 5
                  CALL RANDOM_NUMBER ( Q(L,i,j,k) )
                  FF(L,i,j,k) = 0.5_rp * Q(L,i,j,k) + 1.0_rp
               END DO
            END DO
         END DO
      END DO
!
! Allocation on device
      NbSliceZ = Nz / L_SliceZ
!
      ALLOCATE (PQ(5,Nx,Ny,L_SliceZ,0:1), PF(5,Nx,Ny,L_SliceZ,0:1) )
!
      inew = 1
      iold = 0
!
!========================
! Pre-processing slice #1
!========================
      ik = 1
      kBeg(iold) = (ik-1) * L_SliceZ + 1
      kEnd(iold) =  ik    * L_SliceZ
!
! Association on the device
      length = 8 * SIZE(Q(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
      CALL ACC_MAP_DATA (Q(1,1,1,kBeg(iold) ), PQ(1,1,1,1,iold), length)
!
      length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
      CALL ACC_MAP_DATA (F(1,1,1,kBeg(iold) ), PF(1,1,1,1,iold), length)
!
! Host To Device for slice #1 of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
!
!========================
! Pre-processing slice #2
!========================
      ik = 2
      kBeg(inew) = (ik-1) * L_SliceZ + 1
      kEnd(inew) =  ik    * L_SliceZ
!
! Association on the device
      length = 8 * SIZE(Q(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
      CALL ACC_MAP_DATA (Q(1,1,1,kBeg(inew) ), PQ(1,1,1,1,inew), length)
!
      length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
      CALL ACC_MAP_DATA (F(1,1,1,kBeg(inew) ), PF(1,1,1,1,inew), length)
!
! Host 2 Device for a slice of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) ) ASYNC(ik)
!
! Computing F(Q)
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1, L_SliceZ
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                     PF(L,i,j,k,iold) = 0.5_rp * PQ(L,i,j,k,iold) + 1.0_rp 
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!$ACC WAIT (ik)
!
!========================================
! Pre-processing slice #3 ... #NbSliceZ-2
!========================================
      DO ik = 3, NbSliceZ
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
!
! End of association
         CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(iold) ) )
         CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(iold) ) )
!
! Change of ptr for computing
         itmp = inew
         inew = iold
         iold = itmp
!
! New slice
         kBeg(inew) = (ik-1) * L_SliceZ + 1
         kEnd(inew) =  ik    * L_SliceZ
!
! Association on the device
         WRITE (6,'(A,I3,A)')'Association of slice #', ik, ' on the device'
!
         length = 8 * SIZE(Q(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
         CALL ACC_MAP_DATA (Q(1,1,1,kBeg(inew) ), PQ(1,1,1,1,inew), length)
!
         length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
         CALL ACC_MAP_DATA (F(1,1,1,kBeg(inew) ), PF(1,1,1,1,inew), length)
!
! Host 2 Device for a slice of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) ) ASYNC(ik)
!
! Computing F(Q)
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1, L_SliceZ
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                     PF(L,i,j,k,iold) = 0.5_rp * PQ(L,i,j,k,iold) + 1.0_rp
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!
!$ACC WAIT (ik)
!
! End of loop over slices
      END DO
!
!==================================
! Post-processing slice #NbSliceZ-1
!==================================
      ik = NbSliceZ-1
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) ) ASYNC(ik)
!
! End of association
      CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(iold) ) )
      CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(iold) ) )
!
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1, L_SliceZ
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                     PF(L,i,j,k,inew) = 0.5_rp * PQ(L,i,j,k,inew) + 1.0_rp
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!
!$ACC WAIT (ik)
!
!================================
! Post-processing slice #NbSliceZ
!================================
      ik = NbSliceZ
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
!
! End of association
      CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(inew) ) )
      CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(inew) ) )
!
! Free ressources on device
      DEALLOCATE (PF, PQ)
!
! Check of F
      checksum = 0.0_rp
      DO k = 1, Nz
         DO j = 1, Ny
            DO i = 1, Nx
               DO L = 1, 5
                  checksum = MAX(checksum,ABS(F(L,i,j,k)-FF(L,i,j,k) ) )
               END DO
            END DO
         END DO
      END DO
      WRITE (6,'(A,ES15.8)') 'Checksum of F :', checksum
      CALL Flush (6)
!
! Free ressources on host
      DEALLOCATE (Q, F, FF)
!
      STOP
      END PROGRAM SMALL

Regards,
Guy.

Hi Mat,

I add halo of ghost cells in order to get closer to the real application I want to accelerate, but the code crashes with the following message :

FATAL ERROR: variable in OpenACC API data map routine was partially present on the GPU

I have identified the line, it’s the second call to ACC_MAP_DATA for array Q but the parts of arrays I want to map have the same size :

      length = 8 * SIZE(Q (1:5,1:Nx,1:Ny,kBeg(inew)-k1add:kEnd(inew)+k2add ) )
      write (6,'(A,2I9)') 'before :', length,
     &         8 * SIZE(PQ(1:5,1:Nx,1:Ny,1:L_SliceZ+k1add+k2add,inew) )
      call flush (6)
      CALL ACC_MAP_DATA (Q(1,1,1,kBeg(inew)-k1add ), PQ(1,1,1,1,inew), length)
      write (6,'(A)') 'after'
      call flush (6)

At the runtime I get

./a.out 
before :  6400000  6400000
FATAL ERROR: variable in OpenACC API data map routine was partially present on the GPU

The only difference with before are the 3 ghost cells left and right of the fourth dimension (so Z direction). I’ve done other tests, with ghosts cells in the X or Y direction or both of them and it works perfectly.
So, please, can you explain to me what’s going wrong when it is in the direction where the slicing takes place ?

Here is the whole example

      PROGRAM SMALL
!
      USE openacc
#ifdef _CUDA
      USE cudafor
#endif
!
      IMPLICIT NONE
!
      INTEGER, PARAMETER  ::  rp = kind(1.0D0)
      INTEGER, PARAMETER  ::  L_SliceZ = 10
      INTEGER, PARAMETER  ::  Nx = 100
      INTEGER, PARAMETER  ::  Ny = 100
      INTEGER, PARAMETER  ::  Nz = 200
!
      INTEGER, PARAMETER  :: k1add  = 3
      INTEGER, PARAMETER  :: k2add  = 3
!
      REAL(rp), PARAMETER ::  one  = 1.0_rp
      REAL(rp), PARAMETER ::  zero = 0.0_rp
      REAL(rp), PARAMETER ::  half = 0.5_rp
!
! temporary scalar variables like loop indexes
      INTEGER  ::  i, j, k, L, length, ik
      REAL(rp)  ::  checksum
!
      INTEGER, DIMENSION(0:1)  ::  kBeg, kEnd
      INTEGER  ::  inew, iold, itmp, NbSliceZ
!
! Arrays of data
      REAL(rp), DIMENSION(:,:,:,:),   ALLOCATABLE  ::  Q , F , FF
      REAL(rp), DIMENSION(:,:,:,:,:), POINTER      ::  PQ, PF
#ifdef _CUDA
      ATTRIBUTES(DEVICE)  ::  PQ, PF
#endif
!
! Dynamic Allocation
      ALLOCATE (Q (5, Nx, Ny, 1-k1add:Nz+k2add), F (5, Nx, Ny, Nz), FF (5, Nx, Ny, Nz) )
!
! Initialization of the data
      DO k = 1-k1add, Nz+k2add
         DO j = 1, Ny
            DO i = 1, Nx
               DO L =  1, 5
                  CALL RANDOM_NUMBER ( Q(L,i,j,k) )
               END DO
            END DO
         END DO
      END DO
!
      DO k = 1, Nz
         DO j = 1, Ny
            DO i = 1, Nx
               DO L =  1, 5
                  FF(L,i,j,k) = (
     &           + Q(L,i  ,j  ,k+1) + Q(L,i  ,j  ,k+2) + Q(L,i  ,j  ,k+3)
     &           + Q(L,i  ,j  ,k-1) + Q(L,i  ,j  ,k-2) + Q(L,i  ,j  ,k-3)
     &           + Q(L,i,j,k) ) * half + one
               END DO
            END DO
         END DO
      END DO
!
! Allocation on device
      NbSliceZ = Nz / L_SliceZ
!
      ALLOCATE (PQ(5,Nx,Ny,L_SliceZ+k1add+k2add,0:1), PF(5,Nx,Ny,L_SliceZ,0:1) )
!
      inew = 1
      iold = 0
!
!========================
! Pre-processing slice #1
!========================
      ik = 1
      kBeg(iold) = (ik-1) * L_SliceZ + 1
      kEnd(iold) =  ik    * L_SliceZ
!
! Association on the device
      length = 8 * SIZE(Q(1:5,1:Nx,1:Ny,kBeg(iold)-k1add:kEnd(iold)+k2add ) )
      CALL ACC_MAP_DATA (Q(1,1,1,kBeg(iold)-k1add ), PQ(1,1,1,1,iold), length)
!
      length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
      CALL ACC_MAP_DATA (F(1,1,1,kBeg(iold) ), PF(1,1,1,1,iold), length)
!
! Host To Device for slice #1 of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(iold)-k1add:kEnd(iold)+k2add ) )
!
!========================
! Pre-processing slice #2
!========================
      ik = 2
      kBeg(inew) = (ik-1) * L_SliceZ + 1
      kEnd(inew) =  ik    * L_SliceZ
!
! Association on the device
      length = 8 * SIZE(Q (1:5,1:Nx,1:Ny,kBeg(inew)-k1add:kEnd(inew)+k2add ) )
      write (6,'(A,2I9)') 'before :', length,
     &         8 * SIZE(PQ(1:5,1:Nx,1:Ny,1:L_SliceZ+k1add+k2add,inew) )
      call flush (6)
      CALL ACC_MAP_DATA (Q(1,1,1,kBeg(inew)-k1add ), PQ(1,1,1,1,inew), length)
      write (6,'(A)') 'after'
      call flush (6)
!
      length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
      CALL ACC_MAP_DATA (F(1,1,1,kBeg(inew) ), PF(1,1,1,1,inew), length)
!
! Host 2 Device for slice #2 of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(inew)-k1add:kEnd(inew)+k2add ) ) ASYNC(ik)
!
! Computing F(Q)
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1+k1add, L_SliceZ + k1add
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                  PF(L,i,j,k-k1add,iold) = (
     &          + PQ(L,i  ,j  ,k+1,iold) + PQ(L,i  ,j  ,k+2,iold) + PQ(L,i  ,j  ,k+3,iold)
     &          + PQ(L,i  ,j  ,k-1,iold) + PQ(L,i  ,j  ,k-2,iold) + PQ(L,i  ,j  ,k-3,iold)
     &          + PQ(L,i  ,j,  k  ,iold) ) * half + one
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!$ACC WAIT (ik)
!
!========================================
! Pre-processing slice #3 ... #NbSliceZ-2
!========================================
      DO ik = 3, NbSliceZ
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) )
!
! End of association
         CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(iold)-k1add ) )
         CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(iold) ) )
!
! Change of ptr for computing
         itmp = inew
         inew = iold
         iold = itmp
!
! New slice
         kBeg(inew) = (ik-1) * L_SliceZ + 1
         kEnd(inew) =  ik    * L_SliceZ
!
! Association on the device
         WRITE (6,'(A,I3,A)')'Association of slice #', ik, ' on the device'
!
         length = 8 * SIZE(Q(1:5,1:Nx,1:Ny,kBeg(inew)-k1add:kEnd(inew)-k1add ) )
         CALL ACC_MAP_DATA (Q(1,1,1,kBeg(inew)-k1add ), PQ(1,1,1,1,inew), length)
!
         length = 8 * SIZE(F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
         CALL ACC_MAP_DATA (F(1,1,1,kBeg(inew) ), PF(1,1,1,1,inew), length)
!
! Host 2 Device for slice #ik of Q
!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg(inew)-k1add:kEnd(inew)+k2add ) ) ASYNC(ik)
!
! Computing F(Q)
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1+k1add, L_SliceZ+k1add
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                  PF(L,i,j,k-k1add,iold) = (
     &          + PQ(L,i  ,j  ,k+1,iold) + PQ(L,i  ,j  ,k+2,iold) + PQ(L,i  ,j  ,k+3,iold)
     &          + PQ(L,i  ,j  ,k-1,iold) + PQ(L,i  ,j  ,k-2,iold) + PQ(L,i  ,j  ,k-3,iold)
     &          + PQ(L,i  ,j,  k  ,iold) ) * half + one
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!
!$ACC WAIT (ik)
!
! End of loop over slices
      END DO
!
!==================================
! Post-processing slice #NbSliceZ-1
!==================================
      ik = NbSliceZ-1
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(iold):kEnd(iold) ) ) ASYNC(ik)
!
! End of association
      CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(iold)-k1add ) )
      CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(iold) ) )
!
!$ACC PARALLEL LOOP COLLAPSE(4) ASYNC(ik)
         DO k = 1, L_SliceZ
            DO j = 1, Ny
               DO i = 1, Nx
                  DO L = 1, 5
                  PF(L,i,j,k-k1add,iold) = (
     &          + PQ(L,i  ,j  ,k+1,iold) + PQ(L,i  ,j  ,k+2,iold) + PQ(L,i  ,j  ,k+3,iold)
     &          + PQ(L,i  ,j  ,k-1,iold) + PQ(L,i  ,j  ,k-2,iold) + PQ(L,i  ,j  ,k-3,iold)
     &          + PQ(L,i  ,j,  k  ,iold) ) * half + one
                  END DO
               END DO
            END DO
         END DO
!$ACC END PARALLEL LOOP
!
!$ACC WAIT (ik)
!
!================================
! Post-processing slice #NbSliceZ
!================================
      ik = NbSliceZ
!
! PF's Device To Host transfer
!$ACC UPDATE SELF (F(1:5,1:Nx,1:Ny,kBeg(inew):kEnd(inew) ) )
!
! End of association
      CALL ACC_UNMAP_DATA (Q(1,1,1,kBeg(inew)-k1add ) )
      CALL ACC_UNMAP_DATA (F(1,1,1,kBeg(inew) ) )
!
! Free ressources on device
      DEALLOCATE (PF, PQ)
!
! Check of F
      checksum = 0.0_rp
      DO k = 1, Nz
         DO j = 1, Ny
            DO i = 1, Nx
               DO L = 1, 5
                  checksum = MAX(checksum,ABS(F(L,i,j,k)-FF(L,i,j,k) ) )
               END DO
            END DO
         END DO
      END DO
      WRITE (6,'(A,ES15.8)') 'Checksum of F :', checksum
      CALL Flush (6)
!
! Free ressources on host
      DEALLOCATE (Q, F, FF)
!
      STOP
      END PROGRAM SMALL

I compile it with :

pgf90 -acc -Mcuda -Minfo=all -Minfo=accel -O3 -Mfixed -Mextend -g -traceback small.F90

Regards,
Guy.

Hi Guy,

Again, using acc_map_data to achieve pipelining is not correct. You’ve made it way more complicated than necessary. Yes, you fixed the collision issue by adding an extra dimension to PQ, but then why not just use Q directly?

The actual error you’re seeing is because you’re not allowed to map the same host variable twice. So to get this to work, you must first unmap Q and F, before remapping them. But this is going to cause you other problems. Yes it’s different sections of the array and perhaps you could work around this by passing the address of the section to another routine where you do the mapping, but this is way too much work when there’s a much simpler solution.

Granted, I don’t understand the motivation as to why you want to use acc_map_data, and perhaps there’s a good reason, but lets first start with the simple solution (below) and then build up from there.

% cat small2.F90
      PROGRAM SMALL
       IMPLICIT NONE
       INTEGER, PARAMETER  ::  rp = kind(1.0D0)
       INTEGER, PARAMETER  ::  L_SliceZ = 10
       INTEGER, PARAMETER  ::  Nx = 100
       INTEGER, PARAMETER  ::  Ny = 100
       INTEGER, PARAMETER  ::  Nz = 200
!
       INTEGER, PARAMETER  :: k1add  = 3
       INTEGER, PARAMETER  :: k2add  = 3
!
       REAL(rp), PARAMETER ::  one  = 1.0_rp
       REAL(rp), PARAMETER ::  zero = 0.0_rp
       REAL(rp), PARAMETER ::  half = 0.5_rp
!
! temporary scalar variables like loop indexes
       INTEGER  ::  i, j, k, L, length, ik, queue
       REAL(rp)  ::  checksum
!
       INTEGER  ::  kBeg, kEnd
       INTEGER  ::  inew, iold, itmp, NbSliceZ
!
! Arrays of data
       REAL(rp), DIMENSION(:,:,:,:),   ALLOCATABLE  ::  Q , F , FF
!
! Dynamic Allocation
       ALLOCATE (Q (5, Nx, Ny, 1-k1add:Nz+k2add), F (5, Nx, Ny, Nz), FF (5, Nx, Ny, Nz) )

! Initialization of the data
       DO k = 1-k1add, Nz+k2add
          DO j = 1, Ny
             DO i = 1, Nx
                DO L =  1, 5
                   CALL RANDOM_NUMBER ( Q(L,i,j,k) )
                END DO
             END DO
          END DO
       END DO
!
!$ACC ENTER DATA CREATE(Q,F)
       DO k = 1, Nz
          DO j = 1, Ny
             DO i = 1, Nx
                DO L =  1, 5
                   FF(L,i,j,k) = (
     &           + Q(L,i  ,j  ,k+1) + Q(L,i  ,j  ,k+2) + Q(L,i  ,j ,k+3)
     &           + Q(L,i  ,j  ,k-1) + Q(L,i  ,j  ,k-2) + Q(L,i  ,j ,k-3)
     &           + Q(L,i,j,k) ) * half + one
                END DO
             END DO
          END DO
       END DO

! Allocation on device
       NbSliceZ = Nz / L_SliceZ

! Update the device halos

       do ik = 1, NbSliceZ
         kBeg = (ik-1) * L_SliceZ + 1
         kEnd =  kBeg + L_SliceZ - 1

!$ACC UPDATE DEVICE (Q(1:5,1:Nx,1:Ny,kBeg-k1add:kEnd+k2add)) async(mod(ik,2)+1)
! Computing F(Q)
!$ACC PARALLEL LOOP COLLAPSE(4) async(mod(ik,2)+1)
          DO k = kBeg, kEnd
             DO j = 1, Ny
                DO i = 1, Nx
                   DO L = 1, 5
                   F(L,i,j,k) = (
     &          + Q(L,i  ,j  ,k+1) + Q(L,i  ,j  ,k+2) + Q(L,i  ,j  ,k+3)
     &          + Q(L,i  ,j  ,k-1) + Q(L,i  ,j  ,k-2) + Q(L,i  ,j  ,k-3)
     &          + Q(L,i  ,j,  k  ) ) * half + one
                   END DO
                END DO
             END DO
          END DO
!$ACC END PARALLEL LOOP
       end do  ! en ik loop

!
!$ACC WAIT
!$ACC UPDATE SELF (F)
!
! Check of F
       checksum = 0.0_rp
       DO k = 1, Nz
          DO j = 1, Ny
             DO i = 1, Nx
                DO L = 1, 5
                   checksum = MAX(checksum,ABS(F(L,i,j,k)-FF(L,i,j,k) ))
                END DO
             END DO
          END DO
       END DO
       WRITE (6,'(A,ES15.8)') 'Checksum of F :', checksum
       CALL Flush (6)
!
!$ACC EXIT DATA DELETE(Q,F)
! Free ressources on host
       DEALLOCATE (Q, F, FF)
!
       STOP
       END PROGRAM SMALL
% pgf90 -fast -Mfixed -Mextend small2.F90 -acc -Minfo=accel -ta=tesla:cuda8.0,cc35
small:
     40, Generating enter data create(f(:,:,:,:),q(:,:,:,:))
     63, Generating update device(q(1:5,1:100,1:100,kbeg-3:kend+3))
     65, Accelerator kernel generated
         Generating Tesla code
         66, !$acc loop gang, vector(128) collapse(4) ! blockidx%x threadidx%x
         67,   ! blockidx%x threadidx%x collapsed
         68,   ! blockidx%x threadidx%x collapsed
         69,   ! blockidx%x threadidx%x collapsed
     65, Generating implicit copyin(q(1:5,1:100,1:100,(ik-1)*10-2:(ik-1)*10+13))
         Generating implicit copyout(f(1:5,1:100,1:100,(ik-1)*10+1:(ik-1)*10+10))
     83, Generating update self(f(:,:,:,:))
     99, Generating exit data delete(q(:,:,:,:),f(:,:,:,:))
% a.out
Checksum of F : 0.00000000E+00
Warning: ieee_inexact is signaling
FORTRAN STOP
  • Mat

Hi Mat,

one again,thank you for your explanations.

To answer your questions, I’d like to use slices of arrays like Q, instead of the whole arrays on the device because of memory consumption.

Some days ago, I encountered a person who has better experience than I using OpenACC and the conclusion of our thoughts was to use the acc_map_data function. Maybe, the problem of twice mapping was forgotten.

So, now, I put aside this function and I start over using your version.

Thanks a lot for your help.

Guy.