The CUDA Program Lifecycle

CUDA Deep Dive · Part 1: CUDA Foundations · B-01
This post is one of a series. See the Series Map for the full plan and where this post fits. Next: B-02 The GPU Hardware Model (coming soon).

One vector addition, followed from the input file to the last freed buffer: nine stages, and the kernel is only one of them.

Who this is for: a software engineer or ML practitioner who knows C++ basics (pointers, stack vs heap, manual memory allocation), may have used PyTorch, and has never written a CUDA kernel.

What this post does: it walks one complete CUDA program through its nine stages, from reading the input to releasing the last buffer, and leaves you with a model of where a GPU program’s time actually goes. Everything later in the series is an improvement to this model.

Start here if you know what a GPU kernel is and roughly why GPUs exist. If you have never seen one, read H-01 Why GPU? CPU vs GPU Architecture first. It explains what this post takes for granted: how CPUs hit the power wall, how graphics chips became programmable, and how that turned into CUDA. This post starts where that one ends.

Writing the kernel is the easy part

H-01 ended with the CPU and the GPU working as a pair. The CPU runs the program. When it reaches a large, uniform task, it hands that task to the GPU and collects the answer afterwards. CUDA calls the two sides host and device. This post is about what that hand-off involves in practice, step by step, and what each step costs.

Consider the simplest job a GPU can do: add two arrays of floats. Each output element is the sum of two inputs. The arithmetic is one line. Yet before the GPU can run that line, your program has to find the input, make room for it, carry it to the GPU, start the computation, bring the answer back, and clean up after itself. The one line of arithmetic is surrounded by eight stages of logistics.

The reason is that a CUDA program on a discrete GPU lives in two memory worlds:

• The host: the CPU and its system memory (RAM).
• The device: the GPU and its own memory, usually dedicated VRAM on the card.

The two are connected by an interconnect, most commonly PCI Express (PCIe). Every byte that crosses it costs time and bandwidth. That is why a very fast kernel does not automatically make the program fast: the program may be spending most of its time moving data across PCIe.

If a picture helps: think of the CPU as an office and the GPU as a factory across town. System RAM is the office’s storeroom, VRAM is the factory’s storeroom, and PCIe is the truck between them. Nothing is built until it has been trucked out, and nothing is used until it has been trucked back.

If there is one sentence to keep for the whole post, it is this:

Optimize the movement and lifetime of your data together with the computation that runs on it.

The nine stages below show why that sentence is true, using the explicit memory-management model that every CUDA program starts with. The optimizations that turn this basic program into a pipeline (asynchronous copies, multiple streams, keeping data on the GPU) are each a change to one of these stages, and we will point at them as we go.

The map: nine stages

Figure B01.1 shows the whole journey before we take it. Nine stages, and only one of them, stage 5, does any arithmetic.

Figure B01.1: A CUDA program is nine stages, and only one of them is the kernel

1. Get hold of the input.
2. Make room for it in host memory and fill that room.
3. Reserve room in device memory.
4. Ship the input across to the device.
5. Launch the kernel, which the GPU runs across many threads.
6. Order the work and wait for it where necessary.
7. Ship the results back to the host.
8. Use the results.
9. Give every resource back.

Stages 4 and 7 are the two copies across PCIe. Keep an eye on them: a faster kernel cannot shrink them.

Stage 1: Getting the input

Every computation starts with data, and the data starts somewhere else: a file on an SSD, a camera, a network socket, a sensor, a previous result, or a synthetic dataset the CPU generates. Suppose we want to process a 4096 x 4096 image stored as 32-bit floats, read from a binary file:

#include <fstream>
#include <vector>
#include <cstddef>

int main() {
    const int width  = 4096;
    const int height = 4096;
    const size_t num_elements = static_cast<size_t>(width) * height;
    std::vector<float> image(num_elements);
    std::ifstream file("input.bin", std::ios::binary);
    if (!file) {
        return 1;
    }
    file.read(reinterpret_cast<char*>(image.data()),
              static_cast<std::streamsize>(num_elements * sizeof(float)));
    if (!file || file.gcount() !=
            static_cast<std::streamsize>(num_elements * sizeof(float))) {
        return 1;
    }
    // The input is now in host memory. The GPU has not seen it.
}

