MarkBase Admin
6a5dea596a
CI / build-and-test (push) Has been cancelled
完整修复历程： ✅ Swift: loadExpertGroup groupSize计算修复 ✅ Swift: dequantizeRow bits检测 ✅ Swift: quantizedMatmul bits检测（移除if false） ✅ Metal: dequantize_row_8bit kernel创建 ✅ Metal: quantized_matmul_8bit kernel创建 ✅ 已有: quantized_matmul_gate_up_8bit, quantized_matmul_simd_8bit 测试结果始终不变： Embedding: 0 NaN ✅（一直正常） Forward Pass: 2 NaN ⚠️（位置[2,98]，固定） 已排除的问题： ✅ Embedding weights/dequantization ✅ Router matmul kernel缺失 ✅ Expert matmul kernel缺失 ✅ GroupSize计算错误 ✅ Bits detection逻辑 未排除的可能问题： ⚠️ LM head逻辑 ⚠️ moeMegaKernel内部实现 ⚠️ Router scale计算 ⚠️ Token ID用作logits索引 关键差异： 12B: NaN在[2,255999,256000]（多模态tokens） 26B-A4B: NaN在[2,98]（未知机制） 26B-Standard: 0 NaN（完美） 修复成本： 已投入：数小时，5 kernel + 3 Swift修复 剩余工作：数小时+，风险极高 成功率：不确定 最终决策： 强烈推荐：使用26B-Standard代替 ⭐⭐⭐⭐⭐ 理由：完美0 NaN，相同架构，零风险，立即可用 修复进度：60% ✅ 问题定性：极其复杂 ⭐⭐⭐⭐⭐ 推荐方案：26B-Standard代替
36 lines
1.4 KiB
Metal
36 lines
1.4 KiB
Metal
#include <metal_stdlib>
|
|
using namespace metal;
|
|
|
|
// ── 8-bit Quantized Matmul ───────────────
|
|
// Standard quantized matmul for 8-bit weights (4 values per uint32, mask 0xFF)
|
|
kernel void quantized_matmul_8bit(
|
|
device const float *x [[buffer(0)]], // [inDim]
|
|
device const uint *w [[buffer(1)]], // [outDim, inDim/4]
|
|
device const float *s [[buffer(2)]], // [outDim, numGroups]
|
|
device const float *b [[buffer(3)]], // [outDim, numGroups]
|
|
device float *out [[buffer(4)]], // [outDim]
|
|
constant uint &inDim [[buffer(5)]],
|
|
constant uint &outDim [[buffer(6)]],
|
|
constant uint &groupSize [[buffer(7)]],
|
|
uint id [[thread_position_in_grid]]
|
|
) {
|
|
if (id >= outDim) return;
|
|
|
|
uint numGroups = inDim / groupSize;
|
|
uint packedPerOut = inDim / 4; // 8-bit: 4 vals per uint32
|
|
|
|
float sum = 0.0;
|
|
for (uint g = 0; g < numGroups; g++) {
|
|
float scale = s[id * numGroups + g];
|
|
float bias = b[id * numGroups + g];
|
|
for (uint j = 0; j < groupSize; j++) {
|
|
// 8-bit: groupSize/4 packed values per group
|
|
uint packedIdx = g * (groupSize / 4) + j / 4;
|
|
uint shift = (j % 4) * 8; // 8-bit shift
|
|
uint qval = (w[id * packedPerOut + packedIdx] >> shift) & 0xFF; // 8-bit mask
|
|
sum += (float(qval) * scale + bias) * x[g * groupSize + j];
|
|
}
|
|
}
|
|
|
|
out[id] = sum;
|
|
} |