cudaStreamDestroy()不同步/阻塞？

Question

我在一台Win7 x64机器上使用了CUDA 4.2的象限NVS 295。我从“数据自动化系统C方案编制手册”中读到：

"...Streams是通过调用cudaStreamDestroy()来释放的。

for (int i = 0; i < 2; ++i)
cudaStreamDestroy(stream[i]);

cudaStreamDestroy()在销毁流并将控制返回到主机线程之前，等待给定流中所有前面的命令完成。“

这是真的吗？我编写了一个小代码，其中我或多或少地做了以下工作(我只放伪代码)：

//transfer input buffer to device
cudaMemcpyToArrayAsync( ... , stream[1]);

//launch kernel
my_kernel <<<dimGrid, dimBlock, 0, stream[1]>>> (...);

//transfer from device to host
cudaMemcpyAsync(.., cudaMemcpyDeviceToHost, stream[1]);

//Destroy stream. In theory this should block the host until everything on the stream is completed!
ret = cudaStreamDestroy(stream[1]); 

在这个例子中，cudaStreamDestroy()调用似乎立即返回到主机，即不等待cudaMemcpyAsync()调用和其他strem指令完成。如果我输入一个"cudaStreamSynchronize(stream1);“在破坏流之前调用，一切都进行得很顺利，但速度较慢。我做错什么了？

非常感谢您的回复！

Answer 1

我不确定您正在查看的文档的版本，但它与我的不一样。我的CUDA 4.2文档说：

破坏并清理流指定的异步流。

如果当调用cudaStreamDestroy()时，设备仍在流流中工作，则该函数将立即返回，一旦设备完成流中的所有工作，与流相关的资源将自动释放。

根据我的经验，这正是它所做的。驱动程序等待直到流是空的，并销毁它。但是cudaStreamDestroy不会阻止调用线程。

您可以通过运行以下示例来确认这一点：

#include <stdio.h>
#include <assert.h>
#include <unistd.h>

__global__ void kernel(int * inout, const int N)
{
    int gid = threadIdx.x + blockIdx.x * blockDim.x;
    int gstride = gridDim.x * blockDim.x;

   for (; gid < N; gid+= gstride) inout[gid] *= 2;
}

#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true)
{
   if (code != cudaSuccess) 
   {
      fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
      if (abort) exit(code);
   }
}

int main(void)
{
    const int N = 2<<20, sz = N * sizeof(int);

    int * inputs, * outputs, * _inout;

    gpuErrchk( cudaMallocHost((void **)&inputs, sz) );
    gpuErrchk( cudaMallocHost((void **)&outputs, sz) );
    gpuErrchk( cudaMalloc((void **)&_inout, sz) );

    for(int i=0; i<N; i++) { inputs[i] = i; outputs[i] = 0; }

    cudaStream_t stream[2];
    for (int i = 0; i < 2; i++)
        gpuErrchk( cudaStreamCreate(&stream[i]) );

    gpuErrchk( cudaMemcpyAsync(_inout, inputs, sz, cudaMemcpyHostToDevice, stream[1]) );

    kernel<<<128, 128, 0, stream[1]>>>(_inout, N);
    gpuErrchk(cudaPeekAtLastError());

    gpuErrchk( cudaMemcpyAsync(outputs, _inout, sz, cudaMemcpyDeviceToHost, stream[1]) );

    for(int i = 0; i < 2; i++)
        gpuErrchk( cudaStreamDestroy(stream[i]) );

    sleep(1); // remove the sleep and see what happens....

    for(int i = 0; i < N; i++)
        assert( (2 * inputs[i]) == outputs[i] );

    cudaDeviceReset();

    return 0;
}

如果没有sleep()，代码就会失败，因为GPU还没有完成，但是有了它，assert就会通过。请注意，sleep正在做一些与在cudaStreamDestroy调用之前使用显式流同步原语不同的事情，即使结果是相同的。如果流被销毁时不是空的，则结果检查永远不会通过。

Answer 2

CUDA流只是设备任务的执行队列。所有接受流的函数只向队列中添加新任务，而不等待执行结果。cudaStreamDestroy是一个特殊的任务，它意味着流需要被销毁，然后所有以前的设备任务都完成了。词句

"cudaStreamDestroy()等待给定流中所有前面的命令完成，然后销毁流并将控制返回到主机线程。

这意味着在设备代码完成之前无法销毁该流。