Evaluating Application Characteristics for GPU Portability Layer Selection

Este artigo apresenta um estudo analisando aplicações heterogêneas representativas de grandes experimentos de Física de Altas Energias para identificar características fundamentais que influenciam o desempenho através de várias camadas de portabilidade de GPU, orientando assim os desenvolvedores na seleção da tecnologia mais adequada para suas necessidades específicas.

O essencial

Imagine que você é um chef tentando cozinhar um banquete massivo. Você tem três tipos diferentes de fornos de alta potência em sua cozinha: um feito pela NVIDIA, um pela AMD e um pela Intel. Cada forno cozinha os alimentos de forma diferente, usa botões diferentes e exige receitas diferentes para funcionar da melhor maneira.

Se você escrever sua receita especificamente para o forno NVIDIA (usando uma linguagem chamada CUDA), você não pode simplesmente colocar essa mesma receita nos fornos AMD ou Intel. Você teria que reescrever tudo. Isso é um problema porque você nem sempre sabe qual forno terá em sua cozinha amanhã.

Para resolver isso, o artigo discute "camadas de portabilidade". Pense nelas como tradutores universais ou adaptadores inteligentes. Eles permitem que você escreva uma receita mestre que o tradutor converte na linguagem específica que cada forno entende. O artigo analisa vários desses tradutores (como Kokkos, SYCL, OpenMP e Alpaka) para ver qual é o melhor ajuste para diferentes tipos de tarefas culinárias.

Aqui está o que os autores descobriram quando testaram esses tradutores com "receitas" reais de experimentos de física de altas energias (como os usados para estudar partículas subatômicas):

1. O Problema do "Tempo de Inicialização"

Ligar uma GPU (um forno) não é instantâneo. Leva alguns milissegundos para acordar e ficar pronta.

• O Problema: Alguns tradutores são lentos para iniciar o processo de cozimento. Por exemplo, o Kokkos pode adicionar um atraso significativo ao usar fornos AMD. Se sua tarefa de cozimento for muito curta (como ferver um ovo por 10 segundos), e o tradutor levar 5 segundos apenas para ligar o fogão, você terá desperdiçado metade do seu tempo.
• A Lição: Se suas tarefas forem minúsculas e rápidas, evite tradutores que tornem a inicialização lenta.

2. O Problema da "Cozinha Lotada"

Em um laboratório de física real, a GPU não trabalha sozinha. Ela faz parte de um sistema maior onde muitas pessoas (threads) estão tentando usar o forno ao mesmo tempo.

• O Problema: Alguns tradutores são ruins em lidar com multidões. O Kokkos, por exemplo, tem uma regra que diz: "Apenas uma pessoa pode falar com o forno por vez", o que causa um congestionamento se vários chefs tentarem lançar tarefas simultaneamente. O SYCL é um pouco inconsistente; às vezes ele deixa todos cozinharem ao mesmo tempo, e às vezes força eles a esperarem na fila, dependendo de qual versão do tradutor você está usando.
• A Lição: Se sua aplicação precisa de muitas pessoas trabalhando ao mesmo tempo, você precisa de um tradutor que saiba gerenciar uma cozinha movimentada sem bloquear as portas.

3. O Problema da "Compatibilidade de Ferramentas"

Receitas de física costumam usar ferramentas especiais (bibliotecas como ROOT ou Eigen) que ajudam com matemática e dados.

• O Problema: Algumas dessas ferramentas não se dão bem com os tradutores. Por exemplo, uma ferramenta matemática popular chamada Eigen frequentemente quebra quando usada com o compilador NVIDIA, no qual muitos tradutores se baseiam. Além disso, tentar usar dois compiladores diferentes (um para a CPU e outro para a GPU) no mesmo projeto é como tentar construir uma casa com dois conjuntos de plantas diferentes que não combinam — torna a construção (build) do software um pesadelo.
• A Lição: Antes de escolher um tradutor, verifique se suas ferramentas favoritas cabem dentro dele.

4. O Problema do "Arranjo de Móveis"

