Skip to content

Support pre-Volta GPUs (sm_61 / Pascal) - #19

Open
gdevenyi wants to merge 5 commits into
FlashML-org:mainfrom
gdevenyi:pascal-sm61-support
Open

Support pre-Volta GPUs (sm_61 / Pascal)#19
gdevenyi wants to merge 5 commits into
FlashML-org:mainfrom
gdevenyi:pascal-sm61-support

Conversation

@gdevenyi

@gdevenyi gdevenyi commented Aug 22, 2026

Copy link
Copy Markdown

FreeToken cannot currently be installed or run on a pre-Volta GPU. This branch gets it serving on a GTX 1080 Ti (sm_61, Pascal). It does not change what any supported GPU does.

Result

Serving Qwen/Qwen2.5-0.5B-Instruct on a GTX 1080 Ti, with default sampling:

$ curl -s localhost:8099/v1/chat/completions -d '{"model":"Qwen2.5-0.5B-Instruct",
    "messages":[{"role":"user","content":"Name three primary colours."}],"max_tokens":48}'
{"choices":[{"message":{"role":"assistant",
  "content":"Three primary colours are red, blue, and yellow."},
  "finish_reason":"stop"}],
 "usage":{"prompt_tokens":34,"completion_tokens":12,"total_tokens":46}}

Test suite on that card: 66 failed and 1270 passed becomes 33 failed and 1327 passed, with 29 skipped.

What was blocking

Five separate issues. Each is a construct that NVIDIA introduced after Pascal.

blocker fix
1 __grid_constant__ (compute_70 and up) and ld.global.L1::no_allocate (sm_70 and up) in the JIT kernels. nvcc and ptxas reject the whole translation unit, so the offload gather, the index kernels and the KV store never build. Gate both on __CUDA_ARCH__ < 700, as __nanosleep in the same header already is.
2 tanh.approx.f32 (sm_75 and up) in _fast_tanh. A constexpr_function gate and a libdevice fallback, following what e4m3_compat.py already does.
3 tl.atomic_add in moe_align. Triton lowers every atomic to a scoped and ordered PTX encoding, all of which need sm_70. Route to the existing atomic-free staged moe_align_block_size_triton.
4 Attention tiles overflow a 48 KB shared-memory budget. Extend the existing budget ladder and let the compiler arbitrate.
5 tl.atomic_add and tl.atomic_max in the top-k and top-p threshold search. A torch sort-based fallback.

Both codegen hints in the first row are semantically empty. The parameter is passed identically without the annotation, and the load returns the same bytes without the cache hint.

Testing

I cannot re-test this on main myself. The machine I have access to holds two sm_89 cards and no pre-Volta GPU, so I can only confirm that the gates compile and that the sm_89 paths still behave. The GTX 1080 Ti figures above stand as originally measured.

Anyone with a Pascal or Maxwell card who can re-run the suite would be worth more to this PR than anything further I can do.

Note that a pre-Turing card also needs #26, because the pinned CUDA 13 torch build carries no cubin below sm_75.

gdevenyi and others added 5 commits August 22, 2026 00:58
Two device-side constructs in the JIT kernels are sm_70+ and make nvcc/ptxas
reject the whole translation unit on Pascal (sm_6x), so the offload gather, the
index kernels and the KV store never build there:

  - `__grid_constant__` on kernel parameters (compute_70+)
  - the `L1::no_allocate` modifier on `ld.global` (sm_70+)

Both are codegen hints with no semantic content: the parameter is passed
identically without the annotation, and the load returns the same bytes without
the cache hint. Gate each behind `__CUDA_ARCH__ < 700` so every sm_70+ device
pass and the host pass emit exactly what they did before. `__nanosleep` in the
same header was already guarded this way.

Also stop a failed CUDA call in the pinned-tensor extension from poisoning the
process: the runtime keeps a failure in the per-thread last-error slot, so the
next unrelated `C10_CUDA_CHECK` reported it instead of its own result. This was
visible as an unrelated `torch.empty` on CUDA raising "invalid argument" after
`host_device_ptr` rejected unregistered memory.

