💻 computer science

Non-Uniform L2 Cache Latency Across the Streaming Multiprocessors of an NVIDIA L40

Este artículo revela que la latencia de acierto de la caché L2 en las GPU NVIDIA L40 y Blackwell es no uniforme y depende altamente del multiprocesador de flujo (SM) físico específico que emite la carga, permitiendo un primitivo estable a nivel de usuario para la autolocalización de kernels, la huella digital del dispositivo y la distribución de trabajo optimizada que reduce el tiempo de ejecución hasta en un 11%.

Autores originales: Faruk Alpay, Baris Basaran

Publicado 2026-06-23

📖 5 min de lectura🧠 Análisis profundo

CC BY 4.0

Autores originales: Faruk Alpay, Baris Basaran

Artículo original bajo licencia CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Esta es una explicación generada por IA del artículo a continuación. No ha sido escrita ni avalada por los autores. Para mayor precisión técnica, consulte el artículo original. Leer descargo de responsabilidad completo

Imagina una biblioteca masiva y de alta velocidad (la GPU) donde miles de bibliotecarios (llamados SMs, o multiprocesadores de flujo) trabajan juntos para buscar libros (datos) en una enorme sala de almacenamiento compartida (el caché L2).

Durante años, los científicos de la computación creyeron que esta biblioteca funcionaba como un sistema perfectamente uniforme: sin importar a qué bibliotecario le preguntaras, el tiempo que tardaba en recoger un libro de la sala de almacenamiento era exactamente el mismo. Pensaban que la sala de almacenamiento era un fondo único y plano de información donde la distancia no importaba.

Este artículo, sin embargo, revela que esta creencia es errónea. Los autores descubrieron que en la GPU NVIDIA L40, el lugar donde se encuentra físicamente un bibliotecario determina qué tan rápido puede obtener un libro.

Aquí está el desglose de sus hallazgos utilizando analogías simples:

1. El descubrimiento de que "la distancia importa"

Imagina que la sala de almacenamiento no es un suelo plano, sino un gran almacén con pasillos.

La visión antigua: Todos pensaban que todos los bibliotecarios estaban situados exactamente a la misma distancia de los estantes. Si le preguntabas al Bibliotecario A o al Bibliotecario Z, el tiempo de espera era idéntico (unos 279 "ticks" del reloj).
La nueva realidad: Los autores midieron el tiempo de espera de cada uno de los 142 bibliotecarios. Descubrieron que algunos bibliotecarios están parados justo al lado de los estantes y obtienen los libros en 222 ticks. Otros están en el extremo lejano del almacén y tardan 339 ticks.
El resultado: Esa es una diferencia del 52% en velocidad. Es como si una persona corriera hacia la puerta principal mientras otra tiene que correr hacia la salida trasera solo para conseguir el mismo paquete.

2. El patrón de "imagen especular"

Cuando los autores mapearon quién era rápido y quién era lento, no vieron ruido aleatorio. Vieron un patrón perfecto.

Los 142 bibliotecarios se dividen en dos mitades idénticas (como dos alas de un edificio).
El Bibliotecario #1 en la primera ala tiene exactamente el mismo perfil de velocidad que el Bibliotecario #1 en la segunda ala.
Esta "simetría de espejo" coincide con el plano físico real del chip de la computadora, lo que demuestra que esto no es un error en el software, sino un hecho de la construcción física del hardware.

3. El efecto de la "huella digital"

Debido a que cada bibliotecario tiene una velocidad única basada en su ubicación física, la biblioteca tiene una identidad secreta.

Autolocalización: Si eres un bibliotecario (un programa de computadora) y preguntas: "¿Cuánto tiempo me toma obtener un libro?", puedes saber instantáneamente en qué lugar específico del almacén estás parado. Los autores construyeron una herramienta que puede identificar tu lugar específico con una preccesión del 99%.
Huella digital del dispositivo: Incluso si tienes dos GPUs L40 nuevas e idénticas (dos bibliotecas idénticas), son ligeramente diferentes. Una biblioteca puede tener al Bibliotecario #5 parado un poco más cerca de los estantes que el Bibliotecario #5 de la otra biblioteca debido a pequeñas diferencias de fabricación. Los autores pudieron distinguir las dos máquinas idénticas el 100% de las veces simplemente midiendo estas diminutas diferencias de velocidad. Es como poder distinguir a dos gemelos idénticos por la forma en que caminan.

4. ¿Es esto un riesgo de seguridad?

Los autores son cuidadosos al aclarar lo que esto significa para la seguridad.

