La Pila de Software CUDA

Tres Capas de CUDA

La palabra "CUDA" está sobrecargada. Dependiendo del contexto, puede referirse a una arquitectura de dispositivo, un modelo de programación o una plataforma de software. Este artículo se enfoca en la plataforma de software — las capas de código que se ubican entre tu script de Python (o programa en C++) y el hardware de la GPU. Comprender esta pila importa porque cada sesión de depuración de rendimiento eventualmente llega a una de estas capas, y saber en cuál estás mirando determina si necesitas actualizar un driver, cambiar una biblioteca o reescribir un kernel.

La pila está organizada en tres capas, cada una construyéndose sobre la anterior:

• Capa de driver: nvidia.ko (el módulo del kernel de Linux) y libcuda.so (la API del Driver CUDA). Esta es la frontera entre el sistema operativo y el silicio de la GPU.
• Capa de runtime: libcudart.so (la API de Runtime CUDA), nvcc (el compilador) y nvrtc (el compilador en tiempo de ejecución). Esta es la capa con la que la mayoría de los desarrolladores de CUDA C++ interactúan directamente.
• Capa de bibliotecas: cuBLAS, cuDNN, cuFFT, cuSPARSE y otras. Kernels preconstruidos y optimizados manualmente para operaciones comunes. Esta es la capa a la que PyTorch, TensorFlow y JAX hacen llamadas.

El siguiente diagrama muestra cómo se apilan estas capas (consulta el Modal GPU Glossary sobre la plataforma de software CUDA para contexto adicional):

Your Code (Python / C++ CUDA)
        │
        ▼
┌───────────────────────────────────────────────┐
│  Library Layer                                │
│  cuBLAS (matmul) · cuDNN (conv, attn) · ...   │
├───────────────────────────────────────────────┤
│  Runtime Layer                                │
│  libcudart.so · nvcc · nvrtc                  │
├───────────────────────────────────────────────┤
│  Driver Layer                                 │
│  libcuda.so · nvidia.ko (kernel module)       │
├───────────────────────────────────────────────┤
│  Hardware                                     │
│  GPU (SMs, memory, interconnect)              │
└───────────────────────────────────────────────┘

Cada capa tiene una responsabilidad bien definida. La capa de driver gestiona la GPU como dispositivo de hardware — asignando memoria, planificando trabajo, manejando interrupciones. La capa de runtime proporciona una API de programación más ergonómica y la cadena de compilación que convierte archivos fuente .cu en código ejecutable de GPU. La capa de bibliotecas ofrece kernels precompilados y probados en batalla para las operaciones que dominan las cargas de trabajo de deep learning y computación científica — multiplicaciones de matrices, convoluciones, FFTs, etcétera.

La separación no es meramente conceptual. Cada capa se distribuye como un conjunto distinto de bibliotecas compartidas con su propio versionado. Una actualización de driver (nvidia.ko + libcuda.so) no requiere recompilar tu código CUDA; una actualización de cuDNN no requiere un nuevo driver. Este desacoplamiento es una de las razones por las que CUDA ha sido tan duradero — permite a NVIDIA mejorar cualquier capa de forma independiente, y permite a los usuarios fijar versiones específicas cuando la estabilidad importa.

La Capa de Driver

La capa de driver es la capa de software más baja en la pila CUDA — se comunica directamente con el hardware de la GPU y media todo acceso a él. Todo lo que está por encima (la API de runtime, las bibliotecas, tus kernels) eventualmente pasa por esta capa.

nvidia.ko: el módulo del kernel de Linux. Cuando se carga, registra la GPU como un dispositivo en el sistema operativo, gestiona la asignación de memoria de GPU a nivel físico, planifica el lanzamiento de kernels en los procesadores de comandos de la GPU y maneja las interrupciones de hardware. Sin este módulo cargado ( lsmod | grep nvidia te lo dirá), la GPU es invisible para el sistema operativo — es simplemente un dispositivo PCIe desconocido. En sistemas con múltiples GPUs, una sola instancia del módulo las gestiona a todas.

libcuda.so (API del Driver CUDA): la biblioteca compartida en espacio de usuario que se comunica con nvidia.ko mediante llamadas al sistema ioctl . Proporciona funciones de bajo nivel para cada paso de la programación GPU: crear un contexto CUDA (análogo a un proceso en la CPU), asignar memoria de dispositivo, cargar kernels compilados desde PTX o binario SASS, establecer argumentos del kernel y lanzar kernels. La API del driver es verbosa y explícita — gestionas contextos, módulos y memoria manualmente. La mayoría de los desarrolladores nunca la usan directamente; la API de runtime (libcudart.so) la envuelve con una interfaz más amigable. Pero entender que la API del driver existe es importante porque es la capa donde se aplica la compatibilidad de versiones: la versión del driver CUDA determina la versión máxima del toolkit CUDA que tu sistema soporta.

