Echo: Merging Host Device Buffers to Avoid Redundant Data Movement on Unified Memory SoCs

arXiv:2609.05635v2 [cs.DC] 30 Sep 2026

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs Yuheng Zhu∗

Yanbo Zhao∗

yzhu63@ncsu.edu North Carolina State University Raleigh, North Carolina, USA

yzhao62@ncsu.edu North Carolina State University Raleigh, North Carolina, USA

Jiajia Li

Man-Ki Yoon

jiajia.li@ncsu.edu North Carolina State University Raleigh, North Carolina, USA

man-ki.yoon@ncsu.edu North Carolina State University Raleigh, North Carolina, USA

Abstract GPU applications on unified-memory (UMA) edge platforms often inherit a discrete-GPU memory abstraction in which they allocate one buffer for the CPU, another for the GPU, and copy data between them before and after GPU execution. On UMA hardware these buffers reside in the same physical DRAM pool, so the copies consume bandwidth, time, and energy. Despite the growing adoption of UMA platforms, this pattern remains common because production software stacks, libraries, and samples were written for portability across discrete GPUs. However, removing these copies is not as simple as merging the two buffers, because the original program may rely on the two buffers being distinct, or on the copy itself ordering CPU and GPU accesses. Echo removes these copies only when it preserves the data values and access ordering on which the program depends. It does so along two complementary paths: sourcelevel rewriting and binary deployment. Echo-SR is a compiletime LLVM transformation that proves safety and rewrites accepted pairs in place. Echo-DR is a profile-guided binary optimizer for existing applications that profiles and validates stable allocation/copy patterns and, at runtime, intercepts the matching calls to unify profile-matched pairs while preserving the ordering effects of removed copies. Across seven benchmarks on three NVIDIA Jetson platforms, Echo’s benefit grows with the fraction of baseline time spent on copies. Copy-dominated workloads speed up by up to 7.05×, closed-source end-to-end applications speed up by up to 1.40×. On the five source-available Orin benchmarks, both automatic paths recover ≥99% of the gain of a manually optimized reference.

1

Introduction

Heterogeneous systems-on-chip (SoCs) that integrate a CPU and a GPU on the same die with a shared physical DRAM, known as unified-memory architectures (UMA), are increasingly prevalent in edge inference, robotics, autonomous ∗ Contributed equally.

systems, and mobile computing [3, 26]. Sharing one physical DRAM between processors removes the off-chip data path that limits discrete CPU–GPU systems, while reducing power, cost, and board area. Yet most accelerator code on UMA platforms continues to manage memory as if the GPU were discrete. Applications ported from discrete-GPU (dGPU) platforms or built on top of legacy libraries and frameworks still preserve the discrete-GPU programming model: allocate a host buffer, allocate a device buffer, copy host-to-device (H2D) before GPU execution, and copy device-to-host (D2H) afterward. Because both buffers reside in the same physical DRAM, these copies are redundant. It consumes unnecessary DMA bandwidth, CPU and GPU cycles, and energy, as it only transfers data between two virtual addresses backed by the same memory, without crossing any off-chip device boundary. The root cause is not a hardware limitation but inertia. Widely-used libraries, manufacturer-provided samples, and developer workflows were designed for discrete GPUs and remain largely unchanged even as UMA platforms became common. Portability concerns further reinforce it, since code that must also run on discrete-GPU systems follows the same legacy pattern by default. Among accelerator programming models, CUDA (Compute Unified Device Architecture) is the most widely deployed on UMA edge platforms. NVIDIA Jetson, the dominant commercial UMA SoC family, exposes its compute capability through CUDA, and the production inference stacks built on top of it (e.g., TensorRT, cuDNN, cuBLAS) follow the same host-device separate-buffer convention by default. For example, NVIDIA’s TensorRT samples [22] ship a BufferManager class that allocates host and device buffers separately and copies data between them explicitly, and downstream developers often reuse the same structure. We conduct a static survey of 34 open-source Jetson GPU repositories (§2.4) and find that 19 of them (56%) still contain the legacy pattern, rising to 10 of 11 (91%) among inferencespecific projects. Together, these numbers show that the GPU software ecosystem has not fully adapted to UMA edge

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

ptr = cudaMallocManaged(N);

host = malloc(N); cudaMalloc(&dev, N);

produce_input(ptr);

kernel<<<G,B>>>(dev); cudaMemcpy(host,dev,N,D2H); consume_output(host); free(host); cudaFree(dev);

Standard dGPU-style

Static/Dynamic Transformation

produce_input(host); cudaMemcpy(dev,host,N,H2D);

kernel<<<G,B>>>(ptr); cudaDeviceSynchronize(); consume_output(ptr);

CPU

GPU cudaMemcpy(d,h,N,H2D)

malloc(N)

cudaMemcpy(h,d,N,D2H)

cudaMalloc(N)

DejaVu’s Static/Dynamic Transformation

CPU

