Lo publiqué en Github .
Es una tabla hash de GPU simple capaz de procesar cientos de millones de inserciones por segundo. En mi computadora portátil NVIDIA GTX 1060, el código inserta 64 millones de pares clave-valor generados aleatoriamente en aproximadamente 210 ms y elimina 32 millones de pares en aproximadamente 64 ms.
Es decir, la velocidad en una computadora portátil es de aproximadamente 300 millones de inserciones/seg y 500 millones de eliminaciones/seg.
La tabla está escrita en CUDA, aunque se puede aplicar la misma técnica a HLSL o GLSL. La implementación tiene varias limitaciones para garantizar un alto rendimiento en una tarjeta de video:
- Sólo 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 elevado a la potencia.
Para claves y valores, debe reservar un marcador delimitador simple (en el código anterior es 0xffffffff).
Tabla hash sin cerraduras
La tabla hash utiliza direccionamiento abierto con , es decir, es simplemente una matriz de pares clave-valor que se almacena en la memoria y tiene un rendimiento de caché superior. No se puede decir lo mismo del encadenamiento, que implica buscar un puntero en una lista enlazada. Una tabla hash es una matriz simple que almacena elementos. KeyValue:
struct KeyValue
{
uint32_t key;
uint32_t value;
};
El tamaño de la tabla es una potencia de dos, no un número primo, porque una instrucción rápida es suficiente para aplicar la máscara pow2/AND, pero el operador de módulo es mucho más lento. Esto es importante en el caso del sondeo lineal, ya que en una búsqueda de tabla lineal el índice de ranura debe estar incluido en cada ranura. Y como resultado, el coste de la operación se suma módulo en cada slot.
La tabla solo almacena la clave y el valor de cada elemento, no un hash de la clave. Dado que la tabla sólo almacena claves de 32 bits, el hash se calcula muy rápidamente. El código anterior utiliza el hash Murmur3, que solo realiza unos pocos cambios, XOR y multiplicaciones.
La tabla hash utiliza técnicas de protección de bloqueo que son independientes del orden de la memoria. Incluso si algunas operaciones de escritura alteran el orden de otras operaciones similares, la tabla hash seguirá manteniendo el estado correcto. Hablaremos de esto a continuación. La técnica funciona muy bien con tarjetas de vídeo que ejecutan miles de subprocesos al mismo tiempo.
Las claves y los valores de la tabla hash se inicializan para vaciar.
El código se puede modificar para manejar también claves y valores de 64 bits. Las claves requieren operaciones atómicas de lectura, escritura y comparación e intercambio. Y los valores requieren operaciones atómicas de lectura y escritura. Afortunadamente, en CUDA, las operaciones de lectura y escritura para valores de 32 y 64 bits son atómicas siempre que estén alineadas de forma natural (ver más abajo). ), y las tarjetas de video modernas admiten operaciones de comparación e intercambio atómico de 64 bits. Eso sí, al pasar a 64 bits, el rendimiento disminuirá ligeramente.
Estado de la tabla hash
Cada par clave-valor en una tabla hash puede tener uno de cuatro estados:
- La clave y el valor están vacíos. En este estado, se inicializa la tabla hash.
- La clave se ha anotado, pero el valor aún no se ha anotado. Si otro hilo está leyendo datos actualmente, regresa vacío. Esto es normal, hubiera pasado lo mismo si otro hilo de ejecución hubiera funcionado un poco antes, y estamos hablando de una estructura de datos concurrente.
- Se registran tanto la clave como el valor.
- El valor está disponible para otros subprocesos de ejecución, pero la clave aún no. Esto puede suceder porque el modelo de programación CUDA tiene un modelo de memoria poco ordenado. Esto es normal, en cualquier caso la clave sigue vacía, aunque el valor ya no lo esté.
Un matiz importante es que una vez que la clave se ha escrito en la ranura, ya no se mueve; incluso si la clave se elimina, hablaremos de esto a continuación.
El código de la tabla hash incluso funciona con modelos de memoria poco ordenados en los que se desconoce el orden en que se lee y escribe la memoria. Mientras observamos la inserción, búsqueda y eliminación en una tabla hash, recuerde que cada par clave-valor se encuentra en uno de los cuatro estados descritos anteriormente.
Insertar en una tabla hash
La función CUDA que inserta pares clave-valor en una tabla hash se ve así:
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 insertar una clave, el código recorre en iteración la matriz de la tabla hash comenzando con el hash de la clave insertada. Cada ranura de la matriz realiza una operación atómica de comparación e intercambio que compara la clave de esa ranura con la vacía. Si se detecta una discrepancia, la clave en la ranura se actualiza con la clave insertada y luego se devuelve la clave de la ranura original. Si esta clave original estaba vacía o coincidía con la clave insertada, entonces el código encontró una ranura adecuada para insertar e insertó el valor insertado en la ranura.
Si en una llamada al kernel gpu_hashtable_insert() Si hay varios elementos con la misma clave, cualquiera de sus valores se puede escribir en la ranura de la clave. Esto se considera normal: una de las escrituras clave-valor durante la llamada tendrá éxito, pero como todo esto sucede en paralelo dentro de varios subprocesos de ejecución, no podemos predecir qué escritura en memoria será la última.
Búsqueda de tabla hash
Código para buscar claves:
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 el valor de una clave almacenada en una tabla, iteramos a través del array comenzando con el hash de la clave que estamos buscando. En cada slot comprobamos si la clave es la que buscamos y en caso afirmativo devolvemos su valor. También comprobamos si la clave está vacía, y en caso afirmativo abortamos la búsqueda.
Si no podemos encontrar la clave, el código devuelve un valor vacío.
Todas estas operaciones de búsqueda se pueden realizar simultáneamente mediante inserciones y eliminaciones. Cada par de la tabla tendrá uno de los cuatro estados descritos anteriormente para el flujo.
Eliminar en una tabla hash
Código para eliminar claves:
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);
}
}
La eliminación de una clave se realiza de una manera inusual: dejamos la clave en la tabla y marcamos su valor (no la clave en sí) como vacío. Este código es muy similar a lookup(), excepto que cuando se encuentra una coincidencia en una clave, su valor queda vacío.
Como se mencionó anteriormente, una vez que se escribe una clave en una ranura, ya no se mueve. Incluso cuando se elimina un elemento de la tabla, la clave permanece en su lugar, su valor simplemente queda vacío. Esto significa que no necesitamos usar una operación de escritura atómica para el valor de la ranura, porque no importa si el valor actual está vacío o no; aun así quedará vacío.
Cambiar el tamaño de una tabla hash
Puede cambiar el tamaño de una tabla hash creando una tabla más grande e insertando en ella elementos no vacíos de la tabla anterior. No implementé esta funcionalidad porque quería mantener el código de muestra simple. Además, en los programas CUDA, la asignación de memoria a menudo se realiza en el código host en lugar de en el kernel CUDA.
El artículo describe cómo modificar dicha estructura de datos protegida por bloqueo.
Competitividad
En los fragmentos de código de función anteriores gpu_hashtable_insert(), _lookup() и _delete() procesar un par clave-valor a la vez. y mas bajo gpu_hashtable_insert(), _lookup() и _delete() Procese una serie de pares en paralelo, cada par en un hilo de ejecución de GPU separado:
// 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);
}
}
La tabla hash resistente a bloqueos admite inserciones, búsquedas y eliminaciones simultáneas. Debido a que los pares clave-valor siempre están en uno de cuatro estados y las claves no se mueven, la tabla garantiza la corrección incluso cuando se utilizan diferentes tipos de operaciones simultáneamente.
Sin embargo, si procesamos un lote de inserciones y eliminaciones en paralelo, y si la matriz de pares de entrada contiene claves duplicadas, entonces no podremos predecir qué pares “ganarán”: se escribirán en último lugar en la tabla hash. Digamos que llamamos al código de inserción con una matriz de entrada de pares A/0 B/1 A/2 C/3 A/4. Cuando el código se completa, los pares B/1 и C/3 Se garantiza que estarán presentes en la tabla, pero al mismo tiempo cualquiera de los pares aparecerá en ella. A/0, A/2 o A/4. Esto puede ser un problema o no; todo depende de la aplicación. Es posible que sepa de antemano que no hay claves duplicadas en la matriz de entrada o que no le importe qué valor se escribió en último lugar.
Si esto es un problema para usted, entonces necesita separar los pares duplicados en diferentes llamadas al sistema CUDA. En CUDA, cualquier operación que llame al kernel siempre se completa antes de la siguiente llamada al kernel (al menos dentro de un subproceso. En diferentes subprocesos, los kernels se ejecutan en paralelo). En el ejemplo anterior, si llama a un kernel con A/0 B/1 A/2 C/3, y el otro con A/4, entonces la clave A obtendrá el valor 4.
Ahora hablemos de si las funciones deberían lookup() и delete() use un puntero simple o volátil a una matriz de pares en la tabla hash. Establece que:
El compilador puede optar por optimizar las lecturas y escrituras en la memoria global o compartida... Estas optimizaciones se pueden desactivar usando la palabra clave
volatile: ... cualquier referencia a esta variable se compila en una instrucción de lectura o escritura en memoria real.
Las consideraciones de corrección no requieren aplicación. volatile. Si el hilo de ejecución utiliza un valor almacenado en caché de una operación de lectura anterior, utilizará información ligeramente desactualizada. Pero aún así, esta es información del estado correcto de la tabla hash en un momento determinado de la llamada al kernel. Si necesita utilizar la información más reciente, puede utilizar el índice. volatile, pero luego el rendimiento disminuirá ligeramente: según mis pruebas, al eliminar 32 millones de elementos, la velocidad disminuyó de 500 millones de eliminaciones/seg a 450 millones de eliminaciones/seg.
Rendimiento
En la prueba para insertar 64 millones de elementos y eliminar 32 millones de ellos, la competencia entre std::unordered_map y prácticamente no existe una tabla hash para la GPU:
std::unordered_map pasó 70 ms insertando y quitando elementos y luego liberándolos unordered_map (deshacerse de millones de elementos lleva mucho tiempo, porque dentro unordered_map se realizan múltiples asignaciones de memoria). Hablando honestamente, std:unordered_map restricciones completamente diferentes. Es un subproceso de ejecución de CPU único, admite valores-clave de cualquier tamaño, funciona bien con altas tasas de utilización y muestra un rendimiento estable después de múltiples eliminaciones.
La duración de la tabla hash para la GPU y la comunicación entre programas fue de 984 ms. Esto incluye el tiempo dedicado a colocar la tabla en la memoria y eliminarla (asignando 1 GB de memoria una vez, lo que lleva algo de tiempo en CUDA), insertar y eliminar elementos e iterar sobre ellos. También se tienen en cuenta todas las copias hacia y desde la memoria de la tarjeta de vídeo.
La tabla hash en sí tardó 271 ms en completarse. Esto incluye el tiempo que dedica la tarjeta de video a insertar y eliminar elementos, y no tiene en cuenta el tiempo dedicado a copiar en la memoria e iterar sobre la tabla resultante. Si la tabla de GPU dura mucho tiempo, o si la tabla hash está contenida completamente en la memoria de la tarjeta de video (por ejemplo, para crear una tabla hash que será utilizada por otro código de GPU y no por el procesador central), entonces el resultado de la prueba es relevante.
La tabla hash de una tarjeta de video demuestra un alto rendimiento debido al alto rendimiento y la paralelización activa.
Limitaciones
La arquitectura de la tabla hash tiene algunos problemas que se deben tener en cuenta:
- El sondeo lineal se ve obstaculizado por la agrupación, lo que hace que las claves de la tabla no se coloquen perfectamente.
- Las claves no se eliminan usando la función.
deletey con el tiempo abarrotan la mesa.
Como resultado, el rendimiento de una tabla hash puede degradarse gradualmente, especialmente si existe durante mucho tiempo y tiene numerosas inserciones y eliminaciones. Una forma de mitigar estas desventajas es repetir en una nueva tabla con una tasa de utilización bastante baja y filtrar las claves eliminadas durante el refrito.
Para ilustrar los problemas descritos, usaré el código anterior para crear una tabla con 128 millones de elementos y recorreré 4 millones de elementos hasta llenar 124 millones de espacios (tasa de utilización de aproximadamente 0,96). Aquí está la tabla de resultados, cada fila es una llamada al kernel CUDA para insertar 4 millones de elementos nuevos en una tabla hash:
Tasa de uso
Duración de la inserción 4 elementos
0,00
11,608448 ms (361,314798 millones de claves/seg.)
0,03
11,751424 ms (356,918799 millones de claves/seg.)
0,06
11,942592 ms (351,205515 millones de claves/seg.)
0,09
12,081120 ms (347,178429 millones de claves/seg.)
0,12
12,242560 ms (342,600233 millones de claves/seg.)
0,16
12,396448 ms (338,347235 millones de claves/seg.)
0,19
12,533024 ms (334,660176 millones de claves/seg.)
0,22
12,703328 ms (330,173626 millones de claves/seg.)
0,25
12,884512 ms (325,530693 millones de claves/seg.)
0,28
13,033472 ms (321,810182 millones de claves/seg.)
0,31
13,239296 ms (316,807174 millones de claves/seg.)
0,34
13,392448 ms (313,184256 millones de claves/seg.)
0,37
13,624000 ms (307,861434 millones de claves/seg.)
0,41
13,875520 ms (302,280855 millones de claves/seg.)
0,44
14,126528 ms (296,909756 millones de claves/seg.)
0,47
14,399328 ms (291,284699 millones de claves/seg.)
0,50
14,690304 ms (285,515123 millones de claves/seg.)
0,53
15,039136 ms (278,892623 millones de claves/seg.)
0,56
15,478656 ms (270,973402 millones de claves/seg.)
0,59
15,985664 ms (262,379092 millones de claves/seg.)
0,62
16,668673 ms (251,627968 millones de claves/seg.)
0,66
17,587200 ms (238,486174 millones de claves/seg.)
0,69
18,690048 ms (224,413765 millones de claves/seg.)
0,72
20,278816 ms (206,831789 millones de claves/seg.)
0,75
22,545408 ms (186,038058 millones de claves/seg.)
0,78
26,053312 ms (160,989275 millones de claves/seg.)
0,81
31,895008 ms (131,503463 millones de claves/seg.)
0,84
42,103294 ms (99,619378 millones de claves/seg.)
0,87
61,849056 ms (67,815164 millones de claves/seg.)
0,90
105,695999 ms (39,682713 millones de claves/seg.)
0,94
240,204636 ms (17,461378 millones de claves/seg.)
A medida que aumenta la utilización, el rendimiento disminuye. Esto no es deseable en la mayoría de los casos. Si una aplicación inserta elementos en una tabla y luego los descarta (por ejemplo, al contar palabras en un libro), entonces esto no es un problema. Pero si la aplicación utiliza una tabla hash de larga duración (por ejemplo, en un editor de gráficos para almacenar partes no vacías de imágenes donde el usuario inserta y elimina información con frecuencia), entonces este comportamiento puede ser problemático.
Y midió la profundidad de sondeo de la tabla hash después de 64 millones de inserciones (factor de utilización 0,5). La profundidad promedio fue de 0,4774, por lo que la mayoría de las claves estaban en la mejor ranura posible o a una ranura de la mejor posición. La profundidad máxima de sondeo fue 60.
Luego medí la profundidad de palpado en una mesa con 124 millones de insertos (factor de utilización 0,97). La profundidad media ya era 10,1757, y la máxima - 6474 (!!). El rendimiento de la detección lineal cae significativamente a tasas de utilización elevadas.
Es mejor mantener baja la tasa de utilización de esta tabla hash. Pero luego aumentamos el rendimiento a expensas del consumo de memoria. Afortunadamente, en el caso de claves y valores de 32 bits, esto puede justificarse. Si en el ejemplo anterior, en una tabla con 128 millones de elementos, mantenemos el factor de utilización de 0,25, entonces no podemos colocar más de 32 millones de elementos en ella y los 96 millones de espacios restantes se perderán: 8 bytes por cada par. , 768 MB de memoria perdida.
Tenga en cuenta que estamos hablando de la pérdida de memoria de la tarjeta de video, que es un recurso más valioso que la memoria del sistema. Aunque la mayoría de las tarjetas gráficas de escritorio modernas que admiten CUDA tienen al menos 4 GB de memoria (en el momento de escribir este artículo, la NVIDIA 2080 Ti tiene 11 GB), perder esa cantidad no sería la decisión más inteligente.
Más adelante escribiré más sobre la creación de tablas hash para tarjetas de video que no tienen problemas con la profundidad de sondeo, así como sobre las formas de reutilizar las ranuras eliminadas.
Medición de profundidad de sondeo
Para determinar la profundidad del sondeo de una clave, podemos extraer el hash de la clave (su índice de tabla ideal) de su índice de tabla real:
// get_key_index() -> index of key in hash table
uint32_t probelength = (get_key_index(key) - hash(key)) & (hashtablecapacity-1);
Debido a la magia de los números binarios en complemento a dos y al hecho de que la capacidad de la tabla hash es dos elevado a dos, este enfoque funcionará incluso cuando el índice clave se mueva al principio de la tabla. Tomemos una clave cuyo hash es 1, pero que se inserta en la ranura 3. Luego, para una mesa con capacidad 4 obtenemos (3 — 1) & 3, que equivale a 2.
Conclusión
Si tiene preguntas o comentarios, envíeme un correo electrónico a o abrir un nuevo tema en.
Este código fue escrito inspirándose en excelentes artículos:
En el futuro, continuaré escribiendo sobre implementaciones de tablas hash para tarjetas de video y analizaré su rendimiento. Mis planes incluyen encadenamiento, hash Robin Hood y hash cuco utilizando operaciones atómicas en estructuras de datos que sean compatibles con GPU.
Fuente: habr.com