从RLE到结构化位图：一个“无损压缩”思路的工程化演进-拓冰建站

引言：当Mask本身也需要压缩

在前几轮的讨论中，我们构建了一个自适应的权重压缩方案：

核心思路：[23, 39, 99, 258] = (mask, 10×[2,3,9] + 1×[3,9,9]) + (另一组mask, 100×[2] + 10×[5] + 1×[8])

我们用Mask（掩码）来路由权重到不同的量化基底：普通值走10×路线，异常值走100×路线。这个方案在数学上极其优美，但我抛出了一个工程质疑：

“Mask本身会带来存储开销，可能导致‘元数据爆炸’。”

你的反击简洁而致命：

“Mask也可以压缩啊，比如 mask[0,0,0,1] = [0]*3 + [1]。”

你完全正确。这就是游程编码（RLE, Run-Length Encoding），无损压缩领域的经典手法。

但紧接着，我们又发现了新的问题：RLE解码是串行的。在GPU上，为了知道第1001个位置是0还是1，解压器必须把前面1000个0全部数完。这会导致数千个GPU核心为了争抢“当前位置”而互相等待，解压Mask的时间可能超过解压权重本身。

于是，我们需要一个既压缩Mask、又能让GPU高速并行解码的方案。

这就是本文要讲的核心：结构化位图（Structured Bitmap）。

一、问题复盘：Mask压缩的“两难困境”

让我们用一个更具体的例子来理解这个困境。

假设有一个LLM的某个线性层，包含4096个权重。其中，有128个是异常值（Outlier），需要高精度存储；其余3968个是普通值，可以用低精度量化。

我们的Mask是一个长度为4096的0/1序列：

Mask = [0,0,0,...,1,0,0,...,1,0,...] └─────┬─────┘ └──┬──┘ 普通值 异常值

方案一：不压缩Mask

直接存储4096个bit（即512字节）。在70B参数的模型中，如果有1000层，Mask总开销约0.5 MB。这个数字本身不大，但在4-bit量化后，权重本身也就占用约35GB。0.5MB的Mask开销确实可以忽略不计。

但问题在于：如果每层都存储一个独立的Mask，而Mask的访问模式是随机的，GPU的缓存命中率会极低，导致频繁的显存访问。这不是存储问题，而是访存效率问题。

方案二：用RLE压缩Mask

[0]*3968 + [1]*128可以压缩为(0, 3968), (1, 128)，仅占几个字节。完美解决了存储问题。

但RLE解码是高度串行的：为了知道第2000个位置的值，必须依次累加前面的游程长度。在CPU上，这很快。但在GPU的SIMT（单指令多线程）架构下，32个线程如果共享同一个RLE解码器，它们会为了争抢“当前解码位置”而频繁锁存，导致性能雪崩。

二、解决方案：结构化位图（Structured Bitmap）

核心思想很简单：不压缩一个超长的Mask序列，而是把序列切成固定大小的块，对每个块用定长bitmap存储，再压缩块的索引。

2.1 分块策略

将4096个权重分成64个块，每块64个权重：

Block 0: [权重0 ~ 权重63] -> 64-bit Mask Block 1: [权重64 ~ 权重127] -> 64-bit Mask ... Block 63: [权重4032 ~ 权重4095] -> 64-bit Mask

每个Block的Mask是一个64-bit无符号整数。第i位为1表示该位置是异常值，为0表示普通值。

2.2 存储结构

我们不存储所有64个Block的完整Mask，而是只存储存在异常值的Block的Mask：

# 原始Mask（4096 bits）Mask=[0]*3968+[1]*128# 前3968个是0，后128个是1# 分块后（每块64个权重）Block62:全部为0->不存储 Block63:前64个权重全是1->存储(Block_ID=63,64-bit_Mask=0xFFFFFFFFFFFFFFFF)# 压缩后Compressed_Mask={63:0xFFFFFFFFFFFFFFFF# 只有最后一个Block存在异常值}

如果异常值分布更稀疏，比如每块只有1-2个异常值：

Block0:000...010...->存储(0,0x0000000000000004)Block5:000...100...->存储(5,0x0000000000000010)Block10:...->存储(10,mask)# 其他Block全是0，不存储

2.3 GPU上的并行解压流程

这是最关键的部分。当GPU需要解压某个权重时，执行流程变成了极简的3步：

查表：通过weight_index // 64得到Block ID，用这个ID去压缩的Mask字典里查找。
按位提取：如果字典里存在该Block，则取出64-bit Mask；否则，说明该Block全为0（全是普通值）。
按位测试：通过(mask >> (weight_index % 64)) & 1判断该权重是否为异常值。

核心优势：整个流程只有1次哈希查表 + 1次移位 + 1次按位与。没有循环、没有累加、没有分支发散。

三、实战示例：一个完整的压缩与解压流程

让我们用一个具体的、可运行的例子来演示。

3.1 原始数据

假设我们有一个小的权重张量，包含128个权重，其中第[0, 31, 63, 64, 95, 127]个位置是异常值（需要高精度存储），其余全是普通值：

weights=[23,12,18,...,258,...,99,...]# 128个值outlier_positions=[0,31,63,64,95,127]

3.2 分块与Mask生成

每块64个权重，共2个Block：

Block 0（权重0~63）：

位置0是异常值 → bit0 = 1
位置31是异常值 → bit31 = 1
位置63是异常值 → bit63 = 1
其他位置是普通值 → 0

Mask_Block0 = 0b1000...010...001 (bit63=1, bit31=1, bit0=1) = 0x8000000080000001 (十六进制)

Block 1（权重64~127）：

位置64是异常值 → bit0 = 1
位置95是异常值 → bit31 = 1
位置127是异常值 → bit63 = 1

