arXiv CSopen access

Python in the front, party in the Backline: compiling quantum workloads across CPUs, GPUs, and FPGAs

1

Python in the front, party in the Backline: compiling quantum workloads across CPUs, GPUs, and FPGAs

arXiv:2609.09270v1 [quant-ph] 8 Sep 2026

Joseph K. L. Lee, Mehrdad Malekmohammadi, Hong-Sheng Zheng, Shuli Shu, Cheick Doumbia, Kalman Szenes, Mehran Zamani Abnili, Thomas Ainsworth, Matthew Seymour, Thomas Germain, Leonhard Neuhaus, Josh Izaac, and Lee J. O’Riordan

Abstract—Moving from quantum research and development to production-grade, fault-tolerant quantum workload execution remains one of the most significant challenges facing quantum platform builders. While Python frameworks have enabled an easy entry point for quantum algorithm design, the low-latency requirements for real-time quantum error correction (QEC) demand performance that traditional interpreted environments cannot provide. FPGAs and ASICs play a central role at these layers, but their specialized programming models make development rigid and time-consuming. CPUs, GPUs, and other accelerators introduce a different challenge: as infrastructure becomes increasingly heterogeneous, programming across different devices and their associated abstractions becomes more complex. Designing abstractions that allow researchers to write workloads in high-level languages that map to low-latency execution across diverse distributed target platforms will enable the development of key infrastructure for utility-scale quantum systems. For this, we introduce Backline, a heterogeneous compilation and runtime framework built within PennyLane and Catalyst. Backline allows us to design and build quantum-classical workloads for highperformance and low-latency devices, with compilation directly from a Python interface through MLIR. We demonstrate the compilation and execution of several quantum workloads with low-latency data movement across a mix of CPUs, GPUs, and FPGAs, efficiently handled for both local and distributed remote hardware targets, all from a vendor-agnostic Python frontend. With an AMD VPK120 FPGA board as the controller, issuing each round from its hardware-handshake engine, we measured median steady-state round-trip latencies over RoCE v2 of 2.305 µs to an AMD Ryzen Threadripper PRO CPU and 4.5 µs to an AMD Instinct MI210 GPU across 106 −1 rounds per path, demonstrating microsecond-scale synchronous co-processing. Index Terms—Quantum, RDMA, GPU, HPC, Compiler

I. I NTRODUCTION Across many domains of scientific computing and research software, writing performant and expressive systems is a complex endeavour. Allowing domain experts to work at the right layer of abstraction enables them to efficiently design, prototype, and build out workloads that move their research forward. A quantum computing algorithm researcher should not need to understand memory access optimizations to implement their algorithms, as such low-level details can and should be abstracted away. Similarly, an HPC performance engineer should not need to understand the complex mathematical formalism of the special unitary group, only that the data is handled efficiently on the underlying hardware. All authors are affiliated with Xanadu Quantum Technologies Inc., Toronto, Ontario M5G 2C8, Canada (e-mail: [email protected]).

In the field of quantum system design, we have been able to avoid such targeted complexities in the past with noisy intermediate-scale quantum (NISQ) machines, since fast feedforward from measurement outcomes to subsequent corrections was never a part of the execution paradigm. For fault-tolerant quantum hardware environments we now need to take this into consideration, as the application of gates based on branched pathways is a fundamental part of the execution paradigm. Quantum error correction (QEC) for fault-tolerant quantum computing (FTQC) has extremely tight communication and feedback loops. Additionally, as we require quantum devices to be programmable and controllable in such low-latency environments, FPGAs and ASICs become paramount components of the stack. Writing software for devices at these layers remains hard and rigid, as ASICs are often built to serve a specific singular purpose, and resynthesizing FPGA bitstreams can be a timeconsuming process for large workloads [1]. As QEC requires researchers and quantum control-system developers to interact with these devices, appropriate software abstractions can help hide unnecessary implementation complexity. At the same time, researchers designing algorithms and applications often have heavy computational requirements, which on the spectrum of quantum workloads can be compute-bound, communicationbound, and everything in between [2]. Since FTQC systems remain under active development across all quantum hardware platforms, there remains significant R&D for software, hardware, QEC, and architecture, with many identifying various requirements across the stack for both applications and infrastructure [3]. For fast progress, we need stacks that support this R&D and co-design end-to-end with layered interfaces for everyone involved, from application developers to system engineers. This means supporting users who are prototyping FTQC workloads with QEC in high-level languages (e.g. Python), and allowing them to progressively compile and optimize these workloads onto local and distributed heterogeneous devices matching the real-world environments and components in which large-scale FTQC systems will operate. As these systems scale, developing software that spans both distributed computing and quantum computing becomes increasingly complex, and this complexity is a recognized challenge [3]–[7]. Within fast-moving technical fields, two disparate approaches are often taken: i) aim to lock down and standardize the stack to allow multiple parties to unify around a common base, or ii) move fast with less-than-optimal designs to unblock important ecosystem research work [8].

2

While both have their merits, stability and performance can devices can happen on a given node, and remote memory often be at odds, requiring a nuanced approach in research access via RDMA [22], [23] where communication between and design, and standardizing a fast-moving field too early can devices across different nodes is enabled by efficient deviceentrench suboptimal solutions [9], [10]. Platforms for QEC to-NIC communication. While programming with low-level infrastructure (both hardware and software) remain highly abstractions is common in HPC ecosystem design, application bespoke, with limited access to interface layers allowing developers should not need to understand, or even be aware of, development against multiple hardware variants. While efforts these complex system layers. Given that quantum workloads like NVQLink [11] are a move in the right direction, the will largely run on mixtures of commodity, HPC, and bespoke lack of variety in target platforms is a recognized constraint. hardware across various component types, it is necessary to Indeed, having multiple vendor-agnostic platforms would be create well-designed abstractions for the stack across all layers. For programming quantum devices, the use of high-level valuable to the community [12]. Therefore, we favour a movingdesign specification that evolves with the field, rather than a Python programming APIs backed by MLIR/LLVM compiler rigid standard, unlocking the best of both approaches above. infrastructure is becoming the de facto design of choice to We approach these systems with “best effort” compatibility, abstract away the complexity of the executing environment similar to the LLVM ecosystem’s approach to stability and from the user-facing layers [24]–[28]. In PennyLane, a program compatibility between releases [13], letting us move quickly is written in Python, and compiled by the Catalyst compiler with researcher and ecosystem needs while still offering through an MLIR and LLVM pipeline to progressively lower the classical and quantum program to device-native operations. compatibility guarantees across changes. Recent trends in software development for quantum ecosys- While MLIR and LLVM are expansive compilation tools, tems follow the early days of building expansive machine- they complement the Python-first designs, helping to keep learning pipelines, with increasing development, abstraction, the complex system and compilation layers abstracted away and sophistication across the stack. As workload complexity ex- from quantum application developers. This compilation-first pands, the need for managed device interactions and distributed approach unlocks many abilities that have limited exposure environments is becoming central to the ecosystem’s design. in NISQ-first stacks, such as preservation of program strucThe need for communication across multiple hardware and ture [29], complex targeted decompositions, and hierarchical system-level components, and doing so with minimal overheads, optimization, as well as integration of existing classical code is at the core of many such ecosystems. For example, consider in the same program. Given the separation between high-level algorithmic research machine learning inference workloads, where disaggregated computing environments assign dedicated devices to dedicated for optimal program representations, and high-performance lowfunctions and behaviours [14], [15]. In this scenario, a variety latency workload execution, marrying both across a variety of of device types with unique responsibilities are treated as device types can be a daunting challenge. Our manuscript will resources, and are engaged according to behaviour that matches address this exact scenario. We will explore how to provide accessible high-level programmability while maintaining besta given workload’s needs. Frameworks like TensorFlow, PyTorch, and JAX [16]–[18] in-class low-latency performance for quantum workloads have enabled rapid progress in abstracting away much of through a heterogeneous compilation and runtime layer we call the related machinery and configuration needed to scale up Backline. The manuscript is structured as follows: and scale out ML workloads. In each of these, Python is the frontend, with support for both eager (Python interpreter) • Section II provides a snapshot of fault-tolerant quantum and compiled workloads, allowing execution on local and computing, hardware constraints, and R&D requirements. remote CPUs, GPUs, TPUs, and custom devices. As these • Section III describes the open-source software PennyLane frameworks follow a familiar NumPy-API design for writing for quantum programming and Catalyst for compiling the user algorithms, applications, and end-to-end workloads, they resulting programs to heterogeneous hardware devices. have also expanded the infrastructure to support extensibility • Section IV focuses on the design of Backline’s interopand customization across the ecosystem. Kernel-development erability between heterogeneous devices and low-latency tooling like Pallas [19] and Triton [20] expand the above SDKs communication abstractions from PennyLane through to allow explicit definition of complex accelerator kernels, Catalyst. directly from Python. The usability and performance of these • Section V demonstrates usage of Backline abstractions via tools, relative to bare C-style extensions, have proven to be PennyLane and Catalyst through examples of low-latency advantageous for compute-heavy workloads as well as for workloads running across asymmetric heterogeneous deproductivity [21]. As workloads expressed in such frameworks vices. scale, additional consideration must be given to data locality • Section VI examines the applicability of the infrastructure and availability for all devices in the system. Many frameworks for both current and future FTQC workloads. utilize an explicit data-placement model, allowing data to be II. FTQC PLATFORMS AND QEC manually migrated between devices as dictated by the workflow needs. To reduce bottlenecks in such workloads, this migration In the current state of quantum computing infrastructure, one needs to happen as efficiently as possible. Key technologies active research topic of high importance is the development of to support this are the commonly used direct memory access QEC protocols for encoding, workload execution, and decoding (DMA) abstractions, where direct local communication between of errors in real-time.

3

For Xanadu’s Aurora system [30], the error correction 2) Synchronous co-processing tier [1 µs–100 µs]: A soft realdecoder computation was implemented on an FPGA with a time regime for supportive, high-importance tasks such as latency of 64 ns. The end-to-end feed-forward latency within backup decoding and auxiliary processing that can tolerate which the decoder needed to provide corrections back to the some variation in execution time. Coherent memory access system was slightly below the available 1000 ns budget due and synchronous low-latency kernels on FPGAs, CPUs, to additional latencies coming from analog, analog-to-digital, and GPUs can serve this regime, depending on the and serial communication delays imposed by the hardware computational requirements. selected for the experiment. More optimal hardware choices 3) Co-located runtime tier [100 µs–1 s]: For batched or such as using zero-latency analog-to-digital converters and numerical assistive workloads, as well as tight compilation optimized serial transceivers can bring this additional latency loops. Execution here should not sit on the critical path below 100 ns. As this example relied on a small-scale code, for quantum workloads, and can draw on externally increasing the complexity of the QEC schema will further available libraries and tooling, such as high-throughput extend the execution times of the decoder, potentially exceeding communication collectives and optimized linear-algebra the latency budgets for real-time corrections. There is currently implementations. significant ongoing work to explore hardware implementations 4) Distributed workload tier [>1 s]: Standard data-processing that enable faster decoding in small-time windows [31], [32]. workloads, where start-up costs can be paid to allow subseConsidering the decoding loop of the Aurora system, where quent communication and compute to overlap effectively. an inner and outer decoder are paired to ensure logical This covers pre- and post-processing, as well as noncorrectness of the system within a given number of clock cycles, critical calculations on local or remote systems. having the ability to prototype such protocols can help unlock The tier that a system is required to hit is set by the QEC an accessible R&D path for system-level research. Because cycle rate of the underlying qubit modality, since the decoder these quantum hardware ecosystems are highly complex, throughput must keep pace with the incoming syndrome data; building infrastructure that supports the iterative design and otherwise the computation suffers an exponential backlog testing of these protocols is essential. This infrastructure slowdown [33]. The required timescales span several orders of enables teams to understand bottlenecks, application viability, magnitude across platforms, where tighter modalities such as and overall system-level requirements for real-world quantum photonic and superconducting systems sit around 1 µs, while workloads. Consequently, a unified end-to-end stack (program- looser ones such as ion traps and neutral atoms sit closer to ming, compilation, orchestration, runtime execution) becomes 1 ms. For photonic systems, Tier 1 is a hard requirement for paramount for effective hardware-software co-design, ensuring decoding. That budget binds the primary decoder while the that classical control architectures and quantum processors looser tiers exist for the classical work that does not need to are iteratively tailored to mutually support one another’s keep pace with the cycle rate. constraints. For workloads that have a hard deadline on the arrival of However, this co-design process must account for highly a result, strategies such as latency hiding do not provide non-uniform system requirements. The various classical tasks any benefit: if the result is unavailable, the outcome can surrounding a quantum device operate under vastly different be fatal to the workload. Running programs on real-world timing deadlines, and conflating them obscures which specific quantum hardware falls into this category, with error-correction classical hardware is actually needed at each layer. For the subroutines being essential to ensure the functional operation purpose of quantum system hardware, we can loosely categorize of the device. Moreover, iterative decoders such as belief latencies into different regions of utility. It is important to propagation (BP), are inherently probabilistic, which means that separate two distinct axes: i) what a given class of classical there is a given likelihood that the decoder does not converge hardware can deliver, and ii) what a given quantum modality in the operation timescale budget. This leads to a stall of the actually requires. We first lay out the capabilities and what is system until the correction can be determined and applied. achievable, and then map the per-modality requirements onto While an FPGA decoder is expected to correct the majority it. of such errors, some small tail of events (e.g. those beyond the Following the resource analyses of real-time decoding and P99.999 percentile) cannot be corrected in a timely manner hybrid quantum-classical execution [2], [33], we define the and introduce the problems mentioned above. These tail-errors following tiers of latency-bounded operation for an FTQC are too complex to converge with a fixed decoder design, and system. These tiers describe the class of hardware that can may, therefore, benefit from an additional auxiliary processor service a given latency band, with the modality then dictating equipped with a more sophisticated algorithm. Such a decoder which tier must be met, and looser tiers enabling more complex is not required on every round, and, therefore, does not need to computations at lower cost and development effort: meet the real-time control tier budget above. Indeed, it belongs 1) Real-time control tier [1 ns–1 µs]: The tightest latency in Tier 2, which is reachable with commodity hardware. These regime, servicing hard real-time tasks such as QEC decoders could run on a variety of heterogeneous devices (CPU, decoders and associated adaptive control hardware, where GPU, or even another FPGA), and assist the system to handle each computation must complete within a deterministic, low-likelihood tail-errors [34], preventing the decoding backlog worst-case-bounded budget [33]. ASICs, FPGAs, and from building up to unresolvable amounts [35] and allowing embedded real-time coprocessor cores are the primary execution to continue. targets at this scale. To minimize stalls, this auxiliary hardware should i) have

