OpenMP Runtime Routines for Memory Allocation

Intel compilers support a number of OpenMP runtime routines for performing memory allocations. These routines are shown in the table below.

Note that the three routines omp_target_alloc_device, omp_target_alloc_host, and omp_target_alloc_shared are Intel extensions to the OpenMP specification.

The following examples use the above OpenMP memory allocation routines. Compare those to the ones using map clauses.

For more information about memory allocation, see:

• Data Parallel C++, by James Reinders et al.

• SYCL 2020 Specification

• oneAPI Level Zero Specification

• The SYCL part of this guide

Using the map Clause

The first example uses map clauses to allocate memory on a device and copy data between the host and the device.

In the following example, arrays A, B, and C are allocated in system memory by calling the C/C++ standard library routine, malloc.

The target construct on line 58 is the main kernel that computes the values of array C on the device. The map(tofrom: C[0:length) clause is specified on this target construct since the values of C need to be transferred from the host to the device before the computation, and from the device to the host at the end of the computation. The map(to: A[0:length], B[0:length]) is specified for arrays``A`` and B since the values of these arrays need to be transferred from the host to the device, and the device only reads these values. Under the covers, the map clauses cause storage for the arrays to be allocated on the device and data to be copied from the host to the device, and vice versa.

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define iterations 100
#define length     64*1024*1024

int main(void)
{
  size_t bytes = length*sizeof(double);
  double * __restrict A;
  double * __restrict B;
  double * __restrict C;
  double scalar = 3.0;
  double nstream_time = 0.0;

  // Allocate arrays on the host using plain malloc()

  A = (double *) malloc(bytes);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using plain malloc().\n");
     exit(1);
  }

  B = (double *) malloc(bytes);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using plain malloc().\n");
     exit(1);
  }

  C = (double *) malloc(bytes);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using plain malloc().\n");
     exit(1);
  }

  // Initialize the arrays

  #pragma omp parallel for
  for (size_t i=0; i<length; i++) {
      A[i] = 2.0;
      B[i] = 2.0;
      C[i] = 0.0;
  }

  // Perform the computation

  nstream_time = omp_get_wtime();
  for (int iter = 0; iter<iterations; iter++) {
      #pragma omp target teams distribute parallel for \
        map(to: A[0:length], B[0:length]) \
        map(tofrom: C[0:length])
      for (size_t i=0; i<length; i++) {
          C[i] += A[i] + scalar * B[i];
      }
  }
  nstream_time = omp_get_wtime() - nstream_time;

  // Validate and output results

  double ar = 2.0;
  double br = 2.0;
  double cr = 0.0;
  for (int iter = 0; iter<iterations; iter++) {
      for (int i=0; i<length; i++) {
          cr += ar + scalar * br;
      }
  }

  double asum = 0.0;
  #pragma omp parallel for reduction(+:asum)
  for (size_t i=0; i<length; i++) {
      asum += fabs(C[i]);
  }

  free(A);
  free(B);
  free(C);

  double epsilon=1.e-8;
  if (fabs(cr - asum)/asum > epsilon) {
      printf("Failed Validation on output array\n"
             "       Expected checksum: %lf\n"
             "       Observed checksum: %lf\n"
             "ERROR: solution did not validate\n", cr, asum);
      return 1;
  } else {
      printf("Solution validates\n");
      double avgtime = nstream_time/iterations;
      printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
  }

  return 0;
}

Compilation command:

icpx -fiopenmp -fopenmp-targets=spir64 test_target_map.cpp

Run command:

OMP_TARGET_OFFLOAD=MANDATORY ZE_AFFINITY_MASK=0.0 LIBOMPTARGET_DEBUG=1 ./a.out

The map clauses on the target construct inside the iterations loop cause data (values of A, B, C) to be transferred from the host to the device at the beginning of each target region, and cause data (values of C) to be transferred from the device to the host at the end of each target region. These data transfers incur a significant performance overhead. A better approach using map clauses would be to put the whole iterations loop inside a target data construct with the map clauses. This causes the transfers to occur once at the beginning of the iterations loop, and another time at the end of the iterations loop. The modified example using target data and map clauses is shown below.

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define iterations 100
#define length     64*1024*1024