Lo que SÍ ES: Una forma de que un programa sepa dónde está dentro de la computadora. Es como una persona que entra en una habitación y se da cuenta de: "Ah, estoy parado en la esquina cerca de la ventana".
Lo que NO ES: No es una forma de robar secretos de otros programas. Los autores enfatizan que esto solo mide la propia velocidad del programa. No puede espiar lo que otros programas están haciendo ni robar sus datos. Es una herramienta de "autolocalización", no una herramienta de "espionaje".

5. El beneficio práctico: Programación más inteligente

¿Por qué es esto importante para el rendimiento?

Imagina que tienes una lista de 100 tareas por hacer.
La forma antigua: Asignas las tareas al azar. Algunas van a los bibliotecarios que están lejos (lentos), y otras a los que están cerca (rápidos). Todo el trabajo espera a que las personas más lentas terminen.
La nueva forma: Ahora que tenemos el mapa, podemos asignar el trabajo pesado a los bibliotecarios que están más cerca de los estantes.
El resultado: Al hacer esto, los autores demostraron que podían terminar el trabajo un 11% más rápido para tareas que dependen fuertemente del caché. Sin embargo, si la tarea requiere buscar datos en la memoria principal (un almacenamiento diferente y más lento), este truco no ayuda.

6. No es solo un chip

Los autores también probaron esto en un chip más nuevo y diferente (el RTX 5090). Encontraron que la misma regla de "la distancia importa" se aplica también allí, aunque el patrón es ligeramente diferente. Esto demuestra que la vieja idea de un caché "uniforme" es probablemente errónea para muchas computadoras de alto rendimiento modernas.

Resumen

El artículo rompe la ilusión de que todas las partes del caché de una GPU son iguales. Revela que la ubicación física dicta la velocidad. Al mapear estas diferencias de velocidad ocultas, podemos:

Identificar exactamente dónde se está ejecutando un programa.
Distinguir entre dos máquinas idénticas.
Acelerar tareas específicas asignándolas a los lugares físicos más rápidos.

Resulta que el chip de la computadora no es un campo plano y uniforme; es un paisaje con colinas y valles, y conocer el mapa te hace más rápido.

Resumen Técnico: Latencia de Caché L2 No Uniforme en NVIDIA L40

Planteamiento del Problema
La GPU NVIDIA L40, basada en la implementación completa de AD102, expone una caché L2 de 96 MiB que convencionalmente se modela en herramientas de rendimiento y literatura académica como un pool uniforme con una latencia de acierto (hit latency) única para todo el dispositivo. Esta abstracción asume que cualquier multiprocesador de flujo (SM) que acceda a la caché L2 incurre en el mismo costo. Sin embargo, la realidad física de la caché L2 involucra múltiples segmentos (slices) conectados a controladores de memoria, accedidos a través de una red en chip (NoC). Si bien trabajos previos en las arquitecturas V100, A100 y H100 establecieron que la latencia del NoC es no uniforme, esos estudios enfatizaron que la latencia promedio por núcleo era aproximadamente uniforme, con la no uniformidad apareciendo principalmente como pasos discretos a través de los límites de partición. Seguía siendo una cuestión abierta si una GPU de centro de datos de un solo die y una sola partición (como la L40 de la arquitectura Ada Lovelace) exhibe una estructura comparable y, de ser así, si la latencia promedio por SM es uniforme o dependiente de la ubicación.

Metodología
Los autores desarrollaron una metodología de microbenchmarking a nivel de usuario para medir la latencia de acierto en L2 desde cada SM físico sin requerir privilegios elevados o contadores de rendimiento de hardware.

Diseño de la Sonda: La medición utiliza una cadena de carga dependiente de un solo hilo y un solo warp (ld.global.cg) para asegurar un único viaje de ida y vuelta a través del tejido de la caché. Se emplearon dos formas de buffer: una forma de "línea única" para sondear segmentos específicos de 128 B y una forma de "cadena" (permutación aleatoria de 2 MiB) para medir la distancia promedio a todos los segmentos.
Aislamiento de SM: Para eliminar la contención en el chip, los bloques se serializan mediante un contador de turnos global. Cada bloque reclama un turno mediante atomicAdd, espera, realiza un calentamiento (warm-up) y mide sus cargas dependientes de forma aislada.
Cobertura: La sonda cubre los 142 SM físicos de la L40 y 256 sondas de segmento. Una sonda similar se ejecutó en una NVIDIA RTX 5090 (Blackwell GB202) para probar la generalización.
Modelado: Los datos de latencia se analizan utilizando un modelo de acceso a caché no uniforme (NUCA) aditivo: $L = \mu + a(sm) + b(slice)$ , donde $a(sm)$ representa el término de ubicación del SM y $b(slice)$ representa el término de ubicación del segmento.

