Outrageously Fast Training on Volta: Hand-Written Tensor-Core Kernels for Language Models on the V100
Year: 2026
Type: article
Abstract: The software stack for training transformers has moved on from NVIDIA Volta. FlashAttention 2 and 3 do not support sm_70, Volta has no bf16, and Triton's matrix-multiply primitive does not use Volta's tensor cores. As a result, PyTorch on a V100 falls back to slower attention and to unfused elementwise kernels, even though V100s are now among the cheapest datacenter GPUs available. We present hand-written CUDA and Triton kernels for training language models on a single V100: a variable-length causal attention kernel built directly on the mma.sync.m8n8k4 tensor-core instruction with fp32 accumulation, a fused Canon-convolution and rotary-embedding kernel, and a fused cross-entropy head that never materializes the logits. Our attention is 1.9x (forward) and 2.5x (backward) faster than PyTorch's memory-efficient attention. Every kernel we measure for precision (attention, the Canon+RoPE forward and the LM head) matches or beats the stock path's error against an fp64 reference. End to end, pretraining a 143.7M-parameter model runs 1.52x faster than a torch.compile baseline at equal micro-batch. Because the kernels also need less memory, they fit the production micro-batch the baseline cannot, which raises the gain to 1.59x: 59.2k tokens/s, 50% of the card's real tensor-core peak. We also report the optimizations that did not work, with measurements, and an fp16 overflow that ends training runs and that Volta, lacking bf16, cannot sidestep.
Author: Jones, Joseph Daniel
Language: English
Subfield: Hardware and Architecture
Field: Computer Science
Domain: Physical Sciences
Open Access status: green
OpenAlex ID: W7225455616