feat: port TQ3_0 KV cache from llama-turboquant by carlosfundora · Pull Request #2 · PrismML-Eng/llama.cpp

carlosfundora · 2026-04-01T18:50:56Z

TurboQuant 3-bit (3.5 bpw) KV cache compression combined with PrismML's Q1_0 GPU inference. Ported the TQ3_0 implementation from llama-turboquant, tested on ROCm gfx1030.

…CUDA) Adds two 1-bit quantization types: - Q1_0: block size 32, ~1.5 bpw - Q1_0_g128: block size 128, ~1.125 bpw Backend support: CPU (x86 SSE/AVX + ARM NEON), Metal, CUDA. Kernel implementations follow Q4_0 as boilerplate, adapted for 1-bit sign-based dequantization. CUDA MMQ kernels included but disabled (cuBLAS fallback used for prompt processing) pending accuracy debugging. Made-with: Cursor

[cuda] Fix mmq/mma path

TurboQuant 3-bit (3.5 bpw) KV cache compression: - Per-block WHT rotation with 4-centroid MSE codebook - QJL residual signs for error correction - GPU kernels: vec_dot, MMVQ, convert, set-rows, cpy - CPU: quantize/dequantize with WHT butterfly transform - Flash attention auto-disabled for TQ3_0 K cache Combined with PrismML's Q1_0 GPU inference, this enables 1-bit weights + 3-bit KV cache on a single build.

khosravipasha · 2026-04-02T22:26:28Z

Thanks this is pretty cool. How does it work? It is good?

Our main focus right now is getting our changes in llama.cpp so might not have time to look into details yet, but love to see if speed/quality output if we have tried it.
What the vram usage with long context after this change?

carlosfundora · 2026-04-03T06:35:21Z

It works great. I have SGLang nearly wired up for 1-bit support and TurboQuant as well.

carlosfundora · 2026-04-03T06:37:44Z

VRAM usage was reduced by roughly 35%.

khosravipasha · 2026-04-03T10:46:56Z

Oh how does it work with SGlang for 1-bit, was it easy to add support there?

carlosfundora · 2026-04-04T03:37:35Z

So so, patience and preparation was key. I also crafted a few agents to run methodical research and debugging on a very detailed and slow scale during smoke tests. It ran well for chat in SGLang, I bechmarked it and moved on to implement P-EAGLE though, so I've been training heads for the models all day. I haven't yet tried it on any coding tasks but I'm excited to see how they do.

If you guys would be kind enough to release a 1/2 B or .3-.6 range 1-bit quant that would be amazing and make it much easier for me to rapidly work on creating PRs for advanced speculative decoding architectures. 🧠🤌

khosravipasha · 2026-04-04T18:45:32Z

@carlosfundora Sounds exciting, yeah good ideas, lets chat more on the discord-server next week (I think you were there right?)

rosmur · 2026-04-05T06:49:23Z

excellent! Will this work on Apple Silicon? If yes, I can report back with memory footprint improvement

carlosfundora · 2026-04-06T04:23:56Z

@carlosfundora Sounds exciting, yeah good ideas, lets chat more on the discord-server next week (I think you were there right?)

Yes, I commented on there today. looking forward to it.

Copilot

Pull request overview

This PR introduces a new GGML quantization type (GGML_TYPE_TQ3_0) intended for TurboQuant-style KV-cache compression, wiring it through GGML core traits, CUDA/HIP paths, and CLI/tooling so it can be selected as a KV cache type (tested on ROCm gfx1030 per description).

Changes:

Add GGML_TYPE_TQ3_0 type definition/traits and (de)quantization hooks in GGML core + CPU integration.
Add CUDA kernels/support for writing (SET_ROWS) and using (MMVQ vecdot) TQ3_0 KV cache blocks.
Expose tq3_0 via CLI and bench tooling; disable flash-attention when type_k == TQ3_0.

Reviewed changes

