Portable Acceleration of Learning With Errors KEMs for Post-Quantum Cryptography

Resumo Técnico: Aceleração Portátil de Mecanismos de Encapsulamento de Chave com Erros de Aprendizado para Criptografia Pós-Quântica

Declaração do Problema
A transição para a Criptografia Pós-Quântica (PQC) exige implementações capazes de atender às demandas computacionais de aplicações do mundo real. Entre os candidatos a PQC, os Mecanismos de Encapsulamento de Chave (KEMs) baseados em Aprendizado com Erros (LWE) oferecem bases de segurança robustas, mas incorrem em um overhead computacional substancial. Esse overhead é impulsionado principalmente por operações de matriz-vetor de larga escala e pela geração de números aleatórios criptograficamente seguros. Embora as Unidades de Processamento Gráfico (GPUs) sejam candidatas naturais para acelerar essas cargas de trabalho paralelas, as implementações de GPU existentes dependem amplamente de otimizações específicas de CUDA, criando um bloqueio de fornecedor (vendor lock-in) e limitando a portabilidade entre ecossistemas de aceleradores heterogêneos (por exemplo, NVIDIA e AMD). Além disso, tentativas anteriores de portabilidade usando OpenACC mostraram suporte limitado entre diferentes fornecedores. O desafio central abordado é alcançar uma aceleração de alto desempenho para KEMs densos baseados em LWE, mantendo a portabilidade ao nível de código-fonte entre diferentes arquiteturas de GPU sem sacrificar a capacidade de manter componentes críticos, como a geração de números aleatórios, residentes no dispositivo.

Metodologia
Os autores apresentam uma implementação de GPU portátil de um KEM baseado em LWE puro usando offloading de alvo OpenMP (OpenMP Target offloading). A metodologia envolve três estratégias técnicas primárias:

1. Estratégia de Offloading Portátil: A implementação utiliza diretivas de Alvo OpenMP para gerenciar o movimento de dados e a execução de kernels. Para garantir a residência na GPU e minimizar a comunicação entre host e dispositivo, os autores empregam o gerenciamento de memória persistente no dispositivo (target enter data map(alloc:...)). Isso mantém as matrizes principais, criptogramas e buffers temporários residentes na GPU durante as fases de Geração de Chave (KeyGen), Encapsulamento e Decapsulamento.
2. Geração de Números Aleatórios (RNG) Heterogênea: Um componente crítico do fluxo de trabalho é a geração de números aleatórios criptograficamente seguros. Os autores estenderam a biblioteca existente RNGonGPU, originalmente projetada para CUDA, para suportar o Heterogeneous-computing Interface for Portability (HIP). Ao integrar uma camada adaptadora compatível com HIP ao fluxo de trabalho de Alvo OpenMP, eles permitem a geração de números aleatórios totalmente residente na GPU tanto em aceleradores NVIDIA quanto AMD. Isso é alcançado usando #pragma omp target data use_device_ptr para expor buffers de dispositivo gerenciados pelo OpenMP ao backend de RNG externo baseado em HIP, evitando a necessidade de gerar aleatoriedade no host e copiá-la para o dispositivo.
3. Paralelização de Kernels: Os kernels computacionais são paralelizados usando construtos teams distribute parallel for do OpenMP. A geração de chaves envolve a geração paralela de vetores secretos e vetores de erro, seguida por produtos matriz-vetor densos. A criptografia e a descriptografia são processadas em lotes (batched), processando bits de mensagem ou slots de criptograma independentes simultaneamente para maximizar o paralelismo ao nível de thread e o reuso de buffers. As verificações de igualidade da transformação de Fujisaki-Okamoto (FO) também são realizadas inteiramente na GPU, retornando apenas um único sinal de status ao host.

Configuração Experimental
O estudo avaliou a implementação em quatro arquiteturas de aceleradores distintas:

