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