💻 computer science

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

Este artigo revela que a latência de acerto (hit latency) do cache L2 nas GPUs NVIDIA L40 e Blackwell é não uniforme e altamente dependente do Streaming Multiprocessor (SM) físico específico que emite a carga, permitindo um primitivo estável de nível de usuário para autolocalização de kernel, fingerprinting de dispositivo e distribuição de trabalho otimizada que reduz o makespan em até 11%.

Autores originais: Faruk Alpay, Baris Basaran

Publicado 2026-06-23

📖 5 min de leitura🧠 Leitura aprofundada

CC BY 4.0

Autores originais: Faruk Alpay, Baris Basaran

Artigo original sob licença CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Esta é uma explicação gerada por IA do artigo abaixo. Não foi escrita nem endossada pelos autores. Para precisão técnica, consulte o artigo original. Ler aviso legal completo

Imagine uma biblioteca massiva e de alta velocidade (a GPU) onde milhares de bibliotecários (chamados de SMs, ou Streaming Multiprocessors) trabalham juntos para buscar livros (dados) de uma enorme sala de armazenamento compartilhada (o Cache L2).

Durante anos, cientistas da computação acreditaram que esta biblioteca funcionava como um sistema perfeitamente uniforme: não importava qual bibliotecário você chamasse, o tempo que levava para pegar um livro no depósito era exatamente o mesmo. Eles pensavam que a sala de armazenamento era um único pool plano de informações, onde a distância não importava.

Este artigo, no entanto, revela que essa crença está errada. Os autores descobriram que, na GPU NVIDIA L40, onde um bibliotecário está posicionado fisicamente determina a rapidez com que ele consegue pegar um livro.

Aqui está a divisão de suas descobertas usando analogias simples:

1. A Descoberta de que "A Distância Importa"

Pense na sala de armazenamento não como um chão plano, mas como um grande armazém com corredores.

A Visão Antiga: Todos pensavam que cada bibliotecário estava exatamente à mesma distância das prateleiras. Se você pedisse ao Bibliotecário A ou ao Z, o tempo de espera era idêntico (cerca de 279 "ticks" do relógio).
A Nova Realidade: Os autores mediram o tempo de espera para cada um dos 142 bibliotecários. Eles descobriram que alguns bibliotecários estão parados bem ao lado das prateleiras e pegam os livros em 222 ticks. Outros estão no extremo oposto do armazém e levam 339 ticks.
O Resultado: Isso é uma diferença de 52% na velocidade. É como se uma pessoa corresse até a porta da frente enquanto outra tem que correr até a saída dos fundos apenas para pegar o mesmo pacote.

2. O Padrão de "Imagem Espelhada"

Quando os autores mapearam quem era rápido e quem era lento, eles não viram ruído aleatório. Eles viram um padrão perfeito.

Os 142 bibliotecários são divididos em duas metades idênticas (como duas alas de um edifício).
O Bibliotecário nº 1 na primeira ala tem exatamente o mesmo perfil de velocidade que o Bibliotecário nº 1 na segunda ala.
Esta "simetria de espelho" corresponde ao projeto físico real do chip do computador, provando que isso não é um erro no software — é um fato da construção física do hardware.

3. O Efeito "Impressão Digital"

Como cada bibliotecário tem uma velocidade única baseada em seu local físico, a biblioteca possui uma identidade secreta.

Autolocalização: Se você é um bibliotecário (um programa de computador) e pergunta: "Quanto tempo levo para pegar um livro?", você pode descobrir instantaneamente em qual lugar exato do armazém você está parado. Os autores construíram uma ferramenta que pode identificar seu local específico com 99% de precisão.
Impressão Digital do Dispositivo (Device Fingerprinting): Mesmo que você tenha duas GPUs L40 novas e idênticas (duas bibliotecas idênticas), elas são ligeiramente diferentes. Uma biblioteca pode ter o Bibliotecário nº 5 parado um pouco mais perto das prateleiras do que o Bibliotecário nº 5 da outra biblioteca, devido a minúsculas diferenças de fabricação. Os autores puderam distinguir as duas máquinas idênticas 100% das vezes apenas medindo essas pequenas diferenças de velocidade. É como ser capaz de distinguir dois gêmeos idênticos pela maneira como eles caminham.

4. Isso é um Risco de Segurança?

Os autores são cuidadosos ao esclarecer o que isso significa para a segurança.

