Asynchronous Queues and Streams

Namespaces
namespace	tinker::g
	Global handles for the GPU runtime libraries.

Variables
int	tinker::g::q0
	Default OpenACC async queue. More...
int	tinker::g::q1
	Default OpenACC sync queue. More...
int	tinker::g::qpme
	OpenACC async queue for PME. More...
bool	tinker::use_pme_stream
	Logical flag for use of a separate CUDA stream for PME. More...
cudaStream_t	tinker::g::s0
	CUDA stream for the default OpenACC async queue. More...
cudaStream_t	tinker::g::s1
	CUDA stream for the default OpenACC sync queue. More...
cublasHandle_t	tinker::g::h0
	CUDA BLAS handle using s0. More...
cublasHandle_t	tinker::g::h1
	CUDA BLAS handle using s1. More...

Detailed Description

Asynchronous Queues and Streams

This is a snippet of the OpenACC code.

#pragma acc parallel loop
for (int i = 0; i < limit; ++i) {
   array[i] = i;
}

By default, the host thread will wait until the parallel loop finishes. If you want the host thread to proceed without waiting for the parallel loop to finish, you may add async,

#pragma acc parallel loop async(queue)
// or #pragma acc parallel loop async
for (int i = 0; i < limit; ++i) {
   array[i] = i;
}

where queue is an optional hardwired integer or an integer variable. A special integer constant value acc_async_sync is defined by the OpenACC standard that can be used in the async directive as a queue number, to achieve an synchronous/blocking behavior. Implementations may be different on different platforms though, on the CUDA platform every OpenACC queue is built on top of a CUDA stream.

Variable Documentation

h0

cublasHandle_t tinker::g::h0

extern

CUDA BLAS handle using s0.

h1

cublasHandle_t tinker::g::h1

extern

CUDA BLAS handle using s1.

q0

int tinker::g::q0

extern

Default OpenACC async queue.

q1

int tinker::g::q1

extern

Default OpenACC sync queue.

qpme

int tinker::g::qpme

extern

OpenACC async queue for PME.

s0

cudaStream_t tinker::g::s0

extern

CUDA stream for the default OpenACC async queue.

s1

cudaStream_t tinker::g::s1

extern

CUDA stream for the default OpenACC sync queue.

use_pme_stream

bool tinker::use_pme_stream

extern

Logical flag for use of a separate CUDA stream for PME.