Feat/sm75 cuda128 - #157
Closed
JohnieBraaf wants to merge 2 commits into
Closed
Conversation
added 2 commits
August 24, 2026 22:02
- pyproject.toml: swap cu130 index to cu128 for torch + sglang-kernel; swap flashinfer[cu13] to flashinfer[cu12]; add CUDA 12 classifier - moe/nvfp4_backends.py: add explicit sm_75 guard in select_nvfp4_backend() with clear warning directing users to Q4_0 / hybrid backend - utils/arch.py: add default_compute_dtype() — returns float16 on sm_75 (Turing has no BF16 ALUs), bfloat16 on sm_80+ (Ampere+) - setup.py: default TORCH_CUDA_ARCH_LIST starts at 7.5 (was implicit 8.0+) - moe/configs/triton_3_5_1/: add NVIDIA_GeForce_RTX_2080_Ti configs for E=256/N=2048 (DeepSeek-V4-Flash), E=128/N=768 (Qwen3-30B-A3B), E=64/N=2048; all use num_stages=2 (cp.async requires sm_80+, Turing uses synchronous shared-memory loads) - flake.nix: NixOS unstable dev shell with cudaPackages_12_8, gcc12, NCCL, TORCH_CUDA_ARCH_LIST=7.5;..., NCCL_P2P_LEVEL=NVL - scripts/install-sm75.sh: bootstrap install for CUDA 12.8 environment - docs/install.md: add Turing GPU support section with format support table, hardware requirements, NixOS/manual install, run commands, expected perf, and known limitations Hardware target: 4x RTX 2080 Ti 22GB modded + 2x NVLink + AMD 3975WX + 256GB RAM + NixOS unstable. Model target: DeepSeek-V4-Flash Q4_0 via --moe-backend hybrid --tensor-parallel-size 4 --dtype float16. Validated baseline (PocketLLM on identical hardware): ~3.7 tok/s decode, ~401 tok/s prefill @ 32K-64K context.
Wires the Marlin WNA16 kernel port from weicj/vLLM-2080Ti-Definitive into FreeToken as an optional build-time extension for sm_75 (RTX 2080 Ti) GPUs. Key design: - scripts/fetch_marlin_sm75.py: fetches ops.cu, marlin_template.h, kernel.h, generate_kernels.py + transitive includes from vLLM-2080Ti-Definitive at build time; runs generate_kernels.py with arch=7.5 to produce sm75_kernel_*.cu (stages=2) and kernel_selector.h - python/freetoken/kernel/csrc/marlin_wna16/marlin_sm75_ext.cpp: pybind11 wrapper exposing marlin_mm() as freetoken.kernel._marlin_sm75.marlin_mm - setup.py: CUDAExtension for freetoken.kernel._marlin_sm75; compiled with -gencode arch=compute_75,code=sm_75 only; fetch is skipped gracefully on network failure or FREETOKEN_SKIP_MARLIN_SM75=1 - python/freetoken/moe/marlin_sm75.py: Python dispatch layer; is_marlin_sm75_applicable() guards cc==(7,5) + int32 weights + FP16/INT8 activation + extension present; fused_experts_marlin_sm75() handles gate+up projection, activation split, down projection, topk weighted sum - python/freetoken/moe/fused.py: fused_experts_impl() routes to fused_experts_marlin_sm75() when w1_scales/w2_scales are provided and is_marlin_sm75_applicable() returns True; falls through to Triton otherwise Why stages=2: Turing (sm_75) lacks cp.async (async shared-memory copy, requires sm_80+). The 4-stage Marlin pipeline uses cp.async for prefetching; on sm_75 this is replaced with a 2-stage synchronous pipeline. generate_kernels.py sets stages=2 for all sm_75 instantiations. Runtime ops.cu already guards this with TORCH_CHECK(stages=2 for major==7 && minor==5). Supported formats on sm_75 via this extension: AWQ-INT4 (kU4), GPTQ-INT4 (kU4B8), AWQ-INT8 (kU8B128) All with FP16 activation + FP16 output only. FP8 activation (kFE4M3fn) is NOT built for sm_75 (requires sm_89+).
Author
|
sorry wrong repo |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.