DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

arXiv:2609.05635v1 [cs.DC] 4 Sep 2026

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs Yuheng Zhu∗

Yanbo Zhao∗

yzhu63@ncsu.edu North Carolina State University USA

yzhao62@ncsu.edu North Carolina State University USA

Jiajia Li

Man-Ki Yoon

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

man-ki.yoon@ncsu.edu North Carolina State University 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, so the copies consume bandwidth, time, and energy without moving data across a physical boundary. 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. DejaVu removes these copies only when the program does not depend on the effects above. It does so along two complementary paths, depending on whether source is available. DejaVu-SR is a compile-time LLVM transformation that proves safety and rewrites accepted pairs in place. DejaVuDR is a profile-guided binary optimizer for closed-source deployments that profiles and validates stable allocation/copy patterns and, at runtime, intercepts the matching calls to coalesce profile-matched pairs while preserving the ordering effects of removed copies. Across seven benchmarks on three NVIDIA Jetson platforms, DejaVu’s benefit grows with the fraction of baseline time spent on copies. Copy-dominated workloads speed up by up to 6.9×, closed-source end-to-end applications speed up by 1.10–1.14×. Both source and binary paths achieve ≥99% of the performance achievable by manual optimization.

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, 24]. 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 [20] 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 platforms. We therefore take CUDA on Jetson as our concrete

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 DejaVu. (a) & (b) DejaVu detects and unifies separate host and device allocations into a single managed allocation and eliminates redundant copies on UMA platforms. (c) Speedup (bar) achieved by DejaVu on NVIDIA Jetson AGX Orin, overlaid with 𝑓𝑐𝑜𝑝𝑦 (line), the fraction of total runtime spent on memory copies before transformation. 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 [21] inference pipeline running FCN-ResNet50 [8, 12] on Jetson AGX Orin [17], 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 DejaVu, a system that automatically detects and safely eliminates these redundant copies on UMA platforms. It does so both when source is available and when it is not, trading a universal soundness guarantee for deployability in the latter case. DejaVu-SR (Static Replacement) is a compile-time LLVM pass for source-available code. It identifies host/device pairs whose unification provably cannot change the program’s observable behavior, then transforms each pair into a single unified allocation and replaces copies with explicit synchronization where needed (§3.2). DejaVuDR (Dynamic Replacement) is a binary-level optimizer for deployments where no source is available to rewrite or the copy path is buried inside a precompiled runtime. Because it has no source-level view, DR cannot soundly verify the static safety conditions for all inputs. Instead, it conservatively

applies guarded, profile-matched allocation coalescing for stable binary deployments and provides profile-conditional correctness: the optimization holds as long as runtime execution matches the profile (§3.3). We evaluate DejaVu on seven benchmarks and a negative control: three CUDA-library micro-benchmarks (cuBLAS [14], cuFFT [16], cuDNN [15]), two canonical TensorRT inference pipelines (FCN-ResNet50 [8, 12], FCN-ResNet101), and two end-to-end applications (YOLOv5 [27], DeepSORT [22]). 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. With DejaVu, copydominated micro-benchmarks achieve up to 6.9× speedup, canonical inference pipelines 1.28–1.47×, and closed-source end-to-end applications 1.10–1.14× without source changes. On the five source-available benchmarks where both SR and DR apply, each reaches ≥99% of manually optimized performance. Beyond latency, DejaVu also yields energy savings. On Orin, per-iteration energy reduction ranges from 9% to 86% and scales with the copy fraction. The evaluation also exposes two deployment boundaries. First, copies inside a closed-source runtime such as TensorRT that never cross the intercepted API are invisible to DR and limit the achievable speedup. Second, managed-memory performance varies by platform, so coalescing is not always profitable. 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. DejaVu-SR, a sound LLVM transformation that coalesces allocation pairs and eliminates the associated copies only after proving the transformation is safe.

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

3. DejaVu-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 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. 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. DejaVu 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 pagefault 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 [28] 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

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

Baseline (ms)

DejaVu (ms)

Reduction

Preprocess H2D copy Inference D2H copy Postprocess

0.336 ± 0.003 0.316 ± 0.011 11.602 ± 0.049 5.313 ± 0.047 0.002 ± 0.000

0.329 ± 0.001 0.001 ± 0.000 11.647 ± 0.041 0.010 ± 0.001 0.002 ± 0.000

+2.1% 99.7% −0.4% 99.8% 0%

Total

17.568 ± 0.107

11.989 ± 0.040

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

cudaMemcpy calls inherited from dGPU-style code, not from page faults or data migration. 2.3

Measured Opportunity

We profile FCN-ResNet50 [8, 12] inference on Jetson AGX Orin using Nsight Systems [18]. Table 1 breaks down perstage latency before and after applying DejaVu-DR to the unmodified binary. Copies account for 32% of baseline iteration time. DejaVu 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.97 ± 0.02 ms, only 0.02 ms faster than DejaVu. DejaVu 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, 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

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

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 DejaVu 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 coalesced requires either DejaVu’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 DejaVu

3.1

Overview

DejaVu 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. We first define a candidate pair as a host–device allocation whose associated copies are potentially redundant: Definition 1 (Candidate Pair). A host allocation h (malloc, calloc, or realloc) and a device allocation d (cudaMalloc) form a candidate pair if: 1. h and d allocate the same size 𝑁 . 2. At least one cudaMemcpy transfers data between them. A candidate pair (h, d) that satisfies the following safety conditions is safe to unify into a single (cudaMallocManaged) allocation: Definition 2 (Safety Conditions). Let (h, d) be a candidate pair. Merging them into a single pointer p = cudaMallocManaged(N) is safe if it satisfies the following sufficient conditions: S1. No write to h occurs between the H2D copy and the next kernel launch or synchronization point. S2. No read from or write to h occurs between the last kernel launch and the D2H copy. S3. h and d must be used only as handles to the data they point to, never as observable values themselves (because once unified, they become the same pointer). Hence, all their uses must be visible to SR, so neither escapes into opaque code. S4. No other pointer in the function aliases h. S5. p is allocated before the first use of h or d and freed after the last use of both.

These conditions are stated here as a compact contract for the transformation. §3.2.2 explains why each check is needed: S1–S2 preserve the isolation originally provided by copies, S3 prevents the pointers’ addresses from being observed or escaping the analysis, S4 makes aliases visible to the analysis, and S5 preserves allocation lifetime. After redirecting all references to h and d to p, each cudaMemcpy between them becomes a self-copy, i.e., cudaMemcpy(p,p,N), which is a no-op that can be safely eliminated. SR verifies all five conditions at compile time on LLVM IR and guarantees soundness (§3.2). DR has no source-level view and therefore cannot soundly discharge those conditions for all inputs. Instead, it conservatively optimizes only stable CUDA runtime API patterns that match the profile and leaves unmatched executions in the original pass-through path. 3.2

