Utiliser la structure comme support de tampon

Dans mon implémentation OpenCL actuelle, je voulais gagner du temps avec les arguments, éviter de les transmettre à chaque fois que je voulais utiliser un tampon dans un kernel et avoir une liste d’arguments plus courte pour mon kernel.

J’ai donc créé une structure (espace de travail) qui contient le pointeur sur la mémoire tampon de la mémoire du périphérique. La structure agit comme un object avec une variable membre à laquelle vous souhaitez accéder au fil du temps et que vous souhaitez conserver en vie pendant toute l’exécution. Je n’ai jamais eu de problème avec le processeur graphique AMD ni même avec le processeur. Mais Nvidia cause beaucoup de problèmes avec ça. Cela semble toujours être un problème d’alignement, ne jamais atteindre le tampon droit, etc.

Voici quelques codes pour vous aider, voir question ci-dessous:

La structure définit sur l’hôte:

#define SRC_IMG 0 // (float4 buffer) Source image #define LAB_IMG 1 // (float4 buffer) LAB image // NOTE: The size of this array should be as much as the last define + 1. #define __WRKSPC_SIZE__ 2 // Structure defined on host. struct Workspace { cl_ulong getPtr[__WRKSPC_SIZE__]; }; struct HostWorkspace { cl::Buffer srcImg; cl::Buffer labImg; }; 

La structure définie sur le périphérique:

 typedef struct __atsortingbute__(( packed )) gpuWorkspace { ulong getPtr[__WRKSPC_SIZE__]; } gpuWorkspace_t; 

Notez que sur device, j’utilise ulong et sur hôte, j’utilise cl_ulong comme indiqué ici. OpenCL: utilisation de struct comme argument du kernel .

Ainsi, une fois que cl :: Buffer pour l’image source ou l’image LAB est créé, je les enregistre dans un object HostWorkspace. Ainsi, tant que cet object n’est pas publié, la référence à cl :: Buffer est conservée. Un tampon existe donc pour tout le projet sur l’hôte. et de facto sur l’appareil.

Maintenant, je dois nourrir le périphérique, donc j’ai un kernel simple qui initie mon espace de travail de périphérique comme suit:

 __kernel void Workspace_Init(__global gpuWorkspace_t* wrkspc, __global float4* src, __global float4* LAB) { // Get the ulong pointer on the first element of each buffer. wrkspc->getPtr[SRC_IMG] = &src[0]; wrkspc->getPtr[LAB_IMG] = &LAB[0]; } 

où wrkspc est un tampon alloué avec struct Workspace , et src + LAB sont simplement des allocations de mémoire tampon en tant qu’images de tableau 1D.

Et ensuite, dans n’importe quel kernel, si je veux utiliser src ou LAB, je fais comme suit:

 __kernel void ComputeLABFromSrc(__global gpuWorkspace_t* wrkSpc) { // ============================================================= // Get pointer from work space. // ============================================================= // Cast back the pointer of first element as a normal buffer you // want to use along the execution of the kernel. __global float4* srcData = ( __global float4* )( wrkSpc->getPtr[SRC_IMG] ); __global float4* labData = ( __global float4* )( wrkSpc->getPtr[LAB_IMG] ); // Code kernel as usual. } 

Quand j’ai commencé à utiliser cela, j’avais 4-5 images qui fonctionnaient bien, avec une structure différente comme celle-ci:

 struct Workspace { cl_ulong imgPtr; cl_ulong labPtr; }; 

où chaque image avait son propre pointeur.

À un moment donné, j’atteins plus d’images et j’ai eu un problème. Donc, je cherche en ligne, et j’ai trouvé une recommandation que la structure sizeof () pourrait être différente entre périphérique / hôte, donc je le change en un tableau unique du même temps, et cela fonctionne bien jusqu’à 16 éléments.

Donc, je cherche plus, et j’ai trouvé une recommandation sur l’ atsortingbut ((emballé)), que je mets sur la structure de l’appareil (voir ci-dessus). Mais maintenant, j’atteins 26 éléments. Lorsque je vérifie la taille de la structure sur le périphérique ou sur l’hôte, la taille est de 208 (éléments * sizeof (cl_ulong) == 26 * 8). Mais j’ai toujours un problème similaire à mon modèle précédent, mon pointeur est lu ailleurs au milieu de l’image précédente, etc.

Je me suis donc demandé si quelqu’un essayait jamais un modèle similaire (peut-être avec une approche différente) ou avait des conseils pour avoir un modèle “solide” avec cela.

Notez que tous les kernelx sont bien codés, j’ai un bon résultat lors de l’exécution sur AMD ou sur la CPU avec le même code. Le seul problème concerne Nvidia.

N’essayez pas de stocker les valeurs de pointeur côté GPU au-delà des limites du kernel. Ils ne sont pas garantis pour restr les mêmes. Toujours utiliser des indices. Et si un kernel utilise un tampon spécifique, vous devez le passer comme argument à ce kernel.

Références:

  1. La spécification OpenCL 1.2 (pour autant que je sache, nvidia n’implémente pas de norme plus récente) ne définit pas le comportement des transtypages pointeur sur entier, ni inversement.
  2. La section 6.9p indique: “Les arguments relatifs aux fonctions du kernel déclarées être une structure ou une union ne permettent pas aux objects OpenCL d’être transmis en tant qu’éléments de la structure ou de l’union.” C’est exactement ce que vous essayez de faire: transmettre une structure de tampons à un kernel.
  3. La section 6.9a stipule: “Les arguments des fonctions du kernel dans un programme ne peuvent pas être déclarés en tant que pointeur sur un ou plusieurs pointeurs.” – C’est essentiellement ce que vous essayez de subvertir en transformant vos pointeurs en un entier et en un autre. (point 1) Vous ne pouvez pas “tromper” OpenCL pour que ceci soit bien défini en contournant le système de types.

Comme je le suggère dans le fil de commentaire ci-dessous, vous devrez utiliser des index pour enregistrer des positions dans un object tampon. Si vous souhaitez stocker des positions sur différentes régions de la mémoire, vous devez soit unifier plusieurs tampons en un et sauvegarder un index dans ce tampon géant, soit enregistrer une valeur numérique identifiant le tampon auquel vous faites référence.