6.11. target enter data and target exit data Constructs

The structured data construct (target data) provides persistent data on a device for subsequent target constructs as shown in the target data examples above. This is accomplished by creating a single target data region containing target constructs.

The unstructured data constructs allow the creation and deletion of data on the device at any appropriate point within the host code, as shown below with the target enter data and target exit data constructs.

The following C++ code creates/deletes a vector in a constructor/destructor of a class. The constructor creates a vector with target enter data and uses an alloc modifier in the map clause to avoid copying values to the device. The destructor deletes the data (target exit data) and uses the delete modifier in the map clause to avoid copying data back to the host. Note, the stand-alone target enter data occurs after the host vector is created, and the target exit data construct occurs before the host data is deleted.

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

/*
* name:        target_unstructured_data.1
* type:        C++
* version:     omp_4.5
*/
class Matrix
{

  Matrix(int n) {
    len = n;
    v = new double[len];
    #pragma omp target enter data map(alloc:v[0:len])
  }

  ~Matrix() {
    // NOTE: delete map type should be used, since the corresponding
    // host data will cease to exist after the deconstructor is called.

    #pragma omp target exit data map(delete:v[0:len])
    delete[] v;
  }

  private:
  double* v;
  int len;

};

The following C code allocates and frees the data member of a Matrix structure. The init_matrix function allocates the memory used in the structure and uses the target enter data directive to map it to the target device. The free_matrix function removes the mapped array from the target device and then frees the memory on the host. Note, the stand-alone target enter data occurs after the host memory is allocated, and the target exit data construct occurs before the host data is freed.

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

/*
 * name:        target_unstructured_data.1
 * type:        C
 * version:     omp_4.5
 */
#include <stdlib.h>
typedef struct {
  double *A;
  int N;
} Matrix;

void init_matrix(Matrix *mat, int n)
{
  mat->A = (double *)malloc(n*sizeof(double));
  mat->N = n;
  #pragma omp target enter data map(alloc:mat->A[:n])
}

void free_matrix(Matrix *mat)
{
  #pragma omp target exit data map(delete:mat->A[:mat->N])
  mat->N = 0;
  free(mat->A);
  mat->A = NULL;
}

The following Fortran code allocates and deallocates a module array. The initialize subroutine allocates the module array and uses the target enter data directive to map it to the target device. The finalize subroutine removes the mapped array from the target device and then deallocates the array on the host. Note, the stand-alone target enter data occurs after the host memory is allocated, and the target exit data construct occurs before the host data is deallocated.

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

! name:        target_unstructured_data.1
! type:        F-free
! version:     omp_4.5
module example
  real(8), allocatable :: A(:)

  contains
    subroutine initialize(N)
      integer :: N

      allocate(A(N))
      !$omp target enter data map(alloc:A)

    end subroutine initialize

    subroutine finalize()

      !$omp target exit data map(delete:A)
      deallocate(A)

    end subroutine finalize
end module example