Why Pageable Memory Is Slow
Staging buffers behind ordinary malloc.
Why Pageable Memory Is Slow is a free CUDA Academy lesson on CoddyKit — lesson 1 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.
Where Your Data Lives
Normal C++ buffers from malloc sit in pageable host memory. It feels free, but the GPU cannot copy from it directly, and that small detail costs you speed.
What Pageable Means
Memory is pageable when the operating system may move its pages around or swap them to disk at any moment. Great for flexibility, tricky for hardware that needs fixed addresses.
The GPU Speaks DMA
The GPU pulls data using DMA, a direct hardware transfer that needs a physical address that will not move. Pageable pages break that promise, so they cannot be used directly.
The Hidden Staging Step
To work around this, the driver copies your pageable buffer into a hidden staging buffer first, then sends that to the GPU. You pay for an extra copy you never wrote. 😮
Two Copies, Not One
So a single cudaMemcpy from pageable memory is really two copies: host to staging, then staging to device. That doubled work is exactly why pageable transfers feel sluggish.
An Ordinary malloc
Here is the buffer everyone reaches for first. It works, but every transfer from h_data quietly goes through that staging detour.
float* h_data = (float*)malloc(N * sizeof(float));
cudaMemcpy(d_data, h_data, bytes, cudaMemcpyHostToDevice);Why the OS Cares
The OS keeps memory pageable so it can overcommit RAM and serve many programs at once. That convenience is what blocks the GPU from reading your pages directly.
Bandwidth You Lose
Pageable transfers often reach only half of your link's peak bandwidth. The staging copy eats CPU cycles and memory traffic that real transfers could have used.
It Blocks Overlap Too
Because the staging copy is synchronous, pageable transfers cannot truly overlap with kernels. You lose the async tricks that make streams worthwhile.
When It Still Hurts
For one tiny copy nobody notices. But in a loop that ships data every iteration, the staging tax compounds and quietly dominates your runtime.
The Fix Is Coming
The cure is asking CUDA for a buffer that cannot be paged out, called pinned memory. The GPU reads it directly, no staging, no doubled copy.
Quick Check
Why is a transfer from ordinary pageable memory slower than it needs to be?
Recap
Pageable buffers can move, so the GPU cannot DMA them directly. The driver stages them first, doubling the copy. The fix ahead is pinned memory. 🚀
Frequently asked questions
Is the “Why Pageable Memory Is Slow” lesson free?
Yes — the full text of “Why Pageable Memory Is Slow” 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 “Why Pageable Memory Is Slow”?
Staging buffers behind ordinary malloc. 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 1 of 4, so you can start here or from the beginning and move at your own pace.
How long does the “Why Pageable Memory Is Slow” 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