{"id":551,"date":"2026-08-21T08:19:25","date_gmt":"2026-08-21T13:19:25","guid":{"rendered":"https:\/\/davidwdrell.net\/wordpress\/?p=551"},"modified":"2026-08-22T05:23:45","modified_gmt":"2026-08-22T10:23:45","slug":"cuda-zero-copy-on-jetson","status":"publish","type":"post","link":"https:\/\/davidwdrell.net\/wordpress\/?p=551","title":{"rendered":"Cuda Zero-copy on Jetson"},"content":{"rendered":"

My image processing library was written for desktop machines with discrete GPUs. On a Jetson, most of what it was doing with memory was wasted work. Yes you can just re-compile GPU code for Jetson and it runs, but it does not run fast.<\/p>\n

Measured on AGX Orin, JetPack 6, 4112\u00d73008 image planes.<\/em><\/p>\n

I have a C++ image processing library that I have been building for years on Windows and Linux desktops with big NVIDIA GPUs. Now the same library has to run on a Jetson AGX Orin inside a microscope. It compiled and it produced correct images on the first try, but it was far slower than it should have been. The problem was not the kernels. It was the memory copies.<\/p>\n

On a desktop, host memory and device memory are separate physical chips at opposite ends of a PCIe link, so every Cuda wrapper I have ever written has the same pattern: cudaMalloc, cudaMemcpy up, launch the kernel, cudaMemcpy down, cudaFree. That is the correct pattern on a desktop.<\/p>\n

On a Jetson it is not. The CPU and the GPU share the same physical DRAM. Those two copies move bytes from DRAM back to DRAM and end where they started, and each one costs you a synchronization as well.<\/p>\n

Fuse the kernels first<\/h2>\n

Before you do any of this, look at whether you can do less work. Fusing kernels is always the first option. Every kernel you eliminate is a pair of copies you do not do make.<\/p>\n

But you cannot fuse everything. NPP calls, library primitives and genuinely stand-alone kernels will still be there when you are done, and every one of them has a copy-in and copy-out. That is what the rest of this article is about.<\/p>\n

One memory, but not one kind of memory<\/h2>\n

This part confused me for a while, and if you skip it the rest of the article will look like it contradicts itself. There is only one physical DRAM on a Jetson. That is true. But host memory and device memory are still different things, because the difference is not where the bytes are. It is who is allowed to address them, and how they are cached.<\/p>\n

The CPU and the GPU have separate page tables. Memory you get from new[] or malloc exists only in the CPU’s tables, and it is pageable, meaning the operating system is free to move it around underneath you. The GPU has no entry for it and cannot safely touch it. So when you hand that pointer to a kernel, the driver has no choice but to copy the bytes into a region the GPU is allowed to address.<\/p>\n

That is what your cudaMemcpy is actually doing on a Jetson. It is not carrying data across a bus to another chip. It is moving it from pages the GPU cannot see into pages it can, in the same DRAM. cudaMalloc gives you the second kind directly, but then the CPU cannot dereference it. Try it and you get a segfault, even though the bytes are physically sitting right there.<\/p>\n

cudaHostAlloc with the mapped flag is the way out. Those pages are pinned, so the OS cannot relocate them, and they are mapped into both page tables. One set of bytes, two valid addresses, nothing to copy.<\/p>\n

Three allocators, one DRAM:<\/p>\n\n\n\n\n\n\n\n
Allocation<\/th>\nCPU can dereference<\/th>\nGPU can dereference<\/th>\nCached in GPU L2<\/th>\n<\/tr>\n<\/thead>\n
new[] \/ malloc (pageable)<\/td>\nyes<\/td>\nno<\/td>\nn\/a<\/td>\n<\/tr>\n
cudaMalloc (device)<\/td>\nno<\/td>\nyes<\/td>\nyes<\/td>\n<\/tr>\n
cudaHostAlloc(…Mapped)<\/td>\nyes<\/td>\nyes<\/td>\nno<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n

Look at that last column, because it is the one that catches people out. On Tegra, memory from cudaMalloc is cached in the GPU’s L2 the way you would expect. Mapped host memory is not. It is kept coherent with the CPU instead, and that means every GPU access goes all the way out to DRAM.<\/p>\n

For most image work this costs you nothing. A resize or a multiply reads each pixel once, so there is no reuse for a cache to capture and zero-copy wins outright. But it is a disaster for a kernel where every block hits the same few bytes, which is exactly what a device-side atomic counter is. In device memory those hits are absorbed by L2. In mapped memory each one is a separate trip to DRAM.<\/p>\n

So the rule I ended up with is simple: reuse decides the allocator<\/strong>. If each byte is touched once, map it and take the zero-copy path. If a small region is hit over and over by many blocks, leave it in device memory so the cache can do its job. Almost all of an image pipeline is the first case, which is why this works so well.<\/p>\n

1. Detect the memory model at run time<\/h2>\n

The obvious thing to do is add a build flag for Jetson. Do not do this. An aarch64 CPU does not necessarily mean an integrated GPU, and if you branch at compile time you can no longer test both paths on one machine. Ask the driver instead, once, at startup.<\/p>\n

enum class MemoryModel { Unified, Discrete, NoDevice };\n\nMemoryModel Detect(int device)\n{\n    int n = 0;\n    if (cudaGetDeviceCount(&n) != cudaSuccess || n == 0) return MemoryModel::NoDevice;\n\n    cudaDeviceProp p{};\n    if (cudaGetDeviceProperties(&p, device) != cudaSuccess) return MemoryModel::NoDevice;\n\n    return (p.integrated && p.canMapHostMemory) ? MemoryModel::Unified\n                                                : MemoryModel::Discrete;\n}<\/code><\/pre>\n
Include the NoDevice case. Someone will load a TIFF on a build machine with no GPU in it, and if you route every allocation through Cuda you have just broken them.<\/p>\n
2. Change the allocator, not the kernels<\/h2>\n
The GPU can only read host memory in place if that memory was allocated so it can be mapped. This is the one real change to your data path: everywhere an image plane or a tensor was allocated with new[], allocate it mapped instead.<\/p>\n
cudaSetDeviceFlags(cudaDeviceMapHost);   \/\/ before your first allocation\n\n\/\/ Unified  : cudaHostAlloc(&p, n, cudaHostAllocMapped);   \/\/ kernels read it in place\n\/\/ Discrete : cudaHostAlloc(&p, n, cudaHostAllocDefault);  \/\/ pinned, no bounce buffer\n\/\/ NoDevice : ::operator new[](n);<\/code><\/pre>\n
Put this behind an RAII handle that remembers which allocator owns the block, so the buffer knows how to free itself. Note that the desktop path gets faster too, because pinned memory removes the driver’s own staging copy. This is not a Jetson-only branch that nobody maintains.<\/p>\n
Leave the raw pointer public. Mapped memory is ordinary memory as far as the CPU is concerned, so libtiff, your copy constructors and all of your existing CPU loops keep working exactly as they did.<\/p>\n
3. Put one object at the Cuda boundary<\/h2>\n
Now you can delete the plumbing. Every wrapper gets the same object, and that object decides per pointer whether the GPU can read the buffer where it sits or whether it has to stage a copy.<\/p>\n
void AddConstant_f32(float *img, int w, int h, float c)\n{\n    \/\/ deleted:  cudaMalloc(&dev, w * h * 4);\n    \/\/ deleted:  cudaMemcpy(dev, img, w * h * 4, cudaMemcpyHostToDevice);\n\n    DeviceView<float> v(img, (size_t)w * h, Access::InOut);\n\n    nppiAddC_32f_C1IR_Ctx(c, v.device(), w * 4, {w, h}, NppCtxFor(v.stream()));\n\n    \/\/ deleted:  cudaMemcpy(img, dev, w * h * 4, cudaMemcpyDeviceToHost);\n    \/\/ deleted:  cudaFree(dev);\n\n    v.commit();\n}<\/code><\/pre>\n
I converted 41 entry points this way and the library lost about 2300 lines of code. Three things matter here.<\/p>\n
The function signature does not change, so everything that calls your library gets the speedup without being touched. Unregistered pointers must fall back to staging: call cudaPointerGetAttributes, and if the pointer did not come from your allocator, copy it the old way. If you ship a shared library you cannot assume anything about the caller’s buffers. And use the devicePointer the driver gives you back rather than assuming the host address and the device address are the same number.<\/p>\n
Finally, commit() has to synchronize on both paths. Under zero-copy the kernel is writing directly into the caller’s buffer, so the caller must not read it until the stream has drained. This is the one place where zero-copy changes the rules on you.<\/p>\n
4. Some things still need real device memory<\/h2>\n
In practice the reuse rule from earlier turns into two categories, and only one of them is about speed.<\/p>\n
The first is contended atomics. Some of my kernels have every block atomically updating one scalar, a mean or a min or a max. With the image mapped and those scalars in device memory the kernel ran in 2.31 ms. With the scalars mapped as well it took 452 ms, which is 196 times slower<\/strong>. Same kernel, same image, the only difference is where three floats live. Histogram bins and the CLAHE tile lookup table belong in the same category.<\/p>\n
The second reason has nothing to do with speed. NPP scratch buffers and similar library workspace are never read by the CPU at all, so there is nothing to gain by mapping them. Leave them alone. Between the two categories I had 17 sites out of about 137 that stayed device resident.<\/p>\n
Hand those buffers out from a small pool with exact size buckets instead of calling cudaMalloc and cudaFree around each use. Be aware that cudaFree and cudaFreeHost are device-wide synchronization points: the calling thread blocks until the device drains. I measured a cudaFreeHost call taking 28.2 ms to return with a 24 ms kernel resident on another stream. That is not a fixed cost, it is however much outstanding work there happens to be, and image destructors run on whatever thread your consumer is using.<\/p>\n
5. Prove that nothing changed<\/h2>\n
Two things can go wrong here without any test failing. The output can change slightly, and a buffer you thought was zero-copy can quietly stage anyway.<\/p>\n
For the first, hash the output of every wrapper before you convert it, and compare the hash bit for bit afterwards. Do not eyeball images. For the second, count how many times your accessor decided to stage, and assert that the count is zero for buffers that should have been mapped. Without that counter, one forgotten new[] turns the whole pipeline back into a copying pipeline and nothing tells you.<\/p>\n
What I measured<\/h2>\n
Per call, median, 4112\u00d73008 image planes. Geometric mean 10.5\u00d7<\/strong> over 42 operations.<\/p>\n\n\n\n\n\n\n\n\n\n\n
Operation<\/th>\nBefore<\/th>\nAfter<\/th>\nSpeedup<\/th>\n<\/tr>\n<\/thead>\n
Resize f32 \u00d70.125<\/td>\n26.3 ms<\/td>\n0.35 ms<\/td>\n75\u00d7<\/td>\n<\/tr>\n
MultiplyImages f32<\/td>\n72.8 ms<\/td>\n2.67 ms<\/td>\n27\u00d7<\/td>\n<\/tr>\n
addHSV<\/td>\n181.9 ms<\/td>\n9.39 ms<\/td>\n19\u00d7<\/td>\n<\/tr>\n
Remap chain, 7 primitives<\/td>\n224.5 ms<\/td>\n14.0 ms<\/td>\n16\u00d7<\/td>\n<\/tr>\n
CLAHE 8u (compute bound)<\/td>\n161.0 ms<\/td>\n123.4 ms<\/td>\n1.3\u00d7<\/td>\n<\/tr>\n
Flood fill (compute bound)<\/td>\n95.1 ms<\/td>\n76.3 ms<\/td>\n1.2\u00d7<\/td>\n<\/tr>\n<\/tbody>\n<\/table>\n
The spread is worth paying attention to. The operations that were a single pass over an image plane were almost entirely copy overhead, and they got 10 to 75 times faster. The two that are genuinely compute bound barely moved. If everything in your suite improves by the same amount, you measured something other than what you think.<\/p>\n
Also worth knowing: 145 synchronous cudaMemcpy calls came out of the library, and the library ended up smaller than it was before.<\/p>\n
TensorRT buffers<\/h2>\n
If you use TensorRT, look at its buffer manager. It keeps a host buffer and a device buffer for every binding and copies between them on every inference. On a Jetson you can allocate one mapped buffer per binding, bind the device pointer you get from cudaHostGetDevicePointer, and make the copy functions return immediately. Same predictions, one less round trip per frame. I got 11 percent off my inference path doing this.<\/p>\n
Things that bit me<\/h2>\n
cudaMallocManaged looks easier. It is not.<\/strong> On Orin, concurrentManagedAccess is 0. Touching a managed buffer from another thread while the GPU might be using it is not a slow path, it is a segfault. I reproduced this before giving up on it. Mapped host memory has no such restriction.<\/p>\n
Shutdown order now matters.<\/strong> Every mapped buffer has to be freed before the Cuda context goes away. If a buffer outlives the context, log it and leak it. Do not call cudaFreeHost into a dead runtime. A leak at process exit is harmless; that call is undefined behavior.<\/p>\n
 <\/p>\n","protected":false},"excerpt":{"rendered":"
My image processing library was written for desktop machines with discrete GPUs. On a Jetson, most of what it was […]<\/p>\n","protected":false},"author":1,"featured_media":0,"comment_status":"open","ping_status":"","sticky":false,"template":"","format":"standard","meta":{"_monsterinsights_skip_tracking":false,"_uf_show_specific_survey":0,"_uf_disable_surveys":false,"site-sidebar-layout":"default","site-content-layout":"","ast-site-content-layout":"default","site-content-style":"default","site-sidebar-style":"default","ast-global-header-display":"","ast-banner-title-visibility":"","ast-main-header-display":"","ast-hfb-above-header-display":"","ast-hfb-below-header-display":"","ast-hfb-mobile-header-display":"","site-post-title":"","ast-breadcrumbs-content":"","ast-featured-img":"","footer-sml-layout":"","theme-transparent-header-meta":"","adv-header-id-meta":"","stick-header-meta":"","header-above-stick-meta":"","header-main-stick-meta":"","header-below-stick-meta":"","astra-migrate-meta-layouts":"default","ast-page-background-enabled":"default","ast-page-background-meta":{"desktop":{"background-color":"var(--ast-global-color-5)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"tablet":{"background-color":"","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"mobile":{"background-color":"","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""}},"ast-content-background-meta":{"desktop":{"background-color":"var(--ast-global-color-4)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"tablet":{"background-color":"var(--ast-global-color-4)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""},"mobile":{"background-color":"var(--ast-global-color-4)","background-image":"","background-repeat":"repeat","background-position":"center center","background-size":"auto","background-attachment":"scroll","background-type":"","background-media":"","overlay-type":"","overlay-color":"","overlay-opacity":"","overlay-gradient":""}},"footnotes":""},"categories":[1],"tags":[],"class_list":["post-551","post","type-post","status-publish","format-standard","hentry","category-uncategorized"],"_links":{"self":[{"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/posts\/551","targetHints":{"allow":["GET"]}}],"collection":[{"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/posts"}],"about":[{"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/types\/post"}],"author":[{"embeddable":true,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/users\/1"}],"replies":[{"embeddable":true,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=%2Fwp%2Fv2%2Fcomments&post=551"}],"version-history":[{"count":5,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/posts\/551\/revisions"}],"predecessor-version":[{"id":561,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=\/wp\/v2\/posts\/551\/revisions\/561"}],"wp:attachment":[{"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=%2Fwp%2Fv2%2Fmedia&parent=551"}],"wp:term":[{"taxonomy":"category","embeddable":true,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=%2Fwp%2Fv2%2Fcategories&post=551"},{"taxonomy":"post_tag","embeddable":true,"href":"https:\/\/davidwdrell.net\/wordpress\/index.php?rest_route=%2Fwp%2Fv2%2Ftags&post=551"}],"curies":[{"name":"wp","href":"https:\/\/api.w.org\/{rel}","templated":true}]}}