Resultados Clave

Variación Significativa de la Latencia: Contrario al modelo uniforme, la latencia de acierto en L2 en la L40 no es constante. Oscila entre 222.5 y 339.2 ciclos (un rango del 52%), con una variación media de 57.2 ciclos entre los SM. El ruido de fondo es insignificante (<0.01 ciclos), confirmando que la variación es física.
Simetría de Doble Vía: El término de ubicación del SM exhibe una fuerte simetría de doble vía. Dividir los 142 SM en el identificador 72 revela dos mitades (72 y 70 SM) con perfiles de latencia casi idénticos (correlación $r=0.999$ ). Esto se alinea con el diseño de piso (floorplan) de AD102 de dos clústeres GPC simétricos.
Gradiente Continuo vs. Pasos Discretos: A diferencia de la A100/H100, que muestran un paso de latencia discreto en el límite de partición, la L40 exhibe un gradiente continuo de latencia a través de los SM sin un hueco bimodal.
Ajuste del Modelo Aditivo: El modelo aditivo explica el 87% de la variación de la latencia ( $R^2=0.87$ ). Añadir un único término de interacción de rango 1 aumenta esto a 0.98, lo que indica que el mapa de latencia es una geometría de baja dimensión que sigue el diseño de los GPC.
Generalización y Especificidad:
- Trans-Arquitectura: La RTX 5090 también muestra no uniformidad (dispersión del 46%), pero con una simetría más débil ( $r=0.80$ ) y una latencia absoluta más lenta.
- Huella Digital del Dispositivo: Dos L40 físicamente distintas e idénticas en modelo son distinguibles con un 100% de precisión basándose en sus mapas por SM, a pesar de tener latencias medias casi idénticas (diferencia de 0.28 ciclos). El mapa por SM actúa como una identidad de hardware estable.
Estabilidad: El mapa de latencia por SM permanece sin cambios tras una hora de utilización al 100% en ambos dispositivos, confirmando que la topología es una propiedad del silicio y no un estado transitorio.

Contribuciones y Significado
El artículo reclama las siguientes contribuciones:

Mapeo de Topología: Proporciona el primer mapa de latencia L2 por SM para la L40, revelando una dispersión del 52% que invalida la suposición de caché uniforme para cargas de trabajo limitadas por latencia.
Primitiva de Autolocalización: Los autores demuestran que una sola sonda de latencia L2 permite a un kernel identificar su propia ubicación física con un 99.2% de precisión (o 92% en un conjunto de dispositivos mixtos) e identificar el die específico del dispositivo con un 100% de precisión. Esto se plantea como una primitiva de "autolocalización", no como un ataque a un kernel víctima, ya que no extrae datos secretos y depende de la propia ubicación del kernel.
Implicaciones de Rendimiento: Al distribuir el trabajo limitado por latencia de acuerdo con el mapa de cada SM (colocando el trabajo en los SM más rápidos), los autores lograron una reducción del 10.9% en el tiempo de ejecución (makespan) para cargas de trabajo residentes en L2. Este beneficio desaparece en regímenes limitados por DRAM, resaltando la naturaleza dependiente del régimen del impacto de la topología.
Reinterpretación de la Residencia: El artículo reinterpreta los controles de residencia de un solo hilo existentes (límites de capacidad, escalado de etiquetas de línea y resultados de persistencia en L2) como promedios sobre esta distribución no uniforme, en lugar de constantes.

Alcance y Limitaciones
Los autores limitan explícitamente el alcance de sus afirmaciones de seguridad. La técnica permite la autolocalización (un kernel leyendo su propia ubicación) y la huella digital del dispositivo. No infiere la ubicación de un "kernel víctima" co-residente, ni extrae datos secretos; tales capacidades requerirían un canal de contención o interferencia no medido en este estudio. El estudio se limita a dos L40 y una RTX 5090, y el mapeo de %smid al diseño físico se infiere de la simetría y el recuento de GPC en lugar de ser confirmado independientemente mediante fotografía de die. Los conteos de ciclos absolutos son específicos del diseño del microbenchmark y no se traducen directamente al rendimiento de las aplicaciones donde existen solicitudes independientes y uso de L1.