What Vectors Are & Why GPUs Love Them

The smallest workload that exposes the entire GPU model

A vector is an array. Two vectors of the same length, added: w[i] = u[i] + v[i] for every i. One arithmetic op per element, no dependencies between elements — the textbook definition of embarrassingly parallel.

'Embarrassing' here is praise. Embarrassingly parallel workloads have:

No cross-element communication — each output depends only on its own inputs.
Trivially load-balanced work — every element costs the same.
Linear memory access — fully coalesced, full bandwidth.
Zero synchronization — no __syncthreads needed.

That's why we start here. Every later workload (matrix add, GEMM, attention, softmax) inherits this skeleton and adds complications: 2-D indexing, shared-memory reuse, reductions, masks. Master the skeleton first.

Why GPUs love it specifically: a vector add saturates DRAM bandwidth without ever needing the L2 cache. It's the simplest test of 'is my memory subsystem alive?' — if your hand-rolled kernel hits 80% of theoretical bandwidth here, your toolchain and launch geometry are healthy.

Code

The skeleton — one line of math, infinite scaling·cuda

// CUDA
__global__ void vector_add(
    const float* __restrict__ A,
    const float* __restrict__ B,
    float*       __restrict__ C,
    size_t N)
{
    size_t i = (size_t)blockIdx.x * blockDim.x + threadIdx.x;
    if (i < N) C[i] = A[i] + B[i];
}

// __restrict__ tells nvcc the three pointers don't alias.
// That unlocks reorderings that would otherwise be illegal.

Saturating bandwidth — back-of-envelope expectation·python

# RTX 4090: peak DRAM ~ 1.0 TB/s
# Vector add of N FP32 elements:
#   reads:  2 * N * 4 bytes (A + B)
#   writes: 1 * N * 4 bytes (C)
#   total:  3 * N * 4 bytes
# Time at peak BW: bytes / BW

N = 1 << 26   # 64M elements ≈ 768 MB total traffic
bytes_ = 3 * N * 4
time_s = bytes_ / 1.0e12
print(f'Vector add {N:,}: ~{time_s*1000:.2f} ms at peak BW')
# A healthy kernel reaches 80–90% of this. Below 50% → bandwidth-bound
# but launch geometry is wrong: too few blocks, bad coalescing, etc.

Exercise

On any GPU you have access to, run a 64M-element vector add (FP32) and time just the kernel (not the H↔D copies). Divide bytes touched (3 × N × 4) by your measured time to get achieved bandwidth in GB/s. Compare to your card's spec peak. The number you get is your starting reference for every later optimization.

What Vectors Are & Why GPUs Love Them

The smallest workload that exposes the entire GPU model

Code

External links

Exercise

Progress

Comments 0