Memory Copy Share (!!"#$ )

GPU

cudaFree(ptr); cudaMallocManaged(N)

Optimized version

(a)

Unified Memory

(b)

(c)

Figure 1. Overview of Echo. (a) & (b) Echo unifies separate host and device buffers into one managed allocation and removes redundant copies while preserving the required synchronization. (c) Speedup (bar) achieved by Echo on NVIDIA Jetson AGX Orin, overlaid with 𝑓𝑐𝑜𝑝𝑦 (line), the fraction of total runtime spent on memory copies before transformation. platforms. We therefore take CUDA on Jetson as our concrete study platform. However, the analysis and transformation we develop are not inherently platform-specific. Figure 1(a) contrasts the legacy dGPU-style pattern with what UMA-aware code should do instead. The cost of leaving this redundancy in place is not negligible. On a representative TensorRT [23] inference pipeline running FCN-ResNet50 [8, 12] on Jetson AGX Orin [19], we measure that over 30% of per-iteration wall-clock time is spent on H2D and D2H copies (see Table 1) that are avoidable on a UMA platform. Eliminating this redundancy requires solving two subproblems: (1) pattern detection, identifying a single logical buffer that has been split into a separate host and device allocation linked by copies; and (2) transformation, replacing the allocation pair with a single buffer shared by both host and device, and removing or bypassing the associated copies while preserving the synchronization they enforce. Despite how mechanical this may appear, neither step is a naive findand-replace. The transformation is safe only when no other part of the program relies on the host and device pointers being distinct, or on the copy enforcing an order between host and GPU accesses. Proving that automatically is hard, and harder still in production software, where the copy path may lie inside a precompiled library with no source to analyze. We present Echo, a system that automatically detects and eliminates redundant copies on UMA platforms. It combines source-level safety analysis with profile-guided specialization for binary deployments. Echo-SR (Static Replacement) couples an LLVM transformation with sufficient conditions for sharing storage while preserving data values, access order, and lifetime. Its correctness argument covers pairs satisfying these conditions in the supported execution model (§3.2). Echo-DR (Dynamic Replacement) applies allocation unification to application binaries without recompilation. It identifies recurring allocation/copy patterns, validates host accesses during GPU-use windows, and calibrates the complete-loop benefit before enabling a plan. Its deployment contract combines these observations with

assumptions about buffer values, pointer use, and synchronization (§3.3). We evaluate Echo on seven benchmarks and a negative control: three CUDA-library micro-benchmarks (cuBLAS [15], cuFFT [18], cuDNN [17]), two canonical TensorRT inference pipelines (FCN-ResNet50 [8, 12], FCN-ResNet101), and two application workloads (YOLOv5 detection [14] and DeepSORT feature extraction [24]). We run the evaluation across three NVIDIA Jetson platforms (Orin Nano, AGX Orin, and AGX Thor), covering two GPU architectures (Ampere and Blackwell) and varying memory bandwidths and compute capabilities. Figure 1(c) summarizes the representative result on AGX Orin. On Orin, copy-dominated micro-benchmarks achieve up to 7.05× speedup, and closed-source applications improve by 1.17–1.40×. Over the complete benchmark loop, Echo raises YOLOv5 throughput from 194.5 to 272.4 images/s and DeepSORT feature extraction from 94.9 to 106.8 frame jobs/s. On the five source-available Orin benchmarks, both automatic paths closely match the performance of manual optimization. The shorter primary timing regions also lower the corresponding estimated rail energy. In summary, we make the following contributions:

1. A deployment characterization of the discrete-GPU separatebuffer pattern in widely-used Jetson CUDA software, including its prevalence and its relevance to production inference stacks. 2. Echo-SR, a sound LLVM transformation that coalesces allocation pairs and eliminates the associated copies only after proving the transformation is safe. 3. Echo-DR, a profile-guided binary optimizer that brings the same allocation coalescing to unmodified, closedsource binaries by matching stable, API-visible allocation/copy patterns. 4. An evaluation across three Jetson platforms and seven benchmarks showing that speedup and energy savings scale with copy fraction, that binary-only applications

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

can benefit without recompilation, and that platformspecific managed-memory behavior defines a deployment boundary.

2

Background and Motivation

2.1

Unified Memory Architecture

Modern edge SoCs integrate CPU and GPU on the same die with a shared DRAM pool. Figure 1(b) contrasts this unified memory architecture (UMA) with the discrete-GPU model. On a discrete GPU, malloc allocates in host DRAM, cudaMalloc allocates in device DRAM, and cudaMemcpy transfers data across the PCIe bus. Therefore, the cost is inherent to the physical separation. On a UMA platform such as NVIDIA’s Jetson AGX Orin, both malloc and cudaMalloc allocate from the same 64 GB LPDDR5 DRAM. The two calls return pointers to different virtual-address regions, but both reside in the same DRAM pool. A cudaMemcpy between them copies data from one region to another within the same physical memory. This transfer moves no data across any bus, yet still consumes DMA bandwidth and CPU and GPU cycles. On UMA platforms, duplicate host buffers also consume the same physical memory pool available to the GPU [16]. When both host and device copies are resident, these duplicate allocations can increase the peak shared-memory footprint and reduce the capacity available for model parameters, intermediate tensors, and GPU workspaces. CUDA provides cudaMallocManaged, which returns a single pointer accessible by both CPU and GPU. On UMA platforms like Jetson, managed memory maps directly to shared DRAM without triggering page faults or data migration. However, most CUDA code in practice still follows the legacy malloc + cudaMalloc + cudaMemcpy pattern, as most code was originally written for or ported from discrete-GPU environments. Echo targets this pattern. 2.2

CUDA Unified Virtual Memory

CUDA provides a related but distinct mechanism called Unified Virtual Memory (UVM), which offers a single-pointer programming model for CPU–GPU memory. On discrete GPUs, UVM relies on demand paging with hardware page-fault handling, which introduces migration overhead proportional to access patterns [1, 5]. UVMBench [7] measures an average 34.2% kernel slowdown under UVM compared to explicit data management across their discrete-GPU benchmark suite, and UVMDiscard [29] identifies cases where UVM migrates data that is overwritten before being read. These overheads stem from physical page migration across a PCIe or NVLink bus. On integrated UMA platforms, there is no page migration because CPU and GPU share the same physical DRAM. The overhead targeted in our work is different: it arises from virtual-address duplication and explicit cudaMemcpy calls inherited from dGPU-style code, not from page faults or data migration. Consequently, managed-memory accesses avoid

Table 1. Per-stage latency breakdown for FCN-ResNet50 inference on Jetson AGX Orin (5 independent runs, mean ± std. dev.). Echo eliminates H2D and D2H copies, reducing iteration time by 31.8% (1.47× speedup). Stage

Baseline (ms)

Echo (ms)

Reduction

Preprocess H2D copy Inference D2H copy Postprocess

0.335 ± 0.001 0.295 ± 0.009 11.456 ± 0.015 5.248 ± 0.012 0.002 ± 0.000

0.326 ± 0.002 0.001 ± 0.000 11.479 ± 0.011 0.007 ± 0.000 0.002 ± 0.000

2.6% 99.8% −0.2% 99.9% 0%

Total

17.336 ± 0.010

11.815 ± 0.011

31.8%

Table 2. Prevalence of the dGPU-style allocation pattern across 34 open-source Jetson repositories. Copy-only: uses only explicit cudaMemcpy. Mixed: combines explicit copies with some managed or zero-copy paths. Project Domain (#)

Explicit copy

Candidate pattern

Copy only

Mixed

Inference (11) Robotics (13) Computer Vision (8) Signal Proc. (1) Other (1)

10 8 3 1 1

10 5 2 1 1

2 5 1 0 0

8 3 2 1 1

Total (34)

23

19

8

15

the interconnect migration cost that makes UVM expensive on discrete GPUs. On platforms that provide coherent shared-memory mappings, kernel accesses can be comparable to ordinary allocations, although caching and coherence policies remain platform-dependent. 2.3

Measured Opportunity

We profile FCN-ResNet50 [8, 12] inference on Jetson AGX Orin using Nsight Systems [20]. Table 1 breaks down perstage latency before and after applying Echo-DR to the unmodified binary. Copies account for 32% of baseline iteration time. Echo eliminates them almost entirely, reducing iteration time by 31.8% (1.47×). To establish that this result is close to the practical optimum, we also measured a hand-rewritten version that manually replaces the malloc+cudaMalloc pair with cudaMallocManaged and removes the redundant copies. Over the same five independent runs, this manually optimized variant achieves 11.786±0.008 ms, only 0.029 ms faster than Echo. Echo therefore achieves close to the practical optimum on this benchmark. 2.4

Prevalence of the Target Pattern

The malloc–cudaMalloc–cudaMemcpy pattern is not an artifact of our benchmarks. To assess its prevalence, we surveyed 34 open-source repositories designed for or deployed on Jetson platforms, spanning inference, robotics, computer vision,

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

and signal processing, collected from a variety of sources such as the community projects list and tutorials. For each, we performed a static source scan for the co-occurrence of malloc (or new), cudaMalloc, and cudaMemcpy with HostToDevice or DeviceToHost direction. Of the 34 repositories surveyed (Table 2), 19 (56%) contain the candidate malloc–cudaMalloc–cudaMemcpy pattern that Echo targets. Among inference-specific projects (a primary deployment scenario for edge AI), the prevalence rises to 10 of 11 (91%). Among the repositories with explicit copies, eight use them exclusively with no UMA-aware alternative, while the remaining fifteen combine explicit copies with some managed or zero-copy paths (See Copy-only vs. Mixed columns in Table 2). The static scan identifies syntactic co-occurrence of the pattern, not semantic optimizability. Whether a given pair can be unified requires either Echo’s source-level safety analysis or the guarded runtime profiling used by DR. Nonetheless, the survey confirms that the target pattern is widespread in practice.

3

Design of Echo

3.1

Overview

Echo automatically detects the dGPU-style allocation pattern on UMA platforms and coalesces each such pair into a single buffer, removing the copies that become redundant, but only when it can do so without changing the program’s behavior. It provides two approaches: SR, which proves this at compile time when source is available, and DR, which conservatively optimizes only guarded, profile-matched patterns at runtime for unmodified binaries.

Table 3. Notation for §3.2. Symbol

Meaning

ℎ ⊥𝑑 ℎ ⊥𝑑 →𝑝 𝑃, 𝑃 ′ [[𝑃 ]]

candidate pair: host buffer ℎ, device buffer 𝑑, both 𝑁 bytes unifying ℎ and 𝑑 into one managed buffer 𝑝 program before and after ℎ ⊥ 𝑑 → 𝑝 externally visible behavior of 𝑃 : I/O and exit status, excluding allocation, copy, and synchronization calls byte 𝑖 of ℎ and 𝑑, 0 ≤ 𝑖 < 𝑁 partner byte: ℎ[𝑖 ] = 𝑑 [𝑖 ], 𝑑 [𝑖 ] = ℎ[𝑖 ] events of matching executions of 𝑃 and 𝑃 ′ bytes read and written by event 𝑒 events that read or write a byte in 𝑋 happens-before in 𝑃 and in 𝑃 ′ linking copies: events that copy ℎ[𝑖 ] → 𝑑 [𝑖 ] or 𝑑 [𝑖 ] → ℎ[𝑖 ] events that restore the orderings provided by copy 𝑐 accesses SR attributes to ℎ or 𝑑

ℎ[𝑖 ], 𝑑 [𝑖 ] 𝑥¯ 𝐸, 𝐸 ′ rd(𝑒 ), wr(𝑒 ) Acc(𝑋 ) ≺, ≺ ′ L sync(𝑐 ) VisSR (ℎ, 𝑑 )

A linking copy moves bytes between ℎ and 𝑑 at the same offset, from ℎ[𝑖] to 𝑑 [𝑖] or back. L is the set of linking-copy events. Given a candidate pair, SR merges ℎ and 𝑑 into one managed buffer 𝑝 and deletes their linking copies while keeping any ordering those copies imposed. The following definition describes only this rewrite. §3.2.2 gives the conditions under which it is safe. Definition 2 (Unification ℎ ⊥ 𝑑 → 𝑝). Unifying each accepted candidate pair ℎ ⊥ 𝑑 in 𝑃 into its own managed buffer 𝑝 transforms 𝑃 into 𝑃 ′ by four rules.

When source code is available, SR decides at compile time whether a host/device allocation pair can share one managed buffer without changing the program’s behavior. We define the candidate pair and the transformation (§3.2.1), give sufficient conditions for its safety (§3.2.2), and show that they preserve observable behavior (§3.2.3). Table 3 summarizes the notation.

R1 Allocation. Replace malloc(N) and cudaMalloc(&d, N) with one cudaMallocManaged(&p, N). R2 Renaming. Replace every use of ℎ and 𝑑 with 𝑝, so byte 𝑖 of each buffer, ℎ[𝑖] and 𝑑 [𝑖], becomes 𝑝 [𝑖]. R3 Copies. Replace every linking copy 𝑐 ∈ L with sync(𝑐), a synchronization that accesses no memory and keeps each happens-before ordering 𝑎 ≺ 𝑐 ≺ 𝑏 of 𝑃 as 𝑎 ≺′ 𝑏 in 𝑃 ′ .1 R4 Deallocation. Replace free(h) and cudaFree(d) with one cudaFree(p) placed after both original frees and after all GPU work that accesses 𝑝 completes.

3.2.1 Transformation. SR starts by identifying candidate pairs in the program. A candidate pair is the dGPU-style idiom of Figure 1(a): ℎ is a host staging buffer for 𝑑, and a full-buffer H2D copy fills 𝑑 from ℎ, so byte ℎ[𝑖] corresponds to 𝑑 [𝑖]. Copies between corresponding bytes, which we call linking copies, exist only to keep the two buffers consistent. Once ℎ and 𝑑 share storage on UMA, they become redundant.

Matching executions. Let 𝐸 and 𝐸 ′ be the events of executions of 𝑃 and 𝑃 ′ on the same input and external choices. The retained events 𝐸ˆ are the events of 𝐸 other than linking copies and the allocations and frees of ℎ and 𝑑. Each retained event has a matching event in 𝐸 ′ , and we use the same name for both. The events of sync(𝑐) are new in 𝐸 ′ . R3 then guarantees

3.2 SR: Source-Level Provably Safe Unification

Definition 1 (Candidate pair). Buffers ℎ and 𝑑 form a candidate pair, written ℎ ⊥ 𝑑, if the original program 𝑃 contains the following calls, all with the same size 𝑁 : h = malloc(N); cudaMalloc(&d, N); cudaMemcpy(d, h, N, H2D);

// host buffer // device buffer // full-buffer H2D copy

(OP)

∀𝑎, 𝑏 ∈ 𝐸ˆ :

𝑎 ≺ 𝑏 ⇒ 𝑎 ≺′ 𝑏,

1 After R2 the data movement of 𝑐 is redundant, but the orderings it provides

are not, for example when 𝑐 blocks the host or orders work across streams. For a synchronous D2H copy, sync(𝑐 ) can be cudaDeviceSynchronize(). It is empty when the remaining events already imply every 𝑎 ≺ 𝑐 ≺ 𝑏.

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

1 void case_S1(int N) { float* h = (float*)malloc(N); 2 float* d; cudaMalloc(&d, N); 3 fill(h, data, N); 4 cudaMemcpy(d, h, N, H2D); 5 h[0] = 999.0f; 6 kernel<<<G,B>>>(d, N); 7 cudaMemcpy(h, d, N, D2H); 8 free(h); cudaFree(d); 9 10 }

S1: live device value overwritten

1 void case _S2(int N) { 2 float* h = (float*)malloc(N); 3 float* d; cudaMalloc(&d, N); 4 fill(h, data, N); 5 cudaMemcpy(d, h, N, H2D); 6 kernel<<<G,B>>>(d, N); 7 cudaMemcpy(h, d, N, D2H); 8 use(h[0]); 9 free(h); cudaFree(d); 10 }

S2: preserve D2H completion

1 void case _S3(int N) { float* h = (float*)malloc(N); 2 float* d; cudaMalloc(&d, N); 3 fill(h, data, N); 4 if (h == d) reportSame(); 5 cudaMemcpy(d, h, N, H2D); 6 kernel<<<G,B>>>(d, N); 7 cudaMemcpy(h, d, N, D2H); 8 free(h); cudaFree(d); 9 10 }

S3: pointer identity is observable

1 void case _S4(int N) { float* h = (float*)malloc(N); 2 float* d; cudaMalloc(&d, N); 3 fill(h, data, N); 4 cudaMemcpy(d, h, N, H2D); 5 free(h); 6 kernel<<<G,B>>>(d, N); 7 cudaDeviceSynchronize(); 8 cudaFree(d); 9 10 }

S4: keep p alive through GPU use

Figure 2. Four examples motivating SR’s conditions for replacing host allocation ℎ and device allocation 𝑑 with one managed allocation 𝑝. S1: after H2D, the kernel still needs the value in 𝑑; writing ℎ[0] would overwrite it in 𝑝. S2: the synchronous D2H copy completes the relevant GPU work before use(h[0]), so removing it requires an equivalent completion wait. S3: unification changes the observable test h == d from false to true. S4: although the original program frees ℎ before the kernel uses 𝑑, 𝑝 must remain live until both original lifetimes and the GPU work end. that is, every ordering among retained events of 𝑃 survives in 𝑃 ′ . Allocations and frees are excluded because R1 and R4 move them, and S4 constrains them instead. 3.2.2 Safety Conditions. Unification is not always safe, because the program may see different values, orderings, pointer identities, or lifetimes once ℎ and 𝑑 share storage (Figure 2). SR therefore checks each candidate pair against the conditions of Definition 3 and unifies only the pairs that satisfy all of them. S1 relies on liveness, which we define first. Liveness. Let rd(𝑒) and wr(𝑒) be the bytes that an event 𝑒 reads and writes. A byte 𝑥 is live after an event 𝑒 ∈ 𝐸 if another event 𝑟 can still read the value 𝑥 holds after 𝑒, that is, 𝑟 reads 𝑥, 𝑟 is not ordered before 𝑒, and no write to 𝑥 is ordered between 𝑒 and 𝑟 . Formally, live(𝑒, 𝑥) ≜ ∃𝑟 ≠ 𝑒. 𝑥 ∈ rd(𝑟 ) ∧ 𝑟 ̸≺ 𝑒 ∧ ¬∃𝑤 . 𝑥 ∈ wr(𝑤) ∧ 𝑒 ≺ 𝑤 ≺ 𝑟 . Because 𝑟 ̸≺ 𝑒 also admits reads unordered with 𝑒, which may run after 𝑒, such reads keep 𝑥 live too. Definition 3 (Safe unification). We call ℎ ⊥ 𝑑 → 𝑝 safe iff S1–S4 hold for every execution of 𝑃 and its matching execution of 𝑃 ′ . S1 Value. In 𝑃 ′ , a write event 𝑒 to a byte 𝑥 of ℎ or 𝑑 also ¯ the byte at the same offset in the overwrites its partner 𝑥, other buffer, so the old value of 𝑥¯ must never be read after 𝑒.

be ordered in 𝑃 ′ . Since R3 keeps every ordering of 𝑃 in 𝑃 ′ , it suffices that 𝑎 and 𝑏 are ordered in 𝑃. For example, a host read of ℎ[𝑖] that overlaps a kernel writing 𝑑 [𝑖] is harmless in 𝑃 but becomes a race in 𝑃 ′ . S3 Pointer. In 𝑃 ′ , ℎ and 𝑑 become the same pointer 𝑝, so the program must not be able to tell them apart. Let Uses(ℎ) be the operations that use ℎ or a pointer derived from it. Uses(ℎ) ∪ Uses(𝑑) ⊆ {load, store, copy, kernel arg, free} Thus ℎ and 𝑑 serve only as addresses and are never compared (as in h == d in Figure 2), cast to integers, or allowed to escape to memory. In addition, no remaining operation may assume that ℎ and 𝑑 do not overlap, such as a copy between them at different offsets. S4 Lifetime. The single buffer 𝑝 replaces two buffers with different lifetimes, so it must stay allocated from the first access to either buffer until the last. With Acc(𝑋 ) denoting the events that access a byte of 𝑋 , ∀𝑎 ∈ Acc(ℎ ∪ 𝑑) ∩ 𝐸ˆ : alloc(𝑝) ≺′ 𝑎 ≺′ free(𝑝) In the S4 example of Figure 2, ℎ is freed before the kernel uses 𝑑, so 𝑝 must outlive both. S4 fails when R1 and R4 cannot place alloc(p) and free(p) this way, for example when one buffer is allocated inside a conditional branch that does not dominate all its accesses. 3.2.3

Soundness.

¯ ∧ 𝑥¯ ∉ rd(𝑒) ¬ live(𝑒, 𝑥)

Theorem 1 (Soundness of unification). Let 𝑃 be race-free, let every allocation succeed, and let no event of 𝑃 other than a linking copy read an uninitialized byte of ℎ or 𝑑. If ℎ ⊥ 𝑑 → 𝑝 is safe, then [[𝑃 ′ ]] ⊆ [[𝑃]], where [[𝑃]] is the externally visible behavior of 𝑃.

¯ rules out any other read that could still Here, ¬live(𝑒, 𝑥) ¯ and 𝑥¯ ∉ rd(𝑒) rules out 𝑒 itself see the old value of 𝑥, ¯ which in 𝑃 ′ could return 𝑒’s own write. S1 reading 𝑥, skips linking copies, because after such a copy 𝑥 and 𝑥¯ hold the same value, so merging them loses nothing. S2 Order. An access to ℎ[𝑖] and an access to 𝑑 [𝑖] touch the same byte 𝑝 [𝑖] in 𝑃 ′ , so if either one writes, they must

Proof sketch. Fix an execution of 𝑃 ′ . We build a matching execution of 𝑃 that follows the same order and reinserts the deleted linking copies, and show that every read returns the same value in both. Two conflicting accesses in 𝑃 ′ either conflict in 𝑃, where race freedom orders them, or touch ℎ[𝑖] and 𝑑 [𝑖], where S2 orders them. (OP) in §3.2.1 carries both orders to 𝑃 ′ , so 𝑃 ′ is race-free and its reads do not depend

ˆ ∀𝑥 ∈ wr(𝑒) ∩ (ℎ ∪ 𝑑) : ∀𝑒 ∈ 𝐸,

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

on how unordered events interleave. Along this order, whenever ℎ[𝑖] or 𝑑 [𝑖] will still be read in 𝑃, 𝑝 [𝑖] holds its value in 𝑃 ′ . A write to 𝑥 preserves this because, by S1, the value of 𝑥¯ it overwrites is dead. A linking copy, which runs only in 𝑃, preserves it because its source value is already in 𝑝 [𝑖]. S3 makes the shared address unobservable, and S4 keeps 𝑝 allocated for every access. Every read therefore returns the same value in 𝑃 and 𝑃 ′ , so both programs take the same branches and make the same external calls, which gives [[𝑃 ′ ]] ⊆ [[𝑃]]. The full proof is in the supplementary appendix. □ 3.2.4 Static Analysis and Rewriting. Theorem 1 requires S1–S4 over all accesses to ℎ and 𝑑, but SR can check only the accesses its analysis attributes to them, VisSR (ℎ, 𝑑). Its checks therefore imply S1–S4 only under the coverage premise (C)

Acc(ℎ ∪ 𝑑) ⊆ VisSR (ℎ, 𝑑).

SR secures (C) conservatively. S3 rules out pointer uses that the analysis cannot follow, and SR declines a pair whose pointers flow across compilation units. SR runs these checks and the rewrite as a pass on LLVM IR. For each candidate pair, it checks S1–S4 on the accesses in VisSR (ℎ, 𝑑) and declines the pair whenever a condition cannot be decided statically, for example because of a dynamic size or complex control flow. It rewrites each accepted pair by R1–R4 (Definition 2), so Theorem 1 applies to every pair that SR unifies. SR also skips pairs smaller than a threshold 𝜃 min , because their copies cost too little to matter. 3.3 DR: Binary-Level Pattern-Guided Optimization In many production deployments, source code is unavailable or the critical copy paths are buried inside precompiled libraries such as TensorRT. DR applies the unification ℎ ⊥ 𝑑 → 𝑝 at runtime by interposing on libc and CUDA runtime calls in the unmodified binary. We describe which safety conditions DR checks and which it assumes (§3.3.1), then its three-phase pipeline (§3.3.2). 3.3.1 Checked Conditions and Assumptions. DR applies the unification ℎ ⊥ 𝑑 → 𝑝 (Definition 2) at runtime. However, unlike SR, DR cannot establish S1–S4 for every execution. It sees API calls but not memory loads and stores. DR therefore checks what it can observe and assumes the rest. Phase 2 checks, on profiled executions, that the host does not touch ℎ while the GPU uses the pair (§3.3.2), and Phase 3 keeps 𝑝 alive until both original frees (S4). DR assumes the other conditions, for example that the GPU never overwrites a host value that is still needed (S1) and that the program never compares ℎ and 𝑑 (S3). Its guarantee therefore holds only for later runs that behave like the profiled ones. 3.3.2 Profile-Guided Unification. DR selects the pairs to unify by profiling the unmodified binary, and unifies them in later runs (Figure 3). Phases 1 and 2 run once to produce

a unification plan, and Phase 3 applies this plan to each later run. Phase 1: Profile. Using LD_PRELOAD, DR intercepts libc and CUDA runtime APIs for allocation, copy, synchronization, and deallocation. For each allocation or copy, it records the size 𝑁 and a call-site identity 𝛽, a 64-bit FNV-1a hash of the top 𝑑 return addresses on the call stack (e.g., 𝑑=16). We write 𝛽ℎ for the host allocation site, 𝛽𝑑 for the device allocation site, and 𝛽𝑐 for the copy site. Multiple stack frames distinguish sites behind common wrapper functions (§4.1.2). A size threshold 𝜃 min (e.g., 200 KB) excludes small buffers. This run keeps the original allocations and copies. The runtime writes all recorded events to a trace file, which Phase 2 consumes as input. Phase 2: Analyze. Phase 2 turns the trace into a unification plan. It selects the pairs to unify, records how to keep the orderings their copies provide, rejects pairs that the host touches while the GPU uses them, and enables the plan only if it reduces latency. Pairing. The analyzer groups copies by (𝛽ℎ , 𝛽𝑑 , 𝑁 ) and selects pairs that recur at least 𝑘 min times (e.g., 𝑘 min =2) at the dominant steady-state copy site. Each selected pair ℎ ⊥ 𝑑 is linked by copies in L. Orderings. A synchronous D2H copy does two things. It moves data, and it makes the host wait until the GPU work that produces the data is done. Skipping the copy removes the first but must keep the second, which is the role of sync(𝑐) in R3. By default, Phase 3 replaces the copy with cudaDeviceSynchronize. If the program already waits for that GPU work before the copy, for example through an earlier synchronization with no kernel launched in between, the plan records it and Phase 3 adds no wait. In both cases the host reads the result only after the GPU has produced it, as (OP) in §3.2.1 requires. Validation. A separate run keeps ℎ and 𝑑 distinct and checks that the host does not touch ℎ while the GPU uses the pair, which is the host side of S1 and S2. SR discharges these conditions statically, but DR sees only API calls, not memory loads and stores. It therefore checks them at runtime with a page-protection trick delimited by API calls. The API calls mark the interval to check, which we call the pair’s GPU-use window, and page protection catches any host access inside it. A window normally opens at the pair’s H2D copy and closes at its D2H copy. Some pairs have only one of them, for example an input buffer that the GPU only reads or an output buffer that the GPU only writes. For these pairs, DR takes the missing boundary from the round of GPU work they belong to, which starts with an H2D copy and ends with a D2H copy or a synchronization. An asynchronous D2H copy closes a window only when a later synchronization completes it. While a window is open, DR marks the pages of ℎ inaccessible with PROT_NONE, so any host access to them

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

fill(h, input, N); cudaMemcpy(d, h, N, H2D);

// ③ H2D copy

kernel<<<G,B>>>(d, N);

// ④ GPU compute

cudaMemcpy(h, d, N, D2H); process(h, N);

// ⑤ D2H copy

free(h); cudaFree(d);

// ⑥ free pair

// trace.jsonl {malloc, β=0x3F7A, ptr=0xA000, n=3MB} {cudaMalloc, β=0x8C21, ptr=0xB000, n=3MB}

1 2 3 4 5 6 7 8 9 10 11 12 13 14

{cudaMemcpy, H2D, β_src=0x3F7A, β_dst=0x8C21, n=3MB}

{cudaMemcpy, D2H, β_src=0x8C21, β_dst=0x3F7A, n=3MB} {free, β=0x3F7A} {cudaFree, β=0x8C21}

}

LD_PRELOAD

libc

Intercept

CUDA Runtime APIs Allocations:

malloc, free, calloc , realloc , cudaMalloc, cudaFree

Call_site hash (β) 0x7fff00 0x800000 0x800100 0x800200 0x800e00 d-frame return addr

Group by (βsrc, βdst, n)

FNV-1a

Candidate Pairing

64-bit hash β

• Match allocation/copy sites and sizes • Include H2D and outputonly pairs

…

Copying: cudaMemcpy, cudaMemcpyAsync

Synchronization: cudaDeviceSynchronize, cudaStreamSynchronize

Host-access validation

• Separate run with original buffers • Reject pairs with protectedwindow accesses

Profitability calibration Deployment plan

Key Invariant

Record

copy-skip uses pointer identity (dst == src ∧ isUvmPtr(dst)) only, unmatched calls unchanged

JSONL event per API call Fields: API, β, ptr, size

// optimized behavior cudaMallocManaged(N) -> p return p;

// ① [REPLACE] // ② [REPLACE]

fill(h, input, N); [NO-OP]

// ③ [SKIP]

kernel<<<G,B>>>(d, N);

// ④ GPU compute

[NO-OP] + cudaDeviceSync() process(p);

// ⑤ [SKIP]

free(h); cudaFree(p);}