As GPUs adoram layouts simples e planos. Elas preferem que os dados sejam organizados como uma fileira organizada de caixas. No entanto, os dados de física costem vir em formas complexas e bagunçadas (como uma pilha de malas de tamanhos diferentes).

• O Problema: Os tradutores tentam corrigir essa bagunça envolvendo os dados em recipientes especiais. Embora isso torne o código portátil, adiciona "overhead" — como colocar cada item em uma mala antes de movê-lo, mesmo que você precise mover apenas uma meia. Isso reduz a velocidade. Além disso, nenhum dos tradutores é muito bom em lidar com dados "irregulares" (linhas de comprimentos diferentes), o que é muito comum na física.
• A Lição: Se seus dados forem complexos e bagunçados, o tradutor pode te atrasar tentando arrumá-los.

5. O Problema das "Ferramentas Especializadas"

Às vezes você precisa de uma ferramenta específica, como um Gerador de Números Aleatórios (RNG) ou uma Transformada Rápida de Fourier (FFT).

• O Problema: Cada fabricante de forno tem suas próprias versões especializadas e super rápidas dessas ferramentas. Os tradutores universais muitas vezes não incluem essas versões especializadas ou usam suas próprias versões mais lentas. Embora você possa forçar o tradutor a usar a ferramenta nativa do forno, isso quebra a "portabilidade", porque essa ferramenta só funciona naquele forno específico.
• A Lição: Se você depende fortemente dessas ferramentas específicas, pode ter que escolher entre velocidade (usar a ferramenta nativa do forno) ou portabilidade (usar a ferramenta genérica do tradutor).

6. Os Problemas de "Tempo de Construção" e "Dia de Mudança"

• Construindo a Receita: Alguns tradutores tornam o "tempo de cozimento" (tempo de compilação) muito mais longo. Para projetos enormes, usar certos tradutores pode fazer com que o processo de build leve horas em vez de minutos.
• Mudando a Cozinha: Se você construir seu software para um forno específico (ex: NVIDIA V100), ele pode não funcionar em um mais novo (ex: NVIDIA A100). Alguns tradutores exigem que você construa uma versão separada para cada tipo de forno que possa encontrar. Isso cria um pesadelo logístico enorme para distribuir o software para diferentes laboratórios.

O Veredito Final

O artigo conclui que não existe um tradutor "perfeito".

• Kokkos é ótimo para algumas coisas, mas tem dificuldades com concorrência e tempos de inicialização em certos hardwares.
• SYCL é poderoso, mas pode ser inconsistente dependendo da versão do compilador.
• OpenMP e outros têm seus próprios pontos fortes e fracos em relação a como lidam com memória e diferentes hardwares.

A Conclusão: Você não pode simplesmente escolher um tradutor porque ele é popular. Você tem que olhar para sua "receita" específica (sua aplicação). Se seu código é curto e rápido, escolha um tradutor com baixo tempo de inicialização. Se seu código é complexo e usa muitas ferramentas, escolha um que se dê bem com essas ferramentas.

Os autores também observam que essas tecnologias estão evoluindo rapidamente, como novos modelos de fornos saindo todos os anos. O que funciona melhor hoje pode mudar amanhã, então os desenvolvedores precisam continuar monitorando o cenário. No futuro, novos padrões podem tornar essas escolhas mais fáceis, mas por enquanto, testes cuidadosos são a única maneira de encontrar o ajuste ideal.

Resumo técnico

Resumo Técnico: Avaliando Características de Aplicação para Seleção de Camada de Portabilidade de GPU

Declaração do Problema

A computação de Alta Energia (HEP - High Energy Physics) depende cada vez mais de GPUs para desempenho, mas o cenário de hardware está se diversificando além da NVIDIA para incluir AMD e Intel. Embora linguagens específicas de fornecedores como CUDA ofereçam alto desempenho, elas restringem a implantação a hardwares específicos, reduzindo a portabilidade. Por outro lado, camadas de portabilidade como Kokkos, SYCL, HIP, OpenMP, Alpaka e std::par prometem execução em diversas arquiteturas de CPU e GPU. No entanto, essas camadas não são funcionalmente equivalentes; elas exibem características distintas no que diz respeito a desempenho, suporte a concorrência e compatibilidade com bases de código existentes. O problema central abordado é a falta de um framework claro para desenvolvedores selecionarem a camada de portabilidade mais apropriada com base em características específicas da aplicação, estruturas de kernel e ambientes de execução.

