mirror of
https://github.com/Niko1221/Strata.git
synced 2026-11-07 18:35:04 -10:00
metal: the Unsloth UD-Q4/5/6_K_XL formats - K-quant / Q8_0 / Q5_1 routed experts, F32 GEMV + router, Q8_0 HC read with F32 inject
- Q5_1 (type 7): ref_quant (dequant, random blocks), quant.metal fmt/deq, units.metal Dec (w = d q + m, the min
through the unit's own x sums) / X, mma.metal DW
- decode experts (moe.metal): moe_gu_{12,13,14,8}_* (Q4_K, Q5_K, Q6_K, Q8_0), moe_down_{7,8}_* (Q5_1, Q8_0)
- prompt experts (moe_mm.metal): W4 decoders Q4_K / Q5_K / Q6_K / Q8_0 / Q5_1; mm_gu3_{12,13,14,8}_*,
mm_down3_{7,8}_*; mm_gu2 K-quant tiles with TK 128 (16_32_128_2/4, 16_16_128_2, 32_32_128_2), Q8_0 TK >= 32;
mm_down2_{7,8}_* (TK >= 32)
- F32 (type 0) as a format: gemv_0_*, gemvc_0_*, dq_f16_0, router_route_0
- gr.metal: gr_down_f<SG, KS, F, FI> with the inject rows in their own format: gr_down_q8f_<S>_<K> = Q8_0 down + F32
inject (gr_down_<S>_<K>, gr_down_q8_<S>_<K> unchanged)
- parity: expert_parity, moe_mm_parity v2/v3 (every gate/up x down pair), gemv_parity / dense_mm_parity F32,
fuse_parity router_route F32 (bitwise), gr_parity GR_Q8=1 GR_INJ_F32=1
This commit is contained in:
Niko1221 committed
1 parent
af6738eb74
commit
945c0d2424
17 files changed
+222
-40
No files matched your search
@@ -61,7 +61,7 @@ kernel void gemm_f16(device half* A [[buffer(0)]], device half* B [[buffer(1)]],
|
||||
|
||||
#define DQ(F) template [[host_name("dq_f16_" #F)]] kernel decltype(dq_f16<F>) dq_f16<F>;
|
||||
#define DQ_(F) DQ(F)
|
||||
DQ_(T_BF16) DQ_(T_Q8_0) DQ_(T_Q4_0) DQ_(T_Q5_0) DQ_(T_IQ4_NL) DQ_(T_IQ4_XS) DQ_(T_Q2_0) DQ_(T_IQ3_S) DQ_(T_Q3_K)
|
||||
DQ_(T_F32) DQ_(T_BF16) DQ_(T_Q8_0) DQ_(T_Q4_0) DQ_(T_Q5_0) DQ_(T_IQ4_NL) DQ_(T_IQ4_XS) DQ_(T_Q2_0) DQ_(T_IQ3_S) DQ_(T_Q3_K)
|
||||
DQ_(T_Q4_K) DQ_(T_Q5_K) DQ_(T_Q6_K)
|
||||
template [[host_name("gemm_f16_64_64_4")]] kernel decltype(gemm_f16<64, 64, 4>) gemm_f16<64, 64, 4>;
|
||||
template [[host_name("gemm_f16_128_64_4")]] kernel decltype(gemm_f16<128, 64, 4>) gemm_f16<128, 64, 4>;
|
||||
|
||||
@@ -95,7 +95,7 @@ kernel void gemv(device const uchar* W [[buffer(0)]], device const float* xp [[b
|
||||
GV(F, 8, 2, 4, 1) GV(F, 4, 2, 4, 1) GV(F, 4, 1, 4, 1) GV(F, 32, 1, 4, 4) GV(F, 32, 2, 4, 2) \
|
||||
GV(F, 32, 1, 8, 8) GV(F, 16, 1, 8, 4) GV(F, 32, 1, 8, 2) GV(F, 8, 1, 8, 1) GV(F, 1, 1, 4, 1) GV(F, 1, 2, 4, 1) GV(F, 2, 1, 4, 1) GV(F, 2, 2, 4, 1) GV(F, 1, 1, 2, 1) GV(F, 4, 4, 4, 1)
|
||||
#ifndef GEMV_TYPES
|
||||
#define GEMV_TYPES GV_SET(T_BF16) GV_SET(T_Q8_0) GV_SET(T_Q4_0) GV_SET(T_Q5_0) GV_SET(T_IQ4_NL) GV_SET(T_IQ4_XS) \
|
||||
#define GEMV_TYPES GV_SET(T_F32) GV_SET(T_BF16) GV_SET(T_Q8_0) GV_SET(T_Q4_0) GV_SET(T_Q5_0) GV_SET(T_IQ4_NL) GV_SET(T_IQ4_XS) \
|
||||
GV_SET(T_Q2_0) GV_SET(T_IQ3_S) GV_SET(T_Q3_K) GV_SET(T_Q4_K) GV_SET(T_Q5_K) GV_SET(T_Q6_K)
|
||||
#endif
|
||||
GEMV_TYPES
|
||||
@@ -208,6 +208,7 @@ kernel void gemv_cat(device const uchar* W [[buffer(0)]], device const float* xp
|
||||
#define GVC_(F, L, R, S) template [[host_name("gemvc_" #F "_" #L "_" #R "_" #S)]] kernel decltype(gemv_cat<F, L, R, S>) gemv_cat<F, L, R, S>;
|
||||
#define GVC(F, L, R, S) GVC_(F, L, R, S)
|
||||
GVC(T_BF16, 32, 1, 4) GVC(T_BF16, 32, 1, 8) GVC(T_BF16, 16, 1, 4) GVC(T_BF16, 32, 2, 4)
|
||||
GVC(T_F32, 32, 1, 4) GVC(T_F32, 32, 1, 8) GVC(T_F32, 16, 1, 4) GVC(T_F32, 32, 2, 4)
|
||||
|
||||
// gemv_x: gemv (KS = 1) with the window's prepared x (the format's plane, NTOK rows, and its per-16 sums) staged in
|
||||
// threadgroup memory once per threadgroup, so the per-unit x reads of every token come from threadgroup memory
|
||||
|
||||
+26
-17
@@ -53,7 +53,26 @@ kernel void gr_norm(device float* R [[buffer(0)]], device const float* bo [[buff
|
||||
|
||||
// rows: 0 .. hc_lr - 1 = w_down, hc_lr .. hc_lr + hc - 1 = w_inject (both bf16 rows of hc * n_embd).
|
||||
// SG simdgroups split the row's units KS ways; LPR = 32 (a simdgroup per row); RS = SG / KS rows per threadgroup.
|
||||
template <int SG, int KS, int F>
|
||||
// one row's K-split partial over xn (units of format F)
|
||||
template <int F>
|
||||
static inline void grd_acc(device const uchar* w, int nin, int ks, int KS, ushort lane, device const float* xn,
|
||||
thread float* acc, threadgroup const uint* notab) {
|
||||
constexpr int UNIT = fmt<F>::UNIT;
|
||||
const int U = nin / UNIT, u0 = ks * U / KS, u1 = (ks + 1) * U / KS;
|
||||
for (int u = u0 + lane; u < u1; u += 32) {
|
||||
float q[UNIT];
|
||||
float4 sc;
|
||||
Dec<F>::run(w, u, q, sc, notab);
|
||||
for (int t = 0; t < GR_NTOK; ++t) {
|
||||
float xs[UNIT], xsum[2] = {0.0f, 0.0f};
|
||||
for (int k = 0; k < UNIT; k += 16) load16(xn + t * nin + UNIT * u + k, xs + k);
|
||||
acc[t] += unit_combine<UNIT>(q, sc, xs, xsum);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
// FI: the inject rows' format (the Unsloth files: Q8_0 down, F32 inject); FI = F is the original kernel
|
||||
template <int SG, int KS, int F, int FI>
|
||||
kernel void gr_down_f(device const uchar* w_down [[buffer(0)]], device const uchar* w_inject [[buffer(1)]],
|
||||
device const float* xn [[buffer(2)]], device float* lo [[buffer(3)]],
|
||||
device float* inject [[buffer(4)]], constant GrArgs& a [[buffer(5)]],
|
||||
@@ -67,22 +86,11 @@ kernel void gr_down_f(device const uchar* w_down [[buffer(0)]], device const uch
|
||||
const int nrows = int(a.hc_lr + a.hc);
|
||||
const int rr = min(row, nrows - 1);
|
||||
const int nin = int(a.hc * a.n_embd);
|
||||
constexpr int UNIT = fmt<F>::UNIT;
|
||||
const uint rb = uint(nin / fmt<F>::QK * fmt<F>::BB);
|
||||
device const uchar* w = rr < int(a.hc_lr) ? w_down + ulong(rr) * rb : w_inject + ulong(rr - int(a.hc_lr)) * rb;
|
||||
const int U = nin / UNIT, u0 = ks * U / KS, u1 = (ks + 1) * U / KS;
|
||||
const uint rb = uint(nin / fmt<F>::QK * fmt<F>::BB), rbi = uint(nin / fmt<FI>::QK * fmt<FI>::BB);
|
||||
float acc[8];
|
||||
for (int t = 0; t < 8; ++t) acc[t] = 0.0f;
|
||||
for (int u = u0 + lane; u < u1; u += 32) {
|
||||
float q[UNIT];
|
||||
float4 sc;
|
||||
Dec<F>::run(w, u, q, sc, notab);
|
||||
for (int t = 0; t < GR_NTOK; ++t) {
|
||||
float xs[UNIT], xsum[2] = {0.0f, 0.0f};
|
||||
for (int k = 0; k < UNIT; k += 16) load16(xn + t * nin + UNIT * u + k, xs + k);
|
||||
acc[t] += unit_combine<UNIT>(q, sc, xs, xsum);
|
||||
}
|
||||
}
|
||||
if (rr < int(a.hc_lr)) grd_acc<F>(w_down + ulong(rr) * rb, nin, ks, KS, lane, xn, acc, notab);
|
||||
else grd_acc<FI>(w_inject + ulong(rr - int(a.hc_lr)) * rbi, nin, ks, KS, lane, xn, acc, notab);
|
||||
for (int t = 0; t < GR_NTOK; ++t) {
|
||||
const float v = simd_sum(acc[t]);
|
||||
if (lane == 0) red[(ks * RS + rs) * 8 + t] = v;
|
||||
@@ -163,8 +171,9 @@ kernel void gr_up_f(device const uchar* w_up [[buffer(0)]], device const float*
|
||||
}
|
||||
}
|
||||
|
||||
#define GRD(S, K) template [[host_name("gr_down_" #S "_" #K)]] kernel decltype(gr_down_f<S, K, T_BF16>) gr_down_f<S, K, T_BF16>; \
|
||||
template [[host_name("gr_down_q8_" #S "_" #K)]] kernel decltype(gr_down_f<S, K, T_Q8_0>) gr_down_f<S, K, T_Q8_0>;
|
||||
#define GRD(S, K) template [[host_name("gr_down_" #S "_" #K)]] kernel decltype(gr_down_f<S, K, T_BF16, T_BF16>) gr_down_f<S, K, T_BF16, T_BF16>; \
|
||||
template [[host_name("gr_down_q8_" #S "_" #K)]] kernel decltype(gr_down_f<S, K, T_Q8_0, T_Q8_0>) gr_down_f<S, K, T_Q8_0, T_Q8_0>; \
|
||||
template [[host_name("gr_down_q8f_" #S "_" #K)]] kernel decltype(gr_down_f<S, K, T_Q8_0, T_F32>) gr_down_f<S, K, T_Q8_0, T_F32>;
|
||||
GRD(8, 8) GRD(8, 4) GRD(4, 4) GRD(4, 2) GRD(16, 16) GRD(16, 8)
|
||||
template [[host_name("gr_up")]] kernel decltype(gr_up_f<T_BF16>) gr_up_f<T_BF16>;
|
||||
template [[host_name("gr_up_q8")]] kernel decltype(gr_up_f<T_Q8_0>) gr_up_f<T_Q8_0>;
|
||||
@@ -38,6 +38,26 @@ template <> struct DW<T_Q2_0> {
|
||||
p1 = p0 + 8;
|
||||
}
|
||||
};
|
||||
template <> struct DW<T_Q5_1> {
|
||||
static inline void run(device const uchar* row, int u, int k0, thread float* w, thread int& p0, thread int& p1,
|
||||
threadgroup const uint* tab) {
|
||||
device const uchar* b = row + u * 24;
|
||||
float q[32];
|
||||
const float d = deq<T_Q5_1>(b, 0, q, tab), m = rdh(b + 2);
|
||||
for (int j = 0; j < 32; ++j) w[j] = q[j] * d + m;
|
||||
p0 = 32 * u - k0;
|
||||
p1 = p0 + 16;
|
||||
}
|
||||
};
|
||||
template <> struct DW<T_F32> {
|
||||
static inline void run(device const uchar* row, int u, int k0, thread float* w, thread int& p0, thread int& p1,
|
||||
threadgroup const uint*) {
|
||||
device const float* p = (device const float*)(row + 64 * u);
|
||||
for (int j = 0; j < 16; ++j) w[j] = p[j];
|
||||
p0 = 16 * u - k0;
|
||||
p1 = p0 + 8;
|
||||
}
|
||||
};
|
||||
template <> struct DW<T_BF16> {
|
||||
static inline void run(device const uchar* row, int u, int k0, thread float* w, thread int& p0, thread int& p1,
|
||||
threadgroup const uint*) {
|
||||
|
||||
@@ -309,6 +309,8 @@ GU_SET(T_IQ2_XS)
|
||||
GU_SET(T_IQ2_S)
|
||||
GU_SET(T_IQ3_S)
|
||||
GU_SET(T_IQ4_XS) // the IQ3_S pack's gate/up in one layer
|
||||
GU_SET(T_Q4_K) GU_SET(T_Q5_K) GU_SET(T_Q6_K) GU_SET(T_Q8_0) // the UD-Q4/5/6_K_XL packs
|
||||
DN_SET(T_Q5_1) DN_SET(T_Q8_0)
|
||||
DN_SET(T_Q2_0)
|
||||
DN_SET(T_Q2_0R)
|
||||
DN_SET(T_IQ4_NL)
|
||||
@@ -365,6 +365,13 @@ kernel void mm_down2(device const uchar* W [[buffer(0)]], device const half* HS
|
||||
#define MMDN2(F, M, N, K, S) MMDN2_(F, M, N, K, S)
|
||||
#define MM_GU2_SET(F) MMGU2(F, 32, 64, 32, 2) MMGU2(F, 32, 32, 32, 2) MMGU2(F, 32, 64, 16, 2) MMGU2(F, 64, 32, 16, 2) MMGU2(F, 16, 32, 32, 2) MMGU2(F, 16, 64, 16, 2) MMGU2(F, 16, 128, 32, 4) MMGU2(F, 32, 64, 16, 4) MMGU2(F, 16, 32, 32, 1) MMGU2(F, 16, 64, 32, 4) MMGU2(F, 16, 64, 64, 2) MMGU2(F, 32, 64, 64, 4) MMGU2(F, 64, 64, 64, 4) MMGU2(F, 16, 32, 64, 2) MMGU2(F, 32, 64, 32, 4) MMGU2(F, 64, 64, 32, 4) MMGU2(F, 16, 64, 32, 2)
|
||||
#define MM_DN2_SET(F) MMDN2(F, 32, 64, 32, 2) MMDN2(F, 32, 32, 32, 2) MMDN2(F, 32, 64, 16, 2) MMDN2(F, 64, 32, 16, 2) MMDN2(F, 16, 32, 32, 2) MMDN2(F, 16, 64, 16, 2) MMDN2(F, 16, 128, 32, 4) MMDN2(F, 32, 64, 16, 4) MMDN2(F, 16, 32, 32, 1) MMDN2(F, 16, 64, 32, 4) MMDN2(F, 16, 64, 64, 2) MMDN2(F, 32, 64, 64, 4) MMDN2(F, 64, 64, 64, 4) MMDN2(F, 16, 32, 64, 2) MMDN2(F, 32, 64, 32, 4) MMDN2(F, 64, 64, 32, 4) MMDN2(F, 16, 64, 32, 2)
|
||||
#define MM_GU2_U32(F) MMGU2(F, 16, 64, 32, 2) MMGU2(F, 16, 64, 64, 2) MMGU2(F, 32, 64, 32, 2) MMGU2(F, 16, 32, 32, 2) \
|
||||
MMGU2(F, 64, 64, 32, 4)
|
||||
#define MM_DN2_U32(F) MMDN2(F, 16, 64, 32, 2) MMDN2(F, 16, 64, 64, 2) MMDN2(F, 16, 32, 32, 2) MMDN2(F, 32, 32, 32, 2) \
|
||||
MMDN2(F, 16, 128, 32, 4) MMDN2(F, 32, 64, 32, 2) MMDN2(F, 64, 64, 32, 4)
|
||||
// K-quants: the staged K chunk must hold whole units as DW places them (128 wide): TK 128
|
||||
#define MM_GU2_KQ(F) MMGU2(F, 16, 32, 128, 2) MMGU2(F, 16, 16, 128, 2) MMGU2(F, 32, 32, 128, 2) MMGU2(F, 16, 32, 128, 4)
|
||||
MM_GU2_KQ(T_Q4_K) MM_GU2_KQ(T_Q5_K) MM_GU2_KQ(T_Q6_K) MM_GU2_U32(T_Q8_0) MM_DN2_U32(T_Q5_1) MM_DN2_U32(T_Q8_0)
|
||||
// IQ4_XS (UNIT 32): the K chunk must be >= 32
|
||||
MMGU2(T_IQ4_XS, 16, 64, 32, 2) MMGU2(T_IQ4_XS, 16, 64, 64, 2) MMGU2(T_IQ4_XS, 32, 64, 32, 2) MMGU2(T_IQ4_XS, 16, 32, 32, 2)
|
||||
MMGU2(T_IQ4_XS, 64, 64, 32, 4)
|
||||
@@ -499,6 +506,77 @@ template <> struct W4<T_IQ4_NL> { // 32 a block: d, 16 bytes; low nibbles are
|
||||
|
||||
// W8<F>: the runs k .. k + 3 and k + 16 .. k + 19 of a row (k % 32 < 16): one 32-group, so the per-group fields (scale,
|
||||
// sign word) are read once. The generic one is two W4.
|
||||
// ---- the K-quant / legacy-quant W4 decoders (the UD-Q4/5/6_K_XL packs' experts): values k .. k + 3 (k % 4 == 0)
|
||||
// of a row in ref_quant's arithmetic (float), then fp16
|
||||
template <> struct W4<T_Q8_0> { // 32 a block: d, int8 qs[32]
|
||||
static inline half4 run(device const uchar* row, int k, threadgroup const uint*) {
|
||||
device const uchar* b = row + (k >> 5) * 34;
|
||||
const float d = rdh(b);
|
||||
device const char* q = (device const char*)(b + 2 + (k & 31));
|
||||
return half4(float4(float(q[0]), float(q[1]), float(q[2]), float(q[3])) * d);
|
||||
}
|
||||
};
|
||||
template <> struct W4<T_Q5_1> { // 32 a block: d, m, qh (32 bits), qs[16] (low nibbles 0..15, high 16..31)
|
||||
static inline half4 run(device const uchar* row, int k, threadgroup const uint*) {
|
||||
device const uchar* b = row + (k >> 5) * 24;
|
||||
const float d = rdh(b), m = rdh(b + 2);
|
||||
const uint hb = rd32(b + 4);
|
||||
const int j = k & 31, o = j & 15, hi = j >> 4;
|
||||
float4 w;
|
||||
for (int c = 0; c < 4; ++c) {
|
||||
const uint qv = b[8 + o + c];
|
||||
const int x = int(hi ? (qv >> 4) : (qv & 15)) | int(((hb >> (j + c)) & 1) << 4);
|
||||
w[c] = float(x) * d + m;
|
||||
}
|
||||
return half4(w);
|
||||
}
|
||||
};
|
||||
template <> struct W4<T_Q4_K> { // 256 a block: d, dmin, scales[12], qs[128] (64-chunks: low nibbles then high)
|
||||
static inline half4 run(device const uchar* row, int k, threadgroup const uint*) {
|
||||
device const uchar* b = row + (k >> 8) * 144;
|
||||
const int j = k & 255, c = j >> 6, hi = (j >> 5) & 1, l = j & 31;
|
||||
float s, mm;
|
||||
scale_min_k4(2 * c + hi, b + 4, s, mm);
|
||||
const float d1 = rdh(b) * s, m1 = rdh(b + 2) * mm;
|
||||
device const uchar* q = b + 16 + 32 * c + l;
|
||||
float4 w;
|
||||
for (int i = 0; i < 4; ++i) w[i] = d1 * float(hi ? (q[i] >> 4) : (q[i] & 15)) - m1;
|
||||
return half4(w);
|
||||
}
|
||||
};
|
||||
template <> struct W4<T_Q5_K> { // 256 a block: d, dmin, scales[12], qh[32], qs[128]
|
||||
static inline half4 run(device const uchar* row, int k, threadgroup const uint*) {
|
||||
device const uchar* b = row + (k >> 8) * 176;
|
||||
const int j = k & 255, c = j >> 6, hi = (j >> 5) & 1, l = j & 31;
|
||||
float s, mm;
|
||||
scale_min_k4(2 * c + hi, b + 4, s, mm);
|
||||
const float d1 = rdh(b) * s, m1 = rdh(b + 2) * mm;
|
||||
device const uchar* qh = b + 16 + l;
|
||||
device const uchar* ql = b + 48 + 32 * c + l;
|
||||
const uint bit = 1u << (2 * c + hi);
|
||||
float4 w;
|
||||
for (int i = 0; i < 4; ++i)
|
||||
w[i] = d1 * float(int(hi ? (ql[i] >> 4) : (ql[i] & 15)) + ((qh[i] & bit) ? 16 : 0)) - m1;
|
||||
return half4(w);
|
||||
}
|
||||
};
|
||||
template <> struct W4<T_Q6_K> { // 256 a block: ql[128], qh[64], int8 scales[16], d; halves of 128
|
||||
static inline half4 run(device const uchar* row, int k, threadgroup const uint*) {
|
||||
device const uchar* b = row + (k >> 8) * 210;
|
||||
const int j = k & 255, n = j >> 7, jj = j & 127, qd = jj >> 5, l = jj & 31;
|
||||
device const uchar* ql = b + 64 * n + l + ((qd & 1) ? 32 : 0);
|
||||
device const uchar* qh = b + 128 + 32 * n + l;
|
||||
const float d = rdh(b + 208);
|
||||
const float ds = d * float(((device const char*)(b + 192 + 8 * n))[l / 16 + 2 * qd]);
|
||||
float4 w;
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
const int q = int(((qd >> 1) ? (ql[i] >> 4) : (ql[i] & 15)) | (((qh[i] >> (2 * qd)) & 3) << 4)) - 32;
|
||||
w[i] = ds * float(q);
|
||||
}
|
||||
return half4(w);
|
||||
}
|
||||
};
|
||||
|
||||
template <int F> struct W8 {
|
||||
static inline void run(device const uchar* row, int k, threadgroup const uint* tab, thread half4& a, thread half4& b) {
|
||||
a = W4<F>::run(row, k, tab);
|
||||
@@ -741,8 +819,8 @@ kernel void mm_down3(device const uchar* W [[buffer(0)]], device const half* HS
|
||||
MMGU3(F, 16, 32, 16, 2) MMGU3(F, 32, 32, 32, 2) MMGU3(F, 16, 64, 32, 1) MMGU3(F, 16, 64, 32, 4) MMGU3(F, 16, 32, 16, 4) MMGU3(F, 16, 32, 32, 5) MMGU3(F, 16, 32, 16, 5) MMGU3(F, 16, 16, 32, 4) MMGU3(F, 16, 16, 32, 8) MMGU3(F, 16, 16, 16, 8)
|
||||
#define MM3_DN(F) MMDN3(F, 16, 64, 32, 2) MMDN3(F, 16, 32, 32, 4) MMDN3(F, 32, 64, 32, 2) MMDN3(F, 16, 64, 64, 2) \
|
||||
MMDN3(F, 16, 32, 16, 2) MMDN3(F, 32, 32, 32, 2) MMDN3(F, 16, 64, 32, 1) MMDN3(F, 16, 64, 32, 4) MMDN3(F, 16, 64, 16, 2) MMDN3(F, 16, 32, 16, 4) MMDN3(F, 16, 64, 32, 5) MMDN3(F, 16, 64, 16, 4)
|
||||
MM3_GU(T_IQ4_XS) MM3_GU(T_Q2_0) MM3_GU(T_IQ3_XXS) MM3_GU(T_IQ2_XXS) MM3_GU(T_IQ2_XS) MM3_GU(T_IQ2_S) MM3_GU(T_IQ3_S)
|
||||
MM3_DN(T_Q2_0) MM3_DN(T_IQ4_NL)
|
||||
MM3_GU(T_Q4_K) MM3_GU(T_Q5_K) MM3_GU(T_Q6_K) MM3_GU(T_Q8_0) MM3_GU(T_IQ4_XS) MM3_GU(T_Q2_0) MM3_GU(T_IQ3_XXS) MM3_GU(T_IQ2_XXS) MM3_GU(T_IQ2_XS) MM3_GU(T_IQ2_S) MM3_GU(T_IQ3_S)
|
||||
MM3_DN(T_Q2_0) MM3_DN(T_IQ4_NL) MM3_DN(T_Q5_1) MM3_DN(T_Q8_0)
|
||||
|
||||
// mm_gu4: mm_gu3 with each decoded right operand (TN x TK weights of gate and of up) applied to MS 16-row token
|
||||
// tiles (a tile of 16 * MS rows from the plan, tm = 16 * MS): the decode, which bounds mm_gu3, is shared by MS times
|
||||
|
||||
@@ -13,7 +13,9 @@
|
||||
#define T_Q4_0 2
|
||||
#define T_Q3_K 11
|
||||
#define T_BF16 30
|
||||
#define T_F32 0
|
||||
#define T_Q5_0 6
|
||||
#define T_Q5_1 7
|
||||
#define T_Q8_0 8
|
||||
#define T_Q4_K 12
|
||||
#define T_Q5_K 13
|
||||
@@ -44,9 +46,11 @@ template <> struct fmt<T_Q4_K> { enum { QK = 256, BB = 144, UNIT = 32, TAB =
|
||||
template <> struct fmt<T_Q5_K> { enum { QK = 256, BB = 176, UNIT = 32, TAB = 1 }; };
|
||||
template <> struct fmt<T_Q6_K> { enum { QK = 256, BB = 210, UNIT = 32, TAB = 1 }; };
|
||||
template <> struct fmt<T_BF16> { enum { QK = 16, BB = 32, UNIT = 16, TAB = 1 }; };
|
||||
template <> struct fmt<T_F32> { enum { QK = 16, BB = 64, UNIT = 16, TAB = 1 }; };
|
||||
template <> struct fmt<T_Q8_0> { enum { QK = 32, BB = 34, UNIT = 32, TAB = 1 }; };
|
||||
template <> struct fmt<T_Q4_0> { enum { QK = 32, BB = 18, UNIT = 32, TAB = 1 }; };
|
||||
template <> struct fmt<T_Q5_0> { enum { QK = 32, BB = 22, UNIT = 32, TAB = 1 }; };
|
||||
template <> struct fmt<T_Q5_1> { enum { QK = 32, BB = 24, UNIT = 32, TAB = 1 }; };
|
||||
|
||||
static inline uint rd32(device const uchar* p) {
|
||||
device const ushort* s = (device const ushort*)p;
|
||||
@@ -76,11 +80,13 @@ template <> inline void load_tab<T_Q2_0R>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q8_0>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q4_0>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q5_0>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q5_1>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q3_K>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q4_K>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q5_K>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_Q6_K>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_BF16>(threadgroup uint*, ushort, ushort) {}
|
||||
template <> inline void load_tab<T_F32>(threadgroup uint*, ushort, ushort) {}
|
||||
|
||||
static inline void load_signs(threadgroup uint* tab, int at, ushort tid, ushort nth) {
|
||||
threadgroup uchar* s = (threadgroup uchar*)(tab + at);
|
||||
@@ -283,6 +289,21 @@ template <> inline float deq<T_Q5_0>(device const uchar* b, int, thread float* q
|
||||
return rdh(b);
|
||||
}
|
||||
|
||||
// Q5_1: 32 = 1 unit; q = the 5-bit code (no offset; the min m is applied through the x sums), returns d
|
||||
template <> inline float deq<T_Q5_1>(device const uchar* b, int, thread float* q, threadgroup const uint*) {
|
||||
const uint hb = rd32(b + 4);
|
||||
device const ushort* p = (device const ushort*)(b + 8);
|
||||
for (int j = 0; j < 8; ++j) {
|
||||
const ushort v = p[j];
|
||||
const int j0 = 2 * j, j1 = 2 * j + 1;
|
||||
q[j0] = float(int(v & 15) | int(((hb >> j0) & 1) << 4));
|
||||
q[j0 + 16] = float(int((v >> 4) & 15) | int(((hb >> (j0 + 16)) & 1) << 4));
|
||||
q[j1] = float(int((v >> 8) & 15) | int(((hb >> j1) & 1) << 4));
|
||||
q[j1 + 16] = float(int(v >> 12) | int(((hb >> (j1 + 16)) & 1) << 4));
|
||||
}
|
||||
return rdh(b);
|
||||
}
|
||||
|
||||
// sum over groups of LPR consecutive lanes (LPR = 1, 2, 4, 8, 16 or 32); the result is valid in every lane
|
||||
template <int LPR>
|
||||
static inline float lane_sum(float v) {
|
||||
|
||||
@@ -101,3 +101,4 @@ kernel void router_route(device const uchar* W [[buffer(0)]], device const float
|
||||
}
|
||||
|
||||
template [[host_name("router_route_30")]] kernel decltype(router_route<T_BF16>) router_route<T_BF16>;
|
||||
template [[host_name("router_route_0")]] kernel decltype(router_route<T_F32>) router_route<T_F32>;
|
||||
@@ -85,6 +85,18 @@ template <> struct X<T_Q6_K> {
|
||||
xsum[1] = s16[(p0 >> 4) + 4];
|
||||
}
|
||||
};
|
||||
// Q5_1: 32 contiguous raw x, the two per-16 sums formed from them (the GEMVs that read Q5_1 - the routed down - get
|
||||
// h without prepared sums)
|
||||
template <> struct X<T_Q5_1> {
|
||||
template <typename P> static inline void load(P xp, P, int u, thread float* xs, thread float* xsum) {
|
||||
load16(xp + 32 * u, xs);
|
||||
load16(xp + 32 * u + 16, xs + 16);
|
||||
float a = 0.0f, b = 0.0f;
|
||||
for (int j = 0; j < 16; ++j) { a += xs[j]; b += xs[16 + j]; }
|
||||
xsum[0] = a;
|
||||
xsum[1] = b;
|
||||
}
|
||||
};
|
||||
// Q3_K: 32 contiguous, sums per 16
|
||||
template <> struct X<T_Q3_K> {
|
||||
template <typename P> static inline void load(P xp, P s16, int u, thread float* xs, thread float* xsum) {
|
||||
@@ -179,6 +191,14 @@ static inline void bytes2(ushort w, thread float* q) {
|
||||
q[1] = float(w & 0xFF00);
|
||||
}
|
||||
|
||||
template <> struct Dec<T_Q5_1> { // w = d q + m: sum = d (q . x) + m (sum x)
|
||||
static inline void run(device const uchar* row, int u, thread float* q, thread float4& sc, threadgroup const uint* tab) {
|
||||
device const uchar* b = row + u * 24;
|
||||
const float d = deq<T_Q5_1>(b, 0, q, tab), m = rdh(b + 2);
|
||||
sc = float4(d, d, -m, -m);
|
||||
}
|
||||
};
|
||||
|
||||
template <> struct Dec<T_Q4_K> {
|
||||
static inline void run(device const uchar* row, int u, thread float* q, thread float4& sc, threadgroup const uint*) {
|
||||
device const uchar* b = row + (u >> 3) * 144;
|
||||
@@ -262,6 +282,17 @@ template <> struct Dec<T_Q3_K> {
|
||||
}
|
||||
};
|
||||
|
||||
template <> struct Dec<T_F32> { // 16 floats a unit (the Unsloth files' router / alpha / beta / inject rows)
|
||||
static inline void run(device const uchar* row, int u, thread float* q, thread float4& sc, threadgroup const uint*) {
|
||||
device const float4* p = (device const float4*)(row + 64 * u);
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
const float4 v = p[k];
|
||||
q[4 * k] = v.x; q[4 * k + 1] = v.y; q[4 * k + 2] = v.z; q[4 * k + 3] = v.w;
|
||||
}
|
||||
sc = float4(1.0f, 1.0f, 0.0f, 0.0f);
|
||||
}
|
||||
};
|
||||
|
||||
template <> struct Dec<T_BF16> {
|
||||
static inline void run(device const uchar* row, int u, thread float* q, thread float4& sc, threadgroup const uint*) {
|
||||
device const uint4* p = (device const uint4*)(row + 32 * u);
|
||||
|
||||
+16
-3
@@ -18,7 +18,7 @@ namespace ref {
|
||||
const char* type_name(int t) {
|
||||
switch (t) {
|
||||
case F32: return "F32"; case F16: return "F16"; case BF16: return "BF16";
|
||||
case Q4_0: return "Q4_0"; case Q5_0: return "Q5_0"; case Q8_0: return "Q8_0";
|
||||
case Q4_0: return "Q4_0"; case Q5_0: return "Q5_0"; case Q5_1: return "Q5_1"; case Q8_0: return "Q8_0";
|
||||
case Q3_K: return "Q3_K"; case Q4_K: return "Q4_K"; case Q5_K: return "Q5_K"; case Q6_K: return "Q6_K";
|
||||
case IQ2_XXS: return "IQ2_XXS"; case IQ2_XS: return "IQ2_XS"; case IQ2_S: return "IQ2_S";
|
||||
case IQ3_XXS: return "IQ3_XXS"; case IQ3_S: return "IQ3_S"; case IQ4_NL: return "IQ4_NL";
|
||||
@@ -30,7 +30,7 @@ const char* type_name(int t) {
|
||||
int block_elems(int t) {
|
||||
switch (t) {
|
||||
case F32: case F16: case BF16: return 1;
|
||||
case Q4_0: case Q5_0: case Q8_0: case IQ4_NL: return 32;
|
||||
case Q4_0: case Q5_0: case Q5_1: case Q8_0: case IQ4_NL: return 32;
|
||||
case Q2_0: return 64;
|
||||
default: return 256;
|
||||
}
|
||||
@@ -39,7 +39,7 @@ int block_elems(int t) {
|
||||
int block_bytes(int t) {
|
||||
switch (t) {
|
||||
case F32: return 4; case F16: case BF16: return 2;
|
||||
case Q4_0: return 18; case Q5_0: return 22; case Q8_0: return 34; case IQ4_NL: return 18; case Q2_0: return 18;
|
||||
case Q4_0: return 18; case Q5_0: return 22; case Q5_1: return 24; case Q8_0: return 34; case IQ4_NL: return 18; case Q2_0: return 18;
|
||||
case Q3_K: return 110; case Q4_K: return 144; case Q5_K: return 176; case Q6_K: return 210;
|
||||
case IQ2_XXS: return 66; case IQ2_XS: return 74; case IQ2_S: return 82; case IQ3_XXS: return 98;
|
||||
case IQ3_S: return 110; case IQ4_XS: return 136;
|
||||
@@ -134,6 +134,18 @@ static void deq_block(int t, const uint8_t* b, float* y) {
|
||||
}
|
||||
return;
|
||||
}
|
||||
case Q5_1: { // ggml dequantize_row_q5_1: y = x * d + m
|
||||
const float d = f16_to_f32(rd16(b)), m = f16_to_f32(rd16(b + 2));
|
||||
uint32_t h;
|
||||
std::memcpy(&h, b + 4, 4);
|
||||
const uint8_t* qs = b + 8;
|
||||
for (int j = 0; j < 16; ++j) {
|
||||
const int x0 = (qs[j] & 15) | (((h >> j) & 1) << 4), x1 = (qs[j] >> 4) | (((h >> (j + 16)) & 1) << 4);
|
||||
y[j] = (float)x0 * d + m;
|
||||
y[j + 16] = (float)x1 * d + m;
|
||||
}
|
||||
return;
|
||||
}
|
||||
case IQ4_NL: {
|
||||
const float d = f16_to_f32(rd16(b));
|
||||
for (int j = 0; j < 16; ++j) {
|
||||
@@ -417,6 +429,7 @@ void random_blocks(int t, uint8_t* dst, int rows, int cols, uint64_t seed, float
|
||||
case Q8_0: put16(b, f32_to_f16(scale * u / 64)); break;
|
||||
case Q4_0: case IQ4_NL: put16(b, f32_to_f16(scale * u / (t == Q4_0 ? 4 : 64))); break;
|
||||
case Q5_0: put16(b, f32_to_f16(scale * u / 8)); break;
|
||||
case Q5_1: put16(b, f32_to_f16(scale * u / 16)); put16(b + 2, f32_to_f16(-scale * u)); break;
|
||||
case IQ4_XS: put16(b, f32_to_f16(scale * u / 64 / 16)); break;
|
||||
case Q4_K: case Q5_K: put16(b, f32_to_f16(scale * u / 32 / 8)); put16(b + 2, f32_to_f16(scale * u / 32 / 2)); break;
|
||||
case Q6_K: put16(b + 208, f32_to_f16(scale * u / 32 / 64)); break;
|
||||
|
||||
@@ -11,7 +11,7 @@ namespace ref {
|
||||
|
||||
// GGML type ids (stable, as in native_mmvq.hpp)
|
||||
enum : int {
|
||||
F32 = 0, F16 = 1, Q4_0 = 2, Q5_0 = 6, Q8_0 = 8, Q3_K = 11, Q4_K = 12, Q5_K = 13, Q6_K = 14,
|
||||
F32 = 0, F16 = 1, Q4_0 = 2, Q5_0 = 6, Q5_1 = 7, Q8_0 = 8, Q3_K = 11, Q4_K = 12, Q5_K = 13, Q6_K = 14,
|
||||
IQ2_XXS = 16, IQ2_XS = 17, IQ3_XXS = 18, IQ3_S = 21, IQ2_S = 22, IQ4_NL = 20, IQ4_XS = 23, BF16 = 30, Q2_0 = 42,
|
||||
};
|
||||
|
||||
|
||||
@@ -20,7 +20,7 @@ double median(std::vector<double> v) { std::sort(v.begin(), v.end()); return v[v
|
||||
float h16(float v) { return ref::f16_to_f32(ref::f32_to_f16(v)); }
|
||||
int unit_of(int t) {
|
||||
switch (t) {
|
||||
case ref::Q2_0: case ref::IQ3_S: case ref::BF16: return 16;
|
||||
case ref::Q2_0: case ref::IQ3_S: case ref::BF16: case ref::F32: return 16;
|
||||
default: return 32;
|
||||
}
|
||||
}
|
||||
@@ -58,7 +58,7 @@ int main(int argc, char** argv) {
|
||||
d.grid[0] = (n_out + tn - 1) / tn; d.grid[1] = (T + tm - 1) / tm; d.tg[0] = 128; ds.push_back(d);
|
||||
}
|
||||
};
|
||||
const std::vector<int> types = {ref::BF16, ref::Q8_0, ref::Q4_0, ref::Q5_0, ref::IQ4_NL, ref::IQ4_XS, ref::Q2_0,
|
||||
const std::vector<int> types = {ref::F32, ref::BF16, ref::Q8_0, ref::Q4_0, ref::Q5_0, ref::IQ4_NL, ref::IQ4_XS, ref::Q2_0,
|
||||
ref::IQ3_S, ref::Q3_K, ref::Q4_K, ref::Q5_K, ref::Q6_K};
|
||||
if (!bench) {
|
||||
int fails = 0;
|
||||
|
||||
@@ -36,7 +36,7 @@ struct MoeArgs {
|
||||
};
|
||||
|
||||
int type_of(const std::string& s) {
|
||||
for (int t : {ref::Q2_0, ref::IQ3_XXS, ref::IQ2_XXS, ref::IQ2_XS, ref::IQ2_S, ref::IQ3_S, ref::IQ4_NL, ref::IQ4_XS})
|
||||
for (int t : {ref::Q2_0, ref::IQ3_XXS, ref::IQ2_XXS, ref::IQ2_XS, ref::IQ2_S, ref::IQ3_S, ref::IQ4_NL, ref::IQ4_XS, ref::Q4_K, ref::Q5_K, ref::Q6_K, ref::Q8_0, ref::Q5_1})
|
||||
if (s == ref::type_name(t)) return t;
|
||||
std::fprintf(stderr, "unknown type %s\n", s.c_str());
|
||||
std::exit(2);
|
||||
|
||||
@@ -67,16 +67,17 @@ int main() {
|
||||
}
|
||||
// ---- router_route
|
||||
{
|
||||
for (int rty : {ref::BF16, ref::F32}) {
|
||||
const int NE = 512;
|
||||
const size_t rbr = ref::row_bytes(ref::BF16, N);
|
||||
const size_t rbr = ref::row_bytes(rty, N);
|
||||
mtl::Buffer Wr = cx.alloc(rbr * NE), lg1 = cx.alloc(4 * MAXT * NE), lg2 = cx.alloc(4 * MAXT * NE),
|
||||
id1 = cx.alloc(4 * MAXT * 10), id2 = cx.alloc(4 * MAXT * 10), w1 = cx.alloc(4 * MAXT * 10),
|
||||
w2 = cx.alloc(4 * MAXT * 10), pl1 = cx.alloc(4 * P_INTS), pl2 = cx.alloc(4 * P_INTS), cnt = cx.alloc(16);
|
||||
ref::random_blocks(ref::BF16, Wr.as<uint8_t>(), NE, N, 9, 0.05f);
|
||||
ref::random_blocks(rty, Wr.as<uint8_t>(), NE, N, 9, 0.05f);
|
||||
mtl::Pipeline pr = cx.pipeline(lib, "moe_route");
|
||||
for (int nt = 1; nt <= 4; ++nt) {
|
||||
mtl::Pipeline pg = cx.pipeline(lib, "gemv_30_32_1_4_1", mtl::Constants().i(0, nt)),
|
||||
pf = cx.pipeline(lib, "router_route_30", mtl::Constants().i(0, nt));
|
||||
mtl::Pipeline pg = cx.pipeline(lib, "gemv_" + std::to_string(rty) + "_32_1_4_1", mtl::Constants().i(0, nt)),
|
||||
pf = cx.pipeline(lib, "router_route_" + std::to_string(rty), mtl::Constants().i(0, nt));
|
||||
const size_t s16 = (size_t)4 * 4 * MAXT * N;
|
||||
GemvArgs ga{(uint32_t)N, (uint32_t)NE, (uint32_t)nt, (uint32_t)rbr, (uint32_t)N, (uint32_t)NE};
|
||||
MoeArgs ma{};
|
||||
@@ -111,7 +112,9 @@ int main() {
|
||||
cnt.as<uint32_t>()[0] == 0;
|
||||
}
|
||||
fails += !ok;
|
||||
std::printf("router_route ntok %d: logits, ids, w, plan bitwise (2 runs) %s\n", nt, ok ? "PASS" : "FAIL");
|
||||
std::printf("router_route %s ntok %d: logits, ids, w, plan bitwise (2 runs) %s\n", ref::type_name(rty), nt,
|
||||
ok ? "PASS" : "FAIL");
|
||||
}
|
||||
}
|
||||
}
|
||||
// ---- gr_norm_comb
|
||||
|
||||
@@ -28,7 +28,7 @@ int rows_per_tg(Cfg c) { return (c.S / c.K) * (32 / c.L) * c.R; }
|
||||
int unit_of(int t) {
|
||||
switch (t) {
|
||||
case ref::Q2_0: case ref::IQ3_S: case ref::IQ3_XXS: case ref::IQ2_XXS: case ref::IQ2_XS: case ref::IQ2_S:
|
||||
case ref::BF16: return 16;
|
||||
case ref::BF16: case ref::F32: return 16;
|
||||
default: return 32;
|
||||
}
|
||||
}
|
||||
@@ -163,7 +163,7 @@ int main(int argc, char** argv) {
|
||||
}
|
||||
if (!bench) {
|
||||
int fails = 0;
|
||||
for (int t : {ref::BF16, ref::Q8_0, ref::Q4_0, ref::Q5_0, ref::IQ4_NL, ref::IQ4_XS, ref::Q2_0, ref::IQ3_S,
|
||||
for (int t : {ref::F32, ref::BF16, ref::Q8_0, ref::Q4_0, ref::Q5_0, ref::IQ4_NL, ref::IQ4_XS, ref::Q2_0, ref::IQ3_S,
|
||||
ref::Q3_K, ref::Q4_K, ref::Q5_K, ref::Q6_K}) {
|
||||
for (int n_in : {2560, 640, 10240}) {
|
||||
if (n_in % ref::block_elems(t)) continue;
|
||||
|
||||
@@ -39,7 +39,9 @@ int main(int argc, char** argv) {
|
||||
const int SETS = bench ? 48 : 1;
|
||||
const bool q8 = std::getenv("GR_Q8") != nullptr; // the Q8_0 form of the HC matrices (gr_down_q8_*, gr_up_q8)
|
||||
const int wty = q8 ? ref::Q8_0 : ref::BF16;
|
||||
const size_t wd_bytes = (size_t)LR * NIN * 2, wi_bytes = (size_t)HC * NIN * 2, wu_bytes = (size_t)NIN * LR * 2;
|
||||
// GR_INJ_F32 (with GR_Q8): the inject rows as F32 (the Unsloth files), read by gr_down_q8f_*
|
||||
const bool injf = q8 && std::getenv("GR_INJ_F32") != nullptr;
|
||||
const size_t wd_bytes = (size_t)LR * NIN * 2, wi_bytes = (size_t)HC * NIN * 4, wu_bytes = (size_t)NIN * LR * 2;
|
||||
const size_t set_bytes = wd_bytes + wi_bytes + wu_bytes + NIN * 4;
|
||||
mtl::Buffer W = cx.alloc(set_bytes * SETS);
|
||||
ref::Rng rng(3);
|
||||
@@ -57,7 +59,8 @@ int main(int argc, char** argv) {
|
||||
for (int m = 0; m < 3; ++m) {
|
||||
std::vector<float> f(ns[m]);
|
||||
ref::dequant(ref::BF16, b + offs[m], f.data(), ns[m]);
|
||||
ref::quantize_q8_0(f.data(), ns[m], b + offs[m]);
|
||||
if (m == 1 && injf) std::memcpy(b + offs[m], f.data(), 4 * ns[m]);
|
||||
else ref::quantize_q8_0(f.data(), ns[m], b + offs[m]);
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -78,7 +81,7 @@ int main(int argc, char** argv) {
|
||||
int fails = 0;
|
||||
for (int nt : ntoks) {
|
||||
mtl::Pipeline pn = cx.pipeline(lib, "gr_norm"),
|
||||
pd = cx.pipeline(lib, std::string(q8 ? "gr_down_q8_" : "gr_down_") + grd, mtl::Constants().i(1, nt)),
|
||||
pd = cx.pipeline(lib, std::string(injf ? "gr_down_q8f_" : q8 ? "gr_down_q8_" : "gr_down_") + grd, mtl::Constants().i(1, nt)),
|
||||
pu = cx.pipeline(lib, q8 ? "gr_up_q8" : "gr_up", mtl::Constants().i(1, nt));
|
||||
for (int apply = 0; apply < 2; ++apply) {
|
||||
GrArgs a{N, HC, LR, (uint32_t)nt, (uint32_t)apply, 1e-6f};
|
||||
@@ -122,7 +125,7 @@ int main(int argc, char** argv) {
|
||||
const uint8_t* b = W.as<uint8_t>();
|
||||
std::vector<float> wd((size_t)LR * NIN), wi((size_t)HC * NIN), wu((size_t)NIN * LR);
|
||||
ref::dequant(wty, b, wd.data(), wd.size());
|
||||
ref::dequant(wty, b + wd_bytes, wi.data(), wi.size());
|
||||
ref::dequant(injf ? ref::F32 : wty, b + wd_bytes, wi.data(), wi.size());
|
||||
ref::dequant(wty, b + wd_bytes + wi_bytes, wu.data(), wu.size());
|
||||
const float* wn = (const float*)(b + wd_bytes + wi_bytes + wu_bytes);
|
||||
Err eR, elo, einj, emix;
|
||||
|
||||
@@ -23,7 +23,7 @@ struct MmPlanArgs { uint32_t n_ent, topk, n_expert, tm, n_blocks; };
|
||||
struct MmArgs { uint64_t gate_off, up_off, down_off, gu_stride, d_stride; uint32_t gu_row, d_row, n_embd, n_ff; };
|
||||
double median(std::vector<double> v) { std::sort(v.begin(), v.end()); return v[v.size() / 2]; }
|
||||
int type_of(const std::string& s) {
|
||||
for (int t : {ref::Q2_0, ref::IQ3_XXS, ref::IQ2_XXS, ref::IQ2_XS, ref::IQ2_S, ref::IQ3_S, ref::IQ4_NL, ref::IQ4_XS})
|
||||
for (int t : {ref::Q2_0, ref::IQ3_XXS, ref::IQ2_XXS, ref::IQ2_XS, ref::IQ2_S, ref::IQ3_S, ref::IQ4_NL, ref::IQ4_XS, ref::Q4_K, ref::Q5_K, ref::Q6_K, ref::Q8_0, ref::Q5_1})
|
||||
if (s == ref::type_name(t)) return t;
|
||||
std::exit(2);
|
||||
}
|
||||
|
||||
Reference in new issue
Block a user