CUDA 8.0, GTX 1080, warum ist Vektoraddition langsamer als Matrixmultiplikation um 5x?

Question

Ich verwende die neueste CUDA 8.0 mit GTX 1080 und laufe Proben, um die Geschwindigkeit zu testen. (Ich weiß, dass sie nicht die optimale Geschwindigkeit reflektieren, aber ich möchte nur horizontal vergleichen.)

In 0_Simple/matrixMul, die Geschwindigkeit durch den Code gemessen wird, was ergibt:

Performance= 1029.91 GFlop/s, Time= 0.127 msec, Size= 131072000 Ops, WorkgroupSize= 1024 threads/block 

Dann lief ich 0_Simple/vectorAdd, und kopiere den Geschwindigkeitstestcode von der obigen Probe. d.h .:

// Measure speed 
    cudaEvent_t start; 
    cudaEventCreate(&start); 
    cudaEvent_t stop; 
    cudaEventCreate(&stop); 

    cudaEventRecord(start, NULL); 
    int nIter = 300; 
    for (int i = 0; i < nIter; i++) { 
     vectorAdd<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, numElements); 
    } 
    cudaEventRecord(stop, NULL); 

    cudaEventSynchronize(stop); 
    float msecTotal = 0.0f; 
    cudaEventElapsedTime(&msecTotal, start, stop); 
    float msecPerAdd = msecTotal/nIter; 
    double flopsPerAdd = numElements; 
    double gigaFlops = (flopsPerAdd * 1.0e-9f)/(msecPerAdd/1000.0f); 
    printf("Performance= %.2f GFLOPS, Time= %.3f ms, Size= %.0f Ops\n", gigaFlops, msecPerAdd, flopsPerAdd); 

ich vergrößert auch die numElements50000-67108864. Das Geschwindigkeitsergebnis lautet:

Performance= 19.85 GFLOPS, Time= 3.380 ms, Size= 67108864 Ops 

was fast 5x langsamer ist.

Ich weiß, dass Beispielcode möglicherweise suboptimal ist, also könnte mir jemand sagen, warum der vectorAdd-Code so langsam ist, und wie man es optimiert?

ich verwende CUDA 8.0 und GTX 1080

Answer 1

Anders als Matrixmultiplikation, Vektoraddition ist eine Speicherbandbreite gebunden Betrieb. Die korrekte Messung der Leistung besteht in der Messung der Bandbreite des globalen Speicherzugriffs. Für die Vektoraddition enthält es 2 Eingabe- und 1 Ausgabevektoren und kann wie folgt berechnet werden.

Sie könnten es mit der Bandbreite einer einfachen D2D-Kopie vergleichen, um zu sehen, ob Sie den Höchstwert erreicht haben.

Answer 2

Edit: hat Code für GPU und CPU

ich folgendes Experiment durchgeführt

*** Cuda Kernel

__global__ void add(float *a, float *c , size_t N) 
{ 
    int tid = blockIdx.x*blockDim.x + threadIdx.x;  
    if(tid < N) a[tid] += c[tid]; 
} 

*** CPU naive Version (Baseline)

void naiveAdd(float *h_a, float *h_c, size_t N) 
{ 
    for (size_t i=0; i<N; i++) 
     h_a[i] += h_c[i]; 
} 

für den Leistungsvergleich zum Hinzufügen von zwei Vektoren in der gleichen Maschine (740M GPU, 4. Gen i7) mit dif verschiedene Versionen von CUDA GCC.

auf CUDA 8.0 und GCC 5 die GPU-Versionen waren langsamer

auf CUDA 7.5 und 4.8 GCC die GPU-Versionen waren schneller

Mit diesem Beweis kann ich feststellen, dass es besser ist, vermeiden Sie CUDA 8.0 für den Moment.