mirror of
https://github.com/Niko1221/Strata.git
synced 2026-11-07 18:35:04 -10:00
prefill: STRATA_PF_FUSED_KQ=1 - the fused native-expert path takes UD-Q4_K_XL (gfx11): gate/up Q4_K / Q5_K, down Q5_1 / Q8_0
The formats carry a minimum, so the activation blocks' spare bytes now hold the int8 code sums per 32 values (written by every path, read by none else). Layers MMQ does not cover on HIP (Q4_K) join the fused layout through MmqPlan::fo. Opt-in (default output bit-identical). Test pairs plus an int8-rounded model check; QA_SAME probe in the attention harness.
This commit is contained in:
Niko1221 committed
1 parent
65a48f0d57
commit
7440806987
4 files changed
+235
-49
No files matched your search
@@ -97,6 +97,17 @@ int run(int fmt, int64_t ctx, int64_t nq, int reps) { // fmt 1 int8, 0 fp16
|
||||
std::sort(v.begin(), v.end());
|
||||
std::copy(v.begin(), v.end(), sel);
|
||||
}
|
||||
// QA_SAME=1 (S23 probe): every query reads the same first-w cells - the best case of any K/V sharing between queries
|
||||
// (all loads hit the caches); QA_SAME=2: query i reads a window of w cells at i / 4 * 64 (neighbours share ~90%)
|
||||
if (const char* sm = std::getenv("QA_SAME")) {
|
||||
const int mode = std::atoi(sm);
|
||||
for (int64_t i = 0; i < nq; ++i) {
|
||||
const int64_t w = steps[i * k::kStepCount + k::kStepWidth], nkv = ctx - nq + i + 1;
|
||||
if (w == nkv) continue;
|
||||
const int64_t base = mode == 1 ? 0 : std::min<int64_t>(i / 4 * 64, nkv - w);
|
||||
for (int64_t c = 0; c < w; ++c) ids[i * cap + c] = (int32_t) (base + c);
|
||||
}
|
||||
}
|
||||
k::QsaAttnPools pl;
|
||||
if (fmt == 1) { pl.k_q = up(kq); pl.v_q = up(vq); pl.k_scale = up(ks); pl.v_scale = up(vs); }
|
||||
else { pl.k_pool = up(kh); pl.v_pool = up(vh); }
|
||||
|
||||
+139
-25
@@ -62,15 +62,28 @@ constexpr int GU_ROWS_K = 2560, D_ROWS_K = 640; // K of gate/up (n_embd) and o
|
||||
// the formats (ggml type ids)
|
||||
constexpr int T_IQ2_XXS = GGML_TYPE_IQ2_XXS, T_IQ2_XS = GGML_TYPE_IQ2_XS, T_IQ2_S = GGML_TYPE_IQ2_S,
|
||||
T_IQ3_XXS = GGML_TYPE_IQ3_XXS, T_IQ3_S = GGML_TYPE_IQ3_S, T_IQ4_XS = GGML_TYPE_IQ4_XS,
|
||||
T_IQ4_NL = GGML_TYPE_IQ4_NL, T_Q2_0 = GGML_TYPE_Q2_0;
|
||||
T_IQ4_NL = GGML_TYPE_IQ4_NL, T_Q2_0 = GGML_TYPE_Q2_0,
|
||||
// Aurora (opt-in, HIP gfx11 only): Unsloth UD-Q4_K_XL's experts - gate/up Q4_K (one layer Q5_K), down Q5_1
|
||||
// (five layers Q8_0). Their weights are w = d_w q + m_w (a minimum), so a sub-block's dot is
|
||||
// sum w x = d_x (d_w sum q a + m_w sum a):
|
||||
// the activation blocks carry sum a per 32 values in their eight spare bytes (quant_act_nat_kernel and
|
||||
// the H epilogue write it for every path; no other format reads it)
|
||||
T_Q4_K = GGML_TYPE_Q4_K, T_Q5_K = GGML_TYPE_Q5_K, T_Q5_1 = GGML_TYPE_Q5_1, T_Q8_0 = GGML_TYPE_Q8_0;
|
||||
static_assert(sizeof(block_iq2_xxs) == 66 && sizeof(block_iq2_xs) == 74 && sizeof(block_iq2_s) == 82 &&
|
||||
sizeof(block_iq3_xxs) == 98 && sizeof(block_iq3_s) == 110 && sizeof(block_iq4_xs) == 136 &&
|
||||
sizeof(block_iq4_nl) == 18 && sizeof(block_q2_0) == 18, "the block layouts this file decodes");
|
||||
sizeof(block_iq4_nl) == 18 && sizeof(block_q2_0) == 18 && sizeof(block_q4_K) == 144 &&
|
||||
sizeof(block_q5_K) == 176 && sizeof(block_q5_1) == 24 && sizeof(block_q8_0) == 34,
|
||||
"the block layouts this file decodes");
|
||||
|
||||
// block bytes, scale per 16 values, codebook bytes in shared memory
|
||||
__host__ __device__ constexpr int block_bytes(int t) {
|
||||
return t == T_IQ2_XXS ? 66 : t == T_IQ2_XS ? 74 : t == T_IQ2_S ? 82 : t == T_IQ3_XXS ? 98 : t == T_IQ3_S ? 110
|
||||
: t == T_IQ4_XS ? 136 : 18;
|
||||
: t == T_IQ4_XS ? 136 : t == T_Q4_K ? 144 : t == T_Q5_K ? 176 : t == T_Q5_1 ? 24 : t == T_Q8_0 ? 34 : 18;
|
||||
}
|
||||
// the formats with a minimum (w = d q + m), and the raw words a sub-block's load stage holds
|
||||
__host__ __device__ constexpr bool has_min(int t) { return t == T_Q4_K || t == T_Q5_K || t == T_Q5_1; }
|
||||
__host__ __device__ constexpr int raw_words(int t) {
|
||||
return t == T_Q4_K ? 10 : t == T_Q5_K ? 18 : t == T_Q5_1 ? 6 : t == T_Q8_0 ? 9 : 5;
|
||||
}
|
||||
__host__ __device__ constexpr bool per16(int t) { return t == T_IQ2_XS || t == T_IQ2_S; }
|
||||
__host__ __device__ constexpr int grid_bytes(int t) {
|
||||
@@ -107,9 +120,16 @@ __global__ void quant_act_nat_kernel(const float* __restrict__ x, int64_t nblk,
|
||||
a1 = fmaxf(a1, __shfl_xor_sync(0xffffffffu, a1, o));
|
||||
}
|
||||
uint8_t* out = xa + w * AB;
|
||||
out[lane] = (uint8_t) (int8_t) (a0 > 0.0f ? __float2int_rn(v0 * (127.0f / a0)) : 0);
|
||||
out[32 + lane] = (uint8_t) (int8_t) (a1 > 0.0f ? __float2int_rn(v1 * (127.0f / a1)) : 0);
|
||||
if (lane == 0) *(float4*) (out + 64) = make_float4(a0 / 127.0f, a1 / 127.0f, 0.0f, 0.0f);
|
||||
const int c0 = a0 > 0.0f ? __float2int_rn(v0 * (127.0f / a0)) : 0, c1 = a1 > 0.0f ? __float2int_rn(v1 * (127.0f / a1)) : 0;
|
||||
out[lane] = (uint8_t) (int8_t) c0;
|
||||
out[32 + lane] = (uint8_t) (int8_t) c1;
|
||||
int s0 = c0, s1 = c1; // the codes' sums per 32 values (the formats with a minimum read them)
|
||||
#pragma unroll
|
||||
for (int o = 16; o > 0; o >>= 1) {
|
||||
s0 += __shfl_xor_sync(0xffffffffu, s0, o);
|
||||
s1 += __shfl_xor_sync(0xffffffffu, s1, o);
|
||||
}
|
||||
if (lane == 0) *(float4*) (out + 64) = make_float4(a0 / 127.0f, a1 / 127.0f, (float) s0, (float) s1);
|
||||
}
|
||||
|
||||
#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ >= 800
|
||||
@@ -180,7 +200,12 @@ __device__ __forceinline__ void table16(uint32_t q4, const uint32_t (&t)[4], uin
|
||||
|
||||
// Raw bytes of sub-block `ib` of the block at `bp` (IQ: the 256-value super-block, ib 0..7; Q2_0: the 64-value block,
|
||||
// ib 0..1; IQ4_NL: the 32-value block).
|
||||
template <int T> __device__ __forceinline__ void load_unit(const uint8_t* bp, int ib, uint32_t (&w)[5]) {
|
||||
// 32 bits at a 2-byte aligned address (one load when it is 4-byte aligned)
|
||||
__device__ __forceinline__ uint32_t ld32a(const uint8_t* p) {
|
||||
return ((uintptr_t) p & 3) == 0 ? *(const uint32_t*) p : ld32(p);
|
||||
}
|
||||
template <int T, int NW> __device__ __forceinline__ void load_unit(const uint8_t* bp, int ib, uint32_t (&w)[NW]) {
|
||||
static_assert(NW >= raw_words(T), "raw words");
|
||||
if constexpr (T == T_IQ2_XXS) {
|
||||
w[0] = ld32(bp + 2 + 8 * ib); w[1] = ld32(bp + 6 + 8 * ib); w[2] = ld16(bp);
|
||||
} else if constexpr (T == T_IQ2_XS) {
|
||||
@@ -202,16 +227,69 @@ template <int T> __device__ __forceinline__ void load_unit(const uint8_t* bp, in
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 4; ++k) w[k] = ld32(bp + 2 + 4 * k);
|
||||
w[4] = ld16(bp);
|
||||
} else if constexpr (T == T_Q4_K || T == T_Q5_K) {
|
||||
// qs: 32 bytes of the sub-block pair (ib / 2), the nibble chosen in convert; w[8] = d | dmin << 16; w[9]: the
|
||||
// scale/min bytes (get_scale_min_k4's: ib < 4: sc[ib], sc[ib + 4]; else sc[ib + 4], sc[ib - 4], sc[ib]), ib << 24
|
||||
constexpr int QS = T == T_Q4_K ? 16 : 48;
|
||||
const uint8_t* q = bp + QS + 32 * (ib >> 1);
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 8; ++k) w[k] = ld32a(q + 4 * k);
|
||||
w[8] = ld32a(bp);
|
||||
const uint8_t* sc = bp + 4;
|
||||
w[9] = ib < 4 ? (uint32_t) sc[ib] | ((uint32_t) sc[ib + 4] << 8)
|
||||
: (uint32_t) sc[ib + 4] | ((uint32_t) sc[ib - 4] << 8) | ((uint32_t) sc[ib] << 16);
|
||||
w[9] |= (uint32_t) ib << 24;
|
||||
if constexpr (T == T_Q5_K) {
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 8; ++k) w[10 + k] = ld32a(bp + 16 + 4 * k); // qh
|
||||
}
|
||||
} else if constexpr (T == T_Q5_1) { // d, m, qh, qs[16]
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 4; ++k) w[k] = ld32a(bp + 8 + 4 * k);
|
||||
w[4] = ld32a(bp + 4);
|
||||
w[5] = ld32a(bp);
|
||||
} else if constexpr (T == T_Q8_0) { // d, qs[32]
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 8; ++k) w[k] = ld32(bp + 2 + 4 * k);
|
||||
w[8] = ld16(bp);
|
||||
} else { // Q2_0
|
||||
w[0] = ld32(bp + 2 + 8 * ib); w[1] = ld32(bp + 6 + 8 * ib); w[2] = ld16(bp);
|
||||
}
|
||||
}
|
||||
|
||||
// The sub-block as 32 int8 (q[0..7], natural order) and its scales (s0: values 0-15, s1: 16-31).
|
||||
template <int T>
|
||||
__device__ __forceinline__ void convert(const uint32_t (&w)[5], const uint8_t* grid, const uint32_t (&kv)[4],
|
||||
template <int T, int NW>
|
||||
__device__ __forceinline__ void convert(const uint32_t (&w)[NW], const uint8_t* grid, const uint32_t (&kv)[4],
|
||||
uint32_t (&q)[8], float& s0, float& s1) {
|
||||
if constexpr (T == T_IQ2_XXS) {
|
||||
if constexpr (T == T_Q4_K || T == T_Q5_K) {
|
||||
const int ib = (int) (w[9] >> 24), sh = 4 * (ib & 1);
|
||||
uint32_t sc, mn;
|
||||
if (ib < 4) { sc = w[9] & 63; mn = (w[9] >> 8) & 63; }
|
||||
else {
|
||||
sc = ((w[9] & 0xF) | (((w[9] >> 8) & 0xC0) >> 2)); // sc[ib + 4] & 15 | (sc[ib - 4] >> 6) << 4
|
||||
mn = (((w[9]) >> 4) & 0xF) | (((w[9] >> 16) & 0xC0) >> 2); // sc[ib + 4] >> 4 | (sc[ib] >> 6) << 4
|
||||
}
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 8; ++k) {
|
||||
q[k] = (w[k] >> sh) & 0x0F0F0F0Fu;
|
||||
if constexpr (T == T_Q5_K) q[k] |= ((w[10 + k] >> ib) & 0x01010101u) << 4;
|
||||
}
|
||||
s0 = half_at(w[8]) * (float) sc;
|
||||
s1 = -(half_at(w[8] >> 16) * (float) mn);
|
||||
} else if constexpr (T == T_Q5_1) {
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
const uint32_t lo = ((w[4] >> (4 * k)) & 0xF) * 0x00204081u, hi = ((w[4] >> (16 + 4 * k)) & 0xF) * 0x00204081u;
|
||||
q[k] = (w[k] & 0x0F0F0F0Fu) | ((lo & 0x01010101u) << 4);
|
||||
q[4 + k] = ((w[k] >> 4) & 0x0F0F0F0Fu) | ((hi & 0x01010101u) << 4);
|
||||
}
|
||||
s0 = half_at(w[5]);
|
||||
s1 = half_at(w[5] >> 16);
|
||||
} else if constexpr (T == T_Q8_0) {
|
||||
#pragma unroll
|
||||
for (int k = 0; k < 8; ++k) q[k] = w[k];
|
||||
s0 = s1 = half_at(w[8]);
|
||||
} else if constexpr (T == T_IQ2_XXS) {
|
||||
const uint2* g = (const uint2*) grid;
|
||||
#pragma unroll
|
||||
for (int l = 0; l < 4; ++l) {
|
||||
@@ -581,10 +659,10 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
: blob + geo.down_off + (size_t) (rbase + ur) * geo.d_row;
|
||||
auto unit = [&](int s) -> const uint8_t* {
|
||||
if (GU) return wrow + (s >> 2) * BS;
|
||||
return WT == T_IQ4_NL ? wrow + (2 * s + uj) * BS : wrow + s * BS;
|
||||
return (WT == T_IQ4_NL || WT == T_Q5_1 || WT == T_Q8_0) ? wrow + (2 * s + uj) * BS : wrow + s * BS;
|
||||
};
|
||||
auto sub = [&](int s) { return GU ? 2 * (s & 3) + uj : uj; };
|
||||
auto put = [&](const uint32_t (&raw)[5], int buf) {
|
||||
auto put = [&](const uint32_t (&raw)[raw_words(WT)], int buf) {
|
||||
uint32_t q[8];
|
||||
float s0, s1;
|
||||
convert<WT>(raw, sgrid, kv, q, s0, s1);
|
||||
@@ -598,7 +676,7 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
const int r = row0 + min(tid, nrows - 1); // rows past the tile's end repeat its last one
|
||||
srow[tid] = GU ? src[r] : r;
|
||||
}
|
||||
uint32_t raw[5] = {0, 0, 0, 0, 0};
|
||||
uint32_t raw[raw_words(WT)] = {};
|
||||
load_unit<WT>(unit(0), sub(0), raw);
|
||||
put(raw, 0);
|
||||
if (NS > 1) load_unit<WT>(unit(1), sub(1), raw);
|
||||
@@ -615,21 +693,23 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; ++i) acc[mt][nt][i] = 0.0f;
|
||||
uint4 bq[2][4];
|
||||
float2 bx[2];
|
||||
auto fetch_b = [&](int s, uint4 (&q)[2][4], float2 (&x)[2]) {
|
||||
float2 bx[2], bsum[2]; // the activation scales, and (the formats with a minimum) the codes' sums per 32 values
|
||||
auto fetch_b = [&](int s, uint4 (&q)[2][4], float2 (&x)[2], float2 (&sm)[2]) {
|
||||
#pragma unroll
|
||||
for (int nt = 0; nt < 2; ++nt) {
|
||||
const uint4* p = reinterpret_cast<const uint4*>(brow[nt] + s * AB);
|
||||
q[nt][0] = p[0]; q[nt][1] = p[1]; q[nt][2] = p[2]; q[nt][3] = p[3];
|
||||
x[nt] = *reinterpret_cast<const float2*>(brow[nt] + s * AB + 64);
|
||||
if constexpr (has_min(WT)) sm[nt] = *reinterpret_cast<const float2*>(brow[nt] + s * AB + 72);
|
||||
else sm[nt] = make_float2(0.0f, 0.0f);
|
||||
}
|
||||
};
|
||||
if (on) fetch_b(0, bq, bx);
|
||||
if (on) fetch_b(0, bq, bx, bsum);
|
||||
for (int s = 0; s < NS; ++s) {
|
||||
__syncthreads(); // stage s's weights are in buffer s & 1
|
||||
uint4 nbq[2][4];
|
||||
float2 nbx[2];
|
||||
if (on && s + 1 < NS) fetch_b(s + 1, nbq, nbx);
|
||||
float2 nbx[2], nbsum[2];
|
||||
if (on && s + 1 < NS) fetch_b(s + 1, nbq, nbx, nbsum);
|
||||
if (on) {
|
||||
const int bf = s & 1;
|
||||
#pragma unroll
|
||||
@@ -660,7 +740,10 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
} else {
|
||||
const nw_i8 d = nw_wmma(A1, B1, nw_wmma(A0, B0, m));
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 8; ++i) acc[mt][nt][i] = fmaf(w0[i] * dx, dotf(d[i]), acc[mt][nt][i]);
|
||||
for (int i = 0; i < 8; ++i) {
|
||||
acc[mt][nt][i] = fmaf(w0[i] * dx, dotf(d[i]), acc[mt][nt][i]);
|
||||
if constexpr (has_min(WT)) acc[mt][nt][i] = fmaf(w1[i] * dx, h ? bsum[nt].y : bsum[nt].x, acc[mt][nt][i]);
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -673,6 +756,7 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
#pragma unroll
|
||||
for (int nt = 0; nt < 2; ++nt) {
|
||||
bx[nt] = nbx[nt];
|
||||
bsum[nt] = nbsum[nt];
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) bq[nt][j] = nbq[nt][j];
|
||||
}
|
||||
@@ -703,13 +787,18 @@ native_w11_kernel(const Batch b, const NativeGeom geo, const Tables tb, const ui
|
||||
const float inv = am > 0.0f ? 127.0f / am : 0.0f;
|
||||
uint8_t* o = out + (size_t) (row0 + r) * (10 * AB) + (size_t) fb * AB;
|
||||
uint32_t wd[8] = {};
|
||||
int csum = 0;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 32; ++j)
|
||||
wd[j >> 2] |= (uint32_t) (uint8_t) (int8_t) __float2int_rn(hs[r][32 * hh + j] * inv) << (8 * (j & 3));
|
||||
for (int j = 0; j < 32; ++j) {
|
||||
const int c = __float2int_rn(hs[r][32 * hh + j] * inv);
|
||||
csum += c;
|
||||
wd[j >> 2] |= (uint32_t) (uint8_t) (int8_t) c << (8 * (j & 3));
|
||||
}
|
||||
uint4* o4 = reinterpret_cast<uint4*>(o + 32 * hh);
|
||||
o4[0] = make_uint4(wd[0], wd[1], wd[2], wd[3]);
|
||||
o4[1] = make_uint4(wd[4], wd[5], wd[6], wd[7]);
|
||||
*reinterpret_cast<float*>(o + 64 + 4 * hh) = am / 127.0f;
|
||||
*reinterpret_cast<float*>(o + 72 + 4 * hh) = (float) csum; // (the formats with a minimum)
|
||||
}
|
||||
}
|
||||
} else {
|
||||
@@ -815,10 +904,29 @@ const DevInfo& dev_info() {
|
||||
#endif
|
||||
}
|
||||
|
||||
// opt-in on top of STRATA_PF_FUSED=1: STRATA_PF_FUSED_KQ=1 takes UD-Q4_K_XL's formats (the output differs from the FP16
|
||||
// expert path's at rounding level; KL-gated, see aurora_s23.md)
|
||||
bool kq_on() {
|
||||
static const bool on = [] {
|
||||
const char* v = std::getenv("STRATA_PF_FUSED_KQ");
|
||||
return v != nullptr && v[0] == '1';
|
||||
}();
|
||||
return on;
|
||||
}
|
||||
bool gu_covered(int t) {
|
||||
return t == T_IQ2_XXS || t == T_IQ2_XS || t == T_IQ2_S || t == T_IQ3_XXS || t == T_IQ3_S || t == T_IQ4_XS;
|
||||
return t == T_IQ2_XXS || t == T_IQ2_XS || t == T_IQ2_S || t == T_IQ3_XXS || t == T_IQ3_S || t == T_IQ4_XS
|
||||
#if defined(__HIPCC__)
|
||||
|| ((t == T_Q4_K || t == T_Q5_K) && kq_on()) // the gfx11 kernels only
|
||||
#endif
|
||||
;
|
||||
}
|
||||
bool d_covered(int t) {
|
||||
return t == T_Q2_0 || t == T_IQ4_NL
|
||||
#if defined(__HIPCC__)
|
||||
|| ((t == T_Q5_1 || t == T_Q8_0) && kq_on())
|
||||
#endif
|
||||
;
|
||||
}
|
||||
bool d_covered(int t) { return t == T_Q2_0 || t == T_IQ4_NL; }
|
||||
|
||||
template <int T, bool GU>
|
||||
void launch(int ww, unsigned grid, const Batch& b, const NativeGeom& g, const Tables& tb, const void* act,
|
||||
@@ -886,11 +994,17 @@ void experts_native(const Batch& b, const NativeGeom& g, int n_expert, int64_t n
|
||||
case T_IQ2_S: STRATA_NW_GU(T_IQ2_S); break;
|
||||
case T_IQ3_XXS: STRATA_NW_GU(T_IQ3_XXS); break;
|
||||
case T_IQ3_S: STRATA_NW_GU(T_IQ3_S); break;
|
||||
case T_Q4_K: STRATA_NW_GU(T_Q4_K); break;
|
||||
case T_Q5_K: STRATA_NW_GU(T_Q5_K); break;
|
||||
default: STRATA_NW_GU(T_IQ4_XS); break;
|
||||
}
|
||||
#undef STRATA_NW_GU
|
||||
if (g.d_type == T_Q2_0) native_w11_kernel<T_Q2_0, false><<<g_d, NW_THREADS, 0, s>>>(b, g, tb, ha8, src, nullptr, dm);
|
||||
else native_w11_kernel<T_IQ4_NL, false><<<g_d, NW_THREADS, 0, s>>>(b, g, tb, ha8, src, nullptr, dm);
|
||||
#define STRATA_NW_D(T) native_w11_kernel<T, false><<<g_d, NW_THREADS, 0, s>>>(b, g, tb, ha8, src, nullptr, dm)
|
||||
if (g.d_type == T_Q2_0) STRATA_NW_D(T_Q2_0);
|
||||
else if (g.d_type == T_Q5_1) STRATA_NW_D(T_Q5_1);
|
||||
else if (g.d_type == T_Q8_0) STRATA_NW_D(T_Q8_0);
|
||||
else STRATA_NW_D(T_IQ4_NL);
|
||||
#undef STRATA_NW_D
|
||||
ck(cudaGetLastError(), "experts_native");
|
||||
return;
|
||||
}
|
||||
|
||||
+11
-2
@@ -559,6 +559,8 @@ constexpr size_t MMQ_TAIL = 4096;
|
||||
struct MmqPlan {
|
||||
bool any = false, fallback = true;
|
||||
std::vector<char> layer; // per layer: MMQ
|
||||
std::vector<char> fo; // per layer: no MMQ here but the native fused kernels cover its formats
|
||||
// (gfx11: UD-Q4_K_XL's Q4_K / Q5_K and Q5_1 / Q8_0 experts, STRATA_PF_FUSED=1)
|
||||
size_t gu_max = 0, d_max = 0;
|
||||
};
|
||||
const MmqPlan& mmq_plan() {
|
||||
@@ -569,11 +571,17 @@ const MmqPlan& mmq_plan() {
|
||||
const bool on = mmq::built() && (env == nullptr || std::atoi(env) != 0);
|
||||
const int64_t layers = lay.native ? (int64_t) lay.fmt.size() : lay.n_layers;
|
||||
p.layer.assign((size_t) std::max<int64_t>(layers, 0), 0);
|
||||
p.fo.assign(p.layer.size(), 0);
|
||||
p.fallback = !on || layers <= 0;
|
||||
for (int64_t l = 0; on && l < layers; ++l) {
|
||||
const int gt = lay.native ? lay.fmt[(size_t) l].gu_type : 42, dt = lay.native ? lay.fmt[(size_t) l].d_type : 42;
|
||||
// #420: a tile on every GPU for these shapes (gate+up: 1280 rows, down: N rows), else the FP16 path
|
||||
if (!mmq::fits(gt, 1280) || !mmq::fits(dt, N)) { p.fallback = true; continue; }
|
||||
if (!mmq::fits(gt, 1280) || !mmq::fits(dt, N)) {
|
||||
p.fallback = true;
|
||||
// the fused path's buffers (Xq, H) exist for it; a chunk it does not take (a small one) keeps the FP16 path
|
||||
if (on && lay.native && fused::native_supported(gt, dt)) { p.fo[(size_t) l] = 1; p.any = true; }
|
||||
continue;
|
||||
}
|
||||
p.layer[(size_t) l] = 1;
|
||||
p.any = true;
|
||||
p.gu_max = std::max(p.gu_max, mmq::matrix_bytes(gt, 1280, N));
|
||||
@@ -1758,7 +1766,8 @@ bool Prefill::run(const int64_t* tokens, int64_t n, int64_t pos0, std::string& e
|
||||
const bool use_mmq = mmq_plan().any && mmq_plan().layer[(size_t) l];
|
||||
const int mmq_gt = lay.native ? lay.fmt[(size_t) l].gu_type : 42;
|
||||
const int mmq_dt = lay.native ? lay.fmt[(size_t) l].d_type : 42;
|
||||
const bool fused_nat = use_mmq && stream_all && lay.native && fused::native_supported(mmq_gt, mmq_dt);
|
||||
const bool fused_only = mmq_plan().any && mmq_plan().fo[(size_t) l];
|
||||
const bool fused_nat = (use_mmq || fused_only) && stream_all && lay.native && fused::native_supported(mmq_gt, mmq_dt);
|
||||
const bool fused_l = (use_mmq && stream_all && !lay.native && fused::enabled()) || fused_nat;
|
||||
size_t n_order = 0; // the routed experts (the debug report; unknown when fused)
|
||||
if (fused_l) {
|
||||
|
||||
@@ -57,6 +57,8 @@ struct Pair {
|
||||
ggml_type gu, d;
|
||||
const char* name;
|
||||
};
|
||||
// MMQ covers this pair (the HIP build has no Q4_K / Q5_K / Q5_1 MMQ: those pairs are checked against the reference only)
|
||||
bool mmq_ok(const Pair& p) { return mmq::supported((int) p.gu) && mmq::supported((int) p.d); }
|
||||
|
||||
struct Geo {
|
||||
size_t gu_row, d_row, up_off, down_off, bytes;
|
||||
@@ -80,13 +82,18 @@ struct Geo {
|
||||
void fill_matrix(uint8_t* m, ggml_type t, int64_t rows, int64_t cols, std::mt19937& rng) {
|
||||
const size_t bs = ggml_type_size(t), nb = (size_t) rows * (size_t) (cols / ggml_blck_size(t));
|
||||
std::uniform_int_distribution<int> byte(0, 255);
|
||||
const float lo = t == GGML_TYPE_IQ3_S ? 0.0004f : t == GGML_TYPE_IQ4_XS ? 0.00005f
|
||||
const float lo = t == GGML_TYPE_Q4_K || t == GGML_TYPE_Q5_K ? 0.00004f : t == GGML_TYPE_Q5_1 ? 0.0006f
|
||||
: t == GGML_TYPE_Q8_0 ? 0.0001f : t == GGML_TYPE_IQ3_S ? 0.0004f : t == GGML_TYPE_IQ4_XS ? 0.00005f
|
||||
: t == GGML_TYPE_IQ4_NL ? 0.0002f : t == GGML_TYPE_Q2_0 ? 0.004f : 0.0008f;
|
||||
std::uniform_real_distribution<float> sc(lo, 6.0f * lo);
|
||||
for (size_t i = 0; i < nb * bs; ++i) m[i] = (uint8_t) byte(rng);
|
||||
for (size_t b = 0; b < nb; ++b) {
|
||||
const ggml_fp16_t h = ggml_fp32_to_fp16(sc(rng));
|
||||
std::memcpy(m + b * bs, &h, 2);
|
||||
if (t == GGML_TYPE_Q4_K || t == GGML_TYPE_Q5_K || t == GGML_TYPE_Q5_1) { // the minimum's scale, either sign
|
||||
const ggml_fp16_t h2 = ggml_fp32_to_fp16(sc(rng) * ((rng() & 1) ? 1.0f : -1.0f) * (t == GGML_TYPE_Q5_1 ? 4.0f : 1.0f));
|
||||
std::memcpy(m + b * bs + 2, &h2, 2);
|
||||
}
|
||||
}
|
||||
}
|
||||
std::vector<uint8_t> make_blob(const Pair& p, const Geo& g, std::mt19937& rng) {
|
||||
@@ -271,6 +278,16 @@ Err compare(const std::vector<float>& a, RA row_a, const std::vector<float>& b,
|
||||
return out;
|
||||
}
|
||||
|
||||
// the kernels' int8 rounding per 32 values (quant_act_nat_kernel / the H epilogue), as floats: codes * (amax / 127)
|
||||
void quant32(const float* in, float* out, int n) {
|
||||
for (int b = 0; b < n; b += 32) {
|
||||
float am = 0.0f;
|
||||
for (int j = 0; j < 32; ++j) am = std::max(am, std::fabs(in[b + j]));
|
||||
const float inv = am > 0.0f ? 127.0f / am : 0.0f, d = am / 127.0f;
|
||||
for (int j = 0; j < 32; ++j) out[b + j] = (float) std::nearbyint(in[b + j] * inv) * d;
|
||||
}
|
||||
}
|
||||
|
||||
std::vector<float> download(const Dev& d, size_t n) {
|
||||
std::vector<float> h(n);
|
||||
ck(cudaMemcpy(h.data(), d.p, n * 4, cudaMemcpyDeviceToHost), "download");
|
||||
@@ -310,27 +327,30 @@ void reference_part(const Pair& p, cudaStream_t s) {
|
||||
// (the uploads above are on the legacy stream and the work below on a non-blocking one: a cudaMemcpy from pageable
|
||||
// memory may return before its DMA lands - on Windows it can wait in the queue until the next synchronous call)
|
||||
ck(cudaDeviceSynchronize(), "uploads");
|
||||
const bool mq = mmq_ok(p);
|
||||
mmq::Context ctx;
|
||||
MmqBufs mb(p, rows, E);
|
||||
run_mmq(p, geo, ctx, mb, r, x_dev.as<float>(), src_dev.as<int32_t>(), blob, rows, s);
|
||||
MmqBufs* mbp = mq ? new MmqBufs(p, rows, E) : nullptr;
|
||||
if (mq) run_mmq(p, geo, ctx, *mbp, r, x_dev.as<float>(), src_dev.as<int32_t>(), blob, rows, s);
|
||||
FusedBufs fb(T, rows, E);
|
||||
ck(cudaMemcpy(fb.ids.p, ids.data(), ids.size() * 4, cudaMemcpyHostToDevice), "ids");
|
||||
ck(cudaMemset(fb.dm.p, 0xff, (size_t) rows * N * 4), "sentinel"); // an unwritten row is NaN
|
||||
ck(cudaDeviceSynchronize(), "uploads");
|
||||
run_fused(p, geo, fb, x_dev.as<float>(), T, E, {0, 1, 20, E}, blob, s);
|
||||
ck(cudaStreamSynchronize(s), "sync");
|
||||
const std::vector<float> y_mmq = download(mb.dm, (size_t) rows * N), y_f = download(fb.dm, (size_t) rows * N);
|
||||
const std::vector<float> y_mmq = mq ? download(mbp->dm, (size_t) rows * N) : std::vector<float>((size_t) rows * N, 0.0f),
|
||||
y_f = download(fb.dm, (size_t) rows * N);
|
||||
delete mbp;
|
||||
const std::vector<int32_t> slot = download_i(fb.slot, (size_t) rows), fsrc = download_i(fb.src, (size_t) rows);
|
||||
for (int64_t i = 0; i < rows; ++i)
|
||||
if (fsrc[(size_t) slot[(size_t) i]] != (int32_t) (i / K)) throw std::runtime_error("group: slot/src disagree");
|
||||
|
||||
// the reference, an expert at a time on all cores
|
||||
std::vector<float> ref((size_t) rows * N);
|
||||
std::vector<float> ref((size_t) rows * N), ref2((size_t) rows * N); // ref2: the same with the kernels' int8 rounding
|
||||
std::vector<std::thread> th;
|
||||
const int nth = std::max(1, std::min(16, (int) std::thread::hardware_concurrency()));
|
||||
for (int w = 0; w < nth; ++w)
|
||||
th.emplace_back([&, w] {
|
||||
std::vector<float> gu, dn, h(FF);
|
||||
std::vector<float> gu, dn, h(FF), h2(FF), hq(FF), xq(N);
|
||||
for (int e = w; e < E; e += nth) {
|
||||
if (r.cnt[(size_t) e] == 0) continue;
|
||||
dequant(p, geo, host[(size_t) e], gu, dn);
|
||||
@@ -349,6 +369,20 @@ void reference_part(const Pair& p, cudaStream_t s) {
|
||||
for (int f = 0; f < FF; ++f) a += (double) wd[f] * h[(size_t) f];
|
||||
ref[(size_t) i * N + o] = (float) a;
|
||||
}
|
||||
quant32(xr, xq.data(), N);
|
||||
for (int f = 0; f < FF; ++f) {
|
||||
double gt = 0, up = 0;
|
||||
const float *wg = gu.data() + (size_t) f * N, *wu = gu.data() + (size_t) (FF + f) * N;
|
||||
for (int k = 0; k < N; ++k) { gt += (double) wg[k] * xq[(size_t) k]; up += (double) wu[k] * xq[(size_t) k]; }
|
||||
h2[(size_t) f] = (float) (gt / (1.0 + std::exp(-gt)) * up);
|
||||
}
|
||||
quant32(h2.data(), hq.data(), FF);
|
||||
for (int o = 0; o < N; ++o) {
|
||||
double a = 0;
|
||||
const float* wd = dn.data() + (size_t) o * FF;
|
||||
for (int f = 0; f < FF; ++f) a += (double) wd[f] * hq[(size_t) f];
|
||||
ref2[(size_t) i * N + o] = (float) a;
|
||||
}
|
||||
}
|
||||
}
|
||||
});
|
||||
@@ -357,17 +391,27 @@ void reference_part(const Pair& p, cudaStream_t s) {
|
||||
auto pair = [](size_t i) { return (int64_t) i; };
|
||||
auto mrow = [&](size_t i) { return (int64_t) r.row_of[i]; };
|
||||
auto frow = [&](size_t i) { return (int64_t) slot[i]; };
|
||||
const Err em = compare(y_mmq, mrow, ref, pair, (size_t) rows);
|
||||
const Err em = mq ? compare(y_mmq, mrow, ref, pair, (size_t) rows) : Err{};
|
||||
const Err ef = compare(y_f, frow, ref, pair, (size_t) rows);
|
||||
const Err efm = compare(y_f, frow, y_mmq, mrow, (size_t) rows);
|
||||
std::printf(" MMQ vs FP32 reference: rel RMS %.3e worst row max rel %.3e\n", em.rms, em.worst);
|
||||
const Err efm = mq ? compare(y_f, frow, y_mmq, mrow, (size_t) rows) : Err{};
|
||||
const Err e2 = compare(y_f, frow, ref2, pair, (size_t) rows);
|
||||
std::printf(" fused vs the int8-rounded model (same activation / H rounding, double sums): rel RMS %.3e worst %.3e\n",
|
||||
e2.rms, e2.worst);
|
||||
if (e2.rms > 2e-3 || e2.worst > 3e-2)
|
||||
throw std::runtime_error(std::string(p.name) + ": the fused path differs from its own arithmetic's model");
|
||||
if (mq) std::printf(" MMQ vs FP32 reference: rel RMS %.3e worst row max rel %.3e\n", em.rms, em.worst);
|
||||
std::printf(" fused vs FP32 reference: rel RMS %.3e worst row max rel %.3e\n", ef.rms, ef.worst);
|
||||
std::printf(" fused vs MMQ : rel RMS %.3e worst row max rel %.3e\n", efm.rms, efm.worst);
|
||||
if (mq) std::printf(" fused vs MMQ : rel RMS %.3e worst row max rel %.3e\n", efm.rms, efm.worst);
|
||||
double zmax = 0;
|
||||
for (int k = 0; k < K; ++k)
|
||||
for (int o = 0; o < N; ++o)
|
||||
zmax = std::max(zmax, (double) std::fabs(y_f[(size_t) slot[(size_t) (ZERO * K + k)] * N + o]));
|
||||
if (zmax != 0) throw std::runtime_error(std::string(p.name) + ": the all-zero token's outputs are not zero");
|
||||
if (!mq) { // no MMQ here: the int8 activations' own error (the IQ pairs' fused path sits at ~1e-2 rel RMS)
|
||||
if (ef.rms > 0.03 || ef.worst > 0.25)
|
||||
throw std::runtime_error(std::string(p.name) + ": the fused path's error against the reference is too large");
|
||||
return;
|
||||
}
|
||||
if (ef.rms > 1.5 * em.rms || ef.worst > 2.0 * em.worst)
|
||||
throw std::runtime_error(std::string(p.name) + ": the fused path's error is not comparable to MMQ's");
|
||||
}
|
||||
@@ -401,8 +445,10 @@ void timing_part(const Pair& p, int T, cudaStream_t s) {
|
||||
Dev x_dev(x.size() * 4), src_dev((size_t) rows * 4);
|
||||
ck(cudaMemcpy(x_dev.p, x.data(), x.size() * 4, cudaMemcpyHostToDevice), "x");
|
||||
ck(cudaMemcpy(src_dev.p, r.src.data(), r.src.size() * 4, cudaMemcpyHostToDevice), "src");
|
||||
const bool mq = mmq_ok(p);
|
||||
mmq::Context ctx;
|
||||
MmqBufs mb(p, rows, E);
|
||||
MmqBufs* mbp = mq ? new MmqBufs(p, rows, E) : nullptr;
|
||||
MmqBufs& mb = *mbp; // (only used when mq)
|
||||
FusedBufs fb(T, rows, E);
|
||||
ck(cudaMemcpy(fb.ids.p, ids.data(), ids.size() * 4, cudaMemcpyHostToDevice), "ids");
|
||||
ck(cudaDeviceSynchronize(), "uploads"); // (see reference_part)
|
||||
@@ -420,7 +466,7 @@ void timing_part(const Pair& p, int T, cudaStream_t s) {
|
||||
};
|
||||
auto mmq_run = [&] { run_mmq(p, geo, ctx, mb, r, x_dev.as<float>(), src_dev.as<int32_t>(), blob, rows, s); };
|
||||
auto fused_run = [&] { run_fused(p, geo, fb, x_dev.as<float>(), T, E, {0, E}, blob, s); };
|
||||
const bool all_routed = std::count(r.cnt.begin(), r.cnt.end(), 0) == 0;
|
||||
const bool all_routed = mq && std::count(r.cnt.begin(), r.cnt.end(), 0) == 0;
|
||||
auto direct_run = [&] {
|
||||
run_mmq(p, geo, ctx, mb, r, x_dev.as<float>(), src_dev.as<int32_t>(), blob, rows, s, true, stride);
|
||||
};
|
||||
@@ -446,7 +492,7 @@ void timing_part(const Pair& p, int T, cudaStream_t s) {
|
||||
}
|
||||
}
|
||||
// ~1 s of both first (the clocks ramp up), then alternating rounds; the medians
|
||||
const bool no_mmq = std::getenv("S20_NOMMQ") != nullptr; // (debug: the fused runs alone)
|
||||
const bool no_mmq = !mq || std::getenv("S20_NOMMQ") != nullptr; // (debug: the fused runs alone)
|
||||
for (float spent = 0; spent < 1000.0f;) spent += (no_mmq ? 0.0f : once(mmq_run)) + once(fused_run);
|
||||
std::vector<float> tm, tf, td;
|
||||
for (int rep = 0; rep < 8; ++rep) {
|
||||
@@ -467,10 +513,11 @@ void timing_part(const Pair& p, int T, cudaStream_t s) {
|
||||
std::sort(tf2.begin(), tf2.end());
|
||||
std::printf(" (fused back to back: median %.3f ms, min %.3f; alternating min %.3f; MMQ min %.3f)\n", tf2[4],
|
||||
tf2[0], tf[0], tm[0]);
|
||||
const std::vector<float> y_mmq = download(mb.dm, (size_t) rows * N), y_f = download(fb.dm, (size_t) rows * N);
|
||||
const std::vector<float> y_mmq = mq ? download(mb.dm, (size_t) rows * N) : std::vector<float>((size_t) rows * N, 0.0f),
|
||||
y_f = download(fb.dm, (size_t) rows * N);
|
||||
const std::vector<int32_t> slot = download_i(fb.slot, (size_t) rows);
|
||||
const Err efm = compare(y_f, [&](size_t i) { return (int64_t) slot[i]; }, y_mmq,
|
||||
[&](size_t i) { return (int64_t) r.row_of[i]; }, (size_t) rows);
|
||||
const Err efm = mq ? compare(y_f, [&](size_t i) { return (int64_t) slot[i]; }, y_mmq,
|
||||
[&](size_t i) { return (int64_t) r.row_of[i]; }, (size_t) rows) : Err{};
|
||||
std::printf("%s - timing part: %d tokens x top %d over %d experts, one layer: MMQ path %.3f ms, fused %.3f ms "
|
||||
"(%.2fx); fused vs MMQ rel RMS %.3e worst row %.3e\n",
|
||||
p.name, T, K, E, t_mmq, t_f, t_mmq / t_f, efm.rms, efm.worst);
|
||||
@@ -560,18 +607,19 @@ void timing_part(const Pair& p, int T, cudaStream_t s) {
|
||||
for (int f = 0; f < FF; ++f) a += (double) wd[f] * h[(size_t) f];
|
||||
ref[j * N + o] = (float) a;
|
||||
}
|
||||
std::copy_n(y_mmq.begin() + (ptrdiff_t) ((size_t) r.row_of[i] * N), N, ym.begin() + (ptrdiff_t) (j * N));
|
||||
if (mq) std::copy_n(y_mmq.begin() + (ptrdiff_t) ((size_t) r.row_of[i] * N), N, ym.begin() + (ptrdiff_t) (j * N));
|
||||
std::copy_n(y_f.begin() + (ptrdiff_t) ((size_t) slot[i] * N), N, yf.begin() + (ptrdiff_t) (j * N));
|
||||
}
|
||||
auto id = [](size_t j) { return (int64_t) j; };
|
||||
const Err em = compare(ym, id, ref, id, pick.size()), ef = compare(yf, id, ref, id, pick.size());
|
||||
const Err em = mq ? compare(ym, id, ref, id, pick.size()) : Err{}, ef = compare(yf, id, ref, id, pick.size());
|
||||
std::printf(" sample of %zu pairs vs FP32 reference: MMQ rel RMS %.3e worst %.3e, fused rel RMS %.3e worst "
|
||||
"%.3e\n", pick.size(), em.rms, em.worst, ef.rms, ef.worst);
|
||||
if (ef.rms > 1.5 * em.rms || ef.worst > 2.0 * em.worst)
|
||||
if (!mq ? (ef.rms > 0.03 || ef.worst > 0.25) : (ef.rms > 1.5 * em.rms || ef.worst > 2.0 * em.worst))
|
||||
throw std::runtime_error(std::string(p.name) + ": the fused path's error is not comparable to MMQ's (timing)");
|
||||
}
|
||||
cudaEventDestroy(a);
|
||||
cudaEventDestroy(b);
|
||||
delete mbp;
|
||||
if (efm.rms > 0.05) throw std::runtime_error(std::string(p.name) + ": the paths disagree at the real shape");
|
||||
}
|
||||
} // namespace
|
||||
@@ -582,6 +630,7 @@ int main(int argc, char** argv) {
|
||||
_putenv_s("STRATA_PF_FUSED", "1");
|
||||
#else
|
||||
setenv("STRATA_PF_FUSED", "1", 1);
|
||||
setenv("STRATA_PF_FUSED_KQ", "1", 1); // (the HIP build's Q4_K / Q5_K / Q5_1 / Q8_0 pairs)
|
||||
#endif
|
||||
int n = 0;
|
||||
if (cudaGetDeviceCount(&n) != cudaSuccess || n == 0) { std::printf("no CUDA device: skipped\n"); return 77; }
|
||||
@@ -590,7 +639,10 @@ int main(int argc, char** argv) {
|
||||
const std::vector<Pair> pairs = {
|
||||
{GGML_TYPE_IQ2_S, GGML_TYPE_Q2_0, "IQ2_S / Q2_0"}, {GGML_TYPE_IQ2_XXS, GGML_TYPE_Q2_0, "IQ2_XXS / Q2_0"},
|
||||
{GGML_TYPE_IQ2_XS, GGML_TYPE_IQ4_NL, "IQ2_XS / IQ4_NL"}, {GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ4_NL, "IQ3_XXS / IQ4_NL"},
|
||||
{GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_NL, "IQ3_S / IQ4_NL"}, {GGML_TYPE_IQ4_XS, GGML_TYPE_Q2_0, "IQ4_XS / Q2_0"}};
|
||||
{GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_NL, "IQ3_S / IQ4_NL"}, {GGML_TYPE_IQ4_XS, GGML_TYPE_Q2_0, "IQ4_XS / Q2_0"},
|
||||
// (HIP only) UD-Q4_K_XL's: 47 layers Q4_K / Q5_1, 4 Q4_K / Q8_0, and the one Q5_K layer
|
||||
{GGML_TYPE_Q4_K, GGML_TYPE_Q5_1, "Q4_K / Q5_1"}, {GGML_TYPE_Q4_K, GGML_TYPE_Q8_0, "Q4_K / Q8_0"},
|
||||
{GGML_TYPE_Q5_K, GGML_TYPE_Q5_1, "Q5_K / Q5_1"}, {GGML_TYPE_Q5_K, GGML_TYPE_Q8_0, "Q5_K / Q8_0"}};
|
||||
const char* only = nullptr; // --only=NAME: the pairs whose name contains NAME
|
||||
for (int i = 1; i < argc; ++i)
|
||||
if (std::strncmp(argv[i], "--only=", 7) == 0) only = argv[i] + 7;
|
||||
@@ -611,14 +663,14 @@ int main(int argc, char** argv) {
|
||||
cudaStream_t s = nullptr;
|
||||
ck(cudaStreamCreateWithFlags(&s, cudaStreamNonBlocking), "stream");
|
||||
for (const Pair& p : pairs) {
|
||||
if (only && std::string(p.name).find(only) == std::string::npos) continue;
|
||||
if (only ? std::string(p.name).find(only) == std::string::npos : p.gu == GGML_TYPE_Q4_K || p.gu == GGML_TYPE_Q5_K) continue;
|
||||
if (!fused::native_supported((int) p.gu, (int) p.d))
|
||||
throw std::runtime_error(std::string(p.name) + ": not covered by the native kernels");
|
||||
if (ref) reference_part(p, s);
|
||||
}
|
||||
if (timing) {
|
||||
// the IQ2_XS pack's layers (34 of 48 IQ2_S / Q2_0, 11 IQ2_XXS / Q2_0), the IQ3_S pack's most common, IQ3_XXS's
|
||||
for (const Pair& p : {pairs[0], pairs[1], pairs[4], pairs[3]}) {
|
||||
for (const Pair& p : only ? pairs : std::vector<Pair>{pairs[0], pairs[1], pairs[4], pairs[3]}) {
|
||||
if (only && std::string(p.name).find(only) == std::string::npos) continue;
|
||||
for (int T : chunks) timing_part(p, T, s);
|
||||
}
|
||||
|
||||
Reference in new issue
Block a user