Análise Microarquitetural da SpMV (CSR) em Nós Heterogêneos: Escalabilidade e Gargalos de Memória entre CPU e GPU

1. Introdução

O principal obstáculo para o desempenho ótimo da SpMV está na natureza esparsa dos dados. Para viabilizar o armazenamento e evitar o desperdício de memória com zeros estruturais, matrizes esparsas são compactadas em estruturas como o formato Compressed Sparse Row (CSR). Embora altamente eficiente no consumo de espaço, o formato CSR obriga a ocorrência de acessos indiretos e irregulares à memória durante o cálculo. Essa irregularidade destrói a previsibilidade espacial dos dados, tornando ineficazes os mecanismos de pré-busca (prefetching) dos processadores. Consequentemente, o algoritmo sofre com uma baixíssima intensidade aritmética, tornando-se severamente limitado pela largura de banda da memória principal (o gargalo conhecido como Memory Wall), em vez de ser limitado pela capacidade de processamento aritmético da arquitetura [Vuduc et al. 2002, Williams et al. 2009].

Para mitigar esse impedimento físico e acelerar a resolução computacional, a adoção de estratégias de paralelização em sistemas de memória compartilhada tornou-se indispensável. O padrão OpenMP consolidou-se como uma ferramenta robusta para explorar o paralelismo em arquiteturas multicore [Dagum and Menon 1998]. No entanto, a distribuição eficiente da carga de trabalho na SpMV não é trivial: como as linhas de uma matriz esparsa possuem concentrações discrepantes de elementos não-nulos, divisões de trabalho estáticas e lineares frequentemente geram um desbalanceamento de carga severo (load imbalance), mantendo núcleos de processamento ociosos.

Diante deste cenário, este trabalho tem como objetivo implementar, avaliar e contrastar o desempenho da operação SpMV utilizando uma linha de base sequencial frente a diferentes estratégias de paralelização via OpenMP (incluindo escalonamento Estático, Guiado e particionamento customizado por elementos Não-Nulos) em um nó de cluster CPU de 24 núcleos. Adicionalmente, o estudo explora o descarregamento de processamento (offloading) para aceleradores gráficos (GPU), avaliando as limitações impostas pelo barramento PCI-Express e pelos acessos desalinhados de memória. A análise de desempenho é balizada por métricas de vazão (throughput), pelo ganho de aceleração (SpeedUp) [Amdahl 1967] e estruturada sob os limites teóricos do modelo Roofline [Williams et al. 2009].

2. Referencial Teórico

2.1 A Multiplicação Matriz-Vetor Esparsa (SpMV)

A Multiplicação Matriz-Vetor Esparsa (SpMV) é definida pela operação \(y = Ax\), onde \(A\) é uma matriz esparsa de dimensões \(M \times N\) contendo uma proporção massiva de elementos nulos, e \(x\) e \(y\) são vetores densos. Matrizes esparsas surgem nativamente em simulações numéricas, métodos de elementos finitos e na resolução de grandes sistemas lineares. Devido à sua baixa intensidade aritmética — visto que realiza estritamente duas operações de ponto flutuante (uma multiplicação e uma adição) por elemento não-nulo —, a SpMV opera em um regime severamente limitado pela largura de banda da memória principal [Vuduc et al. 2002, Williams et al. 2009].

2.2 O Formato Compressed Sparse Row (CSR)

Para viabilizar o armazenamento de matrizes esparsas e evitar o desperdício de espaço com zeros estruturais, utiliza-se o formato Compressed Sparse Row (CSR). O formato compacta os dados utilizando três vetores unidimensionais [Vuduc et al. 2002]:

• data: armazena de forma contínua os valores reais não-nulos (Non-Zero Elements - NNZ) da matriz.
• col: armazena os índices das colunas originais correspondentes a cada elemento presente no vetor data.
• row_ptr: vetor de ponteiros com tamanho \(M + 1\) que delimita os intervalos de início e término dos elementos de cada linha nos vetores data e col.

O cálculo da SpMV utilizando o CSR processa cada linha \(i\) por meio do somatório:

\[ y[i] = \sum_{j=\text{row\_ptr}[i]}^{\text{row\_ptr}[i+1]-1} \text{data}[j] \cdot x[\text{col}[j]] \]

Apesar da alta eficiência de compactação, o acesso indireto ao vetor multiplicador \(x\) por meio de \(\text{col}[j]\) gera leituras irregulares e dispersas. Esse comportamento destrói a localidade espacial e temporal dos dados, tornando ineficazes os mecanismos de pré-busca baseados em hardware [Vuduc et al. 2002].

2.3 O Modelo Roofline e Análise de Desempenho

Proposto por Williams et al. [2009], o modelo Roofline é uma ferramenta analítica visual concebida para delimitar o desempenho máximo alcançável de algoritmos em arquiteturas de hardware modernas. O modelo relaciona o desempenho computacional (em GFlops/s) com a largura de banda da memória (em GB/s) por meio da métrica de intensidade operacional (operações de ponto flutuante por byte trafegado). O teto diagonal do modelo evidencia o limite físico da largura de banda da memória principal, caracterizando o gargalo do Memory Wall ao qual a SpMV está submetida [Vuduc et al. 2002, Williams et al. 2009].

