引言当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# 原始Mask4096 bitsMask[0]*3968[1]*128# 前3968个是0后128个是1# 分块后每块64个权重Block62:全部为0-不存储 Block63:前64个权重全是1-存储(Block_ID63,64-bit_Mask0xFFFFFFFFFFFFFFFF)# 压缩后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个BlockBlock 0权重0~63位置0是异常值 → bit0 1位置31是异常值 → bit31 1位置63是异常值 → bit63 1其他位置是普通值 → 0Mask_Block0 0b1000...010...001 (bit631, bit311, bit01) 0x8000000080000001 (十六进制)Block 1权重64~127位置64是异常值 → bit0 1位置95是异常值 → bit31 1位置127是异常值 → bit63 1Mask_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只存储非全零的Blockmasks:{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个权重硬件友好度原始FP16256 字节0无需解压高直接计算无压缩Mask16 字节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个Warp64线程来处理。同一个Warp内的线程共享同一个Mask值通过移位操作各自提取自己的位// 一个Warp32线程处理半个Block32个权重 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)查表。六、总结从“压缩”到“路由”的范式转换我们最初的疑问是“统一位宽是浪费的能否让每个权重使用适合自己的位宽”我们设计了Mask路由 多基底量化的方案。然后我们发现Mask本身也需要压缩于是引入了RLE压缩。但RLE在GPU上串行解析太慢于是我们升级为结构化位图。最终方案的核心可以用一句话概括将Mask视为一种“路由表”用64-bit定长块存储用按位运算实现O(1)并行查表。这个演进的启示是压缩不能只考虑存储空间必须考虑解压速度。在GPU上一个慢速的解压器可能会抵消压缩带来的所有带宽收益。结构化是GPU友好的前提。定长块、固定位宽、无分支——这些“土气”的工程约束是算法在硬件上落地的基础。信息的价值是不均匀的。有些bit比如异常值的路由信息值得用更多位来保存有些bit比如普通值的完整精度可以被压缩到极致。回到最初的例子[23, 39, 99, 258]。如果采用我们最终的结构化位图方案23, 39, 99走10×基底存储为[2,3,9]和[3,9,9]仅占4-bit × 6 24 bits。258走100×基底存储为[2,5,8]占8-bit × 3 24 bits。Mask用2-bit编码两种基底存储为[01, 01, 01, 10]占8 bits。总计56 bits比原始64-bit节省了12.5%。对于大规模的LLM如70B参数这种优化叠加结构化稀疏和层间共享后保守估计可以将模型体积压缩到原来的30%-40%同时保持95%以上的原始精度——且解压速度接近直接读取FP16。这不是科幻这是正在发生的工程实践。