Summe 3D-Matrix cuda

Question

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

Werte in dem Speicher mit so etwas wie abgebildet werden:

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 das Cuda Kernel ok?

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]; 
} 

Q2: Ist dies der schnellere Weg, um die Berechnung durchzuführen?

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

Q1: Test mit Matrizen, wo Sie die Antwort wissen

Anmerkung: Sie könnten Probleme haben, wenn sehr große Matrizen. Verwenden Sie eine While-Schleife mit geeigneten Schritten. Cuda by Example ist wie immer das Nachschlagewerk.

Ein Beispiel für die Implementierung einer verschachtelten Schleife finden Sie hier: For nested loops with CUDA. Dort ist eine while-Schleife implementiert.

marina.k hat Recht bezüglich der Race Condition. Das würde Ansatz 1 bevorzugen, da atomare Operationen den Code tendenziell verlangsamen.

Answer 2

Der erste Kernel ist in Ordnung. Aber wir haben keinen Zugriff auf die Matrix B und C.

Wie für die zweite Kernel-Funktion. Sie haben Datenrennen, weil nicht nur ein Thread eine Fähigkeit hat, in A[idx*width+idy] Adressen zu schreiben. Sie benötigen zusätzliche Synchronisation wie AttomicAdd

Wie für allgemeine Frage: Ich denke, dass Experimente zeigen, dass es besser ist. Es hängt von typischen Matrixgrößen ab, die Sie haben. Denken Sie daran, dass die maximale Thread-Blockgröße auf Fermi < 1024 und wenn Matrizen groß sind, Sie viele Thread-Blöcke gem. Normalerweise ist es langsamer (um viele Thread-Blöcke zu haben).

Answer 3

wirklich einfach in ArrayFire:

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