ArctyrEX: Accelerated Encrypted Execution
of General-Purpose Applications

All encryption schemes that exhibit homomorphic properties enable meaningful computation directly on ciphertext data without revealing the underlying plaintext. The two variants of homomorphic encryption that support functionally complete sets of operations include leveled homomorphic encryption (LHE) and FHE. In both cases, ciphertexts are encoded as tuples of high-degree polynomials and adding or multiplying ciphertexts takes the form of polynomial addition or multiplication. These polynomials typically range from degree $2^{10}$ to $2^{17}$ and the coefficients are integers modulo $q$, which is a product of primes and typically hundreds of bits in length. In the encrypted domain, addition increases the ciphertext noise slightly, while multiplication is significantly more noisy. An unfortunate consequence of this ciphertext noise (which is necessary for security) is that the noise magnitude increases during each homomorphic arithmetic operation, and eventually the noise will corrupt the underlying plaintext message and prevent successful decryption with high probability. LHE can mitigate noise for a finite number of operations using a technique called modulus switching, with larger encryption parameters allowing higher noise tolerance. However, larger parameters entail slower computation and higher memory consumption, which limits scalability for very deep circuits.

FHE solves the scalability issues inherent to LHE and allows for unbounded, arbitrary computation on encrypted data. First realized by Gentry in 2009 (Gentry, 2009), bootstrapping is a noise mitigation technique that can be applied an infinite number of times, unlike modulus switching. In fact, any LHE scheme can be converted to an FHE scheme with the inclusion of bootstrapping. Nevertheless, the bootstrapping procedure itself is costly in terms of latency and remains a key bottleneck of all FHE constructions. Depending on the cryptosystem and chosen parameters, bootstrapping can take anywhere from several milliseconds (Chillotti et al., 2020a) to minutes (Gentry et al., 2012a). Therefore, the only way to achieve feasible FHE for general-purpose computation is to accelerate and optimize the bootstrapping mechanism.

Both the DM (Ducas and Micciancio, 2015) and CGGI cryptosystems (Chillotti et al., 2020a) possess bootstrapping routines that can be evaluated in up to tens of milliseconds on a CPU using modern open-source implementations such as Concrete (Chillotti et al., 2020b) and OpenFHE (Al Badawi et al., 2022), which is faster than other FHE cryptosystems. Also, while other schemes encrypt vectors of integers and floating point numbers, CGGI and DM are typically used to encrypt individual bits into a single ciphertext. Due to this encoding, the core encrypted operations take the form of Boolean gates, which are more flexible in terms of general computation than arithmetic operations over integers (e.g., encrypted comparisons are easily implemented using encrypted bits).

As discussed, to support unlimited computation depths, the FHE scheme must periodically invoke a bootstrapping operation to reset the amount of noise in the ciphertext. In the case of CGGI, which evaluates Boolean gates, bootstrapping must be performed every gate. As a result, evaluating a single homomorphic gate requires on the order of 2,000 polynomial multiplications (Chillotti et al., 2020a), which are typically accomplished using the Discrete Fourier Transform (DFT). While this is an efficient algorithm for a single polynomial multiplication, even a small application could require billions of DFTs. For example, the computation graph for a single inner product of two vectors comprising 16 encrypted 16-bit numbers contains nearly 25,000 encrypted logic gates. Evaluating this circuit results in over 75 million invocations of the DFT. DM (Ducas and Micciancio, 2015) was the first cryptosystem to introduce a functional bootstrap that refreshes noise while simultaneously evaluating a non-linear operation on the encrypted bits. In fact, this bootstrap is a necessary component of the computation for logic gates such as NAND.

CGGI improves upon this construction and generalizes it for all logic gates, including an encrypted MUX that is capable of obliviously choosing between two encrypted bits dependent on the underlying value of an encrypted selector bit. For all gates except the inverter, which is noiseless and linear, the bootstrapping operation comprises the majority of the gate’s latency (Jiang et al., 2022). In turn, the core bottleneck of bootstrapping is the numerous polynomial multiplications between encrypted secret key components and input ciphertexts. Most FHE libraries perform these high-degree polynomial multiplications as element-wise multiplications in the DFT domain, which is asymptotically faster than textbook polynomial multiplication (Chen et al., 2014) (Cormen et al., 2022). Both the number theoretic transform (NTT) and fast fourier transform (FFT) can facilitate the forward and inverse domain conversions for these purposes. However, the NTT is typically preferred over the FFT as it operates directly over integers. Moreover, FFT requires additional type conversions between integers and floating point numbers as FHE ciphertexts contain integer coefficients. As a result, FFT introduces small computation errors due to its reliance on floating point arithmetic.

