Repository navigation
IQ1_XS, IQ1_XXS, IQ1_XXXS: three quant types below IQ1_S - #61
danielhanchen wants to merge 134 commits into
Conversation
|
End-to-end verification on real Kimi-K3 files is done, so the predicted numbers in the description are now measured ones. Three full builds, 276
Every value matches what the description predicted, sizes to the byte ( Two controls make that meaningful:
wikitext-2, 568 chunks, n_ctx 512, KLD against the BF16 master. |
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 60ea1fa1d4
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| // Below IQ1_S: the same 8-d ternary codebook with a narrower index. Fork-local ids, | ||
| // deliberately placed after every upstream type so a future upstream addition does not | ||
| // collide. See block_iq1_xs / block_iq1_xxs / block_iq1_xxxs in ggml-common.h. | ||
| GGML_TYPE_IQ1_XS = 43, // 1.4375 bpw, 1024-entry grid |
There was a problem hiding this comment.
When this fork rebases onto an upstream release that adds the next ggml_type, upstream will allocate 43 because its current count is 43. gguf_writer::write_tensor_meta serializes this raw enum value without a fork namespace, so files emitted now as IQ1_XS will then be read as the unrelated upstream type or rejected, making the stated collision guarantee false. Obtain stable upstream IDs or add a versioned, namespaced encoding before emitting these formats.
Useful? React with 👍 / 👎.
| 0x010101010101ffff, 0x010101010101ff01, 0x01010101010101ff, 0x0101010101010101, | ||
| GGML_TABLE_END() | ||
|
|
||
| // Generated by scripts/iqxs_gen_tables.py. Do not edit by hand. |
There was a problem hiding this comment.
The repo-wide find and rg search finds neither scripts/iqxs_gen_tables.py nor another source for these generated tables. Without the referenced generator, maintainers cannot reproduce the facility-location selection or safely regenerate and compare the mirrored CPU/GPU tables, leaving thousands of opaque constants that cannot be meaningfully reviewed or maintained. Include the generator and its reproducible inputs or derive both representations from one checked-in source.
AGENTS.md reference: AGENTS.md:L16-L18
Useful? React with 👍 / 👎.
| case LLAMA_FTYPE_MOSTLY_IQ1_XS: name = LLAMA_FTYPE_PREFIX "IQ1_XS - 1.4375 bpw"; break; | ||
| case LLAMA_FTYPE_MOSTLY_IQ1_XXS: name = LLAMA_FTYPE_PREFIX "IQ1_XXS - 1.3125 bpw"; break; | ||
| case LLAMA_FTYPE_MOSTLY_IQ1_XXXS: name = LLAMA_FTYPE_PREFIX "IQ1_XXXS - 1.1875 bpw"; break; |
There was a problem hiding this comment.
When a GGUF omits the optional general.file_type metadata, the loader infers it from type_max, but the switch in llama_model_loader has no cases for any of these three new tensor types. Because their enum values exceed every existing type, such a model takes the default branch and is exposed through llama_model_ftype() and the model description as guessed F32; an --only-copy pass can also persist that incorrect guessed value. Add the corresponding GGML_TYPE_IQ1_* mappings alongside the existing IQ1_S mapping.
Useful? React with 👍 / 👎.
| IQ1_XS = 43 | ||
| IQ1_XXS = 44 | ||
| IQ1_XXXS = 45 |
There was a problem hiding this comment.
Files produced by the new quantizer store general.file_type values 42-44, but gguf.constants.LlamaFileType still ends at MOSTLY_Q2_0 = 41 despite its invariant that all values match llama_ftype. Python consumers therefore cannot represent these file types and LlamaFileType(42), (43), or (44) raises ValueError. Add MOSTLY_IQ1_XS, MOSTLY_IQ1_XXS, and MOSTLY_IQ1_XXXS to that enum as well as adding the tensor-type constants here.
Useful? React with 👍 / 👎.
60ea1fa to
c2974d7
Compare
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: 665dc7b938
ℹ️ About Codex in GitHub
Your team has set up Codex to review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
Codex can also answer questions or update the PR. Try commenting "@codex address that feedback".
| GGML_TYPE_IQ1_XS = 64, // 1.4375 bpw, 1024-entry grid | ||
| GGML_TYPE_IQ1_XXS = 65, // 1.3125 bpw, 512-entry grid | ||
| GGML_TYPE_IQ1_XXXS = 66, // 1.1875 bpw, 256-entry grid |
There was a problem hiding this comment.
On a SYCL build, models containing these types are still accepted for GGML_OP_MUL_MAT because ggml_backend_sycl_device_supports_op() returns true for every type except Q1_0. However, get_row_rounding(), the SYCL conversion dispatch, and MMVQ have no cases for these enums: split-buffer loading reaches the default GGML_ABORT, while single-device inference selects MMVQ based only on ggml_is_quantized() and then aborts on the unsupported type. Add the corresponding SYCL implementations or explicitly prevent these tensors and operations from being placed on SYCL.
Useful? React with 👍 / 👎.
|
Codex usage limits have been reached for code reviews. Please check with the admins of this repo to increase the limits by adding credits. |
|
Merged master (via #80, the upstream sync to Why it needed redoingUpstream's MMQ refactor, ggml-org ggml-org#24127 plus five follow-ups, replaced the tile API this branch was written against. The portThe three types keep IQ1_S's tile geometry exactly:
Blackwell only overrides the FP4 types and falls through to Ampere, so it needs nothing. The config tables are the part that is easy to miss and they are not optional. Without them For the record, the IQ1_XS deviation from IQ1_S is the block layout and nothing else: - constexpr int threads_per_row = MMQ_ITER_K / (4 * QR1_S);
+ constexpr int threads_per_row = MMQ_ITER_K / (4 * QR1_XS);
- const int qh = bxi->qh[kqsx];
+ const int qh = bxi->qh[kqsx];
+ const int nib = (bxi->sc[kqsx/2] >> (4*(kqsx%2))) & 0x0F;
- const int grid = iq1s_grid_gpu[qs[l] | (((qh >> (3*l)) & 0x07) << 8)];
+ const int grid = iq1_xs_grid_gpu[qs[l] | (((qh >> (2*l)) & 0x03) << 8)];
- const float d1q = __half2float(bxi->d) * (((qh >> 11) & 0x0E) + 1);
- const float delta = -1.0f + IQ1S_DELTA - (qh & 0x8000) * (2.0f*IQ1S_DELTA/0x8000);
+ const float d1q = __half2float(bxi->d) * (2*(nib & 0x07) + 1);
+ const float delta = nib & 0x08 ? -1.0f - IQ1S_DELTA : -1.0f + IQ1S_DELTA;10-bit index (2 high bits per group, for the 1024-entry grid), 3-bit scale and sign from the Also in this merge, resolved additively against upstream
The id choice survives the sync untouched. Upstream now uses TestsB200, CUDA 13, sm_100: 11 MUL_MAT cases each for Separately, all four published Kimi-K3 quants were loaded end to end from a fresh clone and a fresh download and answered correctly on 4x B200: |
6090cbc to
36fe8e1
Compare
|
Codex usage limits have been reached for code reviews. Please check with the admins of this repo to increase the limits by adding credits. |
|
Rebased onto master now that #80 has landed. 41 files, +2334/-25, one commit. It had ballooned to 1358 files because I merged the sync branch in here to redo the MMQ port, while this PR's base was still the un-synced master, so GitHub attributed all 718 upstream commits to this PR. That is resolved: the branch is now just the narrow-IQ1 change on top of current master. The source is byte-identical to the tree I tested. Rebuilding the branch changed only 11 MUL_MAT cases each for One note for whoever picks this up next. #80 was squash-merged, so master has the upstream content but not the upstream commits as ancestors: |
Three sub-IQ1_S data types: IQ1_S's algorithm with the codebook the mapping step may choose from cut to a subset of iq1s_grid, and the index stored in fewer bits. IQ1_XS id 64 46 B / 256 1024-entry grid 10-bit index 1.4375 bpw IQ1_XXS id 65 42 B / 256 512-entry grid 9-bit index 1.3125 bpw IQ1_XXXS id 66 38 B / 256 256-entry grid 8-bit index 1.1875 bpw Dequantization is IQ1_S's rule throughout, y[j] = d*(2*ls + 1)*(grid[idx][j] + delta). GGML_TYPE_COUNT is 67; 42..63 are left for upstream, which already uses 42 for Q2_0, so a stock build rejects these files rather than misreading them. CPU, CUDA dequant, MMVQ and MMQ. The MMQ path targets the current ggml_cuda_mmq_util_funcs / ggml_cuda_mmq_config API and registers the types in every per-architecture config table; Blackwell falls through to Ampere and needs none. scripts/iqxs_gen_tables.py regenerates all six codebook tables in ggml-common.h byte-identically from the three committed .npy grids, so the tables are reproducible rather than pasted. Tested on B200, CUDA 13, sm_100. test-backend-ops -o MUL_MAT passes 2/2 backends both default and with GGML_CUDA_FORCE_MMQ=1, 11 cases per type matching IQ1_S. Note the default run never enters the new MMQ kernels, since small n routes to MMVQ.
|
Noting here why this cannot go into the nightly pin set as it stands, since it is a branch-base issue rather than anything wrong with the change.
Opened #91 with the same commit on a Keeping this open for the design review. Numbers from the merged nightly tree are in #91. |
Adds IQ1_XS at 1.4375 bpw, IQ1_XXS at 1.3125 and IQ1_XXXS at 1.1875 to the nightly. Pins #91 rather than #61: same commit content, but on an upstream base, so it merges onto the base tag instead of dragging two months of squash-synced master in behind it. Replayed resolve on b10356 with this set: all four pins merge, tag b10356-mix-1c79028, and the resulting tree hash matches the tree that was built and tested (cdc8f949).
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
|
expect |
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
Adds three sub-IQ1_S quantization types on top of the vendored ggml, for running very large MoE models (122B-A10B class) on machines that cannot hold an IQ1_S copy: 03 unslothai/llama.cpp#61, ggml/ portion only -- IQ1_XS (1.4375 bpw), IQ1_XXS (1.3125), IQ1_XXXS (1.1875). Type enum, grids, ggml/CPU/CUDA kernels + MMQ template instances. 04 metal-iq1-narrow.patch (user-supplied; the PR ships no Metal support) -- mul_mv / mul_mv_id / mul_mm / mul_mm_id / get_rows + dequantize for the three types, i.e. the same op coverage IQ1_S has. Kept on its own branch on purpose: both patches are third-party and expect to land in ggml-org/llama.cpp eventually, at which point this branch is deleted and a plain re-vendor picks the types up instead. The base branch carries only the upstream sync and the two qwencpp-local ggml patches. Not exercised end to end: no IQ1_XS/XXS/XXXS GGUF was available, so the types are verified only by construction -- registered in ggml's type traits, in the CPU/CUDA/Metal dispatch tables, and compiling under CUDA. Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…org#31357) Requested in the ggml-org#31125 review. The cases duplicate coverage that is already in the file and add CI time. Test-only, nothing changes in the kernel, the docs or the CMake option. Assisted-by: Claude
…gml-org#31075) * metal : faster MXFP4 mat-vec kernel Decode the fp4 values with the fp16 bit trick instead of a threadgroup lookup table, and load the packed nibbles as aligned 32-bit words. The kernel no longer uses threadgroup memory. * metal : launch contiguous unary and binary ops with full threadgroups The contiguous unary path and the row-broadcast binary path used one thread per threadgroup. Use up to 256 threads per threadgroup and bound-check the tail. * metal : split GLU rows across threadgroups A GLU row ran on one threadgroup, so single-row decode used a single threadgroup for the whole op. Split each row over threadgroups of up to 256 threads.
* opencl: add iq2_s AOS mv * opencl: add iq2_s SOA mv and tiled mm * opencl: add iq2_s noshuffle gemm/gemv for adreno * opencl: add iq2_s dp4a for adreno * opencl: optimize iq2_s adreno decode and narrow batches Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com> * opencl: add iq2_xs AOS mv * opencl: add iq2_xs SOA mv and tiled mm * opencl: add iq2_xs noshuffle gemm/gemv for adreno * opencl: add iq2_xs dp4a for adreno * opencl: optimize iq2_xs adreno decode and narrow batches Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com> * opencl: add iq2_xxs AOS mv * opencl: add iq2_xxs SOA mv and tiled mm * opencl: add iq2_xxs noshuffle gemm/gemv for adreno * opencl: add iq2_xxs dp4a for adreno * opencl: optimize iq2_xxs adreno decode and narrow batches Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com> --------- Co-authored-by: Hongqiang Wang <wangh@qti.qualcomm.com>
* common: select A100 cores for SpaceMiT default threads # Conflicts: # tests/CMakeLists.txt * ggml: export SpaceMiT build marker * common: scope SpacemiT thread detection macro to llama-common * common: detect A100 cores in all Linux RISC-V builds * tests: link barrier test against llama-common * tests: drop SpacemiT default thread test * common: follow platform check order
* gguf-py : add Limite arch Assisted-by: Claude Opus 5.5 Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org> * convert : add Limite support Assisted-by: Claude Opus 5.5 Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org> * llama : add Limite support Assisted-by: Claude Opus 5.5 Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org> --------- Signed-off-by: Giuseppe Scrivano <giuseppe@scrivano.org>
* HIP/CUDA: enable sparse attn on CDNA * Use __ballot for __ballot_sync on old ROCm versions * Ignore new large head fattn mma variant
Assisted-by: Codex
Ask the browser to keep IndexedDB and localStorage data under storage pressure, so conversations and settings are less likely to be evicted. Assisted-by: Claude (Opus 5.5 Medium)
…checkbox colors (ggml-org#31303) MCP server groups in the settings Tools tab now get a group checkbox like the built-in, server and browser groups, toggling the server-scoped tools key. Indeterminate checkboxes now use the checked state colors. Assisted-by: Claude (Opus 5.5 Medium)
…gml-org#31307) * ui: reveal the collapsed sidebar from its logo and toggle it there The collapsed sidebar now shows only the logo; hovering it reveals the icon strip until the pointer leaves. The expanded sidebar drops its top-right collapse button and the logo collapses it on hover instead. Assisted-by: Claude (Opus 5.5 Medium) * ui: fade the peeked sidebar in at full width and keep the toggle slide Hovering the logo of the collapsed sidebar now fades in the fully open sidebar, floating over the page at its full width, instead of resizing it. Opening or closing with the toggle still resizes and slides: the sidebar docks to strip width while it slides, then floats out of sight, and it stays hidden under the pointer after a toggle close. The unused icon-only strip is removed from the sidebar actions. Assisted-by: Claude (Opus 5.5 Medium) * ui: stack the sidebar above the chat tabs on desktop The peeked sidebar floats over the page, but the chat tab bar shared its z-index and came later in the page, so the tabs were drawn on top of it. Assisted-by: Claude (Opus 5.5 Medium)
ChatFormActions gated the mic and stop buttons on a `canSubmit` prop that ChatForm never passed, so it was always false and typing text did not swap them for the submit button. Use `canSend` instead and also hide the skip-reasoning button while there is content to send. Assisted-by: Claude (Opus 5.5 Medium)
) * ui: land at the bottom of a conversation every time it opens Opening a long conversation often stopped short of its end: - the press grace counted from page load (pressedAt = 0), so a chat that landed within 400 ms of a refresh did not follow the rows realizing after the jump - the click that opened a chat (sidebar row, tab) started the same grace, so navigating by click had the same problem - any scroll event away from the bottom stopped following, including the ones the browser fires itself while the chat settles - the chat screen outlives its routes, so going back to the last landed conversation through the new-chat screen skipped the landing entirely scrollToBottom() now clears the grace, only an upward scroll away from the bottom stops following, and every navigation lands anew. Navigations into a chat also disable SvelteKit's scroll reset, which moved the previous conversation to its top for a few frames. Where overflow-anchor is supported, a 1px pin after the messages is the only scroll anchor while following, so the browser keeps the bottom in place itself; WebKit keeps the resize observer. Assisted-by: Claude (Opus 5.5 Medium) * ui: let nested agentic content scrollers chain to the chat Reasoning, shell output and code blocks used overscroll-behavior: contain, so the wheel stopped dead at their edges instead of scrolling the conversation. Assisted-by: Claude (Opus 5.5 Medium) * ui: stop following the stream as soon as the reader scrolls up A trackpad or smooth wheel scrolls in steps smaller than the at-bottom threshold, so each step up was re-stuck and snapped back on the next growth while a response streamed. Any upward move that is not content shrinking now releases at once. The attach effect also read the hook's own state, so every flip re-attached and re-measured, which undid such a release. Assisted-by: Claude (Opus 5.5 Medium)
…gml-org#31350) Assisted-by: Claude Co-authored-by: Zeshawn Zahid <zzshawnn@gmail.com>
* llama: extend DSA and sdpark to new model * fix ci
* skills: specify rule about ggml_build_forward_expand * Update skills/code-review/SKILL.md Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* WIP OCP FP8 E4M3 support * Fix python type-checker * Materialize input_scales also for FP8 when converting * Fix conversion of https://huggingface.co/RedHatAI/Qwen3.8-27B-NVFP4 1. Per row FP8 wrongly fell to NVFP4 conversion, as NVFP4 checked on dims alone. Also check on dtype for FP8 2. Need to reshape FP8 QKV projections scales in the same way that weights are reshaped * Allow remaining FP8 tensors to be dequantized during conversion * Allow tensor-strategy for FP8-only compressed-tensor ckpts * Fix and test gguf-py behavior for FP8_E4M3 * Disallow QKV-fusion for FP8 temporarily * Fix conversion of fused qkvz & FP8 for Qwen3-Next * Avoid fusing FP8 for --fuse-gate-up-exps as we may have different ws * Emit warnings when not fusing FP8 during conversion * More FP8 conversion fixes * Add supported for strided copies to/from FP8 in CPU * remove unneeded scratch buffer from FP8 cpu cpy/dup path * add explanatory comments and TODOs * Add dedicated encoder and decoder tests for FP8 * Whitelist ggml_op_get_rows in metal instead of blacklisting required to fix failing CI, and white-listing should be more future-proof than adding FP8 to the black-list https://github.com/ggml-org/llama.cpp/actions/runs/38433916498/job/115422880633?pr=27512 * Fix more conversion bugs * Defer weight materialization during conversion * Blacklist fp8 in metal backend instead of white-list * Adjust indents * Fix lint + execute tests * Add fast CPU kernel for F8_E4M3 * Conversion fixes * Fix remaining gaps * Enable tied weights with shared scales * Make mandatory-F32 rules apply to FP8 during conversion * Fix python type checks https://github.com/ggml-org/llama.cpp/actions/runs/39765670052/job/119780379595?pr=27512 * Fix kaleidai unit-test failures https://github.com/ggml-org/llama.cpp/actions/runs/39765670125/job/119780383150 * Relax numeric precision constraints for pre-check * Leave fast-FP8 CPU on RISC-V for follow-up We already exclude arm from the fast-path * Allow output scale tensors to be located in any gguf shard * Move test_conversion_fp8.py test execution to behavior-check.yaml * Apply batched suggestions from code review Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> * Move mixed compressed-tensor logic into dequant_fp8 + add unit-tests for it * Install all requirements also on Windows for behavior-check Fixes https://github.com/ggml-org/llama.cpp/actions/runs/40535852590/job/122300758223?pr=27512 * Unmap lazy tensors before clean-up in Windows Should fix https://github.com/ggml-org/llama.cpp/actions/runs/40539797824/job/122314782459?pr=27512 * LF line endings https://github.com/ggml-org/llama.cpp/actions/runs/40539798168/job/122314791118?pr=27512 * Extract FP8 tests from llama-arch to its own test-fixture See ggml-org#27512 (comment) * Update gguf-py/gguf/scripts/gguf_convert_endian.py Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co> --------- Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@huggingface.co>
gfx117X is more like RDNA4, with 128b WMMA. But they're not identical. Adjust the defines accordingly to have a new RDNA4_M and use it where it makes sense. Assisted-by: Codex Assisted-by: GitHub Copilot Co-authored-by: Mario Limonciello <mario.limonciello@amd.com>
* ggml : add generic batched tensor copy API
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* llama : use batched tensor copy API in context io
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml : add batched tensor copy API
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* llama : batch device io tensor copies
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml : add CUDA support for batched tensor APIs
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml : remove async tensor_n APIs and add null buffer hooks
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml : fix batched tensor 2D detection
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* ggml : group batched tensor operations by buffer type
Assisted-by: pi:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* cont : minor indentantion
* cont : const correctness
* tests : add mock buffer type tests for batched tensor API
Restructure test-backend-tensor to run a unified test set on each backend buffer type and on an ad-hoc mock buffer type. The mock set exercises the dispatch logic in ggml_backend_tensor_{set,get,copy}_n: hook invocation, per-tensor fallback, cpy_tensor_n all-or-nothing, and buft grouping.
Assisted-by: dsh:llama.cpp/ggml-org/DeepSeek-V4-Flash-Vision-Exp-GGUF (Mac Studio, M2 Ultra, 192GB)
* cont : add TODO
* opencl: add q2_0 AOS mv Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: add q2_0 SOA mv and tiled mm Co-authored-by: Shaofei Qi <shaoqi@qti.qualcomm.com> Assisted-by: Codex * opencl: add q2_0 noshuffle gemm/gemv for adreno Assisted-by: Codex * opencl: add q2_0 dp4a for adreno * opencl: optimize q2_0 adreno decode and narrow batches --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
…org#31368) * memory: refuse rollbacks past the snapshots of the last ubatch The recurrent memory and the DSV4 cache accepted any rollback up to n_rs_seq, whether or not the last ubatch wrote that snapshot. Each now records the deepest rollback its last ubatch guarantees and seq_rm refuses anything past it: - recurrent: n - 1 for a ubatch of n <= n_rs_seq + 1 tokens, 0 for a longer one, which only guarantees the final state - DSV4: min(n, n_rs_seq), the snapshots are rebuilt from the state before the ubatch The dead rollback block of the rollback test becomes the assertion it asked for, and rollback_bound checks the exact bound of each memory after ubatches of 1, 2, n_rs_seq + 1 and n_rs_seq + 2 tokens. * tests: address review from ggerganov Detect architectures with llama_arch_from_string in comp_padding and arch_reserves_single_seq, as in rollback_bound.
* kv-cache : keep rotation constants on the GPU Assisted-by: Codex * tests : verify resident KV rotation buffers Assisted-by: Codex * kv-cache : document remaining multi-device rotation work Assisted-by: Codex * kv-cache : support per-layer rotation buffers Assisted-by: Codex * kv-cache : fix hybrid SWA rotation inputs and mixed placement check Assisted-by: Codex * models : use per-layer rotation in mistral-large4 and dflash * tests: drop the rotation buffer tests --------- Co-authored-by: Pascal <admin@serveurperso.com>
…31394) The duplicated output looks up token_embd in its buffer context, so token_embd must exist first. With the old order a load where input and output share a buffer type (CPU-only, -ngl 0) held the table twice: gemma-4-E2B-it Q4_K_M, no mmap, CPU model buffer 3163.73 -> 2947.73 MiB.
This commit extracts all apple jobs except for macos-latest-arm64 into a new daily workflow. The motivation for this is to reducd the load on the macos runners.
* cmake : embed an app manifest in all Windows exes Signed-off-by: Adrien Gallouët <angt@huggingface.co> * Remove LLAMA_WIN_APP_CMAKE Signed-off-by: Adrien Gallouët <angt@huggingface.co> --------- Signed-off-by: Adrien Gallouët <angt@huggingface.co>
Adds three quantization types below IQ1_S:
IQ1_XSat 1.4375 bpw,IQ1_XXSat 1.3125 andIQ1_XXXSat 1.1875. IQ1_S at 1.5625 bpw is currently the floor of both llama.cpp and ik_llama.cpp, so there is nothing lower to port and this had to be built.Targets
master. These are plain ggml types with no architecture coupling: they touch the quant tables, the CPU and CUDA kernels and thellama-quantizeftype list, and nothing model-specific. They work for any model, exactly like IQ1_S does. The Kimi-K3 numbers below are just the model I measured them on, because it is the one large enough for the size question to matter.Design
IQ1_S spends 1.3750 of its 1.5625 bpw on an 11-bit index into a 2048-entry ternary grid, which is 88% of the whole budget. That index is the only field worth shrinking, so these three keep every other field with its exact IQ1_S meaning and narrow the index to 10, 9 and 8 bits.
Per 32-element sub-block each type stores four index-high fields, a 3-bit scale and a delta sign. That is 12, 8 and 4 bits respectively, which packs byte-exactly:
IQ1_XXS packs perfectly: its single index-high bit per group leaves exactly room for the scale and sign in the same byte.
Dequantisation is
y[j] = d*(2*ls+1)*(grid[j] + delta)throughout, identical todequantize_row_iq1_s. Only the index reassembly differs.Codebooks
Each grid is a subset of ggml's own
iq1s_grid, selected by imatrix-weighted greedy facility-location over real expert weights. Nothing is trained from scratch.That matters for a specific reason:
ggml_vec_dot_iq1_s_q8_Kreaches the grid through_mm256_maddubs_epi16, an int8 dot product, so a shippable codebook has to be integer-valued. Reusing existing entries inherits that property instead of asserting it. A float codebook looked about 10% better in isolation and came out 18.7 to 20.7% worse once constrained to integers, so this is not a hypothetical concern.The generator and its three grid index files are committed as
scripts/iqxs_gen_tables.pyandscripts/iq1_narrow_grids/, so all six tables inggml-common.hregenerate byte-identically from what is in the tree rather than being 1792 unauditable constants.The neighbour tables are built by ggml's own
iq2xs_init_implover the restricted grid rather than reimplemented, withnwant = 3matching IQ1_S.iq2_data[]widens from 4 slots to 7.d *= 1.125fin the encoder is left exactly as it is.Type ids
Upstream ggml-org master already uses 42 for
Q2_0, 28 forGGML_FTYPE_MOSTLY_Q2_0and 41 forLLAMA_FTYPE_MOSTLY_Q2_0. Sitting on those would mean a stock build reads a file written here as Q2_0, which is block 32 at 18 B/block against 46/42/38 B at block 256, so it would be a wrong-size reinterpretation rather than a clean unknown-type rejection.So the fork-local types start clear of upstream and leave room for it to grow:
GGML_TYPE_IQ1_XS/XXS/XXXSGGML_TYPE_COUNT67)GGML_FTYPE_MOSTLY_IQ1_XS/XXS/XXXSLLAMA_FTYPE_MOSTLY_IQ1_XS/XXS/XXXSOne consequence worth stating: the gap leaves holes in
type_traits[GGML_TYPE_COUNT], sogguf.cpp'stype >= GGML_TYPE_COUNTbound check no longer rejects every invalid id on its own.Scope
CPU scalar and CUDA. Non-x86 routes to scalar through the existing
arch-fallback.h.get_rowsis not implemented on CUDA, which matches master, where no k-quant has it. MMQ mirrors IQ1_S rather than routing through dequantise.No SIMD kernel, on any architecture. IQ1_S has hand-written AVX2, AVX-512, NEON, POWER and s390x
vec_dotkernels; these three get none and fall through to scalar everywhere. Measured on an AVX-512 box, 20000 dots of 8192 weights, L1-resident: IQ1_S 0.012 s against IQ1_XS 0.078 s, IQ1_XXS 0.105 s and IQ1_XXXS 0.097 s, i.e. 6.5x to 8.8x slower per weight. Since the point of these types is fitting a large model into less memory on CPU, that is a real cost and worth knowing before choosing one. The x86 IQ1_S kernel adapts almost verbatim (only the index reassembly and scale/sign extraction differ), so this is a gap rather than a blocker.Verification
All numbers below are from this branch on
master, not carried over from an earlier base.dequantize(quantize(x))vs reference, per typeexact_frac = 1.000000exact_frac = 1.000000, max_abs_diff 0test-backend-ops -o MUL_MAT -b CUDA0GGML_CUDA_FORCE_MMQ=1test-backend-ops -o MUL_MAT_ID -b CUDA0scripts/iqxs_gen_tables.pyExactness is
memcmp, not RMSE. The CPU gate compares the decoded grid index, scale code, delta sign and the fp16dbitwise, over seven synthetic distributions plus three real Kimi-K3 tensors with a real imatrix. Bit equality is reachable on the CUDA side becauseIQ1S_DELTAis 0.125, so bothq[j] + (-1 +/- delta)andgrid[j] +/- deltaare exact in float.Two notes on what the shipped tests do and do not cover, since both are easy to over-read:
nroutes to MMVQ, so the default run never exercises the newload_tiles_iq1_x*skernels.test-quantize-fnsreachesggml_quantize_initfor the three types and stops there, because every assertion in it is behindif (from_float && to_float)and these setfrom_float = NULL. That is also true of IQ1_S on master, so it is not a defect, but the pass carries no information and the round-trip and sweep rows above are what actually cover the formats.test-quant-type-selectionis unchanged at 9/11 models; the two failures are pre-existing snapshot drift onQwen3.5-27Bin ftypes this PR does not touch, confirmed identical with the branch's own change reverted.The regression row matters as much as the new ones: widening the shared
iq2_dataandiq2xs_init_implpaths is exactly the kind of change that can silently perturb neighbouring types.Quantization error is monotone in grid width, 0.5410 / 0.5446 / 0.5589 / 0.5793 for S / XS / XXS / XXXS.
Quality
Measured on Kimi-K3, wikitext-2, 568 chunks, n_ctx 512, KLD against the BF16 master, varying only the 276
ffn_*_expstensors:These have been built as real files in the new formats at exactly those sizes, 14 shards each, and re-measured. They reproduce the numbers above, which were originally obtained by emitting the same restricted-grid weights while still storing an 11-bit index. Agreement between the two is a round-trip check on the pack and unpack code as much as it is a quality result.
Be aware of what the quality-per-byte says: below IQ1_S is the start of a cliff, not a continuation of the curve. Fitting a log-linear frontier through the two measured neighbours predicts noticeably better than what these three actually deliver, and the miss grows as the grid narrows. IQ1_XS costs about 13.7% perplexity for 39.6 GiB on this model. These are useful when the size target is hard, not a free lunch.
Imatrix
All three register in
tensor_requires_imatrix, so a missing imatrix is reported as a clean error rather than reaching theGGML_ASSERT(quant_weights)in the encoder. There is no non-imatrix path for sub-2-bit IQ types anywhere, because the lattice search is defined over per-column importance.One pre-existing caveat inherited from upstream, not introduced here:
tensor_requires_imatrixreturns false for token-embedding and output tensors before it looks at the type, so selecting one of these types for those tensors specifically, via--output-tensor-typeor--token-embedding-typewith no imatrix, reaches the assert instead of the clean error. This reproduces identically withIQ1_Son stock master (--output-tensor-type iq1_s ... Q8_0aborts with rc 134), so fixing it belongs in its own change.Known gaps
vec_doton any architecture, as measured above.gguf-py/gguf/quants.pyhas no Python dequantizer, sogguf.quants.dequantize()raisesNotImplementedErrorfor these types. Same as this fork'sQ1_0.