Deep Dive: Ejecución de un Kernel CUDA en GPU NVIDIA

La ejecución de un programa paralelo en una GPU moderna es un proceso complejo que abarca múltiples capas de abstracción, desde el código fuente de alto nivel hasta las instrucciones de máquina específicas del hardware. Este "deep dive" revela cómo un simple kernel de suma vectorial en CUDA se traduce en una orquestación de compiladores, drivers de sistema operativo, y mecanismos de hardware de la GPU para lograr un paralelismo masivo. El problema fundamental que resuelve esta arquitectura es la gestión eficiente de miles de hilos de ejecución y el movimiento de datos entre la memoria del host y la del dispositivo, optimizando la latencia y el throughput en un entorno de computación heterogénea.

La relevancia actual de este tema radica en la omnipresencia de las GPUs en cargas de trabajo de alto rendimiento, desde la inteligencia artificial hasta la simulación científica. Comprender el camino crítico de un kernel es esencial para optimizar el rendimiento y diagnosticar cuellos de botella, especialmente cuando se busca maximizar la utilización de recursos en arquitecturas de GPU cada vez más complejas. La evolución de las arquitecturas de GPU y las herramientas de software asociadas, como CUDA, ha permitido abstraer gran parte de esta complejidad, pero un conocimiento profundo de los mecanismos subyacentes sigue siendo crucial para los ingenieros que operan a escala de hyperscaler.

Arquitectura del Sistema

El proceso comienza con nvcc, un compilador driver que coordina múltiples etapas. El código host se compila con un compilador C++ estándar, mientras que el código del dispositivo (__global__ functions) pasa por cicc (un compilador basado en LLVM) para generar PTX (Parallel Thread Execution), una ISA virtual. Luego, ptxas transforma el PTX en SASS (Streaming Assembler), la ISA nativa de la GPU. Finalmente, fatbinary empaqueta el SASS y el PTX (para compatibilidad futura y JIT) en un archivo fatbin, que se incrusta en el ejecutable final del host como una sección ELF (.nv_fatbin).

En tiempo de ejecución, un constructor oculto (__attribute__((__constructor__))) registra el fatbin con el runtime de CUDA, mapeando los punteros de función del host a los nombres "mangled" de los kernels del dispositivo. Cuando se invoca un kernel (ej. vadd<<<...>>>), un "host launch stub" generado por el compilador empaqueta los argumentos del kernel en un buffer en la memoria del host. Este stub llama a __cudaLaunch, que a su vez interactúa con la librería libcuda.so.1 (driver de usuario-modo) para iniciar la ejecución. libcuda se comunica con el driver de kernel (nvidia.ko) a través de ioctls, creando un "contexto" para la GPU.

La carga del módulo del kernel es "lazy" por defecto (desde CUDA 12.2), lo que significa que el cubin SASS se carga en la memoria de la GPU solo en la primera invocación del kernel. La comunicación entre CPU y GPU se realiza a través de un "channel" que incluye un pushbuffer (donde el driver escribe comandos GPU, o "methods") y un GPFIFO (un ring buffer de punteros a spans del pushbuffer). El driver llena el pushbuffer con métodos, incluyendo el "Queue Meta Data" (QMD) que describe la configuración del lanzamiento (grid/block dimensions, SASS start address, constant bank para argumentos), y avanza el cursor GP_PUT. Finalmente, el driver "toca el timbre" (doorbell) escribiendo en un registro MMIO mapeado, lo que notifica al "host engine" de la GPU que hay trabajo nuevo. El "host engine" lee el GP_PUT, procesa el GPFIFO y entrega el QMD al "compute work distributor" de la GPU. Este distribuidor asigna bloques a los Streaming Multiprocessors (SMs), que a su vez distribuyen "warps" a sus "warp schedulers". Cada scheduler selecciona warps elegibles basándose en "scoreboard barriers" y "stall counts" codificados en las instrucciones SASS por ptxas, ocultando la latencia de memoria y computación. Las unidades de carga/almacenamiento de los SMs realizan "request coalescing" para optimizar el acceso a la caché L1, L2 y VRAM GDDR6X. Una vez completada la ejecución, la GPU publica un "completion semaphore" que la CPU espera, y el motor de copia de la GPU transfiere los resultados (directamente desde la caché L2 si es posible) de vuelta a la memoria del host.

__global__ void vadd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) c[i] = a[i] + b[i];
}

// ... en main()
vadd<<<4096, 256>>>(da, db, dc, n);

#include <dlfcn.h>
#include <sys/mman.h>
// ... (struct Map, map_count)
void* mmap(void* addr, size_t length, int prot, int flags, int fd, off_t offset) {
    // ... original mmap call ...
    if (ret != MAP_FAILED && fd != -1 && map_count < 128) {
        // ... readlink and check for /dev/nvidia ...
        if (strstr(path, "/dev/nvidia")) {
            // ... store map info ...
        }
    }
    return ret;
}

Fundamentos Teóricos

El concepto de una ISA virtual como PTX, que luego se compila a una ISA nativa (SASS), es un patrón bien establecido en la informática, similar a la JVM de Java o el LLVM IR. Esto permite la portabilidad del código fuente a través de diferentes generaciones de hardware GPU sin necesidad de recompilación completa para cada arquitectura. La idea de "warp scheduling" y la ocultación de latencia mediante el cambio rápido de contexto entre múltiples hilos ligeros es una aplicación directa de principios de diseño de procesadores tolerantes a la latencia, donde en lugar de predecir ramas o ejecutar fuera de orden (como en CPUs), se explota el paralelismo masivo y la capacidad de cambiar de hilo rápidamente para mantener las unidades de ejecución ocupadas. Los "scoreboard barriers" y "stall counts" son una forma de "static scheduling" asistido por el compilador, un enfoque que contrasta con el "dynamic scheduling" de CPUs out-of-order, y que se elige en GPUs por su simplicidad y eficiencia energética, dado el gran número de hilos. Este diseño se alinea con los principios de la arquitectura SIMT (Single Instruction, Multiple Thread) popularizada por NVIDIA, que a su vez tiene raíces en el procesamiento vectorial y paralelo de los años 70 y 80, como se discute en trabajos fundamentales sobre arquitecturas de computación paralela.