Copilot reviewed 20 out of 20 changed files in this pull request and generated 10 comments.

Show a summary per file

File	Description
tools/llama-bench/llama-bench.cpp	Adds `tq3_0` string → type mapping for benchmarking.
src/llama-context.cpp	Forces flash-attention off when using TQ3_0 K-cache.
ggml/src/ggml.c	Registers TQ3_0 type traits and enables chunk quantization dispatch.
ggml/src/ggml-quants.h	Adds TQ3_0 quantize/dequantize API declarations.
ggml/src/ggml-quants.c	Implements TQ3_0 reference quantize + dequantize + quantize wrapper and row validation.
ggml/src/ggml-cuda/vecdotq.cuh	Adds fused TQ3_0×Q8_1 vecdot for MMVQ.
ggml/src/ggml-cuda/set-rows.cu	Enables `SET_ROWS` into TQ3_0 buffers (KV updates).
ggml/src/ggml-cuda/mmvq.cu	Wires TQ3_0 into MMVQ type switches.
ggml/src/ggml-cuda/ggml-cuda.cu	Marks additional ops/types as CUDA-supported (incl. TQ3_0).
ggml/src/ggml-cuda/cpy-utils.cuh	Adds device quantization helper for TQ3_0 blocks.
ggml/src/ggml-cuda/convert.cu	Adds CUDA dequantization kernel for TQ3_0 → fp16/fp32.
ggml/src/ggml-cuda/common.cuh	Adds CUDA type-traits for TQ3_0 (qk/qr/qi).
ggml/src/ggml-cpu/quants.h	Declares CPU quantize entrypoint for TQ3_0.
ggml/src/ggml-cpu/quants.c	Implements CPU quantize wrapper calling reference quantizer.
ggml/src/ggml-cpu/ops.cpp	Allows TQ3_0 through quantized op switch cases.
ggml/src/ggml-cpu/ggml-cpu.cpp	Tightens CPU op support checks for `MUL_MAT` and `FLASH_ATTN_EXT`.
ggml/src/ggml-cpu/ggml-cpu.c	Registers TQ3_0 CPU type-traits (from_float).
ggml/src/ggml-common.h	Defines `block_tq3_0` layout and constants.
ggml/include/ggml.h	Adds `GGML_TYPE_TQ3_0` to the public enum/API.
common/arg.cpp	Exposes `tq3_0` as an allowed KV cache type via CLI.

💡 Add Copilot custom instructions for smarter, more guided reviews. Learn how to get started.

khosravipasha · 2026-04-06T19:08:43Z

Good new our first CPU PR just got merged int llama.cpp master branch now, if you are still working on this please rebase with PrismML's master (just pulled the main llama.cpp)

Changes: Q1_0_g128 naming is gone now, the original Q1_0 with group size 32 was deleted and Q1_0_g128 was renamed to Q1_0 now by default has group size 128.

https://github.com/PrismML-Eng/llama.cpp/tree/master

This one only has generic cpu (slow), and ARM NEON path, planning to gather the best x86 kernels from here and to send a PR there (and tag all the contributers).

khosravipasha · 2026-04-10T06:59:44Z

Now few more of our backends are merged in main llama.cpp.
I have more free time to help with testing now to see how it affects evals/etc.
Also remember you mentioned a sglang option.
Might want to rebase this from the master branch (generic cpu, metal, vulkan backend are already in the master branch, cuda and x86 pending)

khosravipasha · 2026-04-14T00:29:06Z

@carlosfundora
Just revamped the prism branch by grabbing the llama.cpp master and applying our two pending PRs (x86 and cuda backned)

Do you still want to merge this? There is some outdates stuff too, for example we removed Q1_0_g128 since they did not like the name, and removed the old Q1_0 (group size 32), now Q1_0 (group size 128) is the only new quantization that is added (and many backends aleady merged into llama.cpp)>

Which part does the Turboquant related things? I am curous to try, but need to rebase this and fix some confilicts.

