INVESTIGATION_L2.md

# Investigation: "why is torch faster in the L2 regime?"

Triggered by wall-clock runs showing our ReLU ~3–6% behind `torch.relu` for
L2-resident sizes. Result: **the premise is mostly false for floats (measurement
noise) and real-but-small for int8 (a launch-parallelism difference, not ALU).**

## Finding 0 — wall-clock micro-benchmarking is unreliable here

The same call (`opt.ops.relu(out, x)`, fp16 6M) measured **3544 GB/s** in one run
and **662 GB/s** in another — a **5× swing** — because L2-resident kernels are ~7 µs
and the RTX 4090's clocks boost/throttle by ±2× under sustained vs bursty load.
Wall-clock also includes Python wrapper overhead (`empty_like` + op dispatch).
**Use nsys/ncu GPU kernel durations (clock-paired by interleaving) or lock clocks
(`sudo nvidia-smi -lgc`).** All conclusions below use nsys GPU durations with ours
and torch launches interleaved 1:1 (same clock state).

## Finding 1 — fp16/bf16: torch is NOT faster

nsys, fp16 6M, 350 launches each, interleaved:

| kernel | avg | median | stddev |
|--------|---:|---:|---:|
| torch `vectorized_elementwise<vec4>` | 7226 ns | 7232 | 78.0 |
| ours `relu_vec<__half>` | **7151 ns** | **7137** | **55.7** |

Ours is **marginally faster** (1.01×) with lower variance. The apparent "torch
wins" was wall-clock noise. Same for bf16.

## Finding 2 — int8: torch genuinely ~6.6% faster, and why

nsys, int8 16.7M (33 MB ws, L2-resident), interleaved:

| kernel | block | grid | threads | elem/thread | GPU time |
|--------|:---:|:---:|:---:|:---:|:---:|
| torch | 128 | 16384 | 2.1M | ~8 | **9383 ns** |
| ours | 512 | 2048 | 1.05M | 16 | 10005 ns |

torch launches **2× more threads / 8× more blocks** (128-thread blocks, ~8 elem
per thread) vs our 512-thread blocks at 16 elem/thread (int4). In the L2 regime
that extra memory-level parallelism wins.

**Ruled out — ALU.** Hypothesis "our scalar per-byte relu is the cost" was tested
by switching int8 to `__vmaxs4` (SIMD signed-byte max, 4 bytes/instruction). It
moved int8 only 10192→10005 ns (~1.8%) — confirming int8 is memory-bound, not
ALU-bound. (`__vmaxs4` kept anyway: correct, slightly faster, no downside.)

**Confirmed — launch parallelism.** A granularity sweep (bytes/thread × block, int8
@16.7M) ranks **8 B/thread blk256 (3497 GB/s)** and **16 B/thread blk128 (3496)**
above our 16 B/thread blk512 (3447) — i.e. smaller blocks and/or less work per
thread, matching torch's config direction. Recovers ~1.5%; the residual is torch's
finely-tuned TensorIterator heuristic (128-thread, thread_work_size grid-stride).

## Conclusion

- For **fp16/bf16/fp32** there is **no real L2 deficit** — at the GPU level our
  kernel matches or slightly beats `torch.relu`.
- For **int8** torch is ~6.6% faster in the L2 regime due to a **more parallel
  launch** (more/smaller blocks, fewer elements/thread), not compute. Config tuning
  closes ~1.5%; fully matching torch means replicating its launch heuristic.
- **Methodological takeaway:** never rank ~µs L2-resident kernels by wall clock on
  a boosting consumer GPU; use clock-paired nsys/ncu durations or lock clocks.

## Reproduce
```
# clock-paired kernel durations
nsys profile -o p --trace=cuda,nvtx python -c "<interleave torch.relu & ours>"
nsys stats --report cuda_gpu_kern_sum p.nsys-rep
nsys stats --report cuda_gpu_trace --format csv p.nsys-rep   # grid/block
# int8 granularity sweep
cd bench && nvcc -arch=sm_89 -O3 int8opt.cu -o int8opt && ./int8opt
```