int main(void)
{
  size_t bytes = length*sizeof(double);
  double * __restrict A;
  double * __restrict B;
  double * __restrict C;
  double scalar = 3.0;
  double nstream_time = 0.0;

  // Allocate arrays on the host using plain malloc()

  A = (double *) malloc(bytes);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using plain malloc().\n");
     exit(1);
  }

  B = (double *) malloc(bytes);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using plain malloc().\n");
     exit(1);
  }

  C = (double *) malloc(bytes);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using plain malloc().\n");
     exit(1);
  }

  // Initialize the arrays

  #pragma omp parallel for
  for (size_t i=0; i<length; i++) {
      A[i] = 2.0;
      B[i] = 2.0;
      C[i] = 0.0;
  }

  // Perform the computation

  nstream_time = omp_get_wtime();
  #pragma omp target data  map(to: A[0:length], B[0:length]) \
    map(tofrom: C[0:length])
  {
     for (int iter = 0; iter<iterations; iter++) {
         #pragma omp target teams distribute parallel for
         for (size_t i=0; i<length; i++) {
             C[i] += A[i] + scalar * B[i];
         }
     }
  }
  nstream_time = omp_get_wtime() - nstream_time;

  // Validate and output results

  double ar = 2.0;
  double br = 2.0;
  double cr = 0.0;
  for (int iter = 0; iter<iterations; iter++) {
      for (int i=0; i<length; i++) {
          cr += ar + scalar * br;
      }
  }

  double asum = 0.0;
  #pragma omp parallel for reduction(+:asum)
  for (size_t i=0; i<length; i++) {
      asum += fabs(C[i]);
  }

  free(A);
  free(B);
  free(C);

  double epsilon=1.e-8;
  if (fabs(cr - asum)/asum > epsilon) {
      printf("Failed Validation on output array\n"
             "       Expected checksum: %lf\n"
             "       Observed checksum: %lf\n"
             "ERROR: solution did not validate\n", cr, asum);
      return 1;
  } else {
      printf("Solution validates\n");
      double avgtime = nstream_time/iterations;
      printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
  }

  return 0;
}

omp_target_alloc

Next, the example above is modified to use device allocations instead of map clauses. Storage for arrays A, B, and C is directly allocated on the device by calling the OpenMP runtime routine omp_target_alloc. The routine takes two arguments: the number of bytes to allocate on the device, and the number of the device on which to allocate the storage. The routine returns a device pointer that references the device address of the storage allocated on the device. If the call to omp_target_alloc returns NULL, then this indicates that the allocation was not successful.

To access the allocated memory in a target construct, the device pointer returned by a call to omp_target_alloc is listed in an is_device_ptr clause on the target construct. This ensures that there is no data transfer before and after kernel execution since the kernel operates on data that is already on the device.

At the end of the program, the runtime routine omp_target_free is used to deallocate the storage for A, B, and C on the device.

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define iterations 100
#define length     64*1024*1024

