Eu postei no Github .
É uma tabela hash de GPU simples capaz de processar centenas de milhões de inserções por segundo. No meu laptop NVIDIA GTX 1060, o código insere 64 milhões de pares de valores-chave gerados aleatoriamente em cerca de 210 ms e remove 32 milhões de pares em cerca de 64 ms.
Ou seja, a velocidade em um laptop é de aproximadamente 300 milhões de inserções/seg e 500 milhões de exclusões/seg.
A tabela está escrita em CUDA, embora a mesma técnica possa ser aplicada ao HLSL ou GLSL. A implementação possui diversas limitações para garantir alto desempenho em uma placa de vídeo:
- Apenas chaves de 32 bits e os mesmos valores são processados.
- A tabela hash tem um tamanho fixo.
- E esse tamanho deve ser igual a dois elevado à potência.
Para chaves e valores, você precisa reservar um marcador delimitador simples (no código acima é 0xffffffff).
Tabela hash sem bloqueios
A tabela hash usa endereçamento aberto com , ou seja, é simplesmente uma matriz de pares de valores-chave armazenados na memória e com desempenho de cache superior. O mesmo não pode ser dito do encadeamento, que envolve a busca por um ponteiro em uma lista encadeada. Uma tabela hash é um array simples que armazena elementos KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
O tamanho da tabela é uma potência de dois, não um número primo, porque uma instrução rápida é suficiente para aplicar a máscara pow2/AND, mas o operador de módulo é muito mais lento. Isto é importante no caso de sondagem linear, pois em uma consulta de tabela linear o índice de slot deve ser encapsulado em cada slot. E como resultado, o custo da operação é adicionado módulo em cada slot.
A tabela armazena apenas a chave e o valor de cada elemento, não um hash da chave. Como a tabela armazena apenas chaves de 32 bits, o hash é calculado muito rapidamente. O código acima usa o hash Murmur3, que executa apenas algumas mudanças, XORs e multiplicações.
A tabela hash usa técnicas de proteção de bloqueio que são independentes da ordem da memória. Mesmo que algumas operações de gravação perturbem a ordem de outras operações, a tabela hash ainda manterá o estado correto. Falaremos sobre isso abaixo. A técnica funciona muito bem com placas de vídeo que executam milhares de threads simultaneamente.
As chaves e valores na tabela hash são inicializados como vazios.
O código também pode ser modificado para lidar com chaves e valores de 64 bits. As chaves requerem operações atômicas de leitura, gravação e comparação e troca. E os valores requerem operações atômicas de leitura e gravação. Felizmente, no CUDA, as operações de leitura e gravação para valores de 32 e 64 bits são atômicas, desde que estejam naturalmente alinhadas (veja abaixo). ) e placas de vídeo modernas suportam operações atômicas de comparação e troca de 64 bits. Obviamente, ao passar para 64 bits, o desempenho diminuirá ligeiramente.
Estado da tabela hash
Cada par de valores-chave em uma tabela hash pode ter um dos quatro estados:
- Chave e valor estão vazios. Neste estado, a tabela hash é inicializada.
- A chave foi anotada, mas o valor ainda não foi escrito. Se outro thread estiver lendo dados no momento, ele retornará vazio. Isso é normal, a mesma coisa teria acontecido se outro thread de execução tivesse funcionado um pouco antes, e estivéssemos falando de uma estrutura de dados simultânea.
- Tanto a chave quanto o valor são registrados.
- O valor está disponível para outros threads de execução, mas a chave ainda não está. Isso pode acontecer porque o modelo de programação CUDA possui um modelo de memória pouco ordenado. Isto é normal; em qualquer caso, a chave ainda está vazia, mesmo que o valor não esteja mais assim.
Uma nuance importante é que uma vez que a chave foi gravada no slot, ela não se move mais - mesmo que a chave seja excluída, falaremos sobre isso a seguir.
O código da tabela hash funciona até mesmo com modelos de memória pouco ordenados, nos quais a ordem em que a memória é lida e gravada é desconhecida. Ao examinarmos a inserção, a pesquisa e a exclusão em uma tabela hash, lembre-se de que cada par de valores-chave está em um dos quatro estados descritos acima.
Inserindo em uma tabela hash
A função CUDA que insere pares de valores-chave em uma tabela hash é semelhante a esta:
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 uma chave, o código percorre a matriz da tabela hash começando com o hash da chave inserida. Cada slot na matriz executa uma operação atômica de comparação e troca que compara a chave desse slot com a vazia. Se for detectada uma incompatibilidade, a chave no slot será atualizada com a chave inserida e, em seguida, a chave do slot original será retornada. Se esta chave original estivesse vazia ou correspondesse à chave inserida, o código encontrou um slot adequado para inserção e inseriu o valor inserido no slot.
Se em uma chamada do kernel gpu_hashtable_insert() existem vários elementos com a mesma chave, então qualquer um de seus valores pode ser gravado no slot da chave. Isso é considerado normal: uma das gravações de valor-chave durante a chamada será bem-sucedida, mas como tudo isso acontece em paralelo em vários threads de execução, não podemos prever qual gravação na memória será a última.
Pesquisa de tabela hash
Código para pesquisar chaves:
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 encontrar o valor de uma chave armazenada em uma tabela, iteramos pelo array começando com o hash da chave que procuramos. Em cada slot verificamos se a chave é a que procuramos e, em caso afirmativo, devolvemos o seu valor. Também verificamos se a chave está vazia e, em caso afirmativo, abortamos a pesquisa.
Se não conseguirmos encontrar a chave, o código retornará um valor vazio.
Todas essas operações de pesquisa podem ser realizadas simultaneamente por meio de inserções e exclusões. Cada par na tabela terá um dos quatro estados descritos acima para o fluxo.
Excluindo em uma tabela hash
Código para exclusão de chaves:
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 exclusão de uma chave é feita de uma forma incomum: deixamos a chave na tabela e marcamos seu valor (não a chave em si) como vazio. Este código é muito semelhante ao lookup(), exceto que quando uma correspondência é encontrada em uma chave, seu valor fica vazio.
Conforme mencionado acima, uma vez que uma chave é gravada em um slot, ela não é mais movida. Mesmo quando um elemento é excluído da tabela, a chave permanece no lugar, seu valor simplesmente fica vazio. Isso significa que não precisamos usar uma operação de gravação atômica para o valor do slot, porque não importa se o valor atual está vazio ou não - ele ainda ficará vazio.
Redimensionando uma tabela hash
Você pode alterar o tamanho de uma tabela hash criando uma tabela maior e inserindo nela elementos não vazios da tabela antiga. Não implementei essa funcionalidade porque queria manter o código de exemplo simples. Além disso, em programas CUDA, a alocação de memória geralmente é feita no código host, e não no kernel CUDA.
o artigo descreve como modificar essa estrutura de dados protegida por bloqueio.
Competitividade
Nos trechos de código de função acima gpu_hashtable_insert(), _lookup() и _delete() processar um par de valores-chave por vez. E mais baixo gpu_hashtable_insert(), _lookup() и _delete() processar uma matriz de pares em paralelo, cada par em um thread de execução de 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 tabela hash resistente a bloqueio oferece suporte a inserções, pesquisas e exclusões simultâneas. Como os pares chave-valor estão sempre em um dos quatro estados e as chaves não se movem, a tabela garante a correção mesmo quando diferentes tipos de operações são usados simultaneamente.
No entanto, se processarmos um lote de inserções e exclusões em paralelo e se a matriz de pares de entrada contiver chaves duplicadas, não seremos capazes de prever quais pares “ganharão” – serão gravados por último na tabela hash. Digamos que chamamos o código de inserção com uma matriz de pares de entrada A/0 B/1 A/2 C/3 A/4. Quando o código for concluído, pares B/1 и C/3 têm a garantia de estar presentes na mesa, mas ao mesmo tempo qualquer um dos pares aparecerá nela A/0, A/2 ou A/4. Isso pode ou não ser um problema – tudo depende da aplicação. Você pode saber antecipadamente que não há chaves duplicadas na matriz de entrada ou pode não se importar com qual valor foi escrito por último.
Se isso for um problema para você, será necessário separar os pares duplicados em diferentes chamadas de sistema CUDA. No CUDA, qualquer operação que chame o kernel sempre é concluída antes da próxima chamada do kernel (pelo menos dentro de um thread. Em threads diferentes, os kernels são executados em paralelo). No exemplo acima, se você chamar um kernel com A/0 B/1 A/2 C/3, e outro com A/4, então a chave A receberá o valor 4.
Agora vamos falar sobre se as funções devem lookup() и delete() use um ponteiro simples ou volátil para uma matriz de pares na tabela hash. Afirma que:
O compilador pode optar por otimizar leituras e gravações na memória global ou compartilhada... Essas otimizações podem ser desabilitadas usando a palavra-chave
volatile: ... qualquer referência a esta variável é compilada em uma instrução de leitura ou gravação de memória real.
Considerações de correção não requerem aplicação volatile. Se o thread de execução usar um valor armazenado em cache de uma operação de leitura anterior, ele usará informações um pouco desatualizadas. Mas ainda assim, esta é uma informação do estado correto da tabela hash em um determinado momento da chamada do kernel. Se precisar usar as informações mais recentes, você pode usar o índice volatile, mas o desempenho diminuirá um pouco: de acordo com meus testes, ao excluir 32 milhões de elementos, a velocidade diminuiu de 500 milhões de exclusões/s para 450 milhões de exclusões/s.
Desempenho
No teste de inserção de 64 milhões de elementos e exclusão de 32 milhões deles, a competição entre std::unordered_map e praticamente não há tabela hash para a GPU:
std::unordered_map gastou 70 ms inserindo e removendo elementos e depois liberando-os unordered_map (livrar-se de milhões de elementos leva muito tempo, porque dentro unordered_map múltiplas alocações de memória são feitas). Honestamente falando, std:unordered_map restrições completamente diferentes. É um thread de execução de CPU único, suporta valores-chave de qualquer tamanho, funciona bem em altas taxas de utilização e mostra desempenho estável após múltiplas exclusões.
A duração da tabela hash para a GPU e a comunicação entre programas foi de 984 ms. Isso inclui o tempo gasto colocando a tabela na memória e excluindo-a (alocando 1 GB de memória uma vez, o que leva algum tempo em CUDA), inserindo e excluindo elementos e iterando sobre eles. Todas as cópias de e para a memória da placa de vídeo também são levadas em consideração.
A própria tabela hash levou 271 ms para ser concluída. Isso inclui o tempo gasto pela placa de vídeo inserindo e excluindo elementos e não leva em consideração o tempo gasto copiando para a memória e iterando na tabela resultante. Se a tabela GPU durar muito tempo, ou se a tabela hash estiver contida inteiramente na memória da placa de vídeo (por exemplo, para criar uma tabela hash que será usada por outro código GPU e não pelo processador central), então o resultado do teste é relevante.
A tabela hash para uma placa de vídeo demonstra alto desempenho devido ao alto rendimento e paralelização ativa.
Contras:
A arquitetura da tabela hash tem alguns problemas que você deve conhecer:
- A análise linear é dificultada pelo agrupamento, o que faz com que as chaves na tabela não sejam posicionadas de forma perfeita.
- As chaves não são removidas usando a função
deletee com o tempo eles bagunçam a mesa.
Como resultado, o desempenho de uma tabela hash pode degradar gradualmente, especialmente se ela existir por muito tempo e tiver inúmeras inserções e exclusões. Uma maneira de mitigar essas desvantagens é fazer o novo hash em uma nova tabela com uma taxa de utilização bastante baixa e filtrar as chaves removidas durante o novo hash.
Para ilustrar os problemas descritos, usarei o código acima para criar uma tabela com 128 milhões de elementos e percorrer 4 milhões de elementos até preencher 124 milhões de slots (taxa de utilização de cerca de 0,96). Aqui está a tabela de resultados, cada linha é uma chamada do kernel CUDA para inserir 4 milhões de novos elementos em uma tabela hash:
Taxa de uso
Duração da inserção 4 elementos
0,00
11,608448 ms (361,314798 milhões de chaves/seg.)
0,03
11,751424 ms (356,918799 milhões de chaves/seg.)
0,06
11,942592 ms (351,205515 milhões de chaves/seg.)
0,09
12,081120 ms (347,178429 milhões de chaves/seg.)
0,12
12,242560 ms (342,600233 milhões de chaves/seg.)
0,16
12,396448 ms (338,347235 milhões de chaves/seg.)
0,19
12,533024 ms (334,660176 milhões de chaves/seg.)
0,22
12,703328 ms (330,173626 milhões de chaves/seg.)
0,25
12,884512 ms (325,530693 milhões de chaves/seg.)
0,28
13,033472 ms (321,810182 milhões de chaves/seg.)
0,31
13,239296 ms (316,807174 milhões de chaves/seg.)
0,34
13,392448 ms (313,184256 milhões de chaves/seg.)
0,37
13,624000 ms (307,861434 milhões de chaves/seg.)
0,41
13,875520 ms (302,280855 milhões de chaves/seg.)
0,44
14,126528 ms (296,909756 milhões de chaves/seg.)
0,47
14,399328 ms (291,284699 milhões de chaves/seg.)
0,50
14,690304 ms (285,515123 milhões de chaves/seg.)
0,53
15,039136 ms (278,892623 milhões de chaves/seg.)
0,56
15,478656 ms (270,973402 milhões de chaves/seg.)
0,59
15,985664 ms (262,379092 milhões de chaves/seg.)
0,62
16,668673 ms (251,627968 milhões de chaves/seg.)
0,66
17,587200 ms (238,486174 milhões de chaves/seg.)
0,69
18,690048 ms (224,413765 milhões de chaves/seg.)
0,72
20,278816 ms (206,831789 milhões de chaves/seg.)
0,75
22,545408 ms (186,038058 milhões de chaves/seg.)
0,78
26,053312 ms (160,989275 milhões de chaves/seg.)
0,81
31,895008 ms (131,503463 milhões de chaves/seg.)
0,84
42,103294 ms (99,619378 milhões de chaves/seg.)
0,87
61,849056 ms (67,815164 milhões de chaves/seg.)
0,90
105,695999 ms (39,682713 milhões de chaves/seg.)
0,94
240,204636 ms (17,461378 milhões de chaves/seg.)
À medida que a utilização aumenta, o desempenho diminui. Isto não é desejável na maioria dos casos. Se um aplicativo inserir elementos em uma tabela e depois os descartar (por exemplo, ao contar palavras em um livro), isso não será um problema. Mas se o aplicativo usar uma tabela hash de longa duração (por exemplo, em um editor gráfico para armazenar partes não vazias de imagens onde o usuário insere e exclui informações com frequência), esse comportamento pode ser problemático.
E mediu a profundidade de sondagem da tabela hash após 64 milhões de inserções (fator de utilização 0,5). A profundidade média foi de 0,4774, então a maioria das chaves estava no melhor slot possível ou a um slot de distância da melhor posição. A profundidade máxima de sondagem foi de 60.
Em seguida, medi a profundidade de sondagem em uma mesa com 124 milhões de pastilhas (fator de utilização 0,97). A profundidade média já era 10,1757, e a máxima - 6474 (!!). O desempenho da detecção linear cai significativamente em altas taxas de utilização.
É melhor manter baixa a taxa de utilização desta tabela hash. Mas então aumentamos o desempenho às custas do consumo de memória. Felizmente, no caso de chaves e valores de 32 bits, isso pode ser justificado. Se no exemplo acima, em uma tabela com 128 milhões de elementos, mantivermos o fator de utilização de 0,25, então não poderemos colocar mais de 32 milhões de elementos nela, e os 96 milhões de slots restantes serão perdidos - 8 bytes para cada par , 768 MB de memória perdida.
Observe que estamos falando da perda de memória da placa de vídeo, que é um recurso mais valioso que a memória do sistema. Embora a maioria das placas gráficas de desktop modernas que suportam CUDA tenham pelo menos 4 GB de memória (no momento em que este artigo foi escrito, a NVIDIA 2080 Ti tinha 11 GB), ainda não seria a decisão mais sábia perder tais quantidades.
Posteriormente escreverei mais sobre a criação de tabelas hash para placas de vídeo que não apresentam problemas de profundidade de teste, bem como maneiras de reutilizar slots excluídos.
Medição de profundidade de sondagem
Para determinar a profundidade de sondagem de uma chave, podemos extrair o hash da chave (seu índice de tabela ideal) do índice de tabela real:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Devido à magia dos números binários em complemento de dois e ao fato de que a capacidade da tabela hash é dois elevado a dois, essa abordagem funcionará mesmo quando o índice chave for movido para o início da tabela. Vamos pegar uma chave com hash 1, mas inserida no slot 3. Então, para uma mesa com capacidade 4, obtemos (3 — 1) & 3, que é equivalente a 2.
Conclusão
Se você tiver dúvidas ou comentários, envie-me um e-mail para ou abra um novo tópico em.
Este código foi escrito inspirado em excelentes artigos:
No futuro, continuarei escrevendo sobre implementações de tabelas hash para placas de vídeo e analisando seu desempenho. Meus planos incluem encadeamento, hash Robin Hood e hash cuco usando operações atômicas em estruturas de dados compatíveis com GPU.
Fonte: habr.com