GPU-resident Execution for Timing-predictable Sensor Processing Systems

GPU-resident Execution for Timing-predictable Sensor Processing Systems Soham Sinha*

Rahul Choudhury*

Shekhar Dwivedi

NVIDIA Bengaluru, India

NVIDIA Santa Clara, CA, USA

NVIDIA Hyderabad, India

Gregory Lee, Gigon Bae, and Anisha Goel

Elena Agostini

Patrick O’Grady, Yosi Levi, and Mahdi Azizian*

arXiv:2610.13809v1 [cs.DC] 9 Oct 2026

NVIDIA Santa Clara, CA, USA

NVIDIA Milan, Italy * {sohams,rahulc,mazizian}@nvidia.com

Abstract—Real-time embedded systems increasingly place GPUaccelerated processing of high-bandwidth sensor streams inside closed-loop systems, but the CPU (the host) remains the primary orchestrator of each iteration. The host observes sensor inputs, submits CUDA work, waits for completion, and commits actuation. This host-mediated path exposes the closed-loop end-to-end latency to I/O interrupts, host scheduler decisions, and CUDA API overheads even after common host-side mitigations such as GPUmemory RDMA or CPU core isolation. This paper presents GPU-resident Execution, a new execution model for real-time sensor processing pipelines that keeps I/O and GPU execution in a long-lived device-resident loop. Sensor inputs arrive in GPU memory, trigger GPU-side processing, and reach GPU-visible output endpoints without any host wakeups. We implement the model in the NVIDIA Holoscan C++ framework and provide reference components for building GPU-resident applications. We evaluate our system on NVIDIA GPUs and SoCs with inference, host-contention, multi-loop, and streaming camera workloads. On an industry-grade 1080p, 60 Hz real-time camera-to-display application, GPU-resident Execution reduces photon-to-display jitter from 4.83 ms to 0.66 ms (7.3×) relative to a CPU-driven GPU-memory RDMA baseline and reduces software-processing jitter from 1.70 ms to 0.01 ms (170×).

NVIDIA Santa Clara, CA, USA

time

(a) CPU-driven CUDA Stream loop

I/O sensor

CPU

read (interrupt/poll)

K1

GPU

write wait on CUDA event (interrupt/poll)

launch kernel

launch kernel

actuator

K2 time

(b) CPU-driven CUDA Graph loop

I/O sensor

CPU

read (interrupt/poll)

launch graph

actuator write wait on CUDA event (interrupt/poll)

K1 → K2

GPU

(c) CPU-free GPU-resident Graph (this work) I/O

sensor

CPU launch once GPU

I. I NTRODUCTION Modern sensor data processing systems increasingly ingest high-bandwidth camera, radar, LiDAR, and IMU streams that demand parallel processing on accelerators [1], [2], [3]. This makes accelerator execution part of the bounded end-to-end response-time required by closed-loop systems. GPUs are now the dominant accelerator, even for such time-critical workloads, in autonomous driving [4], [5], drones [6], [7], robotics [8], medical devices [9], [10], and industrial automation [11], [12]. GPUs are ofered in compact SoC form factors [13], [14], [15], [16] with feature-rich software stacks (CUDA/TensorRT, ROCm/MIOpen), camera drivers, display, and networking support. However, the conventional CPU-driven GPU usage keeps it subordinate to the host, complicating timing predictability. Host-mediated I/O and GPU orchestration remains on the time-critical path. In a CPU-driven workflow, the GPU drives accelerated processing but leaves per-frame control on the host. The CPU observes sensor inputs, submits GPU work, waits for completion, and commits actuation. The top two rows

loop

time

actuator

host idle read (poll)

process write K1 → K2 → · · · (poll)

Fig. 1. Execution models for sensing→accelerated processing→actuation loop.

of Figure 1 capture this host-controlled sequence: whether the host launches kernels one at a time in a CUDA stream or submits them as a captured CUDA Graph [17], per-frame I/O, CUDA operations, and synchronization still route through the CPU. The sensor-to-output-endpoint latency, therefore, includes host orchestration, which runs under Linux and is subject to event-delivery, driver and scheduler timing variation [18], [19], [20], [21], [22]. Prior work has attempted to mitigate this host-side cost in Linux. PREEMPT_RT patches provide threaded interrupt handlers and makes most kernel execution preemptible [23], [24], [19]; kernel-bypass I/O removes bottlenecks from the network stack [25], [26]; and core isolation reduces interference

1

Growth vs. Tegra K1

and scheduler preemption [21]. These techniques improve the CPU cores environment in which the host runs. They do not change the 2560 12x CUDA cores control-plane placement: per-frame I/O and CUDA controls 2048 remain in Linux, and end-to-end timing continues to inherit 8x scheduling delays, kernel overheads, and OS-level uncertainty. This paper presents GPU-resident execution, a new execution 14 model whose design principle is to avoid the unpredictable CPU 4x 512 12 8 components from the closed-loop system entirely at steady-state 6 256 4 192 and place them on the GPU for its lifetime. We developed it 1x with the observation that GPU-memory RDMA [27], [28] and K1 TX2 Xavier Orin Thor GPU-driven I/O [29], [30] can form a sensor/actuator substrate, (2014) (2017) (2018) (2022) (2025) with inputs arriving in GPU-visible memory and outputs leaving from it. A device-side handler observes sensor input, triggers Fig. 2. NVIDIA SoC CPU and CUDA core count growth [13]. the data processing pipeline, and writes output signals, all without a host wakeup (see Figure 1). The CPU remains in the setup and supervisory plane, so CPU orchestration overheads (C1) Unpredictable Host I/O Scheduling: CPU-driven and unpredictability do not contribute to the end-to-end sensor- execution manages sensor I/O events through the Linux I/O to-output latency. stack. First, an interrupt or polled queue reaches a driver. Then, This paper is the first to address sensor I/O determinism the kernel notifies a user thread which is then scheduled before and GPU workload execution predictability as one problem, GPU work is launched, and actuator commands are committed. through a GPU-resident execution model. We implement a This path exposes the time-critical loop to interrupt latency, reference C++ framework and examples for real-time streaming deferred kernel work (threaded IRQs, softirqs), polling delay, applications. The framework treats a sensor pipeline as a and priority mismatch between the I/O event and the real-time GPU-native fixed-execution unit, amenable to GPU timing userspace threads [23], [24], [19], [20], [22]. analysis and verification [31], [32], [33], [34], [35], rather than RTOSs can provide stronger I/O timing guarantees through a sequence of unbounded CUDA API calls. priority-aware I/O scheduling, budgeted interrupt handling, We implement GPU-resident Execution using CUDA Graphs, and domain-isolated OS designs [41], [22], [42], [43], [44], GPU-driven I/O and GPU-memory RDMA, and evaluate it [45], but supporting the same I/O devices requires rewritten on x86, Orin, and Thor platforms with Ampere, Ada, and drivers. In addition, most RTOSs do not support GPU software Blackwell GPUs. Across inference, host-contention and scala- stacks (CUDA/TensorRT, ROCm), camera, advanced display, bility microbenchmarks, GPU-resident Execution consistently and advanced networking drivers. QNX is the only commercial reduces standard deviation, jitter and worst-case latency, against RTOS with CUDA support [46], [47], and even that is limited a core-isolated CPU-driven baseline using NVIDIA Holoscan to a safety-certified CUDA subset. Linux-based mitigations – SDK [36]. On an industry-grade 1920×1080, 60 Hz IMX274 threaded interrupt handlers in PREEMPT_RT [24], [20], kernelcamera-to-display pipeline, our framework reduces photon-to- bypass I/O [26], [25], isolated CPU cores for I/O threads – display jitter by 7.3× over a CPU-driven solution and software- reduce parts of this cost, but the critical event path still passes processing jitter by 170×. through the Linux kernel and scheduler. The rest of the paper is organized as follows. Section II gives (C2) Host Scheduling and Core Isolation Limitations: On required background for the design. Section III presents the Linux, time-critical I/O and ML applications are managed by system design, API and implementation details of GPU-resident middleware runtimes such as ROS2 and Holoscan, and share the Execution. Section IV describes the reference operators and host with non-critical, third-party, and background workloads. IMX274 application, built on our API, which are evaluated Achieving predictable interrupt handling and task release in in §V. Finally, Sections VI to VIII place the work in scope, this environment requires either an analytically complex or discuss related work, and conclude with future directions. intractable scheduling configuration, or CPU core isolation that dedicates cores to specific handlers and threads [21]. Core II. BACKGROUND ML-driven sensor pipelines on Linux-hosted GPU plat- isolation may improve predictability in some setups, but it forms [36], [37], [38], [39], [40] target predictable and strict does not scale to modern robotics, humanoid, and industrial timing guarantees such as sub-millisecond end-to-end latency. workloads that integrate tens to hundreds of sensors and This section discusses the challenges that prevent the CPU- actuators. Embedded SoCs provide only a few CPU cores, as driven execution model from meeting those timing contracts shown in Figure 2 for NVIDIA platforms, whereas integrated reliably and then introduces the sensor I/O and CUDA Graph GPU CUDA core counts already outnumber them by over an order of magnitude (PCIe GPUs even more). This motivates primitives used by GPU-resident Execution. using CUDA cores for real-time compute. A. Challenges (C3) CPU–GPU Coordination Overheads: In CPU-driven CPU-driven real-time ML pipelines struggle primarily with GPU work for ML, each iteration depends on host-mediated three systems design challenges. coordination with the GPU: the CPU enters the CUDA runtime/-

2

driver, submits work, and later observes completion through a stream/event synchronization before the next control decision can execute (see Fig. 1). With CUDA Streams, the host submits the kernel one launch at a time. CUDA Graphs [17] reduces the number of launch calls, but the launch and completion path remains part of the real-time loop. Consequently, CUDA driver queueing and dispatch, the CPU scheduler, and the kernel notification mechanisms – all contribute to timing variation [48], [49], [50] and degrade end-to-end predictability. Together, these challenges motivate GPU-resident Execution, which keeps the sensor data path and per-frame data and control flow on the GPU and removes the host from the critical path. Next, we describe the substrates that make this possible.

