Skip to content

Releases: deepbeepmeep/kernels

GGUF CUDA kernels 1.0.25

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 25 Sep 18:34

GGUF CUDA 1.0.25 speeds up the PTQ1_0 (Bonsai ternary) short-batch tensor-core path used by speculative verification. Tensor cores take the raw base-3 digits and an exact integer correction replaces per-trit sign conversion, each thread extracts only the digits it uses, and activations are quantized in fragment order for 16-byte loads. Outputs are bit-identical to 1.0.24.

On an RTX 5090, PTQ1_0 verification linears take 28% less time and Bonsai speculative decoding runs about 14% faster (same generated text: 179.8 -> 204.6 tok/s with 4 MTP drafts). Q4_K and all other paths are unchanged. On other GPUs (compute capability 8.0+), WanGP keeps measuring the tensor-core path against MMVQ at model load and uses it only where it is faster.

Five wheels cover Python 3.10/PyTorch 2.7.1/CUDA 12.8 and Python 3.11/PyTorch 2.10/CUDA 13 on Windows and Linux (17 and 12 SASS targets respectively, plus PTX), and the Windows AMD HIP build for PyTorch 2.10.0+rocm7.14.0 (gfx1201). The CUDA wheels were rebuilt incrementally (only the changed kernel source was recompiled); the HIP build does not use the changed source, so its 1.0.24 wheel is relabelled.

Every CUDA wheel passed the toolkit-wide SASS/PTX inventory, bit-exact comparison against the 1.0.24 wheel of its stack with the default policy (existing GGUF formats, attention, PTQ1, Prism, fused paths), bit-identity of 259 PTQ1_0/Q4_K short-batch outputs against 1.0.24, and WanGP's short-batch tests. Other NVIDIA architectures are compiled but not hardware-tested; no AMD GPU was available.

Wheel hashes are listed in SHA256SUMS.txt; build and validation records are in llama.cpp/release/1.0.25/.

GGUF CUDA kernels 1.0.24

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 25 Sep 11:07

GGUF CUDA 1.0.24 adds a short-batch INT8 tensor-core linear for Q4_K and PTQ1_0 speculative verification (2-8 activation rows), for GPUs with compute capability 8.0 and newer. On an RTX 5090 it shortens a Qwen3.8 Q4_K_M decoding cycle with 4 MTP drafts by about 9% and a Bonsai PTQ1_0 cycle by about 30%, with unchanged or better accuracy and no additional weight memory.

A runtime policy selects the kernel: auto (default) applies per-shape decisions recorded by the caller and otherwise enables the new path only on compute capability 12.0 (RTX 50 / RTX PRO Blackwell); native keeps the 1.0.23 dispatch; mma forces the new path. LLAMACPP_GGUF_SHORT_BATCH overrides it at import. WanGP measures both kernels on the model's own weights before CUDA graph capture and records per-GPU choices, so other GPUs use the new path only where it is faster.

Five wheels cover Python 3.10/PyTorch 2.7.1/CUDA 12.8 and Python 3.11/PyTorch 2.10/CUDA 13 on Windows and Linux (17 and 12 SASS targets respectively, plus PTX), and the Windows AMD HIP build for PyTorch 2.10.0+rocm7.14.0 (gfx1201), which keeps the existing kernels.

With the new path disabled, all four CUDA wheels reproduce 1.0.23 bit-exactly on the 18 existing GGUF formats and passed the standard GGUF/attention, PTQ1/Prism and fused-path checks on an RTX 5090; each also passed WanGP's short-batch accuracy, graph-replay and selection tests. Other NVIDIA architectures are compiled but not hardware-tested. The HIP wheel was compiled and imported; no AMD GPU was available. The attached source archive and repository tag contain all maintained kernel sources and the refreshed WanGP integration overlay.

Wheel hashes are listed in SHA256SUMS.txt; full build and validation records are in llama.cpp/release/1.0.24/.

