• Es werden nur 32-Bit-Schlüssel und dieselben Werte verarbeitet.
• Die Hash-Tabelle hat eine feste Größe.
• Und diese Größe sollte gleich zwei Grad sein.

Hash-Tabelle ohne Sperren

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Hash-Tabellenstatus

• Der Schlüssel und die Bedeutung sind leer. In diesem Zustand wird die Hash-Tabelle initialisiert.
• Der Schlüssel wurde aufgezeichnet, aber der Wert ist noch nicht. Wenn in diesem Moment ein anderer Ausführungsthread Daten liest, wird ein leerer Wert zurückgegeben. Dies ist normal, dasselbe würde passieren, wenn ein anderer Ausführungsthread etwas früher funktioniert und wir über eine wettbewerbsfähige Datenstruktur sprechen.
• Sowohl der Schlüssel als auch der Wert werden aufgezeichnet.
• Der Wert ist für andere Ausführungsthreads verfügbar, der Schlüssel jedoch noch nicht. Dies kann passieren, weil das CUDA-Programmiermodell ein schlecht geordnetes Speichermodell impliziert. Dies ist normal, auf jeden Fall ist der Schlüssel noch leer, auch wenn der Wert nicht mehr so ​​ist.

In eine Hash-Tabelle einfügen

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

Hash-Tabellensuche

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

Löschen der Hash-Tabelle

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

Ändern Sie die Größe einer Hash-Tabelle

Wettbewerbsfähigkeit

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

Der Compiler kann nach eigenem Ermessen die Lese- und Schreibvorgänge im globalen oder gemeinsam genutzten Speicher optimieren. Diese Optimierungen können mit dem Schlüsselwort deaktiviert werden volatile: ... Jeder Link zu dieser Variablen wird zu einer echten Lese- oder Schreibanweisung im Speicher kompiliert.

Performance

Nachteile

• Clustering stört die lineare Abtastung, weshalb die Schlüssel in der Tabelle alles andere als ideal sind.
• Schlüssel werden nicht mit der Funktion gelöscht deleteund überladen mit der Zeit die Tabelle.

Tiefenmessung erfassen

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

Fazit

• Die einfachste sperrenfreie Hash-Tabelle der Welt
• Eine sperrfreie wartungsfreie Hash-Tabelle