CUDA 5

Presenter Notes

Plan

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

20160507

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.

En realidad:

2.097.152 bloques de  
4 warps de  
32 hilos cada uno.

Presenter Notes

... acá?

GM204

Presenter Notes

... acá? (zoom parte 1)

GM204 SMM

Presenter Notes

... acá? (zoom parte 2)

GM204 SMM

Presenter Notes

1 bloque a 1 SMM

1 procesador virtual en 1 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 SMM libres y les da bloques (batch)

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

Necesito sobrevender los SMM 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 Maxwell (CC 5.2)

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 32 bloques por SMM y los 64 warps por SMM.

Presenter Notes

CUDA occupancy Calulator

CUDA_Occupancy_calculator.xls

Mirar datos importantes en physical limits, por ejemplo granularidad.

Ejemplo sgemm-shared-Volkov() en Maxwell (CC 5.2)

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

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

Entran 8 bloques por SMM.
La limitante es la Cantidad de Registros.
Ocupación (en warps) = 32/64 = 50%.

¿Dónde está el truco?

Presenter Notes

¿Cómo se llenan estas unidades?

Architecture specifications

Presenter Notes

ILP

GM204

Presenter Notes

ILP

The NVIDIA GeForce GTX 980 Review: Maxwell Mark 2, AnandTech, 2014.

Manda hasta dos instrucciones por ciclo.

Los stalls no se producen en la instrucción que ejecuta la operación con latencia (aritmética o memoria), 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 SMM: ¡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

Accesos a Memoria

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 calls.

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 distintos 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