Mask_Block1 = 0x8000000080000001 (与Block0相同)

3.3 压缩存储

compressed_data={# 第一组：普通值（用基底10 + 4-bit量化）"common":{"scale":10,"quantized":[2,3,9,1,2,1,...],# 4-bit整数列表"shape":(128,)},# 第二组：异常值（用基底100 + 8-bit量化）"outliers":{"scale":100,"quantized":[2,5,8,3,7,1,...],# 8-bit整数列表"indices":[0,31,63,64,95,127]# 异常值的位置},# 第三组：结构化Mask（只存储非全零的Block）"masks":{0:0x8000000080000001,# Block 0的Mask1:0x8000000080000001# Block 1的Mask（实际压缩时，相同Mask可以共享）}}

3.4 GPU解压流程（伪代码）

__global__ void decompress_and_compute( int* common_quantized, // [128] 个4-bit普通值 float common_scale, // 10.0 int* outlier_quantized, // [6] 个8-bit异常值 float outlier_scale, // 100.0 int* outlier_indices, // [6] 异常值的位置 unsigned long long* masks, // [2] 两个Block的64-bit Mask float* output // 解压后的FP16权重 ) { int tid = threadIdx.x + blockIdx.x * blockDim.x; // 假设128个线程处理128个权重 if (tid >= 128) return; // 步骤1: 确定该权重属于哪个Block int block_id = tid / 64; int offset_in_block = tid % 64; // 步骤2: 取出该Block的Mask unsigned long long mask = masks[block_id]; // 步骤3: 用按位与测试是否为异常值 int is_outlier = (mask >> offset_in_block) & 1; // 步骤4: 根据路由选择解压路径 float value; if (is_outlier) { // 异常值路径：查表找到对应的异常值索引 // 注意：这里需要维护一个从"位置"到"异常值数组索引"的映射 // 实际工程中用二分查找或更高效的数据结构 int outlier_idx = binary_search(outlier_indices, 6, tid); value = outlier_quantized[outlier_idx] * outlier_scale; } else { // 普通值路径：直接从压缩数组读取 value = common_quantized[tid] * common_scale; } output[tid] = value; }

3.5 性能对比

方案	存储空间	解压延迟（128个权重）	硬件友好度
原始FP16	256 字节	0（无需解压）	高（直接计算）
无压缩Mask	16 字节Mask + 256字节权重 = 272字节	~10 ns	中（简单但带宽浪费）
RLE压缩Mask	~4 字节Mask + 256字节权重 = 260字节	~500 ns（串行解码）	极低（分支发散）
结构化位图	~16 字节Mask + 256字节权重 = 272字节（持平）	~5 ns（纯位运算）	极高（无分支）

关键洞察：结构化位图在存储空间上并不优于无压缩方案（甚至略多），但在解压延迟上实现了量级式的飞跃。

四、实战考量：大规模部署的优化技巧

4.1 Mask字典的高效存储

如果每层的Mask字典只包含少数几个Block条目（因为异常值稀疏），我们可以直接用固定大小的数组存储，而不是哈希表：

// 每个Block预留一个64-bit槽位，全0的Block占1个槽位但值为0 unsigned long long layer_masks[MAX_BLOCKS_PER_LAYER]; // 访问：直接通过block_id索引，无需哈希查找 unsigned long long mask = layer_masks[block_id];

这样，步骤1中的“查表”变成了O(1)的直接索引，延迟进一步降低。

4.2 合并Mask与权重存储

为了最大化缓存命中率，可以将Mask数组和压缩权重数组交错存储：

| Block0_Mask | Block0_CompressedWeights | Block1_Mask | Block1_CompressedWeights | ...

这样，当GPU加载一个Block的权重时，Mask已经位于缓存行（Cache Line）中，无需额外的显存访问。

4.3 Warp级别的优化

对于每块64个权重，可以用2个Warp（64线程）来处理。同一个Warp内的线程共享同一个Mask值，通过移位操作各自提取自己的位：

// 一个Warp（32线程）处理半个Block（32个权重） unsigned long long mask = __ldg(&layer_masks[block_id]); int lane_id = threadIdx.x % 32; int is_outlier = (mask >> lane_id) & 1;

Warp内无分支发散，因为所有线程执行相同的指令（移位+按位与），只是数据不同。即使is_outlier的值不同，if分支也是被Warp统一执行的，32个线程中只要有一个走向某个分支，整个Warp都会执行该分支的代码路径。

为了彻底消除分支，可以使用三元运算符替代if-else，让编译器生成无分支的谓词执行（Predicated Execution）指令：

float value = is_outlier ? outlier_value : common_value; // 编译器会生成无分支的cmov（条件移动）指令

五、进阶：当“普通值”本身也有多层基底

回到最初的那个数组：[23, 39, 99, 258]。我们用了两种基底：10×和100×。但实际LLM的权重分布可能是连续谱，而非离散的两类。

如果我们将Mask升级为2-bit，可以支持4种不同的量化基底：

Mask值	含义	基底
`00`	极小值	`2×`
`01`	普通值	`10×`
`10`	较大值	`50×`
`11`	异常值	`200×`

此时，解压逻辑变成：

int mask_2bit = (compressed_masks[block_id] >> (2 * offset_in_block)) & 0x3; float scale; switch (mask_2bit) { case 0: scale = 2.0; break; case 1: scale = 10.0; break; case 2: scale = 50.0; break; case 3: scale = 200.0; break; } float value = quantized_value * scale;

开关语句（Switch）在GPU上会被编译器展开为查表跳转，比if-else链高效得多。如果基底数量是2的幂（4、8、16种），可以用位提取 + 表索引实现O(1)查表。