6.8. Array Shaping

A pointer variable can be shaped to a multi-dimensional array to facilitate data access. This is achieved by a shape-operator casted in front of a pointer (lvalue expression):
([ s₁][ s₂]…[ s_n]) pointer
where each s_i is an integral-type expression of positive value. The shape-operator can appear in either the motion-clause of the target update directive or the depend clause.

The following example shows the use of the shape-operator in the target update directive. The shape-operator ([nx][ny+2]) casts pointer variable a to a 2-dimentional array of size nx x (ny+2) . The resulting array is then accessed as array sections (such as [0:nx][1] and [0:nx][ny]) in the from or to clause for transferring two columns of noncontiguous boundary data from or to the device. Note the use of additional parentheses around the shape-operator and a to ensure the correct precedence over array-section operations.

//%compiler: clang
//%cflags: -fopenmp

/*
* name: array_shaping.1
* type: C
* version: omp_5.1
*/
#pragma omp begin declare target
  int do_work(double *a, int nx, int ny);
  int other_work(double *a, int nx, int ny);
#pragma omp end declare target

void exch_data(double *a, int nx, int ny);

void array_shaping(double *a, int nx, int ny)
{
   // map data to device and do work
   #pragma omp target data map(a[0:nx*(ny+2)])
   {
      // do work on the device
      #pragma omp target  // map(a[0:nx*(ny+2)]) is optional here
      do_work(a, nx, ny);

      // update boundary points (two columns of 2D array) on the host
      // pointer is shaped to 2D array using the shape-operator
      #pragma omp target update from( (([nx][ny+2])a)[0:nx][1], \
                                      (([nx][ny+2])a)[0:nx][ny] )

      // exchange ghost points with neighbors
      exch_data(a, nx, ny);

      // update ghost points (two columns of 2D array) on the device
      // pointer is shaped to 2D array using the shape-operator
      #pragma omp target update to( (([nx][ny+2])a)[0:nx][0], \
                                    (([nx][ny+2])a)[0:nx][ny+1] )

      // perform other work on the device
      #pragma omp target  // map(a[0:nx*(ny+2)]) is optional here
      other_work(a, nx, ny);
   }
}

The shape operator is not defined for Fortran. Explicit array shaping of procedure arguments can be used instead to achieve a similar goal. Below is the Fortran-equivalent of the above example that illustrates the support of transferring two rows of noncontiguous boundary data in the target update directive.

!!%compiler: gfortran
!!%cflags: -fopenmp

! name: array_shaping.1
! type: F-free
! version: omp_5.2

module m
   interface
      subroutine do_work(a, nx, ny)
         !$omp declare target enter(do_work)
         integer, intent(in) :: nx, ny
         double precision a(0:nx+1,ny)
      end subroutine do_work

      subroutine other_work(a, nx, ny)
         !$omp declare target enter(other_work)
         integer, intent(in) :: nx, ny
         double precision a(0:nx+1,ny)
      end subroutine other_work

      subroutine exch_data(a, nx, ny)
         integer, intent(in) :: nx, ny
         double precision a(0:nx+1,ny)
      end subroutine exch_data
   end interface
end module m

subroutine array_shaping(a, nx, ny)
   use m
   implicit none
   integer, intent(in) :: nx, ny
   double precision a(0:nx+1,ny)

   ! map data to device and do work
   !$omp target data map(a)

      ! do work on the device
      !$omp target      ! map(a) is optional here
      call do_work(a, nx, ny)
      !$omp end target

      ! update boundary points (two rows of 2D array) on the host.
      ! data transferred are noncontiguous
      !$omp target update from( a(1,1:ny), a(nx,1:ny) )

      ! exchange ghost points with neighbors
      call exch_data(a, nx, ny)

      ! update ghost points (two rows of 2D array) on the device.
      ! data transferred are noncontiguous
      !$omp target update to( a(0,1:ny), a(nx+1,1:ny) )

      ! perform other work on the device
      !$omp target      ! map(a) is optional here
      call other_work(a, nx, ny)
      !$omp end target

   !$omp end target data

end subroutine