QServe：用于高效 LLM 服务的 W4A8KV4 量化与系统协同设计

Posted on

QServe：用于高效 LLM 服务的 W4A8KV4 量化与系统协同设计

ArXiv ID: 2405.04532
作者: Yujun Lin, Haotian Tang, Shang Yang, Zhekai Zhang, Guangxuan Xiao, Chuang Gan, Song Han
机构: MIT Han Lab, NVIDIA, MIT-IBM Watson AI Lab
发布日期: 2024-05-07

摘要

量化可以加速大语言模型推理。在 INT8 量化之外，研究社区正在积极探索更低精度如 INT4。然而，现有 INT4 量化技术仅能加速低批量、边缘 LLM 推理，无法在大批量、基于云的 LLM 服务中提供性能提升。本文发现了一个关键问题：现有的 INT4 量化方法在 GPU 上对权重或部分和进行反量化时存在显著的运行时开销（20-90%）。为了解决这一挑战，本文引入 QoQ，一种W4A8KV4量化算法，采用 4 位权重、8 位激活和 4 位 KV 缓存。QServe 在 L40S 上实现比 TensorRT-LLM 在 A100 上更高的吞吐量，服务成本降低3 倍。

问题背景

INT4 量化的困境

1
2
3
4
5
6
7
8
9
10
11
12
13
INT4 量化的性能悖论：

理论加速：
- FP16 → INT4: 4 倍计算密度
- 内存带宽需求：减少 4 倍
- 预期加速：2-4 倍

实际表现（云服务等高吞吐场景）：
- 反量化开销：20-90%
- 实际加速：1.1-1.5 倍
- 某些场景：甚至更慢！

为什么？

反量化开销分析

问题根源：

1
2
3
4
5
6
7
8
9
10
11
12
# 典型 INT4 推理流程
def int4_inference(weights_int4, activations_fp16):
    # 步骤 1: 反量化权重（运行时开销）
    weights_fp16 = dequantize(weights_int4)  # ← 20-30% 时间

    # 步骤 2: 矩阵乘法
    output_fp16 = matmul(weights_fp16, activations_fp16)

    # 步骤 3: 反量化部分和（额外开销）
    output_fp16 = dequantize_partial(output_int32)  # ← 10-20% 时间

    return output_fp16

关键洞察：

  • 反量化操作在 GPU 上是内存密集型
  • 高批量服务场景下，反量化成为瓶颈
  • 需要系统级协同设计，而非单纯算法优化

QServe 方法

整体架构

1
2
3
4
5
6
7
8
9
10
11
12
13
14
┌─────────────────────────────────────────────────────────┐
│                    QServe System                         │
│                                                         │
│  ┌───────────────┐    ┌───────────────┐                │
│  │  QoQ 量化     │    │  系统优化     │                │
│  │  (算法层)     │    │  (系统层)     │                │
│  │               │    │               │                │
│  │ • W4 权重     │    │ • 权重重排序   │                │
│  │ • A8 激活     │    │ • 寄存器并行   │                │
│  │ • KV4 缓存    │    │ • 融合注意力   │                │
│  └───────────────┘    └───────────────┘                │
│                                                         │
│              协同设计 = 算法 × 系统                      │
└─────────────────────────────────────────────────────────┘

W4A8KV4 量化方案

QoQ (Quattuor-Octo-Quattuor)：

组件 精度 理由
Weight 4-bit 权重是内存瓶颈，可低位量化
Activation 8-bit 保持计算质量，避免反量化开销
KV Cache 4-bit KV 缓存是长序列内存瓶颈

内存节省：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
以 LLaMA-7B 为例：

FP16:
- 权重：14 GB
- KV Cache (4K): ~2 GB
- 总计：16 GB

W4A8KV4:
- 权重：3.5 GB (4 倍节省)
- KV Cache (4K): ~0.5 GB (4 倍节省)
- 激活：~2 GB (保持 8-bit)
- 总计：6 GB

内存节省：62.5%

渐进量化 (Progressive Quantization)

核心思想：将量化操作分散到多个阶段，避免运行时集中开销

1
2
3
4
5
6
7
8
9
10
11
# 传统 INT4：运行时反量化
def traditional_int4_gemm(W_int4, A_fp16):
    W_fp16 = dequantize(W_int4)  # 集中开销
    return matmul(W_fp16, A_fp16)

