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 NeuroBrix's Triton engine is written in. One kernel source serves NVIDIA, AMD, Intel and Apple GPUs, which is exactly why we use it. On Volta (sm_70), however, a matrix product written with tl.dot compiles to FMA instructions: ordinary fused multiply-adds on the CUDA cores. The tensor cores of the V100 stay idle. For a model whose time is dominated by matrix products — which is most of them — that is the difference between using the card and using part of it. The fix: m8n8k4 kernels Volta's tensor cores are driven by the mma.sync.m8n8k4 instruction: an 8x8 output tile accumulated from 8x4 by 4x8 operands, per group of threads. nbxtensor carries hand-written kernels built on that instruction for the operations that matter: matrix products, convolutions, flash attention, the fused mixture-of-experts GEMM used by models such as DeepSeek-MoE and Qwen3-30B-A3B. They live next to the portable Triton kernels, not instead of them. Chosen by measurement, per operation A tensor-core kernel is not automatically the faster one. For small shapes, odd shapes or memory-bound operations, the plain Triton kernel can win. So nbxtensor never switches globally. For each operation and shape class, both paths are candidates in the same certification: each candidate is checked against a float64 oracle, the ones that pass are timed on the GPU class, the winner is written into the certificate with its proof. At run time the engine serves whatever the certificate says. Where m8n8k4 measured faster, the V100 uses its tensor cores; where it did not, it does not. The choice is data, recorded and reviewable, not a rule in code. What it adds up to In NeuroBrix's showcase benchmark of 2026-11-08 on Tesla V100 (18 models, two cards, one run alone on the machine, median of 3), the Triton engine was faster than the PyTorch engine in 30 of 36 model-and-card pairs, up to 5.24x on Janus-Pro-7B; the full table, including the 6 pairs where it is slower, is in docs/BENCHMARKS.md. Why V100 owners should care A V100 is not a museum piece. It has fast HBM2 memory, NVLink on many servers, and tensor cores that were designed for exactly the matrix products modern models are made of. When a toolchain leaves those tensor cores idle, the card looks slower than it is. Bringing them back is about using hardware people already own instead of telling them to buy new hardware. Why it matters beyond Volta The pattern generalises: when the portable path leaves hardware on the table, add a specialised kernel as one more candidate, and let certification decide where it is used. The portable kernel stays the default everywhere the specialised one does not win, and every GPU class keeps its own certificate. Try it pip install neurobrix neurobrix run --model TinyLlama-1.1B-Chat-v1.0 --triton --prompt "Hello" --max-tokens 32 github.com/NeuroBrix/nbxtensor github.com/NeuroBrix/neurobrix The way we build it This post is one piece of a consistent method: One engine, zero per-model code. The runtime knows tensors, graphs and plans; it never knows which model it is running. A fix in the engine is a fix for every model. .nbx, one container for every model: graph and weights in one file format. Prism decides from facts: placement chosen from a hardware profile, not from heuristics in code. Certified, not autotuned: every kernel setting checked against a float64 oracle, timed once per GPU class, shared by pull request. Validated through the real entry point: all 49 models on both engines at every release (196 of 196 cells for 0.6.0 on 4x V100), outputs compared byte for byte with judged outputs, runs deterministic from one repetition to the next. The ecosystem NeuroBrix Engine: the inference engine, pip install neurobrix. nbxtensor: the tensor library without PyTorch, pip install nbxtensor. Model Hub: .nbx containers for 49 models at neurobrix.es/models, more with every release. triton-ext: our Triton build for Apple GPUs. Studio: a desktop application for Linux, Windows and macOS, in development. Where we are going Faster than the frameworks (kernel fusion, graph-level simplification), every processor (AMD and Intel GPUs, many-core CPUs, ARM64, accelerators beyond GPUs, multi-GPU across NVLink, PCIe and networks), community-certified hardware, training in nbxtensor, local agents, a local server and a JSON interface for applications, Studio for everyone, nbxtensor 2.0 in Rust, and .nbx as an open model standard. Get involved Star github.com/NeuroBrix/neurobrix and github.com/NeuroBrix/nbxtensor. Try it: pip install neurobrix. Certify your GPU class and share it (how). Open an issue when something breaks, or to request a model. Hocine Benkelaya, NeuroBrix