GGUF CUDA kernels 1.0.23

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 24 Sep 20:11

GGUF CUDA 1.0.23 adds native short-batch FP16/BF16 outputs and optional fused SiLU-and-multiply for packed GGUF projections. Existing paths remain available.

Four wheels cover Python 3.10/PyTorch 2.7.1/CUDA 12.8 and Python 3.11/PyTorch 2.10/CUDA 13 on Windows and Linux. Each native extension contains every architecture supported by its build toolkit (17 and 12 SASS targets respectively, plus PTX).

All four wheels passed existing-format bit-exact, numerical, attention, PTQ1/Prism, fused-path, and CUDA-graph checks on an RTX 5090. Other architectures are compiled but not hardware-tested. The attached source archive and repository tag contain all maintained kernel sources, generated SM120 binary provenance, and the WanGP integration overlay. The AMD HIP wheel remains at 1.0.22.

Wheel hashes are listed in SHA256SUMS.txt; full build and validation records are in llama.cpp/release/1.0.23/.

GGUF CUDA kernels 1.0.22

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 20 Sep 01:40

Adds Prism PTQ1_0 support to the existing packed GGUF MMVQ/MMQ kernels, signed Hadamard transforms and a fused Prism decode extension. Existing quantization APIs and dispatch are preserved.

Includes four CUDA wheels:

Platform Python PyTorch CUDA runtime / build toolkit
Windows x64 3.10 2.7.1 12.8 / 12.8
Linux x64 3.10 2.7.1 12.8 / 12.8
Windows x64 3.11 2.10.0 13.0 / 13.1
Linux x64 3.11 2.10.0 13.0 / 13.1

In the CUDA wheels, all three native extensions contain every GPU target supported by the corresponding toolkit, plus PTX for its highest target. CUDA 12.8 has 17 SASS targets; CUDA 13.1 has 12. Hardware validation used an RTX 5090 on Windows and Ubuntu 22.04 under WSL; other GPUs were compiled and inspected, not hardware-tested. Automatic Prism/GDN fusions and the architecture-specific Q8 async path retain their validated SM120 selection rules.

Each CUDA wheel passed 870 bit-identical existing-format comparisons against its matching 1.0.21 wheel, the standard linear/attention suite, 240 PTQ1 configurations, 81 Hadamard configurations and 360 fused Prism launch configurations. Real Q4 and Bonsai PTQ1 checkpoints passed repeated native-MTP generation with aligned caches on all four stacks. See llama.cpp/release/1.0.22/ in the source for detailed validation and binary/build records.

The source archive includes the updated WanGP GGUF/Prism, GDN, MTP, DSpark and DFlash2 integration, native and Triton sources, vendored dependencies, licenses and reproducible build scripts. Shared Triton kernels compile on use; the archive does not contain a universal Triton cache. The WanGP snapshot is a source overlay, not a standalone application distribution.

The new checkpoint assets are published in DeepBeepMeep/Wan2.1.

Install the wheel matching your existing environment with pip install --no-deps <wheel URL>. SHA256SUMS.txt covers all four wheels and the source archive.

Experimental AMD HIP wheel (PyTorch 2.10)

The Windows AMD gfx1201 wheel targets Python 3.11, PyTorch 2.10.0+rocm7.14.0 and ROCm 7.14.0. It replaces the earlier +hip asset, which incorrectly targeted PyTorch 2.12. The filename now identifies its target, and the package checks the exact PyTorch build before importing native DLLs.

Download the HIP wheel. Install it into the matching ROCm environment with pip install --no-deps. The HIP and CUDA wheels share the llamacpp_gguf_cuda package name; install the appropriate backend.

All three HIP extensions were rebuilt against PyTorch 2.10 and passed import checks. AMD GPU execution and full Qwen inference remain unverified because the build machine has no AMD GPU. The NVIDIA test results above describe the CUDA wheels.

GGUF CUDA Kernels 1.0.21

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 06 Sep 01:59

