Comment utiliser la mémoire locale dans OpenCL?

Question

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

Par exemple, disons que j'ai voulu convertir Apple OpenCL exemple Hello World noyau à quelque chose de les utilisations de la mémoire locale. Comment le feriez-vous? Voici le noyau d'origine de la source:

__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 faire pour que l'utilisation de la mémoire locale, tout autre exemple simple permettra de le faire.

Merci!

Answer 1

Découvrez les exemples dans les kits de développement logiciel NVIDIA ou AMD, ils devraient vous indiquer la bonne direction. La transposition en matrice utiliserait la mémoire locale par exemple.

En utilisant votre noyau carré, vous pouvez organiser 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];
        output[gtid] =  temp[ltid] * temp[ltid];
    }
}
 

Answer 2

Il existe une autre possibilité de le faire si la taille de la mémoire locale est constante. Sans utiliser un pointeur dans la liste des paramètres du noyau, le tampon local peut être déclaré dans le noyau simplement en le déclarant __local:

 __local float localBuffer[1024];
 

Cela supprime le code car moins d'appels clSetKernelArg sont utilisés.

Answer 3

En OpenCL la mémoire locale est destinée à partager des données entre tous les éléments de travail dans un groupe de travail. Et il nécessite généralement de faire une barrière d'appel devant la mémoire locale de données peuvent être utilisées (par exemple, un élément de travail veut lire une mémoire locale de données qui sont écrites par les autres éléments de travail). La barrière est coûteux en matériel. Gardez à l'esprit, de la mémoire locale doit être utilisée en cas de récidive de données en lecture/écriture. Banque de conflit devrait être évitée autant que possible.

Si vous n'êtes pas prudent avec la mémoire locale, vous pouvez vous retrouver avec de mauvaises performances de certains de temps que l'utilisation de la mémoire globale.