LESSON 01 · 2026.04.18 · T4

vector_add — what the first kernel taught

Is compute more expensive, or the copy? The first kernel I ran on T4 returned the proposition "data movement is the bottleneck" as a 17× gap in numbers.

Setup

My local is Apple Silicon so I can't run CUDA directly. I spun up a T4 Spot VM in us-east1-d with GCP credits, booted the Deep Learning VM image, and compiled remotely. The first kernel is a textbook vector add. Simple — but simple is where you learn.

__global__ void vector_add_kernel(const float* a, const float* b, float* c, size_t n) {
  size_t idx = blockIdx.x * blockDim.x + threadIdx.x;
  size_t stride = gridDim.x * blockDim.x;
  for (size_t i = idx; i < n; i += stride) c[i] = a[i] + b[i];
}

The reason for the grid-stride loop is simple: if problem size n exceeds the grid, one thread handles several elements. Launch config becomes less sensitive to n.

Results

Main run (block_size = 256):

best kernel time

3.488 ms

effective bandwidth

230.909 GB/s

theoretical (T4)

320.064 GB/s

efficiency

72.145%

H2D copy

44.276 ms

D2H copy

20.439 ms

Shaking block size didn't move things dramatically.

The real bottleneck

Kernel: 3.4 ms. Copies: 60 ms+ combined. The real bottleneck is not inside the kernel — it's over PCIe.

The first lesson's takeaway is simple: data movement is more expensive than compute. This proposition defines all nine lessons that follow — pinned memory (02), reduction tree (03), tiling (04), fusion (05), Flash Attention (06). All variations on "how do we move less."

Why bytes_moved = n × 4 × 3

Per element we read two floats and write one. 4 bytes × 3 = 12 bytes per element. Dividing by time gives effective bandwidth. It only measures how fast the kernel does "its own work." Copy time isn't included — the next lesson covers that separately.

Why vector add is memory-bound

Arithmetic intensity is 12 bytes per FLOP. T4's roofline knee sits around FP32 peak 8.1 TFLOPS / 320 GB/s ≈ 25 FLOP/byte. Our intensity is far below that — HBM bandwidth is the ceiling. No amount of kernel tuning beats the bandwidth cap.