SR: Static Analysis and Rewriting

SR analyzes and transforms LLVM IR at compile time. It is sound, that is, when it accepts a pair, it guarantees that the transformation preserves program behavior. However, SR is not complete because it is scoped to single function bodies and conservatively skips pairs when sufficient conditions cannot be established (e.g., dynamically sized allocations or pointers that escape to other compilation units). Conservatively skipping a safe pair only misses a performance opportunity, whereas accepting an unsafe pair can silently corrupt output. It also applies a cost filter to discard pairs whose copy size falls below a threshold 𝜃 min , as their copy overhead is negligible. 3.2.1 Transformation. Figure 1(a) shows the canonical before/after form of this transformation. For each (h, d) pair that passes safety verification (described next) and meets the cost criteria, SR replaces both allocations with a single cudaMallocManaged(N) call and uses LLVM’s Replace-AllUses-With (RAUW) to redirect every reference to h and d to the new pointer p. After RAUW, each H2D cudaMemcpy(d, h, N) is also eliminated as otherwise it would be a self-copy, cudaMemcpy(p, p, N), with no data effect. Each D2H copy is replaced with cudaDeviceSynchronize() to preserve the GPU-to-CPU ordering guarantee that the original copy provided. Finally, the original free(h) and cudaFree(d) are replaced by cudaFree(p). 3.2.2 Safety Verification. The core of SR is its safety analysis, which checks that no execution path depends on the original program having distinct host/device buffers or on the ordering enforced by the cudaMemcpy calls. Figure 2 shows a counterexample for each safety check (S1–S5). S1 and S2 check whether the program relies on h and d holding different values at any point. Before unification they are separate buffers, but after unification they become the same buffer, so modifications to one are immediately visible through the other. If the program relies on this isolation,

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

1 void fail_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 }

1 void fail_S2(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 kernel<<<G,B>>>(d, N); 6 float val = h[0]; 7 cudaMemcpy(h, d, N, D2H); 8 free(h); cudaFree(d); 9 10 }

1 float* g_ptr; 2 void fail_S3(int N) { float* h = (float*)malloc(N); 3 float* d; cudaMalloc(&d, N); 4 g_ptr = h; 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 }

S1: write to host before H2D

S2: access host before D2H

S3: pointer escape

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

S4: pointer alias

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

S5: Allocation not dominating use

Figure 2. SR safety checks illustrated with code examples. S1 rejects a host write between the H2D copy and kernel (line 6). S2 rejects a host access (read/write) between the kernel and D2H copy (line 7). S3 rejects a pointer that escapes to a global variable (line 5). S4 rejects a write through an aliased pointer (line 7). S5 rejects an allocation inside a conditional branch that does not dominate its uses (line 5). unifying the buffers would change its behavior, and SR must reject the pair.

input. We show that 𝑃 and 𝑃 ′ produce the same observable behavior.

• S1 rejects pairs where the host writes to h after the H2D copy. In the example shown in Figure 2 (S1, line 6), after the H2D copy, the host modifies h while the GPU kernel still holds a reference to d. With separate buffers, the kernel sees the original snapshot. After unification (h = d), the same write is immediately visible to the kernel, which potentially corrupts its input. • S2 rejects pairs where the host reads or writes h between the kernel and the D2H copy. In the example shown in Figure 2 (S2, line 7), the host reads h before the GPU has completed its write to d. With separate buffers, this read would safely return the original host data, but after unification, it may return a partially written GPU result instead.

Running example. We use the following canonical instance of the pattern throughout this subsection. The first listing is the original program 𝑃, and the second is the transformed program 𝑃 ′ , in which SR has merged h and d into one managed buffer p, deleted the H2D copy, and replaced the D2H copy with cudaDeviceSynchronize(). The two line up line by line, so every reference to h or d in 𝑃 becomes a reference to p in 𝑃 ′ .

S1 and S2 are sound only if SR can see every access to the buffer. Two checks ensure this: • S3 verifies that the host pointer does not escape the function (e.g., no stores to globals, no passing to opaque callees, no returns). If the pointer escapes (S3, line 1 and 5), external code could access the buffer in ways that SR cannot observe: for example, it could access the buffer (silently violating S1 or S2), or compare or cast the pointer itself, which unification would change since the two addresses become identical. • S4 verifies that no other pointer in the function aliases the host buffer. Since S1 and S2 only track accesses through the original pointer, an alias would allow the program to read or write the buffer through a path that SR cannot see (S4, line 4 and 7). Finally, SR checks that the unified buffer is always allocated before it is used:

Original program 𝑃: 1 2 3 4 5 6 7 8 9 10

