Host and device have separate physical memory, so data has to cross the PCIe bus. Memory transfer is often the real bottleneck in a GPU program, and for small problems it costs more than the computation it feeds.

cudaMemcpy(d_a, h_a, bytes, cudaMemcpyHostToDevice);
kernel<<<blocks, threads>>>(d_a);
cudaMemcpy(h_a, d_a, bytes, cudaMemcpyDeviceToHost);

PCIe gen4 x16 gives roughly 25 GB/s in practice. On-device global memory bandwidth is commonly 20 to 40 times that. A kernel that reads its input once, does a little arithmetic, and writes its output back will spend most of its wall-clock time on the bus, which is the arithmetic intensity argument the Roofline model makes, applied to the interconnect instead of to DRAM.

Ordinary malloc memory is pageable, so the driver has to stage transfers through an internal pinned buffer. Allocating pinned memory directly skips that copy and typically doubles the achieved bandwidth:

float *h_a;
cudaMallocHost(&h_a, bytes);   // pinned, page-locked
cudaFreeHost(h_a);

Pinned memory cannot be paged out by the OS, so allocating too much of it degrades the whole system. It is worth using for buffers that are transferred repeatedly, not for everything.

Pinned memory also enables asynchronous copies, which is what makes overlapping transfers with compute possible:

cudaMemcpyAsync(d_a, h_a, bytes, cudaMemcpyHostToDevice, stream);

The strategies that actually help, in order of effectiveness: move less data, keep data resident on the device across several kernels instead of round-tripping, batch many small transfers into one large one, and overlap what remains with computation using streams.