This is ordinary C++ file I/O. When it finishes, the image sits in CPU-accessible memory, and the GPU has done nothing yet.

Two trips, not one. It is easy to assume that once a file is “loaded” the GPU can get at it. It cannot. There are two separate movements of data:

• Storage to host: the operating system and the storage subsystem bring the file’s bytes into system RAM.
• Host to device: your program copies those bytes from RAM into GPU memory.

Each has its own bottleneck. The first is limited by how fast the drive can deliver bytes (or the OS page cache, which may already hold the file). The second is limited by the interconnect. A faster drive speeds up the first trip and does nothing for the second.

How big is the load? The image has 4096 x 4096 = 16,777,216 elements. At 4 bytes each, that is 67,108,864 bytes, or 64 MiB. The arithmetic the GPU will do on it is trivial; moving 64 MiB across a bus is not free. Figure B01.2 puts the two trips and the three memories side by side, and previews stage 3: by the time we copy, the GPU-side buffers already exist, but they are empty.

Figure B01.2: Reading a file into RAM and copying it to the GPU are two separate trips, and cudaMalloc is neither of them

With the input in hand, we need somewhere on the host to keep it, and there are two kinds of “somewhere”.

Stage 2: Host memory, pageable or pinned

Host memory is plain system RAM, but CUDA cares about one property of it: whether the operating system is allowed to move it.

Pageable memory is what every ordinary allocation gives you:

std::vector<float> h_input(num_elements);
std::vector<float> h_output(num_elements);

Here h_input holds the input and h_output will eventually receive the result. (The h_ prefix for host pointers and d_ for device pointers is a convention you will see in almost all CUDA code, and we use it throughout the series.) Pageable memory is convenient: it is what std::vector, new and malloc give you, and the OS manages it normally, which includes being free to page it out or move it.

That freedom is the catch. The GPU’s copy engine moves bytes with DMA (direct memory access), which needs a buffer that stays put at a fixed physical address. When you copy from pageable memory, the CUDA runtime first copies your bytes into a page-locked staging buffer of its own, then transfers from there. Two hops instead of one. The extra hop costs time and, as we will see in stage 4, rules out truly asynchronous transfers.

Pinned memory, also called page-locked memory, is host memory the OS promises not to move:

float* h_input  = nullptr;
float* h_output = nullptr;

cudaMallocHost(&h_input,  num_elements * sizeof(float));
cudaMallocHost(&h_output, num_elements * sizeof(float));

cudaMallocHost allocates ordinary RAM and pins its pages. The copy engine can now read it directly, in one hop, and cudaMemcpyAsync can overlap the copy with other work. Figure B01.3 shows the difference.

Figure B01.3: A copy from pageable memory takes two hops; from pinned memory to one

Pinned memory is still host memory. It is not VRAM, and it does not make the GPU any faster at reading the data once it has arrived. It is also a system resource: pinned pages cannot be swapped out, so pinning very large amounts starves the rest of the system. The rule of thumb is to pin the buffers you transfer often, not every array in the program.

Initialize before you ship. Whichever kind of host memory you choose, the input buffer must hold valid data before it is copied: file I/O, preprocessing, conversion, or CPU computation all happen here. The output buffer does not need initializing if the kernel will write every element before the CPU reads one. The contract is simple: every input element the kernel reads must be valid, and every output element the program consumes must have been written by the computation.

Now we need the matching room on the other side of the bus.

Stage 3: Device memory

In the explicit model, cudaMalloc reserves memory in the device's address space. For a vector addition we need three buffers, two inputs and one output:

float* d_a = nullptr;
float* d_b = nullptr;
float* d_c = nullptr;

cudaMalloc(&d_a, num_elements * sizeof(float));
cudaMalloc(&d_b, num_elements * sizeof(float));
cudaMalloc(&d_c, num_elements * sizeof(float));

You might expect cudaMalloc to return a pointer, the way malloc does. It does not; it fills in the pointer you pass by address, and returns an error code instead (B-04 is where we start checking those codes). Each call reserves enough device memory for one array. On a discrete GPU that memory is on the card, in VRAM, and the pointer you get back refers to the device's address space: ordinary CPU code cannot dereference d_a as if it were a normal array. It is an address in the GPU's memory, and only GPU code can use it.