float *h = malloc(N); float *d; cudaMalloc(&d, N); for (int i = 0; i < T; i++) { produce_input(h); // host writes h cudaMemcpy(d, h, N, H2D); // copy h to d kernel<<<G, B>>>(d); // kernel reads d, writes d cudaMemcpy(h, d, N, D2H); // copy d to h consume_output(h); // host reads h } free(h); cudaFree(d);

Transformed program 𝑃 ′ : 1 2 3 4 5 6 7 8 9 10

float *p = cudaMallocManaged(N); // (cudaMalloc folded into line 1) for (int i = 0; i < T; i++) { produce_input(p); // now writes p // H2D copy deleted; p already holds the input kernel<<<G, B>>>(p); // kernel reads p, writes p cudaDeviceSynchronize(); consume_output(p); // now reads p } cudaFree(p);

• S5 ensures that the allocation unconditionally precedes all its uses. If an allocation sits inside a conditional branch (S5, line 5) but its uses do not, some execution paths may reach a use without having executed the allocation.

Definition 3 (Divergence). Let 𝑃 and 𝑃 ′ run on the same input. A read event 𝑚 is a host load or a kernel load of the pair’s buffer; let 𝑉𝑃 (𝑚) and 𝑉𝑃 ′ (𝑚) be the value it returns in 𝑃 and 𝑃 ′ . We say 𝑃 ′ diverges from 𝑃 if the two executions take different branches, or if 𝑉𝑃 (𝑚) ≠ 𝑉𝑃 ′ (𝑚) for some corresponding read event 𝑚.

3.2.3 Core correctness property. We compare the original program 𝑃 and the transformed program 𝑃 ′ on the same

Theorem 1 (Correctness of SR). Let (h, d) be a candidate pair (Definition 1) accepted by SR that satisfies the safety conditions

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

S1–S5. Then, the transformed program 𝑃 ′ does not diverge from the original program 𝑃 (Definition 3) and therefore computes the same results. Proof. By Definition 3 we must show two things: that 𝑃 and 𝑃 ′ take the same branches, and that every read event returns the same value in both. The only difference between the two programs is the buffer: 𝑃 keeps h and d as two separate buffers, while 𝑃 ′ keeps a single buffer p. Because the buffer is reached only through h and d (S3, S4) and stays allocated throughout (S5), we can track its contents through the pair’s own operations. We follow one loop iteration (lines 4–8). The buffer is read in exactly two places, the kernel load on line 6 and the host read on line 8. • Host writes the input (Line 4). Both programs fill the buffer with the same input: 𝑃 writes h, and 𝑃 ′ writes p. In 𝑃, d still holds old data from the previous iteration, but nothing reads d yet. • H2D copy (Line 5). In 𝑃 the copy moves the input into d. In 𝑃 ′ the copy is deleted, but p already holds the input from line 4. So in both programs the buffer the kernel will read now holds the input, and by S1 the host does not change it before the kernel. • Kernel (Line 6). The kernel reads the same input in both programs (𝑃 from d, 𝑃 ′ from p), so this read event satisfies 𝑉𝑃 (𝑚) = 𝑉𝑃 ′ (𝑚). It writes its result back into the same buffer, and equal inputs give equal results. • D2H copy (Line 7). In 𝑃 the D2H cudaMemcpy blocks the host until the GPU has finished and then moves the result from d into h. In 𝑃 ′ , cudaDeviceSynchronize() blocks the host the same way, and p already holds the result. By S2 the host has not read the buffer earlier, so this is the first point at which it could, and in both programs the buffer now holds the result. • Host reads the output (Line 8). The host reads the buffer (𝑃 from h, 𝑃 ′ from p), which holds the same result in both programs, so this read also satisfies 𝑉𝑃 (𝑚) = 𝑉𝑃 ′ (𝑚). These are the only reads of the buffer, and host loads of all other memory are identical because the transformation leaves it untouched, so every read satisfies 𝑉𝑃 (𝑚) = 𝑉𝑃 ′ (𝑚). The two executions also take the same branches. A branch condition is computed only from read values and from variables that the transformation never touches, so equal reads force equal branch outcomes. In the example the only branch is the loop condition i < T, which uses the counter i and the bound T; both are untouched, so the two executions run the loop the same number of times. A branch that instead depends on the buffer (e.g., sitting inside consume_output, which reads h in 𝑃 and p in 𝑃 ′ ) is decided the same way, since that read returns the same value in both. These two facts close a step-by-step induction over the execution. Suppose the two runs agree up to some point under the correspondence of Definition 3, where every instruction is identical except for the rewritten allocation, copy, and free

calls. Then their next pair of corresponding instructions also agrees: a read returns the same value, a branch takes the same direction, a rewritten call is replaced as specified, and any other instruction is identical and touches only memory the two runs share. The agreement therefore extends to the whole execution, so 𝑃 ′ does not diverge from 𝑃. □ The argument is per pair, and the pairs compose. Each rewrite changes only its own buffer and preserves every read, so coalescing several pairs together still preserves behavior. 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 addresses this by interposing on API calls at the binary level. 3.3.1 Deployment Contract. Because DR sees API calls rather than source-level loads, stores, and control flow, it cannot soundly verify all of the safety conditions of Definition 2. Instead, it relies on a profiled deployment contract, which is a set of preconditions, captured during profiling, that a deployment must satisfy for coalescing to apply. DR checks these preconditions at the API boundary: the runtime allocation/copy structure must match the profiled plan, i.e., the same CUDA/libc call sites, allocation sizes, copy directions, and synchronization structure. Under this contract, DR provides profile-conditional correctness, meaning it preserves behavior for executions that stay within the profiled patterns, and, as with any profile-guided optimization, the guarantee remains conditional on the stability of those patterns. In this sense, DR is a conservative profile-guided runtime specialization, not a proof system. DR enforces the contract through a set of runtime rules. A plan is enabled only when the runtime structure matches the profiled plan; a copy is skipped only when coalescing has made src == dst on a pointer DR created; a skipped D2H copy preserves the original GPU-to-CPU ordering by inserting synchronization when the trace shows none; host and device lifetimes are preserved by deferring free until the matching cudaFree; and calls that do not match the plan pass through to the original implementation. Finally, before enabling a loaded plan, DR runs a short profitability calibration against the profiled baseline latency, marking the plan inactive if the optimized run does not beat it by the configured margin. 3.3.2 Three-Phase Pipeline. DR implements this approach in three phases, which we illustrate in Figure 3. Phase 1: Profile. Using LD_PRELOAD, DR intercepts libc and CUDA runtime APIs related to memory allocation, copy, synchronization, and deallocation. For each allocation or copy call, the runtime records the buffer size 𝑁 and a callsite identity 𝛽, a 64-bit FNV-1a hash of the top 𝑑 return addresses on the call stack (e.g., 𝑑=16). We write 𝛽ℎ for the

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies 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, β=0xF7A, 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

Pass 1: H2D Pairs

64-bit hash β

• n ≥ θmin • Pair # ≥ kmin (steady-state) • Window consistency checks

d-frame return addr

Copying: cudaMemcpy, cudaMemcpyAsync

Synchronization: cudaDeviceSynchronize, cudaStreamSynchronize

// ④ GPU compute

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

// ⑤ [SKIP]

free(h); cudaFree(p);}

// ⑥ [DEFER]

Load Plan Match incoming API calls by β + size

cudaDevice Synchronize

Emit Plan

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

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

// ③ [SKIP]

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

• Check preceding

Key Invariant

Record

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

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

Pass 2: D2H Pairs

Group by (βsrc, βdst, n)

FNV-1a

…

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

Phase 3: Optimize

Phase 2: Analyze

Phase 1: Profile

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

ONE-TIME

Replace

Skip

malloc

cudaMemcpy H2D

cudaMalloc Managed cudaMalloc

return same ptr

Defer free

skip

D2H

skip +

No match

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

cudaDev iceSync

until cudaFree

Passthrough

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

EVERY RUN