int main(void)
{
  int device_id = omp_get_default_device();
  size_t bytes = length*sizeof(double);
  double * __restrict A;
  double * __restrict B;
  double * __restrict C;
  double scalar = 3.0;
  double nstream_time = 0.0;

  // Allocate arrays in device memory

  A = (double *) omp_target_alloc(bytes, device_id);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using omp_target_alloc().\n");
     exit(1);
  }

  B = (double *) omp_target_alloc(bytes, device_id);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using omp_target_alloc().\n");
     exit(1);
  }

  C = (double *) omp_target_alloc(bytes, device_id);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using omp_target_alloc().\n");
     exit(1);
  }

  // Initialize the arrays

  #pragma omp target teams distribute parallel for \
    is_device_ptr(A,B,C)
  for (size_t i=0; i<length; i++) {
      A[i] = 2.0;
      B[i] = 2.0;
      C[i] = 0.0;
  }

  // Perform the computation 'iterations' number of times

  nstream_time = omp_get_wtime();
  for (int iter = 0; iter<iterations; iter++) {
      #pragma omp target teams distribute parallel for \
         is_device_ptr(A,B,C)
      for (size_t i=0; i<length; i++) {
          C[i] += A[i] + scalar * B[i];
      }
  }
  nstream_time = omp_get_wtime() - nstream_time;

  // Validate and output results

  double ar = 2.0;
  double br = 2.0;
  double cr = 0.0;
  for (int iter = 0; iter<iterations; iter++) {
      for (int i=0; i<length; i++) {
          cr += ar + scalar * br;
      }
  }

  double asum = 0.0;
  #pragma omp target teams distribute parallel for reduction(+:asum) \
    map(tofrom: asum) is_device_ptr(C)
  for (size_t i=0; i<length; i++) {
      asum += fabs(C[i]);
  }

  omp_target_free(A, device_id);
  omp_target_free(B, device_id);
  omp_target_free(C, device_id);

  double epsilon=1.e-8;
  if (fabs(cr - asum)/asum > epsilon) {
      printf("Failed Validation on output array\n"
             "       Expected checksum: %lf\n"
             "       Observed checksum: %lf\n"
             "ERROR: solution did not validate\n", cr, asum);
      return 1;
  } else {
      printf("Solution validates\n");
      double avgtime = nstream_time/iterations;
      printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
  }

  return 0;
}

Notes:

• When calling omp_target_alloc, the device number specified must be one of the supported devices, other than the host device. This will be the device on which storage will be allocated.

• Since the arrays A, B, and C are not accessible from the host, the initialization of the arrays, kernel execution, and summation of elements of C all need to be done inside OpenMP target regions.

• A device allocation can only be accessed by the device specified in the omp_target_alloc call, but may be copied to memory allocated on the host or other devices by calling omp_target_memcpy.

omp_target_alloc_device

The Intel extension omp_target_alloc_device is similar to omp_target_alloc. It is also called with two arguments: the number of bytes to allocate on the device, and the number of the device on which to allocate the storage. The routine returns a device pointer that references the device address of the storage allocated on the device. If the call to omp_target_alloc_device returns NULL, then this indicates that the allocation was not successful.

The above omp_target_alloc example can be rewritten using omp_target_alloc_device by simply replacing the call to omp_target_alloc with a call to omp_targer_alloc_device as shown below.

At the end of the program, the runtime routine omp_target_free is used to deallocate the storage for A, B, and C on the device.

  // Allocate arrays in device memory

  A = (double *) omp_target_alloc_device(bytes, device_id);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using omp_target_alloc_device().\n");
     exit(1);
  }

  B = (double *) omp_target_alloc_device(bytes, device_id);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using omp_target_alloc_device().\n");
     exit(1);
  }

  C = (double *) omp_target_alloc_device(bytes, device_id);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using omp_target_alloc_device().\n");
     exit(1);
  }

Note:

• All of the above Notes that apply to omp_target_alloc also apply to omp_target_alloc_device.

omp_target_alloc_host

The above example can also be rewritten by doing a host allocation for A, B, and C. This allows the memory to be accessible to the host and all supported devices.

In the following modified example, the omp_target_alloc_host runtime routine (an Intel extension) is called to allocate storage for each of the arrays A, B, and C. The routine takes two arguments: the number of bytes to allocate, and a device number. The device number must be one of the supported devices, other than the host device. The routine returns a pointer to a storage location in host memory. If the call to omp_target_alloc_host returns NULL, this indicates that the allocation was not successful.

Note the directive requires unified_address is specified at the top of the program. This requires that the implementation guarantee that all devices accessible through OpenMP API routines and directives use a unified address space. In this address space, a pointer will always refer to the same location in memory from all devices, and the is_device_ptr clause is not necessary to obtain device addresses from device pointers for use inside target regions. When using Intel compilers, the requires unified_address directive is actually not needed, since unified address space is guaranteed by default. However, for portability the code includes the directive.

