Tensor-core dots on Volta (sm_70) and Turing (sm_75): a patch series on Triton 3.6

[Peuqui]

Hello,

this is not a request to change Triton's support policy. I know that sm_70 and sm_75 are not supported targets (the README lists compute capability 8.0+; see also #9970 and #10918), and I am not asking for a review. I only want to point people who still run these cards to a fork that works for us, in case it saves someone the work.

Since #5066 (Triton 3.3) every tl.dot below compute capability 8.0 compiles to FMA, and #5099 removed the Volta code. We serve large models on three Tesla V100 and two Quadro RTX 8000 and needed the tensor cores there. The result is 23 commits on top of v3.6.0, the Triton version our PyTorch 2.10 build ships with:

https://github.com/Peuqui/triton/tree/sm7x-tensor-cores

What it does:

• Turing: the MMAv2 m16n8k8 lowering is enabled for sm_75 again, with software pipelining through registers (there is no cp.async before sm_80). Batched (3D) dots use MMAv2 too.
• Volta: MMAv1 (mma.m8n8k4) expressed through linear layouts. The fragment bases live in the linear layout, the LLVM lowering reads the register slots from it, row/col is chosen per operand from the memory order, plus a 16-row shared-memory swizzle and an 8x32 tile for dots with at most 8 rows.
• Fixes needed on the way: warp reductions and delinearize assumed that the lane bits along an axis are contiguous (Volta places rows on lane bits 0, 2 and 4), which broke online softmax; main has reworked the reduction lowering since ([BACKEND] Simplify ReduceOpHelper and reduction lowering #12013). FMA dots with an fp16/bf16 accumulator now sum in fp32, and FMA dots widen int8 operands as integers.
• Pre-sm_80 performance passes: small loop-invariant dot operands kept in registers by a register budget, shared memory double-buffered through registers, and decode chains (elementwise ops between load and dot) staged in the register buffer.

Measurements on our machine (fp16):

GEMM M x 4096 x 4096 at M = 16 / 128 / 512 / 4096, TFLOPS (higher is better), measured at f1c961e (the three commits after it change the decode and rearrangement paths):

• Quadro RTX 8000: fork 8.9 / 51.5 / 75.4 / 69.4; unmodified 3.6.0 (FMA) 3.0 / 6.6 / 6.8 / 7.5; cuBLAS 7.3 / 34.0 / 92.9 / 82.0
• Tesla V100: fork 10.5 / 48.0 / 71.9 / 76.4; unmodified 3.6.0 (FMA) 3.9 / 6.8 / 8.0 / 8.6; cuBLAS 8.1 / 44.8 / 78.8 / 83.8

Causal flash-attention forward, a chunk of 2,048 queries at the end of a 32,768-token context, 12 query heads, 2 KV heads, head dim 256 (ms, lower is better):

• Quadro RTX 8000: fork 14.45 ms (55.3 TFLOPS); hand-written FA2 Turing build 24.4 to 30.8 ms; unmodified 3.6.0 561.9 ms (1.4 TFLOPS)
• Tesla V100: fork 16.79 ms (47.6 TFLOPS); hand-written V100 flash-attention kernel 75.4 to 75.5 ms (10.6 TFLOPS); unmodified 3.6.0 319.4 ms (2.5 TFLOPS)

Tests: python/test/unit/language/test_core.py -k dot at the head of the branch: 283 passed, 0 failed, 2329 skipped on each card. Unmodified 3.6.0 on the V100: 97 failed, 194 passed, 2321 skipped.

Not tested: the lit and C++ unit tests, and anything on sm_80 or newer. The changes are meant to stay below sm_80, but I have not checked that on newer cards.

We run this build in production in our vLLM fork (https://github.com/Peuqui/1Cat-vLLM/tree/fork-main-v4) on three V100 and two RTX 8000: the Triton kernels of DeepSeek-V4 (sparse-attention indexer and prefill) and GLM-5.3 (linear attention) use the tensor cores on both card types, and the DeepSeek-V4 MoE runs on Triton on the V100 stages.

If more people run Triton on these cards and a maintained pre-sm_80 path would help, I would be glad to hear about it. I am not proposing anything upstream at this point.

AI assistance (Claude) was used to write the patches and this text; every number above comes from our measurement logs.