4

access to the relevant data, ii) be primed for immediate execution of the decoder, and iii) be ready to transmit the correction back for application. Realizing these concepts calls for careful consideration of both the user interface and user experience, as well as the underlying infrastructure and protocols needed to support workloads operating under this paradigm. For these abstractions to be practical on real-world quantum systems, they must be supported by a low-latency execution environment, particularly given the communication overhead associated with transferring data. At the same time, the system should retain sufficient programmability to enable a reasonably short turnaround from prototyping to experimental data collection.

Since programmability is the most important enabler of any new technology (if you cannot program it, you cannot use it), defining the right abstractions becomes important for endusers and researchers. Shallow-embedded DSLs (e.g., tracing interpreters) are often favoured for their extensibility over deep-embeddings (e.g. AST “interpreters”) due to the host language and associated libraries being implicitly accessible to provide key functionality as expressed by the user [40]. This usability comes with a trade-off: deeper embeddings provide more compiler-friendly constructs by enabling direct interpretation at the AST-level. To ease the extensibility of our ecosystem, we opt for a dual-mode approach comprising a shallow embedded DSL [41] that is traced to capture the intention of the quantum program, with additional functionality provided by an AST-level rewriter that captures and represents III. P ROGRAMMING AND COMPILING QUANTUM PROGRAMS Python’s structured control-flow [42]. In this way, we treat Python as the language to express the High-level software abstractions are needed to hide the complexity of the underlying quantum hardware. To accommo- user’s intent and to build the compiled program, rather than date the increasing heterogeneity in computing systems, many the interpreted execution environment it is commonly used for. frameworks have been developed to provide this functionality at Within this environment, users can express their workloads various layers of abstraction, with Pythonic frontends coupled for heterogeneous compilation and execution, with our qjit with compiled backends being widely used [24], [25], [28], frontend layer handling the capturing and lowering of the [36]–[38]. Allowing for the user program to be written in program from high-level concepts through MLIR all the way Python and natively compiled on a variety of hardware types down to the final compiled program. The PennyLane tracing (e.g. CPUs, GPUs, TPUs, QPUs, etc.) helps remove execution mechanism is currently based on JAX [18], [43] and, thus, relies overheads from Python interpreters and schedule execution on similar infrastructure to represent the high-level program, on the most effective device target [18], [39]. PennyLane’s before explicit conversion to MLIR, from which Catalyst [25] compiler Catalyst allows for the expression of workloads at the lowers the program to the pre-defined targets. Python layer, with the program progressing through multiple lowering stages from MLIR, onto LLVM IR, and to a final IV. L OW- LATENCY COMPUTING WITH HPC- GRADE binary program representation for the executable workload. HARDWARE Quantum-device-specific calls are provided by an associated While Section III focused on programming abstractions and runtime layer, which is linked against the program for execution compilation, one must also consider a runtime layer for such on the target hardware. For a PennyLane program, compilation compiled workloads, handling the direct tie-in to the hardware, through Catalyst follows the high-level design given by Fig. 1. as well as data migration, synchronization, and related systemlevel functionality. Operating systems often have well-defined abstractions that allow many devices to communicate effectively 1. Frontend: Program rep. capture Python / PennyLane (@qjit) with the operating system kernel, behaving as a mediator in all + JAX Tracing + Autograph AST rewriting communication paths. These layered abstractions allow for an Captured high-level user program effective way to build interoperability between devices, treating Return final the system-provided abstractions as the means to “wire” devices 2. Compiler Middle-End JAX / NumPy arrays MLIR Quantum Dialect & Optimization together in flexible ways. Such flexibility however can come at the cost of performance: Hierarchical lowered and optimized program ensuring that there is always a mediator handling the data as it traverses the system may add some overheads which can be 3. Backend: Code Generation LLVM IR & RT => Native Executable reduced by a more direct communication path. To meet the realtime needs of quantum hardware systems, it becomes important Final compiled program dyn. lib. to design software ecosystems with data placement and mobility 4. Runtime Environment built-in across all abstraction layers for low-latency execution. Catalyst C++ Runtime Of the tiers set out in Section II, Tier 1 is reachable only measurement Loaded and executing program data with bespoke hardware infrastructure sitting beside the quantum device. We therefore focus on programmability within Tier 2, 5. Quantum Execution Target with Tier 3 and Tier 4 also implicitly supported. Hardware backend (physical or simulated) For Tier 2, having easily programmable auxiliary devices to complement the real-time control path allows for more flexible Fig. 1: High-level overview of the PennyLane program lowering development, design, and exploration of key infrastructure, process. such as auxiliary decoding and switching [34]. Building such

5

low-latency infrastructure, as well as the workloads that run on them, relies on many approaches to reduce overheads, with kernel bypass techniques being commonly used. These techniques allow user-space applications to skip over standard processing practices that happen in the Linux kernel, allowing the removal of overheads by more direct data communication between hardware. We can consider the use of remote memory access paradigms, allowing data exchange between classical devices in the most effective manner based on the topology and composition of the target system. The use of such paradigms has recently been both discussed [44] and demonstrated [11] for quantum programming environments, tying into the remote direct-memory access (RDMA) subsystem of the Linux operating system and using RDMA over Converged Ethernet (RoCE) [22] as the underlying network and transport protocol. Through this, costly context switches from the CPU can be avoided, as well as unnecessary memory movement between PCIe-accessible devices and main memory. Instead, PCIe devices communicate their messages directly while minimizing system-level involvement. Assuming a user wishes to read from or write to a remote target’s memory, the accelerators and the RNICs are the only participants in the data exchange, avoiding costly CPU and DRAM accesses that are normally incurred in default execution pathways. This infrastructure is used by most, if not all, HPC systems, as well as for running large-scale machine-learning workloads [45], where both high throughput and low-latency data transmission are essential for reducing time-to-solution. In large-scale quantum systems, using these networks for auxiliary decoding requires advanced scheduling policies to allocate decoder resources efficiently [46]. For most R&D applications, however, a tightly connected mesh of devices is sufficient. Supporting these use cases requires mechanisms for expressing communication paths between devices with different target architectures, as well as runtime layers that abstract data migration from end users.

A. Backline With the above needs identified, we introduce Backline, a high-level data-placement abstraction, heterogeneous compilation layer, and low-latency runtime layer, built into PennyLane and Catalyst. Backline spans three layers: 1) The PennyLane frontend provides a declarative interface for specifying the target devices that participate in a workload, assigning computations to those devices, and selecting the transport protocol used for data exchange. 2) The Catalyst compiler translates this specification into MLIR, compiles program components for their assigned execution environments, and coordinates their deployment and execution. 3) The runtime implements target-specific execution, memory management, and data movement, including registering memory regions, binding the function that a coprocessor runs, and actually moving payloads between the devices. This modular setup enables easy development of new runtime devices with a unified and portable user interface. Figure 2 shows how the various components compose.

Placement – qp.Backline(...) PennyLane frontend Node – target, locality, connectivity Controller drives the QNode, initiates every round CPU | FPGA

Coprocessor runs a coprocessor function per message CPU | GPU

transport – "memcpy" or "rdma" carried into the compiler MLIR dialects executor which device instructions run on, and how it gets there

Catalyst compiler transport how data moves between them, once they are running

lowered to the runtime’s C entry points unified C ABI controller target session life cycle, memory regions, drives each round

runtime coprocessor target binds and runs the coprocessor function for each message

rounds

Fig. 2: The Backline abstractions and how they compose.

B. User interface for workload expression Our aim for the frontend layer is to enable users to prototype a heterogeneous workload locally on a laptop with a CPU, or on a single workstation with a GPU, which can subsequently be moved onto HPC-grade hardware with an RDMA-capable NIC to achieve very low latency with minimal changes to the program, all directly within Python through PennyLane. 1) Data placement abstraction: Data movement is primarily implicit in the model. The user creates a heterogeneous device explicitly through the usual PennyLane interface, and the transfers between the component targets follow from the declaration, rather than being written out by hand. The abstraction at this layer comprises the following components: • A Node is one participant in a (heterogeneous) workload. It carries the information needed to identify and reach that participant: the hardware it executes on, which the compiler pairs with the placement’s transport to resolve a concrete target library; whether it is remote; how its code is deployed there; and where the role requires, its connectivity and arguments to set it up. The following two roles specialize this: • A Controller is the concrete root of the placement model and drives the quantum device: it initiates all workload execution, mediates data transfers, and is directly responsible for returning results to the originating user process once the compiled workload completes. This could be a CPU, or an FPGA, and it executes a QNode. • A Coprocessor is an accelerator that assists the controller, for example as a decoder. This can be a CPU,

6

a GPU, or any platform directly targetable with the supported compiler and backend infrastructure. Each one registers a coprocessor function, which defines what runs there for every message it receives. • A node may run in the host process, in a local executor process, or in an executor on another machine. Executordispatched nodes are compiled for the executor’s target and executed through LLVM ORC, as described in Section IV-C1. • One controller, zero or more coprocessors, and a transport mode together define a Placement, declared with qp.Backline. This is the control plane, i.e. the single place a workload’s topology is defined. qp.Backline creates a heterogeneous device, which binds directly to a PennyLane QNode and carries the placement into the compiler.

1 2

import pennylane as qp import triton.language as tl

3 4 5 6 7 8 9

# (a) Pre-compiled: name a symbol in a library the # runtime can reach. Demos 1, 1a on a CPU; # demos 3, 4 on a GPU. steane_decode = qp.CoprocessorFunction( name="steane_coprocessor", lib_path=STEANE_CPU_DECODER_LIB_PATH)

10 11 12 13 14 15 16 17 18 19 20 21 22

# (b) From a Triton kernel, written as plain Python # each entry is compiled by the helper. # Demo 5, abridged here. def steane_lookup(syndrome): idx = tl.cast(0, tl.uint32) for i in tl.static_range(3): idx |= tl.cast((syndrome >> (8 * i)) & 1, tl.uint32) << i qubit = (tl.cast(STEANE_LUT, tl.uint32) >> (idx * 4)) & 0xF return tl.where(qubit == 0xF, NO_ERROR, tl.cast(qubit, tl.uint64))

23 24 25 26 27 28 29

# decoder_id indexes the tuple at run time, so one # coprocessor serves X and Z. Steane is self-dual, # so the one kernel is given twice. steane_triton = qp.backline.triton_decoder( (steane_lookup, steane_lookup), platform="hip:gfx90a:64")

30

Taken together, this is an implicit memory model for interactions between target devices. Explicit memory interactions are commonplace in ML workloads (e.g. Torch and JAX expose tensor placement directly); here we aim instead for a data model managed by the device mesh itself, so that transfers follow from how its nodes are declared to interact. Concretely, a placement is a single qp.Backline call carrying a qp.Controller, a list of qp.Coprocessor nodes, and a transport mode, with each node holding the fields above. Section V-B1 gives complete placements for each of the demonstrated configurations.

2) Coprocessor functions: A coprocessor function is what a coprocessor applies to each message it receives. It can be defined in several ways, each trading simplicity against control and performance. For the most optimized implementations, the function can be written and pre-compiled directly in a library (e.g. in C++), for either a CPU or GPU target, and referenced by symbol. This gives direct control over the decoder’s inner loop and memory-access patterns.

31 32 33 34 35 36 37

# (c) From a Tanner graph: supply the CSS parity# check matrices only. Demos 2, 2a use the # [[13,1,3]] hypergraph-product code, whose # Hx and Hz are built in the demonstration. bp_osd_decoder = qp.backline.css_bp_decoder( Hx, Hz, postprocess="osd", num_iters=10, platform="hip:gfx90a:64")

38 39 40 41 42

coprocessor = qp.Coprocessor( hardware="gpu", remote=True, coprocessor_fn=bp_osd_decoder) # or coprocessor_fn=steane_decode/steane_triton

Listing 1: The three routes to a coprocessor function: naming a pre-compiled symbol, compiling a user-supplied Triton kernel, or generating a belief-propagation decoder with orderedstatistics post-processing (BP-OSD) decoder from a CSS code’s Tanner graph. All three produce a CoprocessorFunction and are interchangeable at the placement level. The comments name the demonstrations of Section V-B that take each route.

