Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
CUDA
Carlos Bederián, Nicolás Wolovick
FaMAF, Universidad Nacional de Córdoba, Argentina
31 de Julio de 2012
[email protected]
Revisión 3739, 2012-08-02
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Introducción
Manejo de Memoria
Configuración de ejecución
Paralelismo de datos
Comunicación y sincronización
Desempeño y Debugging
Resumen
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Modelo de programación
CPU tradicional + Placa aceleradora.
Dos arquitecturas, dos espacios de memoria,
dos filosofías distintas de programación.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Modelo de programación
CPU tradicional + Placa aceleradora.
host
device
Dos arquitecturas, dos espacios de memoria,
dos filosofías distintas de programación.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Modelo de programación
CPU tradicional + Placa aceleradora.
host
device
Dos arquitecturas, dos espacios de memoria,
dos filosofías distintas de programación.
¿Porqué?
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Potencia de Cálculo
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Ancho de Banda de Memoria
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
No hay que ser ingenuos
• ¿Cuánto de estos GFLOPS y GBps soy capaz de sacarle yo a
mi aplicación?
• ¿Mi aplicación es FLOPS-bound o GBps-bound?
• ¿Podré hacer una corrida de 132 días sin problemas?
• ¿Entra mi problema en la RAM?
• Ahora puedo calcular mucho más ¿Puedo hacer crecer mi
problema?
• ¿Cuál es el costo total de pasarse a GPU?
• ¿Es eficiente energéticamente?
• ¿Evacúa mucho calor? ¿Necesito refrigeración con nitrógeno?
• ¿Qué opina la Iglesia?
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Indicadores positivos
• “New Top500 list: 4x more GPU supercomputers”.
• Intel Xeon Phi.
Contraejemplo: el Top500-nov11 fue liderado por los SPARC64 VIIIfx,
y el Top500-jun12 fue liderado por los Power7.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
C para CUDA
Mezcla de
• Extensiones y restricciones de C/C++.
• Tipos de datos especiales.
• Bibliotecas.
Compilado por nvcc:
• Genera código CPU: x86 64 (lo hace gcc).
• Genera código GPU: PTX.
• Suma código runtime. (Ej: interoperar con el driver)
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Flujo de procesamiento
• Copiar datos host → device.
• Configurar la ejecución.
• Lanzar kernels.
• Copiar datos host ← device.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Introducción
Manejo de Memoria
Configuración de ejecución
Paralelismo de datos
Comunicación y sincronización
Desempeño y Debugging
Resumen
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Memoria CPU-GPU
Tenemos Unified Virtual Addressing–UVA (CUDA≥4.0)
• Punteros adornados para identificar CPU, GPU0, GPU1, . . .
• Elimina el paso por la CPU: direct access, direct transfer.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Memoria: pedido
#include <stdio.h> // printf
#include <assert.h>
#include <cuda.h> // CUDA library
#include <cutil_inline.h> // CUDA error checking
const unsigned int SIZE=1;
int main(void)
{
float *h_heat = NULL; // pointer to a float in host memory
float *d_heat = NULL; // pointer to a float in device memory
h_heat = (float *) calloc(SIZE, sizeof(float));
assert(h_heat);
cutilSafeCall(cudaMalloc(&d_heat, SIZE*sizeof(float)));
...
• cutilSafeCall() revisa errores en llamados a CUDA Lib.
• assert(p1) revisa que el puntero en el host no sea nulo.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Memoria: copias y devolución
...
*h_heat = 42.0f;
cutilSafeCall(cudaMemcpy(d_heat, h_heat, SIZE*sizeof(float),
cudaMemcpyDefault));
*h_heat = 0.0f;
cutilSafeCall(cudaMemcpy(h_heat, d_heat, SIZE*sizeof(float),
printf("Heat: %f\n", *h_heat);
cudaMemcpyDefault));
cutilSafeCall(cudaFree(d_heat));
free(h_heat);
return 0;
}
Esto imprime:
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Memoria: copias y devolución
...
*h_heat = 42.0f;
cutilSafeCall(cudaMemcpy(d_heat, h_heat, SIZE*sizeof(float),
cudaMemcpyDefault));
*h_heat = 0.0f;
cutilSafeCall(cudaMemcpy(h_heat, d_heat, SIZE*sizeof(float),
printf("Heat: %f\n", *h_heat);
cudaMemcpyDefault));
cutilSafeCall(cudaFree(d_heat));
free(h_heat);
return 0;
}
Esto imprime: 42.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Configuración de ejecución
Queremos lanzar hilos incrementando d heat.
__global__ void inc(volatile float *heat) {
*heat = *heat+1.0f;
}
int main(void)
{
...
inc<<<32768,512>>>(d_heat);
cutilCheckMsg("inc kernel failed");
cutilSafeCall(cudaDeviceSynchronize());
...
return 0;
}
Son 32768×512 hilos: paralelismo masivo.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Corrección
Resultado correcto
heat = 42 + 32768 × 512 = 16777258
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Corrección
Resultado correcto
heat = 42 + 32768 × 512 = 16777258
Histograma, 4096 ejecuciones, Tesla C2070
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Corrección
Resultado correcto
heat = 42 + 32768 × 512 = 16777258
Histograma, 4096 ejecuciones, Tesla C2070
¿Corrección?
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Nodeterminismo
Race conditions, lost updates, memory reorders, etc.
Los problemas de la concurrencia son muy notables.
(debería dar ≈ 16777258, pero retorna ≈ 1000.)
Lo que ejecuta: Fermi ISA y traducción
/*0000*/
/*0008*/
/*0010*/
/*0018*/
/*0020*/
/*0028*/
/*0030*/
/*0x00005de428004404*/
/*0x80009de428004000*/
/*0x9000dde428004000*/
/*0x00201f8584000000*/
/*0x00001c005000cfe0*/
/*0x00201f8594000000*/
/*0x00001de780000000*/
MOV R1, c [0x1] [0x100];
MOV R2, c [0x0] [0x20];
MOV R3, c [0x0] [0x24];
LD.E.CV R0, [R2];
FADD R0, R0, 0x3f800;
ST.E.WT [R2], R0;
EXIT;
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Nodeterminismo
Race conditions, lost updates, memory reorders, etc.
Los problemas de la concurrencia son muy notables.
(debería dar ≈ 16777258, pero retorna ≈ 1000.)
Lo que ejecuta: Fermi ISA y traducción
a1 := heat
a1 := a1 + 1.0
a2 := heat
a2 := a2 + 1.0
a3 := heat
a3 := a3 + 1.0
··· ≈ 16Mmás
heat := a1
heat := a2
heat := a3
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Nodeterminismo
Race conditions, lost updates, memory reorders, etc.
Los problemas de la concurrencia son muy notables.
(debería dar ≈ 16777258, pero retorna ≈ 1000.)
Lo que ejecuta: Fermi ISA y traducción
a1 := heat
a1 := a1 + 1.0
a2 := heat
a2 := a2 + 1.0
a3 := heat
a3 := a3 + 1.0
··· ≈ 16Mmás
heat := a1
heat := a2
heat := a3
Es una ejecución muy síncrona.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Solución: incrementos atómicos
__global__ void atomicinc(float *heat) {
atomicAdd(heat, 1.0f);
}
int main(void)
{
...
atomicinc<<<32768,512>>>(d_heat);
...
}
Ahora si
$ ./inc
16777258
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Introducción
Manejo de Memoria
Configuración de ejecución
Paralelismo de datos
Comunicación y sincronización
Desempeño y Debugging
Resumen
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Configuración de ejecución
Los hilos se agrupan en bloques.
Los bloques se agrupan en grillas.
Esquema bidimensional para
definición de hilos:
(threadIdx, blockIdx)
¿Porqué?
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Configuración de ejecución
Los hilos se agrupan en bloques.
Los bloques se agrupan en grillas.
Esquema bidimensional para
definición de hilos:
(threadIdx, blockIdx)
¿Porqué?
Limitar escalabilidad en la
comunicación y sincronización.
Restricción
threadIdx≤1024
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Dimensionalidad de bloques y grillas
Bloques y grillas pueden ser
uni, bi o tridimensionales.
¿Para qué?
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Dimensionalidad de bloques y grillas
Bloques y grillas pueden ser
uni, bi o tridimensionales.
¿Para qué?
Definir fácil el mapa hilo↔dato.
Restricciones
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Dimensionalidad de bloques y grillas
Bloques y grillas pueden ser
uni, bi o tridimensionales.
¿Para qué?
Definir fácil el mapa hilo↔dato.
Restricciones
threadIdx.{x, y} ≤ 1024
threadIdx.z ≤ 64
blockIdx.{x, y , z} ≤ 65536
Con estos límites no puedo tener más
de 65535 × 1024 threads
unidimensionales.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Configuración de bloques y grillas
Definir dimensiones de grillas y bloques.
dim3 grid_dim(2, 2);
dim3 block_dim(4, 2, 2);
my_kernel<<<grid_dim, block_dim>>>(d_array);
dim3 tripla con default values a 1.
<<<·,·>>> sintáxis para configuración de ejecución.
Introducción
Memoria
Configuración
Paralelismo
Comunicación
Desempeño
Resumen
Configuración de bloques y grillas
Definir dimensiones de grillas y
Comentarios de: CUDA (0)
No hay comentarios