ik_llama.cpp 的 MMQ for Q6_0为 6-bit 量化补上矩阵乘加速内核【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址: https://gitcode.com/GitHub_Trending/ik/ik_llama.cppik_llama.cppllama.cpp 的一个 fork专注于 SOTA 量化方案与推理性能改进在 2024 年 11 月通过 PR #115 为Q6_0量化类型补齐了MMQMatrix Multiplication Quantized量化矩阵乘CUDA 内核。本文以该 PR 的讨论记录为主体结合当前仓库中 mmq.cuh、mmq.cu、mmq_id_common.cuh 等源码实现系统讲解 Q6_0 格式的存储结构、MMQ 内核的数据加载load_tiles与流水线原理、dp4a/INT8 MMA 双路径实现以及它和 Q6_K 在精度与速度上的取舍帮助你理解一个量化类型 × GPU 内核组合从补丁到落地的完整链路。说明PR #115 的原始讨论中测试者 Nexesenex 报告了一个经验性结论——在纯量化pure quant的 Sheared Llama 2.7B 上Q6_0 的困惑度PPL比 Q6_K 高出约 0.1%。该数据来自当时社区测试仅作为精度参考不代表本项目对 Q6_0 与 Q6_K 的官方结论。背景为什么需要为 Q6_0 单独写 MMQ 内核量化矩阵乘在 CUDA 上的两种形态在 ggml 的 CUDA 后端里量化权重的矩阵乘法主要有两条路径MMVQMatrix-Matrix / Vector Quantized每个线程独立处理一列适合 token 数很少接近 GEMV矩阵-向量乘的场景实现简单但并行度受限MMQMatrix Multiplication Quantized按 tile 切分矩阵、通过共享内存暂存分块数据、用 dp4a4 路点积累加指令或 INT8 MMA张量核心做高吞吐乘加适合 token 数较多的 GEMV/GEMM 场景。Q6_0是 ggml 里老牌的基础量化格式6-bit 对称量化。在 PR #115 之前CUDA 后端虽然已经支持Q6_0的普通乘法路径却没有专门的 MMQ 内核导致它在批处理场景下拿不到 dp4a/MMA 的加速红利。PR #115 的标题非常简洁——Add MMQ kernel for Q6_0就是补上这块拼图。从源码看 Q6_0 的存储格式要理解 MMQ 内核先看Q6_0的数据布局。在 ggml-common.h 中定义了块结构#define QK6_0 32 typedef struct { ggml_half d; // delta块缩放因子fp16 uint8_t qh[QK6_0/4]; // 每个量化值第 5、6 bit高位 uint8_t qs[QK6_0/2]; // 每个量化值低 4 bitnibble } block_q6_0;每个块包含 32 个 6-bit 量化值d块的缩放因子fp16qs每个值低 4 位32 个值占 16 字节qh每个值的高 2 位bit5/bit632 个值占 8 字节。量化值是无符号的范围 063在反量化时统一减去偏移 32得到 [-32, 31] 的有符号范围。QI6_0宏ggml-common.h定义了每个线程可打包的量化值数量用于内核的向量化展开。正因Q6_0的量化值是6-bit 拆分存储4 位 2 位MMQ 内核无法像Q8_0那样直接按字节取整型而必须先做**位重组bit unpack**把 6-bit 值还原成 int8 后喂给 dp4a 或 MMA 指令——这正是load_tiles_q6_0存在的意义。MMQ for Q6_0 的实现从位重组到张量核心内核调度入口mmq.cu 与 mmq_id.cuPR #115 的改动落地后Q6_0正式进入 MMQ 调度表。在 mmq.cu 中ggml_cuda_op_mul_mat_q的分发逻辑为void ggml_cuda_op_mul_mat_q(ggml_backend_cuda_context ctx, enum ggml_type type, const mmq_args args) { auto stream ctx.stream(); switch (type) { case GGML_TYPE_Q4_0: mul_mat_q_caseGGML_TYPE_Q4_0(ctx, args, stream); break; // ... case GGML_TYPE_Q6_0: mul_mat_q_caseGGML_TYPE_Q6_0(ctx, args, stream); break; // ... } }同时针对 MoEMixture of Experts模型的mul_mat_id路径mmq_id.cu 与 mmq_id_common.cuh 也接入了Q6_0意味着 MoE 模型中Q6_0量化的专家权重同样可以走 MMQ 内核。模板实例由脚本自动生成见 mmq-instance-q6_0.cu#include ../mmq.cuh DECL_MMQ_CASE(GGML_TYPE_Q6_0);以及配套的_id变体 mmq-instance-q6_0_id.cu。核心函数 load_tiles_q6_0把 6-bit 拆成两个 int8 半值MMQ 的关键步骤是load_tiles——把权重 tile 从全局内存搬进共享内存并完成数据类型转换。Q6_0的版本定义在 mmq.cuhtemplate int mmq_y, int nwarps, bool need_check static __device__ __forceinline__ void load_tiles_q6_0( const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { const int kbx threadIdx.x / QI6_0; const int kqsx threadIdx.x % QI6_0; #pragma unroll for (int i0 0; i0 mmq_y; i0 nwarps) { int i i0 threadIdx.y; if (need_check) { i min(i, i_max); } const block_q6_0 * bxi (const block_q6_0 *)(x i*stride) kbx0 kbx; const int ql get_int_b2(bxi-qs, kqsx); // 取回 4 个低 4 位 const int qh get_int_b2(bxi-qh, kqsx%2) 4*(kqsx/2); // 取回对应高 2 位 int qs0 ((ql 0) 0x0F0F0F0F) | ((qh 4) 0x30303030); int qs1 ((ql 4) 0x0F0F0F0F) | ((qh 2) 0x30303030); qs0 __vsubss4(qs0, 0x20202020); // subtract 32转为有符号 qs1 __vsubss4(qs1, 0x20202020); // subtract 32 #ifdef INT8_MMA_AVAILABLE x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 kbx*(2*QI6_0) kqsx 0] qs0; x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 kbx*(2*QI6_0) kqsx QI6_0] qs1; #else x_qs[i*(2*WARP_SIZE 1) kbx*(2*QI6_0) kqsx 0] qs0; x_qs[i*(2*WARP_SIZE 1) kbx*(2*QI6_0) kqsx QI6_0] qs1; #endif } // 缩放因子 d 的加载循环 ... }这段代码蕴含了 Q6_0 特有的两个技巧位重组bit unpack一个 6-bit 量化值q实际是低 4 位存于 qs 高 2 位存于 qh。ql一次取回 4 个值的低 4 位qh取回对应的高 2 位再通过掩码与移位拼出两个 int32qs0、qs1每个 int32 里是 4 个 8-bit 有符号值恰好匹配 dp4a 的 4 路点积输入格式。__vsubss4(qs0, 0x20202020)是 SIMD 减法一次性给 4 个字节都减 32完成 [-32, 31] 的有符号化。共享内存双份存放一个Q6_0值被拆成两个半值写入共享内存0与QI6_0两个偏移这样后续计算阶段可以直接复用与Q8_0相同的流水线只不过数据量翻倍。双路径dp4a 与 INT8 MMAmmq.cuh 里定义了 MMQ 的通用参数#define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. #define MMQ_ITER_K 256 #define MMQ_NWARPS 8MMQ_DP4A_MAX_BATCH_SIZE当 GPU 有 FP16 张量核心时dp4a 路径可承载的最大 batch size默认 64超出后倾向切换到其他路径MMQ_ITER_Kk 维每次迭代处理的元素数256控制共享内存占用与指令流水MMQ_NWARPS每个 block 的 warp 数8决定并行度与共享内存划分。Q6_0的 tile 尺寸通过以下宏映射mmq.cuh#define MMQ_DP4A_TXS_Q8_0 tile_x_sizes{mmq_y*WARP_SIZE*2 mmq_y, mmq_y*WARP_SIZE*2/QI8_0 mmq_y/(QI8_0/2), 0} // ... case GGML_TYPE_Q6_0: return MMQ_DP4A_TXS_Q8_0; // dp4a 路径复用 Q8_0 的 tile 布局 case GGML_TYPE_Q6_0: return MMQ_MMA_TILE_X_K_Q8_0; // INT8 MMA 路径注意这里Q6_0 在 tile 布局上复用了 Q8_0 的尺寸——因为 Q6_0 拆包后每个值占 1 字节同 Q8_0只是块内元素翻倍这让内核的存储规划可以共享现有宏。mmq_get_q8_1_ds_layoutmmq.cuh也为 Q6_0 返回MMQ_Q8_1_DS_LAYOUT_D4每个 32 值一个 fp32 scale保证权重 tile 与激活 tile 的 scale 布局一致。两条路径的选择由INT8_MMA_AVAILABLE宏在编译期决定dp4a 路径无 INT8 MMA或老架构用__dp4a指令做 int8×int8→int32 的 4 路点积tile 走MMQ_DP4A_TXS_Q8_0布局INT8 MMA 路径有 INT8 张量核心的架构走MMQ_MMA_TILE_X_K_Q8_0布局直接用张量核心的 mma 指令做 8-bit 矩阵乘吞吐更高。load_tiles_q6_0内部用#ifdef INT8_MMA_AVAILABLE分别写出两种布局保证两条路径都能拿到正确排布的共享内存数据。MoE 场景mmq_id 的 Q6_0 支持MoE 模型里专家权重是mul_mat_id运算。PR #115 的改动同样覆盖了这一路径mmq_id_common.cuh 提供了独立的load_tiles_q6_0模板参数为mmq_y, need_check并在 mmq_type_traits_id 中注册。这意味着Q6_0量化权重的 MoE 模型在 CUDA 上也能享受 MMQ 的加速而不是退回到逐列MMVQ路径。内核质量检查模板实例与测试MMQ 内核通过模板实例化表驱动生成仓库里既有普通实例 mmq-instance-q6_0.cu也有 MoE 实例 mmq-instance-q6_0_id.cu实例文件由 generate_cu_files.py 自动生成保证新增量化类型的声明与实现一致。在 CPU 参考实现方面Q6_0的量化和反量化基准函数位于 ggml-quants.h 与 ggml-quants.hquantize_row_q6_0_ref、dequantize_row_q6_0CUDA 内核的正确性以这些参考实现为对齐基准。若想验证内核正确性可在仓库根目录用 CMake 构建后运行 test-backend-ops.cpp 等后端算子测试。精度观察Q6_0 与 Q6_K 的取舍PR #115 讨论区里测试者 Nexesenex 在 IK_LLama 上做了实测并给出了一个常被引用的结论2024-11-20 两条重复评论Tested successfully on IK_LLama, PPL is 0.1% above Q6_K on a pure quant of Sheared Llama 2.7b.即在 Sheared Llama 2.7B 的纯量化模型上Q6_0的困惑度PPL比Q6_K高约 0.1%。由于Q6_0是更简单的块结构单 scale、无 per-block 残差分组精度略低于分组的Q6_K属于预期行为而 MMQ 内核的意义在于让Q6_0在同样能跑的基础上获得更快的推理速度给用户多一个精度-速度权衡点。需要强调的是这个 0.1% 是单次社区测试数据不同模型、不同量化方式纯量化 vs 有 imatrix 的量化、不同上下文长度下差异可能不同本项目不保证该数值在所有模型上复现也不代表Q6_0与Q6_K的官方精度对比结论实际使用时建议结合自己的模型用 perplexity 等工具做评估。如何验证与使用构建按 install.md 说明用 CMake 构建带 CUDA 后端的 ik_llama.cpp-DGGML_CUDAONQ6_0MMQ 内核会随模板实例自动编译进libggml-cuda。量化模型可用 quantize 工具把模型量化到Q6_0q6_0或q6_k等类型取决于转换脚本支持。运行用 main 或 llama-server 加载模型观察 CUDA 利用率与 token 吞吐批处理场景如 batched-bench最能体现 MMQ 相对 MMVQ 的吞吐优势。精度评估用 perplexity 对比同模型Q6_0与Q6_K的 PPL 差异。小结PR #115 虽然只是为 Q6_0 增加一个 MMQ 内核背后却是量化推理工程中存储格式 → 位重组 → 共享内存 tile → dp4a/MMA 双路径的完整链条Q6_0的 6-bit 拆包存储决定了内核必须做位重组load_tiles_q6_0通过掩码移位 SIMD 减偏移把 6-bit 值还原为 int8 并双份写入共享内存内核通过INT8_MMA_AVAILABLE同时支持 dp4a 与 INT8 MMA 两条路径并在 tile 布局上复用Q8_0的成熟设计MoE 场景通过mmq_id路径同样获得支持社区实测给出的Q6_0 PPL 比 Q6_K 高约 0.1%是重要的精度参考让用户在速度与精度之间有了更清晰的决策依据。这一改动也体现了 ik_llama.cpp 的典型工作方式在 ggml 既有内核框架内逐个补齐量化类型的 MMQ 支持从而让每种量化格式在 CUDA 上都能吃到同样的加速红利。延伸阅读mmq.cuhMMQ 内核核心头文件含load_tiles_q6_0、tile 尺寸宏与mmq_type_traitsmmq.cuggml_cuda_op_mul_mat_q调度分发mmq_id.cu 与 mmq_id_common.cuhMoE 的mul_mat_idMMQ 路径ggml-common.hblock_q6_0等量化块结构定义template-instances/mmq-instance-q6_0.cuQ6_0 普通 MMQ 实例template-instances/mmq-instance-q6_0_id.cuQ6_0 MoE MMQ 实例。【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址: https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp创作声明：本文部分内容由AI辅助生成（AIGC），仅供参考