3) Transport mode and local prototyping: Prototyping locally requires nothing beyond a Linux machine, and optionally a GPU. By selecting a local transport (memcpy), the heterogeneous compilation and execution path, including placement, serialization, and the per-round request/reply structure, is still exercised. The messages are transferred as local memory copies. Alternatively, the function can be written in Python; the When a user has access to a more advanced RDMA-ready setup, provided helper functions take either custom Triton functions, changing the transport mode to rdma replaces those copies or Tanner graphs for CSS codes, and compile ahead-of-time with one-sided network writes. No change to the program as GPU kernels that the runtime target can load and launch. is required besides updating the transport name and nodes’ This significantly lowers the barrier to prototyping performant locality. 4) Granular control and runtime call: The abstractions GPU decoders directly from Python. Listing 1 shows all three above are implicit: the user declares a placement, and the routes. compiler decides where a round of communication appears. For iterative R&D, this is sometimes more abstraction than Additionally, a coprocessor may register multiple functions, one wants, so we also introduce functionality that exposes the dynamically dispatched depending on the incoming message. runtime directly through qp.runtime_call, which embeds A section of the payload is allocated for a decoder_id, which into the lowered program, at the point of expression, a call to is used to select which decoder function (e.g. for correcting a pre-compiled and linked dynamic library. X-type or Z-type errors) is applied to the incoming syndrome. This provides the simplest accessible abstraction for calling

7

into performance-optimized programs directly from PennyLane. code for the selected target, and the host runtime transfers it The dynamic library is loaded by the runtime executor and to the executor, where ORC links and loads it into the target launched to support the given workload on the target device, process. External symbols in the object are resolved at link which may be either local or remote. It can be used, for time against shared libraries already loaded in the executor. example, to transfer single messages between devices, or to The executor and its supporting runtime libraries are built trigger a handshake, initialization, or setup steps for a new separately for each target system, using a target sysroot when device. Everything from BLAS libraries, to custom compiled cross-compilation is required. functions can be loaded and executed here, with the data 2) The transport dialect: The transport dialect expresses available at the instruction point during workload execution. the instructions that set up and drive the latency-critical communication path. The dialect models a connection-oriented request/reply session between two endpoints and covers the C. Compiler infrastructure whole life cycle, from creating the session and setting up the To ensure that a program can run on multiple hardware data path, to sending rounds and tearing down. platforms, it is often necessary to utilize platform-agnostic Sessions are typed by their role, as programming abstractions and multiple separate compiler !transport.session<controller> or toolchains, each targeting a given platform. The LLVM !transport.session<coprocessor>. To bring up a ecosystem as a whole supports compilation for a wide variety of node’s communication, transport.connect establishes hardware (CPU, GPU, and custom devices), and so is naturally the connection and transport.exchange_keys positioned as the compiler infrastructure base for many vendor exchanges the memory-region information required ecosystems (e.g. AMD’s ROCm compiler hipcc, Intel’s oneAPI by the transport. Each operation also has a DPC++, etc.). As Catalyst is built atop MLIR, and hence LLVM, non-blocking form (transport.connect_async, we can tie into this tooling with minimal effort, allowing a transport.exchange_keys_async) completed by program to be written and compiled on a given system, and transport.await, allowing bring-up steps to overlap where executed on another target architecture and instruction set the placement permits. transport.establish_channel through cross-compilation. then arms the channel for the selected transport. Before Through PennyLane’s qjit interface, users can configure rounds begin, transport.set_message_sizes configures target-specific compilation options and use a Placement the request and reply sizes, and transport.start definition to assign specific parts of their program to different readies the session. On the controller, each round devices. To translate these high-level instructions into the interresolves its session with transport.get_session, mediate representation (IR), we introduce two complementary stages the payload with transport.stage_payload, dialects: the executor dialect handles the movement of code and transmits it with transport.post. The reply (determining which target device a subprogram runs on and is received into a buffer with transport.collect; how it gets there), while the transport dialect handles the where possible, transport.reply_slot exposes the movement of data (specifying the exchanges between those transport-owned reply buffer directly. On the coprocessor, subprograms during execution). transport.set_coprocessor_fn binds the processing 1) The executor dialect: The placement specifies both the function used by the selected target; depending on the target execution process and the machine locality of each program device, this function may process individual messages or component. A component may execute in the host process, in start a persistent service. Finally, transport.stop and a separate executor process on the same machine, or in an transport.destroy tear down the session. executor on another machine. During lowering, components 3) Compilation flow and the QEC pipeline: The frontend assigned to an executor are compiled separately, and the placement is serialized as an attribute on the generated MLIR executor dialect represents the transfer and invocation of the module, from which every pass below reads the nodes, their resulting code. The placement also supplies the compilation targets, and the transport. A Backline placement augments target and indicates whether supporting libraries are resolved Catalyst’s base pipeline with transport passes; if it also requests from the local installation or from a deployed software bundle. implicit QEC, the QEC encoding and physical-lowering passes This remote execution life cycle follows a sequence of are added. Both compose with Catalyst’s existing quantum dialect operations. First, a session is opened against an and classical transformations. The relevant passes run in the executor endpoint (executor.open). Next, a compiled object following order: is shipped to it (executor.send_binary), and entry points within that object are invoked (executor.launch) using • With implicit QEC enabled, the QEC passes operands and results that mirror the original host-side call. lower the program from a logical to a physical Alternatively, if relying on a pre-existing library on the target, qubit representation. convert-quantum-to-qecl executor.call can directly invoke an already-loaded symbol. raises the logical circuit into the qecl dialect, Finally, executor.close releases the session. inject-noise-to-qecl inserts the noise model used The remote execution is implemented using LLVM ORC by the demonstrations, and convert-qecl-to-qecp v2 [47]. A catalyst-executor process exposes the target followed by convert-qecp-to-quantum expands each process through ORC’s Executor Process Control (EPC) logical gate into its encoded physical circuit with a mechanism. Catalyst emits an object file containing machine round of error correction around it. Each round emits a

8

// Bring-up, emitted once into the program // entry. %ct = transport.create { backend_lib = "lib...cpu_verbs_controller.so", config = "dev=mlx5_1;gid=3", key = "gpu-coproc"} -> !transport.session<controller> ① transport.connect %ct { peer = "192.168.1.2", oob_port = 18560 : ui16} : !transport.session<controller> ② transport.exchange_keys %ct : !transport.session<controller> ③ transport.establish_channel %ct "rdma" : !transport.session<controller> ④ transport.set_message_sizes %ct { in_bytes = 8 : i64, out_bytes = 8 : i64} : !transport.session<controller> ⑤ transport.start %ct : !transport.session<controller> // One round, emitted in place of each decode. %s = transport.get_session {key = "gpu-coproc"} : !transport.session<controller> ⑥ transport.stage_payload %s, %syndrome { decoder_id = 0 : i32} : !transport.session<controller>, memref<3xi1> ⑦ transport.post %s : !transport.session<controller> %rep = transport.reply_slot %s : !transport.session<controller> -> memref<1xindex> ⑧ transport.collect %s, %rep : !transport.session<controller>, memref<1xindex>

Listing 2: Transport operations in MLIR for session bring-up and one request/reply round, as the compiler emits them on the controller. Circled numbers mark the steps that also appear in Fig. 3.

qecp.decode_esm_css operation carrying the code’s Tanner graph and the stabilizer type it decodes. • inject-transport-session reads the backline placement and inserts session creation and bring-up into compiler-generated setup functions, with the corresponding operations for teardown. For nodes assigned to executors, these life-cycle operations are placed in separately compiled components. • lower-decode-to-transport runs after bufferization and replaces each bufferized qecp.decode_esm_css operation with a session lookup and one transport round — stage_payload, post, and collect — taking the reply through reply_slot where the target device exposes a transport-owned buffer. • cross-compile-targets then compiles each target component to an object file for its selected compilation target. • dispatch-executor-targets replaces calls to target components assigned to executors with an executor.open/send_binary/launch sequence. • In the final LLVM conversion stage, convert-executor-to-llvm followed by convert-transport-to-llvm lowers the two

dialects to calls into their runtime C interfaces. When implicit QEC is enabled, convert-qecp-to-llvm subsequently lowers any remaining QEC physical operations. The implicit path currently supports the Steane code, selected with qec_code="steane". To use a different QEC code or to configure where in the circuit the rounds execute, the encoding and decoding can be explicitly written in the QNode with qp.backline.decode, which emits the transport round directly and does not pass through lower-decode-to-transport. This is the path the qLDPC demonstrations of Section V-B2 take. D. The runtime layer The compiled program reaches hardware through a shared library interface with a well-defined C ABI. By decoupling the compiled program from the underlying implementation via this standard ABI, users and developers can easily plug in a custom runtime tailored to their own bespoke hardware. 1) Session life cycle and memory: A target device keeps session state behind an opaque handle, and the interface is strictly ordered. For an RDMA-enabled target, each session goes through the following life cycle: first, a session is created from a named target library and role; connect establishes the out-of-band handshake and initializes local endpoint state; exchange_keys exchanges endpoint and memory-region metadata; establish_channel configures data movement between the local and peer memory; then after role-specific configuration, start readies the session for data exchange. After the required rounds, stop and destroy end the session. The details of memory provisioning and queue-pair state transitions depend on the target. Figure 3 shows the sequence of events within a pair of controller and coprocessor sessions. The memory is managed by the target device itself, and could allow for variable sizes and memory kinds, e.g. host RAM, HBM on GPUs, and DDR or BRAM on FPGAs. Regions are organized as ring buffers with a fixed slot count. 2) Network transport and fabric selection: Existing HPCfocused network-layer interfaces such as UCX [48] or the OpenFabrics Interfaces (OFI) [49] provide abstractions around the various low-latency RDMA layers, allowing one to express higher-level application logic without focus on the underlying network technologies. Although these interfaces are generally preferred for their portability and have successfully supported applications across a wide range of HPC systems, their abstractions inevitably introduce some overhead. When performance is measured in nanoseconds to single-digit microseconds, it can be more effective to directly target the appropriate subsystems. For the current state of this project, we opted for directly targeting and compiling against libibverbs [50], [51], the userspace library that has become the common interface to the Linux RDMA subsystem [52] and associated tooling, originally developed around the InfiniBand specification [53]. Given that many of the higher-level frameworks live above this layer (or related subsystems), potential overhead can be eliminated to achieve the fastest communication path. We anticipate this will have the widest impact, with the ability to adapt

9

(e.g. GPU-NIC-NIC-GPU). For systems that have baseline heterogeneity (CPU, GPU, RNIC), these libraries allow wellcreate create optimized application code to execute with best-in-class perforconnect_async + await ① connect mance on the given systems. However, for more asymmetric designs, such as pairing a GPU and RNIC with an FPGA exchange_keys ② exchange_keys as an endpoint, many of these libraries are not immediately accessible, and require bespoke wrapping to achieve the lowest establish_channel ③ establish_channel floor latency in communication. While it is possible to go set_coprocessor_fn ④ set_message_sizes below the libibverbs layer to directly tie into the RNIC drivers, start ⑤ start such as using GPU-initiated RDMA writes, we found the CPU to be sufficient for our latency needs, with results reported in per round Section V-C1. stage_payload (+ decoder_id) Furthermore, these libraries assume a performance model ⑥ post — syndrome that is different from our quantum use-case. Indeed, systems ⑦ built around these libraries achieve exceptional throughput and decode FLOPs by hiding latencies through the overlap of commureply — correction nications with compute, thereby minimizing upfront latency costs while maintaining high device utilization. This works for ⑧ collect offline quantum workloads (e.g. pre-/post-processing), as well Fig. 3: A representative remote setup and one decode as ones with loose timing needs, but is of limited applicability round. inject-transport-session emits the one-time for truly low-latency workloads, where device stalling while controller and coprocessor setup; dotted arrows show the awaiting completion of a task can cause catastrophic build-up out-of-band interactions between them. Asynchronous op- of corrections in the quantum system. In these cases, adding erations used during bring-up depend on the placement. additional endpoints to our networking setup can be considered. 3) Composing a system: It is worth stating the composability lower-decode-to-transport emits the solid request/reply of these layers explicitly, since it is what makes one workload round in place of each implicit decode. Circled numbers portable across different configurations. The transport, the correspond to the controller-side operations in Listing 2. controller platform, the coprocessor platform, their locality, and the definition of the coprocessor function are independent choices, summarized in Table I. Moving a prototype from laptop the associated tooling to work with other (and often more onto HPC-grade hardware means changing entries within the accessible) networking abstractions (e.g. UCX, OFI, ef_vi, table, without changing the workload. etc.) in follow-up work. In the following section we demonstrate a variety of setups, With this infrastructure, and to build the most widely including local and remote endpoints, and controllers and compatible platform, we opted for RoCE v2 [22], which coprocessors across multiple architectures, including the use supports transmission using Ethernet packets via compatible of an FPGA as a controller. This is to emphasize that on a real RNICs and related FPGA infrastructure [54]. In addition, with quantum hardware system, the controller may not necessarily the libibverbs layer, this allows us to immediately target any be a CPU host processor. Indeed, qubit control, measurement compatible RNIC via the upstream providers supported in the data acquisition, and the primary real-time decoding commonly Linux RDMA subsystem, such as those from NVIDIA Mel- reside on FPGAs, and ultimately on ASICs. This is because lanox (mlx5), AMD Pensando (ionic), Broadcom (bnxt_re), the real-time control tier (Tier 1) requires a deterministic subas well as many others, including the fallback software-defined microsecond budget, which bare CPUs are unable to provide. emulation Soft-RoCE (rxe) for testing. With this wide breadth The key claim that we seek to demonstrate with Backline is of support, we can target any such provider as a testbed that the same Python program can scale from a fully local, for building out quantum computing workloads across many single-process execution to an FPGA controller with a remote provider platforms, as well as developing and testing locally CPU or GPU coprocessor, changing only the placement while using the emulation layer. leaving the circuit untouched. For the simplest prototyping use case, below even the Bringing everything together, an example workload’s life emulation layer, the same session interface is implemented by a cycle through our Backline design is outlined in Fig. 4. local target that replaces one-sided writes with in-process copies between the two roles’ regions. This requires no verbs device V. D EMONSTRATIONS AND EVALUATION at all, and is the mechanism behind the local development path The demonstrations and measurements in this section are of Section IV-B3. all in the accompanying repository1 , which includes detailed Besides portable networking abstraction libraries such as installation and execution instructions, along with a cross-build UCX or OFI, shared memory (SHMEM) libraries such as system used for producing the stack deployed to each remote NVSHMEM and rocSHMEM are also highly relevant. These machine. The demonstrations characterize the performance sit closer to the hardware and are tuned for low-latency 1 https://github.com/PennyLaneAI/backline RDMA communication in symmetric multi-device ecosystems controller