Figure 3. DR three-phase pipeline. Phase 1 (Profile) records allocation and copy events via LD_PRELOAD interception. Phase 2 (Analyze) pairs allocations offline and emits an optimization plan. Phase 3 (Optimize) replaces paired allocations with cudaMallocManaged and skips self-copies at runtime. Phases 1 and 2 incur only one-time overhead during profiling. host allocation site, 𝛽𝑑 for the device allocation site, and 𝛽𝑐 for the copy site. Walking multiple frames avoids hash collisions caused by C++ wrapper functions that share the same innermost return address (§4.1.2). A size threshold 𝜃 min (e.g., 200 KB) filters out small transient buffers at recording time. The runtime writes all recorded events to a trace file, which Phase 2 consumes as input. During this profiling phase, DR checks whether the target program violates the isolation conditions S1 and S2 of Definition 2. Recall that these forbid host access to the buffer between the H2D and D2H copies, the interval in which the separate buffers h and d may hold different values. With a single unified buffer, such an access would observe or produce a different value and change the program’s behavior. SR discharges S1 and S2 statically, but DR sees only API calls and cannot, so it instead checks them at runtime with an API-delimited page-protection trick. Specifically, right after a matched H2D copy has completed, DR marks the host buffer pages inaccessible (PROT_NONE). The protected window is closed at the next boundary, either before the matching D2H copy runs, or after an explicit device synchronization that closes an H2D-only device use interval. Any host read or write while the window is open faults, and that fault is exactly the signal DR is after: it means the program touched the host buffer while the copy was supposed to keep h and d isolated, i.e., an S1 or S2 violation. A pair that faults in any window is rejected, and pairs that survive are retained in the plan. For asynchronous copies, this check is applied only when completion is ordered by synchronization visible in the trace. Otherwise, the pair is rejected.

Phase 2: Analyze. An offline analyzer scans the trace and identifies candidate (h, d) pairs to unify. The analyzer targets only steady-state iterations that repeat the same allocation and copy call sites, to avoid optimizing one-time transfers such as model-weight loading during initialization. The analyzer makes two passes, one for H2D and one for D2H copies, since each direction requires separate matching logic. A triple (𝛽ℎ , 𝛽𝑑 , 𝑁 ) is accepted as a candidate only if: 1. It appears in at least 𝑘 min steady-state iterations (e.g., 𝑘 min =2), confirming it is a recurring copy rather than a one-time transfer. 2. All copies between this host-device pair originate from the same call site 𝛽𝑐 . Multiple distinct copy sites indicate that different code paths copy between the same buffers, which DR cannot unify under this contract. 3. The copy/kernel region for the pair is effectively serialized: the trace uses the default stream or explicit synchronization that establishes the same H2D→kernel→D2H order. Candidate pairs involving multiple streams, events, or overlapped copy/compute are rejected and left out of the plan. 4. The CUDA context was already active when the host allocation occurred. DR replaces malloc with cudaMallocManaged, which requires an active CUDA context. If the context did not exist at that point in the original program, the replacement would fail. Triples that pass all checks are recorded as candidates for Phase 3. Additionally, for each device-to-host copy in the

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

candidate list, the analyzer records whether an explicit synchronization already precedes it in the trace, because eliminating the D2H copy also removes the implicit GPU-to-CPU synchronization it provided. Phase 3 uses this flag to decide whether to insert a compensating cudaDeviceSynchronize. Phases 1 and 2 are a one-time cost. Once the candidate list is produced, it is reused for all subsequent runs without re-profiling. Phase 3: Optimize. On each subsequent run, DR loads the candidate list and intercepts the same allocation and copy call sites identified during profiling. For each accepted triple (𝛽ℎ , 𝛽𝑑 , 𝑁 ), DR applies three rules: Allocation replacement. When a malloc(N) whose call-site hash matches 𝛽ℎ is intercepted, DR calls cudaMallocManaged(N) instead and returns the unified pointer p. When the corresponding cudaMalloc(N) with hash 𝛽𝑑 is intercepted, DR returns the same p. If cudaMalloc arrives before malloc, DR holds p until the paired malloc occurs. Copy skipping. When cudaMemcpy(dst, src, N) is intercepted, DR checks whether dst == src and the pointer belongs to a unified allocation created in the step above. If so, the copy is a self-copy and is skipped. This pointer-identity check is the sole skip criterion; the call-site hash 𝛽𝑐 is deliberately not used because 𝛽 collisions can occur when different cudaMemcpy calls pass through the same wrapper chain. For skipped D2H copies, DR inserts a cudaDeviceSynchronize when the plan indicates no prior explicit synchronization, preserving the GPU-to-CPU ordering guarantee. Deallocation deferral. When free is called on a pointer that DR replaced (now p), deallocation is deferred. The buffer is freed only when the matching cudaFree arrives (originally targeting d, which after unification points to p). 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 implementation. This pass-through behavior keeps DR scoped to the profiled deployment contract.

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 [21]). 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.

Table 3. Benchmark suite. Pairs: number of candidate allocation pairs in the steady-state optimization plan. Per-platform 𝑓copy values are listed in Appendix Table 9. Not all copy directions within a pair are necessarily optimized; for example, 𝐵 M2 ’s D2H copy is rejected for safety (§4.3). ID

Workload

Category

Pairs

Opt.

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

cuBLAS SGEMM (20482 )

micro micro micro canonical canonical end-to-end end-to-end negative

2 2 2 2 2 1 2 0

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

cuFFT OFDM (2048 × 1024) cuDNN Conv (ResNet-50 layer 1) FCN-ResNet50 inference FCN-ResNet101 inference YOLOv5 object detection DeepSORT tracking C1 with manually optimized

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 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. 4.1.3 Benchmarks. Table 3 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 [20]. • Canonical pipelines (𝐵 C1 –𝐵 C2 ) run TensorRT inference (FCN-ResNet50 and FCN-ResNet101) with preprocessing, H2D, inference, D2H, and postprocessing stages. • End-to-end applications (𝐵 E1 –𝐵 E2 ) use unmodified thirdparty codebases (YOLOv5-TensorRT [27], DeepSORTTensorRT [22]). The critical copy paths (i.e., the H2D and D2H copies) reside inside the closed-source TensorRT runtime, so these benchmarks are evaluated with DR 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.3). 4.1.4 Methodology. We compare four methods on each benchmark. (a) The baseline is the original unmodified program using the standard malloc+cudaMalloc+cudaMemcpy pattern. (b) DR and (c) SR are the two DejaVu 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.

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

Baseline

SR

DR

6

7

• Copy fraction (𝑓copy ) is the share of baseline iteration time spent inside cudaMemcpy calls, computed as the GPUside copy time reported by Nsight Systems divided by 𝑡 baseline . Because 𝑓copy measures all copies in the baseline, including any that DejaVu cannot intercept (e.g., TensorRT-internal copies), the Amdahl approximation 𝑆 = 1/(1−𝑓copy ) provides an upper bound on the speedup achievable by copy elimination. • Recovered fraction = (𝑡 baseline −𝑡 DejaVu )/(𝑡 baseline −𝑡 manual ), where 𝑡 DejaVu is the latency of either DR or SR, and 𝑡 manual is the latency of the manually optimized variant. This measures how close DejaVu’s automatic transformation comes to the manual upper bound.

Manually Optimized

B

BM1 SD M

BM2

B S D M

BM3

B S D M

BC1

B S D M

BC2

B S D M

5

0

BE1 BD BE2 BD

1

1

11

13

9

2

3

3

4

2

12

8

14

15

10

4

5

5

16

17

6

18

19

4.2

19

20

21

22

23

24

25

