Comment choisir les dimensions de grille et de bloc pour les noyaux CUDA?

Ceci est une question sur la façon de déterminer la grille CUDA, les tailles de bloc et de filetage. Ceci est une question supplémentaire à celle publiée ici .

Suite à ce lien, la réponse de talonmies contient un extrait de code (voir ci-dessous). Je ne comprends pas le commentaire "valeur généralement choisie par le réglage et les contraintes matérielles".

Je n'ai pas trouvé une bonne explication ou clarification qui explique cela dans la documentation CUDA. En résumé, ma question est de savoir comment déterminer le blocksize(nombre de threads) optimal compte tenu du code suivant:

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);

Il y a deux parties à cette réponse (je l'ai écrite). Une partie est facile à quantifier, l'autre est plus empirique.

Contraintes matérielles:

C'est la partie la plus facile à quantifier. L'annexe F du guide de programmation CUDA actuel répertorie un certain nombre de limites strictes qui limitent le nombre de threads par bloc qu'un lancement de noyau peut avoir. Si vous dépassez l'un d'entre eux, votre noyau ne fonctionnera jamais. Ils peuvent être résumés grossièrement comme suit:

1. Chaque bloc ne peut pas avoir plus de 512/1024 threads au total ( capacité de calcul 1.x ou 2.x et versions ultérieures respectivement)
2. Les dimensions maximales de chaque bloc sont limitées à [512,512,64] / [1024,1024,64] (Compute 1.x / 2.x ou version ultérieure)
3. Chaque bloc ne peut pas consommer plus de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k registres au total (Calculer 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. Chaque bloc ne peut pas consommer plus de 16 Ko / 48 Ko / 96 Ko de mémoire partagée (Compute 1.x / 2.x-6.2 / 7.0)

Si vous restez dans ces limites, tout noyau que vous pouvez compiler avec succès se lancera sans erreur.

L'optimisation des performances:

C'est la partie empirique. Le nombre de threads par bloc que vous choisissez dans les contraintes matérielles décrites ci-dessus peut affecter et affecte les performances du code exécuté sur le matériel. Le comportement de chaque code sera différent et le seul véritable moyen de le quantifier est de procéder à une analyse comparative et à un profilage minutieux. Mais encore une fois, très grossièrement résumé:

1. Le nombre de threads par bloc doit être un multiple arrondi de la taille de la chaîne, qui est de 32 sur tout le matériel actuel.
2. Chaque unité multiprocesseur de streaming sur le GPU doit avoir suffisamment de déformations actives pour masquer suffisamment toutes les différentes latences de mémoire et de pipeline d'instructions de l'architecture et atteindre un débit maximal. L'approche orthodoxe ici est d'essayer d'obtenir une occupation optimale du matériel (ce à quoi la réponse de Roger Dahl fait référence).

Le deuxième point est un sujet énorme dont je doute que quiconque essaiera de le couvrir dans une seule réponse StackOverflow. Il y a des gens qui rédigent des thèses de doctorat autour de l'analyse quantitative des aspects du problème (voir cette présentation de Vasily Volkov de l'UC Berkley et cet article de Henry Wong de l'Université de Toronto pour des exemples de la complexité réelle de la question).

Au niveau d'entrée, vous devez surtout être conscient que la taille de bloc que vous choisissez (dans la plage des tailles de bloc légales définies par les contraintes ci-dessus) peut avoir un impact sur la vitesse d'exécution de votre code, mais cela dépend du matériel. vous avez et le code que vous exécutez. En comparant, vous constaterez probablement que la plupart des codes non triviaux ont un "sweet spot" dans la plage de 128-512 threads par bloc, mais cela nécessitera une analyse de votre part pour trouver où cela se trouve. La bonne nouvelle est que, comme vous travaillez en multiples de la taille de chaîne, l'espace de recherche est très limité et la meilleure configuration pour un morceau de code donné est relativement facile à trouver.

Les réponses ci-dessus montrent comment la taille du bloc peut avoir un impact sur les performances et suggèrent une heuristique commune pour son choix basé sur la maximisation de l'occupation. Sans vouloir fournir le critère de choix de la taille du bloc, il convient de mentionner que CUDA 6.5 (maintenant en version Release Candidate) inclut plusieurs nouvelles fonctions d'exécution pour aider aux calculs d'occupation et lancer la configuration, voir

Astuce CUDA Pro: L'API d'occupation simplifie la configuration de lancement

L'une des fonctions utiles est de cudaOccupancyMaxPotentialBlockSizecalculer de manière heuristique une taille de bloc qui atteint l'occupation maximale. Les valeurs fournies par cette fonction pourraient alors être utilisées comme point de départ d'une optimisation manuelle des paramètres de lancement. Voici un petit exemple.

#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");

}

ÉDITER

Le cudaOccupancyMaxPotentialBlockSizeest défini dans le cuda_runtime.hfichier et est défini comme suit:

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

La signification des paramètres est la suivante

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.

Notez que, à partir de CUDA 6.5, il faut calculer ses propres dimensions de bloc 2D / 3D à partir de la taille de bloc 1D suggérée par l'API.

Notez également que l'API du pilote CUDA contient des API fonctionnellement équivalentes pour le calcul de l'occupation, il est donc possible de l'utiliser cuOccupancyMaxPotentialBlockSizedans le code API du pilote de la même manière que celle indiquée pour l'API d'exécution dans l'exemple ci-dessus.

La taille de bloc est généralement sélectionnée pour maximiser «l'occupation». Recherchez sur occupation CUDA pour plus d'informations. En particulier, consultez la feuille de calcul du calculateur d'occupation CUDA.