Niko1221/Strata
1
0 Fork 0
mirror of https://github.com/Niko1221/Strata.git synced 2026-11-07 14:21:06 -10:00
Code Issues Packages Projects Releases 1 Wiki Activity

mmvq_il_parity: the L2-residency arms (--hot / --gap / --persist / --ab), the clock warm-up and MMVQ_IL_TMIN

Browse Source 
A go/no-go harness for an L2 weight prefetch, with its measurement on the CMP 170HX (GA100,
70 SMs, 32 MiB L2), T = 3.  Harness only; no engine code changes.

mmvq_il_parity --bench gains --hot (one weight copy: the L2-hot ceiling of a prefetch), --gap MB (MB of OTHER
weights read between the timed calls, as the window does between a prefetch and its consumer), --persist (a 20 MiB
persisting L2 set-aside + an access-policy window on the stream) and --ab (every arm per case inside ONE process:
cold / hot / hot+gap / hot+gap+persist / cold - the clocks ramp over the first ~100 ms of a run, so arms in separate
runs are not comparable; a memset gap was tried first and its write-backs made every arm slower than cold), plus a
clock warm-up and MMVQ_IL_TMIN.  Engine row choices (kIlRows_80), cold -> hot, us:
  z   2560x6144: IQ4_XS 16.0 -> 14.7 (-8 %), Q4_K 20.8 -> 17.7 (-15 %), Q5_K 22.3 -> 18.9 (-15 %), Q6_K 27.9 -> 20.7 (-26 %)
  out 6144x2560: IQ4_XS 16.1 -> 13.9 (-14 %), Q4_K 21.9 -> 19.2 (-12 %, multi), Q5_K 23.6 -> 20.5 (-13 %, multi), Q6_K 25.1 -> 20.4 (-19 %)
  ncu (caches kept vs flushed): L2 sector hit rate 100 / 93 / 96 % hot vs 24 / 21 / 16 % cold (Q6_K out, Q6_K z,
  IQ4_XS z), but at ncu's locked clock the duration moves only -4.5 / -10 / -3 %: the IL kernel is issue-bound once
  the lines are in L2 (sm__throughput 47-51 %).
  With 12, 16, 24 or 40 MB of other weights read between the prefetch and the consumer, plain LRU keeps NOTHING
  (hot+gap = cold or 1-2 us worse for every type at every gap), and the 20 MiB set-aside keeps a fraction: Q6_K z
  27.9 -> 26.0, Q6_K out 25.1 -> 24.6, Q4_K / Q5_K unchanged, IQ4_XS z 16.0 -> 17.6 (slower).  So an L2 prefetch
  (z forked 150 us ahead behind the experts' 58 MB, out 25-60 us ahead behind z's 9-13 MB) cannot meet its -15 %
  floor on this card: no-go.

gfx906 (the harness builds there with tests): persist_window's L2 access policy window and the --persist / --ab
set-aside query sit under !STRATA_HIP_GFX906; the harness prints "not available in this build".  The new timing
code ignores the runtime calls' results explicitly ((void)), so the harness adds no new nodiscard warnings to
the gfx906 build.