The CGGI cryptosystem employs different types of ciphertexts, each with different characteristics. The first type, known as LWE ciphertexts, serve as the inputs and outputs of each homomorphic gate evaluation from a user perspective. LWE ciphertexts are the smallest type that CGGI uses; at 128 bits of security, they consist of a single 630-degree polynomial with 32-bit coefficients and an extra 32-bit scalar term. However, these ciphertexts can not be used for nonlinear encrypted operations and are incapable of being used to evaluate a standard encrypted logic gate function (with the trivial NOT gate being the sole exception). Instead, these ciphertexts need to be transformed to TLWE ciphertexts (i.e., Ring-LWE) that are larger in size. Typically, TLWE ciphertexts are composed of a tuple of 1024-degree polynomials with 32-bit coefficients. The third type is TGSW ciphertexts, which are the largest and can conceptually be viewed as an array of TLWE ciphertexts. The bootstrapping key, which is an encryption of the secret key, is composed of this type of ciphertexts. Importantly, TGSW ciphertexts can be multiplied directly with TLWE ciphertexts, which is a necessary step of all bootstrapped gate evaluations. Figure 2 gives a high-level overview of the operations involved in a homomorphic NAND gate. All bootstrapped gates are evaluated in a similar way and only differ in the preliminary linear operations (i.e., the top green box in the figure).

Overall, the CGGI cryptosystem is a good candidate for achieving accelerated general purpose computation on GPUs for a variety of reasons. First, the parameters used by CGGI are often significantly smaller than other FHE schemes, which yields smaller ciphertexts. Thus, multiple ciphertexts can be held in the GPU shared memory simultaneously, which is not always the case for schemes such as CKKS (Cheon et al., 2017), BFV (Fan and Vercauteren, 2012), and BGV (Brakerski et al., 2014) that can utilize ciphertexts on the order of several megabytes (Varia et al., 2015). Notably, certain classes of encrypted operations used for general purpose computation are well-suited for CGGI with binary ciphertexts, but are non-trivial using other cryptosystems that adopt multi-bit encodings. For instance, comparison operations, bitwise manipulations like shifting, and nonlinear functions such as the ReLU activation function in machine learning applications, can be computed directly without the need of costly polynomial approximations (Boura et al., 2019) (Han et al., 2019). Lastly, the requirement of executing hundreds of DFT transforms per bootstrap is particularly well-suited to GPUs due to the parallel nature of FFT and NTT. CGGI can also support multi-bit encodings and employ a special programmable bootstrapping mechanism that evaluates univariate functions. However, only low precision is achievable with realistic parameters and therefore this approach is better suited for specific applications rather than for arbitrary computation. We strictly consider CGGI in gate bootstrapping mode with binary ciphertexts in this work specifically for this reason but note that the methodology proposed is readily extensible to support this scenario.

One of the challenges for achieving efficient encrypted computation with the CGGI cryptosystem involves exploiting circuit-level parallelism at the logic gate level. Essentially, any number of gates with resolved dependencies (e.g., all input wires have been loaded with encrypted ciphertexts) can be executed in parallel as they are entirely independent. For CPU-based systems with a limited number of cores, this parallelism is sufficient to effectively saturate the available CPUs without any significant optimizations at the logic synthesis or application level. However, high-performance computing systems that leverage hundreds of CPU cores or incorporate GPUs require much higher degrees of circuit-level parallelism to achieve high efficiency. For these systems, the characteristics of the underlying Boolean circuit become much more important, therefore avoiding sub-optimal configurations is a critical concern. For example, using the kernels of Figure 4, we present the width of each circuit level for a $10\times 10$ matrix multiplication as well as a logistic regression (LR) inference in Figure 5. The matrix multiplication benchmark represents ideal circuit characteristics for parallel execution as the majority of levels are very wide (the largest being nearly 200,000 gates in width), and the critical path is relatively short. On the other hand, LR inference has over 500 levels (resulting in a much longer critical path) and the width of each level is considerably shorter than those in the matrix multiplication circuit. Another type of circuit configuration ill-suited for systems that can exploit high degrees of parallelism is circuits that adopt cascading. Cascaded circuits typically have a long critical path and each level of the circuit is narrow, limiting the number of gates that can be evaluated in parallel at any given time.