nvidia-smi: una herramienta de línea de comandos construida sobre NVML (libnvml.so) que muestra el estado de la GPU de un vistazo — temperatura, uso de memoria, utilización de GPU, consumo de energía, procesos en ejecución y la versión del driver. Piensa en ella como el comando top o htop para GPUs. Típicamente es la primera herramienta a la que recurres cuando depuras problemas de GPU: "¿Es visible la GPU? ¿Hay algo ejecutándose en ella? ¿Cuánta memoria está libre?"

La Capa de Runtime

La capa de runtime se sitúa por encima del driver y proporciona dos cosas: una API de programación más amigable y la cadena de compilación que convierte el código fuente CUDA en binarios ejecutables de GPU.

libcudart.so (API de Runtime CUDA): envuelve la API del driver con funciones más simples y automáticas. En lugar de crear manualmente contextos CUDA, cargar módulos y establecer argumentos del kernel uno por uno, escribes cudaMalloc para asignar memoria de dispositivo, cudaMemcpy para transferir datos entre host y dispositivo, y la sintaxis de triple chevron kernel<<<grid, block>>>(args) para lanzar kernels. La API de runtime gestiona contextos CUDA automáticamente (uno por hilo por defecto) y maneja la inicialización perezosa, así que no necesitas código de configuración repetitivo. La gran mayoría del código CUDA C++ usa la API de runtime, no la API del driver.

nvcc (NVIDIA CUDA Compiler Driver): compila archivos .cu — código fuente C++ con extensiones CUDA — en código ejecutable. Pero nvcc no es un compilador monolítico. Es un orquestador que divide el trabajo entre el compilador del host y el pipeline de compilación de dispositivo de NVIDIA:

your_kernel.cu (C++ CUDA source)
     │
     ▼ nvcc separates host and device code
     ├── Host code → GCC/Clang → CPU executable
     └── Device code → cicc → PTX → ptxas → SASS
                                              │
                                              ▼
                                    Embedded in the executable
                                    (or .cubin / .fatbin file)

El proceso funciona de la siguiente manera. nvcc lee el archivo .cu y lo separa en dos flujos: código de host (C++ regular — todo lo que está fuera de las funciones __global__ y __device__ ) y código de dispositivo (los kernels de GPU). El código de host se pasa al compilador C++ del sistema — GCC o Clang en Linux, MSVC en Windows — y se compila a un archivo objeto CPU normal. El código de dispositivo pasa por el pipeline propio de NVIDIA: primero a cicc (que produce PTX, una representación intermedia portable), luego a ptxas (que compila PTX a SASS, el código máquina real de GPU para una arquitectura específica). Los dos se enlazan juntos en un único ejecutable, con el binario SASS incrustado como una sección de datos.

nvrtc (Compilador en Tiempo de Ejecución): permite compilar kernels CUDA C++ en tiempo de ejecución en lugar de anticipadamente con nvcc. Pasas una cadena de código fuente CUDA a nvrtc, y devuelve PTX compilado (o SASS) que puedes cargar y lanzar a través de la API del driver. Este es el mecanismo detrás de la compilación just-in-time (JIT) en el ecosistema CUDA. Bibliotecas que generan kernels de GPU dinámicamente — como Triton, XLA y TVM — dependen en gran medida de nvrtc (o de la compilación JIT de PTX en el driver). La compilación JIT añade latencia de arranque, pero permite generar kernels especializados para formas, dtypes o características de hardware específicas en tiempo de ejecución, lo que puede producir mejor rendimiento que la compilación estática.

La Capa de Bibliotecas: cuBLAS y cuDNN

La capa de bibliotecas proporciona kernels de GPU precompilados y altamente optimizados para patrones computacionales comunes. Estas bibliotecas son los verdaderos caballos de batalla detrás de los frameworks de deep learning — cuando llamas a torch.matmul o nn.Conv2d en un tensor GPU, PyTorch despacha a una de estas bibliotecas en lugar de ejecutar un kernel CUDA ingenuo.

cuBLAS (CUDA Basic Linear Algebra Subprograms): la implementación GPU del estándar BLAS. Proporciona multiplicación de matrices ( sgemm para FP32, dgemm para FP64, hgemm para FP16), productos punto, productos matriz-vector, resoluciones triangulares y otras primitivas de álgebra lineal. Cada pase forward de torch.matmul , torch.mm y nn.Linear eventualmente llama a cuBLAS en GPU.

