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.
Un procesador virtual en un procesador físico.
Cada bloque es independiente a todos los otros.
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.
El planificador global le da todo el trabajo al planificador local mientras "entre".
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.
Mirar datos importantes en physical limits, por ejemplo granularidad.
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?
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.
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!
Paralelismo necesario para ocupar todo essssto.
Paralelismo = latencia * throughput
(Vasily Volkov, Better performance at lower occupancy, GTC, 2010.)
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
Hacer un gather interno no molesta. (antes ¿G80, GT200?, si!!!)
1 uint gtid = blockIdx.x*blockDim.x + threadIdx.x;
2 a[gtid+1] = 1.0f
Trae dos líneas de 128 bytes.
En GPU la caché mitiga los accesos desalineados.
Solo se benefician los programas por la localidad espacial.
Está aumentando y aumentando y aumentando.
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.
1 __shared__ float shared[SHARED];
2 float data = shared[threadIdx.x];
¿sgemm-shared-Volkov.cu
tendrá en cuenta estas cosas? ¡Ejercicio!
Table of Contents | t |
---|---|
Exposé | ESC |
Full screen slides | e |
Presenter View | p |
Source Files | s |
Slide Numbers | n |
Toggle screen blanking | b |
Show/hide slide context | c |
Notes | 2 |
Help | h |