How to use CUDA stream in OpenACC

Hello.

I want to use CUDA stream in OpenACC asynchronous kernel.

But how can I do it?

Hi CNJ,

To assign an OpenACC async queue to a particular CUDA stream, use “acc_set_cuda_stream”.

int acc_set_cuda_stream ( int async, void* stream );

Where “async” is the OpenACC queue number and “stream” is the stream created via cudaStreamCreate.

Hope this helps,
Mat

This can also be done with the “async()” clause as well correct?

No, you can’t attach an existing CUDA stream to an async queue via the “async” clause. That states which async queue number to use.

By default, the compiler will create a new CUDA stream for each async queue number.

-Mat

May I hijack this topic to ask a related question?

No, you can’t attach an existing CUDA stream to an async queue via the “async” clause. That states which async queue number to use.

What about a usage case like this one. Is invalid, even If you are providing a correct stream identifier (integer) ?

int a[100];

cudaStream_t stream;
cudaStreamCreate(&stream);
int streamId = 23;
...

acc_set_cuda_stream(streamId, stream);

#pragma acc kernels async(streamId)
for (int i = 0; i < 100; i++)
    a[i] = i;

#pragma acc wait

...

cudaStreamDestroy(stream)

Hi pfarre83876,

Using “acc_set_cuda_stream” as you show above is the correct way of assigning an OpenACC async queue number (i.e. streamID) to an already created CUDA stream.

What would be invalid is trying to set the stream in the “async” clause itself rather than using the API call. I.e. if you tried to use “async(stream)” rather than “async(streamID)”.

-Mat