OpenCL (AMD GCN) Motif global d'accès à la mémoire pour les données vectorisées: stride ou contigu

Question

Je vais améliorer les performances du noyau OCL et vouloir clarifier le fonctionnement des transactions de mémoire et le meilleur accès à la mémoire (et pourquoi). Le noyau est alimenté par des vecteurs de 8 entiers définis comme array: int v [8], ce qui signifie que le vecteur entier doit être chargé dans les GPR avant tout calcul. Donc, je crois que le goulot d'étranglement de ce code est la charge de données initiale.

D'abord, je considère quelques bases théoriques.

Target HW est un Radeon RX 480/580 doté d'un bus mémoire GDDR5 de 256 bits, sur lequel la transaction de lecture/écriture en rafale a une granularité de 8 mots, d'où une transaction de mémoire de 2048 bits ou 256 octets. , Je crois, ce que CL_DEVICE_MEM_BASE_ADDR_ALIGN fait référence à:

Alignment (bits) of base address: 2048. 

Ainsi, ma première question: quel est le sens physique du cacheline 128 octets? Conserve-t-il la partie de données récupérée par simple lecture en rafale mais pas vraiment demandée? Que se passe-t-il avec le reste si nous avons demandé, disons, 32 ou 64 octets - ainsi, les restes dépassent la taille de la ligne de cache? (Je suppose, il sera juste jeté - alors, quelle partie: tête, queue ...?)

Maintenant, revenons à mon noyau, je pense que le cache ne joue pas un rôle important dans mon cas, car une rafale lit 64 entiers -> une transaction de mémoire peut théoriquement alimenter 8 éléments de travail à la fois, il n'y a pas de données supplémentaires à lire et la mémoire est toujours fusionnée.

Mais encore, je peux placer mes données avec deux modèles d'accès différents:

1) contiguë

a[i] = v[get_global_id(0) * get_global_size(0) + i]; 

(Wich en fait perfomed comme)

*(int8*)a = *(int8*)v; 

2) intercalés

a[i] = v[get_global_id(0) + i * get_global_size(0)]; 

Je pense que dans mon cas contigu serait plus rapide car comme mentionné ci-dessus une transaction de mémoire peut complètement bourrer 8 éléments de travail avec des données. Cependant, je ne sais pas, comment le planificateur dans l'unité de calcul fonctionne physiquement: a-t-il besoin de toutes les données pour être prêt pour toutes les voies SIMD ou seulement la première partie pour 4 éléments SIMD parallèles serait suffisante? Néanmoins, je suppose qu'il est assez intelligent pour fournir entièrement des données au moins une CU en premier, dès que les UC peuvent exécuter des flux de commande de manière indépendante. Dans le second cas, nous devons effectuer 8 * global_size/64 transactions pour obtenir un vecteur complet. Donc, ma deuxième question: est-ce que ma supposition est exacte?

Maintenant, la pratique.

En fait, je diviser la tâche entière en deux noyaux parce qu'une partie a moins de pression de registre qu'une autre et peut donc employer plus d'éléments de travail. Donc d'abord je jouais avec pattern comment les données stockées en transition entre les noyaux (en utilisant vload8/vstore8 ou cast to int8 donnaient le même résultat) et le résultat était quelque peu étrange: le noyau qui lit les données de manière contiguë fonctionne environ 10% plus vite CodeXL et par la mesure du temps de l'OS), mais le noyau qui stocke les données de manière contiguë fonctionne étonnamment plus lentement. Le temps global pour deux noyaux est alors à peu près le même. Dans mes pensées, les deux doivent se comporter au moins de la même manière - soit être plus lent ou plus rapide, mais ces résultats inverses semblaient inexplicables.

Et ma troisième question est: quelqu'un peut-il expliquer un tel résultat? Ou peut-être que je fais quelque chose de mal? (Ou complètement faux?)

Answer 1

Jetez un oeil au chapitre 2.1 dans le AMD OpenCL Optimization Guide. Il se concentre principalement sur les anciennes cartes de génération, mais l'architecture GCN n'a pas complètement changé, donc devrait toujours s'appliquer à votre appareil (polaris).

En général, les cartes AMD comportent plusieurs contrôleurs de mémoire auxquels sont réparties les demandes de mémoire de chaque cycle d'horloge. Si vous accédez par exemple à vos valeurs dans une logique majorée en colonne au lieu d'une logique majeure en ligne, votre performance sera pire car les requêtes sont envoyées au même contrôleur de mémoire. (Par colonne majeure, je veux dire qu'une colonne de votre matrice est accédée ensemble par tous les work-items exécutés dans le cycle d'horloge actuel, c'est ce que vous appelez coalesced vs entrelacé). Si vous accédez à une ligne d'éléments (en coalescence) dans un cycle d'horloge unique (ce qui signifie que toutes les valeurs d'accès aux éléments de travail se trouvent dans la même ligne), ces demandes doivent être distribuées aux différents contrôleurs de mémoire. En ce qui concerne l'alignement et la taille des lignes de cache, je me demande si cela aide vraiment à améliorer les performances. Si j'étais dans votre situation, j'essaierais de voir si je peux optimiser l'algorithme lui-même ou si j'accède souvent aux valeurs et il serait logique de les copier dans la mémoire locale. Mais encore une fois, il est difficile de dire sans savoir ce que vos noyaux exécutent.

Cordialement,

Michael

Answer 2

Eh bien, pas vraiment répondu tout ma question, mais certaines informations trouvées dans l'immensité de l'Internet des choses mis chemin ensemble plus clair, du moins pour moi (contrairement mentionné ci-dessus optimisation AMD Guide, qui ne semble pas claire et parfois confus):

«le matériel effectue une coalescent, mais il est compliqué ...
accès mémoire dans une chaîne ne doit pas nécessairement être contigus, mais le nombre de segments de mémoire globale de 32 octets (et de segments de cache l1 de 128 octets) auxquels ils appartiennent est important. le contrôleur de mémoire peut charger 1, 2 ou 4 de ces 32 segments d'octets dans une seule transaction, mais cela est lu dans le cache dans des lignes de cache de 128 octets.
ainsi, si chaque ligne dans une chaîne charge un mot aléatoire dans une plage de 128 octets, alors il n'y a pas de pénalité; c'est 1 transaction et la lecture est à pleine efficacité. mais, si chaque ligne dans une chaîne charge 4 octets avec une foulée de 128 octets, c'est très mauvais: 4096 octets sont chargés mais seulement 128 sont utilisés, ce qui donne un rendement de ~ 3%. »

Donc, pour mon Dans le cas où les données sont lues/stockées alors qu'elles sont toujours contiguës, l'ordre dans lequel les parties des vecteurs sont chargées peut avoir une incidence sur le (re) ordonnancement du flux de commandes par le compilateur.
Je peux également imaginer qu'une architecture GCN plus récente peut effectuer des écritures en cache/coalescées, c'est pourquoi mes résultats sont différents de ceux demandés par ce Guide d'optimisation.