CUDA Streams Performance

Question

Ich lerne gerade CUDA Streams durch die Berechnung eines Skalarprodukts zwischen zwei Vektoren. Die Bestandteile sind eine Kernfunktion, die Vektoren x und y aufnimmt und einen Vektor Ergebnis der Größe gleich der Anzahl der Blöcke zurückgibt, wo jeder Block seine eigene reduzierte Summe beiträgt.

Ich habe auch eine Host-Funktion dot_gpu, die den Kernel aufruft und reduziert den Vektor Ergebnis zum Endprodukt Wert Punkt.

Die synchrone Version genau dies tut:

// copy to device 
copy_to_device<double>(x_h, x_d, n); 
copy_to_device<double>(y_h, y_d, n); 

// kernel   
double result = dot_gpu(x_d, y_d, n, blockNum, blockSize); 

während die Asynchron geht man wie:

double result[numChunks]; 
for (int i = 0; i < numChunks; i++) { 
    int offset = i * chunkSize; 

    // copy to device 
    copy_to_device_async<double>(x_h+offset, x_d+offset, chunkSize, stream[i]); 
    copy_to_device_async<double>(y_h+offset, y_d+offset, chunkSize, stream[i]); 

    // kernel 
    result[i] = dot_gpu(x_d+offset, y_d+offset, chunkSize, blockNum, blockSize, stream[i]); 
} 
for (int i = 0; i < numChunks; i++) { 
    finalResult += result[i]; 
    cudaStreamDestroy(stream[i]); 
} 

Ich schlechtere Leistung bekommen, wenn Streams mit und versuchte, die Gründe zu untersuchen. Ich habe versucht, die Downloads, Kernel-Aufrufe und Uploads zu pumpen, aber ohne Ergebnisse.

// accumulate the result of each block into a single value 
double dot_gpu(const double *x, const double* y, int n, int blockNum, int blockSize, cudaStream_t stream=NULL) 
{ 
double* result = malloc_device<double>(blockNum); 
dot_gpu_kernel<<<blockNum, blockSize, blockSize * sizeof(double), stream>>>(x, y, result, n); 

#if ASYNC 
    double* r = malloc_host_pinned<double>(blockNum); 
    copy_to_host_async<double>(result, r, blockNum, stream); 

    CudaEvent copyResult; 
    copyResult.record(stream); 
    copyResult.wait(); 
#else 
    double* r = malloc_host<double>(blockNum); 
    copy_to_host<double>(result, r, blockNum); 
#endif 

double dotProduct = 0.0; 
for (int i = 0; i < blockNum; i ++) { 
    dotProduct += r[i]; 
} 

cudaFree(result); 
#if ASYNC 
    cudaFreeHost(r); 
#else 
    free(r); 
#endif 

return dotProduct; 
} 

Meine Vermutung ist, dass das Problem innerhalb der ist dot_gpu() Funktionen, die nicht nur den Kernel nicht nennen. Sagen Sie mir, wenn ich das richtig die folgenden Stream Ausführungen verstehen

foreach stream { 
    cudaMemcpyAsync(device[stream], host[stream], ... stream); 
    LaunchKernel<<<...stream>>>(...); 
    cudaMemcpyAsync(host[stream], device[stream], ... stream); 
} 

Der Host führt alle drei Befehle ohne blockiert zu werden, da cudaMemcpyAsync und Kernel-Rückkehr sofort (aber auf der GPU diese ausgeführt werden sequentiell als sie derselben zugeordnet sind, Strom). Der Host geht also zum nächsten Stream (auch wenn stream1 weiß, in welchem ​​Stadium es ist, aber wen interessiert das. Er macht seinen Job auf der GPU, oder?) Und führt die drei Befehle erneut aus, ohne blockiert zu werden. Und so weiter und so weiter. Allerdings blockiert mein Code den Host, bevor er den nächsten Stream irgendwo innerhalb der dot_gpu() Funktion verarbeiten kann. Liegt es daran, dass ich & Sachen freigebe, sowie das vom Kernel zurückgegebene Array auf einen einzigen Wert reduziere?

Answer 1

Ihre objektiviert CUDA-Schnittstelle Unter der Annahme, das tut, was die Funktion und Methode Namen vermuten lassen, gibt es drei Gründe, warum Arbeit von nachfolgenden Aufrufen dot_gpu() möglicherweise nicht überlappen:

1. Ihr Code explizit Blöcke durch eine Ereignisaufzeichnung und wartet dafür.

2. Wenn es nicht bereits für 1. blockieren würde, würde Ihr Code block on the pinned host side allocation and deallocation, wie Sie vermuteten.

3. Wenn Ihr Code bereits für 2 nicht blockiert war, kann die Arbeit von nachfolgenden Aufrufen an dot_gpu() sich immer noch nicht überlappen, je nach Rechenleistung. , auch wenn sie in verschiedene Streams eingereiht sind.

   Auch für Geräte der Rechenfähigkeit und höhere 3,5 the number of streams whose operations can be reordered is limited by the CUDA_​DEVICE_​MAX_​CONNECTIONS environment variable, die standardmäßig bis 8 und kann auf Werte so groß wie 32

eingestellt werden