MarkBase Admin
d3379e23d5
CI / build-and-test (push) Has been cancelled
根本问题确认： ✅ 26B-A4B Router/Expert使用bits=8量化 ✅ inDim = 704*4 = 2816（8-bit: 4 vals/u32） ✅ groupSize = 2816/44 = 64 ⚠️ 现有dequantize_row kernel只支持bits=4 ⚠️ Kernel硬编码：groupSize/8, (inG%8)*4, &0xF ⚠️ 需要8-bit逻辑：groupSize/4, (inG%4)*8, &0xFF 已修复部分： ✅ loadExpertGroup groupSize计算（Line 1247-1251） ✅ 从scales shape正确计算groupSize ⚠️ 但仍需8-bit Metal kernel支持 修复方案对比： 方案A（修改Metal kernels）：数天，极高风险，不确定 ⭐ 方案B（使用26B-Standard）：0分钟，无风险，完美 ⭐⭐⭐⭐⭐ 创建文件： - dequantize_8bit_kernel.metal（示例kernel） - dequantizeRow_analysis.md（函数分析） - 26B_A4B_Deep_Fix_Analysis.md（完整分析） 结论： 技术上可修复，但难度极高（需修改Metal kernels） 强烈推荐使用26B-Standard代替（完美无NaN） 推荐度：方案B ⭐⭐⭐⭐⭐
22 lines
955 B
Metal
22 lines
955 B
Metal
kernel void dequantize_row_8bit(
|
|
device const uint *w [[buffer(0)]], // [nRows, nCols/4]
|
|
device const float *s [[buffer(1)], // [nRows, numGroups]
|
|
device const float *b [[buffer(2)]], // [nRows, numGroups]
|
|
device float *out [[buffer(3)], // [nCols]
|
|
constant uint &nCols [[buffer(4)]],
|
|
constant int &rowIdx [[buffer(5)]],
|
|
constant uint &groupSize [[buffer(6)]],
|
|
uint id [[thread_position_in_grid]]
|
|
) {
|
|
if (id >= nCols) return;
|
|
uint g = id / groupSize;
|
|
uint inG = id % groupSize;
|
|
// For 8-bit: 4 values per uint32
|
|
uint packedIdx = g * (groupSize / 4) + inG / 4;
|
|
uint shift = (inG % 4) * 8; // 8-bit shift
|
|
uint qval = (w[rowIdx * (nCols / 4) + packedIdx] >> shift) & 0xFF; // 8-bit mask
|
|
uint numGroups = nCols / groupSize;
|
|
float scale = s[rowIdx * numGroups + g];
|
|
float bias = b[rowIdx * numGroups + g];
|
|
out[id] = float(qval) * scale + bias;
|
|
} |