奇怪的cudaMemcpyAsync同步行爲

Question

以下代碼用於測試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 ] 

不知道爲什麼...

Answer 1

有（至少）兩件事情在這裏。

你的第一個觀察是：

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），Profiler會顯示cudaMemcpy操作所需的實際「實際」時間，報告爲gputime以及cpu上的「視在時間」，即將啓動操作所需的cpu時間報告爲cputime。請注意，cputime與gputime相比非常小，即它幾乎是瞬間的，所以基於主機的時序方法報告零時間。但他們實際上並非零時間完成，而且剖析器報告了這一點。

如果您使用cudaEvent計時方法，您會看到不同的結果，當然這會更接近您的分析器gputime結果。