// ⑥ [DEFER]

Phase 3: Optimize

Phase 2: Analyze

Phase 1: Profile

1 2 3 4 5 6 7 8 9 10 11 12 13 14

Plan enabled and context matches?

Share allocation

Skip copy

• First allocation creates managed buffer • Matching allocation returns the same pointer

• src == dst and owned managed pointer • Preserve D2H completion

ONE-TIME

Release allocation • Record both original frees • Release after both have occurred

Unmatched API

void infer(const float* input, int N) { float* h = (float*)malloc(N); // ① host buffer float* d; cudaMalloc(&d, N); // ② device buffer

Original API

1 2 3 4 5 6 7 8 9 10 11 12 13 14 15

EVERY RUN

Figure 3. DR’s three-phase pipeline. Profile discovers recurring allocation and copy events; Analyze forms candidate pairs, validates their host-access windows in a separate run, and calibrates the resulting plan; Optimize applies the enabled plan in its calibrated deployment context. Preparation includes discovery, analysis, page validation, and three baseline/optimized calibration pairs. traps. A trap means that the host touched ℎ while ℎ and 𝑑 could hold different values, which may violate S1 or S2, so DR rejects the pair. An unsupported or incomplete window also rejects the pair. Page protection is enforced by hardware, so it also catches accesses through aliases of ℎ. This is the counterpart of (C) for host accesses inside a window. Calibration. A safe plan is not always faster, because the unified buffer and any added synchronization can cost more than the copies they remove. DR therefore performs three trials, each with one run without the plan and one run with it in separate processes, and enables the plan only if the median reduction in complete-loop latency is positive. This is a profitability check, not a safety condition. DR also records the configuration of these runs, which Phase 3 checks before applying the plan. Phase 3: Optimize. Before any unification, DR compares the executable and interposer hashes, hostname, arguments, GPU, and CUDA runtime version with the calibration context. A disabled plan or a mismatch runs 𝑃 unchanged. Otherwise the runtime applies R1–R4 (Definition 2) to each selected pair. Allocation replacement (R1, R2). The first allocation that matches (𝛽, 𝑁 ) creates 𝑝 with cudaMallocManaged, and the partner allocation returns the same 𝑝, in either order. Every use of ℎ and 𝑑 then refers to 𝑝 without rewriting code. Copy skipping (R3). When cudaMemcpy(dst, src, N) is intercepted, DR checks if dst == src == p. If so, the copy is a self-copy and is skipped. A skipped copy that provides an

ordering is replaced by sync(𝑐) (e.g., cudaDeviceSynchronize) from Phase 2. Deallocation (R4). DR records free(h) and cudaFree(d) separately and releases 𝑝 after both, in either order. This keeps the unified buffer valid throughout both original lifetimes. Any API call whose call-site hash and size do not match a planned triple falls back to its original APIs.

4

Evaluation

4.1

Experimental Setup

4.1.1 Platform. Our primary platform is the Jetson AGX Orin Developer Kit (Ampere GPU, SM 8.7, 64 GB LPDDR5, CUDA 12.6, TensorRT 10.3 [23]). We also evaluate on two other Jetson platforms for cross-platform validation: Orin Nano Super (Ampere, SM 8.7, 8 GB, CUDA 12.6) and AGX Thor Developer Kit (Blackwell, SM 11.0, 128 GB, CUDA 13.0, pre-release JetPack 7.0-b128). All experiments use the maximum performance mode available on each platform (nvpmodel -m 0 on Orin and Thor; -m 2 MAXN_SUPER on Nano) with jetson_clocks and 60-second thermal stabilization before measurement. 4.1.2 Implementation. SR is implemented as an LLVM pass plugin. DR is a single shared library injected via LD_ PRELOAD, which forwards each intercepted call to the real implementation via dlsym(...). However, two platformspecific issues required workarounds on Jetson. First, glibc’s

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

Table 4. Benchmark suite. The Pairs column counts candidate allocation rules on Orin before profitability selection. Per-platform 𝑓copy values appear in Figure 7 and Appendix Table 11.

Baseline

Workload

Category

Pairs

Opt.

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2 𝐵 NC

cuBLAS SGEMM (20482 ) cuFFT OFDM (2048 × 1024) cuDNN Conv (ResNet-50 layer 1) FCN-ResNet50 inference FCN-ResNet101 inference YOLOv5 object detection DeepSORT feature extraction 𝐵 C1 with manually optimized

micro micro micro canonical canonical application application negative

2 2 2 2 2 2 2 0

SR & DR SR & DR SR & DR SR & DR SR & DR DR DR DR

4.1.3 Benchmarks. Table 4 lists the seven benchmarks and the negative control, organized into three tiers. • Micro-benchmarks (𝐵 M1 –𝐵 M3 ) isolate a single CUDA library kernel (cuBLAS, cuFFT, cuDNN) with controlled buffer sizes, following the allocation pattern used by NVIDIA’s TensorRT BufferManager class [22]. • Canonical pipelines (𝐵 C1 –𝐵 C2 ) run TensorRT inference (FCN-ResNet50 and FCN-ResNet101) with preprocessing, H2D, inference, D2H, and postprocessing stages. • Application workloads (𝐵 E1 –𝐵 E2 ) use third-party YOLOv5TensorRT [14] detection and DeepSORT-TensorRT [24] feature-extraction code. Although their source code is public, we emulate a closed-source deployment by running the prebuilt binaries distributed by each project as-is. We do not evaluate SR on these two workloads because we treat them as binary-only. We also include a negative control (𝐵 NC ) that already uses cudaMallocManaged to verify that DR finds no candidate pairs when no redundant copies exist. 4.1.4 Methodology. We compare up to four methods, depending on the benchmark. (a) The baseline is the original unmodified program using the standard malloc+cudaMalloc+ cudaMemcpy pattern. (b) DR and (c) SR are the two Echo approaches under evaluation. For the source-available benchmarks (𝐵 M1 –𝐵 M3 , 𝐵 C1 –𝐵 C2 ), we also include a (d) manually optimized variant as a performance upper bound for what the automatic transformation can achieve. Each method is measured over 5 independent process invocations (separate OS-level program launches, not repeated calls within a single process), each executing 100 warmup iterations followed by 100 timed iterations. Each process reports its mean latency over the timed iterations. We compute

Manually Optimized

B M

B BM2 DS M 0

6

7

1

8

2

9

3

4

5

B BM3 DS M 1

2

3

4

5

B BC1 DS M 11

backtrace() internally calls malloc, causing infinite recursion inside the malloc interceptor. Hence, we use a manual frame-pointer walk instead. Second, the first cudaMallocManaged call triggers lazy CUDA initialization, which reenters the intercepted allocation path. A thread-local depth counter prevents deadlock.

SR

BM1 DS 5

ID

DR

12

13

14

15

16

17

18

B BC2 DS M 19

20

21

22

23

24

25

26

BE1 BD 14

16

18

20

22

BE2 BD 5.5

6.0

6.5

7.0

7.5

8.0

8.5

9.0

Iteration latency (ms) Note: x-axis scale differs per panel.

Figure 4. Primary-region latency on Jetson AGX Orin. Box: mean±1𝜎. Whisker: min to p99. Red line: mean. Appendix Table 6 reports the mean of the five per-run p99 values. Each benchmark has its own x-axis scale. 𝐵 E1 and 𝐵 E2 show Baseline and DR as source code is unavailable.

the mean and standard deviation across the 5 per-process means. Some CUDA library kernels are not guaranteed to be bit-exact across process launches because implementation choices and parallel reduction orders may vary. We therefore compare each benchmark’s reported scalar L2 checksum with its same-run baseline using a relative tolerance of 10−5 . All selected runs across methods and platforms pass this smoke check. It is not an element-wise proof of arbitrary outputs. 4.1.5 Metrics. The primary timing region is the latency scope each benchmark uses to compare methods. It includes the benchmark’s host-device transfers, GPU computation, synchronization, and host-processing work, measured either as one wall-clock interval or as the sum of timed phases. The same scope is used for every method within a benchmark. Let 𝑡 method denote its mean latency. We use three metrics throughout the evaluation: • Speedup = 𝑡 baseline /𝑡 method , which represents the primaryregion performance improvement achieved by Echo’s transformation (and the manual optimization).

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

4.2

Per-Benchmark Results

Figure 4 shows primary-region latency on AGX Orin for all seven benchmarks. Full numerical results in Appendix Table 6. The main trend is that Echo improves performance by removing the copy time that it can safely observe and optimize. Therefore, workloads with copy-dominated baselines see the largest speedups, while workloads with small copy fraction see smaller but still useful gains. Micro-benchmarks (𝐵 M1 –𝐵 M3 ). These benchmarks isolate the impact of copy elimination on individual CUDA kernels. 𝐵 M2 (cuFFT) achieves a 7.05× speedup because its baseline primary region is heavily dominated by data movement (𝑓copy = 85.3%). Eliminating redundant copies only leaves the FFT computation and synchronization overhead. As the compute workload increases in 𝐵 M3 (cuDNN) and 𝐵 M1 (cuBLAS), the speedups naturally scale down to 3.37× and 1.69×, closely tracking the available copy fraction. Across all three, both DR and SR match the performance of handoptimized code to within measurement noise (≥99% recovery). Canonical inference pipelines (𝐵 C1 –𝐵 C2 ). In these multistage inference pipelines, copies account for a smaller fraction of total runtime, so the speedup is more modest: 𝐵 C1 (FCN-ResNet50) yields a 1.47× speedup, whereas the deeper 𝐵 C2 (FCN-ResNet101) achieves 1.28×. Because ResNet101 requires more GPU execution time per frame, its data-transfer overhead accounts for a proportionally smaller slice of the iteration (21% vs. 30%). Importantly, Echo scales to these multi-stage pipelines without degrading performance relative to manual optimization. The latency gap between Echo and the manual baseline is less than 0.5%. End-to-end applications (𝐵 E1 –𝐵 E2 ). These benchmarks show the benefit of DR on application binaries without recompilation. 𝐵 E1 (YOLOv5) achieves a 1.40× primary-region speedup on Orin. DR unifies two API-visible host/device allocation pairs, covering the input H2D and output D2H copies. In the profiled iteration, GPU copy time falls by 5.91 ms (Figure 5). The resulting latency recovers 98.4% of the gain achieved by

30%

30%

58%

60%

44% 44%

28% 28%

74%

74%

62%

62%

9% 22%

22%

32%

54%

20%

54%

40%

58%

65%

60%

Overlapped Other

20%

29%

37%

39%

80%

13%

100%

GPU kernel only GPU copy only

28%

Baseline (L) DR (R)

Profiled wall time (% of Baseline)

• Copy fraction (𝑓copy ) is baseline GPU copy-only time divided by baseline complete-iteration wall time, both measured in the same Nsight Systems [20] capture. It excludes CPU-only portions of copy calls and intervals where a copy overlaps a kernel. The Amdahl approximation 𝑆 = 1/(1 − 𝑓copy ) provides an upper bound reference on the speedup achievable by copy elimination. • Recovered fraction = (𝑡 baseline − 𝑡 Echo )/(𝑡 baseline − 𝑡 manual ), where 𝑡 Echo is the latency of either DR or SR, and 𝑡 manual is the latency of the manually optimized variant. This measures how close Echo’s automatic transformation comes to the manual upper bound.

0% BM1 BM2 BM3 BC1 BC2 BE1 BE2 0.59× 0.38× 0.31× 0.68× 0.78× 0.72× 0.89×

Figure 5. Complete-iteration wall time on AGX Orin, divided into GPU kernel-only, GPU copy-only, overlapped, and Other intervals. Other covers time with no GPU kernel or copy running. In each pair, the left bar is Baseline and the hatched right bar is DR, both normalized to Baseline’s total. The multiplier below each pair is DR’s total relative to Baseline (lower is better). DR removes the copy-only intervals and leaves kernel time unchanged, so the remaining kernel and Other intervals limit the gain.

an independent manual implementation of the same transformation. 𝐵 E2 (DeepSORT feature extraction) achieves a 1.17× speedup on Orin, close to the 1.16× copy-removal reference implied by its 13.7% baseline copy fraction. Copies account for a smaller share of this workload than 𝐵 E1 , so eliminating them leaves most of its execution time intact. 4.2.1 Execution-Time Breakdown. The primary-region speedups above show how much Echo reduces latency, but not where the saved time comes from. To answer this, Figure 5 divides the complete-iteration wall time of Baseline and DR into GPU kernel-only, GPU copy-only, overlapping kernel-and-copy, and Other intervals. This breakdown lets us check that the gain comes from the removed copies rather than from changes in kernel execution, and shows which remaining portion limits the gain over the complete iteration. Each bar averages 20 iterations from one Nsight Systems capture per condition, and Baseline is normalized to 100% within each benchmark. Across the Orin workloads, DR removes the GPU copyonly intervals while leaving kernel time nearly unchanged. The remaining intervals therefore determine the gain over the complete iteration. The copy-heavy 𝐵 M3 and 𝐵 M2 gain the most, whereas kernel time limits 𝐵 C1 and 𝐵 C2 , and Other time limits 𝐵 E1 and 𝐵 E2 . Other time also explains why 𝐵 M2 ’s 7.05× primary-region speedup shrinks over the complete

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

Primary-region speedup (×)

Orin

Nano

4.3

Thor

8

6

4

2 *

*

0 BM1

BM2

BM3

BC1

BC2

BE1

BE2

Figure 6. DR speedup across platforms and benchmarks. ∗ denotes 𝐵 M3 and 𝐵 E2 use the baseline fallback selected by the profitability guard.

8

Amdahl S = 1/(1

Observed speedup

7

fcopy)

Orin Nano

6

Thor

5 4 3 2 1

Beyond primary-region latency, Echo also improves the complete application loop. Table 5 reports throughput (𝑄), measured as completed work units per second (e.g., objectdetection images in 𝐵 E1 and segmented images in 𝐵 C2 ) over the complete timed loop. DR raises throughput on every Orin workload. For example, 𝐵 M2 ’s FFT batches increase by 2.64× and 𝐵 C1 ’s images by 1.47×. For the application workloads, YOLOv5 rises from 194.47 to 272.38 images/s, while DeepSORT feature extraction rises from 94.85 to 106.81 frame jobs/s. Echo also saves memory, both in buffer footprint and in copy traffic. To quantify these savings, we record two quantities. The first is the selected buffer capacity (𝑀), which is the total size of the live host and device allocations that DR unifies. The second is the copy payload that DR skips in each complete iteration, which we obtain from runtime counters. As Table 5 shows, 𝑀 is halved in every row, reflecting the removal of duplicate host/device storage. Across the seven workloads, the baseline copies 12.25–51.43 MiB per complete iteration between the unified buffers, including 32 MiB for 𝐵 M2 , 45 MiB for 𝐵 C1 , and 51.43 MiB for 𝐵 E1 , and the runtime counters confirm that DR skips all of these copies. 𝐵 E1 processes four images per iteration, corresponding to 12.86 MiB of avoided copy payload per image. For example, a 20 FPS (frames per second) camera feed with the same engine and tensor shapes would avoid approximately 257 MiB/s of copy payload.

M3* fallback

E2* fallback

0

Throughput and Memory Savings

20

40

60

80

100

Wall-clock copy fraction f copy (%)

Figure 7. Observed DR speedup vs. copy fraction (𝑓copy ) across 21 data points on three platforms. The dashed curve is the Amdahl approximation 𝑆 = 1/(1 − 𝑓copy ).

iteration, because most of its remaining time is spent in Other intervals. 4.2.2 Cross-Platform Validation. Figure 6 shows that the same transformation produces different speedups across platforms. Differences in hardware and software stacks change how each workload’s time is divided among copies, GPU computation, and other work. Consequently, the portion affected by copy elimination varies. Echo accelerates only the memory copy portion of the pipeline, so its speedup is bounded by Amdahl’s law: 𝑆 Amdahl = 1/(1 − 𝑓copy ). Figure 7 plots 𝑓copy against observed speedup for all benchmarks across three platforms (21 data points), with this curve overlaid. The strong fit (𝑅 2 = 0.989) shows that copy elimination, not benchmark identity, explains most of the observed benefit. The abnormal behavior of 𝐵 M3 and 𝐵 E2 on Thor is discussed in Appendix §D.

4.4

Practical Costs and Energy Consumption

DR’s one-time cost is the profile run used to build the optimization plan. Profiling adds less than 1% per-iteration overhead, so the dominant cost is executing 50–100 baseline iterations before optimized runs. For 𝐵 C1 at a 30 FPS operating point, profiling 50 iterations takes about 1.67 s; the 5.52 ms per-iteration latency saving recovers that cost within 156 optimized iterations, or about 5.2 s of operation. Figure 8 shows how Echo affects energy consumption by comparing the primary-region energy of DR with that of the baseline. These memory savings are secondary to latency in our evaluation, but they matter on memory-constrained edge devices. For edge deployments running on battery or solar power, energy consumption is as important as latency, if not more so. Hence, we estimate the energy of each run by multiplying the mean rail power over the timed loop by that run’s primary-region latency, i.e., 𝐸 (mJ) = 𝑃mean (mW) × 𝑡 (ms)/1000. The figure reports the mean and sample standard deviation across five runs. On Orin, estimated energy consumption falls by 3.5–83.2% across the seven benchmarks. On most benchmarks, we found that instantaneous GPU power increases under Echo because the GPU no longer idles during DMA copies, but total energy per iteration still decreases because iterations complete faster.

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

Table 5. Complete-loop throughput and selected buffer capacity on AGX Orin. 𝑄 is mean ± sample SD across five independent runs; 𝑀 counts the live allocations selected for unification. Subscripts B and D denote Baseline and DR. Benchmark

𝑄 𝐵 (work units/s)

𝑄 𝐷 (work units/s)

𝑄 𝐷 /𝑄 𝐵

𝑀𝐵 (MiB)

𝑀𝐷 (MiB)

109.21 ± 0.15 164.53 ± 1.04 199.26 ± 1.03 57.68 ± 0.03 39.86 ± 0.02 194.47 ± 0.63 94.85 ± 0.37

184.65 ± 0.12 434.24 ± 0.29 664.97 ± 2.44 84.63 ± 0.08 51.13 ± 0.02 272.38 ± 0.30 106.81 ± 0.31

1.691× 2.639× 3.337× 1.467× 1.283× 1.401× 1.126×

64.000 64.000 58.188 90.000 90.000 102.869 24.500

32.000 32.000 29.094 45.000 45.000 51.434 12.250

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

Note: Work units are SGEMM jobs for 𝐵 M1 , forward-and-inverse FFT batches for 𝐵 M2 , convolution batches for 𝐵 M3 , images for 𝐵 C1 , 𝐵 C2 , and 𝐵 E1 , and feature-extraction frame jobs for 𝐵 E2 . 𝐵 E1 processes four images per iteration.

5

Orin Energy / iter. (mJ)

1000

-3%

Echo’s two optimization paths have different envelopes: SR proves source-available transformations, while DR provides a guarded contract for stable binary deployments. We summarize the remaining boundaries here.

800 -7%

600 400

-19%

-11%

200 -83%

-12%

-41%

0

Nano -4%

Energy / iter. (mJ)

600

-4%

400

200

-7%

-12% -76%

-8%

-26%

0

Thor -6%

Energy / iter. (mJ)

500 400 -11%

300 200

-8% -5%

100 -46%

0 M1

M2 Baseline

-2%

M3*

Limitations and Future Work

+8%

C1 DR

C2

E1

E2*

Binary deployment contract. DR provides profile-conditional correctness, not SR’s input-universal proof. DR protects the pages of candidate host buffers during host-access windows in which a CPU read or write would violate the allocation unification semantics. If the profiled execution touches those pages, the corresponding pair is rejected from the optimization plan. This strengthens validation beyond API-visible traces, but the remaining limitation is profile coverage. An input encountered later in deployment could trigger unobserved host accesses, allocation pairings, copy orderings, or stream pattern that profiling never witnessed. Page protection checks candidate buffers for unsafe CPU accesses during validation, but is not active in later optimized runs. A new, unseen input could therefore cause a CPU read or write to a unified buffer that DR does not notice, even if its CUDA calls still match the plan. Startup checks and API matching do not detect this kind of change.

SR

Figure 8. Estimated primary-region energy (mJ). Signed percentages show (𝐸 DR − 𝐸 Baseline )/𝐸 Baseline ; negative values indicate lower energy. Each panel has its own y-axis scale. ∗ mark DR’s baseline fallback on Thor 𝐵 M3 /𝐵 E2 , where differences reflect variation between baseline runs. 𝐵 E1 –𝐵 E2 evaluate the binary deployment path and show Baseline and DR.

API coverage. Echo-DR intercepts the CUDA runtime API (e.g., cudaMalloc, cudaMemcpy, cudaMemcpyAsync). Applications relying on the CUDA driver API (e.g., cuMemAlloc, cuMemcpyDtoH), such as DeepStream [21] and trtexec [23], are not covered. Similarly, two-dimensional and pitched copies (cudaMemcpy2D) are currently unsupported. Supporting these would require engineering new interception hooks, but no changes to the core analysis.

The largest reduction is 𝐵 M2 . Its mean rail power rises by 18.8%, but latency falls from 4.27 to 0.61 ms, yielding an 83.2% lower energy estimate. Nano saves 3.9–75.8% (See Appendix Table 10). Thor 𝐵 M3 and 𝐵 E2 use baseline fallback, so their plotted differences reflect variation between baseline runs.

Compile-time visibility. As discussed in §3.2, SR conservatively declines (h,d) pairs with dynamic sizes, crosscompilation-unit pointer flow, or complex control flow. Wholeprogram analysis via link-time optimization (LTO) could widen its scope, but the dominant barrier is dynamic-library boundaries: copies inside separately compiled shared libraries

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

(e.g., TensorRT, cuDNN) are invisible at compile time regardless of analysis granularity. DR complements SR by observing runtime CUDA API calls across those library boundaries. This lets it optimize cases SR cannot reach, but the guarantee it provides is profile-based rather than a compile-time proof. Multi-stream workloads. Our current evaluation covers default-stream or explicitly synchronized copy/kernel regions. Pairs involving multiple streams, events, or overlapped copy/compute are treated as ineligible and execute through the original CUDA runtime path. Optimizing them would require stream-aware dependency analysis and additional synchronization repair. Platform scope. Echo’s transformation principle applies to any UMA platform where software retains a duplicated host/device object model, but our implementation is specific to CUDA. High-end data center architectures such as the Grace Hopper Superchip [25] provide hardware cachecoherency via NVLink-C2C, but this does not itself remove explicit copies between separately allocated buffers. On edge SoCs lacking such interconnects, however, the legacy dGPUstyle code pattern remains pervasive. Adapting Echo to other UMA ecosystems (e.g., Apple Metal, Vulkan on integrated GPUs, AMD APU OpenCL) is a natural direction for future work, provided the implementation preserves the same correctness and safety guarantees.

6

Related Work

Redundant Copies on Other UMA Platforms. The redundant copy problem is not specific to CUDA or Jetson. It arises on any UMA platform where programming abstractions assume separate CPU and GPU memory spaces. Dashti and Fedorova [4] provide a foundational study of memory management methods on integrated CPU–GPU systems, showing that explicit-copy code persists even when zerocopy and unified alternatives are available. On Qualcomm Adreno, Wang et al. [28] identify memory copies between Java layer and native layer as a common OpenCL issue on Snapdragon SoCs and recommend the platform-specific zerocopy extension cl_qcom_ion_host_ptr. Grasso et al. [6] report the same pattern on ARM Mali-T604, where map/unmap operations must be manually substituted to avoid unnecessary transfers. Newer hardware and frameworks have begun to address this problem at the platform level. On AMD MI300A APUs, Wahlgren et al. [27] show that applications using the unified physical memory model match or outperform explicit management while reducing memory costs by up to 44%. On Apple Silicon, MLX-based inference [2] achieves 21–87% higher throughput than llama.cpp by exploiting zero-copy

tensor operations, and Hubner et al. [9] confirm that MTLResourceStorageModeShared buffers eliminate manual transfers entirely. However, these benefits accrue only to applications written for (or ported to) the new APIs. On Jetson, and more broadly across large existing CUDA codebases that follow the dGPU-style pattern, the legacy pattern persists, and rewriting is impractical for closed-source and binary-only deployments. Echo-DR addresses this gap by optimizing existing applications without requiring source modification or framework migration. Compiler and Runtime Copy Optimization. Standard compiler optimizations (mem2reg, GVN, dead-store elimination) remove redundant memory operations within a single address space [10], but they optimize individual load/store instructions and are unaware of the cross-address-space redundancy introduced by CUDA’s allocation and copy API conventions. Similarly, compiler memcpy optimizations reason about byte copies between pointers in one memory model; they do not decide that a host allocation and a device allocation should become one object, nor do they repair the ordering effect of a removed D2H copy. Li et al. [11] apply LLVMbased data-reuse analysis to OpenMP GPU offloading under unified memory, reducing implicit data transfers through compiler–runtime collaboration. However, their work targets OpenMP’s pragma-based implicit data management, not the explicit malloc–cudaMalloc–cudaMemcpy pattern in native CUDA code. The Grace Hopper Superchip [25] supports coherent access to shared system memory via NVLink-C2C, but explicit copies between separate allocations remain a software choice. This hardware feature is unavailable on edge SoCs like Jetson. Source-Level Refactoring. Nejadfard and Sang [13] present a Clang LibTooling-based refactoring tool that automatically replaces malloc/cudaMalloc with cudaMallocManaged in CUDA source code. Their tool operates at the AST level and performs syntactic pattern matching without analyzing control flow, data dependencies, or pointer aliasing, and therefore cannot verify whether unifying a specific pair is semantically safe, as Echo-SR does. Echo-SR formulates the rewrite on LLVM IR, within the full-buffer H2D scope, with sufficient conditions for preserving live values, access order, pointer uses, and allocation lifetimes. Source refactoring also cannot reach copies hidden in closed-source runtimes such as TensorRT, whereas EchoDR can apply a guarded API-level optimization to stable binary-only deployments. We do not include the LibTooling tool as an experimental baseline because it performs no safety analysis. Applying it to our benchmarks might produce similar allocation replacement but without checking the safety conditions, making the comparison about analysis capability rather than runtime performance. To the best of our knowledge, Echo is the first system to combine a source-available path that proves semantic safety

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

for this allocation unification transformation with a guarded binary path for stable, API-visible legacy CUDA patterns.

7

Conclusion

Echo is a system that automatically detects and eliminates redundant memory copies on unified-memory architectures, where dGPU-style code still persists and wastes bandwidth by copying data between buffers that already share the same physical DRAM. Echo’s static and dynamic approaches verify safety conditions to ensure that unifying separate host and device allocations into a single shared allocation does not change program semantics, and then eliminate the associated copies. This principle applies broadly to any UMA platform where programming models assume separate CPU and GPU memory spaces, and Echo’s implementation for CUDA on Jetson serves as a case study demonstrating the potential for significant performance gains with minimal engineering effort.

References [1] Tyler Allen and Rong Ge. 2021. In-depth analyses of unified virtual memory system for GPU accelerated computing. In Proceedings of 2021 International Conference for High Performance Computing, Networking, Storage and Analysis (St. Louis, Missouri) (SC ’21). Association for Computing Machinery, New York, NY, USA, Article 64, 15 pages. doi:10.1145/3458817.3480855 [2] Wayner Barrios. 2026. Native LLM and MLLM Inference at Scale on Apple Silicon. arXiv preprint arXiv:2601.19139 (2026). [3] Jiasi Chen and Xukan Ran. 2019. Deep Learning With Edge Computing: A Review. Proc. IEEE 107, 8 (2019), 1655–1674. doi:10.1109/JPROC. 2019.2921977 [4] Mohammad Dashti and Alexandra Fedorova. 2017. Analyzing memory management methods on integrated CPU-GPU systems. SIGPLAN Not. 52, 9 (June 2017), 59–69. doi:10.1145/3156685.3092256 [5] Debashis Ganguly, Ziyu Zhang, Jun Yang, and Rami Melhem. 2019. Interplay between Hardware Prefetcher and Page Eviction Policy in CPUGPU Unified Virtual Memory. In Proceedings of the 46th ACM/IEEE Annual International Symposium on Computer Architecture (ISCA). 224– 235. [6] Ivan Grasso, Petar Radojkovic, Nikola Rajovic, Isaac Gelado, and Alex Ramirez. 2014. Energy Efficient HPC on Embedded SoCs: Optimization Techniques for Mali GPU. In Proceedings of the 28th IEEE International Parallel and Distributed Processing Symposium (IPDPS). 123–132. doi:10. 1109/IPDPS.2014.24 [7] Yongbin Gu, Wenxuan Wu, Yunfan Li, and Lizhong Chen. 2020. UVMBench: A Comprehensive Benchmark Suite for Researching Unified Virtual Memory in GPUs. CoRR abs/2007.09822 (2020). arXiv:2007.09822 https://arxiv.org/abs/2007.09822 [8] Kaiming He, Xiangyu Zhang, Shaoqing Ren, and Jian Sun. 2016. Deep Residual Learning for Image Recognition. In Proceedings of 2015 IEEE Conference on Computer Vision and Pattern Recognition (CVPR). IEEE. doi:10.1109/CVPR.2016.90 [9] Paul Hübner, Andong Hu, Ivy Peng, and Stefano Markidis. 2025. Apple vs. Oranges: Evaluating the Apple Silicon M-Series SoCs for HPC Performance and Efficiency. In Proceedings of 2025 IEEE International Parallel and Distributed Processing Symposium Workshops (IPDPSW). 45–54. doi:10.1109/IPDPSW66978.2025.00013 [10] C. Lattner and V. Adve. 2004. LLVM: a compilation framework for lifelong program analysis & transformation. In Proceedings of 2004 International Symposium on Code Generation and Optimization. 75–86.

doi:10.1109/CGO.2004.1281665 [11] Lingda Li, Hal Finkel, Martin Kong, and Barbara Chapman. 2018. Manage OpenMP GPU Data Environment Under Unified Address Space. In Proceedings of the 14th International Workshop on OpenMP (IWOMP), Bronis R. de Supinski, Pedro Valero-Lara, Xavier Martorell, Sergi Mateo Bellido, and Jesus Labarta (Eds.). Springer International Publishing, Cham, 69–81. [12] Jonathan Long, Evan Shelhamer, and Trevor Darrell. 2015. Fully Convolutional Networks for Semantic Segmentation. In Proceedings of 2015 IEEE Conference on Computer Vision and Pattern Recognition (CVPR). IEEE. doi:10.1109/CVPR.2015.7298965 [13] Kian Nejadfard and Janche Sang. 2024. Developing a Clang LibToolingBased Refactoring Tool for CUDA GPU Programming. Journal of Software Engineering and Applications 17, 2 (2024), 89–108. doi:10. 4236/jsea.2024.172005 [14] noahmr. 2026. yolov5-tensorrt: YOLOv5 Inference with TensorRT. GitHub. https://github.com/noahmr/yolov5-tensorrt. [15] NVIDIA. 2026. cuBLAS: CUDA Basic Linear Algebra Subroutines. CUDA Toolkit Documentation. https://docs.nvidia.com/cuda/cublas/. [16] NVIDIA. 2026. CUDA for Tegra. CUDA Toolkit Documentation, §4.1 Memory Selection. https://docs.nvidia.com/cuda/cuda-for-tegraappnote/. [17] NVIDIA. 2026. cuDNN: CUDA Deep Neural Network Library. NVIDIA Developer. https://developer.nvidia.com/cudnn. [18] NVIDIA. 2026. cuFFT: CUDA Fast Fourier Transform. CUDA Toolkit Documentation. https://docs.nvidia.com/cuda/cufft/. [19] NVIDIA. 2026. Jetson AGX Orin Series. NVIDIA Embedded Systems. https://www.nvidia.com/en-us/autonomous-machines/ embedded-systems/jetson-orin/. [20] NVIDIA. 2026. Nsight Systems. NVIDIA Developer Tools. https: //developer.nvidia.com/nsight-systems. [21] NVIDIA. 2026. NVIDIA DeepStream SDK Developer Guide. https: //developer.nvidia.com/deepstream-sdk. [22] NVIDIA. 2026. TensorRT C++ Samples. GitHub. https://github.com/ NVIDIA/TensorRT/tree/main/samples. [23] NVIDIA. 2026. TensorRT: High-Performance Deep Learning Inference. NVIDIA Developer. https://developer.nvidia.com/tensorrt. [24] RichardoMrMu. 2026. deepsort-tensorrt: DeepSORT with TensorRT in C++. GitHub. https://github.com/RichardoMrMu/deepsort-tensorrt. [25] Gabin Schieffer, Jacob Wahlgren, Jie Ren, Jennifer Faj, and Ivy Peng. 2024. Harnessing Integrated CPU-GPU System Memory for HPC: a first look into Grace Hopper. In Proceedings of the 53rd International Conference on Parallel Processing (Gotland, Sweden) (ICPP ’24). Association for Computing Machinery, New York, NY, USA, 199–209. doi:10.1145/3673038.3673110 [26] Vivienne Sze, Yu-Hsin Chen, Tien-Ju Yang, and Joel S. Emer. 2017. Efficient Processing of Deep Neural Networks: A Tutorial and Survey. Proc. IEEE 105, 12 (2017), 2295–2329. doi:10.1109/JPROC.2017.2761740 [27] Jacob Wahlgren, Gabin Schieffer, Ruimin Shi, Edgar A. León, Roger Pearce, Maya Gokhale, and Ivy Peng. 2025. Dissecting CPU-GPU Unified Physical Memory on AMD MI300A APUs. In Proceedings of 2025 IEEE International Symposium on Workload Characterization (IISWC). 368–380. doi:10.1109/IISWC66894.2025.00038 [28] Hongqiang Wang, Jay Yun, and Alex Bourd. 2018. OpenCL Optimization and Best Practices for Qualcomm Adreno GPUs. In Proceedings of the International Workshop on OpenCL (IWOCL) (Oxford, United Kingdom) (IWOCL ’18). Association for Computing Machinery, New York, NY, USA, Article 16, 8 pages. doi:10.1145/3204919.3204935 [29] Weixi Zhu, Guilherme Cox, Jan Vesely, Mark Hairgrove, Alan L. Cox, and Scott Rixner. 2022. UVM Discard: Eliminating Redundant Memory Transfers for Accelerators. In Proceedings of 2022 IEEE International Symposium on Workload Characterization (IISWC). 27–38. doi:10.1109/ IISWC55918.2022.00013

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

A

Proof of SR Soundness

Proof of Theorem 1. Fix an execution of 𝑃 ′ with events 𝐸 ′ . We build the matching execution 𝐸 of 𝑃 alongside it and show that every retained read returns the same value in both. Branches and external calls depend only on read values, so the two executions then take the same branches and make the same external calls, which gives [[𝑃 ′ ]] ⊆ [[𝑃]]. Step 1: 𝑃 ′ is race-free. Take 𝑎, 𝑏 ∈ 𝐸 ′ that access the same byte, one of them a write. Events of sync(𝑐) access no memory, and S4 orders the allocation and free of 𝑝 against every access, so we may assume both are retained. If they access a byte outside 𝑝, or ℎ[𝑖] and ℎ[𝑖], or 𝑑 [𝑖] and 𝑑 [𝑖] in 𝑃, then 𝑃 being race-free orders them by ≺. If they access ℎ[𝑖] and 𝑑 [𝑖], S2 orders them by ≺. In both cases (OP) orders them by ≺′ . The values that 𝑃 ′ reads therefore do not depend on how its unordered events interleave, and we may follow any total order 𝜏 ′ of 𝐸 ′ consistent with ≺′ . Step 2: a common order. Drop the sync(𝑐) events and the allocation and free of 𝑝 from 𝜏 ′ , and insert each event of 𝐸 \ 𝐸ˆ (linking copies and the allocations and frees of ℎ and 𝑑) right after its last ≺-predecessor. Call the result 𝜏. By (OP), 𝜏 is consistent with ≺, and since 𝑃 is race-free, reading along 𝜏 gives the values of 𝐸. Moreover, if a read 𝑟 of 𝑥 follows 𝑒 in 𝜏 with no write to 𝑥 between them, then 𝑟 ̸≺ 𝑒 and no 𝑤 satisfies 𝑒 ≺ 𝑤 ≺ 𝑟 , so live(𝑒, 𝑥) holds. Contrapositively, S1 guarantees that 𝑥¯ is never read along 𝜏 after a retained write to 𝑥 unless it is written again first. Step 3: invariant. We walk along 𝜏 and maintain the following. For every 𝑖 and every initialized 𝑥 ∈ {ℎ[𝑖], 𝑑 [𝑖]} that will be read along 𝜏 before being written again, 𝑃 holds in 𝑥 the value that 𝑃 ′ holds in 𝑝 [𝑖]. All bytes outside ℎ, 𝑑, and 𝑝 hold the same values in both programs. The invariant holds initially. For the next event 𝑒 of 𝜏: • Read of 𝑥 ∈ ℎ ∪ 𝑑. 𝑥 is initialized, by the premise on uninitialized bytes, and 𝑒 reads it next, so the invariant gives the same value in 𝑃 ′ . Reads of other bytes agree directly. • Retained write to 𝑥 ∈ ℎ ∪ 𝑑. Its operands agree, so both programs store the same value, 𝑃 into 𝑥 and 𝑃 ′ into 𝑝 [𝑖], and the invariant holds for 𝑥. 𝑃 ′ also overwrites the old ¯ which by Step 2 and S1 is not read again before value of 𝑥, ¯ No other being rewritten, and 𝑒 itself does not read 𝑥. byte changes. • Linking copy (only in 𝜏). For each copied 𝑖, 𝑃 gives the ¯ If 𝑦¯ is initialized, the copy destination 𝑦 the value of 𝑦. reads it, so the invariant held for 𝑦¯ and 𝑝 [𝑖] already holds that value. The invariant then holds for 𝑦 although 𝑃 ′ does nothing. If 𝑦¯ is uninitialized, so is 𝑦, and nothing is required. • Allocation or free of ℎ or 𝑑 (only in 𝜏). A new byte is uninitialized, so nothing is required. By S4, 𝑝 is allocated before and freed after every retained access to ℎ or 𝑑.

Baseline

DR

SR

Manually Optimized

B

BM1 DS M 13

14

15

16

17

18

B BM2 DS M 1

2

3

4

5

6

7

8

B BM3 DS M 3

4

5

6

B BC1 DS M 28

30

32

34

36

38

40

B BC2 DS M 42

44

46

48

50

52

54

BE1 BD 22

24

26

28

30

BE2 BD 9

10

11

12

13

14

15

Primary-region latency (ms) Note: x-axis scale differs per panel.

Figure 9. Primary-region latency on Orin Nano, using the five-run measurement and plotting conventions of Figure 4. Each benchmark panel has its own x-axis scale. • sync(𝑐) (only in 𝜏 ′ ). It accesses no memory. • Any other event. It is the same operation in both programs. By S3, ℎ and 𝑑 appear in it only as addresses, so replacing them with 𝑝 changes neither the values it computes nor the bytes it touches. Hence every retained read returns the same value in 𝑃 and 𝑃 ′ . For several disjoint pairs, the invariant and S1–S4 apply to each pair separately, so the argument covers them together. □

B

Cross-Platform Results

Tables 6–11 and Figures 9–10 report the complete DR and SR results for all three platforms. The main text in Section 4 centers on Orin via Figure 4. The numerical Orin table and the Nano/Thor results appear here. Per-Platform Results (Orin, Nano, and Thor). Nano follows the same qualitative trend as Orin, but with smaller gains throughout. Both platforms use Ampere SM 8.7, but Nano has fewer CUDA cores, so computation occupies a larger fraction of each iteration and 𝑓copy is correspondingly lower. Thor’s guard enables five net-positive workloads: 𝐵 M1 , 𝐵 M2 , 𝐵 C1 , 𝐵 C2 , and 𝐵 E1 achieve primary-region speedups of

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

Table 6. Primary-region DR results on AGX Orin over five independent runs, complementing Figure 4. Each p99 entry is the mean of the five per-process p99 values; Δp99 is DR minus Baseline. All latency and p99 values are in milliseconds. SR results appear in §4.2. ID

Baseline (ms)

DR (ms)

Speedup

Recovered

BL p99 (ms)

DR p99 (ms)

Δp99 (ms)

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

9.115 ± 0.011 4.271 ± 0.024 4.976 ± 0.026 17.336 ± 0.010 25.090 ± 0.013 20.569 ± 0.067 7.979 ± 0.054

5.389 ± 0.003 0.606 ± 0.004 1.479 ± 0.005 11.815 ± 0.011 19.557 ± 0.008 14.685 ± 0.016 6.834 ± 0.023

1.69× 7.05× 3.37× 1.47× 1.28× 1.40× 1.17×

103.3% 99.7% 99.5% 99.5% 99.6% 98.4% n/a

9.17 4.41 5.12 17.49 25.24 20.89 8.59

5.41 0.63 1.52 11.95 19.63 14.80 7.45

−3.76 −3.78 −3.59 −5.54 −5.60 −6.09 −1.14

Table 7. DR results on Nano (5-run). Recovered fraction ranges from 97.1% to 103.9% on benchmarks with a manual reference; 𝐵 E2 has no independent manual reference. ID

Baseline (ms)

DR (ms)

Speedup

Managed (ms)

Recovered

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

17.240 ± 0.011 5.063 ± 0.032 7.029 ± 0.009 35.846 ± 0.031 50.502 ± 0.058 29.161 ± 0.047 13.130 ± 0.040

13.204 ± 0.004 1.029 ± 0.010 3.244 ± 0.003 29.745 ± 0.092 44.584 ± 0.138 22.779 ± 0.126 11.808 ± 0.016

1.31× 4.92× 2.17× 1.21× 1.13× 1.28× 1.11×

13.357 ± 0.006 1.011 ± 0.006 3.234 ± 0.003 29.664 ± 0.044 44.408 ± 0.087 22.616 ± 0.042 n/a

103.9% 99.6% 99.7% 98.7% 97.1% 97.5% n/a

Table 8. DR guarded-deployment results on Thor (5-run). Asterisks denote measured baseline fallback after a non-positive profitability calibration; these rows are not raw transformed performance. ID

Baseline (ms)

DR (ms)

Speedup

Managed (ms)

Recovered

𝐵 M1 𝐵 M2 𝐵 M3 ∗ 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2 ∗

3.301 ± 0.007 0.459 ± 0.004 1.036 ± 0.021 3.836 ± 0.015 5.724 ± 0.031 10.112 ± 0.077 1.750 ± 0.006

3.030 ± 0.004 0.289 ± 0.015 1.025 ± 0.014 3.295 ± 0.008 5.152 ± 0.014 9.740 ± 0.111 1.745 ± 0.003

1.09× 1.59× 1.01× 1.16× 1.11× 1.04× 1.00×

3.019 ± 0.007 0.271 ± 0.012 1.102 ± 0.049 3.297 ± 0.015 5.162 ± 0.023 9.689 ± 0.094 n/a

96.1% 90.3% n/a 100.4% 101.8% 87.9% n/a

1.09×, 1.59×, 1.16×, 1.11×, and 1.04×, respectively. The profitability decision uses three paired complete-loop trials and enables a plan when their median relative gain is strictly positive. For 𝐵 M3 , the candidate plan increases full-loop latency by 3.7%, 16.7%, and 17.2% in the three paired calibration trials. The guard therefore selects baseline fallback. 𝐵 E2 also falls back after a small negative median full-loop gain (−0.65%). The near-1× values on these two rows are independent baseline-process variation, not transformed execution. For the enabled canonical workloads, DR and the manual reference agree closely: 𝐵 C1 recovers 100.4% and 𝐵 C2 101.8% of the measured manual-reference gain. Thor’s native 𝐵 E2 TensorRT engine supports at most 16 feature crops per frame, whereas the measured frames contain 23–32 detections and the Orin/Nano engine supports 128. We therefore label 𝐵 E2 as a capped feature-extraction job and do not compare its

absolute Thor throughput with the uncapped Orin/Nano workload.

SR Results on Nano and Thor. SR produces the same analysis decisions on all platforms because it operates on platform-independent LLVM IR. Table 9 shows the primary-region latency of SR on Nano and Thor for the five source-available benchmarks. On Nano, SR closely matches DR (Table 7), and their speedups differ by at most 0.02×. On Thor, the two approaches also agree on 𝐵 M1 , 𝐵 C1 , and 𝐵 C2 . 𝐵 M2 runs faster under SR (1.71× versus 1.59×), and its latency (0.268 ms) is within noise of the manual reference (0.271 ms in Table 8). 𝐵 M3 is the only regression (0.91×), whereas DR falls back to the baseline for this workload because, unlike guarded DR, SR applies its accepted transformation without a runtime profitability calibration.

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

Baseline

DR

SR

Table 10. Primary-region energy on Nano. For each run, energy is the mean rail power in its measurement window multiplied by that run’s primary-region latency. Savings range from 4–76%.

Manually Optimized

B

BM1 DS M 3.0

3.2

3.4

3.6

3.8

B BM2 DS M 0.20

0.25

0.30

0.35

0.40

0.45

0.50

1.00

1.05

1.10

1.15

1.20

3.2

3.4

3.6

3.8

4.0

0.55

B

BM3 D*S M 0.95

1.25

1.30

B BC1 DS M

5.5

6.0

BL (mW)

DR (mW)

BL (mJ)

DR (mJ)

Saved

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

9741 3268 6213 10242 11188 5081 6304

11840 3894 9935 11820 12185 5744 6473

167.9 16.5 43.7 367.1 565.0 148.2 82.8

156.3 4.0 32.2 351.6 543.2 130.9 76.4

6.9% 75.8% 26.2% 4.2% 3.9% 11.7% 7.7%

Table 11. Primary-region copy fraction across three platforms. 𝑓copy is baseline profiled GPU copy-only time divided by the corresponding baseline primary-region latency, matching the speedup scope used in the Amdahl analysis.

4.2

B BC2 DS M 5.0

ID

6.5

BE1 BD 9.4

9.6

1.700

1.725

9.8

10.0

10.2

10.4

10.6

BE2 D*B 1.750

1.775

1.800

1.825

1.850

Primary-region latency (ms) Note: x-axis scale differs per panel.

Figure 10. Primary-region latency on AGX Thor, using the conventions of Figure 4. D* marks DR’s baseline fallback for 𝐵 M3 and 𝐵 E2 . SR and Manual show transformed executions. Table 9. SR runtime on Nano and Thor (5-run). On Nano, SR and DR remain closely matched. On Thor, 𝐵 M2 reaches 1.71× under SR versus 1.59× under enabled DR. 𝐵 M3 regresses to 0.91× because SR has no deployment-time profitability fallback. Nano

Thor

ID

SR (ms)

Speedup

ID

SR (ms)

Speedup

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2

13.369 ± 0.007 1.027 ± 0.010 3.254 ± 0.004 29.683 ± 0.042 44.518 ± 0.207

1.29× 4.93× 2.16× 1.21× 1.13×

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2

3.032 ± 0.006 0.268 ± 0.001 1.138 ± 0.033 3.307 ± 0.025 5.132 ± 0.010

1.09× 1.71× 0.91× 1.16× 1.12×

Power and Energy on Nano. Table 10 reports the mean rail power and primary-region energy of DR on Nano. On Nano, mean power on the measured rail rises by approximately 3–60% under Echo, but primary-region energy still decreases because the optimized region finishes sooner. Across enabled benchmarks where both apply, DR and SR agree within 3% in primary-region energy. 𝑓copy Across Platforms. Across the 21 platform/benchmark points, the primary-region measurements follow the

ID

Orin 𝑓copy

Nano 𝑓copy

Thor 𝑓copy

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

40.5% 85.3% 67.1% 29.8% 20.1% 28.7% 13.7%

23.3% 78.3% 52.6% 15.7% 11.1% 22.6% 9.7%

7.3% 49.9% 20.9% 8.6% 5.9% 4.2% 1.2%

same Amdahl trend (𝑅 2 = 0.989). The principal deviations are 𝐵 M2 and 𝐵 M3 on Thor, where managed-memory kernel slowdown offsets part or all of the copy-elimination benefit; excluding these two points gives 𝑅 2 = 0.994.

C

Complete-Loop Throughput and Buffer Capacity

Table 5 in Section 4 reports Orin results. Tables 12 and 13 gives the corresponding loop throughput and selected buffer capacity on Nano and Thor. Throughput uses the complete timed loop and each workload’s work unit. Capacity counts the live allocations selected for unification instead of measuring process RSS, because CUDA allocation accounting in RSS is not comparable across the evaluated Tegra generations. The primary-region latency results are shown in Tables 6–8.

D

Managed-Memory Slowdown on Thor

Thor exposes a limit of copy elimination that is not visible from the copy fraction alone. Its 𝐵 M2 plan remains profitable, but 𝐵 M3 ’s guarded DR result is a baseline fallback; the independently transformed SR execution is slower than Baseline at 0.91×. Both pairs pass observed host-access validation, so the different deployment decisions call for a performance explanation rather than another safety check. We examine the paired calibration trials to see whether saved copy time

Echo: Merging Host–Device Buffers to Avoid Redundant Data Movement on Unified-Memory SoCs

Table 12. Complete-loop throughput and selected buffer capacity on Orin Nano. 𝑄 is reported as mean ± sample SD across five independent runs; 𝑀 counts the live allocations selected for unification. Subscripts B and D denote Baseline and DR, respectively. Benchmark 𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2

𝑄 𝐵 (work units/s)

𝑄 𝐷 (work units/s)

𝑄 𝐷 /𝑄 𝐵

𝑀𝐵 (MiB)

𝑀𝐷 (MiB)

57.78 ± 0.06 142.71 ± 0.76 141.27 ± 0.19 27.90 ± 0.02 19.80 ± 0.02 137.17 ± 0.22 61.48 ± 0.19

75.57 ± 0.05 347.55 ± 1.10 305.42 ± 0.23 33.62 ± 0.10 22.43 ± 0.07 175.60 ± 0.96 66.89 ± 0.07

1.308× 2.435× 2.162× 1.205× 1.133× 1.280× 1.088×

64.000 64.000 58.188 90.000 90.000 102.869 24.500

32.000 32.000 29.094 45.000 45.000 51.434 12.250

Table 13. Complete-loop throughput and selected buffer capacity on AGX Thor. 𝑄 is reported as mean ± sample SD across five independent runs; 𝑀 counts the live allocations selected for unification. Subscripts B and D denote Baseline and DR, respectively. Benchmark 𝐵 M1 𝐵 M2 𝐵 M3 ∗ 𝐵 C1 𝐵 C2 𝐵 E1 𝐵 E2 ∗

𝑄 𝐵 (work units/s)

𝑄 𝐷 (work units/s)

𝑄 𝐷 /𝑄 𝐵

𝑀𝐵 (MiB)

𝑀𝐷 (MiB)

300.67 ± 0.63 817.02 ± 3.51 941.95 ± 20.60 260.68 ± 1.00 174.69 ± 0.96 395.56 ± 3.04 321.66 ± 1.53

327.41 ± 0.44 861.52 ± 6.49 953.53 ± 12.65 303.45 ± 0.69 194.09 ± 0.53 410.68 ± 4.72 322.72 ± 1.01

1.089× 1.054× 1.012×∗ 1.164× 1.111× 1.038× 1.003×∗

64.000 64.000 58.188 90.000 90.000 102.869 3.000

32.000 32.000 58.188 45.000 45.000 51.434 3.000

Table 14. Thor profitability calibration. Compute times are medians of three process means; full-loop reductions are medians of the three paired relative reductions. Candidate runs execute the transformed plan before the deployment decision. A negative reduction denotes increased latency.

ID 𝐵 M2 𝐵 M3

Compute (ms) Full-loop Baseline → candidate reduction Decision 0.156 → 0.257 0.745 → 1.174

+6.3% Enable −16.7% Fallback

offsets the changed computation cost, then use five controlled memory-access probes to characterize that cost on the evaluated CUDA 13/R38.2.2 Thor stack. The M2/M3 contrast. Table 14 reports the three paired calibration trials. For 𝐵 M2 , median compute-phase time rises from 0.156 to 0.257 ms, an increase of 65%. Eliminating the copies still reduces complete-loop latency by a median 6.3%, so DR enables the plan. For 𝐵 M3 , compute-phase time rises from 0.745 to 1.174 ms, an increase of 58%; complete-loop latency increases by 16.7%, so the guard retains the baseline. The independent SR measurement has no runtime profitability gate and yields a primary-region speedup of 0.91× on 𝐵 M3 .

Table 15. Memory-access probes on Thor. Ratios are managed-allocation time divided by explicit-allocation time within each probe; the matching Orin probes remain near 1.0×. T1 and T3–T5 use 16 MB buffers. The probes and the calibration in Table 14 were collected in separate campaigns. Probe

Thor observation

T1: Read/write Reads: 16.8×; writes: 1.3×. T2: Buffer size sweep Peak 1.71× at 16 MB; 0.83× at 24 MB. T3: Computation 1.02× at 200 FMA/element. T4: Stride Reaches 10.7× at 128 elements. T5: Repetition About 1.7× over 100 launches.

Access-pattern characterization. The five probes in Table 15 compare managed and explicit allocations while varying access type, buffer size, computation, stride, and repetition. At 16 MB, the read probe slows down by 16.8× on Thor, compared with 1.3× for writes. The size, computation, and stride sweeps show that the penalty depends on the access pattern: it disappears at larger buffer sizes or with sufficient computation, but remains substantial for large strides. It persists over 100 launches. The matching Orin probes remain near 1.0×. The takeaway is that semantic eligibility and net benefit are separate decisions on Thor. Managed allocations make the tested 𝐵 M2 and 𝐵 M3 computation slower, with the impact

Yuheng Zhu, Yanbo Zhao, Jiajia Li, and Man-Ki Yoon

depending on access pattern. Saved copies still outweigh that cost for 𝐵 M2 ; for 𝐵 M3 , they do not, so DR retains the baseline. The probes characterize the measured platform/software stack rather than establish a vendor-confirmed cache

mechanism. Complete-loop calibration matters because copy fraction alone cannot determine which plan should run.

Related documents

• Universal magic state concentration#1143871
• Decoupling Logical Masks from GPU Execution for Dynamic Block-Sparse Attention#1061497
• A Bifurcation Theory for the Equilibria of Modern Hopfield Networks#984657
• Philosophy after Philosophy Vol. 2 — Tonal Being and Meta-Ontology#81739
• Making Array-Based Translation Practical for Modern, High-Performance Buffer Management#2785