…ng#29 Codex post-commit review found: 1. TURBO_D was QK_TURBO3 (now 32) — broke turbo4 C array sizes 2. SET_ROWS kernel turbo3-specific but instantiated for turbo4 3. Tail block drop for non-128 head dims Fixed PrismML-Eng#3 (TURBO_D). Mintplex-Labs#1 and Mintplex-Labs#2 don't affect turbo3+dk128 path. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Complete experiment log: Mintplex-Labs#1 4-mag LUT: 15.1 at 8K (BEST, +38%) Mintplex-Labs#2 Batched extract: 13.7 (+25%) PrismML-Eng#3 Inline FA block: 13.5 (I-cache pressure) PrismML-Eng#4 Deferred norm: 12.9 (loses ILP) PrismML-Eng#5 2-pair half2: 12.0 (ternary overhead) PrismML-Eng#6 Select chain: 11.9 (branches kill) PrismML-Eng#7 Bit-arithmetic: 11.6 (ALU too heavy) PrismML-Eng#8 FMA branchless: 11.4 (ALU still too heavy) PrismML-Eng#9 Named-reg ternary: 10.3 (branches worst) PrismML-Eng#10 Main (8-LUT): 10.95 (baseline) PrismML-Eng#11 Non-vec FA: 10.2 (wrong kernel) Ceiling: 24.5 (no dequant) Apple8 hardware truth: 1 divergent constant read < 7 ALU ops (even with fma) Branches cost MORE than divergent constant reads Array indexing ALWAYS spills on Metal 4 constant addresses is the sweet spot The 4-mag LUT is the dequant-level ceiling on Apple Silicon. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai

…ng#29 Codex post-commit review found: 1. TURBO_D was QK_TURBO3 (now 32) — broke turbo4 C array sizes 2. SET_ROWS kernel turbo3-specific but instantiated for turbo4 3. Tail block drop for non-128 head dims Fixed PrismML-Eng#3 (TURBO_D). Mintplex-Labs#1 and Mintplex-Labs#2 don't affect turbo3+dk128 path. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Complete experiment log: Mintplex-Labs#1 4-mag LUT: 15.1 at 8K (BEST, +38%) Mintplex-Labs#2 Batched extract: 13.7 (+25%) PrismML-Eng#3 Inline FA block: 13.5 (I-cache pressure) PrismML-Eng#4 Deferred norm: 12.9 (loses ILP) PrismML-Eng#5 2-pair half2: 12.0 (ternary overhead) PrismML-Eng#6 Select chain: 11.9 (branches kill) PrismML-Eng#7 Bit-arithmetic: 11.6 (ALU too heavy) PrismML-Eng#8 FMA branchless: 11.4 (ALU still too heavy) PrismML-Eng#9 Named-reg ternary: 10.3 (branches worst) PrismML-Eng#10 Main (8-LUT): 10.95 (baseline) PrismML-Eng#11 Non-vec FA: 10.2 (wrong kernel) Ceiling: 24.5 (no dequant) Apple8 hardware truth: 1 divergent constant read < 7 ALU ops (even with fma) Branches cost MORE than divergent constant reads Array indexing ALWAYS spills on Metal 4 constant addresses is the sweet spot The 4-mag LUT is the dequant-level ceiling on Apple Silicon. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai

…ng#29 Codex post-commit review found: 1. TURBO_D was QK_TURBO3 (now 32) — broke turbo4 C array sizes 2. SET_ROWS kernel turbo3-specific but instantiated for turbo4 3. Tail block drop for non-128 head dims Fixed PrismML-Eng#3 (TURBO_D). Mintplex-Labs#1 and Mintplex-Labs#2 don't affect turbo3+dk128 path. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>