Para avaliar a efetividade das otimizações paralelas frente à linha de base sequencial, a análise de desempenho emprega a métrica de ganho de aceleração (SpeedUp), cujos limites teóricos de escalabilidade em sistemas paralelos são governados pela formulação clássica de Amdahl [1967].

2.4 Vetorização de Hardware e Instruções SIMD

Para maximizar o desempenho aritmético em nível de núcleo único, os processadores modernos incorporam unidades vetoriais capazes de executar o paradigma Single Instruction, Multiple Data (SIMD). Por meio de extensões de conjunto de instruções como o AVX (Advanced Vector Extensions), o processador pode operar simultaneamente sobre múltiplos elementos de dados utilizando um único ciclo de clock [Hager and Wellein 2010]. Em códigos de simulação numérica, a ativação de instruções SIMD combinada com otimizações agressivas de compilação assegura que a linha de base sequencial explore plenamente a capacidade de ponto flutuante do hardware, permitindo isolar rigorosamente o ganho de aceleração atribuído estritamente à paralelização de múltiplos núcleos [Williams et al. 2009, Hager and Wellein 2010].

2.5 O Modelo de Programação Paralela OpenMP

O OpenMP (Open Multi-Processing) é um padrão amplamente estabelecido para a programação paralela em arquiteturas de memória compartilhada [Dagum and Menon 1998]. Ele baseia-se em um modelo de execução do tipo fork-join, onde uma thread mestre cria (fork) uma equipe de threads paralelas para executar blocos de código computacionalmente intensivos delimitados por diretivas de compilação, sincronizando-se novamente ao término da tarefa (join).

Para mitigar problemas de desbalanceamento de carga decorrentes da variação no número de elementos por linha em matrizes esparsas, o OpenMP disponibiliza diferentes políticas de escalonamento (scheduling), tais como o estático (static), onde as iterações são divididas previamente entre as threads, e o guiado (guided), que ajusta dinamicamente o tamanho dos blocos de trabalho atribuídos às threads à medida que a execução avança.

2.6 Descarregamento de Processamento (Offloading) via OpenMP em GPUs

Com a evolução das diretivas OpenMP em suas versões recentes, o padrão expandiu seu escopo para além de arquiteturas de CPU, introduzindo suporte nativo a modelos de computação heterogênea baseados em aceleradores, como as unidades de processamento gráfico (GPUs). Por meio de diretivas de offloading (como target e teams distribute parallel for), o programador pode direcionar blocos computacionais e estruturas de dados da memória principal do hospedeiro (host) para a memória de vídeo dedicada do dispositivo acelerador (device) [Van der Pas et al. 2017].

No entanto, em operações limitadas por largura de banda e que apresentam padrão irregular de acesso, como a SpMV, a eficiência em arquiteturas massivamente paralelas é fortemente desafiada pelas características de hardware das GPUs. Conforme apontado por Bell e Garland [2009], o desempenho torna-se estritamente condicionado à largura de banda do barramento de comunicação (PCI-Express) e à penalidade imposta por acessos desalinhados à memória global da GPU (uncoalesced memory accesses), o que torna o uso direto do formato CSR em aceleradores altamente dependente de estratégias de otimização estrutural.

3. Metodologia

3.1 Conjuntos de Dados e Matrizes Experimentais

Para garantir a representatividade dos testes em cenários práticos de simulação numérica, os experimentos utilizaram matrizes esparsas reais extraídas da SuiteSparse Matrix Collection [Davis and Hu 2011]:

• bcsstk14 e bcsstk18: Matrizes de rigidez estrutural e elementos finitos (grupo HB), caracterizadas por alta simetria.
• G3_circuit: Matriz de grande porte representativa de circuitos integrados (grupo AMD), com alta dispersão.

3.2 Ambiente Experimental e Compilação

• Ambiente de CPU (Cluster Sirius): Processador Intel Core i7-13700 (24 threads lógicas e suporte a AVX2). A submissão foi automatizada por scripts de lote direcionados ao gerenciador Grid Engine (SGE) com alocação dedicada de 24 núcleos (-pe openmp 24). Compilador GCC com -O3 e vetorização SIMD ativa.
• Acelerador Gráfico (GPU Local): NVIDIA GeForce RTX 4050 Laptop (6 GB VRAM, sm_89), Clang LLVM via libomptarget.

3.3 Algoritmos Implementados

A rotina computacional elementar baseia-se na função computeRowSum:

inline double computeRowSum(const std::vector<double>& x, int i) const {
    double sum = 0.0;
    int init = row[i], end = row[i + 1];
    for (int j = init; j < end; ++j) {
        sum += data[j] * x[col[j]];
    }
    return sum;
}

Estratégias de execução paralela implementadas:

// 1. Sequencial Vetorizada (Linha de Base com AVX2)
for (int i = 0; i < nRows; ++i) {
    y[i] = computeRowSum(x, i);
}