Metodologia

O Centro de Excelência em Computação de Alta Energia (HEP-CCE) conduziu um estudo comparativo utilizando aplicações heterogêneas representativas de grandes experimentos de HEP:

• ATLAS: FastCaloSim (simulação de calorímetro parametrizado).
• CMS: Patatrack (reconstrução de pixel heterogênea com mais de 40 kernels, multithreading e pools de memória) e P2R (reconstrução de trajetória de kernel único).
• DUNE: Wire-Cell Toolkit (kernels de rasterização, scatter-add e convolução FFT).
• Geradores de Eventos: Portas para GPU de Sherpa (Pepper) e Madgraph (utilizando código gerado automaticamente).

Esses bancos de teste foram portados de C++ serial para várias tecnologias de portabilidade para identificar características de código que impactam significamente o desempenho e a compatibilidade. O estudo avaliou essas aplicações em múltiplas dimensões, incluindo latência de lançamento de kernel, tratamento de concorrência, compatibilidade de biblioteca, gerenciamento de memória e overhead de compilação.

Principais Descobertas e Resultados

1. Latência de Lançamento de Kernel

O lançamento de kernel não é instantâneo e pode variar de microssegundos a dezenas de microssegundos.

• Kokkos: Observou-se que adiciona um overhead significativo (dezenas de microssegundos), particularmente em GPUs AMD, que já possuem latências de lançamento nativas mais altas do que NVIDIA ou Intel.
• SYCL, HIP, OpenMP, Alpaka: Não exibiram penalidades adicionais grandes em comparação com implementações nativas.
• Especificidades da AMD: Chamar rocRAND em GPUs AMD resultou em latências extremas na primeira chamada (excedendo 1 segundo), sugerindo que kernels de execução curta podem sofrer degradação severa de desempenho se utilizarem camadas de portabilidade que exacerbam o overhead de lançamento.

2. Concorrência e Pools de Threads

Aplicações de HEP frequentemente rodam dentro de frameworks maiores que utilizam multithreading (ex: Intel TBB).

• Kokkos: Demonstrou incompatibilidades significativas com pools de threads externos. Seu backend serial utiliza locks que serializam chamadas, e seu backend de threads proíbe chamadas de threads externas. Lançamentos de kernels concorrentes são atualmente viáveis apenas com backends CUDA/HIP usando APIs específicas de arquitetura, limitando a portabilidade.
• SYCL: As implementações foram inconsistentes. Algumas serializaram chamadas concorrentes de diferentes threads dependendo do fabricante do compilador e da versão. Arquiteturas AMD pareciam carecer de uma solução SYCL concorrente no momento do estudo.

3. Compatibilidade de Bibliotecas Externas e Compiladores