La convención de nombres sigue la tradición BLAS: cublas[S|D|H]gemm donde S = precisión simple (FP32), D = doble (FP64), H = media (FP16). gemm significa General Matrix Multiply:

Los escalares α y β permiten fusionar una escala y acumulación en la propia matmul, evitando un lanzamiento de kernel separado. En deep learning, α = 1 y β = 0 es el caso común (multiplicación de matrices simple), pero la forma general es útil para operaciones como agregar un sesgo o acumular gradientes.

cuDNN (CUDA Deep Neural Network library): proporciona implementaciones GPU de las operaciones que dominan la computación de redes neuronales — convoluciones (con múltiples algoritmos: Winograd, basado en FFT, GEMM implícito), pooling, normalización por lotes, softmax, funciones de activación, y en versiones más recientes, atención multi-cabeza fusionada (incluyendo implementaciones al estilo FlashAttention). Cada nn.Conv2d , F.softmax y nn.BatchNorm2d eventualmente llama a cuDNN en GPU.

¿Por qué precompilados en lugar de compilados al instalar? Porque el rendimiento a este nivel requiere ajuste manual por equipos de arquitectos de GPU que explotan características específicas del hardware — operaciones de matriz a nivel de warp en Tensor Cores, instrucciones de copia de memoria asincrónica, patrones de swizzling de memoria compartida, optimizaciones a nivel de registros — que los compiladores automatizados generalmente no pueden descubrir. Los binarios SASS resultantes son frecuentemente 10–100× más rápidos que una implementación CUDA ingenua de la misma operación. NVIDIA efectivamente amortiza el costo de este ajuste a nivel de experto entre todos los usuarios de la biblioteca.

Compute Capability y Compatibilidad PTX

No todas las GPUs de NVIDIA soportan las mismas características. Los Tensor Cores, la aritmética BF16, FP8 y las instrucciones de copia de memoria aceleradas por hardware se introdujeron cada una en una generación específica de GPU. NVIDIA rastrea esto con Compute Capability — un esquema de versionado que indica al software exactamente qué características de hardware están disponibles en una GPU determinada.

Los hitos principales relevantes para el deep learning moderno:

• SM 7.0 (Volta — V100): introdujo los Tensor Cores de primera generación (multiply-accumulate de matrices en FP16). Esta fue la GPU que hizo práctico el entrenamiento con precisión mixta.
• SM 8.0 (Ampere — A100): añadió soporte para BF16, TF32 (un formato de 19 bits para aceleraciones transparentes de FP32), Tensor Cores de 3ª generación y soporte de dispersión.
• SM 8.9 (Ada Lovelace — RTX 4090): añadió FP8 (formatos E4M3 y E5M2) y Tensor Cores de 4ª generación. FP8 reduce a la mitad la huella de memoria en comparación con FP16.
• SM 9.0 (Hopper — H100): añadió el Tensor Memory Accelerator (TMA) para movimiento masivo de datos asíncrono, operaciones de matrices a nivel de grupo de warps (WGMMA) y el Transformer Engine para cambio dinámico FP8/FP16.

¿Por qué importa esto para la pila de software? Porque SASS — el código máquina final de GPU — es específico de una compute capability . SASS compilado para SM 8.0 no se ejecutará en una GPU SM 7.0, y tampoco usará automáticamente las características de SM 9.0 si lo ejecutas en un H100. PTX, en cambio, es compatible hacia adelante : PTX compilado para SM 7.0 puede ser compilado JIT por el driver a SASS en una GPU SM 9.0. La compensación es la latencia de arranque — la compilación JIT ocurre en tiempo de carga — y optimizaciones potencialmente perdidas que el compilador JIT del driver no conoce.

Por esto existen los fat binaries . Un solo archivo .fatbin puede contener SASS para múltiples compute capabilities más PTX como respaldo. En tiempo de ejecución, el driver selecciona el mejor SASS disponible para la GPU actual, recurriendo a la compilación JIT del PTX si no se encuentra SASS compatible. Esto asegura que el código se ejecute en cualquier GPU compatible con CUDA — con rendimiento óptimo en arquitecturas que tienen SASS precompilado, y rendimiento ligeramente menor con compilación JIT en las demás. También explica por qué las bibliotecas CUDA como cuBLAS son tan grandes: incluyen SASS para cada generación de GPU soportada.

import numpy as np

