mirror of
https://github.com/Niko1221/Strata.git
synced 2026-11-07 22:59:00 -10:00
Decode: hoisted-load, exact-integer interleaved GEMV kernels for Q6_K (STRATA_MMVQ_IL_PF=1)
native_mmvq_il_pf_kernel is native_mmvq_il_kernel with every weight load of a group (4 chunks x R rows) issued before any of the group's dots, in a branch-free main loop over the full groups (the old kernel's four `if (kbx < bpr)` sections each sat behind their own branch, so a warp had one 2-row x 1-block group in flight; ptxas also sinks a load whose only use is behind a branch into it), one pointer set per row with the other three blocks at immediates, and the tail group clamped and predicated. The dots use P2's exact integer forms: __vsubss4(x, 0x20) for Q6_K's [0, 63] codes as (x + 0x60606060) ^ 0x80808080 (2 instructions for ptxas's 8), and (float)(dp * sc) for |dp * sc| < 2^22 as the 1.5 * 2^23 magic add on the FMA pipe instead of the quarter-rate I2F (same for Q4_K / Q5_K's dot1 * sc, dot2 * m, and IQ4_XS's sumi then times (float)(ls - 32), exact under 2^24). The per-lane accumulation order, the v-slot sums and the xor tree are untouched, so every output is bitwise native_mmvq's: mmvq_il_parity checks each variant on every case, T 1-8, rows 1/2/4. The main loop is 441 SASS instructions per 8 row-blocks (55 a row-block, from 89), and ncu's long-scoreboard stall drops from 3.0 to 1.1-1.8 cycles an instruction; the win is then bounded by registers (128 at 4 blocks a SM), so the choice of variant and rows a warp is a measured sm_80 table (kIlPf_80 / kIlPf58_80, variant << 4 | rows, from --emit-table --pf 0,3,4): variant 3 (4 blocks a SM) or 4 (5 blocks a SM: the 2560-row shapes' 320 blocks fit one wave). Only Q6_K is in the table: Q4_K / Q5_K / IQ4_XS are bitwise and measured too but tie today's kernel within noise. Tried and rejected (measured slower): hoisting without P2 (latency hidden but 16% more instructions, issue-bound), a whole-group double buffer (168 registers, 3 blocks a SM) and a per-section ring (spills at 2 rows a warp). mmvq_il_parity --pf a,b,.. (interleaved medians, 5 rounds of 100), cold, engine clocks, T x rows a warp: Q6_K 2560x10240 T2 r2 37.6 -> 30.0 T3 r2 39.6 -> 37.0 T4 r2 45.8 -> 39.6 T5 r2 55.5 -> 48.5 us Q6_K 6144x2560 T2 r2 27.8 -> 19.1 T3 r1 25.1 -> 22.7 T4 r1 27.9 -> 23.8 T5 r2 45.5 -> 35.1 Q6_K 2560x6144 T2 r2 25.3 -> 20.1 T3 r2 28.0 -> 23.9 T4 r2 32.5 -> 25.1 T5 r2 35.1 -> 30.9 Engine (cmp_bench3, 3K context, interleaved ref / on / ref / on): dec 125.5 126.7 | 127.0 127.9 | 124.4 126.0 | 127.1 127.9 tok/s (+1.4%); the 11-prompt greedy set 11 / 11. Unset or 0: today's kernels; 2-6 force one variant for every type (A/B; 6 = hoisted loads without P2). HIP: the inline-PTX load helpers (ldg_nc_*, madw) and the Q6_K dot forms only the CUDA-only il_pf kernels use sit under !__HIPCC__, and the traits' TYPE constants are [[maybe_unused]], so this commit carries its own HIP guard and builds alone for gfx1100 and gfx906 with no new warnings in native_mmvq.cu. Co-Authored-By: Claude Fable 5.1 <[email protected]> Co-Authored-By: Claude Opus 5.5 <[email protected]>
This commit is contained in:
linxuhaoandClaude Fable 5.1 committed
1 parent
1f4b81398b
commit
116992ab8a
4 files changed
+715
-44
No files matched your search
@@ -166,5 +166,8 @@ void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, cons
|
||||
int native_mmvq_il_rows_for(int cc, int ggml_type, int ncols, int n_out);
|
||||
// Tests and benchmarks: every native_mmvq_il call takes `rows` a warp (0: the table).
|
||||
void native_mmvq_il_tune(int rows);
|
||||
// Tests and benchmarks: the P1 / P2 kernel variant for every native_mmvq_il call (0: today's kernel, 2-6 as
|
||||
// STRATA_MMVQ_IL_PF; -1: the environment).
|
||||
void native_mmvq_il_pf(int pf);
|
||||
|
||||
} // namespace strata::kernels
|
||||
@@ -650,6 +650,66 @@ __device__ __forceinline__ float q6_q8_dot_impl(int vl, int vh, const int* __res
|
||||
}
|
||||
return d * sumf;
|
||||
}
|
||||
#if !defined(__HIPCC__) // CUDA only: the native_mmvq_il_pf_kernel family below (itself CUDA only) is the only user
|
||||
// q6_q8_dot_impl with the two scales already in registers (the interleaved kernels' hoisted loads, P1): the same
|
||||
// expression in the same order
|
||||
__device__ __forceinline__ float q6_q8_dot_sc(int vl, int vh, const int* __restrict__ u, int sc0, int sc1, float d,
|
||||
const float* __restrict__ d8) {
|
||||
float sumf = 0.0f;
|
||||
const int scs[2] = {sc0, sc1};
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int sc = scs[i];
|
||||
const int vil = (vl >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int vih = ((vh >> (4 * i)) << 4) & 0x30303030;
|
||||
const int vi = __vsubss4(vil | vih, 0x20202020);
|
||||
sumf += d8[i] * (STRATA_DP4A(vi, u[i], 0) * sc);
|
||||
}
|
||||
return d * sumf;
|
||||
}
|
||||
|
||||
// P1: loads the compiler may not sink into a later branch (it moved the Q6_K scale and d loads of the hoisted group
|
||||
// into the dot sections, where each one is a dependent L1 round trip); the same ld.global.nc the __restrict__ loads use
|
||||
__device__ __forceinline__ int ldg_nc_u16(const void* p) { unsigned v; asm volatile("ld.global.nc.u16 %0, [%1];" : "=r"(v) : "l"(p)); return int(v); }
|
||||
__device__ __forceinline__ int ldg_nc_s8(const void* p) { int v; asm volatile("ld.global.nc.s8 %0, [%1];" : "=r"(v) : "l"(p)); return v; }
|
||||
__device__ __forceinline__ int ldg_nc_u8(const void* p) { unsigned v; asm volatile("ld.global.nc.u8 %0, [%1];" : "=r"(v) : "l"(p)); return int(v); }
|
||||
__device__ __forceinline__ int ldg_nc_s32(const void* p) { int v; asm volatile("ld.global.nc.s32 %0, [%1];" : "=r"(v) : "l"(p)); return v; }
|
||||
__device__ __forceinline__ int ldg_nc_b2(const void* p, int i32) { // load_int_b2 pinned: two u16, little-endian
|
||||
const auto* x = static_cast<const uint16_t*>(p) + 2 * i32;
|
||||
return ldg_nc_u16(x) | (ldg_nc_u16(x + 1) << 16);
|
||||
}
|
||||
// P1: base + k * stride as one IMAD.WIDE.U32 (an asm so the front end does not share k * stride between pointers and
|
||||
// turn each into a 64-bit add pair)
|
||||
__device__ __forceinline__ const void* madw(const void* base, int k, int stride) {
|
||||
const void* r;
|
||||
asm("mad.wide.u32 %0, %1, %2, %3;" : "=l"(r) : "r"(k), "r"(stride), "l"(base));
|
||||
return r;
|
||||
}
|
||||
// P2: (float) p for |p| < 2^22 as an add on the FMA pipe instead of the quarter-rate I2F: p + 1.5 * 2^23
|
||||
// as float bits is exactly 1.5 * 2^23 + p (unit ulp), and the subtraction is exact (Sterbenz), so the value is the
|
||||
// conversion's, +0 included
|
||||
__device__ __forceinline__ float i2f_small(int p_plus_magic) { return __int_as_float(p_plus_magic) - 12582912.0f; }
|
||||
constexpr int kI2fMagic = 0x4B400000;
|
||||
// P2: __vsubss4(x, 0x20202020) for bytes of x in [0, 63] (Q6_K's codes, so nothing saturates): x | 0x80 is x + 0x80 per
|
||||
// byte, minus 0x20 never borrows (0x60..0x9f), and the xor takes the 0x80 back: the same int8 bytes, 2 instructions
|
||||
// instead of ptxas's 8 for the saturating subtract sm_80 lacks
|
||||
__device__ __forceinline__ int sub32_bytes(int x) { return (x + 0x60606060) ^ 0x80808080; }
|
||||
// q6_q8_dot_sc with P2's integer forms: the dp4a inputs are the same 32-bit values, (dp * sc) the same int, and
|
||||
// i2f_small its exact float (|dp * sc| <= 4 * 32 * 127 * 128 = 2.1M < 2^22), so every FFMA sees the bits q6_q8_dot_impl's does
|
||||
__device__ __forceinline__ float q6_q8_dot_fast(int vl, int vh, const int* __restrict__ u, int sc0, int sc1, float d,
|
||||
const float* __restrict__ d8) {
|
||||
float sumf = 0.0f;
|
||||
const int scs[2] = {sc0, sc1};
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
const int vil = (vl >> (4 * i)) & 0x0f0f0f0f;
|
||||
const int vih = ((vh >> (4 * i)) << 4) & 0x30303030;
|
||||
const int vi = sub32_bytes(vil | vih);
|
||||
sumf += d8[i] * i2f_small(STRATA_DP4A(vi, u[i], 0) * scs[i] + kI2fMagic);
|
||||
}
|
||||
return d * sumf;
|
||||
}
|
||||
#endif // !__HIPCC__
|
||||
|
||||
__device__ __forceinline__ float q6_q8_dot(const Q6KBlock* __restrict__ w,
|
||||
const Q81Block* __restrict__ x, int iqs) {
|
||||
@@ -835,6 +895,7 @@ constexpr int MAX_NCOLS = 8;
|
||||
// every column bitwise equal to it.
|
||||
struct Q5KTraits {
|
||||
using Block = Q5KBlock;
|
||||
[[maybe_unused]] static constexpr int TYPE = 13; // ggml type
|
||||
static constexpr int DIV = QK, T = QI / VDR, KBY = QK / Q8K, BPI = VDR * WARPS * WARP / QI;
|
||||
__device__ static int kqs(int tid) { return VDR * (tid % (QI / VDR)); }
|
||||
struct W { int vl[2], vh[2]; uint16_t aux[2]; half2 dm; int bq8_offset; };
|
||||
@@ -876,6 +937,7 @@ struct Q5KTraits {
|
||||
};
|
||||
struct Q4KTraits {
|
||||
using Block = Q4KBlock;
|
||||
[[maybe_unused]] static constexpr int TYPE = 12; // ggml type
|
||||
static constexpr int DIV = QK, T = QI / VDR, KBY = QK / Q8K, BPI = VDR * WARPS * WARP / QI;
|
||||
__device__ static int kqs(int tid) { return VDR * (tid % (QI / VDR)); }
|
||||
struct W { int v[2]; uint16_t aux[2]; half2 dm; int bq8_offset; };
|
||||
@@ -946,6 +1008,7 @@ struct Q20Traits {
|
||||
};
|
||||
struct Q3KTraits {
|
||||
using Block = Q3KBlock;
|
||||
[[maybe_unused]] static constexpr int TYPE = 11; // ggml type
|
||||
static constexpr int DIV = 256, T = 16, KBY = 8, BPI = WARPS * WARP / 16;
|
||||
__device__ static int kqs(int tid) { return tid % 16; }
|
||||
struct W { int vl, vh; float d; const uint8_t* scales; int scale_offset, bq8_offset; };
|
||||
@@ -972,6 +1035,7 @@ struct Q3KTraits {
|
||||
};
|
||||
struct Q6KTraits {
|
||||
using Block = Q6KBlock;
|
||||
[[maybe_unused]] static constexpr int TYPE = 14; // ggml type
|
||||
static constexpr int DIV = 256, T = 32, KBY = 8, BPI = WARPS * WARP / 32;
|
||||
__device__ static int kqs(int tid) { return tid % 32; }
|
||||
struct W { int vl, vh; float d; const int8_t* scales; int bq8_offset; };
|
||||
@@ -999,6 +1063,7 @@ struct Q6KTraits {
|
||||
};
|
||||
struct IQ4XSTraits {
|
||||
using Block = IQ4XSBlock;
|
||||
[[maybe_unused]] static constexpr int TYPE = 23; // ggml type
|
||||
static constexpr int DIV = 256, T = 8, KBY = 8, BPI = 4 * WARPS * WARP / 32;
|
||||
__device__ static int kqs(int tid) { return 4 * (tid % 8); }
|
||||
struct W { int2 v[4]; int ls; float dw; };
|
||||
@@ -1556,10 +1621,141 @@ struct IlIQ4XS {
|
||||
out[c] = d * sumi;
|
||||
}
|
||||
}
|
||||
// P1 (native_mmvq_il_pf_kernel): the block's bytes as loaded, the table lookup and the scale decode deferred to
|
||||
// applyr so the loads of a whole group issue before any of them is consumed; the same integers and floats as apply
|
||||
static constexpr bool HAS_PF = true;
|
||||
struct RP { const uint8_t *qs, *sl, *blk; }; // a row's blocks for this lane: its qs ints, its scales_l byte, the block
|
||||
__device__ static RP rowptrs(const F::Block* __restrict__ row, int iqs) {
|
||||
const auto* b = reinterpret_cast<const uint8_t*>(row);
|
||||
return {b + 8 + 4 * iqs, b + 4 + iqs / 8, b};
|
||||
}
|
||||
struct WR { int q[4]; int sl, sh, d; };
|
||||
// the pointers of block kbx (RP::at), and the loads of block kbx + dv from them (dv a compile-time offset, so the
|
||||
// group's loads share one pointer set and differ by an immediate)
|
||||
__device__ static RP at(const RP& rp, int kbx) {
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
return {static_cast<const uint8_t*>(madw(rp.qs, kbx, S)), static_cast<const uint8_t*>(madw(rp.sl, kbx, S)),
|
||||
static_cast<const uint8_t*>(madw(rp.blk, kbx, S))};
|
||||
}
|
||||
__device__ static WR loadr(const RP& rp, int dv) {
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
const auto* qs = reinterpret_cast<const int*>(rp.qs + dv * S);
|
||||
WR r;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) r.q[j] = ldg_nc_s32(qs + j);
|
||||
r.sl = ldg_nc_u8(rp.sl + dv * S);
|
||||
r.sh = ldg_nc_u16(rp.blk + dv * S + 2);
|
||||
r.d = ldg_nc_u16(rp.blk + dv * S);
|
||||
return r;
|
||||
}
|
||||
// the activations: position-major (every int of the 8 at a stride of nb), so the kernel's block pointers are not used
|
||||
__device__ static int xb(int) { return 0; }
|
||||
__device__ static int xp(int) { return 0; }
|
||||
template <int NC, bool FAST, class X>
|
||||
__device__ static void applyr(const WR& rr, const X& x, const int*, const float*, int kby, int iqs, float (&out)[NC]) {
|
||||
F::W r;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) r.v[j] = iq4_table_lookup(rr.q[j]);
|
||||
r.ls = ((rr.sl >> (iqs & 0x04)) & 0x0f) | (((rr.sh >> (iqs / 2)) & 0x03) << 4);
|
||||
r.dw = __half2float(__ushort_as_half((unsigned short) rr.d));
|
||||
if constexpr (!FAST) {
|
||||
apply<NC>(r, x, kby, iqs, out);
|
||||
} else {
|
||||
// P2: sumi (|sumi| <= 32 * 127 * 127 = 516K < 2^22) through i2f_small, then times (float)(ls - 32): the product
|
||||
// |sumi * (ls - 32)| < 2^24 is exact in fp32, so it is the float of the int apply() converts; the same d * sumi
|
||||
const int b = kby + iqs / 4;
|
||||
int u0[4][NC], u1[4][NC];
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
x.u(b, j, u0[j]);
|
||||
x.u(b, j + 4, u1[j]);
|
||||
}
|
||||
float d8[NC];
|
||||
x.scales(b, d8);
|
||||
const float fls = float(r.ls - 32);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
int sumi = 0;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < 4; ++j) {
|
||||
sumi = STRATA_DP4A(r.v[j].x, u0[j][c], sumi);
|
||||
sumi = STRATA_DP4A(r.v[j].y, u1[j][c], sumi);
|
||||
}
|
||||
const float fsumi = i2f_small(sumi + kI2fMagic) * fls;
|
||||
const float d = r.dw * d8[c];
|
||||
out[c] = d * fsumi;
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
struct IlQ4K {
|
||||
using F = Q4KTraits;
|
||||
static constexpr bool PM = false;
|
||||
static constexpr bool HAS_PF = true; // P1: F::W holds every byte of the block already (no pointer into it)
|
||||
struct RP { const uint8_t *qs, *sc, *blk; int iqs; }; // this lane's ql ints, its scale words, the block (dm)
|
||||
__device__ static RP rowptrs(const F::Block* __restrict__ row, int iqs) {
|
||||
const auto* b = reinterpret_cast<const uint8_t*>(row);
|
||||
const int bq8_offset = 2 * ((iqs / 2) / 4), jm = (bq8_offset / 2) & 1;
|
||||
return {b + 16 + 16 * bq8_offset + 4 * ((iqs / 2) % 4), b + 4 + 2 * jm, b, iqs};
|
||||
}
|
||||
using WR = F::W;
|
||||
__device__ static RP at(const RP& rp, int kbx) {
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
return {static_cast<const uint8_t*>(madw(rp.qs, kbx, S)), static_cast<const uint8_t*>(madw(rp.sc, kbx, S)),
|
||||
static_cast<const uint8_t*>(madw(rp.blk, kbx, S)), rp.iqs};
|
||||
}
|
||||
__device__ static WR loadr(const RP& rp, int dv) { // F::load's bytes and aux arithmetic, block dv past rp's
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
const auto* ql = reinterpret_cast<const int*>(rp.qs + dv * S);
|
||||
const auto* sc = rp.sc + dv * S;
|
||||
WR r;
|
||||
r.bq8_offset = 2 * ((rp.iqs / 2) / 4);
|
||||
r.v[0] = ldg_nc_s32(ql);
|
||||
r.v[1] = ldg_nc_s32(ql + 4);
|
||||
const uint32_t s0 = ldg_nc_u16(sc), s2 = ldg_nc_u16(sc + 4), s4 = ldg_nc_u16(sc + 8);
|
||||
const int dm = ldg_nc_s32(rp.blk + dv * S);
|
||||
const int j = r.bq8_offset / 2;
|
||||
const uint32_t hi = uint32_t(-int32_t(j >= 2));
|
||||
r.aux[0] = uint16_t(((s0 & 0x3f3f) & ~hi) | ((((s4 >> 0) & 0x0f0f) | ((s0 & 0xc0c0) >> 2)) & hi));
|
||||
r.aux[1] = uint16_t(((s2 & 0x3f3f) & ~hi) | ((((s4 >> 4) & 0x0f0f) | ((s2 & 0xc0c0) >> 2)) & hi));
|
||||
{ half2 h; std::memcpy(&h, &dm, 4); r.dm = h; }
|
||||
return r;
|
||||
}
|
||||
__device__ static int xb(int iqs) { return 2 * ((iqs / 2) / 4); } // the lane's q8_1 block and int within a 256-chunk
|
||||
__device__ static int xp(int iqs) { return (iqs / 2) % 4; }
|
||||
template <int NC, bool FAST, class X>
|
||||
__device__ static void applyr(const WR& r, const X& x, const int* qb, const float* db, int kby, int iqs, float (&out)[NC]) {
|
||||
if constexpr (!FAST) {
|
||||
apply<NC>(r, x, kby, iqs, out);
|
||||
} else { // P2: q4_q8_dot_impl with (dot1 * sc), (dot2 * m) through i2f_small (|dot1 * sc| <= 15240 * 63, |dot2 * m| <= 1016 * 63)
|
||||
int u[4][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
x.qat(qb + i * X::QB, u[2 * i]);
|
||||
x.qat(qb + i * X::QB + x.qp(4), u[2 * i + 1]);
|
||||
x.dat(db + i * X::DB, 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 = STRATA_DP4A(v1i, u[2 * i + 1][c], STRATA_DP4A(v0i, u[2 * i][c], 0));
|
||||
const int dot2 = STRATA_DP4A(0x01010101, u[2 * i + 1][c], STRATA_DP4A(0x01010101, u[2 * i][c], 0));
|
||||
sumf_d += d8[i][c] * i2f_small(dot1 * sc[i] + kI2fMagic);
|
||||
sumf_m += d8[i][c] * i2f_small(dot2 * m[i] + kI2fMagic);
|
||||
}
|
||||
out[c] = dm4f.x * sumf_d - dm4f.y * sumf_m;
|
||||
}
|
||||
}
|
||||
}
|
||||
template <int NC, class X>
|
||||
__device__ static void apply(const F::W& r, const X& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[4][NC];
|
||||
@@ -1583,6 +1779,78 @@ struct IlQ4K {
|
||||
struct IlQ5K {
|
||||
using F = Q5KTraits;
|
||||
static constexpr bool PM = false;
|
||||
static constexpr bool HAS_PF = true; // P1: as IlQ4K
|
||||
struct RP { const uint8_t *qs, *qh, *sc, *blk; int iqs; };
|
||||
__device__ static RP rowptrs(const F::Block* __restrict__ row, int iqs) {
|
||||
const auto* b = reinterpret_cast<const uint8_t*>(row);
|
||||
const int bq8_offset = 2 * ((iqs / 2) / 4), jm = (bq8_offset / 2) & 1;
|
||||
return {b + 48 + 16 * bq8_offset + 4 * ((iqs / 2) % 4), b + 16 + 4 * ((iqs / 2) % 4), b + 4 + 2 * jm, b, iqs};
|
||||
}
|
||||
using WR = F::W;
|
||||
__device__ static RP at(const RP& rp, int kbx) {
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
return {static_cast<const uint8_t*>(madw(rp.qs, kbx, S)), static_cast<const uint8_t*>(madw(rp.qh, kbx, S)),
|
||||
static_cast<const uint8_t*>(madw(rp.sc, kbx, S)), static_cast<const uint8_t*>(madw(rp.blk, kbx, S)), rp.iqs};
|
||||
}
|
||||
__device__ static WR loadr(const RP& rp, int dv) { // F::load's bytes and aux arithmetic, block dv past rp's
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
const auto* ql = reinterpret_cast<const int*>(rp.qs + dv * S);
|
||||
const auto* qh = reinterpret_cast<const int*>(rp.qh + dv * S);
|
||||
const auto* sc = rp.sc + dv * S;
|
||||
WR r;
|
||||
r.bq8_offset = 2 * ((rp.iqs / 2) / 4);
|
||||
r.vl[0] = ldg_nc_s32(ql);
|
||||
r.vl[1] = ldg_nc_s32(ql + 4);
|
||||
r.vh[0] = ldg_nc_s32(qh) >> r.bq8_offset;
|
||||
r.vh[1] = ldg_nc_s32(qh + 4) >> r.bq8_offset;
|
||||
const uint32_t s0 = ldg_nc_u16(sc), s2 = ldg_nc_u16(sc + 4), s4 = ldg_nc_u16(sc + 8);
|
||||
const int dm = ldg_nc_s32(rp.blk + dv * S);
|
||||
const int j = r.bq8_offset / 2;
|
||||
const uint32_t hi = uint32_t(-int32_t(j >= 2));
|
||||
r.aux[0] = uint16_t(((s0 & 0x3f3f) & ~hi) | ((((s4 >> 0) & 0x0f0f) | ((s0 & 0xc0c0) >> 2)) & hi));
|
||||
r.aux[1] = uint16_t(((s2 & 0x3f3f) & ~hi) | ((((s4 >> 4) & 0x0f0f) | ((s2 & 0xc0c0) >> 2)) & hi));
|
||||
{ half2 h; std::memcpy(&h, &dm, 4); r.dm = h; }
|
||||
return r;
|
||||
}
|
||||
__device__ static int xb(int iqs) { return 2 * ((iqs / 2) / 4); }
|
||||
__device__ static int xp(int iqs) { return (iqs / 2) % 4; }
|
||||
template <int NC, bool FAST, class X>
|
||||
__device__ static void applyr(const WR& r, const X& x, const int* qb, const float* db, int kby, int iqs, float (&out)[NC]) {
|
||||
if constexpr (!FAST) {
|
||||
apply<NC>(r, x, kby, iqs, out);
|
||||
} else { // P2: q5_q8_dot_impl with (dot1 * sc), (dot2 * m) through i2f_small (|dot1 * sc| <= 31496 * 63 < 2^22)
|
||||
int u[4][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
x.qat(qb + i * X::QB, u[2 * i]);
|
||||
x.qat(qb + i * X::QB + x.qp(4), u[2 * i + 1]);
|
||||
x.dat(db + i * X::DB, 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 = STRATA_DP4A(v0i, u[2 * i][c], STRATA_DP4A(v1i, u[2 * i + 1][c], 0));
|
||||
const int dot2 = STRATA_DP4A(0x01010101, u[2 * i][c], STRATA_DP4A(0x01010101, u[2 * i + 1][c], 0));
|
||||
sumf_d += d8[i][c] * i2f_small(dot1 * sc[i] + kI2fMagic);
|
||||
sumf_m += d8[i][c] * i2f_small(dot2 * m[i] + kI2fMagic);
|
||||
}
|
||||
out[c] = dm5f.x * sumf_d - dm5f.y * sumf_m;
|
||||
}
|
||||
}
|
||||
}
|
||||
template <int NC, class X>
|
||||
__device__ static void apply(const F::W& r, const X& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[4][NC];
|
||||
@@ -1608,6 +1876,7 @@ struct IlQ5K {
|
||||
struct IlQ3K {
|
||||
using F = Q3KTraits;
|
||||
static constexpr bool PM = false;
|
||||
static constexpr bool HAS_PF = false; // the Q2_0 pack's type: keeps native_mmvq_il_kernel (F::W points into the block)
|
||||
template <int NC, class X>
|
||||
__device__ static void apply(const F::W& r, const X& x, int kby, int iqs, float (&out)[NC]) {
|
||||
int u[4][NC];
|
||||
@@ -1662,8 +1931,210 @@ struct IlQ6K {
|
||||
out[c] = q6_q8_dot_impl(r.vl, r.vh, uc, r.scales, r.d, dc);
|
||||
}
|
||||
}
|
||||
// P1: the two scale bytes loaded with the rest (F::W keeps a pointer into the block, which would put their loads
|
||||
// in the compute phase); q6_q8_dot_sc is q6_q8_dot_impl's expression
|
||||
static constexpr bool HAS_PF = true;
|
||||
// a row's four lane-adjusted byte pointers (ql, qh, scales, d): a block's loads are then one IMAD.WIDE each from kbx
|
||||
struct RP { const uint8_t *ql, *qh, *sc, *d; };
|
||||
__device__ static RP rowptrs(const F::Block* __restrict__ row, int iqs) {
|
||||
const auto* b = reinterpret_cast<const uint8_t*>(row);
|
||||
const int scale_offset = 8 * (iqs / 16) + (iqs % 16) / 4;
|
||||
return {b + 4 * iqs, b + 128 + 4 * (8 * (iqs / 16) + iqs % 8), b + 192 + scale_offset, b + 208}; // load_int_b2's int iqs
|
||||
}
|
||||
struct WR { int vl, vh, sc0, sc1, d; };
|
||||
__device__ static RP at(const RP& rp, int kbx) {
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
return {static_cast<const uint8_t*>(madw(rp.ql, kbx, S)), static_cast<const uint8_t*>(madw(rp.qh, kbx, S)),
|
||||
static_cast<const uint8_t*>(madw(rp.sc, kbx, S)), static_cast<const uint8_t*>(madw(rp.d, kbx, S))};
|
||||
}
|
||||
__device__ static WR loadr(const RP& rp, int dv) { // block dv past rp's: immediates off its four pointers
|
||||
constexpr int S = int(sizeof(F::Block));
|
||||
WR r;
|
||||
r.vl = ldg_nc_b2(rp.ql + dv * S, 0);
|
||||
r.vh = ldg_nc_b2(rp.qh + dv * S, 0);
|
||||
r.sc0 = ldg_nc_s8(rp.sc + dv * S);
|
||||
r.sc1 = ldg_nc_s8(rp.sc + dv * S + 4);
|
||||
r.d = ldg_nc_u16(rp.d + dv * S);
|
||||
return r;
|
||||
}
|
||||
__device__ static int xb(int iqs) { return 4 * (iqs / 16) + (iqs % 16) / 8; }
|
||||
__device__ static int xp(int iqs) { return iqs % 8; }
|
||||
template <int NC, bool FAST, class X>
|
||||
__device__ static void applyr(const WR& r, const X& x, const int* qb, const float* db, int kby, int iqs, float (&out)[NC]) {
|
||||
const int vh = r.vh >> (2 * ((iqs % 16) / 8));
|
||||
int u[2][NC];
|
||||
float d8[2][NC];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < 2; ++i) {
|
||||
x.qat(qb + 2 * i * X::QB, u[i]);
|
||||
x.dat(db + 2 * i * X::DB, d8[i]);
|
||||
}
|
||||
const float d = __half2float(__ushort_as_half((unsigned short) r.d));
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) {
|
||||
const int uc[2] = {u[0][c], u[1][c]};
|
||||
const float dc[2] = {d8[0][c], d8[1][c]};
|
||||
out[c] = FAST ? q6_q8_dot_fast(r.vl, vh, uc, r.sc0, r.sc1, d, dc) : q6_q8_dot_sc(r.vl, vh, uc, r.sc0, r.sc1, d, dc);
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
// P1: native_mmvq_il_kernel with every weight load of a group (the 4 virtual warps x R rows) issued
|
||||
// before any of the group's dots - the four `if (kbx < bpr)` sections of the kernel below each sit behind their own
|
||||
// branch, so a warp had one 2-row x 1-block group in flight (ncu: 432 GB/s on the qkv Q6_K matrix). The addresses are
|
||||
// clamped to the row's last block (a tail re-reads it; its dot is skipped), and with DB the next group's loads are
|
||||
// issued before this group's dots (two register sets, ping-pong, no copies). The per-lane accumulation order is the
|
||||
// kernel below's: acc[v] += block(4m + v) in ascending m, ((acc0 + acc1) + acc2) + acc3, warp_sum - the same bits.
|
||||
// STRATA_MMVQ_IL_PF=1: hoisted group loads (4 blocks a SM); =3: + P2's exact integer forms (the traits' applyr<FAST>);
|
||||
// =4: as 3 at 5 blocks a SM; =2 / =5: as 3 / 4 with the ring (DB 2: a section's registers reloaded with the next group's
|
||||
// blocks right after its dots); unset / 0: the kernel below.
|
||||
template <typename I, int NC, int R, int DB, int MINB, bool FAST>
|
||||
__launch_bounds__(4 * WARP, MINB)
|
||||
__global__ void native_mmvq_il_pf_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;
|
||||
using WR = typename I::WR;
|
||||
constexpr int SUB = WARP / F::T; // weight blocks a chunk covers
|
||||
const int lane = int(threadIdx.x) & (WARP - 1);
|
||||
const auto x = q8_1_cols<NC, I::PM>(xq, xd, n_in / Q8K);
|
||||
using X = std::remove_const_t<decltype(x)>;
|
||||
const int bpr = n_in / F::DIV;
|
||||
const int kqs = F::kqs(lane);
|
||||
const int sub = lane / F::T;
|
||||
const int nm = (bpr + SUB - 1) / SUB;
|
||||
for (int grp = int(blockIdx.x) * 4 + (int(threadIdx.x) >> 5); grp * R < n_out; grp += int(gridDim.x) * 4) {
|
||||
const int row0 = grp * R;
|
||||
typename I::RP wr[R];
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) wr[i] = I::rowptrs(w + std::size_t(min(row0 + i, n_out - 1)) * bpr, kqs); // a partial last group recomputes its last row
|
||||
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;
|
||||
// the groups whose four chunks are whole blocks for every lane run with no branch at all (ptxas sinks a load
|
||||
// whose only use sits behind a branch into it, which is what put the scale loads back into the dot sections);
|
||||
// the tail group clamps its loads and predicates its dots
|
||||
const int m_full = bpr / SUB;
|
||||
// a full group's 4 x R blocks: one pointer set per row at the group's first block, the other three at immediates
|
||||
// (v * SUB blocks on); a clamped (tail) group addresses each block on its own
|
||||
auto load_group = [&](int m0, bool clamp, WR (&wv)[4][R]) {
|
||||
if (!clamp) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) {
|
||||
const typename I::RP p0 = I::at(wr[i], SUB * m0 + sub);
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v) wv[v][i] = I::loadr(p0, v * SUB);
|
||||
}
|
||||
} else {
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v) {
|
||||
const int kbx = min(SUB * (m0 + v) + sub, bpr - 1);
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) wv[v][i] = I::loadr(I::at(wr[i], kbx), 0);
|
||||
}
|
||||
}
|
||||
};
|
||||
// the lane's activation ints and scales of the group's first chunk; a section's are an immediate on
|
||||
const int* xq0 = x.qbase(I::xb(kqs) + sub * F::KBY, I::xp(kqs));
|
||||
const float* xd0 = x.dbase(I::xb(kqs) + sub * F::KBY);
|
||||
auto dot_group = [&](int m0, bool pred, const WR (&wv)[4][R]) {
|
||||
const int* qg = static_cast<const int*>(madw(xq0, SUB * m0 * F::KBY, X::QB * 4));
|
||||
const float* dg = static_cast<const float*>(madw(xd0, SUB * m0 * F::KBY, X::DB * 4));
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v) {
|
||||
const int kbx = SUB * (m0 + v) + sub;
|
||||
if (!pred || kbx < bpr) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) {
|
||||
float o[NC];
|
||||
I::template applyr<NC, FAST>(wv[v][i], x, qg + v * SUB * F::KBY * X::QB, dg + v * SUB * F::KBY * X::DB,
|
||||
kbx * F::KBY, kqs, o);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[v][i][c] += o[c];
|
||||
}
|
||||
}
|
||||
}
|
||||
};
|
||||
int m0 = 0;
|
||||
if constexpr (DB == 0) {
|
||||
for (; m0 + 4 <= m_full; m0 += 4) {
|
||||
WR wv[4][R];
|
||||
load_group(m0, false, wv);
|
||||
dot_group(m0, false, wv);
|
||||
}
|
||||
} else if constexpr (DB == 2) {
|
||||
// a ring: after a section's dots its registers take the next group's blocks of that section, so the warp
|
||||
// always has about a group's worth of loads in flight at the register cost of one group (the loads sit
|
||||
// behind a uniform branch whose consumers are in the next iteration, so nothing can sink)
|
||||
WR wv[4][R];
|
||||
if (m0 + 4 <= m_full) load_group(0, false, wv);
|
||||
for (; m0 + 4 <= m_full; m0 += 4) {
|
||||
const bool more = m0 + 8 <= m_full;
|
||||
typename I::RP pn[R];
|
||||
if (more) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) pn[i] = I::at(wr[i], SUB * (m0 + 4) + sub);
|
||||
}
|
||||
const int* qg = static_cast<const int*>(madw(xq0, SUB * m0 * F::KBY, X::QB * 4));
|
||||
const float* dg = static_cast<const float*>(madw(xd0, SUB * m0 * F::KBY, X::DB * 4));
|
||||
#pragma unroll
|
||||
for (int v = 0; v < 4; ++v) {
|
||||
const int kbx = SUB * (m0 + v) + sub;
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) {
|
||||
float o[NC];
|
||||
I::template applyr<NC, FAST>(wv[v][i], x, qg + v * SUB * F::KBY * X::QB, dg + v * SUB * F::KBY * X::DB,
|
||||
kbx * F::KBY, kqs, o);
|
||||
#pragma unroll
|
||||
for (int c = 0; c < NC; ++c) acc[v][i][c] += o[c];
|
||||
}
|
||||
if (more) {
|
||||
#pragma unroll
|
||||
for (int i = 0; i < R; ++i) wv[v][i] = I::loadr(pn[i], v * SUB);
|
||||
}
|
||||
}
|
||||
}
|
||||
} else { // two register sets: a group's loads issue before the previous group's dots; no branch inside the loop
|
||||
WR wa[4][R], wb[4][R];
|
||||
if (m0 + 4 <= m_full) load_group(0, false, wa);
|
||||
for (; m0 + 12 <= m_full; m0 += 8) {
|
||||
load_group(m0 + 4, false, wb);
|
||||
dot_group(m0, false, wa);
|
||||
load_group(m0 + 8, false, wa);
|
||||
dot_group(m0 + 4, false, wb);
|
||||
}
|
||||
if (m0 + 8 <= m_full) {
|
||||
load_group(m0 + 4, false, wb);
|
||||
dot_group(m0, false, wa);
|
||||
dot_group(m0 + 4, false, wb);
|
||||
m0 += 8;
|
||||
} else if (m0 + 4 <= m_full) {
|
||||
dot_group(m0, false, wa);
|
||||
m0 += 4;
|
||||
}
|
||||
}
|
||||
if (m0 < nm) {
|
||||
WR wv[4][R];
|
||||
load_group(m0, true, wv);
|
||||
dot_group(m0, true, wv);
|
||||
}
|
||||
#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 NC, int R>
|
||||
__launch_bounds__(4 * WARP)
|
||||
__global__ void native_mmvq_il_kernel(const typename I::F::Block* __restrict__ w, const int* __restrict__ xq,
|
||||
@@ -1717,23 +2188,57 @@ __global__ void native_mmvq_il_kernel(const typename I::F::Block* __restrict__ w
|
||||
}
|
||||
}
|
||||
|
||||
int g_il_pf = -1; // native_mmvq_il_pf (tests, benchmarks): the variant for every launch; -1: STRATA_MMVQ_IL_PF
|
||||
int il_pf_env() {
|
||||
static const int env = [] { const char* v = std::getenv("STRATA_MMVQ_IL_PF"); return v ? std::atoi(v) : 0; }();
|
||||
return env;
|
||||
}
|
||||
int il_pf_table(int type, int ncols, int n_out); // the sm_80 table below: pf << 4 | rows, 0 = today's kernel and table
|
||||
// the variant of one launch (g_il_pf, else STRATA_MMVQ_IL_PF: 0 / unset today's kernel, 1 the measured table, 2-6 that
|
||||
// variant for every type and shape), and with the table its rows a warp (0: the rows table's)
|
||||
int il_pf_for(int type, int ncols, int n_out, int* rows) {
|
||||
*rows = 0;
|
||||
if (g_il_pf >= 0) return g_il_pf;
|
||||
const int env = il_pf_env();
|
||||
if (env != 1) return env;
|
||||
const int t = il_pf_table(type, ncols, n_out);
|
||||
*rows = t & 15;
|
||||
return t >> 4;
|
||||
}
|
||||
// one (type, columns, rows) launch: today's kernel, or the variant il_pf_for names for a type that has them
|
||||
template <typename I, int NC, int R>
|
||||
void launch_il_one(unsigned blocks, const typename I::F::Block* w, const int* xq, const float* xd, float* y, int n_in, int n_out,
|
||||
cudaStream_t s, int pf) {
|
||||
if constexpr (I::HAS_PF) {
|
||||
switch (pf) {
|
||||
case 6: native_mmvq_il_pf_kernel<I, NC, R, 0, 4, false><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); return;
|
||||
case 2: native_mmvq_il_pf_kernel<I, NC, R, 2, 4, true><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); return;
|
||||
case 3: native_mmvq_il_pf_kernel<I, NC, R, 0, 4, true><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); return;
|
||||
case 4: native_mmvq_il_pf_kernel<I, NC, R, 0, 5, true><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); return;
|
||||
case 5: native_mmvq_il_pf_kernel<I, NC, R, 2, 5, true><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); return;
|
||||
default: break;
|
||||
}
|
||||
}
|
||||
native_mmvq_il_kernel<I, NC, R><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out);
|
||||
}
|
||||
|
||||
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) {
|
||||
void launch_il(const void* weights, const void* x_il, float* y, int n_in, int n_out, int ncols, cudaStream_t s, int pf) {
|
||||
const auto* w = static_cast<const typename I::F::Block*>(weights);
|
||||
const unsigned blocks = unsigned(((n_out + R - 1) / R + 3) / 4);
|
||||
if (ncols == 1) { // 143-dense: one column reads the plain q8_1 blocks (x_il is then the caller's x_q8_1)
|
||||
native_mmvq_il_kernel<I, 1, R><<<blocks, 4 * WARP, 0, s>>>(w, static_cast<const int*>(x_il), nullptr, y, n_in, n_out);
|
||||
launch_il_one<I, 1, R>(blocks, w, static_cast<const int*>(x_il), nullptr, y, n_in, n_out, s, pf);
|
||||
return;
|
||||
}
|
||||
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;
|
||||
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;
|
||||
#define STRATA_IL_CASE(N) case N: launch_il_one<I, N, R>(blocks, w, xq, xd, y, n_in, n_out, s, pf); break;
|
||||
STRATA_IL_CASE(2) STRATA_IL_CASE(3) STRATA_IL_CASE(4)
|
||||
#undef STRATA_IL_CASE
|
||||
// 147-verify: 5..8 columns (the cp = 8 copy) at 1 or 2 rows a warp; launch_il_rows maps 4 rows to 2 there
|
||||
#define STRATA_IL_CASE(N) case N: if constexpr (R <= 2) native_mmvq_il_kernel<I, N, R><<<blocks, 4 * WARP, 0, s>>>(w, xq, xd, y, n_in, n_out); break;
|
||||
#define STRATA_IL_CASE(N) case N: if constexpr (R <= 2) launch_il_one<I, N, R>(blocks, w, xq, xd, y, n_in, n_out, s, pf); break;
|
||||
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 1 <= ncols <= 8");
|
||||
@@ -1742,11 +2247,14 @@ void launch_il(const void* weights, const void* x_il, float* y, int n_in, int n_
|
||||
|
||||
template <typename I>
|
||||
void launch_il_rows(int r, const void* weights, const void* x_il, float* y, int n_in, int n_out, int ncols, cudaStream_t s) {
|
||||
int pf_rows = 0;
|
||||
const int pf = il_pf_for(I::F::TYPE, ncols, n_out, &pf_rows);
|
||||
if (pf_rows != 0) r = pf_rows; // the pf table measured its rows with its kernel
|
||||
if (ncols > 4 && r == 4) r = 2; // 5..8 columns: 1 or 2 rows a warp
|
||||
switch (r) {
|
||||
case 1: launch_il<I, 1>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 4: launch_il<I, 4>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
default: launch_il<I, 2>(weights, x_il, y, n_in, n_out, ncols, s); break;
|
||||
case 1: launch_il<I, 1>(weights, x_il, y, n_in, n_out, ncols, s, pf); break;
|
||||
case 4: launch_il<I, 4>(weights, x_il, y, n_in, n_out, ncols, s, pf); break;
|
||||
default: launch_il<I, 2>(weights, x_il, y, n_in, n_out, ncols, s, pf); break;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -3054,6 +3562,20 @@ constexpr IlRows kIlRows_80[] = {
|
||||
{14, {{0, 0, 2, 2, 2}, {0, 2, 2, 2, 2}, {0, 2, 4, 2, 2}}}, // Q6_K
|
||||
{11, {{1, 2, 2, 2, 2}, {1, 1, 2, 2, 2}, {1, 2, 2, 2, 2}}}, // Q3_K
|
||||
};
|
||||
// sm_80 P1/P2: the hoisted-load / exact-integer kernel variant and its rows a warp per cell, pf << 4 | rows
|
||||
// (0x32 = variant 3 at 2 rows a warp; 0 = today's kernel with kIlRows_80), read off mmvq_il_parity --bench --emit-table
|
||||
// --pf 0,3,4 on the CMP 170HX with the same 3% rule against native_mmvq; STRATA_MMVQ_IL_PF=1 takes it. Every variant is
|
||||
// bitwise native_mmvq (the tool checks every case, T 1-8, rows 1/2/4, each variant).
|
||||
struct IlPf { int type; uint8_t v[3][5]; };
|
||||
struct IlPf58 { int type; uint8_t v[4][5]; };
|
||||
// Only Q6_K: the same variants are bitwise and measured for Q4_K / Q5_K / IQ4_XS too, but there they tie today's kernel
|
||||
// within the run-to-run noise (their emitted cells flip between 0 and a variant), so those types keep kIlRows_80.
|
||||
constexpr IlPf kIlPf_80[] = {
|
||||
{14, {{0x31, 0x42, 0x32, 0x42, 0x42}, {0x31, 0x41, 0x32, 0x32, 0x41}, {0x31, 0x41, 0x32, 0x32, 0x32}}}, // Q6_K
|
||||
};
|
||||
constexpr IlPf58 kIlPf58_80[] = { // (rows 4 runs as 2 at 5+ columns, so the emitted 0x34 cells are written 0x32)
|
||||
{14, {{0x31, 0x42, 0x32, 0x32, 0x32}, {0x31, 0x31, 0x32, 0x32, 0x32}, {0x31, 0x31, 0x02, 0x32, 0x02}, {0x31, 0x31, 0x31, 0x31, 0x31}}}, // Q6_K
|
||||
};
|
||||
// Turing (sm_75): read off mmvq_il_parity --bench --emit-table on an RTX 2080 Ti 22 GB (SM 1650 MHz), the same 3% rule.
|
||||
// The plain multi-column kernels are bound by dp4a throughput on Turing from 2 columns on (IQ4_XS 2560x10240 at 4
|
||||
// columns 81.1 -> 37.4 us, the Q5_K head 1671 -> 1045 us); bitwise equal to native_mmvq. No one-column table here:
|
||||
@@ -3220,9 +3742,22 @@ int il_cc() {
|
||||
return cc[dev];
|
||||
}
|
||||
int il_rows(int type, int ncols, int n_out) { return native_mmvq_il_rows_for(il_cc(), type, ncols, n_out); }
|
||||
int il_pf_table(int type, int ncols, int n_out) {
|
||||
if (il_cc() != 80 || ncols < 2 || ncols > 8) return 0; // only the CMP 170HX measured one
|
||||
const int cls = n_out < 2048 ? 0 : n_out < 4096 ? 1 : n_out < 8192 ? 2 : n_out < 12288 ? 3 : 4;
|
||||
if (ncols <= 4) {
|
||||
for (const IlPf& t : kIlPf_80)
|
||||
if (t.type == type) return t.v[ncols - 2][cls];
|
||||
} else {
|
||||
for (const IlPf58& t : kIlPf58_80)
|
||||
if (t.type == type) return t.v[ncols - 5][cls];
|
||||
}
|
||||
return 0;
|
||||
}
|
||||
int g_tune_rows = 0; // native_mmvq_il_tune (tests, benchmarks): rows a warp for every shape (0: the table)
|
||||
} // namespace
|
||||
void native_mmvq_il_tune(int rows) { g_tune_rows = rows; }
|
||||
void native_mmvq_il_pf(int pf) { g_il_pf = pf; }
|
||||
|
||||
namespace {
|
||||
// sm_80 and newer (measured on sm_86 and sm_120), Volta (sm_70) and Pascal (sm_60, sm_61), each with its own table
|
||||
@@ -3298,6 +3833,7 @@ void native_mmvq_il(int ggml_type, const void* weights, const void* x_q8_1, cons
|
||||
native_mmvq(ggml_type, weights, x_q8_1, y, n_in, n_out, ncols, stream);
|
||||
}
|
||||
void native_mmvq_il_tune(int) {}
|
||||
void native_mmvq_il_pf(int) {}
|
||||
#endif // !__HIPCC__
|
||||
|
||||
} // namespace strata::kernels
|
||||
@@ -41,6 +41,52 @@ template <int NC, bool PM> struct Q81Il {
|
||||
}
|
||||
}
|
||||
}
|
||||
// P1 (native_mmvq_il_pf_kernel): the same loads from a pointer the caller strength-reduces - qat(qbase(b, p) + k * QB)
|
||||
// is u(b + k, p), dat(dbase(b) + k * DB) is scales(b + k)
|
||||
static constexpr int QB = PM ? CP : 8 * CP, DB = CP;
|
||||
__device__ __forceinline__ int qp(int dp) const { return PM ? dp * nb * CP : dp * CP; } // ints from p to p + dp
|
||||
__device__ __forceinline__ const int* qbase(int b, int p) const { return q + std::size_t(PM ? p * nb + b : b * 8 + p) * CP; }
|
||||
__device__ __forceinline__ const float* dbase(int b) const { return d + std::size_t(b) * CP; }
|
||||
__device__ __forceinline__ void qat(const int* s, int (&o)[NC]) const { // __ldg: s comes from an asm (madw), so the
|
||||
if constexpr (CP == 2) { // front end no longer knows it is global
|
||||
const int2 a = __ldg(reinterpret_cast<const int2*>(s));
|
||||
o[0] = a.x;
|
||||
o[1] = a.y;
|
||||
} else {
|
||||
const int4 a = __ldg(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 = __ldg(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 dat(const float* s, float (&o)[NC]) const {
|
||||
if constexpr (CP == 2) {
|
||||
const float2 a = __ldg(reinterpret_cast<const float2*>(s));
|
||||
o[0] = a.x;
|
||||
o[1] = a.y;
|
||||
} else {
|
||||
const float4 a = __ldg(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 = __ldg(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;
|
||||
}
|
||||
}
|
||||
}
|
||||
__device__ __forceinline__ void scales(int b, float (&o)[NC]) const {
|
||||
const float* s = d + std::size_t(b) * CP;
|
||||
if constexpr (CP == 2) {
|
||||
@@ -73,6 +119,12 @@ struct Q81One {
|
||||
__device__ __forceinline__ void scales(int b, float (&o)[1]) const {
|
||||
o[0] = __low2float(*reinterpret_cast<const half2*>(x + std::size_t(b) * 36));
|
||||
}
|
||||
static constexpr int QB = 9, DB = 9; // a 36-byte block in ints / floats
|
||||
__device__ __forceinline__ int qp(int dp) const { return dp; }
|
||||
__device__ __forceinline__ const int* qbase(int b, int p) const { return reinterpret_cast<const int*>(x + std::size_t(b) * 36 + 4) + p; }
|
||||
__device__ __forceinline__ const float* dbase(int b) const { return reinterpret_cast<const float*>(x + std::size_t(b) * 36); }
|
||||
__device__ __forceinline__ void qat(const int* s, int (&o)[1]) const { o[0] = __ldg(s); }
|
||||
__device__ __forceinline__ void dat(const float* s, float (&o)[1]) const { o[0] = __low2float(__ldg(reinterpret_cast<const half2*>(s))); }
|
||||
};
|
||||
|
||||
// The activations a kernel for NC columns reads: one column's q8_1 blocks, 2+ columns' interleaved copy (nb blocks a
|
||||
|
||||
+117
-37
@@ -1,7 +1,13 @@
|
||||
// src/kernels/mmvq_il_parity.cpp - native_mmvq_il (2-4 columns from the interleaved q8_1 copy, fork F4) against
|
||||
// native_mmvq's multi-column kernels, bitwise, for every rows-a-warp choice; --bench times both per call.
|
||||
//
|
||||
// build/mmvq_il_parity [--bench [--emit-table] [--hot] [--gap MB] [--persist] [--ab]]
|
||||
// build/mmvq_il_parity [--bench [--emit-table] [--hot] [--gap MB] [--persist] [--ab] [--pf a,b,..] [--rounds N]]
|
||||
//
|
||||
// --pf a,b,..: the P1 / P2 kernel variants (native_mmvq_il_pf: 0 = today's kernel, 6 hoisted group loads, 3 + P2's exact
|
||||
// integer forms, 4 = 3 at 5 blocks a SM, 2 / 5 = 3 / 4 with a ring) are each bit-checked against native_mmvq, and with
|
||||
// --bench timed interleaved per (case, T, rows): N rounds
|
||||
// (default 5) of 100 calls each, variant after variant, the median of the rounds printed per variant (PFBENCH lines; the
|
||||
// table cells take the first listed variant). Without --pf every call takes STRATA_MMVQ_IL_PF (unset: today's kernel).
|
||||
//
|
||||
// --hot times one copy of the weights (they sit in L2 after the first call: the ceiling of an L2
|
||||
// prefetch); --gap MB reads MB megabytes of OTHER weights (native_mmvq over other copies) before every timed call, as
|
||||
@@ -23,6 +29,7 @@
|
||||
#include <cuda_runtime.h>
|
||||
|
||||
#include <algorithm>
|
||||
#include <array>
|
||||
#include <cmath>
|
||||
#include <cstdint>
|
||||
#include <cstdio>
|
||||
@@ -35,9 +42,11 @@
|
||||
namespace {
|
||||
using namespace strata::kernels;
|
||||
|
||||
struct Cell { int type, T, n_out; double ref; double il[5]; };
|
||||
struct Cell { int type, T, n_out, pf; double ref; double il[5]; }; // pf: the variant the il[] times ran (-1: the environment's)
|
||||
std::vector<Cell> g_cells;
|
||||
bool g_hot = false, g_persist = false, g_ab = false;
|
||||
std::vector<int> g_pf; // --pf: the variants to check and time; empty: the environment's
|
||||
int g_rounds = 5;
|
||||
std::size_t g_gap_bytes = 0;
|
||||
std::size_t g_gap_now = 0; // the gap in force for the current arm (--ab) or g_gap_bytes
|
||||
// time `reps` calls of `fn` on `s`; with --gap each call is preceded by `gap` (other weights' reads) and timed on its own
|
||||
@@ -195,6 +204,7 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
|
||||
std::vector<uint32_t> a((std::size_t) T * c.n_out), b(a.size());
|
||||
ck(cudaMemcpy(a.data(), yref, ybytes, cudaMemcpyDeviceToHost), "read ref");
|
||||
double best_ref = 0, best_il[5] = {0, 0, 0, 0, 0};
|
||||
std::vector<std::array<double, 5>> pf_il(g_pf.size(), std::array<double, 5>{0, 0, 0, 0, 0}); // per --pf variant
|
||||
int table_r = 0;
|
||||
struct Arm { const char* name; bool hot; std::size_t gap; bool persist; };
|
||||
const Arm arms_ab[] = {{"cold", false, 0, false}, {"hot", true, 0, false}, {"hot+gap", true, g_gap_bytes, false},
|
||||
@@ -208,34 +218,60 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
|
||||
g_gap_now = arms[ai].gap;
|
||||
if (arms[ai].persist != g_persist) { g_persist = arms[ai].persist; persist_window(s, dw[0], wbytes, g_persist); }
|
||||
}
|
||||
const std::vector<int> pfs = g_pf.empty() ? std::vector<int>{-1} : g_pf;
|
||||
for (int r : {0, 1, 2, 4}) {
|
||||
native_mmvq_il_tune(r);
|
||||
cudaMemset(yil, 0xff, ybytes);
|
||||
native_mmvq_il(c.type, dw[0], xq, xil, yil, c.n_in, c.n_out, T, s);
|
||||
ck(cudaStreamSynchronize(s), "il");
|
||||
ck(cudaMemcpy(b.data(), yil, ybytes, cudaMemcpyDeviceToHost), "read il");
|
||||
long long diff = 0, nonfinite = 0;
|
||||
for (std::size_t i = 0; i < a.size(); ++i) {
|
||||
float fa;
|
||||
std::memcpy(&fa, &a[i], 4);
|
||||
if (!std::isfinite(fa)) ++nonfinite;
|
||||
if (a[i] != b[i]) ++diff;
|
||||
for (int pf : pfs) {
|
||||
native_mmvq_il_pf(pf);
|
||||
cudaMemset(yil, 0xff, ybytes);
|
||||
native_mmvq_il(c.type, dw[0], xq, xil, yil, c.n_in, c.n_out, T, s);
|
||||
ck(cudaStreamSynchronize(s), "il");
|
||||
ck(cudaMemcpy(b.data(), yil, ybytes, cudaMemcpyDeviceToHost), "read il");
|
||||
long long diff = 0, nonfinite = 0;
|
||||
for (std::size_t i = 0; i < a.size(); ++i) {
|
||||
float fa;
|
||||
std::memcpy(&fa, &a[i], 4);
|
||||
if (!std::isfinite(fa)) ++nonfinite;
|
||||
if (a[i] != b[i]) ++diff;
|
||||
}
|
||||
if (diff != 0 || nonfinite != 0) {
|
||||
++bad;
|
||||
std::printf("FAIL %-26s T=%d rows=%d pf=%d: %lld bits differ, %lld non-finite of %zu\n", c.name, T, r, pf, diff,
|
||||
nonfinite, a.size());
|
||||
}
|
||||
}
|
||||
const bool used = r != 0 || native_mmvq_il_supported(c.type, T, c.n_out);
|
||||
if (diff != 0 || nonfinite != 0) {
|
||||
++bad;
|
||||
std::printf("FAIL %-26s T=%d rows=%d: %lld bits differ, %lld non-finite of %zu\n", c.name, T, r, diff, nonfinite, a.size());
|
||||
}
|
||||
(void) used;
|
||||
native_mmvq_il_pf(-1);
|
||||
if (bench && r != 0) {
|
||||
best_il[r] = time_us([&](int i) { native_mmvq_il(c.type, dw[g_hot ? 0 : i % copies], xq, xil, yil, c.n_in, c.n_out, T, s); }, 200, s, gap);
|
||||
auto il_call = [&](int i) { native_mmvq_il(c.type, dw[g_hot ? 0 : i % copies], xq, xil, yil, c.n_in, c.n_out, T, s); };
|
||||
if (g_pf.empty()) {
|
||||
best_il[r] = time_us(il_call, 200, s, gap);
|
||||
} else { // the variants interleaved, round after round; the median of each one's rounds
|
||||
std::vector<std::vector<double>> t(g_pf.size());
|
||||
for (int round = 0; round < g_rounds; ++round)
|
||||
for (std::size_t k = 0; k < g_pf.size(); ++k) {
|
||||
native_mmvq_il_pf(g_pf[k]);
|
||||
t[k].push_back(time_us(il_call, 100, s, gap));
|
||||
}
|
||||
native_mmvq_il_pf(-1);
|
||||
std::printf("PFBENCH %-26s T=%d rows=%d:", c.name, T, r);
|
||||
for (std::size_t k = 0; k < g_pf.size(); ++k) {
|
||||
std::sort(t[k].begin(), t[k].end());
|
||||
const double med = t[k][t[k].size() / 2];
|
||||
if (k == 0) best_il[r] = med;
|
||||
pf_il[k][r] = med;
|
||||
std::printf(" pf%d %.1f us (min %.1f max %.1f)", g_pf[k], med, t[k].front(), t[k].back());
|
||||
}
|
||||
std::printf("\n");
|
||||
}
|
||||
}
|
||||
}
|
||||
native_mmvq_il_tune(0);
|
||||
if (bench) {
|
||||
best_ref = time_us([&](int i) { native_mmvq(c.type, dw[g_hot ? 0 : i % copies], xq, yref, c.n_in, c.n_out, T, s); }, 200, s, gap);
|
||||
(void) table_r;
|
||||
g_cells.push_back({c.type, T, c.n_out, best_ref, {0, best_il[1], best_il[2], 0, best_il[4]}});
|
||||
if (g_pf.empty()) g_cells.push_back({c.type, T, c.n_out, -1, best_ref, {0, best_il[1], best_il[2], 0, best_il[4]}});
|
||||
for (std::size_t k = 0; k < g_pf.size(); ++k)
|
||||
g_cells.push_back({c.type, T, c.n_out, g_pf[k], best_ref, {0, pf_il[k][1], pf_il[k][2], 0, pf_il[k][4]}});
|
||||
std::printf("BENCH %-26s T=%d %-16s multi %.1f us | il rows1 %.1f rows2 %.1f rows4 %.1f us\n", c.name, T,
|
||||
arms[ai].name, best_ref, best_il[1], best_il[2], best_il[4]);
|
||||
}
|
||||
@@ -256,33 +292,37 @@ void emit_table() {
|
||||
std::printf("// measured on this card by mmvq_il_parity --bench --emit-table\n");
|
||||
// per cell: the rows count whose WORST ratio to native_mmvq over the cell's measured shapes is the best, taken only
|
||||
// when that worst ratio is under 0.97 (else 0); a class with no measured shape takes its lower neighbour's value
|
||||
auto pick_cell = [&](int type, int nc, int cls, int prev) {
|
||||
// pf_sel: the cells of that variant (the environment's when --pf was not given, else 0 = today's kernel); -2: every
|
||||
// listed variant, the answer then pf << 4 | rows (the kIlPf form), 0 when nothing beats native_mmvq by 3%
|
||||
auto pick_cell = [&](int type, int nc, int cls, int prev, int pf_sel) {
|
||||
double best_worst = 1e9;
|
||||
int best = -1;
|
||||
bool any = false;
|
||||
for (int r : {1, 2, 4}) {
|
||||
double worst = 0;
|
||||
bool have = false;
|
||||
for (const Cell& c : g_cells) {
|
||||
if (c.type != type || c.T != nc) continue;
|
||||
const int cc = c.n_out < 2048 ? 0 : c.n_out < 4096 ? 1 : c.n_out < 8192 ? 2 : c.n_out < 12288 ? 3 : 4;
|
||||
if (cc != cls || c.il[r] <= 0) continue;
|
||||
have = true;
|
||||
worst = std::max(worst, c.il[r] / c.ref);
|
||||
for (int pf : (pf_sel == -2 ? (g_pf.empty() ? std::vector<int>{-1} : g_pf) : std::vector<int>{pf_sel}))
|
||||
for (int r : {1, 2, 4}) {
|
||||
double worst = 0;
|
||||
bool have = false;
|
||||
for (const Cell& c : g_cells) {
|
||||
if (c.type != type || c.T != nc || c.pf != pf) continue;
|
||||
const int cc = c.n_out < 2048 ? 0 : c.n_out < 4096 ? 1 : c.n_out < 8192 ? 2 : c.n_out < 12288 ? 3 : 4;
|
||||
if (cc != cls || c.il[r] <= 0) continue;
|
||||
have = true;
|
||||
worst = std::max(worst, c.il[r] / c.ref);
|
||||
}
|
||||
any = any || have;
|
||||
if (have && worst < best_worst) { best_worst = worst; best = pf_sel == -2 ? (std::max(pf, 0) << 4 | r) : r; }
|
||||
}
|
||||
any = any || have;
|
||||
if (have && worst < best_worst) { best_worst = worst; best = r; }
|
||||
}
|
||||
if (!any) return prev;
|
||||
return best_worst < 0.97 ? best : 0;
|
||||
};
|
||||
const int pf_today = g_pf.empty() ? -1 : 0;
|
||||
for (int ti = 0; ti < 5; ++ti) {
|
||||
std::printf(" {%d, {", types[ti]);
|
||||
for (int nc = 2; nc <= 4; ++nc) {
|
||||
int prev = 0;
|
||||
std::printf("{");
|
||||
for (int cls = 0; cls < 5; ++cls) {
|
||||
prev = pick_cell(types[ti], nc, cls, prev);
|
||||
prev = pick_cell(types[ti], nc, cls, prev, pf_today);
|
||||
std::printf("%d%s", prev, cls < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}%s", nc < 4 ? ", " : "");
|
||||
@@ -296,19 +336,49 @@ void emit_table() {
|
||||
int prev = 0;
|
||||
std::printf("{");
|
||||
for (int cls = 0; cls < 5; ++cls) {
|
||||
prev = pick_cell(types[ti], nc, cls, prev);
|
||||
prev = pick_cell(types[ti], nc, cls, prev, pf_today);
|
||||
std::printf("%d%s", prev, cls < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}%s", nc < 8 ? ", " : "");
|
||||
}
|
||||
std::printf("}}, // %s\n", names[ti]);
|
||||
}
|
||||
if (!g_pf.empty()) { // the joint (variant, rows) table over the listed variants: kIlPf_80 / kIlPf58_80
|
||||
std::printf("// variant << 4 | rows over --pf (kIlPf_<cc>), 2-4 columns:\n");
|
||||
for (int ti = 0; ti < 5; ++ti) {
|
||||
std::printf(" {%d, {", types[ti]);
|
||||
for (int nc = 2; nc <= 4; ++nc) {
|
||||
int prev = 0;
|
||||
std::printf("{");
|
||||
for (int cls = 0; cls < 5; ++cls) {
|
||||
prev = pick_cell(types[ti], nc, cls, prev, -2);
|
||||
std::printf("0x%02x%s", prev, cls < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}%s", nc < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}}, // %s\n", names[ti]);
|
||||
}
|
||||
std::printf("// 5-8 columns (kIlPf58_<cc>):\n");
|
||||
for (int ti = 0; ti < 5; ++ti) {
|
||||
std::printf(" {%d, {", types[ti]);
|
||||
for (int nc = 5; nc <= 8; ++nc) {
|
||||
int prev = 0;
|
||||
std::printf("{");
|
||||
for (int cls = 0; cls < 5; ++cls) {
|
||||
prev = pick_cell(types[ti], nc, cls, prev, -2);
|
||||
std::printf("0x%02x%s", prev, cls < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}%s", nc < 8 ? ", " : "");
|
||||
}
|
||||
std::printf("}}, // %s\n", names[ti]);
|
||||
}
|
||||
}
|
||||
std::printf("// one column (kIl1RowsShared):\n");
|
||||
for (int ti = 0; ti < 5; ++ti) {
|
||||
std::printf(" {%d, {", types[ti]);
|
||||
int prev = 0;
|
||||
for (int cls = 0; cls < 5; ++cls) {
|
||||
prev = pick_cell(types[ti], 1, cls, prev);
|
||||
prev = pick_cell(types[ti], 1, cls, prev, pf_today);
|
||||
std::printf("%d%s", prev, cls < 4 ? ", " : "");
|
||||
}
|
||||
std::printf("}}, // %s\n", names[ti]);
|
||||
@@ -354,12 +424,22 @@ int check_tables() {
|
||||
int main(int argc, char** argv) {
|
||||
const bool bench = argc > 1 && std::string(argv[1]) == "--bench";
|
||||
const bool emit = bench && argc > 2 && std::string(argv[2]) == "--emit-table";
|
||||
for (int i = 2; i < argc; ++i) {
|
||||
for (int i = 1; i < argc; ++i) {
|
||||
const std::string a = argv[i];
|
||||
if (a == "--hot") g_hot = true;
|
||||
else if (a == "--ab") g_ab = true;
|
||||
else if (a == "--persist") g_persist = true;
|
||||
else if (a == "--gap" && i + 1 < argc) g_gap_bytes = (std::size_t) std::atoi(argv[++i]) << 20;
|
||||
else if (a == "--rounds" && i + 1 < argc) g_rounds = std::atoi(argv[++i]);
|
||||
else if (a == "--pf" && i + 1 < argc) {
|
||||
const std::string list = argv[++i];
|
||||
for (std::size_t p = 0; p < list.size();) {
|
||||
std::size_t q = list.find(',', p);
|
||||
if (q == std::string::npos) q = list.size();
|
||||
g_pf.push_back(std::atoi(list.substr(p, q - p).c_str()));
|
||||
p = q + 1;
|
||||
}
|
||||
}
|
||||
}
|
||||
cudaStream_t s;
|
||||
cudaStreamCreate(&s);
|
||||
|
||||
Reference in new issue
Block a user