#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.

depend(dependence-type:list)

Establishes scheduling dependences between the target task and sibling tasks that share list items. The dependence-type can be in, out, or inout. Data variables in list are separated by commas.

If dependence-type is in or inout, for each list item that has the same storage location as a list item in a depend clause of a sibling task with the out or inout dependence type, a scheduling dependence for the target task on the sibling task is created.
If dependence-type is out or inout, for each list item that has the same storage location as a list item in a depend clause of a sibling task with the in, out, or inout dependence type, a scheduling dependence for the target task on the sibling task is created.

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.

is_device_ptr(list)

Indicates that the data variables in list are device pointers that exist within the device data environment. Supported types are pointers, references to pointers, arrays, and references to arrays.

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.

nowait

Eliminates the implicit barrier so the parent task can make progress even if the target task is not yet completed. By default, an implicit barrier exists at the end of the target construct, which ensures the parent task cannot continue until the target task is completed.

private(list)

reduction(reduction-identifier:list)

Specifies that for each data variable in list, a private copy is created and initialized based on the reduction-identifier. At the end of the region, the original data variable is updated with the values of the private copies using a combiner based on the reduction-identifier.

Scalar variables are supported in list. Items in list are separated by commas. reduction-identifier is one of the following operators: +, -, *, &, |, ^, &&, and ||.

Limitations: The reduction clause on the omp target directive is supported only in the form of combined constructs, such as omp target parallel for or omp target teams distribute parallel for.

Usage

To enable the omp target directive to execute the target region, you must specify the -qsmp and -qoffload 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.

A listed data variable cannot appear in both a data-sharing clause and the map clause on the same target construct.

Examples

int main()
{ 
  int x = 1;
  #pragma omp target map(tofrom: 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