Mirror the Python hanzo_research SDK verbs + semantics:
experiment(hypothesis, predict) -> log(text) -> conclude(verdict, because, value),
plus record / snapshot / report and query / totals. The Verdict enum
(Proven|Refuted|Inconclusive) rejects any other value at compile time; a
refutation is a first-class sealed result. A stated hypothesis with no verdict
seals Inconclusive, never silently a proof.
Zero-config provenance: the first entry method posts the run in-flight and stamps
git sha/branch/dirty, the commit narrative since the experiment's last recorded
run, the crate versions resolved from Cargo.lock, and the host. The caller
supplies none. Records serialize to the identical wire shape as the Python SDK,
proven by an equality test against the golden shape.
Per-org Bearer auth only (no client-forged identity headers); the api key is
redacted from Debug; query values are percent-encoded; artifacts carry a client
sha256 the server verifies.
Deps stay within the resolved lock (ureq 2.12, base64 0.22, sha2 0.10, serde);
hex and percent-encoding are hand-rolled. Blocking client for a bench tool.
Tests: shape identity vs Python golden, running-record empty-log, finish
defaults, attempt shape, content-addressing, key redaction, and an end-to-end
mock-server run of a kernel-perf refutation. Example: examples/kernel_perf.rs.
ROCm qmmq asymmetric min-bias precompute (hanzo-rocm-kernels)
------------------------------------------------------------
The int8-WMMA prefill GEMM (qmmq_core<WTYPE>) carries a min-bias for the
sub-block-min k-quants: out -= dmin_w * m_g * d_x * sum_k(q_x). The per-block
activation scale d_x was recomputed inside the per-output inner loop (one extra
multiply per output element per K-block). Fold it once per activation block at
quantize time: quantize_q8_1 now stores xs = d_x * sum(q_x) (f32, llama
block_q8_1 `s`) instead of the raw int sum, so the epilogue reads
acc -= xs * (dmin*m) -- the d_x hoisted out of the hot loop. xs changes
i32 -> f32 across the quantize kernels, qmmq_core, the DEFINE_QMMQ/MoE launchers
and the rocm_backend binding.
Bit-exact by construction: __half2float(dh) * (float)qsum is exactly the product
the old epilogue formed from the stored f16 scale, so every per-block bias term
is IEEE-identical. Verified on evo (gfx1151) with Qwen3-4B-Q4_K_M greedy decode:
output byte-identical pre/post (401 chars). rocprofv3 SQ_INSTS_VALU per dispatch,
648/648 dispatches: qmmq_q4k_f16 -6.6% (the asymmetric target); qmmq_q6k_f16
+0.01% (symmetric control -- the !SYMM branch elides, unchanged as required).
Symmetric weight types never read xs; decode matvec (rows==1) is a separate
kernel and is untouched.
Vulkan int8 per-token KV-quant paged attention (opt-in)
-------------------------------------------------------
Symmetric per-token int8 KV cache (KIVI / KVQuant): each cached per-head K/V
vector is quantized to signed int8 with one f32 scale (amax/127), codes packed
4-per-u32 with a co-located scale word, so a token vector is self-describing and
paged by slot. Two kernels: reshape_and_cache_q8 (amax reduce -> scale -> pack on
write) and paged_attn_q8 (dequantize-on-read code*token_scale, same online-softmax
as paged_attn). Opt-in via --pa-cache-type int8; default Auto is byte-unchanged.
Verified on evo (gfx1151, RADV): the U32 int8 cache path is active with no
fallback; int8 decode is argmax-identical to f16 at 128 greedy tokens and coherent
with exact ordered 12-item recall at 256 (information-faithful under the lossy KV).
The engine's hanzo-paged-attn::quant CPU oracle (6 tests) gates the packing.
Committed SPIR-V regenerated for glslc-less hosts.
Port the full FlashAttention-3 Hopper forward kernel set (arXiv:2407.08608):
head dims 64/128/256/512, fp16 + bf16, dense and GQA-packed, causal/local/
varlen, with the warp-specialized producer/consumer pipeline, TMA async copy
and wgmma of the reference hopper/ kernels.
- Kernels compile ONLY under the `cuda` feature, cross-compiled for sm_90a
(nvcc -gencode arch=compute_90a,code=sm_90a) on any CUDA host. Without the
feature the crate is a pure-Rust stub, so ROCm/Metal/Vulkan/CPU/non-Hopper
CUDA builds are unaffected. Un-excluded as a stub-by-default workspace member.
- cuda_compat.h: force-included shim restoring the unversioned
PFN_cuTensorMapEncodeTiled/Im2col typedefs CUTLASS 3.5 needs under CUDA 13;
touches no kernel. All 49 sources compile clean for sm_90a on nvcc 13.0.
- Numerical-equivalence harness: FA3 vs f32 CPU oracle within a scale-relative
bound across head dims/dtypes/causal/GQA/varlen, #[ignore]d for Hopper.
- publish.yml publishes the crate after hanzo-ml.
fp8 (e4m3) left as a documented hook. Perf UNVALIDATED (no Hopper silicon).
Add kernel_mul_mm_q{4,6}_K_bf16_half: stage the bf16 activation tile as half so the
K-quant prefill GEMM outer product runs at the GPU half-precision simdgroup rate. The
weight tile is already half and a bf16 activation widens to half exactly, so the operand
values match the f32-tile bf16 kernel -- only the matmul tile precision changes (half
products vs f32). Two new template params (B, simdgroup_B8x8) default to float, so every
existing mul_mm instantiation stays byte-identical.
Opt-in via HANZO_METAL_MM_F16 (ml fwd gates it); the default remains the bit-exact
f32-tile bf16 kernel. Bit-exact to that path on synthetic data (metal_mm_q{4,6}k_bf16_half_bounded:
max_abs=0) and behaviourally identical greedy output on qwen2.5-3b + qwen3-4b.
Measured dbc M4 Max, Qwen3-4B Q4_K_M pp512: fp16 1474 vs f32-tile 1413 t/s (+4.3% pure
wall-clock; +5.5% engine-internal prefill). Does not overtake llama.cpp -fa 1 at 1642 t/s
(0.90-0.935x) -- the matmul lever converts but the residual prefill gap is elsewhere.
The f16 coopmat matrix-core prefill GEMM (mul_mm_q{4,6}k_coopmat) was gated on
m % 16 == 0, so any prompt whose token count is not a multiple of 16 -- nearly
every real prompt -- fell through to the ~3.4x slower int8 dp4a tile. Pad the
output-row allocation up to a 16-multiple (the trailing <16 scratch rows are fed
by stage()'s zero-padded activation and never read back, count stays m*nout) so
the 16x16 coopmat store runs for any m. This is llama mul_mm.comp's partial-tile
store expressed as an allocation instead of a divergent store path. coopmat is
now the default for all dense (woff==0, m>1) Q4_K/Q6_K prefill.
Measured on gfx1151 (RADV), 506-token prefill, per-op GPU timestamps:
Q4_K GEMM 498.9ms (dp4a) -> 142.4ms (coopmat) 3.5x
Q6_K GEMM 106.0ms (dp4a) -> 35.9ms (coopmat) 3.0x
prefill block total 687ms -> 253ms 2.7x
Bit-exact vs the column oracle at ragged m (vulkan_q4k_coopmat_matches_default
extended to m in {2,17,31,33,127,129}; mul_mm_q6k_coopmat_and_dp4a_match_column
gates both the coopmat default and the dp4a A/B fallback at m in {37,40,100}).
Coherence byte-identical to the dp4a path on qwen2 (zen-eco-4b).
hanzo-ml 0.11.83 + hanzo-nn 0.11.39 + hanzo-rocm-kernels 0.11.35.
ROCm single-query decode had no fused attention: it fell to eager repeat_kv
(Tensor::cat -> copy2d) + QKᵀ bmm + f32-softmax + P@V bmm + a per-layer
.contiguous() copy of the WHOLE active KV cache -- ~10 dispatches and the 653
copyBufferRectAligned/token that dominated ROCm decode glue.
flash_decode_{f16,bf16} (flash.hip): one block per (batch, q_head), 128 threads,
GQA implicit (kv_head = q_head/(Hq/Hkv), no repeat_kv), KV cache read IN PLACE via
(batch,head,seq) strides (head dim contiguous, no cache copy), tiled online softmax
in f32. RocmStorage::flash_attn_decode launches it from strided k/v layouts;
hanzo_nn::attention::rocm_flash_attn_decode is the Tensor wrapper; the engine wires
it as rocm_decode_attn (mirrors vulkan_decode_attn's gating).
Correct vs an f64 eager oracle across GQA ratios, tile-boundary L, and the strided
cache layout (test flash_decode_matches_eager, rel 2.4e-4..3.9e-4). copyBufferRect
33950->5760/run (-83%); decode QKᵀ/softmax bmms 3708->216 (prefill only).
Extend the fused add_rmsnorm op + reduce.hip kernel to F16, the ROCm decode
residual dtype. QRmsNorm::forward_of_sum previously fell back to a separate
badd_f16 + rmsnorm_f16 (barrier between) for F16 since the fused op was f32-only;
the fused kernel merges them into one dispatch, one HBM round-trip.
add_rmsnorm<T> now stores the sum rounded to T and reduces over the reloaded
rounded value: bit-identical to a separate add (rounds to T) then rmsnorm
(squares the stored T), the CUDA rms_norm_of_sum contract. F32 path is bit-
unchanged (round exact); F16 matches the badd_f16 + rmsnorm_f16 fallback bit-for-
bit (test add_rmsnorm_f16_bit_identical_to_separate, GPU byte-equality, 5 shapes).
Add kernel_mul_mm_q{4,6}_K_bf16: src1 is read as bfloat and widened to f32 on
stage (so the simdgroup math is identical to the _f32 mm) and the f32 accumulator
is narrowed to bfloat on store via an overloaded epilogue, keeping the _f32 kernel
byte-identical. This removes the bf16<->f32 activation round-trip GgufMatMul wraps
around the projection and halves the src1 read + dst write traffic. fwd routes a
bf16 activation against a Q4_K/Q6_K weight to it -- the multi-token twin of the
bf16 decode matvec. Bit-identical to the f32-then-round-bf16 path
(metal_mm_q{4,6}k_bf16_matches_f32: max_abs=0 over full and partial tiles).
The Metal prefill attention materializes several `contiguous(x.transpose(1, 2))`
copies per layer (q after RoPE, the attention output before o_proj, and the K/V
cache append). Their innermost dim (head_dim) is contiguous but the per-block
source stride is non-uniform across the outer axis, so `copy_strided_src` could
not express them as a single copy2d and fell back to the generic strided kernel,
which recomputes a full multi-dimensional index divide for every element.
Add copy3d: the natural three-dimensional extension of copy2d. Each thread
derives its source offset from the grid position directly (three fused multiply-
adds, no division) and the contiguous innermost dim maps to grid x so adjacent
lanes touch contiguous memory. `copy_strided_src` reduces a strided layout to at
most three non-unit dims with a contiguous inner dim and routes it here; anything
else keeps the generic path. Bit-exact by construction (a contiguity copy is an
identity on values); `metal_copy3d_transpose_bit_exact` asserts max_abs == 0.0
against the CPU oracle over the live prefill shapes, a 3D transpose, and the
four-non-unit-dim fallback.
Measured on dbc (M4 Max, Qwen3-4B Q4_K_M, pp512): the slow strided copy per
transpose drops 127.7us -> 19.7us (6.5x); prefill 1317.3 -> 1435.5 t/s (+9.0%,
TTFT 388.7 -> 356.7ms), 0.836x -> 0.911x vs llama.cpp Metal. Decode
non-regressed at 130.2 t/s. The mul_mm GEMM is untouched (already at ggml
parity); the gain is pure layout-glue elimination. Coherent ("Paris").
The quantized decode matvec required an f32 activation, so GgufMatMul wrapped every
projection in a bf16->f32->bf16 round-trip. Measured (METAL_PROFILE_OPS) that pair was
~12-15% of decode GPU time and, decode being a serialized dependency chain, on the
critical path. Give kernel_mul_mv_q4_K/q6_K a bf16 activation path: templatize the impl
on the src1/dst type (T=float is bit-identical, keeps the f32 host kernel + MoE mmv_fn),
add bfloat host kernels, and thread a src1_bf16 flag through call_quantized_matmul_mv_t so
a bf16 activation reads/writes bf16 with f32 accumulation. fwd_mv allocates a bf16 dst and
selects the bf16 kernel when the activation is bf16 and the weight is Q4_K/Q6_K.
Bit-exact: bf16 matvec (f32 accumulate, bf16 store) == f32 matvec on the same bf16-rounded
activation rounded to bf16 -- exactly what GgufMatMul.forward_raw computes. New tests
metal_matvec_q4k/q6k_bf16_matches_f32 assert max_abs = 0. f32 matvec tests unchanged.
Qwen3-4B Q4_K_M, M4 Max: decode 107.8 -> 131.2 t/s; coherent. Engine wiring in a
companion hanzo-quant commit (GgufMatMul::forward override).
The MSL kernel_mul_mm_id_iq*_f32 templates were already present in quantized.metal
(verbatim ggml mul_mat_id port); the gap was 100% Rust dispatch -- there was no
call_mul_mm_id at all, so Metal MoE prefill (t>1) fell to the generic per-expert loop
(and glm-dsa IQ1_S prefill bailed 'Failed to create pipeline').
- metal-kernels: new call_mul_mm_id (all wired types incl the 9 i-quants), mirroring
call_mul_mv_id's arg block + [t,topk,n] output. Classic ggml grid
(ceil(t/32), ceil(n/64), n_experts) x 128 threads; nb01 = expert_bytes/n (the mm impl
indexes weight rows explicitly); threadgroup mem = 8192 + t*sizeof(ushort2) rowid map.
- hanzo-ml metal.rs: indexed_moe_forward dispatches mul_mm_id for t>1 (expert-grouped
GEMM), keeps mul_mv_id for t==1 (decode). Same low buffer churn on both paths.
- mod.rs: Metal MoE arm no longer t==1-only -> prefill computes on the Metal GPU.
- Gate: metal_iquant_tests::metal_moe_prefill_iquant_matches_cpu (IQ1_S + IQ2_XXS + IQ3_S,
multi expert/token/tile shapes, gate/up + down input) vs CPU f64 ref.
metal-kernels 0.11.35->0.11.36, ml 0.11.47->0.11.48.
The MSL kernels (get_rows/mul_mv/mul_mm/mul_mv_id/mul_mm_id + dequantize_iq* +
codebook grids) were already present in quantized.metal (verbatim ggml-metal
port); the gap was 100% Rust dispatch. from_data bailed 'IQ1_S is not supported
on the Metal backend', blocking GLM-5.2 (UD-IQ1_S) from loading on Apple Metal.
Wire the whole i-quant family (IQ1_S, IQ1_M, IQ2_XXS, IQ2_XS, IQ2_S, IQ3_XXS,
IQ3_S, IQ4_NL, IQ4_XS), no new MSL:
- hanzo_metal_kernels::GgmlDType +9 variants (names mirror ml's GgmlDType 1:1)
- the 3 dispatch fns (mv_t decode / mm_t prefill / mul_mv_id MoE): per-type
(nth0,nth1,align) + kernel-name arms; matvec binds threadgroup grid scratch
- mod.rs Metal load arm; metal.rs dequantize() readback + TryFrom mapping
IQ4_NL/IQ4_XS were half-wired (load+dequant, no matmul); now complete.
Gate (dbc M4 Max, tests/metal_iquant_tests.rs, 14/14): dequant round-trip
Metal-vs-CPU bit-identical (max_abs=0.0); native matvec vs CPU f64 ref max_rel
~1e-7 (all 9 types); prefill mm ~2e-4. Integration: TinyLlama-1.1B-IQ1_S loads
on Metal + generates coherent English @ 264 T/s (no NaN/garbage).
metal-kernels 0.11.34->0.11.35, ml 0.11.46->0.11.47.
- declare quantized_lfm2 in models/mod.rs (file existed, module never wired; broke the quantized-lfm2 example on cargo test)
- trainer.rs metal arm: anyhow::anyhow! instead of String.into() (E0277 on --features metal)
- clippy -D warnings sweep: .clamp() in dsv4_qat, FRAC_1_SQRT_2 in hanzo-3d test, doc-list indents, rename misnamed run_unary_benchmark, cfg-gated allows for rocm-only field/dim reads, drop redundant trim
- workspace clippy allows: needless_range_loop + too_many_arguments (quant-kernel idioms)
- rust-ci: install cc on the clippy runner (image ships without a C toolchain)
First all-Rust training loop in the stack: trains a Qwen3 DSpark speculative-draft
model on real cached target hidden states (DeepSpec v2 cache) with hanzo-ml autograd
+ hanzo-nn AdamW, and saves a checkpoint in the exact layout the engine's
qwen3_dspark.rs loader expects (verified by a key+shape load-check).
- cache.rs: DeepSpec v2 reader (idx records, bf16->f32), mini300 cache.
- model.rs: trainable DSpark draft (fc/hidden_norm/N Qwen3 layers/norm/markov head),
frozen embed (streamed on demand) + f16 lm_head. Frozen-head CE via a surrogate
backward so the [hidden,vocab] head grad is never formed (memory+compute win;
exact-gradient unit test). Output-norm small-init well-conditions the start (~ln V).
- main.rs: CLI training loop, checkpoint save + config.json + load-check gate,
MemAvailable safety guard (never starves a co-resident job), --freeze-markov and
seq-cap knobs for memory-constrained hosts.
- 5 unit tests: cache record parse, markov bias, attention mask, frozen-head grad identity.
CPU/f32 MVP. Full 5-layer + tv/confidence losses + GPU are follow-ons.
The DSL crate was published to crates.io with repository=github.com/hanzoai/ml but its source
lived only in a standalone local repo -- the source link resolved to nothing. Home it in the ml
tree as excluded crates (matching hanzo-kernels/hanzo-metal-kernels/hanzo-rocm-kernels), so the
published crate's source is where it claims to be. README + LICENSE + runnable example included.
Extend the fused F32 hd512 flash-decode kernel to trailing query blocks of 1..=8 rows (grid.z plane per row; per-row causal end = kv_len-q_len+row+1 and per-row ds4 window clamp). S=1 byte-identical to 0.11.35/0.11.39. Rust API unchanged (q_len inferred from q shape [n_head, q_len, 512]). Tests: 10-case matrix (q_len x kv_len straddling the window x GQA x sink) vs CPU reference, max|delta| ~1e-7. Consumer: hanzo-engine V4 MTP/EAGLE verify forwards (previously fell to the eager 3-pass, making speculation net-negative).
0.11.33 called .device_ptr_mut() (the varlen dst/softmax_lse writes for the DS4 hdim512
+ all fwd kernels) but only imported DevicePtr, so it never compiled against cudarc 0.19
-- device_ptr_mut lives on the separate DevicePtrMut trait (cudarc core.rs:1229). Shipped
broken because publish.yml uses --no-verify (no CUDA compile gate on CI). One-line import
fix unblocks EVERY CUDA flash-attn build (was: no method named device_ptr_mut on CudaSlice).
0.11.36 on crates.io predates the Vulkan decode cleanup (805aa517) + docs; bump to
0.11.37 so the shipped decode path (column dp4a optimum, no dead-end env knobs) is in
a published release. Publishes via the now-fixed CI (ubuntu-latest + set +e skip).
DeepSeek-V4's hyper-connection combine matrix is tiny ([.,.,nh,nh], nh=4) but the
20-iteration row/col Sinkhorn alternation, as elementwise tensor ops, was ~120 tiny
kernel launches per call — a swarm of 16-thread kernels that dominated V4 decode
(nsys: bdiv/fast_sum/affine ~20% of GPU time, ~10.7k tiny kernels/token).
- hanzo-kernels/reduce.cu: hc_sinkhorn_f32 — one block per matrix, softmax + full
iteration in shared memory, ONE launch. Semantics match the reference exactly
(softmax(src)+eps; col-norm denom=colsum+eps; (iters-1)×(row,col)).
- hanzo-ml/hc_sinkhorn.rs: HcSinkhorn CustomOp1 (cpu_fwd reference + cuda_fwd) +
Tensor::hc_sinkhorn(iters, eps). f32 (the tensor-op path ran bf16 — low precision
for a 4×4 normalization; f32 is closer to the ds4.c reference).
Measured on GB10 (81GB IQ2 V4-Flash): decode 7.3 -> 9.8 t/s. Output coherent + correct
('The capital of France is Paris'). Bumps hanzo-kernels 0.11.34, hanzo-ml 0.11.36.
vulkan_prefill_gemm_max_rows returns usize::MAX for types with a native matmul_q*_gpu
(Q4_0/Q8_0/Q4K/Q5K/Q6K) so prefill always uses the quantized GEMM and never the f32-dequant
fallback. The dequant path materialized a fresh ~100-235MB f32 BO per weight (upload_f32 ->
raw_buffer -> amdgpu_bo_alloc); under the deferred single batch none free until end-of-forward
flush, so a dense prefill re-expanded the whole model to ~32GB of fresh BOs and GEM_CREATE
blocked forever on the 64GB-carveout/31GB-GTT Strix Halo UMA (gdb-confirmed kernel ioctl block).
Decode (rows==1, native matvec) and ROCm (int8-WMMA qmmq prefill) never hit it. release 0.11.31
Composes the outstanding CUDA work onto main (rebased onto swarm 0.11.29 dsv4_qat):
- moe_route CUDA fast path -- completes the backend abstraction (ROCm+CUDA fast paths + ml-op
fallback, symmetric); bit-exact (cuda_moe_route_numeric nbad_id=0), flat GB10 decode (re-bench
on 256-expert DeepSeek-V4).
- DECOMPLECT: drop dead KSIGNS_IQ2XS_D CUDA decode table (gen_iquant_grids.py DECODE_SKIP;
unpack_ksigns reconstructs it in-kernel; Rust ksigns stays for the CPU oracle).
Wave-5 completeness batch (all bit-exact on GB10, 13/13):
- IQ1_S int8-WMMA qmmq PREFILL: filled the real iq1s_grid_gpu[2048] (from llama.cpp,
GPU-packed; was a zeroed [512] stub) + DS4 ds-layout (delta bias rides the activation sum);
IQ1_S/MoE prefill now native. max_rel 5.6e-4 (f16 ds-rounding floor).
- TQ1_0/TQ2_0 native dp4a DECODE (+ MoE-decode): ternary {-1,0,1} -> signed int8 -> dp4a,
no bias-sum. TQ2_0 clean 2-bit; TQ1_0 base-3 fits-but-unpack-bound (honest). ~400-2700x
over the (CPU-round-trip) dequant fallback. llama.cpp has NO CUDA ternary path -- pioneering.
- IQ4_XS now in the fused-MoE numeric gate (red's [LOW]): run_moe is dtype-parametrized,
cuda_moe covers IQ2_XXS + IQ4_XS (the dominant MoE-quant) -- a MoE-wiring regression can't slip.
IQ4_XS red verdict was SHIP (decode bit-exact vs ggml/ROCm/CPU). Occupancy tune + CUDA-graph
levers closed as evidence-backed non-levers (47 GB/s = GB10 VALU wall; CUDA graphs already on
for Qwen3 -- the gap is kernel VALU + op-fusion, not dispatch).
DeepSeek-V4 i-quant MoE GGUFs carry I32 `ffn_gate_tid2eid` expert-routing maps (GGML_TYPE_I32
= 26) that the loader rejected ("unknown dtype for tensor 26 / Critical failure loading model
part 0"), blocking the whole model. These are auxiliary index tensors, NOT float/quant weights
-- the engine computes routing itself and loads weights by name, never requesting them (grep
confirms no `tid2eid` reference in the engine model code). So the GGUF header reader now skips
the integer GGML types (24..=27) rather than hard-failing: their header fields are still consumed
(reader stays aligned), they're just not recorded as loadable tensors. Minimal + correct -- vs a
full GgmlDType::I32 (which would need ~10 exhaustive-match arms + a QuantizedType impl for a tensor
that's never used). New gguf_i32_skip_test (set HANZO_TEST_GGUF) reads the header + asserts the
integer tensors are skipped + the weights remain. ml 0.11.24 (pure-Rust loader; kernels unchanged).
IQ4_XS (the dominant tensor type in Unsloth UD-IQ2_M MoE quants, previously dequant-capped
at 0.66 T/s) now has native dp4a decode + int8-WMMA qmmq prefill. 11/11 GB10 tests; decode
115-148x + prefill 171x over the dequant fallback; bit-exact prefill.
Reconciles the parallel-swarm 0.11.20 (unified-memory) with the published i-quant
0.11.20 (decode) + 0.11.21 (qmmq prefill). Clean auto-merge (unified-memory touches the
allocation path, i-quant the matmul dispatch -- disjoint regions of cuda.rs); only the
hanzo-ml version collided. Bumped to 0.11.22 = first version carrying BOTH feature sets.
# Conflicts:
# hanzo-ml/Cargo.toml
Wave 3: native int8-WMMA qmmq prefill for the 5 IQ2/IQ3 codebooks (the prefill twin of
the dp4a decode shipped in 0.11.20) -- dense + expert-grouped MoE. Bit-exact (max_rel
~1.6e-7) + 292-335x over the dequant fallback; the int8 tensor cores process 128 prefill
tokens in LESS time than 1 decode token (full weight-load amortization). IQ1_S/M stay
dequant prefill (distinct GPU grid / no kernel) but keep native decode. Also folds in the
CPU i-quant MoE matmul block-count fix (was a panic on the CPU per-expert path).
Native dp4a i-quant decode for IQ2_XXS/XS/S, IQ3_XXS/S, IQ1_S/M on CUDA -- dense
(69-425x over the dequant fallback) + fused MoE (bank-resident, on-device gather,
replacing the per-token whole-bank DtoH). 9/9 numeric gates on GB10, red-reviewed.
Native Vulkan decode shaders for IQ2_XXS/IQ2_XS/IQ2_S/IQ3_XXS/IQ3_S/IQ1_S/IQ1_M
(were dequant->f32). Each mirrors the ROCm/CPU qdec codebook decode bit-faithfully;
non-u32-aligned blocks repack to a padded stride on upload. Vulkan native decode
coverage 10 -> 17 types. Full suite 21/21 bit-exact on RADV gfx1151 (max_rel ~1e-6).
Fixes silent MoE routing corruption on CUDA for all .topk()-based models (Qwen3.6-35B-A3B,
deepseek3, glm4_moe, gpt_oss, granite): the fused indexed-MoE now derives topk from
elem_count/batch (rank-agnostic, matches the flat-buffer kernel) + forces ids contiguous.
Paired with the engine cuda_topk all-rows fix. Red-reviewed SHIP; 35B-A3B coherent on GB10,
ROCm/dense unaffected. Only hanzo-ml changed; rocm-kernels stays 0.11.16.
qdp4a<DW_IQ2_XXS/XS/S, IQ3_XXS/S, IQ1_S/M>: decode the grid coord to int8 and sudot4
against the q8_1 activation (llama's vec_dot_iq*_q8_1 on the unified dp4a core),
replacing the scalar codebook float-MAC. IQ3_XXS decode 30->68 T/s (== llama-HIP 67.6),
prefill 4x; bit-exact (dp4a-vs-scalar + vs-CPU nbad=0). Sign via dp4a(g,u)-2*dp4a(g&mask,u)
(no __vsub4 on ROCm 7.13). rocm-kernels + ml 0.11.11->0.11.12.
16 new native ROCm decode families (Q2_K Q3_K Q5_K Q4_1 Q5_0 Q5_1 Q8_1 IQ2_XXS/XS/S
IQ3_XXS/S IQ1_S/M TQ1_0 IQ4_NL) on the unified core: one qdec + traits row each, ONE
qmmq_capable predicate (11 prefill-capable / 11 decode-only). Every one of the 22
RocmQuantType variants decodes bit-exact vs CPU to_float (nbad=0); zero regression on
the 6 pre-existing. rocm-kernels + ml bumped 0.11.10->0.11.11.
Routes MoE prefill (rows>1) through a fused expert-grouped iu8 WMMA GEMM (qmmq_core<WTYPE,MOE>) instead of the per-slot dp4a matvec: tokens grouped by expert, each expert weight staged once and amortized over its token tile (llama mul_mat_id strategy). Decode keeps the capture-clean dp4a matvec. Qwen3-30B-A3B-Q4_K_M prefill 363->936 T/s (2.58x, 0.34x->0.89x of llama-HIP), decode 41->50 (+22%). Bit-exact nbad=0 (Q4_K/Q6_K/Q8_0 MoE oracle); dense path codegens byte-identical (if constexpr MOE elides). HANZO_MOE_QMMQ_FALLBACK=1 to force the matvec path.
Hanzo Dev
z