Racionalizar lo que está sucediendo en mi simple núcleo opencl en lo que respecta a la memoria global

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 núcleo de arriba es una adición de vector realizada diez veces por bucle. He usado la guía de programación y el desbordamiento de pila para descubrir cómo funciona la memoria global, pero aún no puedo entender mirando mi código si estoy accediendo a la memoria global de la misma manera. Lo estoy accediendo de manera contigua y supongo de manera alineada. ¿La tarjeta se carga de 128 kb de memoria global para las matrices A, B y C? ¿Luego carga los fragmentos de 128 kb para cada matriz una vez por cada 32 índices GID procesados? (4*32 = 128) Parece que entonces no estoy desperdiciando ningún ancho de banda de memoria global, ¿verdad?

Por cierto, el Compute Profiler muestra una eficiencia GLD y GST de 1.00003, lo que parece extraño, pensé que sería 1.0 si todas mis tiendas y cargas estuvieran unidas. ¿Cómo está por encima de 1.0?

Answer 1

Sí, su patrón de acceso a la memoria es bastante óptimo. Cada mediowarp accede a 16 palabras consecutivas de 32 bits. Además, el acceso está alineado con 64 bytes, ya que los buffers en sí están alineados y el inicio de cada mediowarp es un múltiplo de 16. Por lo tanto, cada mediowarp generará una transacción de 64 bytes. Por lo tanto, no debe desperdiciar el ancho de banda de memoria a través de accesos sin coal.

Dado que solicitó ejemplos en su última pregunta, modifiquemos este código para otro (un patrón de acceso menos óptimo (ya que el bucle realmente 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, vamos cómo funciona esto en el cálculo 1.3 (GT200) Hardware

Para las escrituras a A, esto generará un patrón ligeramente unópico (siguiendo los medios mide identificados por su rango de identificación 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 (cuanto menos duplicado el ancho de acceso para los extraños mide no ayuda mucho porque genera más accesos, lo que no es más rápido que desperdiciar más bytes, por así decirlo).

Para las lecturas de B, los hilos acceden solo a elementos incluso de la matriz, por lo que para cada mediowarp todos los accesos se encuentran en un bloque alineado de 128 bytes (el primer elemento está en el límite de 128b, ya que para ese elemento el GID es un 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 Access se extiende sobre todo el bloque 128B, por lo que esto hará una transferencia de 128b por cada mediowarp, nuevamente cinturando la mitad del ancho de banda.

Las lecturas de C generan uno de los peores escenarios de los casos, donde cada hilo índice en su propio bloque 128B, por lo que cada hilo necesita su propia transferencia, que una mano es un escenario de serialización (aunque no tan malo como normalmente, ya que el hardware debería poder superponer las transferencias). Lo que peor es el hecho de que esto transferirá un bloque 32B para cada hilo, desperdiciando 7/8 del ancho de banda (accedemos a 4B/hilo, 32B/4B = 8, por lo que solo se utiliza 1/8 del ancho de banda). Dado que este es el punto de acceso de acceso de ingenuos MatrixTranspose, es muy aconsejable hacer aquellos que usan la memoria local (hablando de la experiencia).

Calcule 1.0 (G80)

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

Compute 2.0 (Fermi)

Aquí, cada acceso a la memoria crea transacciones 128B (tantas como necesarias para recopilar todos los datos, por lo que 16x128b en el peor de los casos), sin embargo, se almacenan en caché, lo que hace que sea 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 cacheline 128B se transferirá como máximo una vez. Profundamos asumir una ejecución serializada de los medianos, por lo que tenemos una ocupación de caché determinista.

Los accesos a B siempre transferirán los bloques 128B (no hay otros índices de hilo en la memoria de correspondencia). El acceso a C generará transferencias 128B por hilo (peor acceso de acceso posible).

Para los accesos a A es el siguiente (tratarlos 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 matrices grandes, los accesos a un desgastará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 destrozarán el caché bastante bien

Para el Profiler, supongo que las eficiencias de más de 1.0 son simplemente resultados de las inexactitudes de puntos flotantes.

Espero que ayude