deep analysis: 26B-A4B根本问题 - Metal kernel需支持bits=8

根本问题确认：
✅ 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 ⭐⭐⭐⭐⭐

Sources/MarkBase/Metal/dequantize_8bit_kernel.metal

@@ -0,0 +1,22 @@
+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;
+}