26

27

18

19

20

21

22

23

24

25

26

5.5

6.0

6.5

7.0

7.5

8.0

Iteration latency (ms)

8.5

9.0

Note: x-axis scale differs per panel.

Figure 4. Per-iteration latency on Jetson AGX Orin. Each row uses an independent x-axis scale. Box: mean±1𝜎. Whisker: min to p99. Red line: mean. 𝐵 E1 –𝐵 E2 show only Baseline and DR because their copy paths are inside the closed-source TensorRT runtime (§4.2).

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 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 verify correctness by comparing the relative L2 norm of each benchmark’s primary output tensor against the baseline. A method passes if the relative error is below 10−5 . All runs across all methods and platforms pass the check. 4.1.5 Metrics. Let 𝑡 method denote the mean wall-clock iteration latency of a given method. We use three metrics throughout the evaluation: • Speedup = 𝑡 baseline /𝑡 method , which represents the overall performance improvement achieved by DejaVu’s transformation (and the manual optimization).

Per-Benchmark Results

Figure 4 shows per-iteration latency on AGX Orin for all seven benchmarks (full numerical results in Appendix Table 4; Nano and Thor in Figures 8 and 9). The main trend is that DejaVu 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 end-to-end applications see smaller but still useful gains when copies are a smaller or partially hidden part of the application. Micro-benchmarks (𝐵 M1 –𝐵 M3 ). These benchmarks isolate the impact of copy elimination on individual CUDA kernels. 𝐵 M2 (cuFFT) achieves a 6.91× speedup because its baseline iteration is heavily dominated by data movement (𝑓copy = 86%). Eliminating redundant copies leaves only the 0.7 ms FFT kernel. As the compute workload increases in 𝐵 M3 (cuDNN) and 𝐵 M1 (cuBLAS), the speedups naturally scale down to 3.40× and 1.70×, closely tracking the available copy fraction. Across all three, both DR and SR match the performance of hand-optimized code to within measurement noise (≥99% recovery). 𝐵 M1 ’s recovered fraction slightly exceeds 100% (103.3%) on both Orin and Nano, but the absolute difference is only 0.12 ms, comparable to the run-to-run standard deviation of 0.04 ms. We attribute this to minor differences in CUDA driver page-placement decisions, which can vary between execution paths even when both kernels use cudaMallocManaged. 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, DejaVu scales to these multi-stage pipelines without degrading performance relative to manual optimization; the latency gap between DejaVu and the manual baseline is less than 0.5%.

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

8

Micro Canonical End-to-end

Amdahl S = 1/(1

7

6 5 4 3 2 1 0

R 2 = 0.98 (20 points, 3 platforms)

8

Observed speedup

Speedup (×)

7

Orin Nano Thor

Nano

6

Thor

5 4 3 2

M2 (Blackwell L2)

1

BM1

BM2

BM3

BC1

BC2

BE1

Figure 5. DR speedup across platforms and benchmarks. Speedup scales with copy fraction and varies across platforms due to differences in memory bandwidth and compute throughput. 𝐵 E1 is absent on Thor (codebase does not compile under the pre-release JetPack). SR achieves equivalent speedups on source-available benchmarks (see Figure 4).

End-to-end applications (𝐵 E1 –𝐵 E2 ). These benchmarks show the benefit of DR in settings where source-level rewriting is unavailable. Both use unmodified third-party code with critical copies inside the closed-source TensorRT runtime, so neither SR nor manual optimization can reach the same path. 𝐵 E2 (DeepSORT) achieves a 1.14× speedup, nearly reaching its 1.16× Amdahl bound, confirming that DR intercepts and eliminates the majority of API-visible copies. 𝐵 E1 (YOLOv5) yields a 1.10× speedup. Although the TensorRT engine internally executes multiple proprietary memory transfers that bypass standard CUDA runtime interposition and inflate the theoretical Amdahl bound to 1.42×, DR still optimizes the application’s primary API-visible bottleneck. By automatically converting the 19.7 MB host-to-device input upload to unified memory, DejaVu achieves a 10% performance gain on a binary-only inference pipeline without requiring source-code access or recompilation. 4.2.1 Cross-Platform Validation. Figure 5 shows DR speedup across all three platforms. The same optimization yields different gains because 𝑓copy changes with each platform’s compute and memory bandwidth. Nano spends more time in compute than Orin, so copies form a smaller share of the iteration. Thor’s higher memory bandwidth similarly reduces baseline copy time, leaving less room for copy elimination to improve wall-clock latency. DejaVu accelerates only the memory copy portion of the pipeline, so its speedup is bounded by Amdahl’s law: 𝑆 Amdahl = 1/(1 − 𝑓copy ). Figure 6 plots 𝑓copy against observed speedup for all benchmarks across three platforms (20 data points), with this curve overlaid. The strong fit (𝑅 2 = 0.98) shows that copy elimination, not benchmark identity, explains most of the observed benefit. The deviations are also

M3 (Blackwell L2)

0

BE2

fcopy)

Orin

20

R1 (Partial)

40

60

80

100

Wall-clock copy fraction f copy (%)

Figure 6. Observed DR speedup vs. copy fraction (𝑓copy ) across 20 data points on three platforms. The dashed curve is the Amdahl approximation 𝑆 = 1/(1 − 𝑓copy ). informative. 𝐵 E1 on Orin and Nano falls below the curve because DR can intercept only the API-visible subset of copies in the closed-source TensorRT runtime. On Thor, 𝐵 M2 and 𝐵 M3 fall below the speedup predicted by their 𝑓copy values, and 𝐵 M3 is the only raw enabled-plan regression (0.91×). We found that, on this pre-release Blackwell platform, some kernels operating on managed-memory buffers run slower than the same kernels on explicitly allocated buffers, with up to 105% kernel-time overhead. Compute-heavy benchmarks such as 𝐵 C1 and 𝐵 C2 are not affected. We report this raw result to show that, on this platform, enabling the plan can be slower. In deployment, DR’s profitability guard disables such an unprofitable plan, so subsequent executions remain on the original (pass-through) path. Appendix D characterizes the effect in detail. Excluding the four explained outliers, the remaining 16 points fit the Amdahl curve with 𝑅 2 = 0.996. 4.3

Boundary Validation

We validate boundary behavior separately for the two optimization paths because their guarantees differ. For SR, all 19 regression tests covering S1–S5 pass, and all benchmarks with D2H pairs produce correct output after synchronization insertion (L2 checksum < 10−5 ). On real benchmarks, the analyzer also rejects unsafe patterns, such as 𝐵 M2 ’s D2H-only output buffer. For DR, the negative control (𝐵 NC ) confirms that the profiler finds zero candidate pairs when the binary already uses cudaMallocManaged. Full details, including percheck test cases and the 𝐵 M2 rejection analysis, are presented in Appendix C. We also test DR under four workload-divergence scenarios. When the workload matches the profile, DR delivers full optimization. When it diverges completely (different binary), no call-site and size matches trigger, so every CUDA call passes through to the baseline runtime with no crash and no data corruption. The plan is also stable across input variations and reduced profiling iterations, because the allocation pattern is structural rather than statistical. These tests

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