O que É: Uma forma de um programa saber onde ele está dentro do computador. É como uma pessoa entrando em uma sala e percebendo: "Ah, estou parado no canto, perto da janela".
O que NÃO É: Não é uma forma de roubar segredos de outros programas. Os autores enfatizam que isso apenas mede a própria velocidade do programa. Não é possível espiar o que outros programas estão fazendo ou roubar seus dados. É uma ferramenta de "autolocalização", não uma ferramenta de "espionagem".

5. O Benefício Prático: Escalonamento Mais Inteligente

Por que isso importa para o desempenho?

Imagine que você tem uma lista de 100 tarefas para fazer.
O Jeito Antigo: Você atribui as tarefas aleatoriamente. Algumas vão para os bibliotecários que estão longe (lentos), e outras para aqueles que estão perto (rápidos). Todo o trabalho espera pelos mais lentos terminarem.
O Novo Jeito: Agora que temos o mapa, podemos atribuir o trabalho pesado aos bibliotecários que estão mais próximos das prateleiras.
O Resultado: Ao fazer isso, os autores mostraram que conseguiram concluir o trabalho 11% mais rápido para tarefas que dependem fortemente do cache. No entanto, se a tarefa exigir a busca de dados na memória principal (um armazenamento diferente e mais lento), este truque não ajuda.

6. Não é Apenas Um Chip

Os autores testaram isso em um chip mais novo e diferente (o RTX 5090) também. Eles descobriram que a mesma regra de "a distância importa" se aplica lá também, embora o padrão seja ligeiramente diferente. Isso prova que a antiga ideia de um cache "uniforme" provavelmente está errada para muitos computadores de alto desempenho modernos.

Resumo

O artigo destrói a ilusão de que todas as partes do cache de uma GPU são iguais. Ele revela que a localização física dita a velocidade. Ao mapear essas diferenças de velocidade ocultas, podemos:

Identificar exatamente onde um programa está sendo executado.
Distinguir entre duas máquinas idênticas.
Acelerar tarefas específicas ao atribuí-las aos locais físicos mais rápidos.

Acontece que o chip do computador não é um campo plano e uniforme; é uma paisagem com colinas e vales, e conhecer o mapa torna você mais rápido.

Resumo Técnico: Latência de Cache L2 Não Uniforme na NVIDIA L40

Definição do Problema
A GPU NVIDIA L40, baseada na implementação completa do AD102, expõe um cache L2 de 96 MiB que é convencionalmente modelado em ferramentas de desempenho e literatura acadêmica como um pool uniforme com uma latência de hit única para todo o dispositivo. Essa abstração assume que qualquer SM (Streaming Multiprocessor) que acesse o cache L2 incorre no mesmo custo. No entanto, a realidade física do cache L2 envolve múltiplas fatias (slices) conectadas a controladores de memória, acessadas via uma rede no chip (NoC). Embora trabalhos anteriores nas arquiteturas V100, A100 e H100 tenham estabelecido que a latência da NoC é não uniforme, esses estudos enfatizaram que a latência média por núcleo era aproximadamente uniforme, com a não uniformidade aparecendo principalmente como degraus discretos através das fronteiras de partição. Permanecia uma questão em aberto se uma GPU de datacenter de matriz única e partição única da arquitetura Ada Lovelace (como a L40) exibe uma estrutura comparável e, se sim, se a latência média por SM é uniforme ou dependente do posicionamento.

Metodologia
Os autores desenvolveram uma metodologia de microbenchmarking em nível de usuário para medir a latência de hit de L2 de cada SM físico sem a necessidade de privilégios elevados ou contadores de desempenho de hardware.

Design da Sonda: A medição utiliza uma cadeia de carga dependente de thread única e warp único (ld.global.cg) para garantir uma única viagem de ida e volta pelo fabric de cache. Dois formatos de buffer foram empregados: um formato de "linha única" para sondar fatias específicas de 128 B e um formato de "corrente" (permutação aleatória de 2 MiB) para medir a distância média para todas as fatias.
Isolamento de SM: Para eliminar a contenção no chip, os blocos são serializados usando um contador de turno global. Cada bloco reivindica um turno via atomicAdd, aguarda, realiza o aquecimento (warm-up) e cronometra suas cargas dependentes em isolamento.
Cobertura: A sonda cobre todos os 142 SMs físicos na L40 e 256 sondas de fatia. Uma sonda semelhante foi executada em uma NVIDIA RTX 5090 (Blackwell GB202) para testar a generalização.
Modelagem: Os dados de latência são analisados usando um modelo de Acesso a Cache Não Uniforme (NUCA) aditivo: $L = \mu + a(sm) + b(slice)$ , onde $a(sm)$ representa o termo de posicionamento do SM e $b(slice)$ representa o termo de posicionamento da fatia.

