Gaußsche Eliminierung Parallelität

Question

Ich habe erfolgreich ein einzelnes Thread-Programm in CUDA für die Gaußsche Eliminierung implementiert und möchte Parallelität erreichen. Bis zu diesem Punkt ist die parallele Code wie folgt aussieht:

__global__ void ParallelGaussian(double* A) 
{ 
    int index = threadIdx.x; 
    int stride = blockDim.x; 

    if (index < ROWS) //Skip additional threads 
    { 
     for (unsigned int r = index; r < ROWS; r += stride) 
     { 
      //Forward elimination to reduce to row echelon form 
      for (unsigned int k = r + 1; k < ROWS; ++k) 
      { 
       double c = -A[(ROWS + 1) * k + r]/A[(ROWS + 1) * r + r]; 
       for (unsigned int j = r; j < ROWS + 1; ++j) 
       { 
        if (r == j) 
         A[(ROWS + 1) * k + j] = 0.0; 
        else 
         A[(ROWS + 1) * k + j] += c * A[(ROWS + 1) * r + j]; 
       } 
      } 
     } 
    } 
} 

Wie wir den Code auf der GPU sehen kann, wird die 1D-Array (Matrix) in eine untere Dreiecksmatrix zu transformieren und dann auf der CPU werde ich mit zurück weiter Substitution, um das Endergebnis zu erhalten. Bei diesem Ansatz wird keine Schwenkung durchgeführt, da diese nicht vollständig benötigt wird, sondern tatsächlich die numerische Stabilität des Algorithmus verbessert.

Starten des Kernels mit einem einzelnen Faden und einem Block arbeitet und transformiert die Matrix in Stufenform:

ParallelGaussian << < 1, 1 >> >(dev_a); 

Wenn ich jedoch die Anzahl der Threads erhöhen möchten, wie

ParallelGaussian << < 1, 32 >> >(dev_a); 

wird die untere Dreiecksmatrix nicht erzeugt. Das Hinzufügen von __synctreads() - Aufrufen in den Code, um die Threads in einem Block zu synchronisieren, verbessert die Situation nicht, was auch immer, und ich kann nicht herausfinden, warum.

Answer 1

Betrachten Sie Ihre innere Schleife. Jeder Thread greift auf A zu, und da k und j von r bis zum Ende der Matrix laufen, besteht die Möglichkeit, dass mehrere Threads den gleichen A[(ROWS + 1) * k + j] Wert ändern.

Sie haben möglicherweise auch einige Threads, die auf A[(ROWS + 1) * r + j] zugreifen, während andere Threads diesen Wert aktualisieren.

Eine mögliche Lösung besteht darin, jeden Thread in einzelnen Ergebnisfeldern zu speichern und diese am Ende zu kombinieren. Dies ist speicherintensiv.

Eine andere wäre, dies so zu restrukturieren, dass nur ein Thread auf einen bestimmten Wert schreibt und diese Werte in einer neuen Matrix speichert (so dass Sie keinen Wert ändern, der von einem anderen Thread benötigt wird).