What is cudaMemcpy?

The blocking copy between host and device, whose direction argument is the first thing beginners get wrong.

The argument order is destination, source, bytes, direction, and it is the same order as memcpy, which is the only mercy in the signature. The two pointers live in different address spaces: one came from malloc or a std::vector, the other from cudaMalloc and points into global memory. Nothing in the type system separates them, both are float*, so the naming convention does the work instead. Day 5 puts h_ on every host pointer and d_ on every device one, which makes a wrong argument visible in the call rather than in a crash three lines later.

The direction has to agree with the pointers you actually passed, and the runtime will not save you when it does not. NVIDIA's wording is flat: "Calling cudaMemcpy() with dst and src pointers that do not match the direction of the copy results in an undefined behavior." Undefined, not an error code, so a swapped cudaMemcpyHostToDevice sometimes returns cudaSuccess and sometimes turns up later as invalid argument from a call that had nothing to do with it. Checking the return value of every call is the only way to bound how far the blame can travel, which is why error checking arrives on day 5 before the kernel does.

It also blocks, and that matters more than it looks. "The cudaMemcpy API is synchronous. That is, it does not return until the copy has completed." A device-to-host copy is therefore a synchronization point, and in a program with no explicit cudaDeviceSynchronize() it is usually the first call that can report what the kernel did. That is why a beginner's first fault appears to come from the copy back. It did not. The copy is just where the bill arrived.

Every copy also crosses PCIe, which is roughly an order of magnitude slower than the GPU's own memory, so on a small problem the two copies cost more than the arithmetic they feed. That is most of the answer to "why is my GPU slower than my CPU", and day 9 puts the copy time and the kernel time side by side for this exact program. When the copies are the bottleneck rather than the setup, pinned memory and cudaMemcpyAsync let them overlap with compute, and unified memory removes the explicit call at the cost of page faults you did not schedule.

Measured

On a Tesla T4, driver 595.84, CUDA 12.6 (V12.6.85), built with nvcc -O3 -arch=sm_75, day 5 copied two input arrays of 611 floats to the device, launched 3 blocks of 256 threads, and copied the result back:

n = 611, 256 threads per block, 3 blocks, 768 threads
157 threads have no element to add
all 611 elements match the CPU reference

Nothing on that run is timed, deliberately. A host clock around a first CUDA program measures context creation and two PCIe copies rather than the kernel, so day 5 proves the program is right and day 9 works out what each part costs.

Code

From code/day05-vector-add/vector_add.cu. The comments are in the shipped file because the ordering here is the lesson.

// Destination, source, bytes, direction. The direction has to agree with
// the pointers; the runtime documents a mismatch as undefined behaviour.
CUDA_CHECK(cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice));
CUDA_CHECK(cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice));

vectorAdd<<<blocks, kThreadsPerBlock>>>(d_a, d_b, d_out, kElems);

CUDA_CHECK(cudaGetLastError());
CUDA_CHECK(cudaDeviceSynchronize());

// A device-to-host cudaMemcpy is itself a sync point, so the synchronize
// above is not needed for correctness. It is there to attribute a kernel
// failure to the kernel's line instead of to this one.
CUDA_CHECK(cudaMemcpy(h_out.data(), d_out, bytes, cudaMemcpyDeviceToHost));

Drop the cudaDeviceSynchronize() and the program is still correct. Drop it and remove the checks, and a kernel fault reports itself from the last line, which is how a copy gets blamed for a bug in a kernel.

Diagram

timeline-host-device: two lanes, host above and device below. cudaMalloc returns a device address with no data crossing. Two host-to-device copies cross as solid arrows and the host lane is blocked for their whole width. The launch returns to the host immediately while the device lane keeps working. The device-to-host copy blocks the host again, and the host lane resumes only when the device lane is clear.

Alt text: "A host and device timeline for day 5's vector add. The host blocks for the whole of each cudaMemcpy, returns immediately from the kernel launch, and blocks again on the copy back, which is where a kernel failure is usually reported."

Related terms

Where you meet this

Day 5, CUDA vector addition, end to end with error checking, the lesson that owns this term and breaks the program three ways on purpose.
Day 9, why your GPU code looks slower than your CPU, where the copies get timed against the kernel.
invalid argument, the catch-all a bad pointer or a bad direction usually lands on.
out of memory, when the cudaMalloc before the copy failed and was not checked.
How to set up CUDA, if nvcc is not yet compiling this program at all.

Sources

CUDA Runtime API, memory management, including the undefined-behaviour wording for a mismatched direction: https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__MEMORY.html (checked 2026-08-29)
CUDA Programming Guide, on cudaMemcpy being synchronous and on what a launch does not report: https://docs.nvidia.com/cuda/cuda-programming-guide/02-basics/intro-to-cuda-cpp.html (checked 2026-08-29)

Byline

Written by: pending. Reviewed by: pending. Written on: pending. Last checked on: pending. This entry stays a draft until a named author and a different named reviewer sign it.