mirror of
https://github.com/Niko1221/Strata.git
synced 2026-11-07 22:59:00 -10:00
Verify window: the dense projections of 2+ tokens read an interleaved copy of their activations
The quantizer also writes the tokens' q8_1 values interleaved (a position of every token in one 16-byte word, block-major for the K-quants and Q8_0, position-major for IQ4_XS and IQ3_S). A warp then takes 2-4 rows, decodes each weight block once and applies it to every token with the single-token dot's arithmetic, walking the chunks of the exact layout's four warps with a sum per warp (IQ3_S: row_dot's one sum) added in the same order before the xor tree, so every value is bitwise the multi-column kernel's. Matrices of fewer than 2048 rows and single tokens keep the old kernels. Used for the mixer's projections and the head. Per call, 4 tokens (mmvq_il_parity --bench): IQ4_XS 2560x10240 40.7 -> 29.3 us, Q6_K 68.7 -> 47.4, Q8_0 60.9 -> 43.1, IQ3_S 2560x6144 48.8 -> 21.3; 2 tokens 10-35% faster. Fixed text, 2 alternating rounds: IQ3_XXS code +3.8%, thinking +2.1%; UD-Q4_K_XL within noise. Window hashes identical to the previous build (both models, thinking and code; --spec-split, --pcie-frac 0.3, --second-gpu -1). Claude-Session: https://claude.ai/code/session_01YDx8gSYSbts2783qXv2Nxz
This commit is contained in:
Eddoursul committed
1 parent
70e73db4bc
commit
738e16602d
10 files changed
+766
-14
No files matched your search
@@ -372,6 +372,10 @@ if(STRATA_ENABLE_CUDA)
|
||||
add_executable(qsa_select_parity src/kernels/qsa_select_parity.cu)
|
||||
target_link_libraries(qsa_select_parity PRIVATE strata_kernels CUDA::cudart)
|
||||
add_test(NAME qsa_select_parity COMMAND qsa_select_parity)
|
||||
# The MMVQ of 2..8 tokens from interleaved activations against native_mmvq: bitwise (synthetic matrices).
|
||||
add_executable(mmvq_il_parity src/kernels/mmvq_il_parity.cu)
|
||||
target_link_libraries(mmvq_il_parity PRIVATE strata_kernels CUDA::cudart)
|
||||
add_test(NAME mmvq_il_parity COMMAND mmvq_il_parity)
|
||||
|
||||
# Plan v0.3 P3: native MMVQ for 1..8 tokens (bitwise selftest + --bench timing); synthetic data, no model.
|
||||
if(STRATA_BUILD_TESTS)
|
||||
|
||||
@@ -194,6 +194,7 @@ private:
|
||||
float *ple_ = nullptr, *emb_ = nullptr, *R_ = nullptr, *mixed_ = nullptr, *bo_ = nullptr;
|
||||
float *inj_ = nullptr, *inj2_ = nullptr, *lo_ = nullptr, *rs_ = nullptr, *xn_ = nullptr;
|
||||
uint8_t* xq_ = nullptr; // T columns of q8_1
|
||||
uint8_t* xil_ = nullptr; // their interleaved copy (2+ tokens)
|
||||
float *qkv_L_ = nullptr, *h_L_ = nullptr, *gate_L_ = nullptr, *beta_L_ = nullptr; // per GDN layer
|
||||
float *z_ = nullptr, *y_ = nullptr, *y_dummy_ = nullptr;
|
||||
float *qfull_ = nullptr, *qcur_ = nullptr, *kcur_ = nullptr, *vcur_ = nullptr, *idx_raw_L_ = nullptr;
|
||||
|
||||
@@ -22,6 +22,9 @@ void quantize_q8_1_rows(const float* x, int64_t n_rows, int64_t n_cols, void* y,
|
||||
|
||||
/// y[c][r] = W[r] . x[c] for `ncols` columns of q8_1 activations (x stride n_in/32 blocks per column).
|
||||
void iq_mmvq(int ggml_type, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, void* stream);
|
||||
/// iq_mmvq for 2..8 columns from their interleaved copy (native_quantize_q8_1_il), bitwise: IQ3_S.
|
||||
bool iq_mmvq_il_supported(int ggml_type) noexcept;
|
||||
void iq_mmvq_il(int ggml_type, const void* w, const void* x_il, float* y, int n_in, int n_out, int ncols, void* stream);
|
||||
|
||||
/// Dequantize `n` contiguous values (n a multiple of 256) to fp16 / fp32.
|
||||
void iq_dequant_f16(int ggml_type, const void* src, int64_t n, uint16_t* dst, void* stream);
|
||||
|
||||
@@ -109,4 +109,17 @@ std::size_t native_mmvq_weight_bytes(int ggml_type, int n_in, int n_out);
|
||||
void native_mmvq(int ggml_type, const void* weights, const void* x_q8_1, float* y,
|
||||
int n_in, int n_out, int ncols, void* stream);
|
||||
|
||||
// Plan v0.3 P6: 2..8 columns also INTERLEAVED, so one 16-byte load reads a position of every column: CP = 2 (2
|
||||
// columns), 4 (3-4) or 8 (5-8), zero-padded; nb = n_in / 32 blocks; int32 bm[nb][8][CP] (block-major), int32
|
||||
// pm[8][nb][CP] (position-major), float d[nb][CP] (each block's scale as the dots read it, __low2float of ds).
|
||||
constexpr int native_q8_1_il_cp(int ncols) { return ncols <= 2 ? 2 : ncols <= 4 ? 4 : 8; }
|
||||
std::size_t native_q8_1_il_bytes(int n_in, int ncols);
|
||||
// native_quantize_q8_1 that also writes the interleaved copy to x_il (native_q8_1_il_bytes); 2 <= ncols <= 8.
|
||||
void native_quantize_q8_1_il(const float* x, void* x_q8_1, void* x_il, int n_in, int ncols, void* stream);
|
||||
// native_mmvq for 2..8 columns quantized by native_quantize_q8_1_il, bitwise native_mmvq's output: kernels that read
|
||||
// the interleaved copy for IQ4_XS, Q4_K, Q5_K, Q6_K, Q8_0 and IQ3_S matrices of at least 2048 rows, native_mmvq
|
||||
// otherwise.
|
||||
void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, const void* x_il, float* y, int n_in,
|
||||
int n_out, int ncols, void* stream);
|
||||
|
||||
} // namespace strata::kernels
|
||||
+29
-14
@@ -216,6 +216,7 @@ bool Verifier::init(const WeightTable& wt, const ModelGeometry& g, SessionState&
|
||||
inj_ = b.take<float>(T * HC); inj2_ = b.take<float>(T * HC);
|
||||
lo_ = b.take<float>(T * (uint64_t) g.hc_lr); rs_ = b.take<float>(T * HC); xn_ = b.take<float>(T * HC * N);
|
||||
xq_ = b.take<uint8_t>(strata::kernels::native_q8_1_bytes(max_in, (int) T));
|
||||
xil_ = b.take<uint8_t>(strata::kernels::native_q8_1_il_bytes(max_in, (int) T));
|
||||
qkv_L_ = b.take<float>(nG * T * C); h_L_ = b.take<float>(nG * T * C);
|
||||
gate_L_ = b.take<float>(nG * T * HV); beta_L_ = b.take<float>(nG * T * HV);
|
||||
z_ = b.take<float>(T * ZV); y_ = b.take<float>(T * ZV); y_dummy_ = b.take<float>(T * ZV);
|
||||
@@ -474,6 +475,16 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
gr_read_group(0, pending, inj2_, inj_);
|
||||
stamp(l, 1, cs);
|
||||
float* xm = mixed_ + tb * N;
|
||||
// the mixer's projections: a group of 2+ tokens is quantized with the interleaved copy the multi-token kernels
|
||||
// read
|
||||
auto quant = [&](const float* x, int64_t n_in) {
|
||||
if (n >= 2) native_quantize_q8_1_il(x, xq_, xil_, (int) n_in, n, cs);
|
||||
else native_quantize_q8_1(x, xq_, (int) n_in, n, cs);
|
||||
};
|
||||
auto proj = [&](const WeightRef* w, float* y, int64_t n_in, int64_t n_out) {
|
||||
if (n >= 2) native_mmvq_il(w->native_type, w->native_data, xq_, xil_, y, (int) n_in, (int) n_out, n, cs);
|
||||
else native_mmvq(w->native_type, w->native_data, xq_, y, (int) n_in, (int) n_out, n, cs);
|
||||
};
|
||||
try {
|
||||
if (!is_qsa_layer(g, l)) {
|
||||
// ======================= GDN =======================
|
||||
@@ -493,18 +504,18 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
float* hb = h_L_ + (size_t) gi * MT * C;
|
||||
float* gate = gate_L_ + (size_t) gi * MT * HV;
|
||||
float* beta = beta_L_ + (size_t) gi * MT * HV;
|
||||
native_quantize_q8_1(xm, xq_, (int) N, n, cs);
|
||||
native_mmvq(wqkv->native_type, wqkv->native_data, xq_, qkv + (size_t) tb * C, (int) N, (int) C, n, cs);
|
||||
quant(xm, N);
|
||||
proj(wqkv, qkv + (size_t) tb * C, N, C);
|
||||
gdn_conv_l2_multi(conv, qkv, (const float*) wc->data, hb, (int) C, (int) (2 * HK), EPS, n, cs, tb);
|
||||
gdn_ab_multi(xm, (const uint16_t*) wa->data, (const uint16_t*) wb->data, (const float*) wdt->data,
|
||||
(const float*) wsa->data, gate + (size_t) tb * HV, beta + (size_t) tb * HV, (int) N, (int) HV,
|
||||
n, cs);
|
||||
native_mmvq(wg->native_type, wg->native_data, xq_, z_ + (size_t) tb * ZV, (int) N, (int) ZV, n, cs);
|
||||
proj(wg, z_ + (size_t) tb * ZV, N, ZV);
|
||||
// the recurrence from the untouched state over tokens [0, te); outputs only for this group's
|
||||
gdn_step_norm_multi(state, hb, (int) C, gate, beta, z_, (const float*) wnm->data, EPS, y_, (int) HK,
|
||||
(int) HV, te, nullptr, cs, tb);
|
||||
native_quantize_q8_1(y_ + (size_t) tb * ZV, xq_, (int) ZV, n, cs);
|
||||
native_mmvq(wout->native_type, wout->native_data, xq_, bo_ + tb * N, (int) ZV, (int) N, n, cs);
|
||||
quant(y_ + (size_t) tb * ZV, ZV);
|
||||
proj(wout, bo_ + tb * N, ZV, N);
|
||||
} else {
|
||||
// ======================= QSA =======================
|
||||
const int64_t qi = qsa_idx[(size_t) l];
|
||||
@@ -532,10 +543,10 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
(int) NH, cs);
|
||||
};
|
||||
float* idx_raw = idx_raw_L_ + (size_t) qi * MT * ID;
|
||||
native_quantize_q8_1(xm, xq_, (int) N, n, cs);
|
||||
quant(xm, N);
|
||||
bf16_gemv_fp32_mmvf_multi(xm, (const uint16_t*) wik->data, idx_raw + tb * ID, N, ID, n, cs);
|
||||
native_mmvq(wk->native_type, wk->native_data, xq_, kcur_ + tb * NKV * HD, (int) N, (int) (NKV * HD), n, cs);
|
||||
native_mmvq(wv->native_type, wv->native_data, xq_, vcur_ + tb * NKV * HD, (int) N, (int) (NKV * HD), n, cs);
|
||||
proj(wk, kcur_ + tb * NKV * HD, N, NKV * HD);
|
||||
proj(wv, vcur_ + tb * NKV * HD, N, NKV * HD);
|
||||
if (batched) norm_rope_tokens(kcur_ + tb * NKV * HD, wkn, (int) NKV, (int) HD, pos_ + tb * NH);
|
||||
else for (int t = tb; t < te; ++t) norm_rope(kcur_ + t * NKV * HD, wkn, (int) NKV, (int) HD, pos_ + t * NH);
|
||||
if (grp == 0) copy_from_mapped(tail_snap_ + (size_t) qi * TS, st.idx_tail, TS, cs);
|
||||
@@ -550,8 +561,7 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
native_qsa_indexer_append_multi(idx_raw + tb * ID, step_ + tb * kStepCount + kStepPos, (int) kStepCount, n,
|
||||
0, (const float*) wikn->data, EPS, ib, s, st.max_cells,
|
||||
(float) qsa_freq_base(), cs);
|
||||
native_mmvq(wq->native_type, wq->native_data, xq_, qfull_ + tb * NH * 2 * HD, (int) N, (int) (NH * 2 * HD),
|
||||
n, cs);
|
||||
proj(wq, qfull_ + tb * NH * 2 * HD, N, NH * 2 * HD);
|
||||
// the query halves of every token's heads (rows of 2 * HD: query, then gate)
|
||||
if (cudaMemcpy2DAsync(qcur_ + tb * NH * HD, (size_t) HD * 4, qfull_ + tb * NH * 2 * HD, (size_t) HD * 2 * 4,
|
||||
(size_t) HD * 4, (size_t) (n * NH), cudaMemcpyDeviceToDevice, cs) != cudaSuccess) {
|
||||
@@ -579,8 +589,8 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
else
|
||||
for (int t = tb; t < te; ++t)
|
||||
qsa_gate_apply_f32(attn_ + t * NH * HD, qfull_ + t * NH * 2 * HD, s, attn32_ + t * NH * HD, cs);
|
||||
native_quantize_q8_1(attn32_ + tb * NH * HD, xq_, (int) (NH * HD), n, cs);
|
||||
native_mmvq(wo->native_type, wo->native_data, xq_, bo_ + tb * N, (int) (NH * HD), (int) N, n, cs);
|
||||
quant(attn32_ + tb * NH * HD, NH * HD);
|
||||
proj(wo, bo_ + tb * N, NH * HD, N);
|
||||
}
|
||||
} catch (const std::exception& e) {
|
||||
err = "verify layer " + std::to_string(l) + ": " + e.what();
|
||||
@@ -772,8 +782,13 @@ bool Verifier::record_window(int T, cudaStream_t cs, std::string& err) {
|
||||
}
|
||||
if (head_ != nullptr && head_->loaded()) {
|
||||
try {
|
||||
native_quantize_q8_1(head_mixed_, xq_, (int) N, T, cs);
|
||||
native_mmvq(head_->type(), head_->weights(), xq_, head_logits_, (int) N, (int) n_vocab_, T, cs);
|
||||
if (T >= 2) {
|
||||
native_quantize_q8_1_il(head_mixed_, xq_, xil_, (int) N, T, cs);
|
||||
native_mmvq_il(head_->type(), head_->weights(), xq_, xil_, head_logits_, (int) N, (int) n_vocab_, T, cs);
|
||||
} else {
|
||||
native_quantize_q8_1(head_mixed_, xq_, (int) N, T, cs);
|
||||
native_mmvq(head_->type(), head_->weights(), xq_, head_logits_, (int) N, (int) n_vocab_, T, cs);
|
||||
}
|
||||
} catch (const std::exception& e) {
|
||||
err = std::string("verify head: ") + e.what();
|
||||
return false;
|
||||
|
||||
@@ -216,6 +216,60 @@ __device__ __forceinline__ float vec_dot_iq3_s_q8_1(const void* __restrict__ vbq
|
||||
return d * sumi;
|
||||
}
|
||||
|
||||
// vec_dot_iq3_s_q8_1 for several activation columns: the grid values, masks and scale decoded once (`load`), each
|
||||
// column's integer sums and float expression as above (`apply`; X::u(b, p, u) gives int p of q8_1 block b of every
|
||||
// column, X::scales(b, d) their __low2float scales).
|
||||
struct IQ3SCols {
|
||||
uint32_t g[8], gm[8]; // grid words, and their negative values' bytes
|
||||
int sc;
|
||||
float dw;
|
||||
__device__ __forceinline__ void load(const void* __restrict__ vbq, const int& kbx, const int& iqs) {
|
||||
const block_iq3_s* bq3 = (const block_iq3_s*) vbq + kbx;
|
||||
const int2 qs_packed = make_int2(get_int_b2(bq3->qs, iqs + 0), get_int_b2(bq3->qs, iqs + 1));
|
||||
const uint8_t* qs = (const uint8_t*) &qs_packed;
|
||||
const int qh = bq3->qh[iqs / 2];
|
||||
const int signs_packed_32 = get_int_b2(bq3->signs, iqs / 2);
|
||||
const uint8_t* signs_packed_8 = (const uint8_t*) &signs_packed_32;
|
||||
#pragma unroll
|
||||
for (int l0 = 0; l0 < 8; l0 += 2) {
|
||||
const uint32_t gx = iq3s_grid[qs[l0 + 0] | ((qh << (8 - l0)) & 0x100)];
|
||||
const uint32_t gy = iq3s_grid[qs[l0 + 1] | ((qh << (7 - l0)) & 0x100)];
|
||||
const uint2 m = sign_masks(signs_packed_8[l0 / 2]);
|
||||
g[l0] = gx;
|
||||
gm[l0] = gx & m.x;
|
||||
g[l0 + 1] = gy;
|
||||
gm[l0 + 1] = gy & m.y;
|
||||
}
|
||||
sc = 1 + 2 * ((bq3->scales[iqs / 4] >> ((iqs << 1) & 0x04)) & 0x0F);
|
||||
dw = __half2float(bq3->d);
|
||||
}
|
||||
template <int NC, class X>
|
||||
__device__ __forceinline__ void apply(const X& x, const int& b, float (&out)[NC]) const {
|
||||
int sp[NC], sn[NC];
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) sp[c] = sn[c] = 0;
|
||||
#pragma unroll
|
||||
for (int l = 0; l < 8; ++l) {
|
||||
int u[NC];
|
||||
x.u(b, l, u);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
sp[c] = ggml_cuda_dp4a((int) g[l], u[c], sp[c]);
|
||||
sn[c] = ggml_cuda_dp4a((int) gm[l], u[c], sn[c]);
|
||||
}
|
||||
}
|
||||
float d8[NC];
|
||||
x.scales(b, d8);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
int sumi = sp[c] - 2 * sn[c];
|
||||
sumi *= sc;
|
||||
const float d = dw * d8[c];
|
||||
out[c] = d * sumi;
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
__device__ __forceinline__ float vec_dot_iq1_m_q8_1(const void* __restrict__ vbq, const block_q8_1* __restrict__ bq8_1,
|
||||
const int& kbx, const int& iqs) {
|
||||
const block_iq1_m* bq1 = (const block_iq1_m*) vbq + kbx;
|
||||
|
||||
@@ -5,6 +5,7 @@
|
||||
// in iq_dot.cuh.
|
||||
#include "strata/kernels/iq_kernels.hpp"
|
||||
#include "iq_dot.cuh"
|
||||
#include "q8_1_il.cuh"
|
||||
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
@@ -32,6 +33,45 @@ __global__ void __launch_bounds__(128) mmvq_kernel(const uint8_t* __restrict__ w
|
||||
}
|
||||
}
|
||||
|
||||
// mmvq_kernel for NC columns from their interleaved (position-major) copy: a warp takes R rows, decodes each call's
|
||||
// weights once and keeps row_dot's lane-strided sum per (row, column), so every value is bitwise mmvq_kernel's.
|
||||
template <int NC, int R>
|
||||
__global__ void __launch_bounds__(128) mmvq_il_iq3_s_kernel(const uint8_t* __restrict__ w, size_t row_bytes,
|
||||
const int* __restrict__ xq, const float* __restrict__ xd,
|
||||
float* __restrict__ y, int n_in, int n_out) {
|
||||
using F = Fmt<21>;
|
||||
const int lane = threadIdx.x & 31;
|
||||
const int row0 = (blockIdx.x * 4 + (threadIdx.x >> 5)) * R;
|
||||
if (row0 >= n_out) return;
|
||||
const Q81Il<NC, true> x{xq, xd, n_in / 32};
|
||||
const int nb = n_in / F::qk;
|
||||
float acc[R][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i)
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[i][c] = 0.0f;
|
||||
for (int k = lane; k < nb * F::ipb; k += 32) {
|
||||
const int kbx = k / F::ipb, iqs = F::step * (k % F::ipb);
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) {
|
||||
const int row = min(row0 + i, n_out - 1); // a partial last group recomputes its last row
|
||||
IQ3SCols wv;
|
||||
wv.load(w + (size_t) row * row_bytes, kbx, iqs);
|
||||
float o[NC];
|
||||
wv.apply<NC>(x, kbx * (F::qk / 32) + iqs / 2, o);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[i][c] += o[c];
|
||||
}
|
||||
}
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i)
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
const float s = warp_sum(acc[i][c]);
|
||||
if (lane == 0 && row0 + i < n_out) y[(size_t) c * n_out + row0 + i] = s;
|
||||
}
|
||||
}
|
||||
|
||||
// ---------------------------------------------------------------- grouped native experts
|
||||
// Two kernels a call. The first computes a group's gate and up rows for one chunk of 32 rows of n_ff - one q8_1 block
|
||||
// of h - SwiGLU and that block for each entry; the second the down rows. A warp takes several rows at once, one
|
||||
@@ -531,6 +571,27 @@ void iq_mmvq(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n
|
||||
check("iq_mmvq");
|
||||
}
|
||||
|
||||
bool iq_mmvq_il_supported(int t) noexcept { return t == 21; }
|
||||
|
||||
void iq_mmvq_il(int t, const void* w, const void* x_il, float* y, int n_in, int n_out, int ncols, void* stream) {
|
||||
if (t != 21 || ncols < 2 || ncols > 8 || n_in % 256 != 0) {
|
||||
std::fprintf(stderr, "iq_mmvq_il: type %d, %d columns are not supported\n", t, ncols);
|
||||
std::exit(1);
|
||||
}
|
||||
constexpr int R = 2;
|
||||
const Q81IlParts parts = q8_1_il_parts(x_il, n_in, ncols);
|
||||
const size_t rb = iq_row_bytes(t, n_in);
|
||||
const auto* W = (const uint8_t*) w;
|
||||
const unsigned blocks = (unsigned) (((n_out + R - 1) / R + 3) / 4);
|
||||
cudaStream_t s = (cudaStream_t) stream;
|
||||
switch (ncols) {
|
||||
#define STRATA_IL(N) case N: mmvq_il_iq3_s_kernel<N, R><<<blocks, 128, 0, s>>>(W, rb, parts.pm, parts.d, y, n_in, n_out); break;
|
||||
STRATA_IL(2) STRATA_IL(3) STRATA_IL(4) STRATA_IL(5) STRATA_IL(6) STRATA_IL(7) STRATA_IL(8)
|
||||
#undef STRATA_IL
|
||||
}
|
||||
check("iq_mmvq_il");
|
||||
}
|
||||
|
||||
void iq_dequant_f16(int t, const void* src, int64_t n, uint16_t* dst, void* stream) {
|
||||
if (n % 256 != 0 || !is_iq(t)) { std::fprintf(stderr, "iq_dequant_f16: bad arguments\n"); std::exit(1); }
|
||||
dequant_flat_kernel<__half><<<(unsigned) (n / 256), 32, 0, (cudaStream_t) stream>>>(t, src, (__half*) dst);
|
||||
|
||||
@@ -25,6 +25,7 @@
|
||||
|
||||
#include "strata/kernels/native_mmvq.hpp"
|
||||
#include "strata/kernels/iq_kernels.hpp"
|
||||
#include "q8_1_il.cuh"
|
||||
|
||||
#include <cuda_fp16.h>
|
||||
#include <cuda_runtime.h>
|
||||
@@ -1076,6 +1077,274 @@ void launch_multi(const void* weights, const void* x_q8_1, float* y, int n_in, i
|
||||
}
|
||||
}
|
||||
|
||||
// ============================ plan v0.3 P6: ncols = 2..8 from interleaved activations ============================
|
||||
//
|
||||
// `native_quantize_q8_1_il` also writes the columns interleaved (native_mmvq.hpp), so one load reads the same int of
|
||||
// every column. A warp takes R rows and walks their blocks in the chunks of the exact layout above - its lane l's chunk
|
||||
// m is what thread 32 (m % 4) + l does in iteration m / 4 - with a sum per virtual warp, added in warp order before the
|
||||
// xor tree; each weight block is decoded once (the traits' `load`) and every column takes the single-column dot's
|
||||
// arithmetic, so each value is bitwise the multi-column kernel's. Formats whose lanes read consecutive ints of one q8_1
|
||||
// block read the block-major copy; IQ4_XS, whose lanes read the same int of consecutive blocks, the position-major one.
|
||||
//
|
||||
// Per format: the traits' blocks and `load`, and `apply` for NC columns (kby = the columns' first q8_1 block of the
|
||||
// weight block): the single-column `apply`'s expression per column.
|
||||
struct IlIQ4XS {
|
||||
using F = IQ4XSTraits;
|
||||
static constexpr bool PM = true;
|
||||
template <int NC>
|
||||
__device__ static void apply(const F::W& r, const Q81Il<NC, PM>& x, int kby, int iqs, float (&out)[NC]) {
|
||||
const int b = kby + iqs / 4;
|
||||
int sumi[NC];
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) sumi[c] = 0;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
int u0[NC], u1[NC];
|
||||
x.u(b, j, u0);
|
||||
x.u(b, j + 4, u1);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
sumi[c] = __dp4a(r.v[j].x, u0[c], sumi[c]);
|
||||
sumi[c] = __dp4a(r.v[j].y, u1[c], sumi[c]);
|
||||
}
|
||||
}
|
||||
float d8[NC];
|
||||
x.scales(b, d8);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
const int s = sumi[c] * (r.ls - 32);
|
||||
const float d = r.dw * d8[c];
|
||||
out[c] = d * s;
|
||||
}
|
||||
}
|
||||
};
|
||||
struct IlQ4K {
|
||||
using F = Q4KTraits;
|
||||
static constexpr bool PM = false;
|
||||
template <int NC>
|
||||
__device__ static void apply(const F::W& r, const Q81Il<NC, PM>& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[4][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int b = kby + r.bq8_offset + i;
|
||||
x.u(b, (iqs / 2) % 4, u[2 * i]);
|
||||
x.u(b, (iqs / 2) % 4 + 4, u[2 * i + 1]);
|
||||
x.scales(b, d8[i]);
|
||||
}
|
||||
const uint8_t* sc = reinterpret_cast<const uint8_t*>(r.aux);
|
||||
const uint8_t* m = sc + 2;
|
||||
const float2 dm4f = __half22float2(r.dm);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
float sumf_d = 0.0f;
|
||||
float sumf_m = 0.0f;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int v0i = (r.v[0] >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int v1i = (r.v[1] >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int dot1 = __dp4a(v1i, u[2 * i + 1][c], __dp4a(v0i, u[2 * i][c], 0));
|
||||
const int dot2 = __dp4a(0x01010101, u[2 * i + 1][c], __dp4a(0x01010101, u[2 * i][c], 0));
|
||||
sumf_d += d8[i][c] * (dot1 * sc[i]);
|
||||
sumf_m += d8[i][c] * (dot2 * m[i]);
|
||||
}
|
||||
out[c] = dm4f.x * sumf_d - dm4f.y * sumf_m;
|
||||
}
|
||||
}
|
||||
};
|
||||
struct IlQ5K {
|
||||
using F = Q5KTraits;
|
||||
static constexpr bool PM = false;
|
||||
template <int NC>
|
||||
__device__ static void apply(const F::W& r, const Q81Il<NC, PM>& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[4][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int b = kby + r.bq8_offset + i;
|
||||
x.u(b, (iqs / 2) % 4, u[2 * i]);
|
||||
x.u(b, (iqs / 2) % 4 + 4, u[2 * i + 1]);
|
||||
x.scales(b, d8[i]);
|
||||
}
|
||||
const uint8_t* sc = reinterpret_cast<const uint8_t*>(r.aux);
|
||||
const uint8_t* m = sc + 2;
|
||||
const float2 dm5f = __half22float2(r.dm);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
float sumf_d = 0.0f;
|
||||
float sumf_m = 0.0f;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int vl0i = (r.vl[0] >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int vl1i = (r.vl[1] >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int vh0i = ((r.vh[0] >> i) << 4) & 0x10101010;
|
||||
const int vh1i = ((r.vh[1] >> i) << 4) & 0x10101010;
|
||||
const int v0i = vl0i | vh0i;
|
||||
const int v1i = vl1i | vh1i;
|
||||
const int dot1 = __dp4a(v0i, u[2 * i][c], __dp4a(v1i, u[2 * i + 1][c], 0));
|
||||
const int dot2 = __dp4a(0x01010101, u[2 * i][c], __dp4a(0x01010101, u[2 * i + 1][c], 0));
|
||||
sumf_d += d8[i][c] * (dot1 * sc[i]);
|
||||
sumf_m += d8[i][c] * (dot2 * m[i]);
|
||||
}
|
||||
out[c] = dm5f.x * sumf_d - dm5f.y * sumf_m;
|
||||
}
|
||||
}
|
||||
};
|
||||
struct IlQ6K {
|
||||
using F = Q6KTraits;
|
||||
static constexpr bool PM = false;
|
||||
template <int NC>
|
||||
__device__ static void apply(const F::W& r, const Q81Il<NC, PM>& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[2][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
x.u(kby + r.bq8_offset + 2 * i, iqs % 8, u[i]);
|
||||
x.scales(kby + r.bq8_offset + 2 * i, d8[i]);
|
||||
}
|
||||
int vi[2], sc[2];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
sc[i] = r.scales[4 * i];
|
||||
const int vil = (r.vl >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int vih = ((r.vh >> (4 * i)) << 4) & 0x30303030;
|
||||
vi[i] = __vsubss4(vil | vih, 0x20202020);
|
||||
}
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
float sumf = 0.0f;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) sumf += d8[i][c] * (__dp4a(vi[i], u[i][c], 0) * sc[i]);
|
||||
out[c] = r.d * sumf;
|
||||
}
|
||||
}
|
||||
};
|
||||
// Q8_0: its traits keep the block pointer (`small_q8_dot` loads it per column); here the codes and scale load once.
|
||||
struct IlQ80 {
|
||||
using F = SmallTraits<Q80Block, 8>;
|
||||
static constexpr bool PM = false;
|
||||
struct W { int v[2]; float d0; };
|
||||
__device__ static W load(const Q80Block* __restrict__ w, int iqs) {
|
||||
W r;
|
||||
r.v[0] = load_int_b2(w->qs, iqs);
|
||||
r.v[1] = load_int_b2(w->qs, iqs + 1);
|
||||
r.d0 = w->d;
|
||||
return r;
|
||||
}
|
||||
template <int NC>
|
||||
__device__ static void apply(const W& r, const Q81Il<NC, PM>& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u0[NC], u1[NC];
|
||||
x.u(kby, iqs, u0);
|
||||
x.u(kby, iqs + 1, u1);
|
||||
float d8[NC];
|
||||
x.scales(kby, d8);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
int sumi = 0;
|
||||
sumi = __dp4a(r.v[0], u0[c], sumi);
|
||||
sumi = __dp4a(r.v[1], u1[c], sumi);
|
||||
const float d0 = r.d0;
|
||||
const float d1 = d8[c];
|
||||
out[c] = d0 * d1 * float(sumi);
|
||||
}
|
||||
}
|
||||
};
|
||||
template <typename I> struct IlLoad {
|
||||
using W = typename I::F::W;
|
||||
__device__ static W load(const typename I::F::Block* w, int kqs) { return I::F::load(w, kqs); }
|
||||
};
|
||||
template <> struct IlLoad<IlQ80> {
|
||||
using W = IlQ80::W;
|
||||
__device__ static W load(const Q80Block* w, int kqs) { return IlQ80::load(w, kqs); }
|
||||
};
|
||||
|
||||
template <typename I, int NC, int R>
|
||||
__launch_bounds__(4 * WARP)
|
||||
__global__ void native_mmvq_il_kernel(const typename I::F::Block* __restrict__ w, const int* __restrict__ xq,
|
||||
const float* __restrict__ xd, float* __restrict__ y, int n_in, int n_out) {
|
||||
using F = typename I::F;
|
||||
constexpr int SUB = WARP / F::T; // weight blocks a chunk covers
|
||||
const int lane = int(threadIdx.x) & (WARP - 1);
|
||||
const int row0 = (int(blockIdx.x) * 4 + (int(threadIdx.x) >> 5)) * R;
|
||||
if (row0 >= n_out) return;
|
||||
const Q81Il<NC, I::PM> x{xq, xd, n_in / Q8K};
|
||||
const int bpr = n_in / F::DIV;
|
||||
const int kqs = F::kqs(lane);
|
||||
const int nm = (bpr + SUB - 1) / SUB;
|
||||
float acc[4][R][NC];
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v)
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i)
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[v][i][c] = 0.0f;
|
||||
for (int m0 = 0; m0 < nm; m0 += 4) {
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v) {
|
||||
const int kbx = SUB * (m0 + v) + lane / F::T;
|
||||
if (kbx < bpr) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) {
|
||||
const int row = min(row0 + i, n_out - 1); // a partial last group recomputes its last row
|
||||
const typename IlLoad<I>::W wv = IlLoad<I>::load(w + std::size_t(row) * bpr + kbx, kqs);
|
||||
float o[NC];
|
||||
I::template apply<NC>(wv, x, kbx * F::KBY, kqs, o);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[v][i][c] += o[c];
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i)
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
float s = acc[0][i][c];
|
||||
s += acc[1][i][c];
|
||||
s += acc[2][i][c];
|
||||
s += acc[3][i][c];
|
||||
s = warp_sum(s);
|
||||
if (lane == 0 && row0 + i < n_out) y[std::size_t(c) * n_out + row0 + i] = s;
|
||||
}
|
||||
}
|
||||
|
||||
template <typename I, int R>
|
||||
void launch_il(const void* weights, const void* x_il, float* y, int n_in, int n_out, int ncols, cudaStream_t s) {
|
||||
const Q81IlParts parts = q8_1_il_parts(x_il, n_in, ncols);
|
||||
const int* xq = I::PM ? parts.pm : parts.bm;
|
||||
const float* xd = parts.d;
|
||||
const auto* w = static_cast<const typename I::F::Block*>(weights);
|
||||
const unsigned blocks = unsigned(((n_out + R - 1) / R + 3) / 4);
|
||||
switch (ncols) {
|
||||
#define STRATA_IL_CASE(N) case N: native_mmvq_il_kernel<I, N, R><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); break;
|
||||
STRATA_IL_CASE(2) STRATA_IL_CASE(3) STRATA_IL_CASE(4) STRATA_IL_CASE(5) STRATA_IL_CASE(6) STRATA_IL_CASE(7)
|
||||
STRATA_IL_CASE(8)
|
||||
#undef STRATA_IL_CASE
|
||||
default: throw std::invalid_argument("native_mmvq_il requires 2 <= ncols <= 8");
|
||||
}
|
||||
}
|
||||
|
||||
// native_quantize_q8_1_kernel plus the interleaved copy (native_mmvq.hpp): the same values in the same bytes
|
||||
__launch_bounds__(QUANT_THREADS, 1)
|
||||
__global__ void native_quantize_q8_1_il_kernel(const float* __restrict__ x, Q81Block* __restrict__ y,
|
||||
int8_t* __restrict__ bm, int8_t* __restrict__ pm, float* __restrict__ dl,
|
||||
int n_in, int n_total, int cp) {
|
||||
const int i = int(blockIdx.x) * QUANT_THREADS + int(threadIdx.x);
|
||||
if (i >= n_total) return; // n_in is a multiple of 32: only whole warps return.
|
||||
const float xi = x[i];
|
||||
const float amax = warp_max(fabsf(xi));
|
||||
const float sum = warp_sum(xi);
|
||||
const float d = amax / 127.0f;
|
||||
const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
|
||||
const half2 ds = make_half2(d, sum);
|
||||
y[i / Q8K].qs[i % Q8K] = q;
|
||||
if (i % Q8K == 0) y[i / Q8K].ds = ds;
|
||||
const int c = i / n_in, e = i - c * n_in, b = e / Q8K, p = (e % Q8K) / 4, nb = n_in / Q8K;
|
||||
bm[(std::size_t(b * 8 + p) * cp + c) * 4 + e % 4] = q;
|
||||
pm[(std::size_t(p * nb + b) * cp + c) * 4 + e % 4] = q;
|
||||
if (e % Q8K == 0) dl[std::size_t(b) * cp + c] = __low2float(ds);
|
||||
}
|
||||
|
||||
void validate_shape(int n_in, int ncols, int block_elems = Q8K) {
|
||||
if (n_in <= 0 || n_in % block_elems != 0) {
|
||||
throw std::invalid_argument("native MMVQ requires n_in > 0 and divisible by its block element count");
|
||||
@@ -1489,4 +1758,62 @@ void native_mmvq(int ggml_type, const void* weights, const void* x_q8_1, float*
|
||||
}
|
||||
}
|
||||
|
||||
std::size_t native_q8_1_il_bytes(int n_in, int ncols) {
|
||||
validate_shape(n_in, ncols);
|
||||
return std::size_t(n_in / Q8K) * std::size_t(native_q8_1_il_cp(ncols)) * (16 * sizeof(int) + sizeof(float));
|
||||
}
|
||||
|
||||
void native_quantize_q8_1_il(const float* x, void* x_q8_1, void* x_il, int n_in, int ncols, void* stream) {
|
||||
validate_shape(n_in, ncols);
|
||||
if (ncols < 2) throw std::invalid_argument("native_quantize_q8_1_il requires 2 <= ncols <= 8");
|
||||
validate_pointer(x);
|
||||
validate_pointer(x_q8_1);
|
||||
validate_pointer(x_il);
|
||||
validate_stream(stream);
|
||||
const int nb = n_in / Q8K, cp = native_q8_1_il_cp(ncols);
|
||||
auto* bm = static_cast<int8_t*>(x_il);
|
||||
auto* pm = bm + std::size_t(nb) * 8 * cp * 4;
|
||||
auto* dl = reinterpret_cast<float*>(pm + std::size_t(nb) * 8 * cp * 4);
|
||||
const int n_total = n_in * ncols;
|
||||
const unsigned blocks = unsigned((std::size_t(n_total) + QUANT_THREADS - 1) / QUANT_THREADS);
|
||||
native_quantize_q8_1_il_kernel<<<blocks, QUANT_THREADS, 0, static_cast<cudaStream_t>(stream)>>>(
|
||||
x, static_cast<Q81Block*>(x_q8_1), bm, pm, dl, n_in, n_total, cp);
|
||||
launch_check();
|
||||
}
|
||||
|
||||
void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, const void* x_il, float* y, int n_in,
|
||||
int n_out, int ncols, void* stream) {
|
||||
// below ~2048 rows the warps are too few for one row each: the multi-column kernel spreads a row over four (and
|
||||
// for Q8_0 2 columns of fewer than 4096 rows)
|
||||
const bool il = ggml_type == 23 || ggml_type == 12 || ggml_type == 13 || ggml_type == 14 || ggml_type == 8;
|
||||
if (ncols < 2 || n_out < 2048 || (ggml_type == 8 && ncols == 2 && n_out < 4096) ||
|
||||
(!il && !iq_mmvq_il_supported(ggml_type))) {
|
||||
native_mmvq(ggml_type, weights, x_q8_1, y, n_in, n_out, ncols, stream);
|
||||
return;
|
||||
}
|
||||
if (!il) {
|
||||
iq_mmvq_il(ggml_type, weights, x_il, y, n_in, n_out, ncols, stream);
|
||||
return;
|
||||
}
|
||||
validate_shape(n_in, ncols, ggml_type == 8 ? Q8K : QK);
|
||||
validate_pointer(weights);
|
||||
validate_pointer(x_il);
|
||||
validate_pointer(y);
|
||||
validate_stream(stream);
|
||||
const auto s = static_cast<cudaStream_t>(stream);
|
||||
// rows a warp: 2, and 4 for Q4_K and Q6_K from 8192 rows (the faster on the model files' shapes, mmvq_il_parity
|
||||
// --bench)
|
||||
const bool wide = n_out >= 8192;
|
||||
switch (ggml_type) {
|
||||
case 23: launch_il<IlIQ4XS, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 12: wide ? launch_il<IlQ4K, 4>(weights, x_il, y, n_in, n_out, ncols, s)
|
||||
: launch_il<IlQ4K, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 13: launch_il<IlQ5K, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 14: wide ? launch_il<IlQ6K, 4>(weights, x_il, y, n_in, n_out, ncols, s)
|
||||
: launch_il<IlQ6K, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 8: launch_il<IlQ80, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
}
|
||||
launch_check();
|
||||
}
|
||||
|
||||
} // namespace strata::kernels
|
||||
@@ -0,0 +1,78 @@
|
||||
// src/kernels/cuda/q8_1_il.cuh - the interleaved copy of 2..8 q8_1 activation columns (native_mmvq.hpp) as the
|
||||
// multi-column dot products read it, for native_mmvq.cu and iq_kernels.cu.
|
||||
#pragma once
|
||||
|
||||
#include "strata/kernels/native_mmvq.hpp"
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <cstddef>
|
||||
|
||||
namespace strata::kernels {
|
||||
namespace {
|
||||
|
||||
// NC columns of the block-major (PM false) or position-major copy: u(b, p, o) = int p of q8_1 block b of every column,
|
||||
// scales(b, o) = their scales
|
||||
template <int NC, bool PM> struct Q81Il {
|
||||
static constexpr int CP = native_q8_1_il_cp(NC);
|
||||
const int* __restrict__ q;
|
||||
const float* __restrict__ d;
|
||||
int nb; // blocks a column
|
||||
__device__ __forceinline__ void u(int b, int p, int (&o)[NC]) const {
|
||||
const int* s = q + std::size_t(PM ? p * nb + b : b * 8 + p) * CP;
|
||||
if constexpr (CP == 2) {
|
||||
const int2 a = *reinterpret_cast<const int2*>(s);
|
||||
o[0] = a.x;
|
||||
o[1] = a.y;
|
||||
} else {
|
||||
const int4 a = *reinterpret_cast<const int4*>(s);
|
||||
o[0] = a.x;
|
||||
if constexpr (NC > 1) o[1] = a.y;
|
||||
if constexpr (NC > 2) o[2] = a.z;
|
||||
if constexpr (NC > 3) o[3] = a.w;
|
||||
if constexpr (NC > 4) {
|
||||
const int4 e = *reinterpret_cast<const int4*>(s + 4);
|
||||
o[4] = e.x;
|
||||
if constexpr (NC > 5) o[5] = e.y;
|
||||
if constexpr (NC > 6) o[6] = e.z;
|
||||
if constexpr (NC > 7) o[7] = e.w;
|
||||
}
|
||||
}
|
||||
}
|
||||
__device__ __forceinline__ void scales(int b, float (&o)[NC]) const {
|
||||
const float* s = d + std::size_t(b) * CP;
|
||||
if constexpr (CP == 2) {
|
||||
const float2 a = *reinterpret_cast<const float2*>(s);
|
||||
o[0] = a.x;
|
||||
o[1] = a.y;
|
||||
} else {
|
||||
const float4 a = *reinterpret_cast<const float4*>(s);
|
||||
o[0] = a.x;
|
||||
if constexpr (NC > 1) o[1] = a.y;
|
||||
if constexpr (NC > 2) o[2] = a.z;
|
||||
if constexpr (NC > 3) o[3] = a.w;
|
||||
if constexpr (NC > 4) {
|
||||
const float4 e = *reinterpret_cast<const float4*>(s + 4);
|
||||
o[4] = e.x;
|
||||
if constexpr (NC > 5) o[5] = e.y;
|
||||
if constexpr (NC > 6) o[6] = e.z;
|
||||
if constexpr (NC > 7) o[7] = e.w;
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// The copy's parts inside the buffer native_q8_1_il_bytes sizes: bm, pm (int32) and d (float).
|
||||
struct Q81IlParts {
|
||||
const int* bm;
|
||||
const int* pm;
|
||||
const float* d;
|
||||
};
|
||||
inline Q81IlParts q8_1_il_parts(const void* x_il, int n_in, int ncols) {
|
||||
const std::size_t words = std::size_t(n_in / 32) * 8 * std::size_t(native_q8_1_il_cp(ncols));
|
||||
const int* bm = static_cast<const int*>(x_il);
|
||||
return {bm, bm + words, reinterpret_cast<const float*>(bm + 2 * words)};
|
||||
}
|
||||
|
||||
} // namespace
|
||||
} // namespace strata::kernels
|
||||
@@ -0,0 +1,196 @@
|
||||
// src/kernels/mmvq_il_parity.cu - native_mmvq_il (2..8 columns from the interleaved q8_1 copy) against native_mmvq,
|
||||
// bitwise, on synthetic matrices of the model files' dense formats and shapes; the interleaving quantizer's q8_1 blocks
|
||||
// against native_quantize_q8_1's. --bench times both inside CUDA graphs (weights cycled past the L2).
|
||||
#include "strata/kernels/native_mmvq.hpp"
|
||||
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <chrono>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
#include <cstdlib>
|
||||
#include <cstring>
|
||||
#include <functional>
|
||||
#include <random>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
using namespace strata::kernels;
|
||||
|
||||
#define CK(x) do { const cudaError_t e_ = (x); if (e_ != cudaSuccess) { std::printf("%s: %s\n", #x, cudaGetErrorString(e_)); std::exit(1); } } while (0)
|
||||
|
||||
namespace {
|
||||
|
||||
struct Fmt { int type; const char* name; int block_elems, block_bytes, d_off, dmin_off; };
|
||||
// the dense formats of the IQ3_XXS and UD-Q4_K_XL files; d_off: the fp16 block scale (random bytes elsewhere)
|
||||
const Fmt kFmts[] = {
|
||||
{23, "IQ4_XS", 256, 136, 0, -1}, {12, "Q4_K", 256, 144, 0, 2}, {13, "Q5_K", 256, 176, 0, 2},
|
||||
{14, "Q6_K", 256, 210, 208, -1}, {8, "Q8_0", 32, 34, 0, -1}, {21, "IQ3_S", 256, 110, 0, -1},
|
||||
};
|
||||
|
||||
cudaStream_t g_s;
|
||||
std::mt19937 g_rng(29);
|
||||
|
||||
void* dev_weights(const Fmt& f, int n_in, int n_out) {
|
||||
const size_t bytes = native_mmvq_weight_bytes(f.type, n_in, n_out);
|
||||
std::vector<uint8_t> h(bytes);
|
||||
for (auto& b : h) b = (uint8_t) g_rng();
|
||||
std::uniform_int_distribution<int> e(0x1c00, 0x2400); // fp16 scales in ~[0.004, 0.016]
|
||||
for (size_t o = 0; o < bytes; o += (size_t) f.block_bytes) {
|
||||
const uint16_t d = (uint16_t) e(g_rng);
|
||||
std::memcpy(&h[o + f.d_off], &d, 2);
|
||||
if (f.dmin_off >= 0) { const uint16_t m = (uint16_t) e(g_rng); std::memcpy(&h[o + f.dmin_off], &m, 2); }
|
||||
}
|
||||
void* p;
|
||||
CK(cudaMalloc(&p, bytes));
|
||||
CK(cudaMemcpy(p, h.data(), bytes, cudaMemcpyHostToDevice));
|
||||
CK(cudaDeviceSynchronize()); // a pageable copy may return before its DMA lands
|
||||
return p;
|
||||
}
|
||||
|
||||
struct Act {
|
||||
float* x = nullptr;
|
||||
void *q_ref = nullptr, *q = nullptr, *il = nullptr;
|
||||
int n_in = 0;
|
||||
};
|
||||
Act make_act(int n_in) {
|
||||
Act a;
|
||||
a.n_in = n_in;
|
||||
std::vector<float> h((size_t) 8 * n_in);
|
||||
std::normal_distribution<float> nd(0.0f, 1.0f);
|
||||
for (auto& v : h) v = nd(g_rng);
|
||||
for (size_t i = 0; i < h.size(); i += 97) h[i] *= 20.0f; // outliers set some blocks' scales
|
||||
for (int j = 0; j < 32; ++j) h[(size_t) 3 * n_in + 64 + j] = 0.0f; // an all-zero block
|
||||
CK(cudaMalloc(&a.x, h.size() * 4));
|
||||
CK(cudaMemcpy(a.x, h.data(), h.size() * 4, cudaMemcpyHostToDevice));
|
||||
CK(cudaMalloc(&a.q_ref, native_q8_1_bytes(n_in, 8)));
|
||||
CK(cudaMalloc(&a.q, native_q8_1_bytes(n_in, 8)));
|
||||
CK(cudaMalloc(&a.il, native_q8_1_il_bytes(n_in, 8)));
|
||||
CK(cudaDeviceSynchronize());
|
||||
return a;
|
||||
}
|
||||
|
||||
template <typename T>
|
||||
std::vector<T> fetch(const void* d, size_t n) {
|
||||
std::vector<T> h(n);
|
||||
CK(cudaMemcpy(h.data(), d, n * sizeof(T), cudaMemcpyDeviceToHost));
|
||||
return h;
|
||||
}
|
||||
|
||||
int check_shape(const Fmt& f, const Act& a, int n_out) {
|
||||
const int n_in = a.n_in;
|
||||
void* w = dev_weights(f, n_in, n_out);
|
||||
float *y0, *y1;
|
||||
CK(cudaMalloc(&y0, (size_t) 8 * n_out * 4));
|
||||
CK(cudaMalloc(&y1, (size_t) 8 * n_out * 4));
|
||||
int bad = 0;
|
||||
for (int nc = 2; nc <= 8; ++nc) {
|
||||
native_quantize_q8_1(a.x, a.q_ref, n_in, nc, g_s);
|
||||
native_quantize_q8_1_il(a.x, a.q, a.il, n_in, nc, g_s);
|
||||
CK(cudaMemsetAsync(y0, 0xff, (size_t) 8 * n_out * 4, g_s));
|
||||
CK(cudaMemsetAsync(y1, 0xee, (size_t) 8 * n_out * 4, g_s));
|
||||
native_mmvq(f.type, w, a.q_ref, y0, n_in, n_out, nc, g_s);
|
||||
native_mmvq_il(f.type, w, a.q, a.il, y1, n_in, n_out, nc, g_s);
|
||||
CK(cudaStreamSynchronize(g_s));
|
||||
const size_t qb = native_q8_1_bytes(n_in, nc);
|
||||
if (fetch<uint8_t>(a.q_ref, qb) != fetch<uint8_t>(a.q, qb)) {
|
||||
++bad;
|
||||
std::printf(" %s %d->%d, %d columns: the q8_1 blocks differ\n", f.name, n_in, n_out, nc);
|
||||
}
|
||||
const auto r0 = fetch<uint32_t>(y0, (size_t) nc * n_out), r1 = fetch<uint32_t>(y1, (size_t) nc * n_out);
|
||||
size_t diff = 0;
|
||||
for (size_t i = 0; i < r0.size(); ++i) diff += r0[i] != r1[i];
|
||||
if (diff) {
|
||||
++bad;
|
||||
std::printf(" %s %d->%d, %d columns: %zu of %zu values differ\n", f.name, n_in, n_out, nc, diff, r0.size());
|
||||
}
|
||||
}
|
||||
cudaFree(w);
|
||||
cudaFree(y0);
|
||||
cudaFree(y1);
|
||||
return bad;
|
||||
}
|
||||
|
||||
double time_calls(const std::function<void(int)>& fn, int reps = 48) {
|
||||
cudaGraph_t g;
|
||||
CK(cudaStreamBeginCapture(g_s, cudaStreamCaptureModeThreadLocal));
|
||||
for (int r = 0; r < reps; ++r) fn(r);
|
||||
CK(cudaStreamEndCapture(g_s, &g));
|
||||
cudaGraphExec_t ex;
|
||||
CK(cudaGraphInstantiate(&ex, g, 0));
|
||||
cudaGraphDestroy(g);
|
||||
const auto t0 = std::chrono::steady_clock::now();
|
||||
while (std::chrono::duration<double>(std::chrono::steady_clock::now() - t0).count() < 0.3) { // clocks up
|
||||
CK(cudaGraphLaunch(ex, g_s));
|
||||
CK(cudaStreamSynchronize(g_s));
|
||||
}
|
||||
cudaEvent_t a, b;
|
||||
CK(cudaEventCreate(&a));
|
||||
CK(cudaEventCreate(&b));
|
||||
const int iters = 20;
|
||||
CK(cudaEventRecord(a, g_s));
|
||||
for (int i = 0; i < iters; ++i) CK(cudaGraphLaunch(ex, g_s));
|
||||
CK(cudaEventRecord(b, g_s));
|
||||
CK(cudaEventSynchronize(b));
|
||||
float ms = 0;
|
||||
CK(cudaEventElapsedTime(&ms, a, b));
|
||||
cudaGraphExecDestroy(ex);
|
||||
cudaEventDestroy(a);
|
||||
cudaEventDestroy(b);
|
||||
return ms * 1000.0 / iters / reps;
|
||||
}
|
||||
|
||||
void bench_shape(const Fmt& f, const Act& a, int n_out) {
|
||||
const int n_in = a.n_in;
|
||||
const double mb = (double) native_mmvq_weight_bytes(f.type, n_in, n_out) / 1e6;
|
||||
const int copies = std::max(4, (int) (24.0 / mb) + 1);
|
||||
std::vector<void*> w((size_t) copies);
|
||||
for (auto& p : w) p = dev_weights(f, n_in, n_out);
|
||||
float* y;
|
||||
CK(cudaMalloc(&y, (size_t) 8 * n_out * 4));
|
||||
std::printf("%-6s %5d -> %5d (%5.2f MB), us (GB/s):", f.name, n_in, n_out, mb);
|
||||
for (int nc : {2, 3, 4, 8}) {
|
||||
native_quantize_q8_1(a.x, a.q_ref, n_in, nc, g_s);
|
||||
native_quantize_q8_1_il(a.x, a.q, a.il, n_in, nc, g_s);
|
||||
const double t0 = time_calls([&](int r) { native_mmvq(f.type, w[(size_t) (r % copies)], a.q_ref, y, n_in, n_out, nc, g_s); });
|
||||
const double t1 = time_calls([&](int r) {
|
||||
native_mmvq_il(f.type, w[(size_t) (r % copies)], a.q, a.il, y, n_in, n_out, nc, g_s);
|
||||
});
|
||||
std::printf(" %d: %5.1f -> %5.1f (%3.0f)", nc, t0, t1, mb * 1e3 / t1);
|
||||
}
|
||||
std::printf("\n");
|
||||
for (auto p : w) cudaFree(p);
|
||||
cudaFree(y);
|
||||
}
|
||||
|
||||
} // namespace
|
||||
|
||||
int main(int argc, char** argv) {
|
||||
const bool bench = argc > 1 && std::strcmp(argv[1], "--bench") == 0;
|
||||
CK(cudaStreamCreateWithFlags(&g_s, cudaStreamNonBlocking));
|
||||
Act a2560 = make_act(2560), a6144 = make_act(6144), a640 = make_act(640);
|
||||
if (bench) {
|
||||
for (const Fmt& f : kFmts) {
|
||||
bench_shape(f, a2560, 10240);
|
||||
bench_shape(f, a2560, 6144);
|
||||
bench_shape(f, a6144, 2560);
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
int bad = 0, cases = 0;
|
||||
for (const Fmt& f : kFmts) {
|
||||
const std::pair<Act*, int> shapes[] = {{&a2560, 10240}, {&a2560, 6144}, {&a6144, 2560}, {&a2560, 12288},
|
||||
{&a2560, 2050}, {&a2560, 4097}, {&a2560, 8192}, {&a2560, 512}};
|
||||
for (const auto& sh : shapes) {
|
||||
bad += check_shape(f, *sh.first, sh.second);
|
||||
++cases;
|
||||
}
|
||||
if (f.block_elems == 32) { // Q8_0 also at the shared expert's down shape
|
||||
bad += check_shape(f, a640, 2560);
|
||||
++cases;
|
||||
}
|
||||
}
|
||||
std::printf("mmvq_il_parity: %d shape/format cases x 2..8 columns: %s\n", cases, bad ? "FAILED" : "bitwise equal");
|
||||
return bad ? 1 : 0;
|
||||
}
|
||||
Reference in new issue
Block a user