Allocation is not transfer. This is the most common misunderstanding in beginner CUDA code, so it gets its own paragraph. cudaMalloc puts nothing into the memory it reserves. After the three calls above, d_a, d_b and d_c exist and contain garbage. Allocation decides where data can live; a copy, or a kernel that generates data on the device, decides what is actually there. Reserving the space does not fill it.

Why three separate buffers? The kernel reads d_a and d_b and writes d_c. Separate buffers make those roles explicit and make it impossible to overwrite an input by accident. Later, when memory is tight, you will see in-place computation and buffer reuse, which are fine when the algorithm's dependencies permit it. For now, clarity wins.

Allocate for the workload, not for one operation. Suppose an image-processing pipeline runs five kernels over the same image. Reallocating the image buffer before each kernel would be wasteful: allocation has a cost, and allocation inside a loop makes the loop hard to pipeline. Allocate once, reuse across the pipeline, free at the end. Stage 9 comes back to this, because it turns out to be a performance habit, not just a tidiness habit.

With room reserved on both sides, the data can be copied across.

Stage 4: Host to device

cudaMemcpy moves bytes between the two memories. The last argument says which direction:

cudaMemcpy(d_a, h_a.data(), num_elements * sizeof(float), cudaMemcpyHostToDevice);
cudaMemcpy(d_b, h_b.data(), num_elements * sizeof(float), cudaMemcpyHostToDevice);

The argument order is destination first, then source, like memcpy. On a discrete GPU the bytes travel over PCIe, usually driven by a DMA engine rather than by the CPU, although the exact path depends on the hardware and the memory type.

What happens during the call. Conceptually: the runtime identifies source and destination, the runtime and driver set up the transfer, the bytes move, and the call completes according to its synchronization guarantee, which for the plain cudaMemcpy means "when the copy is finished". With pageable memory the staging hop from stage 2 happens inside this call. With pinned memory the engine reads your buffer directly.

Estimating the trip. A useful first approximation is that transfer time is the bytes moved divided by the effective bandwidth:

T_transfer ≈ bytes transferred / effective bandwidth

Suppose a program moves 1 GiB and the link delivers an effective 24 GiB/s. Then the transfer takes about 1/24 of a second, roughly 41.7 ms. Take 24 GiB/s as a round number for an effective PCIe rate; the real figure depends on the link generation, the lane count and whether the host buffer is pinned. The estimate also ignores the fixed per-call overhead, which matters for small copies. What it gets right is the order of magnitude, and the lesson: in many programs, moving the data takes longer than computing on it. Our 64 MiB image at the same rate takes about 2.6 ms to cross, each way. Keep that number. Stage 7 uses it again.

Synchronous copies. The plain cudaMemcpy above is a blocking call: it does not return to your CPU thread until the copy is complete. That makes the program easy to reason about. After the call returns, the data is on the device and the next line can use it.

Asynchronous copies, a preview. CUDA also offers cudaMemcpyAsync:

cudaMemcpyAsync(d_a, h_a, num_elements * sizeof(float), cudaMemcpyHostToDevice, stream);

The extra argument is a stream, a queue of GPU work that executes in order. An asynchronous copy returns to the CPU immediately, which lets the CPU do other work while the bytes move and, with a second stream, lets a copy overlap a kernel. This is where stage 2’s choice comes back: the overlap only happens from pinned memory. From a pageable buffer the runtime still has to do the staging hop on the CPU before the DMA engine can start, so the “asynchronous” call does most of its work before it returns. The price of real asynchrony is that “the call returned” no longer means “the copy is finished”. You must keep the source buffer alive until the copy completes, and you must make sure nothing reads the destination before the copy is done. B-05 is entirely about getting this right. In this post we stay with blocking copies.

The data is now in device memory. The next stage puts the GPU to work on it.

Stage 5: Launching the kernel

A kernel is a function the GPU executes, with many threads running it at the same time. Here is vector addition:

__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}

__global__ marks the function as a kernel: callable from the host, executed on the device. The body has no loop, which surprises everyone coming from C. The repetition lives in the launch, not in the function. Each thread runs this body once, and the only thing that differs between threads is the value of i.

Which element am I? Each thread computes its global index from three built-in variables:

int i = blockIdx.x * blockDim.x + threadIdx.x;

