J’ai une bonne compréhension de la façon d’allouer et de copier de la mémoire linéaire avec cudaMalloc() et cudaMemcpy() . Cependant, lorsque je souhaite utiliser les fonctions CUDA pour allouer et copier des masortingces 2D ou 3D, je suis souvent troublé par les différents arguments, en particulier en ce qui concerne les pointeurs pitchés qui sont toujours présents dans les tableaux 2D / 3D. La documentation est utile pour donner quelques exemples sur la façon de les utiliser, mais cela suppose que je sois familier avec la notion de padding et de pitch, ce que je ne suis pas.
Je finis généralement par peaufiner les divers exemples que je trouve dans la documentation ou ailleurs sur le Web, mais le débogage aveugle qui suit est assez pénible, alors ma question est la suivante:
Qu’est-ce qu’un pitch? Comment l’utiliser? Comment allouer et copier des tableaux 2D et 3D dans CUDA?
Voici une explication sur le pointeur aigu et le remplissage dans cuda.
Tout d’abord, commençons par la raison de l’existence d’une mémoire non linéaire. Lors de l’allocation de mémoire avec cudaMalloc, le résultat est comme une allocation avec malloc, nous avons un bloc de mémoire contigu de la taille spécifiée et nous pouvons y mettre tout ce que nous voulons. Si nous voulons atsortingbuer un vecteur de 10000 float, nous faisons simplement:
float* myVector; cudaMalloc(&myVector,10000*sizeof(float));
puis accédez à cet élément de myVector par indexation classique:
float element = myVector[i];
et si nous voulons accéder à l’élément suivant, nous faisons simplement:
float next_element = myvector[i+1];
Cela fonctionne très bien car accéder à un élément juste à côté du premier est (pour des raisons que je ne connais pas et que je ne souhaite pas être pour le moment) bon marché.
Les choses deviennent un peu différentes lorsque nous utilisons notre mémoire en tant que tableau 2D. Disons que notre vecteur 10000 float est en fait un tableau 100×100. Nous pouvons l’allouer en utilisant la même fonction cudaMalloc, et si nous voulons lire la dernière ligne, nous faisons:
float* myArray; cudaMalloc(&myArray,10000*sizeof(float)); int row[100]; // number of columns for (int j=0; j<100; ++j) row[j] = myArray[i*100+j];
Nous devons donc lire la mémoire de myArray + 100 * i sur myArray + 101 * i-1. Le nombre d'opérations d'access à la mémoire nécessaires dépend du nombre de mots mémoire pris par cette ligne. Le nombre d'octets dans un mot mémoire dépend de l'implémentation. Pour minimiser le nombre d'access à la mémoire lors de la lecture d'une seule ligne, nous devons nous assurer que nous commençons la ligne au début d'un mot, nous devons donc remplir la mémoire pour chaque ligne jusqu'au début d'une nouvelle.
Une autre raison du remplissage des tableaux est le mécanisme bancaire de cuda, concernant l'access à la mémoire partagée. Lorsque la masortingce est dans la mémoire partagée, elle est divisée en plusieurs banques de mémoire. Deux threads cuda peuvent y accéder simultanément, à condition qu'ils n'accèdent pas à la mémoire appartenant à la même banque de mémoire. Comme nous voulons généralement traiter chaque ligne en parallèle, nous pouvons nous assurer que nous pouvons y accéder de manière simulée en remplissant chaque ligne au début d'une nouvelle banque.
Maintenant, au lieu d’allouer le tableau 2D avec cudaMalloc, nous utiliserons cudaMallocPitched:
size_t pitch; float* myArray; cudaMallocPitch(&myArray,&pitch,100*sizeof(float),100);//width in bytes by height
Notez que le pitch est la valeur de retour de la fonction: cudaMallocPitch vérifie ce qu'il doit être sur votre système et renvoie la valeur appropriée. Ce que cudaMallocPitch fait est le suivant:
À la fin, nous avons généralement alloué plus de mémoire que nécessaire car chaque ligne a maintenant la taille de la hauteur, et non la taille de w * sizeof (float).
Mais maintenant, lorsque nous voulons accéder à l'élément suivant dans une colonne, nous devons faire:
float next_column_element = myArray[(j+1)*pitch+i];
Le décalage en octets entre deux colonnes successives ne peut plus être déduit de la taille de notre tableau, c'est pourquoi nous voulons garder la hauteur renvoyée par cudaMallocPitch. Et comme le pitch est un multiple de la taille du padding (généralement le plus grand de la taille des mots et de la taille de la banque), cela fonctionne très bien. Yay.
Maintenant que nous soaps comment créer et accéder à un élément unique dans un tableau créé par cudaMallocPitch, nous pourrions vouloir en copier une partie entière depuis et vers une autre mémoire, linéaire ou non.
Disons que nous voulons copier notre tableau dans un tableau 100x100 alloué sur notre hôte avec malloc:
float* host_memory = (float*)malloc(100*100*sizeof(float));
Si nous utilisons cudaMemcpy, nous copierons toute la mémoire allouée avec cudaMallocPitch, y compris les octets remplis entre chaque ligne. Ce que nous devons faire pour éviter de bourrer la mémoire est de copier chaque ligne une par une. Nous pouvons le faire manuellement:
for (size_t i=0;i<100;++i) { cudaMemcpy(host_memory[i*100],myArray[pitch*i], 100*sizeof(float),cudaMemcpyDeviceToHost); }
Ou nous pouvons dire à l'API cuda que nous ne voulons que la mémoire utile de la mémoire que nous avons allouée avec des octets de remplissage pour sa commodité, donc si elle pouvait gérer son propre problème automatiquement, ce serait très bien, merci. Et voici cudaMemcpy2D:
cudaMemcpy2D(host_memory,100*sizeof(float)/*destination pitch*/,myArray,pitch, 100*sizeof(float)/*width*/,100/*heigth*/,cudaMemcpyDeviceToHost);
Maintenant, la copie sera faite automatiquement. Il va copier le nombre d'octets spécifié en largeur (ici: 100xsizeof (float)), le premier temps (ici: 100), en sautant les octets de pitch à chaque fois qu'il passe à une ligne suivante. Notez que nous devons toujours fournir la hauteur de la mémoire de destination car elle pourrait aussi être remplie. Ici, ce n'est pas le cas, donc la hauteur est égale à la hauteur d'un tableau non-rembourré: c'est la taille d'une ligne. Notez également que le paramètre width de la fonction memcpy est exprimé en octets, mais que le paramètre de hauteur est exprimé en nombre d'éléments. C'est à cause de la façon dont la copie est faite, comme si j'avais écrit la copie manuelle ci-dessus: la largeur est la taille de chaque copie le long d'une ligne (elemnts contigus en mémoire) et la hauteur est le nombre de fois que cette opération doit être accompli. (Ces incohérences dans les unités, en tant que physicien, me gênent beaucoup.)
Les tableaux 3D ne sont pas différents des tableaux 2D, aucun remplissage supplémentaire n'est inclus. Un tableau 3D est simplement un tableau 2D classique de lignes remplies. C'est pourquoi, lors de l'allocation d'un tableau 3D, vous n'obtenez qu'un seul pitch représentant la différence en nombre d'octets entre les points successifs d'une ligne. Si vous souhaitez accéder à des points successifs le long de la dimension de profondeur, vous pouvez multiplier en toute sécurité le pas par le nombre de colonnes, ce qui vous donne le slicePitch.
Le cuda api pour accéder à la mémoire 3D est légèrement différent de celui de la mémoire 2D, mais l'idée est la même:
Dans la réponse de Hephaestos
Si nous utilisons cudaMemcpy, nous copierons toute la mémoire allouée avec cudaMallocPitch, y compris les octets remplis entre chaque ligne. Ce que nous devons faire pour éviter de bourrer la mémoire est de copier chaque ligne une par une. Nous pouvons le faire manuellement:
for (size_t i=0;i<100;++i) { cudaMemcpy(host_memory[i*100],myArray[pitch*100], 100*sizeof(float),cudaMemcpyDeviceToHost); }
Ici, l’adresse "Source mémoire" doit être myArray[i*pitch] plutôt que myArray[pitch*100] .
Dans la réponse d’Ernest_Galbrun
float next_column_element = myArray[(j+1)*pitch+i];
besoin d’être
float next_column_element = *((float*)((char*)myArray + (j+1) * pitch) + i);
Et comme l’a souligné @RobertCrovella,
float next_column_element = myArray[(j+1)*pitch/sizeof(float)+i];
n’est pas une bonne façon non plus.