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;
}