1

我在 Win7 x64 机器上的 Quadro NVS 295 上使用 CUDA 4.2。从 CUDA C Programming Manual 我读到：

“...通过调用 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(stream[1]);” 在销毁流之前调用，一切顺利，但速度较慢。那么，我做错了什么？

非常感谢您的回复！

4

2 回答 2

3

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

销毁并清理 stream 指定的异步流。

如果调用 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，即使结果相同。如果流在销毁时不为空，则结果检查永远不会通过。

于 2012-06-11T18:46:43.563 回答
2

CUDA 流只是设备任务的执行队列。所有接受流的函数只将新任务添加到队列中，而不等待执行结果。cudaStreamDestroy是一项特殊任务，这意味着需要销毁流，然后完成所有先前的设备任务。话

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

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

于 2012-06-11T17:46:10.690 回答