Likewise, not all encrypted gates have the same execution time. For instance, NOT gates are significantly faster than other gates because they don’t require any bootstrapping, while MUX gates are approximately twice as expensive as standard gates (like AND and OR gates). Efficient FHE circuit generation should take into account these differences in gate efficiency.

The conversion process from a C program to an equivalent FHE algorithm can be completed in two distinct steps borrowed from modern hardware design paradigms: high-level synthesis (HLS) followed by logic or register transfer level (RTL) synthesis. While any HLS tool can be used for this purpose, we employ the Google XLS framework (XLS, 2020), which is a fast and efficient open-source HLS engine that can be used to rapidly generate synthesizable Verilog code. This Verilog code serves as an intermediate representation and describes the circuit functionality, which is then transformed by a logic synthesis tool to generate the actual Boolean netlist.

We utilize the open-source Yosys Open Synthesis Suite to facilitate this process and perform crucial circuit-level optimizations (Wolf, ). However, all existing logic synthesis tools, including Yosys, are tailored specifically for physical hardware development and optimize for several constraints that are not relevant to virtual FHE circuits (such as minimizing area or reducing clock cycle latency). The most relevant factors for optimal FHE circuit generation are minimizing the critical path delay, which is luckily a goal shared with actual hardware development, and prioritizing gates that run efficiently in the encrypted domain. Similarly to techniques introduced by the Google Transpiler (Gorantala et al., 2021), we can configure the logic synthesis tool to choose FHE-friendly gates by encoding the relative costs of each gate type as a function of area. For instance, we assign the multiplexer gate to be twice as big as the standard two-input gates to reflect the fact that the latency of the MUX is twice as slow as a standard gate.

Where prior work has adopted generic synthesis scripts for generating netlists for FHE evaluation (Gorantala et al., 2021; Gouert and Tsoutsos, 2020), our synthesis flow: (1) reduces the time required to generate the netlist relative to the Yosys generic synthesis script, and (2) results in more efficient circuits for FHE. The core optimizations that we utilize with Yosys include functional and word-size reduction, removing redundant logic, and omitting unreachable branches in decision trees. Compared to the baseline Google XLS logic optimizations, we observe a reduction of about 40% in the overall size of the circuit for a dot product of two vectors with length 500. However, we note that the Google XLS logic optimizer is more lightweight and can process the encrypted circuit about twice as fast. We emphasize that this process is a one-time cost; after the circuit is processed, it can be executed using an arbitrary number of inputs.

After the Boolean netlist has been generated by the frontend compiler, and before encrypted computation can be carried out, we need to translate each gate to the encrypted domain. This process involves traversing the circuit, which is represented as a directed-acyclic graph (DAG), and mapping each node to the equivalent CGGI gate function. All wires become ciphertext data, the inputs are loaded with encryptions provided by the client, while the outputs are communicated back to the client for decryption after circuit evaluation.

The intuitive approach for providing the computing party with an executable FHE circuit is to simply generate code that invokes the encrypted gate functions using the underlying backend directly one after the other. This approach works well for small programs where performance is not critical, but is ill-suited for non-trivial programs. For complex programs, the generated FHE code can grow to millions of lines in length, as each logic gate in the circuit would require 2-3 lines of code on average.

Our key observation is that it is more efficient to avoid code generation entirely and incorporate a scheduler that traverses the DAG and distributes each gate to additional workers that exclusively run the corresponding FHE logic gate function. In this approach, gates that are ready to be evaluated can be distributed across a set of workers to exploit the circuit-level parallelism inherent in all applications. We remark that a similar methodology is employed by the Google FHE transpiler (Gorantala et al., 2021) and is referred to as interpreter mode. However, their strategy of distributing gates one at a time is not feasible when the workers constitute GPUs. Previous GPU-centric CGGI implementations as well as our proposed implementation (described in Section 5) can execute one homomorphic logic gate per streaming multiprocessor (SM) concurrently. In the case of an NVIDIA A100 GPU, $108$ homomorphic logic gates are the least number required to achieve $100\%$ device utilization at any given time. Thus, only one SM could be engaged if gates are assigned one at a time, resulting in extremely inefficient evaluation. Further, interpreter mode creates new ciphertext objects and allocates more memory as needed throughout circuit evaluation. While this technique is suitable for CPU workers, it therefore creates a prohibitive bottleneck on GPUs as memory allocation and ciphertext transfers between the host and the device are costly.

