|
| View previous topic :: View next topic |
| Author |
Message |
xray
Joined: 21 Jan 2010
Posts: 71
|
 Posted: Fri Mar 01, 2013 3:19 am Post subject: Overlapping kernel and data execution |
 |
|
Hi,
with the OpenACC async clauses, I can execute kernels and update data asynchronously. Now, I want to use different streams (=integer expressions).
1) I have heard that using integer=0 as argument for the async clause, PGI interprets it as synchronous behavior (only integer values > 0 are asynchronous). Is that true?
2) With CUDA streams, it is said that if a call blocks, it blocks all other calls of the same type behind it (even in other streams) - where the call is either of type kernel or memcopy. I assume that is a hardware issue. So, is it also true for OpenACC?
Example code:
| Code: | #pragma acc kernels async(1)
...
#pragma acc update async(1)
...
#pragma acc update async(2) |
Issuing these operations, does it mean that the second update can only be executed after the first update? The reason would be: stream 1 is executed first since an operations was first issued here (and not in stream 2). So, the kernel execution will start, afterwards the first update. And since the second update has to wait until the first update did finish, it will be executed at last. In this case, everything actually serialized.
Second example
| Code: | #pragma acc kernels async(1)
...
#pragma acc update async(1)
...
#pragma acc kernels async(2)
...
#pragma acc update async(2) |
If I understand it right, here, only kernel execution in stream 2 and update in stream 1 could be really overlapped. Correct?
Bye, Sandra |
|
| Back to top |
|
 |
xray
Joined: 21 Jan 2010
Posts: 71
|
| Posted: Tue Mar 05, 2013 3:08 am Post subject: |
|
|
| Any news on the issue (for Nvidia GPUs)? |
|
| Back to top |
|
|
mkcolg
Joined: 30 Jun 2004
Posts: 4996
Location: The Portland Group Inc.
|
| Posted: Fri Mar 08, 2013 3:28 pm Post subject: |
|
|
Hi Sandra,
Sorry for the late reply. I needed some clarification from engineering. Here's their response:
| Quote: | | In PGI 12.x, async(0) is interpreted as synchronous. This is in conflict with OpenACC V2.0, so PGI 13.x uses async(-1) to mean synchronous. Any other integer value is allowed, including negative values. The current OpenACC runtime only uses 8 CUDA streams, so the values are mapped down to between 0:7. The mapping is a little more complex than taking just the lower 3 bits, but values that differ in only the lower 3 bits will get different streams. |
As for your second question, the OpenACC 1.0 Spec states:
| Quote: | | Two asynchronous activities with the same argument value will be executed on the device in the order they are encountered by the host process. Two asynchronous activities with different handle values may be executed on the device in any order relative to each other. If there are two or more host threads executing and sharing the same accelerator device, two asynchronous activities with the same argument value will execute on the device one after the other, though the relative order is not determined. |
One thing to add is that "async" is asynchronous to the host code and doesn't block until a "wait" directive is encountered. Hence, in the case of:
| Code: | #pragma acc kernels async(1)
...
#pragma acc update async(1)
...
#pragma acc update async(2) |
The second update could be executed before the first.
For the second example, the only guarantee is that the stream 1's kernel will launch before stream 1's update and stream 2's kernel will launch before stream 2's update.
Hope this helps,
Mat |
|
| Back to top |
|
|
|
|
You cannot post new topics in this forum
You cannot reply to topics in this forum
You cannot edit your posts in this forum
You cannot delete your posts in this forum
You cannot vote in polls in this forum
|
Powered by phpBB © 2001, 2002 phpBB Group
|
Search
Memberlist
Register
Profile
Log in
