Comment puis-je utiliser la mémoire locale à OpenCL?

Question

Je joue avec OpenCL récemment, et je suis capable d'écrire des noyaux simples qui utilisent uniquement la mémoire globale. Maintenant, je voudrais commencer à utiliser la mémoire locale, mais je ne peux pas sembler comprendre comment utiliser get_local_size() et get_local_id() pour calculer un « morceau » de la production à la fois.

Par exemple, disons que je voulais convertir OpenCL d'Apple Bonjour exemple du monde du noyau à quelque chose les utilise de la mémoire locale. Comment feriez-vous? Voici la source du noyau 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 facilement être transformé en quelque chose qui montre comment utiliser la mémoire locale, tout autre exemple simple fera.

Answer 1

Découvrez les échantillons dans les SDKs NVIDIA ou AMD, ils vous devez pointer dans la bonne direction. Matrix Transpose utiliserait la mémoire locale par exemple.

Utilisation de votre noyau d'élévation au carré, vous pouvez mettre en scène les données dans un tampon intermédiaire. N'oubliez pas de passer dans 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];
    }
}

Answer 2

Il y a une autre possibilité de le faire, si la taille de la mémoire locale est constante. Sans l'aide d'un pointeur dans la liste des paramètres des noyaux, la mémoire tampon locale peut être déclarée dans le noyau juste en déclarant __local:

__local float localBuffer[1024];

supprime le code en raison de moins d'appels clSetKernelArg.

Answer 3

Dans la mémoire locale OpenCL est destiné à partager des données sur tous les éléments de travail dans un groupe de travail. Et il faut généralement faire un appel barrière avant que les données de la mémoire locale peuvent être utilisés (par exemple, un élément de travail veut lire les données de la mémoire locale qui est écrit par les autres éléments de travail). Barrière est coûteuse en matériel. Gardez à l'esprit, la mémoire locale doit être utilisé pour les données répétées en lecture / écriture. conflit bancaire devrait être évité autant que possible.

Si vous ne faites pas attention à la mémoire locale, vous pouvez vous retrouver avec de moins bonnes performances temps que l'utilisation de la mémoire globale.