Como Projetar Chips de GPU
Chapter 7 Streaming Multiprocessor Design

Capítulo 7: Design do Streaming Multiprocessador no Design de GPU

O streaming multiprocessador (SM) é o bloco fundamental de construção das arquiteturas de GPU da NVIDIA. Cada SM contém um conjunto de núcleos CUDA que executam instruções em um estilo SIMT (Single Instruction, Multiple Thread). O SM é responsável por gerenciar e agendar warps, lidar com divergência de ramificação e fornecer acesso rápido à memória compartilhada e caches. Neste capítulo, exploraremos a microarquitetura do SM, incluindo seus pipelines, mecanismos de agendamento de warps, design do arquivo de registros e organização da memória compartilhada e do cache L1.

Microarquitetura e Pipelines do SM

O SM é um processador altamente paralelo e encanado projetado para executar eficientemente centenas de threads concorrentemente. A Figura 7.1 mostra um diagrama de bloco simplificado de um SM na arquitetura NVIDIA Volta.

                                 Cache de Instruções
                                         |
                                         v
                                Agendador de Warps
                                         |
                                         v
                               Unidade de Despacho (4 warps)
                                 |   |   |   |
                                 v   v   v   v
                               Núcleo CUDA (FP64/FP32/INT)
                               Núcleo CUDA (FP64/FP32/INT)
                               Núcleo CUDA (FP64/FP32/INT)
                               ...
                               Tensor Core
                               Tensor Core
                               ...
                               Unidade de Carga/Armazenamento
                               Unidade de Carga/Armazenamento
                               ...
                               Unidade de Função Especial
                                         ^
                                         |
                                Arquivo de Registros (64 KB)
                                         ^
```Memória Compartilhada / Cache L1 (96 KB)

Figura 7.1: Diagrama de bloco simplificado de um SM na arquitetura NVIDIA Volta.

Os principais componentes do SM incluem:

  1. Instruction Cache: Armazena instruções acessadas com frequência para reduzir a latência e melhorar o throughput.

  2. Warp Scheduler: Seleciona warps prontos para execução e os despacha para as unidades de execução disponíveis.

  3. Dispatch Unit: Busca e decodifica instruções para até 4 warps por ciclo e as despacha para as unidades de execução apropriadas.

  4. CUDA Cores: Unidades de execução programáveis que suportam uma ampla gama de operações de inteiro e ponto flutuante. Cada SM na Volta contém 64 CUDA Cores.

  5. Tensor Cores: Unidades de execução especializadas projetadas para acelerar cargas de trabalho de aprendizado profundo e IA. Cada SM na Volta contém 8 Tensor Cores.

  6. Load/Store Units: Lidam com operações de memória, incluindo carregamentos e armazenamentos na memória global, memória compartilhada e caches.

  7. Special Function Units: Executam operações transcendentais e outras operações matemáticas complexas.

  8. Register File: Fornece acesso rápido a registradores privados do thread. Cada SM na Volta possui um registro de 64 KB.

  9. Memória Compartilhada / Cache L1: Um espaço de memória configurável que pode ser usado como um cache gerenciado por software (memória compartilhada) ou como um cache de dados L1 gerenciado por hardware.

O pipeline SM é projetado para maximizar o throughput, permitindo a execução concorrente de vários warps e ocultando a latência da memória. A Figura 7.2 ilustra uma visão simplificada do pipeline SM.

    Instruction Fetch
            |
            v
    Instruction Decode
            |
            v
    Operand Collection
            |
            v
    Execution (CUDA Cores, Tensor Cores, Load/Store Units, Special Function Units)
            |
            v
    Writeback

Figura 7.2: Pipeline SM simplificado.

As etapas do pipeline são as seguintes:

  1. Instruction Fetch: O warp scheduler seleciona um warp pronto para exe

1. **Busca de Instrução**: A próxima instrução é buscada do cache de instruções para essa warpe.

2. **Decodificação de Instrução**: A instrução buscada é decodificada para determinar o tipo de operação, operandos e registradores de destino.

3. **Coleta de Operandos**: Os operandos necessários para a instrução são coletados do arquivo de registradores ou da memória compartilhada.

4. **Execução**: A instrução é executada na unidade de execução apropriada (Núcleo CUDA, Núcleo de Tensor, Unidade de Carga/Armazenamento ou Unidade de Função Especial).

5. **Gravação de Resultado**: O resultado da execução é gravado de volta no arquivo de registradores ou na memória compartilhada.

Para alcançar um alto desempenho, o SM emprega várias técnicas para maximizar a utilização de recursos e ocultar a latência:

- **Emissão Dupla**: O SM pode emitir duas instruções independentes por warp em um único ciclo, permitindo um maior paralelismo em nível de instrução.
- **Unidades de Execução Pipeline**: As unidades de execução são pipelineadas, permitindo que o SM inicie uma nova operação em uma unidade antes que a operação anterior tenha sido concluída.
- **Ocultação de Latência**: O SM pode alternar entre warps em uma base de ciclo a ciclo, permitindo ocultar a latência de acessos à memória e operações de longa duração, executando instruções de outras warps.

O Exemplo 7.1 mostra um kernel CUDA simples que realiza a adição elemento a elemento de dois vetores.

```cpp
__global__ void vectorAdd(int *a, int *b, int *c, int n) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (tid < n) {
        c[tid] = a[tid] + b[tid];
    }
}