Resultos Principais

Variação Significativa de Latência: Contrário ao modelo uniforme, a latência de hit de L2 na L40 não é constante. Ela varia de 222,5 a 339,2 ciclos (uma amplitude de 52%), com uma variação média de 57,2 ciclos entre os SMs. O piso de ruído é insignificante (<0,01 ciclos), confirmando que a variação é física.
Simetria Dupla: O termo de posicionamento do SM exibe uma forte simetria de dois lados. A divisão dos 142 SMs no identificador 72 revela duas metades (72 e 70 SMs) com perfis de latência quase idênticos (correlação $r=0,999$ ). Isso se alinha ao floorplan do AD102 de dois clusters GPC simétricos.
Gradiente Contínuo vs. Degraus Discretos: Diferente da A100/H100, que mostram um degrau de latência discreto na fronteira da partição, a L40 exibe um gradiente contínuo de latência entre os SMs, sem um hiato bimodal.
Ajuste do Modelo Aditivo: O modelo aditivo explica 87% da variação de latência ( $R^2=0,87$ ). Adicionar um único termo de interação de rank-1 eleva isso para 0,98, indicando que o mapa de latência é uma geometria de baixa dimensão que segue o layout do GPC.
Generalização e Especificidade:
- Cross-Architecture: A RTX 5090 também mostra não uniformidade (dispersão de 46%), mas com uma simetria mais fraca ( $r=0,80$ ) e latência absoluta mais lenta.
- Fingerprinting de Dispositivo: Duas L40s fisicamente distintas e de modelo idêntico são distinguíveis com 100% de precisão com base em seus mapas por SM, apesar de terem latências médias quase idênticas (diferença de 0,28 ciclo). O mapa por SM atua como uma identidade de hardware estável.
Estabilidade: O mapa de latência por SM permanece inalterado após uma hora de utilização de 100% em ambos os dispositivos, confirmando que a topologia é uma propriedade do silício e não um estado transitório.

Contribuições e Significância
O artigo reivindica as seguintes contribuições:

Mapeamento de Topologia: Fornece o primeiro mapa de latência de L2 por SM para a L40, revelando uma dispersão de 52% que invalida a suposição de cache uniforme para cargas de trabalho limitadas por latência.
Primitiva de Autolocalização: Os autores demonstram que uma única sonda de latência de L2 permite que um kernel identifique seu próprio posicionamento físico com 99,2% de precisão (ou 92% em um conjunto de dispositivos mistos) e identifique o die específico do dispositivo com 100% de precisidade. Isso é enquadrado como uma primitiva de "autolocalização", não um ataque a um kernel vítima, pois não extrai dados secretos e depende do próprio posicionamento do kernel.
Implicações de Desempenho: Ao distribuir o trabalho limitado por latência de acordo com o mapa por SM (posicionando o trabalho em SMs mais rápidos), os autores alcançaram uma redução de 10,9% no makespan para cargas de trabalho residentes em L2. Esse benefício desaparece em regimes limitados por DRAM, destacando a natureza regime-específica do impacto da topologia.
Reinterpretação de Residência: O artigo reinterpreta os controles existentes de residência de thread única (fronteiras de capacidade, escalonamento de tag de linha e resultados de persistência de L2) como médias sobre esta distribuição não uniforme, em vez de constantes.

Escopo e Limitações
Os autores limitam explicitamente o escopo de suas reivindicações de segurança. A técnica permite a autolocalização (um kernel lendo seu próprio posicionamento) e o fingerprinting de dispositivo. Ela não infere o posicionamento de um kernel "vítima" co-residente separado, nem extrai dados secretos; tais capacidades exigiriam um canal de contenção ou interferência não medido neste estudo. O estudo é limitado a duas L40s e uma RTX 5090, e o mapeamento de %smid para o floorplan físico é inferido a partir da simetria e contagem de GPCs, em vez de ser confirmado independentemente por fotografia de die. As contagens de ciclos absolutos são específicas do design do microbenchmark e não se traduzem diretamente para o throughput da aplicação onde ocorrem requisições independentes e uso de L1.