• Solo se procesan claves de 32 bits y los mismos valores.
• La tabla hash tiene un tamaño fijo.
• Y este tamaño debe ser igual a dos en grado.

Mesa hash sin cerraduras

struct KeyValue
{
    uint32_t key;
    uint32_t value;
};

Estado de la tabla hash

• La clave y el significado están vacíos. En este estado, la tabla hash se inicializa.
• La clave ha sido registrada, pero el valor aún no. Si otro hilo de ejecución está leyendo datos en ese momento, entonces devuelve un valor vacío. Esto es normal, lo mismo sucedería si otro hilo de ejecución funcionara un poco antes, y estamos hablando de una estructura de datos competitiva.
• Se registran tanto la clave como el valor.
• El valor está disponible para otros hilos de ejecución, pero la clave aún no lo está. Esto puede suceder porque el modelo de programación CUDA implica un modelo de memoria mal ordenado. Esto es normal; en cualquier caso, la clave todavía está vacía, incluso si el valor ya no es tal.

Insertar en una tabla 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);
    }
}

Búsqueda de tabla 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);
        }
}

Eliminación de tabla 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);
    }
}

Cambiar el tamaño de una tabla hash

Competitividad

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

El compilador puede, a su discreción, optimizar las operaciones de lectura y escritura en la memoria global o compartida ... Estas optimizaciones pueden deshabilitarse usando la palabra clave volatile: ... cualquier enlace a esta variable se compila en una instrucción real de lectura o escritura en la memoria.

Actuación

desventajas

• El agrupamiento interfiere con el sondeo lineal, por lo que las claves en la tabla están lejos de ser ideales.
• Las teclas no se eliminan mediante la función deletey con el tiempo desordenan la tabla.

Medición de profundidad de detección

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

Conclusión

• La tabla de hash sin bloqueo más simple del mundo
• Una tabla de hash sin bloqueo y sin espera