GPU Acceleration of Learning With Errors KEMs Using OpenACC for Post-Quantum Cryptography∗ Tiziana Liberati1,3 , Nitin Shukla2 , Matteo Barbieri1,4,5 , Gabriella Bettonte1 , Elisabetta Boella1 , Simone Rizzo1 , Daniele Gregori1 , and Marco Pedicini3 E4 Computer Engineering SpA, via Martiri della Libertà 66, 42019 Scandiano (Italy) {tiziana.liberati, matteo.barbieri, gabriella.bettonte, elisabetta.boella, simone.rizzo, daniele.gregori}@e4company.com 2 SuperComputing Applications and Innovation Department, Cineca, via Magnanelli 6/3, 40033 Bologna (Italy) [email protected] 3 Dipartimento di Matematica e Fisica, Roma Tre University, Largo San Leonardo Murialdo 1, 00146 Rome (Italy) [email protected] 4 Dipartimento di Fisica e Astronomia “G. Galilei”, Università di Padova, via F. Marzolo 8, 35131 Padova (Italy) 5 INFN, Sezione di Padova, Italy
arXiv:2606.01211v1 [cs.CR] 31 May 2026
1
Abstract. Shor’s algorithm proved that asymmetric cryptographic protocols based on the integer factorization and discrete logarithm problems are no longer safe in a world with large-scale quantum computers. As a result, Post-Quantum Cryptography (PQC) has been developed over the last few years, seeking cryptographic primitives resistant to quantum attacks. One of the main hard problems underlying PQC schemes is the Learning with Errors (LWE) problem, which is significantly more computationally intensive than its classical predecessors. In this work, we present a Key Encapsulation Mechanism (KEM) based on plain LWE and develop a GPU-oriented implementation using OpenACC. We evaluate the performance of our accelerated application in terms of both time-to-solution and energy-to-solution, considering bare-metal and containerized executions across multiple NVIDIA GPU models and generations. Our implementation achieves significant acceleration across all tested GPU platforms. In particular, on the NVIDIA Grace Hopper Superchip, it attains up to a 208× speedup over a multithreaded CPU baseline and enables the execution of problem sizes that are impractical on CPU architectures due to memory and synchronization constraints. Energy consumption analysis also shows ≈ 2× better efficiency when using the Superchip compared to systems equipped with x86-based CPUs and NVIDIA H100 GPUs. These results highlight the effectiveness of GPU acceleration for computationally demanding LWE-based cryptographic workloads. Keywords: Post-Quantum Cryptography · Key Encapsulation Mechanism · Learning with Errors · Parallel Computing · OpenACC acceleration.
1
Introduction
Shor’s algorithm [13] solves the integer factorization and the discrete logarithm problems in polynomial time with a quantum computer, threatening widely used public-key cryptosystems, such as Rivest–Shamir–Adleman (RSA) and Elliptic-Curve Diffie-Hellman. In response, the National Institute of Standards and Technology (NIST) has initiated standardization of Post-Quantum Cryptography (PQC) protocols. In this realm, lattice-based schemes are widely studied due to their security and flexibility, particularly those based on Learning With Errors (LWE) for key encapsulation mechanisms (KEMs) [10]. Albeit being secure from known quantum attacks, they introduce significant computational overhead from large-scale matrix operations and randomness generation [9]. To address these limitations, PQC schemes have been accelerated on GPUs leveraging their parallel processing capabilities, particularly in lattice-based constructions. Implementations of LWE-based ∗
Pre-print version of the manuscript submitted to NUMTA2026
2
T. Liberati et al.
schemes such as FrodoKEM, NewHope, and Kyber [7], as well as other lattice-based cryptographic primitives [12,7], have demonstrated substantial speedups through optimized memory access patterns, efficient data transfer, and data-parallel execution strategies. However, most existing approaches rely on CUDA-specific optimizations. In this work, we propose a KEM based on plain LWE and investigate a directive-based heterogeneous programming approach using OpenACC. This method enables existing CPU code bases to leverage GPUs with minimal modifications, thereby improving developer productivity while preserving code readability and long-term maintainability. Although OpenACC currently provides more limited portability than alternative heterogeneous programming models, it offers a favorable trade-off between programmability and performance [14]. Moreover, the maturity of the NVIDIA compiler software stack now enables high performance, in some cases approaching or even exceeding CUDA [16]. Furthermore, thanks to the straightforward interoperability between OpenACC and CUDA [5], we effortlessly integrate a CUDA-based modular Advanced Encryption Standard (AES)-based random number generator (RNGonGPU6 ) to produce noise directly on the GPU, reducing CPU-side overhead during the main KEM operations and thus mitigating one of the primary bottlenecks in LWE-based cryptosystems. In addition, we introduce an intra-operation batching strategy that exploits fine-grained parallelism within KEM operations by supporting bit-level parallel execution. Unlike inter-instance batching used in prior work, this method improves resource utilization within a single KEM instance while mitigating kernel-level inefficiencies. Overall, our implementation achieves speedups of up to 208× compared to the CPU OpenMP version, mainly due to the parallelization of the matrix–vector multiplications. Importantly, it also enables the efficient handling of problem sizes that are impractical to process on CPUs. Besides comparing against the CPU implementation, this work also evaluates the performance of the proposed approach across multiple GPU generations and models, analyzing how architectural characteristics influence execution time and energy consumption. Thus, the main contributions of this paper are: – An open-source implementation of a plain LWE-based KEM for CPU and an OpenACC accelerated version, enabling reproducible experiments and benchmarks. – A set of GPU-oriented optimizations, applicable to a broad class of LWE-based cryptographic schemes, along with an evaluation and discussion of the main optimization strategies used to address the primary performance bottlenecks. – A comprehensive performance, energy-efficiency, and scalability study across different parameter sets, problem size, and NVIDIA GPU architectures. This paper is organized as follows. Section 2 presents implementation details of the post-quantum cryptosystem, with emphasis on its parallelization using OpenMP on multicore architectures and on the choice of valid parameter sets. Section 3 describes the strategy adopted to extend the implementation with GPU support. Section 4 reports performance, energy consumption, and scalability results for the proposed implementation across multiple NVIDIA GPU architectures. Finally, Section 5 summarizes the main findings and discusses their implications for high-performance GPU software development. All performance evaluations and benchmark reproductions in this work are based on the implementation hosted in the LWE-KEM Git repository7 .
2
Implementation and Parallelization of the Post-Quantum Cryptosystem
This section introduces the mathematical and cryptographic foundations underlying the proposed plain LWE KEM. We first briefly review the security assumptions of lattice-based cryptography, which motivate the large adoption of lattice problems in modern post-quantum cryptographic schemes. We then 6 7
RNGonGPU library Git repository: https://github.com/Alisah-Ozcan/RNGonGPU/ LWE-KEM Git repository: https://github.com/TizianaLiberati/LWE-KEM/.
GPU Acceleration of Plain LWE KEMs
3
introduce the notion of Key Encapsulation Mechanisms (KEMs) together with the underlying publickey encryption construction on which KEMs rely. Subsequently, we present a KEM based on the LWE problem and the main algorithms composing the scheme. We also discuss the parameter choices adopted in this work and the validation methodology used to ensure correct decapsulation behavior across both CPU and GPU implementations. Finally, we consider the OpenMP multicore parallelization and highlight the bottlenecks of the scheme; motivating the adoption of GPU acceleration strategies that will be analyzed later on. 2.1
Mathematical Framework
The proposed KEM is based on the Learning With Errors (LWE) problem. Intuitively, LWE can be viewed as solving noisy linear system modulo q: the small error term e hides the linear relation between the public matrix A and the secret s, making the recovery of the secret computationally hard. Formally, in the LWE setting, given a random matrix A ∈ Zn×n , a secret vector s ∈ Znq , and a small q n error vector e ∈ Zq , where n is the lattice dimension, one observes samples of the form b = As + e
(mod q).
(1)
The security of the scheme relies on the difficulty of either recovering s from such samples (search-LWE) or distinguishing valid samples from random values (decision-LWE) [11]. LWE is a central hardness assumption in lattice-based cryptography due to its worst-case to average-case reductions [11], which relate its difficulty to lattice problems such as GapSVP, BDD and SIVP. These hardness assumptions are the foundation on which many post-quantum protocols are built. During the NIST standardization process for PQC, two main primitives were considered: digital signatures and key encapsulation mechanisms. In this work, we focus on KEM constructions, which are a type of protocols that enable secure establishment of shared symmetric keys over untrusted channels. A KEM can be constructed from an underlying public-key encryption (PKE) scheme as shown in Figures 1 and 2. It is composed of three main algorithms: KeyGen, which generates the secret and the public key (sk and pk) (Algorithm 1); Encrypt, which encrypts the message m (a pseudorandom bit string) into a ciphertext c (Algorithm 2); Decrypt which recovers the encrypted message m from the ciphertext (Algorithm 3). The KEM extends this PKE protocol through Encaps (Algorithm 4) and Decaps (Algorithm 5). Encaps generates both the ciphertext and the shared secret key, while Decaps decrypts the ciphertext, recomputes it deterministically, and verifies its correctness through re-encryption. This mechanism is implemented using the Fujisaki-Okamoto (FO) transform which upgrades the scheme to IND-CCA2 security [6]. If the recomputed ciphertext matches the received one, the shared secret is derived from the valid pair (Figure 3); otherwise, implicit rejection outputs a fallback pseudorandom value from the secret key and ciphertext, ensuring constant-time decapsulation without leaking information about ciphertext validity. In our implementation, this appears as: (i) explicit re-encryption verification, (ii) hash-based fallback secret derivation, and (iii) a constant-time decapsulation path independent of ciphertext validity. 2.2
Correctness and Parameter Validation
The correctness of the proposed KEM is determined by the choice of the LWE parameters, namely the lattice dimension n, the modulus q, and the adopted noise distributions. Intuitively, decryption succeeds only when the accumulated noise remains sufficiently small compared to the modulus; otherwise, the recovered value may exceed the decoding threshold and lead to an incorrect message reconstruction. In this work, the secret key distribution is fixed to a binomial distribution with parameter η = 3, a standard and well-studied choice in lattice-based cryptography [10]. The remaining parameters (n, q, σ, k) are varied to balance decryption correctness and computational hardness. Table 1 displays
4
T. Liberati et al.
Algorithm 1 KeyGen()
Algorithm 2 Encrypt(pk, m)
Algorithm 3 Decrypt(sk, c)
1: A ← sample(ρ) ∈ Z 2: s ← χn η 3: e ← χn η 4: t ← As + e (mod q) 5: return pk = (A, t), sk = s
1: e2 ← Uk 2: r ← χn η 3: e1 ← χn η 4: u ← AT r + e1 (mod q) 5: v ← tT r + e2 + m (mod q) 6: return c = (u, v)
1: µ ← v − sT u (mod q) 2: if µj ≤ q/4 or µj ≥ 3q/4 then 3: mj ← 0 4: else 5: mj ← q/2 6: end if 7: return m
n×n q
Fig. 1: Pseudo-code of the underlying LWE-based public key encryption scheme, showing key generation, encryption, and decryption algorithm.
Algorithm 4 Encaps(pk)
Algorithm 5 Decaps(sk, c)
1: m ← {0, 1} 2: c ← Encrypt(pk, m) 3: K ′ ← H(H(pk) ∥ m) 4: K ← H(K ′ ∥ H(c)) 5: return (c, K)
1: m′ ← Decrypt(sk, c) 2: c′ ← Encrypt(pk, m′ ) 3: K ′ ← H(H(pk) ∥ m′ ) 4: if c′ = c then 5: K ← H(K ′ ∥ H(c)) 6: else 7: K ← H(z ∥ H(c)) 8: end if 9: return K
256
Fig. 2: Pseudo-code of the LWE-based key encapsulation mechanism, showing encapsulation and decapsulation using the Fujisaki–Okamoto transform.
Table 1: Selected LWE parameters and estimated hardness (i.e., security level) n
q
512 3329 1024 12289 2048 12289 4096 12289 8192 65537 16384 65537 32768 65537 65536 65537 131072 65537
Gaussian σ Centered uniform k Binomial η Estimated (e, r, e1 ) (e2 ) (s) hardness 2.3 2.3 2.3 1.7 1.2 1.2 1.0 1.0 1.0
3 3 3 3 2 2 2 1 1
3 3 3 3 3 3 3 3 3
141 257 559 1230 2323 5070 10990 23699 50855
the parameter sets considered in our evaluation, including lattice dimension, modulus, noise parameters, and the corresponding approximate hardness. The selected noise distributions follow common practices in LWE-based schemes, providing wellunderstood hardness properties while preserving reliable decryption behavior, and their approximate hardness was assessed using the fast LWE parameter selection tool by Biasoli et al. [4], under simplified LWE model with binary secrets and Gaussian noise. The reported values in Table 1 should therefore be interpreted as indicative hardness levels, rather than exact KEM bit-security claims.
GPU Acceleration of Plain LWE KEMs
Alice (receiver)
5
Bob (sender) KeyGen
Decapsulation key
Encapsulation key
Decaps
Encaps ciphertext
Alice's shared secret
Bob's shared secret
K'
K
K' = K when decapsulation succeds
Fig. 3: Key establishment using a KEM
Overall, the chosen parameter sets are intended to ensure reliable decapsulation in our experiments while spanning a range of increasing, approximate hardness levels. 2.3
Bottlenecks and Parallelization Approach
Our baseline implementation is written in C++ and parallelized using OpenMP to exploit multicore architectures, addressing the high computational cost of lattice-based post-quantum cryptographic schemes. For LWE-based KEMs, the dominant contribution to the overall computational cost arises from the matrix–vector multiplications in Equation (1). This operation is executed in all three KEM phases (KeyGen Algorithm 1, Encaps Algorithm 4, and Decaps Algorithm 5) and dominates runtime due to the quadratic memory traffic associated with the public matrix A in Equation (1). Additional costs arise from dot-product reductions, matrix transposition, and the generation of random noise. In our plain LWE-KEM additional overhead is introduced by processing 256 independent message bits separately in both encapsulation and decapsulation. As the problem size increases, these costs lead to high latency and limited scalability, especially in server-side settings with many concurrent key exchanges, where the CPU becomes a throughput bottleneck. At the same time, these operations (matrix generation, matrix–vector multiplication, dot products, per-bit ciphertext processing, and noise generation) are inherently data-parallel making them well suited for parallel execution. Therefore, our optimization strategy begins with shared-memory parallelism using OpenMP and is later extended to GPU offloading via OpenACC, to better exploit large-scale thread parallelism and memory bandwidth.
3
GPU Porting and Optimization of LWE-KEM
This section presents our OpenACC-based GPU implementation of the proposed plain LWE-KEM. We begin from a naive OpenACC baseline, in which the relevant computational loops are offloaded to the GPU while random numbers are generated on the host and subsequently transferred to the device. Building upon this baseline, we progressively introduce a series of refinements, including intraoperation batching, device-resident random number generation, explicit data management, and kernel parallelization strategies. Before presenting these techniques, we motivate the main design choices and identify the primary performance bottleneck. The section concludes with a detailed performance analysis. All performance results reported in this section were obtained on a Booster node of CINECA’s Leonardo supercomputer, whose hardware and software specifications are provided in Appendix A.
6
T. Liberati et al.
3.1
Design Rationale
The optimization strategy is driven by the bandwidth asymmetry between host and device memory systems. PCIe Gen4 provides approximately 32 GB/s bandwidth, whereas NVIDIA A100 global memory bandwidth reaches approximately 1.5 TB/s. This yields an effective ratio of approximately 50:1. Three architectural principles follow from this observation: 1. Device-resident random number generation: all randomness is generated directly on the GPU, eliminating bulk host-to-device transfer. 2. Explicit data management: OpenACC enter data and exit data directives are used for deterministic control over host-device transfers. 3. Batched kernel execution: all 256 component encryptions are fused into a single kernel launch to amortize invocation overhead. In addition, the public matrix A ∈ Zn×n and the associated working buffers are managed to q minimize unnecessary host-device communication and temporary allocation overhead during largescale executions. 3.2
Intra-Operation Batching
A significant inefficiency in the naive OpenACC baseline originates from sequential processing of 256 independent encryption instances. Each instance triggers separate memory transfers and kernel launches. The baseline execution model is characterized as follows: CPU noise generation → Host-to-device transfer → GPU encryption and repeated independently for each of the 256 message bits. This design introduces excessive synchronization and launch overhead. To mitigate this limitation, an intra-operation batching strategy is introduced. All noise vectors are precomputed and transferred in a single host-to-device operation. CPU seed generation → Host-to-device transfer → Noise generation on GPU → Single batched GPU kernel for 256 encryptions For n = 4096, where n is the lattice rank, the batched noise representation includes 256 × n elements. This corresponds to approximately 8 MB per encapsulation. The GPU kernel is executed once per batch, reducing launch overhead from 256 invocations to a single invocation. The resulting execution exposes 256 × n independent work items that can be executed concurrently by the OpenACC implementation.. 3.3
Device-Resident Random Number Generation
For n = 4096, each encapsulation requires approximately 2.1×106 random samples. Transferring these as 32-bit integers over PCIe Gen4 (∼16 GB/s effective) introduces approximately 0.5 ms of latency per call, which is comparable to the GPU compute time itself. Moreover, CPU-sequential generation places noise generation on the critical path, leaving the device idle while randomness is produced. To eliminate both the transfer and serialization bottlenecks,we employ the RNGonGPU library for fully on-device pseudorandom sampling. Since RNGonGPU implements an AES-256 Deterministic Random Bit Generator (DRBG) compliant with NIST SP 800-90A, it provides cryptographic-grade randomness together with forward secrecy and backtracking resistance. Consequently, only a 64-bit seed is transferred from host to device at initialization, while all subsequent pseudorandom expansion is performed on-device. This reduce the per-encapsulation host-to-device transfer volume by approximately six orders of magnitude (from ∼8 MB to 8 bytes). As a result, this produces a throughput on
GPU Acceleration of Plain LWE KEMs
7
the NVIDIA A100 that exceeds 109 samples/s, compared to ∼ 107 samples/s for sequential OpenSSL, results in a two-order-of-magnitude advantage. Furthermore, RNGonGPU generates identical pseudorandom sequences across GPU architectures (Volta, Ampere, Hopper, ...), a property that is necessary for the correct functioning of the KEM. The library is integrated into OpenACC parallel regions via #pragma acc routine. Per-thread generator state is derived deterministically from a master seed, with one unique sub-seed per output element. 3.4
Explicit Data Management
Reliance on CUDA Unified Memory is avoided due to its non-deterministic page-migration behavior under irregular access patterns. Device allocations are instead established once at program initialization via #pragma acc enter data create(...) directives. All subsequent kernels operate on these device-resident buffers. The present clause is used to enforce consistency and suppress implicit data movement. Host-device synchronization is performed only at two well-defined points: prior to CPU-side cryptographic hashing (device → host) and after CPU-side message generation (host → device). This approach ensures that transfers are initiated only when necessary and with strictly bounded scope. 3.5
Kernel Parallelization Strategy
To maximize performance, all the operations described in Subsection 2.3 are executed directly on the GPU. In the following, we briefly describe the main optimization strategies adopted in our implementation and the phases of the scheme that were primarily parallelized. Additional implementation details and kernel-level optimizations are available in the public repository accompanying this work. We recall that, in KEM constructions, the encapsulation and decapsulation procedures rely on the encryption and decryption algorithms as subroutines for the underlying PKE scheme. Consequently, the optimization of the PKE phases (key generation, encryption, decryption) directly impact the KEM performance. In the following, we describe how these phases were parallelized to improve the overall performance of the scheme. Key Generation: Secret vector sampling from the centered binomial distribution χn3 is fully parallelized, since each component can be generated independently and directly mapped to GPU threads. The computation of the public vector t = As + e assigns one row per thread. The corresponding inner products are evaluated sequentially within each thread in order to preserve deterministic generation of Ai,j directly on the on the GPU from a seed. A fully vectorized implementation of this loop was intentionally avoided to prevent redundant recomputation overhead. Encryption and Decryption: The encapsulation phase processes all 256 message bits within a single GPU kernel. Parallelism is exposed hierarchically by assigning independent message encryptions across OpenACC gangs; while the ciphertext components are computed concurrently across vector lanes. For a lattice dimension of n = 4096, the total thread count per launch is more than a million concurrent GPU threads (256 × 4096 = 1 048 576). The decapsulation phase follows an analogous strategy, where the computation of µ = v − s · u (mod q) is parallelized using the same hierarchical decomposition and reduction-based accumulation scheme via an identical gang/vector decomposition with a parallel reduction. 3.6
Performance Analysis
System-level behavior was profiled using NVIDIA Nsight Systems 2025.2 with CUDA, NVTX, and OpenACC tracing enabled, while kernel-level microarchitectural characteristics were obtained using NVIDIA Nsight Compute with full metric coverage, including memory subsystem activity, warp
8
T. Liberati et al.
(a) Nsight Systems timeline of the naive LWE-KEM implementation (N = 10, n = 1024) on the NVIDIA A100. Execution is dominated by host-side computation and PCIe data transfers, resulting in severe GPU starvation and a single-iteration time of approximately 2.787 s.
(b) Nsight Systems timeline of the optimized LWE-KEM implementation (N = 10, n = 1024) on the NVIDIA A100. Host-device communication is restricted to an 8-byte seed transfer, enabling sustained GPU activity (∼98% of iteration time) and reducing the single-iteration time to 8.24 ms.
Fig. 4: Comparison of LWE-KEM execution timelines on the NVIDIA A100
scheduling behavior, and execution efficiency. This dual-level profiling enables a consistent separation between system-wide execution structure and low-level kernel behavior. The workload, i.e. the number of keys generated, was fixed at N = 10 KEM instances with lattice dimension n = 1024, ensuring a controlled and representative comparison between implementations under identical input conditions. As shown in Figure 4, a substantial performance improvement is observed in the optimized implementation. Total execution time is reduced from 27.869 s in the naive OpenACC baseline to 0.082 s, corresponding to a speedup of approximately 338×. This improvement cannot be attributed to a single factor but emerges from the combined effects of reduced host-device communication, improved kernel-level parallelism, and more efficient utilization of GPU resources. In the naive implementation, execution is largely dominated by host-side processing. Matrix generation and randomness construction are performed on the CPU prior to kernel invocation, during which the GPU remains idle for significant periods. In addition, for each encapsulation, large data
GPU Acceleration of Plain LWE KEMs
9
structures, including public matrices and noise vectors, are transferred from host to device. These repeated transfers introduce substantial latency due to PCIe bandwidth constraints and lead to a fragmented execution timeline, where short kernel executions are interleaved with comparatively long CPU-bound phases. In contrast, the optimized implementation shifts the majority of computation to the device. Only a small seed value is transferred from host to GPU, while matrix elements and random samples are generated directly on-device. As a result, repeated bulk transfers are eliminated, and kernel execution proceeds in a more continuous manner. Nsight Systems traces confirm sustained GPU activity with reduced idle intervals, indicating improved overlap of computation and memory operations at the system level. Further insights are obtained from Nsight Compute metrics. The naive implementation exhibits low occupancy, with a limited number of active warps per streaming multiprocessor. This restricts the ability to hide memory latency effectively. Moreover, memory accesses are frequently uncoalesced, resulting in inefficient use of global memory bandwidth. Warp stall analysis indicates that the dominant source of inefficiency arises from memory dependency stalls, consistent with a memory-bound execution profile. The optimized implementation improves both occupancy and memory efficiency. A larger number of independent threads is exposed through batching and restructuring of computation, leading to higher warp activation and improved latency hiding. Memory access patterns are also significantly improved, with a higher proportion of coalesced transactions and better cache utilization. Consequently, effective memory throughput increases and global memory pressure is reduced. A change in stall behavior is also observed. In the naive implementation, stalls are primarily attributed to memory latency. In the optimized version, the dominant stall reasons shift towards synchronization-related delays, suggesting that memory is no longer the primary bottleneck. Overall, the execution transitions from a memory-bound regime in the naive case to a more compute- and synchronization-limited regime in the optimized implementation.
4
Results and discussion
In this section, we present the results obtained by running the complete KEM workflow across OpenMP and OpenACC, focusing on the scalability between different GPUs. The performance of the GPUaccelerated LWE-KEM implementation was evaluated against the multicore CPU version, corresponding to the OpenACC and OpenMP branches of our repository. The systems used to gather performance data for this section are detailed in Appendix A. All performance experiments in this section, unless specified otherwise, were conducted on a containerized environment based on the NVIDIA HPC SDK image version 24.3. The Dockerfile used to build the container is available inside the main branch of the LWE-KEM Git repository. All time-to-solution measurements (hereafter used synonymously with execution time) were performed using the C++ steady_clock from the chrono library with timing calls embedded in the application. All experiments were conducted using the validated parameter sets discussed in Section 2, with identical parameters for the OpenMP and OpenACC implementations. Since both implementations perform equivalent modular arithmetic and noise sampling, any decryption failure is due to the chosen parameters rather than to hardware-specific effects. 4.1
CPU and GPU scalability tests
As established in Section 2, the dominant computational kernels scale quadratically with the LWE dimension n, making them the primary contributors to runtime for larger problem sizes. This behavior is consistent with the performance trends reported in Table 2. As n grows, the computational workload increases rapidly, allowing the GPU implementation to better leverage its massive parallelism and high memory bandwidth. Consequently, the OpenACC implementation exhibits superior scalability compared to the OpenMP implementation, achieving higher speedups for larger problem sizes. Both
10
T. Liberati et al.
Table 2: Performance comparison between the multicore OpenMP implementation (72 threads) and the OpenACC implementation on the GH200 for N = 100 iterations. The table reports the average execution times of the KeyGen, Encaps, and Decaps phases (in µs), as well as the total workflow runtime (in s). The performance gain of OpenACC over OpenMP is quantified as Speedup = TOpenMP /TOpenACC , which increases with n.
OpenMP n
KeyGen
Encaps
Decaps
OpenACC Total KeyGen Encaps Decaps Total Speedup
32 881.83 274.04 147.49 0.13 128 7914.52 674.41 552.98 0.92 256 30232.35 1490.97 1375.53 3.32 512 116923.74 4064.75 3958.97 12.52 1024 477595.27 12902.04 12750.04 50.37 2048 2001914.9 46668.94 46803.2 209.63 4096 7742431.87 171860.34 172880.33 809.01
3199.23 3212.22 3218.47 3305.69 3636 4622.99 8576.53
5740.13 5952.51 6258.3 7176.59 8121.1 11175.6 19668.7
5575.59 5579.62 5582.86 5997.84 5785.05 6755.38 10551.8
1.46 1.48 1.51 1.65 1.76 2.26 3.89
0.09 0.62 2.2 7.57 28.61 92.68 208.13
the OpenMP and OpenACC benchmarks were executed on the same GH200 compute node to ensure a fair comparison under identical hardware and software conditions. The OpenMP experiments were performed using 72 threads, corresponding to all physical cores available on the Grace CPU. Figure 5 presents the scaling behavior for a fixed problem size of n = 4096. The execution time increases approximately linearly with respect to the workload parameter N for both implementations. The OpenMP-based CPU implementation (blue) consistently exhibits higher execution times, whereas the OpenACC-based GPU implementation (green) maintains substantially lower values across all tested workloads. The roughly constant separation between the curves suggests that the relative speedup remains stable as N increases. These results indicate that the GPU implementation consistently outperforms the CPU implementation, with both implementations exhibiting approximately linear scaling with respect to the workload size.
Scaling Behaviour for Fixed n = 4096
Execution Time (s)
104
OpenMP (CPU) OpenACC (GPU)
103 102 101 100 101
102 Workload Size N
103
Fig. 5: Scaling behavior for fixed n = 4096, showing execution time as a function of workload size N for OpenMP (CPU, blue) and OpenACC (GPU, green) implementations.
GPU Acceleration of Plain LWE KEMs
4.2
11
GPU performance
We assessed the performance of the optimized GPU implementation in terms of time-to-solution and energy-to-solution on x86 + Hopper and GH200 nodes featuring the most recent NVIDIA GPU models available to us. For these tests we used n = 32768 and N = 1000, which correspond to the largest problem sizes that our GPUs could accommodate. On both systems, the code was compiled using the NVHPC 25.7 compiler suite. CUDA 12.9 was used for the RNG library together with GCC 8.5.0 on the x86 + Hopper system and GCC 11.3.0 on GH200. Each simulation was repeated 20 times to obtain sufficient statistics, with runs executed in batches and separated by 300-second pre- and post-run sleep intervals to ensure the system returned to idle between simulations. The top row of Figure 6 shows the resulting time-to-solution distributions over 20 runs on the x86 + Hopper system (left) and the GH200 node (right). The reported times do not include the pre- and post-run sleep intervals. On the x86 + Hopper system, the average execution time was 1119.9 ± 2.6 s, while on Grace Hopper it decreased to 943.2 ± 2.4 s, yielding a speedup of 1.19×. This improvement is likely due to the Grace Hopper Superchip unified memory design and higher memory bandwidth. Energy-to-solution was computed by summing the energy consumed by the GPU and the CPU during each simulation, excluding the sleep intervals. For both systems, GPU energy consumption was estimated from the power measurements provided by nvidia-smi. Following the methodology described in [3,5], GPU power was sampled in the background from user space throughout each simulation and logged to a text file at a frequency of 1 Hz. The GPU energy was then obtained by numerically integrating the power measurements over the simulation time. For the x86 + Hopper system, CPU energy consumption was measured using perf stat -a -e in combination with a one-second sleep interval, following the approach adopted in [1,2]. On GH200, CPU energy was estimated through system-level power monitoring interfaces. CPU power was obtained by reading the file /sys/class/hwmon/hwmon3/device/power1_average once per second. This file reports the output of the sensor placed on the Grace Power Socket 0, averaged over a time window of 50 ms. CPU energy was then computed by numerically integrating the measured power over time. As for the GPU measurements, polling of the CPU power data was performed in the background from user space during the simulations. The distribution of energyto-solution over 20 simulations for both systems is shown in the bottom row of Figure 6. On the x86 + Hopper system, a simulation consumed 420.9 ± 1.0 kJ, whereas the same simulation on GH200 required only 235.9±0.5 kJ, corresponding to an energy saving factor of 1.78×. The reduction in energy consumption, which exceeds the observed time-to-solution improvement, is primarily attributable to the substantially lower CPU power consumption of the ARM-based Grace architecture. 4.3
Cross-Generational GPU Architectural Scaling
Figure 7 presents the total time-to-solution of the LWE-KEM implementation for the largest configuration (n = 32768, N = 1000) across three NVIDIA GPU generations: Volta (V100S), Ampere (A100), and Hopper (GH200). The OpenMP-based CPU implementation was not able to complete this configuration within a 24-hour time limit on the GH200 node, due to cache contention, NonUniform Memory Access (NUMA) penalties, and OpenMP barrier overhead. All GPU experiments were executed in containers to ensure reproducible and consistent comparisons across heterogeneous platforms. Execution times were 3500.08 s for V100S, 2634.69 s for A100, and 1740.39 s for GH200, yielding speedups of 1.33× and 2.01× for A100 and GH200, respectively, relative to the V100S baseline. To better understand these gains, the results were analyzed with respect to three main hardware characteristics: HBM memory bandwidth, cache hierarchy, and host–device communication architecture. Table 3 summarizes the relevant architectural parameters of the evaluated platforms. Nsight Compute profiling (Subsection 3.6) indicates that the LWE-KEM kernel is memory-bound, with stalls primarily driven by long-latency memory dependencies. Consequently, cross-generational performance is expected to scale predominantly with HBM bandwidth rather than computational throughput. The V100S provides 1134 GB/s (HBM2), the A100 provides approximately 1555 GB/s
T. Liberati et al. 0.200
0.40
0.175
0.35
0.150
0.30
0.125
0.25
# sim [arb.un.]
# sim [arb.un.]
12
0.100 0.075
0.15
0.050
0.10
0.025
0.05
0.000 900
950
1000 1050 1100 Time-to-solution [s]
0.00 220
1150
0.200
0.40
0.175
0.35
0.150
0.30
0.125
0.25
# sim [arb.un.]
# sim [arb.un.]
0.20
0.100 0.075
0.025
0.05 1000 1050 1100 Time-to-solution [s]
1150
420
460
260
300 340 380 Energy-to-solution [kJ]
420
460
0.15 0.10
950
300 340 380 Energy-to-solution [kJ]
0.20
0.050
0.000 900
260
0.00 220
Fig. 6: Top row: distribution of time-to-solution (TTS) over 20 simulations for runs with n = 32768 and N = 1000 on the x86 + Hopper system (left) and the GH200 node (right). Bottom row: distribution of energy-to-solution (ETS) over 20 simulations for the same runs and systems. The red dashed lines indicate the average TTS and ETS values for each system.
Table 3: Key architectural characteristics of the GPUs evaluated in the cross-generational comparison. Parameter
V100S
A100
GH200
Microarchitecture Volta Ampere Hopper HBM type HBM2 HBM2e HBM3 HBM bandwidth 1134 GB/s ∼1555 GB/s ∼3352 GB/s L2 cache 6 MB 40 MB 60 MB SM count 80 108 132 CPU–GPU interconnect PCIe 3.0 ×16 PCIe 4.0 ×16 NVLink 4 Interconnect bandwidth ∼16 GB/s ∼32 GB/s 900 GB/s
(HBM2e), and the GH200 provides approximately 3352 GB/s (HBM3), giving bandwidth scaling factors of 1.37× and 2.96× relative to the V100S. The measured speedup on A100 (1.33×) closely matches this scaling, indicating that memory throughput remains the primary constraint. The GH200 achieves
GPU Acceleration of Plain LWE KEMs
Walltime [s]
4000
13
1.0x
3000
1.33x 2.01x
2000 1000 0 V100S
A100
GH200
Fig. 7: Execution times and relative speedups for the different GPU models evaluated in our crossgenerational comparison (n = 32768, N = 1000).
2.01×, which is below the expected 2.96×, suggesting that additional non-bandwidth-limited effects emerge at higher performance levels. In particular, on-demand matrix element generation introduces increasing arithmetic overhead as bandwidth improves, shifting part of the workload away from a purely memory-bound regime, while CPU-side costs from the Fujisaki–Okamoto transform impose a fixed serial component that limits end-to-end scaling under Amdahl’s law [8]. For n = 32768, the working set is approximately 100 MB, comprising large noise buffers and ciphertext data, which exceeds the L2 cache capacity across all architectures (6 MB on V100S, 40 MB on A100, and 60 MB on GH200). As a result, capacity misses to HBM are unavoidable, although larger caches on A100 and GH200 reduce traffic for frequently reused structures such as the public vector t and tiled regions of noise buffers, improving effective bandwidth use. The public matrix is generated on demand via a seed-expansion function rather than stored explicitly, avoiding multi-gigabyte allocations but introducing O(n2 ) arithmetic overhead from per-element computation. This cost becomes increasingly significant as memory bandwidth improves, contributing to sublinear scaling on newer architectures. Additionally, host-device synchronization associated with cryptographic hashing persists across generations, with interconnect improvements (PCIe 3.0, PCIe 4.0, and NVLink 4) reducing but not eliminating transfer overhead. Overall, the cross-generational scaling of LWE KEM performance is governed by the interplay of three factors: (i) HBM bandwidth improvements, which drive the primary gains for the memorybound GPU kernels; (ii) growing relative compute overhead from on-demand matrix generation, which attenuates bandwidth-driven scaling at higher performance levels; and (iii) fixed CPU-side costs from the Fujisaki-Okamoto transform, which establish a diminishing-returns ceiling via Amdahl’s law. These results highlight that for memory-bound crypto workloads with significant non-GPU components, bandwidth improvements give meaningful but eventually saturating gains, and full utilization of nextgeneration architectures requires co-optimization of the full heterogeneous pipeline.
5
Conclusion and Future work
This work presents a KEM based on plain LWE and describes our strategy to accelerate the algorithm using OpenACC. Unlike structured variants such as Ring-LWE and Module-LWE, plain LWE avoids additional algebraic structure beyond standard linear algebra over finite fields. While this choice significantly increases computational requirements, it also relies on more conservative security assumptions. The proposed implementation leverages the inherent parallelism of operations such as matrix–vector products, combining intra-operation batching, explicit data management, and on-device randomness
14
T. Liberati et al.
generation to reduce host-device communication overhead and expose large-scale fine-grained parallelism on different GPU architectures. In addition, by integrating the RNGonGPU library, we generate deterministic, cryptographically secure pseudorandom strings directly on the GPU, further limiting data transfers. Indeed, profiling results (Figure 4a) show that our initial naive implementation without RNGonGPU and the rest of the optimizations is dominated by CPU-side noise generation and hostdevice coordination overheads. In the enhanced implementation (Figure 4b) we managed to reduce memory-transfer overheads, enabling more sustained GPU activity while shifting the performance bottleneck towards synchronization and coordination of host-side operations. Node-to-node benchmarks comparing the OpenMP-parallelized CPU reference implementation with the accelerated GPU code show speedups of up to 208× on the GH200 platform, depending on the value of n (Table 2). In particular, the GPU implementation begins to outperform the multicore CPU baseline for n ≥ 256, where kernel parallelism and memory throughput are sufficient to amortize datamanagement overheads (Table 2). In addition, the proposed approach enabled the execution of large parameter configurations that are impractical on CPUs due to memory contention, synchronization overhead, and limited scalability. We further evaluated the accelerated implementation in terms of execution time and energy consumption across different GPU models and architectures. In particular, the GH200 platform achieves up to a 1.78× reduction in energy consumption compared to the x86 + Hopper configuration, whereas the reduction in execution time is more modest, reaching at most 1.19×. A cross-generation analysis across Volta, Ampere, and Hopper GPUs confirms that memory bandwidth remains the primary limiting factor for plain LWE workloads, while CPU-side operations increasingly constrain scalability as parameters grow. Overall, our results show that the proposed LWE-based KEM and its directive-based GPU implementation provide a practical and efficient path toward scalable post-quantum cryptography on modern heterogeneous platforms, possibly shifting the practicality boundary of conservative plain LWE cryptography. Future work will investigate OpenMP portability across heterogeneous accelerator architectures, particularly AMD and Intel GPUs, to assess its effectiveness on different architectures, as C/C++ OpenACC support for AMD GPUs still lacks mature compilers. We will also extend the implementation to multi-GPU and distributed HPC environments for larger workloads and higher throughput. Additionally, further security and side-channel analyses will be conducted to evaluate robustness in practical deployments. Disclosure of Interests. The authors have no competing interests to declare that are relevant to the content of this article. T.L. developed the OpenMP-parallelized CPU implementation. T.L. and N.S. contributed equally to the GPU porting and optimization of the application. S.R. conducted the GPU experiments in containerized environments. E.B. carried out the energy measurements. All authors contributed to the conception of the study, the analysis and interpretation of the results, and the writing and revision of the manuscript. During the preparation of this article, artificial intelligence tools (ChatGPT and Grammarly) were used exclusively to assist with writing and language editing. All scientific content and analyses are the responsibility of the authors.
A
Testbed Details
Our work was developed and tested on four compute nodes containing Volta-, Ampere-, and Hoppergeneration NVIDIA GPUs. The Volta-based platform consists of two AMD EPYC 7282 processors (SMT enabled, 2.8 GHz), 512 GB DDR4 memory (3200 MT/s), and two NVIDIA V100S GPUs with 32 GB HBM2 (High Bandwidth Memory) each connected through PCIe 3.0 ×16. The system runs Red Hat Enterprise Linux 9.4 with NVIDIA driver version 545.23.08 and utilizes Podman 5.2.2 for containerized execution. The Ampere-based platform, hosted on the Booster partition of the Leonardo supercomputer [15], features an Intel Xeon Platinum 8358 processor (SMT disabled, 2.6 GHz), 512 GB DDR4 memory, and four NVIDIA A100 GPUs with 64 GB HBM2e each connected through PCIe 4.0
GPU Acceleration of Plain LWE KEMs
15
×16. The platform operates on Red Hat Enterprise Linux 8.8 and executes workloads through SingularityPRO 4.3.1. The x86 Hopper platform integrates dual AMD EPYC 9554 processors (SMT disabled, 3.10 GHz), 512 GB DDR5 memory (4800 MT/s), and one NVIDIA H100 PCIe GPU with 80 GB HBM3 connected through PCIe 5.0 ×16. It runs Red Hat Enterprise Linux 8.6. Finally, the Grace Hopper platform employs the NVIDIA GH200 Grace Hopper Superchip, combining an ARM Neoverse V2 CPU (SMT disabled, 3.47 GHz), 480 GB LPDDR5 memory, and a Hopper GPU with 96 GB HBM3 memory. The CPU and GPU are connected through NVLink 4, providing up to 900 GB/s bidirectional bandwidth. This system runs Ubuntu 22.04.5 with Docker 29.0.2 for containerized execution. Across all evaluated systems, the Linux performance CPU frequency governor was enforced to ensure stable measurements.
References 1. Almerol, J.L., et al.: Accelerating gravitational n-body simulations using the risc-v-based tenstorrent wormhole. In: SC25-W: Workshops of the International Conference for High Performance Computing, Networking, Storage and Analysis. pp. 1729–1735 (2025). https://doi.org/10.1145/3731599.3767528 2. Almerol, J.L., et al.: Assessing performance and porting strategies for gravitational n-body simulations on the risc-v-based tenstorrent wormhole™. Astronomy and Computing 56, 101121 (2026). https://doi.org/ https://doi.org/10.1016/j.ascom.2026.101121 3. Amati, G., et al.: Experience on clock rate adjustment for energy-efficient gpu-accelerated real-world codes. p. 245–257. Springer-Verlag, Berlin, Heidelberg (2025). https://doi.org/10.1007/978-3-032-07612-0_19, https://doi.org/10.1007/978-3-032-07612-0_19 4. Biasioli, B., et al.: A tool for fast and secure LWE parameter selection: the FHE case. Cryptology ePrint Archive, Paper 2024/1895 (2024) 5. Boella, E., et al.: Accelerating the particle-in-cell code ecsim with openacc (2026) 6. Fujisaki, E., et al.: Secure integration of asymmetric and symmetric encryption schemes. J. Cryptol. 26, 80–101 (2013). https://doi.org/10.1007/s00145-011-9114-1, https://doi.org/10.1007/s00145-011-9114-1 7. Gupta, N., et al.: Pqc acceleration using gpus: Frodokem, newhope, and kyber. IEEE Transactions on Parallel and Distributed Systems 32(3), 575–586 (2021). https://doi.org/10.1109/TPDS.2020.3025691 8. Gustafson, J.L.: Amdahl’s Law, pp. 53–60. Springer US, Boston, MA (2011). https://doi.org/10.1007/ 978-0-387-09766-4_77 9. Kumar, M.: Post-quantum cryptography algorithm’s standardization and performance analysis. Array 15, 100242 (08 2022). https://doi.org/10.1016/j.array.2022.100242 10. National Institute of Standards and Technology: FIPS 203: Module-Lattice-Based Key-Encapsulation Mechanism Standard. Tech. Rep. FIPS 203, U.S. Department of Commerce, National Institute of Standards and Technology (2024) 11. Regev, O.: On lattices, learning with errors, random linear codes, and cryptography. J. ACM 56(6) (2009). https://doi.org/10.1145/1568318.1568324 12. Shen, S., et al.: High-throughput gpu implementation of dilithium post-quantum digital signature. IEEE Transactions on Parallel and Distributed Systems 35(11), 1964–1976 (2024). https://doi.org/10.1109/ TPDS.2024.3453289 13. Shor, P.: Algorithms for quantum computation: discrete logarithms and factoring. In: Proceedings 35th Annual Symposium on Foundations of Computer Science. pp. 124–134 (1994). https://doi.org/10.1109/ SFCS.1994.365700 14. Stack, M., et al.: Openacc acceleration of an agent-based biological simulation framework. Computing in Science & Engineering 24(5), 53–63 (2022). https://doi.org/10.1109/MCSE.2022.3226602 15. Turisini, M., et al.: Leonardo: A pan-european pre-exascale supercomputer for hpc and ai applications (2023) 16. Vignolo, A., et al.: A tale of two codes: Cuda vs openacc for mass-zero constrained dynamics. The International Journal of High Performance Computing Applications 39(4), 502–518 (2025). https://doi.org/ 10.1177/10943420251331673