coprocessor

10

Fig. 4: Life cycle of a heterogeneous workload in Backline across a host, FPGA controller, and GPU coprocessor. Numbers indicate execution order. Steps 0-2 occur once per session. Steps 3-11 are one correction round and repeat for every decode. Steps 12-14 return results and tear down.

Axis

Choices

Transport

memcpy – in-process copies; no fabric hardware required rdma – RoCE v2 over libibverbs, any supported RNIC provider

Controller

CPU – Any CPU (for rdma, with connection to a NIC) FPGA – Application on APU, RDMA through e.g. Xilinx ERNIC engine

Coprocessor

CPU – per-message callback

Coprocessor function

C++ library – most control, referenced by symbol Triton kernel – written in Python, AOT-compiled for the GPU

GPU – launch-once persistent kernel

Tanner graph – Hx/Hz as input, BP-OSD decoder generated Locality

Local in-process – no executor Local out-of-process – a local executor Remote machine – an executor and a bundle deployed

TABLE I: The configuration axes of a Backline placement. Each row can be varied independently without changing the QNode.

of our benchmark setup while illustrating the breadth of the abstractions introduced in Section IV, spanning configurations from a laptop with no specialized hardware to an FPGA controller driving a GPU decoder across a fabric, with the decoder function either pre-compiled, written directly in Python, or generated from a code’s parity-check matrices. A. Hardware setup The demonstrations are arranged in four hardware configurations with increasing setup requirements, allowing users to test progressively more capabilities. The first configuration only requires a single Linux CPU machine and nothing else. For demonstrations using local executors, the entire PennyLane and Catalyst program is compiled and executed on the local machine; for those using remote executors, the local machine instead handles compilation and orchestration, while the target components execute on one or more remote machines. The second configuration additionally requires a verbs device on the same machine; kernel Soft-RoCE is sufficient for this purpose, making an RDMA NIC optional. The third adds the server with an RDMA NIC and a GPU. For our demonstrations, we relied on an AMD Threadripper PRO 9975WX CPU with an AMD

11

Instinct MI210 GPU (PCIe 4.0) on an ASRock WRX90 WS EVO motherboard carrying a ConnectX-7 RNIC (PCIe 5.0). The fourth configuration adds an FPGA board, an AMD Xilinx VPK120 board [55] running the ERNIC IP block [54] with a customized reference design2 , whose processing system is a pair of aarch64 Arm Cortex-A72 cores. The FPGA board is connected to the server via the RDMA NIC using a 100 Gb Direct Attach Copper cable. For the FPGA board running the controller node, we implemented a hardware-handshake (HWHS) engine in the programmable logic of the FPGA adjacent to the ERNIC RoCE engine. This allows us to avoid issuing work requests from the processing system that would place the application processing unit (APU), Linux scheduler, and repeated memory-mapped I/O (MMIO) transactions on the latency-critical path. Session setup remains a software control-plane operation, but the per-round operations of posting a work-queue entry (WQE), ringing the send-queue doorbell, detecting the reply, and measuring the round-trip latency are all performed by a finite-state machine inside the HWHS engine. Figure 5 gives an overview of the HWHS architecture and its connection to the ERNIC. The APU configures the engine through AXI4-Lite, while the HWHS orchestrator handles request posting, doorbells, reply observation, and round-trip time (RTT) measurement in programmable logic. WQEs, queuepair state, and reply data are kept close to the orchestrator so that repeated round-trip operations do not return to host software.

1

import pennylane as qp

2 3

STEANE_CPU_DECODER_LIB_PATH = "/path/to/ libsteane_coprocessor_cpu.so"

4 5 6

steane_decode = qp.CoprocessorFunction( name="steane_coprocessor", lib_path= STEANE_CPU_DECODER_LIB_PATH)

7 8 9 10 11 12

ctrl = qp.Controller( hardware="cpu", device=qp.device("null.qubit", wires=3)) coproc = qp.Coprocessor( hardware="cpu", coprocessor_fn=steane_decode)

13 14 15 16 17

dev = qp.Backline(controller=ctrl, coprocessors=[coproc], transport="memcpy", qec_code="steane")

18 19 20 21 22 23 24 25 26 27 28

@qp.qjit(capture=True) @qp.set_shots(10) @qp.qnode(dev, mcm_method="one-shot") def ghz(): qp.Hadamard(0) qp.CNOT([0, 1]) qp.CNOT([1, 2]) return qp.sample([qp.measure(0), qp.measure(1), qp.measure(2)])

29 30

print("samples:", ghz())

Listing 3: The smallest complete program: local CPU to local CPU over the memcpy transport. Asking for qec_code="steane" is what turns the three logical gates into encoded circuits with a correction round around each, and each of those rounds into a message to the coprocessor.

ERNIC RoCE engine

B. Demonstrations inbound RDMA write

WQE fetch

Shared SQ/Data BRAM (AXI/SmartConnect)

direct doorbells

SmartConnect + AXI BRAM controller port A

WQE write

HWHS orchestrator post, doorbell, RTT

WQE RAM QP/WQE state

native port B

Local Reply BRAM port A: AXI write port B: native monitor

free-running counter + trace RAM HWHS engine

AXI4-Lite config/status

AXI4-Lite

PS/APU software control plane

Fig. 5: Architecture of the VPK120 hardware-handshake engine, which connects the APU control plane, ERNIC, local request/reply storage, and hardware orchestrator. To reduce system-related overheads the server treats each CPU chiplet as an individual NUMA domain, via the NPS4 (NUMA Per Socket) BIOS option, and the RNIC and GPU are placed in PCIe slots on the same domain so that traffic between them stays within one chiplet. The controller process is pinned to a core on that domain and runs at real-time priority. 2 Board images and setup instructions are available at https://github.com/ PennyLaneAI/backline-vpk120.

The demonstrations span four choices: where each node runs and on what kind of processor, which transport carries a round between them, whether the error-correction round is inserted by the compiler or written out by the user, and where the decode function comes from. Each of these choices is explained below; to illustrate, Listing 3 first shows our smallest demo in full: a logical GHZ state on three qubits, corrected by a pre-compiled decoder in the same process. A few points worth noting: 1) Both the controller and coprocessor are purely local, so no remote execution (or configuration) is required. 2) qec_code="steane" turns the three logical gates into encoded circuits with a correction round around each, and each of those rounds into a message to the coprocessor. 3) coprocessor_fn=steane_decode picks a precompiled "steane_coprocessor" function to perform the per-message decoding. 4) transport="memcpy" means syndrome and reply messages are transferred purely locally; no RDMA device is required. 1) Placement: As defined in Section IV-B, a Placement contains one controller, zero or more coprocessors, and the transport between them. The nodes. Across the demonstrations the controller is a local CPU, a remote CPU, or a remote FPGA, and the coprocessor is a local CPU or a remote GPU. Listing 4 shows four of the placements we showcase in our demos.

12

1 2 3 4 5 6 7 8 9

# demo 1 -- both roles are local ctrl = qp.Controller( hardware="cpu", device=qp.device("null.qubit", wires=3)) coproc = qp.Coprocessor( hardware="cpu", coprocessor_fn=steane_decode) dev = qp.Backline(controller=ctrl, coprocessors=[coproc], transport="memcpy", qec_code="steane")

10 11 12 13 14 15 16 17 18 19 20 21 22

1 2 3 4 5 6 7 8 9 10

# demo 1a -- the same pair, over a verbs device ctrl = qp.Controller( hardware="cpu", device=qp.device("null.qubit", wires=3), init_args={...}) coproc = qp.Coprocessor( hardware="cpu", coprocessor_fn=steane_decode, endpoint=qp.Endpoint(...), # soft-RoCE init_args={...}) dev = qp.Backline(controller=ctrl, coprocessors=[coproc], transport="rdma", qec_code="steane")

# (a) implicit - demos 1, 1a, 3, 4, 5 # a logical circuit; the compiler inserts # the encoding and the decode round dev = qp.Backline(..., qec_code="steane")

11

# (b) explicit - demo 2 z_syn, x_syn = extract_syndromes() corr_z = qp.backline.decode(x_syn, decoder_id=0) corr_x = qp.backline.decode(z_syn, decoder_id=1) apply_correction(corr_x, qp.X) apply_correction(corr_z, qp.Z)

12 13 14 15 16 17 18 19 20 21 22

# (c) runtime calls - demo 2a # what (b) is made of session = qp.runtime_call( GET_SESSION, ROLE_CONTROLLER, "gpu-coproc") qp.runtime_call(STAGE_PAYLOAD, session, packed, PACKED_BYTES, decoder_id) qp.runtime_call(POST, session, WORK_ITEM) _status, correction = qp.runtime_call( COLLECT, session, PACKED_BYTES, out_bytes=PACKED_BYTES)

23 24 25 26 27 28 29 30 31 32 33 34 35 36 37

# demo 3 -- remote CPU controller, remote GPU ctrl = qp.Controller( hardware="cpu", remote=True, executor_options={...}, init_args={...}) coproc = qp.Coprocessor( hardware="gpu", remote=True, coprocessor_fn="gpu_steane_launcher", endpoint=qp.Endpoint(...), # fabric address executor_options={...}, init_args={...}) dev = qp.Backline(controller=ctrl, coprocessors=[coproc], transport="rdma", qec_code="steane")

38 39 40 41 42 43 44

# demo 4 -- the controller moves to the FPGA; # everything else is demo 3 unchanged ctrl = qp.Controller( hardware="fpga", remote=True, executor_options={...}, init_args={...})

Listing 4: Four placements for the same program. Demo 1 and demo 1a differ in the transport alone; demo 1a to demo 3 changes the machine and the device class together; demo 3 to demo 4 changes only the controller. The qp.Backline line is the placement itself, carrying the two nodes and the transport between them. The ellipses in the constructors abbreviate testbed-specific values declared once in the repository’s placement.py for easy definition and re-use, which are shared by every demonstration.

A node is characterized by the machine its compiled code runs on and what it loads there. The following are arguments used to configure a node: remote decides whether the code is shipped instead of run in-process, executor_options describes how to reach the machine over a network, endpoint specifies the coprocessor’s address on the RDMA fabric, init_args carries target-specific setup, and hardware is paired with the placement’s transport to resolve a concrete target device library. The transport. The two transports of Section IV-B3 divide the demonstrations. Demo 1 uses memcpy, which needs no verbs device at all and is the one to run first; every other demonstration uses rdma. Demo 1a is the useful step between them, being demo 1 over a Soft-RoCE device: the verbs path is genuine while the transport beneath it is still software, so

Listing 5: The same round at three levels. In (b) the syndrome is extracted by the user and the correction applied to live qubits, with decoder_id choosing which of the coprocessor’s decoders serves it. In (c) that one call is replaced by the four it expands to.

it establishes that the path works. Given a real RDMA NIC, the same placement runs over that NIC’s loopback on one machine, or over a wire between two. 2) Implicit and explicit decoding: With qec_code="steane" given on the device, the user defines only a logical circuit, and points the compiler at the passes and libraries that do the rest: error correction is applied automatically within the stack, with the encoding introduced at compile time by the QEC lowering passes and the decoding performed on the coprocessor. Each logical gate is expanded into its physical circuit with a round of error correction around it — extract the stabilizers, decode the syndrome, apply the correction — none of which the user needs to specify. Alternatively, users can express these steps explicitly in their Python program. The user can encode the logical circuit by hand, in effect writing the physical circuit to be compiled and executed through Backline, then measure the syndromes, decode them, and apply the correction explicitly. An explicit call to qp.backline.decode sends a syndrome to the coprocessor and blocks until the correction comes back. Where several decoders are registered, a decoder_id (which is part of the payload) selects between them. decode expands into a series of local runtime calls in the controller’s process. The same operations are accessible directly through the qp.runtime_call interface of Section IV-B: get_session finds the controller’s open session with the named coprocessor, stage_payload writes the packed syndrome into the outgoing slot, post hands that slot to the transport, and collect blocks until the correction lands in the reply slot. Demo 2a reproduces demo 2’s result with the exchange written out this way. The three levels differ only in how much is left to the compiler. 3) The decoder function: The demonstrations exercise all three routes of Section IV-B2: a pre-compiled decoder for the [[7, 1, 3]] Steane code, used on both a CPU and a GPU

13

coprocessor; a Steane decoder written in Python as a Triton a median of 4.5 µs, a mean of 4.448 µs, a standard deviation kernel; and a belief-propagation decoder with ordered-statistics of 0.159 µs, and a P99.999 of 5.255 µs. A full description of post-processing [56] generated from the parity-check matrices the FPGA timing setup and further experiments is given in of the [[13, 1, 3]] hypergraph-product code, run for 10 iterations. Appendix A. The baseline measurements suggest that this framework is Several decoders may be registered on one coprocessor and selected per round by the decoder_id travelling with the suitable for the synchronous co-processing tier described in payload, which is how a single coprocessor serves the X and Section II with room to spare. In particular, the CPU path is not only 2.2 µs faster in the median but also over seven times Z stabilizer types with a decoder each. Table II lists the demonstrations available in the repository, tighter in the spread. Against a budget that has to be met on their respective choices on the above, and the route each one every round rather than on average, this spread in particular is decisive in whether a decoder is admissible, and on that takes to its decoder. measure the commodity CPU path is the stronger of the two. Controller local CPU local CPU remote CPU remote CPU remote CPU remote FPGA remote FPGA

