Your first kernel

01 · The whole program

Four steps: allocate, copy, launch, copy back

Day 01 showed you a single line of kernel code with a promise attached. Here's the full ritual around it. A CUDA program is ordinary C++ where some functions run on the GPU; the host orchestrates four moves — allocate device memory, copy inputs over, launch the kernel, copy the result back:

__global__ void vecadd(const float* x, const float* y,
                              float* out, int n){
    int i = blockIdx.x * blockDim.x + threadIdx.x;
    if (i < n) out[i] = x[i] + y[i];
}

int main(){
    int n = 1 << 20;  // ~1M elements → 12 MB total
    float *x, *y, *out;
    // 1. allocate on the device (note: cudaMalloc, not malloc)
    cudaMalloc(&x, n * sizeof(float));
    cudaMalloc(&y, n * sizeof(float));
    cudaMalloc(&out, n * sizeof(float));

    // 2. host → device
    cudaMemcpy(x, hx, n * sizeof(float), cudaMemcpyHostToDevice);
    cudaMemcpy(y, hy, n * sizeof(float), cudaMemcpyHostToDevice);

    // 3. launch: 4096 blocks × 256 threads = 1,048,576 threads
    vecadd<<<(n + 255) / 256, 256>>>(x, y, out, n);

    // 4. device → host
    cudaMemcpy(hout, out, n * sizeof(float), cudaMemcpyDeviceToHost);
}

Day 01's mysterious line C[i] = A[i] + B[i], now in context: three cudaMallocs, two copies in, a launch, one copy back. The kernel itself is four lines. Everything interesting is in how it's launched and indexed.

02 · `__global__` and the launch

The function that runs everywhere at once

__global__ marks a function the device executes and the host launches. The triple-chevron syntax <<<blocks, threads>>> is not C++ — it's CUDA's launch configuration, saying "run this function once per thread, arranged in this many blocks of this many threads." With <<<4096, 256>>> the GPU starts a million threads. They all execute the same function; what makes each useful is the indexing line.

blockIdx, blockDim, and threadIdx are built-in variables — day 03's grid → block → thread hierarchy, now as arithmetic. global_id = block_index × threads_per_block + thread_in_block. One million threads, each computing its own offset into the arrays, each adding one pair of floats. Day 01's paradox — how does one program become a million workers? — resolves into one line of index math.

03 · Bounds and errors

The two lines between working and debugging until 3am

First, if (i < n). Launch configurations round up — (1,048,576 + 255) / 256 is exact here, but for n = 1,000,000 you'd launch 1,048,832 threads and the last 832 must not touch memory that isn't theirs. Unchecked, out-of-bounds reads return garbage and out-of-bounds writes corrupt neighbors — day 14 makes this a design rule.

Second, error checking. CUDA API calls return error codes that are easy to ignore, and kernel launches fail silently — the error surfaces later, possibly at an unrelated call. Production code wraps both:

#define CUDA_CHECK(call) do {
    cudaError_t err = (call);
    if (err != cudaSuccess)
        fprintf(stderr, "CUDA error %s at %s:%d\n",
               cudaGetErrorString(err), __FILE__, __LINE__);
} while (0)

CUDA_CHECK(cudaMalloc(&x, bytes));
vecadd<<<grid, block>>>(x, y, out, n);
CUDA_CHECK(cudaGetLastError());          // catches launch config errors
CUDA_CHECK(cudaDeviceSynchronize());     // surfaces kernel runtime errors

The synchronize matters because launches are asynchronous: the host moves on while the device works. Runtime errors inside the kernel — a bad index, an illegal address — only become visible when something waits for the GPU. Every mysterious "it only crashes at the memcpy" story traces back to this.

04 · What to notice

Even hello-world has honest numbers

Time this kernel with day 18's tools and the profile is exactly what days 10–11 predict: 12 MB moved (two arrays read, one written), so the bandwidth floor is 12 MB ÷ 504 GB/s ≈ 24 microseconds — and a well-written launch will land close to it. But notice the other number: kernel launch overhead alone is ~5–10 µs. Tiny kernels aren't bandwidth-bound or compute-bound; they're launch-bound, and no amount of day 15 heroics helps. That's the real argument for day 23's fusion — fewer launches is sometimes the whole optimization.

Next: day 14 — launch configuration in practice: grid-stride loops, awkward sizes, 2D grids, and how production kernels pick their block dimensions.