Skip to content

perf(bitnet): CUDA port of the BitLinear ternary matmul kernel - #322

Merged
inureyes merged 1 commit into
mainfrom
feat/bitnet-cuda-kernel
Jun 17, 2026
Merged

inureyes merged 1 commit into
mainfrom
feat/bitnet-cuda-kernel

Conversation

@inureyes

Copy link
Copy Markdown
Member

The BitLinear ternary matmul (BitNet b1.58) was a Metal-only mx.fast.metal_kernel, so on the CUDA backend both bitnet-b1.58-2b-4t and the 4bit variant aborted at warmup with [metal_kernel] No Metal back-end.

Port the kernel to mx.fast.cuda_kernel (the CUDA analogue of metal_kernel), the same pattern as #319 (the fused decode-MoE CUDA port):

  • One warp per (batch, out/4) row group; the simd_sum reduction over in_features becomes a __shfl_down_sync warp reduction.
  • The 2-bit-packed ternary unpacking (4 output rows per byte, value = bits − 1) is unchanged.
  • bitlinear_matmul selects the cuda_kernel port when metal::is_available() is false; no_cuda.cpp / no_metal.cpp stub the unused side, so both link on either backend.

Validation (GB10 / DGX Spark, CUDA 13.0)

Both bitnet variants now produce coherent output (previously an abort):

Hello! I'm just a computer program, so I don't have feelings, but I'm here and ready to help you. How can I assist...

Only src/lib/mlxcel-core/cpp/mlx_cxx_kernels.cpp changes.

The BitLinear ternary matmul (BitNet b1.58) was a Metal-only
mx.fast.metal_kernel, so on the CUDA backend it aborted at warmup with
"[metal_kernel] No Metal back-end" for both bitnet-b1.58-2b-4t and the
4bit variant.

Port it to mx.fast.cuda_kernel: one warp per (batch, out/4) row group,
the simd_sum reduction over in_features becomes __shfl_down_sync, and the
2-bit-packed ternary unpacking (4 output rows per byte) is unchanged.
bitlinear_matmul selects the cuda_kernel port when metal::is_available()
is false.

Verified on GB10 (DGX Spark, CUDA 13.0): both bitnet variants now produce
coherent output (previously an abort).
@inureyes inureyes added area:inference Generation, sampling, decoding (incl. speculative, DRY) priority:high High priority type:bug Bug fixes, error corrections, or issue resolutions labels Jun 17, 2026
@inureyes
inureyes merged commit f1547ac into main Jun 17, 2026
5 checks passed
@inureyes
inureyes deleted the feat/bitnet-cuda-kernel branch June 17, 2026 04:02
inureyes added a commit that referenced this pull request Jun 17, 2026
…large (#324)

bitnet re-benched on the fixed binary (passes, merged into CSV); deepseek-v3 reclassified as too-large/capacity (671B ~350GB, incomplete download). Counts 135/8/3/2.
inureyes added a commit that referenced this pull request Jun 17, 2026
The two bitnet rows were measured separately when their CUDA ternary
kernel landed (#322), so they carried a different hardware label
(NVIDIA_GB10_CUDA13.0_122GB vs NVIDIA_GB10_122GB) and an unprefixed
model_path (models/... vs ./models/.../) from the other 145 rows. Same
machine and run, just label drift. Normalized both columns so all 147
rows are uniform. Measured numbers are unchanged.
@inureyes inureyes added this to the 0.3 milestone Jun 21, 2026
@inureyes inureyes self-assigned this Aug 31, 2026
@inureyes inureyes added the status:done Completed label Sep 6, 2026
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

area:inference Generation, sampling, decoding (incl. speculative, DRY) priority:high High priority status:done Completed type:bug Bug fixes, error corrections, or issue resolutions

Projects

None yet

Development

Successfully merging this pull request may close these issues.

1 participant