Energy / iter (mJ)

Energy / iter (mJ)

Energy / iter (mJ)

Orin 15%

600 400 200 0

400 300 200 100 0

25% 9%

30%

BM1

86%

69%

BM2

BM3

12%

BC1

Nano

BC2

14% 14%

BM1

7% 79%

49%

BM2

BM3

7%

BC1

Thor

BC2

BE1

BE2

9% 23%

45%

BM1

5

17%

50 0

BE2

8%

150 100

BE1

BM2

Baseline

1%

-5%

BM3

BC1

DR

BC2

BE2

SR

Figure 7. Per-iteration energy (mJ). Percentages show DR savings over baseline. Each subplot has its own y-axis. 𝐵 E1 is absent on Thor (pre-release JetPack incompatibility); SR is absent on 𝐵 E1 –𝐵 E2 (closed-source TensorRT).

are boundary checks for the profiled deployment contract; they do not claim that DR proves arbitrary binary semantics beyond the eligibility rules in §3.3.1. 4.4

mean over a dedicated profiling run for each condition. Figure 7 shows results for all three platforms. Energy savings on Orin range from 9% (𝐵 E1 ) to 86% (𝐵 M2 ) and scale with copy fraction. On most benchmarks, we found that instantaneous GPU power typically increases by 6–28% under DejaVu because the GPU no longer idles during DMA copies, but total energy per iteration still decreases because iterations complete faster. On 𝐵 M2 , the pattern reverses. Instantaneous power decreases by 1–7% under DR because the eliminated copies were themselves drawing substantial DMA engine power, producing the largest energy savings (79–86%) across all benchmarks. Nano follows the same trend with smaller savings, consistent with its lower 𝑓copy . On Thor, 𝐵 M3 shows a small energy increase due to the Blackwell L2 overhead.

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 0.88 s; the 5.58 ms per-iteration latency saving recovers that cost within 156 optimized iterations, or about 5.2 s of operation. Allocation coalescing also reduces memory footprint by removing duplicate buffers. The savings match the eliminated host/device pair sizes: 32 MB on 𝐵 M1 and 𝐵 M2 , 29 MB on 𝐵 M3 , 24 MB on 𝐵 C1 and 𝐵 C2 , 19 MB on 𝐵 E1 , and 2–3 MB on 𝐵 E2 . 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 compute per-iteration energy as 𝐸 = 𝑃mean × 𝑡, where 𝑃mean is the mean GPU power from Jetson’s tegrastats utility and 𝑡 is the mean iteration latency. Power is sampled via tegrastats at 100 ms intervals; 𝑃mean is the

Limitations and Future Work

DejaVu’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. 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 coalescing 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. When a runtime execution does not match the profiled plan, DR simply falls back to the original path, and any plan that turns out to be unprofitable is switched off after calibration. API coverage. DejaVu-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 [19] and trtexec [21], 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. 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 (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. Tis

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

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. DejaVu’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 [23] provide hardware cachecoherency via NVLink-C2C, largely obviating software-level copy elimination for many allocations. On edge SoCs lacking such interconnects, however, the legacy dGPU-style code pattern remains pervasive. Adapting DejaVu 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. [26] 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. [25] 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. DejaVu-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 [23] eliminates explicit data transfers for system-allocated memory via hardware-level page-table sharing over NVLink-C2C, which is a platform-specific hardware feature not available 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 DejaVu-SR does. DejaVu-SR operates on optimized LLVM IR and checks alias relationships, control-flow dominance, and pointer-escape properties before transforming each pair. Source refactoring also cannot reach copies hidden in closed-source runtimes such as TensorRT, whereas DejaVu-DR can apply a guarded APIlevel 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, DejaVu is the first system to combine a source-available path that proves semantic safety for this allocation-coalescing transformation with a guarded binary path for stable, API-visible legacy CUDA patterns.

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

7

Conclusion

DejaVu 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. DejaVu’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 DejaVu’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] NVIDIA. 2026. cuBLAS: CUDA Basic Linear Algebra Subroutines. CUDA Toolkit Documentation. https://docs.nvidia.com/cuda/cublas/. [15] NVIDIA. 2026. cuDNN: CUDA Deep Neural Network Library. NVIDIA Developer. https://developer.nvidia.com/cudnn. [16] NVIDIA. 2026. cuFFT: CUDA Fast Fourier Transform. CUDA Toolkit Documentation. https://docs.nvidia.com/cuda/cufft/. [17] NVIDIA. 2026. Jetson AGX Orin Series. NVIDIA Embedded Systems. https://www.nvidia.com/en-us/autonomous-machines/ embedded-systems/jetson-orin/. [18] NVIDIA. 2026. Nsight Systems. NVIDIA Developer Tools. https: //developer.nvidia.com/nsight-systems. [19] NVIDIA. 2026. NVIDIA DeepStream SDK Developer Guide. https: //developer.nvidia.com/deepstream-sdk. [20] NVIDIA. 2026. TensorRT C++ Samples. GitHub. https://github.com/ NVIDIA/TensorRT/tree/main/samples. [21] NVIDIA. 2026. TensorRT: High-Performance Deep Learning Inference. NVIDIA Developer. https://developer.nvidia.com/tensorrt. [22] RichardoMrMu. 2026. deepsort-tensorrt: DeepSORT with TensorRT in C++. GitHub. https://github.com/RichardoMrMu/deepsort-tensorrt. [23] 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 [24] 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 [25] 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 [26] 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 [27] Xinyu Wang. 2026. tensorrtx: Implementation of popular deep learning networks with TensorRT. GitHub. https://github.com/wang-xinyu/ tensorrtx. [28] 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

Baseline

SR

DR

Manually Optimized

Baseline

B

BM3

B S D M

BC1

B S D M

13

14

15

16

17

18

19

M

20

BM2 1

2

3

3

4

28

4

5

30

5

6

6

32

7

7

34

8

8

36

38

BM3

B S D M

BC1

B S D M

BC2

B S D M

BE2

B D

9

40

42

B M

42

BE2

B D

44

30

46

32 10

48 34

11

50 36

12

52 38

13

54 40 14

Iteration latency (ms)

56 42 15

B S D M

0.0

BC2 SD BE1 BD

Manually Optimized

BM1 SD

M

BM2

DR

B

BM1 SD B S D M

SR

58 44 16

3.0

0.2

3.5

0.4

1.0

3.0

4.0

0.6

0.8