Coprocessor local CPU local CPU remote GPU remote GPU remote GPU remote GPU remote GPU

Transport memcpy rdmaa rdma rdma rdma rdma rdma

Decode implicit implicit explicit runtime implicit implicit implicit

Decode fn pre-compiled pre-compiled Triton BP-OSD Triton BP-OSD pre-compiled pre-compiled Triton Steane

TABLE II: All demonstrations included in the repository, as combinations of the choices of Sections V-B1 to V-B3. Decode is how the round is expressed: inserted by the compiler (implicit), written with qp.backline.decode (explicit), or written out as runtime calls (runtime). Decode fn is where the decoder comes from: a symbol in a pre-compiled CPU or GPU library, generated from the code’s parity-check matrices (Triton BP-OSD), or directly through a Triton kernel (Triton Steane). a over kernel Soft-RoCE, so no RDMA NIC is required.

2.6

round-trip time (µs)

1 1a 2 2a 3 4 5

CPU echo

2.4

GPU echo 5.0 4.5 0.0

0.2

0.4

0.6

0.8

1.0 ×106

samples

(a) Round-trip time against sample number. CPU echo mean=2.312 µs std= 0.02137 µs med= 2.305 µs min= 2.285 µs max= 2.59 µs t0 = 4.64 µs P99= 2.42 µs P99.9= 2.445 µs P99.999= 2.475 µs

105 104 103 102

C. Measurements 1) Baseline: To accurately measure the overheads, we built an end-to-end latency evaluation benchmark (bench.py) to determine the floor latency of the ecosystem using the hardware configuration specified in Section V-A. For the baseline case, the FPGA issues a 16-byte payload (including an 8-byte syndrome) as a one-sided RDMA write to a given target (CPU or GPU memory). Through the hardware-handshake engine, the programmable logic (PL) directly posts that write and detects the reply, so there is no host involvement in the round trip. To ensure a fast response, the target coprocessor polls on the expected memory buffer, then in the CPU case echoes the reply back directly through a one-sided write to the FPGA, or in the GPU case signals the CPU, which then issues the one-sided write back. The FPGA controller times the round trip in its own clock domain. Figure 6 shows the timing and distribution of 106 − 1 round trips initiated from the FPGA controller to and from a CPU and GPU coprocessor device, as a time series and as a distribution. Both panels show the steady state, with a single initial warm-up round trip excluded.3 . From the data, the CPU path produced a median of 2.305 µs, a mean of 2.312 µs, a standard deviation of 21.4 ns, and a P99.999 tail at 2.475 µs; the GPU path measured 3 The first warm-up round trip costs 4.64 µs on the CPU path and 9.27 µs on the GPU path, in both cases far outside the steady-state distribution that follows it. Full results are shown in Appendix A

samples

101 100 2.30

2.35

2.40

2.45

2.50

2.55

2.60

GPU echo mean=4.448 µs std= 0.1591 µs med= 4.5 µs min= 4.2 µs max= 5.285 µs t0 = 9.27 µs P99= 4.985 µs P99.9= 5.125 µs P99.999= 5.255 µs

105 104 103 102 101 100 4.2

4.4

4.6

4.8

5.0

5.2

5.4

5.6

round-trip time (µs)

(b) Round-trip time distribution.

Fig. 6: Baseline steady-state round-trip latency over 106 − 1 rounds from FPGA controller to and from CPU or GPU coprocessor. Timing is measured by the FPGA clock domain, with the first warm-up sample excluded from all plots. Statistical quantities are given in the legend of each distribution plot, with the excluded warm-up value given by t0 . 2) Decoder runtime: Using the same benchmarking script and setup, we measure the end-to-end round-trip timing from FPGA controller to a CPU or GPU coprocessor device including a decoder execution. Table III summarizes the time it takes for the three decoders of Section V-B3, alongside an ‘echo’ row giving the baseline of Section V-C1 on the same path. The BP-OSD decoder is a naive implementation; there are

14

many optimizations possible, including using hardware-friendly normalized min-sum to approximate more expensive operations on GPUs, which could be implemented directly within the Triton kernel. The point of the Backline platform is to enable quick turnarounds for such experimentation.

coprocessor 4.5 µs over 106 − 1 rounds (Fig. 6). Both sit inside the synchronous co-processing tier of Section II, which is the regime that backup decoding and auxiliary real-time processing occupy.

Device Coprocessor function

Mean Median

P99

P99.9

Max

A. Choosing a coprocessor

CPU CPU

2.312 2.311

2.420 2.410

2.445 2.445

2.590 2.545

Echo Steane decode

2.305 2.305

An enterprise-grade GPU is not necessarily a prerequisite for real-time QEC. From our measurements, the CPU path is GPU Echo 4.448 4.500 4.985 5.125 5.285 not only faster, but substantially more predictable. Against a GPU Steane decode 4.412 4.405 4.945 5.105 5.265 deterministic budget, the tail is a crucial factor in deciding GPU qLDPC BP-OSD 31.109 31.075 32.410 32.680 33.260 whether an external decoder is admissible, and on that measure decode the CPU path is stronger by a wide margin. This also lowers the TABLE III: Steady-state software-loop RTT statistics in mi- hardware cost of entry for users who want to begin prototyping. croseconds. Round 0 is excluded. This is not an argument that GPUs are unnecessary; the device should follow what the workload requires. Where the From the results, we observe that for very simple decoders, decode is cheap relative to the round trip, a CPU coprocessor is the decoding time is virtually free compared to the commu- the better endpoint, since the GPU would be paying a transport nication overheads. For more complex decoders, the decoder penalty for parallelism it does not need. Where the decode may take longer than the transport overhead. For prototyping itself dominates (e.g. for larger decoders, batched processing, and fast R&D, defining the decoder directly in Python through or neural-network based decoders), the transport difference Triton provides the least resistance, but for best performance, stops being the deciding term and the GPU’s throughput may users should write a tuned and optimized version directly as a be what matters. Backline enforces neither choice, leaving it CPU function or GPU kernel. Backline is designed to facilitate to the user to try both. both modes, and easy transition from one to the other. VI. C ONCLUSIONS AND OUTLOOK

B. Fabric is a moving target

RoCE v2 gave us a latency floor in the single-digit miA quantum computer is not purely quantum. A fault-tolerant quantum computer requires an amalgamation of and complex croseconds, and that floor is what puts the synchronous cointeraction between quantum and classical systems, which processing tier within reach. Similar experiments [57] with encode, decode, measure, and react on the timescale of other HPC-grade networking also establish a comparable floor. the quantum hardware itself. The difficulty of building and This is just the starting point. Emerging interconnects are programming such systems is now a first-order constraint on converging on improved performance in every direction, for the field. The solution to this is not to harden a single fast example UALink [58] for scaling up accelerator-to-accelerator path, but to give researchers a way to express a workload once links within a rack, and Ultra Ethernet [59] for scaling out and place it across whichever classical devices the problem fabrics with transport optimizations surrounding both small latency-sensitive messages and bulk data throughput. We expect requires and that are available to them. Backline is our answer to that. It adds a heterogeneous these to be important for FTQC infrastructure, enabling both compilation and runtime layer to PennyLane and Catalyst high performance and flexible topologies. The abstractions in Section IV were designed with exactly that spans the required stack: a declarative frontend, in which a placement is defined by a controller, its coprocessors, this flexibility in mind. The transport dialect describes only and the transport between them; an operational compilation the life cycle of a session; a new fabric is a new runtime target layer, in which the executor and transport dialects turn behind the same C ABI, and reaches the user as a new value that placement into explicit sessions, staging and launching for a single transport= keyword. operations as well as lowering them alongside the other QEC Currently the dialects and runtimes do not abstract over the passes in MLIR; and a runtime layer, whose unified ABI message sizes. The message size is fixed, and every round in the makes swapping target devices and transport modes seamless. session uses those exact byte counts. This is sufficient for the Section V exercised these layers against real hardware. The workloads shown here today, where the syndrome has a width same GHZ program ran on a single workstation over memcpy, known at compile time and a fixed slot allows registering the on that workstation over rdma, across a remote CPU controller memory once and staging it, without allocation on the critical and GPU coprocessor, and finally from a Xilinx VPK120 path. However, this may be unsuitable for cases where the controller to that same GPU, with the differences between payload varies from round to round, such as a decoder that configurations confined solely to the placement. returns a variable-length correction, or a controller batching a The baseline measurements establish the achievable latency number of rounds that is only known at execution time. The regime: driven from the FPGA controller under the hardware- current limitation of the transport layer is therefore support handshake engine, the steady-state round trip to a CPU for variable payload sizes without giving up the pre-registered, coprocessor measures a median of 2.305 µs and to a GPU allocation-free fast path; we leave this to future work.

15

Fig. 7: An open architecture pathway for FTQC and adjacent quantum systems. Blue denotes what Backline supports today: the application and compilation layers it is built into, the orchestration and runtime layers it supplies, and real-time QEC as the workload driving them. Pink denotes next steps: high-throughput data pre- and post-processing, quantum sensing platforms, and reaching further down through control and data acquisition to the physical instruments.

16

C. Scaling up With hardware system scale-ups happening to accommodate the increasing demand for AI, quantum software layers can likely benefit from integration with such systems. Rackscale designs such as AMD Helios (MI455X GPUs, EPYC CPUs, Pensando RNICs) should be implicitly compatible with the Backline programming abstractions, since a rack of accelerators behind an RDMA-capable NIC is the same object that our placement model already describes, only with more coprocessors in it. This admits direct multi-device targeting and compilation from a single program. We see this being beneficial not just for simulation, but for full-system emulation of FTQC control stacks whose hardware designs are approaching reality.

researchers already use. Getting a decoder from a Python prototype onto real low-latency hardware should be a matter of changing where it runs, with little or no need to rewrite the program. We envision Backline as a platform for researchers and engineers across the quantum computing stack. Quantum architecture and algorithm researchers can explore new designs for applications and error correction with minimal friction. Quantum hardware developers can use a cohesive, widely adopted environment to prototype control systems. Systems engineers and classical computing vendors can rely on a userfriendly abstraction to seamlessly orchestrate heterogeneous computing environments. ACKNOWLEDGEMENTS

D. Scaling down Access to enterprise-grade hardware is frequently the binding constraint for a research group, and local development is what makes iteration on workloads, applications, and SDK layers fast. We deliberately kept a continuous path from a laptop to the full system. The same program runs over memcpy on one machine, over Soft-RoCE [60] against a software RDMA device with no NIC present at all, over verbs on a workstation with a commodity RNIC, or over an FPGA’s hardware engine; moving between these requires only a change to the placement rather than to the workload. Another part of enabling this accessible approach is keeping the framework open, Python-native, and infrastructure-agnostic. Backline is built and shipped alongside the open-source PennyLane and Catalyst codebases, with the frontend API, the MLIR dialects, and the runtime targets all available and extensible. We consider the hardware side to be extensible on the same terms: there is no restriction on the infrastructure or vendors for any of the CPU, GPU, and FPGA components, while open-source RNIC designs [61], [62] make it possible for research groups to implement and test different FPGA setups. Prototyping and co-designing QEC codes, decoders, and transport is entirely possible within Backline. E. Down to silicon The real-time control tier of Section II does not end at the FPGA. As we noted in Section IV-D3, qubit control, data acquisition, and the primary decoder sit on FPGAs and ultimately on ASICs, since this tier requires a bounded, submicrosecond worst-case latency that general-purpose hosts cannot provide. Figure 7 summarizes in blue what Backline enables today and, in pink, the remaining steps needed to cover the whole quantum system in the near future. Platforms for QEC infrastructure remain highly bespoke, and as we argued at the outset the field is unlikely to be served by freezing a standard around today’s hardware. What it does need is vendor-agnostic options, and interfaces that evolve with the field while remaining stable enough to build against. Backline is a stride in that direction: an architecture for workloads spanning CPUs, GPUs, and FPGAs, with transport abstractions exposed at compile time and implemented under a runtime interface, reached from a Python frontend that quantum

