CUDA 5

Presenter Notes

Plan

  • Como planifica bloques y grillas.
    • Factores limitantes de la concurrencia.
    • Inner&outer scheduler.
  • Accesos a memoria:
    • Global.
    • Compartida.

Presenter Notes

¿Cómo se planifica esto ...

 1 #define N (1<<28)
 2 #define BLOCK_SIZE 128
 3 
 4 __global__ void ma4(void) {
 5     unsigned int gtid = blockIdx.x*blockDim.x + threadIdx.x;
 6     d[gtid] = a[gtid]*b[gtid]+c[gtid];
 7 }
 8 
 9 int main(void)
10 {
11     ma4<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
12 }

268.435.456 hilos divididos en

2.097.152 bloques de  
128 hilos cada uno.

Presenter Notes

... acá?

Full GK110

Presenter Notes

... acá? (zoom)

GK110 SMX

Presenter Notes

1 bloque a 1 SMX

Un procesador virtual en un procesador físico.

Cada bloque es independiente a todos los otros.

  • No hay comunicación entre ellos (bah, si, en la global por atomics).
  • Se pueden ejecutar en cualquier orden (concurrentemente también).
  • Escalabilidad trivial.

Automatic scalability

Presenter Notes

Planificación de dos niveles

Planificador global: busca SMX libres y les da bloques (batch)

¿Relación 1 a 1 ó n a 1?

Necesito sobrevender los SMX para que siempre estén ocupados.

Planificador local: cicla entre los bloques asignado (preemptive dynamic scheduing via scoreboarding).

Nicolás Wolovick:

Por fuera tengo un SLURM, por dentro el planificador de un sistema operativo.

Rob Farber:

In toto, the abstraction of a thread block and replication of SM hardware work in concert to transparently provide unlimited and efficient scalability. The challenge for the CUDA programmer is to express their application kernels in such a way to exploit this parallelism and scalability.

Presenter Notes

Límites duros del planificador local

El planificador global le da todo el trabajo al planificador local mientras "entre".

Límites

Technical specifications

Ejemplo, ma4() en full Kepler (CC 3.5)

Configuración de ejecución: <<<2097152,128>>>.
Registros: 16
ShMem: 0 KiB

Ni los registros, ni la shmem, ni los hilos por bloque son el limitante.
El límite son los 16 bloques por SMX y los 64 warps por SMX.

Presenter Notes

CUDA occupancy Calulator

CUDA_Occupancy_calculator.xls

Mirar datos importantes en physical limits, por ejemplo granularidad.

Ejemplo sgemm_shared_Volkov() en full Kepler (CC 3.5)

Para máxima performance: N=1024, B=32, U=8.

Configuración de ejecución: <<<(32,32),(32,4)>>> = <<<1024,128>>>.
Registros (full unroll): 49
ShMem: 8 KiB

Entran 6 bloques por SMX.
La limitante es la ShMem.
Ocupación (en warps) = 24/64 = 38%.

¿Dónde está el truco?

Presenter Notes

¿Cómo se llenan estas unidades?

Architecture specifications

Presenter Notes

ILP

GK110 warp scheduler

Manda hasta dos instrucciones por ciclo.

Los stalls no se producen en la instrucción que ejecuta la operación con latencia (arit o mem), sino en la instrucción que depende de esta.

Reordering a nivel SASS, ILP masivo.

Presenter Notes

ILP, otra forma de ocultar latencia

Dos extremos

  • (tradicional ) Pocos bloques, bloques grandes, cada hilo poco trabajo.
  • (Volkov style) Muchos bloques, bloques pequeños, cada hilo mucho trabajo.

Que se busca maximizar

  • Tradicional: ocultamiento de la latencia a través de TLP de warps del mismo bloque.
  • Volkov y MAGMA style: ocultamiento de latencia a través de ILP de warps de varios bloques.
    • Mucho trabajo por hilo => usar muchos regs, brindar mucho ILP, y en algunos casos (sgemm_shared_Volkov) reutilizar más la información de la ShMem.

(Creo) que la mayor ganancia está en mezclar bloques dentro de un SMX: ¡expone aun más paralelismo!

Presenter Notes

Little's Law

Paralelismo necesario para ocupar todo essssto.

Paralelismo = latencia * throughput

Volkov Little's Law

Presenter Notes

Ejemplos en Fermi

Paralelismo aritmético

Paralelismo de memoria

(Vasily Volkov, Better performance at lower occupancy, GTC, 2010.)

Presenter Notes

Como funciona el acceso a memoria

Lo ideal

Un warp lee 128 bytes consecutivos y alineados.
32 hilos, 4 bytes cada uno.

1 uint gtid = blockIdx.x*blockDim.x + threadIdx.x;
2 a[gtid] = 1.0f

Acceso alineado perfecto

Permutaciones dentro de una línea

Hacer un gather interno no molesta. (antes ¿G80, GT200?, si!!!)

Acceso alineado con permutación

Presenter Notes

Accesos desalineados

1 uint gtid = blockIdx.x*blockDim.x + threadIdx.x;
2 a[gtid+1] = 1.0f

Acceso desalineado

Trae dos líneas de 128 bytes.

La caché

En GPU la caché mitiga los accesos desalineados.
Solo se benefician los programas por la localidad espacial.

  • Suaviza las rugosidades en la performance de acceso a memoria.
  • Mitiga los register spillings, stack frames, function call.

Está aumentando y aumentando y aumentando.

Presenter Notes

¿Deshabilitar la cache?

Por ahi no tiene sentido traer cosas a caché.

Mejor deshabilitar la L1 y dejar la L2 que tiene granularidad de 32 bytes.

Fermi: siempre usa L1, a menos que pidamos -Xptxas -dlcm=cg.
Kepler: no usa L1 para accesos a memoria, pero si para stack y reg. spill, a menos que -Xptxas -dlcm=ca.

"He visto mejoras de performance deshabilitando la caché L1 en Fermi.", NW, circa 2010.

Presenter Notes

Memoria Compartida

Organización

  • Dividida en 32 bancos entremezclados de palabras de 32 bits.
  • Cada thread en un warp puede leer en paralelo de un banco distinto.
  • Si más de un hilo lee la misma palabra de 32 bits de un banco, el resultado se difunde.

Acceso ideal

1 __shared__ float shared[SHARED];
2 float data = shared[threadIdx.x];

Conflicto de bancos

  • Dos hilos dentro de un warp acceden a distintas palabras de 32 bits en el mismo banco.
  • Los hilos en conflicto se serializan.

Presenter Notes

¿Conflictos? No, Si, No

Presenter Notes

¿Conflictos? No, No (bcast), No (bcast)

¿sgemm-shared-Volkov.cu tendrá en cuenta estas cosas? ¡Ejercicio!

Presenter Notes

Bibliografía

Presenter Notes