Question about the device() clause

Discuss the OpenMP 4.0 API Specifications released July 2013 (Read Only)
Forum rules
The OpenMP Forums are now closed to new posts. Please visit Stack Overflow if you are in need of help:
Posts: 1
Joined: Wed Dec 10, 2014 12:21 am
Location: Ioannina, Greece

Question about the device() clause

Post by dimako »

Hello everybody,

although there was an old question posted regarding incorrect values within a device() clause, I have two more questions which are crucial:
Consider a target region within a target data region.
(1) If both directives contain their own device() clause, may they refer to different devices? The specifications impose no restriction but this makes it almost impossible for the compiler to construct the correct data environments.
(2) Assuming that the answer to question (1) is NO, the question still stands: according to the specifications if there is an if() clause in the inner target directive and the condition evaluates to false, the target region will be executed by the host. Thus, the outer target data specifies the data envirnment for a particular device but the inner target will be executed by the host. The question is then what would be the data environment for the host.

Here is an example:

Code: Select all

int i = 0, j;
#pragma omp target data map(tofrom: i) device(D)
  #pragma omp target device(D)   // If there was no device() clause here, would the *default* device be implied??
  #pragma omp target if(0) map(from: j) device(D)
    j = i;
  printf("i = %d, j = %d\n", i, j);
What would be printed? In the second target, what is the data environment for execution by the host? E.g. is variable i refering to the "original" item or to the "corresponding" item (are there "corresponding" items for the host?)?

I would appreciate any clarifications on the above.
Best regards,

Posts: 54
Joined: Fri May 16, 2008 9:27 am

Re: Question about the device() clause

Post by james »

1) separate device clauses refer to separate devices there is no reason to not allow the programmer to send a computation to a a different device than what the surrounding data region is bound to. The spec does not even forbid a data construct for device 1 from being contained inside of a data construct for device 2. Again there is no reason and there may even be legitimate programs that use this feature. Constructing the data environment for a device is simpler to understand if we look at the unstructured data directives see 4.1 version of the spec.

#pragma omp target enter data map(a[0:n]) device ( 1 ) // add the object a[0:n] to the data environment for device 1.

#pragma omp target enter data map(a[0:n] device( 2 ) // add the object a[0:n] to the data environment for device 2.

// at this point the data environments for the two devices contain the same object the compiler and runtime have no problem building this or maintaining.

// here is the problem
#pragma omp target device ( 1 )
a[1] = 1;

#pragma omp target device ( 2 )
a[1] = 2;

// this is a data race since two devices have both modified the object however there was not synchronization between them. Adding a target update from device 1 and to device 2 would synchronize the memories and allow the program to be conformant.

If the programmer writes they get what they ask for:

#pragma omp target data map( to: a[0:n] ) device ( 1 ) // checks and adds a[0:n] to data environment for device ( 1 )
#pragma omp target map( tofrom: a[0:n] ) device ( 2 ) // checks and adds a[0:n] to data environment for device ( 2 )
} // posibly moves data back from device ( 2 ) to host potentially exposing a race if device 1 is not updated before it uses the object.

#pragma omp target map( to: a[0:n] ) device ( 1 ) // checks and finds a[0:n] in the data environment for device ( 1 ) not data motion occurs likely an error but the data environments built by the system are correct.


Simply put there is not restriction because the compiler and runtime can always build the data environment that is requested, however the programmer must ensure that the data environments remain consistent and race free to get a conforming program.

2) NO the answer to (1) is not NO, it is the users responsibility to insure that the if clause on the target data and the target constructs for the same device evaluate to the same value. If they are different the user will get what they asked for which means they are responsible for all required data motion need to get correct results from the program

Posts: 4
Joined: Mon Dec 15, 2014 2:55 am

Re: Question about the device() clause

Post by psxlover »

So in the case of 2) if the "if" clause defaults to false the target area would run as if "i" was implicitly declared as tofrom since there is no "i" in the hosts data environment?