`cudaHostGetDevicePointer` on *unregistered* memory is unspecified -- newer
arches let UVA degenerate it to identity, Pascal validates registration and
returns cudaErrorInvalidValue. Widen the test to accept either, and keep
asserting the invariant that actually matters: it must never return a different
nonzero alias. The in-contract case stays covered by
test_host_bank_pin_registers_and_translates.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01KQy3DziJN9peJ7nA8L59ns
`_fast_tanh` emits `tanh.approx.f32` as inline PTX. That instruction is sm_75+,
and ptxas rejects the whole module without it ("Feature 'tanh' requires .target
sm_75 or higher"), so every GELU_TANH activation failed to compile on Pascal.

Gate it on a `constexpr_function` reading the compilation target, the same idiom
e4m3_compat.py uses to branch fp8e4nv below sm_89. Being compile-time, this adds
no kernel parameter and leaves the cache key untouched: sm_75+ still takes the
single-instruction path, folds the branch away, and emits the PTX it did before.
Older cards compute the same tanh through libdevice, which this file already
depends on for `libdevice.erf`.

Adds tests/kernels/test_activation.py: the four *_and_mul kernels against torch
references, plus a saturation case at +-1e4 that would catch a fallback that
overflows to nan at the tails. gelu_tanh had no direct coverage before -- it was
only reached transitively through the MoE tests.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01KQy3DziJN9peJ7nA8L59ns
`moe_align.py` ranks its scatter with `tl.atomic_add`. Triton lowers every
atomic to a scoped *and* ordered PTX encoding -- `atom.global.gpu.<sem>.add` --
and both the `.gpu` scope and every memory order (`.relaxed` included, the
weakest it can emit) arrived with sm_70. ptxas therefore rejects the entire
module on Pascal, and no `sem=`/`scope=` argument avoids it. Verified by
probing all four sems and all three scopes on an sm_61 device: every one fails
to assemble.

Rather than emulate atomics, route around them. `moe_impl.py` already carries a
second implementation -- the staged `moe_align_block_size_triton`, a
counts/cumsum/binary-search chain across five launches with no atomic anywhere.
It produces the same three buffers, so pre-sm_70 dispatches there. Checked
against a reference of the documented contract on sm_61 across decode and
prefill shapes (1..1024 rows, 8..257 experts, block 16..64): identical padding,
every token placed exactly once inside its own expert's region.

sm_70+ is untouched -- same branch, same kernel, same launch as before. The
staged path costs five launches instead of one, which is the right trade
against not running.

Not attempted: emulating the atomic with unqualified inline PTX
(`atom.global.add.u32` is sm_20+). It assembles and gives correct ranks, but
`tl.inline_asm_elementwise` is contracted for *pure* elementwise ops -- when the
operand tensor is narrower than the thread block, the layout replicates elements
across threads and a side-effecting instruction executes once per replica. A
128-thread block over an 8-wide tensor adds 16x. That silently miscounts at
exactly the decode widths that matter (numel = batch x topk), so it is not a
safe basis for expert dispatch.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01KQy3DziJN9peJ7nA8L59ns
… budget

`_select_extend_tile` already shrank the prefill tile when a device's opt-in shared
memory could not hold it, but two gaps let it overflow anyway:

  - head_dim <= 128 returned (128, 64) before any budget check, so the most common
    head sizes never consulted the device at all;
  - the ladder stopped one rung short of what a 48KB budget needs (pre-Volta gets
    48KB per block with no opt-in at all).

Carry the ladder down to (16, 16) -- the floor, since `tl.dot` needs N >= 16 -- and
budget-check every rung including the largest.

The byte estimate is also only a lower bound: what triton allocates depends on how it
schedules the pipeline, and the split kernel wants about twice the q/k/v tile bytes.
A head_dim 64 model passed the estimate and still needed 64KB. Rather than carry a
per-kernel fudge factor, `_select_extend_tiles` now returns the descending ladder and
`_launch_first_fitting_tile` walks it, letting the compiler have the final say:
OutOfResources is raised during launch setup, after the compile and before any GPU
work, so retrying smaller is side-effect free and triton caches each compile.

The split-k decode kernel had no sizing at all: BLOCK_N=32 with num_stages=2
hardcoded, whose pipelined k/v tiles want 128KB at head_dim 512 -- fine on A100/H100,
over budget on consumer Ampere/Ada, never mind Pascal. Add `_select_decode_tile`,
stepping 32/2 -> 32/1 -> 16/1.

Both selectors only ever shrink a config that does not fit, and the existing
tile-selection test still pins the exact choice for every datacenter and consumer
budget it covered before, so any device that fits the current default keeps it.
`smem_optin == 0` (unavailable) keeps the prior choice. Devices that previously
raised OutOfResources here may now run.

head_dim 512 extend has no fitting configuration on a 48KB device rather than a
merely slower one -- 16x16 q/k/v tiles alone want 48KB. Its test now skips with that
reason, gated on the device's actual budget so it still runs everywhere it can.

On sm_61: tests/kernels/test_triton_attention.py goes from 9 failed to 35 passed,
2 skipped.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01KQy3DziJN9peJ7nA8L59ns
The triton top-p/top-k threshold search reduces through `tl.atomic_add` and
`tl.atomic_max`. Triton lowers every atomic to a scoped, ordered PTX encoding, and
both the scope and every memory order it can emit arrived with sm_70, so ptxas
rejects those kernels on Pascal. No unit test covered them, so this surfaced only
when serving: `top_k` and `top_p` are set by default for most models, and the first
real request killed the scheduler with a raw ptxas dump.

Sorting needs no atomics. `kernel/torch_sampling.py` reproduces the same selection
with torch ops -- top-k by the k-th largest value, top-p by the shortest descending
prefix that reaches p, both keeping the elements the kernel path keeps, including
ties and the crossing element. `softmax` has no atomics and is re-exported from the
triton module unchanged, so only the threshold search changes.

The draw uses cumsum + searchsorted rather than torch.multinomial: it stays
capturable in a CUDA graph and never syncs to the host, matching what the kernel path
guarantees. Seeding mirrors `triton.sampling._gen_u`, including its plain-torch.rand
path while a stream is capturing.

Selected in `sample_impl` only when flashinfer is absent *and* the device is
pre-sm_70, so every currently supported GPU keeps the kernel path untouched.

Full-vocabulary sorting is slower than the bracketed histogram search it replaces.
On a card that cannot run the kernels at all, sampling is not the bottleneck.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01KQy3DziJN9peJ7nA8L59ns
@apcameron

apcameron commented Aug 24, 2026

Copy link
Copy Markdown

I can confirm that this PR works on my Tesla P40 which is a Pascal Based GPU with 24GB VRAM on my computer with 128GB RAM
I used the following
hf download openai/gpt-oss-120b --local-dir models/openai/gpt-oss-120b
ft serve --model ./models/openai/gpt-oss-120b --host 0.0.0.0 --port 5001

ft shell --server 192.168.0.100:5001
FreeToken shell -> gpt-oss-120b @ http://192.168.0.100:5001 (/help for commands, /exit to quit)

what is the value to PI to 100 digits

Thinking...
| Thus π = 3.1415926535 8979323846 2643383279 5028841971 6939937510 5820974944 5
| 923078164 0628620899 8628034825 3421170679
|
| We can present as a single string without spaces. Also note that the 100th dig
| it after decimal is 9? Let's verify: The 100th digit after decimal is the last
| digit of the sequence above: 9. Yes.
|
| Thus answer. Also maybe mention that it's irrational, etc. Provide reference.
End Think
Here is π (pi) written out to 100 decimal places:

[
\pi = 3.!1415926535;8979323846;2643383279;5028841971;6939937510;5820974944
;5923078164;0628620899;8628034825;3421170679
]

(Spaces are added only for readability; the 100‑digit fractional part is the str
ing after the decimal point.)

Why 100 digits?
π is an irrational, transcendental number, so its decimal expansion never termia
tes or repeats. Computing many digits is a classic test for numerical algorithm
and computer hardware. The first 100 digits are often quoted for reference or fo
r high‑precision calculations in scientific work.
$

[done] gpt-oss-120b (medium) | ↓77 ↑497 4.7 tok/s | cache 1255 lru 27.2% | kv 57

@erichstuntebeck

Copy link
Copy Markdown

This PR also fixes a Turing (sm_75) crash, which the title undersells — worth noting since
the PR reads as Pascal-only scope.

decode_paged_attention launches _decode_grouped_stage1_kernel with a hardcoded
BLOCK_N=32 / num_stages=2 and no shared-memory check, while the extend/prefill sibling has
been smem-aware via _select_extend_tile all along. On a 64 KiB opt-in device that overflows
once next_power_of_2(head_dim) > 256. _select_decode_tile() here is exactly what was
missing.

Verified on a Tesla T4, running upstream's own decode parametrization
(test_decode_triton_attention_matches_reference: (256, 8, 3) and (512, 2, None)) against
three trees, each with a fresh TRITON_CACHE_DIR and a never-warmed cache volume:

tree head_dim=256, kv=8 head_dim=512, kv=2
main @ 2757bb5 OK FAILOutOfResources: Required: 81920, Hardware limit: 65536
#24 @ 35668da ("support Turing GPUs") OK FAIL — identical, Required: 81920
this PR @ c8e4c08 OK OKout(2, 16, 512) finite=True

That middle row is the surprising one: the PR named "fix(cuda): support Turing GPUs" does not
fix this, because its attention.py hunk only touches _select_extend_tile — the extend path.
gh pr diff 24 | grep -i decode returns nothing. Anyone triaging a Turing decode crash would
reasonably assume #24 covers it.

Boundary on main (sm_75), showing it is next_power_of_2(head_dim) that matters, not head_dim:

device Tesla T4 sm_75  smem_optin=65536
  OK    head_dim=128 (BLOCK_D=128) num_kv_heads=2
  OK    head_dim=256 (BLOCK_D=256) num_kv_heads=2
  FAIL  head_dim=320 (BLOCK_D=512) num_kv_heads=2  Required: 81920, limit: 65536
  FAIL  head_dim=384 (BLOCK_D=512) num_kv_heads=2  Required: 81920, limit: 65536
  FAIL  head_dim=512 (BLOCK_D=512) num_kv_heads=2  Required: 81920, limit: 65536
  FAIL  head_dim=512 (BLOCK_D=512) num_kv_heads=8  Required: 69632, limit: 65536

One detail that may be useful for the tile ladder: on this device BLOCK_N is the only lever
that matters. Taking main and changing only num_stages=2 → 1 leaves the failure
byte-identical (still Required: 81920); changing only BLOCK_N=32 → 16 makes it pass, and
matches upstream's _reference_paged_attention at upstream tolerance
(max_abs_diff=1.562e-02, atol=rtol=2e-2). Turing has no cp.async, so num_stages buys
nothing there regardless.

Practical impact: this is not an exotic geometry. gemma-4-26B-A4B's full-attention layers are
key_length=512, head_count=16, head_count_kv=2 — the (512, 2) row above — so serving that
checkpoint on any ≤64 KiB opt-in card dies in decode on main today. I have been carrying a
narrower local patch (BLOCK_N 32→16 below a 99 KiB cutoff); this PR's ladder is the better
form of it and I would rather it land than carry mine.

Reproduces only on devices with ≤64 KiB opt-in shared memory (sm_75 and below) — 81920
fits comfortably under Ampere+/sm_89's ~99 KiB, so the same script passes on A100/H100/RTX 4090.

Environment: FreeToken 0.1.2 from source · Tesla T4 16 GiB sm_75 · driver 580.178.04 ·
CUDA 13.0.88 · torch 2.11.0+cu130 · triton 3.6.0 · Ubuntu 24.04.4.

@CodingInCarhartts

Copy link
Copy Markdown

I tested this on a GTX 1060 6GB, Pascal sm_61, and can confirm openai/gpt-oss-20b serves end to end with this branch.

Hardware was a GTX 1060 6GB with about 5.1 GiB free VRAM after the Windows desktop, i9-9900K, 32 GB RAM, Docker Desktop with WSL2. The container was running driver 582.28, CUDA 12.6, Python 3.12.14, torch 2.11.0+cu126, and Triton 3.6.0.

I tested the merge of gdevenyi:pascal-sm61-support at c8e4c08e6548 plus pyproject-cu126-group at fb2f80d9c52f from #26, on top of main@0ab982f10905.

Model was openai/gpt-oss-20b using the MXFP4 safetensors with default context settings. KV auto-sized to 8205 tokens.

Kernel tests

tests/kernels finished with:

88 passed
15 skipped
1 failed

One failure is:

test_minimax_m3_sparse.py::test_prefill_index_score_and_topk

Triton OutOfResources
Required: 98304
Hardware limit: 49152

kernel needs 96 KiB of shared memory per block. Pascal only has 48 KiB, so it cannot fit on sm_61.

gpt-oss does not use that path, so it does not prevent serving.

Serving

End-to-end serving works.

Auto-selection picked:

attention=triton
moe=offload
dtype=bfloat16

CUDA graph capture also succeeded at batch sizes 1, 2, and 4.

The main thing I found is that --moe-backend hybrid matters a lot on Pascal.

ft bench bw measured:

CPU STREAM:       45.6 GB/s
PCIe H2D: 12.6 GB/s
mxfp4 CPU-MoE: 20.9 GB/s
PCIe gather: 9.0 GB/s

CPU-MoE is about 2.32x faster than PCIe gather

The problem is that auto still selected offload because the bench profile was written inside the ephemeral container and did not persist between runs.

Measured difference:

Metric offload hybrid
Startup prefill warmup [80,128] 1228 s 58 s
TTFT, first ~1.5k-token prompt 1158 s 46 s
TTFT, same prompt with warm cache n/a 1.3 s
Steady decode 4.4 to 4.7 tok/s 7.3 to 7.6 tok/s

Offload made the first request look basically hung. Hybrid made the same setup usable.

VRAM after loading was:

5535 MiB / 6144 MiB

That leaves about 0.4 GiB of headroom. Host RAM usage for the container was about 11.4 GiB.

With hybrid and a warm Triton cache, repeat TTFT landed around 1.3 to 2.9 seconds.

fp16

I also tried forcing:

--dtype float16

That does not boot gpt-oss-20b on this card.

It OOMs during weight materialization:

sharded_tensors -> get_tensor
RuntimeError: CUDA driver error: out of memory

Auto-selected bf16 loads successfully with the same amount of free VRAM.

I reproduced the fp16 OOM twice.

PR #26

One separate issue I hit while testing #26.

The group-scoped uv source does not appear to control torch's normal [project.dependencies] edge.

Running:

uv sync --no-default-groups --group cu126 --python 3.12

resolved torch from PyPI instead of the cu126 index, which pulled the cu130 build. The nvcc/torch toolchain check then failed.

For testing I worked around it by pointing torch's source entry directly at the cu126 index.

Reproduction

Docker base:

nvidia/cuda:12.6.0-devel-ubuntu22.04

The devel image is required because _pinned_tensor links against the CUDA runtime API and needs CUDA_HOME and nvcc during the build.

ninja is also needed at runtime for the JIT index kernel.

Setup was:

git clone <fork>
cd FreeToken
git checkout <branch>
git merge <pr-26-branch>

uv sync --no-default-groups --group cu126 --python 3.12
uv pip install pytest ninja

For Pascal I would persist /root/.cache/freetoken, run:

ft bench bw

once, then either let auto use the saved profile or explicitly serve with:

ft serve --moe-backend hybrid

All numbers above are from this one machine at temperature 0. Repeated runs produced byte-identical token counts.

@KodeMunkie

Copy link
Copy Markdown

I have 6x Tesla P4 (sm_61) in a single box, and an external P100 (sm_60) so I'll also test this PR as soon as I get chance (too excited for Qwen 3.8 Flash MoE at the moment, so that's taking priority).

Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

5 participants