We acknowledge support from AMD for providing GPU hardware to support this project. Additionally, we thank Jose M. Monsalve-Diaz, Paul Hartke, Yasuko Eckert, and Muhammad Osama for their discussions on the software side, as well as Rip Sohan, Rowan Lyons, Chris Neely, and Mike Crowley for assistance on the hardware side. We also acknowledge significant help from the Xanadu software team members Ali Asadi and David Ittah for software contributions, discussions and reviews, Jason Selby and Isaac De Vlugt for manuscript review, as well as IT team members Nabil Dib, Andy Yen, Andrew Hodder, Sanchit Bapat, Seun Shoga, and Raymond Baksh for infrastructure setup and configurations. R EFERENCES [1] M. Gort and J. H. Anderson, “Combined architecture/algorithm approach to fast fpga routing,” IEEE Transactions on Very Large Scale Integration (VLSI) Systems, vol. 21, no. 6, pp. 1067–1079, 2013. [2] P. Rao, D. Trenev, J. Gonthier, T. Patti, S. Stern, T. Takeshita, Y. Alexeev, C. Lin, S. McArdle, J. Lietz, K. Klymko, E. Rrapaj, N. Tubman, K. Svore, P. Komar, and E. Kyoseva, “Performance model for hybrid quantum-classical workflows,” 2026. [Online]. Available: https://arxiv.org/abs/2607.15426 [3] A. Shehata, B. Austin, T. Beck, L. Burgholzer, A. Chernoguzov, S. Churchill, A. Delgado, Y. Eckert, J. Heckey, K. Kissell, K. Klymko, J. Moles, T. Naughton, L. J. O’Riordan, C. O. Pauyac, G. Prawiroatmodjo, E. Rrapaj, J. Schindler, L. Schulz, S. Stern, T. Takeshita, M. Tsuji, A. Wennersteen, T. Humble, and M. Schulz, “Quantum-hpc software stacks and the openqse reference architecture: A survey,” 2026. [Online]. Available: https://arxiv.org/abs/2604.20912 [4] M. Mohseni, A. Scherer, K. G. Johnson, O. Wertheim, M. Otten, N. Anand, N. A. Aadit, Y. Alexeev, G. Ben-Shach, K. M. Bresniker, K. Y. Camsari, B. Chapman, S. Chatterjee, S. Chowdhury, G. A. Dagnew, T. Dvir, A. Esposito, F. Fahim, M. Ferguson, M. Fiorentino, A. Gajjar, K. Gratsea, G. Gyawali, C. Heiter, A. H. Z. Kavaki, A. Khalid, X. Kong, B. Kulchytskyy, E. Kyoseva, R. Li, P. A. Lott, I. L. Markov, R. F. McDermott, L. Morais, G. Pedretti, P. Rao, E. Rieffel, A. Silva, J. Sorebo, P. Spentzouris, Z. Steiner, B. Torosov, D. Venturelli, R. J. Visser, Z. Webb, X. Zhan, Y. Cohen, P. Ronagh, A. Ho, R. G. Beausoleil, and J. M. Martinis, “How to build a quantum supercomputer: Scaling from hundreds to millions of qubits,” 2026. [Online]. Available: https://arxiv.org/abs/2411.10406 [5] J. Kim, S. Lee, B. Johnston, and J. S. Vetter, “Iris: A performanceportable framework for cross-platform heterogeneous computing,” IEEE Transactions on Parallel and Distributed Systems, vol. 35, no. 10, pp. 1796–1809, 2024. [6] Y. L. Fur, E. Egger, H.-Y. Hu, V. Russo, W. J. Zeng, and R. LaRose, “Opportunities and challenges in scaling quantum error detection on hardware,” 2026. [Online]. Available: https://arxiv.org/abs/2605.02861 [7] S. Seelam, J. M. Chow, A. Córcoles, S. Sheldon, T. Mittal, A. Kandala, S. Dague, I. Hincks, H. Horii, B. Johnson, M. Le, H. Jamjoom, and J. M. Gambetta, “Reference architecture of a quantum-centric supercomputer,” 2026. [Online]. Available: https://arxiv.org/abs/2603.10970

17

[8] A. Bhatia, F. Khomh, B. Adams, and A. E. Hassan, “An empirical study of self-admitted technical debt in machine learning software,” ACM Trans. Softw. Eng. Methodol., Dec. 2025, just Accepted. [Online]. Available: https://doi.org/10.1145/3785001 [9] M. Henning, “The rise and fall of corba,” Commun. ACM, vol. 51, no. 8, pp. 52–57, Aug. 2008. [Online]. Available: https://doi.org/10.1145/1378704.1378718 [10] A. L. Russell, “The internet that wasn’t,” IEEE Spectrum, vol. 50, no. 8, pp. 39–43, 2013. [11] S. A. Caldwell, M. Khazraee, E. Agostini, T. Lassiter, C. Simpson, O. Kahalon, M. Kanuri, J.-S. Kim, S. Stanwyck, M. Li, J. Olle, C. Chamberland, B. Howe, B. Schmitt, J. G. Lietz, A. McCaskey, J. Ye, A. Li, A. B. Magann, C. I. Ostrove, K. Rudinger, R. Blume-Kohout, K. Young, N. E. Miller, Y. Xu, G. Huang, I. Siddiqi, J. Lange, C. Zimmer, and T. Humble, “Platform architecture for tight coupling of high-performance computing with quantum processors,” 2025. [Online]. Available: https://arxiv.org/abs/2510.25213 [12] Y. Chen, C.-Y. Luan, P. Zheng, X. Zeng, J.-Y. Hou, Z. Fu, Y. Jin, F. Wang, G. Yang, and D. Lv, “Real-time quantum error correction system stack: Architecture, algorithms, and engineering practice,” 2026. [Online]. Available: https://arxiv.org/abs/2605.30765 [13] L. developers. (2026) LLVM Developer Policy. Accessed: 2026-07-24. [Online]. Available: https://llvm.org/docs/DeveloperPolicy.html [14] M. Copik, M. Chrapek, L. Schmid, A. Calotoiu, and T. Hoefler, “Software resource disaggregation for hpc with serverless computing,” in 2024 IEEE International Parallel and Distributed Processing Symposium (IPDPS), 2024, pp. 139–156. [15] C. Guo, J. Xu, and M. Zukerman, “Disaggregated architectures and the redesign of data center ecosystems: Scheduling, pooling, and infrastructure trade-offs,” IEEE Communications Magazine, vol. 64, no. 6, pp. 111–117, 2026. [16] M. Abadi, P. Barham, J. Chen, Z. Chen, A. Davis, J. Dean, M. Devin, S. Ghemawat, G. Irving, M. Isard, M. Kudlur, J. Levenberg, R. Monga, S. Moore, D. G. Murray, B. Steiner, P. Tucker, V. Vasudevan, P. Warden, M. Wicke, Y. Yu, and X. Zheng, “Tensorflow: a system for large-scale machine learning,” in Proceedings of the 12th USENIX Conference on Operating Systems Design and Implementation, ser. OSDI’16. USA: USENIX Association, 2016, p. 265–283. [17] A. Paszke, S. Gross, F. Massa, A. Lerer, J. Bradbury, G. Chanan, T. Killeen, Z. Lin, N. Gimelshein, L. Antiga, A. Desmaison, A. Köpf, E. Yang, Z. DeVito, M. Raison, A. Tejani, S. Chilamkurthy, B. Steiner, L. Fang, J. Bai, and S. Chintala, PyTorch: an imperative style, highperformance deep learning library. Red Hook, NY, USA: Curran Associates Inc., 2019. [18] R. Frostig, M. Johnson, and C. Leary, “Compiling machine learning programs via high-level tracing,” 2018. [Online]. Available: https://mlsys.org/Conferences/doc/2018/146.pdf [19] The jax authors, “Pallas: A jax extension for writing low-level, highperformance custom kernels,” https://docs.jax.dev/en/latest/pallas/index. html, accessed: 2026. [20] P. Tillet, H. T. Kung, and D. Cox, “Triton: an intermediate language and compiler for tiled neural network computations,” in Proceedings of the 3rd ACM SIGPLAN International Workshop on Machine Learning and Programming Languages, ser. MAPL 2019. New York, NY, USA: Association for Computing Machinery, 2019, p. 10–19. [Online]. Available: https://doi.org/10.1145/3315508.3329973 [21] Y. Chen and C. Xia, “From naive cuda to triton: A systematic evaluation of ai-era hpc operator development,” in Proceedings of the 40th ACM International Conference on Supercomputing Workshops, ser. ICS Workshops ’26. New York, NY, USA: Association for Computing Machinery, 2026, p. 60–64. [Online]. Available: https://doi.org/10.1145/3774895.3812198 [22] C. Guo, H. Wu, Z. Deng, G. Soni, J. Ye, J. Padhye, and M. Lipshteyn, “Rdma over commodity ethernet at scale,” in Proceedings of the 2016 ACM SIGCOMM Conference, ser. SIGCOMM ’16. New York, NY, USA: Association for Computing Machinery, 2016, p. 202–215. [Online]. Available: https://doi.org/10.1145/2934872.2934908 [23] A. Kalia, M. Kaminsky, and D. G. Andersen, “Design guidelines for high performance RDMA systems,” in 2016 USENIX Annual Technical Conference (USENIX ATC 16). Denver, CO: USENIX Association, Jun. 2016, pp. 437–450. [Online]. Available: https: //www.usenix.org/conference/atc16/technical-sessions/presentation/kalia [24] V. Bergholm, J. Izaac, M. Schuld, C. Gogolin, S. Ahmed, V. Ajith, M. S. Alam, G. Alonso-Linaje, B. AkashNarayanan, A. Asadi, J. M. Arrazola, U. Azad, S. Banning, C. Blank, T. R. Bromley, B. A. Cordier, J. Ceroni, A. Delgado, O. D. Matteo, A. Dusko, T. Garg, D. Guala, A. Hayes, R. Hill, A. Ijaz, T. Isacsson, D. Ittah,

S. Jahangiri, P. Jain, E. Jiang, A. Khandelwal, K. Kottmann, R. A. Lang, C. Lee, T. Loke, A. Lowe, K. McKiernan, J. J. Meyer, J. A. Montañez-Barrera, R. Moyard, Z. Niu, L. J. O’Riordan, S. Oud, A. Panigrahi, C.-Y. Park, D. Polatajko, N. Quesada, C. Roberts, N. Sá, I. Schoch, B. Shi, S. Shu, S. Sim, A. Singh, I. Strandberg, J. Soni, A. Száva, S. Thabet, R. A. Vargas-Hernández, T. Vincent, N. Vitucci, M. Weber, D. Wierichs, R. Wiersema, M. Willmann, V. Wong, S. Zhang, and N. Killoran, “Pennylane: Automatic differentiation of hybrid quantum-classical computations,” 2022. [Online]. Available: https://arxiv.org/abs/1811.04968 [25] D. Ittah, A. Asadi, E. O. Lopez, S. Mironov, S. Banning, R. Moyard, M. J. Peng, and J. Izaac, “Catalyst: a python jit compiler for auto-differentiable hybrid quantum programs,” Journal of Open Source Software, vol. 9, no. 99, p. 6720, 2024. [Online]. Available: https://doi.org/10.21105/joss.06720 [26] M. B. Healy, R. Jokar, S. Thomas, V. R. Pascuzzi, K. Barton, T. A. Alexander, R. Elkabetz, B. C. Donovan, H. Horii, and M. Hillenbrand, “Design and architecture of the ibm quantum engine compiler,” in 2024 IEEE International Conference on Quantum Computing and Engineering (QCE), vol. 01, 2024, pp. 866–872. [27] P. Hopf, E. Ochoa, Y. Stade, D. Rovara, N. Quetschlich, I. A. Florea, J. Izaac, R. Wille, and L. Burgholzer, “Integrating quantum software tools with(in) mlir,” in Proceedings of the Supercomputing Asia and International Conference on High Performance Computing in Asia Pacific Region, ser. SCA/HPCAsia ’26. New York, NY, USA: Association for Computing Machinery, 2026, p. 42–54. [Online]. Available: https://doi.org/10.1145/3773656.3773658 [28] J.-S. Kim, A. McCaskey, B. Heim, M. Modani, S. Stanwyck, and T. Costa, “Cuda quantum: The platform for integrated quantum-classical computing,” in Proceedings of the 60th Annual ACM/IEEE Design Automation Conference, ser. DAC ’23. IEEE Press, 2025, p. 1–4. [Online]. Available: https://doi.org/10.1109/DAC56929.2023.10247886 [29] D. Ittah, J. Fraser, J. Izaac, and O. D. Matteo, “Constant-time hybrid compilation of shor’s algorithm with quantum just-in-time compilation,” 2025. [Online]. Available: https://arxiv.org/abs/2504.12449 [30] H. Aghaee Rad, T. Ainsworth, R. N. Alexander, B. Altieri, M. F. Askarani, R. Baby, L. Banchi, B. Q. Baragiola, J. E. Bourassa, R. S. Chadwick, I. Charania, H. Chen, M. J. Collins, P. Contu, N. D’Arcy, G. Dauphinais, R. De Prins, D. Deschenes, I. Di Luch, S. Duque, P. Edke, S. E. Fayer, S. Ferracin, H. Ferretti, J. Gefaell, S. Glancy, C. Gonz’alez-Arciniegas, T. Grainge, Z. Han, J. Hastrup, L. G. Helt, T. Hillmann, J. Hundal, S. Izumi, T. Jaeken, M. Jonas, S. Kocsis, I. Krasnokutska, M. V. Larsen, P. Laskowski, F. Laudenbach, J. Lavoie, M. Li, E. Lomonte, C. E. Lopetegui, B. Luey, A. P. Lund, C. Ma, L. S. Madsen, D. H. Mahler, L. Mantilla Calder’on, M. Menotti, F. M. Miatto, B. Morrison, P. J. Nadkarni, T. Nakamura, L. Neuhaus, Z. Niu, R. Noro, K. Papirov, A. Pesah, D. S. Phillips, W. N. Plick, T. Rogalsky, F. Rortais, J. Sabines-Chesterking, S. Safavi-Bayat, E. Sazhaev, M. Seymour, K. Rezaei Shad, M. Silverman, S. A. Srinivasan, M. Stephan, Q. Y. Tang, J. F. Tasker, Y. S. Teo, R. B. Then, J. E. Tremblay, I. Tzitrin, V. D. Vaidya, M. Vasmer, Z. Vernon, L. F. S. S. M. Villalobos, B. W. Walshe, R. Weil, X. Xin, X. Yan, Y. Yao, M. Zamani Abnili, and Y. Zhang, “Scaling and networking a modular photonic quantum computer,” Nature, vol. 638, no. 8052, pp. 912–919, Feb 2025. [Online]. Available: https://doi.org/10.1038/s41586-024-08406-9 [31] M. Wang, A. Li, and F. Mueller, “Fully parallelized bp decoding for quantum ldpc codes can outperform bp-osd,” in 2026 IEEE International Symposium on High Performance Computer Architecture (HPCA), 2026, pp. 1–14. [32] D. Báscones, A. S. Maan, V. Savin, and F. Garcia-Herrero, “A scalable fpga architecture for real-time decoding of quantum ldpc codes using gari,” 2026. [Online]. Available: https://arxiv.org/abs/2605.01035 [33] F. Battistel, C. Chamberland, K. Johar, R. W. J. Overwater, F. Sebastiano, L. Skoric, Y. Ueno, and M. Usman, “Real-time decoding for fault-tolerant quantum computing: progress, challenges and outlook,” Nano Futures, vol. 7, no. 3, p. 032003, aug 2023. [Online]. Available: https://doi.org/10.1088/2399-1984/aceba6 [34] R. Toshio, K. Kishi, J. Fujisaki, H. Oshima, S. Sato, and K. Fujii, “Decoder switching: Breaking the speed-accuracy tradeoff in real-time quantum error correction,” 2025. [Online]. Available: https://arxiv.org/abs/2510.25222 [35] B. M. Terhal, “Quantum error correction for quantum memories,” Rev. Mod. Phys., vol. 87, pp. 307–346, Apr 2015. [Online]. Available: https://link.aps.org/doi/10.1103/RevModPhys.87.307 [36] A. Javadi-Abhari, M. Treinish, K. Krsulich, C. J. Wood, J. Lishman, J. Gacon, S. Martiel, P. D. Nation, L. S. Bishop, A. W. Cross, B. R.