operations as a reusable graph. Launching the graph replaces many host-side CUDA calls with a single call, but in CPUdriven runtimes that call is still issued once per sensor frame. Indeed, most embedded SDKs using GPUs, including Isaac ROS [39] and Holoscan SDK [60], use either CPU-driven CUDA Streams or repeated host-launched CUDA Graphs. Two CUDA primitives can eliminate this per-frame host relaunch. (1) A CUDA Graph conditional WHILE node [61] reexecutes a body subgraph without leaving the device. (2) Device graph tail-self-launch repeatedly relaunches the same graph from the device [62]. Both mechanisms enable a frameprocessing loop to run entirely on the device. Unlike any prior work, GPU-resident Execution integrates these primitives with GPU-visible I/O into an execution model for real-time sensor processing pipelines.

B. Sensor I/O Removing the host from the real-time I/O path requires that sensor data reach the GPU and actuator commands leave it, directly. In standard Linux, both directions traverse the kernel or a user thread. Kernel-bypass frameworks such as DPDK [26] and U-Net [25] map NIC queues directly into user space, polling receive queues from a dedicated CPU core and reducing per-packet latency to single-digit microseconds on 100 Gb/s connections. GPU-memory RDMA [27], [28] extends bypass to the data path by letting a remote NIC write directly into device memory over PCIe at microsecond-scale latency [51], [52]. Both mechanisms remove kernel mediation and CPU copies but leave control on the CPU. 1) Holoscan Sensor Bridge: Holoscan Sensor Bridge (HSB) [53], [54] is an open-source external FPGA bridge that delivers sensor data to the host. It packetizes sensor output as RDMA over Converged Ethernet (RoCEv2) traffic and provides a unified RDMA/RoCE interface for native sensor links such as MIPI, SPI, and I2C. ConnectX SmartNICs place frame payloads directly into pre-allocated device memory via GPU-memory RDMA, bypassing the host network stack. HSB, therefore, enables a GPU data-plane for sensor and actuator data, but the CPU still controls I/O events and frame reception. We preserve HSB’s RDMA data plane and move those decisions into a GPU-resident data-ready handler (§IV-B) via GPU-driven I/O, placing both the data path and the steady-state control path on the device. 2) GPU-driven I/O: GPU-driven I/O frameworks [29], [30], [55], [56] go one step further by exposing NIC doorbell registers to device kernels which can poll for receive completions, post send requests, and ring doorbells without CPU involvement [57], [58], [59]. Our design combines GPUmemory RDMA with GPU-driven I/O so the sensor and actuator path no longer needs a host-side I/O thread.

III. S YSTEM D ESIGN AND I MPLEMENTATION GPU-resident Execution expresses a sensor processing application as a directed acyclic graph (DAG) of GPU-resident operators. Each operator, representing a basic functional component of the application, reads device-visible inputs, runs one or more CUDA kernels, and writes device-visible outputs. Data-flow edges connect the output ports of one operator to the input ports of another. GPU-resident Execution encapsulates this operator DAG into a single GPU-resident CUDA Graph that the host launches once. Next, we describe the GPUresident Execution framework and GPU-resident Graph before discussing the API and implementation details. A. Framework Architecture and Lifecycle

Sensor + HSB

NIC / RoCE

WHILE

Operators

Output

CPU

launch (once) VRAM write data_ready commit

idle

time

frame

loop

Fig. 3. Lifecycle overview (data-flow: solid, control-flow: dashed).

A GPU-resident application has an initialization phase and a steady-state phase. The host drives the initialization phase, and the device runs the steady-state phase. Figure 3 shows the lifecycle overview of a GPU-resident application. During initialization, the host configures the sensor and output C. CUDA Graphs endpoints, allocates device buffers, constructs the GPU-resident GPU-visible I/O places sensor data and actuator endpoints CUDA Graph from the operator DAG, and launches it once. on the device, but uninterrupted processing for new sensor data The steady-state path is then executed entirely on the GPU. must also execute without host mediation. Advanced CUDA Each iteration checks for a sensor frame, runs the operator Graphs features provide those GPU-native control mechanisms. DAG when input data is ready, and commits the result to the CUDA Graphs [17] allow an application to capture a output endpoint. The host remains outside the time-critical sequence of kernel launches, memory copies and other CUDA steady-state path.

3

Algorithm 1 High-level GPU-resident application lifecycle

subgraph of CUDA kernels (see §III-C3). Next, the controller kernel determines whether to execute the operator DAG. It sets the predicate of the subsequent IF conditional node from the handler’s data-ready flag. This IF node gates the operator DAG. When its predicate is true, the DAG executes before the end-marker marks the completion of an iteration. CUDA Graph dependencies enforce this overall execution order without an additional framework scheduler and timing overheads. During initialization, GPU-resident Execution flattens the operator DAG in topological order. Serialization satisfies all data-flow dependencies and avoids parallel kernel launches inside the captured graph, which would introduce timing Algorithm 1 makes the phase boundary explicit. The host variability from the on-device work scheduler. GPU-resident performs the first two steps once, while the GPU repeats the Execution captures each operator’s CUDA work in this order on sensor-to-output loop. To implement this loop, GPU-resident a dedicated stream to form the true branch of the IF conditional Execution places a data-ready handler and a controller before node, as depicted in Figure 4. the application operator DAG. The handler checks for new input, 1) GPU-resident Controller: The controller kernel is the the controller decides whether to execute the operator DAG, device-side decision point in the GPU-resident CUDA Graph. and an end-marker completes the iteration. The next subsection This single-thread controller translates control flags into the explains how these components form the GPU-resident Graph. WHILE and IF predicates. It runs after the data-ready handler in every WHILE iteration. Algorithm 2 shows pseudocode for B. GPU-resident Graph The application creates a GPU-resident CUDA Graph from the controller kernel. its operator DAG. We provide two mechanisms for keeping this Algorithm 2 Controller kernel in WHL CUDA Graph graph resident on the device: CUDA Graph WHILE conditional 1: d ← dataReady; t ← tearDown ▷ atomic loads node (WHL) [61] and device graph tail launch (DVG) [62]. 2: ifPred ← (d ̸= 0) Their performance is similar (see §V-B4), but they differ 3: if d = 0 then in graph-node support and device-resident execution control 4: sleep for δ µs 5: if t = 1 then mechanics. Figure 4 shows the internal structure of our default 6: whilePred ← 0; ifPred ← 0 WHL mechanism. 1: host: configure endpoints, device buffers, and operator DAG 2: host: construct and launch the GPU-resident CUDA Graph once 3: device: execute the steady-state loop below 4: while application is active do 5: check for device-visible sensor input 6: if input is not ready then 7: sleep for a configurable interval; continue 8: end if 9: invoke the operator DAG 10: commit the result to the output endpoint 11: end while

false

tearDown is a host-mapped control flag through which the host requests termination of the GPU-resident graph. When execution flow set, the controller clears both predicates, skipping the operator true device flag control DAG and terminating the graph. Because the flag is polled data-ready framework node/kernel once per iteration, termination takes effect within one loop tear-down: 0 handler user-supplied node/kernel boundary in the worst case. condition flag The dataReady flag carries input availability through a controller kernel steady-state iteration. The data-ready handler sets it to 1 Operator DAG Op A when input is available, allowing the controller to release the operator DAG through the IF node. If it remains 0, the true input-ready check IF Op C cond. (0/1) node controller clears the IF predicate and sleeps for a configurable interval δ µs. After the IF node’s true branch completes, the Op B next false end-marker resets dataReady to 0 and sets the resultReady iteration flag. resultReady is an optional host-mapped completion end-marker flag for debugging or readback and does not participate in kernel steady-state control. For debugging, the host may map the dataReady flag and control it. Fig. 4. Internal structure of the GPU-resident CUDA Graph. 2) Device Graph Tail Launch: GPU-resident The root of the GPU-resident CUDA Graph is a WHILE Execution also supports the device-launchable CUDA conditional node. It keeps the graph active in the device until Graphs [62] mechanism to keep the graph resident its condition becomes false, at which point the graph terminates. on the device. A device graph is instantiated with As long as the WHILE condition is true, the application remains cudaGraphInstantiateFlagDeviceLaunch and uploaded active in device-resident steady-state mode, and the graph once to the device. A running kernel can launch this graph via executes the WHILE body. The data-ready handler is the first cudaStreamGraphTailLaunch, without host involvement. node in the body and checks for new input via a user-supplied GPU-resident Execution builds two such graphs. Graph M WHILE

node

check root cond. (0/1)

terminate

4

(control) holds the data-ready handler followed by a singlethread dispatcher kernel that replaces the WHL controller (Algorithm 2). Graph W (workload) sequentially stitches the captured operator DAG, the end-marker kernel, and a single-thread tail-relaunch kernel. The host fires Graph M once via cudaGraphLaunch; every subsequent iteration boundary is a device-side tail launch, iterating as in the WHILE design. The dispatcher operates on the same flags as the WHL controller: tearDown stops the execution, dataReady = 1 tail-launches Graph W, and dataReady = 0 sleeps for a configurable interval then tail-launches Graph M. At the end of Graph W, a tail-relaunch kernel unconditionally tail-launches Graph M, closing the loop. All flag contracts and the data-ready handler are reused unchanged. Because device-launchable graphs may not contain conditional nodes, this mechanism cannot support conditional subgraphs that CUDA-based libraries increasingly adopt [63]. Therefore, it is not default in GPU-resident Execution, but it provides a path to low-overhead and flexible on-device dynamic scheduling in the future. Both mechanisms expose the same application-facing API described next.

declare named device-side ports and their storage requirements. Each port uses one of two device-memory modes. With internally-managed memory, the executor allocates one shared buffer per edge, maps it to both endpoints, and manages its lifetime. With externally-managed memory, an endpoint supplies a preallocated device pointer that is mapped to both endpoints. If both endpoints supply pointers, the source pointer is used. During graph construction, operators are visited in topological order. It calls setup() to declare ports and bind parameters, then resolves each edge to a device buffer. It calls compute() once in the same order. This method obtains port pointers through device_memory("name") and issues CUDA operations on cuda_stream(). These CUDA operations are recorded into a CUDA Graph as a subgraph which is inserted in the IF node’s true branch (see Fig. 4). 3) Data-ready Handler Interface: A data-ready handler performs device-side input detection at the start of every GPU-resident iteration. Its GPUResidentOperator subgraph may read sensor metadata, poll a NIC completion queue, or otherwise inspect device-visible input state. When input is available, a kernel signals it via the dataReady flag. During graph construction, handler operators issue their CUDA operations on data_ready_handler_cuda_stream(). The handler subgraph is recorded and inserted as the first node in the WHILE body (see Fig. 4). Reference handlers are described in §IV-B.