1.5

3.5

5.5

1.7

1.8

1.0

1.2

2.0

4.0

5.0

4.5

4.5

5.0

6.0

6.5

1.9

2.0

5.0

1.4

1.6

2.5

3.0

5.5

6.0

6.5

7.0

7.5

8.0

2.1

Iteration latency (ms)

2.2

2.3

Note: x-axis scale differs per panel.

Note: x-axis scale differs per panel.

Figure 8. Per-benchmark results on Jetson Orin Nano. Same conventions as Figure 4. Nano shares Orin’s Ampere SM 8.7 GPU but has fewer CUDA cores, so compute takes a larger share of iteration time and 𝑓copy is lower (Table 9). Speedups are smaller than on Orin.

A

Cross-Platform Results

Tables 4–9 and Figures 8–9 report the complete DR and SR results for all three platforms. The main text (§4.2) 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 raw enabled-plan measurements show five netpositive benchmarks and one regression (𝐵 M3 : 0.91×). The 𝐵 M3 regression is caused by the Blackwell L2 cache readprefetch divergence (§D). Compute-dominated inference workloads (𝐵 C1 : 1.24×, 𝐵 C2 : 1.09×) are unaffected. 𝐵 C1 ’s recovered fraction (114.7%) exceeds 100% because DR slightly outperforms the manually optimized variant on Thor (3.433 ms vs. 3.538 ms). We attribute this to the same CUDA driver pageplacement variability observed on Orin for 𝐵 M1 (§4.2); the absolute difference (0.1 ms) is within the manual variant’s run-to-run standard deviation (0.153 ms).

Figure 9. Per-benchmark raw enabled-plan results on Jetson AGX Thor (Blackwell SM 11.0). Same conventions as Figure 4. 𝐵 E1 is absent because its codebase does not compile under the pre-release JetPack on Thor. On 𝐵 M2 and 𝐵 M3 , the optimized candles are closer to baseline due to a Blackwell-specific managed-memory L2 read-prefetch divergence (§D); 𝐵 M3 regresses to 0.91× in this raw run. SR Results on Nano and Thor. SR produces the same analysis decisions on all platforms because it operates on platform-independent LLVM IR. SR on Thor was compiled with LLVM 18 (Ubuntu 24.04); the New Pass Manager API is compatible across LLVM 15–18 without source changes. Power and Energy on Nano. Across the benchmarks where both apply, DR and SR agree within 2% in per-iteration energy. Power on the measured rail typically rises by 6– 28% under DejaVu, but total energy still decreases because iterations finish sooner. 𝑓copy Across Platforms. Despite the cross-platform variation in 𝑓copy , all 20 data points follow the same Amdahl curve (§4.2.1, 𝑅 2 = 0.98).

B

Memory Footprint and Break-Even

Memory footprint. Each eliminated pair removes one duplicate buffer from resident set size (RSS), as the host and device allocations collapse into a single managed buffer. 𝐵 M1 and 𝐵 M2 (two 17 MB pairs each) save 32 MB. 𝐵 M3 (input 4.8 MB + output 25.7 MB) saves 29 MB. 𝐵 C1 and 𝐵 C2 (input

DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs

Table 4. DR numerical results on Orin (5-run), complementing the DR candles in Figure 4. p99 columns quantify tail latency improvement. SR runtime numbers are reported separately in §4.2. ID

Baseline (ms)

DR (ms)

Speedup

Recovered

BL p99

DR p99

Δp99

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

9.222 ± 0.039 4.316 ± 0.012 5.047 ± 0.034 17.568 ± 0.107 25.461 ± 0.067 20.954 ± 0.162 7.779 ± 0.016

5.409 ± 0.004 0.624 ± 0.004 1.485 ± 0.005 11.989 ± 0.040 19.935 ± 0.020 18.963 ± 0.155 6.850 ± 0.124

1.70× 6.91× 3.40× 1.47× 1.28× 1.10× 1.14×

103.3% 99.5% 99.6% 99.6% 99.6% n/a n/a

9.33 4.46 5.26 18.26 25.75 22.16 8.44

5.46 0.68 1.59 12.25 20.33 19.62 7.51

−3.86 −3.78 −3.66 −6.01 −5.42 −2.54 −0.93

Table 5. DR results on Nano (5-run). Recovered fraction ranges from 97.6% to 102.9%, within 3% of Orin across all benchmarks. ID

Baseline (ms)

DR (ms)

Speedup

Managed (ms)

Recovered

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

18.717 ± 0.011 6.753 ± 0.019 8.310 ± 0.015 37.214 ± 0.141 52.895 ± 0.156 34.962 ± 0.259 13.680 ± 0.026

13.349 ± 0.007 1.437 ± 0.006 3.282 ± 0.002 29.818 ± 0.040 45.503 ± 0.062 32.204 ± 0.068 12.517 ± 0.037

1.40× 4.70× 2.53× 1.25× 1.16× 1.09× 1.09×

13.502 ± 0.020 1.409 ± 0.015 3.268 ± 0.008 29.642 ± 0.050 45.318 ± 0.066 n/a n/a

102.9% 99.5% 99.7% 97.7% 97.6% n/a n/a

Table 6. DR raw enabled-plan results on Thor (5-run). Highlighted rows are affected by the Blackwell L2 read-prefetch divergence (§D). 𝐵 E1 absent: its third-party codebase does not compile under the pre-release JetPack on Thor. ID

Baseline (ms)

DR (ms)

Speedup

Managed (ms)

Recovered

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

3.413 ± 0.008 0.546 ± 0.015 1.098 ± 0.006 4.245 ± 0.018 5.905 ± 0.027 1.763 ± 0.003

3.108 ± 0.007 0.325 ± 0.007 1.206 ± 0.073 3.433 ± 0.011 5.396 ± 0.024 1.753 ± 0.009

1.10× 1.68× 0.91× 1.24× 1.09× 1.01×

3.092 ± 0.006 0.299 ± 0.007 1.277 ± 0.036 3.538 ± 0.153 5.373 ± 0.019 n/a

95.1% 89.5% 60.2% 114.7% 95.7% n/a

Table 7. SR runtime on Nano and Thor (5-run). On Nano, SR speedups match DR within 3%. On Thor, SR can exceed DR by a wider margin on sub-millisecond benchmarks due to zero runtime overhead (e.g., 𝐵 M2 : SR 2.05× vs. DR 1.68×). Highlighted rows are affected by the Blackwell L2 readprefetch divergence (§D). Nano

Thor

ID

SR (ms)

Speedup

ID

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2

13.501 ± 0.012 1.420 ± 0.011 3.279 ± 0.010 29.738 ± 0.100 45.501 ± 0.046

1.39× 4.76× 2.53× 1.25× 1.16×

𝐵 M1 𝐵 M2 𝐵 M3 𝐵 C1 𝐵 C2

SR (ms)

Speedup

3.082 ± 0.003 0.267 ± 0.006 1.140 ± 0.017 3.449 ± 0.023 5.405 ± 0.027

1.11× 2.05× 0.96× 1.23× 1.09×

3 MB + output 21 MB each) save 24 MB each. 𝐵 E1 saves 19 MB (one pair) and 𝐵 E2 saves 2–3 MB (small feature-extraction output).