The pointer returned by a call to omp_target_alloc_host can be used to access the storage from the host and all supported devices. No map clauses and no is_device_ptr clauses are needed on a target construct to access the memory from a device since a unified address space is used.

At the end of the program, the runtime routine omp_target_free is used to deallocate the storage for A, B, and C.

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#pragma omp requires unified_address

#define iterations 100
#define length     64*1024*1024

int main(void)
{
  int device_id = omp_get_default_device();
  size_t bytes = length*sizeof(double);
  double * __restrict A;
  double * __restrict B;
  double * __restrict C;
  double scalar = 3.0;
  double nstream_time = 0.0;

  // Allocate arrays in host memory

  A = (double *) omp_target_alloc_host(bytes, device_id);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using omp_target_alloc_host().\n");
     exit(1);
  }

  B = (double *) omp_target_alloc_host(bytes, device_id);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using omp_target_alloc_host().\n");
     exit(1);
  }

  C = (double *) omp_target_alloc_host(bytes, device_id);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using omp_target_alloc_host().\n");
     exit(1);
  }

  // Initialize the arrays

  #pragma omp parallel for
  for (size_t i=0; i<length; i++) {
      A[i] = 2.0;
      B[i] = 2.0;
      C[i] = 0.0;
  }

  // Perform the computation

  nstream_time = omp_get_wtime();
  for (int iter = 0; iter<iterations; iter++) {
      #pragma omp target teams distribute parallel for
      for (size_t i=0; i<length; i++) {
          C[i] += A[i] + scalar * B[i];
      }
  }
  nstream_time = omp_get_wtime() - nstream_time;

  // Validate and output results

  double ar = 2.0;
  double br = 2.0;
  double cr = 0.0;
  for (int iter = 0; iter<iterations; iter++) {
      for (int i=0; i<length; i++) {
          cr += ar + scalar * br;
      }
  }

  double asum = 0.0;
  #pragma omp parallel for reduction(+:asum)
  for (size_t i=0; i<length; i++) {
      asum += fabs(C[i]);
  }

  omp_target_free(A, device_id);
  omp_target_free(B, device_id);
  omp_target_free(C, device_id);

  double epsilon=1.e-8;
  if (fabs(cr - asum)/asum > epsilon) {
      printf("Failed Validation on output array\n"
             "       Expected checksum: %lf\n"
             "       Observed checksum: %lf\n"
             "ERROR: solution did not validate\n", cr, asum);
      return 1;
  } else {
      printf("Solution validates\n");
      double avgtime = nstream_time/iterations;
      printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
  }

  return 0;
}

Notes:

• When calling omp_target_alloc_host, the device number specified must be one of the supported devices, other than the host device.

• Since the arrays A, B, and C are accessible from the host and device, the initialization of the arrays and summation of elements of C may be done either on the host (outside of a target construct) or on the device (inside a target construct).

• Intel^® Data Center GPU Max Series does not support atomic operations (or algorithms that use atomic operations, such as some reductions) on host allocations (i.e., memory allocated via omp_target_alloc_host). Use atomic operations on memory allocated via omp_target_alloc_device, instead.

omp_target_alloc_shared

The above example is modified so that shared allocations are used instead of host allocations. The omp_target_alloc_shared runtime routine is called to allocate storage for each of arrays A, B, and C. The routine takes two arguments: the number of bytes to allocate on the device, and a device number. The device number must be one of the supported devices, other than the host device. The routine returns a pointer to a storage location in shared memory. If the call to omp_target_alloc_shared returns NULL, then this indicates that the allocation was not successful.

Note the requires unified_address directive is specified at the top of the program, for portability.

The pointer returned by a call to omp_target_alloc_shared can be used to access the storage from the host and all supported devices. No map clauses and no is_device_ptr clauses are needed on a target construct to access the memory from a device since a unified address space is used.