# QoQ 渐进量化：预计算 + 融合
def qoq_gemm(W_int4, A_int8):
    # 权重已预转换为特殊格式
    # 激活保持 INT8 到计算最后
    # 使用 INT8 Tensor Core
    return int8_gemm(W_int4, A_int8)

量化感知训练替代方案：

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
def smoothquant_based_quantization(model, calibration_data):
    """
    基于 SmoothQuant 的渐进量化
    无需训练，只需校准数据
    """
    # 步骤 1: 收集激活统计
    activation_stats = collect_activation_stats(model, calibration_data)

    # 步骤 2: 计算平滑因子
    smooth_factors = compute_smooth_factors(
        model.weights,
        activation_stats,
        alpha=0.5
    )

    # 步骤 3: 应用平滑并量化
    for layer in model.layers:
        # 平滑激活难度
        layer.smooth(smooth_factors)
        # 量化权重为 W4
        layer.quantize_weights(bits=4)
        # 量化 KV 缓存为 KV4
        layer.quantize_kv_cache(bits=4)

    return model

SmoothAttention 技术

问题：4-bit KV 量化导致注意力精度下降

解决方案：结合 SmoothQuant 思想平滑 KV 缓存

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
def smooth_attention_kv_quantize(K, V, smooth_factor):
    """
    SmoothAttention: 平滑 KV 缓存后量化

    Args:
        K: 键向量 [seq_len, dim]
        V: 值向量 [seq_len, dim]
        smooth_factor: 平滑因子 [dim]
    """
    # 平滑变换
    K_smooth = K / smooth_factor
    V_smooth = V * smooth_factor

    # 量化到 4-bit
    K_int4 = quantize_per_channel(K_smooth, bits=4)
    V_int4 = quantize_per_channel(V_smooth, bits=4)

    return K_int4, V_int4

系统级优化

1. 计算感知的权重重排序

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
def compute_aware_weight_reorder(weights_int4):
    """
    重排序权重以优化内存访问模式

    目标：
    - 提高内存合并效率
    - 减少 bank conflict
    - 优化 L2 缓存命中率
    """
    # 分析权重访问模式
    access_pattern = analyze_access_pattern()

    # 重排序以匹配 GPU 内存层次
    reordered = reorder_for_memory_hierarchy(
        weights_int4,
        access_pattern
    )

    return reordered

2. 寄存器级并行

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
__global__ void qoq_gemm_kernel(W_int4, A_int8, C_fp16) {
    // 寄存器级并行隐藏反量化延迟
    __shared__ float shared_W[Block_Size];
    __shared__ float shared_A[Block_Size];

    // 预取到寄存器
    float4 reg_W = preload_to_registers(W_int4);
    float4 reg_A = preload_to_registers(A_int8);

    // 计算与数据加载重叠
    #pragma unroll
    for (int i = 0; i < K; i += 4) {
        // 计算当前块
        float acc = mad(reg_W, reg_A);

        // 预取下一块（异步）
        async_load_next_block();

        C[tid] = acc;
    }
}

3. 融合注意力内核

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
20
21
22
23
24
25
26
27
28
29
30
31
32
33
34
35
__global__ void fused_attention_kv4_kernel(Q, K_int4, V_int4, O) {
    // 融合 QK^T、Softmax、PV 计算
    // 减少全局内存访问

    __shared__ float shared_Q[Tile_Size];
    __shared__ float shared_K[Tile_Size];
    __shared__ float shared_V[Tile_Size];

    // 加载 Q
    load_Q_to_shared(Q, shared_Q);
    __syncthreads();

    float acc[Head_Dim];
    float max_val = -INF;
    float sum_exp = 0;

    // 流式处理 KV 块
    for (int kv_block = 0; kv_block < num_kv_blocks; kv_block++) {
        // 加载并反量化 K
        load_and_dequant_K(K_int4, shared_K);

        // 计算注意力分数
        float score = dot_product(shared_Q, shared_K);

        // Online Softmax
        update_max_and_sum(max_val, sum_exp, score);

        // 累加输出
        load_and_dequant_V(V_int4, shared_V);
        acc += exp(score - max_val) * shared_V;
    }

    // 归一化输出
    O[tid] = acc / sum_exp;
}

实验结果

推理加速

LLaMA-3-8B

