Wie wähle ich Raster- und Blockdimensionen für CUDA-Kernel aus?

Dies ist eine Frage zum Bestimmen der CUDA-Raster-, Block- und Gewindegrößen. Dies ist eine zusätzliche Frage zu der hier geposteten .

Nach diesem Link enthält die Antwort von talonmies ein Code-Snippet (siehe unten). Ich verstehe den Kommentar "Wert, der normalerweise durch Tuning und Hardwareeinschränkungen gewählt wird" nicht.

Ich habe keine gute Erklärung oder Klarstellung gefunden, die dies in der CUDA-Dokumentation erklärt. Zusammenfassend ist meine Frage, wie man das Optimum blocksize(Anzahl der Threads) unter Verwendung des folgenden Codes bestimmt:

const int n = 128 * 1024;
int blocksize = 512; // value usually chosen by tuning and hardware constraints
int nblocks = n / nthreads; // value determine by block size and total work
madd<<<nblocks,blocksize>>>mAdd(A,B,C,n);

Diese Antwort besteht aus zwei Teilen (ich habe sie geschrieben). Ein Teil ist leicht zu quantifizieren, der andere ist empirischer.

Hardwareeinschränkungen:

Dies ist der leicht zu quantifizierende Teil. Anhang F des aktuellen CUDA-Programmierhandbuchs enthält eine Reihe von harten Grenzwerten, die die Anzahl der Threads pro Block begrenzen, die ein Kernelstart haben kann. Wenn Sie einen dieser Werte überschreiten, wird Ihr Kernel niemals ausgeführt. Sie können grob zusammengefasst werden als:

1. Jeder Block kann nicht mehr als 512/1024 Fäden insgesamt ( Compute Capability 1.x oder 2.x und später jeweils)
2. Die maximalen Abmessungen jedes Blocks sind auf [512,512,64] / [1024,1024,64] begrenzt (Berechnen Sie 1.x / 2.x oder höher).
3. Jeder Block kann insgesamt nicht mehr als 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k-Register verbrauchen (Compute 1.0,1.1 / 1.2,1.3 / 2.x- / 3.0 / 3.2 / 3.5-5.2 / 5,3 / 6-6,1 / 6,2 / 7,0)
4. Jeder Block kann nicht mehr als 16 KB / 48 KB / 96 KB gemeinsam genutzten Speicher belegen (Compute 1.x / 2.x-6.2 / 7.0).

Wenn Sie innerhalb dieser Grenzen bleiben, wird jeder Kernel, den Sie erfolgreich kompilieren können, ohne Fehler gestartet.

Leistungsoptimierung:

Dies ist der empirische Teil. Die Anzahl der Threads pro Block, die Sie innerhalb der oben beschriebenen Hardwareeinschränkungen auswählen, kann und wirkt sich auf die Leistung des auf der Hardware ausgeführten Codes aus. Das Verhalten jedes Codes ist unterschiedlich und die einzige Möglichkeit, ihn zu quantifizieren, besteht in einem sorgfältigen Benchmarking und Profiling. Aber noch einmal sehr grob zusammengefasst:

1. Die Anzahl der Threads pro Block sollte ein rundes Vielfaches der Warp-Größe sein, die auf der gesamten aktuellen Hardware 32 beträgt.
2. Jede Streaming-Multiprozessoreinheit auf der GPU muss über genügend aktive Warps verfügen, um die unterschiedliche Latenz der Speicher- und Befehlspipeline der Architektur ausreichend zu verbergen und einen maximalen Durchsatz zu erzielen. Der orthodoxe Ansatz besteht darin, eine optimale Hardware-Belegung zu erreichen (worauf sich Roger Dahls Antwort bezieht).

Der zweite Punkt ist ein großes Thema, von dem ich bezweifle, dass irgendjemand versuchen wird, es in einer einzigen StackOverflow-Antwort zu behandeln. Es gibt Leute, die Doktorarbeiten über die quantitative Analyse von Aspekten des Problems schreiben (siehe diese Präsentation von Vasily Volkov von der UC Berkley und diese Arbeit von Henry Wong von der University of Toronto für Beispiele, wie komplex die Frage wirklich ist).

Auf der Einstiegsebene sollten Sie sich vor allem darüber im Klaren sein, dass die von Ihnen gewählte Blockgröße (innerhalb des durch die obigen Einschränkungen definierten Bereichs zulässiger Blockgrößen) einen Einfluss darauf haben kann und hat, wie schnell Ihr Code ausgeführt wird, dies hängt jedoch von der Hardware ab Sie haben und den Code, den Sie ausführen. Beim Benchmarking werden Sie wahrscheinlich feststellen, dass der meiste nicht triviale Code einen "Sweet Spot" in den 128-512 Threads pro Blockbereich aufweist, aber es erfordert eine Analyse von Ihrer Seite, um herauszufinden, wo sich dieser befindet. Die gute Nachricht ist, dass der Suchraum sehr begrenzt ist und die beste Konfiguration für einen bestimmten Code relativ leicht zu finden ist, da Sie in Vielfachen der Warp-Größe arbeiten.

Die obigen Antworten zeigen auf, wie sich die Blockgröße auf die Leistung auswirken kann, und schlagen eine gemeinsame Heuristik für die Auswahl auf der Grundlage der Belegungsmaximierung vor. Ohne das Kriterium für die Auswahl der Blockgröße angeben zu wollen, ist zu erwähnen, dass CUDA 6.5 (jetzt in der Release Candidate-Version) mehrere neue Laufzeitfunktionen enthält, die bei der Belegungsberechnung und der Startkonfiguration helfen (siehe)

