Publiqueino en Github .
É unha simple táboa hash de GPU capaz de procesar centos de millóns de insercións por segundo. No meu portátil NVIDIA GTX 1060, o código insire 64 millóns de pares clave-valor xerados aleatoriamente nuns 210 ms e elimina 32 millóns de pares nuns 64 ms.
É dicir, a velocidade dun portátil é de aproximadamente 300 millóns de insercións/s e 500 millóns de eliminacións/s.
A táboa está escrita en CUDA, aínda que a mesma técnica pódese aplicar a HLSL ou GLSL. A implementación ten varias limitacións para garantir un alto rendemento nunha tarxeta de vídeo:
- Só se procesan claves de 32 bits e os mesmos valores.
- A táboa hash ten un tamaño fixo.
- E este tamaño debe ser igual a dous á potencia.
Para claves e valores, cómpre reservar un marcador delimitador sinxelo (no código anterior é 0xffffffff).
Mesa de hash sen pechaduras
A táboa hash usa o enderezo aberto con , é dicir, é simplemente unha matriz de pares clave-valor que se almacena na memoria e ten un rendemento de caché superior. Non se pode dicir o mesmo do encadeamento, que implica buscar un punteiro nunha lista enlazada. Unha táboa hash é unha matriz sinxela que almacena elementos KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
O tamaño da táboa é unha potencia de dous, non un número primo, porque unha instrución rápida é suficiente para aplicar a máscara pow2/AND, pero o operador do módulo é moito máis lento. Isto é importante no caso de sondaxe lineal, xa que nunha busca de táboa lineal o índice de slots debe estar envolto en cada slot. E como resultado, o custo da operación engádese módulo en cada slot.
A táboa só almacena a clave e o valor de cada elemento, non un hash da chave. Dado que a táboa só almacena claves de 32 bits, o hash calcúlase moi rapidamente. O código anterior usa o hash Murmur3, que só realiza algunhas quendas, XOR e multiplicacións.
A táboa hash usa técnicas de protección de bloqueo que son independentes da orde da memoria. Aínda que algunhas operacións de escritura perturben a orde doutras operacións deste tipo, a táboa hash aínda manterá o estado correcto. Disto falaremos a continuación. A técnica funciona moi ben con tarxetas de vídeo que executan miles de fíos simultáneamente.
As claves e os valores da táboa hash inícianse para baleirar.
O código tamén se pode modificar para manexar claves e valores de 64 bits. As claves requiren operacións atómicas de lectura, escritura e comparación e intercambio. E os valores requiren operacións de lectura e escritura atómicas. Afortunadamente, en CUDA, as operacións de lectura e escritura para valores de 32 e 64 bits son atómicas sempre que estean aliñadas de forma natural (ver a continuación). ), e as tarxetas de vídeo modernas admiten operacións de comparación e intercambio atómicos de 64 bits. Por suposto, ao pasar a 64 bits, o rendemento diminuirá lixeiramente.
Estado da táboa hash
Cada par clave-valor nunha táboa hash pode ter un dos catro estados:
- A clave e o valor están baleiros. Neste estado, a táboa hash iníciase.
- A clave foi anotada, pero o valor aínda non foi escrito. Se outro fío está lendo datos actualmente, volve baleiro. Isto é normal, ocorrería o mesmo se outro fío de execución funcionara un pouco antes, e estamos a falar dunha estrutura de datos concorrente.
- Rexístrase tanto a clave como o valor.
- O valor está dispoñible para outros fíos de execución, pero a clave aínda non está. Isto pode ocorrer porque o modelo de programación CUDA ten un modelo de memoria pouco ordenado. Isto é normal; en calquera caso, a chave aínda está baleira, aínda que o valor xa non o sexa.
Un matiz importante é que unha vez que se escribiu a chave no slot, xa non se move, aínda que se elimine a chave, falaremos sobre isto a continuación.
O código de táboa hash funciona incluso con modelos de memoria pouco ordenados nos que se descoñece a orde na que se le e escribe a memoria. Mentres miramos a inserción, a busca e a eliminación nunha táboa hash, lembre que cada par clave-valor está nun dos catro estados descritos anteriormente.
Inseríndose nunha táboa hash
A función CUDA que insire pares clave-valor nunha táboa hash ten o seguinte aspecto:
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);
}
}
Para inserir unha clave, o código itera a través da matriz da táboa hash comezando polo hash da clave inserida. Cada slot da matriz realiza unha operación de comparación e intercambio atómico que compara a chave dese slot coa baleira. Se se detecta unha falta de coincidencia, a chave do slot actualízase coa chave inserida e, a continuación, devólvese a clave do slot orixinal. Se esta chave orixinal estaba baleira ou coincidía coa clave inserida, entón o código atopou un espazo axeitado para a inserción e inseriu o valor inserido no slot.
Se nunha chamada de núcleo gpu_hashtable_insert() hai varios elementos coa mesma chave, entón calquera dos seus valores pódese escribir na ranura da chave. Isto considérase normal: unha das escrituras de clave-valor durante a chamada terá éxito, pero dado que todo isto ocorre en paralelo dentro de varios fíos de execución, non podemos predecir que escritura de memoria será a última.
Busca de táboas hash
Código para buscar claves:
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);
}
}
Para atopar o valor dunha clave almacenada nunha táboa, iteramos a través da matriz comezando polo hash da clave que buscamos. En cada slot, comprobamos se a clave é a que buscamos e, de ser así, devolvemos o seu valor. Tamén comprobamos se a chave está baleira, e se é así, abortamos a busca.
Se non atopamos a chave, o código devolve un valor baleiro.
Todas estas operacións de busca pódense realizar ao mesmo tempo mediante insercións e eliminacións. Cada par da táboa terá un dos catro estados descritos anteriormente para o fluxo.
Eliminando nunha táboa hash
Código para borrar claves:
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);
}
}
A eliminación dunha chave realízase dun xeito inusual: deixamos a chave na táboa e marcamos o seu valor (non a chave en si) como baleiro. Este código é moi parecido a lookup(), agás que cando se atopa unha coincidencia nunha chave, o seu valor está baleiro.
Como se mencionou anteriormente, unha vez que se escribe unha clave nun slot, xa non se move. Mesmo cando se elimina un elemento da táboa, a clave permanece no seu lugar, o seu valor simplemente queda baleiro. Isto significa que non necesitamos usar unha operación de escritura atómica para o valor do slot, porque non importa se o valor actual está baleiro ou non; aínda quedará baleiro.
Cambiar o tamaño dunha táboa hash
Podes cambiar o tamaño dunha táboa hash creando unha táboa máis grande e inserindo nela elementos non baleiros da táboa antiga. Non implementei esta funcionalidade porque quería manter o código de mostra sinxelo. Ademais, nos programas CUDA, a asignación de memoria adoita facerse no código do host e non no núcleo CUDA.
O artigo describe como modificar unha estrutura de datos protexida por bloqueo.
Competitividade
Nos fragmentos de código de función anteriores gpu_hashtable_insert(), _lookup() и _delete() procesar un par clave-valor á vez. E máis baixo gpu_hashtable_insert(), _lookup() и _delete() procesar unha matriz de pares en paralelo, cada par nun fío de execución GPU separado:
// 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);
}
}
A táboa hash resistente ao bloqueo admite insercións, buscas e eliminacións simultáneas. Dado que os pares clave-valor están sempre nun dos catro estados e as chaves non se moven, a táboa garante a corrección mesmo cando se usan diferentes tipos de operacións simultaneamente.
Non obstante, se procesamos un lote de insercións e eliminacións en paralelo, e se a matriz de pares de entrada contén claves duplicadas, non poderemos predecir cales son os pares que "gañarán": escribiranse en último lugar na táboa hash. Digamos que chamamos ao código de inserción cunha matriz de pares de entrada A/0 B/1 A/2 C/3 A/4. Cando se completa o código, emparejase B/1 и C/3 están garantidos para estar presentes na táboa, pero ao mesmo tempo aparecerá nela calquera das parellas A/0, A/2 ou A/4. Isto pode ou non ser un problema - todo depende da aplicación. Quizais saibas de antemán que non hai chaves duplicadas na matriz de entrada ou pode que non che importe que valor se escribiu o último.
Se isto é un problema para ti, debes separar os pares duplicados en chamadas de sistema CUDA diferentes. En CUDA, calquera operación que chame ao núcleo sempre se completa antes da seguinte chamada ao núcleo (polo menos dentro dun fío. En distintos fíos, os núcleos execútanse en paralelo). No exemplo anterior, se chamas a un núcleo con A/0 B/1 A/2 C/3, e o outro con A/4, entón a chave A obterá o valor 4.
Agora imos falar sobre se as funcións deberían lookup() и delete() use un punteiro simple ou volátil a unha matriz de pares na táboa hash. Indica que:
O compilador pode optar por optimizar as lecturas e as escrituras na memoria global ou compartida... Estas optimizacións pódense desactivar usando a palabra clave
volatile: ... calquera referencia a esta variable compílase nunha instrución de lectura ou escritura de memoria real.
As consideracións de corrección non requiren aplicación volatile. Se o fío de execución usa un valor almacenado na caché dunha operación de lectura anterior, entón utilizará información lixeiramente desactualizada. Pero aínda así, esta é información do estado correcto da táboa hash nun momento determinado da chamada do núcleo. Se precisa utilizar a información máis recente, pode utilizar o índice volatile, pero entón o rendemento diminuirá lixeiramente: segundo as miñas probas, ao eliminar 32 millóns de elementos, a velocidade diminuíu de 500 millóns de eliminacións/seg a 450 millóns de eliminacións/seg.
Produtividade
Na proba para inserir 64 millóns de elementos e eliminar 32 millóns deles, a competencia entre std::unordered_map e non hai practicamente ningunha táboa hash para a GPU:
std::unordered_map pasou 70 ms inserindo e eliminando elementos e despois liberándoos unordered_map (Desfacerse de millóns de elementos leva moito tempo, porque dentro unordered_map realízanse múltiples asignacións de memoria). Sinceramente, std:unordered_map restricións completamente diferentes. É un único fío de execución da CPU, admite valores clave de calquera tamaño, funciona ben a altas taxas de utilización e mostra un rendemento estable despois de varias eliminacións.
A duración da táboa hash para a comunicación entre programas e GPU foi de 984 ms. Isto inclúe o tempo dedicado a colocar a táboa na memoria e borrala (asignar 1 GB de memoria unha vez, o que leva algún tempo en CUDA), inserir e eliminar elementos e iterar sobre eles. Tamén se teñen en conta todas as copias desde e cara á memoria da tarxeta de vídeo.
A propia táboa hash tardou 271 ms en completarse. Isto inclúe o tempo empregado pola tarxeta de vídeo insertando e eliminando elementos, e non ten en conta o tempo dedicado a copiar na memoria e iterar sobre a táboa resultante. Se a táboa GPU vive durante moito tempo, ou se a táboa hash está contida enteiramente na memoria da tarxeta de vídeo (por exemplo, para crear unha táboa hash que será usada por outro código GPU e non polo procesador central), entón o resultado da proba é relevante.
A táboa hash para unha tarxeta de vídeo demostra un alto rendemento debido ao alto rendemento e á paralelización activa.
Limitacións
A arquitectura da táboa hash ten algúns problemas que debes ter en conta:
- A sondaxe lineal vese dificultada pola agrupación, o que fai que as claves da táboa non se coloquen perfectamente.
- As teclas non se eliminan mediante a función
deletee co paso do tempo desordenan a mesa.
Como resultado, o rendemento dunha táboa hash pode degradarse gradualmente, especialmente se existe durante moito tempo e ten numerosas insercións e eliminacións. Unha forma de mitigar estas desvantaxes é rehacer nunha táboa nova cunha taxa de utilización bastante baixa e filtrar as claves eliminadas durante a repetición.
Para ilustrar os problemas descritos, usarei o código anterior para crear unha táboa con 128 millóns de elementos e recorrer 4 millóns de elementos ata cubrir 124 millóns de espazos (taxa de utilización duns 0,96). Aquí está a táboa de resultados, cada fila é unha chamada ao núcleo CUDA para inserir 4 millóns de elementos novos nunha táboa hash:
Taxa de uso
Duración da inserción 4 elementos
0,00
11,608448 ms (361,314798 millóns de teclas/s)
0,03
11,751424 ms (356,918799 millóns de teclas/s)
0,06
11,942592 ms (351,205515 millóns de teclas/s)
0,09
12,081120 ms (347,178429 millóns de teclas/s)
0,12
12,242560 ms (342,600233 millóns de teclas/s)
0,16
12,396448 ms (338,347235 millóns de teclas/s)
0,19
12,533024 ms (334,660176 millóns de teclas/s)
0,22
12,703328 ms (330,173626 millóns de teclas/s)
0,25
12,884512 ms (325,530693 millóns de teclas/s)
0,28
13,033472 ms (321,810182 millóns de teclas/s)
0,31
13,239296 ms (316,807174 millóns de teclas/s)
0,34
13,392448 ms (313,184256 millóns de teclas/s)
0,37
13,624000 ms (307,861434 millóns de teclas/s)
0,41
13,875520 ms (302,280855 millóns de teclas/s)
0,44
14,126528 ms (296,909756 millóns de teclas/s)
0,47
14,399328 ms (291,284699 millóns de teclas/s)
0,50
14,690304 ms (285,515123 millóns de teclas/s)
0,53
15,039136 ms (278,892623 millóns de teclas/s)
0,56
15,478656 ms (270,973402 millóns de teclas/s)
0,59
15,985664 ms (262,379092 millóns de teclas/s)
0,62
16,668673 ms (251,627968 millóns de teclas/s)
0,66
17,587200 ms (238,486174 millóns de teclas/s)
0,69
18,690048 ms (224,413765 millóns de teclas/s)
0,72
20,278816 ms (206,831789 millóns de teclas/s)
0,75
22,545408 ms (186,038058 millóns de teclas/s)
0,78
26,053312 ms (160,989275 millóns de teclas/s)
0,81
31,895008 ms (131,503463 millóns de teclas/s)
0,84
42,103294 ms (99,619378 millóns de teclas/s)
0,87
61,849056 ms (67,815164 millóns de teclas/s)
0,90
105,695999 ms (39,682713 millóns de teclas/s)
0,94
240,204636 ms (17,461378 millóns de teclas/s)
A medida que aumenta a utilización, o rendemento diminúe. Isto non é desexable na maioría dos casos. Se unha aplicación insire elementos nunha táboa e despois descartaos (por exemplo, ao contar palabras nun libro), non é un problema. Pero se a aplicación usa unha táboa hash de longa duración (por exemplo, nun editor de gráficos para almacenar partes non baleiras de imaxes onde o usuario insire e elimina con frecuencia información), entón este comportamento pode ser problemático.
E mediu a profundidade de sondaxe da táboa hash despois de 64 millóns de insercións (factor de utilización 0,5). A profundidade media era de 0,4774, polo que a maioría das chaves estaban na mellor ranura posible ou a unha distancia da mellor posición. A profundidade máxima de sondaxe foi de 60.
A continuación, medii a profundidade de sondaxe nunha mesa con 124 millóns de insercións (factor de utilización 0,97). A profundidade media xa era de 10,1757, e a máxima - 6474 (!!). O rendemento da detección lineal cae significativamente a altas taxas de utilización.
É mellor manter baixa a taxa de utilización desta táboa hash. Pero entón aumentamos o rendemento a costa do consumo de memoria. Afortunadamente, no caso de claves e valores de 32 bits, isto pódese xustificar. Se no exemplo anterior, nunha táboa con 128 millóns de elementos, mantemos o factor de utilización de 0,25, entón non podemos colocar máis de 32 millóns de elementos nela e perderanse os 96 millóns de slots restantes: 8 bytes por cada par. , 768 MB de memoria perdida.
Teña en conta que estamos a falar da perda de memoria da tarxeta de vídeo, que é un recurso máis valioso que a memoria do sistema. Aínda que a maioría das tarxetas gráficas de escritorio modernas que admiten CUDA teñen polo menos 4 GB de memoria (no momento da escritura, a NVIDIA 2080 Ti ten 11 GB), aínda non sería a decisión máis acertada perder tales cantidades.
Máis adiante escribirei máis sobre a creación de táboas hash para tarxetas de vídeo que non teñan problemas coa profundidade de sondaxe, así como sobre formas de reutilizar os slots eliminados.
Medición da profundidade de sondaxe
Para determinar a profundidade de sondaxe dunha chave, podemos extraer o hash da chave (o seu índice de táboa ideal) do seu índice de táboa real:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Debido á maxia dos números binarios do complemento a dous e ao feito de que a capacidade da táboa hash é de dous a potencia de dous, este enfoque funcionará mesmo cando o índice clave se move ao principio da táboa. Tomemos unha chave que haxa a 1, pero que está inserida no slot 3. Despois, para unha táboa con capacidade 4, obtemos (3 — 1) & 3, que equivale a 2.
Conclusión
Se tes preguntas ou comentarios, envíame un correo electrónico a ou abrir un novo tema en.
Este código foi escrito inspirado en excelentes artigos:
No futuro, seguirei escribindo sobre implementacións de táboas hash para tarxetas de vídeo e analizarei o seu rendemento. Os meus plans inclúen o encadeamento, o hash Robin Hood e o hash de cuco mediante operacións atómicas en estruturas de datos compatibles coa GPU.
Fonte: www.habr.com