查询默认流时如何理解 cudaStreamQuery 行为？答案

【问题标题】：How to understand cudaStreamQuery behavior when querying default stream?查询默认流时如何理解 cudaStreamQuery 行为？
【发布时间】：2018-12-22 12:45:49
【问题描述】：

从CUDA manual，我知道cudaStreamQuery 在异步流上运行：

查询异步流的完成状态。

我编写了一个简单的程序来测试它在“默认流”上的行为：

#include <cstdlib>
#include <iostream>

#define cudaSafeCall(call)  \
        do {\
            cudaError_t err = call;\
            if (cudaSuccess != err) \
            {\
                std::cerr << "CUDA error in " << __FILE__ << "(" << __LINE__ << "): " \
                    << cudaGetErrorString(err) << '\n';\
                exit(EXIT_FAILURE);\
            }\
        } while(0)

int main(void)
{
  int N = 1<<20;
  float *x, *d_x;
  x = (float*)malloc(N*sizeof(float));

  cudaSafeCall(cudaMalloc(&d_x, N*sizeof(float))); 

  cudaSafeCall(cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyDefault));

  cudaSafeCall(cudaStreamQuery(0));

  cudaSafeCall(cudaFree(d_x));

  free(x);

  return 0;
}

构建并运行它：

$ nvcc sync_test.cu -o sync_test
$ ./sync_test
CUDA error in sync_test.cu(25): device not ready

在cudaMemcpy和cudaStreamQuery之间加上cudaStreamSynchronize(0);后，错误消失了：

cudaSafeCall(cudaMemcpy(d_x, x, N*sizeof(float), cudaMemcpyDefault));
cudaSafeCall(cudaStreamSynchronize(0));
cudaSafeCall(cudaStreamQuery(0));

我有点困惑：
a) cudaMemcpy() 是与主机码同步的，所以cudaMemcpy 返回后，该流中的复制工作应该完成，为什么cudaStreamQuery(0) 在第一种情况下返回错误？
b) 如果cudaStreamQuery 仅适用于异步流，而不适用于默认流，为什么cudaStreamQuery 在第二种情况下不返回错误？

【问题讨论】：

标签： cuda

【解决方案1】：

对于示例代码中的特定情况，您的陈述 a) 并不十分准确：

a)cudaMemcpy()与宿主码同步，所以cudaMemcpy返回后，该流中的复制工作应该完成，

cudaMemcpy 的 documentation 声明：

此函数在大多数用例中表现出synchronous 行为。

如果我们点击synchronous 链接，我们将被带到一个进一步定义确切行为的页面。您可能想阅读整页，它不是很长。我从同步的定义中摘录两点：

对于从可分页主机内存到设备内存的传输，在启动复制之前会执行流同步。一旦可分页缓冲区被复制到暂存内存以便 DMA 传输到设备内存，该函数将返回，但到最终目的地的 DMA 可能尚未完成。（强调）

和：

对于从设备到可分页或固定主机内存的传输，该函数仅在复制完成后返回。

您的代码适合第一个摘录（上面的第 2 项）。该传输类型（主机到设备）可以在复制完全完成之前返回。

如果您在cudaMemcpy 示例中颠倒了指针的顺序，从而颠倒了副本的方向，您将陷入第二段摘录（上面的第 4 项）。在这种情况下，可以保证传输完成，并且错误报告确实会从您的代码中消失。

这种行为是可能的，并在文档中说明。

【讨论】：