Memory transfer moves data between host and device across the PCIe bus and is often the real bottleneck in a GPU program. For small problems it costs more than the computation it feeds, since host and device have separate physical memory that must be explicitly moved.

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.