#pragma omp target data
Purpose
The omp target data directive maps variables to a device data environment, and defines the lexical scope of the data environment that is created. The omp target data directive can reduce data copies to and from the offloading device when multiple target regions are using the same data.
Syntax
.-+---+------. | '-,-' | V (1) | >>-#--pragma--omp target data----clause-----+------------------><
- You must specify at least one map or use_device_ptr clause.
Parameters
- if([target 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.
- 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 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.
- 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 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.
- If a list item does not exist in the device data environment,
a new item is created in the device data environment.
- The map-type-modifier is always. If this modifier
is present, 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 the termination of the construct.
- At most one map clause can appear on the directive.
- use_device_ptr(list)
- Converts the data variables in list into device pointers to the corresponding variables in the device data environment. References to variables in list throughout the target data region must only be to the address of those variables. You can use only direct access variables; object members and indirectly accessed variables are not supported by the use_device_ptr clause.
Usage
The omp target data directive takes effect only if you specify both the -qsmp and -qoffload compiler options.
Rules
The target data construct creates a device data environment, for the duration of the execution of its lexical scope by the encountering thread, during which the storage, and values, of the mapped list items are present on the device.
Example
double *array = (double*)malloc(sizeof(double)*N);
// Target data region
#pragma omp target data map(from: array[0:N])
{
// The first target region
#pragma omp target map(tofrom: array[0:N])
{
for (int i=0; i<N; i++)array[i] = double(i) / N;
}
// The second target region
#pragma omp target map(tofrom: array[0:N])
{
for (int i=0; i<N; i++) array[i] = 1.0 / array[i];
}
}
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 data region is encountered, corresponding storage for the array is created on the device. The device storage is also not initialized.
When the first 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, while the array storage on the device is initialized.
The same happens when the second 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 end of the target data lexical scope is encountered. This causes the values of the array storage on the device to be copied back to the host, and the device storage to be deallocated. After the lexical scope of the target data region is executed, only the host storage of the array exists, and it contains the values that were calculated on the device.
In contrast, see the following example that doesn't contain a target data construct.
double *array = (double*)malloc(sizeof(double)*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];
}
for (int i=0; i<N; i++) sum += array[i];- Allocate storage on the device corresponding to the array.
- Copy the values of the array from the host to the device.
- Execute the region on the device.
- Copy the values of the array from the device to the host.
- Deallocate the device storage of the array.