// 2. OpenMP Estatico / Guiado
#pragma omp parallel for num_threads(threads) schedule(guided)
for (int i = 0; i < nRows; ++i) {
    y[i] = computeRowSum(x, i);
}

// 3. Particionamento Balanceado por Elementos Nao-Nulos (NNZ)
#pragma omp parallel num_threads(nThreads) {
    int tid = omp_get_thread_num();
    for (int i = threadRowStart[tid]; i < threadRowStart[tid + 1]; ++i) {
        y[i] = computeRowSum(x, i);
    }
}

// 4. Descarregamento para GPU (Target Offloading)
#pragma omp target teams distribute parallel for \
    map(to: ptrData[0:nnz], ptrCol[0:nnz], ptrRow[0:rows+1], ptrX[0:cols]) \
    map(from: ptrY[0:rows])
for (int i = 0; i < rows; ++i) {
    double sum = 0.0;
    for (int j = ptrRow[i]; j < ptrRow[i+1]; ++j) {
        sum += ptrData[j] * ptrX[ptrCol[j]];
    }
    ptrY[i] = sum;
}

O código-fonte completo está disponível publicamente no repositório do projeto [Viana 2026].

4. Avaliação Experimental

4.1 Impacto da Vetorização e Linha de Base Sequencial

A linha de base sequencial foi avaliada no cluster Sirius comparando a versão padrão sem otimizações e a versão otimizada com AVX e -O3. A ativação de SIMD elevou drasticamente a largura de banda e a eficiência de ponto flutuante em todas as matrizes, assegurando que o núcleo único opere próximo ao limite do hardware.

(a) Desempenho Sequencial (GB/s)

(b) Escalabilidade CPU OpenMP

4.2 Escalabilidade e Desempenho Paralelo em CPU (OpenMP)

A avaliação no cluster Sirius explorou as estratégias Static, Guided e Balanceado por NNZ ($1$ a $24$ threads). O escalonamento Guided demonstrou superioridade ao redistribuir dinamicamente blocos menores no final da execução, mitigando com eficiência o load imbalance inerente ao formato CSR.

4.3 Avaliação em GPU e Análise de SpeedUp

O descarregamento via OpenMP na NVIDIA GeForce RTX 4050 obteve vazões entre $3.11\text{ GB/s}$ e $5.18\text{ GB/s}$. Essa restrição severa decorre do custo de transferência bidirecional via PCI-Express e da penalidade imposta por acessos desalinhados à memória global (uncoalesced memory accesses) no CSR.

(a) Banda Efetiva GPU (GB/s)

(b) SpeedUp GPU vs. Sequencial

(a) SpeedUp OpenMP CPU

(b) Modelo Roofline Comparativo

4.4 Análise do Modelo Roofline

A projeção analítica sob o modelo Roofline ratifica que a SpMV opera estritamente na região de Memory Bound, com intensidade operacional característica abaixo de \(0.2\text{ Flops/byte}\). Como os pontos experimentais alinham-se à rampa inclinada de ambas as arquiteturas, conclui-se que o desempenho do algoritmo é condicionado exclusivamente pela largura de banda da memória e não pelo poder de processamento aritmético bruto (GFlops).

5. Trabalhos Relacionados

A otimização da SpMV e a caracterização de seus limites físicos têm sido amplamente investigadas na literatura de computação de alto desempenho. No trabalho seminal de Vuduc et al. [2002], foram estabelecidos limites teóricos de desempenho para a SpMV em formatos compactados, demonstrando que a operação é limitada por largura de banda e propondo otimizações em múltiplos níveis de memória cache e registradores.

Com a popularização de arquiteturas massivas, diversos estudos analisaram o impacto de formatos esparsos em GPUs [Bell and Garland 2009]. Em particular, Bell e Garland evidenciaram que o formato CSR sofre com divergência de threads e acessos não coalescidos à memória global (uncoalesced memory accesses), propondo alternativas como o ELLPACK e variações por linha (CSR-Vector). De forma complementar, o modelo analítico Roofline [Williams et al. 2009] consolidou a SpMV como caso clássico de saturação de banda. Este trabalho dialoga com essas referências ao confrontar o comportamento do CSR nativo em nós dedicados de CPU frente ao descarregamento via #pragma omp target em aceleradores gráficos de consumo.

6. Considerações Finais

Este trabalho apresentou uma análise microarquitetural da SpMV em formato CSR em ambientes heterogêneos. A vetorização SIMD com -O3 provou-se essencial para saturar o núcleo único. Em sistemas multicore, o OpenMP obteve ganhos expressivos com escalonamento Guided, mitigando o desbalanceamento de carga até atingir saturação por largura de banda entre 12 e 16 threads. Em contraste, o offloading em GPU obteve desempenho inferior à baseline sequencial devido à sobrecarga de barramento PCI-Express e acessos dispersos. A validação via modelo Roofline confirmou analiticamente o regime Memory-Bound da aplicação em ambas as arquiteturas, ressaltando que a vazão de dados sobrepõe-se à capacidade computacional bruta.