#pragma omp target enter data

Purpose

The omp target enter data directive maps variables to a device data environment. The omp target enter data directive can reduce data copies to and from the offloading device when multiple target regions are using the same data, and when the lexical scope requirement of the omp target data construct is not appropriate for the application.

Syntax

Note: You must specify at least one map clause.

Parameters

clause is any of the following clauses:

if([target enter data:] scalar-expression)

When an if clause is present and the if clause expression evaluates to zero, the device is the host.

At most one if clause can appear on the directive.

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 generated 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 generated task on the sibling task is created.

device(integer-expression)

Creates the data environment on the device with the designated ID. The integer-expression must evaluate to a non-negative integer value that is less than the value of omp_get_num_devices().

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 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, this new item is initialized with the value of the original list item in list in the host data environment.
- If map-type is alloc, the initial value of the list item in the device data environment is undefined.
If the list item was already present in the device data environment when the construct is encountered, the allocation count of the item in the device environment is decremented by one.

The map-type-modifier is always. If this modifier is present, the value of the original list item is always copied to the device environment if the map-type is to, regardless of whether a new item was created in the device data environment for the list item.

nowait

Enables the target enter data construct to perform asynchronously with respect to the encountering thread. By default, the encountering thread must wait for the completion of the construct.

Usage

The omp target enter data directive takes effect only if you specify both the -qsmp and -qoffload compiler options.

Rules

The target enter data construct adds mapped list items to the device data environment.

Example

double *array = (double*)malloc(sizeof(double)*N);
#pragma omp target enter data map(alloc: array[0:N])
#pragma omp target map(tofrom: array[0:N])
{  
  #pragma omp target map(tofrom: array[0:N])
  {
     for (int i=0; i<N; i++)array[i] = double(i) / N;
  }
  #pragma omp target map(tofrom: array[0:N])
  {
     for (int i=0; i<N; i++) array[i] = 1.0 / array[i];
  }
}
#pragma omp target update from(array[0:N])
for (int i=0; i<N; i++) sum += array[i];

At first, storage for an array is allocated in the host environment, but this storage is not initialized. When the target enter data construct is encountered, corresponding storage for the array on the device is created. This device storage is also not initialized.

When the first nested target region is encountered, the OpenMP runtime will check whether storage corresponding to the array already exists on the device. No further action, with respect to the device storage, is taken because the corresponding storage already exists on the device when the target region is encountered. The device storage of the array is then initialized, and the target region completes. Upon completion, the OpenMP runtime will recognize that the storage of the array should remain on the device, and so no copy-back from device to host will occur. At this point, the array storage on the host remains uninitialized, and the array storage on the device is initialized.

The same happens when the second nested target region is encountered. No copies between host and device will occur because the storage is already present on the device when the target construct is encountered, and because that storage is to remain on the device after completion of the target region. After execution of this target region is completed, the array storage on the host remains uninitialized and the array storage on the device has been altered.

Finally, the target update construct is encountered. This causes the values of the array storage on the device to be copied back to the host. Now, both host and device have the storage for the array, and both of these copies contain the same values.