139 votes

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

Il s'agit d'une question sur la manière de déterminer la taille des grilles, des blocs et des threads CUDA. Il s'agit d'une question complémentaire à celle posée aquí .

En suivant 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 en fonction des contraintes de réglage et de matériel".

Je n'ai pas trouvé de bonne explication ou de clarification à ce sujet dans la documentation CUDA. En résumé, ma question est de savoir comment déterminer la valeur optimale des blocksize (nombre de threads) pour le 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);

176voto

talonmies Points 41460

Cette réponse comporte deux parties (je l'ai écrite). L'une 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 énumère un certain nombre de limites strictes qui limitent le nombre de threads par bloc qu'un noyau peut lancer. Si vous dépassez l'une de ces limites, votre noyau ne fonctionnera jamais. Ces limites peuvent être grossièrement résumées comme suit :

  1. Chaque bloc ne peut avoir plus de 512/1024 threads au total ( Capacité de calcul 1.x ou 2.x et plus)
  2. Les dimensions maximales de chaque bloc sont limitées à [512,512,64]/[1024,1024,64] (Compute 1.x/2.x ou ultérieur)
  3. Chaque bloc ne peut consommer plus de 8k/16k/32k/64k/32k/64k/32k/64k/32k/64k registres au total. (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. Chaque bloc ne peut consommer plus de 16kb/48kb/96kb 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.

Optimisation des performances :

C'est la partie empirique. Le nombre de threads par bloc que vous choisissez dans le cadre des contraintes matérielles décrites ci-dessus peut avoir et a un effet sur les performances du code qui s'exécute sur le matériel. Le comportement de chaque code sera différent et le seul moyen de le quantifier est de procéder à une analyse comparative et à un profilage minutieux. Mais encore une fois, il s'agit d'un résumé très approximatif :

  1. Le nombre de threads par bloc doit être un multiple rond 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 warps actifs pour masquer toutes les différentes latences de la mémoire et du pipeline d'instructions de l'architecture et atteindre un débit maximal. L'approche orthodoxe consiste à essayer d'obtenir une occupation optimale du matériel (ce que l'on appelle le "hardware occupancy"). Réponse de Roger Dahl se réfère).

Le deuxième point est un sujet très vaste et je doute que quelqu'un essaie de le couvrir dans une seule réponse StackOverflow. Il y a des gens qui écrivent des thèses de doctorat sur l'analyse quantitative de certains aspects du problème (voir cette présentation par Vasily Volkov de l'Université de Californie à Berkley et ce document par Henry Wong de l'Université de Toronto pour des exemples illustrant la complexité de la question).

Au niveau débutant, vous devez surtout savoir que la taille de bloc que vous choisissez (dans la fourchette des tailles de bloc légales définies par les contraintes ci-dessus) peut avoir et a un impact sur la vitesse d'exécution de votre code, mais cela dépend du matériel dont vous disposez et du code que vous exécutez. En effectuant des analyses comparatives, vous constaterez probablement que la plupart des codes non triviaux ont un "point idéal" dans la plage de 128 à 512 threads par bloc, mais cela nécessitera une certaine analyse de votre part pour trouver ce point. La bonne nouvelle est que, comme vous travaillez en multiples de la taille de la chaîne, l'espace de recherche est très limité et la meilleure configuration pour un morceau de code donné est relativement facile à trouver.

47voto

JackOLantern Points 6965

Les réponses ci-dessus soulignent l'impact de la taille du bloc sur les performances et suggèrent une heuristique commune pour son choix, basée sur la maximisation de l'occupation. Sans vouloir fournir des les pour choisir 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 faciliter les calculs d'occupation et la configuration du lancement, voir

CUDA Pro Tip : L'API d'occupation simplifie la configuration du lancement

L'une des fonctions utiles est cudaOccupancyMaxPotentialBlockSize qui calcule de manière heuristique une taille de bloc permettant d'atteindre l'occupation maximale. Les valeurs fournies par cette fonction pourraient ensuite ê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");

}

EDIT

En cudaOccupancyMaxPotentialBlockSize est défini dans le cuda_runtime.h 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 qu'à partir de CUDA 6.5, vous devez calculer vos 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, de sorte qu'il est possible d'utiliser cuOccupancyMaxPotentialBlockSize dans le code de l'API du pilote de la même manière que pour l'API d'exécution dans l'exemple ci-dessus.

9voto

Roger Dahl Points 8326

La taille des blocs est généralement sélectionnée pour maximiser l'"occupation". Faites une recherche sur CUDA Occupancy pour plus d'informations. En particulier, consultez la feuille de calcul CUDA Occupancy Calculator.

Prograide.com

Prograide est une communauté de développeurs qui cherche à élargir la connaissance de la programmation au-delà de l'anglais.
Pour cela nous avons les plus grands doutes résolus en français et vous pouvez aussi poser vos propres questions ou résoudre celles des autres.

Powered by:

X