oneAPI GPU Optimization Guide

Developer Guide

Contents

Clauses: is_device_ptr, use_device_ptr, has_device_addr, use_device_addr

The OpenMP clauses
is_device_ptr
,
use_device_ptr
,
has_device_addr
, and
use_device_addr
 can be used to convey information about variables referenced in
target
,
target data
, or
dispatch
 constructs. These clauses are described as follows.

is_device_ptr

The
is_device_ptr
 clause appears on a
target
 or
dispatch
 directive. It indicates that the list items are device pointers. So each list item is privatized inside the construct and the new list item is initialized to the device address to which the original list item refers.
  • In C, each list item should be of type pointer or array.
  • In C++, each list item should be of type pointer, array, reference to pointer, or reference to array.
  • In Fortran, each list item should be of type C_PTR.
The following C/C++ example illustrates the use of the
is_device_ptr
 clause. The
omp_target_alloc_device
 routine allocates memory on the device and returns a device pointer for that memory which is saved in the host variable
arr_device
. On the
target
 directive, we use the
is_device_ptr(arr_device)
 clause to indicate that
arr_device
 points to device memory. So inside the
target
 construct
arr_device
 is privatized and initialized to the device address to which
arr_device
 refers.
//==============================================================
// Copyright © 2022 Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
// clang-format off
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define N 100

int main(void)
{
  int *arr_host = NULL;
  int *arr_device = NULL;

  arr_host = (int *) malloc(N * sizeof(int));
  arr_device = (int *) omp_target_alloc_device(N * sizeof(int),
					       omp_get_default_device());

  #pragma omp target is_device_ptr(arr_device) map(from: arr_host[0:N])
  {
    for (int i = 0; i < N; ++i) {
      arr_device[i] = i;
      arr_host[i] = arr_device[i];
    }
  }

  printf ("%d, %d, %d \n", arr_host[0], arr_host[N/2], arr_host[N-1]);
}

use_device_ptr

The
use_device_ptr
 clause appears on a
target data
 directive. It indicates that each list item is a pointer to an object that has corresponding storage on the device or is accessible on the device.
If a list item is a pointer to an object that is mapped to the device, then references to the list item in the construct are converted to references to a device pointer that is local to the construct and that refers to the device address of the corresponding object.
If the list item does not point to a mapped object, it must contain a valid device address, and the list item references are converted to references to a local device pointer that refers to this device address.
Each list item must be a pointer for which the value is the address of an object that has corresponding storage in the device data environment or is accessible on the target device.
In C, each list item should be of type pointer or array.
In C++, each list item should be of type pointer, array, reference to pointer, or reference to array.
In Fortran, each list item should be of type C_PTR.
The following C/C++ example illustrates the use of the
use_device_ptr
 clause. The
omp_target_alloc_device
 routine is called three times to allocate memory on the device. The addresses of the memory allocated is saved in the pointer variables
A
,
B
, and
C
 on the host. We use the
use_device_ptr(A, B,C)
 clause on the
target data
 directive to indicate that
A
,
B
, and
C
 contain valid device addresses.
//==============================================================
// Copyright © 2022 Intel Corporation
//
// SPDX-License-Identifier: MIT
// =============================================================
// clang-format off
#include <stdio.h>
#include <stdlib.h>
#include <stdint.h>
#include <math.h>
#include <omp.h>

#define length 65536

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 ar;
  double br;
  double cr;
  double asum;

  // 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);
  }

  #pragma omp target data use_device_ptr(A,B,C)
  {
      // Initialize the arrays

      #pragma omp target teams distribute 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

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

      // Validate and output results

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

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

  } // end target data

  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. Checksum = %lf\n", asum);
  }

  return 0;
}

has_device_addr

The
has_device_addr
 clause appears on a
target
 directive. It indicates that the list items already have valid device addresses, and therefore may be directly accessed from the device.