• blockIdx.x is which block this thread belongs to, counting across the grid.
• blockDim.x is how many threads a block holds.
• threadIdx.x is this thread's position inside its block.

Think of the launch as a grid of blocks, each block a group of threads. With 256 threads per block, thread 5 of block 2 computes i = 2 * 256 + 5 = 517 and handles exactly one element: c[517] = a[517] + b[517]. The check if (i < n) protects the last block. When the array length does not divide evenly by the block size, the final block contains threads whose index lands beyond the array, and those threads must do nothing.

Configuring the launch. The host decides how many threads to launch with the triple-chevron syntax:

int threads_per_block = 256;
int blocks = (num_elements + threads_per_block - 1) / threads_per_block;

vectorAdd<<<blocks, threads_per_block>>>(d_a, d_b, d_c, num_elements);

The expression for blocks is integer ceiling division: enough blocks to cover every element, rounding up. For 1,048,576 elements (that is 1 << 20, the size used in the complete program below) and 256 threads per block, it comes to exactly 4,096 blocks: 1,048,576 threads, one per element. Figure B01.4 traces one of them.

Get Agcshubham’s stories in your inbox

Join Medium for free to get updates from this writer.

Remember me for faster sign in

Figure B01.4: <<<4096, 256>>> launches 1,048,576 threads, the GPU hands the blocks to its SMs a few at a time, and each thread computes exactly one index

The launch describes a grid. The GPU decides how to run it. A GPU is made of processing units called streaming multiprocessors, or SMs. Each block is handed to one SM when that SM has room. An SM runs several blocks at once, but a grid of 4,096 blocks is far more than the SMs can hold at one time, so the blocks go through in waves: when a block finishes, the next one in line takes its place. B-02 explains the hardware behind this.

What the GPU touches while it runs. Even this tiny kernel uses several resources. Global memory (VRAM) holds the arrays. Caches absorb some repeated accesses. Registers hold each thread’s i and temporaries. Shared memory, a programmer-managed scratchpad per block, goes unused here but is central to later posts. And the SMs execute the instructions, running threads in fixed groups of 32 that CUDA calls warps. Vector addition lives almost entirely in registers and global memory. A tiled matrix multiply (P-06) depends on shared memory. B-03 maps this hierarchy level by level.

Launching is not finishing. A kernel launch does not wait for the GPU. The CPU reaches the <<<...>>> line, submits the work, and moves on to the next statement immediately, while the GPU may not even have started. So when the launch expression returns, d_c does not necessarily hold results. The CPU has handed the work to the GPU and moved on. The GPU is still running. Before anyone uses d_c, something has to wait.

Stage 6: Synchronization and ordering

CUDA separates submitting work from completing it. That separation is what makes overlap possible, and it is also why synchronization is your responsibility.

vectorAdd<<<blocks, threads_per_block>>>(d_a, d_b, d_c, num_elements);
// Right here, the GPU may still be running vectorAdd.

Device-wide synchronization. The simplest tool, and the heaviest:

cudaDeviceSynchronize();

This blocks the host thread until every piece of device work submitted so far has completed. It is the right tool when you genuinely need a device-wide checkpoint: before timing, while debugging, before the program exits. It is the wrong habit after every kernel, because each call stops everything and removes any overlap you might have had.

Stream ordering. Work submitted to the same stream executes in the order you submitted it. If a kernel and then a device-to-host copy go into the same stream, the copy cannot start until the kernel has finished, and a blocking copy waits for the copy itself. So the sequence “launch, then blocking cudaMemcpy back" is already correctly synchronized, with no cudaDeviceSynchronize in between: the blocking copy does the waiting for you. Figure B01.5 shows the two clocks and where the waiting actually happens.

Figure B01.5: The launch returns before the kernel has run; the blocking copy afterwards is what actually waits

Events. A CUDA event is a marker you drop into a stream. The host can wait on it, one stream can wait on another stream’s event, and two events can measure the elapsed GPU time between them. Events let you coordinate without stopping the whole device, and they are the correct way to time a kernel, since a CPU-side stopwatch around a launch only captures the launch call itself, not the work. B-05 covers events as a dependency tool; B-06 covers them as a measuring tool.

When the kernel has finished, the result is in d_c, in device memory. The CPU still cannot read it from there.