At the end of the program, the runtime routine omp_target_free is used to deallocate the storage for A, B, and C.

  // Allocate arrays in shared memory

  A = (double *) omp_target_alloc_shared(bytes, device_id);
  if (A == NULL){
     printf(" ERROR: Cannot allocate space for A using omp_target_alloc_shared().\n");
     exit(1);
  }

  B = (double *) omp_target_alloc_shared(bytes, device_id);
  if (B == NULL){
     printf(" ERROR: Cannot allocate space for B using omp_target_alloc_shared().\n");
     exit(1);
  }

  C = (double *) omp_target_alloc_shared(bytes, device_id);
  if (C == NULL){
     printf(" ERROR: Cannot allocate space for C using omp_target_alloc_shared().\n");
     exit(1);
  }

Notes:

• When calling omp_target_alloc_shared, the device number specified must be one of the supported devices, other than the host device.

• Since the arrays are accessible from the host and device, the initialization and verification may be done either on the host or on the device (inside a target construct).

• Concurrent access from host and device to memory allocated via omp_target_alloc_shared is not supported.

omp_target_memcpy

The following example shows how the runtime routine omp_target_memcpy may be used to copy memory from host to device, and from device to host. First arrays h_A, h_B, and h_C are allocated in system memory using plain malloc, and then initialized. Corresponding arrays d_A, d_B, and d_C are allocated on the device using omp_target_alloc.

Before the start of the target construct on line 104, the values in h_A, h_B, and h_C are copied to d_A, d_B, and d_C by calling omp_target_memcpy. After the target region, new d_C values computed on the device are copied to h_C by calling omp_target_memcpy.

#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define iterations 100
#define length     64*1024*1024

int main(void)
{
  int device_id = omp_get_default_device();
  int host_id = omp_get_initial_device();
  size_t bytes = length*sizeof(double);
  double * __restrict h_A;
  double * __restrict h_B;
  double * __restrict h_C;
  double * __restrict d_A;
  double * __restrict d_B;
  double * __restrict d_C;
  double scalar = 3.0;
  double nstream_time = 0.0;

  // Allocate arrays h_A, h_B, and h_C on the host using plain malloc()

  h_A = (double *) malloc(bytes);
  if (h_A == NULL){
     printf(" ERROR: Cannot allocate space for h_A using plain malloc().\n");
     exit(1);
  }

  h_B = (double *) malloc(bytes);
  if (h_B == NULL){
     printf(" ERROR: Cannot allocate space for h_B using plain malloc().\n");
     exit(1);
  }

  h_C = (double *) malloc(bytes);
  if (h_C == NULL){
     printf(" ERROR: Cannot allocate space for h_C using plain malloc().\n");
     exit(1);
  }

  // Allocate arrays d_A, d_B, and d_C on the device using omp_target_alloc()

  d_A = (double *) omp_target_alloc(bytes, device_id);
  if (d_A == NULL){
     printf(" ERROR: Cannot allocate space for d_A using omp_target_alloc().\n");
     exit(1);
  }

  d_B = (double *) omp_target_alloc(bytes, device_id);
  if (d_B == NULL){
     printf(" ERROR: Cannot allocate space for d_B using omp_target_alloc().\n");
     exit(1);
  }

  d_C = (double *) omp_target_alloc(bytes, device_id);
  if (d_C == NULL){
     printf(" ERROR: Cannot allocate space for d_C using omp_target_alloc().\n");
     exit(1);
  }

  // Initialize the arrays on the host

  #pragma omp parallel for
  for (size_t i=0; i<length; i++) {
      h_A[i] = 2.0;
      h_B[i] = 2.0;
      h_C[i] = 0.0;
  }

  // Call omp_target_memcpy() to copy values from host to device

  int rc = 0;
  rc = omp_target_memcpy(d_A, h_A, bytes, 0, 0, device_id, host_id);
  if (rc) {
     printf("ERROR: omp_target_memcpy(A) returned %d\n", rc);
     exit(1);
  }

  rc = omp_target_memcpy(d_B, h_B, bytes, 0, 0, device_id, host_id);
  if (rc) {
     printf("ERROR: omp_target_memcpy(B) returned %d\n", rc);
     exit(1);
  }

  rc = omp_target_memcpy(d_C, h_C, bytes, 0, 0, device_id, host_id);
  if (rc) {
     printf("ERROR: omp_target_memcpy(C) returned %d\n", rc);
     exit(1);
  }

  // Perform the computation

  nstream_time = omp_get_wtime();
  for (int iter = 0; iter<iterations; iter++) {
      #pragma omp target teams distribute parallel for \
        is_device_ptr(d_A,d_B,d_C)
      for (size_t i=0; i<length; i++) {
          d_C[i] += d_A[i] + scalar * d_B[i];
      }
  }
  nstream_time = omp_get_wtime() - nstream_time;

  // Call omp_target_memcpy() to copy values from device to host

  rc = omp_target_memcpy(h_C, d_C, bytes, 0, 0, host_id, device_id);
  if (rc) {
     printf("ERROR: omp_target_memcpy(A) returned %d\n", rc);
     exit(1);
  }

  // Validate and output results

  double ar = 2.0;
  double br = 2.0;
  double cr = 0.0;
  for (int iter = 0; iter<iterations; iter++) {
      for (int i=0; i<length; i++) {
          cr += ar + scalar * br;
      }
  }

  double asum = 0.0;
  #pragma omp parallel for reduction(+:asum)
  for (size_t i=0; i<length; i++) {
      asum += fabs(h_C[i]);
  }

  free(h_A);
  free(h_B);
  free(h_C);
  omp_target_free(d_A, device_id);
  omp_target_free(d_B, device_id);
  omp_target_free(d_C, device_id);

  double epsilon=1.e-8;
  if (fabs(cr - asum)/asum > epsilon) {
      printf("Failed Validation on output array\n"
             "       Expected checksum: %lf\n"
             "       Observed checksum: %lf\n"
             "ERROR: solution did not validate\n", cr, asum);
      return 1;
  } else {
      printf("Solution validates\n");
      double avgtime = nstream_time/iterations;
      printf("Checksum = %lf; Avg time (s): %lf\n", asum, avgtime);
  }

  return 0;
}

