AVX2: Speed up large batch size prompt processing of IQ models - #27402
Conversation
github-actions
Bot
added
the
ggml
|
Updated with latest numbers for performance, looks quite good across the board. Also noted that I ran the test ops on machines with VNNI for completeness |
ggerganov
left a comment
ggerganov
left a comment
There was a problem hiding this comment.
If I understand correctly, this is not really repacking of the weights - they stay the same, we just add some on-the-fly repacking to the work buffers. If this is correct, then the repack.cpp should not be modified. Most likely you need to reorganize the code similar to llamafile and not modify the repack logic at all.
|
@ggerganov it's not exactly repacking but it is doing similar work to other stuff already in repack.cpp, namely the row interleaving I guess the question is, does repack only apply to load-time interleaving or also transient interleaving? If it makes more sense to move it to avoid calling it something it isn't, I can do that, will need to duplicate some existing work around architecture mapping but it's definitely doable |
Repacking is just for load-time. It uses an extra buffer type and modifies the contents of the weight tensors. Runtime-only optimizations like in this PR should be implemented separately similar to |
|
@ggerganov okay it's all now self contained in the |
|
@ggerganov let me know if you need anything else from me (more explanation, more tests, etc) all sanity checks look good |
|
Please rebase on latest |
809f6da to
9f6857f
Compare
|
Rebased and squashed |
|
Add yourself to the |
|
Okay here's what I ran, let me know if it was right: |
|
I have a few other CPUs I can test on and also can ask people to help test if you need more confirmation, just say the word :) |
|
You will also need Also, add yourself to the CODEOWNERS for the new |
|
Okay I'll try to run that when I get the chance! I did add myself to the CODEOWNERS file, did I do it right? |
|
So I ran and it returned all OK, but it ran on all my cores for about 19 hours... Do I need to also run MUL_MAT_ID? haha.. |
|
@ggerganov if you need me to run the other i'll do it overnight tonight, not sure what the difference is between them but i'm guessing the |
42a8aba to
1d7caa0
Compare
|
also rebased on master again, hence the force-push |
| @@ -1131,7 +1131,7 @@ GGML_TABLE_END() | |||
| #define NGRID_IQ1S 2048 | |||
| #define IQ1S_DELTA 0.125f | |||
| #define IQ1M_DELTA 0.125f | |||
| #if defined(GGML_COMMON_IMPL_C) | |||
| #if defined(GGML_COMMON_IMPL_C) || defined(GGML_COMMON_IMPL_CPP) | |||
There was a problem hiding this comment.
noticed that technically this change removes the iq1s_grid_gpu from any area that has GGML_COMMON_IMPL_CPP, however obviously none of them actually use it currently, just wanted to call attention to it in case it raises any questions about the grid that they have access to switching from this change
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb.
fa27ca6 to
1ed17b4
Compare
|
Moved IQP mul_mat_id test and rebased on latest |
|
@CISC I just woke up so I trusted Claude to investigate so take this with a grain of salt, but it seems to think it's actually a CI ccache issue, not a bug in the code: Fix is allegedly to clear the ccache for those jobs and re run, or add |
Yes, most likely, it's a bit weird, will try deleting caches. |
Seems to have worked: |
|
Great thanks for handling it! |
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
|
Hi, I just coincidentally had the same idea and implemented something similar for k-quants. I put a proposal for merging the 2 approaches into a single code path on my PR here: #27851 @bartowski1182 if you want to take a look. |
Port upstream PR ggml-org#27402 (grid IQ batched GEMM) onto the bmoe/expert-ready-hook line: - new iqp.cpp/iqp.h: decode 8 src0 rows at a time into per-thread int8 panels (block_iqp_x8) and run an integer gemm against all src1 columns, for the grid IQ types (IQ1_S/M, IQ2_XXS/XS/S, IQ3_XXS/S, IQ4_XS/NL) - up to 8-10x faster CPU prefill at large batch sizes, ~2x on CPU MoE layers - ggml-cpu.c: dispatch mul_mat and mul_mat_id to the iqp path after the src1->q8_K barrier (mirrors the upstream placement: the iqp path consumes the q8_K rows from the work buffer); the per-expert mul_mat_id dispatch comes after the expert-ready hook so the streamer's residency block still fires on both paths - graph_plan reserves one scratch panel per thread for both ops - ggml-common.h: widen iq1s_grid table guard to C++ translation units - tests: 10-row-batch IQP coverage on all grid IQ types for MUL_MAT and MUL_MAT_ID Verified: test-backend-ops 883/883 MUL_MAT + MUL_MAT_ID on CPU (x86 AVX2) including the new 10-row batch cases; CUDA0 topk/fusion 416/416 unchanged; bmoe byte-identity gates 13/16 (pre-existing streaming-init failures only).
…gml-org#25294 ggml-org#26003 ggml-org#26414 ggml-org#27402 ggml-org#20596 ggml-org#22671) and the CPU-kernel cherry-picks Co-Authored-By: Claude Fable 5 <noreply@anthropic.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…org#27402) * Batched gemm for grid IQ quants Style updates and a bit more performance Clean up comments Move code around Vectorize IQ panel decode, lower threshold for speedup IQ panel: single-source gather layout, gate bias, vectorize interleave Add ggml_gemm_iqp_8x8_q8_K_p4 kernel, remove gather buffer Move IQ panel code out of repack into iqp.cpp, clean up comments Another comment sweep * Add myself as iqp.* codeownder * Remove ggml_cpu_iqp_scratch_offset and ggml_cpu_iqp_src1_conv_size * Renaming and moving * The other half of renaming and moving * Move macros and ggml_cpu_iqp_mul_mat_id_min_batch definition * Update ggml/src/ggml-cpu/iqp.h Co-authored-by: Georgi Gerganov <ggerganov@gmail.com> * Add iqp_rows work buffer * Revert "Add iqp_rows work buffer" This reverts commit 425542991eee1b01fa3844bf87fc4f205ddbfccb. * Add NUMA fallback * Add 10 row batch tests for IQP coverage on all grid IQ types * Swap assert for return false in support check * Move IQP mul_mat_id test --------- Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
4 participants
Overview
IQ quants are particularly slow on CPU at large batch sizes (what you'd see for imatrix and perplexity)
Part of this comes from the fact that during a 512-token batch, every weight in the model is decoded 512 times from the associated lookup table
This change introduces a GEMM panel,
block_iqp_x8, which is 8 weight rows x 256 columns:Now, instead of decoding each weight 512 times at batch 512, we decode 8 rows at a time into a cache-sized
int8tile and run integer GEMM over it.For dense models, this is a huge increase, up to 10x
For MoE it's more modest, more like 2x, but still quite good
PPL shifts a tiny bit, only ~0.24%, aka margin of error
Design details from Claude
Four design details make it work:
Interleaving 8 rows. qs[sb128 + g32 + row4 + k] holds column sb16 + g*4 + k. So one 32-byte load pulls 4 consecutive columns for each of 8 rows — exactly the operand shape _mm256_dpbusd_epi32 wants against a broadcast int32 of 4 activations. Eight output rows advance per instruction.
Sub-blocks of 16, not 32. iq2_xs/iq2_s take a group of 32 weights' scale from two nibbles, one per half. One scale per 32 couldn't represent them exactly, so 16 is used uniformly for all eight types → one kernel.
Split scale = float × integer (this is commit 4f7e113). Every supported type's sub-block scale factors as (d · 2⁻ᵏ) · small_int — e.g. iq2_xs is d·(0.5+ls)·0.25 = (d/8)·(2ls+1). So dfac holds the float, iscales the integer. The kernel then applies sub-block scales in integer (mullo_epi32), accumulates the whole 256-weight super-block exactly in int32, and touches float exactly once per (row, super-block). That's 16 float roundings → 1. A float64 reference harness measured the result as 12–22% more accurate than upstream vec_dot (the previous float-scale version had been 2.6× less accurate).
The bias trick. VNNI's dpbusd needs unsigned × signed, but both operands here are signed. Instead of fixing that with _mm256_sign_epi8 per operand, activations are fed as y ^ 0x80 (i.e. y + 128), and the resulting spurious 128·Σw term — a per-row, per-super-block constant — is precomputed into bias at decode time and subtracted with one int32 op. Non-VNNI AVX2 builds fall back to the sign trick with bias disabled, because the maddubs int16 accumulator would overflow with unsigned operands.
Additional information
On dense models, this only kicks in when batch >=
328 columns, less than that was found to have lower performance, set withGGML_IQP_MIN_BATCHinggml/src/ggml-cpu/iqp.h(edit: lowered to 8 with a better kernel)for MoE, it instead gates per-expert, and only works when needs to see
168 columns to route through the new panel, otherwise old performance is better, set withGGML_IQP_MIN_BATCH_IDinggml/src/ggml-cpu/iqp.h. (edit: also lowered to 8 with better kernel)Can be disabled with
GGML_NO_IQ_PANEL=1Confirmed with tests that setting
GGML_NO_IQ_PANEL=1returns original performance and PPL completely, as well as when using a lower batch size (-ub 16). At-ub 16, there is still an ~80% speed up in performance on denseBenchmark numbers
I ran PPL against master and this PR to get speed and numbers on
--chunks 50for Qwen3.6-27B and Qwen3.6-35B-A3B on EPYC 9654 using 24 threadsCreated pure
IQ1_S,IQ1_M,IQ2_XXS,IQ2_XS,IQ2_S,IQ3_XXS,IQ3_S, andIQ4_XS. Made pure to make sure each tensor type is fully exercised.Updated figures:
Speed:
Perplexity:
I ran at various chunk lengths for each ubatch to save time, all were run with these settings:
I also tested:
On a Ryzen 9 7950X3D and an Intel Core Ultra 7 358H since they have VNNI support, both gave OK
On the Ryzen 9 I got these numbers:
original speeds/ppl just for posterity
These are the most extremely differences because it's at a big batch size (512), lower batch sizes get smaller increases
Note, since some of these are extremely long running even at only 50 chunks, the performance numbers may vary slightly, but the gains were seen repeatedly.
Requirements