Stage 7: Device to host

The return trip is the same call with the direction reversed:

cudaMemcpy(h_c.data(), d_c, num_elements * sizeof(float), cudaMemcpyDeviceToHost);

Once this blocking copy returns, h_c holds the sums and the host can use them.

Do you always need the trip back? No, and this is where a lot of real performance lives. Suppose an image pipeline does preprocessing, then classification, both on the GPU. The preprocessed image is an intermediate result. If it stays on the device, the classification kernel reads it from VRAM and only the final answer, which may be a single pass/fail value, comes back to the CPU. The naive version copies the intermediate down to the CPU and back up again, for no reason other than habit. Figure B01.6 counts the cost.

Figure B01.6 : The kernels are unchanged, but the naive pipeline moves 192 MiB over PCIe and the GPU-resident one moves 64 MiB plus 4 bytes

At our round-number 24 GiB/s, that is about 7.8 ms of transfers against 2.6 ms, three times the PCIe traffic for the same answer, and the kernels did not change at all. The rule that falls out: keep data on the GPU for as long as the pipeline allows, and bring back only what the host actually needs. Every deep learning framework you have used does exactly this; a training step passes activations from layer to layer without ever visiting the CPU. We build that pattern by hand in T-01 and I-01.

Stage 8: Consuming the result

With the output in host memory, the program does whatever it was for: writes a file, displays an image, returns a prediction to another service, runs CPU postprocessing, sends bytes over a network, or logs a metric. For our vector addition:

std::cout << "Result[0] = " << h_c[0] << '\n';

If the inputs were all 1.0 and 2.0, the output is 3.0. A real program checks more than element zero, ideally every element against a CPU reference, because a successful allocation, copy and launch tell you nothing about whether the arithmetic was right. “It ran” and “it is correct” are different claims; H-04 is about the gap between them.

One more time, because this is the stage where the mistakes show up: consume only completed results. With blocking copies that is automatic. The moment you move to asynchronous copies, you must wait on the right event or stream before touching the destination buffer.

Stage 9: Cleanup

When the computation is done and the buffers are no longer needed, release them:

cudaFree(d_a);
cudaFree(d_b);
cudaFree(d_c);

and, for pinned host memory:

cudaFreeHost(h_input);
cudaFreeHost(h_output);

Ordinary containers like std::vector free themselves when they go out of scope.

Match the pair. Every allocation API has one release API: cudaMalloc with cudaFree, cudaMallocHost with cudaFreeHost, new[] with delete[], malloc with free. Crossing them (passing a pinned pointer to free, or a device pointer to cudaFreeHost) is undefined behavior at best and a crash at worst.

Do not free what is still in use. A kernel or an asynchronous copy can still depend on a buffer after the host has submitted the operation. Free only after the work that uses the buffer has completed, which with multiple streams and reused buffers takes some care.

Cleanup is performance engineering. In an iterative workload (a training loop, a video stream), allocating and freeing large buffers every iteration adds avoidable overhead and defeats pipelining. Allocate once, reuse across iterations, free when the loop ends. Larger programs go further with CUDA memory pools and caching allocators, which is exactly what PyTorch’s allocator does on your behalf.

The complete program

Here are all nine stages in one file. It uses blocking copies and the default stream, so the order of events is exactly the order of the lines:

#include <cuda_runtime.h>
#include <iostream>
#include <vector>