C. Application, Operator, and Handler API The GPU-resident graph above is constructed from three user-facing abstractions. GPU-resident Execution exposes them through a C++/CUDA API integrated into the NVIDIA Holoscan SDK [36], [64]. An Application defines the operator DAG and controls its lifecycle. An operator is Holoscan’s basic functional unit. GPU-resident Execution represents it as a GPUResidentOperator, which implements a processing or I/O stage. A GPUResidentHandler defines an operator subgraph that detects new device-visible input. The following interfaces describe these roles before Listing 1 combines them in an example. 1) Application Interface: The application’s compose() method instantiates operators with make_operator() and connects them with add_flow(). A handler’s compose() method uses the same API to form a DAG of input-detection operators. The application’s gpu_resident() accessor exposes register_data_ready_handler(), which attaches this handler to the application. Composition and handler registration occur during host-side initialization. After composition, run_async() constructs and launches the GPU-resident Graph once, beginning device-resident steadystate streaming. While it runs, the host can query its state with is_launched() and request shutdown with tear_down(). 2) Operator Interface: A GPUResidentOperator performs only CUDA operations in GPU-resident Execution. These CUDA operations are captured from the host once, and are repeated as part of the GPU-resident Graph, as explained in §III-B. Following the Holoscan convention, setup(OperatorSpec&) declares an operator’s configuration. GPU-resident Execution extends OperatorSpec with device_input("name", n_bytes) and device_output("name", n_bytes) methods to

// processing operator class MyKernelOp : public GPUResidentOperator { void setup(OperatorSpec& s) override { s.device_input ("in", N*sizeof(float)); s.device_output("out", N*sizeof(float)); } void compute() override { auto *in = (float*)device_memory("in"); auto *out = (float*)device_memory("out"); my_kernel<<<G,B,0,cuda_stream()>>>(in,out); } }; class GpuResidentApp : public Application { void compose() override { auto sense = make_operator<SensorOp>("src"); auto k1 = make_operator<MyKernelOp>("k1"); auto output = make_operator<OutputOp>("sink"); add_flow(sense, k1); add_flow(k1, output); } }; // data-ready handler class MyDRH : public GPUResidentHandler { void compose() override { add_operator(make_operator<HsbPollOp>("drh")); } }; int main() { auto app = make_application<GpuResidentApp>(); auto drh = make_handler<MyDRH>(); app->gpu_resident().register_data_ready_handler(drh); auto fut = app->run_async(); // GPU-resident Graph launched, CPU free for other work. app->gpu_resident().tear_down(); fut.wait(); }

Listing 1. GPU-resident application example

D. Static SM Affinity for GPU-resident Graphs Each GPU-resident application runs inside a CUDA context assigned a full GPU or a static CUDA MPS partition [65]. The

5

partition fixes the streaming multiprocessor (SM) budget availIV. R EFERENCE O PERATORS , H ANDLERS AND able to the graph for its lifetime, and all data-ready, controller, A PPLICATIONS operator, and egress kernels execute in that context on each This section presents reference implementations of the frameiteration. Predictable GPU-resident execution also requires work’s ingress, processing, and egress components and then locked GPU clocks during deployment (e.g., nvidia-smi combines them in a real-time camera-to-display application. -lock-gpu-clocks [66]), so the steady-state loop runs on a fixed resource at a stable frequency. This forms a key practical A. GPU-resident Operator Implementations basis for timing predictability. We do not currently consider These reference operators show that the API in §III-C shared L2 cache, memory bandwidth, and other device-side supports custom CUDA kernels and library-backed imageinterference, whose control would require stronger resource processing and inference stages. They configure image dimenpartitioning [67], [68], akin to Intel RDT [69]. sions, tensor layouts, auxiliary device buffers, library contexts, and kernel launch parameters before their CUDA operations are captured into a CUDA Graph. E. Timing Predictability by Design Bayer Demosaic. Bayer is the raw color-filter-array layout The timing consequence of the design is straightforward: produced by the camera, and demosaicing reconstructs a fullonce the graph has been launched, every steady-state decision color image from it. This operator wraps NVIDIA Performance that releases work for a sensor sample is made inside the Primitives (NPP) demosaicing functions [70]. It takes image device-resident graph. We use CUDA Graph dependencies to size, pixel type, Bayer layout, and alpha as setup parameters, establish a single device-side chain from ingress to egress. and sets the final in and out port sizes during initialization. When the data-ready handler marks input ready, the controller Graph construction then allocates the corresponding deviceupdates the IF predicate before that IF node is reached in the side input and output buffers. For a 1920×1080 raw image same loop iteration. The processing body is therefore released with 16-bit Bayer input and four RGBA output channels, the without a host wakeup, a host-side completion event, or a ports are 4.1 MB and 16.6 MB. During capture, the operator binds the NPP stream context to cuda_stream() and calls per-frame cudaGraphLaunch call by the host. nppiCFAToRGB_* or nppiCFAToRGBA_* kernels. The NPP This gives a design-level timing contract rather than a new demosaic operation is captured as a CUDA Graph node, so GPU scheduler. The deployed loop must have a fixed graph the host does not resubmit it for each frame. topology, preallocated device buffers, locked clocks, and fixed Format Conversion. The format conversion operator supcompute resources assigned to the CUDA context. Under that ports fixed device-only packed and planar layouts. It fixes deployment model, the steady-state response path has only two input and output pixel formats, image dimensions, optional kinds of costs: device-side execution inside the graph and the resize parameters, plane strides, and plane offsets before launch, fixed graph transitions between stages. and allocates any resize or channel-reorder scratch buffers If a sample arrives just after the handler has checked for during initialization. During capture it calls the corresponding input, the critical-instant wait is the time until the next handler NPP conversion on cuda_stream(), so color conversion and check: one controller pass plus the larger of the configurable optional resize remain inside the graph body. data-not-ready sleep interval and the operator kernels’ worstTensorRT Inference. The inference operator places model case execution time. After the next handler observes the sample, inference in the same captured body as hand-written kernels. It the body release and commit remain inside that same graph derives input and output port sizes from the model configuration, iteration. constructs a TensorRT-backed inference context, and connects This timing model makes overload behavior explicit. If a the context’s input and output buffers to executor-owned devicesensor or actuator runs faster than the body can sustainably port addresses. Once configured, TensorRT exposes the model process, GPU-resident Execution does not hide the mismatch as a reusable CUDA Graph, which the operator’s compute() behind an implicit queueing policy. The application defines launches on cuda_stream() during capture. the over-rate behavior at the data-ready handler and egress Camera Image Operators. Together, these operators reboundary: it may process the newest available sample and drop construct and correct a raw camera frame before demosaicing. older ones, suppress duplicate processing when the loop polls 1) Camera-frame reconstruction. This operator reassembles faster than the sensor produces data, or repeat the latest output a MIPI Camera Serial Interface 2 (CSI-2) packet stream when an actuator commits faster than new processed results into a dense 16-bit Bayer frame for RAW8, RAW10, and are available. The timing model then applies to the samples RAW12 sensor formats. The input port is sized to the CSI that this application policy elects to process. byte length and the output port to one 16-bit Bayer value per Thus the main timing claim is not that GPU work is pixel. 2) Raw-image signal processing. This operator applies inherently deterministic, but that the per-frame host scheduler, sensor-level corrections in-place, by mapping both in and interrupt, and CUDA-launch terms are absent from the steady- out to a single internal device buffer. Initialization derives state response path; the remaining terms are device-side graph Bayer-pattern channel offsets, computes the optical-black shift and operator costs measured in the evaluation. from the pixel format, and allocates a three-channel 256-bin

6

histogram with a white-balance gain vector that stores the perchannel scale factors. During capture, the operator runs four kernels on the same stream: black-level correction, histogram accumulation, white-balance gain computation, and the final per-pixel correction pass.

IMX274 camera

HSB FPGA

ConnectX-7

GPU CQ + receive ring

configure CPU launch once

B. Holoscan Sensor Bridge (HSB) Data-ready Handlers

GPU-resident CUDA Graph DOCA data-ready

image pipeline

display operator display buffer

monitor

