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

Sintesi Tecnica: Accelerazione Portabile di Meccanismi di Incapsulamento delle Chiavi basati su Learning With Errors per la Crittografia Post-Quantistica

Problematica
La transizione alla Crittografia Post-Quantistica (PQC) richiede implementazioni capaci di soddisfare le richieste computazionali delle applicazioni del mondo reale. Tra i candidati PQC, i Meccanismi di Incapsulamento delle Chiavi (KEM) basati su Learning With Errors (LWE) offrono solide fondamenta di sicurezza, ma comportano un overhead computazionale sostanziale. Questo overhead è guidato principalmente da operazioni matrice-vettore su larga scala e dalla generazione di numeri casuali crittograficamente sicuri. Le Unità di Elaborazione della Grafica (GPU) sono candidati naturali per accelerare questi carichi di lavoro paralleli, tuttavia le attuali implementazioni GPU si affidano ampiamente a ottimizzazioni specifiche per CUDA, creando un vincolo al fornitore (vendor lock-in) e limitando la portabilità attraverso ecosistemi di acceleratori eterogenei (ad esempio, NVIDIA e AMD). Inoltre, precedenti tentativi di portabilità utilizzando OpenACC hanno mostrato un supporto cross-vendor limitato. La sfida centrale affrontata è il raggiungimento di un'accelerazione ad alte prestazioni per KEM basati su LWE densi, mantenendo al contempo la portabilità a livello di sorgente tra diverse architetture GPU senza sacrificare la capacità di mantenere componenti critici, come la generazione di numeri casuali, residenti sul dispositivo.

Metodologia
Gli autori presentano un'implementazione portabile su GPU di un KEM basato su LWE puro utilizzando l'offloading OpenMP Target. La metodologia si basa su tre strategie tecniche primarie:

1. Strategia di Offloading Portabile: L'implementazione utilizza le direttive OpenMP Target per gestire il movimento dei dati e l'esecuzione dei kernel. Per garantire la residenza sulla GPU e minimizzare la comunicazione host-device, gli autori impiegano una gestione della memoria persistente sul dispositivo (target enter data map(alloc:...)). Ciò mantiene le matrici principali, i ciphertext e i buffer temporanei residenti sulla GPU durante le fasi di KeyGen, Encapsulation e Decapsulation.
2. Generazione di Numeri Casuali (RNG) Eterogenea: Un componente critico del flusso di lavoro è la generazione di numeri casuali crittograficamente sicuri. Gli autori hanno esteso la libreria esistente RNGonGPU, originariamente progettata per CUDA, per supportare l'interfaccia HIP (Heterogeneous-computing Interface for Portability). Integrando uno strato di adattamento compatibile con HIP al flusso di lavoro OpenMP Target, essi consentono la generazione di numeri casuali interamente residente sulla GPU sia su acceleratori NVIDIA che AMD. Ciò viene ottenuto tramite #pragma omp target data use_device_ptr per esporre i buffer di dispositivo gestiti da OpenMP al backend RNG esterno basato su HIP, evitando la necessità di generare la casualità sull'host e copiarla sul dispositivo.
3. Parallelizzazione dei Kernel: I kernel computazionali sono parallelizzati utilizzando le costrutti OpenMP teams distribute parallel for. La generazione delle chiavi prevede la generazione parallela di vettori segreti e vettori di errore, seguita da prodotti matrice-vettore densi. La cifratura e la decifratura sono processate in batch, elaborando simultaneamente bit di messaggio indipendenti o slot di ciphertext per massimizzare il parallelismo a livello di thread e il riutilizzo dei buffer. Anche i controlli di uguaglianza della trasformata di Fujisaki-Okamoto (FO) sono eseguiti interamente sulla GPU, restituendo al solo host un singolo flag di stato.

Configurazione Sperimentale
Lo studio ha valutato l'implementazione attraverso quattro distinte architetture di acceleratori:

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

