Entremezclar hilos en un solo core para evitar stalls: ocultamiento de latencia.
Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.
CPU baja latencia de un hilo a costa del throughput.
El trabajo puede ser secuencial o paralelo.
GPU tiene alta latencia, pero alto thoroughput.
Asume trabajo en paralelo.
Se usa el área de la caché, branch prediction, scheduler (OoOE), prefetcher, etc.
Amortizar aun más el área de Fetch/Decode:
SIMD adentro de la GPU.
blend
de SSE4.1.¿Cómo lo hace?:
internal masks, branch syncronization stack, instruction markers.
W.Fung, I.Sham, G.Yuan, T.Aamodt, Dynamic Warp Formation and Scheduling for Efficient GPU Control Flow, MICRO40, 2007.
PC
Writing programs that operate on SIMD Lanes in this highly independent MIMD mode is like writing programs that use lots of virtual address space on a computer with a smaller physical memory. Both are correct, but they may run so slowly that the programmer could be displeaced with the result.
Hennessy, Patterson, CAAQA5, p.303.
Cambio de contexto es prácticamente gratis.
Pero tengo que poder almacenar el contexto (PC, regs, flags) en algún lugar muy cercano.
Tensión entre:
Cada SM (streaming multiprocessor) de la Tesla C2070/75 (Fermi) tiene 32768 registros de 32 bits, o sea 128 KiB.
Tesla C2075:
El nudo está en la caché.
Elimina la caché, pero da 6x de ancho de banda de memoria.
Agrega caché manual (shared memory), caché L1 y caché L2.
¡Cuiadado! La caché tiene propósitos distintos que en la CPU.
Registros: 128 KB × 15 ~= 2 MB.
Caché L1: [16 KB × 15, 48 KB × 15 ] ~= [0.25 MB, 0.75 MB].
Caché L2: 0.75 MB.
Necesito 1024GFLOP / (150GBps/4) = 27 FLOP/word (operaciones de punto flotante por lectura de fp32
).
¿Cómo lograr esto? Ya veremos más adelante:
FMA
toma entre 18 y 22 ciclos.PCIe 3.0 x16
15.75 GB/s.Microprocessor (single core) performance scales roughly as the square root of its complexity, where transistor count is often used as a proxy to quantify complexity.
(M. Själander, M. Martonosi, S. Kaxiras, Power-Efficient Computer Architectures, Synthesis Lectures on Computer Architecture, 2014)
Ali Bakhoda et.al, Analyzing CUDA Workloads Using a Detailed GPU Simulator, ISPASS, 2009.
Ali Bakhoda et.al, Analyzing CUDA Workloads Using a Detailed GPU Simulator, ISPASS, 2009.
Además de la memoria global que pueden ver todos los cores, hay memoria local.
Caché L1 read-only, broadcast a todas las ALUs si piden el mismo dato.
Caché L1 read-only, con un interleaving para que las lecturas espaciales sean óptimas con el 4D blocking scheme.
Caché L1 non-coherent de memoria global. Write policy depende de si es memoria local o global.
... evict-on-write for global memory, writeback for local memory ...
Memoria de lecto-escritura tipo scratchpad highly-banked (32-way) para lecturas paralelas y broadcast para lecturas al mismo banco.
Hardware!
Un simple spinlock entra en deadlock.
Volta los soluciona con libre avance de los hilos dentro del warp
A. ElTantawy, T.M. Aamodt, MIMD synchronization on SIMT architectures, MICRO16, 2016.
Un CTA llamando a bar.sync
hace que todos los warps del CTA sincronicen en esa instrucción.
Asigna bloques a shader cores mientras HAYA LUGAR.
¿Qué pasa con la const
y la tex
?
Algunas arquitecturas aceptan bloques de kernels distintos que ejecuten en el mismo shader.
Ejecutar dos o más programas distintos en los mismos shader cores.
Operaciones de enteros y fp32
.
En igual cantidad.
Throughput 1.
Unidades de pedido de memoria. En cantidad menor a ALU.
fp64
Menor cantidad que fp32
.
Acá está el ratio fp64/fp32 = {1/2, 1/8, 1/32} dependiendo de la arquitectura.
Operaciones trascendentales: sin
, cos
, log2
, exp2
, rcp
, sqrt
, rsqrt
.
En mucho menor cantidad que los dos anteriores.
Unidades de interpolación de texturas. Cantidad similar a SFU.
Si tengo una latencia de 20 ciclos para una instrucción típica y una ALU de 32 de ancho.
¿Cuántos warps necesito para tener un throughput de 32 IPC?
Paralelismo = latencia × throughput
640 = 20 × 32
¡Un montón de paralelismo!
(no necesariamente TLP, también puede ser ILP)
Todos los pedidos de memoria de un warp se juntan para armar un pedido grande.
Permutation crossbar de 32x4=128 bytes.
(GPGPU-Sim v3, 2016)
(GPGPU-Sim v3, 2016)
fp32
y fp64
.fp32
.The Maxwell GPU inside the GeForce GTX 980 is built on the same 28nm manufacturing process used for previous GeForce GPUs, this means that the GPU’s remarkable performance and efficiency doesn’t come from the smaller transistors you get with a die shrink, rather these benefits come from the new GPU architecture.
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 |