# Compute capabilities and their key features
architectures = [
    ("SM 7.0", "Volta (V100)",        2017, "1st-gen Tensor Cores, FP16 MMA"),
    ("SM 7.5", "Turing (RTX 2080)",   2018, "INT8 Tensor Cores, RT Cores"),
    ("SM 8.0", "Ampere (A100)",       2020, "BF16, TF32, 3rd-gen Tensor Cores"),
    ("SM 8.6", "Ampere (RTX 3090)",   2020, "Consumer Ampere, same TC gen"),
    ("SM 8.9", "Ada Lovelace (4090)", 2022, "FP8, 4th-gen Tensor Cores"),
    ("SM 9.0", "Hopper (H100)",       2022, "TMA, WGMMA, Transformer Engine"),
]

print("NVIDIA Compute Capability Timeline")
print("=" * 72)
print(f"{'SM':>6}  {'Architecture':<25} {'Year':<6} {'Key Features'}")
print("-" * 72)
for sm, arch, year, features in architectures:
    print(f"{sm:>6}  {arch:<25} {year:<6} {features}")

print()
print("Compatibility rules:")
print("  SASS: runs ONLY on the exact SM version it was compiled for")
print("  PTX:  forward-compatible (SM 7.0 PTX → JIT-compiled on SM 9.0)")
print("  Fat binary: bundles multiple SASS + PTX fallback → runs everywhere")

Perfilado: nvidia-smi, Nsight y CUPTI

Una vez que tu código se ejecuta en la GPU, la siguiente pregunta es si se está ejecutando bien . ¿Está usando los Tensor Cores? ¿Está limitado por el ancho de banda de memoria? ¿Está gastando más tiempo copiando datos que computando? El ecosistema CUDA proporciona un conjunto de herramientas de perfilado a diferentes niveles de granularidad, y elegir la correcta depende de qué pregunta estás haciendo.

• nvidia-smi: la verificación rápida de salud. Muestra la utilización de GPU, uso de memoria, temperatura, consumo de energía y qué procesos están usando la GPU. Es la primera herramienta que usas — no para análisis profundo de rendimiento, sino para confirmar que la GPU está activa y aproximadamente qué tan cargada está. Ejecuta nvidia-smi dmon para monitoreo continuo.
• NVML (libnvml.so): la biblioteca C detrás de nvidia-smi. Proporciona acceso programático a las mismas métricas de GPU — útil para construir dashboards, sistemas de alertas o registrar la salud de la GPU en producción. Bibliotecas como pynvml la envuelven para Python.
• Nsight Systems: un perfilador a nivel de sistema que traza la línea temporal completa de ejecución — actividad de CPU, lanzamientos de kernels de GPU, transferencias de memoria (H2D, D2H, D2D), llamadas a la API CUDA y sus relaciones temporales. Responde preguntas a nivel macro: ¿dónde se está gastando el tiempo? ¿Está la GPU inactiva entre lanzamientos de kernels? ¿Se superponen las copias de memoria con el cómputo? ¿Es la CPU el cuello de botella? La salida es una visualización de línea temporal donde puedes hacer zoom en intervalos específicos.
• Nsight Compute: un perfilador a nivel de kernel que profundiza en el rendimiento de un solo kernel de GPU. Mide la ocupación (cuántos warps están activos en relación al máximo), utilización del ancho de banda de memoria, throughput de cómputo, mezcla de instrucciones, razones de estancamiento y estadísticas a nivel de warp. Responde preguntas a nivel micro: ¿este kernel está limitado por cómputo o por memoria? ¿Dónde se están estancando los warps? ¿Se está usando la memoria compartida eficientemente?
• CUPTI (CUDA Profiling Tools Interface): la API de bajo nivel sobre la cual se construyen tanto Nsight Systems como Nsight Compute. CUPTI proporciona acceso a contadores de rendimiento de hardware, trazado de actividad y mecanismos de callback. Herramientas de terceros (como el trazado CUDA del PyTorch Profiler) también usan CUPTI internamente.

El flujo de trabajo típico de perfilado va de grueso a fino: comienza con nvidia-smi para verificar que la GPU está activa, usa Nsight Systems para encontrar qué kernels u operaciones de memoria dominan el tiempo de ejecución, y luego usa Nsight Compute para profundizar en los kernels específicos que parecen subóptimos. Saltar directamente a Nsight Compute sin primero entender el panorama a nivel de sistema tiende a producir optimizaciones que no importan — podrías perfeccionar un kernel que solo representa el 2% del tiempo total de ejecución.

Quiz

Pon a prueba tu comprensión de la pila de software CUDA — las capas de driver, runtime y bibliotecas, el pipeline de compilación, las compute capabilities y las herramientas de perfilado.