screen HSB is an external sensor-side FPGA bridge that packetizes native sensor data into RoCE traffic for a ConnectX NIC [53]. A data-ready handler is required to integrate HSB sensor Fig. 5. GPU-resident IMX274 camera-to-display application. stream into a GPU-resident application. This handler detects new frames on the GPU and signals input availability to the controller kernel. Implementing such a handler requires the to the NIC. Thus, the NIC driver remains on the host, but the underlying I/O interface to expose data buffers and completion steady-state control is moved to the GPU. GPU-resident Execution captures the DOCA GPUNetIO state and/or doorbell registers to the GPU. We implement two HSB-facing handlers using GPUDirect RDMA and GPU- kernel launch in the handler’s compute() method on the driven networking with DOCA GPUNetIO [29]. Other sensor data-ready handler CUDA stream. In the streaming phase, the interfaces can use the same abstraction by exposing this contract kernel performs a non-blocking poll of the receive CQ using the DOCA GPUNetIO API [29]. On a miss, it writes data directly or through a bridge like HSB. HSB metadata handler. This first handler detects frames not ready. On a hit, it reads the immediate-data field of the from GPU-visible metadata while a CPU thread continues to Completion Queue Entry (CQE) to identify the completed manage the NIC receive completion queue. HSB connects to receive-buffer slot, then copies that frame to a device buffer a MIPI CSI-2 camera, such as the IMX274 [71]. It receives for processing. It then advances the CQ consumer position, the camera stream and sends each frame as chunked RoCEv2 posts the next receive WQE, and signals data ready. RDMA Write requests. The ConnectX NIC places the payload C. Display Output into a two-slot (or larger, sensor rate-dependent) GPU memory An output operator closes the end-to-end loop by committing buffer via DMA. After each frame, the FPGA sends a metadata the final device-resident result to an externally visible action. packet (frame counter, status, timestamps) as an RDMA Write Our display operator demonstrates this egress pattern via with Immediate [72]. This generates a completion event [73] on rendering an image stream to a monitor, analogous to a the ConnectX NIC [74]. In RDMA terminology, a Completion network transmit or GPIO write. During initialization, it fixes Queue (CQ) reports finished work requests, while a Queue the display mode (resolution, frame rate, etc.), creates the Pair (QP) is the send/receive endpoint that posts them. In the display swapchain, and obtains device pointers to the scanout canonical CPU-driven path, a CPU thread polls the receive buffers using the Linux DRM API [75]. It does not create a CQ, copies metadata to host memory, selects the frame on the Vulkan, OpenGL, or other graphics context, avoiding a costly CPU, and releases downstream processing. graphics/compute context switch [76] on every frame. During The GPU-resident handler, instead, uses a CUDA kernel that steady-state, the captured graph issues a gamma-corrected polls two metadata records directly on the GPU and selects display kernel. Then, the rendering to a monitor happens in the slot with the newer frame counter. On a new frame, the one of the two modes. Front-buffer rendering mode writes handler sets the dataReady flag and triggers the rest of the one buffer while the display engine continuously scans it out. pipeline. This removes the device-to-host metadata copy and Continuous-flip mode alternates two buffers, scanning out one CPU frame-selection decision from steady state. A lightweight while rendering into the other. CPU thread still manages the receive CQ, but frame detection Other actuators may follow the same structure when they and application execution proceed on the GPU. The next expose a GPU-visible doorbell. Initialization fixes the queue, handler removes the remaining host role. register, or buffer mapping. The captured compute() body HSB DOCA GPUNetIO handler. This second handler writes the payload and rings the doorbell after processing removes the remaining per-frame CPU I/O role. The GPU polls completes. Supporting an egress protocol requires vendorreceive completions and reposts receive descriptors through the specific integration and is part of ongoing work. NVIDIA DOCA GPUNetIO API [29]. During initialization, the host uses DOCA GPUNetIO to allocate the CQ rings, QP D. IMX274 Camera-to-Display Application work queues, doorbell record, and receive-frame ring in GPU This application combines the ingress handlers, processing memory. The ConnectX driver exposes controls that bind these operators, and display egress into a complete externally GPU memory allocations to the NIC. The host then exports observable sensor-to-output loop. Figure 5 shows its hardware the QP as a device-visible handle and connects it to the HSB and GPU-resident graph structure. The IMX274 camera [71] FPGA peer. A setup CUDA kernel pre-posts receive Work is the sensor endpoint in the pipeline, and a monitor is the Queue Elements (WQEs) identifying the GPU buffers available output endpoint. The sensor supports 1080p and 4K resolutions

7

at 60 Hz in 10- or 12-bit Bayer RAW format. A HSB FPGA packetizes each frame as RoCEv2 traffic. A ConnectX-7 (CX7) NIC on the IGX Orin + RTX 6000 Ada platform delivers the payload directly into GPU memory. The existing CPU-driven application uses GPUDirect RDMA for host-bypassed data movement, but the host still handles HSB-CX7 I/O events, launches CUDA work, and orchestrates each frame. We build two GPU-resident configurations using the HSB metadata and DOCA GPUNetIO handlers described above. For each camera frame, the graph reconstructs and corrects the raw image, demosaics and converts it, then commits it to the display, using the previously described operators [77].

periodic ticks), and rcu_nocbs (offload RCU callbacks), with IRQs offloaded to other cores. We use box-and-whisker plots to show latency distributions. Each box spans the middle 50% of observations (25th–75th percentiles), the center line marks the median, the diamond marks the mean, and the whiskers extend to the observed minimum and maximum. 1) Single-stage Inference Pipeline: The first experiment runs a single TensorRT inference operator as the processing workload in a closed loop, resembling a small ML perception pipeline. Five TensorRT image-classification CNNs (ShuffleNetV2-x0.5 [80], MobileNetV2 [81], EfficientNetB0 [82], ResNet-50 [83], and VGG-16 [84]) are tested. The two V. E VALUATION execution models launch the same inference CUDA Graph: the We evaluate GPU-resident Execution in two steps. First, a set Isolated CPU-driven CUDA Graph configuration (CPU-CG of microbenchmarks isolates where predictability is gained: one hereafter) captures a TensorRT engine once and host-launches it inference pipeline, the same loop under host contention, many on every iteration per new input, while the GPU-resident execoncurrent loops, and the fixed framework cost of keeping the cution runs the same model with GPUResidentInferenceOp, loop resident on the GPU. Second, an industry-grade IMX274 with a data-ready handler to iterate over new input on the device. camera-to-display case study tests whether those gains carry Both configurations preload 100 ImageNet images in device through an end-to-end pipeline. memory and round-robin over them during measurement to A. Setup simulate a camera stream delivered over GPU-memory RDMA. We evaluated GPU-resident Execution on several platforms For every measurement, we record 50,000 iterations after a and report representative configurations. All microbenchmarks 5,000 iteration warmup. Figure 6 shows the per-iteration latency distributions. The in §V-B run on a 24-core/48-hardware-thread AMD 7965WX experiment is intentionally quiet, and the two modes have x86_64 workstation (Linux kernel 6.17 PREEMPT_RT) with comparable tail latencies. The benefit of GPU-resident appears an RTX A6000 Ampere (84 SMs, 48 GB VRAM) GPU. We mainly in lower standard deviation and a shorter absolute jitter repeat an experiment on IGX Thor with an integrated GPU (minimum-to-maximum span). For MobileNetV2, EfficientNet(iGPU: 20 SMs, 128 GB RAM, 14-core Armv9-A CPU, Linux B0, and ResNet-50, GPU-resident Execution reduces standard 6.8 PREEMPT_RT). The real-time IMX274 streaming camera deviation and absolute finishing jitter by 1.4×–2.4×. Launchapplication (explained in §IV-D) case study in §V-C runs ing these models as CUDA Graphs per iteration adds variation on IGX Orin with a 12-core Armv8 CPU (Linux 5.15 with to the CPU-CG timing, even though the launcher is running PREEMPT_RT), plus an RTX 6000 Ada (142 SMs, 48 GB on an isolated core. The gains are smaller for ShuffleNetV2 VRAM) GPU. and VGG-16, where model-level behavior dominates the a) Metrics: We focus on latency tails and variation as our timing predictability metrics. A key metric is absolute latency cost [85], and the variance is similar between the finishing jitter [78], [79]: for job k with release rk and finish fk , two modes. GPU-resident Execution remains subject to deviceresponse time Rk = fk −rk and jitter J = maxk Rk −mink Rk . side microarchitectural effects such as runtime cache residency Microbenchmarks treat each loop iteration as one job with zero and memory-system contention, but CPU-driven baselines add release jitter (rk at iteration start, fk at completion), so J host-side overheads on top. The next experiments explore more of these host-side effects. captures end-to-end variation. 2) Host-side Contention: In a realistic deployment, the B. Microbenchmarks CPU not only runs the real-time loop but also telemetry The microbenchmarks demonstrate that GPU-resident Execu- logging, health monitoring, non-critical services, and third-party tion improves timing predictability over CPU-driven baselines processes. We emulate this co-running activity on non-isolated against different sources of timing variation. From a quiet cores with buffered file I/O (20× disk I/O) and memorysystem to host contention and multi-loop scaling, CPU-driven intensive workloads (stress-ng [86] malloc, vm). The CPUexecution and its mitigations remain sensitive to host-side CG launcher remains pinned to a fully isolated core, so the effects, while GPU-resident stays consistently robust. At the experiment tests interference through shared OS kernel, driver, and memory subsystem paths, rather than from CPU scheduling. end, we measure the fixed framework overhead. For CPU-driven baselines, we use Holoscan SDK [36] with Inference. We test against the same five inference models CUDA Graphs rather than individual CUDA stream launches from the previous experiment. We also use the same iteration because Graphs are already the optimized, lower-overhead host- regime. Figure 7 shows the latency distributions for the two driven option. To avoid any host-side scheduling noise, the configurations in this experiment. The boxes capture each host process is pinned to an isolated CPU core configured with model’s typical latency range, which remains similar between isolcpus (exclude scheduler tasks), nohz_full (suppress the two configurations. However, the CPU-CG upper whiskers

8

360

430

340

420

835

410

300

400

ResNet-50

VGG-16 1945

1400

1925 815

320

Isolated CPU-CG

GPU-resident

EffNet-B0

1380

1360

y-axis: log scale

1 maximum latency: 98.8 ms (beyond axis range)

0.1

0 0 6 et V2 et-B sNet-5 VGG-1 Net ffleN EffN Re Shu Mobile

1905 795

GPU-resident

10

1885

Fig. 7. Single-stage inference under host-side background workloads on RTX A6000.

Fig. 6. Single-stage inference latency.

extend to the observed maxima and reveal long-tail responsetime outliers. Across four models, CPU-CG jitter expands to 1.05 ms to 2.15 ms, up to an order of magnitude larger than GPU-resident, while EfficientNet-B0 reaches 97.99 ms, catching a rare host-side stall, although its median stays low at 832 µs. GPU-resident, in contrast, remains close to the quiet experiment with 9 µs to 86 µs jitter across the five models. This CPU-CG tail behavior was reproducible across runs, not a one-off measurement artifact. More broadly, the larger jitter and corresponding worst-case latencies indicate that a process, even on an isolated core, is not immune to host-side contention. Such latency anomalies did not appear in GPU-resident Execution across multiple runs. Core isolation avoids scheduler preemption of the CPUCG process. However, every iteration still issues CUDA host API calls which have no timing guarantees in CUDA’s userand kernel-space drivers. In addition, background memory and I/O pressure can affect these calls even on isolated cores, for example via TLB shootdowns induced by stress-ng, and delayed driver or IRQ activity. Because of the timing-predictable design of GPU-resident Execution, a device-resident Graph is immune to such uncontrollable timing effects on the host. Emulated I/O. Next, we design a lightweight closedloop emulated I/O experiment and put it under the same background workloads. To emulate a sensor/actuator pair, two SCHED_FIFO (99 priority) CPU threads with small, locked memory footprints are pinned to isolated cores: the sensor thread releases a 1 kHz frame token and timestamps it, while the actuator thread busy-polls the completion and timestamps the response. The processing component between the sensor and actuator can be CPU-driven or GPU-resident. The latency is the difference between the actuator’s and sensor’s timestamps. For this experiment, the CPU-driven configuration is designed as a kernel-bypass GPU-memory-RDMA-style deployment: the input token is visible in a page-locked (pinned) buffer, but host-side control still polls it, launches a CUDA Graph, waits for completion, and emits the output token. The GPUresident configuration uses the same emulated endpoints, but the data-ready handler polls the device-mapped and pinned input token from the GPU, and the final graph node writes the output token back to host-visible memory, modeling GPU-memory RDMA plus GPU-driven I/O. Unlike the previous experiments, latency here is observed by external threads. For processing workload, we use six 512×512 fp32 matrix multiplication

kernels that run under a millisecond on RTX A6000. For this lightweight workload, each configuration runs for 1,000,000 loop iterations (first 1,000 dropped as warmup). Figure 8a shows the latency distribution and variation metrics. Compared with CPU-CG, GPU-resident reduces standard deviation from 14.5 µs to 9.3 µs, p99.9 − p50 from 25.7 µs to 16.5 µs, and jitter from 239 µs to 189 µs. The improvements are not as large as the inference experiments, but GPU-resident Execution achieves this without reserving any host cores. Next, we examine whether the benefits of GPU-resident carry over to other platforms. IGX Thor. We repeat the above two contention experiments on IGX Thor to test whether the result carries to an SoC platform. We also add a CPU-side mitigation with CPU-CG that busy-waits for CUDA completion (cudaDeviceScheduleSpin) instead of using default blocking semantics. Busy-waiting can reduce wakeup latency, but increases CPU utilization and can starve lower-priority work. For the small emulated-I/O workload, Fig. 8b shows that busy-waiting can make CPU-CG behave similarly to GPUresident. Still, CPU remains busy polling for events, whereas CPU core counts are limited in SoCs (see Fig. 2) and do not scale with more sensors, requiring more CPU cores to run the

Isol. CPU-CG (blocking) GPU-resident

Isol. CPU-CG (busy-wait) (b) IGX Thor 1250

latency (μs)

latency (μs)

(a) RTX A6000 1000

800 700

latency (μs)

300

−21% 239 189

200 100

y-axis: log scale

1100

900

−36%

−36% 14.5

0

ev

stdd

25.7 16.5 9.3

−p 50

p 99.9

950 800

1000

latency (μs)

latency (μs)

MobileNetV2

latency (ms)

Isolated CPU-driven CUDA Graph ShuffleNet

y-axis: log scale

100 10 1

r

jitte

p 50 dev std p 99.9−

er

jitt

Fig. 8. Matrix multiplication under host-side contention: latency distributions (top) and variation metrics (bottom).

9

Isol. CPU-CG (blocking) GPU-resident

Isol. CPU-CG (busy-wait)

Isol. CPU-CG (6 cores)

Isol. CPU-CG (10 cores) GPU-resident (a) p99.9 (μs) (c) stddev (μs) (b) max (ms)

y-axis: log scale

100

10

1000

10

latency

latency (ms)

Isol. CPU-CG (2 cores)

1

30

200 100

1

y-axis log scale

20 y-axis log scale

0.1

0.1

ShuffleNet

MobileNetV2

EffNet-B0

10

10

ResNet-50

0 1

5

10 15 20

1

5

10 15 20

1

5

10 15 20

number of concurrent loops

Fig. 9. Single-stage inference latency under host-side background workloads on IGX Thor. Isol. CPU-CG (6 cores) GPU-resident

WHILE node latency (μs) [log scale]

Isol. CPU-CG (2 cores) Isol. CPU-CG (10 cores)

Fig. 11. Concurrent loop latency tail and variation.

latency (μs)

1250 1000 750 500 250 0

Device Graph Tail Launch

(a) maximum

(b) standard deviation 10

1000

5

500

2

200

1

100

64

128

256

512

64

128

256

512

matrix dimension 1

5

10

15

20

number of concurrent loops

Fig. 12. WHILE node vs. device graph tail launch.

Fig. 10. Concurrent loop with MPS static partitioning.

from the box, and outliers are omitted. The figure shows that GPU-resident latency remains stable as the concurrentpolling threads. For the inference workloads, however, Fig. 9 loop count increases, whereas no CPU-CG configuration shows that busy-waiting does not remove the costly CUDA maintains this stability. Figure 11 reveals the p99.9 and max calls, runtime paths, and other host overheads, so the CPU-CG latencies (both log scale), and standard deviation. For GPUlatency tail remains exposed to host contention. Together, these resident Execution, all three metrics rise from one to five loops, results show why a mitigation in CPU-driven execution is not meaning that device-side interference does not disappear in a robust substitute: it can tune around one surface of the shared GPU-resident Execution. Cache, memory-system, and other CPU–GPU stack, but the effective solution depends on the device-side effects still need better hardware controls [67]. application, workload, and other factors. Indeed, the mitigations However, after this initial rise, the tail metrics remain relatively aim to practically avoid the unpredictable components of the stable: from 5 to 20 loops, GPU-resident Execution’s p99.9 and surface, but do not address them fundamentally. The GPU- maximum latency stay within 2×, and stddev stays within 2.5×. resident execution paradigm prevents any host-side interference CPU-CG compounds those same device-side effects with hostby design, with a device-resident steady-state real-time path. side work-queue scheduling and runlist submission [87]; its 3) Scaling to Concurrent Loops: The next experiment tests max latency and stddev spike across loop and core counts. the scalability of GPU-resident Execution against concurrent Relative to its own five-loop values, CPU-CG worst-case sensor processing loops, running matrix multiplications kernels. latency grows up to 30×, p99.9 up to 4×, and stddev up to We co-run 1, 5, 10, 15, and 20 such loops, each assigned a 16×. At 10 concurrent loops, compared with the CPU-CG disjoint 4-SM static CUDA MPS partition [65], the smallest configuration using 10 isolated cores, GPU-resident Execution partition on the A6000. As before, each CPU-CG loop is also reserves no CPU cores and still reduces standard deviation pinned to an isolated CPU core. We vary the isolated-core from 29.3 µs to 7.6 µs, p99.9 from 1113.6 µs to 50.2 µs, and count to 2/6/10, a realistic limit for embedded SoCs where maximum latency from 1850.3 µs to 404.5 µs. This shows that remaining cores serve the OS and other system tasks. To GPU-resident Execution makes concurrent loops more stable demonstrate timing predictability, scaling to concurrent loops even when each CPU-CG loop receives an isolated core. Thus, must keep the latency of individual loops stable. We use the GPU-resident demonstrates a more scalable and predictable same measurement regime as in the emulated I/O experiment. design foundation for multi-sensor processing applications with concurrent GPU-resident loops. Figure 10 shows the per-iteration latency distributions for 4) Framework Overheads: The final experiments focus on each measurement using Tukey box plots. Their whiskers reach the framework costs. First, we show that the two GPU-resident the furthest observations within 1.5× the interquartile range iteration driver mechanisms, WHILE node (WHL) and Device

