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