Grundlegendes zu CUDA-Rasterdimensionen, Blockdimensionen und Thread-Organisation (einfache Erklärung) [geschlossen]

Wie sind Threads organisiert, um von einer GPU ausgeführt zu werden?

Hardware

Wenn ein GPU-Gerät beispielsweise über 4 Multiprozessor-Einheiten verfügt und diese jeweils 768 Threads ausführen können, werden zu einem bestimmten Zeitpunkt nicht mehr als 4 * 768 Threads tatsächlich parallel ausgeführt (wenn Sie mehr Threads geplant haben, warten sie Sie sind dran).

Software

Threads sind in Blöcken organisiert. Ein Block wird von einer Mehrfachverarbeitungseinheit ausgeführt. Die Threads eines Blocks können mithilfe von 1Dimension (x) -, 2Dimensions (x, y) - oder 3Dim-Indizes (x, y, z) identifiziert (indiziert) werden. In unserem Beispiel ist jedoch in jedem Fall x y z <= 768 (es gelten andere Einschränkungen) zu x, y, z siehe Anleitung und Gerätefähigkeit).

Wenn Sie mehr als diese 4 * 768-Threads benötigen, benötigen Sie natürlich mehr als 4 Blöcke. Blöcke können auch 1D, 2D oder 3D indiziert werden. Es gibt eine Warteschlange von Blöcken, die darauf warten, in die GPU eingegeben zu werden (in unserem Beispiel verfügt die GPU über 4 Multiprozessoren und nur 4 Blöcke werden gleichzeitig ausgeführt).

Nun ein einfacher Fall: Verarbeiten eines 512x512-Bildes

Angenommen, ein Thread soll ein Pixel (i, j) verarbeiten.

Wir können Blöcke mit jeweils 64 Threads verwenden. Dann brauchen wir 512 * 512/64 = 4096 Blöcke (also 512x512 Threads = 4096 * 64)

Es ist üblich, die Threads in 2D-Blöcken mit blockDim = 8 x 8 (die 64 Threads pro Block) zu organisieren (um die Indizierung des Bildes zu vereinfachen). Ich nenne es lieber threadsPerBlock.

dim3 threadsPerBlock(8, 8);  // 64 threads

und 2D gridDim = 64 x 64 Blöcke (die benötigten 4096 Blöcke). Ich nenne es lieber numBlocks.

dim3 numBlocks(imageWidth/threadsPerBlock.x,  /* for instance 512/8 = 64*/
              imageHeight/threadsPerBlock.y); 

Der Kernel wird folgendermaßen gestartet:

myKernel <<<numBlocks,threadsPerBlock>>>( /* params for the kernel function */ );       

Endlich: Es wird so etwas wie eine "Warteschlange mit 4096 Blöcken" geben, in der ein Block darauf wartet, einem der Multiprozessoren der GPU zugewiesen zu werden, damit seine 64 Threads ausgeführt werden.

Im Kernel wird das von einem Thread zu verarbeitende Pixel (i, j) folgendermaßen berechnet:

uint i = (blockIdx.x * blockDim.x) + threadIdx.x;
uint j = (blockIdx.y * blockDim.y) + threadIdx.y;

Angenommen, eine 9800GT-GPU:

• Es verfügt über 14 Multiprozessoren (SM)
• Jeder SM hat 8 Thread-Prozessoren (AKA-Stream-Prozessoren, SP oder Kerne).
• erlaubt bis zu 512 Threads pro Block
• Warpsize ist 32 (was bedeutet, dass jeder der 14x8 = 112 Thread-Prozessoren bis zu 32 Threads planen kann)

https://www.tutorialspoint.com/cuda/cuda_threads.htm

Ein Block kann nicht mehr aktive Threads als 512 haben, daher __syncthreadskann nur eine begrenzte Anzahl von Threads synchronisiert werden. dh wenn Sie Folgendes mit 600 Threads ausführen:

func1();
__syncthreads();
func2();
__syncthreads();

dann muss der Kernel zweimal ausgeführt werden und die Ausführungsreihenfolge lautet:

1. func1 wird für die ersten 512 Threads ausgeführt
2. func2 wird für die ersten 512 Threads ausgeführt
3. func1 wird für die verbleibenden Threads ausgeführt
4. func2 wird für die verbleibenden Threads ausgeführt

Hinweis:

Der Hauptpunkt ist __syncthreadseine blockweite Operation, bei der nicht alle Threads synchronisiert werden.

Ich bin mir nicht sicher, wie viele Threads __syncthreadsgenau synchronisiert werden können, da Sie einen Block mit mehr als 512 Threads erstellen und den Warp die Planung übernehmen lassen können. Nach meinem Verständnis ist es genauer zu sagen: func1 wird zumindest für die ersten 512 Threads ausgeführt.

Bevor ich diese Antwort bearbeitet habe (im Jahr 2010), habe ich gemessen, dass 14x8x32-Threads mit synchronisiert wurden __syncthreads.

Ich würde mich sehr freuen, wenn jemand dies erneut testen würde, um genauere Informationen zu erhalten.