10

camera

TABLE I F RAMEWORK OVERHEADS B REAKDOWN Edge DRH → CTRL CTRL kernel CTRL → BODY

body → end end → DRH

display

clock

MIPI

IGX

HSB GPIO

Graph Component Worst-case (µs) handler-to-controller transition 0.82 controller kernel 1.67 body-release transition 3.69 body-completion transition 1.37 next-iteration transition 5.67

Pin0

Pin1

frame led trigger start (T )

photon-to-display latency

Photodetector trigger (TPD)

LED

Fig. 13. HSB-native photon-to-display latency tool. CPU-driven (no GDR) CPU-driven (w/ GDR)

(a) Photon-to-display Latency 40

GPU-resident

(b) CUDA Processing Latency 3.0

2.6

J=22

2.5 30 J=4.83

20

J=0.66

10 0

Latency (ms)

Avg. Latency (ms)

Graph Tail Launch (DVG), have similar performance. Then, we measure the WHL mechanism overheads. WHL uses CUDA Graph WHILE conditional nodes, whereas DVG achieves the same behavior by tail-launching device graphs. Fig. 12 shows that the two mechanisms are equivalent on a workload of 10 matrix multiplication operators. Across matrix dimensions, max latency differs by at most 6% and standard deviation by at most 3 µs, with neither mechanism consistently ahead. We support both mechanisms but make WHL the default because DVG does not support advanced CUDA Graph features such as conditional nodes, which many CUDA-based libraries increasingly use [63]. We characterize the overhead of GPU-resident Execution on RTX A6000 with an empty processing kernel experiment in Table I. These platform-specific fixed costs can be measured once per new platform for timing analysis.

2.0

1.0 0.5

24.0

16.2

12.8

1.7

1.5 1.0 0.3

0.3 0.01

0.0 avg

max

jitter

Fig. 14. IMX274 camera-to-display case-study latency.

