十年匠心定制 · 商业建站与技术教学双线并行 咨询热线:400-886-1026 service@lmnt.cn
ARTICLE DETAIL

资讯详情

深耕网站建设与运营推广的一线实战洞察。

ik_llama.cpp PR 461 深度解析:IQ2_K_R4 / IQ3_K_R4 / IQ4_K_R4 / IQ5_K_R4 的 CUDA 实现与部署实战

ik_llama.cpp PR 461 深度解析:IQ2_K_R4 / IQ3_K_R4 / IQ4_K_R4 / IQ5_K_R4 的 CUDA 实现与部署实战 人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载导读本文围绕 ik_llama.cpp 的 PR 461 展开该 PR 为行交错row-interleaved量化类型IQ2_K_R4、IQ3_K_R4、IQ4_K_R4、IQ5_K_R4补上了 CUDA 端 GEMM/GEMV 实现解决了CPU 端性能更好的 R4 量化无法在 GPU 上直接推理的痛点。读完本文你将掌握这些 R4 量化类型的来龙去脉与格式定义、PR 461 的 CUDA 实现策略dequantize cuBLAS及其在源码中的落点、GGML_CUDA_IQK_FORCE_BF16编译选项的适用场景以及如何在 CUDA 环境下编译、量化/重打包并实际运行 R4 量化模型。背景为什么需要给IQX_K_R4补 CUDA 实现IQX_K与IQX_K_R4的关系ik_llama.cpp 的核心特色之一是在llama.cpp主线上游的基础上引入了大量 SOTA 级的新量化类型参见 README.md 中 IQK quants 一节。其中IQX_K如IQ2_K、IQ3_K、IQ4_K、IQ5_K是K 系列 i-quants 的行交错变体在相同 bpwbits per weight下量化质量优于对应的 i-quants、k-quants 或 legacy quantsIQX_K_R4是IQX_K的行交错重打包repacked row-interleaved版本在 CPU 上拥有更好的矩阵乘法GEMM/GEMV性能。这两组类型的 GGML 类型 ID 与模型格式 IDftype都可以在头文件中直接查到在 ggml/include/ggml.h 中GGML_TYPE_IQ2_K_R4 337、GGML_TYPE_IQ3_K_R4 338、GGML_TYPE_IQ4_K_R4 339、GGML_TYPE_IQ5_K_R4 340在 include/llama.h 中LLAMA_FTYPE_MOSTLY_IQ2_K_R4 338、IQ3_K_R4 339、IQ4_K_R4 340、IQ5_K_R4 341注释except 1d tensors表明 1 维张量保持原始类型不参与重打包。痛点CPU 快、GPU 无解PR 461 的描述点出了当时的现实困境IQX_K_R4在 CPU 上性能更好但没有对应的 CUDA GEMM/GEMV 实现因此无法直接用于 GPU 推理于是quant cookers量化作者只能发布IQX_K非 R4版本以便 GPU 用户使用而纯 CPU 用户为了享受更好的 CPU 性能又需要自己把模型重打包repack成 R4此外社区成员 ubergarm 发布的IQK_X_R4量化模型同样无法用于 GPU 推理。也就是说同一份模型不得不在 CPU 与 GPU 两个阵营之间二选一。PR 461 的目标正是消除这种不便为IQ2_K_R4、IQ3_K_R4、IQ4_K_R4、IQ5_K_R4提供 CUDA 实现让 R4 模型可以直接在 NVIDIA GPU 上运行。PR 461 的 CUDA 实现策略与源码落点第一阶段dequantize cuBLASPR 461 明确说明当前阶段 GEMM 采用dequantize cuBLAS的方式实现即先把 R4 量化权重反量化成浮点再交给 cuBLAS 执行矩阵乘法作者表示后续可能补充量化 GEMMMMQmatrix-matrix multiplication with quantized data。从当前仓库源码看该实现已经落地并进一步演进为完整的量化内核quantized kernels体系MMQ量化 GEMM实例在 ggml/src/ggml-cuda/template-instances/mmq-instance-iq2_k_r4.cu 等文件中为IQ2_K_R4/IQ3_K_R4/IQ4_K_R4/IQ5_K_R4各自生成了load_tiles_*与mmq_type_traits特化。以IQ2_K_R4为例load_tiles_iq2_k_r4会按block_iq2_k_r4布局取出 4 位量化值bxi-qs、符号/额外位bxi-extra与缩放bxi-scales、bxi-d经iq2nl_values/iq2k_table查找表还原后送入vec_dot_q8_0_16_q8_1_mma或vec_dot_q8_0_16_q8_1_dp4a完成点积GEMV量化矩阵乘向量实例在 ggml/src/ggml-cuda/template-instances/mmvq-instance-iq2_k_r4.cu 中定义了vec_dot_iq2_k_r4_q8_1通过ggml_cuda_dp4a与iq2k_table查找表实现 R4 权重与Q8_1激活的点积并由iqk_mul_mat_vec_q_cudaGGML_TYPE_IQ2_K_R4, ...实例化mul_mat_vec_iq2_k_r4_q8_1_cuda类型分发在 ggml/src/ggml-cuda/mmq.cu 的mul_mat_q_case与 mmq_id.cu 的mul_mat_q_case_id中四个 R4 类型均被分发到对应的量化内核反量化转换在 ggml/src/ggml-cuda/convert.cu 中四个 R4 类型也都注册了 CUDA 反量化路径这正是dequantize cuBLAS兜底方案所依赖的转换入口。从源码结构可以推断PR 461 提交时的 dequantize cuBLAS 方案后来在 README 所述的 PR 557 等后续工作中被完整的量化 GEMM/GEMV 内核所增强——README 明确记录All quantization types now have quantized matrix multiplication CUDA kernels所有量化类型现在都有 CUDA 量化矩阵乘法内核。对 MoE 专家并行mul_mat_id的支持仓库中还可以看到IQX_K_R4同时出现在mmq_id.cu的mul_mat_q_case_id分发路径中说明这些 R4 类型也支持 MoE 模型的专家张量expert tensors量化矩阵乘法这与 PR 461 描述中提及的 DeepSeek-V3/R1 等大 MoE 模型场景是吻合的。关键编译选项GGML_CUDA_IQK_FORCE_BF16为什么需要 bf16PR 461 描述中给出了一个重要警告fp16 在 cuBLAS 路径下可能导致数值不稳定和乱码输出这一现象在 DeepSeek 系列模型如 DeepSeek-V3/R1的IQX_K_R4量化上被观察到。因此作者提供了编译选项GGML_CUDA_IQK_FORCE_BF16开启后当没有可用的 MMQ 量化内核时强制以bf16精度配合 cuBLAS 执行矩阵乘法作者默认不开启因为这会降低 prompt processingPP性能而就作者所知 bf16 仅对 DeepSeek 模型是必须的。该选项在源码中的落点CMake 选项定义于 ggml/CMakeLists.txtoption(GGML_CUDA_IQK_FORCE_BF16 ggml: use bf16 cuBLAS when no MMQ kernel is available OFF)编译宏注入于 ggml/src/CMakeLists.txt 与 ggml/src/CMakeLists.txt运行时行为实现在 ggml/src/ggml-cuda.cu当定义了GGML_CUDA_IQK_FORCE_BF16且src0为量化类型时先把src0/src1转换为nv_bfloat16再通过cublasGemmEx(..., CUDA_R_16BF, ..., CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)完成矩阵乘法最后把结果转回 fp32。编译命令示例针对 DeepSeek-V3/R1 这类需要 bf16 的模型推荐使用cmake -B ./build -DCMAKE_BUILD_TYPERelease -DGGML_CUDAON -DGGML_CUDA_IQK_FORCE_BF161 cmake --build ./build --config Release -j对普通模型不需要 bf16则可不加-DGGML_CUDA_IQK_FORCE_BF16或显式置0以获得更快的 prompt processing。社区实测反馈bf16 未必更慢PR 461 的对话中还记录了一条来自社区的反馈discussions/477 - DeepSeek-R1-0528 ik quants 中亦有引用有用户在特定硬件与特定量化组合下报告FORCE_BF161反而带来了速度提升这与作者的预期相反。可见 bf16 的影响与硬件配置GPU 架构、Tensor Core 行为等和具体量化类型密切相关建议用户在自己的硬件上实测对比。如何在 CUDA 环境下使用 R4 量化模型方案一直接使用已发布的 R4 模型社区如 ubergarm发布过IQK_X_R4量化模型。在 PR 461 之前这类模型只能用于 CPU 推理现在则可以直接在 NVIDIA GPU 上运行。使用方式与普通模型一致例如./build/bin/llama-cli -m /path/to/model-IQ3_K_R4.gguf -p Hello -n 64具体二进制名称以构建产物为准如llama-cli/llama-server等。方案二自己量化或重打包如果手头只有IQX_K非 R4模型可以使用 examples/quantize/quantize.cpp 提供的工具进行量化或重打包。从该文件中的 ftype 表可以看到四个目标类型{ IQ2_K_R4, LLAMA_FTYPE_MOSTLY_IQ2_K_R4, IQ2_K repacked, }, { IQ3_K_R4, LLAMA_FTYPE_MOSTLY_IQ3_K_R4, IQ3_K repacked, }, { IQ4_K_R4, LLAMA_FTYPE_MOSTLY_IQ4_K_R4, IQ4_K repacked, }, { IQ5_K_R4, LLAMA_FTYPE_MOSTLY_IQ5_K_R4, IQ5_K repacked, },典型用法把已有的IQ2_K模型重打包为 R4./build/bin/llama-quantize /path/to/model-IQ2_K.gguf /path/to/model-IQ2_K_R4.gguf IQ2_K_R4PR 461 对话中还提到-rtrrepack-to-row-interleaved选项会在加载模型时把所有留在内存中的张量重打包为行交错格式。需要特别注意的是 README 的警告混合 CPU/GPU 推理 MoE 模型且部分专家留在 CPU 时不要随意使用-rtr因为并非所有量化类型都有 CUDA 行交错实现如 k-quantsQ2_K/Q3_K/Q4_K/Q5_K/Q6_K就没有 CUDA R4 实现一旦重打包这些张量的矩阵乘法将永远在 CPU 上执行反而拖慢 prompt processing。哪些模型适合用 R4 CUDAPR 461 明确点名的场景是DeepSeek-V3/R1这类大 MoE 模型——它们体量大、bpw 敏感R4 的低 bpw 高保真特性在这里收益明显但使用时必须注意GGML_CUDA_IQK_FORCE_BF16的取舍见上文。普通模型则默认不需要 bf16。与 R4 家族其他类型的关联IQX_K_R4只是 ik_llama.cpp R4 行交错家族的一部分。从 examples/quantize/quantize.cpp 的 ftype 表可以看到完整的家族包括IQ1_S_R4、IQ1_M_R4、IQ2_XXS_R4、IQ2_XS_R4、IQ2_M_R4、IQ2_BN_R4、IQ3_XXS_R4、IQ3_S_R4、IQ4_NL_R4、IQ4_KS_R4、IQ5_KS_R4、Q2_K_R4、Q3_K_R4、Q4_K_R4、Q5_K_R4、Q6_K_R4、Q8_K_R8、Q8_KV_R8等覆盖 1.5~8 bpw 的多个档位。README 的 IQK quants 一节记录了各类型的初始实现 PR其中IQ2_K_R4、IQ3_K_R4、IQ4_K_R4、IQ5_K_R4的 CUDA 实现正是 PR 461。PR 461 作者还预告了后续会单独为IQ2_KS_R4、IQ4_KS_R4、IQ5_KS_R4补充 CUDA 实现对话中确认IQ2_KS_R4当时尚未完成且作者一度考虑对其 packing 做破坏性变更这也是推迟 CUDA 实现的原因。从当前仓库看IQ4_KS_R4、IQ5_KS_R4的 CUDA 实现已在 README 中被记为 PR 462 与 PR 493IQ1_S_R4、IQ1_M_R4也分别由 PR 492/494 覆盖说明该 PR 家族的 CUDA 支持已基本补齐。延伸对话中的量化技术讨论要点PR 461 的评论区还沉淀了不少关于量化方法论的讨论可作为理解该项目量化思路的参考KTtrellis量化作者开发的IQX_KT系列QTIP/exl3 风格的 trellis quants在 CUDA 上性能不错、bpw 更低如iq2_kt约 2.125 bpw低于iq2_ks的 2.1875且 PPL 更低但 CPU 上 TGtoken generation很慢且量化耗时约为IQK的 5 倍启发式而非最优解作者明确表示量化实现并不追求 RMSE 最小化的证明最优解——因为 CPU 上精确求解代价过高且经验表明更低的 RMSE 未必带来更好的可观察模型质量因此采用针对 trellis 输出值调优的启发式fudge factors对话中讨论的1.05f、1.01f等反量化缩放系数见dequantize_block_iq2_kt等实现是作者实验期间为快速调参而保留的经验系数本应吸收进行缩放中差异很小评估口径作者建议使用ln(PPL(Q)/PPL(f16))与 KLD 直接相关来做不同框架/语料间的苹果对苹果比较因为该比值对 PPL 计算方式不敏感。这些讨论帮助我们理解R4 的 CUDA 支持只是该仓库量化生态的一环其背后是作者围绕 bpw、PPL、CPU/GPU 性能之间的持续权衡。总结PR 461 为IQ2_K_R4、IQ3_K_R4、IQ4_K_R4、IQ5_K_R4打通了 NVIDIA GPU 推理路径实现策略以dequantize cuBLAS起步源码落点见 convert.cu、ggml-cuda.cu后续演进为完整的量化 GEMM/GEMV 内核见 mmq.cu、mmq_id.cu 与template-instances/下的实例文件DeepSeek-V3/R1 等模型在 cuBLAS 路径下需要-DGGML_CUDA_IQK_FORCE_BF161编译以保证数值稳定代价是 prompt processing 性能下降该开关对性能的影响与硬件相关建议实测配合 examples/quantize/quantize.cpp 的IQ2_K_R4/IQ3_K_R4/IQ4_K_R4/IQ5_K_R4ftype用户可以自行量化或重打包模型让同一份模型同时享受 CPU 与 GPU 两侧的优势但需注意-rtr在混合 CPU/GPU 推理 MoE 模型时的副作用。至此CPU 用户要 R4、GPU 用户要 IQX_K的割裂局面在IQX_K_R4上得到了解决发布一份 R4 模型CPU 与 CUDA 用户均可直接使用。赞分享人工智能大模型推理引擎本地部署模型量化模型优化【免费下载链接】ik_llama.cppllama.cpp fork with additional SOTA quants and improved performance项目地址https://gitcode.com/GitHub_Trending/ik/ik_llama.cpp点击查看免费下载相关推荐ik_llama.cpp 的 CUDA MMQ 内核为 IQ2_K_R4~IQ5_K_R4 量化消除反量化开销ik_llama.cpp 的 CUDA MMQ 内核为 IQ2_K_R4~IQ5_K_R4 量化消除反量化开销 本文围绕 ik_llama.cpp 仓库中 P人工智能大模型推理引擎本地部署模型量化模型优化ik_llama.cpp CUDA 量化加速实战IQ2_KS / IQ2_K / IQ2_K_R4 的 Prompt Processing 优化解析ik_llama.cpp CUDA 量化加速实战IQ2_KS / IQ2_K / IQ2_K_R4 的 Prompt Processing 优化解析 导读 本人工智能大模型推理引擎本地部署模型量化模型优化深入解析 Q6_0 的 MMQ CUDA Kernel从 ik_llama.cpp PR 114 的失败尝试到完整实现深入解析 Q6_0 的 MMQ CUDA Kernel从 ik_llama.cpp PR 114 的失败尝试到完整实现 导读 本文以 ik_llama.cpp人工智能大模型推理引擎本地部署模型量化模型优化创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考
返回列表