• Eigen: Muitas versões são incompatíveis com o nvcc da NVIDIA, afetando camadas de portabilidade que dependem dele (Kokkos, Alpaka). Também é atualmente incompatível com OpenMP offload.
• ROOT: Não pode ser totalmente compilado com o nvc++ da NVIDIA. A compilação híbrida (usando g++ e nvc++) é possível, mas complica as regras de build e exige gerenciamento cuidadoso de alocação de memória (dados de GPU devem ser alocados pelo nvc++).
• Complexidade de Build: Suportar múltiplas camadas de portabilidade em um único projeto requer regras de build complexas (uso de #ifdef, separação de arquivos) para evitar conflitos de cabeçalho entre camadas como Kokkos e CUDA.

4. Estruturas de Dados e Transferências de Memória

• Layouts Orientados a Objetos vs. Planos: Estruturas de dados complexas e orientadas a objetos, comuns em HEP, mapeiam mal para GPUs, que preferem layouts planos ou Estruturas de Arrays (SoA).
• Overhead de Portabilidade: Camadas como Kokkos e SYCL fornecem abstrações (ex: Kokkos::Views, sycl::buffers) que aumentam o overhead de alocação e transferência, especialmente para objetos pequenos.
• Inicialização: A inicialização padrão de memória em Kokkos::Views pode ser um gargalo de desempenho; isso pode ser mitigado usando Kokkos::ViewAllocateWithoutInitializing.
• Unified Shared Memory (USM): Transferências automáticas de USM (usadas por std::par) foram observadas como mais lentas que transferências explícitas e podem não ser acionadas no momento ideal.
• Arrays Irregulares (Jagged Arrays): Nenhuma das APIs suporta nativamente vetores multidimensionais de tamanhos variados (jagged arrays), exigindo criação manual, redução de dimensão ou preenchimento (padding).

5. RNGs, FFTs e Atômicos

• Bibliotecas Nativas: Bibliotecas específicas de fornecedores (cuRAND, rocRAND, oneMKL) oferecem desempenho ajustado, mas carecem de portabilidade. Embora o Kokkos forneça implementações de RNG/FFT multiplataforma, elas podem não igualar o desempenho nativo.
• Contribuições do HEP-CCE: O grupo desenvolveu uma biblioteca header-only que envolve RNGs nativos de dispositivo com uma API uniforme e backends nativos de hardware para o oneMKL, permitendo que o SYCL chame implementações nativas entre diferentes fornecedores.
• Atômicos: Operações atômicas mapeiam mal para o hardware de GPU. O desempenho varia drasticamente conforme o compilador e o hardware (ex: atomic<int> na NVIDIA com OpenMP/clang13 foi 35x mais lento que CUDA). Atômicos de inteiros geralmente apresentam desempenho inferior aos de float/double.

6. Tempo de Compilação

• Overhead do Compilador: O nvc++ é geralmente 2x–3x mais lento que o g++. O Kokkos adiciona um overhead de 10%–20%.
• Código Pesado em Templates: Madgraph, utilizando código auto-gerado fortemente baseado em templates, experimentou tempos de compilação com o backend icpx do SYCL mais de 1000x mais lentos do que outros backends para kernels grandes.

7. Provisionamento de Runtime

• Kokkos: Requer ser construído com identificadores de hardware específicos. Um binário construído para uma GPU (ex: V100) pode não rodar em outra (ex: A100). Suportar bibliotecas compartilhadas e múltiplas arquiteturas cria uma matriz complexa de configurações de build, desafiando a distribuição de binários.
• Alpaka: Requer compilação separada para cada compilador de backend.
• SYCL/OneAPI: Binários multiplataforma estão emergindo, mas frequentemente limitados a arquiteturas específicas ou exigem identificação em tempo de compilação.
• OpenMP: Algumas implementações (LLVM Clang) permitem listar múltiplas arquiteturas em tempo de compilação e selecionar dinamicamente o backend em tempo de execução.

Significância e Conclusão

O artigo conclui que nenhuma solução única de portabilidade de GPU é ideal para todos os cenários. A escolha da camada de portabilidade deve ser impulsionada por uma análise cuidadosa das características específicas da aplicação, incluindo o tempo de execução do kernel, requisitos de concorrência, complexidade de estrutura de dados e dependências de bibliotecas externas.

Os autores enfatizam que o ecossistema está evoluindo rapidamente. Tanto as camadas de portabilidade (Kokkos, SYCL, etc.) quanto os compiladores nativos (clang, nvcc, hipcc) estão melhorando, com tendências voltadas para o aumento do suporte a recursos e desempenho otimizado. O artigo defende o monitoramento frequente dessas tecnologias para garantir a melhor correspondência entre bases de código e camadas de portabilidade. Adicionalmente, o artigo observa que os padrões emergentes do C++ (como o P2300 para execução assíncrona) podem simplificar essas escolhas no futuro, embora o suporte amplo de hardware ainda esteja a anos de distância.

Em última análise, selecionar uma camada de portabilidade sem compreender totalmente a natureza do código, suas dependências externas e o ambiente de execução corre o risco de causar problemas significativos tanto em desempenho quanto na flexibilidade de implantação.