Programación CUDA Flashcards
(13 cards)
Device
- Coprocesador de la CPU (host)
- Posee su propia memoria DRAM
- Ejecuta “muchos” threads en paralelo
Diferencias entre threads GPU y CPU
- Los threads de GPU son extremadamente livianos (muy poco overhead para crearlos)
- La GPU necesita 1000s de threads para alcanzar la eficiencia
Algoritmo básico
- Instrucciones en el host
- Envío de datos de host a device
- Se procesa en device
- Se recuperan los datos
- Continúan instrucciones en el host
Array de threads y SPMT
- Un kernel es ejecutado por un array de threads todos corriendo el mismo código (SPMT)
- Cada thread tiene un ID que se puede utilizar para computar direcciones de memoria y tomar decisiones de control
Bloques de threads
- Permite dividir un array en multiples bloques
- Threads de un mismo bloque cooperan via shared, operaciones atomic y barrier de sync.
- Threads de diferentes bloques no pueden cooperar directamente (atomic en global)
- Los bloques son ejecutados en cualquier orden y en diferentes MPs
Tipos de funciones
__device__: device a device
__global__: host a device
__host__: host a host
Lanzamiento de kernels
- Para invocar un kernel es necesario determinar “una configuración de ejecución”
- Las llamadas a kernel son asincrónicas. Es necesario usar sincronizaciones explícitas
Keywords
- threadIdx
- blockIdx
- blockDim
- gridDim
Todos se eacceden con .x/.y/.z
Generalmente int idX = threadIdx.x + blockDim.x * blockIdx.x
Algunas funciones
- cudaDeviceSynchronize()
- _ _syncthreads
- cudaMalloc(*var, size)
- cudaFree(*var)
- cudaMemcpy(*dest, *orig, bytes, tipo)
Acceso Coalesced
- El acceso a memoria global es por segmentos de 128bytes
- Si no se usan todos los datos del segmento, se desperdicia ancho de banda
- El acceso no alineado es más costoso que el alineado
- Cada solicitud de acceso de un warp se parte en varias solicitudes. Si todas las solicitudes acceden al mismo segmento, se hace una sola y se usan todos los datos
- Si los accesos están distribuidos entre distinto segmentos se realizan multiples solicitudes y hay datos transferidos que no son usados
Memoria compartida
- Cientos de veces más rápida que la global
- Hay una en cada MP
- Alcance a nivel de bloque
- El uso y contenido se define explícitamente por el programador
- El puede definir en forma dinámica con extern
- Se suele utilizar como una especie de caché para reducir los accesos a global mem
- Permite evitar accesos no coalesced a mem glob (mediante tiling)
Conflicto de bancos
- Memoria compartida se divide en módulos del mismo tamaño llamados bancos
- Palabras de 32 bits contiguas, están en bancos contiguos
- Bancos pueden ser accedidos simultáneamente a nivel de warp
- Lecturas o escrituras en bancos distintos pueden ser atendidas simultáneamente
- Cada banco puede atender una solicitud por ciclo por warp
- Si dos solicitudes caen en el mismo banco, se produce un conflicto
- No hay conflicto si todos los hilos acceden a distintos bancos, o si todos acceden a la misma palabra (broadcast)
Recomendaciones
- Evitar divergencia entre hilos de un warp
- Número de bloques mayor al número de MPs
- Número de hilos por bloque múltiplo de 32