We propose a novel methodology for efficient evaluation of encrypted circuits on both CPU and GPU devices. The host thread parses the intermediate representation (IR) generated by the frontend, and generates a set of nodes stored using XLS data structures (XLS, 2020). Each node contains an opcode, which defines the operation performed by the circuit gate, and its input operands, which are pointers to other XLS nodes.

The IR thus defines a sequence of gates which can be processed sequentially to generate a valid execution of the circuit. In order to introduce parallelism, we transform this ordered set of XLS nodes into a vector of circuit gates. In addition to logic gates, we also create gates which compute encrypted constant values and augment the frontend IR with gates which copy an input ciphertext into another one. These copies are used to integrate the retrieval of encrypted results as part of the circuit itself instead of having to extract individual ciphertexts after waiting for circuit termination.

To derive a parallel execution of the circuit, we first dispatch all gates into multiple waves. Each wave must be processed in-order, but all entries of a wave can be processed concurrently. We now detail the topological sort algorithm that we use to build the list of waves.

For each entry of the vector of gates, we compute its successors (gates depending on it), and count its predecessors (gates on which it depends). To assign gates to different waves, we create a FIFO of ready gates, which are gates with no remaining dependencies. We start by adding all gates with no input dependencies into this FIFO. Until the FIFO is empty, we remove the first entry $n$ from the FIFO, and do the following :

• •

  Assign $n$ to the first wave if there are no input dependencies, or compute the max index of the wave of all predecessors, and add $n$ to the next wave. All predecessors have an index or $n$ would not be in the ready FIFO.

• •

  We decrement the predecessor count of all successors of $n$. Any of these successors reaching a null predecessor count are put in the ready FIFO.

This algorithm terminates even if the circuit is not connected. As the IR can be processed sequentially one node after the other, there cannot be cycles in the circuit and all gates will be given an index. Since nodes are assigned to waves with indexes that are strictly greater than their predecessors, all entries in a wave are independent and can be processed concurrently, as long as the waves are processed in order. We thus automatically derived a parallel execution from the IR, based on the fact that the IR was a valid sequential execution and used node operands to compute dependencies. This algorithm has a linear complexity because each node is taken exactly once from the FIFO, and we decrement counters as many times as there are wires in the circuit. Partitioning the circuit into such waves provides concurrency which can be exploited to efficiently use a single GPU device. For multiple processing units, we dispatch waves over the different devices. A simple solution to dispatch a wave with $N$ gates over $K$ devices that consists of splitting it into roughly $N/K$ gates per device, is illustrated in Figure 6.

Let us consider a wave with 43875 gates, composed of 2125 AND gates, 25000 OR gates and 16750 NOT gates. On two devices we could have 1356, 10465 and 769 gates of type AND, OR and NOT for device 0; and have 769, 14535 and 6633 gates of these types on device 1. This represents a total of 21938 gates on device 0, and 21937 on device 1, but we measured that device 0 and 1 respectively need 2.21 ms and 2.87 ms to process their portion of the wave. This 30% load imbalance is explained by the fact that AND and OR gates take the same time to process (about 0.19 us per gate), while it also takes 0.19 us to process 1024 NOT gates, which are non-bootstrapped. Equally dividing the gates between devices is therefore not a satisfactory approach, and only results in a 1.7x speedup with two devices.

We could consider the disparities between the different types of gates to evenly divide the load between the different devices based on performance models, but it would require an extra training phase per gate type. This may be tedious and not reliable when combining multiple gates with different compute or memory bandwidth intensity. In practice, the number of gates is usually large enough that a simpler but effective solution is to assign the same number of identical gates on all processing units. In our previous example, this results in putting 1064 and 1063 AND gates respectively on devices 0 and 1, and putting 12500 OR gates and 8375 NOT gates on both devices 0 and 1. We then measure 2.54 ms of work on both devices, with a negligible difference of less than 0.2 us, which corresponds to a perfect balancing.

For each wave, ArctyrEX implements this strategy using a hash-table which associates a vector to each of the opcodes encountered in the wave. Each entry of the wave is then appended to the list which corresponds to its opcode. Considering that there are only 8 types of standard logic gates currently supported, and that this number would not grow significantly, appending an entry roughly has a constant complexity. This phase therefore also has an overall linear complexity. In Section 5 we will show that building such vectors of identical gates makes it straightforward to implement batched kernels which obtain much higher performance.

In this Section, we have shown how ArctyrEX converts the frontend IR into a well-balanced parallel workload. Provided CPUs with sufficient processing power, nothing prevents us from assigning them parts of the waves too. This could be done using performance models based on per-gate performance models, or more simply based on the respective peak performance of the different types of processing units. Our methodology is therefore suitable to address hybrid systems combining CPUs and GPUs.

