1. Muchos núcleos simples que corren en paralelo.
2. Empaquetar muchos cores y que corran el mismo flujo de instrucciones.
   • SIMD implícito.
   • Instrucciones escalares, el hardware "vectoriza".
3. Entremezclar hilos en un solo core para evitar stalls: ocultamiento de latencia.
   • De memoria.
   • De operaciones largas.

(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 las ALU:
SIMD adentro de la GPU.

• SSE: hw SIMD, sw SIMD.
  • Recordar instrucciones como blend de SSE4.1.
• GPU: hw SIMD, sw MIMD.
  • Grupos de ALUs comparten instruction stream.
  • NVIDIA: warp, grupos de 32 hilos.
  • ATI: wavefront, grupos de 64 hilos.

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

Ilusión de que cada hilo tiene su propio 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)

(Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.)

(Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.)

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:

• tamaño del contexto
• capacidad de ocultar la latencia.

Ejemplo

Cada SM (streaming multiprocessor) de la Tesla C2070/75 (Fermi) tiene 32768 registros de 32 bits, o sea 128 KiB.

(Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.)

(Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.)

• Las aplicaciones memory-bound, no tienen chance.
  • Hay que conocer la arithmetic intensity del problema (CAAQA5).

• El tamaño del contexto por core en cada SM.
  • G80: 8192×4/16 = 2 KB , GT200: 16384×4/24 = 2.66 KB , GF100:32768×4/32 = 4 KB, GK100: 65536×4/192 = 1.33 KB.
• El número máximo de hilos por SM:
  • G80: 768, GT200: 1024, GF100: 1536, GK100: 2048.
  • Notar que una C2075 puedo tener 15×1536 = 23040 hilos corriendo!

Tesla C2075:

• 15 SM.
• 32 lanes per SM.
• 1 FMA precisión simple por lane (2 floating point ops).
• 1.15 GHz clock frequency.

15×32×2×1.15 = 1.104 TFLOPS precisión simple

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.

Notar la pirámide inversa en la jerarquía (GF100)

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.

Ejemplo, C2075

• 1 TFLOPS
• 150 GBps

Necesito 1024GFLOP / (150GBps/4) = 27 FLOP/word (operaciones de punto flotante por lectura de fp32).

¿Cómo lograr esto? Ya veremos más adelante:

• Vasily Volkov, Better Performance at Lower Occupancy, GTC, 2010.

• Registros: 8.1 TB/s, 1 ciclo.
• Shared: 1.6 TB/s, ~6 ciclos.
• Principal: 150 GB/s, ~200,400 ciclos.

Para comparar

• Una instrucción FMA toma entre 18 y 22 ciclos.
• Ancho de banda de 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)

Ejemplo de la arquitectura bit.LITTLE de ARM

Las GPUs son muchísimos núcleos pequeños

• Gran throughput.
• pero con gran latencia.

• Muchos núcleos.
• Caché L2.
• Memoria paralela.

(Ali Bakhoda et.al, Analyzing CUDA Workloads Using a Detailed GPU Simulator, ISPASS, 2009.)

• Comunicación de programas y datos a través de PCIe. Es un acelerador u offloader.
• Muchos núcleos, del orden de decenas.
• Todos los núcleos ejecutan en mismo programa.
• Entre 2 y 8 canales de comunicación de memoria.
• Memoria global entre 0.5 GiB y 12 GiB.
• La red de interconexión de las memorias intenta maximizar el paralelismo. No hay necesariamente una localidad.
• Los memory controllers juntan pedidos de memoria y cuando están listos avisan a los cores (async access).

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

Memoria constante / de parámetros ~8 KiB

Caché L1 read-only, broadcast a todas las ALUs si piden el mismo dato.

Memoria de texturas ~16 KiB

Caché L1 read-only, con un interleaving para que las lecturas espaciales sean óptimas con el 4D blocking scheme.

Memoria local / Cache L1 ~16 KiB

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 compartida ~32 KiB

Memoria de lecto-escritura tipo scratchpad highly-banked (32-way) para lecturas paralelas y broadcast para lecturas al mismo banco.

• El mismo código para todas las ALUs del shader core.
• Átomo: paquetes de 32 lanes: warp.
• Agrupan átomos en CTA (cooperative threads array), aka block.
• 1 CTA (varios warps) -> 1 core.
• 1 core -> varios CTA.

Planificador local (por core)

1. Elegir circularmente (round robin) un warp listo para ejecutar de algún CTA.
2. Se lo corre ininterrumpidamente hasta terminar o ...
3. Si hay una dependencia no resuelta, switch de banco de registros (zero overhead).
4. GOTO 1

Barrera

Un CTA llamando a bar.sync hace que todos los warps del CTA sincronicen en esa instrucción.

Asigna CTAs a shader cores mientras HAYA LUGAR.

Hay lugar si no se supera:

• Cantidad de CTAs por shader.
• Cantidad de warps por shader.
• Cantidad de registros por shader.
• Cantidad de memoria compartida por shader.

¿Qué pasa con la const y la tex?

Multiprogramación

Algunas arquitecturas aceptan CTAs de kernels distintos que ejecuten en el mismo shader.

Ejecutar dos o más programas distintos en los mismos shader cores.

ALU

Operaciones de enteros y fp32.
En igual cantidad.
Throughput 1.

LD/ST

Unidades de pedido de memoria. En cantidad menor a ALU.

ALU fp64

Menor cantidad que fp32.
Acá está el ratio fp64/fp32 = {1/2, 1/8, 1/32} dependiendo de la arquitectura.

SFU

Operaciones trascendentales: sin, cos, log2, exp2, rcp, sqrt, rsqrt. En mucho menor cantidad que los dos anteriores.

TEX

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)

• 550 mm² die size.
• 40 nm litografía.
• 3000 M transitores.
• 512 cores.
  • 32 lanes por SM (streaming multiprocessor). 16 SM.
• Dual-issue warp scheduler (ILP).
• Full IEEE-754-2008 fp32 y fp64.
• Direcciones 64 bit.
• Espacio unificado de direcciones CPU-GPU.
• Memoria ECC.
• Caché L2 y L1 (configurable con la shared).
• Operaciones atómicas rápidas y en fp32.
• Hasta 6GB DDR5, 6x1GB, 6 puertos, ancho 384 bits (6x64 bits).

• 398 mm2 die size.
• 28 nm litografía.
• 5200 M transitores.
• 165 W de TDP.
• 4 SIMT Core Clusters, 4 controladores de memoria.
• 2048 "cores".
  • 32 lanes por warp.
  • 4 bloques de procesamiento.
  • 16 SMM.
• Dual-issue warp scheduler.
• fp64/fp32 = 1/32
• Sin memoria ECC.
• Caché L2 2 MiB
• 4 canales de memoria.

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.

• Kayvon Fatahalian, CMU 15-869, Graphics and Imaging Architectures, 1-Course Introduction, 2011.
• Kayvon Fatahalian, From Shader Code to a Teraflop: How GPU Shader Cores Work, 2011.
• Hennessy, Patterson, Computer Architecture: A Quantitative Approach, Fifth Edition, Morgan Kaufmann, 2011.
• J. Owens, et.al, A Survey of General-Purpose Computation on Graphics Hardware, Computer Graphics Forum, 26(1), 2007.
• Ali Bakhoda et.al, Analyzing CUDA Workloads Using a Detailed GPU Simulator, ISPASS, 2009.

• ISA.
• CUDA1.