feat: OSCAR2 quantized KV-cache (draft) - #234
Closed
giveen wants to merge 445 commits into
Closed
Conversation
…default OFF Replaces the prior approach (auto-enable rotation for non-turbo quantized types) with explicit per-side opt-in via LLAMA_ATTN_ROT_K_OVERRIDE and LLAMA_ATTN_ROT_V_OVERRIDE. Default behavior: rotation OFF on both sides across all KV types. Background: TheTom/turboquant_plus#88 surfaced that asymmetric q8_0/turbo* configs were missing the upstream activation rotation from ml-explore/llama.cpp#21038. Multiple iterations (broad enable, per-side gating with turbo skip) tried to find a smart default that balances quality across model families. Empirical PPL+KLD testing on 7 model families (gemma-4 26B-A4B / 31B / E2B, Qwen2.5-7B, Qwen3.5-2B, Mistral-Small-24B, phi-4) showed the optimal rotation policy is highly model-and-quant specific. No single default is correct everywhere, including within the same architecture family (gemma-4 26B-A4B Q8, 31B Q8, and E2B Q4_K_L showed three distinct optima). phi-4 V-side rotation crashes with graph-node hash overflow, ruling out any default-on policy that touches V rotation across model families. Default OFF avoids regressing any tested model. Env-knob opt-in lets power users tune for their specific config based on documented per-model findings (see README/docs follow-up). LLAMA_ATTN_ROT_DISABLE remains as a no-op alias for historical scripts. Co-Authored-By: tturney@psyguard.ai
…disable=false) Previous state: db3595a added LLAMA_ATTN_ROT_K_OVERRIDE / _V_OVERRIDE per-side opt-in knobs but kept attn_rot_disable defaulting to TRUE for legacy LLAMA_ATTN_ROT_DISABLE compatibility. The override branches included `&& !attn_rot_disable` guards, so when LLAMA_ATTN_ROT_DISABLE is unset (default true) the per-side env knobs were silently no-ops. Users could not opt into rotation without also setting LLAMA_ATTN_ROT_DISABLE=0. Fix: flip attn_rot_disable default to false. Rotation is still OFF by default because attn_rot_k/v default to false. LLAMA_ATTN_ROT_DISABLE=1 still acts as a hard lock-out that blocks the per-side overrides for users who want a single switch to guarantee no rotation. Caught while running the cross-format KLD matrix for the rotation/PPL investigation paper — V-only override appeared to silently fail. Confirmed with logs that attn_rot_v stayed 0 even with LLAMA_ATTN_ROT_V_OVERRIDE=1 until this default flip.
Optimizations found via automated kernel optimization (33 experiments): - nthreads_KQ=1 + nthreads_V/=8 for better occupancy - Warp shuffle KQ scores (eliminates shared memory for reduction) - Precomputed scaled V centroids per block - __expf fast-math softmax - __launch_bounds__ occupancy 2 - Shmem KQ LUT: precompute Q×centroid in shared memory Also includes: - Auto-asymmetric KV: detect GQA ratio ≥6:1, upgrade K to q8_0 (fixes catastrophic PPL on Qwen2.5 symmetric turbo3) - HIP -Wnodiscard fix: (void) casts on cudaMemcpyToSymbol/FromSymbol
…nt (TheTom#78) Post-attention V-padded reshape in build_attn was using hparams.n_head_kv(il), but cur returned from build_attn_mha has shape (n_embd_head * n_head, n_tokens) — n_head is the Q-head count. On GQA models where n_head != n_head_kv (e.g. Qwen2.5-0.5B with head_dim=64 padded → 128, n_head=14, n_head_kv=2), the reshape element count fails the assertion in ggml_reshape_3d and the process aborts. Symptom: GGML_ASSERT(ggml_nelements(a) == ne0*ne1*ne2) at ggml.c:3656. Reported and diagnosed by @bingh0 in TheTom#78. Verified locally on Qwen2.5-7B (head_dim=128, no padding, regression check passes) and on AMD MI300X with Qwen2.5-0.5B (head_dim=64, was crashing pre-fix). Three sites fixed (lines 2285, 2412, 2532 — same idiom in three build_attn overloads). Closes TheTom#78. Likely also closes TheTom#108 (speculative decoding hits the same assertion). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fix: build_attn V-padded reshape uses Q-head count, not KV-head count (TheTom#78)
Cherry-pick signalnine PR TheTom#53: auto-asymmetric GQA + turbo VEC FA opts
fix(kv-cache): per-side env-knob control for upstream attn rotation (default OFF)
Two well-meaning fixes both added `spirv-headers` to the package.nix
function pattern arglist (line 19 and line 22 on current HEAD), causing
a hard parse-time failure on any nix evaluation:
error: duplicate formal function argument 'spirv-headers'
at .devops/nix/package.nix:22:3:
21| shaderc,
22| spirv-headers,
| ^
23| useBlas ?
Drops the second occurrence. The remaining single declaration is what
the rest of the file actually references (line 19 binds the input;
`vulkanBuildInputs` and `nativeBuildInputs` consume it once each).
Reported by @cguentherTUChemnitz in TheTom#81 (originally), then re-confirmed
by @alanscodelog on the current tip after two prior fix attempts both
landed the same line.
Verified: fixed file parses clean via nix-instantiate; injecting the
duplicate back reproduces the exact error message above.
Closes TheTom#81.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fix(nix): remove duplicate spirv-headers function arg (TheTom#81)
Mirror of @apollosenvy's turbo3_0 Vulkan SET_ROWS port (PR TheTom#33 + TheTom#87) to the other two turbo types. Reported by @dpblnt in TheTom#50 with a clean matrix on RX 9060 XT showing turbo3 V works on Vulkan but turbo2/turbo4 V abort with: pre-allocated tensor (cache_v_l*) in a buffer (Vulkan0) that cannot run the operation (SET_ROWS) at llama_context::sched_reserve() time, before any compute runs. Mechanical port across 4 files: - vulkan-shaders/types.glsl: block_turbo2_0 + block_turbo4_0 struct declarations matching the C side (ggml-common.h). - vulkan-shaders/copy_to_quant.comp: SET_ROWS quantize main() blocks for turbo2 (4 centroids, 2-bit pack, no signs byte) and turbo4 (16 centroids, 4-bit nibble pack, no signs byte). WHT setup and reduction structure identical to turbo3 (QK = 128 across all three). Centroid + midpoint tables ported from CENTROIDS_2BIT and CENTROIDS_4BIT in ggml-turbo-quant.c. - vulkan-shaders/vulkan-shaders-gen.cpp: turbo2_0 and turbo4_0 added to the set_rows iteration list at line ~789. - ggml-vulkan.cpp: SET_ROWS pipeline registrations + supports_op switch + dispatch element-count all extended with TURBO2_0 and TURBO4_0 cases. ## Verified on llvmpipe Vulkan (CPU software, AMD MI300X cloud droplet) Patched ggml-vulkan.cpp temporarily during repro to allow llvmpipe (normally filtered out as eCpu); patch reverted before commit. The SET_ROWS abort is a backend-capability check at graph build time so it fires regardless of GPU vs CPU Vulkan backend. | ctk / ctv | tg16 (t/s) | status | |-------------------|-----------:|---------------| | q4_0 / q4_0 | 17.68 | baseline | | q4_0 / turbo3 | 5.91 | already worked| | q4_0 / turbo4 | 6.14 | was aborting | | q4_0 / turbo2 | 5.65 | was aborting | llvmpipe perf numbers are not meaningful (CPU-emulated Vulkan); they are reported here only to confirm the abort is gone and the kernels run end-to-end without divergence. ## Needs GPU validation Cannot validate GPU shader correctness on the droplet (MI300X SR-IOV VF does not expose itself to RADV/amdvlk on cloud). Specifically: - Subgroup shuffle / ballot behavior on real GPU subgroup sizes - Shader compilation under non-llvmpipe Vulkan drivers - PPL / quality on the actual quantization math @dpblnt @apollosenvy if either of you has cycles, would appreciate a quick rebuild on RDNA Vulkan (gfx1100/gfx1200) to confirm: 1. The SET_ROWS abort that triggered TheTom#50 is gone 2. Output coherence on turbo4 V (not garbage tokens) 3. PPL stays in the expected ballpark vs the CUDA / Metal implementations of the same quants Closes TheTom#50. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
…s-turbo24 vulkan: add SET_ROWS support for turbo2_0 and turbo4_0 (TheTom#50)
…869) (ggml-org#22267) * server: clamp n_discard to non-negative at JSON parse boundary (CVE-2026-21869) A negative n_discard from client JSON causes heap-buffer-overflow in update_slots() context-shift loop (CWE-787, CVSS 8.8). Clamp to 0 at ingress; n_discard=0 already triggers auto-discard (n_left/2). Ref: GHSA-8947-pfff-2f3c * cont : cleaner * cont : cleanerer * cont : cleanest --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…d-clamp security: cherry-pick CVE-2026-21869 (n_discard heap-buffer-overflow in server)
When users set --cache-type-k turbo3 (or turbo2/4) without
explicitly setting --cache-type-v, V defaults to F16 and the
(TURBOx_0, F16) pair hits ggml_cuda_flash_attn_ext_vec without a
matching FATTN_VEC_CASE → GGML_ABORT("fatal error") at fattn.cu:348.
The reverse direction (F16 K + turbo V) was already instantiated
for all three turbo variants. This adds the matching
(turbo K + F16 V) pairs.
Reported on Radeon 8060S (gfx1151) with Qwen2.5-7B-Instruct-Q4_K_M
running --cache-type-k turbo3 --cache-type-v f16. Not GPU-specific
— same crash on any CUDA/HIP target with that flag combo.
Files added:
- fattn-vec-instance-turbo2_0-f16.cu
- fattn-vec-instance-turbo3_0-f16.cu
- fattn-vec-instance-turbo4_0-f16.cu
Files updated:
- fattn.cu (3 dispatch entries)
- fattn-vec.cuh (6 extern decls)
- CMakeLists.txt (3 entries in non-FA_ALL_QUANTS list)
Mac/Metal build verified. CUDA/HIP build needs validation on a
target with the toolchain (compile-only, no behavior change for
existing instantiated combos).
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
) When --kl-divergence-base is set, line 527 sizes the log_probs vector with int * int. For Qwen3-class models (vocab=151,936) at n_ctx=16384, the product is 16384 * 151940 = 2,489,610,240 which overflows INT32_MAX (2,147,483,647). The wrapped negative value sign-extends to a giant size_t when passed to vector::resize, exceeds vector::max_size, and throws std::length_error. Reproduced on M5 Max with Qwen3.6-35B-A3B Q8_0: llama-perplexity -m <model> -f wiki.test.raw -c 16384 -fa 1 \ --kl-divergence-base /tmp/baseline.dat perplexity: saving all logits to /tmp/baseline.dat perplexity: tokenizing the input .. perplexity: tokenization took 316 ms perplexity: calculating perplexity over 18 chunks ... libc++abi: terminating due to uncaught exception of type std::length_error: vector Smaller-vocab models (vocab <= ~128K) and shorter context (n_ctx <= 8192) do not trip the overflow, which is why this only surfaces on Qwen3 family at the standard 16K PPL bench depth. Same size_t cast already exists at line 514 for the parallel logits.reserve allocation; this brings line 527 to the same convention. After fix the same command runs cleanly to completion. Validated A/B on Qwen3.6-35B-A3B Q8_0 at 16K context against wikitext-2-raw: Cache | PPL | KL Div | Top-1 agree -------+------------------+-----------------+------------- f16 | 5.3513 +/- 0.032 | (baseline) | - q8_0 | 5.3508 +/- 0.032 | 0.0014 +/- 0e-5 | 98.708% turbo3 | 5.3907 +/- 0.032 | 0.0121 +/- 1e-4 | 95.293%
Commit e69af78 added 3 new dispatch entries to fattn.cu for the (turbo2/3/4, F16) mixed-KV combinations and the matching template-instance .cu files, but only updated ggml/src/ggml-cuda/CMakeLists.txt. The parallel list in ggml/src/ggml-hip/CMakeLists.txt was missed, so the HIP build links without those instantiations and fails: ld.lld: error: undefined symbol: void ggml_cuda_flash_attn_ext_vec_case<64, TURBO3_0, F16>(...) void ggml_cuda_flash_attn_ext_vec_case<128, TURBO3_0, F16>(...) void ggml_cuda_flash_attn_ext_vec_case<256, TURBO3_0, F16>(...) (and same for TURBO2_0, TURBO4_0) clang++: error: linker command failed with exit code 1 Surfaced first by LocalAI's hipblas-turboquant build job (mudler/LocalAI#9740 CI). Fix is mechanical: mirror the 3 new entries from the CUDA CMakeLists into the HIP CMakeLists, paired next to their existing f16-X counterparts.
…ache Wholesale sync of 384 upstream commits since merge-base 7fc1c4e (2026-04-22). Headline upstream feature: MTP / Multi-Token Prediction (ggml-org#22673) + spec-decoding stack (ggml-org#22838 parallel drafting, ggml-org#22227 spec-simple checkpoints, ggml-org#19493 server spec checkpointing, plus 5 spec bug-fixes). 11 conflicts resolved across CUDA fattn / Metal / Vulkan / common: ggml/src/ggml-cuda/fattn-mma-f16.cuh RDNA config matrix: union TQ's (640, 512) entries with upstream's expanded (112..576) RDNA matrix. Took upstream's new sentinel fallback (no ampere fallback for RDNA). ggml/src/ggml-cuda/fattn.cu - Extended hoisted ncols2_max to include 640 head-dim. - Volta: dropped TQ's local ncols2_max redefinition in favor of upstream's hoisted version (with 640 added). - WMMA gate: union exclusions (40, 72, 192, 512, 576, 640). - Preserved TQ's RDNA4 vector-kernel branch for TurboQuant cache types (renamed inner gqa_ratio_eff_rdna4 to avoid shadowing); took upstream's restructured MFMA/CDNA path verbatim. ggml/src/ggml-cuda/ggml-cuda.cu Supported-op switch: union TQ's GGML_OP_TURBO_WHT case with upstream's GGML_OP_ADD/SUB/MUL/DIV FP16 cases. ggml/src/ggml-metal/ggml-metal-device.h Kept TQ's get_pipeline_turbo_wht declaration; took upstream's new get_pipeline_mul_mv_ext(lib, const ggml_tensor * op, ...) signature (replaces split tsrc0/tsrc1 args). ggml/src/ggml-metal/ggml-metal-device.cpp Kept TQ's get_pipeline_turbo_wht implementation; took upstream's new get_pipeline_mul_mv_ext signature — body already uses op-> for tsrc0/tsrc1/ne12/r2/r3. ggml/src/ggml-metal/ggml-metal-ops.cpp Preserved TQ's is_tq_weight rotate→matmul→unrotate path with original hardcoded dispatch shape. Updated non-TQ fallback to upstream's pipeline- param dispatch (pipeline.nr0 / nr1 / nsg + (ne11+nr1-1)/nr1 shape). ggml/src/ggml-vulkan/* (3 files) Upstream-wholesale via `git checkout --theirs`. Upstream architecturally refactored FA from compile-time DATA_A_* variants to runtime FaTypeK/FaTypeV spec-constant switches. TQ's TURBO3_0 GLSL path is DEFERRED — Vulkan TURBO3_0 support needs re-implementation against the new architecture in a follow-up PR. Mac mini + M5 Max have no Vulkan; no in-house validation path for an immediate re-adaptation. common/arg.cpp --spec-default: took upstream's new struct shape (params.speculative.types vector + params.speculative.ngram_mod.{n_match,n_min,n_max}). common/speculative.cpp Low-acceptance reset: took upstream's sinfo.n_low / sinfo.i_last (variables moved into sinfo struct). NOT-CONFLICTED upstream additions that touch TQ-adjacent surface (auto-merged clean, but worth eyes during review): - src/llama-memory-recurrent.{cpp,h} (MTP rollback API) - src/llama-memory-hybrid.{cpp,h} (recall feedback_llama_memory_types + feedback_layer0_hybrid_trap) - src/llama-graph.cpp, src/llama-kv-cache.cpp, src/llama-context.cpp - tools/server/server-context.cpp (+~1100 lines: MTP + parallel drafting + spec checkpointing) - src/models/qwen35*.cpp, qwen3next.cpp, delta-net-base.cpp (entirely new in upstream — MTP draft-head integration) Known-deferred follow-ups: 1. Vulkan TURBO3_0 re-implementation against runtime spec-constant FA arch 2. PR ggml-org#21245 QKV refactor helpers — landed; TQ models not migrated to use them. Migrate in a focused follow-up; do not bundle here. Validation gate (pending): M2 mini PPL/decode comparison @ Qwen2.5-7B-Q8_0 K=q8_0/V=turbo4 asymmetric ctx 2048 + 16384 — see PR body. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
…ed pipeline
Regression introduced by the upstream b9190 merge. Upstream moved ne12/ne13/r2/r3
from kernel args to function constants (FC_MUL_MM + 2..5) in kernel_mul_mm.
get_pipeline_mul_mm was updated to set them; get_pipeline_mul_mm_tq_rotated was
not, so the TQ-rotated kernel templates (which reuse the same kernel_mul_mm
body) read those constants as zero, producing wrong tensor offsets → NaN/inf
outputs.
Symptoms caught by extended M2 mini validation:
- test-backend-ops MTL0: ALL MUL_MAT(type_a=tq3_1s|tq4_1s, ...) FAIL with
ERR = inf > 0.000500000 (and one NaN at index 136). 279 TQ MUL_MAT tests
pass on TQ tip 5aeb2fd.
- Qwen2.5-7B-TQ4_1S PPL: 146/146 chunks "nan" both dense and asym KV
(vs 6.7530 / 6.7887 baseline).
- Qwen2.5-7B-TQ4_1S decode bench tg128: 7.47 → 3.95 t/s = -47% regression,
pp128 variance ±0.27 → ±40.85 indicating dispatch chaos.
Fix:
- Compute ne12/ne13/r2/r3 in get_pipeline_mul_mm_tq_rotated identical to
get_pipeline_mul_mm.
- Set FC_MUL_MM + 2..5 alongside the existing bc_inp/bc_out constants.
- Include ne12/ne13/r2/r3 in the pipeline cache name so different tensor
shapes don't collide on a single compiled pipeline (cache poisoning).
The MUL_MAT_ID variant (get_pipeline_mul_mm_id_tq_rotated) mirrors
get_pipeline_mul_mm_id which only sets bc_inp — kernel_mul_mm_id is a
different template and doesn't need ne12/ne13/r2/r3, so no change there.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
CI failure (every -Werror build job — ubuntu-latest-cuda, ggml-ci-x64-cpu-*,
arm64-cpu-*, macOS-latest-{x64,arm64,arm64-webgpu}, ubuntu-22-{hip-quality-check,musa},
android-arm64, ubuntu-cpu x64-22.04):
ggml/src/ggml-turbo-quant.c:247:15: error: no previous prototype for
'turbo_cpu_fwht_inverse' [-Werror=missing-prototypes]
Pre-existing issue on the TQ branch — the function is defined GGML_API but
has no prototype declaration. M5 Max + M2 mini builds don't use -Werror so it
slid through local validation. Upstream CI does.
Fix: forward-declare the function near the top of the .c file. Matches the
extern declaration already used by tests/test-turbo-quant.c. No semantic change.
The other GGML_API symbol in this file (turbo3_cpu_wht_group_size) is a
variable, not a function — -Wmissing-prototypes does not apply.
Flagged by @pacak on PR TheTom#146 (ubuntu-latest-cuda CI). Restores green CI on
all -Werror build jobs without affecting any runtime path.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
Two CI build failures both pre-existing on TQ tip, exposed by upstream-policy CI on PR TheTom#146: 1. ubuntu-22-hip-quality-check — fattn-common.cuh:1312-1313 in TQ's HIP hip_f16_alloc destructor calls hipStreamSynchronize / hipFree without consuming the return value. HIP's recent runtime declares these [[nodiscard]] and the HIP quality build uses -Werror: error: ignoring return value of type 'hipError_t' declared with 'nodiscard' attribute [-Werror,-Wunused-value] Fix: (void) cast both calls. We're in a destructor and can't propagate errors anyway; intent is fire-and-forget cleanup. Matches the idiom used in upstream code for the same situation. 2. ubuntu-22-musa — turbo-quant.cuh uses cudaMemcpyToSymbol in InnerQ calibration; MUSA's vendor header (vendors/musa.h) aliases every other cudaMemcpy* variant but missed cudaMemcpyToSymbol. Result on MUSA build: error: use of undeclared identifier 'cudaMemcpyToSymbol'; did you mean 'musaMemcpyToSymbol'? Fix: add the missing alias next to the other cudaMemcpy* defines. Mirrors the same alias already present in vendors/hip.h:143. Both are TQ-only paths (HIP f16 alloc was added in 0757ff4 2026-04-18; MUSA was never built against TQ in-tree). M5 Max + M2 mini local Metal builds unaffected by either change. Flagged by @pacak on PR TheTom#146 (ubuntu-latest-cuda + cross-vendor CI fails). Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
Cherry-picked from upstream ggml-org/llama.cpp@87589042c (merged 2026-05-17). option(LLAMA_BUILD_WEBUI ... ON) always leaves the deprecated flag DEFINED, so the compat-block guard `AND NOT DEFINED LLAMA_BUILD_UI` never fires. tools/ui/CMakeLists.txt then ORs both flags, so passing only the new `-DLLAMA_BUILD_UI=OFF` was silently ignored. Removes the deprecated options and simplifies the compat block + UI gate to a single flag. Fixes the nix-sandbox build failure reported by @arch-fan and @pacak on PR TheTom#146 — both hit the resulting xxd.cmake crash when an empty tools/ui/dist/index.html was produced by failed npm + HF Bucket provisioning. After this cherry-pick, `-DLLAMA_BUILD_UI=OFF` alone works as documented. Co-Authored-By: TheTom <tturney@psyguard.ai> Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Round 2 of CI fixes addressing the remaining red jobs on the b9190 sync
PR. All were pre-existing TQ-tip bugs exposed by upstream CI's -Werror
policy (M5 Max + M2 mini local builds don't use -Werror).
1. ggml/src/ggml-cuda/fattn-mma-f16.cuh — fall back to ampere config
(not zero-sentinel) in get_config_rdna
----------------------------------------------------------------
Reverts the round-1 conflict choice. Round 1 took upstream's new
sentinel `fattn_mma_config(32, 1, 0, 0, 0, 0, 0, false)` for the
RDNA fallback. Template instances like
fattn-mma-f16-instance-ncols1_1-ncols2_16.cu do constexpr arithmetic
on the returned config (np = nwarps * cols_per_warp / ncols, etc).
nwarps=0 from the sentinel propagates to np=0, triggering compile-
time div/mod-by-zero at lines 1265/1371/1375/1512/1519/1572. HIP
quality build is -Werror,-Wdivision-by-zero so it errors out.
TQ-tip behavior (delegate to ampere) returns a valid config —
restore it. Keeps all (640, 512) RDNA entries unioned in round 1.
2. ggml/src/ggml-cuda/vendors/musa.h — add cudaMemcpyFromSymbol alias
----------------------------------------------------------------
turbo-quant.cuh InnerQ calibration uses both cudaMemcpyToSymbol AND
cudaMemcpyFromSymbol. Round-1 fix added _ToSymbol; _FromSymbol was
missed. Mirrors vendors/hip.h line 142.
3. src/llama-kv-cache.cpp — [[maybe_unused]] stubs + remove unused `il`
----------------------------------------------------------------
The non-CUDA stub block (g_innerq_finalized, g_innerq_scale_inv_host,
turbo_innerq_needs_tensor_update, turbo_innerq_mark_tensor_updated)
are declared static but every consumer is gated by #ifdef GGML_USE_CUDA,
so the file-local copies look unused on non-CUDA builds. Annotate
with [[maybe_unused]]. Also drops two `const uint32_t il = layer.il;`
locals in the state-save k/v writer loops where `il` was unreferenced —
dead-code from a removed logging pass.
4. scripts/xxd.cmake — defensive quote of ${hex_data}
----------------------------------------------------------------
Belt-and-suspenders for the LLAMA_BUILD_UI nix-sandbox failure. The
primary fix is the cherry-pick of upstream PR ggml-org#23190 (previous
commit), which makes -DLLAMA_BUILD_UI=OFF actually work. This patch
makes the underlying xxd.cmake robust: when an empty UI source file
slips through, produce a 0-length .hpp instead of crashing with
cmake's cryptic "string sub-command LENGTH requires two arguments"
error. Worth proposing upstream as a follow-up.
Local Metal build green on M5 Max with all four fixes applied.
Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
Co-Authored-By: tturney@psyguard.ai
arch-fan's next nix-sandbox build (after PR ggml-org#23190 cherry-pick + earlier empty-input defensive quote) hit a different xxd.cmake failure: scripts/xxd.cmake:10 (file): file failed to open for reading (No such file or directory): /build/source/build/tools/ui/dist/bundle.js Empty-file case (LENGTH error) was already handled by quoting the variable. This is the sibling case: file READ itself fails when the UI provisioning flow leaves an asset missing entirely (npm absent AND HF Bucket download blocked → some assets created empty, some not created at all). Fix: early-return with a valid 0-byte symbol when ${INPUT} doesn't exist. Also unify the empty-content path to emit {0} instead of {} (zero-element array initializer is C++ extension, not portable). Verified end-to-end on M5 Max by reproducing arch-fan's exact conditions: build/tools/ui/dist/ removed, PATH stripped of npm, LLAMA_USE_PREBUILT_UI=OFF. Without the fix, build crashes on bundle.js.hpp generation. With the fix, all four .hpp files generate as 0-byte symbols, llama-ui target completes cleanly, server builds with LLAMA_UI_DEFAULT_ENABLED=0 (no embedded UI but no crash) — exactly upstream's intended graceful degradation. No effect on normal builds with UI assets present (regenerated all 4 .hpp files at original 26MB / 2.5MB / 34KB / 1.4KB sizes, byte- identical to pre-fix output). Worth proposing upstream as defensive hardening for the xxd helper. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
sync: upstream master b9190 + MTP/spec stack (DO NOT MERGE — tester review)
Short README calling out what this fork adds vs upstream ggml-org/llama.cpp: TurboQuant KV-cache and weight quantization types (turbo3, turbo4, TQ3_1S, TQ4_1S) and their CUDA / HIP-ROCm / Metal / Vulkan kernel integrations. Points at TheTom/turboquant_plus for the codec design and papers, and TheTom/tqkit for cross-backend bench results. Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com> Co-Authored-By: tturney@psyguard.ai
Restore the upstream llama.cpp README so users have the standard build / model / bindings docs, with a short fork-specific intro at the top calling out the TurboQuant additions and pointing at TheTom/turboquant_plus for the codec design.
Cross-checked against commit log diff vs upstream + paper corpus: - add turbo2 KV-cache format (was already in codebase but missed) - add auto-asymmetric K/V compression policy (PR TheTom#53) - add Boundary V (layer-aware experimental V compression for turbo2-V) - add Sparse V dequantization (on all Metal targets) - add turbo VEC flash attention +9% decode on CUDA - add V2.1 fused Metal TQ kernels - add TurboFlash Apple10 known-limitation caveat - add Vulkan SET_ROWS support for turbo2/turbo4 - note turbo block size moved from 32 to 128 Direct paper links inline (TheTom/turboquant_plus papers) instead of vague references. Drop the tqkit cross-reference per scope.
Move fork README to a professional structure suitable for downstream projects evaluating the integration: - Title + tagline + badges (license / status / paper corpus) - Production deployments section (LocalAI, AtomicChat, others) - Status table (branch, commits ahead, upstream tracking) - Quantization types table with bits / notes / paper links - Compression policies (asymmetric K/V, Boundary V, sparse V) - Backend coverage matrix with per-backend notes + caveats - Model-family support + operational fixes the fork carries - Quick start + usage with concrete invocations - Citation pointer to the TurboQuant+ paper corpus Upstream llama.cpp README preserved verbatim below the fork section so users still have the standard build / model / bindings docs.
- Add Chronara.io (quantum-safe fintech infrastructure) as a production user of this fork. - Link AtomicChat (https://atomic.chat/).
Restructure the KV-cache usage section to make the asymmetric-turbo
pattern unmissable and to guide users to escalate compression instead
of starting at maximum and walking back:
- Add prominent "Start light, then compress" callout — some model
families are more delicate (small models, certain MoE configs,
quant-sensitive instruction-tuned variants).
- Reframe the recommendations table as a 5-step ladder, from safest
('f16'-K + 'turbo3'-V, first-contact) through recommended default
('q8_0'-K + 'turbo3'-V, asymmetric turbo sweet spot) to aggressive
V and MoE-aware aggressive, ending with the discouraged symmetric-K
config and its failure-mode citation.
- Add concrete CLI examples for steps 1, 2, 3.
- Add closing reminder that the frontier is per-model — walk back a
step if quality drops.
github-actions
Bot
added
ggml
examples
Apple Metal
testing
model
CUDA
conversion
labels
Author
OSCAR2 KV Cache BenchmarkDate: 2026-07-26 (updated from 2026-07-22) Verified 4K Baseline (2026-07-26)
Commands: # f16 baseline
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 4096 \
--cache-type-k f16 --cache-type-v f16 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"
# oscar2
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 4096 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"Context Scaling — Verified 2026-07-26 (Qwen3.6-27B, oscar2/oscar2)
Generation speed is flat across all context sizes — no degradation from 4K to 1.16M. VRAM AnalysisKV cache cost: ~10 MiB per 1,024 tokens (both K+V combined)
Maximum oscar2 context on 32 GB RTX 5090: 1,160,000 tokens (1.16M) Commands: # 128K
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 131072 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"
# 256K
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 262144 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"
# 512K
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 524288 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"
# 1M
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 1048576 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"
# 1.16M (max)
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 1160000 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--temp 0 --top-p 0.95 --min-p 0.05 -n 128 --single-turn \
-p "The capital of France is Paris. The capital of Germany is Berlin. The capital of Italy is Rome. The capital of Spain is"Context Limits: f16 vs oscar2 (Qwen3.6-27B)
Maximum f16 context: ~170K on 32 GB card # Test the OOM boundary for f16
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-cli \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 196608 \
--cache-type-k f16 --cache-type-v f16 \
--temp 0 -n 1 --prompt "x" --single-turnoscar2 vs turbo2 (at 256K context)
oscar2 delivers 6.8× more context than turbo2 in 32 GB (1.16M vs ~170K tokens). Compression Ratios
OSCAR2 compresses slightly less aggressively than turbo2 (2.25 vs 2.12 bits/elem) because of the fp16 zero-point metadata, but it enables symmetric INT2 compression for both K and V, unlike turbo2's asymmetric policy which forces K to q8_0. Optimization Progression
*The 68.0 t/s vs f16 64.9 was measured on a slightly different build. Current f16 baseline is 64.3 t/s. Generation speed is now 45% faster than f16 at large contexts, reversing the original speed penalty. This is because the tensor-core MMA path benefits from reduced memory traffic: oscar2's INT2 blocks transfer 7× less data from global memory than f16, and the on-chip dequant overhead is hidden behind shared memory latency while tensor cores run at full throughput. How the Tensor-Core Path Works
Files changed for the tensor-core path
Known Issues Fixed
Rotation Model ConversionTo generate Hadamard rotations for any GGUF model and bake them in: # One-step: generate + bake
python3 oscar-rotation/generate_and_bake_rot.py \
--base model.gguf --out model-rot-kv.gguf \
--method hadamard
# Standalone Hadamard generator
python3 oscar-rotation/generate_hadamard_rot.py \
--head-dim 256 --num-layers 48 --output-dir .
# Bake existing .pt rotation files
python3 oscar-rotation/export_rot_kv_gguf.py \
--base model.gguf --out model-rot-kv.ggufSupports per-layer head dimensions (Gemma-4 SWA: 40 layers at 256, 8 at 512). Server Test (llama-server + oscar2)The server loads with oscar2 at up to 1.16M context and produces coherent responses. # Start server with max context
CUDA_VISIBLE_DEVICES=0 ./build/bin/llama-server \
-m /mnt/storage/models/oscar-rotations/qwen3.6-27b-q5kxl-hadamard.gguf \
-ngl 99 -fa on -c 1160000 \
--cache-type-k oscar2 --cache-type-v oscar2 \
--host 127.0.0.1 --port 18080Server metrics at 512K (single request):
|
… enable prefill parallelism Changes: - fattn-common.cuh: remove +m from VEC oscar2 K dequant (mean is position-dependent bias — wrong if VEC ever dispatched for oscar2) - fattn-mma-f16.cuh: add restore_mean template param for V tile load (correct for MMA path when re-enabled) - fattn-oscar2.cuh: fix nbatch_fa = D instead of K->ne[1] to enable multi-warp parallelism during prefill via flash_attn_combine_results (+21% pp4096 throughput) - set-rows.cu: add LLAMA_KV_CLIP_RATIO percentile clipping to oscar2 quantizer (matches q2_0 behavior, matches vLLM VLLM_OSCAR_K_CLIP_RATIO) Assisted-by: deepseek-v4-flash
Author
2026-07-26 — nbatch_fa fix (prefill parallelism)Change: Before: Results (Qwen3.6-27B Q5_K_XL, oscar2 KV)
Decode unchanged (position count < D → Other changes in this build
|
- OSCAR2_BUGS.md: Re-evaluate all 20 bugs (B1-B20) against current code. Mark B1/B2/B5/B8/B10/B11/B13 as FIXED, B3 as PARTIALLY FIXED, B4/B6 as still open, B7/B12/B16/B18/B19/B20 as NOT A BUG. Update VEC path status: fundamental Hadamard domain mismatch, not fixable without VEC kernel rewrite. Add B17 verification (struct width correct). - OSCAR-PORT-STATUS.md: Update verification table with per-bug footnotes. Rewrite Known Issues 1-5 with accurate current status. Mark rotation fallback (issue 3) as FIXED. Document VEC path as fundamentally broken (not just D>256). Assisted-by: Claude Code
GPT 5.6 SOL identified that the dedicated FA kernel drops the stored K mean from the KQ dot product. The code comment claims mean 'doesn't affect softmax' but this is incorrect: mean varies per K token, and the missing term mean(K_block) * sum(Q_block) varies per token and per Q position, directly corrupting attention score rankings. This is now the leading hypothesis for garbled oscar2 output. The V path correctly handles its mean via VKQ_mean accumulation, proving the design pattern exists. Also corrects Issue 1 analysis: the normalized Hadamard is orthonormal (H_norm^T @ H_norm = I), not scaled-by-128. Assisted-by: Claude Code
P0 — Garbled output fixes: - B21: Add per-block K mean to KQ dot (fattn-oscar2.cuh). The stored block_oscar2.m was omitted from the KQ dot product with an incorrect comment that it 'doesn't affect softmax'. The mean varies per K token so the missing term mean(K_block) * sum(Q_block) is token-dependent. Fix: add m_k * Q_had_DC * sqrt(128) per block, only thread 0 holds the DC component after Hadamard transform. - Disable VEC path for oscar2 (fattn.cu). VEC kernel lacks inverse Hadamard, set_rows-stored Hadamard values read as natural domain. Removed all VEC oscar2 template instantiations. P1 — Quality: - Remove LLAMA_KV_V_ROT gate on Gemma4 (gemma4.cpp). Align with Qwen3 pattern: V rotation unconditional when attn_v_rot tensor exists. - Auto-detect rotation for q2_0 NO_HADAMARD (llama-kv-cache.cpp). When q2_0 KV cache used with rotation tensors, auto-set the env var. - INT2-without-rotation warning (llama-kv-cache.cpp). Warn when oscar2 or q2_0 is selected without any rotation tensors in the model. P2 — Hardening: - Add nb11 stride assert (fattn-oscar2.cuh). Catches pre-rotated K or zero-padded K stride mismatches. - Fix B3 dst_ptr index for batch>1 (fattn-oscar2.cuh). Replace ne01.z (batch) with ne01.x (ncols) in output indexing. P3 — Cleanup: - Remove unused P_BR_DEV table (fattn-oscar2.cuh). P_br is CPU-only, GPU set_rows does not apply bit-reversal permutation. Assisted-by: Claude Code
- gemma4.cpp: Replace TURBO_ROTATION_RT #include with runtime Sylvester construction of normalized Hadamard matrix for any power-of-2 head dim >= 64. D=256 and D=512 no longer abort. - OSCAR-PORT-STATUS.md: Update limitation note for D>128. Assisted-by: Claude Code
Documents that CPU quant path stores natural-domain values (no Hadamard, optional P_br) while GPU set_rows stores Hadamard-domain values (no P_br). Blocks from one path must not cross to the other. Assisted-by: Claude Code
B21 fix verified working on Qwen3.6 models at D=128 and D=512. Gemma4 failures are a pre-existing ISWA cache issue affecting ALL quantized types, not specific to oscar2 FA kernel. Assisted-by: Claude Code
Author
|
So basically where I stand is , chatting works just fine, but when given a coding task or agentic tasks it starts outputting crap, this is why this is staying in draft till i resolve this. |
B21 fix verified for short completions and thinking/reasoning. Instruction-style prompts show systematic token bias with oscar2. Hypothesis: fp16 m_k precision interacting with float32 Q_had + attn_scale. Suggested fix: apply mean correction in Hadamard domain. Assisted-by: Claude Code
Merge OSCAR2_BUGS.md, OSCAR2_PERF.md, and docs/oscar-baseline-commands.md into OSCAR-PORT-STATUS.md as appendices. Added table of contents, status overview section, and refreshed stale model paths in baseline commands. Delete the absorbed files. Assisted-by: Claude Code
- convert.cu: Remove k_dequantize_oscar2_ih + dequantize_oscar2_ih_cuda (zero callers, KV-cache-only type never needs general fp16 conversion). Replace dispatch case with nullptr + explanatory comment. Fix F32 fallthrough bug (was falling into BF16 handler). - fattn-oscar2.cuh: Update stale 'P_br(H) domain' comments to 'Hadamard domain' — P_br is CPU-only, not applied on GPU. - gemma4.cpp: Remove unused #include '../turbo-rotation-data.h' (replaced by runtime Sylvester Hadamard generator). - oscar-rotation/kld_oscar2_vs_f16.py: Remove TODO-stub with no implementation. Assisted-by: Claude Code
This was referenced Jul 30, 2026
Merged
TheTom
force-pushed
the
feature/turboquant-kv-cache
branch
from
August 2, 2026 03:02
8a891f4 to
28c68fe
Compare
Owner
|
Heads-up: |
Author
once the rebase is done, i'm going to completely redo this work. |
giveen
pushed a commit
to giveen/llama-cpp-turboquant
that referenced
this pull request
Aug 3, 2026
Rebuild of PR TheTom#234 on the post-rebase turboquant-kv-cache lineage: - GGML_TYPE_OSCAR2 (ID 48): INT2 per-128-elem quant with fp16 sigma/mean - CUDA: set_rows quantizer (Hadamard-domain encode), dedicated FA kernel (fattn-oscar2.cuh), MMA-F16 tile loader, vec/MMA/tile dispatch, cpy/convert/ get_rows support, D=512 (Gemma-4) MMA instances - CPU fallback: oscar2 quant/dequant + vec_dot in ggml-quants.c/ggml-cpu - llama core: attn_k_rot/attn_v_rot GGUF rotation tensors, per-layer head-dim guard, SWA pre-scan, V rotation undo (R^T), qwen3/qwen3moe/gemma4 wiring - Tooling: oscar-rotation/ scripts, tools/oscar-calib, llama-bench parser, run_oscar_tests.sh, test-backend-ops oscar2 cases, docs Skipped vs PR TheTom#234: q2_0/q2_preh KV types (not in base, no upstream q2_0 KV support either), HP sink+recent buffer system (q2_0-coupled, mixed-FA kernel cannot compile against the standard block_q2_0 and was never validated with oscar2), VEC-on-Blackwell disable (regressed hsk=72+sinks FA tests; VEC is not broken on this machine), chat parser workarounds, noise/docs churn. Also fixes a merge-induced bug: the FA dispatch switch lost its BEST_FATTN_KERNEL_TILE case label, silently dropping tile-kernel execution for hsk=72/sinks configs (garbage attention, ERR~1 in test-backend-ops). Bumped the per-kernel dynamic-smem tracking cap from 64 to 4096. Known: oscar2 FA test-backend-ops cases are impossible as written (harness feeds natural-domain data, kernel expects Hadamard-domain), matching PR TheTom#234. Assisted-by: omp
giveen
added a commit
to giveen/llama-cpp-turboquant
that referenced
this pull request
Aug 3, 2026
Rebuild of PR TheTom#234 on the post-rebase turboquant-kv-cache lineage: - GGML_TYPE_OSCAR2 (ID 48): INT2 per-128-elem quant with fp16 sigma/mean - CUDA: set_rows quantizer (Hadamard-domain encode), dedicated FA kernel (fattn-oscar2.cuh), MMA-F16 tile loader, vec/MMA/tile dispatch, cpy/convert/ get_rows support, D=512 (Gemma-4) MMA instances - CPU fallback: oscar2 quant/dequant + vec_dot in ggml-quants.c/ggml-cpu - llama core: attn_k_rot/attn_v_rot GGUF rotation tensors, per-layer head-dim guard, SWA pre-scan, V rotation undo (R^T), qwen3/qwen3moe/gemma4 wiring - Tooling: oscar-rotation/ scripts, tools/oscar-calib, llama-bench parser, run_oscar_tests.sh, test-backend-ops oscar2 cases, docs Skipped vs PR TheTom#234: q2_0/q2_preh KV types (not in base, no upstream q2_0 KV support either), HP sink+recent buffer system (q2_0-coupled, mixed-FA kernel cannot compile against the standard block_q2_0 and was never validated with oscar2), VEC-on-Blackwell disable (regressed hsk=72+sinks FA tests; VEC is not broken on this machine), chat parser workarounds, noise/docs churn. Also fixes a merge-induced bug: the FA dispatch switch lost its BEST_FATTN_KERNEL_TILE case label, silently dropping tile-kernel execution for hsk=72/sinks configs (garbage attention, ERR~1 in test-backend-ops). Bumped the per-kernel dynamic-smem tracking cap from 64 to 4096. Known: oscar2 FA test-backend-ops cases are impossible as written (harness feeds natural-domain data, kernel expects Hadamard-domain), matching PR TheTom#234. Assisted-by: omp
giveen
added a commit
to giveen/llama-cpp-turboquant
that referenced
this pull request
Aug 5, 2026
Rebuild of PR TheTom#234 on the post-rebase turboquant-kv-cache lineage: - GGML_TYPE_OSCAR2 (ID 48): INT2 per-128-elem quant with fp16 sigma/mean - CUDA: set_rows quantizer (Hadamard-domain encode), dedicated FA kernel (fattn-oscar2.cuh), MMA-F16 tile loader, vec/MMA/tile dispatch, cpy/convert/ get_rows support, D=512 (Gemma-4) MMA instances - CPU fallback: oscar2 quant/dequant + vec_dot in ggml-quants.c/ggml-cpu - llama core: attn_k_rot/attn_v_rot GGUF rotation tensors, per-layer head-dim guard, SWA pre-scan, V rotation undo (R^T), qwen3/qwen3moe/gemma4 wiring - Tooling: oscar-rotation/ scripts, tools/oscar-calib, llama-bench parser, run_oscar_tests.sh, test-backend-ops oscar2 cases, docs Skipped vs PR TheTom#234: q2_0/q2_preh KV types (not in base, no upstream q2_0 KV support either), HP sink+recent buffer system (q2_0-coupled, mixed-FA kernel cannot compile against the standard block_q2_0 and was never validated with oscar2), VEC-on-Blackwell disable (regressed hsk=72+sinks FA tests; VEC is not broken on this machine), chat parser workarounds, noise/docs churn. Also fixes a merge-induced bug: the FA dispatch switch lost its BEST_FATTN_KERNEL_TILE case label, silently dropping tile-kernel execution for hsk=72/sinks configs (garbage attention, ERR~1 in test-backend-ops). Bumped the per-kernel dynamic-smem tracking cap from 64 to 4096. Known: oscar2 FA test-backend-ops cases are impossible as written (harness feeds natural-domain data, kernel expects Hadamard-domain), matching PR TheTom#234. Assisted-by: omp
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
20 participants
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.
OSCAR2 — Quantized KV-cache with Hadamard-domain attention
DRAFT — only Qwen3.6-27B validated so far.
Status
--cache-type-k oscar2 --cache-type-v oscar2oscar-rotation/)What OSCAR2 adds
AI disclosure
Usage
See
oscar-rotation/README.mdfor the rotation pipeline.