the offload model
When a CPU teams up with an accelerator like a GPU, who is in charge, and how does work actually get to the accelerator? The offload model is the standard answer: the CPU (the host) stays the boss and runs the main program, and it offloads selected heavy, parallel chunks to the device (the GPU or other accelerator), which does the number-crunching and hands results back. The host orchestrates; the device accelerates.
Concretely, an offload has a recurring shape, and the host drives every step. First the host allocates memory on the device (the device usually has its own separate memory). Then it copies the input data from host memory across the connecting bus into device memory. Then it launches the kernel - tells the device to run the parallel computation across many threads. The device runs, the host typically waits or does other work, and finally the host copies the results back from device memory to host memory. Picture mailing a big batch of forms to a fast processing center: you have to package and ship them there (copy in), they process the whole batch in parallel, and they mail the results back (copy out). Crucially the host and device often do not share one address space, so a host pointer is meaningless on the device and vice versa - the copies are not optional bookkeeping, they are how data crosses the gap.
It matters because this round-trip is the dominant cost model for heterogeneous computing, and it explains a constant practical lesson: the transfer is slow relative to the compute. A kernel that runs in microseconds is pointless if you spend milliseconds shipping its data over the bus, so the offload model only pays off when the chunk of work is big enough, parallel enough, and reused enough to amortize the copies - or when you keep data resident on the device across many kernels to avoid copying at all. 'Unified' or 'shared' memory schemes try to hide the copy by letting both sides see one address space, but the data still has to physically move, so they ease the programming, not the laws of bandwidth.
host: cudaMalloc(&d_a, bytes); // 1. allocate on device cudaMemcpy(d_a, a, bytes, H2D); // 2. copy input host -> device kernel<<<grid, block>>>(d_a, d_c); // 3. launch (device computes) cudaMemcpy(c, d_c, bytes, D2H); // 4. copy result device -> host // steps 2 and 4 cross the bus and often cost more than step 3
The four-step offload: allocate, copy in, launch, copy out. The two copies cross the host-device bus and frequently dominate the total time.
Host and device usually have separate memories, so a host pointer is meaningless on the device - the copies are mandatory, not bookkeeping. The frequent mistake is timing only the kernel and ignoring the transfer; for small or one-shot work, the copy-in/copy-out can dwarf the computation and make offloading a net loss.