Comment et quand dois-je utiliser le pointeur avec l’API cuda?

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.

Mémoire linéaire vs mémoire rembourrée

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

Alignement de mots

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.

Conflits bancaires

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:

  1. Allouer la première rangée.
  2. Vérifiez si le nombre d'octets alloués le rend correctement aligné ( c'est- à- dire qu'il est un multiple de 128).
  3. Sinon, allouez d'autres octets pour atteindre le multiple suivant de 128. Le pas est alors le nombre d'octets alloués pour une seule ligne, y compris les octets supplémentaires (octets de remplissage).
  4. Répétez pour chaque ligne.

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

Copie des données vers / depuis la mémoire lancée

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

Traiter avec des tableaux 3D

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:

  • Lorsque vous utilisez cudaMalloc3D, vous recevez une valeur de hauteur que vous devez conserver avec soin pour un access ultérieur à la mémoire.
  • Lors de la copie d'un bloc de mémoire 3D, vous ne pouvez pas utiliser cudaMemcpy sauf si vous copiez une seule ligne. Vous devez utiliser tout autre type d'utilitaire de copie fourni par l'utilitaire cuda qui prend en compte le pitch.
  • Lorsque vous copiez vos données vers / depuis la mémoire linéaire, vous devez fournir une hauteur à votre pointeur même si cela n'a pas d'importance: cette hauteur est la taille d'une ligne, exprimée en octets.
  • Les parameters de taille sont exprimés en octets pour la taille de la ligne et en nombre d'éléments pour la dimension de colonne et de profondeur.

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

comme dans http://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html#group__CUDART__MEMORY_1g32bd7a39135594788a542ae72217775c

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.