CUDA

Presenter Notes

Resumen

  • Tres conceptos claves de la arquitectura GPU.
  • Jerarquía de Memoria.
  • CUDA.
  • Ejemplo.

Nicolás Wolovick, 20140520.

Presenter Notes

It's been almost three years since GPU computing broke into the mainstream of HPC with the introduction of NVIDIA’s CUDA API in September 2007. Adoption of the technology since then has proceeded at a surprisingly strong and steady pace. Many organizations that began with small pilot projects a year or two ago have moved on to enterprise deployment, and GPU accelerated machines are now rep- resented on the TOP500 list starting at position two. The relatively rapid adoption of CUDA by a community not known for the rapid adoption of much of anything is a noteworthy signal. Contrary to the accepted wisdom that GPU computing is more difficult, I believe its success thus far signals that it is no more complicated than good CPU programming. Further, it more clearly and succinctly expresses the parallelism of a large class of problems leading to code that is easier to maintain, more scalable and better positioned to map to future many-core architectures. Vincent Natol “Kudos for CUDA”, HPCwire (2010)

(Hennessy, Patterson, CAAQA5, L-48)

Presenter Notes

Tres conceptos claves de la arquitectura GPU

Presenter Notes

Tres conceptos claves

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

Presenter Notes

Muchos núcleos simples: shader core

1 shader

Presenter Notes

Comparación con CPU

Shader vs CPU

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.

Presenter Notes

Muchos shaders

16 shaders

Se usa el área de la caché, branch prediction, scheduler (OOE), prefetcher, etc.

Presenter Notes

Muchos shaders, comparten instrucciones

sharing instruction stream

Amortizar aun más las ALU:
SIMD adentro de la GPU.

Presenter Notes

Hard SIMD no implica soft SIMD

  • SSE: hw SIMD, sw SIMD.
    • Recordar instrucciones como blend de SSE4.1.
  • GPUs: 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 sincronization stack, instruction markers.

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)

Presenter Notes

Ejecución condicional

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

Presenter Notes

Muchos hilos para ocultar latencia

Latency hiding

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

Presenter Notes

Muchos hilos para ocultar latencia

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.

Presenter Notes

Máximo ocultamiento de latencia

Context size and latency

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

Presenter Notes

Mínimo ocultamiento de latencia

Context size and latency

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

Presenter Notes

Límites del latency-hiding

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

Arithmetic instensity chart

  • 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!

Presenter Notes

Si hacemos todo bien

Tesla C2075:

  • 15 SM.
  • 32 cores per SM.
  • 1 FMA fp32 per core (2 floating point ops).
  • 1.15 GHz clock frequency.

15×32×2×1.15 = 1.104 TFLOPS f32

Presenter Notes

Jerarquía de Memoria

Presenter Notes

Estilo CPU

CPU memory hierarchy

El nudo está en la caché.

Presenter Notes

Estilo GPU

GPU memory hierarchy

Elimina la caché, pero da 6x de ancho de banda de memoria.

Presenter Notes

Estilo GPU moderno

Modern GPU memory hierarchy

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.

Presenter Notes

Intensidad aritmética y performance pico

Ejemplo, C2075

  • 1 TFLOP
  • 150 GBps

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

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

Presenter Notes

Memoria: throughput y latencia

  • Registros: 8.1 GB/s, 1 ciclo.
  • Shared: 1.6 GB/s, ¿6? ciclos.
  • Principal: 150 MB/s, ¿200,400? ciclos.

Para comparar: una instrucción FMA toma entre 18 y 22 ciclos.

Presenter Notes

CUDA

Presenter Notes

Orígenes de CUDA

Brook for GPUs: Stream Computing on Graphics Hardware

Ian Buck, Tim Foley, Daniel Horn, Jeremy Sugerman, Kayvon Fatahalian, Mike Houston, and Pat Hanrahan Computer Science Department Stanford University To appear at SIGGRAPH 2004

Figure 1: Programming model for current programmable graphics hardware.

Ian Buck es "VP Accelerated Computing at NVIDIA".

Presenter Notes

Modelo SIMT

Single Instruction, Multiple Thread

A medio camino entre SMT y SIMD

  • Se programa como SIMD, pero permite divergencia en el flujo de control y de datos.
  • Se programa como SMT, pero todos los hilos ejecutan el mismo código.

Abstracción del hardware, es independiente del:

  • Tamaño del warp.
  • Cantidad de cores por SM.
  • Cantidad de SMs.
  • Ratio fp32 vs. fp64 units.
  • Cantidad de special units (funciones trascendentes).

Permite ejecutar un subconjunto interesante del lenguaje desde una G80 a una GK100.
Puede sufrir problemas de performance.

Presenter Notes

Modelo de paralelismo de CUDA

CUDA Parallelism

CUDA languaje virtualiza el hardware:

  • Thread: procesador escalar virtualizado (PC, registros, pila).
  • Block: multiprocesador virtualizado (hilos, memoria compartida).

