PDF de programación - CUDA

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

CUDA

Carlos Bederián, Nicolás Wolovick

FaMAF, Universidad Nacional de Córdoba, Argentina

19 de Febrero de 2013

[email protected]

Revisión 3869, 2013-02-19

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Motivación

Arquitectura

Manejo de Memoria

Configuración de ejecución

Paralelismo de datos

Comunicación y sincronización

Desempeño y Debugging

Resumen

Motivación

Arch

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.

Motivación

Arch

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.

Motivación

Arch

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é?

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Potencia de Cálculo

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Ancho de Banda de Memoria

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

En realidad. . .

Es una solución de compromiso con respecto a las CPUs.

En números crudos:

• 10x de throughput.
• 0.1x de latencia.

Hay muchos trade-offs detrás de esto.

Cuidado, la eficiencia es un tema complicado.
Eficiencia para HPL

• CPU Xeon E ˜95 %.
• GPU Fermi ˜65 %.
• GPU Kepler ˜78 %.

Motivación

Arch

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/clang).
• Genera código GPU: PTX.
• Suma código runtime. (Ej: interoperar con el driver)

Motivación

Arch

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.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Motivación

Arquitectura

Manejo de Memoria

Configuración de ejecución

Paralelismo de datos

Comunicación y sincronización

Desempeño y Debugging

Resumen

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Streaming Processor – SM

• Procesador SIMD de 32 lanes, aka warp.
• Memoria local (de contextos): muchos
warps ejecutando el mismo programa.

• Memoria compartida: comunicación entre

warps.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

SM – Memory latency hidding

Jerarquía de memoria

Nivel
Registros
Shared
Global

BW

8.1 TB/s
1.3 TB/s
177 GB/s

Latencia
1 ciclo

¿6? ciclos
400 ciclos

Ocultamiento de latencia

• Antes de esperar por mem&ops: ctx switch.
• Ctx switch muy rápido. (solo hay que mover un ptr.)
• +ctx memory, +warps ejecutando, +ocultamiento de latencia.

Superposición de comunicación y computación,
o de computación y computación!

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

SM – identificadores, divergencia de control

Cada lane tiene un id [0 . . . 31]. Cada warp tiene un id [0 . . . ].
Puedo identificar cada uno de los lanes globalmente.

¡Flujo divergente de control?

if (0 == lane%2) {

ISETP.LT.U32.AND P0, pt, R7, 0x4, pt;

@!P0 IADD R4, R2, R4;

++a;
} else {
++b;

}

Dos soluciones

• Instrucciones predicadas.
• Re-ejecución.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

SM – planificador externo

Alimentar el SM con warps.

• Acomodar todos los warps posibles en un SM.

• Cantidad de memoria de contexto (regs).
• Cantidad de memoria compartida (shared).
• Límite en la cantidad de warps.
• Límite en la cantidad de bloques.

• Cuando todos los warps terminan, seleccionar un nuevo

conjunto.

Planificador de dos niveles.

1. Interno: preemptive usando ctx&shared mem.
2. Externo: batch de warps.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Características dadas por la implementación

• Grado de concurrencia en SM está dado por #registros y

shared mem que usa el programa.

• La memoria compartida tiene el tiempo de vida del warp

batch.

• Solo se pueden sincronizar dentro de un warp batch.

La clave es entender como funciona el planificador de dos niveles.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Automatic Scalability

Si la aplicación se divide en warp batchs independientes;

+SMs ⇒ +más rápido.

¡Como no va a escalar si no hay comunicación ni sincronización!

El fabricante vende #SMs según el bolsillo del cliente.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Warps y memoria global

El caso ideal

No cambia nada si se pierde secuencialidad

Accesos desalineados pero secuenciales

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Motivación

Arquitectura

Manejo de Memoria

Configuración de ejecución

Paralelismo de datos

Comunicación y sincronización

Desempeño y Debugging

Resumen

Motivación

Arch

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.

Motivación

Arch

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 uint 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(h heat) revisa que el puntero en el host no sea nulo.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Memoria: ida y vuelta

...
*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:

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Memoria: ida y vuelta

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

Motivación

Arch

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.

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Corrección

Resultado correcto

heat = 42 + 32768 × 512 = 16777258

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

Corrección

Resultado correcto

heat = 42 + 32768 × 512 = 16777258

Histograma, 4096 ejecuciones, Tesla C2070

Motivación

Arch

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?

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

No-determinismo

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;

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

No-determinismo

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

··· ≈ 16M más

heat := a1

heat := a2

heat := a3

Motivación

Arch

Memoria

Configuración

Paralelismo

Comunicación

Desempeño

Resumen

No-determinismo

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