Table 8. Per-iteration energy on Nano (VDD_CPU_GPU_ CV sensor). Savings range from 7–79%, consistent with Orin trends but uniformly lower due to lower 𝑓copy . ID

BL (mW)

DR (mW)

BL (mJ)

DR (mJ)

Saved

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

7820 2760 4440 6784 7891 3729 4082

9384 2732 5682 7290 8394 3762 4143

146.4 18.6 36.9 252.5 417.4 130.4 55.8

125.3 3.9 18.6 217.4 381.9 121.2 51.9

14.4% 78.9% 49.5% 13.9% 8.5% 7.1% 7.1%

Break-even. Since profiling overhead is less than 1%, the dominant one-time cost is the profile run itself, which executes 50–100 iterations at baseline speed. For 𝐵 C1 at a 30 FPS operating point, profiling 50 iterations takes approximately 0.88 seconds. The per-iteration latency saving of 5.58 ms means that the profiling cost is recovered within 156 optimized iterations, or about 5.2 seconds of operation.

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

Table 9. Wall-clock copy fraction across three platforms. 𝑓copy varies with both GPU compute throughput and memory bandwidth: Orin’s stronger GPU makes compute faster relative to copies, raising 𝑓copy ; Thor’s much higher memory bandwidth shrinks copy time disproportionately, lowering 𝑓copy .

C

ID

Orin 𝑓copy

Nano 𝑓copy

Thor 𝑓copy

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

41.0% 86.2% 69.5% 30.2% 21.0% 29.3% 13.8%

28.8% 78.8% 59.6% 20.5% 14.3% 25.3% 10.7%

8.2% 57.0% 22.6% 9.8% 7.6% — 3.1%

Safety and Boundary Validation

M2 safety rejection. 𝐵 M2 (cuFFT) uses separate input and output buffers with the following per-iteration pattern: memcpy(host_in, signal, N); // refill input cudaMemcpy(dev_in, host_in, N, H2D);// upload cufftExec(dev_in, dev_out, FORWARD); // FFT cudaMemcpy(host_out, dev_out, N, D2H);

The input pair (host_in, dev_in) is safe. The output pair (host_out, dev_out) is an output-only pair: it is linked by a D2H copy, but there is no corresponding H2D transfer that establishes device input for the shared buffer. DejaVu’s current optimization rule does not handle this case, so both DR and SR leave it unchanged. DR’s analyzer finds no H2D copy site linking this pair, and SR finds no H2D cudaMemcpy connecting the two allocations within the function body. D2H synchronization insertion. Eliminating a D2H cudaMemcpy removes the implicit GPU→CPU synchronization it provided: kernel<<<...>>>(dev); // baseline: cudaMemcpy(host, dev, N, D2H) -- implicit sync // DejaVu: cudaDeviceSynchronize(); -- explicit sync

When no explicit cudaDeviceSynchronize precedes the copy in the trace, the optimizer inserts one. In all five benchmarks with D2H pairs (𝐵 M1 , 𝐵 M3 , 𝐵 C1 , 𝐵 C2 , 𝐵 E2 ), a compensating sync is inserted. All produce correct output (L2 checksum < 10−5 , recovered fraction ≥97%). Workload robustness. T1 is a control using the same program for profiling and optimization. T2 profiles 𝐵 C1 but optimizes a different binary (𝐵 M1 ), testing cross-program divergence. T3 varies the input image. T4 profiles with only 10 iterations instead of 50. When the workload matches the profile (T1, T3, T4), DR recovers the same optimization plan and the same speedup. When it diverges (T2), no call-site hash matches and DR runs at baseline speed with no crash or corruption.

Table 10. Thor-only managed-memory effect. Ratios compare managed to explicit allocation; Orin stays at 1.0× in all probes. Probe

Thor result

Read/write Reads slow down 16.8×; writes do not. Size Peak at 16 MB (1.71×); gone by 24 MB. Compute Overhead vanishes at ≥200 FMA/element. Stride Large strides reach 10.7× at ≥128 elements. Repeat Stable over 100 launches; not a warmup effect.

D

Blackwell Managed-Memory Overhead

With a forced optimization, 𝐵 M3 (cuDNN convolution) exhibits a 0.91× regression on Thor (Blackwell, SM 11.0). We report this raw behavior to characterize the hardware effect; in deployment mode, DR’s profitability guard disables this unprofitable plan and uses the pass-through path. Nsight Systems [18] CUPTI profiling indicates that the cuDNN kernel uses the same launch configuration (registers, shared memory, grid/block) on both Orin and Thor. On Orin, managed and explicit buffers produce the same kernel execution time (ratio 1.00×). On Thor, however, the same kernel runs 1.80× slower when operating on cudaMallocManaged buffers compared to cudaMalloc buffers. We isolate the effect with a self-contained CUDA program and characterize it through five targeted microbenchmarks (Table 10). The microbenchmark evidence (Table 10) is consistent with a difference in how Blackwell’s L2 handles managedmemory reads. The slowdown is read-only and size-dependent: it appears once the buffer exceeds 128 KB and disappears again once the working set exceeds L2 capacity, where misses dominate for both allocation types. It also disappears when the kernel becomes compute-bound (≥50 FLOP/byte), at which point the memory system is no longer the bottleneck. Orin shows no such size dependence, which argues against a simple L2-capacity explanation. This overhead appears in bandwidth-sensitive kernels with non-sequential access patterns (𝐵 M2 : +105% kernel time, 𝐵 M3 : +79.6%), but we do not observe a measurable penalty on the evaluated TensorRT inference pipelines (𝐵 C1 : 1.24×, 𝐵 C2 : 1.09×) or on the compute-bound micro-benchmark 𝐵 M1 (1.10×). JetPack 7.0-b128 is a pre-release SDK, and the L2 policy may evolve in future releases.

Related documents

• DejaVu: Unifying Memory Allocations to Eliminate Redundant Copies on Unified-Memory SoCs#972687
• Random access and high dimensional integrated quantum memory#1128514
• Transparent Interconnection of Lots of Links (TRILL) Operations, Administration, and Maintenance (OAM) Framework#399053
• Refreshing a Protected Quantum State Extends Its Survival Time on IBM Kingston: A Delay Sweep with Matched Controls#166490
• Dense auto-hetero associative memories applied to noisy communication channels#1116936