#pragma omp target

Purpose

The omp target directive instructs the compiler to generate a target task, that is, to map variables to a device data environment and to execute the enclosed block of code on that device.

Use the omp target directive to define a target region, which is a block of computation that operates within a distinct data environment and is intended to be offloaded onto a parallel computation device during execution.

Syntax


                          .-+---+------.   
                          | '-,-'      |   
                          V            |   
>>-#--pragma--omp target----+--------+-+-----------------------><
                            '-clause-'


>>-block-------------------------------------------------------><

Parameters

clause is any of the following clauses:

defaultmap(tofrom:scalar)

Changes the default implicit mapping rule from firstprivate to tofrom for scalar variables. For more information about implicit mapping, see the Rules section.

device(exp)

Creates the data environment on the device of ID exp. exp is an integer expression that evaluates to a non-negative integer value less than the value of omp_get_num_devices().

firstprivate(list)

Declares the data variables in list to be private to the target task and shared by every thread team that runs the region. A new item is created for each list item that is referenced by the target task. Each new data variable is initialized with the value of the original variable at the time the target construct is encountered. Data variables in list are separated by commas.

if([target:]exp)

When the if clause is specified and the scalar expression that is represented by exp evaluates to zero, the target region is executed by the host device in the host data environment.

map([[map-type-modifier[,]]map-type:]list)

Specifies the data variables in list to be explicitly mapped from the original variables in the host data environment to the corresponding variables in the device data environment of the device specified by the construct.

The map-type can be to, from, tofrom, or alloc.

If a list item does not exist in the device data environment, a new item is created in the device data environment.
- If map-type is to or tofrom, this new item is initialized with the value of the original list item in list in the host data environment.
- If map-type is from or alloc, the initial value of the list item in the device data environment is undefined.
On exit from the target region, if storage for the list item was created in the device data environment when this construct was first encountered, the list item is deallocated from the device data environment. Furthermore, if map-type is from or tofrom, the original list item is updated with the current value of the corresponding list item in the device data environment before the list item in the device data environment is deallocated.
If the list item exists in the device data environment when the construct is encountered, the allocation count of the item in the device environment changes as follows:
- It is incremented by one at the start of the construct
- It is decremented by one at the end of the construct.

The map-type-modifier is always. If this modifier is specified, the following rules apply:

If map-type is to or tofrom, the value of the original list item is always copied to the device environment, regardless of whether a new item was created in the device data environment for the list item.
If map-type is from or tofrom, the value of the list item is always copied from the device environment to the original list item, regardless of whether the device list item will be deallocated at termination of the construct.

private(list)

Usage

To enable the omp target directive to execute the target region, you must specify the -qoffload and -qsmp options to offload the region to the device environment. If a target region cannot be successfully offloaded to a device, the target region is executed within the host environment.

Rules

Nesting of target regions, either dynamically or statically, is not allowed.

General mapping rules are as follows:

Pointers are mapped as zero-length array sections with zero base for both explicit and implicit mapping.
If a zero-length array section that is derived from a pointer variable is mapped, that variable is initialized with the address of the corresponding storage location on the device. If the corresponding storage does not exist, that is, it has not been mapped before, the pointer variable is initialized to NULL.

The data environment of a target region is defined by the implicit and explicit mapping of variables between the host and device:

Implicit mapping: The compiler determines which variables must be mapped to, from, or both to and from the device data environment. Scalar variables that are not explicitly mapped are implicitly mapped as firstprivate if defaultmap(tofrom:scalar) is not specified.
Explicit mapping: You can use the map clause on the target region to explicitly list variables to be mapped to, from, or both to and from the device data environment.

Examples

int main()
{ 
  int x = 1;
  #pragma omp target map(toform: x)
  x = x + 1 // The copy of x on the device has a value of 2.
  printf("After the target region is executed, x = %d\n", x);
  return 0;
}

The integer x is declared in the host environment, and its initial value is set to 1 on the host. The target region is declared with explicit map type tofrom of x, so the storage for x is allocated on the device and the device copy of x is initialized to 1. Within the target region, the value of the copy of x on the device is incremented by 1. At the end of the target region, x is mapped back to the host environment according to the map type tofrom, and the host prints the value of x to be 2.

Voice your opinion on getting help information

Ask IBM compiler experts a technical question in the IBM XL compilers forum