• Hanya kunci 32-bit dan nilai yang sama diproses.
• Tabel hash memiliki ukuran tetap.
• Dan ukuran ini harus sama dengan dua derajat.

Meja hash tanpa kunci

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Status Tabel Hash

• Kuncinya dan artinya kosong. Dalam keadaan ini, tabel hash diinisialisasi.
• Kunci telah direkam, tetapi nilainya belum. Jika utas eksekusi lainnya sedang membaca data pada saat itu, maka itu mengembalikan nilai kosong. Ini normal, hal yang sama akan terjadi jika utas eksekusi lainnya bekerja sedikit lebih awal, dan kita berbicara tentang struktur data kompetitif.
• Baik kunci maupun nilainya dicatat.
• Nilai tersedia untuk utas eksekusi lainnya, tetapi kuncinya belum. Ini bisa terjadi karena model pemrograman CUDA menyiratkan model memori yang tidak tertata dengan baik. Ini normal, dalam hal apapun, kuncinya masih kosong, bahkan jika nilainya tidak lagi seperti itu.

Masukkan ke dalam tabel hash

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

Pencarian Tabel Hash

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

Penghapusan tabel hash

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

Ubah ukuran tabel hash

Daya saing

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

Kompiler dapat, atas kebijakannya sendiri, mengoptimalkan operasi baca dan tulis ke memori global atau bersama ... Optimalisasi ini dapat dinonaktifkan menggunakan kata kunci volatile: ... tautan apa pun ke variabel ini dikompilasi menjadi instruksi baca atau tulis nyata dalam memori.

Performa

kerugian

• Clustering mengganggu linear probing, karena itu kunci dalam tabel jauh dari ideal.
• Tombol tidak dihapus menggunakan fungsi deletedan seiring waktu mengacaukan tabel.

Pengukuran Kedalaman Sensing

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

Kesimpulan

• Tabel Hash Bebas-Kunci Paling Sederhana di Dunia
• Meja Hash Bebas-Tunggu-Bebas-Kunci