Target Offload
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