CUDA Pro-Tipp: Die Belegungs-API vereinfacht die Startkonfiguration

Eine der nützlichen Funktionen ist cudaOccupancyMaxPotentialBlockSizedie heuristische Berechnung einer Blockgröße, die die maximale Belegung erreicht. Die von dieser Funktion bereitgestellten Werte könnten dann als Ausgangspunkt für eine manuelle Optimierung der Startparameter verwendet werden. Unten ist ein kleines Beispiel.

#include <stdio.h>

/************************/
/* TEST KERNEL FUNCTION */
/************************/
__global__ void MyKernel(int *a, int *b, int *c, int N) 
{ 
    int idx = threadIdx.x + blockIdx.x * blockDim.x; 

    if (idx < N) { c[idx] = a[idx] + b[idx]; } 
} 

/********/
/* MAIN */
/********/
void main() 
{ 
    const int N = 1000000;

    int blockSize;      // The launch configurator returned block size 
    int minGridSize;    // The minimum grid size needed to achieve the maximum occupancy for a full device launch 
    int gridSize;       // The actual grid size needed, based on input size 

    int* h_vec1 = (int*) malloc(N*sizeof(int));
    int* h_vec2 = (int*) malloc(N*sizeof(int));
    int* h_vec3 = (int*) malloc(N*sizeof(int));
    int* h_vec4 = (int*) malloc(N*sizeof(int));

    int* d_vec1; cudaMalloc((void**)&d_vec1, N*sizeof(int));
    int* d_vec2; cudaMalloc((void**)&d_vec2, N*sizeof(int));
    int* d_vec3; cudaMalloc((void**)&d_vec3, N*sizeof(int));

    for (int i=0; i<N; i++) {
        h_vec1[i] = 10;
        h_vec2[i] = 20;
        h_vec4[i] = h_vec1[i] + h_vec2[i];
    }

    cudaMemcpy(d_vec1, h_vec1, N*sizeof(int), cudaMemcpyHostToDevice);
    cudaMemcpy(d_vec2, h_vec2, N*sizeof(int), cudaMemcpyHostToDevice);

    float time;
    cudaEvent_t start, stop;
    cudaEventCreate(&start);
    cudaEventCreate(&stop);
    cudaEventRecord(start, 0);

    cudaOccupancyMaxPotentialBlockSize(&minGridSize, &blockSize, MyKernel, 0, N); 

    // Round up according to array size 
    gridSize = (N + blockSize - 1) / blockSize; 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Occupancy calculator elapsed time:  %3.3f ms \n", time);

    cudaEventRecord(start, 0);

    MyKernel<<<gridSize, blockSize>>>(d_vec1, d_vec2, d_vec3, N); 

    cudaEventRecord(stop, 0);
    cudaEventSynchronize(stop);
    cudaEventElapsedTime(&time, start, stop);
    printf("Kernel elapsed time:  %3.3f ms \n", time);

    printf("Blocksize %i\n", blockSize);

    cudaMemcpy(h_vec3, d_vec3, N*sizeof(int), cudaMemcpyDeviceToHost);

    for (int i=0; i<N; i++) {
        if (h_vec3[i] != h_vec4[i]) { printf("Error at i = %i! Host = %i; Device = %i\n", i, h_vec4[i], h_vec3[i]); return; };
    }

    printf("Test passed\n");

}

BEARBEITEN

Das cudaOccupancyMaxPotentialBlockSizeist in der cuda_runtime.hDatei definiert und wie folgt definiert:

template<class T>
__inline__ __host__ CUDART_DEVICE cudaError_t cudaOccupancyMaxPotentialBlockSize(
    int    *minGridSize,
    int    *blockSize,
    T       func,
    size_t  dynamicSMemSize = 0,
    int     blockSizeLimit = 0)
{
    return cudaOccupancyMaxPotentialBlockSizeVariableSMem(minGridSize, blockSize, func, __cudaOccupancyB2DHelper(dynamicSMemSize), blockSizeLimit);
}

Die Bedeutung für die Parameter ist die folgende

minGridSize     = Suggested min grid size to achieve a full machine launch.
blockSize       = Suggested block size to achieve maximum occupancy.
func            = Kernel function.
dynamicSMemSize = Size of dynamically allocated shared memory. Of course, it is known at runtime before any kernel launch. The size of the statically allocated shared memory is not needed as it is inferred by the properties of func.
blockSizeLimit  = Maximum size for each block. In the case of 1D kernels, it can coincide with the number of input elements.

Beachten Sie, dass ab CUDA 6.5 die eigenen 2D / 3D-Blockabmessungen aus der von der API vorgeschlagenen 1D-Blockgröße berechnet werden müssen.

Beachten Sie auch, dass die CUDA-Treiber-API funktional äquivalente APIs für die Belegungsberechnung enthält, sodass die Verwendung cuOccupancyMaxPotentialBlockSizeim Treiber-API-Code auf die gleiche Weise möglich ist, wie sie im obigen Beispiel für die Laufzeit-API gezeigt wurde.

Die Blockgröße wird normalerweise ausgewählt, um die "Belegung" zu maximieren. Suchen Sie nach CUDA-Belegung, um weitere Informationen zu erhalten. Siehe insbesondere die Tabelle zum CUDA-Belegungsrechner.