---
titulo: "Práctica P4a: Taller de escritura de kernels CUDA elementales"
modulo: M4
ec: EC4
contenido: "4.1, 4.2, 4.4"
semanas: "11"
horas: 4
tipo: "Taller en laboratorio"
asignatura: "Sistemas Paralelos"
sigla: TEL-420
docente: "Ing. Elias Cassal Baldiviezo"
institucion: "Universidad Autónoma Juan Misael Saracho — Facultad de Ciencias Integradas de Yacuiba / F.I.R.N.T."
---

# P4a · Taller de escritura de kernels CUDA elementales y jerarquía de ejecución

**Contenido del programa:** §20.4 · contenidos 4.1, 4.2 y 4.4  
**Horas de laboratorio:** 4 · **Modalidad:** individual  

---

## 1. Propósito y objetivos

Familiarizarse con el modelo de programación masivamente paralelo de NVIDIA CUDA en C/C++, comprendiendo la descomposición jerárquica de hilos (*Grids*, *Blocks* y *Threads*), los calificadores de función (`__global__`, `__device__`, `__host__`), la gestión de memoria en el dispositivo (`cudaMalloc`, `cudaMemcpy`, `cudaFree`) y el cálculo de índices globales unidimensionales para operaciones vectoriales elementales.

---

## 2. Fundamentación técnica del modelo SIMT

La arquitectura de GPU se basa en el paradigma **SIMT (*Single Instruction, Multiple Threads*)**: miles de hilos ligeros ejecutan el mismo código binario (kernel) sobre diferentes elementos de datos.

### 2.1 Mapeo de índices globales
Para un vector unidimensional procesado por bloques de $B$ hilos (`threadsPerBlock`), el identificador global del elemento asignado a un hilo es:

```c
int idx = blockIdx.x * blockDim.x + threadIdx.x;
```

Si el tamaño del vector $N$ no es múltiplo exacto de $B$, la grilla debe dimensionarse con techo entero:
$$G = \left\lceil \frac{N}{B} \right\rceil = \frac{N + B - 1}{B}$$
y dentro del kernel debe incluirse una condición de guarda obligatoria: `if (idx < N)`.

---

## 3. Ejercicios guiados

### Ejercicio 1: Suma de vectores en GPU (`vec_add.cu`)
1. Implementar el kernel con calificador `__global__ void vec_add(const float *A, const float *B, float *C, int n)`.
2. Asignar memoria en host con `malloc` y en GPU con `cudaMalloc`.
3. Transferir los datos de host a device con `cudaMemcpy(..., cudaMemcpyHostToDevice)`.
4. Lanzar el kernel configurando `threadsPerBlock = 256` y calcular `blocksPerGrid = (n + 255) / 256` mediante la sintaxis de triple chevron:
   ```c
   vec_add<<<blocksPerGrid, threadsPerBlock>>>(d_A, d_B, d_C, n);
   ```
5. Comprobar errores de lanzamiento con `cudaGetLastError()` y sincronizar el dispositivo con `cudaDeviceSynchronize()`.
6. Transferir el resultado de vuelta al host con `cudaMemcpy(..., cudaMemcpyDeviceToHost)` y validar contra la suma secuencial de la CPU.

### Ejercicio 2: Producto punto / escalar de vectores
Implementar una versión preliminar de producto punto donde cada hilo calcula la multiplicación elemento a elemento `prod[idx] = A[idx] * B[idx]`.
- Analizar por qué la reducción de la suma acumulada de todos los elementos representa un desafío en GPU y requiere sincronización a nivel de bloque (`__syncthreads()`).

---

## 4. Entregables
1. Código fuente `vec_add.cu` compilable con `nvcc -O2 -arch=sm_75 vec_add.cu -o vec_add` (o arquitectura correspondiente a la GPU de laboratorio).
2. Reporte de verificación para $N = 10^7$ elementos comparando el tiempo de ejecución del kernel con el tiempo de la CPU.