CUDA runtime planifica en el hardware:

  • Non-preemptive. Los hilos dentro de un bloque se lanzan y ejecutan hasta que se terminan. (No hay critical sections, semaphores, a la manera de SistOp.).
  • Los bloques son independientes.

Presenter Notes

Escalabilidad automática

Esto permite ejecutar en hardware con diferente número de SMs, o bien en algún momento con ¡SM remotos!.

Notar que las placas baratas (GTX 635, 1 SMX) tienen menos SMs. El programa corre igual que en una mediana (GTX 680, 8 SMX), o que una full (GTX 780 Ti, 15 SMX) pero con menor performance.

Presenter Notes

Ejemplo: grilla 2D de bloques de hilos 2D

Este esquema facilita el paralelismo de datos para grillas 2D y 3D.
No siempre este esquema se adapta a lo que necesitamos.

Presenter Notes

CUDA

Aumenta "C" con unas poquitas cosas.

Abstracciones clave

  • Esquema fork-join como en OpenMP.
    • Pero el fork y el join tienen costo 0.
  • Jerarquía de dos niveles de hilos.
    • Bloques de hilos.
    • Grillas de bloques.
  • Sincronización de barrera muy barata dentro de un bloque.
    • Por fuera no hay (no escalaría!).
    • Usar fork-join para sincronizar.
  • Memoria compartida dentro de un bloque.

Presenter Notes

Esquema fork-join

Presenter Notes

Jerarquía de hilos

hilo ∈ bloque ∈ grilla

Bloque como unidad de cooperación

  • Sincronización.
  • Memoria compartida.

Límite de hilos/bloque: G80: 512, GT200: 512, GF100: 1024, GK100: 1024.

Scoping y lifetime de variables

Notar el scope. Este nuevo nivel aumenta los dolores de cabeza.

Presenter Notes

"C" para CUDA

Calificadores de funciones

1 __host__ int add_cpu() { } // CPU function
2 __global__ void add_gpu() { } // callable kernel from host
3 __device__ float add_partial_gpu() { } // GPU function

Calificadores de variables

1 __device__ float a[256]; // device global variable
2 __constant__ float my_constant_array[32]; // constant global array
3 __shared__ float my_shared_array[32]; // device block-shared array
4 __global__ add_gpu float(const float* __restrict__ a); // non-aliased pointer

Tipos

1 typedef struct {float x, float y} float2; // aligned
2 // also float3, float4, double2, uint2, uint3, etc.
3 typedef uint3 dim3; // initial sections, trailing values default to 1.

Configuración de la ejecución

1 dim3 grid_dim(100, 50); // 5000 thread blocks
2 dim3 block_dim(4, 8, 8); // 256 threads per block
3 my_kernel <<< grid_dim, block_dim >>> (...); // Launch kernel

Presenter Notes

"C" para CUDA (cont'd)

Variables y funciones built-in

1 dim3 gridDim; // Grid dimension
2 dim3 blockDim; // Block dimension
3 uint3 blockIdx; // Block index
4 uint3 threadIdx; // Thread index
5 int warpSize;
6 void __syncthreads(); // Intra-block thread synchronization

Muchas más cosas:

· Memory fence.
· Funciones de texturas (interpolación 1D, 2D y 3D) y superficies.
· Funciones de tiempo (clock).
· Funciones atómicas.
· Funciones de votación de warps.
· Funciones de shuffle de warps.
· Aserciones (assert).
· Salida estándar (printf).
· Pedido de memoria dinámica (malloc).
· Funciones matemáticas.

Presenter Notes

Compute Capability (CC)

No todas las placas implementan esta montaña de cosas.

  • 1.0 (G80): sin doble precisión, sin atomics.
  • 1.1 (G86): atomics.
  • 1.2 (GT216): más atomics, atomics en shared, warp vote.
  • 1.3 (GT200 Tesla): doble precisión.
  • 2.0 (GF100 Fermi): 3D grids, float atomics, predicated synchthreads, surfaces, printf, assert.
  • 2.1 (GF104, aka aborto'e mono): 3 half-warp core groups (48), ILP.
  • 3.0 (GK104 Kepler): 2^32-1 grid.x, shared vs. L1 más configurable, ILP, warp shuffle, 16 kernels concurrentes.
  • 3.5 (GK110 full Kepler): 255 registros por hilo, paralelismo dinámico, ldg, 64-bit atomics, funnel shifts, 32 kernels concurrentes.
  • 5.0 (GM104, half Maxwell): (software) unified memory, crypto instructions.

Presenter Notes

CC chart, Feature Support

CUDA C Programming Guide (6.0).
Página 180

Presenter Notes

CC chart, Technical Specs

CUDA C Programming Guide (6.0).
Página 181

Notar el salto fenomenal de cores por SM(X) en Kepler.

Presenter Notes

CC chart, artithmetic throughput