• NVIDIA: A100-SXM-64GB (PCIe 4.0, HBM2e) e GH200 Grace Hopper (NVLink-C2C, HBM3).
• AMD: MI300X (PCIe 5.0, HBM3) e MI300A (APU Unificada, HBM3 compartilhada).

O desempenho foi medido contra um baseline de CPU multicore OpenMP e comparado entre plataformas usando dimensões de rede ($n$) variando até 16.384. As métricas incluíram tempo de execução, escalabilidade, energia por solução e perfilamento de tempo de execução detalhado usando ferramentas de fornecedores (Nsight Systems e ROCm Systems Profiler).

Principais Resultados

• Aceleração vs. CPU: O offloading de Alvo OpenMP demonstrou uma aceleração substancial sobre o baseline de CPU multicore. Para grandes dimensões de rede (ex: $n=4096$), os ganhos de velocidade excederam duas ordens de magnitude (120x), principalmente porque a carga de trabalho é limitada por memória (memory-bound), e as GPUs oferecem largura de banda de memória e paralelismo superiores para operações de streaming de matrizes.
• Desempenho Cross-Platform:
  • A NVIDIA GH200 alcançou o melhor desempenho geral, atribuído à sua alta largura de banda HBM3 (4 TB/s) e ao interconector NVLink-C2C, que minimizou o overhead de comunicação CPU-GPU.
  • A AMD MI300X seguiu de perto, mostrando desempenho próximo ao ótimo para dimensões de rede menores e forte escalabilidade, embora tenha exibido tempos de execução ligeiramente maiores para $n$ maiores em comparação com a GH200.
  • A NVIDIA A100 mostrou lentidão moderada em relação à GH200, limitada por sua menor largura de banda HBM2e e overhead de comunicação PCIe.
  • A AMD MI300A exibiu a escalabilidade mais fraca e os maiores tempos de execução. O perfilamento revelou que, apesar de possuir memória unificada de alta largura de banda, o subsistema HBM3 compartilhado criou contenção entre as operações do lado da CPU (hashing SHA3, derivação de chave FO) e os kernels da GPU, reduzindo a largura de banda efetiva em aproximadamente 38% em comparação com a MI300X discreta.
• Eficiência Energética: A NVIDIA GH200 foi considerada aproximadamente 2,5x mais eficiente energeticamente do que a AMD MI300X para grandes dimensões de rede, impulsionada tanto pelos tempos de execução mais curtos quanto pelo menor consumo de potência média.
• Portabilidade: A base de código única executou com sucesso em todas as plataformas testadas sem modificações de código específicas de arquitetura, validando a viabilidade do Alvo OpenMP para aceleração de PQC portátil.

Significância e Alegações
O artigo afirma que a aceleração de GPU portátil é uma estratégia viável para reduzir o overhead computacional de PQC sem incorrer em bloqueio de fornecedor. O estudo demonstra que:

1. A portabilidade ao nível de código-fonte é alcançável: O Alvo OpenMP, combinado com um backend de RNG integrado via HIP, permite que uma única base de código aproveite ecossistemas de GPU heterogêneos de forma eficaz.
2. A arquitetura de memória dita o desempenho: Para cargas de trabalho LWE limitadas por memória, a capacidade de computação de pico é menos crítica do que a largura de banda da memória, a organização do sistema de memória e a eficiência do interconector host-dispositivo. Os resultados destacam que arquiteturas de memória unificada (como a MI300A) não garantem automaticamente desempenho superior se o tráfego da CPU e da GPU competir pelos mesmos recursos de memória.
3. Implantação Prática: Ao evitar dependências específicas de CUDA, esta abordagem facilita a implantação de infraestruturas criptográficas resistentes ao quantum em diversos ambientes de hardware.

Os autores concluem que, embora o Alvo OpenMP forneça um offloading de código-fonte eficaz, o desempenho em tempo de execução permanece dependente da maturidade das ferramentas de compilador, do comportamento das bibliotecas de dispositivo interoperáveis (especificamente em relação à alocação dinâmica de memória no adaptador de RNG HIP) e da hierarquia de memória específica do hardware alvo.