memory - gpus - Racionalizando lo que está sucediendo en mi kernel OpenCL simple con respecto a la memoria global

nvidia gpu

Question

const char programSource[] =
        "__kernel void vecAdd(__global int *a, __global int *b, __global int *c)"
        "{"
        "    int gid = get_global_id(0);"
        "for(int i=0; i<10; i++){"
        "    a[gid] = b[gid] + c[gid];}"
        "}";

El kernel anterior es una adición vectorial hecha diez veces por ciclo. He usado la guía de programación y el desbordamiento de la pila para descubrir cómo funciona la memoria global, pero todavía no puedo averiguar mirando mi código si estoy accediendo a la memoria global de una buena manera. Estoy accediendo de forma contigua y estoy adivinando de forma alineada. ¿Carga la tarjeta trozos de 128kb de memoria global para las matrices a, b y c? ¿Carga entonces los trozos de 128kb para cada matriz una vez por cada 32 índices de gid procesados? (4 * 32 = 128) Parece que no estoy desperdiciando ningún ancho de banda de memoria global ¿verdad?

Por cierto, el perfilador de cómputo muestra una eficiencia de gld y gst de 1.00003, lo cual parece extraño, pensé que solo sería 1.0 si todas mis tiendas y cargas se fusionaran. ¿Cómo es por encima de 1.0?

Answer 1

Sí, su patrón de acceso a la memoria es bastante óptimo. Cada semirred es acceder a 16 palabras consecutivas de 32 bits. Además, el acceso está alineado en 64byte, ya que los buffers mismos están alineados y el índice de inicio para cada halfwarp es un múltiplo de 16. Así que cada halfwarp generará una transacción de 64Byte. Por lo tanto, no debe perder ancho de banda de memoria a través de accesos no integrados.

Como solicitó ejemplos en su última pregunta, modifiquemos este código por otro (patrón de acceso menos óptimo (dado que el ciclo en realidad no hace nada, lo ignoraré):

kernel void vecAdd(global int* a, global int* b, global int* c)
{
   int gid = get_global_id(0);
   a[gid+1] = b[gid * 2] + c[gid * 32];
}

Al principio, deja ver cómo funciona esto en el hardware de cómputo 1.3 (GT200)

Para las escrituras en a, esto generará un patrón poco óptimo (siguiendo los medios identificadores identificados por su rango de id y el patrón de acceso correspondiente):

   gid  | addr. offset | accesses     | reasoning
  0- 15 |     4- 67    | 1x128B       | in aligned 128byte block
 16- 31 |    68-131    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access
 32- 47 |   132-195    | 1x128B       | in aligned 128byte block
 48- 63 |   196-256    | 1x64B, 1x32B | crosses 128B boundary, so no 128B access

Básicamente, estamos desperdiciando aproximadamente la mitad de nuestro ancho de banda (el ancho de acceso inferior al doble para los halfwarps impares no ayuda mucho porque genera más accesos, que no es más rápido que perder más bytes, por así decirlo).

Para las lecturas de b, los hilos solo acceden a elementos pares de la matriz, por lo que para cada semicanal, todos los accesos se encuentran en un bloque alineado de 128 bytes (el primer elemento está en el límite 128B, ya que para ese elemento el gid es múltiplo de 16 => el índice es un múltiplo de 32, para elementos de 4 bytes, lo que significa que el desplazamiento de la dirección es un múltiplo de 128B). El patrón de acceso se extiende a lo largo de todo el bloque 128B, por lo que hará una transferencia de 128B por cada semirremolque, y reducirá la mitad del ancho de banda.

Las lecturas de c generan uno de los peores escenarios, donde cada subproceso se indexa en su propio bloque 128B, por lo que cada subproceso necesita su propia transferencia, que por un lado es un poco un escenario de serialización (aunque no tan malo como normalmente, ya que el hardware debería poder superponerse a las transferencias). Lo que es peor es el hecho de que esto transferirá un bloque 32B para cada hilo, desperdiciando 7/8 del ancho de banda (accedemos a 4B / thread, 32B / 4B = 8, por lo que solo se utiliza 1/8 del ancho de banda). Dado que este es el patrón de acceso de matrices maternas ingenuas, es muy recomendable hacer ésas que usan memoria local (hablando desde la experiencia).

Compute 1.0 (G80)

Aquí el único patrón que creará un buen acceso es el original, todos los patrones del ejemplo crearán un acceso completamente sin cola, desperdiciando 7/8 del ancho de banda (transferencia / hilo 32B, ver arriba). Para el hardware G80, cada acceso donde el enésimo subproceso de una semirredida no acceda al enésimo elemento crea tales accesos no reforzados

Compute 2.0 (Fermi)

Aquí, cada acceso a la memoria crea transacciones de 128B (tantas como sea necesario para recopilar todos los datos, por lo que 16x128B en el peor de los casos), sin embargo, estos se almacenan en caché, lo que hace menos obvio dónde se transferirán los datos. Por el momento, supongamos que el caché es lo suficientemente grande como para contener todos los datos y no hay conflictos, por lo que cada caché de 128B se transferirá a lo sumo una vez. Además, asumiremos una ejecución serializada de los halfwarps, por lo que tenemos una ocupación de caché determinista.

Los accesos a b seguirán transfiriendo bloques de 128B (no hay otros índices de subprocesos en la memoria correspondiente). El acceso a c generará transferencias de 128B por subproceso (el peor patrón de acceso posible).

Para accesos a a es el siguiente (tratándolos como lecturas por el momento):

   gid  | offset  | accesses | reasoning
  0- 15 |   4- 67 |  1x128B  | bringing 128B block to cache
 16- 31 |  68-131 |  1x128B  | offsets 68-127 already in cache, bring 128B for 128-131 to cache
 32- 47 | 132-195 |    -     | block already in cache from  last halfwarp
 48- 63 | 196-259 |  1x128B  | offsets 196-255 already in cache, bringing in 256-383

Entonces, para arreglos grandes, los accesos a a no perderán casi ningún ancho de banda teóricamente. Para este ejemplo, la realidad, por supuesto, no es tan buena, ya que los accesos a c arruinarán bastante bien el caché.

Para el generador de perfiles, supongo que las eficiencias superiores a 1.0 son simplemente resultados de inexactitudes flotantes.

Espero que ayude