Latest packed GGUF CUDA kernels with precompiled RTX50xx (SM120) async attention, plus universal matrix and standard attention builds for the architectures supported by each CUDA toolkit.

Platform Python PyTorch CUDA runtime
Windows x86-64 3.11 2.10.0 13.0
Linux x86-64 3.11 2.10.0 13.0
Windows x86-64 3.10 2.7.1 12.8
Linux x86-64 3.10 2.7.1 12.8

Install the matching wheel with python -m pip install --no-deps <wheel-or-URL>.

Included

  • Architecture-aware packed MMVQ/MMQ dispatch and direct FP16/BF16 activation quantization.
  • Graph-safe activation padding, reusable workspaces and Q8/dense paged attention.
  • Four precompiled SM120 async programs per CUDA major, covering FP16/BF16 prefill and grouped decode/verification. Shapes and lengths remain runtime inputs. These use cp.async, not TMA.
  • The async programs need no runtime Triton compilation. WanGP vllm still uses shared Triton kernels elsewhere; the included WanGP source overlay selects the native SM120 module. Other GPU architectures retain the shared implementation, and legacy/cg retain native/PyTorch paths.
  • Complete native sources, vendored dependencies/licenses, AOT compiler inputs, scripts, tests, source overlay and validation evidence.

CUDA 13.1 builds contain SASS for 12 architectures, starting at SM75; CUDA 12.8 builds contain 17, including pre-SM75 targets. Both include PTX for the highest toolkit target. Availability on older hardware also depends on PyTorch. Linux runtime library paths are package-relative.

Validation

Each wheel passed 280 nonzero linear configurations (including real Qwen Q4/Q6 weight slices), embeddings, 12 standard attention cases, 3 padding cases and 20 compiled SM120 configurations with mutated CUDA-graph inputs. Another 24 compiled-path/backend-isolation checks passed. Repeated real Qwen3.8 Q4 inference passed on all four stacks; PyTorch 2.7 MTP was also checked with CUDA_LAUNCH_BLOCKING=1. Hardware testing was on RTX5090 only; other fatbin architectures were inspected.

The Windows 2.7 vllm test used an isolated Triton 3.3.1 installation; the existing environment also passed the CG path. The native async programs themselves were tested without importing Triton. The source overlay includes the corresponding PyTorch 2.7 startup/device fixes.

See the validation report and build instructions. Timing figures in the report are validation observations, not a controlled before/after speedup claim.

The source ZIP includes the exact AOT compiler input and an archive-provenance record. SHA256SUMS.txt covers the four wheels and source ZIP.

Source commit: 2808bfb073bd91e4fe3ef83712f600b8d642579b.

Nunchaku V1.2.0

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 15 Jan 11:35

official kernels https://github.com/nunchaku-tech/nunchaku without dependencides on diffusers, peft & transformers as they caused some comatibility issues
No that these kernels contain both int4 and fp4 support but fp4 will work only with RTX50xx / sm120+ GPUs

Light2xv NVFP4 Kernels

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 13 Jan 13:53

Light2xv NVFP4 Kernels compiled for Pytorch 2.9.1 Cuda 13.
These kernels will work only with RTX50xx / sm120+ GPUs
Source: https://github.com/ModelTC/LightX2V

GGUF Kernels

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 10 Mar 17:59

Precompiled GGUF CUDA kernels for WanGP.

Latest version: 1.0.11.

Windows and Linux wheels are available for Python 3.11 / PyTorch 2.10 and Python 3.10 / PyTorch 2.7.

Flash2

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 02 Feb 22:43

Flash 2 compiled kernels

Nunchaku Linux Kernels

Choose a tag to compare

@deepbeepmeep deepbeepmeep released this 02 Jan 10:25
ea98d85

official kernels https://github.com/nunchaku-tech/nunchaku without dependencides on diffusers, peft & transformers as they caused some comatibility issues