Since we have covered how gates are batched for distribution to different processing units, we now describe how we can access data across the entire system.

NVIDIA GPUs have a distinct memory hierarchy that differs in key ways from traditional CPUs. Inside a streaming multiprocessor (SM), there is a fast on-chip piece of memory partitioned between an L2 cache, and a resource called shared memory. These on-chip memories are much faster than global memory as they are part of the SM itself. In fact, shared memory latency is roughly 100x lower than un-cached global memory latency, provided efficient memory access patterns. Shared memory is allocated per thread block, so all threads in the block have access to the same shared memory. In the case of the A100 GPU, the combined capacity of on-chip memory per SM is 192 kB. Global memory is the largest memory (40 or 80 GB for the A100) and resides off-chip, making it the slowest aside from accessing memory on the host (NVIDIA, 2020).

Bootstrapping keys have a relatively large size, of approximately 100 MB for 110 bits of security. They cannot fit into GPU L2 caches, but are used for the majority of encrypted gate evaluations. Because of this, we replicate them in the global memory of all devices so that each can access the evaluation keys directly. We note that these keys are constant, and can be accessed concurrently within a device.

Ciphertexts are processed during circuit evaluation, and may be used simultaneously on different devices, or accessed from the host. We thus store them in pinned host memory, which is memory allocated with cudaMallocHost()). Ciphertexts are then cached into shared memory which is much faster and is located close to the GPU SMs. These kernels indeed have an extremely high arithmetic intensity and the PCI-e bandwidth consumption is limited, and the large amount of concurrency overlaps transfers with computation. This was verified experimentally by profiling a kernel that processes 1024 gates using the ncu tool. We observed that it only consumed 19.96 MB/s of “system memory” bandwidth, which is orders of magnitude lower than the available PCI-e bandwidth. Using pinned host memory to load the input ciphertexts is thus efficient enough, in spite of its simplicity.

A similar strategy is to use managed memory (also called unified memory, and allocated using cudaMallocManaged). Contrary to pinned memory where devices access host memory directly through the PCI-e bus, managed memory is kept coherent across the entire machine by the means of paging mechanisms. When a page fault occurs, the CUDA driver automatically fetches a valid copy of the page where the fault occurred. Subsequent accesses to the same page will occur at the speed of the memory embedded on the device, until the page is evicted from the device.

Both managed memory and pinned host memory incur a significant overhead per allocation, so that we do not allocate all ciphertexts individually, but group these thanks to pooled memory allocators. This pooling mechanism may introduce false sharing issues, but effectively amortizes allocation overhead, which remains noticeable with pinned host memory, but is several orders of magnitude lower than the time required to evaluate the circuit. Memory pages allocated with managed memory and modified concurrently by multiple devices may bounce from one device to another, and have a severe impact on performance.

In practice, we observe similar performance for an encrypted dot product over 8 GPUs with both strategies. With pinned memory, circuit evaluation takes 14.8 s, compared to 15.1 s with managed memory. Allocating 1 GB of pinned host memory however takes 0.4 s, but is negligible with managed memory. Due to the expected page faults when using managed memory on multiple devices, we observe some slightly imperfect parallelism, while it is flawless with pinned memory. ArctyrEX therefore allows user to store ciphertexts either in host pinned or managed memory, for example depending on the amount of system memory which can limit the availability of pinned memory. All experiments presented in the rest of this paper use pinned host memory.

A strawman approach to assign tasks to multiple workers involves having a single producer thread and a set of worker/consumer threads. When using multiple GPUs, each worker thread will consume an assigned batch from the producer and outsource the computation to a dedicated GPU. This approach is quite simple to implement, but requires numerous synchronizations between CPU threads, which negatively impacts scalability by introducing idle periods on the GPUs when CPU threads fail to provide them computation.

