Getting the V100 tensor cores back from Triton
A lot of serious GPU capacity still sits in V100s: university clusters, second-hand servers, cloud instances priced down. They have tensor cores. Triton, on that architecture, does not use them. This post is about how nbxtensor — the tensor library under NeuroBrix's Triton engine — turns them back on, and why it only does so where a measurement says it helps. The gap Triton is the kernel language…
Tesla V100 GPUs still possess tensor cores, which are crucial for certain matrix products in modern models. However, NeuroBrix's Triton engine, written in kernel language, does not utilize these tensor cores by default. The reason is that a matrix product written with tl.dot, a common operation in most models, compiles to FMA instructions on Volta (sm_70) and does not fully leverage the tensor cores.
Instead, the tensor cores remain idle, leading to lower performance. To address this issue, nbxtensor, the tensor library under NeuroBrix's Triton engine, introduces hand-written kernels that utilize the mma.sync.m8n8k4 instruction. These kernels are designed specifically for matrix products, convolutions, flash attention, and the fused mixture-of-experts GEMM used in models like DeepSeek-MoE and Qwen3-30B-A3B.
The decision to enable these kernels is not global but rather operation and shape class-specific. For each operation and shape class, both the portable Triton kernel and the specialized m8n8k4 kernel are evaluated against a float64 oracle. The kernel that passes the oracle and achieves the best timing on the GPU class is selected and recorded in a certification.
At runtime, the engine serves the kernel recommended by the certification. This approach ensures that the V100 GPUs are used to their full potential, utilizing their tensor cores when it benefits performance, while maintaining compatibility with other GPU architectures and classes. The result is improved performance for models dominated by matrix products, without the need for users to purchase new hardware.
This strategy can be generalized to other GPU classes, with specialized kernels serving as additional candidates during certification.
Written by urgent.news from Dev.to's reporting — not their text. Machine-written — may contain errors; check the original before relying on it.
