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