Exemplo 7.1: Kernel CUDA para adição de vetores.

Neste exemplo, cada thread no kernel calcula a soma dos elementos correspondentes dos vetores de entrada a e b e armazena o resultado no vetor de saída c. O SM executa esse kernel atribuindo cada thread a um Núcleo CUDA e agendando warps de threads para serem executadas nos Núcleos disponíveis. As unidades de carga/armazenamento são usadas para buscar os dados de entrada da memória global e escrever os resultados de volta.

Escalonamento de Warps e Tratamento de Divergência

EfAqui está a tradução em português do arquivo Markdown, com os comentários traduzidos e o código não traduzido:

Uma programação eficiente de warps é crucial para maximizar o desempenho do SM. O escalonador de warps é responsável por selecionar os warps prontos para execução e enviá-los para as unidades de execução disponíveis. O objetivo principal do escalonador de warps é manter as unidades de execução ocupadas, garantindo que haja sempre warps disponíveis para executar.

O SM emprega um mecanismo de escalonamento de warps em dois níveis:

  1. Escalonamento de Warps: O escalonador de warps seleciona os warps prontos para execução com base em uma política de escalonamento, como round-robin ou primeiro-a-chegar-primeiro-a-ser-atendido. Os warps selecionados são então enviados para as unidades de execução disponíveis.

  2. Escalonamento de Instruções: Dentro de cada warp, o SM programa as instruções com base em suas dependências e na disponibilidade das unidades de execução. O SM pode emitir várias instruções independentes do mesmo warp em um único ciclo para maximizar o paralelismo em nível de instrução.

A Figura 7.3 ilustra o mecanismo de escalonamento de warps em dois níveis.

    Warp Pool
    Warp 1 (Pronto)
    Warp 2 (Esperando)
    Warp 3 (Pronto)
    ...
    Warp N (Pronto)
        |
        v
    Warp Scheduler
        |
        v
    Dispatch Unit
        |
        v
    Execution Units

Figura 7.3: Mecanismo de escalonamento de warps em dois níveis.

Um dos principais desafios no escalonamento de warps é lidar com a divergência de ramificações. No modelo de execução SIMT, todos os threads em um warp executam a mesma instrução em sincronia. No entanto, quando um warp encontra uma instrução de ramificação (por exemplo, uma declaração if-else), alguns threads podem seguir o caminho do if, enquanto outros seguem o caminho do else. Essa situação é chamada de divergência de ramificação.

Para lidar com a divergência de ramificação, o SM emprega uma técnica chamada predicação. Quando um warp encontra uma ramificação divergente, o SM executa ambos os caminhos da ramificação sequencialmente, mascarando os threads que não seguem cada caminho. Os resultados são então combinados usando registradores de predição para garantir que cada thread receba o resultado correto.

O Exemplo 7.2 mostra um kernel CUDA com uma ramificação divergente.Aqui está a tradução em português do arquivo markdown, com a observação de não traduzir o código, apenas os comentários:

__global__ void divergentKernel(int *data, int *result) {
    int tid = blockIdx.x * blockDim.x + threadIdx.x;
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }
}

Exemplo 7.2: Kernel CUDA com um ramo divergente.

Neste exemplo, a condição do ramo data[tid] > 0 pode fazer com que alguns threads em uma warp sigam o caminho do if, enquanto outros seguem o caminho do else. O SM lida com essa divergência executando ambos os caminhos sequencialmente e mascarando os threads inativos em cada caminho.

A Figura 7.4 ilustra o processo de predição para uma warp com threads divergentes.

    Warp (32 threads)
    Thread 1: data[1] = 5, result[1] = 10
    Thread 2: data[2] = -3, result[2] = -9
    ...
    Thread 32: data[32] = 7, result[32] = 14

    Ramo Divergente:
    if (data[tid] > 0) {
        result[tid] = data[tid] * 2;
    } else {
        result[tid] = data[tid] * 3;
    }

    Predição:
    Passo 1: Executar o caminho do if com máscara
        Thread 1: result[1] = 10
        Thread 2: (mascarado)
        ...
        Thread 32: result[32] = 14

    Passo 2: Executar o caminho do else com máscara
        Thread 1: (mascarado)
        Thread 2: result[2] = -9
        ...
        Thread 32: (mascarado)

    Resultado Final:
    Thread 1: result[1] = 10
    Thread 2: result[2] = -9
    ...
    Thread 32: result[32] = 14

