Ce qui se passe la rationalisation dans mon noyau simple, OpenCL en ce qui concerne la mémoire globale

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

Le noyau ci-dessus est une addition vectorielle fait dix fois par boucle. Je l'ai utilisé le guide de programmation et de débordement pile pour comprendre comment fonctionne la mémoire globale, mais je ne peux toujours pas comprendre en regardant mon code si je suis accès à la mémoire globale d'une bonne façon. J'accède dans une manière contiguë et je devine que d'une manière alignée. Est-ce que la charge carte des morceaux de 128Ko de mémoire globale pour les tableaux a, b, et c? Est-il alors charger les morceaux de 128Ko pour chaque tableau une fois pour tous les 32 indices gid traités? (4 * 32 = 128) Il semble que je ne suis pas en train de perdre tout droit de la bande passante de la mémoire globale?

BTW, le profileur montre un calcul efficacité GLD et gst de 1,00003, ce qui semble bizarre, je pensais que ce serait juste 1.0 si tous mes magasins et les charges ont été coalescées. Comment est-il supérieur à 1,0?

Answer 1

Oui, votre modèle d'accès à la mémoire est à peu près optimale. Chaque halfwarp accède à 16 mots de 32 bits consécutifs. De plus, l'accès est aligné 64 octets, étant donné que les tampons eux-mêmes sont alignées et startindex pour chaque halfwarp est un multiple de 16. Donc, chaque halfwarp générera une transaction 64 octets. Donc, vous ne devriez pas perdre bande passante mémoire par accès uncoalesced.

Depuis que vous avez demandé des exemples dans votre dernière question permet de modifier ce code pour d'autres (modèle d'accès moins optimal (puisque la boucle ne fait pas vraiment tout ce que je vais ignorer):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

permet d'abord se comment fonctionne sur Compute 1.3 (GT200) matériel

Pour les écritures à un cela va générer un modèle légèrement unoptimal (suivant les halfwarps identifiés par leur gamme d'identité et le modèle d'accès correspondant):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

Donc, fondamentalement, nous perdons environ la moitié de notre bande passante (moins de doubler la largeur d'accès pour les halfwarps bizarres ne contribue pas beaucoup, car il génère plus accès, ce qui est plus rapide alors perdre plus d'octets pour ainsi dire).

Pour le lit à partir de b fils accéder uniquement même des éléments de la matrice, de sorte que pour chaque halfwarp tous les accès sont situés dans un bloc aligné 128byte (le premier élément est à la frontière 128B, étant donné que pour cet élément du GID est un multiple de 16 => l'indice est un multiple de 32, 4 éléments d'octet, cela signifie que l'adresse de décalage est un multiple de 128B). Les étirements accesspattern sur tout le bloc 128B, donc cela fera un transfert 128B pour chaque halfwarp, waisting encore la moitié de la bande passante.

Le lit de c génèrent un des pires scénarios de cas, où chaque indices de fil dans son propre bloc 128B, de sorte que chaque fil a besoin de son propre transfert, dont l'un est d'une part un peu d'un scénario de sérialisation (bien que pas tout à fait aussi mauvais comme normaly, étant donné que le matériel devrait pouvoir chevaucher les transferts). Ce qui est pire est le fait que cela va transférer un bloc 32B pour chaque fil, perdre 7/8 de la bande passante (nous avons accès 4B / fil, 32B / 4B = 8, de sorte que 1/8 de la bande passante est utilisée). Puisque c'est le accesspattern de matrixtransposes naïf, il est fortement conseillé de faire ceux qui utilisent la mémoire locale (parlant de l'expérience).

Compute 1,0 (G80)

Voici le seul modèle qui va créer un bon accès est l'original, tous les modèles dans l'exemple crée un accès complètement uncoalesced, perdre 7/8 de la bande passante (transfert 32B / fil, voir ci-dessus). Pour chaque accès G80 matériel où le fil nième dans un halfwarp n'a pas accès à l'élément n-ième crée ces accès uncoalesced

Compute 2.0 (Fermi)

Ici, chaque accès à la mémoire crée 128B transactions (autant que nécessaire pour recueillir toutes les données, de sorte 16x128B dans le pire des cas), mais ceux-ci sont mises en cache, ce qui rend moins évident lorsque les données seront transférées. Pour le moment, laisse supposer le cache est assez grand pour contenir toutes les données et il n'y a pas de conflit, de sorte que chaque cacheline 128B sera transféré au plus une fois. Permet de prendre furthermoe une exécution sérialisée des halfwarps, nous avons donc une occupation de cache déterministe.

b à Accède transférera encore toujours 128B blocs (pas d'autres indices de fil dans le memoryarea) lui correspondant. L'accès à c va générer des transferts 128B par fil (pire possible accesspattern).

Pour les accès à un, il est le suivant (les traiter comme lit pour le moment):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Donc, pour de grands tableaux les accès à une gaspillent théoriquement presque pas de bande passante. Pour cet exemple, la réalité est bien sûr pas tout à fait aussi bon, puisque les accès à c plantera le cache assez bien

Pour le profileur je suppose que les gains d'efficacité sur 1.0 sont des résultats simplement de flotter inaccurencies point.

L'espoir qui aide