Views
No views yet
kernels ecosystem.t:h[t] = ( sum_{j = max(0, t-cf+1)}^{t} x[j] ) / min(t+1, cf)cf frames, count-normalized so sequence-start frames
divide by the real window size rather than a zero-diluted cf. It fuses the eager
cumsum → pad → subtract → divide → cast (≈5 full [B, D, T] passes + an fp32 round-trip)
into a single launch, using a two-cumsum identity s[t] = cumsum(x)[t] − cumsum(x delayed by cf)[t].1from kernels import get_kernel
2
3k = get_kernel("futo-org/causal-trailing-mean", version=1)
4h = k.causal_trailing_mean(x, cf) # x: [B, D, T] contiguous, any float dtype -> same shape/dtypeversion=1 pins the v1 build; omit it to track main (latest).x.dtype on store. Because it
accumulates in fp32 it is at least as accurate as the eager op in low precision. Parity vs a
fp32 eager reference: max|Δ| ≈ 1e-7 (fp32), 1e-3 (bf16), 2e-4 (fp16).T ≤ 4096; otherwise it falls back to an autograd-safe eager reference (also used on CPU).
eager_causal_trailing_mean is exported for that path.torch-cuda (the Triton kernel — the op only accelerates CUDA). The exported
eager_causal_trailing_mean covers CPU / grad-enabled / oversized-T in-process.cumsum → pad → subtract → divide → cast, ~5 full
[B,D,T] passes + an fp32 round-trip), measured with triton.testing.do_bench under
torch.no_grad(), bf16, cf=100, on an NVIDIA RTX PRO 6000 Blackwell:shape [B, D, T] | eager (ms) | kernel (ms) | speedup |
|---|---|---|---|
| 32 × 512 × 256 | 0.221 | 0.056 | 4.0× |
| 32 × 512 × 512 | 0.203 | 0.062 | 3.3× |
| 32 × 512 × 1024 | 0.828 | 0.127 | 6.5× |
| 64 × 512 × 2048 | 4.510 | 0.402 | 11.2× |
| 16 × 256 × 3000 | 0.545 | 0.137 | 4.0× |
T as the fused single-pass launch
displaces more redundant memory traffic); ~670 GB/s at the top end. Reproduce with
triton.testing.do_bench(lambda: causal_trailing_mean(x, cf)) vs the exported
eager_causal_trailing_mean.kernel-builder; Apache-2.0.