Each list item must have a valid device address for the device data environment. It can be on any type, including an array section.
The
has_device_addr
 clause is especially useful in Fortran, because it can be used with list items of any type (not just C_PTR) to indicate that the list items have device addresses.
The following Fortran example illustrates the use of the
has_device_addr
 clause. In the example, the three arrays
A
,
B
, and
C
 are allocated on the device. When the arrays are referenced in a
target
 region, we use the
has_device_addr(A, B, C)
 clause to indicate that
A
,
B
, and
C
 already have device addresses.
  !=============================================================
  ! Copyright © 2022 Intel Corporation
  !
  ! SPDX-License-Identifier: MIT
  !=============================================================
program main
  use iso_fortran_env
  use omp_lib
  implicit none

  integer, parameter :: iterations=1000
  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 allocate allocator(omp_target_device_mem_alloc)
  allocate(A(length))

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

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

  !
  ! Initialize the arrays

  !$omp target teams distribute parallel do has_device_addr(A, B, C)
  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 has_device_addr(A, B, C)
     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) has_device_addr(C)
  do i = 1, length
     asum = asum + abs(C(i))
  end do

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

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

end program main

use_device_addr

The
use_device_addr
 clause appears on a
target data
 directive. It indicates that each list item already has corresponding storage on the device or is accessible on the device.
If a list item is mapped, then references to the list item in the construct are converted to references to the corresponding list item. If a list item is not mapped, it is assumed to be accessible on the device.
A list item may be an array section.
Just like
has_device_addr
, the
use_device_addr
 clause is especially useful in Fortran, because it can be used with list items of any type (not just C_PTR) to indicate that the list items have device addresses.
The following Fortran example illustrates the use of the
use_device_addr
 clause. In the example,
array_d
 is mapped to the device with the
alloc
 map-type, so storage is allocated for
array_d
 on the device and no data transfer between the host and the device occurs. We use the
use_device_addr(array_d)
 clause on the
target data
 directive to indicate that
array_d
 has corresponding storage on the device.
  !=============================================================
  ! Copyright © 2022 Intel Corporation
  !
  ! SPDX-License-Identifier: MIT
  !=============================================================
program target_use_device_addr

  use omp_lib
  use iso_fortran_env, only : real64
  implicit none

  integer, parameter :: N1 = 1024
  real(kind=real64), parameter :: aval = real(42, real64)
  real(kind=real64), allocatable :: array_d(:), array_h(:)
  integer :: i,err

  ! Allocate host data
  allocate(array_h(N1))

  !$omp target data map (from:array_h(1:N1)) map(alloc:array_d(1:N1))
  !$omp target data use_device_addr(array_d)
  !$omp target
      do i=1, N1
        array_d(i) = aval
        array_h(i) = array_d(i)
     end do
  !$omp end target
  !$omp end target data
  !$omp end target data

  ! Check result
  write (*,*) array_h(1), array_h(N1)
  if (any(array_h /= aval)) then
    err = 1
  else
    err = 0
  end if

  deallocate(array_h)
  if (err == 1) then
    stop 1
  else
    stop 0
  end if

end program target_use_device_addr
The following table summarizes the properties of the clauses described in this section.
Clause
On which directive
Type of list item
Description
is_device_ptr
target, dispatch


C/C++: Pointer,
array, or reference
Fortran: C_PTR

Indicates that list item is a device pointer (has valid device address).
use_device_ptr
target data


C/C++: Pointer,
array, or reference
Fortran: C_PTR

Indicates that list item is a pointer to an object that has corresponding storage on device or is accessible on device.
has_device_addr
target
Any type (may be array section)
Indicates that list item has a valid device address.
use_device_addr
target data
Any type (may be array section)
Indicates that list item has corresponding storage on device or is accessible on the device.

Product and Performance Information

1

Performance varies by use, configuration and other factors. Learn more at www.Intel.com/PerformanceIndex.