Complete experiment log: Mintplex-Labs#1 4-mag LUT: 15.1 at 8K (BEST, +38%) Mintplex-Labs#2 Batched extract: 13.7 (+25%) PrismML-Eng#3 Inline FA block: 13.5 (I-cache pressure) PrismML-Eng#4 Deferred norm: 12.9 (loses ILP) PrismML-Eng#5 2-pair half2: 12.0 (ternary overhead) PrismML-Eng#6 Select chain: 11.9 (branches kill) PrismML-Eng#7 Bit-arithmetic: 11.6 (ALU too heavy) PrismML-Eng#8 FMA branchless: 11.4 (ALU still too heavy) PrismML-Eng#9 Named-reg ternary: 10.3 (branches worst) PrismML-Eng#10 Main (8-LUT): 10.95 (baseline) PrismML-Eng#11 Non-vec FA: 10.2 (wrong kernel) Ceiling: 24.5 (no dequant) Apple8 hardware truth: 1 divergent constant read < 7 ALU ops (even with fma) Branches cost MORE than divergent constant reads Array indexing ALWAYS spills on Metal 4 constant addresses is the sweet spot The 4-mag LUT is the dequant-level ceiling on Apple Silicon. Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai

* vulkan: add TQ4_1S weight compression support Adds Vulkan shader support for TQ4_1S (4-bit WHT-rotated weight compression with 16 Lloyd-Max centroids, 32-element blocks). Shaders: - dequant_tq4_1s.comp: standalone dequant with WHT inverse via subgroupShuffleXor (32-thread workgroup, 5-stage butterfly) - mul_mat_vec_tq4_1s.comp: specialized MUL_MAT_VEC with inline activation pre-rotation (forward RHT on activation, centroid*scale dequant without inverse RHT) - copy_from_quant.comp: TQ4_1S dequant path with full WHT inverse - copy_to_quant.comp: TQ4_1S SET_ROWS quantization path with forward RHT, dual half-block RMS scales, 16-centroid quantization - types.glsl: block_tq4_1s struct (d0, d1, qs[16]) - dequant_funcs.glsl: TQ4_1S centroid*scale dequant (no RHT) Pipeline wiring (ggml-vulkan.cpp): - MUL_MAT, SET_ROWS, CPY supports_op - pipeline_dequant, pipeline_set_rows, pipeline_cpy_quant_f32 - Specialized MUL_MAT_VEC with forced subgroup workgroup size Tests: - test_set_rows_tq4_1s: SET_ROWS round-trip validation * vulkan: add fused mul_mat_vec kernel for TQ4_1S Adds a specialised MUL_MAT_VEC shader for TQ4_1S weights so the per-decode-step matrix-vector product no longer has to dequant the full weight tensor to f16 and then go through the generic matmul path. The kernel pre-rotates the activation via a forward Walsh-Hadamard Transform in shared memory and dot-products against the raw centroid*scale stored weights, folding the inverse-WHT on the weight side into the activation by the symmetry H = H^T. Math: w[k] = sign[k] * INV_SQRT32 * (H @ stored)[k] sum_k w[k] * a[k] = INV_SQRT32 * sum_j stored[j] * (H @ (sign * a))[j] Portability choices: - Workgroup size is pinned to 32 threads regardless of the DMMV_WG_SIZE bucket the rest of the mul_mat_vec family picks for the current architecture. The butterfly operates on 32-element blocks with one element per thread; that contract is fixed by the quantization format, not by the GPU. Earlier revisions used `gl_WorkGroupSize.x` as the stride unit, which silently skipped half the work on Intel drivers that force the subgroup to 16 (tests passed via NMSE tolerance while real inference output was garbage). - Butterfly implementation is shared memory only. A subgroup-shuffle variant (`subgroupShuffleXor`) was prototyped and measured on Intel Arc A380 with Mesa Xe HPG: it ran ~60-85 %% slower than the explicit shared-memory butterfly, because Mesa emulates subgroup shuffles via LDS and ends up doing the same LDS traffic with extra driver overhead. The shared-memory butterfly is correct on every device regardless of subgroup-op support, is the fastest path on every device we can actually measure, and leaves the `pipeline_dequant_mul_mat_vec_f32_f32[w][TQ4_1S]` slot uniform across all DMMV_WG_SIZE buckets. - Reduction is the shared-memory tree reduction (no subgroupAdd), for the same reason: on Intel Arc the subgroupAdd is also LDS-backed and the hybrid reduction path was measurably slower. Future vendor-specific heuristics can switch to the hybrid or pure-subgroup reduction variants on NVIDIA / AMD RDNA if hardware subgroup ops turn out to beat the LDS roundtrip there; the existing reduction modes in `mul_mat_vec_base.glsl` already provide the necessary variants. - NUM_ROWS is 8 so the butterfly cost amortises across 8 output rows per workgroup. Each thread holds one position of each of the 8 weight blocks and pairs them with the shared rotated activation. - `mul_mm` and `flash_attn_cm2` shader generation is skipped for TQ4_1S because it is a weight-only format that never reaches the coopmat2 matmul or the KV cache flash-attention paths. Tests: - `test-backend-ops` MUL_MAT tolerance tightened from 2.0 to 0.01 NMSE so real defects can't hide behind a loose check. - Added Gemma-4 E2B, Qwen, Phi and Llama dimensional coverage (k in {1536, 2048, 2304, 3072, 4096}, m in {256, 1152, 1536, 2048, 5120, 6144}, n in {1..8, 16, 64, 256}). 148 MUL_MAT test cases total. Verification (Intel Arc A380, 6 GB VRAM, Vulkan ANV / Mesa Xe HPG, `llama-bench -p 512 -n 128 -r 3` and `llama-perplexity -c 512 --chunks 20 wiki.test.raw`): | Model | Config | Size | Reduction | PPL Δ | pp512/Q8 | tg128/Q8 | |---------------|---------|----------:|----------:|-------:|---------:|---------:| | Qwen2.5-1.5B | I | 1570→1082 | -31.1% | +4.66% | 53.9% | 107.5% | | Phi-3.5-mini | I | 3873→2839 | -26.7% | +5.36% | 57.6% | 52.8% | | Llama-3.2-3B | hybrid | 3263→2147 | -34.2% | +2.03% | 82.4% | 84.2% | | Llama-3.2-3B | premium | 3263→2577 | -21.0% | +0.98% | 71.3% | 67.3% | Qwen2.5-1.5B is faster than its own Q8_0 baseline with Config I: the compressed model fits in less VRAM, and on a small model the TQ4_1S compute cost is offset by the reduced memory traffic. All four models produce coherent output end-to-end and the reductions line up with the TurboQuant paper's validation matrix (§5.8). The remaining gap to Q8_0 on the bigger models is compute-bound on the A380; it closes further on GPUs with more raw throughput. * vulkan: restructure TQ4_1S inner loop for cross-row smem reuse Splits the dequant+accumulate phase into two sub-loops: 1. Pre-compute w_vals[n] for all NUM_ROWS rows (centroid lookup + scale multiply, reads from weight buffer only). 2. Read the rotated activation from shared memory ONCE per column, then FMA across all rows in a tight register loop. This is the Vulkan analogue of the 'hot loop load dedup' from the CUDA kernel (PR PrismML-Eng#57 optimisation Mintplex-Labs#2). It makes the shared memory read explicitly loop-invariant across rows, which helps compilers that don't auto-hoist LDS loads out of unrolled loops. Measured effect on Intel Arc A380 (Llama-3.2-3B premium, llama-bench tg128, r=5): 15.50 -> 15.78 t/s (+1.8%, within noise but not a regression). The structure is cleaner regardless and should benefit architectures with higher LDS latency.

khosravipasha · 2026-05-05T06:32:11Z

I think will close this PR for now, as there the main Q1_0 backends are now in main llama.cpp.
The turboquant related changes might be there soon too.
(I am planning to rebase again with main llama.cpp soon sometime this week and start sending our PRs for Q2_0 format for ternary models) Feel free to send a new PR after that.

…he kernel Implements the re-triage's #2 target (elementwise glue = ~2.2ms of nearly- pure dispatch tax): env-gated GGML_GDN_FUSED_BA=1 passes RAW beta/alpha to the GDN kernel and computes sigmoid(beta) and ssm_a*softplus(alpha+ssm_dt) inline (formulas match the Metal unary kernels exactly: softplus with the x>20 guard, sigmoid 1/(1+exp(-x))). Kills sigmoid+add+softplus+mul = 4 dispatches per GDN layer (~190/token). ssm_dt/ssm_a ride as srcs 6/7 with an op_params[1] flag -> FC constant; qwen35 arch only for now (qwen3next/ qwen35moe have an interleaved mixed_ba layout needing separate handling). Correctness: token-identical on both checkpoints, and the FULL THREE-FLAG STACK (FUSED_BA + STATE_INPLACE + RECURRENT_STATE_F16) is byte-identical through a 512-token coherent ternary generation. Benchmarks still PENDING on the GPU-contention clearing (same as prior commit). Expected from the calibrated model: ~+3-4% on top of the inplace win. Metal-only, same caveat family as the sibling flags: the CPU GDN op ignores op_params[1] and would compute garbage with raw inputs -- requires full offload, which the guard's fused_gdn_ar/ch requirement effectively implies for this experiment.

Merge PR PrismML-Eng#2 from ciru-ai/fpx-vulkan-fp3-speed and credit ciru-ai for the ROCmFPX FP3 Vulkan matvec/dequant speed path.

Opus 4.8 response to /tmp/diagnostic-brief-perf.md saved at /tmp/diagnostic-brief-perf-response.md. Read-only; src untouched. Top line: kernel is staging-bound on prologue, not tensor-core-bound. Three ranked fixes: Rank 1 (H6): Stage A coalesce + single int load (also fixes sign-extension-or latent correctness bug for negative activations), kill in-loop weight decode (4x mblk redundancy). Predicted 209 -> 300-380 from Stage A alone, 450-600 with pre-expansion. Rank 2 (H2): swap kk-outer/f-inner to f-outer/kk-inner so only one acc_i[4] is live per warp (halves regs). Predicted 550-700. Rank 3 (H1): cp.async + double-buffered sA/sB SMEM. Predicted 800-900 (last mile to cuBLAS parity, not first thing to fix). Opus revised pass criterion: two-step gate. Step 1: >= 550 tok/s after minimum patch (PrismML-Eng#1+PrismML-Eng#2+PrismML-Eng#3). Step 2: >= 850 tok/s after full Hopper-style pipeline. Original single-number 856 criterion is not realistic from a hand kernel without cp.async; revise as above. Recorded decision: apply minimum patch as next concrete experiment (Phase 3.6). Hold the .cu source commits until measured lift confirms the kernel is past the 550 ground. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>

…ismML-Eng#3 no-op) Recorded measured performance from the two Opus minimum-patch edits: Patch PrismML-Eng#1 (Stage A coalesce + single int load): Cold-cache pp512 -> 470 +/- 18 tok/s (was 209 +/- 22; +2.25x) Above Opus's predicted 300-380 ceiling for that step. Patch PrismML-Eng#3 (Stage D frag-outer / kk-inner): Cold-cache pp512 -> 466 +/- 23 tok/s No measurable additional lift. Confirms Opus stacking order: H2/H1 lifts were predicated on H6 (pre-expansion) being in place. Without it, register-pressure relief just buys more warps blocked on Stage B global reads. Honest framing of current state in the ledger: ~470 tok/s == ~52% of the cuBLAS reference (899 tok/s) with the kernel running cleanly, decode (M=1) untouched, 86/86 unit tests, sanitizer clean, and 3.5 GB VRAM-resident throughout. Patch PrismML-Eng#2 (pre-expand to int8 in global VRAM) is rejected on the principle that 8x VRAM (27 GB resident) defeats the Bonsai 1-bit intelligence-density thesis. Open path: 4x mblk redundancy elimination in Stage B without VRAM blow-up -- separate Opus brief in flight. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>

This commit closes the experiment. The Blackwell int8 tensor-core MMQ kernel for Q1_0/Q2_0 weights ships at 692 tok/s cold-cache pp512 (79.5% of cuBLAS reference 871), env-gated behind GGML_BLACKWELL_Q1 defaulting to off. Decode (M=1, tg32) unchanged at 40 tok/s. Memory footprint exactly Bonsai's 3.5 GB packed throughout. 86/86 unit tests pass with env on AND env off. compute-sanitizer clean. Two negative results are part of the shipped artifact, documented so future engineers don't re-investigate the same dead ends: - Patch PrismML-Eng#2 (pre-expand weights to int8 in global VRAM) rejected on the principle that 8x VRAM (27 GB resident) defeats the Bonsai 1-bit intelligence-density thesis. - Phase 3.9 cp.async + KG=2 attempt failed the 750-tok/s gate (550 tok/s measured, 20% regression vs the 692 baseline), isolated to SMEM-doubling occupancy hit on sm_121a sync mma. Branch head: gb10-blackwell @ e45ce1b. Co-authored-by: factory-droid[bot] <138933559+factory-droid[bot]@users.noreply.github.com>

khosravipasha and others added 6 commits March 2, 2026 10:49

[cuda] fix mmq/mma path

43e67bd

Merge pull request #1 from PrismML-Eng/mmq

bc8122e

[cuda] Fix mmq/mma path

simplify dequantize_row_q1_0_g128

1b0fadf

add slim release workflow for prism

1179bfc

github-actions Bot added Nvidia GPU ggml examples labels Apr 1, 2026

khosravipasha requested a review from Copilot April 6, 2026 08:03

Copilot started reviewing on behalf of khosravipasha April 6, 2026 08:04 View session

Copilot AI reviewed Apr 6, 2026

View reviewed changes

claudlos mentioned this pull request Apr 7, 2026

vulkan: add Q1_0_g128 (1-bit ternary) shader support #9

Closed

khosravipasha force-pushed the prism branch from ba7e817 to 6df15a6 Compare April 13, 2026 22:28

khosravipasha closed this May 5, 2026

carlosfundora deleted the feature/tq3_0-kv-cache branch June 13, 2026 12:50

bong-water-water-bong pushed a commit to bong-water-water-bong/llama.cpp that referenced this pull request Jul 15, 2026

Merge ROCmFPX FP3 Vulkan path

553504b

Merge PR PrismML-Eng#2 from ciru-ai/fpx-vulkan-fp3-speed and credit ciru-ai for the ROCmFPX FP3 Vulkan matvec/dequant speed path.

Uh oh!

Conversation

carlosfundora commented Apr 1, 2026

Uh oh!

khosravipasha commented Apr 2, 2026

Uh oh!

carlosfundora commented Apr 3, 2026

Uh oh!

carlosfundora commented Apr 3, 2026

Uh oh!

khosravipasha commented Apr 3, 2026

Uh oh!

carlosfundora commented Apr 4, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

khosravipasha commented Apr 4, 2026

Uh oh!

rosmur commented Apr 5, 2026

Uh oh!

carlosfundora commented Apr 6, 2026

Uh oh!

Copilot AI left a comment

Choose a reason for hiding this comment

Pull request overview

Reviewed changes

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

khosravipasha commented Apr 6, 2026

Uh oh!

khosravipasha commented Apr 10, 2026

Uh oh!

khosravipasha commented Apr 14, 2026

Uh oh!

khosravipasha commented May 5, 2026

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

4 participants

carlosfundora commented Apr 4, 2026 •

edited

Loading