GPU Baseline QServe 加速比
A100 TensorRT-LLM FP16 QServe W4A8KV4 1.2×
L40S TensorRT-LLM FP16 QServe W4A8KV4 1.4×
L4 TensorRT-LLM FP16 QServe W4A8KV4 1.5×

Qwen1.5-72B

GPU Baseline QServe 加速比
A100 TensorRT-LLM FP16 QServe W4A8KV4 2.4×
L40S TensorRT-LLM FP16 QServe W4A8KV4 3.5×

关键发现：

  • 大模型加速更明显（72B 达 3.5×）
  • L40S（消费级）受益更多
  • 高批量场景优势更大

跨 GPU 对比

1
2
3
4
5
6
7
8
9
吞吐量对比 (tokens/s，Qwen-72B，batch=32)：

GPU          | TensorRT-LLM | QServe | 提升
-------------|--------------|--------|-----
A100 80GB    | 125          | 300    | 2.4×
L40S 48GB    | 85           | 298    | 3.5×
L4 24GB      | 45           | 68     | 1.5×

关键：QServe 在 L40S 上超越 A100 + TensorRT-LLM

成本效益

配置 每小时成本 吞吐量 每 token 成本
A100 + TensorRT-LLM $4.00 125 tok/s $0.032
L40S + QServe $1.20 298 tok/s $0.004

成本降低: 8× 每 token 成本，3× 总体拥有成本

精度评估

基准 FP16 QServe 差异
WikiText2 PPL 12.45 12.52 +0.56%
MMLU 65.3% 64.8% -0.5%
GSM8K 45.2% 44.8% -0.4%
HumanEval 32.5% 32.1% -0.4%

结论：精度损失可忽略不计（<1%）

批量大小影响

1
2
3
4
5
6
7
8
9
10
11
12
13
加速比 vs 批量大小 (Qwen-72B, L40S):

加速比
  │
  │           ● 3.5×
  │       ●
  │   ●
  │ ●
  └─────────────────────────
    1    8    16   32   64
            批量大小

高批量场景受益更多

消融实验

各组件贡献

配置 A100 吞吐量 相对性能
完整 QServe 300 tok/s 100%
- 权重重排序 265 tok/s 88%
- 寄存器并行 248 tok/s 83%
- 融合注意力 220 tok/s 73%
- W4A8 (只用 W4) 180 tok/s 60%
基线 (FP16) 125 tok/s 42%

量化位宽选择

配置 精度 吞吐量 推荐度
W4A8KV4 64.8% 300 tok/s ⭐⭐⭐⭐⭐
W4A4KV4 62.1% 320 tok/s ⭐⭐⭐
W4A16KV4 65.1% 250 tok/s ⭐⭐⭐⭐
W8A8KV8 65.2% 180 tok/s ⭐⭐

实践指南

使用 QServe

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
# 安装
pip install qserve

# 量化模型
qserve-quantize \
    --model meta-llama/Llama-2-7b \
    --output llama-2-7b-qserve \
    --bits 4 \
    --calibration-data pile:512

# 运行服务
qserve-serve \
    --model llama-2-7b-qserve \
    --host 0.0.0.0 \
    --port 8000 \
    --batch-size 32

部署配置

1
2
3
4
5
6
7
8
9
10
11
12
13
14
15
16
17
18
19
# qserve_config.yaml
model:
  path: "llama-2-7b-qserve"

quantization:
  weight_bits: 4
  activation_bits: 8
  kv_cache_bits: 4
  group_size: 128

serving:
  max_batch_size: 32
  max_sequence_length: 4096
  gpu_memory_utilization: 0.9

optimization:
  enable_weight_reorder: true
  enable_register_parallel: true
  enable_fused_attention: true

总结

QServe 通过算法 - 系统协同设计，实现了高效 LLM 服务：

核心贡献：

  1. 发现并解决了反量化开销问题
  2. 提出 W4A8KV4 量化方案
  3. 系统级优化（重排序、并行、融合）
  4. 在消费级 GPU 上超越数据中心 GPU

实际影响：

  • 3 倍服务成本降低
  • 使大模型服务更经济
  • 开源实现支持多种模型

资源

评分: 4.5/5.0 ⭐⭐⭐⭐⭐

推荐度: 强烈推荐。系统协同设计典范，成本效益显著。

© 2026 Generative AI Discovery All Rights Reserved.
Theme by hiero