#pragma omp target exit data
Purpose
The omp target exit data directive unmaps variables from a device data environment. The omp target exit data directive can limit the amount of device memory when you use the omp target enter data construct to map items to the device data environment.
Syntax
.-+---+--. | '-,-' | V | >>-#--pragma--omp target exit data----clause-+-----------------><
Parameters
- if([target exit 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 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 from, release, or delete.
- If a list item does not exist in the device data environment,
then this construct does nothing with respect to that list item.
- If map-type is from or release, the allocation count of the item in the device environment is decremented by one.
- If map-type is delete, the item's allocation count in the device data environment is set to zero.
- If this construct deallocates a list item from the device data environment, and the map-type is from, then the values for that list item are copied from the device data environment to the host prior to deallocation of the list item on the device.
- If a list item does not exist in the device data environment,
then this construct does nothing with respect to that list item.
- The map-type-modifier is always. If this modifier is present, the value of the original list item is always assigned the value of the corresponding list item on the device environment if the map-type is from, regardless of whether a new item was created in the device data environment for the list item.
- nowait
- Enables the target exit 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 exit data directive takes effect only if you specify both the -qsmp and -qoffload compiler options.
Rules
The target exit data construct decrements the allocation count of mapped list items in the device data environment. When the allocation count of an item reaches zero, it is deallocated from the device data environment and its value might be copied back to the host 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 exit data map(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 is created on the device. This device storage is also not initialized.
When the first nested target region is encountered, the OpenMP runtime will check whether the 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, 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 exit data construct is encountered. The array storage in the device data environment will be deallocated. This causes the values of the array storage on the device to be copied back to the host because the map-type is from.
Now, only the host has the storage for this array, and the values in the array are the same as those calculated by the offload device.
