PDF de programación - CUDA

Imágen de pdf CUDA

CUDAgráfica de visualizaciones

Publicado el 22 de Junio del 2018
415 visualizaciones desde el 22 de Junio del 2018
955,4 KB
73 paginas
Creado hace 11a (02/08/2012)
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
  • Links de descarga
http://lwp-l.com/pdf12052

Comentarios de: CUDA (0)


No hay comentarios
 

Comentar...

Nombre
Correo (no se visualiza en la web)
Valoración
Comentarios...
CerrarCerrar
CerrarCerrar
Cerrar

Tienes que ser un usuario registrado para poder insertar imágenes, archivos y/o videos.

Puedes registrarte o validarte desde aquí.

Codigo
Negrita
Subrayado
Tachado
Cursiva
Insertar enlace
Imagen externa
Emoticon
Tabular
Centrar
Titulo
Linea
Disminuir
Aumentar
Vista preliminar
sonreir
dientes
lengua
guiño
enfadado
confundido
llorar
avergonzado
sorprendido
triste
sol
estrella
jarra
camara
taza de cafe
email
beso
bombilla
amor
mal
bien
Es necesario revisar y aceptar las políticas de privacidad