💻 computer science

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

Dit artikel presenteert een draagbare GPU-implementatie van een op Learning With Errors (LWE) gebaseerd Key Encapsulation Mechanism (KEM) met behulp van OpenMP Target offloading, waarmee wordt aangetoond dat een enkele broncodebasis aanzienlijke prestatieversnelling en energie-efficiëntie kan bereiken over zowel NVIDIA- als AMD-acceleratoren heen, terwijl vendor lock-in wordt vermeden.

Oorspronkelijke auteurs: Tiziana Liberati, Nitin Shukla, Simone Rizzo, Elisabetta Boella, Matteo Barbieri, Gabriella Bettonte, Daniele Gregori, Marco Pedicini

Gepubliceerd 2026-07-13

📖 4 min leestijd☕ Koffiepauze-leesvoer

CC BY 4.0

Oorspronkelijke auteurs: Tiziana Liberati, Nitin Shukla, Simone Rizzo, Elisabetta Boella, Matteo Barbieri, Gabriella Bettonte, Daniele Gregori, Marco Pedicini

Oorspronkelijk artikel gelicentieerd onder CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ Dit is een AI-gegenereerde uitleg van het onderstaande artikel. Het is niet geschreven of goedgekeurd door de auteurs. Raadpleeg het oorspronkelijke artikel voor technische nauwkeurigheid. Lees de volledige disclaimer

Stel je voor dat je een superveilige digitale kluis probeert te bouwen die niet gekraakt kan worden, zelfs niet door een futuristische "quantum"-supercomputer. Om dit te doen, moet je een enorme, rommelige wiskundige puzzel oplossen die "Learning With Errors" (LWE) wordt genoemd. Het is alsof je probeert een specifieke naald in een hooiberg te vinden, maar de hooiberg bestaat uit miljoenen piepkleine, magnetische deeltjes die trillen en de meer je zoekt, hoe meer de magneten heen en weer wiebelen.

Het probleem? Het oplossen van deze puzzels duurt eeuwig op een gewone computer. Het is alsof je een berg zand korrel voor korrel verplaatst met een klein lepeltje. De auteurs van dit artikel vroegen zich af: "Wat als we een enorme, supersnelle bouwploeg (een GPU) gebruiken om dat zand te verplaatsen?"