__global__ void vectorAdd(const float* a, const float* b, float* c, int n) {
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) {
        c[i] = a[i] + b[i];
    }
}
int main() {
    const int N = 1 << 20;                 // 1,048,576 elements
    const size_t bytes = N * sizeof(float);
    // Stages 1 and 2: host memory, initialized.
    std::vector<float> h_a(N, 1.0f);
    std::vector<float> h_b(N, 2.0f);
    std::vector<float> h_c(N);
    // Stage 3: device memory.
    float* d_a = nullptr;
    float* d_b = nullptr;
    float* d_c = nullptr;
    cudaMalloc(&d_a, bytes);
    cudaMalloc(&d_b, bytes);
    cudaMalloc(&d_c, bytes);
    // Stage 4: host to device.
    cudaMemcpy(d_a, h_a.data(), bytes, cudaMemcpyHostToDevice);
    cudaMemcpy(d_b, h_b.data(), bytes, cudaMemcpyHostToDevice);
    // Stage 5: launch.
    const int threads = 256;
    const int blocks  = (N + threads - 1) / threads;
    vectorAdd<<<blocks, threads>>>(d_a, d_b, d_c, N);
    // Stages 6 and 7: the blocking copy waits for the kernel, then copies.
    cudaMemcpy(h_c.data(), d_c, bytes, cudaMemcpyDeviceToHost);
    // Stage 8: validate and consume.
    int wrong = 0;
    for (int i = 0; i < N; ++i) {
        if (h_c[i] != 3.0f) {
            ++wrong;
        }
    }
    std::cout << "Result[0] = " << h_c[0] << '\n';
    std::cout << (N - wrong) << " of " << N << " elements correct\n";
    // Stage 9: release.
    cudaFree(d_a);
    cudaFree(d_b);
    cudaFree(d_c);
    return wrong == 0 ? 0 : 1;
}

Save it as lifecycle.cu, compile it with the CUDA compiler, and run it:

nvcc lifecycle.cu -o lifecycle
./lifecycle

Expected output:

Result[0] = 3
1048576 of 1048576 elements correct

Two things are deliberately missing. There is no error checking: every CUDA call returns a status code, and a real program checks each one (B-04 introduces the discipline, and the macro that makes it bearable). And there is no cleanup on early exit: a production program would release what it had already acquired, which in C++ means RAII wrappers around the CUDA allocations. The example stays plain so the lifecycle is visible in the order it happens.

Where the time goes

The lifecycle gives us a frame for performance. As a first model, total time is the sum of the stages:

T_total = T_input + T_alloc + T_H2D + T_kernel + T_D2H + T_output + T_other

where T_input is host-side input preparation and file reading, T_alloc is allocation overhead, T_H2D and T_D2H are the two transfers, T_kernel is execution, T_output is host-side consumption, and T_other is everything else. This is a conceptual model, not a measurement recipe: in real programs stages overlap, and a synchronizing call can hide one stage’s time inside another’s, so adding up independently measured stages can double count. B-06 shows how to measure without double counting; here we only need the shape.

A worked example. Suppose a program’s stages run one after another with these times: host preparation 1 ms, input transfer 5 ms, kernel 2 ms, output transfer 5 ms. Total: 13 ms.

Now suppose you optimize the kernel and halve its time, from 2 ms to 1 ms. The program takes 12 ms. Speedup: 13/12, about 1.08x. The kernel got twice as fast and the program got 8 percent faster.

Instead, suppose the program no longer needs to copy its big output tensor back to the CPU, because the next stage runs on the GPU too. Drop the 5 ms output transfer and the program takes 8 ms, a 1.63x speedup, with the kernel untouched. Figure B01.7 draws the three cases.

Figure B01.7: Halving the kernel saves 1 ms of 13; removing the output copy saves 5

This is Amdahl’s law from H-01, applied to one CUDA program: the part you did not optimize limits the speedup you can get. It is also why the first measurement of any GPU program should be end-to-end, transfers included. A kernel-only number can look excellent while the program as a whole has barely moved. B-06 makes “say what your timing includes” a habit.

What to optimize first. The process is the same every time: measure the whole program; find out how much goes to input, transfers, kernels and output; identify the dominant stage; optimize that stage without adding overhead elsewhere; measure again. If the kernel dominates, the Optimization Toolkit (Part 3) is for you: access patterns, occupancy, instruction throughput, algorithms. If data movement dominates, look at transfer volume, pinned buffers, buffer reuse, asynchronous copies, and keeping intermediates on the device (B-05). If allocation dominates a loop, reuse buffers or use a memory pool.

Beyond the basic lifecycle

Once the explicit lifecycle is clear, the standard optimizations stop being tricks and become edits to a specific stage:

