P4b · Taller de optimización de transferencias Host-Device y memoria compartida

Contenido del programa: §20.4 · contenidos 4.5 y 4.6 Horas de laboratorio: 4 · Modalidad: individual

1. Propósito y objetivos

Dominar la jerarquía de memoria en arquitecturas GPU NVIDIA, diferenciando el acceso a memoria global de alta latencia (~400-800 ciclos de reloj) frente a la memoria compartida en chip (Shared Memory, ~20-30 ciclos de reloj), implementando una reducción paralela en árbol optimizada y minimizando el cuello de botella del bus PCIe mediante transferencias asíncronas y streams.

2. Fundamentación técnica

2.1 Memoria Compartida (__shared__)

La memoria compartida reside físicamente en cada multiprocesador de transmisión (Streaming Multiprocessor, SM) y es compartida exclusivamente por todos los hilos pertenecientes al mismo bloque.

• Permite reutilizar datos leídos de la memoria global evitando lecturas repetitivas al bus DRAM.
• Requiere barreras de sincronización explícitas con __syncthreads() para asegurar que todos los hilos del bloque hayan escrito sus datos antes de que comience la fase de lectura.

2.2 Reducción en árbol libre de divergencia de warps

Al sumar elementos dentro de un bloque, una reducción inocente con if (tid % (2 * stride) == 0) provoca una severa divergencia de warps (warp divergence), donde la mitad de los hilos quedan inactivos mientras la otra mitad computa. La técnica optimizada reindexa los hilos de forma contigua:

for (unsigned int s = blockDim.x / 2; s > 0; s >>= 1) {
    if (tid < s) {
        sdata[tid] += sdata[tid + s];
    }
    __syncthreads();
}

3. Ejercicios guiados

Ejercicio 1: Reducción paralela de suma con __shared__

1. Escribir un kernel reducir_bloque que cargue datos de la memoria global a un arreglo __shared__ float sdata[256].
2. Implementar la reducción en árbol contigua para que el hilo 0 de cada bloque calcule la suma parcial de su bloque.
3. El hilo 0 escribe el resultado parcial en un arreglo intermedio en memoria global.
4. Medir la ganancia de velocidad frente a una reducción atómica global (atomicAdd).

Ejercicio 2: Mediciones de tiempo de transferencia PCIe vs. Cómputo

Utilizando eventos de CUDA (cudaEvent_t, cudaEventRecord, cudaEventElapsedTime):

1. Medir el tiempo invertido en cudaMemcpy (Host -> Device y Device -> Host).
2. Medir el tiempo puro de ejecución del kernel en GPU.
3. Calcular la fracción del tiempo total absorbida por el bus PCIe. Comprobar experimentalmente por qué transferir arreglos pequeños destruye la ventaja de usar una GPU.

4. Entregables

1. Código fuente reduccion_shared.cu con comentarios detallados.
2. Gráfica comparativa de tiempo total vs. tiempo de transferencia PCIe para tamaños de vector desde hasta elementos.