Inspiration
Modern attention research is a Hopper/Blackwell party — FlashAttention-3/4, TMA, warp specialization, fp8. But the GPUs most people can actually get for free (a Kaggle Tesla T4) are Turing parts from 2018: fp16 tensor cores, 320 GB/s, 64 KB shared memory, no cp.async, no bf16. We wanted to find out how much of the modern playbook survives the trip back to sm_75 — and how fast a Transformer layer can go there without giving up numerical equivalence. The correctness bar made it interesting: every output element must satisfy $|u_i - r_i| \le 0.002$ or $|u_i - r_i| \le 0.02,|r_i|$ against the fp32 baseline, across 14 shapes from batch $1$ to $10^4$ and sequences from $32$ to $10^5$ tokens.
What it does
A drop-in replacement for the organisers' baseline Transformer layer that runs a 7.11x geometric-mean speedup across all 13 harness-runnable appendix shapes on a Tesla T4, with every case passing the accuracy criterion — measured entirely by the organisers' unmodified benchmark. Per-regime highlights: batch-1 latency drops $2.9,\text{ms} \to 0.30,\text{ms}$ ($9.5\times$, whole forward replayed as one CUDA graph), the $S=1024$ case hits $17\times$, and appendix case 14 — where the baseline would materialise a $B \cdot H \cdot S^2 = 32 \cdot 16 \cdot (10^5)^2$ fp32 score tensor $\approx 20$ TB, ungradable on any hardware — runs out of core and passes against a validated fp32 proxy over all $3.28\times10^9$ output elements.
How we built it
Three measured passes, each aimed at what profiling of the last one exposed:
Restructure the graph. Fuse Q/K/V into one GEMM; replace the materialised $[B,H,S,S]$ score tensor with memory-efficient SDPA (CUTLASS); pre-cast weights once and run GEMMs on fp16 tensor cores while the residual stream, LayerNorm statistics, and softmax stay fp32. The tolerance analysis drove the precision split: outputs are $\approx \mathcal{N}(0,1)$ after the final LayerNorm, so the exposed band is $|r| \lesssim 0.1$ where only the absolute budget applies. → $3.56\times$ geomean. Custom Triton kernels + launch elimination. At $d=128$ even the GEMMs are bandwidth-bound — arithmetic intensity $\frac{KN}{K+N} \approx 96$ flops/byte against the T4's $\approx 203$ ridge — so the add→LayerNorm→cast chain between them (nine kernels per block) dominated. One fused Triton kernel per residual+LayerNorm site ($2.62 \to 1.40$ ms on the reference shape alone), manual CUDA-graph capture for launch-bound shapes, and a data-driven per-shape dispatch table. → $7.11\times$ geomean. Research-driven experiments. Surveyed what transfers to sm_75, built a FlashAttention-2-style Triton attention kernel (fp32 accumulate, exact online softmax, direct $[B,S,H\cdot d_h]$ output), and ran a 25-trial × 2-seed accuracy stress program — all measured on a second T4 session, with an AI-assisted workflow (Claude Code) whose every commit links its session.
Challenges we ran into
No GPU where the code was written. Development ran in CPU-only sessions; we validated kernels numerically through Triton's numpy interpreter and bit-exact CPU harness runs before every T4 measurement session. Worst-element error is an extreme-value statistic. Max error grows roughly like $\sigma\sqrt{2\ln N}$ with sample count $N$, so official 5-trial passes lied about the margin: case 7 ($d{=}32$) failed a 25-trial stress at $0.00208 > 0.002$, and case 6 ($1.64\times10^8$ elements/trial) failed at $0.002074$. Fix: shape dispatch — $d < 64$ runs fp32 ($\sim 2\times10^{-6}$ error, still $3.9\times$). CUDA graphs overwrote our lazily-built weight cache (allocations captured inside the graph region were clobbered on replay) — fixed by forcing cache builders outside compiled/captured regions. Two parallel workstreams collided: a second session independently built the same LN fusion; merging meant reconciling duplicate kernels, flags, and test suites without losing either side's measurements.
Accomplishments that we're proud of
$7.11\times$ geomean on the organisers' harness with zero failed elements, and a $1.000\times$ measured noise floor backing every claim. The fused-LN win was replicated independently — two implementations, two sessions, same result. Making the ungradable case 14 pass out of core over $3.28$ B elements ($\max|u-r| = 0.00102$). Publishing our negative results with data: the hand-written attention kernel loses to CUTLASS SDPA on every shape (7.47 vs 3.47 ms on case 11), and torch.compile was retired once the custom kernels beat it everywhere.
What we learned
Measure, don't assume — our attention kernel and torch.compile both looked good on paper and lost on silicon (without cp.async, Triton's pipeliner can't hide load latency the way hand-pipelined CUTLASS does). Fusion wins came as much from eliminating kernel launches as from saving DRAM bytes — our traffic-only cost model under-predicted the LN fusion by an order of magnitude. And accuracy margins must be treated as distributions, not point estimates: stress-test beyond the official trials or get burned by $\sqrt{2\ln N}$.
What's next for Transformerkernel: a 7.1x Transformer layer for the Tesla T4
Price the seed-robust fp32 dispatch for case 6 (--fp16-max-elements) and settle case 7's default-vs-compiled head-to-head — the follow-up notebook (F1–F4) is already written. Tune the attention kernel's known weaknesses ($\exp_2$ softmax, tile/warp retuning, single-phase causal loop) — or let its documented loss stand as the honest answer for sm_75. Explore INT8 $QK^\top$ (SageAttention-style) on Turing's unused INT8 tensor cores — the one quantised-attention idea that could physically run here, deliberately deferred until the fp16 margin story was nailed down.
Log in or sign up for Devpost to join the conversation.