Co-Authored-By: Claude Opus 5.5 <[email protected]>
This commit is contained in:
linxuhaoandClaude Opus 5.5 committed 2026-10-22 20:53:14 +00:00
1 parent 68c487c32f
commit 1f4b81398b
1 file changed
+122 -33
Download Patch File Download Diff File Expand all files Collapse all files
src/kernels/mmvq_il_parity.cpp
+122 -33
View File
@@ -1,7 +1,15 @@
// 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]]
// build/mmvq_il_parity [--bench [--emit-table] [--hot] [--gap MB] [--persist] [--ab]]
//
// --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
// the window's GEMVs do between a prefetch and its consumer (a memset gap was tried first: its write-backs drain
// during the timed call and make every arm slower than cold); --persist puts the timed weights in a persisting L2
// set-aside (cudaLimitPersistingL2CacheSize + an access-policy window on the stream) so the gap's reads stream past
// them. --ab runs the arms cold / hot / hot+gap / hot+gap+persist / cold per case in ONE process (the clocks ramp
// over the first ~100 ms of a run, so arms in separate runs are not comparable). MMVQ_IL_TMIN / MMVQ_IL_TMAX bound T.
//
// --emit-table (with --bench) prints the rows-a-warp table this card measured, in the source form of kIlRows in
// native_mmvq.cu: per cell the rows count (1/2/4) that takes at least 3% off native_mmvq's time, else 0.
@@ -29,6 +37,58 @@ using namespace strata::kernels;
struct Cell { int type, T, n_out; double ref; double il[5]; };
std::vector<Cell> g_cells;
bool g_hot = false, g_persist = false, g_ab = false;
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
template <class F, class G>
double time_us(F fn, int reps, cudaStream_t s, G gap) {
cudaEvent_t e0, e1;
(void) cudaEventCreate(&e0);
(void) cudaEventCreate(&e1);
float total = 0;
for (int rep = 0; rep < 2; ++rep) { // rep 0 warms, rep 1 times
total = 0;
if (g_gap_now == 0) {
(void) cudaEventRecord(e0, s);
for (int i = 0; i < reps; ++i) fn(i);
(void) cudaEventRecord(e1, s);
(void) cudaEventSynchronize(e1);
(void) cudaEventElapsedTime(&total, e0, e1);
} else {
for (int i = 0; i < reps; ++i) {
gap();
(void) cudaEventRecord(e0, s);
fn(i);
(void) cudaEventRecord(e1, s);
(void) cudaEventSynchronize(e1);
float ms = 0;
(void) cudaEventElapsedTime(&ms, e0, e1);
total += ms;
}
}
}
(void) cudaEventDestroy(e0);
(void) cudaEventDestroy(e1);
return 1000.0 * total / reps;
}
// --persist: the matrix `w` (wbytes) persisting in L2 for every kernel on `s`; off: the window cleared
void persist_window(cudaStream_t s, void* w, std::size_t wbytes, bool on) {
#if !defined(STRATA_HIP_GFX906)
cudaStreamAttrValue v{};
v.accessPolicyWindow.base_ptr = on ? w : nullptr;
v.accessPolicyWindow.num_bytes = on ? wbytes : 0;
v.accessPolicyWindow.hitRatio = 1.0f;
v.accessPolicyWindow.hitProp = cudaAccessPropertyPersisting;
v.accessPolicyWindow.missProp = cudaAccessPropertyStreaming;
if (cudaStreamSetAttribute(s, cudaStreamAttributeAccessPolicyWindow, &v) != cudaSuccess) {
std::printf("--persist: the access policy window was refused: %s\n", cudaGetErrorString(cudaGetLastError()));
}
#else
(void) s; (void) w; (void) wbytes; // (no L2 access policy window in the gfx906 compat layer)
std::printf("--persist: not available in this build (%s)\n", on ? "on" : "off");
#endif
}
struct Case {
const char* name;
int type, n_in, n_out, block_elems, block_bytes, scale_at[2];
@@ -98,15 +158,20 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
const uint16_t h = sane_half(rng);
std::memcpy(&w[k * (std::size_t) c.block_bytes + (std::size_t) at], &h, 2);
}
const int copies = bench ? (int) std::max<std::size_t>(1, std::min<std::size_t>(24, (192u << 20) / wbytes)) : 1;
// --hot: one copy (timed hot); with --gap, 1 + the copies the gap reads (dw[1..]) - the timed call always reads dw[0]
const int gap_copies = g_gap_bytes > 0 ? (int) std::min<std::size_t>(23, (g_gap_bytes + wbytes - 1) / wbytes) : 0;
const int copies = bench && (!g_hot || g_ab) ? (int) std::max<std::size_t>(1, std::min<std::size_t>(24, (192u << 20) / wbytes)) : 1 + gap_copies;
std::vector<void*> dw(copies, nullptr);
for (int i = 0; i < copies; ++i) {
if (!ck(cudaMalloc(&dw[i], wbytes), "malloc w") || !ck(cudaMemcpy(dw[i], w.data(), wbytes, cudaMemcpyHostToDevice), "copy w"))
return 1;
}
int bad = 0;
if (bench && g_persist && !g_ab) persist_window(s, dw[0], wbytes, true);
g_gap_now = g_gap_bytes;
const int t_max = std::getenv("MMVQ_IL_TMAX") ? std::atoi(std::getenv("MMVQ_IL_TMAX")) : 8; // 147-verify: 5-8 columns
for (int T = 1; T <= t_max; ++T) {
const int t_min = std::getenv("MMVQ_IL_TMIN") ? std::atoi(std::getenv("MMVQ_IL_TMIN")) : 1;
for (int T = t_min; T <= t_max; ++T) {
std::vector<float> x((std::size_t) T * c.n_in);
std::normal_distribution<float> nd(0.f, 1.f);
for (auto& v : x) v = nd(rng);
@@ -124,10 +189,25 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
if (T >= 2) native_q8_1_interleave(xq, xil, c.n_in, T, s);
native_mmvq(c.type, dw[0], xq, yref, c.n_in, c.n_out, T, s);
ck(cudaStreamSynchronize(s), "ref");
auto gap = [&]() { // --gap: other weights' reads between the timed calls (the window's traffic)
for (int j = 1; j <= gap_copies; ++j) native_mmvq(c.type, dw[j], xq, yref, c.n_in, c.n_out, 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};
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},
{"hot+gap+persist", true, g_gap_bytes, true}, {"cold2", false, 0, false}};
const Arm arms_one[] = {{g_hot ? "hot" : "cold", g_hot, g_gap_bytes, g_persist}};
const Arm* arms = g_ab ? arms_ab : arms_one;
const int n_arms = g_ab ? 5 : 1;
for (int ai = 0; ai < (bench ? n_arms : 1); ++ai) {
if (bench) {
g_hot = arms[ai].hot;
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); }
}
for (int r : {0, 1, 2, 4}) {
native_mmvq_il_tune(r);
cudaMemset(yil, 0xff, ybytes);
@@ -148,45 +228,21 @@ int run_case(const Case& c, bool bench, cudaStream_t s) {
}
(void) used;
if (bench && r != 0) {
cudaEvent_t e0, e1;
cudaEventCreate(&e0);
cudaEventCreate(&e1);
const int reps = 200;
for (int rep = 0; rep < 2; ++rep) { // rep 0 warms, rep 1 times
cudaEventRecord(e0, s);
for (int i = 0; i < reps; ++i) native_mmvq_il(c.type, dw[i % copies], xq, xil, yil, c.n_in, c.n_out, T, s);
cudaEventRecord(e1, s);
cudaEventSynchronize(e1);
}
float ms;
cudaEventElapsedTime(&ms, e0, e1);
best_il[r] = 1000.0 * ms / reps;
cudaEventDestroy(e0);
cudaEventDestroy(e1);
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);
}
}
native_mmvq_il_tune(0);
if (bench) {
cudaEvent_t e0, e1;
cudaEventCreate(&e0);
cudaEventCreate(&e1);
const int reps = 200;
for (int rep = 0; rep < 2; ++rep) {
cudaEventRecord(e0, s);
for (int i = 0; i < reps; ++i) native_mmvq(c.type, dw[i % copies], xq, yref, c.n_in, c.n_out, T, s);
cudaEventRecord(e1, s);
cudaEventSynchronize(e1);
}
float ms;
cudaEventElapsedTime(&ms, e0, e1);
best_ref = 1000.0 * ms / reps;
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]}});
std::printf("BENCH %-26s T=%d multi %.1f us | il rows1 %.1f rows2 %.1f rows4 %.1f us\n", c.name, T, best_ref,
best_il[1], best_il[2], best_il[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]);
}
} // arms
cudaFree(dx); cudaFree(xq); cudaFree(xil); cudaFree(yref); cudaFree(yil);
}
if (bench && g_persist) persist_window(s, dw[0], wbytes, false);
for (void* p : dw) cudaFree(p);
return bad;
}
@@ -298,10 +354,43 @@ 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) {
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;
}
cudaStream_t s;
cudaStreamCreate(&s);
#if !defined(STRATA_HIP_GFX906)
if (g_persist || g_ab) {
int dev = 0, l2 = 0, maxp = 0, maxw = 0;
cudaGetDevice(&dev);
cudaDeviceGetAttribute(&l2, cudaDevAttrL2CacheSize, dev);
cudaDeviceGetAttribute(&maxp, cudaDevAttrMaxPersistingL2CacheSize, dev);
cudaDeviceGetAttribute(&maxw, cudaDevAttrMaxAccessPolicyWindowSize, dev);
std::size_t want = std::min<std::size_t>(20u << 20, (std::size_t) maxp), got = 0;
const cudaError_t e = cudaDeviceSetLimit(cudaLimitPersistingL2CacheSize, want);
cudaDeviceGetLimit(&got, cudaLimitPersistingL2CacheSize);
std::printf("--persist: L2 %d MiB, max persisting %d MiB, max window %d MiB; set-aside asked %zu MiB -> %s, granted %zu MiB\n",
l2 >> 20, maxp >> 20, maxw >> 20, want >> 20, cudaGetErrorString(e), got >> 20);
}
#endif
if (bench) std::printf("bench arms: %s%s%s\n", g_hot ? "hot (one copy) " : "cold (192 MB rotation) ",
g_gap_bytes ? ("gap " + std::to_string(g_gap_bytes >> 20) + " MB ").c_str() : "", g_persist ? "persist" : "");
int bad = 0;
bad += check_tables();
if (bench) { // ~150 ms of launches so the clocks are up before the first timed case
void* wu = nullptr; float* yu = nullptr; void* xu = nullptr;
if (cudaMalloc(&wu, 6144u * 2560u) == cudaSuccess && cudaMalloc((void**) &yu, 2560 * 4) == cudaSuccess &&
cudaMalloc(&xu, native_q8_1_bytes(6144, 1)) == cudaSuccess) {
(void) cudaMemset(wu, 0, 6144u * 2560u); (void) cudaMemset(xu, 0, native_q8_1_bytes(6144, 1));
for (int i = 0; i < 6000; ++i) native_mmvq(12, wu, xu, yu, 6144, 2560, 1, s);
(void) cudaStreamSynchronize(s);
(void) cudaFree(wu); (void) cudaFree(yu); (void) cudaFree(xu);
}
}
const char* only = std::getenv("MMVQ_IL_ONLY"); // 143-dense: only the cases whose name contains this
for (const Case& c : CASES)
if (only == nullptr || std::strstr(c.name, only) != nullptr) bad += run_case(c, bench, s);