Nossa postagem anterior seguiu um kernel de soma de vetores do nvcc até os warps. Desta vez, seguimos uma instrução SASS crítica (um carregamento global) pelo hardware de uma RTX4090. O kernel CUDA investigado tem duas linhas: __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];}. O SASS compilado mostra: /*0080*/ IMAD.WIDE R4, R6, R7, c[0x0][0x168] ; // &b[i] /*00a0*/ LDG.E R4, [R4.64] ; // b[i]. Essas instruções carregam elementos do vetor b da memória global para os registradores.
Uma LDG.E solicita quatro bytes em cada um dos 32 lanes. Atendê-la requer quatro setores de 32 bytes, uma linha de cache, uma tradução de endereço, uma travessia do crossbar, uma das trinta e seis fatias L2 e, quando falha em todos os lugares, uma ativação e quatro leituras de coluna em um chip DRAM. Nosso warp vive em uma das quatro subpartições do SM junto com outros onze warps residentes. A instrução é emitida para a unidade de carga/armazenamento (LSU), que envia o opcode, a máscara de lanes ativas, os endereços calculados e o número do registrador para o coalescer.
O coalescer determina que são necessárias 4 solicitações de setor contíguas para os 128 bytes solicitados pelo warp. Ao entrar no cache L1, a solicitação é enviada ao cache endereçado virtualmente, associativo por conjuntos de 4 vias. A busca compara tags em conjuntos de quatro slots. Como este é um carregamento pela primeira vez, ele falha e desce mais. Um acerto no L1 retorna em cerca de 15,4 ns — 40 ciclos. A memória virtual requer tradução antes de sair do SM. A TLB do SM contém dezesseis traduções recentes. As falhas custam cerca de 4,4 ns — onze ciclos.
Após a tradução, uma solicitação por linha de 128 bytes segue pelo crossbar até uma das 36 fatias L2 de 2 MiB, escolhida por uma função de seu endereço físico. Todas as fatias podem atender em paralelo, dando uma largura de banda agregada 36 vezes a de uma única fatia.
Fonte: Hacker News · Resumido por HeadlinesBriefing