Figura 7.4: Processo de predição para uma warp com threads divergentes.

Usando a predição, o SM pode lidar com a divergência de ramos sem a necessidade de instruções de ramificação explícitas ou divergência de fluxo de controle. No entanto, ramos divergentes ainda podem impactar o desempenho, pois o SM deve executar ambos os caminhos sequencialmente, reduzindo o paralelismo efetivo.

Arquivo de Registros e Coletores de Operandos

O arquivo de registros é um componente crítico do SM, fornecendo acesso rápido a registros privados de thread. Cada SM tem um grande arquivo de registros para suportar os muitos threads ativos e permitir uma troca de contexto eficiente entre warps.Aqui está a tradução em português do arquivo Markdown, com o código mantido no original:

Na arquitetura NVIDIA Volta, cada SM (Streaming Multiprocessor) possui um arquivo de registros de 64 KB, organizado em 32 bancos de 2 KB cada. O arquivo de registros é projetado para fornecer alto largura de banda e baixa latência de acesso para suportar o grande número de threads concorrentes.

Para minimizar os conflitos de banco e melhorar o desempenho, o SM emprega uma técnica chamada de coleta de operandos. Os coletores de operandos são unidades especializadas que coletam os operandos dos bancos do arquivo de registros e os entregam às unidades de execução. Ao usar os coletores de operandos, o SM pode reduzir o impacto dos conflitos de banco e melhorar a utilização das unidades de execução.

A Figura 7.5 mostra um diagrama simplificado do arquivo de registros e dos coletores de operandos em um SM.

    Arquivo de Registros (64 KB)
    Banco 1 (2 KB)
    Banco 2 (2 KB)
    ...
    Banco 32 (2 KB)
        |
        v
    Coletores de Operandos
        |
        v
    Unidades de Execução

Figura 7.5: Arquivo de registros e coletores de operandos em um SM.

Os coletores de operandos funcionam coletando operandos de múltiplas instruções e múltiplas warps, permitindo que o SM emita instruções de diferentes warps para as unidades de execução em um único ciclo. Isso ajuda a esconder a latência dos acessos ao arquivo de registros e melhora o desempenho geral do SM.

O Exemplo 7.3 mostra um kernel CUDA que realiza um produto escalar de dois vetores.

__global__ void dotProduct(float *a, float *b, float *result, int n) {
    // Cada thread calcula uma soma parcial do produto escalar usando seu índice atribuído
    __shared__ float partialSum[256];
    int tid = threadIdx.x;
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    partialSum[tid] = 0;
 
    while (i < n) {
        partialSum[tid] += a[i] * b[i];
        i += blockDim.x * gridDim.x;
    }
 
    __syncthreads();
 
    for (int s = blockDim.x / 2; s > 0; s >>= 1) {
        if (tid < s) {
            partialSum[tid] += partialSum[tid + s];
        }
        __syncthreads();
    }
 
    if (tid == 0) {
        result[blockIdx.x] = partialSum[0];
    }
}

Neste exemplo, cada thread calcula uma soma parcial do produto escalar usando seu índice atribuído.Aqui está a tradução em português desse arquivo Markdown. Para o código, não traduzi o código, apenas os comentários.

Elementos dos vetores de entrada. As somas parciais são armazenadas no array de memória compartilhada partialSum. Depois que todos os threads calcularam suas somas parciais, é realizada uma redução paralela para somar as somas parciais e obter o resultado final do produto escalar.

O coletor de operandos desempenha um papel crucial neste exemplo, coletando eficientemente os operandos para os acessos à memória compartilhada e as operações aritméticas. Ele ajuda a evitar conflitos de banco e melhora a utilização das unidades de execução.

Conclusão

O multiprocessador de streaming é a unidade computacional principal nas arquiteturas modernas de GPU. Seu design se concentra em maximizar o throughput e esconder a latência de memória por meio de uma combinação de multithreading de grão fino, execução SIMT e coleta eficiente de operandos.

Componentes-chave do SM incluem o planejador de warp, que seleciona warps para execução; a pilha SIMT, que lida com divergência e convergência de ramificação; o arquivo de registros e os coletores de operandos, que fornecem acesso rápido a registradores privados de thread; e a memória compartilhada e o cache L1, que permitem o compartilhamento e a reutilização de dados de baixa latência.

À medida que as arquiteturas de GPU continuam a evoluir, pesquisas em áreas como o tratamento de divergência de ramificação, o escalonamento de warp e o design do arquivo de registros serão cruciais para melhorar o desempenho e a eficiência de futuras GPUs. Técnicas inovadoras, como formação dinâmica de warp, compactação de blocos de thread e caches de reutilização de operandos, têm o potencial de melhorar significativamente as capacidades do SM e permitir novos níveis de desempenho em cargas de trabalho de computação paralela.

Chapter 6 Gpu Performance Metrics and AnalysisChapter 8 Interconnect and on Chip Networks