Performance Considerations

In the above examples (using the map clause, omp_target_alloc, omp_target_alloc_device, omp_target_alloc_host, omp_target_alloc_shared, omp_target_memcpy), the main kernel is the target construct that computes the values of array C. To get more accurate timings, this target construct is enclosed in a loop, so the offload happens iterations number of times (where iterations = 100). The average kernel time is computed by dividing the total time taken by the iterations loop by 100.

  // Perform the computation 'iterations' number of times

  nstream_time = omp_get_wtime();
  for (int iter = 0; iter<iterations; iter++) {
      #pragma omp target teams distribute parallel for \
         is_device_ptr(A,B,C)
      for (size_t i=0; i<length; i++) {
          C[i] += A[i] + scalar * B[i];
      }
  }
  nstream_time = omp_get_wtime() - nstream_time;

LIBOMPTARGET_DEBUG=1 output shows that all the above examples have the same ND_range partitioning.

Target LEVEL0 RTL --> Allocated a device memory 0xff00000020200000
Libomptarget --> omp_target_alloc returns device ptr 0xff00000020200000
Libomptarget --> Call to omp_target_alloc for device 0 requesting 536870912 bytes
Libomptarget --> Call to omp_get_num_devices returning 1
Libomptarget --> Call to omp_get_initial_device returning 1
Libomptarget --> Checking whether device 0 is ready.
Libomptarget --> Is the device 0 (local ID 0) initialized? 1
Libomptarget --> Device 0 is ready to use.

The following table shows the average times taken by the kernel in the various versions when running on the particular GPU used (1-stack only).

The above performance numbers show that the map version is the slowest version (0.183604 seconds). This is because of the data transfers that occur at the beginning and end of each kernel launch. The main kernel is launched 100 times. At the beginning of each kernel launch, storage for arrays A, B and C is allocated on the device, and the values of these arrays are copied from the host to the device. At the end of the kernel, the values of array C are copied from the device to the host. Putting the whole iterations loop inside a target data construct with map clauses reduced the runtime to 0.012757 seconds, because the transfers occur once at the launch of the first kernel in the iterations loop, and again after the last kernel in that loop.