18

Johnson, and J. M. Gambetta, “Quantum computing with qiskit,” 2024. [Online]. Available: https://arxiv.org/abs/2405.08810 [37] R. Seidel, S. Bock, R. Zander, M. Petrič, N. Steinmann, N. Tcholtchev, and M. Hauswirth, “Qrisp: A framework for compilable high-level programming of gate-based quantum computers,” 2024. [Online]. Available: https://arxiv.org/abs/2406.14792 [38] M. Koch, A. Lawrence, K. Singhal, S. Sivarajah, and R. Duncan, “Guppy: Pythonic quantum-classical programming,” in Informal Proceedings of the Fourth International Workshop on Programming Languages for Quantum Computing (PLanQC ’24), 2024. [Online]. Available: https://github.com/quantinuum/guppylang [39] J. Ansel, E. Yang, H. He, N. Gimelshein, A. Jain, M. Voznesensky, B. Bao, P. Bell, D. Berard, E. Burovski, G. Chauhan, A. Chourdia, W. Constable, A. Desmaison, Z. DeVito, E. Ellison, W. Feng, J. Gong, M. Gschwind, B. Hirsh, S. Huang, K. Kalambarkar, L. Kirsch, M. Lazos, M. Lezcano, Y. Liang, J. Liang, Y. Lu, C. K. Luk, B. Maher, Y. Pan, C. Puhrsch, M. Reso, M. Saroufim, M. Y. Siraichi, H. Suk, S. Zhang, M. Suo, P. Tillet, X. Zhao, E. Wang, K. Zhou, R. Zou, X. Wang, A. Mathews, W. Wen, G. Chanan, P. Wu, and S. Chintala, “Pytorch 2: Faster machine learning through dynamic python bytecode transformation and graph compilation,” in Proceedings of the 29th ACM International Conference on Architectural Support for Programming Languages and Operating Systems, Volume 2, ser. ASPLOS ’24. New York, NY, USA: Association for Computing Machinery, 2024, p. 929–947. [Online]. Available: https://doi.org/10.1145/3620665.3640366 [40] V. Jovanovic, A. Shaikhha, S. Stucki, V. Nikolaev, C. Koch, and M. Odersky, “Yin-yang: concealing the deep embedding of dsls,” in Proceedings of the 2014 International Conference on Generative Programming: Concepts and Experiences, ser. GPCE 2014. New York, NY, USA: Association for Computing Machinery, 2014, p. 73–82. [Online]. Available: https://doi.org/10.1145/2658761.2658771 [41] J. Svenningsson and E. Axelsson, “Combining deep and shallow embedding for edsl,” in Proceedings of the 2012 Conference on Trends in Functional Programming - Volume 7829, ser. TFP 2012. Berlin, Heidelberg: Springer-Verlag, 2012, p. 21–36. [Online]. Available: https://doi.org/10.1007/978-3-642-40447-4_2 [42] D. Moldovan, J. Decker, F. Wang, A. Johnson, B. Lee, Z. Nado, D. Sculley, T. Rompf, and A. B. Wiltschko, “Autograph: Imperative-style coding with graph-based performance,” in SysML, 2019. [43] J. Bradbury, R. Frostig, P. Hawkins, M. J. Johnson, Y. Katariya, C. Leary, D. Maclaurin, G. Necula, A. Paszke, J. VanderPlas, S. Wanderman-Milne, and Q. Zhang, “JAX: composable transformations of Python+NumPy programs,” 2018. [Online]. Available: http://github.com/jax-ml/jax [44] V. Klimis, “A low-latency control fabric for distributed quantum error correction using rdma: A timed event structure model,” in 2025 IEEE Conference on Pervasive and Intelligent Computing (PICom), 2025, pp. 244–249. [45] A. Gangidi, R. Miao, S. Zheng, S. J. Bondu, G. Goes, H. Morsy, R. Puri, M. Riftadi, A. J. Shetty, J. Yang, S. Zhang, M. J. Fernandez, S. Gandham, and H. Zeng, “Rdma over ethernet for distributed training at meta scale,” in Proceedings of the ACM SIGCOMM 2024 Conference, ser. ACM SIGCOMM ’24. New York, NY, USA: Association for Computing Machinery, 2024, p. 57–70. [Online]. Available: https://doi.org/10.1145/3651890.3672233 [46] S. Maurya, A. Molavi, A. Albarghouthi, and S. Tannu, “A case for elastic quantum error correction decoders,” in Proceedings of the 21st European Conference on Computer Systems, ser. EUROSYS ’26. New York, NY, USA: Association for Computing Machinery, 2026, p. 514–531. [Online]. Available: https://doi.org/10.1145/3767295.3803584 [47] L. authors, “ORC Design and Implementation,” Online documentation, 2026. [Online]. Available: https://llvm.org/docs/ORCv2.html [48] P. Shamis, M. G. Venkata, M. G. Lopez, M. B. Baker, O. Hernandez, Y. Itigin, M. Dubman, G. Shainer, R. L. Graham, L. Liss, Y. Shahar, S. Potluri, D. Rossetti, D. Becker, D. Poole, C. Lamb, S. Kumar, C. Stunkel, G. Bosilca, and A. Bouteiller, “Ucx: An open source framework for hpc network apis and beyond,” in 2015 IEEE 23rd Annual Symposium on High-Performance Interconnects, 2015, pp. 40–43. [49] P. Grun, S. Hefty, S. Sur, D. Goodell, R. D. Russell, H. Pritchard, and J. M. Squyres, “A brief introduction to the openfabrics interfaces - a new network api for maximizing high performance application efficiency,” in 2015 IEEE 23rd Annual Symposium on High-Performance Interconnects, 2015, pp. 34–39. [50] B. Woodruff, “Introduction to the InfiniBand core software,” in Proceedings of the Linux Symposium, vol. 2, 2005, pp. 279–290. [Online]. Available: https://www.kernel.org/doc/ols/2005/ ols2005v2-pages-279-290.pdf

[51] G. Kerr, “Dissecting a small InfiniBand application using the verbs API,” 2011. [Online]. Available: https://arxiv.org/abs/1105.1827 [52] Linux RDMA Community, “rdma-core: Userspace RDMA core components and libibverbs,” https://github.com/linux-rdma/rdma-core, 2024. [53] InfiniBand Trade Association, “Infiniband architecture specification,” InfiniBand Trade Association, Specification Release 1.3, 2015. [Online]. Available: https://www.infinibandta.org/ibta-specification/ [54] “AMD Embedded RDMA Enabled NIC LogiCORE IP,” https://docs. amd.com/r/en-US/pg332-ernic, accessed: 2026-06-12. [55] “AMD VPK120 evaluation board user guide,” https://docs.amd.com/r/ en-US/ug1568-vpk120-eval-bd, accessed: 2026-06-12. [56] J. Roffe, D. R. White, S. Burton, and E. T. Campbell, “Decoding across the quantum low-density parity-check code landscape,” Phys. Rev. Res., vol. 2, no. 4, p. 043423, 2020. [57] L. Lao, Q. Wang, Y. Liu, Y. Liu, H. Wang, Y. Chen, Y. Zhao, Z. Wu, W. Zhang, Y. Dong, Y. Liu, M. Lai, and J. Wu, “Real-time decoding of quantum error correction codes using high-performance computing,” 2026. [Online]. Available: https://arxiv.org/abs/2608.03948 [58] UALink Consortium, “UALink: Ultra Accelerator Link,” https:// ualinkconsortium.org, accessed: Aug, 2026. [59] Ultra Ethernet Consortium, “Ultra Ethernet Specification,” https:// ultraethernet.org, accessed: Aug, 2026. [60] Red Hat, “Configuring Soft-RoCE,” https://docs.redhat. com/en/documentation/red_hat_enterprise_linux/9/html/ configuring_infiniband_and_rdma_networks/configuring-roce_ configuring-infiniband-and-rdma-networks#configuring-soft-roce_ configuring-roce, accessed: Aug, 2026. [61] M. J. Heer, B. Ramhorst, Y. Zhu, L. Liu, Z. Hu, J. Dann, and G. Alonso, “Roce balboa: Service-enhanced data center rdma for smartnics,” 2025. [Online]. Available: https://arxiv.org/abs/2507.20412 [62] B. Ramhorst, D. Korolija, M. J. Heer, J. Dann, L. Liu, and G. Alonso, “Coyote v2: Raising the level of abstraction for data center fpgas,” in Proceedings of the ACM SIGOPS 31st Symposium on Operating Systems Principles, ser. SOSP ’25. New York, NY, USA: Association for Computing Machinery, 2025, p. 639–654. [Online]. Available: https://doi.org/10.1145/3731569.3764845

19

A PPENDIX A T EST E NVIRONMENT AND FPGA EXPERIMENTS A. Reference test environment 1) CPU/GPU server: The primary server used for our demonstrations and benchmarks is described in Table IV. The GPU and RNIC occupy the same NUMA domain to avoid an inter-domain hop on the data path. The executor and its memory allocations are bound to that domain with numactl; the CPU polling thread is pinned to a fixed core and requests real-time scheduling. TABLE IV: Reference server used for demonstrations and benchmarking. Component

Reference configuration

Platform CPU

ASRock WRX90 workstation AMD Ryzen Threadripper PRO 9975WX, 32 cores / 64 threads AMD Instinct MI210, CDNA2 / gfx90a, PCIe 4.0 NVIDIA ConnectX-7 / mlx5, PCIe 5.0, RoCE v2 NPS4; RNIC, GPU, process, and buffers on one domain

GPU RNIC NUMA layout

For benchmarking, we use the FPGA as controller, and the CPU or GPU on this server as the coprocessor to measure the round-trip timing. In the GPU case, the GPU executes a persistent kernel to poll on its memory for message arrival, then writes the response message to host memory; a separate CPU thread polls the host memory and then writes the message back to the controller. The CPU and GPU echo responders return the received value without decoding and therefore allow us to measure the transport floor. The cpu-steane and gpu-steane responders use pre-compiled Steane lookup decoder implementations. The gpu-qldpc responder runs a Triton-generated BP-OSD decoder on the MI210 GPU. 2) FPGA controller: The controller for our benchmarks is described in Table V. Both the ARM processing system (PS) and the programmable logic (PL) participate in bringing up a session, but their responsibilities are separated, as summarized in Table VI. TABLE V: Reference FPGA controller used for demonstrations and benchmarking. Component

Reference configuration

Platform Device

AMD VPK120 evaluation board AMD Versal Premium XCVP1202

Processing system RDMA engine HWHS clock

Arm (aarch64), PetaLinux 2024.2 AMD (Xilinx) ERNIC, RoCE v2 200 MHz; one recorded cycle is 5 ns

B. ERNIC and HWHS implementation 1) ERNIC software and hardware split: The Embedded RDMA Enabled NIC (ERNIC) is an AMD soft-core IP that implements FPGA-side transport and packet processing for RDMA over RoCE v2 in programmable logic on an FPGA. The PS is responsible for host and device memory allocation, establishing and programming queue-pair (QP)

TABLE VI: VPK120 controller partition in the current design. Region

Responsibility

PS

PetaLinux, session setup, QP/MR creation, configuration, readout WQE consumption, doorbells, RoCE v2 transmit/receive data path Per-round sequencing, WQE construction, reply detection, RTT timestamping SQ/data BRAM, dual-port reply URAM, RTT trace RAM

ERNIC PL HWHS PL On-chip memory

contexts, and handling error exceptions. Hardware then implements the zero-copy data transport without software intervention. The ERNIC Core manages packet encapsulation and decapsulation, sequence validation, Direct Memory Access, and hardware acknowledgements/re-transmissions. Applicationlayer logic handling timing and dispatch are handled by the Hardware-Handshake (HWHS) module’s orchestrator (hh_orchestrator), which interfaces directly with the ERNIC’s payload buses and doorbell registers. 2) HWHS state machine: The application-layer logic is governed by the single-issue hardware-handshake orchestrator’s state machine in Fig. 8. The state machine optimizes round-trip latency through the pre-posting of work descriptors. The edge labels in Fig. 8 denote the condition or completion event that triggers each transition; ! denotes negation. The two entry paths distinguish a software-initiated round (sw-loop) from a hardware-stimulus round (hw-loop). Along the common path, completion signals advance the FSM after WQE posting, doorbell delivery, and ERNIC transmission. The reply sequence number determines whether the observed data belong to the current round, while the timeout path prevents an unsuccessful transaction from stalling the engine indefinitely. The JOIN state synchronizes reply detection with the background pre-post of the next WQE: prepost_done indicates that the descriptor is ready, while demo_act selects whether the completed correction is returned through DAT_TX. Otherwise, the FSM proceeds directly to NEXT, where the round status is committed. This overlap keeps descriptor preparation outside the steady-state critical path and allows reply handling to remain entirely in PL. The overall design consumes only 16% of the device’s LUTs and 9% of its flip-flops, leaving room to insert additional logic modules on the same FPGA and to leverage its low-latency fabric interconnects. On-chip memory usage is intentionally larger to support cycle-accurate RTT profiling and demonstrations. • Send Queue (SQ) & Queue-Pair Context & Demo Syndromes: The majority of Block RAM blocks (Emb_mem_gen_0/Emb_mem_gen_1 and ERNIC Core) are mapped to hold the hardware SQ descriptor rings, and completion queues. This guarantees immediate access to RDMA descriptors without incurring off-chip access latencies. • Reply Buffers: Dedicated dual-ported UltraRAMs store incoming data, enabling memory polling at full memory bandwidth to minimize poll miss costs. • RTT Trace & Debug Space: To support latency profiling,

