Programación CUDA

Question 1

Device

Answer 1

• Coprocesador de la CPU (host)
• Posee su propia memoria DRAM
• Ejecuta “muchos” threads en paralelo

Question 2

Diferencias entre threads GPU y CPU

Answer 2

• Los threads de GPU son extremadamente livianos (muy poco overhead para crearlos)
• La GPU necesita 1000s de threads para alcanzar la eficiencia

Question 3

Algoritmo básico

Answer 3

1. Instrucciones en el host
2. Envío de datos de host a device
3. Se procesa en device
4. Se recuperan los datos
5. Continúan instrucciones en el host

Question 4

Array de threads y SPMT

Answer 4

• 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

Question 5

Bloques de threads

Answer 5

• 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

Question 6

Tipos de funciones

Answer 6

__device__: device a device
__global__: host a device
__host__: host a host

Question 7

Lanzamiento de kernels

Answer 7

• 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

Question 8

Keywords

Answer 8

• threadIdx
• blockIdx
• blockDim
• gridDim

Todos se eacceden con .x/.y/.z
Generalmente int idX = threadIdx.x + blockDim.x * blockIdx.x

Question 9

Algunas funciones

Answer 9

• cudaDeviceSynchronize()
• _ _syncthreads
• cudaMalloc(*var, size)
• cudaFree(*var)
• cudaMemcpy(*dest, *orig, bytes, tipo)

Question 10

Acceso Coalesced

Answer 10

• 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

Question 11

Memoria compartida

Answer 11

• 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)

Question 12

Conflicto de bancos

Answer 12

• 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)

Question 13

Recomendaciones

Answer 13

• 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