C. Case Study: IMX274 Camera-to-display Pipeline We evaluate the IMX274 camera-to-display application the IMX274 frame-start event directly. For this rolling-shutter from §IV-D on IGX Orin [14] with an RTX 6000 Ada. The sensor, a HSB event sequencer drives GPIO Pin0 high at frame connected display is a 240 Hz G-SYNC monitor. This setup start, illuminating an LED in the camera’s field of view during is representative of industry-grade real-time streaming sensor the current frame readout. After the frame is sent, processed, pipelines in medical devices [88], surgical robotics [89], [90], and displayed on the monitor, a photodiode mounted on the medical imaging and endoscopy [91], [92], and broadly, in display drives GPIO Pin1 when the illuminated region appears. industrial automation domains. The LED-trigger timestamp (TLED ) and photodiode timestamp We measure the performance of three configurations in (TPD ) are latched by the same FPGA clock, so the reported this case study. 1) CPU-driven (no GDR) (GDR means photon-to-display latency is the hardware-timestamped interval GPUDirect RDMA [27] which is NVIDIA’s implementation of TPD − TLED . GPU-memory RDMA) keeps both the data and control path on Result. The left panel of Figure 14 compares photon-toLinux on the CPU; 2) CPU-driven (w/ GDR) uses GDR for display latency (average with min–max whiskers). The CPUthe frame data path but still uses CPU-side receive completion, only Linux receiver has an average of 24 ms and a 22 ms frame selection, CUDA launch, and display handoff. Vast jitter. Keeping the frame data path in GPU memory with GDR majority of industry and research deployments use these two reduces that span to 4.83 ms, but it is still affected by host mechanisms [91], [90], [93], [94]. Both CPU-driven baselines orchestration variation, which is problematic in surgical robotics are developed with the CPU-driven Holoscan SDK [36]. 3) and medical devices. The GPU-resident graph eliminates the GPU-resident is our proposed framework with GDR and source of variance, achieving sub-millisecond jitter (0.66 ms), DOCA GPUNetIO for GPU-driven I/O. The HSB metadata a 7.3× reduction relative to CPU-driven (w/ GDR) and 33× handler has similar latency but relies on CPU-side polling, relative to CPU-driven (no GDR). The mean also moves below making it vulnerable to host contention; we omit it since the to 12.81 ms for GPU-resident versus 16.2 ms for CPU-driven gains are clear even without added contention. (w/ GDR) and 24 ms for CPU-driven (no GDR). Crucially Measurement. We measure 1080p photon-to-display latency for time-sensitive domains, the worst-case photon-to-display with a HSB-native closed-loop tool, shown in Fig. 13. The latency drops to 13.35 ms for GPU-resident, compared with camera stream enters the HSB over MIPI, and the correspond- 18.83 ms for CPU-driven (w/ GDR) and 35 ms for CPUing frame payload is forwarded over the sensor bridge and driven (no GDR). The right panel of Figure 14 isolates the ConnectX-7 NIC to the IGX. The configurations receive frames software processing interval after frame arrival into the platform. through either CPU memory (config 1) or GPUDirect (config 2 The CPU-driven (w/ GDR) pipeline averages 0.96 ms, reaches and 3). Independently of the data path, the HSB FPGA observes 2.60 ms at the maximum, and has J=1.70 ms. The GPU-

11

resident path averages 0.30 ms, reaches only 0.31 ms at the maximum, and has J=0.01 ms. The processing path is therefore 3.2× faster on average, 8.4× lower at the maximum, and 170× lower in absolute jitter. Overall, GPU-resident control removes the millisecond-scale latency tails with host-driven execution for this streaming camera application.

the GPU and schedule work on device. Their mechanisms range from task queues and GPU-resident runtimes to CUDA device graph launch [62]. Persistent kernels require modification to CUDA kernels, which GPU-resident does not need. These systems share our goal of CPU bypass, but none integrates sensing, actuation, and display I/O as first-class primitives. In addition, most make dynamic on-device scheduling part of the application runtime, affecting predictability and analysis. Prior real-time GPU work spans driver-level schedulers and managers [100], [101], [102], [103], [104], [105], middleware and pipeline schedulers for accelerator-backed sensor workloads [106], [107], [108], [109], [110]. These systems improve predictability by regulating launches from the host, GPU context arbitration, or task release in the host pipeline. In contrast, GPU-resident Execution introduces a new paradigm that stops treating the GPU as a host-subordinate client and avoids host-mediated orchestration in the time-critical path. It establishes the GPU as an independent real-time accelerator [111] by moving the end-to-end path, including the I/O, onto the device. GPU resource partitioning and reservation mechanisms [65], [67], [112], [113] support this shift by exposing dedicated SM partitions for concurrent timingpredictable workloads. Scheduling a taskset therefore becomes a design-time partitioning problem: for each GPU-resident loop, the response-time bounds the captured graph body on its partition; for multiple loops, a partition configuration must satisfy every loop’s timing requirement.

VI. D ISCUSSION

Scope. The primary significance of the results in §V is to demonstrate the structural benefits of GPU-resident Execution, not just the quantitative improvements. A CPU-driven loop, even if tuned with CUDA Graphs, GPUDirect RDMA, MPS, and core isolation, is susceptible to host latency tails. Mitigations such as core isolation try to bypass the host-side challenges that we identified, while GPU-resident Execution removes the host altogether from steady-state. Removing this source of variance produces the orders-of-magnitude reductions in latency jitter observed in the streaming camera application. For predictable latency, GPU-resident Execution mandates static device-memory layout, fixed clocks, no steady-state CPU interaction, and an execution working set that fits in device memory. These properties already match the common case for certifiable sensor pipelines whose memory footprint and rate are known at design time. The configurable data-not-ready sleep interval trades polling latency for GPU utilization and power. Furthermore, concurrent loops must be confined to a configured SM partition per GPU-resident graph application. Unmanaged compute- or graphics-context coexistence can affect predictability. For this reason, our case study uses a VIII. C ONCLUSION AND F UTURE W ORK GPU-resident display operator, avoiding a graphics context. Real-time ML-driven sensor processing needs both preIn our design, the CPU supervises the non-critical plane, while the GPU runs the time-critical sensor processing loop. dictable end-to-end latency and the advanced software stacks Broadly, our deliberately straightforward device-resident design available on Linux-hosted GPU platforms. This paper presented demonstrates empirically verified timing predictability on GPU-resident execution, which keeps a steady-state closedloop system on the device for its lifetime, with GPU-native industry-grade applications that no prior work has achieved. GPU-resident graphs as asynchronous tasks. Each GPU- ingress, egress, and control. By combining an operator and resident graph behaves as an asynchronous task that runs on application API, device-side data-ready handlers, GPU-visible a fixed SM budget for the application lifetime. This work I/O, and a long-lived CUDA Graph, the design removes perconsiders independent tasks, including the concurrent loops frame host wakeups and CUDA launches from the critical evaluated on disjoint SM partitions. Within one graph, the path. The design yields a predictable system design with current handler abstraction supports a sensor fusion pipeline device-side execution and graph-transition costs, rather than by polling multiple sensors with parallel CUDA threads before Linux kernel, drivers, or CUDA driver-dispatch latency tails. issuing a combined dataReady signal, although our evaluation Across microbenchmarks and an industry-grade camera-tocurrently uses one sensor per handler. Furthermore, the frame- display application, experiments show that this architectural work does not define communication or synchronization among shift tightens latency distributions and reduces end-to-end jitter distinct GPU-resident tasks or between GPU-resident and CPU by orders of magnitude on NVIDIA SoCs and GPUs. Future work will focus on sense–think–act loops [114] with tasks. Designing and analyzing predictable communication broader GPU-driven I/O support, including HSB-style FPGAamong such asynchronous tasks remains future work. backed adapters for CAN/SPI-connected motor control and direct control of SPI, GPIO, and I2C. This work lays the VII. R ELATED W ORK groundwork for deploying GPU-native action models such as The sensor-I/O background and GPU mechanisms behind GR00T N1 [115] inside real-time control loops. As the broader this work are already covered in §II. We now place our work trend in generative AI and robotics shifts from producing in the context of the closest adjacent research. pixels [116] to producing actions [115], our architecture Persistent kernels [95] and graphs [96], Whippletree [97], provides a timing-predictable systems foundation for GPUPERKS [98], and CharminG [99] keep execution resident on native sensor processing systems.

12

R EFERENCES [1] L. Chen, P. Wu, K. Chitta, B. Jaeger, A. Geiger, and H. Li, “End-toend autonomous driving: Challenges and frontiers,” IEEE Transactions on Pattern Analysis and Machine Intelligence, vol. 46, no. 12, pp. 10 164–10 183, 2024. [2] A. Roychoudhury, S. Khorshidi, S. Agrawal et al., “Perception for humanoid robots,” Current Robotics Reports, vol. 4, pp. 127–140, 2023. [3] K. Lee, H.-S. Cho, and J. Lee, “RT-BEV: An open-source framework for real-time BEV perception,” in IEEE Real-Time Systems Symposium (RTSS), 2023. [4] NVIDIA Corporation, “Mercedes-Benz and NVIDIA to build softwaredefined architecture for automated driving,” NVIDIA Newsroom, 2020. [5] ——, “Hyundai motor, kia and NVIDIA expand strategic partnership for next-generation autonomous driving technology,” NVIDIA Newsroom, 2025. [Online]. Available: https://nvidianews.nvidia.com/ news/hyundai-motor-kia-autonomous-driving [6] N. Tijtgat, W. Van Ranst, B. Volckaert, T. Goedemé, and F. De Turck, “Embedded real-time object detection for a UAV warning system,” in IEEE International Conference on Computer Vision Workshops (ICCVW), 2017, pp. 2110–2118. [7] S. Hossain and D.-J. Lee, “Deep learning-based real-time multipleobject detection and tracking from aerial imagery via a flying robot with GPU-based embedded devices,” Sensors, vol. 19, no. 15, 2019. [8] Boston Dynamics, “Spot Core I/O,” https://bostondynamics.com/ products/payload/spot-core-io/, 2025. [9] NVIDIA Corporation, “AI platform to bring real-time sensing to medical devices,” https://blogs.nvidia.com/blog/ clara-holoscan-real-time-ai-medical-devices/, 2022. [10] F. A. Ahmed, M. Yousef, M. A. Ahmed et al., “Deep learning for surgical instrument recognition and segmentation in robotic-assisted surgeries: A systematic review,” Artificial Intelligence Review, vol. 58, 2025. [11] NVIDIA Corporation, “Build a real-time visual inspection pipeline with NVIDIA TAO and NVIDIA DeepStream,” NVIDIA Technical Blog, 2025. [12] Y. Liu, C. Zhang, and X. Dong, “A survey of real-time surface defect inspection methods based on deep learning,” Artificial Intelligence Review, vol. 56, pp. 12 131–12 170, 2023. [13] NVIDIA Corporation, “Jetson modules, support, ecosystem, and lineup,” https://developer.nvidia.com/embedded/jetson-modules, 2025. [14] ——, “NVIDIA IGX Orin,” https://www.nvidia.com/en-us/ edge-computing/products/igx/, 2026. [15] ——, “NVIDIA IGX Thor,” https://www.nvidia.com/en-us/ edge-computing/products/igx/, 2026. [16] Advanced Micro Devices, Inc., “AMD Zynq UltraScale+ MPSoC,” https://www.amd.com/en/products/adaptive-socs-and-fpgas/soc/ zynq-ultrascale-plus-mpsoc.html, 2026, eG and EV variants include Arm Mali-400 MP2 GPU. [17] NVIDIA Corporation, “Getting started with CUDA graphs,” https:// developer.nvidia.com/blog/cuda-graphs/, 2019. [18] D. B. de Oliveira, D. Casini, and T. Cucinotta, “Operating system noise in the linux kernel,” IEEE Transactions on Computers, vol. 72, no. 1, pp. 196–207, 2022. [19] D. B. de Oliveira, D. Casini, R. S. de Oliveira, and T. Cucinotta, “Demystifying the real-time Linux scheduling latency,” in Euromicro Conference on Real-Time Systems (ECRTS), 2020. [20] Y. Zhang and R. West, “Process-aware interrupt scheduling and accounting,” in IEEE Real-Time Systems Symposium (RTSS), 2006. [21] M. M. Madden, “Challenges using Linux as a real-time operating system,” in AIAA SciTech Forum, 2019, NASA Technical Reports Server, Document ID 20200002390. [22] E. Missimer, K. Missimer, and R. West, “Mixed-criticality scheduling with I/O,” in Euromicro Conference on Real-Time Systems (ECRTS), 2016. [23] Linux Kernel Documentation, “Real-time linux: Theory of operation,” https://docs.kernel.org/core-api/real-time/theory.html, 2026. [24] F. Reghenzani, G. Massari, and W. Fornaciari, “The real-time linux kernel: A survey on PREEMPT_RT,” ACM Computing Surveys, vol. 52, no. 1, 2019. [25] T. von Eicken, A. Basu, V. Buch, and W. Vogels, “U-Net: A userlevel network interface for parallel and distributed computing,” in ACM Symposium on Operating Systems Principles (SOSP), 1995, pp. 40–53.

