Summe 3D-Matrix cuda

Question

Ich muss Berechnungen anstellen: A[x][y] = Summe{von z=0 bis z=n}{B[x][y][z]+C[x][y][z]}, wobei Matrix A die Dimensionen [Höhe][Breite] und Matrix B,C die Dimensionen [Höhe][Breite][n] hat.

Die Werte werden dem Speicher etwa mit der folgenden Formel zugeordnet:

index = 0;
for (z = 0; z<n; ++z)
    for(y = 0; y<width; ++y)
        for(x = 0; x<height; ++x) {
            matrix[index] = value;
            index++;
        }

Q1: Ist dieser Cuda-Kernel in Ordnung?

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;

for(z=0; z<n; z++){
    A[idx*width+idy] += B[idx*width+idy+z*width*height] + C[idx*width+idy+z*width*height];
}

F2: Kann die Berechnung so schneller durchgeführt werden?

idx = blockIdx.x*blockDim.x + threadIdx.x;
idy = blockIdx.y*blockDim.y + threadIdx.y;
idz = blockIdx.z*blockDim.z + threadIdx.z;

int  stride_x = blockDim.x * gridDim.x;
int  stride_y = blockDim.y * gridDim.y;
int  stride_z = blockDim.z * gridDim.z;

while ( idx < height && idy < width && idz < n ) {
    atomicAdd( &(A[idx*width+idy]), B[idx*width+idy+idz*width*height] + C[idx*width+idy+idz*width*height] );
    idx += stride_x;
    idy += stride_y;
    idz += stride_z;
}

Answer 1

Der erste Kernel ist in Ordnung. Aber wir haben den Zugriff auf die Matrix nicht koaliert B y C .

Wie bei der zweiten Kernel-Funktion. Sie haben Daten Rennen, weil nicht nur ein Thread hat eine Fähigkeit, in schreiben A[idx*width+idy] Adresse. Sie benötigen eine zusätzliche Synchronisation wie AttomicAdd

Was die allgemeine Frage betrifft: Ich denke, dass Experimente zeigen, dass es besser ist. Es hängt von den typischen Matrixgrößen ab, die Sie haben. Denken Sie daran, dass die maximale Größe der Thread-Blöcke bei Fermi < 1024 ist und dass Sie bei großen Matrizen viele Thread-Blöcke benötigen. Normalerweise ist es langsamer (wenn man viele Thread-Blöcke hat).

Answer 2

Ganz einfach in ArrayFire :

array A = randu(nx,ny,nz);
array B = sum(A,2); // sum along 3rd dimension
print(B);

Answer 3

Q1: Testen Sie es mit Matrizen, bei denen Sie die Antwort kennen

Bemerkung: Bei der Verwendung sehr großer Matrizen können Probleme auftreten. Verwenden Sie eine while-Schleife mit geeigneten Inkrementen. Cuda by Example ist wie immer das Referenzbuch.

Ein Beispiel für die Implementierung einer verschachtelten Schleife finden Sie hier: Für verschachtelte Schleifen mit CUDA . Dort ist eine while-Schleife implementiert.

marina.k hat Recht, was die Rennbedingungen angeht. Das würde Ansatz eins bevorzugen, da atomare Operationen den Code eher verlangsamen.