The omp_target_alloc and omp_target_alloc_device versions have the best performance (0.002501 and 0.002499 seconds, respectively). In these versions, storage for A, B, and C is allocated directly in device memory, so accesses on the device happen from device-local memory. This is a useful model for applications that use scratch arrays on the device side. These arrays never need to be accessed on the host. In such cases, the recommendation is to allocate the scratch arrays on the device and not worry about data transfers, as illustrated in this example.

The omp_target_alloc_shared version also performs well, but is somewhat slower (0.012491 seconds). In this version, storage for A, B, and C is allocated in shared memory. So the data can migrate between the host and the device. There is the overhead of migration but, after migration, accesses on the device happen from much faster device-local memory. In this version, the initialization of the arrays happens on the host. At the first kernel launch, the arrays are migrated to the device, and the kernels access the arrays locally on the device. Finally, before the host performs the reduction computation, the entire C array is migrated back to the host.

The omp_target_alloc_host version (0.074412 seconds) takes almost 6x more time than the omp_target_alloc_shared version. This is because data allocated in host memory does not migrate from the host to the device. When the kernel tries to access the data, the data is typically sent over a bus, such as PCI Express, that connects the device to the host. This is slower than accessing local device memory. If the device accesses only a small part of an array infrequently, then that array may be allocated in host memory using omp_target_alloc_host. However, if the array is accessed frequently on the device side, then it should be kept in device memory. Keeping the data in host memory and accessing it over the PCI will degrade performance.

Finally, a note regarding data transfers: The amount of data transferred in the map version can be seen in LIBOMPTARGET_DEBUG=1 output by grepping for "Libomptarget --> Moving". Notice that each launch of the main kernel yields the following data transfers:

$ grep "Libomptarget --> Moving" test_target_map.debug
Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a5fc8b010) -> (tgt:0xff00000000200000)
Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a9fc8d010) -> (tgt:0xff00000020200000)
Libomptarget --> Moving 536870912 bytes (hst:0x00007f1a7fc8c010) -> (tgt:0xff00000040200000)
Libomptarget --> Moving 536870912 bytes (tgt:0xff00000000200000) -> (hst:0x00007f1a5fc8b010)

On the other hand, data transfers in the omp_target_alloc_... versions are handled by a lower layer of the runtime system. So grepping for "Libomptarget --> Moving" in LIBOMPTARGET_DEBUG=1 output for these versions will not show the data transfers that took place.

Fortran Examples

The Fortran version of the example using target data and map clauses is shown below.

    program main
    use iso_fortran_env
    use omp_lib
    implicit none

    integer, parameter :: iterations=100
    integer, parameter :: length=64*1024*1024
    real(kind=REAL64), parameter ::  epsilon=1.D-8
    real(kind=REAL64), allocatable ::  A(:)
    real(kind=REAL64), allocatable ::  B(:)
    real(kind=REAL64), allocatable ::  C(:)
    real(kind=REAL64) :: scalar=3.0
    real(kind=REAL64) :: ar, br, cr, asum
    real(kind=REAL64) :: nstream_time, avgtime
    integer :: err, i, iter

    !
    ! Allocate arrays on the host using plain allocate

    allocate( A(length), stat=err )
    if (err .ne. 0) then
      print *, "Allocation of A returned ", err
      stop 1
    endif

    allocate( B(length), stat=err )
    if (err .ne. 0) then
      print *, "Allocation of B returned ", err
      stop 1
    endif

    allocate( C(length), stat=err )
    if (err .ne. 0) then
      print *, "Allocation of C returned ", err
      stop 1
    endif

    !
    ! Initialize the arrays

    !$omp parallel do
    do i = 1, length
       A(i) = 2.0
       B(i) = 2.0
       C(i) = 0.0
    end do

    !
    ! Perform the computation

    nstream_time = omp_get_wtime()
    !$omp target data  map(to: A, B) map(tofrom: C)

    do iter = 1, iterations
       !$omp target teams distribute parallel do
       do i = 1, length
          C(i) = C(i) + A(i) + scalar * B(i)
       end do
    end do

    !$omp end target data
    nstream_time = omp_get_wtime() - nstream_time

    !
    ! Validate and output results

    ar = 2.0
    br = 2.0
    cr = 0.0
    do iter = 1, iterations
       do i = 1, length
          cr = cr + ar + scalar * br
       end do
    end do

    asum = 0.0
    !$omp parallel do reduction(+:asum)
    do i = 1, length
       asum = asum + abs(C(i))
    end do

    if (abs(cr - asum)/asum > epsilon) then
       write(*,110) "Failed Validation on output array: Expected =", cr, ", Observed =", asum
    else
       avgtime = nstream_time/iterations
       write(*,120) "Solution validates: Checksum =", asum, ", Avg time (s) =",  avgtime
    endif

