I'm spending the next 8 weeks learning LLM inference from the GPU up: CUDA kernels, fusion, attention, and eventually a small inference engine running on my own kernels. I don't own an NVIDIA GPU, so everything runs on a Tesla T4 rented by the second on Modal.
Day 1 had one job: figure out what this GPU can actually do, so every kernel I write afterward has something to be measured against.
Why start with memory bandwidth?
Generating tokens one at a time (decode) is memory-bound. Every step reads all the model weights and the KV cache from memory, and the math is tiny by comparison. Roughly, tokens per second ≈ memory bandwidth ÷ bytes read per token. So before optimizing anything, I need the real bandwidth number, not the one on the box.
Step 1: Ask the GPU about itself
I wrote a small devquery program using cudaGetDeviceProperties:
- 40 SMs, compute capability 7.5 (Turing)
- 1024 threads per SM, 64K registers per SM (so 64 registers per thread at full occupancy)
- 64 KB shared memory per SM, 4 MB L2
- 256-bit memory bus at 5,001 MHz
Theoretical bandwidth = memory clock × 2 (double data rate) × bus width in bytes = 320 GB/s, matching NVIDIA's spec sheet.
Step 2: The simplest memory-bound kernel there is
A copy kernel does zero math; it only moves bytes. That makes it the cleanest way to measure real bandwidth. I copied 1 GB four different ways, timing the median of 100 runs and verifying every output:
| Version | GB/s | % of theoretical |
|---|---|---|
cudaMemcpy (device to device) | 231 | 72% |
| Scalar: one float per thread | 245 | 77% |
float4: 16 bytes per thread | 234 | 73% |
Grid-stride float4 (best of sweep) | 236 | 74% |
What surprised me
1. Paper bandwidth isn't real bandwidth. The best I could get was 77% of the spec. GDDR6 never hits its theoretical number, and the T4's 70 W power cap holds it back further. So my real ceiling for this GPU is about 245 GB/s, not 320.
2. My naive kernel beat NVIDIA's cudaMemcpy. One thread per element, about a million blocks, 6% faster than the library call. cudaMemcpy has to handle any size and alignment; a dumb kernel with lots of parallelism doesn't.
3. Vectorizing made it slower. The standard advice is "use float4 loads": fewer instructions for the same bytes. Here it was consistently about 5% slower than scalar. I swapped the run order to rule out clock throttling, and the gap held. I don't know why yet. That's tomorrow's job in Nsight Compute.
4. You need fewer threads than I thought. I swept the grid-stride version from 1 to 32 blocks per SM, expecting low block counts to starve memory. Instead, 1 block per SM (about 10K threads) already hit 236 GB/s. The math (Little's law): to sustain about 240 GB/s with about 600 ns of memory latency, you need about 150 KB of loads in flight. 10,240 threads × 16 bytes ≈ 164 KB. Already enough. On an H100, with about 10× the bandwidth, that won't be true, and I want to rerun this sweep there later.
The bug of the day
My first grid-stride kernel hung forever. I'd written the stride as blockIdx.x * blockDim.x instead of gridDim.x * blockDim.x. For block 0, that's a stride of zero, so i += 0 and those threads never finish. A one-word bug, an infinite loop.
Next
- Measure peak FP32 compute with an FMA microbenchmark (and see how badly one accumulator does vs eight)
- Profile scalar vs
float4to explain the slowdown - Build the full roofline sheet for the T4