20

hw_start

TABLE VII: Overall resource utilization and sub-module breakdown on VPK120.

S_IDLE sw_start

DAT_RX

dma_done

Module

ring_done

BRAM

URAM

900,224

450,112

1,800,448

1,341

677

Total Useda

141,652

23,010

168,559

725

220

16%

5%

9%

54%

32% 156

Component breakdown round_done

WAIT_RPL rcv_completion

POLL_RPL

!seq_match

seq_match or timeout

JOIN

tx_done

FF

Available Utilization (%)

RING_SQ

demo_act and prepost_done

LUTRAM

Overall Device Utilization WQE_POST post_done

DAT_TX

LUT

!demo_act and prepost_done

ERNIC Subsystem

73,077

4,254

52,207

172.5

ERNIC Core

67,748

4,254

44,679

165.5

20

HW Handshake

4,147

0

6,368

7

136

Reply RAM

193

0

512

0

8

RTT Trace

275

0

705

0

128

Emb_mem_gen_0

1,811

0

1

224

0

Emb_mem_gen_1

204

0

0

0

64

Emb_mem_gen_2

2,212

0

2

256

0

a Overall utilization also includes 397 I/O pins (57%), 2 DSP blocks in the !reply_poll or timeout

ERNIC Core, and 1 MMCM (8%).

deployed system. Under hw-loop the PL pacer decides when the next Fig. 8: Consolidated FSM transition diagram for the syndrome is issued, and Table IX gives the two cadences hh_orchestrator control unit. In WQE_POST the orches- reported here. Both use the same Fibonacci-LFSR interval trator constructs a 64-byte WQE and computes the local and generator, interval = freq + (LFSR & span), with tap sequence remote offsets inline, posting it to the SQ slot memory if the [32, 22, 2, 1, 0] and a fixed seed for reproducibility. previous descriptor was not already pre-posted. RING_SQ rings For b2b (back-to-back), freq and span are zero, so the the ERNIC’s SQ producer-index doorbell and starts the RTT next syndrome is issued as soon as the previous correctimer. After waiting on an ERNIC completion in WAIT_RPL, tion returns; this is the continuous-service stress case. For POLL_RPL polls memory until a matching sequence-number random-delayed, the masked LFSR produces a random tag written by the remote endpoint indicates valid data; on quiet period, modelling a nearby real-time decoder that handles match or timeout the round-trip time is recorded. DAT_RX and the common cases and calls out to the remote CPU or GPU DAT_TX carry the payload between PL and the remote endpoint decoder only for difficult tail events. The programmed quiet under hw-loop. The next round’s WQE is posted as soon as period falls between rounds and is excluded from the reported the ERNIC completion arrives, and execution synchronizes at RTT. JOIN to confirm that the background pre-post has finished. We verified the generator before use: replaying 10,000 intervals with the deployed parameters (200 MHz clock, freq = 1,000,000, span = 0x0FFFFFFF, seed 1) gives a mean of 128 URAM blocks in the HWHS module form an 677.847 ms against a uniform expectation of 676.089 ms, and addressable debug trace space (1, 048, 576 entries ×16 the distribution across 20 equal-width bins is flat to within bits = 2 MiB). This allows continuous, cycle-accurate RTT sampling error. measurement at line rate for up to one million consecutive RDMA rounds. • Benchmarking syndromes: Reference syndromes C. FPGA experiment results and expected replies are stored exclusively in The main text reports the primary sw-loop calibration in Emb_mem_gen_2, allowing the HWHS engine to Fig. 6 and Table III. As noted there, round 0 takes significantly access them efficiently for hardware-paced decoding longer than the steady-state mean, because it carries onebenchmarks. time endpoint initialization including the WQE and queue-pair 3) Controller modes and cadences: Two settings are varied settings: 4.640 µs for CPU echo and 9.270 µs for GPU echo, independently in the benchmark: what machinery drives the against steady-state medians of 2.305 µs and 4.500 µs. The sequence of rounds, and how often a round is issued. Table VIII complete ordered traces and distributions, including round 0, lists the two controller modes. hw-handshake moves the for the sw-loop experiments are shown in Fig. 9. Since this handshake into PL so that no software sits on the measured single-round initialization effect appears in every experiment, request/reply path, while hw-loop additionally moves the in the results reported subsequently (including Fig. 10 and round scheduling into PL, so the measurement includes the Table X) we omit this round and show only the N − 1 steadytime to move the payload between engines as it would be in a state rounds. NEXT

21

TABLE VIII: Controller modes used by the current demos and measurements. Mode

Round driver

Posting and reply detection

Purpose

hw-handshake-sw-loop (sw-loop) hw-handshake-hw-loop (hw-loop)

Host loop PL pacer

HWHS in PL HWHS in PL

Demo path and host-driven calibration Autonomous QPU-side and cadence experiments

Complete software-loop traces including the warm-up first round

CPU echo: complete distribution

CPU echo: complete ordered trace

105 103 102 101 100

Samples

3.5 3.0 2.5

2.5

3.0

3.5 RTT ( s) GPU echo: complete distribution

4.0

4.5 mean=4.448 µs median=4.500 µs std=0.159 µs min=4.200 µs max=9.270 µs P99=4.985 µs P99.9=5.125 µs P99.99=5.200 µs P99.999=5.260 µs Sample 0: 9.270 µs

104 103 102 101 100

4.0

0.0

0.4 0.6 Sample number GPU echo: complete ordered trace

0.2

0.8

1.0 1e6

median=4.500 µs

Sample 0 warm-up

9 8 RTT ( s)

Samples

104

median=2.305 µs

Sample 0 warm-up

4.5 RTT ( s)

mean=2.312 µs median=2.305 µs std=0.021 µs min=2.285 µs max=4.640 µs P99=2.420 µs P99.9=2.445 µs P99.99=2.465 µs P99.999=2.475 µs Sample 0: 4.640 µs

7 6 5

4

5

6

7 RTT ( s)

8

9

4

0.0

0.2

0.4 0.6 Sample number

0.8

1.0 1e6

Fig. 9: Complete software-loop CPU/GPU echo traces. The left column shows the full distributions and the right column shows sample order. Round 0 is the one-time warm-up round containing the initial WQE and queue-pair settings. TABLE IX: Hardware-pacer cadences retained in the current evaluation. Cadence

Interval

Interpretation

b2b random-delayed

No added wait 5 ms–1347 ms

Continuous-service stress case Infrequent fallback

1) Hardware-loop experiments: Here we use hw-loop to understand the effect of the message-issuing cadence on roundtrip latency; as mentioned, this mimics the sporadic decoding interval of a real setup. For b2b mode, each run executes Nb2b = 1,000,000 rounds and for random-delayed mode, each run executes Nrandom−delayed = 10,000 rounds.4 Table X reports the steady-state traces. CPU and GPU echo isolate the transport, CPU and GPU Steane add the pre-compiled decoder, and qLDPC uses the generated GPU decoder. In b2b mode, the five responders have median RTTs of 2.470, 2.465, 4.575, 4.570, and 32.635 µs, respectively. CPU 4With a mean interval of 676 ms per random-delayed round, collecting

10,000 samples takes 1.88 hours excluding setup and result transfer; collecting 1,000,000 samples would take 188 hours.

Steane adds almost nothing over CPU echo for this lookup implementation, GPU Steane is 0.005 µs faster than GPU echo, and the qLDPC kernel dominates its own end-to-end RTT, consistent with the results reported in the main text. TABLE X: Autonomous hw-loop steady-state benchmark matrix after removing Round 0. All latency values are in microseconds. These runs drive the loop from PL and therefore include the PL-to-endpoint payload movement, so they are not directly comparable with the host-driven sw-loop figures of Table III in the main text. Responder

Cadence

Mean

Median

P99

P99.9

Max

cpu-echo cpu-echo

b2b random-delayed

2.472 2.494

2.470 2.490

2.580 2.630

2.610 2.740

2.630 2.820

cpu-steane cpu-steane

b2b random-delayed

2.471 2.494

2.465 2.485

2.570 2.630

2.605 2.765

2.635 2.985

gpu-echo gpu-echo

b2b random-delayed

4.558 4.543

4.575 4.485

5.115 5.115

5.255 5.320

5.415 5.505

gpu-steane gpu-steane

b2b random-delayed

4.611 4.599

4.570 4.570

5.160 5.170

5.295 5.340

5.445 5.560

gpu-qldpc gpu-qldpc

b2b random-delayed

32.633 32.561

32.635 32.570

33.290 33.215

33.430 33.465

33.550 33.815

22

2) Cadence effect on latency: Figure 10 shows the distributions for all five responders. The two processor classes behave differently, and the difference is between the body of the distribution and its tail. On the CPU path, the whole distribution moves toward higher latency. Both responders are 20 ns slower in the median under random-delayed, while their tails widen further: P99.9 rises by 130 ns for echo and 160 ns for Steane. The identical median shift for echo and Steane places this cost on the long-gap path rather than on the decode itself. On the GPU path, all three distributions move together slightly toward lower latency. Under random-delayed, the median decreases by 90 ns, 0 ns, and 65 ns for echo, Steane, and qLDPC, respectively. Their extreme tails nevertheless become heavier, with P99.9 increasing by 65 ns, 45 ns, and 35 ns. Thus, cadence shifts the body of the CPU and GPU distributions in different directions while widening the tail on both paths. These changes remain small relative to round trips of several microseconds. Two hypothesized mechanisms are compatible with the measured shift: • CPU C-state exit. The 676 ms mean gap is long enough for an idle core to enter a deep state such as C6, which would be consistent with the effect being confined to the CPU responders. Against this, the responder is pinned and polls continuously, so deep-state residency may be limited, and the observed penalty is smaller than a typical full C6 exit. • P-state or frequency ramp. Even with the polling core active, core or uncore frequency may fall during a quiet period. The time to return to the performance state would explain both the small central increase and the broader CPU tail without a deep C-state transition.

CPU Steane

RTT ( s)

3.0 2.8

2.6

2.6

5.5

2000 4000 6000 8000 10000 Steady-state round

GPU echo

b2b (P50 4.570, P99.9 5.255 µs)

5.5 5.0

4.5

4.5

5.5

2000 4000 6000 8000 10000 Steady-state round

GPU Steane

b2b (P50 4.535, P99.9 5.280 µs)

5.0

5.0

4.5

4.5 0

2000 4000 6000 8000 10000 Steady-state round

GPU qLDPC

b2b (P50 32.650, P99.9 33.455 µs)

33

33

32

32 0

2000 4000 6000 8000 10000 Steady-state round

P(RTT x)

random-delayed (P50 2.485, P99.9 2.765 µs)

0

100 10 1 10 2 10 3

empirical CCDF

b2b random-delayed

random-delayed (P50 4.485, P99.9 5.320 µs)

100 10 1 10 2 10 3

2.6 2.7 RTT threshold ( s)

CPU Steane

empirical CCDF

GPU echo

2.6 2.7 2.8 2.9 RTT threshold ( s)

3.0

empirical CCDF

b2b random-delayed

4.50 4.75 5.00 5.25 5.50 RTT threshold ( s)

random-delayed (P50 4.570, P99.9 5.340 µs)

100 10 1 10 2 10 3

GPU Steane

empirical CCDF

b2b random-delayed

4.50 4.75 5.00 5.25 5.50 RTT threshold ( s)

random-delayed (P50 32.570, P99.9 33.465 µs) GPU qLDPC

2000 4000 6000 8000 10000 Steady-state round

2.8

b2b random-delayed

2.5

2000 4000 6000 8000 10000 Steady-state round

GPU qLDPC

CPU echo

2.5

2000 4000 6000 8000 10000 Steady-state round

GPU Steane

0

100 10 1 10 2 10 3

2000 4000 6000 8000 10000 Steady-state round

GPU echo

0 5.5

random-delayed (P50 2.490, P99.9 2.740 µs)

2000 4000 6000 8000 10000 Steady-state round

CPU Steane

0

5.0 0

RTT ( s)

b2b (P50 2.465, P99.9 2.605 µs)

0

2.8

0

RTT ( s)

2000 4000 6000 8000 10000 Steady-state round

CPU echo

P(RTT x)

RTT ( s)

3.0

2.8 2.7 2.6 2.5

P(RTT x)

0

Hardware-paced RTT: b2b versus random-delayed interval, steady-state traces

b2b (P50 2.470, P99.9 2.610 µs)

P(RTT x)

2.8 2.7 2.6 2.5

CPU echo

P(RTT x)

RTT ( s)

23

100 10 1 10 2 10 3

empirical CCDF

b2b random-delayed

32.0

32.5 33.0 33.5 RTT threshold ( s)

Fig. 10: Steady-state b2b (blue) and random-delayed (red) traces, one row per responder. Only the first Nrandom−delayed − 1 rounds are plotted here. The P50 and P99.9 quoted in each panel include only the rounds plotted, so for b2b values come from the truncated subset and can differ slightly from the full-run values in Table X. The first two columns plot round-trip time against round number and the third column overlays the two empirical CCDFs for that responder.