Add oneliner for batch quantization#17
Conversation
| which can be scripted like this if you are lazy (for 65B model): | ||
|
|
||
| ```bash | ||
| for i in models/65B/ggml-model-f16.bin*;do quantized=`echo "$i" | sed -e 's/f16/q4_0/'`; ./quantize "$i" "$quantized" 2 ;done |
There was a problem hiding this comment.
Sed is not necessary, bash, zsh and other modern shells can perform pattern replacement of a variable:
| for i in models/65B/ggml-model-f16.bin*;do quantized=`echo "$i" | sed -e 's/f16/q4_0/'`; ./quantize "$i" "$quantized" 2 ;done | |
| for i in models/65B/ggml-model-f16.bin* ; do ./quantize "$i" "${i/f16/q4_0}" 2 ;done |
There was a problem hiding this comment.
This will generate 'models/65B/ggml-model-q4_0/.bin.2' such paths and will fail with errors, the right command (in bash) should be for i in models/65B/ggml-model-f16.bin* ; do ./quantize "$i" "${i/f16/q4_0}" 2 ;done
There was a problem hiding this comment.
@Player-205 right, updated the suggestion above, thanks
|
Lets put this in a Should be much easier to follow |
|
Note that if the disk space is limited, it is still useful to quantize each file separately so that we could delete each intermediate file in between. |
|
Good point, should have a second parameter for "keep f16" which is on by default |
|
Superseded by #92 |
improve docs and example
QVAC-5545: Use char instead uint8_t for streams
Fix - unsetenv()
* Basic JIT compilation for mul_mat, get_rows, and scale (#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (#20173) * K quant speedup (#20) * Basic JIT compilation for mul_mat, get_rows, and scale (#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
) * ggml: backend-agnostic tensor parallelism * support for GPT-OSS, Qwen 3 MoE * partial Vulkan fix * add support for 4/8 GPUs * unconditional peer access * re-use buffers + ggml contexts * fix output pattern * NCCL support * GGML: HIP: add RCCL support * Remove shfl and AllReduce from backend interface * move allocation workaround out of ggml-alloc.c * 2d tensor set/get support * Fix the seg fault without NCCL * Apply suggestion from JohannesGaessler * support for tensor dims % n_devs != 0 * fix view_offs scaling * arbitrary num. of GPUs/tensor split * fix compilation * better granularity estimate * Support device-specific host buffer types if all underlying backends expose the same type. This allows using pinned memory instead of pageable memory for CUDA. Fix compilation errors. * partial Qwen 3 Next support * Fix qwen3 30b (ggml-org#8) * Fix crash with Qwen-30B-A3B Q4_0 Qwen-30B-A3B Q4_0 has an intermediate dimension of 768. Using a granularity of 256 forces an uneven split between GPUs, which is not supported by the current implementation. * Decide block size based on tensor quantization type * Fix crashes due to KV cache serialization (ggml-org#9) KV cache serialization requires non-zero offsets on the tensor. Add support in the meta backend to set/get a tensor with a non-zero offset. * metal : fix build (ggml-org#7) * static memory allocations, fix usage count * fix tensor granularity * more even memory distribution * use BF16 for allreduce * rebase fixup * better error message for unsupported architectures * Fix device mismatch during scatter of allReduce. (ggml-org#11) There is a mismatch between the dst buffer device and the backend device, causing the use of sync copies * Enable the previous allreduce implementation. It is better in both perf and stability (ggml-org#12) * delay AllReduce for Moe for less I/O * build : clean-up compile warnings * backend : move most of the meta backend API to ggml-backend-impl.h * cont : hide unused public API in the implementation * llama : use llama_device + remove ggml_backend_dev_is_meta() * ggml-backend : remove unused alloc include * minor : remove regex include * ggml : introduce ggml-ext.h for staging new APIs * rebase fixup * fix tests * llama : more robust logic for determining Meta devices (ggml-org#16) * llama : more robust logic for determining Meta devices * cont : fix devs size check Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * cont : fix log type Co-authored-by: Johannes Gäßler <johannesg@5d6.de> --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * disable roundtrip for meta backend * fix arch selection * Qwen 3.5 support * fix Gemma 4 MoE * fix OpenVino, SYCL * fix test-llama-archs for CPU-only builds * Fix Qwen 3.5 MoE * disable meta backend tests for WebGPU * tests : filter CPU-based devices from the Meta backend tests (ggml-org#17) * meta : formatting, naming, indentation (ggml-org#18) * formatting : llama-model.cpp * formatting : ggml-ext.h * formatting : ggml-backend-meta.cpp * meta : add TODO * add documentation * better error messages * fix GPT-OSS --------- Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…gml-org#17 Splits channels into outlier (higher bits) and normal (lower bits): - 2.5-bit: 64/128 at 3b + 64/128 at 2b - 3.5-bit: 64/128 at 4b + 64/128 at 3b This matches Prince Canuma's MLX approach for achieving 4.9× compression. 8 tests, 91% coverage on outlier.py. Co-Authored-By: tturney@psyguard.ai Co-Authored-By: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
) * ggml: backend-agnostic tensor parallelism * support for GPT-OSS, Qwen 3 MoE * partial Vulkan fix * add support for 4/8 GPUs * unconditional peer access * re-use buffers + ggml contexts * fix output pattern * NCCL support * GGML: HIP: add RCCL support * Remove shfl and AllReduce from backend interface * move allocation workaround out of ggml-alloc.c * 2d tensor set/get support * Fix the seg fault without NCCL * Apply suggestion from JohannesGaessler * support for tensor dims % n_devs != 0 * fix view_offs scaling * arbitrary num. of GPUs/tensor split * fix compilation * better granularity estimate * Support device-specific host buffer types if all underlying backends expose the same type. This allows using pinned memory instead of pageable memory for CUDA. Fix compilation errors. * partial Qwen 3 Next support * Fix qwen3 30b (ggml-org#8) * Fix crash with Qwen-30B-A3B Q4_0 Qwen-30B-A3B Q4_0 has an intermediate dimension of 768. Using a granularity of 256 forces an uneven split between GPUs, which is not supported by the current implementation. * Decide block size based on tensor quantization type * Fix crashes due to KV cache serialization (ggml-org#9) KV cache serialization requires non-zero offsets on the tensor. Add support in the meta backend to set/get a tensor with a non-zero offset. * metal : fix build (ruixiang63#7) * static memory allocations, fix usage count * fix tensor granularity * more even memory distribution * use BF16 for allreduce * rebase fixup * better error message for unsupported architectures * Fix device mismatch during scatter of allReduce. (ggml-org#11) There is a mismatch between the dst buffer device and the backend device, causing the use of sync copies * Enable the previous allreduce implementation. It is better in both perf and stability (ggml-org#12) * delay AllReduce for Moe for less I/O * build : clean-up compile warnings * backend : move most of the meta backend API to ggml-backend-impl.h * cont : hide unused public API in the implementation * llama : use llama_device + remove ggml_backend_dev_is_meta() * ggml-backend : remove unused alloc include * minor : remove regex include * ggml : introduce ggml-ext.h for staging new APIs * rebase fixup * fix tests * llama : more robust logic for determining Meta devices (ggml-org#16) * llama : more robust logic for determining Meta devices * cont : fix devs size check Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * cont : fix log type Co-authored-by: Johannes Gäßler <johannesg@5d6.de> --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * disable roundtrip for meta backend * fix arch selection * Qwen 3.5 support * fix Gemma 4 MoE * fix OpenVino, SYCL * fix test-llama-archs for CPU-only builds * Fix Qwen 3.5 MoE * disable meta backend tests for WebGPU * tests : filter CPU-based devices from the Meta backend tests (ggml-org#17) * meta : formatting, naming, indentation (ggml-org#18) * formatting : llama-model.cpp * formatting : ggml-ext.h * formatting : ggml-backend-meta.cpp * meta : add TODO * add documentation * better error messages * fix GPT-OSS --------- Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
) * ggml: backend-agnostic tensor parallelism * support for GPT-OSS, Qwen 3 MoE * partial Vulkan fix * add support for 4/8 GPUs * unconditional peer access * re-use buffers + ggml contexts * fix output pattern * NCCL support * GGML: HIP: add RCCL support * Remove shfl and AllReduce from backend interface * move allocation workaround out of ggml-alloc.c * 2d tensor set/get support * Fix the seg fault without NCCL * Apply suggestion from JohannesGaessler * support for tensor dims % n_devs != 0 * fix view_offs scaling * arbitrary num. of GPUs/tensor split * fix compilation * better granularity estimate * Support device-specific host buffer types if all underlying backends expose the same type. This allows using pinned memory instead of pageable memory for CUDA. Fix compilation errors. * partial Qwen 3 Next support * Fix qwen3 30b (ggml-org#8) * Fix crash with Qwen-30B-A3B Q4_0 Qwen-30B-A3B Q4_0 has an intermediate dimension of 768. Using a granularity of 256 forces an uneven split between GPUs, which is not supported by the current implementation. * Decide block size based on tensor quantization type * Fix crashes due to KV cache serialization (ggml-org#9) KV cache serialization requires non-zero offsets on the tensor. Add support in the meta backend to set/get a tensor with a non-zero offset. * metal : fix build (ruixiang63#7) * static memory allocations, fix usage count * fix tensor granularity * more even memory distribution * use BF16 for allreduce * rebase fixup * better error message for unsupported architectures * Fix device mismatch during scatter of allReduce. (ggml-org#11) There is a mismatch between the dst buffer device and the backend device, causing the use of sync copies * Enable the previous allreduce implementation. It is better in both perf and stability (ggml-org#12) * delay AllReduce for Moe for less I/O * build : clean-up compile warnings * backend : move most of the meta backend API to ggml-backend-impl.h * cont : hide unused public API in the implementation * llama : use llama_device + remove ggml_backend_dev_is_meta() * ggml-backend : remove unused alloc include * minor : remove regex include * ggml : introduce ggml-ext.h for staging new APIs * rebase fixup * fix tests * llama : more robust logic for determining Meta devices (ggml-org#16) * llama : more robust logic for determining Meta devices * cont : fix devs size check Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * cont : fix log type Co-authored-by: Johannes Gäßler <johannesg@5d6.de> --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * disable roundtrip for meta backend * fix arch selection * Qwen 3.5 support * fix Gemma 4 MoE * fix OpenVino, SYCL * fix test-llama-archs for CPU-only builds * Fix Qwen 3.5 MoE * disable meta backend tests for WebGPU * tests : filter CPU-based devices from the Meta backend tests (ggml-org#17) * meta : formatting, naming, indentation (ggml-org#18) * formatting : llama-model.cpp * formatting : ggml-ext.h * formatting : ggml-backend-meta.cpp * meta : add TODO * add documentation * better error messages * fix GPT-OSS --------- Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
…ml-org#18) The turbo KV cache code uses CUDA __device__ constant memory for runtime-initialized values (InnerQ channel scales, calibration accumulators, Q² statistics, extraction buffers). These symbols are per-device in CUDA — each GPU gets its own copy. All initialization functions (turbo_innerq_init, turbo_innerq_init_fattn, turbo_innerq_start_calibration, turbo_innerq_finalize_calibration, turbo_innerq_update_fattn_scales, turbo_q_calibrate_init, turbo_q_calibrate_finalize, turbo_extract_init, turbo_extract_check_done) used cudaMemcpyToSymbol without specifying which device to target, so they only wrote to whichever GPU happened to be active. On multi-GPU setups with --tensor-split, the other GPU(s) would read uninitialized constant memory, causing garbage output or illegal memory access crashes. Standard cache types like q8_0 are unaffected because they use either compile-time constants or read directly from global memory. Fix: save/restore the current device and loop over all devices for every cudaMemcpyToSymbol call in the turbo initialization and calibration paths. Fixes ggml-org#17 Relates to ggml-org#12 Co-authored-by: Claude Opus 4.6 (1M context) <noreply@anthropic.com>
… (setup for ggml-org#18) Manual port of buun's `692cffde1` ("experiment: S3 TCQ codebook from __constant__ to __shared__ memory"). Moves the TURBOQ2_TCQ / TURBOQ3_TCQ codebooks out of __constant__ memory into a cooperatively-loaded __shared__ array in `flash_attn_ext_vec`. __constant__ serializes when threads in a warp hit different 32B cache lines during random codebook lookups; __shared__ gives full 32-bank parallel access. Standalone this commit is near-no-op on perf — its value is setting up ggml-org#18. Yggdrasil divergences from buun's diff: - Naming: buun `turbo{2,3}_tcq` → yggdrasil `turboq{2,3}_tcq`; codebook symbol `d_turbo{2,3}_tcq_codebook_fattn` → `d_turboq{2,3}_tcq_codebook` (no `_fattn` suffix because yggdrasil shares one codebook across encode + decode paths). - V dispatch: yggdrasil's `flash_attn_ext_vec` inlines TURBOQ{2,3}_TCQ V dequant directly inside the V-aggregation loop rather than routing through `dequantize_V_*` (per Phase 3a ggml-org#11 — `727799470`). So the V-side payoff is a one-line `d_turboq{2,3}_tcq_codebook` → `tcq_smem_codebook` swap inside the inline branches at fattn-vec.cuh, instead of buun's `if constexpr` + `_cb<half/float, ne>` wrapper calls. The K-side matches buun: split helper into `_cb` core + thin `__constant__` wrapper, dispatch the `_cb` form for TCQ types and fall through to `vec_dot_KQ` otherwise. - V-helper split (`dequantize_V_turboq{2,3}_tcq_cb` + wrapper) lifted for surface parity even though the VEC kernel bypasses these helpers — they are still reachable via `get_dequantize_V` from non-VEC dispatch paths and keep using `__constant__` there. SMEM footprint: 512 floats (2 KiB) for TCQ3, 256 floats (1 KiB) for TCQ2 — declared per-block via the conditional `tcq_smem_cb_size` constexpr and zeroed to size 1 when neither side is TCQ (no SMEM cost for unrelated kernels). Well under HIP's 64 KiB cap and below the existing turbo_lut alloc for D=256. Composition check: touches only the TCQ helper definitions in fattn-common.cuh and three sites in fattn-vec.cuh (top-of-kernel codebook load + K dispatch + V inline). Disjoint from s55 warp-shuffle cleanup (set-rows.cu), s56 measurement plumbing (set-rows.cu), s57 routing fix (fattn.cu), s58 backtrace + double-buffer (set-rows.cu). Validation: see follow-up smoke + PPL parity cells (session 60). recon/09 INCLUDE-table row ggml-org#17.
…ation
Manual port of buun's `3c4c44e5c` ("simplify: consolidate codebook load loop,
remove hardcoded sizes"). Merges the dual `is_tcq3` / `is_tcq2` codebook-load
branches in `flash_attn_ext_vec` into a single loop driven by `smem_cb_size > 0`,
with the source codebook selected via a constexpr-resolvable ternary on the
hot pointer. Hardcoded `512` and `256` loop bounds collapse into the shared
`tcq_smem_cb_size` constant.
Pure refactor — no algorithmic change. Both branches were behaviorally
identical (cooperative load + __syncthreads); consolidation just dedups the
loop body. Expected PPL: bit-identical to s60 anchor. Expected TPS: zero
delta (compiler should produce the same SASS via constexpr resolution).
Note on starter framing: the s61 starter prompt described buun ggml-org#17+ggml-org#18 as a
"setup + use" pair where ggml-org#18 was the perf-payoff. That model is incorrect.
Buun ggml-org#17 (692cffd, s60) WAS the full __constant__→__shared__ functional
change — it already moved the codebook AND swapped all VEC call sites. Buun
ggml-org#18 (3c4c44e, this commit) is just a code-shape cleanup of the load loop
that landed on the same merge bubble (b2cfd17). The decode-side perf
benefit of buun's design lives entirely in s60.
Source: buun 3c4c44e (merged in b2cfd17).
Mainline forward-sync per [[yggdrasil-mainline-forward-sync]].
… (setup for ggml-org#18) Manual port of buun's `692cffde1` ("experiment: S3 TCQ codebook from __constant__ to __shared__ memory"). Moves the TURBOQ2_TCQ / TURBOQ3_TCQ codebooks out of __constant__ memory into a cooperatively-loaded __shared__ array in `flash_attn_ext_vec`. __constant__ serializes when threads in a warp hit different 32B cache lines during random codebook lookups; __shared__ gives full 32-bank parallel access. Standalone this commit is near-no-op on perf — its value is setting up ggml-org#18. Yggdrasil divergences from buun's diff: - Naming: buun `turbo{2,3}_tcq` → yggdrasil `turboq{2,3}_tcq`; codebook symbol `d_turbo{2,3}_tcq_codebook_fattn` → `d_turboq{2,3}_tcq_codebook` (no `_fattn` suffix because yggdrasil shares one codebook across encode + decode paths). - V dispatch: yggdrasil's `flash_attn_ext_vec` inlines TURBOQ{2,3}_TCQ V dequant directly inside the V-aggregation loop rather than routing through `dequantize_V_*` (per Phase 3a ggml-org#11 — `727799470`). So the V-side payoff is a one-line `d_turboq{2,3}_tcq_codebook` → `tcq_smem_codebook` swap inside the inline branches at fattn-vec.cuh, instead of buun's `if constexpr` + `_cb<half/float, ne>` wrapper calls. The K-side matches buun: split helper into `_cb` core + thin `__constant__` wrapper, dispatch the `_cb` form for TCQ types and fall through to `vec_dot_KQ` otherwise. - V-helper split (`dequantize_V_turboq{2,3}_tcq_cb` + wrapper) lifted for surface parity even though the VEC kernel bypasses these helpers — they are still reachable via `get_dequantize_V` from non-VEC dispatch paths and keep using `__constant__` there. SMEM footprint: 512 floats (2 KiB) for TCQ3, 256 floats (1 KiB) for TCQ2 — declared per-block via the conditional `tcq_smem_cb_size` constexpr and zeroed to size 1 when neither side is TCQ (no SMEM cost for unrelated kernels). Well under HIP's 64 KiB cap and below the existing turbo_lut alloc for D=256. Composition check: touches only the TCQ helper definitions in fattn-common.cuh and three sites in fattn-vec.cuh (top-of-kernel codebook load + K dispatch + V inline). Disjoint from s55 warp-shuffle cleanup (set-rows.cu), s56 measurement plumbing (set-rows.cu), s57 routing fix (fattn.cu), s58 backtrace + double-buffer (set-rows.cu). Validation: see follow-up smoke + PPL parity cells (session 60). recon/09 INCLUDE-table row ggml-org#17.
…ation
Manual port of buun's `3c4c44e5c` ("simplify: consolidate codebook load loop,
remove hardcoded sizes"). Merges the dual `is_tcq3` / `is_tcq2` codebook-load
branches in `flash_attn_ext_vec` into a single loop driven by `smem_cb_size > 0`,
with the source codebook selected via a constexpr-resolvable ternary on the
hot pointer. Hardcoded `512` and `256` loop bounds collapse into the shared
`tcq_smem_cb_size` constant.
Pure refactor — no algorithmic change. Both branches were behaviorally
identical (cooperative load + __syncthreads); consolidation just dedups the
loop body. Expected PPL: bit-identical to s60 anchor. Expected TPS: zero
delta (compiler should produce the same SASS via constexpr resolution).
Note on starter framing: the s61 starter prompt described buun ggml-org#17+ggml-org#18 as a
"setup + use" pair where ggml-org#18 was the perf-payoff. That model is incorrect.
Buun ggml-org#17 (692cffd, s60) WAS the full __constant__→__shared__ functional
change — it already moved the codebook AND swapped all VEC call sites. Buun
ggml-org#18 (3c4c44e, this commit) is just a code-shape cleanup of the load loop
that landed on the same merge bubble (b2cfd17). The decode-side perf
benefit of buun's design lives entirely in s60.
Source: buun 3c4c44e (merged in b2cfd17).
Mainline forward-sync per [[yggdrasil-mainline-forward-sync]].
… (setup for ggml-org#18) Manual port of buun's `692cffde1` ("experiment: S3 TCQ codebook from __constant__ to __shared__ memory"). Moves the TURBOQ2_TCQ / TURBOQ3_TCQ codebooks out of __constant__ memory into a cooperatively-loaded __shared__ array in `flash_attn_ext_vec`. __constant__ serializes when threads in a warp hit different 32B cache lines during random codebook lookups; __shared__ gives full 32-bank parallel access. Standalone this commit is near-no-op on perf — its value is setting up ggml-org#18. Yggdrasil divergences from buun's diff: - Naming: buun `turbo{2,3}_tcq` → yggdrasil `turboq{2,3}_tcq`; codebook symbol `d_turbo{2,3}_tcq_codebook_fattn` → `d_turboq{2,3}_tcq_codebook` (no `_fattn` suffix because yggdrasil shares one codebook across encode + decode paths). - V dispatch: yggdrasil's `flash_attn_ext_vec` inlines TURBOQ{2,3}_TCQ V dequant directly inside the V-aggregation loop rather than routing through `dequantize_V_*` (per Phase 3a ggml-org#11 — `727799470`). So the V-side payoff is a one-line `d_turboq{2,3}_tcq_codebook` → `tcq_smem_codebook` swap inside the inline branches at fattn-vec.cuh, instead of buun's `if constexpr` + `_cb<half/float, ne>` wrapper calls. The K-side matches buun: split helper into `_cb` core + thin `__constant__` wrapper, dispatch the `_cb` form for TCQ types and fall through to `vec_dot_KQ` otherwise. - V-helper split (`dequantize_V_turboq{2,3}_tcq_cb` + wrapper) lifted for surface parity even though the VEC kernel bypasses these helpers — they are still reachable via `get_dequantize_V` from non-VEC dispatch paths and keep using `__constant__` there. SMEM footprint: 512 floats (2 KiB) for TCQ3, 256 floats (1 KiB) for TCQ2 — declared per-block via the conditional `tcq_smem_cb_size` constexpr and zeroed to size 1 when neither side is TCQ (no SMEM cost for unrelated kernels). Well under HIP's 64 KiB cap and below the existing turbo_lut alloc for D=256. Composition check: touches only the TCQ helper definitions in fattn-common.cuh and three sites in fattn-vec.cuh (top-of-kernel codebook load + K dispatch + V inline). Disjoint from s55 warp-shuffle cleanup (set-rows.cu), s56 measurement plumbing (set-rows.cu), s57 routing fix (fattn.cu), s58 backtrace + double-buffer (set-rows.cu). Validation: see follow-up smoke + PPL parity cells (session 60). recon/09 INCLUDE-table row ggml-org#17.
…ation
Manual port of buun's `3c4c44e5c` ("simplify: consolidate codebook load loop,
remove hardcoded sizes"). Merges the dual `is_tcq3` / `is_tcq2` codebook-load
branches in `flash_attn_ext_vec` into a single loop driven by `smem_cb_size > 0`,
with the source codebook selected via a constexpr-resolvable ternary on the
hot pointer. Hardcoded `512` and `256` loop bounds collapse into the shared
`tcq_smem_cb_size` constant.
Pure refactor — no algorithmic change. Both branches were behaviorally
identical (cooperative load + __syncthreads); consolidation just dedups the
loop body. Expected PPL: bit-identical to s60 anchor. Expected TPS: zero
delta (compiler should produce the same SASS via constexpr resolution).
Note on starter framing: the s61 starter prompt described buun ggml-org#17+ggml-org#18 as a
"setup + use" pair where ggml-org#18 was the perf-payoff. That model is incorrect.
Buun ggml-org#17 (692cffd, s60) WAS the full __constant__→__shared__ functional
change — it already moved the codebook AND swapped all VEC call sites. Buun
ggml-org#18 (3c4c44e, this commit) is just a code-shape cleanup of the load loop
that landed on the same merge bubble (b2cfd17). The decode-side perf
benefit of buun's design lives entirely in s60.
Source: buun 3c4c44e (merged in b2cfd17).
Mainline forward-sync per [[yggdrasil-mainline-forward-sync]].
* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * flashattention and matrix multiplication moved to new format * clean up preprocessing * Formatting * remove duplicate constants * Split large shaders into multiple static strings --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
…better shader parameter handling (ggml-org#20173) * K quant speedup (ggml-org#20) * Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17) * scale jit working * preliminary working jit for getrows and mulmat, needs refining * simplified mul_mat preprocessing switch statement * get_rows fixes, mul_mat refinement * formatted + last edits * removed some extraneous prints * fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish * small fix * some changes, working * get_rows and mul_mat jit fixed and working * Update formatting * formatting * Add header --------- Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local> Co-authored-by: Reese Levine <reeselevine1@gmail.com> * Start work on all-encompassing shader library * refactor argmax, set_rows * Refactor all but flashattention, mat mul * no gibberish, all k quants added, merged * vec memory fix * q6_k matching metal on my machine, tests passing * Set tile size for q6_k separately * Separate out fast shaders --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com> * Move towards writeBuffer for params * Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups * Remove extra file * Formatting --------- Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
) * ggml: backend-agnostic tensor parallelism * support for GPT-OSS, Qwen 3 MoE * partial Vulkan fix * add support for 4/8 GPUs * unconditional peer access * re-use buffers + ggml contexts * fix output pattern * NCCL support * GGML: HIP: add RCCL support * Remove shfl and AllReduce from backend interface * move allocation workaround out of ggml-alloc.c * 2d tensor set/get support * Fix the seg fault without NCCL * Apply suggestion from JohannesGaessler * support for tensor dims % n_devs != 0 * fix view_offs scaling * arbitrary num. of GPUs/tensor split * fix compilation * better granularity estimate * Support device-specific host buffer types if all underlying backends expose the same type. This allows using pinned memory instead of pageable memory for CUDA. Fix compilation errors. * partial Qwen 3 Next support * Fix qwen3 30b (ggml-org#8) * Fix crash with Qwen-30B-A3B Q4_0 Qwen-30B-A3B Q4_0 has an intermediate dimension of 768. Using a granularity of 256 forces an uneven split between GPUs, which is not supported by the current implementation. * Decide block size based on tensor quantization type * Fix crashes due to KV cache serialization (ggml-org#9) KV cache serialization requires non-zero offsets on the tensor. Add support in the meta backend to set/get a tensor with a non-zero offset. * metal : fix build (ggml-org#7) * static memory allocations, fix usage count * fix tensor granularity * more even memory distribution * use BF16 for allreduce * rebase fixup * better error message for unsupported architectures * Fix device mismatch during scatter of allReduce. (ggml-org#11) There is a mismatch between the dst buffer device and the backend device, causing the use of sync copies * Enable the previous allreduce implementation. It is better in both perf and stability (ggml-org#12) * delay AllReduce for Moe for less I/O * build : clean-up compile warnings * backend : move most of the meta backend API to ggml-backend-impl.h * cont : hide unused public API in the implementation * llama : use llama_device + remove ggml_backend_dev_is_meta() * ggml-backend : remove unused alloc include * minor : remove regex include * ggml : introduce ggml-ext.h for staging new APIs * rebase fixup * fix tests * llama : more robust logic for determining Meta devices (ggml-org#16) * llama : more robust logic for determining Meta devices * cont : fix devs size check Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * cont : fix log type Co-authored-by: Johannes Gäßler <johannesg@5d6.de> --------- Co-authored-by: Johannes Gäßler <johannesg@5d6.de> * disable roundtrip for meta backend * fix arch selection * Qwen 3.5 support * fix Gemma 4 MoE * fix OpenVino, SYCL * fix test-llama-archs for CPU-only builds * Fix Qwen 3.5 MoE * disable meta backend tests for WebGPU * tests : filter CPU-based devices from the Meta backend tests (ggml-org#17) * meta : formatting, naming, indentation (ggml-org#18) * formatting : llama-model.cpp * formatting : ggml-ext.h * formatting : ggml-backend-meta.cpp * meta : add TODO * add documentation * better error messages * fix GPT-OSS --------- Co-authored-by: Carl Philipp Klemm <carl@uvos.xyz> Co-authored-by: Gaurav Garg <gaugarg@nvidia.com> Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
No description provided.