pub const BLOCKS: &str = r#"
struct lk_w { const signed char* q; const __half* s; const unsigned char* b; long rb; };
// A part's outputs: an `f32` array or a binary16 table, set from either pointer; the host reads an `f32` part's back as
// its pointer, a kernel through `lk_part_out`.
struct lk_part_y {
void* p = nullptr;
bool h = false;
lk_part_y() = default;
__host__ __device__ lk_part_y(std::nullptr_t) {}
__host__ __device__ lk_part_y(float* q) : p(q) {}
__host__ __device__ lk_part_y(__half* q) : p(q), h(true) {}
__host__ operator float*() const { return (float*)p; }
};
// Up to `LK_PARTS` matrices one activation is projected by (the planner's `GEMM_PARTS`), their outputs side by side:
// part `p`'s matrix `w[p]`, outputs `y[p]` (rows from the one-based position `row[p]` names, `pitch[p]` apart, where it
// names one), bias `bias[p]`, and its outputs `end[p − 1] … end[p] − 1` of the launch's.
#define LK_PARTS 6
struct lk_parts {
lk_w w[LK_PARTS];
lk_part_y y[LK_PARTS];
const float* bias[LK_PARTS];
int end[LK_PARTS];
const int* row[LK_PARTS];
long pitch[LK_PARTS];
};
__device__ __forceinline__ lk_out lk_part_out(const lk_parts& parts, int p){
return parts.y[p].h ? lk_out((__half*)parts.y[p].p) : lk_out((float*)parts.y[p].p);
}
// The little-endian f16 at `p`, as an `f32`.
__device__ __forceinline__ float lk_h(const unsigned char* p){ return __half2float(__ushort_as_half((unsigned short)(p[0] | (p[1] << 8)))); }
__device__ __forceinline__ float lk_e8m0_half(unsigned char x){
unsigned int bits = x < 2 ? 0x00200000u << x : ((unsigned int)(x - 1)) << 23;
return __uint_as_float(bits);
}
__device__ __forceinline__ void lk_scale_min_k4(int j, const unsigned char* q, int& sc, int& m){
if(j < 4){ sc = q[j] & 63; m = q[j + 4] & 63; }
else { sc = (q[j + 4] & 0xf) | ((q[j - 4] >> 6) << 4); m = (q[j + 4] >> 4) | ((q[j] >> 6) << 4); }
}
// Word `i` of four, by selects rather than an index, so the four stay in registers.
__device__ __forceinline__ unsigned int lk_pick4(const unsigned int* a, int i){ return i == 0 ? a[0] : i == 1 ? a[1] : i == 2 ? a[2] : a[3]; }
// The IQ4 formats' sixteen non-linear levels (ggml's `kvalues_iq4nl`): one by its nibble, and four to a word.
__device__ __forceinline__ int lk_iq4_level(int n){
const signed char kv[16] = {-127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113};
return kv[n];
}
// The levels of the low nibbles of `q`'s four bytes (`lo`) and of their high nibbles (`hi`), four to a word, from a
// table of sixteen int8 levels four to a word (`t0` levels 0–3 … `t3` 12–15), by byte permutes: each nibble's low three
// bits pick a byte from levels 0–7 and from 8–15, its fourth bit picks between them — no memory read.
__device__ __forceinline__ void lk_table16(unsigned int q, unsigned int t0, unsigned int t1, unsigned int t2, unsigned int t3,
int& lo, int& hi){
unsigned int t[2], sel = 0x32103210u | ((q & 0x88888888u) >> 1);
#pragma unroll
for(int i = 0; i < 2; ++i){
unsigned int s = q >> (16 * i);
t[i] = __byte_perm(__byte_perm(t0, t1, s), __byte_perm(t2, t3, s), sel >> (16 * i));
}
lo = (int)__byte_perm(t[0], t[1], 0x6420);
hi = (int)__byte_perm(t[0], t[1], 0x7531);
}
// IQ4's levels and MXFP4's doubled E2M1 values (ggml's `kvalues_mxfp4`), four to a word.
__device__ __forceinline__ void lk_iq4_word(unsigned int q, int& lo, int& hi){
lk_table16(q, 0xBFAD9881u, 0xF6EADDCFu, 0x26190D01u, 0x71594535u, lo, hi);
}
__device__ __forceinline__ void lk_mxfp4_word(unsigned int q, int& lo, int& hi){
lk_table16(q, 0x03020100u, 0x0C080604u, 0xFDFEFF00u, 0xF4F8FAFCu, lo, hi);
}
// The `n` bytes from `p`, any alignment, as words: aligned word loads joined by funnel shifts. Reads up to three bytes
// past them, which every block-format upload's 32 trailing bytes cover.
template<int N>
__device__ __forceinline__ void lk_words(const unsigned char* p, unsigned int* q){
const unsigned int* a = (const unsigned int*)((size_t)p & ~(size_t)3);
unsigned int sh = 8 * (unsigned int)((size_t)p & 3), w[N + 1];
#pragma unroll
for(int i = 0; i <= N; ++i) w[i] = a[i];
#pragma unroll
for(int i = 0; i < N; ++i) q[i] = __funnelshift_r(w[i], w[i + 1], sh);
}
// An IQ4_XS sub-block's scale: the super-block's `d` times its 6-bit signed scale (low four bits from `sl`, the byte of
// `scales_l` it shares with its neighbour, high two from `sh`, the super-block's `scales_h`), as the CPU dequantizer
// rounds it.
__device__ __forceinline__ float lk_iq4xs_scale(float d, unsigned int sh, unsigned int sl, int ib){
int ls = (int)((sl >> (4 * (ib & 1))) & 0xf) | (int)(((sh >> (2 * ib)) & 3) << 4);
return d * (float)(ls - 32);
}
template<int F>
__device__ __forceinline__ void lk_sub(const unsigned char* __restrict__ row, int sb, signed char* v, float* d, float* c){
if(F == 2 || F == 3){
const unsigned char* b = row + (long)sb * (F == 2 ? 18 : 20);
float dd = lk_h(b), cc = F == 3 ? lk_h(b + 2) : 0.f;
const unsigned char* qs = b + (F == 2 ? 2 : 4);
#pragma unroll
for(int j = 0; j < 16; ++j){ v[j] = (qs[j] & 0xf) - (F == 2 ? 8 : 0); v[j + 16] = (qs[j] >> 4) - (F == 2 ? 8 : 0); }
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 6 || F == 7){
const unsigned char* b = row + (long)sb * (F == 6 ? 22 : 24);
float dd = lk_h(b), cc = F == 7 ? lk_h(b + 2) : 0.f;
const unsigned char* h = b + (F == 6 ? 2 : 4);
unsigned int qh = h[0] | (h[1] << 8) | (h[2] << 16) | ((unsigned int)h[3] << 24);
const unsigned char* qs = h + 4;
#pragma unroll
for(int j = 0; j < 16; ++j){
v[j] = ((qs[j] & 0xf) | (((qh >> j) << 4) & 0x10)) - (F == 6 ? 16 : 0);
v[j + 16] = ((qs[j] >> 4) | ((qh >> (j + 12)) & 0x10)) - (F == 6 ? 16 : 0);
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 39){
const unsigned char* b = row + (long)sb * 17;
const signed char kv[16] = {0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12};
float dd = lk_e8m0_half(b[0]);
#pragma unroll
for(int j = 0; j < 16; ++j){ v[j] = kv[b[1 + j] & 0xf]; v[j + 16] = kv[b[1 + j] >> 4]; }
d[0] = d[1] = dd; c[0] = c[1] = 0.f;
} else if(F == 20 || F == 23){
const unsigned char* b = F == 20 ? row + (long)sb * 18 : row + (long)(sb >> 3) * 136;
int ib = sb & 7;
float dd = F == 20 ? lk_h(b) : lk_iq4xs_scale(lk_h(b), b[2] | (b[3] << 8), b[4 + (ib >> 1)], ib);
const unsigned char* qs = F == 20 ? b + 2 : b + 8 + 16 * ib;
#pragma unroll
for(int j = 0; j < 16; ++j){ v[j] = lk_iq4_level(qs[j] & 0xf); v[j + 16] = lk_iq4_level(qs[j] >> 4); }
d[0] = d[1] = dd; c[0] = c[1] = 0.f;
} else if(F == 12 || F == 13){
const unsigned char* b = row + (long)(sb >> 3) * (F == 12 ? 144 : 176);
int j = sb & 7, p = j >> 1, hi = j & 1, sc, m;
lk_scale_min_k4(j, b + 4, sc, m);
float dd = lk_h(b) * (float)sc, cc = -lk_h(b + 2) * (float)m;
const unsigned char* ql = b + (F == 12 ? 16 : 48) + p * 32;
#pragma unroll
for(int l = 0; l < 32; ++l){
int lv = hi ? ql[l] >> 4 : ql[l] & 0xf;
if(F == 13) lv += (b[16 + l] & (1 << (2 * p + hi))) ? 16 : 0;
v[l] = (signed char)lv;
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 10){
const unsigned char* b = row + (long)(sb >> 3) * 84;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* q = b + 16 + part * 32;
float dd = lk_h(b + 80), dm = lk_h(b + 82);
#pragma unroll
for(int l = 0; l < 32; ++l) v[l] = (signed char)((q[(l & 15) + (l >> 4) * 16] >> (2 * k)) & 3);
#pragma unroll
for(int h = 0; h < 2; ++h){ int sc = b[part * 8 + k * 2 + h]; d[h] = dd * (float)(sc & 0xf); c[h] = -(dm * (float)(sc >> 4)); }
} else if(F == 11){
const unsigned char* b = row + (long)(sb >> 3) * 110;
int j = sb & 7, part = j >> 2, k = j & 3, mbit = 1 << (part * 4 + k);
const unsigned char* q = b + 32 + part * 32;
const unsigned char* p = b + 96;
unsigned int a0 = p[0] | (p[1] << 8) | (p[2] << 16) | ((unsigned int)p[3] << 24);
unsigned int a1 = p[4] | (p[5] << 8) | (p[6] << 16) | ((unsigned int)p[7] << 24);
unsigned int t = p[8] | (p[9] << 8) | (p[10] << 16) | ((unsigned int)p[11] << 24);
unsigned int aux[4] = {(a0 & 0x0f0f0f0fu) | ((t & 0x03030303u) << 4), (a1 & 0x0f0f0f0fu) | (((t >> 2) & 0x03030303u) << 4),
((a0 >> 4) & 0x0f0f0f0fu) | (((t >> 4) & 0x03030303u) << 4), ((a1 >> 4) & 0x0f0f0f0fu) | (((t >> 6) & 0x03030303u) << 4)};
float dd = lk_h(b + 108);
#pragma unroll
for(int l = 0; l < 32; ++l){
int at = (l & 15) + (l >> 4) * 16;
v[l] = (signed char)(((q[at] >> (2 * k)) & 3) - ((b[at] & mbit) ? 0 : 4));
}
#pragma unroll
for(int h = 0; h < 2; ++h){
int is = part * 8 + k * 2 + h;
int sc = (signed char)((lk_pick4(aux, is >> 2) >> (8 * (is & 3))) & 0xff);
d[h] = dd * (float)(sc - 32); c[h] = 0.f;
}
} else if(F == 14){
const unsigned char* b = row + (long)(sb >> 3) * 210;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* ql = b + part * 64;
const unsigned char* qh = b + 128 + part * 32;
const signed char* sc = (const signed char*)(b + 192 + part * 8);
float dd = lk_h(b + 208);
#pragma unroll
for(int l = 0; l < 32; ++l){
int lo = (k & 1) ? ql[l + 32] : ql[l];
lo = (k & 2) ? lo >> 4 : lo & 0xf;
v[l] = (signed char)((lo | (((qh[l] >> (2 * k)) & 3) << 4)) - 32);
}
d[0] = dd * (float)sc[2 * k]; d[1] = dd * (float)sc[2 * k + 1]; c[0] = c[1] = 0.f;
}
}
// Element `l` of sub-block `sb` of a row in format `F`, as `lk_sub` decodes it: its int8 level, and the scale `d` and
// offset `c` of its half — the decoder of one weight a lane reads alone.
template<int F>
__device__ __forceinline__ int lk_level(const unsigned char* __restrict__ row, int sb, int l, float& d, float& c){
if(F == 2 || F == 3){
const unsigned char* b = row + (long)sb * (F == 2 ? 18 : 20);
d = lk_h(b); c = F == 3 ? lk_h(b + 2) : 0.f;
const unsigned char* qs = b + (F == 2 ? 2 : 4);
return (l < 16 ? (qs[l] & 0xf) : (qs[l - 16] >> 4)) - (F == 2 ? 8 : 0);
} else if(F == 6 || F == 7){
const unsigned char* b = row + (long)sb * (F == 6 ? 22 : 24);
d = lk_h(b); c = F == 7 ? lk_h(b + 2) : 0.f;
const unsigned char* h = b + (F == 6 ? 2 : 4);
unsigned int qh = h[0] | (h[1] << 8) | (h[2] << 16) | ((unsigned int)h[3] << 24);
const unsigned char* qs = h + 4;
int q = l < 16 ? ((qs[l] & 0xf) | (((qh >> l) << 4) & 0x10)) : ((qs[l - 16] >> 4) | ((qh >> (l - 4)) & 0x10));
return q - (F == 6 ? 16 : 0);
} else if(F == 39){
const unsigned char* b = row + (long)sb * 17;
const signed char kv[16] = {0, 1, 2, 3, 4, 6, 8, 12, 0, -1, -2, -3, -4, -6, -8, -12};
d = lk_e8m0_half(b[0]); c = 0.f;
return kv[l < 16 ? (b[1 + l] & 0xf) : (b[1 + l - 16] >> 4)];
} else if(F == 20 || F == 23){
const unsigned char* b = F == 20 ? row + (long)sb * 18 : row + (long)(sb >> 3) * 136;
int ib = sb & 7;
d = F == 20 ? lk_h(b) : lk_iq4xs_scale(lk_h(b), b[2] | (b[3] << 8), b[4 + (ib >> 1)], ib); c = 0.f;
const unsigned char* qs = F == 20 ? b + 2 : b + 8 + 16 * ib;
return lk_iq4_level(l < 16 ? (qs[l] & 0xf) : (qs[l - 16] >> 4));
} else if(F == 12 || F == 13){
const unsigned char* b = row + (long)(sb >> 3) * (F == 12 ? 144 : 176);
int j = sb & 7, p = j >> 1, hi = j & 1, sc, m;
lk_scale_min_k4(j, b + 4, sc, m);
d = lk_h(b) * (float)sc; c = -lk_h(b + 2) * (float)m;
const unsigned char* ql = b + (F == 12 ? 16 : 48) + p * 32;
int lv = hi ? ql[l] >> 4 : ql[l] & 0xf;
if(F == 13) lv += (b[16 + l] & (1 << (2 * p + hi))) ? 16 : 0;
return lv;
} else if(F == 10){
const unsigned char* b = row + (long)(sb >> 3) * 84;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* q = b + 16 + part * 32;
float dd = lk_h(b + 80), dm = lk_h(b + 82);
int sc = b[part * 8 + k * 2 + (l >> 4)];
d = dd * (float)(sc & 0xf); c = -(dm * (float)(sc >> 4));
return (q[(l & 15) + (l >> 4) * 16] >> (2 * k)) & 3;
} else if(F == 11){
const unsigned char* b = row + (long)(sb >> 3) * 110;
int j = sb & 7, part = j >> 2, k = j & 3, mbit = 1 << (part * 4 + k);
const unsigned char* q = b + 32 + part * 32;
const unsigned char* p = b + 96;
unsigned int a0 = p[0] | (p[1] << 8) | (p[2] << 16) | ((unsigned int)p[3] << 24);
unsigned int a1 = p[4] | (p[5] << 8) | (p[6] << 16) | ((unsigned int)p[7] << 24);
unsigned int t = p[8] | (p[9] << 8) | (p[10] << 16) | ((unsigned int)p[11] << 24);
unsigned int aux[4] = {(a0 & 0x0f0f0f0fu) | ((t & 0x03030303u) << 4), (a1 & 0x0f0f0f0fu) | (((t >> 2) & 0x03030303u) << 4),
((a0 >> 4) & 0x0f0f0f0fu) | (((t >> 4) & 0x03030303u) << 4), ((a1 >> 4) & 0x0f0f0f0fu) | (((t >> 6) & 0x03030303u) << 4)};
int is = part * 8 + k * 2 + (l >> 4);
int sc = (signed char)((lk_pick4(aux, is >> 2) >> (8 * (is & 3))) & 0xff);
d = lk_h(b + 108) * (float)(sc - 32); c = 0.f;
int at = (l & 15) + (l >> 4) * 16;
return ((q[at] >> (2 * k)) & 3) - ((b[at] & mbit) ? 0 : 4);
} else {
const unsigned char* b = row + (long)(sb >> 3) * 210;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* ql = b + part * 64;
const unsigned char* qh = b + 128 + part * 32;
const signed char* sc = (const signed char*)(b + 192 + part * 8);
d = lk_h(b + 208) * (float)sc[2 * k + (l >> 4)]; c = 0.f;
int lo = (k & 1) ? ql[l + 32] : ql[l];
lo = (k & 2) ? lo >> 4 : lo & 0xf;
return (lo | (((qh[l] >> (2 * k)) & 3) << 4)) - 32;
}
}
__device__ __forceinline__ int lk_pack4(const signed char* v){
return (v[0] & 0xff) | ((v[1] & 0xff) << 8) | ((v[2] & 0xff) << 16) | ((unsigned int)(v[3] & 0xff) << 24);
}
// A word of four bytes from an address 2-byte aligned, or 4-byte aligned.
__device__ __forceinline__ int lk_w16(const unsigned char* p){
const unsigned short* h = (const unsigned short*)p;
return (int)((unsigned int)h[0] | ((unsigned int)h[1] << 16));
}
__device__ __forceinline__ int lk_w32(const unsigned char* p){ return *(const int*)p; }
// The fifth bits of four levels, from four consecutive bits of `qh`, placed at bit 4 of each byte.
__device__ __forceinline__ int lk_bit4(unsigned int b){
return (int)(((b << 4) & 0x10u) | ((b << 11) & 0x1000u) | ((b << 18) & 0x100000u) | ((b << 25) & 0x10000000u));
}
// Sub-block `sb` of a row in format `F`, as `lk_sub` decodes it, its 32 levels packed four to an int for `__dp4a` (`w[i]`
// the levels 4i … 4i+3) — read a word at a time, never a level at a time; the scale and offset of each half as `lk_sub`.
template<int F>
__device__ __forceinline__ void lk_sub4(const unsigned char* __restrict__ row, int sb, int* w, float* d, float* c){
if(F == 2 || F == 3){
const unsigned char* b = row + (long)sb * (F == 2 ? 18 : 20);
float dd = lk_h(b), cc = F == 3 ? lk_h(b + 2) : 0.f;
const unsigned char* qs = b + (F == 2 ? 2 : 4);
#pragma unroll
for(int i = 0; i < 4; ++i){
int q = F == 2 ? lk_w16(qs + 4 * i) : lk_w32(qs + 4 * i);
w[i] = q & 0x0F0F0F0F; w[i + 4] = (q >> 4) & 0x0F0F0F0F;
if(F == 2){ w[i] = __vsub4(w[i], 0x08080808); w[i + 4] = __vsub4(w[i + 4], 0x08080808); }
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 6 || F == 7){
const unsigned char* b = row + (long)sb * (F == 6 ? 22 : 24);
float dd = lk_h(b), cc = F == 7 ? lk_h(b + 2) : 0.f;
const unsigned char* h = b + (F == 6 ? 2 : 4);
unsigned int qh = F == 6 ? (unsigned int)lk_w16(h) : (unsigned int)lk_w32(h);
const unsigned char* qs = h + 4;
#pragma unroll
for(int i = 0; i < 4; ++i){
int q = F == 6 ? lk_w16(qs + 4 * i) : lk_w32(qs + 4 * i);
w[i] = (q & 0x0F0F0F0F) | lk_bit4((qh >> (4 * i)) & 0xF);
w[i + 4] = ((q >> 4) & 0x0F0F0F0F) | lk_bit4((qh >> (16 + 4 * i)) & 0xF);
if(F == 6){ w[i] = __vsub4(w[i], 0x10101010); w[i + 4] = __vsub4(w[i + 4], 0x10101010); }
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 12 || F == 13){
const unsigned char* b = row + (long)(sb >> 3) * (F == 12 ? 144 : 176);
int j = sb & 7, p = j >> 1, hi = j & 1, sc, m;
lk_scale_min_k4(j, b + 4, sc, m);
float dd = lk_h(b) * (float)sc, cc = -lk_h(b + 2) * (float)m;
const unsigned char* ql = b + (F == 12 ? 16 : 48) + p * 32;
#pragma unroll
for(int i = 0; i < 8; ++i){
int q = (lk_w32(ql + 4 * i) >> (4 * hi)) & 0x0F0F0F0F;
if(F == 13) q |= ((lk_w32(b + 16 + 4 * i) >> (2 * p + hi)) & 0x01010101) << 4;
w[i] = q;
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if(F == 10){
const unsigned char* b = row + (long)(sb >> 3) * 84;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* q = b + 16 + part * 32;
float dd = lk_h(b + 80), dm = lk_h(b + 82);
#pragma unroll
for(int i = 0; i < 8; ++i) w[i] = (lk_w32(q + 4 * i) >> (2 * k)) & 0x03030303;
#pragma unroll
for(int h = 0; h < 2; ++h){ int sc = b[part * 8 + k * 2 + h]; d[h] = dd * (float)(sc & 0xf); c[h] = -(dm * (float)(sc >> 4)); }
} else if(F == 11){
const unsigned char* b = row + (long)(sb >> 3) * 110;
int j = sb & 7, part = j >> 2, k = j & 3, bit = part * 4 + k;
const unsigned char* q = b + 32 + part * 32;
const unsigned char* p = b + 96;
unsigned int a0 = (unsigned int)lk_w16(p), a1 = (unsigned int)lk_w16(p + 4), t = (unsigned int)lk_w16(p + 8);
unsigned int aux[4] = {(a0 & 0x0f0f0f0fu) | ((t & 0x03030303u) << 4), (a1 & 0x0f0f0f0fu) | (((t >> 2) & 0x03030303u) << 4),
((a0 >> 4) & 0x0f0f0f0fu) | (((t >> 4) & 0x03030303u) << 4), ((a1 >> 4) & 0x0f0f0f0fu) | (((t >> 6) & 0x03030303u) << 4)};
float dd = lk_h(b + 108);
#pragma unroll
for(int i = 0; i < 8; ++i){
int lo = (lk_w16(q + 4 * i) >> (2 * k)) & 0x03030303;
int missing = (~(lk_w16(b + 4 * i) >> bit) & 0x01010101) << 2;
w[i] = __vsub4(lo, missing);
}
#pragma unroll
for(int h = 0; h < 2; ++h){
int is = part * 8 + k * 2 + h;
int sc = (signed char)((lk_pick4(aux, is >> 2) >> (8 * (is & 3))) & 0xff);
d[h] = dd * (float)(sc - 32); c[h] = 0.f;
}
} else if(F == 14){
const unsigned char* b = row + (long)(sb >> 3) * 210;
int j = sb & 7, part = j >> 2, k = j & 3;
const unsigned char* ql = b + part * 64 + ((k & 1) ? 32 : 0);
const unsigned char* qh = b + 128 + part * 32;
const signed char* sc = (const signed char*)(b + 192 + part * 8);
float dd = lk_h(b + 208);
#pragma unroll
for(int i = 0; i < 8; ++i){
int lo = (lk_w16(ql + 4 * i) >> (4 * (k >> 1))) & 0x0F0F0F0F;
int hi = ((lk_w16(qh + 4 * i) >> (2 * k)) & 0x03030303) << 4;
w[i] = __vsub4(lo | hi, 0x20202020);
}
d[0] = dd * (float)sc[2 * k]; d[1] = dd * (float)sc[2 * k + 1]; c[0] = c[1] = 0.f;
} else if(F == 20 || F == 23){
const unsigned char* b = F == 20 ? row + (long)sb * 18 : row + (long)(sb >> 3) * 136;
int ib = sb & 7;
float dd = F == 20 ? lk_h(b) : lk_iq4xs_scale(lk_h(b), b[2] | (b[3] << 8), b[4 + (ib >> 1)], ib);
const unsigned char* qs = F == 20 ? b + 2 : b + 8 + 16 * ib;
#pragma unroll
for(int i = 0; i < 4; ++i) lk_iq4_word((unsigned int)(F == 20 ? lk_w16(qs + 4 * i) : lk_w32(qs + 4 * i)), w[i], w[i + 4]);
d[0] = d[1] = dd; c[0] = c[1] = 0.f;
} else if(F == 39){
const unsigned char* b = row + (long)sb * 17;
unsigned int q[4];
lk_words<4>(b + 1, q);
#pragma unroll
for(int i = 0; i < 4; ++i) lk_mxfp4_word(q[i], w[i], w[i + 4]);
d[0] = d[1] = lk_e8m0_half(b[0]); c[0] = c[1] = 0.f;
} else {
signed char v[32];
lk_sub<F>(row, sb, v, d, c);
#pragma unroll
for(int i = 0; i < 8; ++i) w[i] = lk_pack4(v + 4 * i);
}
}
// A sub-block's 32 weights (its levels `wv`, four to a word, and each half's scale `d` and offset `c`) dotted with an
// activation's 32 int8 words `h` (and its low words `l` under Q 2), before the activation's scale.
template<int Q>
__device__ __forceinline__ float lk_sub_dot(const int* wv, const float* d, const float* c, const int* h, const int* l){
int dh0 = 0, dh1 = 0, sh0 = 0, sh1 = 0;
#pragma unroll
for(int i = 0; i < 4; ++i){ dh0 = __dp4a(wv[i], h[i], dh0); sh0 = __dp4a(0x01010101, h[i], sh0); }
#pragma unroll
for(int i = 4; i < 8; ++i){ dh1 = __dp4a(wv[i], h[i], dh1); sh1 = __dp4a(0x01010101, h[i], sh1); }
float part = d[0] * (float)dh0 + c[0] * (float)sh0 + d[1] * (float)dh1 + c[1] * (float)sh1;
if(Q == 2){
int dl0 = 0, dl1 = 0, sl0 = 0, sl1 = 0;
#pragma unroll
for(int i = 0; i < 4; ++i){ dl0 = __dp4a(wv[i], l[i], dl0); sl0 = __dp4a(0x01010101, l[i], sl0); }
#pragma unroll
for(int i = 4; i < 8; ++i){ dl1 = __dp4a(wv[i], l[i], dl1); sl1 = __dp4a(0x01010101, l[i], sl1); }
part += 0.0078125f * (d[0] * (float)dl0 + c[0] * (float)sl0 + d[1] * (float)dl1 + c[1] * (float)sl1);
}
return part;
}
// Rows `row … row + W − 1` of `w` dotted with the `M` activation rows `ld` apart (f32 `x` under Q 0 and 3, the int8
// words `xq` (and `xl`) with a scale per 32 in `xs` otherwise), row `r`'s dot with activation row `t` in `out[r·M + t]` —
// one warp, every lane's sums reduced. Each row's dot is the same whatever `W`: the rows only share the activation's
// loads. Under Q 0 each weight is dequantized as the CPU dequantizer rounds it and summed in f64; under Q 3 a lane
// takes half a sub-block, its levels a word at a time and its inputs sixteen bytes at a time, and sums each half's
// `d·Σ level·x + c·Σ x` in f32.
template<int F, int M, int W>
__device__ __forceinline__ void lk_row_dot_f32(lk_w w, long row, int K, const float* __restrict__ x, long ld, float* out){
int lane = threadIdx.x & 31, nb = K >> 5;
const unsigned char* r0 = w.b + row * w.rb;
float acc[W][M];
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) acc[r][t] = 0.f;
for(int u = lane; u < 2 * nb; u += 32){
int sb = u >> 1, hf = u & 1;
#pragma unroll
for(int r = 0; r < W; ++r){
int wv[8]; float d[2], c[2];
lk_sub4<F>(r0 + r * w.rb, sb, wv, d, c);
#pragma unroll
for(int t = 0; t < M; ++t){
const float4* xt = (const float4*)(x + (long)t * ld + sb * 32 + 16 * hf);
float sd = 0.f, sx = 0.f;
#pragma unroll
for(int k = 0; k < 4; ++k){
float4 xv = xt[k];
int lv = wv[4 * hf + k];
sd = fmaf((float)(signed char)lv, xv.x, sd); sd = fmaf((float)(signed char)(lv >> 8), xv.y, sd);
sd = fmaf((float)(signed char)(lv >> 16), xv.z, sd); sd = fmaf((float)(signed char)(lv >> 24), xv.w, sd);
sx += (xv.x + xv.y) + (xv.z + xv.w);
}
acc[r][t] = fmaf(d[hf], sd, fmaf(c[hf], sx, acc[r][t]));
}
}
}
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) out[r * M + t] = lk_warp_sum(acc[r][t]);
}
template<int F, int Q, int M, int W = 1>
__device__ __forceinline__ void lk_row_dot(lk_w w, long row, int K, const float* __restrict__ x,
const signed char* __restrict__ xq, const signed char* __restrict__ xl,
const float* __restrict__ xs, long ld, float* out){
int lane = threadIdx.x & 31, nb = K >> 5;
const unsigned char* r0 = w.b + row * w.rb;
if(Q == 3){
lk_row_dot_f32<F, M, W>(w, row, K, x, ld, out);
} else if(Q == 0){
typedef typename lk_acc<Q>::t A;
A acc[W][M];
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) acc[r][t] = 0;
for(int sb = lane; sb < nb; sb += 32){
#pragma unroll
for(int r = 0; r < W; ++r){
signed char v[32]; float d[2], c[2];
lk_sub<F>(r0 + r * w.rb, sb, v, d, c);
#pragma unroll
for(int t = 0; t < M; ++t){
const float* xt = x + (long)t * ld + sb * 32;
#pragma unroll
for(int i = 0; i < 32; ++i) acc[r][t] += (A)__fadd_rn(__fmul_rn(d[i >> 4], (float)v[i]), c[i >> 4]) * (A)xt[i];
}
}
}
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) out[r * M + t] = (float)lk_warp_sum(acc[r][t]);
} else {
float acc[W][M];
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) acc[r][t] = 0.f;
static_assert(W == 1 || M == 1, "several rows a warp meet one activation row");
for(int sb = lane; sb < nb; sb += 32){
if(W == 1){
// The weight row decoded once, dotted with each activation row in turn.
int wv[8]; float d[2], c[2];
lk_sub4<F>(r0, sb, wv, d, c);
#pragma unroll
for(int t = 0; t < M; ++t){
const int* l = Q == 2 ? (const int*)(xl + (long)t * ld) + sb * 8 : nullptr;
acc[0][t] += xs[(long)t * (ld >> 5) + sb] * lk_sub_dot<Q>(wv, d, c, (const int*)(xq + (long)t * ld) + sb * 8, l);
}
} else {
// The activation's words loaded once, each weight row decoded and dotted in turn.
const int* hp = (const int*)xq + sb * 8;
const int* lp = Q == 2 ? (const int*)xl + sb * 8 : nullptr;
int h[8], l[8];
#pragma unroll
for(int i = 0; i < 8; ++i){ h[i] = hp[i]; l[i] = Q == 2 ? lp[i] : 0; }
float s = xs[sb];
#pragma unroll
for(int r = 0; r < W; ++r){
int wv[8]; float d[2], c[2];
lk_sub4<F>(r0 + r * w.rb, sb, wv, d, c);
acc[r][0] += s * lk_sub_dot<Q>(wv, d, c, h, Q == 2 ? l : nullptr);
}
}
}
#pragma unroll
for(int r = 0; r < W; ++r)
#pragma unroll
for(int t = 0; t < M; ++t) out[r * M + t] = lk_warp_sum(acc[r][t]);
}
}
template<int F, int Q, int M>
__device__ __forceinline__ void lk_dot(lk_w w, long row, int K, const float* __restrict__ x,
const signed char* __restrict__ xq, const signed char* __restrict__ xl,
const float* __restrict__ xs, long ld, float* out){
if(F != 8){ lk_row_dot<F, Q, M>(w, row, K, x, xq, xl, xs, ld, out); return; }
int lane = threadIdx.x & 31;
const signed char* wq = w.q + row * K;
const __half* ws = w.s + row * (K >> 5);
if(Q == 3){
// A lane four levels at a time, each block's `d·Σ level·x` in f32.
float acc[M];
#pragma unroll
for(int t = 0; t < M; ++t) acc[t] = 0.f;
const int* w4 = (const int*)wq;
for(int g = lane; g < (K >> 2); g += 32){
int lv = w4[g];
float d = __half2float(ws[g >> 3]);
#pragma unroll
for(int t = 0; t < M; ++t){
float4 xv = *(const float4*)(x + (long)t * ld + 4 * g);
float sd = (float)(signed char)lv * xv.x;
sd = fmaf((float)(signed char)(lv >> 8), xv.y, sd); sd = fmaf((float)(signed char)(lv >> 16), xv.z, sd);
sd = fmaf((float)(signed char)(lv >> 24), xv.w, sd);
acc[t] = fmaf(d, sd, acc[t]);
}
}
#pragma unroll
for(int t = 0; t < M; ++t) out[t] = lk_warp_sum(acc[t]);
} else if(Q == 0){
typedef typename lk_acc<Q>::t A;
A acc[M];
#pragma unroll
for(int t = 0; t < M; ++t) acc[t] = 0;
for(int k = lane; k < K; k += 32){
A wv = (A)__half2float(ws[k >> 5]) * (A)wq[k];
#pragma unroll
for(int t = 0; t < M; ++t) acc[t] += wv * (A)x[(long)t * ld + k];
}
#pragma unroll
for(int t = 0; t < M; ++t) out[t] = (float)lk_warp_sum(acc[t]);
} else {
float acc[M];
#pragma unroll
for(int t = 0; t < M; ++t) acc[t] = 0.f;
const int* w4 = (const int*)wq;
for(int g = lane; g < (K >> 2); g += 32){
int blk = g >> 3, wv = w4[g];
float wsc = __half2float(ws[blk]);
#pragma unroll
for(int t = 0; t < M; ++t){
float s = xs[(long)t * (ld >> 5) + blk];
int d = __dp4a(wv, ((const int*)(xq + (long)t * ld))[g], 0);
if(Q == 1) acc[t] += (float)d * (wsc * s);
else acc[t] += wsc * (s * (float)d + s * 0.0078125f * (float)__dp4a(wv, ((const int*)(xl + (long)t * ld))[g], 0));
}
}
#pragma unroll
for(int t = 0; t < M; ++t) out[t] = lk_warp_sum(acc[t]);
}
}
// Rows `o … o + W − 1` of a matrix in format `F` (8: Q8_0, split) dotted with the `M` activation rows of `K` values as
// the plain GEMV of that format dots them, row `r`'s dots at `out + r·M` — one warp of a block of at most `WARPS`.
template<int F, int Q, int M, int W = 1, int WARPS = 8>
__device__ __forceinline__ void lk_gemv_dot(lk_w w, long o, int K, const float* __restrict__ x,
const signed char* __restrict__ xq, const signed char* __restrict__ xl,
const float* __restrict__ xs, float* out){
if constexpr(F == 8) lk_q8_rows_dot<Q, M, W, WARPS>(w.q, w.s, o, K, x, xq, xl, xs, out);
else lk_row_dot<F, Q, M, W>(w, o, K, x, xq, xl, xs, K, out);
}
// A weight sub-block's raw words, fetched ahead of their decoding so their loads are in flight while the kernel works
// on what it fetched before: the bytes `lk_sub4` reads for it, by the widest loads the format's alignment allows (16
// bytes in Q4_K, Q5_K and Q8_0's levels, 4 in Q4_1, Q5_1, Q2_K and IQ4_XS, 2 in Q4_0, Q5_0, Q3_K, Q6_K and IQ4_NL, 1 in
// MXFP4).
template<int F> struct lk_raw { };
template<> struct lk_raw<2> { unsigned int q[4]; unsigned short d; };
template<> struct lk_raw<3> { unsigned int q[4], dm; };
template<> struct lk_raw<6> { unsigned int q[4], qh; unsigned short d; };
template<> struct lk_raw<7> { unsigned int q[4], qh, dm; };
template<> struct lk_raw<8> { uint4 q0, q1; unsigned short d; };
template<> struct lk_raw<10> { unsigned int q[8], dm; unsigned short sc; };
template<> struct lk_raw<11> { unsigned int q[8], hm[8], a0, a1, t; unsigned short d; };
template<> struct lk_raw<12> { uint4 h, q0, q1; };
template<> struct lk_raw<13> { uint4 h, qh0, qh1, q0, q1; };
template<> struct lk_raw<14> { unsigned int ql[8], qh[8]; unsigned short sc, d; };
template<> struct lk_raw<39> { unsigned int q[4]; unsigned char e; };
template<> struct lk_raw<20> { unsigned int q[4]; unsigned short d; };
template<> struct lk_raw<23> { unsigned int q[4]; unsigned short d, sh; unsigned char sl; };
__device__ __forceinline__ unsigned short lk_u16(const unsigned char* p){ return *(const unsigned short*)p; }
// Sub-block `ks` of matrix row `row` (`K` values, `nb` sub-blocks), its raw words.
template<int F> __device__ __forceinline__ void lk_fetch(lk_raw<F>& raw, lk_w w, long row, int K, int nb, int ks){
const unsigned char* r = w.b + row * w.rb;
int j = ks & 7, part = j >> 2, k = j & 3;
if constexpr(F == 2 || F == 6){
const unsigned char* b = r + (long)ks * (F == 2 ? 18 : 22);
raw.d = lk_u16(b);
if constexpr(F == 6) raw.qh = (unsigned int)lk_w16(b + 2);
#pragma unroll
for(int i = 0; i < 4; ++i) raw.q[i] = (unsigned int)lk_w16(b + (F == 2 ? 2 : 6) + 4 * i);
} else if constexpr(F == 3 || F == 7){
const unsigned char* b = r + (long)ks * (F == 3 ? 20 : 24);
raw.dm = (unsigned int)lk_w32(b);
if constexpr(F == 7) raw.qh = (unsigned int)lk_w32(b + 4);
#pragma unroll
for(int i = 0; i < 4; ++i) raw.q[i] = (unsigned int)lk_w32(b + (F == 3 ? 4 : 8) + 4 * i);
} else if constexpr(F == 8){
const uint4* q = (const uint4*)(w.q + row * K + (long)ks * 32);
raw.q0 = __ldg(q); raw.q1 = __ldg(q + 1);
raw.d = __half_as_ushort(w.s[row * nb + ks]);
} else if constexpr(F == 10){
const unsigned char* b = r + (long)(ks >> 3) * 84;
#pragma unroll
for(int i = 0; i < 8; ++i) raw.q[i] = (unsigned int)lk_w32(b + 16 + part * 32 + 4 * i);
raw.sc = lk_u16(b + part * 8 + k * 2);
raw.dm = (unsigned int)lk_w32(b + 80);
} else if constexpr(F == 11){
const unsigned char* b = r + (long)(ks >> 3) * 110;
#pragma unroll
for(int i = 0; i < 8; ++i){ raw.q[i] = (unsigned int)lk_w16(b + 32 + part * 32 + 4 * i); raw.hm[i] = (unsigned int)lk_w16(b + 4 * i); }
raw.a0 = (unsigned int)lk_w16(b + 96); raw.a1 = (unsigned int)lk_w16(b + 100); raw.t = (unsigned int)lk_w16(b + 104);
raw.d = lk_u16(b + 108);
} else if constexpr(F == 12 || F == 13){
const unsigned char* b = r + (long)(ks >> 3) * (F == 12 ? 144 : 176);
const uint4* q = (const uint4*)(b + (F == 12 ? 16 : 48) + (j >> 1) * 32);
raw.h = __ldg((const uint4*)b); raw.q0 = __ldg(q); raw.q1 = __ldg(q + 1);
if constexpr(F == 13){ raw.qh0 = __ldg((const uint4*)(b + 16)); raw.qh1 = __ldg((const uint4*)(b + 32)); }
} else if constexpr(F == 14){
const unsigned char* b = r + (long)(ks >> 3) * 210;
#pragma unroll
for(int i = 0; i < 8; ++i){
raw.ql[i] = (unsigned int)lk_w16(b + part * 64 + ((k & 1) ? 32 : 0) + 4 * i);
raw.qh[i] = (unsigned int)lk_w16(b + 128 + part * 32 + 4 * i);
}
raw.sc = lk_u16(b + 192 + part * 8 + 2 * k);
raw.d = lk_u16(b + 208);
} else if constexpr(F == 39){
const unsigned char* b = r + (long)ks * 17;
raw.e = b[0];
#pragma unroll
for(int i = 0; i < 4; ++i) raw.q[i] = (unsigned int)b[1 + 4 * i] | ((unsigned int)b[2 + 4 * i] << 8) | ((unsigned int)b[3 + 4 * i] << 16) | ((unsigned int)b[4 + 4 * i] << 24);
} else if constexpr(F == 20){
const unsigned char* b = r + (long)ks * 18;
raw.d = lk_u16(b);
#pragma unroll
for(int i = 0; i < 4; ++i) raw.q[i] = (unsigned int)lk_w16(b + 2 + 4 * i);
} else if constexpr(F == 23){
const unsigned char* b = r + (long)(ks >> 3) * 136;
raw.d = lk_u16(b); raw.sh = lk_u16(b + 2); raw.sl = b[4 + (j >> 1)];
#pragma unroll
for(int i = 0; i < 4; ++i) raw.q[i] = (unsigned int)lk_w32(b + 8 + 16 * j + 4 * i);
} else static_assert(F != F, "a block format with no raw words");
}
__device__ __forceinline__ float lk_hf(unsigned int bits){ return __half2float(__ushort_as_half((unsigned short)(bits & 0xFFFF))); }
// The sub-block from its fetched words — `lk_sub4`'s levels, scales and offsets, operation for operation.
template<int F> __device__ __forceinline__ void lk_unpack(const lk_raw<F>& raw, int ks, int* w, float* d, float* c){
int j = ks & 7, part = j >> 2, k = j & 3;
if constexpr(F == 2 || F == 3){
float dd, cc = 0.f;
if constexpr(F == 2) dd = lk_hf(raw.d);
else { dd = lk_hf(raw.dm); cc = lk_hf(raw.dm >> 16); }
#pragma unroll
for(int i = 0; i < 4; ++i){
int q = (int)raw.q[i];
w[i] = q & 0x0F0F0F0F; w[i + 4] = (q >> 4) & 0x0F0F0F0F;
if(F == 2){ w[i] = __vsub4(w[i], 0x08080808); w[i + 4] = __vsub4(w[i + 4], 0x08080808); }
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if constexpr(F == 6 || F == 7){
float dd, cc = 0.f;
if constexpr(F == 6) dd = lk_hf(raw.d);
else { dd = lk_hf(raw.dm); cc = lk_hf(raw.dm >> 16); }
unsigned int qh = raw.qh;
#pragma unroll
for(int i = 0; i < 4; ++i){
int q = (int)raw.q[i];
w[i] = (q & 0x0F0F0F0F) | lk_bit4((qh >> (4 * i)) & 0xF);
w[i + 4] = ((q >> 4) & 0x0F0F0F0F) | lk_bit4((qh >> (16 + 4 * i)) & 0xF);
if(F == 6){ w[i] = __vsub4(w[i], 0x10101010); w[i + 4] = __vsub4(w[i + 4], 0x10101010); }
}
d[0] = d[1] = dd; c[0] = c[1] = cc;
} else if constexpr(F == 8){
w[0] = raw.q0.x; w[1] = raw.q0.y; w[2] = raw.q0.z; w[3] = raw.q0.w; w[4] = raw.q1.x; w[5] = raw.q1.y; w[6] = raw.q1.z; w[7] = raw.q1.w;
d[0] = d[1] = lk_hf(raw.d); c[0] = c[1] = 0.f;
} else if constexpr(F == 10){
float dd = lk_hf(raw.dm), dm = lk_hf(raw.dm >> 16);
#pragma unroll
for(int i = 0; i < 8; ++i) w[i] = ((int)raw.q[i] >> (2 * k)) & 0x03030303;
#pragma unroll
for(int h = 0; h < 2; ++h){ int sc = (raw.sc >> (8 * h)) & 0xFF; d[h] = dd * (float)(sc & 0xf); c[h] = -(dm * (float)(sc >> 4)); }
} else if constexpr(F == 11){
int bit = part * 4 + k;
unsigned int a0 = raw.a0, a1 = raw.a1, t = raw.t;
unsigned int aux[4] = {(a0 & 0x0f0f0f0fu) | ((t & 0x03030303u) << 4), (a1 & 0x0f0f0f0fu) | (((t >> 2) & 0x03030303u) << 4),
((a0 >> 4) & 0x0f0f0f0fu) | (((t >> 4) & 0x03030303u) << 4), ((a1 >> 4) & 0x0f0f0f0fu) | (((t >> 6) & 0x03030303u) << 4)};
float dd = lk_hf(raw.d);
#pragma unroll
for(int i = 0; i < 8; ++i){
int lo = ((int)raw.q[i] >> (2 * k)) & 0x03030303;
int missing = (~((int)raw.hm[i] >> bit) & 0x01010101) << 2;
w[i] = __vsub4(lo, missing);
}
#pragma unroll
for(int h = 0; h < 2; ++h){
int is = part * 8 + k * 2 + h;
int sc = (signed char)((lk_pick4(aux, is >> 2) >> (8 * (is & 3))) & 0xff);
d[h] = dd * (float)(sc - 32); c[h] = 0.f;
}
} else if constexpr(F == 12 || F == 13){
int hi = j & 1, pp = j >> 1;
// Scale byte `i` of the head's twelve.
auto sb = [&](int i){ unsigned int wd = i < 4 ? raw.h.y : i < 8 ? raw.h.z : raw.h.w; return (int)((wd >> (8 * (i & 3))) & 0xFF); };
int sc, m;
if(j < 4){ sc = sb(j) & 63; m = sb(j + 4) & 63; }
else { sc = (sb(j + 4) & 0xF) | ((sb(j - 4) >> 6) << 4); m = (sb(j + 4) >> 4) | ((sb(j) >> 6) << 4); }
d[0] = d[1] = lk_hf(raw.h.x) * (float)sc; c[0] = c[1] = -lk_hf(raw.h.x >> 16) * (float)m;
unsigned int q[8] = {raw.q0.x, raw.q0.y, raw.q0.z, raw.q0.w, raw.q1.x, raw.q1.y, raw.q1.z, raw.q1.w};
#pragma unroll
for(int i = 0; i < 8; ++i){
int v = (int)((q[i] >> (4 * hi)) & 0x0F0F0F0Fu);
if constexpr(F == 13){
unsigned int qh[8] = {raw.qh0.x, raw.qh0.y, raw.qh0.z, raw.qh0.w, raw.qh1.x, raw.qh1.y, raw.qh1.z, raw.qh1.w};
v |= (int)(((qh[i] >> (2 * pp + hi)) & 0x01010101u) << 4);
}
w[i] = v;
}
} else if constexpr(F == 14){
float dd = lk_hf(raw.d);
#pragma unroll
for(int i = 0; i < 8; ++i){
int lo = ((int)raw.ql[i] >> (4 * (k >> 1))) & 0x0F0F0F0F;
int hi = (((int)raw.qh[i] >> (2 * k)) & 0x03030303) << 4;
w[i] = __vsub4(lo | hi, 0x20202020);
}
d[0] = dd * (float)(signed char)(raw.sc & 0xFF); d[1] = dd * (float)(signed char)(raw.sc >> 8); c[0] = c[1] = 0.f;
} else if constexpr(F == 39){
#pragma unroll
for(int i = 0; i < 4; ++i) lk_mxfp4_word(raw.q[i], w[i], w[i + 4]);
d[0] = d[1] = lk_e8m0_half(raw.e); c[0] = c[1] = 0.f;
} else if constexpr(F == 20 || F == 23){
float dd;
if constexpr(F == 20) dd = lk_hf(raw.d);
else dd = lk_iq4xs_scale(lk_hf(raw.d), raw.sh, raw.sl, j);
#pragma unroll
for(int i = 0; i < 4; ++i) lk_iq4_word(raw.q[i], w[i], w[i + 4]);
d[0] = d[1] = dd; c[0] = c[1] = 0.f;
} else static_assert(F != F, "a block format with no raw words");
}
"#;Expand description
A weight matrix as the kernels read it, and the sub-block decoders of every format besides Q8_0 (template argument
F, the GGUF type id): a Q8_0 matrix is its int8 values q and block scales s; any other is its blocks b, a row
every rb bytes. lk_sub<F> decodes sub-block sb (weights 32·sb …) of a row into 32 int8 levels and, per half,
a scale and an offset; a weight is scale·level + offset rounded as the CPU dequantizer rounds it (a product, then a
sum); lk_level<F> decodes one element of a sub-block alike. lk_row_dot<F, Q, M> is a row’s dot with M activation
rows (ld apart) under numerics Q, one warp; lk_dot<F, Q, M> the same for any format, a Q8_0 matrix’s split
arrays included; lk_gemv_dot<F, Q, M> a row’s dot with M activation rows of its length as the plain GEMV of its
format takes it, the one dot of a format every kernel shares so a fused launch computes what its parts computed.