Conversely, a more intuitive method involves utilizing a single host thread that will submit work asynchronously to different devices. On each device, we create a pool of CUDA streams, so that we can submit multiple concurrent CUDA kernels on this device. The execution of a single wave therefore consists of taking each individual batched kernel from the wave, selecting a CUDA stream on the device on which our scheduling algorithm assigned the batched kernel (e.g., with a round-robin strategy), and submit the kernel in this CUDA stream. Since waves must be executed in-order, we need to ensure that the execution of a wave does not start until the previous wave has been fully processed. A simple approach would consist of submitting all kernels in a wave on all devices, and then having the host thread wait for the completion of all work on all devices. Waiting for computation to complete from the host however introduces some inefficiency, as devices become idle during the synchronization phase, until the next phase has been submitted. Any potential load imbalance or jitter on a device may also reflect on other devices which could wait longer than expected to get more work. Instead of blocking devices, we have therefore implemented a non-blocking synchronization fence primitive which ensures that the work in a CUDA stream cannot start running until all work submitted previously in all other streams has been done. These fences are implemented by the means of CUDA events which are asynchronously inserted in the CUDA streams. After inserting an event in each stream, we insert a non blocking CUDA operation which synchronizes one of our CUDA streams with all of these events. We then insert another event in that stream, and make sure all other streams wait for that event. Event insertions and dependency declarations between an event and a stream can be performed asynchronously, ahead of time, and therefore do not require the host thread to block during the execution. These event-based synchronizations are implemented using hardware features, which is much more efficient that having the host thread block the entire device. This ensures that successive waves can be executed in order, without ever blocking the submission flow of asynchronous operations, until the very end of the circuit evaluation. With this distribution methodology, we observe a speedup of approximately 12% over the strawman approach for an encrypted dot product benchmark executed on an NVIDIA DGX A100. This may appear to be a moderate improvement, but more than 99% of the circuit evaluation is spent executing CUDA kernels. We therefore have a close to optimal scheduling strategy over multiple devices, which is essential for the scalability of ArctyrEX according to Amdhal’s law. This also indicates that the latency of result retrieval and synchronizations are almost completely hidden.

Due to the relatively small size of TFHE ciphertexts (compared to other FHE schemes), it is possible to process many FHE gate operations at the same time on GPUs over a large number of ciphertexts. Prior works have either launched a separate kernel for every gate evaluation (Dai and Sunar, 2018) or allow for “vectorized” gates (such as performing a bitwise NAND between two 32-bit ciphertext arrays) (NuCypher, 2019). Conversely, we observe that a better approach for general computation is to leverage a kernel capable of executing arbitrary numbers of gates of any supported type. The ArctyrEX backend utilizes a single kernel for each batch of gates that launches with $N$ thread-blocks of 512 threads each, where $N$ indicates the number of gates. This approach is more performant compared to the cuFHE library that initiates host-to-device and device-to-host transfers for each logic gate. This allows each worker in the runtime environment to launch a single kernel for each batch received from the coordinator, avoiding additional kernel launch overheads. Additionally, this technique also allows the GPU to determine the best utilization strategy for the SMs, instead of relying on the user to distribute gates on a per SM basis. Grouping gates into homogeneous gates saves memory bandwidth as we only copy the opcode value once per batched kernel, and the generated code is more regular and requires less registers, increasing the occupancy of our CUDA kernels (NVIDIA, 2023).

Designing batched CUDA kernels which do not require blocking the submitting host thread is also challenging. These kernels indeed need to access buffers with the description of the work to perform, such as the location of the input ciphertexts. We therefore adopt a strategy which consists of assigning such a buffer to each CUDA stream of our pool, and fills them asynchronously from the host using a host callback. As a result, our asynchronous batched kernels consist of 1) selecting a CUDA stream on the target device, 2) submitting a host callback that will update the buffer associated to this stream, and 3) launching a CUDA kernel in this stream which will process the batch described using this buffer. Assigning each CUDA stream a unique buffer requires a limited memory footprint, and ensures there is no concurrent buffer update. This also avoids the use of relatively expensive asynchronous allocations around all asynchronous kernels.

The performance of CGGI bootstrapping is largely determined by the DFT used to perform polynomial multiplication. Both nuFHE (NuCypher, 2019) and cuFHE (Dai and Sunar, 2018) use the NTT for this operation, and employ the same strategy in terms of NTT parameters. We opt to use these parameters as well, since they provide multiple key optimizations that reduce the NTT latency. First, we utilize the modulus $Q=2^{64}-2^{32}+1$, which simplifies modular reduction and supports NTTs up to size $2^{32}$. Lastly, we use the primitive element $g=12037493425763644479$ that allows most multiplications in the NTT algorithm to become bitshifts modulo $Q$.