13

[26] DPDK Project, “Data Plane Development Kit (DPDK),” https://www. dpdk.org/, 2026. [27] NVIDIA Corporation, “GPUDirect RDMA documentation,” https://docs. nvidia.com/cuda/gpudirect-rdma/, 2024. [28] Advanced Micro Devices, Inc., “ROCm gpu-enabled message passing interface,” https://rocm.docs.amd.com/en/develop/how-to/ gpu-enabled-mpi.html, 2026. [29] NVIDIA Corporation, “DOCA GPUNetIO: Real-time GPU-driven network processing,” https://docs.nvidia.com/doca/sdk/DOCA-GPUNetIO/ index.html, 2026. [30] Advanced Micro Devices, Inc., “rocm-xio: A ROCm library for GPUinitiated IO,” https://github.com/ROCm/rocm-xio, 2026. [31] A. Betts and A. F. Donaldson, “Estimating the WCET of GPUaccelerated applications using hybrid analysis,” in Euromicro Conference on Real-Time Systems (ECRTS), 2013, pp. 193–202. [32] K. Berezovskyi, L. Santinelli, K. Bletsas, and E. Tovar, “WCET measurement-based and extreme value theory characterisation of CUDA kernels,” in International Conference on Real-Time Networks and Systems (RTNS), 2014, pp. 279–288. [33] L. Jeanmougin, T. Carle, and C. Rochange, “Bounding the WCET of a GPU thread block with a multi-phase representation of warps execution,” in Euromicro Conference on Real-Time Systems (ECRTS), 2025, pp. 11:1–11:26. [34] G. Li, P. Li, G. Sawaya, G. Gopalakrishnan, I. Ghosh, and S. P. Rajan, “GKLEE: Concolic verification and test generation for GPUs,” in ACM SIGPLAN Symposium on Principles and Practice of Parallel Programming (PPoPP), 2012, pp. 215–224. [35] A. Betts, N. Chong, A. F. Donaldson, S. Qadeer, and P. Thomson, “GPUVerify: A verifier for GPU kernels,” in ACM SIGPLAN Conference on Object-Oriented Programming, Systems, Languages, and Applications (OOPSLA), 2012, pp. 113–132. [36] NVIDIA Corporation, “Holoscan SDK: AI sensor processing platform,” https://docs.nvidia.com/holoscan/sdk-user-guide/, 2026. [37] Open Source Robotics Foundation, “ROS 2 documentation,” https: //docs.ros.org/, 2024. [38] NVIDIA Corporation, “NVIDIA DeepStream SDK,” https://developer. nvidia.com/deepstream-sdk, 2026. [39] ——, “NVIDIA Isaac ROS,” https://developer.nvidia.com/isaac-ros, 2026. [40] Advanced Micro Devices, Inc., “Vitis video analytics SDK (VVAS),” https://www.amd.com/en/developer/resources/ vitis-video-analytics-sdk.html, 2026. [41] M. Danish, Y. Li, and R. West, “Virtual-CPU scheduling in the Quest operating system,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2011, pp. 169–179. [42] BlackBerry QNX, “QNX Neutrino RTOS: System architecture guide,” https://www.qnx.com/developers/docs/7.1/com.qnx.doc.neutrino. sys_arch/topic/intro.html, 2024. [43] Y. Li, E. Missimer, and R. West, “Predictable migration and communication in the Quest-V multikernel,” arXiv:1310.6301, 2014. [44] S. Sinha and R. West, “Towards an integrated vehicle management system in DriveOS,” ACM Transactions on Embedded Computing Systems, vol. 20, no. 5s, 2021. [45] S. Sinha, “Towards a centralized multicore automotive system,” Ph.D. dissertation, Boston University, 2022. [Online]. Available: https://hdl. handle.net/2144/45640 [46] NVIDIA Corporation, “DRIVE OS QNX installation guide,” https://developer.nvidia.com/docs/drive/drive-os/6.0.6/public/ drive-os-qnx-installation/index.html, 2024. [47] ——, “CUDA safety installation guide for DRIVE OS (QNX),” https://docs.nvidia.com/drive/drive-os-5.2.0.0L/drive-qsg-nv/ cuda-safety-installation-guide/, 2024. [48] I. S. Olmedo, N. Capodieci, J. L. Martínez, A. Marongiu, and M. Bertogna, “Dissecting the CUDA scheduling hierarchy: A performance and predictability perspective,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2020. [49] J. Bakita and J. H. Anderson, “Demystifying NVIDIA GPU internals to enable reliable GPU management,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2024. [50] S. Liu, R. Wagle, J. H. Anderson, M. Yang, C. Zhang, and Y. Li, “Autonomy today: Many delay-prone black boxes,” in 36th Euromicro Conference on Real-Time Systems (ECRTS 2024), ser. Leibniz International Proceedings in Informatics (LIPIcs), vol. 298. Schloss Dagstuhl – Leibniz-Zentrum für Informatik, 2024, pp. 12:1–12:27.

[51] NVIDIA Corporation, “Benchmarking GPUDirect RDMA on modern server platforms,” https://devblogs.nvidia.com/parallelforall/ benchmarking-gpudirect-rdma-on-modern-server-platforms/, 2017. [52] A. Sapio, M. Canini, C.-Y. Ho, J. Nelson, P. Sudber et al., “Scaling distributed machine learning with in-network aggregation,” in USENIX Symposium on Networked Systems Design and Implementation (NSDI), 2021. [53] NVIDIA Corporation, “Holoscan Sensor Bridge documentation,” https: //docs.nvidia.com/holoscan/sensor-bridge/, 2026. [54] ——, “Holoscan Sensor Bridge (open-source fpga IP and host software),” https://github.com/nvidia-holoscan/holoscan-sensor-bridge, 2026. [55] E. Agostini, D. Rossetti, and S. Potluri, “Offloading communication control logic in GPU accelerated applications,” in IEEE/ACM Int. Symp. on Cluster, Cloud and Grid Computing (CCGrid), 2017, pp. 248–257. [56] ——, “GPUDirect Async: Exploring GPU synchronous communication techniques for InfiniBand clusters,” Journal of Parallel and Distributed Computing, vol. 114, pp. 28–45, 2018. [57] NVIDIA Corporation, “Inline GPU packet processing with NVIDIA DOCA GPUNetIO,” https://developer.nvidia.com/blog/ inline-gpu-packet-processing-with-nvidia-doca-gpunetio/, 2023. [58] ——, “Unlocking GPU-accelerated RDMA with NVIDIA DOCA GPUNetIO,” https://developer.nvidia.com/blog/ unlocking-gpu-accelerated-rdma-with-nvidia-doca-gpunetio/, 2024. [59] R. Liao et al., “The landscape of GPU-centric communication,” arXiv preprint arXiv:2409.09874, 2024. [60] NVIDIA Corporation, “CUDA stream handling in Holoscan applications,” https://docs.nvidia.com/holoscan/sdk-user-guide/holoscan_cuda_ stream_handling.html, 2026. [61] ——, “Dynamic control flow in CUDA graphs with conditional nodes,” https://developer.nvidia.com/blog/ dynamic-control-flow-in-cuda-graphs-with-conditional-nodes/, 2024. [62] S. Stevenson, “Enabling dynamic control flow in CUDA graphs with device graph launch,” https://developer.nvidia.com/blog/ enabling-dynamic-control-flow-in-cuda-graphs-with-device-graph-launch/, 2022. [63] PyTorch Contributors, “Add support for torch.cond with CUDA graphs,” GitHub Pull Request #168912, 2026. [Online]. Available: https://github.com/pytorch/pytorch/pull/168912 [64] NVIDIA Corporation, “GPU-resident graph execution api,” https://docs. nvidia.com/holoscan/sdk-user-guide/gpu_resident.html, 2026. [65] ——, “CUDA Multi-Process Service: Static SM partitioning,” https: //docs.nvidia.com/deploy/mps/, 2026. [66] ——, “NVIDIA system management interface,” https://docs.nvidia.com/ deploy/nvidia-smi/, 2026. [67] S. Jain, I. Baek, S. Wang, and R. Rajkumar, “Fractional GPUs: Softwarebased compute and memory bandwidth reservation for GPUs,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2019, pp. 29–41. [68] Y. Zhang, H. Yu, C. Han, C. Wang, B. Lu, Y. Li, Z. Jiang, Y. Li, X. Chu, and H. Li, “SGDRC: Software-defined dynamic resource control for concurrent DNN inference on NVIDIA GPUs,” in Proceedings of the 30th ACM SIGPLAN Annual Symposium on Principles and Practice of Parallel Programming, 2025, pp. 267–281. [69] Intel Corporation, “Intel Resource Director Technology (Intel RDT) on 2nd generation Intel Xeon scalable processors,” https://www.intel.com/content/dam/develop/external/us/en/documents/ 180115-intel-rdtcascadelake-serverreferencemanual-806717.pdf, 2019. [70] NVIDIA Corporation, “NVIDIA performance primitives (NPP),” https: //developer.nvidia.com/npp, 2026. [71] Leopard Imaging Inc., “LI-IMX274-MIPI-CS: IMX274 8m 4k camera module with CS lens,” https://leopardimaging.com/product/ robotics-cameras/cis-2-mipi-modules/i-pex-mipi-camera-modules/ rolling-shutter-mipi-cameras/8-49mp-imx274/li-imx274-mipi-cs/, 2026. [72] InfiniBand Trade Association, “InfiniBand architecture specification volume 1,” InfiniBand Trade Association, Tech. Rep., 2015, release 1.3. Defines RDMA Write and RDMA Write with Immediate operations. [73] NVIDIA Corporation, “RDMA aware networks programming user manual,” https://docs.nvidia.com/ rdma-aware-networks-programming-user-manual-1-7.pdf, 2026. [74] ——, “ConnectX-7 SmartNIC,” https://www.nvidia.com/en-us/ networking/ethernet-adapters/connectx-7/, 2024.

