Pinned Memory with cudaMallocHost
Page-locked buffers for fast DMA.
Pinned Memory with cudaMallocHost is a free CUDA Academy lesson on CoddyKit — lesson 2 of 4. You can read the complete lesson below for free — then practise it hands-on in the browser with a built-in code editor and a 24/7 AI tutor. It is part of the CUDA Academy learning path, one of 4 lessons in the course, and your progress syncs across the web and the CoddyKit app.
Meet Pinned Memory
Pinned memory is host memory the operating system promises never to move or swap out. Because its physical address is fixed, the GPU can read it directly.
Also Called Page-Locked
You will see the term page-locked used as a synonym for pinned. Both mean the same thing: pages locked in place so no swapping can ever happen.
Allocate It Right
Ask CUDA for pinned memory with cudaMallocHost. It hands back an ordinary host pointer you can read and write just like any C++ array.
float* h_data;
cudaMallocHost(&h_data, N * sizeof(float));Use It Like Normal
The returned pointer behaves like any host buffer, so your CPU code stays unchanged. You just fill it and read it exactly as before.
for (int i = 0; i < N; ++i) h_data[i] = i * 1.0f;No More Staging
Since the pages cannot move, the driver skips its hidden staging copy. The GPU pulls straight from your buffer, so the doubled copy disappears.
Faster Transfers
The payoff is real bandwidth. Pinned transfers often run near the link's peak speed, sometimes roughly twice as fast as the pageable version.
Free It Correctly
Pinned memory needs its own matching release. Use cudaFreeHost, never the plain free, or you risk a crash and a leak.
cudaFreeHost(h_data);It Is Not Free Real Estate
Pinned pages cannot be swapped, so they hold real RAM hostage. Pin too much and you starve the rest of the system of usable memory.
Allocation Is Slower
Locking pages takes work, so cudaMallocHost is slower to allocate than malloc. Reuse one pinned buffer across many transfers instead of reallocating.
The Real Superpower
Beyond speed, pinned memory is the key that unlocks truly asynchronous copies. Without it, cudaMemcpyAsync cannot overlap with anything.
When to Reach for It
Pin buffers that you transfer often or that feed streams. For a one-off copy of a tiny array, plain malloc is perfectly fine. 🙂
Quick Check
You allocated a buffer with cudaMallocHost. How should you release it?
Recap
cudaMallocHost pins host pages so the GPU reads them directly, giving faster copies and enabling async overlap. Always release with cudaFreeHost. 🔒
Frequently asked questions
Is the “Pinned Memory with cudaMallocHost” lesson free?
Yes — the full text of “Pinned Memory with cudaMallocHost” is free to read here on the web, and the CUDA Academy course includes 4 lessons in total. To practise it interactively (a built-in code editor and a 24/7 AI tutor) and unlock the rest of the CUDA Academy course, upgrade to CoddyKit PRO.
What will I learn in “Pinned Memory with cudaMallocHost”?
Page-locked buffers for fast DMA. You practise CUDA Academy with hands-on code you run directly in the browser, and a 24/7 AI tutor answers your questions as you work through the lesson.
Do I need any experience to start CUDA Academy?
No prior experience is required. CUDA Academy on CoddyKit is structured for beginners through advanced learners; this is — lesson 2 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Pinned Memory with cudaMallocHost” lesson take?
Most CoddyKit lessons take about 5–10 minutes. Each one is bite-sized and interactive, so you make steady progress and pick up exactly where you left off across the web and the app.
Can I write and run code in this CUDA Academy lesson?
Yes. Every CUDA Academy lesson includes a built-in code editor, so you write and run real code right in your browser and get instant AI feedback — no local setup required.
All lessons in this course
- Why Pageable Memory Is Slow
- Pinned Memory with cudaMallocHost
- cudaMemcpyAsync in a Stream
- The Double-Buffering Pipeline