The FHE Basic Linear Algebra Subroutines are benchmarks that form core components of algorithms in a wide variety of fields, such as image processing and machine learning. We focus on three distinct tensor multiplication algorithms on 16-bit encrypted data: a dot product of two vectors of length 500, a matrix-vector multiplication between a vector of length 125 and a $125\times 4$ matrix, and a matrix multiplication between two $10\times 10$ matrices. Additionally, we include a vector addition between two vectors of length 500; this benchmark was executed with a larger wordsize than the previous ones to increase its computational complexity. We compare a $256$-thread CPU execution of these tensor algorithms with our approach running on up to 8 GPUs.

In Figure 7 and Figure 8, the dark-green bars show running time, and the light green bars plot the speedup of GPU vs. CPU . One A100 is $5.9\times$ faster than the CPU reference running on the 256-threaded CPU execution model, and $8$ A100s are $40\times$ faster. We show the latency of these circuits for an increasing number of A100 GPUs and the speedup for all GPU configurations versus a CPU configuration with 256 threads. Our analysis shows a linear speedup with more GPUs, as our design exploits the ample circuit-level parallelism in both synthesis and runtime phases.

Figure 9 shows the linear algebra topologies; the vector addition is more performant as the critical path is approximately $2\times$ shorter and the levels remain wide. Indeed, this is reflected in the execution times in Figure 7, where the vector addition runs nearly $4\times$ faster on 8 GPUs. Both matrix benchmarks have very wide levels and are well-suited for evaluation on multi-GPU systems.

One of the most widely explored use-cases for FHE is privacy-preserving machine learning as a service. We consider two distinct scenarios: one where the model is owned by the cloud server and another where the client or a third party owns the ML model. For the former, the network parameters are kept in plaintext form and can benefit from faster plaintext-ciphertext operations. Conversely, when the cloud doesn’t own the model, the network parameters are encrypted as well as the classification inputs. The majority of existing works demonstrating ML inference with FHE adopt the first scenario as it is generally much more efficient (Folkerts et al., 2021; Bourse et al., 2018; Chou et al., 2018).

Our analysis considers two important machine learning procedures for encrypted classification in the form of logistic regression (LR) inference and neural network (NN) inference. For logistic regression, we perform binary classification for datasets with four attributes, such as the Iris dataset (Fisher, 1936). The key bottleneck in encrypted LR inference is approximating the sigmoid function ($\frac{1}{1+e^{-x}}$), since it is not possible to evaluate it directly. Therefore, we employ a polynomial approximation by evaluating terms of the Maclaurin series. In general, when approximating nonlinear activation functions, there is a trade-off between accuracy and computational complexity. We show this trade-off with ArctyrEX through the use of an approximation that evaluates the first four terms, and one that evaluates the first six. Figure 10 shows diminishing returns when increasing the number of GPUs due to the large critical path and relatively thin circuit levels of the benchmark in Figure 12. Using 8 GPUs with the more accurate sigmoid approximation still outperforms the CPU implementation with 256 threads by $21\times$.

For neural network inference, we employ the same network architecture used by FHE-DiNN (Bourse et al., 2018). The network consists of two fully-connected layers with a sign activation function to classify the MNIST dataset of handwritten digits. We consider two variants of this network that differ in whether or not the network parameters are encrypted and both achieve an accuracy of 96% for MNIST classification. The execution times across both configurations are depicted in Figure 11 while Figure 12 presents the characteristics of these workloads. As expected, the variant where the cloud server does not own the proprietary network (i.e., using encrypted weights) has approximately $2\times$ higher latency because of the increased number of ciphertext-ciphertext operations.

Aside from machine learning and linear algebra benchmarks, we explore another realistic and useful application of FHE enabled by ArctyrEX in the form of private information retrieval (PIR). The ability to search and perform computation across an encrypted database has many useful applications, such as managing a directory of health-care records that must be kept confidential for compliance with standards such as HIPAA (Zhang et al., 2017). We represent the encrypted database as a key/value storage where both keys and values are encrypted. Figure 14 demonstrates the circuit characteristics of a single query for databases of increasing size. For brevity, we show results for databases with 256 and 128 elements, where each value is 64-bits in size. In Figure 13, we observe a linear scaling with increasing database size; this is expected as a query requires comparison operations with each record in the database due to the termination problem. Specifically, the termination problem states that it is impossible to make a decision based on encrypted data as the computing party does not know the underlying value of the ciphertext (Mouris et al., 2018). As such, each element of the database must be visited and the correct entry needs to be chosen through oblivious encrypted multiplexing operations.