110 format (A, F20.6, A, F20.6)
120 format (A, F20.6, A, F10.6)

    deallocate(A)
    deallocate(B)
    deallocate(C)

    end program main

The Fortran version of the example using omp_target_alloc_device is shown below. In this example, allocate directives, with the allocator omp_target_device_mem_alloc, are used to allocate arrays A, B, and C on the device. The use_device_addr(A, B, C) clause is used on the target data directive (line 37) to indicate that the arrays have device addresses, and these addresses should be used in the target region.

    use iso_fortran_env
    use omp_lib
    implicit none

    integer, parameter :: iterations=100
    integer, parameter :: length=64*1024*1024
    real(kind=REAL64), parameter ::  epsilon=1.D-8
    real(kind=REAL64), allocatable ::  A(:)
    real(kind=REAL64), allocatable ::  B(:)
    real(kind=REAL64), allocatable ::  C(:)
    real(kind=REAL64) :: scalar=3.0
    real(kind=REAL64) :: ar, br, cr, asum
    real(kind=REAL64) :: nstream_time, avgtime
    integer :: i, iter

    !
    ! Allocate arrays in device memory

    !$omp allocators allocate(allocator(omp_target_device_mem_alloc): A)
    allocate(A(length))

    !$omp allocators allocate(allocator(omp_target_device_mem_alloc): B)
    allocate(B(length))

    !$omp allocators allocate(allocator(omp_target_device_mem_alloc): C)
    allocate(C(length))

    !
    ! Begin target data

    !$omp target data use_device_addr(A, B, C)

    !
    ! Initialize the arrays

    !$omp target teams distribute parallel do
    do i = 1, length
       A(i) = 2.0
       B(i) = 2.0
       C(i) = 0.0
    end do

    !
    ! Perform the computation

    nstream_time = omp_get_wtime()
    do iter = 1, iterations
       !$omp target teams distribute parallel do
       do i = 1, length
          C(i) = C(i) + A(i) + scalar * B(i)
       end do
    end do
    nstream_time = omp_get_wtime() - nstream_time

    !
    ! Validate and output results

    ar = 2.0
    br = 2.0
    cr = 0.0
    do iter = 1, iterations
       do i = 1, length
          cr = cr + ar + scalar * br
       end do
    end do

    asum = 0.0
    !$omp target teams distribute parallel do reduction(+:asum) &
    !$omp map(tofrom: asum)
    do i = 1, length
       asum = asum + abs(C(i))
    end do

    !
    ! End target data

    !$omp end target data

    if (abs(cr - asum)/asum > epsilon) then
       write(*,110) "Failed Validation on output array: Expected =", cr, ", Observed =", asum
    else
       avgtime = nstream_time/iterations
       write(*,120) "Solution validates: Checksum =", asum, ", Avg time (s) =",  avgtime
    endif

110 format (A, F20.6, A, F20.6)
120 format (A, F20.6, A, F10.6)

    deallocate(A)
    deallocate(B)
    deallocate(C)

    end program main