Je l'ai posté sur Github .
Il s'agit d'une simple table de hachage GPU capable de traiter des centaines de millions d'insertions par seconde. Sur mon ordinateur portable NVIDIA GTX 1060, le code insère 64 millions de paires clé-valeur générées aléatoirement en 210 ms environ et supprime 32 millions de paires en 64 ms environ.
Autrement dit, la vitesse sur un ordinateur portable est d'environ 300 millions d'insertions/s et 500 millions de suppressions/s.
Le tableau est écrit en CUDA, bien que la même technique puisse être appliquée à HLSL ou GLSL. L'implémentation présente plusieurs limitations pour garantir des performances élevées sur une carte vidéo :
- Seules les clés de 32 bits et les mêmes valeurs sont traitées.
- La table de hachage a une taille fixe.
- Et cette taille doit être égale à deux à la puissance.
Pour les clés et les valeurs, vous devez réserver un simple marqueur de délimiteur (dans le code ci-dessus, il s'agit de 0xffffffff).
Table de hachage sans verrous
La table de hachage utilise l'adressage ouvert avec , c'est-à-dire qu'il s'agit simplement d'un tableau de paires clé-valeur stocké en mémoire et offrant des performances de cache supérieures. On ne peut pas en dire autant du chaînage, qui consiste à rechercher un pointeur dans une liste chaînée. Une table de hachage est un simple tableau stockant des éléments KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
La taille de la table est une puissance de deux, pas un nombre premier, car une instruction rapide suffit pour appliquer le masque pow2/AND, mais l'opérateur de module est beaucoup plus lent. Ceci est important dans le cas d'un sondage linéaire, car dans une recherche de table linéaire, l'index d'emplacement doit être enveloppé dans chaque emplacement. Et du coup, le coût de l’opération s’ajoute modulo dans chaque slot.
La table stocke uniquement la clé et la valeur de chaque élément, pas un hachage de la clé. Puisque la table ne stocke que des clés de 32 bits, le hachage est calculé très rapidement. Le code ci-dessus utilise le hachage Murmur3, qui n'effectue que quelques décalages, XOR et multiplications.
La table de hachage utilise des techniques de protection par verrouillage indépendantes de l'ordre de la mémoire. Même si certaines opérations d'écriture perturbent l'ordre d'autres opérations similaires, la table de hachage conservera toujours l'état correct. Nous en parlerons ci-dessous. La technique fonctionne très bien avec les cartes vidéo qui exécutent des milliers de threads simultanément.
Les clés et valeurs de la table de hachage sont initialisées à vides.
Le code peut être modifié pour gérer également les clés et valeurs 64 bits. Les clés nécessitent des opérations atomiques de lecture, d’écriture et de comparaison et d’échange. Et les valeurs nécessitent des opérations de lecture et d'écriture atomiques. Heureusement, dans CUDA, les opérations de lecture-écriture pour les valeurs 32 et 64 bits sont atomiques tant qu'elles sont naturellement alignées (voir ci-dessous). ), et les cartes vidéo modernes prennent en charge les opérations de comparaison et d'échange atomiques 64 bits. Bien entendu, lors du passage au 64 bits, les performances diminueront légèrement.
État de la table de hachage
Chaque paire clé-valeur d'une table de hachage peut avoir l'un des quatre états suivants :
- La clé et la valeur sont vides. Dans cet état, la table de hachage est initialisée.
- La clé a été écrite, mais la valeur n'a pas encore été écrite. Si un autre thread est en train de lire des données, il renvoie alors vide. C'est normal, la même chose se serait produite si un autre thread d'exécution avait fonctionné un peu plus tôt, et nous parlons d'une structure de données concurrente.
- La clé et la valeur sont enregistrées.
- La valeur est disponible pour d’autres threads d’exécution, mais la clé ne l’est pas encore. Cela peut se produire parce que le modèle de programmation CUDA possède un modèle de mémoire peu ordonné. C'est normal, de toute façon, la clé est toujours vide, même si la valeur ne l'est plus.
Une nuance importante est qu'une fois la clé écrite dans l'emplacement, elle ne bouge plus - même si la clé est supprimée, nous en reparlerons ci-dessous.
Le code de la table de hachage fonctionne même avec des modèles de mémoire peu ordonnés dans lesquels l'ordre dans lequel la mémoire est lue et écrite est inconnu. Lorsque nous examinons l'insertion, la recherche et la suppression dans une table de hachage, n'oubliez pas que chaque paire clé-valeur se trouve dans l'un des quatre états décrits ci-dessus.
Insertion dans une table de hachage
La fonction CUDA qui insère des paires clé-valeur dans une table de hachage ressemble à ceci :
void gpu_hashtable_insert(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
uint32_t prev = atomicCAS(&hashtable[slot].key, kEmpty, key);
if (prev == kEmpty || prev == key)
{
hashtable[slot].value = value;
break;
}
slot = (slot + 1) & (kHashTableCapacity-1);
}
}
Pour insérer une clé, le code parcourt le tableau de table de hachage en commençant par le hachage de la clé insérée. Chaque emplacement du tableau effectue une opération de comparaison et d'échange atomique qui compare la clé de cet emplacement à la clé vide. Si une incompatibilité est détectée, la clé de l'emplacement est mise à jour avec la clé insérée, puis la clé de l'emplacement d'origine est renvoyée. Si cette clé d'origine était vide ou correspondait à la clé insérée, alors le code a trouvé un emplacement approprié pour l'insertion et a inséré la valeur insérée dans l'emplacement.
Si dans un appel au noyau gpu_hashtable_insert() il y a plusieurs éléments avec la même clé, alors n'importe laquelle de leurs valeurs peut être écrite dans l'emplacement de clé. Ceci est considéré comme normal : l'une des écritures clé-valeur lors de l'appel réussira, mais comme tout cela se produit en parallèle au sein de plusieurs threads d'exécution, nous ne pouvons pas prédire quelle écriture mémoire sera la dernière.
Recherche de table de hachage
Code de recherche de clés :
uint32_t gpu_hashtable_lookup(KeyValue* hashtable, uint32_t key)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
return hashtable[slot].value;
}
if (hashtable[slot].key == kEmpty)
{
return kEmpty;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
Pour trouver la valeur d'une clé stockée dans une table, nous parcourons le tableau en commençant par le hachage de la clé que nous recherchons. Dans chaque slot, on vérifie si la clé est celle que l'on recherche, et si oui, on renvoie sa valeur. Nous vérifions également si la clé est vide, et si c'est le cas, nous abandonnons la recherche.
Si nous ne trouvons pas la clé, alors le code renvoie une valeur vide.
Toutes ces opérations de recherche peuvent être effectuées simultanément via des insertions et des suppressions. Chaque paire du tableau aura l'un des quatre états décrits ci-dessus pour le flux.
Suppression dans une table de hachage
Code pour supprimer les clés :
void gpu_hashtable_delete(KeyValue* hashtable, uint32_t key, uint32_t value)
{
uint32_t slot = hash(key);
while (true)
{
if (hashtable[slot].key == key)
{
hashtable[slot].value = kEmpty;
return;
}
if (hashtable[slot].key == kEmpty)
{
return;
}
slot = (slot + 1) & (kHashTableCapacity - 1);
}
}
La suppression d'une clé se fait d'une manière inhabituelle : nous laissons la clé dans le tableau et marquons sa valeur (pas la clé elle-même) comme vide. Ce code est très similaire à lookup(), sauf que lorsqu'une correspondance est trouvée sur une clé, sa valeur est vide.
Comme mentionné ci-dessus, une fois qu'une clé est écrite dans un emplacement, elle n'est plus déplacée. Même lorsqu'un élément est supprimé de la table, la clé reste en place, sa valeur devient simplement vide. Cela signifie que nous n'avons pas besoin d'utiliser une opération d'écriture atomique pour la valeur de l'emplacement, car peu importe que la valeur actuelle soit vide ou non - elle deviendra toujours vide.
Redimensionner une table de hachage
Vous pouvez modifier la taille d'une table de hachage en créant une table plus grande et en y insérant des éléments non vides de l'ancienne table. Je n'ai pas implémenté cette fonctionnalité parce que je voulais garder l'exemple de code simple. De plus, dans les programmes CUDA, l'allocation de mémoire se fait souvent dans le code hôte plutôt que dans le noyau CUDA.
L'article décrit comment modifier une telle structure de données protégée par un verrou.
Compétitivité
Dans les extraits de code de fonction ci-dessus gpu_hashtable_insert(), _lookup() и _delete() traiter une paire clé-valeur à la fois. Et plus bas gpu_hashtable_insert(), _lookup() и _delete() traiter un tableau de paires en parallèle, chaque paire dans un thread d'exécution GPU distinct :
// CPU code to invoke the CUDA kernel on the GPU
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);
// GPU code to process numkvs key/values in parallel
void gpu_hashtable_insert_kernel(KeyValue* hashtable, const KeyValue* kvs, unsigned int numkvs)
{
unsigned int threadid = blockIdx.x*blockDim.x + threadIdx.x;
if (threadid < numkvs)
{
gpu_hashtable_insert(hashtable, kvs[threadid].key, kvs[threadid].value);
}
}
La table de hachage résistante au verrouillage prend en charge les insertions, les recherches et les suppressions simultanées. Étant donné que les paires clé-valeur sont toujours dans l'un des quatre états et que les clés ne bougent pas, le tableau garantit l'exactitude même lorsque différents types d'opérations sont utilisés simultanément.
Cependant, si nous traitons un lot d'insertions et de suppressions en parallèle, et si le tableau de paires d'entrée contient des clés en double, alors nous ne pourrons pas prédire quelles paires « gagneront » – elles seront écrites en dernier dans la table de hachage. Disons que nous avons appelé le code d'insertion avec un tableau d'entrée de paires A/0 B/1 A/2 C/3 A/4. Une fois le code terminé, les paires B/1 и C/3 sont garantis d'être présents dans le tableau, mais en même temps, n'importe laquelle des paires y apparaîtra A/0, A/2 ou A/4. Cela peut ou non être un problème – tout dépend de l'application. Vous savez peut-être à l'avance qu'il n'y a pas de clés en double dans le tableau d'entrée, ou vous ne vous souciez peut-être pas de la dernière valeur écrite.
Si cela vous pose un problème, vous devez séparer les paires en double en différents appels système CUDA. Dans CUDA, toute opération qui appelle le noyau se termine toujours avant le prochain appel du noyau (au moins dans un thread. Dans différents threads, les noyaux sont exécutés en parallèle). Dans l'exemple ci-dessus, si vous appelez un noyau avec A/0 B/1 A/2 C/3, et l'autre avec A/4, puis la clé A obtiendra la valeur 4.
Voyons maintenant si les fonctions doivent lookup() и delete() utilisez un pointeur simple ou volatile vers un tableau de paires dans la table de hachage. Stipule que:
Le compilateur peut choisir d'optimiser les lectures et écritures dans la mémoire globale ou partagée... Ces optimisations peuvent être désactivées à l'aide du mot-clé
volatile: ... toute référence à cette variable est compilée en une véritable instruction de lecture ou d'écriture en mémoire.
Les considérations d’exactitude ne nécessitent pas d’application volatile. Si le thread d'exécution utilise une valeur mise en cache provenant d'une opération de lecture antérieure, il utilisera alors des informations légèrement obsolètes. Mais il s’agit quand même d’informations provenant de l’état correct de la table de hachage à un certain moment de l’appel du noyau. Si vous avez besoin d'utiliser les dernières informations, vous pouvez utiliser l'index volatile, mais ensuite les performances diminueront légèrement : d'après mes tests, lors de la suppression de 32 millions d'éléments, la vitesse est passée de 500 millions de suppressions/s à 450 millions de suppressions/s.
Performance
Dans le test d'insertion de 64 millions d'éléments et d'en supprimer 32 millions, la concurrence entre std::unordered_map et il n'y a pratiquement pas de table de hachage pour le GPU :
std::unordered_map passé 70 691 ms à insérer et retirer des éléments puis à les libérer unordered_map (se débarrasser de millions d'éléments prend beaucoup de temps, car à l'intérieur unordered_map plusieurs allocations de mémoire sont effectuées). A dire vrai, std:unordered_map des restrictions complètement différentes. Il s'agit d'un thread d'exécution de processeur unique, prend en charge les valeurs-clés de toute taille, fonctionne bien à des taux d'utilisation élevés et affiche des performances stables après plusieurs suppressions.
La durée de la table de hachage pour le GPU et la communication inter-programmes était de 984 ms. Cela inclut le temps passé à placer la table en mémoire et à la supprimer (en allouant 1 Go de mémoire une fois, ce qui prend un certain temps dans CUDA), à insérer et à supprimer des éléments et à les parcourir. Toutes les copies vers et depuis la mémoire de la carte vidéo sont également prises en compte.
La table de hachage elle-même a pris 271 ms. Cela inclut le temps passé par la carte vidéo à insérer et à supprimer des éléments, et ne prend pas en compte le temps passé à copier en mémoire et à parcourir la table résultante. Si la table GPU dure longtemps ou si la table de hachage est entièrement contenue dans la mémoire de la carte vidéo (par exemple, pour créer une table de hachage qui sera utilisée par un autre code GPU et non par le processeur central), alors le résultat du test est pertinent.
La table de hachage d'une carte vidéo démontre des performances élevées grâce à un débit élevé et une parallélisation active.
Limites
L'architecture des tables de hachage présente quelques problèmes à prendre en compte :
- Le sondage linéaire est gêné par le regroupement, ce qui fait que les clés dans la table ne sont pas parfaitement placées.
- Les clés ne sont pas supprimées à l'aide de la fonction
deleteet avec le temps, ils encombrent la table.
En conséquence, les performances d'une table de hachage peuvent progressivement se dégrader, surtout si elle existe depuis longtemps et comporte de nombreuses insertions et suppressions. Une façon d'atténuer ces inconvénients consiste à effectuer un nouveau hachage dans une nouvelle table avec un taux d'utilisation assez faible et à filtrer les clés supprimées pendant le rehachage.
Pour illustrer les problèmes décrits, j'utiliserai le code ci-dessus pour créer une table avec 128 millions d'éléments et parcourir 4 millions d'éléments jusqu'à ce que j'aie rempli 124 millions d'emplacements (taux d'utilisation d'environ 0,96). Voici le tableau des résultats, chaque ligne est un appel du noyau CUDA pour insérer 4 millions de nouveaux éléments dans une table de hachage :
Taux d'utilisation
Durée d'insertion 4 194 304 éléments
0,00
11,608448 ms (361,314798 millions de clés/sec.)
0,03
11,751424 ms (356,918799 millions de clés/sec.)
0,06
11,942592 ms (351,205515 millions de clés/sec.)
0,09
12,081120 ms (347,178429 millions de clés/sec.)
0,12
12,242560 ms (342,600233 millions de clés/sec.)
0,16
12,396448 ms (338,347235 millions de clés/sec.)
0,19
12,533024 ms (334,660176 millions de clés/sec.)
0,22
12,703328 ms (330,173626 millions de clés/sec.)
0,25
12,884512 ms (325,530693 millions de clés/sec.)
0,28
13,033472 ms (321,810182 millions de clés/sec.)
0,31
13,239296 ms (316,807174 millions de clés/sec.)
0,34
13,392448 ms (313,184256 millions de clés/sec.)
0,37
13,624000 ms (307,861434 millions de clés/sec.)
0,41
13,875520 ms (302,280855 millions de clés/sec.)
0,44
14,126528 ms (296,909756 millions de clés/sec.)
0,47
14,399328 ms (291,284699 millions de clés/sec.)
0,50
14,690304 ms (285,515123 millions de clés/sec.)
0,53
15,039136 ms (278,892623 millions de clés/sec.)
0,56
15,478656 ms (270,973402 millions de clés/sec.)
0,59
15,985664 ms (262,379092 millions de clés/sec.)
0,62
16,668673 ms (251,627968 millions de clés/sec.)
0,66
17,587200 ms (238,486174 millions de clés/sec.)
0,69
18,690048 ms (224,413765 millions de clés/sec.)
0,72
20,278816 ms (206,831789 millions de clés/sec.)
0,75
22,545408 ms (186,038058 millions de clés/sec.)
0,78
26,053312 ms (160,989275 millions de clés/sec.)
0,81
31,895008 ms (131,503463 millions de clés/sec.)
0,84
42,103294 ms (99,619378 millions de clés/sec.)
0,87
61,849056 ms (67,815164 millions de clés/sec.)
0,90
105,695999 ms (39,682713 millions de clés/sec.)
0,94
240,204636 ms (17,461378 millions de clés/sec.)
À mesure que l’utilisation augmente, les performances diminuent. Ce n’est pas souhaitable dans la plupart des cas. Si une application insère des éléments dans un tableau puis les supprime (par exemple, lors du comptage des mots dans un livre), cela ne pose pas de problème. Mais si l'application utilise une table de hachage de longue durée (par exemple, dans un éditeur graphique pour stocker des parties non vides d'images où l'utilisateur insère et supprime fréquemment des informations), alors ce comportement peut être problématique.
Et mesuré la profondeur de sondage de la table de hachage après 64 millions d'insertions (facteur d'utilisation 0,5). La profondeur moyenne était de 0,4774, donc la plupart des clés se trouvaient soit dans le meilleur emplacement possible, soit à un emplacement de la meilleure position. La profondeur maximale de sondage était de 60.
J'ai ensuite mesuré la profondeur de sondage sur une table comportant 124 millions d'inserts (facteur d'utilisation 0,97). La profondeur moyenne était déjà de 10,1757, et le maximum - 6474 (!!). Les performances de détection linéaire diminuent considérablement à des taux d’utilisation élevés.
Il est préférable de maintenir le taux d'utilisation de cette table de hachage à un niveau bas. Mais nous augmentons ensuite les performances au détriment de la consommation de mémoire. Heureusement, dans le cas de clés et valeurs de 32 bits, cela peut être justifié. Si dans l'exemple ci-dessus, dans une table contenant 128 millions d'éléments, nous conservons le facteur d'utilisation de 0,25, alors nous ne pouvons pas y placer plus de 32 millions d'éléments et les 96 millions d'emplacements restants seront perdus - 8 octets pour chaque paire. , 768 Mo de mémoire perdue.
Veuillez noter que nous parlons de la perte de mémoire de la carte vidéo, qui est une ressource plus précieuse que la mémoire système. Bien que la plupart des cartes graphiques de bureau modernes prenant en charge CUDA disposent d'au moins 4 Go de mémoire (au moment de la rédaction, la NVIDIA 2080 Ti en dispose de 11 Go), ce ne serait toujours pas la décision la plus sage de perdre de telles quantités.
Plus tard, j'écrirai davantage sur la création de tables de hachage pour les cartes vidéo qui n'ont pas de problèmes de profondeur de sondage, ainsi que sur les moyens de réutiliser les emplacements supprimés.
Mesure de profondeur de sondage
Pour déterminer la profondeur de sondage d'une clé, nous pouvons extraire le hachage de la clé (son index de table idéal) de son index de table réel :
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
En raison de la magie des nombres binaires complémentaires à deux et du fait que la capacité de la table de hachage est de deux puissance deux, cette approche fonctionnera même lorsque l'index clé est déplacé au début de la table. Prenons une clé hachée à 1, mais insérée dans l'emplacement 3. Ensuite, pour une table de capacité 4, nous obtenons (3 — 1) & 3, ce qui équivaut à 2.
Conclusion
Si vous avez des questions ou des commentaires, envoyez-moi un courriel à ou ouvrez un nouveau sujet dans.
Ce code a été écrit en s'inspirant d'excellents articles :
À l'avenir, je continuerai à écrire sur les implémentations de tables de hachage pour les cartes vidéo et à analyser leurs performances. Mes plans incluent le chaînage, le hachage Robin des Bois et le hachage coucou utilisant des opérations atomiques dans des structures de données compatibles GPU.
Source: habr.com