The Cingulata framework (formerly Armadillo (Carpov et al., 2015)) allows users to map C++ code into a sequence of AND and XOR gates. Cingulata works strictly with binary FHE contexts using the TFHE library (which implements CGGI) and a custom BFV implementation as its backends. Compared to ArctyrEX, Cingulata only supports single-core execution for CGGI and does not offer GPU support. The BFV mode is parallelized on CPUs, but does not support bootstrapping and hence cannot be used for arbitrary general-purpose computation.

E³ is a C++ library that introduces custom encrypted data types for leveraging FHE in general applications (Chielle et al., 2018). It supports a variety of cryptographic backends, including TFHE, Microsoft SEAL, and HElib, encompassing all major FHE schemes. Unlike ArctyrEX, E³ uses a direct mapping to hardcoded FHE functional units and does not offer an optimizing compiler. It also does not support any GPU-friendly cryptographic backends and no parallelization is included.

Google’s FHE Transpiler (Gorantala et al., 2021) and Romeo (Gouert and Tsoutsos, 2020) leverage logic synthesis and optimizations to generate FHE programs for general computation. However, both works employ generic synthesis scripts that include optimizations not relevant to encrypted computation. The FHE Transpiler targets TFHE and the OpenFHE implementations of the CGGI cryptosystem as backends, and can evaluate multiple gates in parallel using interpreter mode. However, it does not support GPUs and its parallelization strategy is not suited for them, yielding very low device utilization. Likewise, Romeo targets TFHE and generates an FHE program instead of interpreting it. This approach, however, does not scale for large programs or HPC systems as described in Section 4. Conversely, ArctyrEX offers a novel dispatch strategy and multigate kernels that can efficiently compute batches of any set of gates.

The cuFHE (Dai and Sunar, 2018) and nuFHE (NuCypher, 2019) constitute the current state-of-the-art for GPU acceleration of the CGGI cryptosystem. The former is a proof-of-concept library that implements high throughput logic gate evaluations on a single NVIDIA GPU. However, cuFHE is not configurable (i.e., only supports 80 bits of security), has non-optimal data transfers and requires frequent high-cost synchronization between GPU and CPU. Each cuFHE gate evaluation requires all ciphertext inputs be copied from the host to the device, and each output is copied back from the device to the host. This approach is impractical for realistic circuit evaluation, as it yields millions of large ciphertext transfers between the CPU and GPU. Lastly, not all cuFHE computations are outsourced to the GPU and the CPU needs to perform certain operations (such as evaluating the homomorphic NOT gate). Unfortunately, this defeats the benefits gained from asynchronous CUDA kernel launches and the CPU execution must block when it reaches a NOT gate until the GPU has finished evaluating all prior gates, instead of continuing to do more meaningful work. Similarly, nuFHE specializes in vectorized gates; for instance, it can evaluate a bitwise AND operation across 64-bit operands. However, this approach is very restrictive in terms of circuit evaluations as typically a circuit level is not composed of one type of gate.

REDcuFHE (Folkerts et al., 2021) enhances cuFHE to add multi-bit plaintext support and multi-GPU support. However, it still suffers from the same synchronization issues as cuFHE, and puts the burden of scheduling and handling communication between multiple GPUs on the programmer. ArctyrEX, on the other hand, handles all scheduling and communication procedures automatically. Lastly, Morshed et al. (Morshed et al., 2020) present a GPU implementation of CGGI that leverages the NVIDIA cuFFT library and incorporates a set of handwritten circuits such as vector addition and matrix multiplication. Table 1 demonstrates that ArctyrEX outperforms (Morshed et al., 2020) by a factor of about $1.5\times$ for a small 32-bit addition and $16\times$ for 32-bit multiplication (which is a significantly larger circuit). Additionally, ArctyrEX evaluates a vector addition with 32 elements of 32-bit integers $4.1\times$ faster and a $16\times 16$ matrix multiplication of 32-bit elements $10.6\times$ faster. We also emphasize that all frameworks in Table 1, aside from ArctyrEX and the Google FHE Transpiler, require developers to write their own circuits by hand, as opposed to automatically generating them.

A final class includes custom ASICs and FGPA implementations of FHE to act as a co-processor to speed up the underlying FHE arithmetic operations. However, many of these designs only support limited parameter sets (Nabeel et al., 2022; Turan et al., 2020; Roy et al., 2019; Samardzic et al., 2021) and usually target other cryptosystems that enable approximate computing (Samardzic et al., 2022; Agrawal et al., 2023) or computation with modular integers (Nabeel et al., 2022). On the other hand, ArctyrEX can support arbitrary parameters and implements the CGGI cryptosystem, which is more suitable for general-purpose computation.