De Grote Ontdekking: Eén Code, Twee Ploegen
Normaal gesproken, als je een bouwploeg wilt inhuren, moet je een specifiek team inhuren (zoals NVIDIA's CUDA) en instructies schrijven die alleen zij begrijpen. Als je een ander team wilt inhuren (zoals AMD), moet je het hele instructieboekje opnieuw schrijven. Dat is duur en irritant.

Dit artikel laat zien dat je één enkele set instructies kunt schrijven (met behulp van iets dat OpenMP Target wordt genoemd) die perfect werkt voor zowel de NVIDIA-ploeg als de AMD-ploeg. Het is alsof je een recept schrijft dat werkt, of je nu kookt in een keuken met gasfornuizen of elektrische fornuizen, zonder dat je één ingrediënt hoeft te veranderen.

De Magische Truc: In de Keuken Blijven
De grootste tijdverspiller bij deze wiskundige puzzels is het heen en weer rennen tussen de hoofdcomputer (de CPU) en de supersnelle GPU. Het is alsof een chef-kok steeds naar de voorraadkast rent voor elk klein snufje zout.

De auteurs hebben ontdekt hoe ze de "zout" (de willekeurige getallen die nodig zijn voor de wiskunde) direct in de GPU-keuken kunnen houden. Ze hebben een hulpmiddel genaamd RNGonGPU verbeterd, zodat het met zowel de NVIDIA- als de AMD-ploegen werkt. Nu kan de GPU alle willekeurige getallen genereren die hij nodig heeft zonder ooit zijn stoel te verlaten. Dit houdt de workflow soepel en snel.

De Race: Wie Wint?
Het team heeft hun nieuwe methode getest op vier verschillende supercomputers:

NVIDIA A100: Een krachtige, standaard supercomputer.
NVIDIA GH200: Een enorme "SuperChip" waarbij de hersenen (CPU) en de spieren (GPU) aan elkaar zijn gelijmd met een supersnelle snelweg (NVLink).
AMD MI300X: Een krachtige, standaard supercomputer met een enorme geheugenbank.
AMD MI300A: Een chip waarbij de hersenen en de spieren exact hetzelfde geheugen delen.

Dit is wat ze vonden:

De Snelheidsboost: Wanneer de wiskundige puzzels enorm groot werden (met een grootte van 4.096), was de GPU-versie 120 keer sneller dan de gewone computerversie. Voor zelfs grotere puzzels (grootte 16.384) was de GPU nog steeds razendsnel, terwijl de gewone computer praktisch in slaap was gevallen.
De Winnaars: De NVIDIA GH200 was de snelste en voltooide de klus in ongeveer 60 seconden. De AMD MI300X was een zeer goede tweede en deed er ongeveer 85 seconden over. Beide machines hebben enorme, snelle geheugenbanken (HBM3) die gegevens zeer snel kunnen uitstromen.
De Verrassende Verliezer: Je zou denken dat de AMD MI300A geweldig zou zijn omdat deze het geheugen deelt met de hersenen en de spieren. Maar het was juist de langzaamste, met een tijd van 114 seconden. Waarom? Omdat de hersenen en de spieren vochten om dezelfde waterkraan. Terwijl de GPU probeerde data te verplaatsen, was de CPU ook bezig met zijn eigen wiskunde, en ze zaten elkaar in de weg. Het is alsof twee mensen tegelijkertijd uit hetzelfde rietje proberen te drinken; niemand krijgt een volledige slok.

Energie en Warmte
Het team heeft ook gecontroleerd hoeveel elektriciteit deze machines gebruikten. De NVIDIA GH200 was niet alleen sneller; hij gebruikte ook minder energie om de klus te klaren. Het kostte ongeveer 9,7 kJ aan energie om de puzzel op te lossen, terwijl de AMD MI300X 26,2 kJ nodig had. Dat betekent dat de NVIDIA-machine ongeveer 2,5 keer energie-efficiënter was.

Wat Dit Betekent
Dit artikel bewijst dat je niet hoeft te kiezen tussen NVIDIA of AMD om supersnelle beveiliging te krijgen. Je kunt één codebasis gebruiken om op beide te draaien. Ze hebben echter ook aangetoond dat het hebben van een snelle chip alleen niet genoeg is; hoe het geheugen is georganiseerd, is minstens zo belangrijk. Als de hersenen en de spieren vechten om hetzelfde geheugen, vertraagt het hele systeem.

Kortom, de auteurs hebben gemeten dat draagbare, door GPU-versnelde beveiliging niet alleen een droom is — het is een realiteit die quantum-bestendige sloten veel sneller en efficiënter kan maken, zolang je de juiste hardwareopstelling kiest om verkeersopstoppingen op de geheugenbaan te voorkomen.

Technische Samenvatting: Draagbare Versnelling van Learning With Errors KEMs voor Post-Quantum Cryptografie

Probleemstelling
De overgang naar Post-Quantum Cryptografie (PQC) vereist implementaties die in staat zijn om te voldoen aan de computationele eisen van real-world toepassingen. Onder de PQC-kandidaten bieden Learning With Errors (LWE)-gebaseerde Key Encapsulation Mechanisms (KEMs) een sterke beveiligingsbasis, maar ze veroorzaken aanzienlijke computationele overhead. Deze overhead wordt primair gedreven door grootschalige matrix-vector operaties en de generatie van cryptografisch veilige willekeurige getallen. Hoewel Graphics Processing Units (GPU's) natuurlijke kandidaten zijn voor het versnellen van deze parallelle workloads, vertrouwen bestaande GPU-implementaties grotendeels op CUDA-specifieke optimalisaties, wat leidt tot vendor lock-in en de portabiliteit over heterogene accelerator-ecosystemen (bijv. NVIDIA en AMD) beperkt. Bovendien hebben eerdere draagbare pogingen met OpenACC een beperkte cross-vendor ondersteuning getoond. De kernuitdaging die wordt aangepakt, is het bereiken van hoogwaardige versnelling voor dense LWE-gebaseerde KEMs, terwijl broncode-portabiliteit behouden blijft over verschillende GPU-architecturen zonder het vermogen op te offeren om kritieke componenten, zoals de generatie van willekeurige getallen, op het device resident te houden.

Methodologie
De auteurs presenteren een draagbare GPU-implementatie van een plain LWE-gebaseerde KEM met behulp van OpenMP Target offloading. De methodologie omvat drie primaire technische strategieën:

Draagbare Offloading Strategie: De implementatie maakt gebruik van OpenMP Target directives om datamovement en kernelexecutie te beheren. Om GPU-residency te waarborgen en host-device communicatie te minimaliseren, maken de auteurs gebruik van persistent device memory management (target enter data map(alloc:...)). Dit houdt belangrijke matrices, ciphertexts en tijdelijke buffers gedurende de KeyGen, Encapsulation en Decapsulation fasen resident op de GPU.
Heterogene Random Number Generation (RNG): Een cruciaal onderdeel van de workflow is de cryptografisch veilige generatie van willekeurige getallen. De auteurs hebben de bestaande RNGonGPU bibliotheek, oorspronkelijk ontworpen voor CUDA, uitgebreid om ondersteuning te bieden voor de Heterogeneous-computing Interface for Portability (HIP). Door een HIP-compatibele adapterlaag te integreren met de OpenMP Target workflow, maken zij volledig GPU-resident random number generation mogelijk op zowel NVIDIA als AMD accelerators. Dit wordt bereikt met behulp van #pragma omp target data use_device_ptr om OpenMP-beheerde device buffers bloot te stellen aan de externe HIP-gebaseerde RNG backend, waardoor het voorkomen wordt dat willekeurheid op de host wordt gegenereerd en vervolgens naar het device wordt gekopieerd.
Kernel Parallelisatie: De computationele kernels worden geparalleliseerd met behulp van OpenMP teams distribute parallel for constructen. Key generation omvat de parallelle generatie van secret vectoren en error vectoren, gevolgd door dense matrix-vector producten. Encryptie en decryptie worden gebatcht, waarbij onafhankelijke berichtbits of ciphertext slots simultaan worden verwerkt om thread-level parallellisme en bufferhergebruik te maximaliseren. De Fujisaki-Okamoto (FO) transform equality checks worden eveneens volledig op de GPU uitgevoerd, waarbij slechts een enkele statusflag naar de host wordt geretourneerd.

Experimentele Opstelling
De studie evalueerde de implementatie over vier verschillende accelerator-architecturen:

NVIDIA: A100-SXM-64GB (PCIe 4.0, HBM2e) en GH200 Grace Hopper (NVLink-C2C, HBM3).
AMD: MI300X (PCIe 5.0, HBM3) en MI300A (Unified APU, gedeelde HBM3).

Prestaties werden gemeten tegenover een multicore OpenMP CPU baseline en vergeleken over platforms met variërende lattice dimensies ( $n$ ) tot 16.384. De gebruikte metrieken waren executietijd, schaalbaarheid, energy-to-solution, en gedetailleerde runtime profiling met vendor tools (Nsight Systems en ROCm Systems Profiler).

Belangrijkste Resultaten

Acceleratie vs. CPU: OpenMP Target offloading demonstreerde substantiële acceleratie ten opzichte van de multicore CPU baseline. Voor grote lattice dimensies (bijv. $n=4096$ ) overstegen de speedups twee ordes van grootte (120x), primair omdat de workload memory-bound is, en GPU's superieure geheugenbandbreedte en parallellisme bieden voor streaming matrix operaties.
Cross-Platform Prestaties:
- NVIDIA GH200 behaalde de beste algehele prestaties, toegeschreven aan de hoge HBM3 bandbreedte (4 TB/s) en de NVLink-C2C interconnect, die CPU-GPU communicatie-overhead minimaliseerden.
- AMD MI300X volgde nauw, waarbij het bijna optimale prestaties liet zien voor kleinere lattice dimensies en sterke schaalbaarheid, hoewel het voor grotere $n$ iets hogere executietijden vertoonde vergeleken met de GH200.
- NVIDIA A100 vertoonde gematigde vertragingen ten opzichte van de GH200, beperkt door de lagere HBM2e bandbreedte en PCIe communicatie-overhead.
- AMD MI300A vertoonde de zwakste schaalbaarheid en hoogste executietijden. Profiling onthulde dat, ondanks het hebben van een high-bandwidth unified memory, het gedeelde HBM3 subsystem contentie creëerde tussen CPU-zijde operaties (SHA3 hashing, FO key derivation) en GPU kernels, wat de effectieve bandbreedte met ongeveer 38% verminderde vergeleken met de discrete MI300X.
Energie-efficiëntie: De NVIDIA GH200 bleek ongeveer 2,5x energie-efficiënter dan de AMD MI300X voor grote lattice groottes, gedreven door zowel kortere executietijden als lager gemiddeld stroomverbruik.
Portabiliteit: De enkele broncode werd succesvol uitgevoerd op alle geteste platforms zonder architectuur-specifieke code-modificaties, wat de haalbaarheid van OpenMP Target voor draagbare PQC acceleratie valideert.

Betekenis en Claims
Het artikel claimt dat draagbare GPU-acceleratie een levensvatbare strategie is om de computationele overhead van PQC te verminderen zonder vendor lock-in te veroorzaken. De studie toont aan dat:

Broncode-portabiliteit bereikbaar is: OpenMP Target, gecombineerd met een HIP-geïntegreerde RNG backend, stelt een enkele codebase in staat om heterogeneous GPU-ecosystemen effectief te benutten.
Geheugenarchitectuur de prestaties dicteert: Voor memory-bound LWE workloads is de piek rekenkracht minder kritiek dan geheugenbandbreedte, geheugenhiërarchie-organisatie en de efficiëntie van de host-device interconnect. De resultaten benadrukken dat unified memory architecturen (zoals de MI300A) niet automatisch superieure prestaties garanderen als CPU en GPU verkeer strijden om dezelfde geheugenbronnen.
Praktische implementatie: Door het vermijden van CUDA-specifieke afhankelijkheden, faciliteert deze aanpak de implementatie van quantum-resistente cryptografische infrastructuren over diverse hardware-omgevingen.

De auteurs concluderen dat hoewel OpenMP Target effectieve broncode-portabiliteit biedt, de runtime prestaties afhankelijk blijven van de volwassenheid van compiler toolchains, het gedrag van interopererende device libraries (met name met betrekking tot dynamische geheugenallocatie in de HIP RNG adapter), en de specifieke geheugenhiërarchie van de doelhardware.

Technische Samenvatting: Draagbare Versnelling van Learning With Errors KEMs voor Post-Quantum Cryptografie

Meer zoals dit