Pourquoi la modification de la taille des blocs et de la grid a-t-elle un tel impact sur l’exécution?

Je travaille actuellement sur un didacticiel cuda permettant de convertir une image RGBA en niveaux de gris. Mais je ne pouvais pas comprendre pourquoi changer les blockSize et gridSize temps X33.

 __global__ void rgba_to_greyscale(const uchar4* const rgbaImage, unsigned char* const greyImage, int numRows, int numCols) { int i = blockIdx.x*numCols + threadIdx.x; float channelSum = .299f * rgbaImage[i].x + .587f * rgbaImage[i].y + .114f * rgbaImage[i].z; greyImage[i]= channelSum; } void your_rgba_to_greyscale(const uchar4 * const h_rgbaImage, uchar4 * const d_rgbaImage, unsigned char* const d_greyImage, size_t numRows, size_t numCols) { const dim3 blockSize(numCols, 1, 1); const dim3 gridSize(numRows, 1 , 1); rgba_to_greyscale<<>>(d_rgbaImage, d_greyImage, numRows, numCols); cudaDeviceSynchronize(); checkCudaErrors(cudaGetLastError()); } 

Quand je définis comme ci-dessus:

 const dim3 blockSize(numCols, 1, 1); const dim3 gridSize(numRows, 1 , 1); 

Je reçois Your code executed in 0.030304 ms

Quand je mets:

  const dim3 blockSize(1, 1, 1); const dim3 gridSize(numRows, numCols , 1); 

et mettre à jour la fonction thread pour fonctionner avec le nouvel index:

 int i = blockIdx.x*numCols + blockIdx.y; 

Je reçois Your code executed in 0.995456 ms .

  1. Je m’attendrais à ce que ce soit l’inverse puisque le gpu peut calculer tous les pixels séparément sur la deuxième division en grid Est-il lié aux problèmes de cohérence de la mémoire cache? Pourquoi ai-je ces résultats?
  2. Quelle est la meilleure théorie sur la grid et la taille de bloc pour ce problème? est-il possible de le calculer au moment de l’exécution?

FYI:

 numRows = 313 numCols =557 

Propriétés techniques:

 #uname -a && /usr/bin/nvidia-settings -v Linux ip-10-16-23-92 3.2.0-39-virtual #62-Ubuntu SMP Thu Feb 28 00:48:27 UTC 2013 x86_64 x86_64 x86_64 GNU/Linux nvidia-settings: version 304.54 (buildmeister@swio-display-x86-rhel47-11) 

Aucune des configurations grid / bloc n’est recommandée. Le premier n’est pas évolutif car le nombre de threads par bloc est limité pour le GPU, il échouera donc pour une taille d’image plus grande. Le second est un mauvais choix car il n’y a qu’un seul thread par bloc, ce qui n’est pas recommandé, car l’occupation du processeur graphique serait très faible. Vous pouvez le vérifier à l’aide du calculateur d’occupation du GPU fourni avec le CUDA Toolkit. La taille de bloc recommandée doit être un multiple de la taille de chaîne du GPU (16 ou 32) en fonction du GPU.

Une approche générale et évolutive de la grid 2D et de la taille de bloc dans votre cas ressemblerait à ceci:

 const dim3 blockSize(16, 16, 1); const dim3 gridSize((numCols + blockSize.x - 1)/blockSize.x, (numRows + blockSize.y - 1)/blockSize.y , 1); 

Vous pouvez modifier la taille du bloc de 16 x 16 à la taille de votre choix, à condition de respecter les limites de l’appareil. Un maximum de 512 threads par bloc est autorisé pour les périphériques de capacité de calcul de 1,0 à 1,3. Pour les périphériques de capacité de calcul 2.0 et suivants, cette limite est fixée à 1024 threads par bloc.

Comme maintenant, la grid et le bloc sont bidimensionnels, l’indexation à l’intérieur du kernel serait modifiée comme suit:

 int i = blockIdx.x * blockDim.x + threadIdx.x; //Column int j = blockIdx.y * blockDim.y + threadIdx.y; //Row int idx = j * numCols + i; //Don't forget to perform bound checks if(i>=numCols || j>=numRows) return; float channelSum = .299f * rgbaImage[idx].x + .587f * rgbaImage[idx].y + .114f * rgbaImage[idx].z; greyImage[idx]= channelSum;