1

以下代码用于测试 cudaMemcpyAsync 的同步行为。

#include <iostream>
#include <sys/time.h>

#define N 100000000

using namespace std;


int diff_ms(struct timeval t1, struct timeval t2) 
{
    return (((t1.tv_sec - t2.tv_sec) * 1000000) +
            (t1.tv_usec - t2.tv_usec))/1000;
}

double sumall(double *v, int n)
{
    double s=0;
    for (int i=0; i<n; i++) s+=v[i];
    return s;
}


int main()
{
    int i;

    cudaStream_t strm;
    cudaStreamCreate(&strm);

    double *h0;
    double *h1;
    cudaMallocHost(&h0,N*sizeof(double));
    cudaMallocHost(&h1,N*sizeof(double));

    for (i=0; i<N; i++) h0[i]=99./N;
    double *d; 
    cudaMalloc(&d,N*sizeof(double));

    struct timeval t1, t2; gettimeofday(&t1,NULL);
    cudaMemcpyAsync(d,h0,N*sizeof(double),cudaMemcpyHostToDevice,strm);
    gettimeofday(&t2, NULL); printf("cuda H->D %d takes: %d ms\n",i, diff_ms(t2, t1)); gettimeofday(&t1, NULL);
    cudaMemcpyAsync(h1,d,N*sizeof(double),cudaMemcpyDeviceToHost,strm);
    gettimeofday(&t2, NULL); printf("cuda D->H %d takes: %d ms\n",i, diff_ms(t2, t1)); gettimeofday(&t1, NULL);

    cout<<"sum h0: "<<sumall(h0,N)<<endl;
    cout<<"sum h1: "<<sumall(h1,N)<<endl;


    cudaStreamDestroy(strm);
    cudaFree(d);
    cudaFreeHost(h0);
    cudaFreeHost(h1);

    return 0;
}

h0/h1 的打印输出提示 cudaMemcpyAsync 与主机同步

sum h0: 99
sum h1: 99

但是，包含 cudaMemcpyAsync 调用的时间差表明它们与主机不同步

cuda H->D 100000000 takes: 0 ms
cuda D->H 100000000 takes: 0 ms

因为 cuda-profiling 结果不支持这一点：

method=[ memcpyHtoDasync ] gputime=[ 154896.734 ] cputime=[ 17.000 ] 
method=[ memcpyDtoHasync ] gputime=[ 141175.578 ] cputime=[ 6.000 ]

不知道为什么...

4

1 回答 1

5

这里（至少）发生了两件事。

您的第一个观察结果是：

sum h0: 99
sum h1: 99

向同一流发出的CUDA 调用将按顺序执行。如果您希望CUDA 调用相互重叠，则必须将它们发布到单独的流中。由于您在同一流中向设备和从设备发出 cuda memcpy，因此它们将按顺序执行。第二个在第一个完成之前不会开始（即使两者都立即排队）。因此数据是完整的（在第一个 cudaMemcpy 之后），并且您观察到两个数组都生成了正确的总和。

您剩余的观察结果也彼此一致。您正在报告：

cuda H->D 100000000 takes: 0 ms
cuda D->H 100000000 takes: 0 ms

这是因为这两个异步调用都会立即将控制权返回给主机线程，并且调用排队等待与主机执行异步执行。然后调用与进一步的主机执行并行进行。由于控制权立即返回给主机，并且您正在使用基于主机的计时方法对操作进行计时，因此它们似乎花费了零时间。

当然，它们实际上并不需要零时间，您的分析器结果表明了这一点。由于 GPU 与 CPU 异步执行（在这种情况下包括 cudaMemcpyAsync），因此分析器显示了 cudaMemcpy 操作所花费的实际“实时”时间，报告为gputimeCPU 上的“表观时间”，即时间量启动操作所需的 CPU 报告为cputime. 请注意，与 gputime 相比，cputime 非常小，即它几乎是瞬时的，因此基于主机的计时方法报告零时间。但它们实际上并不是零时间完成的，分析器报告了这一点。

如果您使用 cudaEvent 计时方法，您当然会看到不同的结果，这将更接近您的分析器gputime结果。

于 2012-12-22T11:54:40.603 回答