• Reuse allocations (stages 3 and 9). Allocate once, reuse across iterations.
• Keep intermediate data on the GPU (stage 7). Copy back only what the host needs.
• Pinned buffers and asynchronous copies (stages 2 and 4). cudaMallocHost plus cudaMemcpyAsync make overlap possible; correct stream dependencies make it safe. See B-05.
• Multiple streams (stage 6). Independent streams let independent work overlap. Streams do not guarantee concurrency: dependencies, hardware resources and bandwidth all have a say. See B-05.
• Fuse compatible kernels (stage 5). One kernel instead of two removes a round trip through global memory and a launch. Fusion can raise register pressure, so measure it. See P-01.
• Managed memory (stages 2 to 4, and 7). cudaMallocManaged lets the runtime migrate pages between host and device for you. Convenient, but the migrations still happen and still cost; you have traded control for simplicity. See B-03.
• Profile the real execution (all stages). Nsight Systems shows the timeline of the whole program: CPU activity, transfers, synchronization, streams. Nsight Compute dissects one kernel. The right optimization is the one the profiler points at, not the one you assumed. See B-06.

Gotchas

1. cudaMalloc does not copy anything. After cudaMalloc(&d_a, bytes), d_a is reserved and full of garbage. The cudaMemcpy in stage 4 is what fills it. A kernel that reads d_a before the copy produces plausible-looking nonsense rather than an error.
2. The launch returns before the kernel runs. vectorAdd<<<blocks, threads>>>(...) hands work to the GPU and returns immediately. Reading d_c from the host right after it, or freeing d_a right after it, races the GPU. Something must wait: a blocking copy, an event, or cudaDeviceSynchronize.
3. A blocking cudaMemcpy is already a sync point. Because same-stream work runs in order, the device-to-host copy in the listing waits for the kernel. Adding cudaDeviceSynchronize before every copy is harmless in this program and harmful in a pipelined one, where it drains the overlap you worked to create.
4. Host pointers and device pointers are not interchangeable. h_a and d_a look identical in C++ and point into different memories. Dereferencing d_a on the CPU, or passing h_a to a kernel, is undefined behavior. The h_ and d_ prefixes are the only thing protecting you.
5. Pinned memory is not device memory. cudaMallocHost gives you RAM that the OS will not move. It speeds up the trip to the GPU; it does not shorten it to zero, and pinning everything slows the whole machine.
6. Pair every allocation with its own release. cudaMalloc with cudaFree, cudaMallocHost with cudaFreeHost, new[] with delete[], malloc with free. Mixing them is undefined behavior.
7. The fast kernel did not make the program fast. In the worked example, halving the kernel improved the program by 8 percent; removing one transfer improved it by 63 percent. Measure the whole program before deciding what to optimize.
8. Summing stage timings can double count. A blocking copy’s measured time includes the kernel it waited for. If you time stages independently and add them up, you can exceed the real total. Say what each measurement includes.

Summary

A CUDA program is a coordinated sequence: get the data, make room on both sides, ship it out, compute, wait, ship it back, use it, clean up. Each stage has its own job and its own cost. Host memory has to be prepared, device memory has to be allocated and then separately filled, the kernel has to read valid data, synchronization has to establish the dependencies, and resources have to live until the work that uses them is done.

The kernel is one stage of nine. The fastest kernel in the world cannot make up for transfers you did not need, allocations you repeated, synchronization you did not need, or host code that was the real bottleneck all along. So, the sentence from the top, one more time: optimize the movement and lifetime of your data together with the computation that runs on it.

Everything that follows in this series (asynchronous pipelines, memory coalescing, shared-memory tiling, kernel fusion, multi-stream execution) is an improvement to one of these nine stages. You now have the frame they fit into.

In This Series

Before this post: H-01 Why GPU? CPU vs GPU Architecture, if you have never seen a GPU kernel. It is the history this post assumes: the power wall, programmable shaders, the arrival of CUDA, and why a GPU is built the way it is. If you already know that history, B-01 is the start of the main sequence.

What comes next: B-02 The GPU Hardware Model (coming soon) explains what the SMs, warps and waves of stage 5 actually are. B-06 Measuring CUDA Programs (coming soon) turns the “where does the time go” model into measured numbers.

Related deep dives: B-05 Streams, Events & Async Execution (coming soon) covers in full the overlap previewed in stages 4 and 6. B-03 Memory Hierarchy Deep Dive (coming soon) maps the memory levels the kernel touched.

CUDA Deep Dive · Part 1: CUDA Foundations · B-01
The Series Map lists every post in the series and the order to read them. Next: B-02 The GPU Hardware Model (coming soon).