The Ant Colony Mental Model

CPU = octopus, GPU = ant colony

Here's the elevator-pitch analogy that makes the rest of this quest click:

The CPU is an octopus — clever, multi-armed, but coloring only a handful of shapes at a time. The GPU is a bustling colony of thousands of ants. Each ant is simpler, yet together they finish the entire coloring book before the octopus colors a single page.

That intuition turns into concrete hardware structure once you map it.

Hardware hierarchy (macro → micro)

GPU — the colony itself. Example: RTX 4090 has 16,384 CUDA cores.
GPC (Graphics Processing Cluster) — city-sized districts. RTX 4090 has 7.
SM (Streaming Multiprocessor) — communal workshops. RTX 4090 has 128 SMs.
Warp — a 32-ant platoon marching in lockstep on one SM.
Thread — the individual ant.

Software / execution chain (your kernel's view)

Kernel — the mission briefing broadcast to thousands of ants.
Grid — the entire campaign (all blocks for this launch).
Block — a single squad whose members can chat (shared memory) and synchronize.
Warp — a 32-ant platoon inside a block, marching in lockstep.
Thread — the individual ant on the frontline.

The blocks-to-SMs assignment is the GPU scheduler's job. You ask for 1,000 blocks; the scheduler hands them out to the 128 SMs as resources free up. That's why output order is non-deterministic: the colony self-schedules.

Code

Map the hierarchy onto a kernel launch·cuda

__global__ void example(int *out, int n) {
    int gid = blockIdx.x * blockDim.x + threadIdx.x;
    //         ^^^^^^^^^   ^^^^^^^^^^^   ^^^^^^^^^^^
    //         which block which block-  thread index
    //         in the grid size           inside block
    if (gid < n) out[gid] = gid;
}

int main() {
    int n = 4096;
    int threadsPerBlock = 256;        // typical sweet spot
    int blocks = (n + threadsPerBlock - 1) / threadsPerBlock;
    example<<<blocks, threadsPerBlock>>>(d_out, n);
    //        ^^^^^^  ^^^^^^^^^^^^^^^^
    //        grid    block (per block thread count)
    cudaDeviceSynchronize();
}

Same hierarchy in Metal — same words, different spellings·metal

kernel void example(
    device int *out [[buffer(0)]],
    constant uint &n [[buffer(1)]],
    uint gid [[thread_position_in_grid]])
{
    if (gid < n) out[gid] = int(gid);
}

// Host side (Swift):
// encoder.dispatchThreadgroups(
//     MTLSize(width: blocks, height: 1, depth: 1),
//     threadsPerThreadgroup: MTLSize(width: 256, height: 1, depth: 1))

Exercise

Take any of the GPUs you have access to and write down: (a) total cores, (b) SM (or threadgroup-equivalent) count, (c) warp / SIMD-group width (always 32), (d) max threads per block. nvidia-smi -q | grep -i 'sm\|warp' on NVIDIA, system_profiler SPDisplaysDataType on macOS. Internalize these as the dimensions of your colony.