【设Q找A，避免陷入细节陷阱】
llama-simple -m ./models/Qwen3-1.7B-Q4_K_M.gguf -n 30 "What is the result of 5/0 in math?"
@ backend 如何进行计算 graph-compute ？
ggml_status llama_context::graph_compute(ggml_cgraph * gf, bool batched) {
// 设置线程池
int n_threads = batched ? cparams.n_threads_batch : cparams.n_threads;
ggml_threadpool_t tp = batched ? threadpool_batch : threadpool;
if (backend_cpu != nullptr) {
auto * reg = ggml_backend_dev_backend_reg(ggml_backend_get_device(backend_cpu));
auto * set_threadpool_fn =
(decltype(ggml_backend_cpu_set_threadpool) *) ggml_backend_reg_get_proc_address(reg,
"ggml_backend_cpu_set_threadpool");
set_threadpool_fn(backend_cpu, tp);
}
for (const auto & set_n_threads_fn : set_n_threads_fns) {
set_n_threads_fn.second(set_n_threads_fn.first, n_threads);
}
// 异步执行计算
// sched（ggml_backend_sched）根据节点的后端（CPU 或 CUDA）分配任务
auto status = ggml_backend_sched_graph_compute_async(sched.get(), gf);
if (status != GGML_STATUS_SUCCESS) {
LLAMA_LOG_ERROR("%s: ggml_backend_sched_graph_compute_async failed with error %d\n", __func__, status);
}
return status;
}
进入异步计算
// sched 后端调度器， 理任务在 CPU、CUDA 等后端的分配
enum ggml_status ggml_backend_sched_graph_compute_async(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
// 清除旧状态
if (!sched->is_reset && !sched->is_alloc) {
ggml_backend_sched_reset(sched);
}
// 为计算图的节点和张量分配内存
// 遍历 graph 的节点，确定每个节点的内存需求，并为它们分配内存
if (!sched->is_alloc) {
if (!ggml_backend_sched_alloc_graph(sched, graph)) {
return GGML_STATUS_ALLOC_FAILED;
}
}
// 执行计算图的节点，完成前向传播
// 调度器将计算图分成多个子任务（splits），根据后端（CPU/CUDA）和依赖关系分配。
// 每个节点调用对应 backend 的 kernel
return ggml_backend_sched_compute_splits(sched) /*包含：*/ {
ggml_backend_graph_compute_async() {
ggml_backend_cpu_graph_compute() {
ggml_graph_compute()
}
} /* 或者 cuda backend (通过gdb定位到)： */ {
ggml_backend_cuda_graph_compute() {
......
evaluate_and_capture_cuda_graph() { // 循环 1097，对于当时 case
......
ggml_cuda_compute_forward(*cuda_ctx, node); {
// 这里是实际的 kernel 调用
// 循环分别执行的 op 是：
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); {
ggml_cuda_mul_mat_vec_f(ctx, src0, src1, nullptr, dst); // or
ggml_cuda_mul_mat_f(ctx, src0, src1, nullptr, dst); // or
ggml_cuda_mul_mat_vec_q(ctx, src0, src1, nullptr, dst); // or
ggml_cuda_mul_mat_batched_cublas(ctx, src0, src1, dst); // or
ggml_cuda_mul_mat_q(ctx, src0, src1, nullptr, dst);{
......
quantize_mmq_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type,
ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); {
// gpu 上的 mul_mat kernel 调用
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4>
<<<num_blocks, block_size, 0, stream>>>
(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
}
}
// 。。。。。。
}
ggml_cuda_op_rope(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_rope(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_set_rows(ctx, dst);
ggml_cuda_op_set_rows(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_soft_max(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_dup(ctx, dst); // GGML_OP_CONT
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_add(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_swiglu(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_add(ctx, dst);
ggml_cuda_mul_mat(ctx, dst->src[0], dst->src[1], dst); // 可能是量化 kernel，或是计算 kernel
ggml_cuda_op_rope(ctx, dst);
// ......
// 对于当前case 会有接近 1000 个cuda kernel 的调用
}
}
}
}
};
}
quantize_mmq_q8_1_cuda 做了什么：
// 在 CUDA 设备上将浮点数数据 x 量化为 Q8_1 格式（8 位整数加缩放因子），并存储到输出缓冲区 vy
// ne00, s01, s02, s03：输入张量的维度和步幅（strides），描述数据布局
// ne0, ne1, ne2, ne3：输出张量的维度
void quantize_mmq_q8_1_cuda(
const float * x, const int32_t * ids, void * vy, const ggml_type type_src0,
const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) {
// 避免未对齐的内存访问，提高 CUDA 性能
// 输入张量的第一个维度（ne00）必须是 4 的倍数，保证数据对齐（便于 SIMD 或 CUDA 线程处理）
GGML_ASSERT(ne00 % 4 == 0);
// 输出张量的第一个维度（ne0）必须是 4*QK8_1 的倍数，QK8_1 是 Q8_1 量化的块大小: 32，确保量化块对齐
GGML_ASSERT(ne0 % (4*QK8_1) == 0);
// CUDA_QUANTIZE_BLOCK_SIZE_MMQ ：每个线程块处理的元素数 128
// 根据输出张量维度 ne0 计算 block y 轴
// 因子 4 和 QK8_1 反映量化块的内存对齐
const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ);
//
const dim3 num_blocks(ne1, block_num_y, ne2*ne3);
// 每个 block 包含 CUDA_QUANTIZE_BLOCK_SIZE_MMQ 个线程
const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1);
// 返回数据布局类型，表示不同的 Q8_1 量化格式
switch (mmq_get_q8_1_ds_layout(type_src0)) {
case MMQ_Q8_1_DS_LAYOUT_D4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D4>
<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
break;
case MMQ_Q8_1_DS_LAYOUT_DS4:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_DS4>
<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
break;
case MMQ_Q8_1_DS_LAYOUT_D2S6:
quantize_mmq_q8_1<MMQ_Q8_1_DS_LAYOUT_D2S6>
<<<num_blocks, block_size, 0, stream>>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2);
break;
default:
GGML_ABORT("fatal error");
break;
}
}
进入 cuda 量化 kernel
llama.cpp/ggml/src/ggml-cuda/quantize.cu
template <mmq_q8_1_ds_layout ds_layout>
static __global__ void quantize_mmq_q8_1(
const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy,
const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03,
const int64_t ne0, const int ne1, const int ne2) {
// D2S6 表示“2 维数据 + 6 位缩放因子变体”
//每个量化块包含 64 个量化值
constexpr int vals_per_scale = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 64 : 32;
// 对于需要求 sum 的 D2S6 布局，每16个一个sum
constexpr int vals_per_sum = ds_layout == MMQ_Q8_1_DS_LAYOUT_D2S6 ? 16 : 32;
// 当前线程处理的数据元素索引，将 cuda grid 映射到输出 tensor 的维度上，确保每个线程处理正确元素
// 每个线程处理 4 个元素
const int64_t i0 = ((int64_t)blockDim.x*blockIdx.y + threadIdx.x)*4;
if (i0 >= ne0) {
return;
}
const int64_t i1 = blockIdx.x;
const int64_t i2 = blockIdx.z % ne2;
const int64_t i3 = blockIdx.z / ne2;
const int64_t i00 = i0;
const int64_t i01 = ids ? ids[i1] : i1;
const int64_t i02 = i2;
const int64_t i03 = i3;
// 从输入 x 加载 4 个浮点数作为 float4 向量
const float4 * x4 = (const float4 *) x;
block_q8_1_mmq * y = (block_q8_1_mmq *) vy;
const int64_t ib0 = blockIdx.z*((int64_t)gridDim.x*gridDim.y*blockDim.x/QK8_1); // first block of channel
const int64_t ib = ib0 + (i0 / (4*QK8_1))*ne1 + blockIdx.x; // block index in channel
const int64_t iqs = i0 % (4*QK8_1); // quant index in block
const float4 xi = i0 < ne00 ? x4[(i03*s03 + i02*s02 + i01*s01 + i00)/4] : make_float4(0.0f, 0.0f, 0.0f, 0.0f);
// 计算 当前线程一个 float4，即 4个浮点数的最大绝对值 amax
float amax = fabsf(xi.x);
amax = fmaxf(amax, fabsf(xi.y));
amax = fmaxf(amax, fabsf(xi.z));
amax = fmaxf(amax, fabsf(xi.w));
// 计算 warp 中计算 32 个线程之间的 amax，即 块级 amax
#pragma unroll
for (int offset = vals_per_scale/8; offset > 0; offset >>= 1) {
amax = fmaxf(amax, __shfl_xor_sync(0xFFFFFFFF, amax, offset, WARP_SIZE));
}
// 某些 layout 需要这个sum 用于额外的放缩和校验，D4 layout 不需要
float sum;
if (ds_layout != MMQ_Q8_1_DS_LAYOUT_D4) {
sum = xi.x + xi.y + xi.z + xi.w;
#pragma unroll
for (int offset = vals_per_sum/8; offset > 0; offset >>= 1) {
sum += __shfl_xor_sync(0xFFFFFFFF, sum, offset, WARP_SIZE);
}
}
/* 与 Q8_1 量化逻辑一致:
1. 找到块内绝对值的最大值（`amax`）。
2. 计算缩放因子：`scale = amax / 127`（`int8_t` 最大值）
3. 量化为：`q[i] = round(x[i] / scale)`
等价于 `q[i] = round(x[i] * (127.0f / amax))`
*/
const float d_inv = 127.0f / amax;
char4 q;
q.x = roundf(xi.x*d_inv);
q.y = roundf(xi.y*d_inv);
q.z = roundf(xi.z*d_inv);
q.w = roundf(xi.w*d_inv);
// 将 y[ib].qs（量化值数组）转换为 char4 指针，便于高效写入
char4 * yqs4 = (char4 *) y[ib].qs;
// 将4个量化值写入正确位置
yqs4[iqs/4] = q;
// 针对其他layout的特殊处理, 见原函数
......
}
Q8_1 量化结果类型：
struct block_q8_1_mmq {
// y 的浮点数据被转换为可以直接复制到共享内存的连续块布局。
// y 的浮点数据首先按每 128 个值分为一组。
// 这些块随后被视为单独的数据值并进行转置。
//
// 为避免共享内存的 bank 冲突，每个块填充 16 字节。
// 填充空间还用于存储块的缩放因子/部分和。
// 缩放因子与量化数据相乘后等于未量化的值。
// 部分和通过对子组值（量化前）求和得到，仅用于性能优化。
//
// 具体存储的数据取决于 x 的数据类型。
union {
float d4[4]; // 每 32 个值 1 个 32 位缩放因子，存储为 d0,d1,d2,d3
half2 ds4[4]; // 每 32 个值 1 个 16 位缩放因子 + 1 个 16 位部分和，存储为 d0,s0,d1,s1,d2,s2,d3,s3
half d2s6[8]; // 每 64 个值 1 个 16 位缩放因子 + 前 96 个值每 16 个值 1 个 16 位部分和，
// 存储为 d0,d1,s1,s2,s3,s4,s5
};
int8_t qs[4*QK8_1]; // 128 个值量化为 8 位整数
};
@ Q8_1 量化
Q8_1 是一种 8 位整数量化格式，将浮点数数据量化为 8 位有符号整数（int8_t），并为每个数据块附加一个 32 位浮点数缩放因子（scale）。量化对象一般是 Weights。
每个数据块包含固定数量的元素，（ggml-common.h 定义 QK8_1=32）。对于 D4 layout 每个块存储：
量化值：类型为
int8_t数组（范围从 -128 到 127，对称量化），元素 32 个，表示压缩后的数据。共 32 bytes缩放因子：1 个
float值 32 bit，用于恢复原始浮点数。共 4 byte
所以每个数据块，大小是 36 bytes。
Q8_1 量化过程
输入浮点数数组 x, 计算：
将数据分组为块, 每块 32 个元素
找到块内绝对值的最大值
max_abs计算缩放因子：
scale = max_abs / 127（int8_t最大值）每个浮点数
x[i]量化为：q[i] = round(x[i] / scale)（int8_t），是量化核心
输出：q（int8_t 数组）和 scale（float）
恢复过程
浮点数恢复：x[i] ≈ q[i] * scale, 就是量化步骤的逆向计算
误差：由于 8 位精度限制，量化会引入小误差，但压缩比高
实例 数据块是8
x = [2.5, -1.8, 3.2, 0.5, -2.7, 1.2, -0.9, 2.1]
找到块内绝对值的最大值 max_abs：max_abs_1 = 3.2（ x[2] = 3.2）
计算缩放因子 scale = max_abs / 127 = 3.2 / 127 ≈ 0.025196850
量化每个浮点数：根据 q[i] = round(x[i] / scale)，结果为 int8_t([-128, 127])。
- q[0] = round(2.5 / 0.025196850) = round(99.250) = 99
- q[1] = round(-1.8 / 0.025196850) = round(-71.438) = -71
- q[2] = round(3.2 / 0.025196850) = round(126.938) = 127
- q[3] = round(0.5 / 0.025196850) = round(19.850) = 20
- q[4] = round(-2.7 / 0.025196850) = round(-107.156) = -107
- q[5] = round(1.2 / 0.025196850) = round(47.638) = 48
- q[6] = round(-0.9 / 0.025196850) = round(-35.719) = -36
- q[7] = round(2.1 / 0.025196850) = round(83.363) = 83
得到量化结果 int8：q = [99, -71, 127, 20, -107, 48, -36, 83 ] 和放缩因子 scale = 0.025196850
- input是 x = [2.5, -1.8, 3.2, 0.5, -2.7, 1.2, -0.9, 2.1]
- output是 q = [99, -71, 127, 20, -107, 48, -36, 83]
验证最大绝对值 x[2] = 3.2 → q[2] = 127，验证 3.2 / 0.025196850 ≈ 127
布局
见 block_q8_1_mmq 定义，它是线性量化的特例。它是一种针对 Transformer 模型优化的量化技术。
衡量量化的性能
每值存储：每个原始 float 在量化后平均占用的存储空间（以 Byte 为单位）。一个数据块有32个元素，量化后 32 个元素占用 36 Bytes， 故每一个 float 占用空间：
36 (bytes) ÷ 32 (个float) = 1.125 字节/值。相比原始 float（4 字节），显著减少存储需求。压缩比：表示量化后数据存储大小与原始数据大小的比率，用百分比表示。压缩比 = (量化后每值存储 ÷ 原始每值存储) =
1.125 ÷ 4 ≈ 0.281 = 28.1%。节省比例：1 - 28.1% = 71.9%，表示节省了约 71.9% 的存储空间。
含有 zero-point 的量化
步骤：
确定数据范围：找到输入浮点数数组的最小值 min_x 和最大值 max_x。
计算缩放因子和零点：
- 缩放因子：scale = (max_x - min_x) / (2^n - 1)，其中 n 是量化位数（如 8 位，2^8 - 1 = 255）。
- 零点 zero-point：zero_point = round(-min_x / scale)，确保量化范围覆盖数据的实际分布。
量化：每个浮点数 x[i] 量化为 q[i] = round(x[i] / scale + zero_point)，并裁剪到目标整数范围（如 [0, 255] 对于 uint8）。
输出：量化后的整数数组 q（如 uint8）和元信息（scale 和 zero_point），用于反量化。
zero-point 作用：
- uint8 运算比 int8 更高效，因为无符号整数运算无需处理符号位。zero_point 确保量化值非负，适配硬件优化的整数运算单元。
- 适配非对称数据分布，映射到整数范围（如 [0, 255]）
- 充分利用整数范围，最小化量化误差
更多量化内容
尝试 llama-quantize 工具：
将 f16 精度的模型量化为 Q4_K_M 量化模型，并使用它：
./llama-quantize ./models/mymodel/ggml-model-f16.gguf ./models/mymodel/ggml-model-Q4_K_M.gguf Q4_K_M
./llama-cli -m ./models/mymodel/ggml-model-Q4_K_M.gguf -cnv -p "You are a helpful assistant"
工具用法详见 llama.cpp/tools/quantize/README.md
KAQ：量化的目的是主要是减少存储压力，计算时，还是需要解量化到原始精度？
一个量化了的 LLM，它的权重总是跟随一个 scale 值，用于恢复到原始数值精度。在推理时，权值加载到内存后，对于每一个值，在内存中先解量化到原始数值精度，然后计算。所以量化的目的主要是减少存储压力。
推理的计算过程中，量化和非量化对比，前者在内存中计算时，多了一步 W_fp = W_int8 * scale 解量化。但是存储空间大大减小，同时降低了贷款需求。
KAQ：量化后的GGUF文件中一定有 scale 值的？
通过解析GGUF文件验证：
import gguf
reader = gguf.GGUFReader("Qwen3-1.7B-Q4_K_M.gguf") # 替换为你的 GGUF 文件路径
for tensor in reader.tensors:
if "weight" in tensor.name: # 过滤权重 tensor
print(f"Tensor: {tensor.name}")
print(f"Shape: {tensor.shape}")
print(f"Quantization type: {tensor.tensor_type}") # 应为 GGUF_TYPE_Q4_K
# Q4_K_M 的 scale 值在 tensor.data 中（每 32 或 64 权重一个 float scale）
# 假设每 32 个权重一块，scale 在数据块中
scale_size = tensor.data.size // (32 * 4 // 8 + 4) # 4 位权重 + float scale
scales = tensor.data.view('f4')[:scale_size] # 提取 scale（float32）
print(f"First 5 scale values: {scales[:5]}")
每个 tensor 的 data 字段包含量化权重（4 位）和 scale（float32）