• Seules les clés 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 degrés.

Table de hachage sans serrures

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

État de la table de hachage

• La clé et le sens sont vides. Dans cet état, la table de hachage est initialisée.
• La clé a été enregistrée, mais la valeur n'est pas encore. Si un autre thread d'exécution lit des données à ce moment, il renvoie une valeur vide. C'est normal, la même chose se produirait si un autre fil d'exécution fonctionnait un peu plus tôt, et nous parlons d'une structure de données compétitive.
• 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 car le modèle de programmation CUDA implique un modèle de mémoire mal ordonné. C'est normal, en tout état de cause, la clé est toujours vide, même si la valeur ne l'est pas.

Insérer dans une table de hachage

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);
    }
}

Recherche de table de hachage

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);
        }
}

Suppression de table de hachage

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);
    }
}

Redimensionner une table de hachage

Compétitivité

// 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);
    }
}

Le compilateur peut, à sa discrétion, optimiser les opérations de lecture et d'écriture dans la mémoire globale ou partagée ... Ces optimisations peuvent être désactivées à l'aide du mot-clé volatile: ... tout lien vers cette variable est compilé en une véritable instruction de lecture ou d'écriture en mémoire.

Performance

désavantages

• Le clustering interfère avec le palpage linéaire, à cause de quoi les clés du tableau sont loin d'être idéales.
• Les clés ne sont pas supprimées à l'aide de la fonction deleteet au fil du temps encombrent la table.

Mesure de la profondeur de détection

// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);

Conclusion

• La table de hachage sans verrou la plus simple au monde
• Une table de hachage sans attente sans verrouillage