Comment choisir les dimensions des grids et des blocs pour les kernelx CUDA?

Ceci est une question sur la façon de déterminer la taille de la grid, du bloc et du filetage CUDA. Ceci est une question supplémentaire à celle affichée ici:

https://stackoverflow.com/a/5643838/1292251

À la suite de 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 réglage et contraintes matérielles”.

Je n’ai pas trouvé de bonne explication ou de clarification qui explique cela dans la documentation CUDA. En résumé, ma question est de savoir comment déterminer la taille de bloc optimale (= nombre de threads) en fonction 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<<>>mAdd(A,B,C,n); 

BTW, j’ai commencé ma question avec le lien ci-dessus parce que cela répond en partie à ma première question. Si ce n’est pas une manière appropriée de poser des questions sur le débordement de stack, excusez-moi ou conseillez-moi.

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 facile à quantifier. L’Annexe F du guide de programmation CUDA actuel énumère un certain nombre de limites matérielles qui limitent le nombre de threads par bloc qu’un lancement de kernel peut avoir. Si vous dépassez un de ces éléments, votre kernel ne fonctionnera jamais. Ils peuvent être sommairement résumés comme suit:

  1. Chaque bloc ne peut pas avoir plus de 512/1024 threads au total ( Compute Capability 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 ultérieur)
  3. Chaque bloc ne peut pas consumr plus de 8k / 16k / 32k / 64k / 32k / 64k / 32k / 64k / 32k / 64k registres totaux (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 pas consumr plus de 16 Ko / 48 Ko / 96 Ko de mémoire partagée (Compute 1.x / 2.x-6.2 / 7.0)

Si vous restz dans ces limites, tout kernel que vous pouvez comstackr 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 les performances du code exécuté sur le matériel. Le comportement de chaque code sera différent et le seul moyen réel de le quantifier est de procéder à une parsing et à un profilage rigoureux. Mais encore une fois, très sommairement résumé:

  1. Le nombre de threads par bloc doit être un multiple de la taille de la chaîne, soit 32 sur tout le matériel actuel.
  2. Chaque unité multiprocesseur en continu sur le processeur graphique doit avoir suffisamment de chaînes actives pour masquer de manière suffisante toute la latence de pipeline de mémoire et d’instructions d’instructions de l’architecture et obtenir un débit maximal. L’approche orthodoxe consiste à essayer d’obtenir une occupation optimale du matériel (à quoi fait référence la réponse de Roger Dahl ).

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

Au niveau de l’entrée, vous devez surtout savoir que la taille de bloc que vous choisissez (dans la plage des tailles de bloc 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 utilisez. En effectuant des parsings comparatives, vous constaterez probablement que la plupart des codes non sortingviaux ont un «point sensible» dans la plage de 128 à 512 threads par bloc, mais vous aurez besoin d’parsings de votre part pour trouver cet emplacement. La bonne nouvelle est que, étant donné que 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.

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 en fonction de 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) comprend plusieurs nouvelles fonctions d’exécution pour faciliter les calculs d’occupation et la configuration du lancement, voir

Astuce CUDA Pro: l’API d’occupation simplifie la configuration du lancement

Une des fonctions utiles est cudaOccupancyMaxPotentialBlockSize qui calcule de manière heuristique une taille de bloc cudaOccupancyMaxPotentialBlockSize l’occupation maximale. Les valeurs fournies par cette fonction pourraient alors être utilisées comme sharepoint départ d’une optimisation manuelle des parameters de lancement. Voici un petit exemple.

 #include  /************************/ /* 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>>(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 

MODIFIER

Le cudaOccupancyMaxPotentialBlockSize est défini dans le fichier cuda_runtime.h et est défini comme suit:

 template __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 parameters 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, il est nécessaire de calculer ses propres dimensions de blocs 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 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.

La taille de bloc est généralement sélectionnée pour maximiser “l’occupation”. Rechercher sur l’occupation CUDA pour plus d’informations. En particulier, consultez la feuille de calcul CUDA Occupancy Calculator.