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 hilos (threadsPerBlock), el identificador global del elemento asignado a un hilo es:

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

Si el tamaño del vector no es múltiplo exacto de , 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 elementos comparando el tiempo de ejecución del kernel con el tiempo de la CPU.