Nicolás Wolovick, $Date: 2012-05-21 09:41:18 -0300 (Mon, 21 May 2012) $, $Revision: 3530 $
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)
(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 (OOE), prefetcher, etc.
Amortizar aun más las ALU:
SIMD adentro de la GPU.
blend
de SSE4.1.¿Cómo lo hace?: internal masks, branch sincronization stack, instruction markers.
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:
Cada SM (streaming multiprocessor) de la Tesla C2070/75 (Fermi) tiene 32768 registros de 32 bits, o sea 128 KB.
(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.)
Tesla C2075:
fp32
per core (2 floating point ops).f32
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 f32
).
¿Cómo lograr esto? Ya veremos más adelante:
Para comparar: una FMA
toma entre 18 y 22 ciclos.
Single Instruction, Multiple Thread
fp32
vs. fp64
units.Permite ejecutar un subconjunto interesante del lenguaje desde una G80 a una GK100.
Puede sufrir problemas de performance.
CUDA virtualiza el hardware:
CUDA runtime planifica en el hardware:
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 (GeForce 510, 1 SM) tienen menos SMs. El programa corre igual que en una full (GTX 580, 16 SM), pero con menor performance.
Este esquema facilita el paralelismo de datos para grillas 2D y 3D.
No siempre este esquema se adapta a lo que necesitamos.
Aumenta C
con unas poquitas cosas.
hilo ∈ bloques ∈ grilla
Calificadores de funciones
1 __host__ int add_cpu() { } // CPU function
2 __global__ void add_gpu() { } // kernel callable 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
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.
No todas las placas implementan esta montaña de cosas.
float
atomics, predicated synchthreads
, surfaces, printf
, assert
.Notar el salto fenomenal de cores por SM(X) en Kepler.
1 #define N (1<<28)
2 float a[N], b[N], c[N], d[N];
3
4 int main(void) {
5 unsigned int i = 0;
6 for(i=0; i<N; ++i)
7 d[i] = a[i]*b[i]+c[i];
8
9 return 0;
10 }
Con millones de hilos no tiene sentido paralelización de tareas.
1 #include <cuda.h>
2
3 #define N (1<<22)
4 #define BLOCK_SIZE 128
5
6 __device__ float a[N], b[N], c[N], d[N];
7
8 __global__ void fma(void) {
9 unsigned int tid = blockIdx.x*blockDim.x + threadIdx.x;
10 d[tid] = a[tid]*b[tid]+c[tid];
11 }
12
13 int main(void)
14 {
15 fma<<<N/BLOCK_SIZE, BLOCK_SIZE>>>();
16 cudaDeviceSynchronize();
17 return 0;
18 }
1 $ nvcc fma.cu -O3 -arch=sm_20 --ptxas-options=-v -o fmacuda
2 ptxas info : Compiling entry function '_Z3fmav' for 'sm_20'
3 ptxas info : Function properties for _Z3fmav
4 0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
5 ptxas info : Used 13 registers, 32 bytes cmem[0], 32 bytes cmem[14]
6 $ ./fmacuda
7 $
Si pongo #define N (1<<23)
tengo más de 65536 bloques en la grilla y excedo el límite de GF100. Falla silenciosamente.
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 |