12.10. Target Offload#

In the OpenMP 5.0 implementation the OMP_TARGET_OFFLOAD environment variable was defined to change default offload behavior. By default the target code (region) is executed on the host if the target device does not exist or the implementation does not support the target device.

In an OpenMP 5.0 compliant implementation, setting the OMP_TARGET_OFFLOAD variable to MANDATORY will force the program to terminate execution when a target construct is encountered and the target device is not supported or is not available. With a value DEFAULT the target region will execute on a device if the device exists and is supported by the implementation, otherwise it will execute on the host. Support for the DISABLED value is optional; when it is supported the behavior is as if only the host device exists (other devices are considered non-existent to the runtime), and target regions are executed on the host.

The following example reports execution behavior for different values of the OMP_TARGET_OFFLOAD variable. A handy routine for extracting the OMP_TARGET_OFFLOAD environment variable value is deployed here, because the OpenMP API does not have a routine for obtaining the value.

Note: The example issues a warning when a pre-5.0 implementation is used, indicating that the OMP_TARGET_OFFLOAD is ignored. The value of the OMP_TARGET_OFFLOAD variable is reported when the OMP_DISPLAY_ENV environment variable is set to TRUE or VERBOSE.

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

/*
* name: target_offload_control.1
* type: C
* version: omp_5.0
*/
#include    <omp.h>
#include  <stdio.h>
#include  <ctype.h>
#include <stdlib.h>
#include <string.h>

typedef enum offload_policy
{MANDATORY, DISABLED, DEFAULT, UNKNOWN, NOTSET} offload_policy_t;


offload_policy_t get_offload_policy()
{
   char *env, *end;
   size_t n;

   env = getenv("OMP_TARGET_OFFLOAD");
   if(env  == NULL) return NOTSET;

   end = env + strlen(env);                //Find trimmed beginning/end
   while (      *env && isspace(*(env  )) ) env++;
   while (end != env && isspace(*(end-1)) ) end--;
   n = (int)(end - env);

                    //Find ONLY string -nothing more, case insensitive
   if      (n == 9 && !strncasecmp(env, "MANDATORY",n)) return MANDATORY;
   else if (n == 8 && !strncasecmp(env, "DISABLED" ,n)) return DISABLED ;
   else if (n == 7 && !strncasecmp(env, "DEFAULT"  ,n)) return DEFAULT  ;
   else                                                 return UNKNOWN  ;
}


int main()
{
   int i;
   int device_num, on_init_dev;

   // get policy from OMP_TARGET_OFFLOAD variable
   offload_policy_t policy = get_offload_policy();

   if(_OPENMP< 201811)
   {
      printf("Warning: OMP_TARGET_OFFLOAD NOT supported, version %d\n",
             _OPENMP );
      printf("         If OMP_TARGET_OFFLOAD is set, "
             "it will be ignored.\n");
   }

   // Set target device number to an unavailable
   // device to test offload policy.
   device_num = omp_get_num_devices() + 1;

   // Policy:
   printf("OMP_TARGET_OFFLOAD Policy:  ");
   if     (policy==MANDATORY)
      printf("MANDATORY-Terminate if dev. not avail\n");
   else if(policy==DISABLED )
      printf("DISABLED -(if supported) Only on Host\n");
   else if(policy==DEFAULT  )
      printf("DEFAULT  -On host if device not avail\n");
   else if(policy==UNKNOWN  )
      printf("OMP_TARGET_OFFLOAD has unknown value\n" );
   else if(policy==NOTSET   )
      printf("OMP_TARGET_OFFLOAD not set\n" );


   on_init_dev = 1;
   // device# out of range--not supported
   #pragma omp target device(device_num) map(tofrom: on_init_dev)
     on_init_dev=omp_is_initial_device();

   if (policy == MANDATORY && _OPENMP >= 201811)
      printf("ERROR: OpenMP implementation ignored MANDATORY policy.\n");

   printf("Target region executed on init dev %s\n",
          on_init_dev ? "TRUE":"FALSE");

   return 0;
}
!!%compiler: gfortran
!!%cflags: -fopenmp

! name: target_offload_control.1
! type: F-free
! version: omp_5.0
module offload_policy
  implicit none
  integer, parameter :: LEN_POLICY=10
contains
  character(LEN_POLICY) function get_offload_policy()
     character(64) :: env
     integer       :: length, i
     env=repeat(' ',len(env))
                                    !policy is blank if not found *
     call get_environment_variable("OMP_TARGET_OFFLOAD",env,length)

     do i = 1,len(env)              !Makes a-z upper case
        if(iachar(env(i:i))>96) env(i:i)=achar(iachar(env(i:i))-32)
     end do

     get_offload_policy = trim(adjustl(env)) !remove peripheral spaces

     if(length==0) get_offload_policy="NOTSET"

     return

  end function

end module

program policy_test

  use omp_lib
  use offload_policy

  integer                :: i, device_num
  logical                :: on_init_dev
  character(LEN_POLICY)  :: policy

  policy = get_offload_policy() !!Get OMP_TARGET_OFFLOAD value

  if (OPENMP_VERSION < 201811) then
     print*,"Warning: OMP_TARGET_OFFLOAD NOT supported by VER.", &
            OPENMP_VERSION
     print*,"         If OMP_TARGET_OFFLOAD is set, it will be ignored."
  endif

     ! Set target device number to an unavailable device
     ! to test offload policy.
  device_num = omp_get_num_devices() + 1

                      !! Report OMP_TARGET_OFFOAD value
  select CASE (policy)
     case("MANDATORY")
          print*,"Policy:  MANDATORY-Terminate if dev. not avail."
     case("DISABLED")
          print*,"Policy:  DISABLED-(if supported) Only on Host."
     case("DEFAULT")
          print*,"Policy:  DEFAULT On host if device not avail."
     case("NOTSET")
          print*,"         OMP_TARGET_OFFLOAD is not set."
     case DEFAULT
          print*,"         OMP_TARGET_OFFLOAD has unknown value."
          print*,"         UPPER CASE VALUE=",policy
  end select


  on_init_dev = .FALSE.
                      !! device# out of range--not supported
  !$omp target  device(device_num) map(tofrom: on_init_dev)
     on_init_dev=omp_is_initial_device()
  !$omp end target

  if (policy=="MANDATORY" .and. OPENMP_VERSION>=201811) then
     print*,"OMP ERROR: ", &
            "OpenMP 5.0 implementation ignored MANDATORY policy."
     print*,"           Termination should have occurred", &
            " at target directive."
  endif

  print*, "Target executed on init dev (T|F): ", on_init_dev

end program policy_test