Auswirkung der Verwendung von seitenfähigem Speicher für die asynchrone Speicherkopie?

Question

In CUDA C Best Practices-Leitfaden Version 5.0, Abschnitt 6.1.2 ist es, dass geschrieben:

Im Gegensatz zu cudaMemcpy(), der asynchronen Transfer Version gepinnt Host-Speicher erfordert (siehe pinned Speicher) und Es enthält ein zusätzliches Argument , eine Stream-ID.

Es bedeutet die cudaMemcpyAsync Funktion sollte fehlschlagen, wenn ich einfachen Speicher verwende.

Aber das ist nicht, was passiert ist.

Nur für Testzwecke, habe ich versucht, das folgende Programm:

Kernel:

__global__ void kernel_increment(float* src, float* dst, int n) 
{ 
    int tid = blockIdx.x * blockDim.x + threadIdx.x; 

    if(tid<n) 
     dst[tid] = src[tid] + 1.0f; 
} 

Main:

int main() 
{ 
    float *hPtr1, *hPtr2, *dPtr1, *dPtr2; 

    const int n = 1000; 

    size_t bytes = n * sizeof(float); 

    cudaStream_t str1, str2; 

    hPtr1 = new float[n]; 
    hPtr2 = new float[n]; 

    for(int i=0; i<n; i++) 
     hPtr1[i] = static_cast<float>(i); 

    cudaMalloc<float>(&dPtr1,bytes); 
    cudaMalloc<float>(&dPtr2,bytes); 

    dim3 block(16); 
    dim3 grid((n + block.x - 1)/block.x); 

    cudaStreamCreate(&str1); 
    cudaStreamCreate(&str2); 

    cudaMemcpyAsync(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice,str1); 
    kernel_increment<<<grid,block,0,str2>>>(dPtr1,dPtr2,n); 
    cudaMemcpyAsync(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost,str1); 

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError())); 

    cudaDeviceSynchronize(); 

    printf("Status: %s\n",cudaGetErrorString(cudaGetLastError())); 

    cudaStreamDestroy(str1); 
    cudaStreamDestroy(str2); 

    cudaFree(dPtr1); 
    cudaFree(dPtr2); 

    for(int i=0; i<n; i++) 
     std::cout<<hPtr2[i]<<std::endl; 

    delete[] hPtr1; 
    delete[] hPtr2; 

    return 0; 
} 

Das Programm gab korrekte Ausgabe. Das Array wurde erfolgreich inkrementiert.

Wie wurde cudaMemcpyAsync ohne Seitensperre ausgeführt? Fehle ich hier etwas?

Answer 1

cudaMemcpyAsync ist im Grunde eine asynchrone Version von cudaMemcpy. Dies bedeutet, dass der aufrufende Host-Thread bei der Ausgabe des Kopieraufrufs nicht blockiert wird. Das ist das grundlegende Verhalten des Anrufs.

Optional, wenn der Anruf in die nicht Standard-Stream gestartet wird, und wenn der Host-Speicher eine festgelegte Zuordnung ist, und das Gerät verfügt über eine kostenlose Kopie Motor DMA, kann der Kopiervorgang passieren, während die GPU führt gleichzeitig eine andere Operation: entweder Kernel-Ausführung oder eine andere Kopie (im Fall einer GPU mit zwei DMA-Kopier-Engines). Wenn alle diese Bedingungen nicht erfüllt sind, ist die Operation auf der GPU funktionell identisch mit einem Standard cudaMemcpy Aufruf, dh. es serialisiert Operationen auf der GPU, und es kann keine gleichzeitige Kopie-Kernel-Ausführung oder gleichzeitige Mehrfachkopien auftreten. Der einzige Unterschied besteht darin, dass die Operation den aufrufenden Host-Thread nicht blockiert.

In Ihrem Beispielcode sind der Host-Quell- und Zielspeicher nicht angeheftet. Die Speicherübertragung kann sich also nicht mit der Kernel-Ausführung überschneiden (dh sie serialisiert Operationen auf der GPU). Die Aufrufe sind immer noch asynchron auf dem Host. Also, was Sie haben, ist funktional äquivalent zu:

cudaMemcpy(dPtr1,hPtr1,bytes,cudaMemcpyHostToDevice); 
kernel_increment<<<grid,block>>>(dPtr1,dPtr2,n); 
cudaMemcpy(hPtr2,dPtr2,bytes,cudaMemcpyDeviceToHost); 

mit der Ausnahme, dass alle Anrufe asynchron auf dem Host, so dass die Host-Thread blockiert am cudaDeviceSynchronize() Anruf anstatt an jeden der Speichertransfer Anrufe.

Dies ist absolut zu erwarten.