Mejora del rendimiento de kernels CUDA con spilling de registros a la memoria compartida (CUDA 13.0)

Visión general

Cuando un kernel de CUDA requiere más registros de hardware de los que están disponibles, el compilador puede derramar (spill) los datos excedentes hacia la memoria local, ubicada en la memoria global del dispositivo. Este spilling puede afectar el rendimiento debido a la mayor latencia y ancho de banda de las accesos a la memoria local. CUDA 13.0 introduce una optimización: spilling de registros hacia la memoria compartida para kernels CUDA. El compilador da prioridad al spilling en la memoria compartida disponible y recurre a la memoria local si es necesario. Esta memoria on‑chip de baja latencia acerca los datos derramados a la SM, reduciendo la presión sobre la caché L2 y mejorando el rendimiento en rutas con alta presión de registros. En CUDA 13.0, PTXAS añade soporte para derramar registros a la memoria compartida en kernels CUDA. Si la característica está habilitada, el compilador intenta derramar primero en la memoria compartida y solo recurre a la memoria local cuando no hay suficiente espacio, preservando la corrección del programa. Esta optimización complementa tamaños de caché L1/L2 grandes y apunta a reducir la latencia de datos en la memoria on‑chip. Anteriormente, los spills se dirigían a la memoria local; este nuevo camino puede generar mejoras significativas en kernels con alta presión de registros, como bucles o secciones ejecutadas con frecuencia. La optimización fue evaluada en kernels CUDA de la biblioteca QUDA, utilizada para cálculos de lattice QCD en GPUs, y mostró mejoras típicas alrededor del 5–10%, atribuidas a la reducción o eliminación de spills hacia la memoria local. La optimización está disponible solo con CUDA 13.0 o posterior y requiere activar explícitamente mediante ensamblaje inline.

Lo que cambia en el compilador

Con CUDA 13.0, PTXAS introduce un mecanismo de activación que redirige spills hacia la memoria compartida cuando hay espacio para ellos. Cuando está activada, la memoria compartida se usa en prioridad para los spills y la memoria local sólo se usa si la memoria compartida se agota. Un kernel compilado con esta función muestra, en salidas de Nsight Compute, el uso de memoria compartida por bloque (por ejemplo, 46080 bytes de smem). El objetivo es mantener los datos derramados cerca de la SM y reducir evicciones de la L2 en regiones críticas de rendimiento. Esta optimización se activa insertando una directiva PTX (pragma) enable_smem_spilling mediante ensamblaje inline dentro de la función, justo después de su declaración. La función es válida solo dentro del ámbito de una función. No debe usarse cuando no se especifican claramente los límites de lanzamiento, ya que la estimación de bloques por SM puede volverse imprecisa, potencialmente limitando la concurrencia. Cuando los límites de lanzamiento están bien definidos y el uso de memoria compartida es estable, habilitar spilling a la memoria compartida puede mejorar el rendimiento.

Características clave

• Soporte PTXAS para derramar registros a la memoria compartida (CUDA 13.0+).
• Spill prioritario hacia la memoria compartida; retorno a la memoria local si no hay espacio suficiente.
• Beneficios de rendimiento potenciales debido a la reducción de latencia y menor presión en la L2 en zonas críticas.
• Activación vía inline assembly PTX: enable_smem_spilling, colocado después de la declaración de la función.
• Válido solo dentro del alcance de una función; recomendado cuando los límites de lanzamiento están claramente definidos.
• Demostrado en cargas como QUDA lattice QCD.
• Disponible solo en CUDA 13.0 y versiones posteriores; toolkits anteriores no soportan spilling en la memoria compartida.
• Si la memoria compartida por bloque suele estar infrautilizada, la optimización puede ser beneficiosa cuando la presión de registros es alta.

Casos de uso comunes

• Kernels con alta presión de registros y bucles ajustados, donde spills a memoria local eran un cuello de botella.
• Cargas en las que gran parte de la memoria compartida por bloque queda sin usar, haciendo factible almacenar spills en la memoria on‑chip.
• Escenarios con ocupación definida por límites de lanzamiento y uso predecible de memoria compartida que toleran el aumento del footprint por bloque.

Setup & instalación (comandos exactos)

• Asegúrese de usar CUDA 13.0 o posterior.
• Active spilling de memoria compartida insertando una pragma PTX en ensamblaje inline después de la declaración de la función: enable_smem_spilling. Esta activación se describe como válida dentro del alcance de una función y debe usarse cuando los límites de lanzamiento estén claramente definidos. Nota: la fuente describe el método de activación y sus restricciones, pero no proporciona comandos externos de instalación aparte de la versión requerida del toolkit CUDA.

// Activar spilling en memoria compartida (CUDA 13.0+)
// Colocar después de la declaración de la función
// Inline PTX pragma: enable_smem_spilling

Quick start (ejemplo mínimo ejecutable)

A continuación se muestran dos esqueletos de kernel mínimos para ilustrar la idea conceptualmente. El primero muestra un kernel simple sin la activación de spilling. El segundo indica dónde y cómo se aplicaría la activación de spilling en el código, de acuerdo con la guía de CUDA 13.0.

// Kernel mínimo (base, sin spilling)
extern "C" __global__ void demo_kernel(float* in, float* out) {
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float v = in[idx];
out[idx] = v * 2.0f;
}

// Kernel mínimo (conceptual: enable_smem_spilling después de la declaración de la función)
extern "C" __global__ void demo_kernel(float* in, float* out) {
// Activar spilling en memoria compartida (CUDA 13.0+)
// Inline PTX pragma: enable_smem_spilling
int idx = threadIdx.x + blockIdx.x * blockDim.x;
float v = in[idx];
out[idx] = v * 2.0f;
}

Nota: la sintaxis exacta para activar mediante inline PTX se describe en la fuente como una pragma PTX colocada tras la declaración de la función. Los ejemplos anteriores ilustran la ubicación y la intención, no una directiva de compilador literal.

Ventajas y desventajas

• Ventajas
• Ganancias de rendimiento potenciales (típicamente 5–10%) en kernels con alta presión de registros, al reducir spills hacia la memoria local.
• Mantiene los spills dentro de la memoria compartida on‑chip, reduciendo latencia y presión en la L2 cuando hay espacio disponible.
• Proporciona un camino de optimización específico para kernels con límites de lanzamiento definidos y uso estable de memoria compartida.
• Desventajas
• Requiere activación explícita; no es automática para todos los kernels.
• No es aplicable si la memoria compartida por bloque ya está saturada o si la presión de registros no es el cuello de botella.
• Puede afectar la ocupación si los límites de lanzamiento no están bien definidos, debido a estimaciones de bloques por SM.
• Disponible solo en CUDA 13.0 y posteriores; los toolkits más antiguos no soportan spills en la memoria compartida.

Alternativas (resumen)

| Enfoque | Pros | Contras |---|---|---| | Spill de registros a la memoria compartida (CUDA 13.0+) | Reduce spills a memoria local; aprovecha memoria on‑chip | Requiere activación explícita; depende de la disponibilidad de memoria compartida |Spill hacia memoria local (comportamiento previo a CUDA 13.0) | Sin necesidad de cambios de código | Latencia mayor; mayor presión en L2 |Optimización del código para reducir la presión de registros | Aplicable a más casos; menor dependencia del hardware | Puede requerir reestructuración significativa |

Pricing o Licencia

La fuente no proporciona información sobre precios o licencias. La funcionalidad se describe como parte de CUDA 13.0.

Referencias

• NVIDIA Dev Blog: How to Improve CUDA Kernel Performance with Shared Memory Register Spilling