Postei no Github um novo projeto chamado A Simple GPU Hash Table .Esta é uma tabela de hash simples para a GPU, capaz de processar centenas de milhões de inserções por segundo. No meu laptop com um 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 do laptop é de aproximadamente 300 milhões de inserções / s e 500 milhões de remoções / s.A tabela está escrita em CUDA, embora a mesma técnica possa ser aplicada ao HLSL ou GLSL. A implementação possui várias limitações que garantem alto desempenho na placa de vídeo:- Somente chaves de 32 bits e os mesmos valores são processados.
- A tabela de hash tem um tamanho fixo.
- E esse tamanho deve ser igual a dois em grau.
Para chaves e valores, você precisa reservar um marcador de delimitação simples (no código acima, é 0xffffffff).Mesa de hash sem fechaduras
A tabela de hash usa endereçamento aberto com detecção linear , ou seja, é apenas uma matriz de pares de valores-chave armazenados na memória e com excelente desempenho de cache. Esse não é o caso do encadeamento, o que significa procurar um ponteiro em uma lista vinculada. Uma tabela de hash é uma matriz simples que armazena elementos KeyValue:struct KeyValue
{
uint32_t key;
uint32_t value;
};
O tamanho da tabela é igual a dois em potência, e não um número primo, porque para usar pow2 / AND-mask, basta uma instrução rápida e o operador do módulo é muito mais lento. Isso é importante no caso da detecção linear, pois em uma pesquisa linear na tabela, o índice do slot deve ser empacotado em cada slot. E, como resultado, o custo da operação é adicionado ao módulo em cada slot.A tabela armazena apenas a chave e o valor de cada item, não o 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 alguns turnos, XORs e multiplicações.A tabela de hash usa uma técnica de proteção de bloqueio que não depende da ordem de colocação da memória. Mesmo que algumas operações de gravação violem a ordem de outras operações, a tabela de hash ainda manterá o estado correto. Falaremos sobre isso abaixo. A técnica funciona muito bem com placas de vídeo nas quais milhares de threads competem.Chaves e valores na tabela de hash são inicializados para esvaziar.O código pode ser modificado para poder processar chaves e valores de 64 bits. As chaves requerem operações atômicas de leitura, gravação e troca (comparar e trocar). 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 aqui) e as placas de vídeo modernas oferecem suporte a operações atômicas de 64 bits comparadas à troca. Obviamente, ao mudar para 64 bits, o desempenho diminuirá um pouco.Status da tabela de hash
Cada par de valores-chave em uma tabela de hash pode ter um dos quatro estados:- A chave e o significado estão vazios. Nesse estado, a tabela de hash é inicializada.
- A chave foi gravada, mas o valor ainda não é. Se outro encadeamento de execução estiver lendo dados nesse momento, ele retornará um valor vazio. Isso é normal, o mesmo aconteceria se outro segmento de execução funcionasse um pouco mais cedo, e estamos falando de uma estrutura de dados competitiva.
- A chave e o valor são registrados.
- O valor está disponível para outros encadeamentos de execução, mas a chave ainda não está. Isso pode acontecer porque o modelo de programação CUDA implica um modelo de memória mal ordenado. Isso é normal; de qualquer forma, a chave ainda está vazia, mesmo que o valor não seja mais esse.
Uma nuance importante é que, assim que a chave é gravada no slot, ela não se move mais - mesmo que a chave seja excluída, falaremos sobre isso abaixo.O código da tabela de hash funciona mesmo com modelos de memória mal ordenados que não sabem a ordem de leitura e gravação na memória. Ao analisar a inserção, pesquisar e excluir na tabela de hash, lembre-se de que cada par de valor-chave está em um dos quatro estados descritos acima.Inserir em uma tabela de hash
Uma função CUDA que insere pares de valores-chave em uma tabela de hash é assim: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 itera sobre a matriz da tabela de hash, começando com o hash da chave inserida. Em cada slot da matriz, uma operação de comparação atômica é realizada com a central, na qual a chave nesse slot é comparada com uma vazia. Se uma incompatibilidade for detectada, a chave no slot é atualizada para a chave inserida e a chave original do slot é retornada. Se essa chave original estava vazia ou correspondeu à chave inserida, o código encontrou um slot adequado para inserção e traz o valor inserido para o slot.Se em uma chamada do kernelgpu_hashtable_insert()Existem vários elementos com a mesma chave, então qualquer um dos seus valores pode ser gravado no slot de chave. Isso é considerado normal: uma das operações de gravação de valor-chave durante a chamada será bem-sucedida, mas como tudo isso acontece em paralelo em vários encadeamentos de execução, não podemos prever qual operação de gravação na memória será a última.Pesquisa de tabela de hash
Código do localizador de 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 da chave armazenada na tabela, iteramos sobre a matriz começando com o hash da chave desejada. Em cada slot, verificamos se a chave é a que estamos procurando e, se for o caso, retornamos seu valor. Também verificamos se a chave está vazia e, se estiver, interrompemos a pesquisa.Se não conseguirmos encontrar a chave, o código retornará um valor vazio.Todas essas operações de pesquisa podem ser executadas competitivamente durante inserções e exclusões. Cada par na tabela terá um dos quatro estados descritos acima para o fluxo.Exclusão de tabela de hash
Código de remoção de chave: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 remoção de uma chave é incomum: deixamos a chave na tabela e marcamos seu valor (não a própria chave) em branco. Esse código é muito semelhante a lookup(), exceto que, quando uma correspondência é encontrada para a chave, ele torna seu valor vazio.Como mencionado acima, uma vez que a chave é gravada no slot, ela não se move mais. Mesmo quando você exclui um item da tabela, a chave permanece no lugar, apenas seu valor fica vazio. Isso significa que não precisamos usar a operação atômica para escrever o valor do slot, porque não importa se o valor atual está vazio ou não - ele ainda ficará vazio.Redimensionar uma tabela de hash
Você pode redimensionar a tabela de hash criando uma tabela maior e inserindo elementos não vazios da tabela antiga nela. Eu não implementei essa funcionalidade porque queria manter o código de exemplo simples. Além disso, nos programas CUDA, a alocação de memória geralmente é feita no código do host e não no núcleo do CUDA.O artigo Tabela de hash sem bloqueio e sem espera descreve como alterar uma estrutura de dados protegida por bloqueio.Competitividade
Nos fragmentos acima de código, as funções gpu_hashtable_insert(), _lookup()e _delete()processo de um par de chaves de valor de cada vez. E abaixo gpu_hashtable_insert(), _lookup()eles _delete()processam uma matriz de pares em paralelo, cada par em um thread de execução da GPU separado:
uint32_t threadblocksize = 1024;
uint32_t gridsize = (numkvs + threadblocksize - 1) / threadblocksize;
gpu_hashtable_insert_kernel<<<gridsize, threadblocksize>>>(hashtable, kvs, numkvs);
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);
}
}
Uma tabela de hash bloqueável suporta inserções, pesquisas e exclusões simultâneas. Como os pares de valores-chave estão sempre em um dos quatro estados e as chaves não se movem, a tabela garante a correção, mesmo ao usar diferentes tipos de operações.No entanto, se processarmos um pacote de inserções e exclusões em paralelo, e se chaves duplicadas estiverem contidas na matriz de entrada de pares, não conseguiremos prever quais pares serão "vencidos" - eles serão gravados na tabela de hash por último. Suponha que tenhamos chamado um código de inserção com uma matriz de entrada de pares A/0 B/1 A/2 C/3 A/4. Quando o código é concluído, os pares B/1têm a C/3garantia de estar presentes na tabela, mas ao mesmo tempo qualquer um dos pares aparecerá nele A/0, A/2ouA/4. Isso pode ou não ser um problema - tudo depende do aplicativo. Você deve saber antecipadamente que não há chaves duplicadas na matriz de entrada ou talvez não seja importante para você qual valor foi gravado pela última vez.Se isso for um problema para você, será necessário dividir os pares duplicados em diferentes chamadas CUDA do sistema. No CUDA, qualquer operação de chamada do kernel sempre termina antes da próxima chamada do kernel (pelo menos dentro do mesmo thread. Em threads diferentes, o kernel é executado em paralelo). Se no exemplo acima, chame um núcleo com A/0 B/1 A/2 C/3e o outro com A/4, a chave Aobterá um valor 4.Agora vamos falar se a função lookup()e delete()usar um ponteiro simples (simples) ou variável (volátil) para uma matriz de pares em uma tabela de hash.A documentação da CUDA afirma que:O compilador pode, a seu critério, otimizar as operações de leitura e gravação na memória global ou compartilhada ... Essas otimizações podem ser desativadas usando a palavra volatile- chave : ... qualquer link para essa variável é compilado em uma instrução real de leitura ou gravação na memória.
As considerações de correção não requerem aplicação volatile. Se o encadeamento de execução usar o valor em cache de uma operação de leitura anterior, isso significa que ele usará informações um pouco desatualizadas. Mas ainda assim, essas são informações do estado correto da tabela de hash em um determinado ponto da chamada do kernel. Se você precisar usar as informações mais recentes, poderá usar o ponteiro volatile, mas o desempenho diminuirá um pouco: de acordo com meus testes, quando você exclui 32 milhões de itens, a velocidade diminui de 500 milhões de exclusões para 450 milhões de exclusões.atuação
No teste para inserir 64 milhões de elementos e excluir 32 milhões deles, std::unordered_mappraticamente não há concorrência entre e a tabela de hash para a GPU:std::unordered_mapGastou 70 691 ms na inserção e exclusão de elementos com liberação subsequente unordered_map(a liberação de milhões de elementos leva muito tempo, porque unordered_mapmuitas alocações de memória são executadas no interior ). Honestamente, existem std:unordered_maplimitações completamente diferentes. Este é um único thread de execução da CPU, suporta valores-chave de qualquer tamanho, funciona bem em altas taxas de utilização e mostra desempenho estável após inúmeras exclusões.A duração da tabela de hash para a comunicação GPU e entre programas foi de 984 ms. Isso inclui o tempo gasto para colocar a tabela na memória e excluí-la (alocação única de 1 GB de memória, que no CUDA leva algum tempo), inserção e exclusão de elementos e também iteração sobre eles. Também foram levadas em consideração todas as cópias de e para a memória da placa de vídeo.A própria tabela de hash levou 271 ms. Isso inclui o tempo gasto pela placa de vídeo para inserir e remover itens e não leva em consideração o tempo necessário para copiar para a memória e iterar sobre a tabela resultante. Se a tabela da GPU permanecer por muito tempo ou se a tabela de hash estiver totalmente contida na memória da placa de vídeo (por exemplo, para criar uma tabela de hash que será usada por outro código da GPU e não pelo processador central), o resultado do teste é relevante.A tabela de hash da placa de vídeo demonstra alto desempenho devido à sua alta largura de banda e paralelização ativa.desvantagens
A arquitetura da tabela de hash tem vários problemas a serem lembrados:- O agrupamento interfere na análise linear, pelo que as chaves da tabela estão longe do ideal.
- As teclas não são excluídas usando a função
deletee, com o tempo, desorganizam a mesa.
Como resultado, o desempenho da tabela de hash pode diminuir gradualmente, especialmente se ele existir por um longo tempo e várias inserções e exclusões forem realizadas nela. Uma maneira de mitigar essas deficiências é refazer uma nova tabela com uma taxa de utilização razoavelmente baixa e filtrar as chaves remotas ao refazer a tarefa.Para ilustrar os problemas descritos, eu uso o código acima para criar uma tabela para 128 milhões de elementos, inserirei ciclicamente 4 milhões de elementos até preencher 124 milhões de slots (a utilização é de cerca de 0,96). Aqui está a tabela de resultados, cada linha é uma chamada para o núcleo CUDA com a inserção de 4 milhões de novos elementos em uma tabela de hash:À medida que a utilização aumenta, a produtividade diminui. Isso é indesejável na maioria dos casos. Se um aplicativo inserir elementos em uma tabela e depois descartá-los (por exemplo, ao contar palavras em um livro), isso não será um problema. Mas se o aplicativo usar uma tabela de hash de longa duração (por exemplo, em um editor de gráficos para armazenar partes não vazias de imagens quando o usuário geralmente inserir e excluir informações), esse comportamento poderá ser problemático.E ele mediu a profundidade da análise da tabela de hash após 64 milhões de inserções (fator de utilização 0,5). A profundidade média foi de 0,4774; portanto, a maioria das teclas estava localizada no melhor slot possível ou em um slot da melhor posição. A profundidade máxima do som foi 60.Depois, medi a profundidade do som na tabela com 124 milhões de inserções (taxa de utilização 0,97). A profundidade média já era 10.1757, e a máxima - 6474 (!!). O desempenho do som linear diminui drasticamente a altas taxas de utilização.É melhor manter essa tabela de hash baixa. Mas aumentamos a produtividade consumindo memória. Felizmente, no caso de chaves e valores de 32 bits, isso pode ser justificado. Se no exemplo acima na tabela para 128 milhões de elementos o coeficiente de utilização de 0,25 estiver armazenado, não podemos colocar mais de 32 milhões de elementos nele 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 sobre a 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 possuam pelo menos 4 GB de memória (no momento da escrita, o NVIDIA 2080 Ti possui 11 GB), a perda desses volumes não será a decisão mais sensata.Posteriormente, escreverei mais sobre a criação de tabelas de hash para placas de vídeo que não apresentam problemas com a profundidade da sonoridade, bem como maneiras de reutilizar slots remotos.Medição da profundidade da detecção
Para determinar a profundidade do som da chave, podemos extrair o hash da chave (seu índice ideal na tabela) do seu índice real da tabela:
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Devido à magia de dois números binários no código adicional e ao fato de a capacidade da tabela de hash ser igual a dois no poder, essa abordagem funcionará mesmo quando o índice de chave for movido para o início da tabela. Pegue uma chave com hash em 1, mas inserida no slot 3. Em uma tabela com capacidade 4, obtemos (3 — 1) & 3o equivalente a 2.Conclusão
Se você tiver perguntas ou comentários, escreva-me no Twitter ou abra um novo tópico no repositório .Este código é inspirado em alguns ótimos artigos:No futuro, continuarei escrevendo sobre implementações de tabelas de hash para placas de vídeo e analisarei seu desempenho. Tenho planos de encadear, hash Robin Hood e cuco usando operações atômicas em estruturas de dados que são convenientes para placas de vídeo.