💻 computer science

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

This paper presents a portable GPU implementation of a Learning With Errors (LWE) based Key Encapsulation Mechanism (KEM) using OpenMP Target offloading, demonstrating that a single source codebase can achieve substantial performance acceleration and energy efficiency across both NVIDIA and AMD accelerators while avoiding vendor lock-in.

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

Published 2026-07-13

📖 4 min read☕ Coffee break read

CC BY 4.0

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

Original paper licensed under CC BY 4.0 (http://creativecommons.org/licenses/by/4.0/). ✨ This is an AI-generated explanation of the paper below. It is not written or endorsed by the authors. For technical accuracy, refer to the original paper. Read full disclaimer

Imagine you're trying to build a super-secure digital vault that can't be cracked even by a futuristic "quantum" supercomputer. To do this, you need to solve a massive, messy math puzzle called "Learning With Errors" (LWE). It's like trying to find a specific needle in a haystack, but the haystack is made of millions of tiny, noisy magnets, and the more you look, the more the magnets wiggle around.

The problem? Solving these puzzles takes forever on a regular computer. It's like trying to move a mountain of sand one grain at a time with a tiny spoon. The authors of this paper asked: "What if we used a giant, super-fast construction crew (a GPU) to move that sand?"

The Big Discovery: One Code, Two Crews
Usually, if you want to use a construction crew, you have to hire a specific team (like NVIDIA's CUDA) and write instructions only they understand. If you want to hire a different team (like AMD), you have to rewrite the whole instruction manual. That's expensive and annoying.

This paper shows that you can write one single set of instructions (using something called OpenMP Target) that works perfectly for both the NVIDIA crew and the AMD crew. It's like writing a recipe that works whether you're cooking in a kitchen with gas stoves or electric ones, without changing a single ingredient.

The Magic Trick: Staying in the Kitchen
The biggest time-waster in these math puzzles is running back and forth between the main computer (the CPU) and the super-fast GPU. It's like a chef running to the pantry to get salt for every single pinch of seasoning.

The authors figured out how to keep the "salt" (random numbers needed for the math) right inside the GPU kitchen. They upgraded a tool called RNGonGPU so it works with both NVIDIA and AMD crews. Now, the GPU can generate all the random numbers it needs without ever leaving its seat. This keeps the workflow smooth and fast.

The Race: Who Wins?
The team tested their new method on four different super-fast computers:

NVIDIA A100: A powerful, standard super-computer.
NVIDIA GH200: A massive "SuperChip" where the brain (CPU) and the muscle (GPU) are glued together with a super-fast highway (NVLink).
AMD MI300X: A powerful, standard super-computer with a huge memory bank.
AMD MI300A: A chip where the brain and muscle share the exact same memory bank.

Here is what they found:

The Speed Boost: When the math puzzles got huge (with a size of 4,096), the GPU version was 120 times faster than the regular computer version. For even bigger puzzles (size 16,384), the GPU was still blazing fast, while the regular computer was practically asleep.
The Winners: The NVIDIA GH200 was the fastest, finishing the job in about 60 seconds. The AMD MI300X was a close second, taking about 85 seconds. Both of these machines have huge, fast memory banks (HBM3) that can pour data out quickly.
The Surprising Loser: You might think the AMD MI300A would be great because it shares memory between the brain and muscle. But it was actually the slowest, taking 114 seconds. Why? Because the brain and the muscle were fighting over the same water hose. While the GPU was trying to move data, the CPU was also trying to do its own math, and they got in each other's way. It's like two people trying to drink from the same straw at the same time; neither gets a full drink.

Energy and Heat
The team also checked how much electricity these machines used. The NVIDIA GH200 didn't just finish faster; it used less energy to do the job. It took about 9.7 kJ of energy to solve the puzzle, while the AMD MI300X needed about 26.2 kJ. That means the NVIDIA machine was roughly 2.5 times more energy-efficient.

What This Means
The paper proves that you don't have to pick a side between NVIDIA and AMD to get super-fast security. You can use one code base to run on both. However, they also showed that just having a fast chip isn't enough; how the memory is organized matters just as much. If the brain and muscle fight over the same memory, the whole system slows down.

In short, the authors measured that portable, GPU-accelerated security is not just a dream—it's a reality that can make quantum-proof locks much faster and more efficient, as long as you pick the right hardware setup to avoid traffic jams in the memory lane.

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

Problem Statement
The transition to Post-Quantum Cryptography (PQC) necessitates implementations capable of meeting the computational demands of real-world applications. Among PQC candidates, Learning With Errors (LWE)-based Key Encapsulation Mechanisms (KEMs) offer strong security foundations but incur substantial computational overhead. This overhead is primarily driven by large-scale matrix-vector operations and the generation of cryptographically secure random numbers. While Graphics Processing Units (GPUs) are natural candidates for accelerating these parallel workloads, existing GPU implementations largely rely on CUDA-specific optimizations, creating vendor lock-in and limiting portability across heterogeneous accelerator ecosystems (e.g., NVIDIA and AMD). Furthermore, previous portable attempts using OpenACC have shown limited cross-vendor support. The core challenge addressed is achieving high-performance acceleration for dense LWE-based KEMs while maintaining source-level portability across different GPU architectures without sacrificing the ability to keep critical components, such as random number generation, resident on the device.

Methodology
The authors present a portable GPU implementation of a plain LWE-based KEM using OpenMP Target offloading. The methodology involves three primary technical strategies:

Portable Offloading Strategy: The implementation utilizes OpenMP Target directives to manage data movement and kernel execution. To ensure GPU residency and minimize host-device communication, the authors employ persistent device memory management (target enter data map(alloc:...)). This keeps major matrices, ciphertexts, and temporary buffers resident on the GPU throughout the KeyGen, Encapsulation, and Decapsulation phases.
Heterogeneous Random Number Generation (RNG): A critical component of the workflow is cryptographically secure random number generation. The authors extended the existing RNGonGPU library, originally designed for CUDA, to support the Heterogeneous-computing Interface for Portability (HIP). By integrating a HIP-compatible adapter layer with the OpenMP Target workflow, they enable fully GPU-resident random number generation on both NVIDIA and AMD accelerators. This is achieved using #pragma omp target data use_device_ptr to expose OpenMP-managed device buffers to the external HIP-based RNG backend, preventing the need to generate randomness on the host and copy it to the device.
Kernel Parallelization: The computational kernels are parallelized using OpenMP teams distribute parallel for constructs. Key generation involves parallel generation of secret vectors and error vectors, followed by dense matrix-vector products. Encryption and decryption are batched, processing independent message bits or ciphertext slots simultaneously to maximize thread-level parallelism and buffer reuse. The Fujisaki-Okamoto (FO) transform's equality checks are also performed entirely on the GPU, returning only a single status flag to the host.

Experimental Setup
The study evaluated the implementation across four distinct accelerator architectures:

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

Performance was measured against a multicore OpenMP CPU baseline and compared across platforms using varying lattice dimensions ( $n$ ) up to 16,384. Metrics included execution time, scalability, energy-to-solution, and detailed runtime profiling using vendor tools (Nsight Systems and ROCm Systems Profiler).

Key Results

Acceleration vs. CPU: OpenMP Target offloading demonstrated substantial acceleration over the multicore CPU baseline. For large lattice dimensions (e.g., $n=4096$ ), speedups exceeded two orders of magnitude (120x), primarily because the workload is memory-bound, and GPUs offer superior memory bandwidth and parallelism for streaming matrix operations.
Cross-Platform Performance:
- NVIDIA GH200 achieved the best overall performance, attributed to its high HBM3 bandwidth (4 TB/s) and the NVLink-C2C interconnect, which minimized CPU-GPU communication overhead.
- AMD MI300X followed closely, showing near-optimal performance for smaller lattice dimensions and strong scalability, though it exhibited slightly higher execution times for larger $n$ compared to the GH200.
- NVIDIA A100 showed moderate slowdowns relative to the GH200, limited by its lower HBM2e bandwidth and PCIe communication overhead.
- AMD MI300A exhibited the weakest scalability and highest execution times. Profiling revealed that despite having high-bandwidth unified memory, the shared HBM3 subsystem created contention between CPU-side operations (SHA3 hashing, FO key derivation) and GPU kernels, reducing effective bandwidth by approximately 38% compared to the discrete MI300X.
Energy Efficiency: The NVIDIA GH200 was found to be approximately 2.5x more energy-efficient than the AMD MI300X for large lattice sizes, driven by both shorter execution times and lower average power consumption.
Portability: The single source code base successfully executed on all tested platforms without architecture-specific code modifications, validating the feasibility of OpenMP Target for portable PQC acceleration.

Significance and Claims
The paper claims that portable GPU acceleration is a viable strategy for reducing the computational overhead of PQC without incurring vendor lock-in. The study demonstrates that:

Source-level portability is achievable: OpenMP Target, combined with a HIP-integrated RNG backend, allows a single codebase to leverage heterogeneous GPU ecosystems effectively.
Memory architecture dictates performance: For memory-bound LWE workloads, peak compute capability is less critical than memory bandwidth, memory-system organization, and host-device interconnect efficiency. The results highlight that unified memory architectures (like the MI300A) do not automatically guarantee superior performance if CPU and GPU traffic contend for the same memory resources.
Practical deployment: By avoiding CUDA-specific dependencies, this approach facilitates the deployment of quantum-resistant cryptographic infrastructures across diverse hardware environments.

The authors conclude that while OpenMP Target provides effective source-level portability, runtime performance remains dependent on the maturity of compiler toolchains, the behavior of interoperating device libraries (specifically regarding dynamic memory allocation in the HIP RNG adapter), and the specific memory hierarchy of the target hardware.

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

More like this