Le prestazioni sono state misurate rispetto a un baseline CPU multicore OpenMP e confrontate tra le piattaforme utilizzando dimensioni del reticolo ($n$) variabili fino a 16.384. Le metriche includevano il tempo di esecuzione, la scalabilità, l'energia per soluzione (energy-to-solution) e un profilo dettagliato del runtime utilizzando strumenti del fornitore (Nseste Systems e ROCm Systems Profiler).

Risultati Chiave

• Accelerazione rispetto alla CPU: L'offloading OpenMP Target ha dimostrato un'accelerazione sostanziale rispetto al baseline CPU multicore. Per grandi dimensioni del reticolo (ad esempio, $n=4096$), l'accelerazione ha superato i due ordini di grandezza (120x), principalmente perché il carico di lavoro è limitato dalla memoria (memory-bound), e le GPU offrono una larghezza di banda della memoria e un parallelismo superiori per le operazioni di streaming delle matrici.
• Prestazioni Cross-Platform:
  • NVIDIA GH200 ha ottenuto le migliori prestazioni complessive, grazie alla sua elevata larghezza di banda HBM3 (4 TB/s) e all'interconnessione NVLink-C2C, che ha minimizzato l'overhead di comunicazione CPU-GPU.
  • AMD MI300X è seguita da vicino, mostrando prestazioni quasi ottimali per dimensioni del reticolo più piccole e una forte scalabilità, sebbene abbia mostrato tempi di esecuzione leggermente superiori per $n$ più grandi rispetto alla GH200.
  • NVIDIA A100 ha mostrato rallentamenti moderati rispetto alla GH200, limitati dalla sua inferiore larghezza di banda HBM2e e dall'overhead di comunicazione PCIe.
  • AMD MI300A ha esibito la scalabilità più debole e i tempi di esecuzione più elevati. Il profiling ha rivelato che, nonostante la presenza di una memoria unificata ad alta larghezza di banda, il sottosistema HBM3 condiviso ha creato contesa tra le operazioni lato CPU (hashing SHA3, derivazione chiavi FO) e i kernel GPU, riducendo la larghezza di banda effettiva di circa il 38% rispetto alla MI300X discreta.
• Efficienza Energetica: La NVIDIA GH200 è stata trovata essere circa 2,5 volte più efficiente dal punto di vista energetico rispetto alla AMD MI300X per grandi dimensioni del reticolo, grazie sia ai tempi di esecuzione più brevi che al minor consumo medio di potenza.
• Portabilità: L'unico codice sorgente ha eseguito con successo su tutte le piattaforme testate senza modifiche al codice specifico dell'architettura, convalidando la fattibilità di OpenMP Target per l'accelerazione portabile della PQC.

Significatività e Rivendicazioni
L'articolo sostiene che l'accelerazione portabile su GPU sia una strategia valida per ridurre l'overhead computazionale della PQC senza incorrere nel vendor lock-in. Lo studio dimostra che:

1. La portabilità a livello di sorgente è realizzabile: OpenMP Target, combinato con un backend RNG integrato in HIP, consente a un singolo codebase di sfruttare efficacemente gli ecosistemi GPU eterogenei.
2. L'architettura della memoria determina le prestazioni: Per i carichi di lavoro LWE limitati dalla memoria, la capacità di calcolo di picco è meno critica della larghezza di banda della memoria, dell'organizzazione del sistema di memoria e dell'efficienza dell'interconnessione host-device. I risultati evidenziano che le architetture a memoria unificata (come la MI300A) non garantiscono automaticamente prestazioni superiori se il traffico CPU e GPU compete per le stesse risorse di memoria.
3. Implementazione pratica: Evitando le dipendenze specifiche di CUDA, questo approccio facilita il dispiegamento di infrastrutture crittografiche resistenti ai quanti attraverso diversi ambienti hardware.

Gli autori concludono che, sebbene OpenMP Target fornisca un'efficace portabilità a livello di sorgente, le prestazioni a runtime rimangono dipendenti dalla maturità degli strumenti di compilazione (toolchain), dal comportamento delle librerie di dispositivo interoperabili (specificamente riguardo all'allocazione dinamica della memoria nell'adattatore RNG HIP) e dalla specifica gerarchia di memoria dell'hardware target.