14

[75] Linux Kernel Documentation, “Direct rendering manager (DRM) internals,” https://docs.kernel.org/gpu/drm-internals.html, 2026. [76] S. Sinha, S. Dwivedi, and M. Azizian, “Towards deterministic end-to-end latency for medical AI systems in NVIDIA Holoscan,” in ACM/IEEE Int. Conf. on Cyber-Physical Systems (ICCPS), 2024, pp. 235–246. [77] NVIDIA Corporation, “Source code for IMX274 GPU-resident application,” https://nvidia-holoscan.github.io/holohub/applications/imx274_ gpu_resident/, 2026. [78] G. C. Buttazzo, Hard Real-Time Computing Systems: Predictable Scheduling Algorithms and Applications, 3rd ed. Springer, 2011. [79] A. Cervin, B. Lincoln, J. Eker, K.-E. Årzén, and G. C. Buttazzo, “The jitter margin and its application in the design of real-time control systems,” in International Conference on Real-Time and Embedded Computing Systems and Applications (RTCSA), 2004, pp. 1–9. [80] N. Ma, X. Zhang, H.-T. Zheng, and J. Sun, “ShuffleNet V2: Practical guidelines for efficient CNN architecture design,” in European Conference on Computer Vision (ECCV), 2018, pp. 116–131. [81] M. Sandler, A. Howard, M. Zhu, A. Zhmoginov, and L.-C. Chen, “MobileNetV2: Inverted residuals and linear bottlenecks,” in IEEE Conference on Computer Vision and Pattern Recognition (CVPR), 2018, pp. 4510–4520. [82] M. Tan and Q. V. Le, “EfficientNet: Rethinking model scaling for convolutional neural networks,” in International Conference on Machine Learning (ICML), 2019, pp. 6105–6114. [83] K. He, X. Zhang, S. Ren, and J. Sun, “Deep residual learning for image recognition,” in IEEE Conference on Computer Vision and Pattern Recognition (CVPR), 2016, pp. 770–778. [84] K. Simonyan and A. Zisserman, “Very deep convolutional networks for large-scale image recognition,” in International Conference on Learning Representations (ICLR), 2015. [85] J. Kosaian and K. V. Rashmi, “Arithmetic-intensity-guided fault tolerance for neural network inference on GPUs,” in International Conference for High Performance Computing, Networking, Storage and Analysis (SC), 2021, pp. 1–14. [86] C. I. King, “stress-ng: Stress test a computer system,” https: //github.com/ColinIanKing/stress-ng, 2026. [87] J. Bakita and J. H. Anderson, “Hardware compute partitioning on NVIDIA GPUs for composable systems,” in 37th Euromicro Conference on Real-Time Systems (ECRTS), ser. Leibniz International Proceedings in Informatics (LIPIcs), vol. 335. Schloss Dagstuhl – Leibniz-Zentrum für Informatik, 2025, pp. 21:1–21:25. [88] Stryker Corporation, “1688 Advanced Imaging Modalities (AIM) 4K platform,” https://www.stryker.com/us/en/endoscopy/systems/ 1688-aim-platform.html, 2026. [89] Intuitive Surgical, Inc., “da Vinci robotic surgical systems,” https://www. intuitive.com/en-us/products-and-services/da-vinci, 2026. [90] Orsi Academy, “Orsi academy brings real-time AI to the OR for a firstof-its-kind augmented reality surgery,” https://www.orsi-online.com/ news/, 2023. [91] Olympus Corporation, “Endoscopic ultrasound (EUS),” https://medical. olympusamerica.com/procedure/endoscopic-ultrasound-eus, 2026. [92] Johnson & Johnson MedTech, “MONARCH platform for roboticassisted bronchoscopy,” https://www.jnjmedtech.com/en-US/products/ robotics/monarch-platform/bronchoscopy/, 2026. [93] Barco NV, “Nexxis: Uncompressed AV-over-IP platform,” https://www. barco.com/en/product/nexxis, 2026. [94] DELTACAST, “NVIDIA GPUDirect RDMA for video capture and processing,” https://www.deltacast.tv/technologies/gpudirect-rdma, 2026. [95] K. Gupta, J. A. Stuart, and J. D. Owens, “A study of persistent threads style GPU programming for GPGPU workloads,” in Innovative Parallel Computing (InPar), 2012. [96] C. Cetre, F. Ferreira, R. Barrere, and D. Gratadour, “Introducing GPU persistent graphs for time-sensitive workflows,” in ACM/SPEC International Conference on Performance Engineering (ICPE), 2025, pp. 181–191. [97] M. Steinberger, M. Kenzel, P. Boechat, B. Kerbl, M. Dokter, and D. Schmalstieg, “Whippletree: Task-based scheduling of dynamic workloads on the GPU,” ACM Trans. Graph. (Proc. SIGGRAPH Asia), vol. 33, no. 6, 2014. [98] L. Zhang, M. Wahib, P. Chen, J. Meng, X. Wang, T. Endo, and S. Matsuoka, “PERKS: A locality-optimized execution model for iterative memory-bound GPU applications,” arXiv preprint arXiv:2204.02064, 2022.

[99] J. Choi, D. F. Richards, and L. V. Kale, “CharminG: A scalable GPUresident runtime system,” in ACM International Symposium on HighPerformance Parallel and Distributed Computing (HPDC), 2021, pp. 261–262. [100] S. Kato, K. Lakshmanan, R. Rajkumar, and Y. Ishikawa, “TimeGraph: GPU scheduling for real-time multi-tasking environments,” in USENIX Annual Technical Conference (USENIX ATC), 2011. [101] G. A. Elliott, B. C. Ward, and J. H. Anderson, “GPUSync: A framework for real-time GPU management,” in IEEE Real-Time Systems Symposium (RTSS), 2013, pp. 33–44. [102] N. Capodieci, R. Cavicchioli, M. Bertogna, and A. Paramakuru, “Deadline-based scheduling for GPU with preemption support,” in IEEE Real-Time Systems Symposium (RTSS), 2018. [103] A. Zou, J. Li, C. D. Gill, and X. Zhang, “RTGPU: Real-time GPU scheduling of hard deadline parallel tasks with fine-grain utilization,” IEEE Trans. Parallel Distrib. Syst., 2023. [104] M. Han, H. Zhang, R. Chen, and H. Chen, “Microsecond-scale preemption for concurrent GPU-accelerated DNN inferences,” in USENIX Symposium on Operating Systems Design and Implementation (OSDI), 2022, pp. 539–558. [105] Y. Wang, C. Liu, D. Wong, and H. Kim, “GCAPS: GPU context-aware preemptive priority-based scheduling for real-time tasks,” in Euromicro Conference on Real-Time Systems (ECRTS), 2024. [106] D. Enright, Y. Xiang, H. Choi, and H. Kim, “PAAM: A framework for coordinated and priority-driven accelerator management in ROS 2,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2024. [107] Y. Xiang and H. Kim, “Pipelined data-parallel CPU/GPU scheduling for

15

multi-DNN real-time inference,” in IEEE Real-Time Systems Symposium (RTSS), 2019, pp. 392–405. [108] W. Kang, K. Lee, J. Lee, I. Shin, and H. S. Chwa, “LaLaRAND: Flexible layer-by-layer CPU/GPU scheduling for real-time DNN tasks,” in IEEE Real-Time Systems Symposium (RTSS), 2021. [109] Y. Wang, M. Karimi, Y. Kim, and H. Kim, “Balancing energy efficiency and real-time performance in GPU scheduling,” in IEEE Real-Time Systems Symposium (RTSS), 2021. [110] F. Kang, K. Lee, and J. Lee, “RT-Swap: Addressing GPU memory bottlenecks for real-time multi-DNN inference,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2024. [111] J. Bakita and J. H. Anderson, “The advantage of the GPU as a real-time AI accelerator,” Real-Time Systems, vol. 61, pp. 275–280, 2025. [112] ——, “Hardware compute partitioning on NVIDIA GPUs,” in IEEE Real-Time and Embedded Technology and Applications Symposium (RTAS), 2023, pp. 54–66. [113] NVIDIA Corporation, “CUDA Green context for spatial partitioning,” https://developer.nvidia.com/blog/ dynamic-load-balancing-in-cuda-with-green-contexts/, 2024. [114] M. Siegel, “The sense-think-act paradigm revisited,” in 1st International Workshop on Robotic Sensing (ROSE), 2003, p. 5. [115] NVIDIA Corporation, “GR00T N1: An open foundation model for generalist humanoid robots,” arXiv:2503.14734; https://arxiv.org/abs/ 2503.14734, 2025. [116] J. Huang, “Every single pixel will be generated soon. not rendered: generated,” Remarks at NVIDIA GTC 2023; https://www.tomshardware.com/pc-components/gpus/ rtx-off-ai-on-jensen-says-well-see-fully-ai-generated-games-in-5-10-years, 2023.

Related documents

• GPU-resident Execution for Timing-predictable Sensor Processing Systems#1298305
• Exploiting Dependency and Parallelism: Real-Time Scheduling and Analysis for GPU Tasks#2742
• WarpGuard: Towards Control-Flow Attestation for Heterogeneous CPU-GPU Execution#370276
• Penrose: Privacy-Preserving Performance Profiling of In-The-Wild GPUs#1062797
• VDCores: Resource Decoupled Programming and Execution for Asynchronous GPU#155230