ISA + CUDA

Presenter Notes

Resumen

  • ISA
  • CUDA.
  • Ejemplos básicos.

Nicolás Wolovick, 20160518.

Presenter Notes

ISA

Presenter Notes

ISA, generalidades

  • Four operand instructions (como AVX2).
  • Predicated instructions (como ARM).
  • Punto flotante de 16, 32 y 64 bits.
  • Funciones trascendentes.
  • (por supuesto operaciones enteras y booleanas).
  • Load/Store para cada una de las 5 memorias: const, texture, shared, local, global.
  • Operaciones atómicas.
  • Sincronización de barrera (con scope acotado).

Presenter Notes

Formato de instrucciones PTX

opcode.type d, a, b, c;

PTX types

(Hennessy, Patterson, CAAQA5, p.298)

Presenter Notes

PTX ISA

CAAQA5 Fig-4.17

(Hennessy, Patterson, CAAQA5, Fig-4.17)

Presenter Notes

PTX ISA

CAAQA5 Fig-4.17

(Hennessy, Patterson, CAAQA5, Fig-4.17)

Presenter Notes

PTX ISA

CAAQA5 Fig-4.17

(Hennessy, Patterson, CAAQA5, Fig-4.17)

Presenter Notes

Capacidades de la ISA

Presenter Notes

Capacidades de la ISA

Presenter Notes

Límites de la ISA

Presenter Notes

Límites de la ISA

Presenter Notes

Límites de la ISA

Presenter Notes

Presenter Notes

PTX es intermedio!

A demostration of the value of the PTX is that the Fermi architecture radically changed the hardware instruction set -- from being memory-oriented like x86 to being register-oriented like MIPS as well as doubling the address size to 64 bits -- without disrupting the NVIDIA software stack.

(Hennessy, Patterson, CAAQA5, p.330)

Presenter Notes

CUDA

Presenter Notes

Kudos for CUDA

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

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 MIMD y SIMD

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

Abstracción del hardware, es independiente del:

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

Permite ejecutar un subconjunto interesante del lenguaje desde una G80 a una GM200.
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

Modelo de Memoria de CUDA

Jerarquía de hilos, bloques, grillas

CUDA hierarchy of threads, blocks, grids

(NVIDIA’s Next Generation CUDA Compute Architecture: Fermi™, NVIDIA, 2009.)

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 ∈ warp ∈ bloque ∈ grilla

Cantidades típicas

32 hilos en 1 warp
32 warps en 1 bloque
1048576 bloques en una grilla.

CUDA hierarchy of threads, blocks, grids

Presenter Notes

Unidades de Cooperación

Warp

  • Ejecución interlocked y ...
  • Comunicación de variables privadas.

(via ballots & warp shuffling)

Bloque

  • Sincronización de barrera.
  • Memoria compartida local shared.
  • Instrucciones atómicas sobre la shared: atomicAdd, CAS.

Grilla

  • Sincronización fork-join por lanzamiento de kernels.
  • Memoria compartida global.
  • Instrucciones atómicas sobre la global: atomicAdd, CAS.

Presenter Notes

Scoping y lifetime de variables

Yo declaro

1 int i;

Esta i puede ser:

  • Local a un hilo.
  • Compartida en un warp.
  • Compartida en un bloque.
  • Compartiga en LA grilla.

¡Compara con OpenMP que solo tenía local y global!

Presenter Notes

Locality = Performance

All performance is from parallelism

Machines are power limited

(efficiency IS performance)

Machines are communication limited

(locality IS performance)

-

(Bill Dally, Efficiency and Programmability: Enablers for ExaScale, SC13)

-

locality=performance

locality=performance

locality=performance

locality=performance

locality=performance

locality=performance

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__ float add_gpu(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)

  • 1.0 (G80): solo float, 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, small Maxwell): (software) unified memory, crypto instructions.
  • ...
  • 7.5 (GP104, small Pascal): ...

Presenter Notes

Ejemplos: map, reduce

Presenter Notes

Map

Presenter Notes

Reduce

Presenter Notes

Bibliografía

Presenter Notes

La clase que viene

  • Ejemplos sencillos de CUDA.
    • Debugging.
    • Herramientas para detectar errores.
    • Profiling.
    • Buenas prácticas de programación.

Presenter Notes