Hi,
I wrote a code which uses many host (openMP) threads per one GPU. Each thread has its own CUDA stream to order it requests. It looks very similar to below code:
#pragma omp parallel for num_threads(STREAM_NUMBER)
for (int sid = 0; sid < STREAM_NUMBER; sid++) {
cudaStream_t stream;
cudaStreamCreate(&stream);
while (hasJob()) {
//... code to prepare job - dData, hData, dataSize etc
cudaError_t streamStatus = cudaStreamQuery(stream);
if (streamStatus == cudaSuccess) {
cudaMemcpyAsync(dData, hData, dataSize, cudaMemcpyHostToDevice, stream);
doTheJob<<<gridDim, blockDim, smSize, stream>>>(dData, dataSize);
else {
CUDA_CHECK(streamStatus);
}
cudaStreamSynchronize();
}
cudaStreamDestroy(stream);
}
And everything were good till I got many small jobs. In that case, from time to time, cudaStreamQuery returns cudaErrorNotReady, which is for me unexpected because I use cudaStreamSynchronize. Till now I were thinking that cudaStreamQuery will always return cudaSuccess if it is called after cudaStreamSynchronize. Unfortunately it appeared that cudaStreamSynchronize may finish even when cudaStreamQuery still returns cudaErrorNotReady.
I changed the code into fallowing and everything works correctly.
#pragma omp parallel for num_threads(STREAM_NUMBER)
for (int sid = 0; sid < STREAM_NUMBER; sid++) {
cudaStream_t stream;
cudaStreamCreate(&stream);
while (hasJob()) {
//... code to prepare job - dData, hData, dataSize etc
cudaError_t streamStatus;
while ((streamStatus = cudaStreamQuery(stream)) == cudaErrorNotReady) {
cudaStreamSynchronize();
}
if (streamStatus == cudaSuccess) {
cudaMemcpyAsync(dData, hData, dataSize, cudaMemcpyHostToDevice, stream);
doTheJob<<<gridDim, blockDim, smSize, stream>>>(dData, dataSize);
else {
CUDA_CHECK(streamStatus);
}
cudaStreamSynchronize();
}
cudaStreamDestroy(stream);
}
So my question… is it a bug or a feature?
EDIT: I asked the same question here: openmp - CUDA 4.0 RC - many host threads per one GPU - cudaStreamQuery and cudaStreamSynchronize behaviour - Stack Overflow