Comment utiliser la mémoire locale dans OpenCL?

J’ai joué avec OpenCL récemment, et je suis capable d’écrire des kernelx simples qui n’utilisent que la mémoire globale. Maintenant, je voudrais commencer à utiliser la mémoire locale, mais je n’arrive pas à comprendre comment utiliser get_local_size() et get_local_id() pour calculer un “morceau” de sortie à la fois.

Par exemple, supposons que je veuille convertir le kernel OpenCL Hello World d’Apple en quelque chose que la mémoire locale utilise. Comment le feriez-vous? Voici la source du kernel d’origine:

 __kernel square( __global float *input, __global float *output, const unsigned int count) { int i = get_global_id(0); if (i < count) output[i] = input[i] * input[i]; } 

Si cet exemple ne peut pas être facilement converti en quelque chose qui montre comment utiliser la mémoire locale, tout autre exemple simple fera l’affaire.

Consultez les exemples dans les SDK NVIDIA ou AMD, ils devraient vous orienter dans la bonne direction. Masortingx Transpose utiliserait par exemple la mémoire locale.

En utilisant votre kernel carré, vous pouvez mettre en scène les données dans un tampon intermédiaire. N’oubliez pas de transmettre le paramètre supplémentaire.

 __kernel square( __global float *input, __global float *output, __local float *temp, const unsigned int count) { int gtid = get_global_id(0); int ltid = get_local_id(0); if (gtid < count) { temp[ltid] = input[gtid]; // if the threads were reading data from other threads, then we would // want a barrier here to ensure the write completes before the read output[gtid] = temp[ltid] * temp[ltid]; } } 

Il existe une autre possibilité pour cela, si la taille de la mémoire locale est constante. Sans utiliser un pointeur dans la liste des parameters des kernelx, le tampon local peut être déclaré dans le kernel simplement en le déclarant __local:

 __local float localBuffer[1024]; 

Cela supprime le code dû à moins d’appels clSetKernelArg.

Dans OpenCL, la mémoire locale est conçue pour partager des données sur tous les éléments de travail d’un groupe de travail. Et il faut généralement faire un appel de barrière avant que les données de la mémoire locale puissent être utilisées (par exemple, un élément de travail veut lire une donnée de la mémoire locale écrite par les autres éléments de travail). La barrière est coûteuse en matériel informatique. Gardez à l’esprit que la mémoire locale doit être utilisée pour lire / écrire des données répétées. Les conflits bancaires doivent être évités autant que possible.

Si vous ne faites pas attention à la mémoire locale, vous risquez d’obtenir des performances moins bonnes que d’utiliser la mémoire globale.