Unified memory is a single allocation addressable from both the host and the device, created with cudaMallocManaged. The driver migrates pages between CPU and GPU on demand, so the explicit cudaMemcpy pairs disappear.

float *a;
cudaMallocManaged(&a, n * sizeof(float));
 
for (int i = 0; i < n; i++) a[i] = 1.0f;   // touched on the host
 
kernel<<<blocks, threads>>>(a, n);          // pages migrate to the device
cudaDeviceSynchronize();                    // required before reading on the host
 
printf("%f\n", a[0]);                       // pages migrate back
cudaFree(a);

This cuts a lot of boilerplate, and it makes pointer-rich data structures such as linked lists and trees usable on the GPU without rebuilding them with device offsets. It is the fastest way to get existing CPU code running on a GPU at all.

The cost is that migration happens through page faults, and a fault is far more expensive than a bulk copy. Code that alternates between host and device access on the same pages can thrash badly enough to be slower than the CPU version. The cudaDeviceSynchronize() before host access is not optional either, since reading a page the GPU still owns is a race.

Two hints help when the access pattern is known ahead of time:

cudaMemPrefetchAsync(a, bytes, deviceId);        // move pages before the kernel needs them
cudaMemAdvise(a, bytes, cudaMemAdviseSetReadMostly, deviceId);

Prefetching recovers most of the gap against explicit copies. The usual approach is to prototype with unified memory, then profile and either add prefetch hints or switch the hot allocations to explicit transfers.