diff --git a/CMakeLists.txt b/CMakeLists.txt index cd60517799..c1968ed849 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -250,6 +250,11 @@ if(STRATA_BUILD_TESTS) add_executable(tuning_header_arch_family_test tests/prefill/tuning_header_arch_family_test.cpp) target_include_directories(tuning_header_arch_family_test PRIVATE "${CMAKE_CURRENT_SOURCE_DIR}/src/prefill") add_test(NAME tuning_header_arch_family_test COMMAND tuning_header_arch_family_test) + + # the decode tuning table's parser and its refusals (plain C++: no GPU) + add_executable(decode_tuning_test tests/core/decode_tuning_test.cpp src/kernels/decode_tuning.cpp) + target_include_directories(decode_tuning_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/include) + add_test(NAME decode_tuning_test COMMAND decode_tuning_test) endif() # ---- strata_plan ---------------------------------------------------------- @@ -349,6 +354,7 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP) src/kernels/cuda/qsa_select.cu src/kernels/cuda/verify_kernels.cu src/kernels/cuda/iq_kernels.cu + src/kernels/decode_tuning.cpp src/kernels/ngram.cpp src/ngram/ple_reader.cpp src/platform/direct_file.cpp) @@ -528,6 +534,11 @@ if(STRATA_ENABLE_CUDA OR STRATA_ENABLE_HIP) target_link_libraries(tune_hipblaslt PRIVATE hip::host roc::hipblas roc::hipblaslt) endif() + if(STRATA_ENABLE_HIP) + # on-device tuning of the decode kernels' block shapes (docs/AMD_HIP_AUTOTUNE.md; tools/hip/autotune.py runs it) + add_executable(tune_decode tools/hip/tune_decode.cpp) + target_link_libraries(tune_decode PRIVATE strata_kernels strata_hip_runtime) + endif() add_executable(strata src/program/generate.cpp) target_include_directories(strata PRIVATE ${_strata_gpu_include_directories}) target_link_libraries(strata PRIVATE strata_engine strata_prefill strata_spec ${_strata_gpu_runtime_target}) diff --git a/docs/AMD_HIP.md b/docs/AMD_HIP.md index 5996177e1e..a08c6b5c53 100644 --- a/docs/AMD_HIP.md +++ b/docs/AMD_HIP.md @@ -50,7 +50,8 @@ the kernel's amdgpu driver (no ROCm install needed): - **Limits for now:** images only through the CPU encoder (`--vision cpu`, 0.1.32). Setup does not offer the tuning (calibration) on AMD yet: its controls are being checked on HIP one at a time (#566). Since 0.1.39 a tuning run by hand (`./setup.sh --calibrate`) is saved for the AMD card it ran on and reused when setup runs again. The Monitor - shows the card's load, VRAM, temperature and power from Linux sysfs (0.1.32). + shows the card's load, VRAM, temperature and power from Linux sysfs (0.1.32). On-device tuning of the +decode kernels (block shapes, draft and CPU settings) is available by hand: [AMD_HIP_AUTOTUNE.md](AMD_HIP_AUTOTUNE.md). The rest of setup is the same as on NVIDIA: the model download, the start script, the server. diff --git a/docs/AMD_HIP_AUTOTUNE.md b/docs/AMD_HIP_AUTOTUNE.md new file mode 100644 index 0000000000..5db588ae69 --- /dev/null +++ b/docs/AMD_HIP_AUTOTUNE.md @@ -0,0 +1,117 @@ +# On-device autotuning for the RX 7900 XTX (gfx1100) + +Strata's AMD backend runs CUDA-shaped kernels whose launch shapes were chosen on an RTX 5070. This adds +two measured layers of tuning for the 7900 XTX, in the spirit of engines that tune kernels on the user's +device instead of shipping one shape for a whole class of hardware. Both layers keep answers exactly the +same. + +| Layer | What is measured | Changes answers? | Cost | +|---|---|---|---| +| Kernels | Rows per thread block of the decode pass's GPU expert kernels (gate/up, down) and `iq_mmvq` | No: bitwise-equal variants only, checked on the machine and again end to end | ~1–3 min + 4 engine starts | +| Engine settings | Draft floor (`--spec-min-p`), CPU expert workers (`--pool-workers`) | No (greedy verification keeps the tokens) | 5–10 min | +| Drafting (opt-in) | MTP window (`--spec`), prompt-lookup match (`--suffix-draft`) | No | +6 engine starts | + +## Run it + +Stop the server first: the timings need the GPU idle and about 1 GB of free VRAM. Then: + +``` +python tools/hip/autotune.py strata-iq3_xxs.json # kernels only +python tools/hip/autotune.py strata-iq3_xxs.json --settings # + draft floor and CPU workers +python tools/hip/autotune.py strata-iq3_xxs.json --draft --prompts my-agent-prompts.json +python tools/hip/autotune.py strata-iq3_xxs.json --off # back to the default kernel shapes +``` + +`--prompts` takes a JSON list of strings, or plain text with prompts separated by a line `---`. Use requests your +agent really sends. Prompt lookup only pays where the answer repeats the context (edits, refactors, quoting), +so the built-in prompts are a weak proxy for an agent workload. + +The script needs the `tune_decode` tool. It builds it in `build-hip` when that build directory exists, or you +can build it yourself: + +``` +cmake --build build-hip --target tune_decode -j2 +``` + +When it finishes, the config's `env` carries `STRATA_DECODE_TUNING` (next to `STRATA_HIPBLASLT_TUNING`), and the +engine log says `decode tuning: N shapes from ...` at start. A report goes to `strata-.autotune.json`, and the +original config is kept once as `strata-.json.bak-autotune`. + +The settings (`--settings`, `--draft`) also work through setup: `./setup.sh --calibrate`, `--calibrate-draft`, +`--calibrate-prompts FILE`. On AMD the PCIe share is left exactly as the config has it, because the HIP expert +paths (mmap, resident CPU experts, the staging buffer) were measured with that share. `tools/calibrate.py +--sweep-pcie` sweeps it anyway. + +## How the kernel layer stays exact + +`native_expert_grouped` (the GPU experts in every verify window) and `iq_mmvq` give one warp to each weight row. +How many rows share a block changes occupancy and scheduling but never the arithmetic: each row is still reduced +by one warp, over the same lanes, in the same order. `tune_decode` checks this anyway: + +1. Every candidate's output is compared **byte for byte** with the engine default's - the real `native_expert_grouped` + / `iq_mmvq` call, no table, with its own launch heuristics - before it may be timed. A difference excludes the + candidate, fails the run (exit 4), and nothing is written. +2. The workload is the verify pass's: windows of 1, 3 and 5 tokens (weighted 0.2/0.3/0.5), `--hit-frac` of each + token's 10 experts on the GPU, and experts drawn from a 768 MB pool. The pool is far larger than the 96 MB + Infinity Cache, because decode streams around a gigabyte of experts per token and a cache-warm benchmark would lie. +3. Candidates are timed interleaved over several rounds, in a rotating order. A winner must beat the default by more + than 3% on the weighted median, then again in a second interleaved confirmation. +4. `autotune.py` then starts the engine without and with the table, twice each, interleaved. It keeps the table + only if the generated tokens are identical and decode is not slower. If two runs *without* the table already + differ (an adaptive expert tier, for example), the token check is reported as not checkable and speed decides. + +## What a table row means + +A row names the rows per block of the **plain one-warp-per-row kernels** (`native_gu_kernel`, `native_down_kernel`, +`mmvq_kernel`). The no-table default for split formats instead runs the multi decode-once kernels, whose rows per +block is their own - so `16` in a table and `16` in the default's launch geometry are different mechanisms, and the +tuner compares whole launches against the engine's default entry, not numbers against numbers. Two consequences: + +- A table row replaces the whole call. Kernels the table leaves at their default (no row, or the default value) run + the default dispatch chain, at the grid the parameter sizes - not necessarily the grid the no-table path would + have chosen. Outputs stay bitwise equal; only the block count can differ. +- The `mmvq` key is `(type, n_in, n_out)` with no token-count axis: the winner is a blend over the 1/3/5-token + verify windows, which is the decode workload. + +The table is refused at engine start, with a message, unless its header matches the running engine: the GPU +architecture, the HIP runtime version, a hash of the compiler that built the kernels, and the kernel selection +space (a version that bumps when the kernel set or the default dispatch changes - a table from an older engine is +stale and the tuner is run again). It is also refused when a dispatch rerouting variable is set +(`STRATA_OLD_IQ_MMVQ`, `STRATA_NO_SUB16_GU`, `STRATA_IQ_STAGE_GRID=0`, `STRATA_IQ_STAGE_GRID_MMVQ=1`, +`STRATA_EXPERT_V2`, `STRATA_EXPERT_V2K`, `STRATA_TSUM`, `STRATA_GROUPED_V1`): the table was measured against the +default dispatch, and a knob that reroutes it makes those measurements meaningless. CUDA builds compile only the +default shapes and ignore the variable. + +A note on the staging knobs, measured on the RX 7900 XTX: the defaults stage the i-quant codebook grids into +shared memory, which upstream chose on an RTX 5070. On this card the codebook already sits in the 96 MB Infinity +Cache, so staging is a wash for the IQ3-class formats (within 0.2% on every shape measured) and a small win when +OFF for IQ2_S gate/up (~2-5% with `STRATA_IQ_STAGE_GRID=0`); mmvq staging is a wash. Measured by the same tool, +byte-checked; if you run an IQ2_S-class model it is worth one paired run on your machine. + +## What to expect + +Honestly: unknown until it runs on your card. The tunable kernels are a slice of each token's time. In the +hybrid setup, CPU-side expert work and the GPU-reach wait are large too, and kernel shapes cannot touch those. +Measure where your time goes before and after: + +``` +"env": {"STRATA_DECODE_TIMING": "1", ...} +``` + +With this set, each request logs one line: windows, average window size, tokens per window, and ms per window +split into verify (GPU-reach wait, per-layer host and CPU expert time), commit and draft. It also logs VRAM hits, +CPU expert entries and PCIe experts per layer-window, followed by the GPU stage table. If "verify" is dominated by +the CPU expert part, the drafting and worker settings matter more than kernel shapes. + +## Drafting notes for Qwen3.8-Flash-Next + +- **MTP** is the model's own draft layer. Its window grows while the draft's confidence is at least + `--spec-min-p`. Every extra row in a window brings more experts the VRAM cache misses onto the CPU. On a + CPU-bound machine a higher floor or a shorter `--spec` often wins, so it is measured, not assumed. +- **Prompt lookup** (`--suffix-draft`) drafts the continuation of an earlier repeat of the context. A built-in + policy learns the measured round cost per window size and lookup acceptance per match length, and takes a lookup + window only when its expected tokens per millisecond beat the MTP's. It fires only when its first token agrees + with the MTP's first guess. +- **DFlash2** (block-diffusion drafters) exists for Qwen3.8-27B, not for Flash-Next. A drafter for Flash-Next would + have to be trained, and it would take VRAM from the expert cache (about 700 experts per GB). This is not part of + this tuning. diff --git a/include/strata/kernels/decode_tuning.hpp b/include/strata/kernels/decode_tuning.hpp new file mode 100644 index 0000000000..fe8a35e116 --- /dev/null +++ b/include/strata/kernels/decode_tuning.hpp @@ -0,0 +1,111 @@ +// include/strata/kernels/decode_tuning.hpp - on-device tuning of the decode kernels' launch shapes (gfx1100). +// +// The decode hot path's GPU experts (`native_expert_grouped`: gate/up, then down) and the i-quant matrix-vector +// kernel (`iq_mmvq`) give one warp to each weight row; how many rows share a thread block is a free choice. It +// changes occupancy and scheduling but not one floating-point operation: every row is still reduced by the same +// warp, over the same lanes, in the same order, so every variant writes bitwise the same output as the default. +// Which one is fastest depends on the GPU, the driver and the compiler, so it is measured on the machine +// (tools/hip/tune_decode.cpp writes the table, tools/hip/autotune.py runs the whole procedure) instead of guessed. +// +// The table (STRATA_DECODE_TUNING=) is plain text: +// +// # comments +// STRATA_DECODE_TUNING_V1 +// gu gate/up of the grouped experts (n_in = n_embd, n_out = n_ff) +// down down of the grouped experts (n_in = n_ff, n_out = n_embd) +// mmvq iq_mmvq +// +// The header must match the running engine exactly - the architecture (e.g. gfx1100), the HIP runtime, the +// compiler that built the kernels (the tuner checked bitwise equality with THAT build), and the kernel selection +// space (below) - or the table is refused with a message and the defaults stay. A shape the table does not list +// keeps its default. CUDA builds compile only the default shapes and ignore the variable (with a message), so +// their behaviour is unchanged. +// +// What a row names: the rows per block of the PLAIN one-warp-per-row kernels (mmvq_kernel / native_gu_kernel / +// native_down_kernel). The no-table default for split types instead runs the multi decode-once kernels, whose +// rows per block is their own - so the number 16 in a table and the number 16 in the default's launch geometry +// are different mechanisms. The tuner compares whole configurations against the engine's true default entry +// (native_expert_grouped / iq_mmvq with no table), byte-checks every candidate against it first, and a table row +// replaces the whole call's launch: the kernels the table leaves at their default run the default chain. +#pragma once + +#include +#include +#include +#include + +namespace strata::kernels { + +enum class DecodeKernel : int { GateUp = 0, Down = 1, Mmvq = 2 }; + +/// Defaults (the layouts before tuning existed): 8 rows per block for the grouped experts, 4 for iq_mmvq. +inline constexpr int kDefaultExpertRows = 8; +inline constexpr int kDefaultMmvqRows = 4; + +/// Which kernel selection space this build offers the tuner. The value rides in the table header; a build whose +/// space differs refuses the table ("run the tuner again"). Bump it when the plain kernels' arithmetic or the +/// default dispatch chain changes enough that an older table no longer describes this build - a compiler or +/// runtime bump alone does not change it, a kernel rework does (the sub-warp/EXACT_N rework is the case that +/// motivated the field: same compiler, same runtime, different kernels). +inline constexpr int kDecodeTuningSpace = 2; + +/// The rows-per-block variants HIP builds compile besides the default (one list: the parser's validity set, the +/// launchers' compiled-in set and the tuner's candidate list all derive from these two macros). CUDA builds +/// compile only the defaults. +#define STRATA_EXPERT_ROWS_VARIANTS(X) X(4) X(16) X(2) +#define STRATA_MMVQ_ROWS_VARIANTS(X) X(2) X(8) X(1) + +/// The rows-per-block values a kernel is compiled for (one warp of 32 lanes per row, at most 512 threads). +bool decode_rows_valid(DecodeKernel k, int rows) noexcept; +/// The candidate list the tuner measures for a kernel, default first. +std::vector decode_rows_candidates(DecodeKernel k); +const char* decode_kernel_name(DecodeKernel k) noexcept; +int decode_default_rows(DecodeKernel k) noexcept; + +struct DecodeTuningRow { + DecodeKernel kernel = DecodeKernel::GateUp; + int type = 0; + int64_t n_in = 0, n_out = 0; + int rows = 0; +}; + +/// Who the kernels were built and run by. A table is only valid for exactly this identity. +struct DecodeIdentity { + std::string arch; ///< e.g. "gfx1100" (feature suffixes such as ":xnack-" stripped); "cuda" on CUDA + long long runtime = 0; ///< hipRuntimeGetVersion / cudaRuntimeGetVersion + std::string toolchain; ///< 16 hex digits: a hash of the device compiler's version string + int space = 0; ///< kDecodeTuningSpace of the build the table was measured on +}; + +class DecodeTuningTable { +public: + /// Parse a table; refuses (returns false, `err` says why) a malformed line, an unknown kernel, a rows value + /// the kernel is not compiled for, a duplicate shape, or a header that does not match `want`. + bool parse(std::istream& in, const DecodeIdentity& want, std::string& err); + bool load(const std::string& path, const DecodeIdentity& want, std::string& err); + /// Rows per block for this shape, or `fallback` when the table does not list it. + int rows(DecodeKernel k, int type, int64_t n_in, int64_t n_out, int fallback) const noexcept; + const std::vector& entries() const noexcept { return rows_; } + bool empty() const noexcept { return rows_.empty(); } + + static std::string header(const DecodeIdentity& id); + static std::string line(const DecodeTuningRow& r); + +private: + std::vector rows_; +}; + +/// A stable 16-hex-digit FNV-1a hash of `s` (the toolchain id). +std::string decode_toolchain_hash(const std::string& s); + +// ---- the running engine (implemented in the GPU translation unit) ---- + +/// This process's identity: the current device's architecture, the runtime version and the kernels' compiler. +DecodeIdentity decode_identity(); +/// Load STRATA_DECODE_TUNING once (later calls are no-ops) and say on stderr what was used. Call it before any +/// graph capture; captured graphs keep the kernels they captured. Safe to skip: the first launch loads it lazily. +void decode_tuning_init(); +/// The active table, or nullptr (none set, refused, or a CUDA build). +const DecodeTuningTable* decode_tuning_active(); + +} // namespace strata::kernels diff --git a/include/strata/kernels/iq_kernels.hpp b/include/strata/kernels/iq_kernels.hpp index 2c38f6f98d..1e18590c4d 100644 --- a/include/strata/kernels/iq_kernels.hpp +++ b/include/strata/kernels/iq_kernels.hpp @@ -23,7 +23,12 @@ size_t iq_row_bytes(int ggml_type, int64_t n) noexcept; void quantize_q8_1_rows(const float* x, int64_t n_rows, int64_t n_cols, void* y, void* stream); /// y[c][r] = W[r] . x[c] for `ncols` columns of q8_1 activations (x stride n_in/32 blocks per column). +/// Rows per thread block: the decode tuning table's choice for this shape (decode_tuning.hpp), else 4. 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); +/// The same with an explicit rows-per-block (the tuner's entry; bitwise equal to iq_mmvq for every value +/// `decode_rows_valid(DecodeKernel::Mmvq, rows)` accepts; CUDA builds compile only 4). +void iq_mmvq_rows(int ggml_type, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, int rows, + 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); @@ -76,5 +81,12 @@ void native_expert_set_mode(int mode, int phase); /// bitwise the same results. Set before graph capture; captured graphs keep the kernels they captured. void iq_set_old_kernels(bool old); bool iq_old_kernels(); +/// The same with explicit rows per block for the gate/up and the down kernel (the tuner's entry; bitwise equal to +/// native_expert_grouped for every value decode_rows_valid accepts; CUDA builds compile only 8). The default entry +/// above reads the decode tuning table (decode_tuning.hpp) and routes here when it names a non-default rows value. +void native_expert_grouped_rows(const NativeExpertLayout& L, const unsigned long long* grp_ptr, + const int32_t* grp_start, const int32_t* n_groups, const int32_t* ent_dst, + const int32_t* ent_tok, int64_t cap_groups, int64_t cap_entries, const void* x_q8_1, + void* scratch, float* out, int gu_rows, int down_rows, void* stream); } // namespace strata::kernels diff --git a/setup.py b/setup.py index ee657ce7d2..26a4c7abae 100644 --- a/setup.py +++ b/setup.py @@ -3384,32 +3384,39 @@ def hardware_key(cfg: dict) -> str: cfg.get("model_name", "?"), ctx, "images" if "--vision" in a else "text"]) -def calibrate_config(cfg_path: Path) -> bool: +def calibrate_config(cfg_path: Path, draft: bool = False, prompts: str | None = None) -> bool: """Measure the engine's hardware-dependent settings on this PC (tools/calibrate.py), write them into the run - config and remember them per PC and model in the settings file, so an update or a reinstall keeps them.""" + config and remember them per PC and model in the settings file, so an update or a reinstall keeps them. + `draft`: also the drafting settings (the MTP window, prompt lookup); `prompts`: a file of your own prompts.""" sys.path.insert(0, str(ROOT / "tools")) import calibrate as CAL cfg = json.loads(cfg_path.read_text(encoding="utf-8-sig")) + hip = cfg.get("backend") == "hip" say() - say(" Tuning Strata for this PC: the output speed is measured with a few engine settings (the PCIe share, the") - say(" draft depth, the CPU threads). It takes about 5-10 minutes; the PC is busy meanwhile.") + say(" Tuning Strata for this PC: the output speed is measured with a few engine settings (" + + ("the draft floor" if hip else "the PCIe share, the draft depth") + ", the CPU threads" + + (", the MTP window, prompt lookup" if draft else "") + ").") + say(f" It takes about {'15-30' if draft else '5-10'} minutes; the PC is busy meanwhile.") try: since = os.path.getsize(cfg["log"]) if cfg.get("log") and os.path.isfile(cfg["log"]) else 0 except OSError: since = 0 try: - res = CAL.run(cfg, say=say) + extra = {"draft": True} if draft else {} # only what was asked for: older stand-ins take neither + if prompts: + extra["prompts"] = CAL.load_prompts(prompts) + res = CAL.run(cfg, say=say, **extra) except Exception as e: # never stops an install: the defaults stay warn(f"the tuning did not finish ({e}): the default settings stay") why = CAL.engine_error(cfg.get("log"), since) # #447: the engine's own reason, not only "see the log" if why: say(f" the engine said: {why}") return False - cfg["args"] = CAL.apply(cfg["args"], res["settings"]) + cfg["args"] = CAL.apply(cfg["args"], res["settings"], backend=cfg.get("backend"), draft=draft) write_config(cfg_path, cfg) st = load_settings() st.setdefault("calibration", {})[hardware_key(cfg)] = {"settings": res["settings"], "tok_s": res["report"].get("tok_s"), - "date": time.strftime("%Y-%m-%d")} + "draft": bool(draft), "date": time.strftime("%Y-%m-%d")} save_settings(st) if res["settings"]: ok("tuned for this PC: " + ", ".join(f"{k} {v}" for k, v in res["settings"].items()) @@ -3428,8 +3435,8 @@ def saved_calibration(cfg: dict) -> dict | None: def setup_calibration(cfg: dict, hip: bool) -> dict | None: """The calibration a (re-)install applies to its new config: the one saved for this PC and model. #566: on Linux HIP too - hardware_key now names the AMD cards, so a `./setup.sh --calibrate` run is matched to its card and - model. Windows HIP keeps the defaults for now (not tried there). Setup still offers the tuning itself on NVIDIA - only: each control is verified on HIP first.""" + model. Windows HIP keeps the defaults for now (not tried there). Setup offers the tuning itself on NVIDIA and + Linux HIP (#744: the HIP run keeps the config's PCIe share); on Windows HIP it is still not offered.""" if hip and WIN: return None return saved_calibration(cfg) @@ -4093,6 +4100,12 @@ def main() -> int: ap.add_argument("--check", action="store_true", help="only check this PC and exit") ap.add_argument("--calibrate", action="store_true", help="tune the engine's settings for this PC (about 5-10 minutes), then start the model") + ap.add_argument("--calibrate-draft", action="store_true", + help="--calibrate, and also the drafting settings: the MTP window and prompt lookup (a restart " + "per value; about 15-30 minutes)") + ap.add_argument("--calibrate-prompts", metavar="FILE", + help="with --calibrate: measure with your own prompts (a JSON list of strings, or text blocks " + "separated by a line ---), e.g. requests your coding agent sends") ap.add_argument("--draft-vocab", choices=list(DRAFT_VOCABS), help="the draft layer's tokens: cjk = with Chinese, Japanese and Korean (default), en = English " "and code only (~110 MiB less VRAM, English answers 1-2%% faster), cyrillic = English, code " @@ -4181,6 +4194,8 @@ def main() -> int: # choice, and asked once when the PC has cards that could share the model run_gpu = start_gpus(a.gpus) or a.gpu port = a.port or 8080 # a new install's port (issue #32: --port for an existing one) + if a.calibrate_draft or a.calibrate_prompts: + a.calibrate = True if have and a.calibrate and not (a.setup or a.model or a.family or a.check): if not a.build: update_installed_engine(a.prebuilt) @@ -4190,7 +4205,8 @@ def main() -> int: for i, c in enumerate(have, 1): say(f" {i}) {json.loads(c.read_text(encoding='utf-8-sig')).get('model_name', c.stem)}") pick_cfg = have[int(ask("Tune which one?", [str(i) for i in range(1, len(have) + 1)], "1", a.yes)) - 1] - if not calibrate_config(pick_cfg): # #447: said again where it is not lost above the start + tuned_ok = calibrate_config(pick_cfg, draft=a.calibrate_draft, prompts=a.calibrate_prompts) + if not tuned_ok: # #447: said again where it is not lost above the start say() warn("this PC is NOT tuned: the tuning failed (the reason is above); the model " + ("keeps" if a.no_start else "starts with") + " the default settings") @@ -4918,11 +4934,12 @@ def main() -> int: elif a.vision_tokens is not None: warn("--vision-tokens: images are off for this model, so it is not used") cfg_path = ROOT / f"strata-{tag.lower()}.json" - cal = setup_calibration(cfg, hip) # #566: Linux HIP too; the tuning is offered on NVIDIA only + cal = setup_calibration(cfg, hip) # #566: Linux HIP too; #744: the tuning itself runs on HIP if cal is not None: sys.path.insert(0, str(ROOT / "tools")) import calibrate as CAL - cfg["args"] = CAL.apply(cfg["args"], cal.get("settings") or {}) + cfg["args"] = CAL.apply(cfg["args"], cal.get("settings") or {}, backend=cfg.get("backend"), + draft=bool(cal.get("draft"))) ok("the settings tuned for this PC earlier are used" + (f" ({cal['date']})" if cal.get("date") else "")) else: # #642: measured counts (a calibration) win over the rule cfg["args"] = recommend_pool_workers(cfg["args"]) @@ -4933,7 +4950,7 @@ def main() -> int: write_setup_config(cfg_path, cfg, adopted if adopted is not None and adopted.name == cfg_path.name else None) script = write_run_script(tag, cfg_path, port, cfg.get("open_browser") is not False) # offered only when someone answers: --yes installs and adopted earlier installs are not held up by it - if cal is None and not hip and not a.no_start and not a.yes and ask( + if cal is None and not (hip and WIN) and not a.no_start and not a.yes and ask( "Tune Strata for this PC now? It measures a few engine settings (about 5-10 minutes; the PC is busy " "meanwhile; later: START-HERE --calibrate)", ["y", "n"], "y", a.yes) == "y": tuned = calibrate_config(cfg_path) diff --git a/src/kernels/cuda/iq_kernels.cu b/src/kernels/cuda/iq_kernels.cu index 7a9e4301b9..fbe459e0a5 100644 --- a/src/kernels/cuda/iq_kernels.cu +++ b/src/kernels/cuda/iq_kernels.cu @@ -8,6 +8,7 @@ #include "strata/kernels/dp4a.hpp" #include "strata/kernels/q8_1_finite.hpp" #include "s26_tsum.cuh" +#include "strata/kernels/decode_tuning.hpp" #include #include @@ -19,6 +20,7 @@ #include #include #include +#include #include namespace strata::kernels { @@ -633,11 +635,14 @@ __device__ __forceinline__ float row_dot(const uint8_t* row, const block_q8_1* x return warp_sum(s); } -template -__global__ void __launch_bounds__(128) mmvq_kernel(const uint8_t* __restrict__ w, size_t row_bytes, - const block_q8_1* __restrict__ x, float* __restrict__ y, int n_in, - int n_out, int ncols) { - const int row = blockIdx.x * 4 + threadIdx.y; +// ROWS rows per block, one warp each (blockDim = 32 x ROWS). The block size only groups whole rows, so every +// variant computes each row with the same warp, lanes and order: outputs are bitwise those of ROWS = 4, the +// default (decode_tuning.hpp; tools/hip/tune_decode.cpp checks it on the machine). +template +__global__ void __launch_bounds__(ROWS * 32) mmvq_kernel(const uint8_t* __restrict__ w, size_t row_bytes, + const block_q8_1* __restrict__ x, float* __restrict__ y, + int n_in, int n_out, int ncols) { + const int row = blockIdx.x * ROWS + threadIdx.y; if (row >= n_out) return; const int lane = threadIdx.x; const int nb = n_in / Fmt::qk; @@ -1208,15 +1213,17 @@ constexpr int GU_ROWS = 8; // rows per block (one warp each) // unused one was cap x (2 n_ff / 8 + n_embd / 8) blocks that started only to return. A smaller gridDim.y strides // instead. Each group's rows are computed by the same warp code in the same order either way, so the results do // not depend on gridDim.y (gridDim.y = cap_groups is the one-row-per-group launch). -template -__global__ void __launch_bounds__(256) native_gu_kernel(const unsigned long long* __restrict__ grp_ptr, +// ROWS rows per block, one warp each (blockDim = 32 * ROWS; default 8). As with mmvq_kernel, the block size never +// changes which warp reduces a row or in what order, so every variant is bitwise equal to the default. +template +__global__ void __launch_bounds__(ROWS * 32) native_gu_kernel(const unsigned long long* __restrict__ grp_ptr, const int32_t* __restrict__ grp_start, const int32_t* __restrict__ n_groups, const int32_t* __restrict__ ent_tok, const block_q8_1* __restrict__ xq, NativeExpertLayout L, float* __restrict__ gate, float* __restrict__ up) { const int warp = threadIdx.x >> 5, lane = threadIdx.x & 31; - const int row = blockIdx.x * GU_ROWS + warp; // 0 .. 2*n_ff + const int row = blockIdx.x * ROWS + warp; // 0 .. 2*n_ff if (row >= 2 * L.n_ff) return; const bool is_up = row >= L.n_ff; const int r = is_up ? row - (int) L.n_ff : row; @@ -1335,15 +1342,15 @@ __global__ void swiglu_entries_kernel(const float* __restrict__ gate, const floa h[i] = (g / (1.0f + __expf(-g))) * up[i]; } -template -__global__ void __launch_bounds__(256) native_down_kernel(const unsigned long long* __restrict__ grp_ptr, +template +__global__ void __launch_bounds__(ROWS * 32) native_down_kernel(const unsigned long long* __restrict__ grp_ptr, const int32_t* __restrict__ grp_start, const int32_t* __restrict__ n_groups, const int32_t* __restrict__ ent_dst, const block_q8_1* __restrict__ hq, NativeExpertLayout L, float* __restrict__ out) { const int warp = threadIdx.x >> 5, lane = threadIdx.x & 31; - const int r = blockIdx.x * 8 + warp; + const int r = blockIdx.x * ROWS + warp; if (r >= L.n_embd) return; const size_t off = L.down_off + (size_t) r * L.d_row; const int nb = (int) (L.n_ff / Fmt::qk), hb = (int) (L.n_ff / 32); @@ -1984,12 +1991,15 @@ void launch_mmvq_multi_nc(dim3 grid, dim3 block, cudaStream_t s, const uint8_t* } } -template +// The multi kernels and the old kernels are 4-rows-to-a-block shapes (mmvq_multi_kernel hardcodes it), so they run +// only in the default instantiation; an explicit-rows call (iq_mmvq_rows, the tuning table) takes the plain kernel. +template void launch_mmvq(const uint8_t* W, size_t rb, const block_q8_1* X, float* y, int n_in, int n_out, int ncols, cudaStream_t s) { - const dim3 grid((unsigned) ((n_out + 3) / 4)), block(32, 4); - if constexpr (!kSplit) mmvq_kernel<<>>(W, rb, X, y, n_in, n_out, ncols); - else if (g_old_kernels) mmvq_kernel<<>>(W, rb, X, y, n_in, n_out, ncols); + const dim3 grid((unsigned) ((n_out + ROWS - 1) / ROWS)), block(32, ROWS); + if constexpr (ROWS != 4 || !kSplit) + mmvq_kernel<<>>(W, rb, X, y, n_in, n_out, ncols); + else if (g_old_kernels) mmvq_kernel<<>>(W, rb, X, y, n_in, n_out, ncols); else if constexpr (kStageIqGrid) { if (!g_stage_grid_mmvq) { launch_mmvq_multi_nc(grid, block, s, W, rb, X, y, n_in, n_out, ncols); @@ -2662,12 +2672,16 @@ void s27_launch_ts(bool ts, const NativeExpertLayout& L, int64_t cap_groups, cud else s27_launch(L, cap_groups, s, grp_ptr, grp_start, n_groups, ent_dst, ent_tok, X, gate, up, h, hq, out, nh); } -template +// The multi kernels decode each weight part once with GU_ROWS rows to a block and the old kernels are the same +// shapes, so they run only in the default instantiation; an explicit-rows call (native_expert_grouped_rows, the +// tuning table) takes the plain kernel, with `grid` sized for ROWS by its caller. +template void launch_gu(dim3 grid, cudaStream_t s, const unsigned long long* grp_ptr, const int32_t* grp_start, const int32_t* n_groups, const int32_t* ent_tok, const block_q8_1* X, const NativeExpertLayout& L, float* gate, float* up) { - if constexpr (!kSplit) native_gu_kernel<<>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); - else if (g_old_kernels) native_gu_kernel<<>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); + if constexpr (ROWS != GU_ROWS || !kSplit) + native_gu_kernel<<>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); + else if (g_old_kernels) native_gu_kernel<<>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); else if constexpr (kStageIqGrid) { if (!g_stage_grid) { if (g_no_sub16_gu) native_gu_multi_kernel<<>>(grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); @@ -2682,12 +2696,15 @@ void launch_gu(dim3 grid, cudaStream_t s, const unsigned long long* grp_ptr, con } } -template +// as launch_gu: the multi kernel (and its n_ff = 640 grid) is the default instantiation's; an explicit-rows call +// takes the plain kernel, with `grid` sized for ROWS by its caller. +template void launch_down(dim3 grid, cudaStream_t s, const unsigned long long* grp_ptr, const int32_t* grp_start, const int32_t* n_groups, const int32_t* ent_dst, const block_q8_1* hq, const NativeExpertLayout& L, float* out) { - if constexpr (!kSplit) native_down_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); - else if (g_old_kernels) native_down_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); + if constexpr (ROWS != 8 || !kSplit) + native_down_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); + else if (g_old_kernels) native_down_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); else if (g_no_sub16_gu) native_down_multi_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); else native_down_multi_kernel<<>>(grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); } @@ -2732,20 +2749,64 @@ void quantize_q8_1_rows(const float* x, int64_t n_rows, int64_t n_cols, void* y, check("quantize_q8_1_rows"); } -void iq_mmvq(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, void* stream) { +namespace { +// The rows-per-block values compiled in. CUDA builds keep only the defaults (their behaviour and build time are +// unchanged); HIP builds also compile the variants the tuner measures (decode_tuning.cpp lists the same values). +#if defined(STRATA_USE_HIP) +#define STRATA_DECODE_VARIANTS 1 +#else +#define STRATA_DECODE_VARIANTS 0 +#endif + +template bool with_mmvq_rows(int rows, F&& f) { + switch (rows) { + case kDefaultMmvqRows: f(std::integral_constant{}); return true; +#if STRATA_DECODE_VARIANTS +#define STRATA_ROWS_CASE(v) case v: f(std::integral_constant{}); return true; + STRATA_MMVQ_ROWS_VARIANTS(STRATA_ROWS_CASE) +#undef STRATA_ROWS_CASE +#endif + default: return false; + } +} + +template bool with_expert_rows(int rows, F&& f) { + switch (rows) { + case kDefaultExpertRows: f(std::integral_constant{}); return true; +#if STRATA_DECODE_VARIANTS +#define STRATA_ROWS_CASE(v) case v: f(std::integral_constant{}); return true; + STRATA_EXPERT_ROWS_VARIANTS(STRATA_ROWS_CASE) +#undef STRATA_ROWS_CASE +#endif + default: return false; + } +} +} // namespace + +void iq_mmvq_rows(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, int rows, + void* stream) { const size_t rb = iq_row_bytes(t, n_in); cudaStream_t s = (cudaStream_t) stream; const auto* W = (const uint8_t*) w; const auto* X = (const block_q8_1*) x_q8_1; - switch (t) { -#define STRATA_MMVQ(T) case T: launch_mmvq(W, rb, X, y, n_in, n_out, ncols, s); break; - STRATA_MMVQ_FMTS(STRATA_MMVQ) + const bool ok = with_mmvq_rows(rows, [&](auto r) { + switch (t) { +#define STRATA_MMVQ(T) case T: launch_mmvq(W, rb, X, y, n_in, n_out, ncols, s); break; + STRATA_MMVQ_FMTS(STRATA_MMVQ) #undef STRATA_MMVQ - default: std::fprintf(stderr, "iq_mmvq: type %d is not supported\n", t); std::exit(1); - } + default: std::fprintf(stderr, "iq_mmvq: type %d is not supported\n", t); std::exit(1); + } + }); + if (!ok) { std::fprintf(stderr, "iq_mmvq: %d rows per block is not compiled in this build\n", rows); std::exit(1); } check("iq_mmvq"); } +void iq_mmvq(int t, const void* w, const void* x_q8_1, float* y, int n_in, int n_out, int ncols, void* stream) { + const DecodeTuningTable* tt = decode_tuning_active(); + const int rows = tt ? tt->rows(DecodeKernel::Mmvq, t, n_in, n_out, kDefaultMmvqRows) : kDefaultMmvqRows; + iq_mmvq_rows(t, w, x_q8_1, y, n_in, n_out, ncols, rows, stream); +} + 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); @@ -3375,6 +3436,18 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long const int32_t* n_groups, const int32_t* ent_dst, const int32_t* ent_tok, int64_t cap_groups, int64_t cap_entries, const void* x_q8_1, void* scratch, float* out, void* stream, int64_t grid_groups) { + // The decode tuning table (decode_tuning.hpp) names the plain gu/down kernels' rows per block: a non-default + // value routes to native_expert_grouped_rows (cap_groups group rows); otherwise this body runs unchanged. + const DecodeTuningTable* tt = decode_tuning_active(); + if (tt) { + const int gu = tt->rows(DecodeKernel::GateUp, L.gu_type, L.n_embd, L.n_ff, kDefaultExpertRows); + const int dn = tt->rows(DecodeKernel::Down, L.d_type, L.n_ff, L.n_embd, kDefaultExpertRows); + if (gu != kDefaultExpertRows || dn != kDefaultExpertRows) { + native_expert_grouped_rows(L, grp_ptr, grp_start, n_groups, ent_dst, ent_tok, cap_groups, cap_entries, + x_q8_1, scratch, out, gu, dn, stream); + return; + } + } if (cap_groups <= 0 || cap_entries <= 0) return; if (L.n_ff % 32 != 0) { std::fprintf(stderr, "native_expert_grouped: n_ff %lld\n", (long long) L.n_ff); std::exit(1); } cudaStream_t s = (cudaStream_t) stream; @@ -3530,4 +3603,141 @@ void native_expert_grouped(const NativeExpertLayout& L, const unsigned long long check("native_expert_grouped/down"); } +// The tuner's entry (decode_tuning.hpp): the same call with the gate/up and down kernels' rows per block explicit. +// The grid is one block row per possible group (cap_groups), the launch the tuner measures; the engine routes here +// only when the table names a non-default rows value, so grid_groups does not apply. A non-default rows value takes +// the launchers' plain one-warp-per-row kernels (the multi decode-once kernels hardcode 4 / 8 rows to a block); a +// value left at the default runs the default chain - bitwise the same output (iq_multi_parity checks it). +void native_expert_grouped_rows(const NativeExpertLayout& L, const unsigned long long* grp_ptr, + const int32_t* grp_start, const int32_t* n_groups, const int32_t* ent_dst, + const int32_t* ent_tok, int64_t cap_groups, int64_t cap_entries, const void* x_q8_1, + void* scratch, float* out, int gu_rows, int down_rows, void* stream) { + if (cap_groups <= 0 || cap_entries <= 0) return; + if (L.n_ff % 32 != 0) { std::fprintf(stderr, "native_expert_grouped: n_ff %lld\n", (long long) L.n_ff); std::exit(1); } + cudaStream_t s = (cudaStream_t) stream; + const size_t f = (size_t) cap_entries * (size_t) L.n_ff * sizeof(float), fa = (f + 255) & ~(size_t) 255; + float* gate = (float*) scratch; + float* up = (float*) ((uint8_t*) scratch + fa); + float* h = (float*) ((uint8_t*) scratch + 2 * fa); + block_q8_1* hq = (block_q8_1*) ((uint8_t*) scratch + 3 * fa); + const auto* X = (const block_q8_1*) x_q8_1; + const dim3 ggu((unsigned) ((2 * L.n_ff + gu_rows - 1) / gu_rows), (unsigned) cap_groups); + const bool gu_ok = with_expert_rows(gu_rows, [&](auto r) { + switch (L.gu_type) { +#define STRATA_GU(T) case T: launch_gu(ggu, s, grp_ptr, grp_start, n_groups, ent_tok, X, L, gate, up); break; + STRATA_GU_FMTS(STRATA_GU) +#undef STRATA_GU + default: std::fprintf(stderr, "native_expert_grouped: gate/up type %d\n", L.gu_type); std::exit(1); + } + }); + if (!gu_ok) { + std::fprintf(stderr, "native_expert_grouped: %d gate/up rows per block is not compiled in this build\n", gu_rows); + std::exit(1); + } + check("native_expert_grouped/gu"); + const long long nh = (long long) cap_entries * L.n_ff; + swiglu_entries_kernel<<<(unsigned) ((nh + 255) / 256), 256, 0, s>>>(gate, up, h, nh); + quantize_q8_1_kernel<<<(unsigned) ((nh + 255) / 256), 256, 0, s>>>(h, hq, nh); + const dim3 gd((unsigned) ((L.n_embd + down_rows - 1) / down_rows), (unsigned) cap_groups); + const bool down_ok = with_expert_rows(down_rows, [&](auto r) { + switch (L.d_type) { +#define STRATA_DOWN(T) case T: launch_down(gd, s, grp_ptr, grp_start, n_groups, ent_dst, hq, L, out); break; + STRATA_D_FMTS(STRATA_DOWN) +#undef STRATA_DOWN + default: std::fprintf(stderr, "native_expert_grouped: down type %d\n", L.d_type); std::exit(1); + } + }); + if (!down_ok) { + std::fprintf(stderr, "native_expert_grouped: %d down rows per block is not compiled in this build\n", down_rows); + std::exit(1); + } + check("native_expert_grouped/down"); +} + +// ---------------------------------------------------------------- the decode tuning table (decode_tuning.hpp) +DecodeIdentity decode_identity() { + DecodeIdentity id; + id.space = kDecodeTuningSpace; +#if defined(STRATA_USE_HIP) + int dev = 0; + hipDeviceProp_t p{}; + if (hipGetDevice(&dev) == hipSuccess && hipGetDeviceProperties(&p, dev) == hipSuccess) { + id.arch = p.gcnArchName; + const size_t colon = id.arch.find(':'); // "gfx1100:sramecc-:xnack-" -> "gfx1100" + if (colon != std::string::npos) id.arch.resize(colon); + } else { + id.arch = "unknown"; + } + int rt = 0; + if (hipRuntimeGetVersion(&rt) == hipSuccess) id.runtime = rt; +#if defined(__VERSION__) + id.toolchain = decode_toolchain_hash(std::string(__VERSION__) + " hip " + std::to_string(HIP_VERSION)); +#else + id.toolchain = decode_toolchain_hash("hip " + std::to_string(HIP_VERSION)); +#endif +#else + id.arch = "cuda"; + int rt = 0; + if (cudaRuntimeGetVersion(&rt) == cudaSuccess) id.runtime = rt; + id.toolchain = decode_toolchain_hash("cuda"); +#endif + return id; +} + +namespace { +struct ActiveTuning { + DecodeTuningTable table; + bool on = false; +}; +const ActiveTuning& active_tuning() { + static const ActiveTuning a = [] { + ActiveTuning t; + const char* path = std::getenv("STRATA_DECODE_TUNING"); + if (path == nullptr || *path == '\0') return t; +#if STRATA_DECODE_VARIANTS + // The table was measured against the default dispatch. A knob that reroutes the decode kernels (the old + // per-column kernels, sub16 off, staging off/on, the experimental expert paths) makes those measurements - + // and the bitwise checks behind them - describe a different build; refuse rather than mis-apply them. + const char* knob = nullptr; + if (g_old_kernels) knob = "STRATA_OLD_IQ_MMVQ"; + else if (g_no_sub16_gu) knob = "STRATA_NO_SUB16_GU"; + else if (!g_stage_grid) knob = "STRATA_IQ_STAGE_GRID"; + else if (g_stage_grid_mmvq) knob = "STRATA_IQ_STAGE_GRID_MMVQ"; + else if (env_on("STRATA_EXPERT_V2")) knob = "STRATA_EXPERT_V2"; + else if (env_on("STRATA_EXPERT_V2K")) knob = "STRATA_EXPERT_V2K"; + else if (env_on("STRATA_TSUM")) knob = "STRATA_TSUM"; + else if (env_on("STRATA_GROUPED_V1")) knob = "STRATA_GROUPED_V1"; + if (knob) { + std::fprintf(stderr, "decode tuning: %s refused (%s reroutes the decode dispatch away from what the table " + "measured); the default kernel shapes stay\n", path, knob); + return t; + } + const DecodeIdentity id = decode_identity(); + std::string err; + if (!t.table.load(path, id, err)) { + std::fprintf(stderr, "decode tuning: %s refused (%s); the default kernel shapes stay\n", path, err.c_str()); + return t; + } + t.on = true; + std::fprintf(stderr, "decode tuning: %zu shapes from %s (%s, runtime %lld)\n", t.table.entries().size(), path, + id.arch.c_str(), id.runtime); + if (std::getenv("STRATA_DECODE_TUNING_VERBOSE")) + for (const DecodeTuningRow& r : t.table.entries()) + std::fprintf(stderr, "decode tuning: %s\n", DecodeTuningTable::line(r).c_str()); +#else + std::fprintf(stderr, "decode tuning: STRATA_DECODE_TUNING is for HIP builds; this CUDA build ignores %s\n", path); +#endif + return t; + }(); + return a; +} +} // namespace + +void decode_tuning_init() { (void) active_tuning(); } + +const DecodeTuningTable* decode_tuning_active() { + const ActiveTuning& a = active_tuning(); + return a.on ? &a.table : nullptr; +} + } // namespace strata::kernels diff --git a/src/kernels/decode_tuning.cpp b/src/kernels/decode_tuning.cpp new file mode 100644 index 0000000000..8705ded73f --- /dev/null +++ b/src/kernels/decode_tuning.cpp @@ -0,0 +1,172 @@ +// src/kernels/decode_tuning.cpp - the decode tuning table (see include/strata/kernels/decode_tuning.hpp). +// Plain C++: no GPU headers, so the parser is tested without a GPU (tests/core/decode_tuning_test.cpp). +#include "strata/kernels/decode_tuning.hpp" + +#include +#include +#include +#include + +namespace strata::kernels { + +namespace { +#define ROW(v) v, +const int kExpertRows[] = {8, STRATA_EXPERT_ROWS_VARIANTS(ROW)}; // default first +const int kMmvqRows[] = {4, STRATA_MMVQ_ROWS_VARIANTS(ROW)}; +#undef ROW + +bool kernel_from_name(const std::string& s, DecodeKernel& k) { + if (s == "gu") { k = DecodeKernel::GateUp; return true; } + if (s == "down") { k = DecodeKernel::Down; return true; } + if (s == "mmvq") { k = DecodeKernel::Mmvq; return true; } + return false; +} + +std::string trim(const std::string& s) { + const size_t a = s.find_first_not_of(" \t\r\n"); + if (a == std::string::npos) return {}; + const size_t b = s.find_last_not_of(" \t\r\n"); + return s.substr(a, b - a + 1); +} +} // namespace + +bool decode_rows_valid(DecodeKernel k, int rows) noexcept { + if (k == DecodeKernel::Mmvq) return std::find(std::begin(kMmvqRows), std::end(kMmvqRows), rows) != std::end(kMmvqRows); + return std::find(std::begin(kExpertRows), std::end(kExpertRows), rows) != std::end(kExpertRows); +} + +std::vector decode_rows_candidates(DecodeKernel k) { + if (k == DecodeKernel::Mmvq) return {std::begin(kMmvqRows), std::end(kMmvqRows)}; + return {std::begin(kExpertRows), std::end(kExpertRows)}; +} + +const char* decode_kernel_name(DecodeKernel k) noexcept { + switch (k) { + case DecodeKernel::GateUp: return "gu"; + case DecodeKernel::Down: return "down"; + case DecodeKernel::Mmvq: return "mmvq"; + } + return "?"; +} + +int decode_default_rows(DecodeKernel k) noexcept { + return k == DecodeKernel::Mmvq ? kDefaultMmvqRows : kDefaultExpertRows; +} + +std::string decode_toolchain_hash(const std::string& s) { + uint64_t h = 1469598103934665603ull; + for (unsigned char c : s) { h ^= c; h *= 1099511628211ull; } + char b[17]; + std::snprintf(b, sizeof(b), "%016llx", (unsigned long long) h); + return b; +} + +std::string DecodeTuningTable::header(const DecodeIdentity& id) { + std::ostringstream o; + o << "STRATA_DECODE_TUNING_V1 " << id.arch << ' ' << id.runtime << ' ' << id.toolchain << ' ' << id.space; + return o.str(); +} + +std::string DecodeTuningTable::line(const DecodeTuningRow& r) { + std::ostringstream o; + o << decode_kernel_name(r.kernel) << ' ' << r.type << ' ' << r.n_in << ' ' << r.n_out << ' ' << r.rows; + return o.str(); +} + +bool DecodeTuningTable::parse(std::istream& in, const DecodeIdentity& want, std::string& err) { + rows_.clear(); + std::string raw; + bool have_header = false; + int lineno = 0; + while (std::getline(in, raw)) { + ++lineno; + const size_t hash = raw.find('#'); + const std::string s = trim(hash == std::string::npos ? raw : raw.substr(0, hash)); + if (s.empty()) continue; + std::istringstream ls(s); + std::string word; + ls >> word; + if (!have_header) { + if (word != "STRATA_DECODE_TUNING_V1") { + err = "line " + std::to_string(lineno) + ": expected the STRATA_DECODE_TUNING_V1 header"; + return false; + } + DecodeIdentity got; + if (!(ls >> got.arch >> got.runtime >> got.toolchain)) { + err = "line " + std::to_string(lineno) + ": the header needs "; + return false; + } + if (!(ls >> got.space)) { + err = "line " + std::to_string(lineno) + ": made before the kernel selection space was tracked - " + "the kernels have changed since; run the tuner again"; + return false; + } + if (got.arch != want.arch) { + err = "made for " + got.arch + ", this GPU is " + want.arch; + return false; + } + if (got.runtime != want.runtime) { + err = "made with HIP runtime " + std::to_string(got.runtime) + ", this one is " + std::to_string(want.runtime); + return false; + } + if (got.toolchain != want.toolchain) { + err = "made for an engine built by another compiler (" + got.toolchain + ", this build is " + + want.toolchain + "): run the tuner again"; + return false; + } + if (got.space != want.space) { + err = "made for kernel selection space " + std::to_string(got.space) + ", this build is " + + std::to_string(want.space) + " (the kernels changed): run the tuner again"; + return false; + } + have_header = true; + continue; + } + DecodeTuningRow r; + long long n_in = 0, n_out = 0; + std::string extra; + if (!kernel_from_name(word, r.kernel) || !(ls >> r.type >> n_in >> n_out >> r.rows) || (ls >> extra)) { + err = "line " + std::to_string(lineno) + ": expected `gu|down|mmvq `"; + return false; + } + r.n_in = n_in; + r.n_out = n_out; + if (r.type < 0 || r.n_in <= 0 || r.n_out <= 0) { + err = "line " + std::to_string(lineno) + ": type and shape must be positive"; + return false; + } + if (!decode_rows_valid(r.kernel, r.rows)) { + err = "line " + std::to_string(lineno) + ": " + std::to_string(r.rows) + " rows per block is not compiled for " + + decode_kernel_name(r.kernel); + return false; + } + for (const DecodeTuningRow& q : rows_) + if (q.kernel == r.kernel && q.type == r.type && q.n_in == r.n_in && q.n_out == r.n_out) { + err = "line " + std::to_string(lineno) + ": shape listed twice"; + return false; + } + rows_.push_back(r); + } + if (!have_header) { + err = "empty table (no STRATA_DECODE_TUNING_V1 header)"; + return false; + } + return true; +} + +bool DecodeTuningTable::load(const std::string& path, const DecodeIdentity& want, std::string& err) { + std::ifstream f(path); + if (!f) { + err = "cannot open " + path; + return false; + } + return parse(f, want, err); +} + +int DecodeTuningTable::rows(DecodeKernel k, int type, int64_t n_in, int64_t n_out, int fallback) const noexcept { + for (const DecodeTuningRow& r : rows_) + if (r.kernel == k && r.type == type && r.n_in == n_in && r.n_out == n_out) return r.rows; + return fallback; +} + +} // namespace strata::kernels diff --git a/src/program/generate.cpp b/src/program/generate.cpp index 3b551f3aaf..f2450f0c1b 100644 --- a/src/program/generate.cpp +++ b/src/program/generate.cpp @@ -49,6 +49,7 @@ #include "strata/kernels/native_qsa_indexer.hpp" #include "strata/kernels/native_rope.hpp" #include "strata/kernels/mrope.hpp" +#include "strata/kernels/decode_tuning.hpp" #include "strata/kernels/kv_q4.hpp" #include "strata/kernels/qsa.hpp" #include "strata/core/native_head.hpp" @@ -2630,6 +2631,9 @@ int main(int argc, char** argv) { strata::kernels::native_qsa_set_enabled(o.native_qsa); strata::kernels::native_qsa_indexer_set_enabled(o.native_qsa_indexer); strata::kernels::native_rope_set_enabled(o.native_rope); + // STRATA_DECODE_TUNING (HIP): the decode kernels' block shapes measured on this GPU (tools/hip/autotune.py); + // loaded here, before any graph captures a kernel that reads it + strata::kernels::decode_tuning_init(); // The vision path: every rope kernel reads a cell's (t, h, w) from this table (strata/kernels/mrope.hpp). It is // the identity until an image request, and it is set here, before any CUDA graph captures a rope kernel. int32_t* d_mrope = nullptr; diff --git a/tests/core/decode_tuning_test.cpp b/tests/core/decode_tuning_test.cpp new file mode 100644 index 0000000000..59ceb86599 --- /dev/null +++ b/tests/core/decode_tuning_test.cpp @@ -0,0 +1,89 @@ +// tests/core/decode_tuning_test.cpp - the decode tuning table: what it accepts, and every refusal (no GPU). +#include "strata/kernels/decode_tuning.hpp" + +#include +#include +#include + +using namespace strata::kernels; + +namespace { +int failures = 0; +void expect(bool ok, const char* what, const std::string& detail = {}) { + if (!ok) { + std::fprintf(stderr, "FAIL: %s %s\n", what, detail.c_str()); + ++failures; + } +} + +DecodeIdentity me() { return DecodeIdentity{"gfx1100", 70253211, "0123456789abcdef", kDecodeTuningSpace}; } + +bool parse(const std::string& text, DecodeTuningTable& t, std::string& err) { + std::istringstream in(text); + return t.parse(in, me(), err); +} +} // namespace + +int main() { + const std::string head = DecodeTuningTable::header(me()) + "\n"; + { + DecodeTuningTable t; + std::string err; + const bool ok = parse("# measured\n" + head + "gu 18 2560 640 16 # gate/up\n down 20 640 2560 4\nmmvq 18 2560 4096 8\n", t, err); + expect(ok, "a valid table parses", err); + expect(t.entries().size() == 3, "three shapes"); + expect(t.rows(DecodeKernel::GateUp, 18, 2560, 640, 8) == 16, "gu lookup"); + expect(t.rows(DecodeKernel::Down, 20, 640, 2560, 8) == 4, "down lookup"); + expect(t.rows(DecodeKernel::Mmvq, 18, 2560, 4096, 4) == 8, "mmvq lookup"); + expect(t.rows(DecodeKernel::GateUp, 21, 2560, 640, 8) == 8, "another type keeps the default"); + expect(t.rows(DecodeKernel::Down, 18, 2560, 640, 8) == 8, "the kernel is part of the key"); + // writing what was read gives the same table + std::string text = head; + for (const DecodeTuningRow& r : t.entries()) text += DecodeTuningTable::line(r) + "\n"; + DecodeTuningTable t2; + expect(parse(text, t2, err) && t2.entries().size() == 3, "header() and line() round-trip", err); + } + { + DecodeTuningTable t; + std::string err; + expect(parse(head, t, err) && t.empty(), "a header alone is a valid, empty table", err); + } + struct Bad { const char* why; std::string text; const char* says; }; + const Bad bad[] = { + {"no header", "gu 18 2560 640 8\n", "header"}, + {"empty file", "", "empty"}, + {"other arch", "STRATA_DECODE_TUNING_V1 gfx1101 70253211 0123456789abcdef 2\n", "gfx1101"}, + {"other runtime", "STRATA_DECODE_TUNING_V1 gfx1100 60000000 0123456789abcdef 2\n", "runtime"}, + {"other compiler", "STRATA_DECODE_TUNING_V1 gfx1100 70253211 fedcba9876543210 2\n", "compiler"}, + {"older table, no space", "STRATA_DECODE_TUNING_V1 gfx1100 70253211 0123456789abcdef\n", "run the tuner again"}, + {"other selection space", "STRATA_DECODE_TUNING_V1 gfx1100 70253211 0123456789abcdef 1\n", "selection space 1"}, + {"short header", "STRATA_DECODE_TUNING_V1 gfx1100\n", "header needs"}, + {"unknown kernel", head + "attn 18 2560 640 8\n", "expected"}, + {"missing field", head + "gu 18 2560 640\n", "expected"}, + {"extra field", head + "gu 18 2560 640 8 9\n", "expected"}, + {"rows not compiled", head + "gu 18 2560 640 32\n", "not compiled"}, + {"mmvq rows not compiled", head + "mmvq 18 2560 640 16\n", "not compiled"}, + {"zero shape", head + "gu 18 0 640 8\n", "positive"}, + {"duplicate", head + "gu 18 2560 640 8\ngu 18 2560 640 4\n", "twice"}, + }; + for (const Bad& b : bad) { + DecodeTuningTable t; + std::string err; + const bool ok = parse(b.text, t, err); + expect(!ok, b.why, "was accepted"); + expect(err.find(b.says) != std::string::npos, b.why, "message: " + err); + } + for (DecodeKernel k : {DecodeKernel::GateUp, DecodeKernel::Down, DecodeKernel::Mmvq}) { + const std::vector c = decode_rows_candidates(k); + expect(!c.empty() && c.front() == decode_default_rows(k), "candidates start with the default"); + for (int r : c) expect(decode_rows_valid(k, r) && r * 32 <= 512, "every candidate is valid"); + } + expect(decode_toolchain_hash("a").size() == 16 && decode_toolchain_hash("a") != decode_toolchain_hash("b"), + "toolchain hash"); + DecodeTuningTable missing; + std::string err; + expect(!missing.load("/nonexistent/decode-tuning.txt", me(), err) && err.find("cannot open") != std::string::npos, + "a missing file is refused"); + if (failures == 0) std::printf("decode_tuning_test: all checks passed\n"); + return failures == 0 ? 0 : 1; +} diff --git a/tools/calibrate.py b/tools/calibrate.py index 48d79c1eee..c03774014f 100644 --- a/tools/calibrate.py +++ b/tools/calibrate.py @@ -11,10 +11,22 @@ The first two are measured through one engine (per-request `strata_tune` keys); the worker count needs a restart per value. Decode speed only: the prompt path streams every expert whatever these settings say. +AMD (a config with "backend": "hip"): the PCIe share is left as the config has it (the HIP expert paths - mmap, +resident CPU experts, the staging buffer - were measured with the share they run with; `sweep_pcie=True` sweeps it +anyway), and the draft floor is swept wider: every extra draft token pulls more VRAM-missed experts onto the CPU. + +Drafting (`draft=True`, setup's --calibrate-draft): two more settings, a restart per value - + --spec the MTP's longest window (with prompt lookup on, the engine allows lookup windows 2 tokens longer) + --suffix-draft prompt lookup's shortest match (0 = MTP only) +None of them changes an answer: verification keeps exactly the tokens plain greedy decoding would produce; they +only decide which guesses are checked. The built-in prompts include one that repeats its input (an edit), where +prompt lookup pays; pass your own (`prompts=`, e.g. requests an agent really sends) for a measurement of YOUR mix. + A setting is kept only when it beats the default by more than MIN_GAIN in an interleaved re-measurement - the adaptive expert tier and the OS make single measurements noisy by a few percent. python tools/calibrate.py strata-q2_0.json # measure and print; setup.py --calibrate also saves it + python tools/calibrate.py strata-iq3_xxs.json --draft --prompts my-agent-prompts.json """ from __future__ import annotations @@ -32,12 +44,44 @@ MIN_GAIN = 0.03 # a setting must beat the default by this much to be kept PCIE_FRACS = (0.0, 0.2, 0.35, 0.55, 0.75) SPEC_MIN_PS = (0.3, 0.5, 0.7) +SPEC_MIN_PS_HIP = (0.3, 0.5, 0.6, 0.7, 0.8) +SPECS = (3, 4, 5) # the MTP's longest window (setup's default 4) +SUFFIX_DRAFTS = (0, 3, 6) # prompt lookup's shortest match (the engine's default 3; 0 = off) MAX_NEW = 128 PROMPTS = ( "Write a Python function that merges two sorted lists into one sorted list, with a docstring and two tests.", "Explain in two paragraphs how a refrigerator moves heat from inside to outside.", "List twelve European capitals with one sentence about each.", ) +# an edit: most of the answer repeats the input, which is what prompt lookup drafts (used with draft=True) +EDIT_PROMPT = ( + "Rename the variable `total` to `running_sum` in this function and return the whole function, nothing else:\n\n" + "def summarize(values, scale=1.0):\n" + " total = 0.0\n" + " count = 0\n" + " for v in values:\n" + " if v is None:\n" + " continue\n" + " total += v * scale\n" + " count += 1\n" + " mean = total / count if count else 0.0\n" + " return {\"total\": total, \"count\": count, \"mean\": mean}\n" +) + + +def load_prompts(path) -> list[str]: + """Prompts from a file: a JSON list of strings, or plain text with prompts separated by a line `---`.""" + text = Path(path).read_text(encoding="utf-8-sig") + if Path(path).suffix.lower() == ".json" or text.lstrip()[:1] in ("[", "{"): + items = json.loads(text) + if not isinstance(items, list) or not all(isinstance(x, str) for x in items): + raise ValueError(f"{path}: expected a JSON list of strings") + out = [x for x in items if x.strip()] + else: + out = [b.strip() for b in text.replace("\r\n", "\n").split("\n---\n") if b.strip()] + if not out: + raise ValueError(f"{path}: no prompts") + return out def chat_ids(tok, text: str) -> list[int]: @@ -119,9 +163,10 @@ def warm_up(self, rounds: int = 2): self.rate() -def run(cfg: dict, say=print, start_engine=None) -> dict: +def run(cfg: dict, say=print, start_engine=None, prompts=None, draft: bool = False, sweep_pcie=None) -> dict: """Measure on the engine `cfg` describes; returns {"settings": {flag: value}, "report": {...}}. - `start_engine(args)` returns a started engine (serve.server.StrataEngine or a stand-in in tests).""" + `start_engine(args)` returns a started engine (serve.server.StrataEngine or a stand-in in tests). + `prompts`: your own prompt texts (default: the built-in ones, plus an edit when `draft`).""" if start_engine is None: from serve.server import StrataEngine, child_env @@ -135,8 +180,10 @@ def start_engine(args): toks[i] = t tok = ST.Tokenizer(toks, (tpath / "merges.txt").read_text(encoding="utf-8").split("\n"), json.loads((tpath / "token_type.json").read_text())) - ids_list = [chat_ids(tok, p) for p in PROMPTS] - return measure(engine_args(cfg), ids_list, start_engine, say, host_worker_extras()) + texts = list(prompts) if prompts else list(PROMPTS) + ([EDIT_PROMPT] if draft else []) + ids_list = [chat_ids(tok, p) for p in texts] + return measure(engine_args(cfg), ids_list, start_engine, say, host_worker_extras(), + backend=cfg.get("backend"), draft=draft, sweep_pcie=sweep_pcie) def engine_args(cfg: dict) -> list[str]: @@ -162,11 +209,15 @@ def engine_error(log: str | None, since: int = 0) -> str | None: return next((x for x in reversed(lines) if x.startswith(("strata", "ERR"))), lines[-1] if lines else None) -def measure(base_args: list[str], ids_list, start_engine, say=print, extra_workers=()) -> dict: +def measure(base_args: list[str], ids_list, start_engine, say=print, extra_workers=(), backend=None, + draft: bool = False, sweep_pcie=None) -> dict: t0 = time.time() - report: dict = {} + hip = backend == "hip" + if sweep_pcie is None: + sweep_pcie = not hip + report: dict = {"backend": backend or "cuda", "draft": bool(draft), "pcie_swept": bool(sweep_pcie)} say(" Loading the model for the measurements ...") - base_args = apply(base_args, {}) # the product defaults: what the measurements must beat + base_args = apply(base_args, {}, backend=backend, draft=draft) # the product defaults the measurements must beat eng = start_engine(base_args) try: info = dict(getattr(eng, "info", {}) or {}) @@ -175,16 +226,21 @@ def measure(base_args: list[str], ids_list, start_engine, say=print, extra_worke d_workers = int(info.get("pool_workers", 0)) or None s = Session(eng, ids_list) s.warm_up() - # 1. the PCIe share, at the default draft floor + + def tune(pcie, minp): + return {"pcie_frac": pcie, "spec_min_p": minp} if sweep_pcie else {"spec_min_p": minp} + # 1. the PCIe share, at the default draft floor (NVIDIA; on AMD the config's own share stays) by_pcie = {} - for f in sorted(set(PCIE_FRACS) | {round(d_pcie, 2)}): - by_pcie[f] = [s.rate({"pcie_frac": f, "spec_min_p": d_minp})] - say(f" PCIe share {f:.2f}: {by_pcie[f][0]:.1f} tok/s") - best_pcie = max(by_pcie, key=lambda k: by_pcie[k][0]) + best_pcie = round(d_pcie, 2) + if sweep_pcie: + for f in sorted(set(PCIE_FRACS) | {round(d_pcie, 2)}): + by_pcie[f] = [s.rate(tune(f, d_minp))] + say(f" PCIe share {f:.2f}: {by_pcie[f][0]:.1f} tok/s") + best_pcie = max(by_pcie, key=lambda k: by_pcie[k][0]) # 2. the draft floor, at that share by_minp = {} - for p in sorted(set(SPEC_MIN_PS) | {round(d_minp, 2)}): - by_minp[p] = [s.rate({"pcie_frac": best_pcie, "spec_min_p": p})] + for p in sorted(set(SPEC_MIN_PS_HIP if hip else SPEC_MIN_PS) | {round(d_minp, 2)}): + by_minp[p] = [s.rate(tune(best_pcie, p))] say(f" draft floor {p:.2f}: {by_minp[p][0]:.1f} tok/s") best_minp = max(by_minp, key=lambda k: by_minp[k][0]) # 3. the winner against the default, interleaved, three times each @@ -193,7 +249,7 @@ def measure(base_args: list[str], ids_list, start_engine, say=print, extra_worke if cand != dflt: for _ in range(3): for k in (dflt, cand): - confirm[k].append(s.rate({"pcie_frac": k[0], "spec_min_p": k[1]})) + confirm[k].append(s.rate(tune(k[0], k[1]))) chosen = pick(confirm, dflt) if cand != dflt else dflt report.update(default={"pcie_frac": dflt[0], "spec_min_p": dflt[1], "pool_workers": d_workers}, pcie_sweep={str(k): v for k, v in by_pcie.items()}, @@ -202,25 +258,53 @@ def measure(base_args: list[str], ids_list, start_engine, say=print, extra_worke finally: close(eng) settings = {} - if chosen != dflt: - settings["--pcie-frac"] = f"{chosen[0]:.2f}" + if chosen != dflt: # (on AMD only the floor can differ: the share is fixed) + if sweep_pcie: + settings["--pcie-frac"] = f"{chosen[0]:.2f}" settings["--spec-min-p"] = f"{chosen[1]:.2f}" base_rate = statistics.median(confirm[chosen]) if confirm.get(chosen) else None - # 4. fewer CPU workers (a restart each), with the chosen settings - if d_workers and len(worker_candidates(d_workers, extra_workers)) > 1: - tuned = with_arg(with_arg(base_args, "--pcie-frac", f"{chosen[0]:.2f}"), "--spec-min-p", f"{chosen[1]:.2f}") - by_workers = {} - for w in worker_candidates(d_workers, extra_workers): - say(f" Measuring with {w} CPU workers (restarts the engine) ...") - e = start_engine(with_arg(tuned, "--pool-workers", None if w == d_workers else str(w))) + tuned = with_arg(base_args, "--spec-min-p", f"{chosen[1]:.2f}") + if sweep_pcie: + tuned = with_arg(tuned, "--pcie-frac", f"{chosen[0]:.2f}") + + def restarts(label, variants, default_key, measure_args): + """One engine start per variant: {key: [tok/s, tok/s]}, and the pick against `default_key`.""" + out = {} + for key in variants: + say(f" Measuring with {label} {key} (restarts the engine) ...") + e = start_engine(measure_args(key)) try: sw = Session(e, ids_list) sw.warm_up(1) - by_workers[w] = [sw.rate(), sw.rate()] - say(f" {w} workers: {statistics.median(by_workers[w]):.1f} tok/s") + out[key] = [sw.rate(), sw.rate()] + say(f" {label} {key}: {statistics.median(out[key]):.1f} tok/s") finally: close(e) - w_best = pick(by_workers, d_workers) + return out, pick(out, default_key) + + # 4. drafting (opt-in): the MTP's window, then prompt lookup's match length + if draft: + d_spec = int(arg_value(tuned, "--spec") or 4) + specs = sorted(set(SPECS) | {d_spec}) + by_spec, best_spec = restarts("MTP window", specs, d_spec, lambda v: with_arg(tuned, "--spec", str(v))) + report["spec"] = {str(k): v for k, v in by_spec.items()} + if best_spec != d_spec: + settings["--spec"] = str(best_spec) + tuned = with_arg(tuned, "--spec", str(best_spec)) + d_sfx = int(arg_value(tuned, "--suffix-draft") or 3) + sfx = sorted(set(SUFFIX_DRAFTS) | {d_sfx}) + by_sfx, best_sfx = restarts("prompt-lookup match", sfx, d_sfx, + lambda v: with_arg(tuned, "--suffix-draft", None if v == 3 else str(v))) + report["suffix_draft"] = {str(k): v for k, v in by_sfx.items()} + if best_sfx != d_sfx: + settings["--suffix-draft"] = str(best_sfx) + tuned = with_arg(tuned, "--suffix-draft", None if best_sfx == 3 else str(best_sfx)) + base_rate = statistics.median(by_sfx[best_sfx]) + # 5. fewer CPU workers (a restart each), with the chosen settings + if d_workers and len(worker_candidates(d_workers, extra_workers)) > 1: + by_workers, w_best = restarts( + "CPU workers", worker_candidates(d_workers, extra_workers), d_workers, + lambda w: with_arg(tuned, "--pool-workers", None if w == d_workers else str(w))) report["workers"] = {str(k): v for k, v in by_workers.items()} if w_best != d_workers: settings["--pool-workers"] = str(w_best) @@ -246,20 +330,34 @@ def close(eng): DEFAULTS = {"--pcie-frac": None, "--spec-min-p": "0.5", "--pool-workers": None} # None: the engine's own choice +HIP_DEFAULTS = {"--spec-min-p": "0.5", "--pool-workers": None} # the PCIe share is the config's +DRAFT_DEFAULTS = {"--spec": "4", "--suffix-draft": None} # setup's --spec, the engine's 3 -def apply(args: list[str], settings: dict) -> list[str]: +def apply(args: list[str], settings: dict, backend=None, draft: bool = False) -> list[str]: """`args` with the calibrated settings; a setting the calibration did not change goes back to the product default (setup's --spec-min-p 0.5, the engine's own PCIe share and worker count), so an older calibration's - values never linger.""" + values never linger. AMD (`backend="hip"`) leaves --pcie-frac as the config has it, unless the settings name + one; the drafting flags are reset only by a calibration that measured them (`draft`) or that names them.""" out = list(args) - for flag, default in DEFAULTS.items(): + flags = dict(HIP_DEFAULTS if backend == "hip" else DEFAULTS) + if backend == "hip" and "--pcie-frac" in settings: + flags["--pcie-frac"] = None + if draft or any(k in settings for k in DRAFT_DEFAULTS): + flags.update(DRAFT_DEFAULTS) + for flag, default in flags.items(): out = with_arg(out, flag, settings.get(flag, default)) return out if __name__ == "__main__": - if len(sys.argv) != 2: - sys.exit("usage: calibrate.py ") - res = run(json.loads(Path(sys.argv[1]).read_text(encoding="utf-8-sig"))) + import argparse + ap = argparse.ArgumentParser(description="measure Strata's hardware-dependent engine settings on this PC") + ap.add_argument("config", help="a strata-*.json run config") + ap.add_argument("--draft", action="store_true", help="also measure the drafting settings (--spec, --suffix-draft)") + ap.add_argument("--prompts", help="your own prompts: a JSON list of strings, or text blocks separated by ---") + ap.add_argument("--sweep-pcie", action="store_true", help="AMD: sweep the PCIe share too") + a = ap.parse_args() + res = run(json.loads(Path(a.config).read_text(encoding="utf-8-sig")), draft=a.draft, + prompts=load_prompts(a.prompts) if a.prompts else None, sweep_pcie=True if a.sweep_pcie else None) print(json.dumps(res, indent=1)) diff --git a/tools/hip/autotune.py b/tools/hip/autotune.py new file mode 100644 index 0000000000..b0cbc70746 --- /dev/null +++ b/tools/hip/autotune.py @@ -0,0 +1,321 @@ +"""On-device autotuning for the AMD (gfx1100) engine: measure, confirm end to end, and save into the run config. + +Two layers, both measured on this PC rather than guessed (docs/AMD_HIP_AUTOTUNE.md): + + 1. kernels - tune_decode times every block shape of the decode pass's GPU expert kernels at THIS model's types + and shapes, keeps a variant only if it is bitwise equal to the default and faster twice in a row, + and writes a table. Then the engine is started with and without the table (interleaved) and the + table is kept only if the generated tokens are identical and decode is not slower. The table goes + into the config's "env" (STRATA_DECODE_TUNING), like the hipBLASLt table. + 2. settings - (--settings / --draft) tools/calibrate.py on top: the draft floor and CPU workers, and with --draft + the MTP window and prompt lookup. Saved through setup's calibration record, so updates keep them. + +Stop the Strata server first: the kernel timings need the GPU idle and ~1 GB of free VRAM. + + python tools/hip/autotune.py strata-iq3_xxs.json # kernels + python tools/hip/autotune.py strata-iq3_xxs.json --settings # kernels, then the engine settings + python tools/hip/autotune.py strata-iq3_xxs.json --draft --prompts my-agent-prompts.json + python tools/hip/autotune.py strata-iq3_xxs.json --off # remove the kernel table from the config +""" +from __future__ import annotations + +import argparse +import json +import os +import shutil +import socket +import statistics +import subprocess +import sys +import threading +import time +from pathlib import Path + +ROOT = Path(__file__).resolve().parents[2] +sys.path.insert(0, str(ROOT / "tools")) +sys.path.insert(0, str(ROOT)) + +ENV_KEY = "STRATA_DECODE_TUNING" +NATIVE_DOWN = (20, 23, 42) # the down formats native_expert_grouped runs (IQ4_NL, IQ4_XS, Q2_0) +IQ_MMVQ = (16, 17, 18, 21, 22, 29) # the dense formats native_mmvq hands to iq_mmvq +MAX_NEW = 128 +MAX_SLOWDOWN = 0.02 # end to end, the table may not be slower than this (it is bitwise exact) + + +def say(*a): + print(*a, flush=True) + + +# ---------------------------------------------------------------- the model's shapes +def shapes_from_tensors(tensors) -> dict: + """The decode kernels' shapes in a model: {"experts": [(gu, down)], "n_embd", "n_ff", "mmvq": [(t, n_in, n_out)]}. + `tensors`: objects with .name, .shape (GGUF order: n_in first) and .type_id.""" + gate, down = {}, {} + mmvq = set() + n_embd = n_ff = None + for t in tensors: + n = t.name + if n.endswith("ffn_gate_exps.weight"): + gate[n.split(".")[1]] = t.type_id + n_embd, n_ff = int(t.shape[0]), int(t.shape[1]) + elif n.endswith("ffn_down_exps.weight"): + down[n.split(".")[1]] = t.type_id + elif len(t.shape) == 2 and t.type_id in IQ_MMVQ and "token_embd" not in n: + mmvq.add((t.type_id, int(t.shape[0]), int(t.shape[1]))) + pairs = sorted({(gate[b], down[b]) for b in gate if b in down and down[b] in NATIVE_DOWN}) + return {"experts": pairs, "n_embd": n_embd, "n_ff": n_ff, "mmvq": sorted(mmvq)} + + +def model_shapes(cfg: dict) -> dict: + from gguf_reader import GGUFFile + args = cfg["args"] + if "--native" not in args: + raise SystemExit("autotune: the config has no --native model (the AMD engine runs native packs only)") + first = Path(args[args.index("--native") + 1]) + files = sorted(first.parent.glob(first.name.replace("00001-of", "*-of"))) if "00001-of" in first.name else [first] + tensors = [] + for f in files: + tensors += GGUFFile(f).tensors + sh = shapes_from_tensors(tensors) + if not sh["experts"]: + raise SystemExit(f"autotune: no native expert tensors found in {first}") + return sh + + +# ---------------------------------------------------------------- the kernel tuner +def find_tuner(cfg: dict, build: bool = True) -> Path | None: + exe = Path(cfg["exe"]) + for p in (exe.parent / "tune_decode", ROOT / "build-hip" / "tune_decode", exe.parent.parent / "build-hip" / "tune_decode"): + if p.exists(): + return p + cache = ROOT / "build-hip" / "CMakeCache.txt" + if build and cache.exists(): + say(" Building the kernel tuner (tune_decode) in build-hip ...") + r = subprocess.run(["cmake", "--build", str(ROOT / "build-hip"), "--target", "tune_decode", "-j", + str(max(1, (os.cpu_count() or 2) // 2))]) + if r.returncode == 0 and (ROOT / "build-hip" / "tune_decode").exists(): + return ROOT / "build-hip" / "tune_decode" + return None + + +def tuner_args(sh: dict, out: Path, quick: bool) -> list[str]: + a = ["--n-embd", str(sh["n_embd"]), "--n-ff", str(sh["n_ff"]), "--out", str(out)] + for gu, dn in sh["experts"]: + a += ["--expert", f"{gu}:{dn}"] + for t, ni, no in sh["mmvq"]: + a += ["--mmvq", f"{t}:{ni}:{no}"] + if quick: + a += ["--rounds", "5", "--reps", "20"] + return a + + +def parse_results(stdout: str) -> list[dict]: + """The tuner's RESULT lines as dicts (`changed`, `default_us`, `chosen_us`, ...).""" + out = [] + for line in stdout.splitlines(): + if not line.startswith("RESULT "): + continue + f = line.split() + r = {"what": " ".join(x for x in f[1:] if "=" not in x)} + for kv in f[1:]: + if "=" in kv: + k, _, v = kv.partition("=") + r[k] = v if k == "chosen" else float(v) + r["changed"] = bool(r.get("changed")) + out.append(r) + return out + + +def table_shapes(path: Path) -> int: + """How many shapes a table changes (lines that are neither comments nor the header).""" + n = 0 + for line in path.read_text().splitlines(): + s = line.split("#", 1)[0].strip() + if s and not s.startswith("STRATA_DECODE_TUNING_V1"): + n += 1 + return n + + +# ---------------------------------------------------------------- end to end +def engine_starter(cfg: dict): + from serve.server import StrataEngine, child_env + + def start(table: str | None): + c = json.loads(json.dumps(cfg)) + env = dict(c.get("env") or {}) + env.pop(ENV_KEY, None) + if table: + env[ENV_KEY] = table + c["env"] = env + e = child_env(c) + if not table: + e.pop(ENV_KEY, None) + return StrataEngine(c["exe"], c["args"], cwd=c.get("cwd"), log=c.get("log"), env=e) + return start + + +def decode_run(eng, ids_list) -> tuple[list[float], list[list[int]]]: + """Greedy decode of every prompt: (tok/s per prompt, the tokens). One warm-up request first.""" + def one(ids): + toks = [t for t in eng.generate(ids, MAX_NEW, {"temperature": 0}, threading.Event()) if t is not None] + ms = (eng.last or {}).get("decode_ms") or 0.0 + return (len(toks) / (ms / 1000.0) if len(toks) > 8 and ms > 0 else 0.0), toks + one(ids_list[0]) + rates, toks = [], [] + for ids in ids_list: + r, t = one(ids) + rates.append(r) + toks.append(t) + return rates, toks + + +def end_to_end(start, table: str, ids_list, pairs: int = 2, say=say, close=None) -> dict: + """Start the engine without and with the table, `pairs` times each, interleaved; compare tokens and speed.""" + if close is None: + from calibrate import close + runs = {"off": [], "on": []} + tokens = {"off": [], "on": []} + for i in range(pairs): + for arm in ("off", "on"): + say(f" End-to-end check {i + 1}/{pairs}: {'with' if arm == 'on' else 'without'} the table (restarts the engine) ...") + eng = start(table if arm == "on" else None) + try: + rates, toks = decode_run(eng, ids_list) + finally: + close(eng) + runs[arm].append(statistics.median(rates)) + tokens[arm].append(toks) + say(f" {arm}: {runs[arm][-1]:.1f} tok/s") + base_det = all(t == tokens["off"][0] for t in tokens["off"]) + same = all(t == tokens["off"][0] for t in tokens["on"]) + off, on = statistics.median(runs["off"]), statistics.median(runs["on"]) + if same: + verdict, keep = "identical tokens", on >= off * (1.0 - MAX_SLOWDOWN) + elif pairs >= 2 and not base_det: + # two runs WITHOUT the table already differ (an adaptive expert tier, for one): tokens cannot judge it + verdict, keep = "not checkable (the default itself varies between runs)", on >= off * (1.0 - MAX_SLOWDOWN) + else: + verdict, keep = "TOKENS DIFFER", False + return {"off": runs["off"], "on": runs["on"], "off_tok_s": round(off, 1), "on_tok_s": round(on, 1), + "gain": round(on / off - 1.0, 4) if off > 0 else None, "tokens": verdict, "keep": keep} + + +def server_running(cfg: dict) -> bool: + port = int(cfg.get("port") or 8080) + with socket.socket() as s: + s.settimeout(0.5) + return s.connect_ex(("127.0.0.1", port)) == 0 + + +def save_config(cfg_path: Path, cfg: dict, table: str | None): + bak = cfg_path.with_name(cfg_path.name + ".bak-autotune") + if not bak.exists(): + shutil.copy2(cfg_path, bak) + env = dict(cfg.get("env") or {}) + if table: + env[ENV_KEY] = table + else: + env.pop(ENV_KEY, None) + if env: + cfg["env"] = env + else: + cfg.pop("env", None) + cfg_path.write_text(json.dumps(cfg, indent=1), encoding="utf-8") + + +def main(argv=None) -> int: + ap = argparse.ArgumentParser(description="on-device autotuning for Strata's AMD (gfx1100) engine") + ap.add_argument("config", help="the strata-*.json run config") + ap.add_argument("--settings", action="store_true", help="also tune the engine settings (draft floor, CPU workers)") + ap.add_argument("--draft", action="store_true", help="--settings, plus the MTP window and prompt lookup") + ap.add_argument("--prompts", help="your own prompts for the end-to-end measurements (JSON list or --- blocks)") + ap.add_argument("--quick", action="store_true", help="fewer rounds, one end-to-end pair (less certain)") + ap.add_argument("--skip-kernels", action="store_true", help="only the settings (with --settings / --draft)") + ap.add_argument("--off", action="store_true", help="remove the kernel table from the config and stop") + ap.add_argument("--force", action="store_true", help="run even if something answers on the server's port") + a = ap.parse_args(argv) + cfg_path = Path(a.config).resolve() + cfg = json.loads(cfg_path.read_text(encoding="utf-8-sig")) + if a.off: + save_config(cfg_path, cfg, None) + say(f" The kernel table is off for {cfg_path.name} (the default shapes; a backup is beside it).") + return 0 + if cfg.get("backend") != "hip": + say(" Note: this config is not marked \"backend\": \"hip\"; the kernel table applies to HIP builds only.") + if server_running(cfg) and not a.force: + say(f" Something answers on port {cfg.get('port') or 8080}: stop the Strata server first (the measurements need " + "the GPU to themselves), or pass --force.") + return 2 + import calibrate as CAL + report = {"config": str(cfg_path), "date": time.strftime("%Y-%m-%d %H:%M")} + + if not a.skip_kernels: + sh = model_shapes(cfg) + say(f" Model: experts {', '.join(f'{g}:{d}' for g, d in sh['experts'])} (n_embd {sh['n_embd']}, n_ff {sh['n_ff']})" + + (f", {len(sh['mmvq'])} i-quant dense shapes" if sh["mmvq"] else "")) + tuner = find_tuner(cfg) + if tuner is None: + say(" The kernel tuner is not built: cmake --build build-hip --target tune_decode (see docs/AMD_HIP_AUTOTUNE.md)") + return 1 + out = cfg_path.with_name(cfg_path.stem + ".decode-tuning.txt") + from serve.server import child_env + env = child_env(cfg) + env.pop(ENV_KEY, None) + say(f" Timing the decode kernels on this GPU ({tuner}) ...") + r = subprocess.run([str(tuner), *tuner_args(sh, out, a.quick)], capture_output=True, text=True, env=env) + sys.stderr.write(r.stderr) + if r.returncode == 3: + say(" Not enough free VRAM for the timings: stop the Strata server (and other GPU programs) first.") + return 2 + if r.returncode == 4: + say(" The tuner found a variant that is NOT bitwise equal to the default - nothing is saved. Please report it.") + return 4 + if r.returncode != 0 or not out.exists(): + say(f" The kernel tuner failed (exit {r.returncode}); the config is unchanged.") + return 1 + results = parse_results(r.stdout) + report["kernels"] = results + n = table_shapes(out) + if n == 0: + say(" The default kernel shapes are already the fastest here: nothing to save.") + save_config(cfg_path, cfg, None) + else: + texts = CAL.load_prompts(a.prompts) if a.prompts else list(CAL.PROMPTS) + [CAL.EDIT_PROMPT] + tok = tokenizer(cfg) + ids_list = [CAL.chat_ids(tok, t) for t in texts] + e2e = end_to_end(engine_starter(cfg), str(out), ids_list, pairs=1 if a.quick else 2) + report["end_to_end"] = e2e + gain = f"{e2e['gain'] * 100:+.1f}%" if e2e["gain"] is not None else "?" + say(f" End to end: {e2e['off_tok_s']} -> {e2e['on_tok_s']} tok/s ({gain}); tokens: {e2e['tokens']}") + if e2e["keep"]: + save_config(cfg_path, cfg, str(out)) + say(f" Kept: {n} kernel shape(s) from {out.name} are in {cfg_path.name}'s env ({ENV_KEY}).") + else: + save_config(cfg_path, cfg, None) + say(" Not kept: " + ("the tokens changed" if e2e["tokens"] == "TOKENS DIFFER" + else "decode was slower with it") + " - the config keeps the default shapes.") + cfg = json.loads(cfg_path.read_text(encoding="utf-8-sig")) + + if a.settings or a.draft: + import setup as S + ok = S.calibrate_config(cfg_path, draft=a.draft, prompts=a.prompts) + report["settings"] = "saved" if ok else "unchanged" + log = cfg_path.with_name(cfg_path.stem + ".autotune.json") + log.write_text(json.dumps(report, indent=1), encoding="utf-8") + say(f" Report: {log.name}") + return 0 + + +def tokenizer(cfg: dict): + import strata_tokenizer as ST + tpath = Path(cfg["tokenizer"]) + vocab = json.loads((tpath / "vocab.json").read_text(encoding="utf-8")) + toks = [None] * len(vocab) + for t, i in vocab.items(): + toks[i] = t + return ST.Tokenizer(toks, (tpath / "merges.txt").read_text(encoding="utf-8").split("\n"), + json.loads((tpath / "token_type.json").read_text())) + + +if __name__ == "__main__": + sys.exit(main()) diff --git a/tools/hip/test_autotune.py b/tools/hip/test_autotune.py new file mode 100644 index 0000000000..7a77c6b940 --- /dev/null +++ b/tools/hip/test_autotune.py @@ -0,0 +1,146 @@ +"""Tests for tools/hip/autotune.py without a GPU: the model's shapes, the tuner's output, the end-to-end verdict and +the config it writes. + + python -m unittest tools.hip.test_autotune +""" +from __future__ import annotations + +import json +import sys +import tempfile +import unittest +from pathlib import Path +from types import SimpleNamespace + +ROOT = Path(__file__).resolve().parents[2] +sys.path.insert(0, str(ROOT / "tools" / "hip")) +sys.path.insert(0, str(ROOT / "tools")) +import autotune as AT # noqa: E402 + + +def T(name, shape, type_id): + return SimpleNamespace(name=name, shape=shape, type_id=type_id) + + +class Shapes(unittest.TestCase): + def test_experts_dense_and_exclusions(self): + ts = [T(f"blk.{i}.ffn_gate_exps.weight", [2560, 640, 512], 18) for i in range(3)] + ts += [T(f"blk.{i}.ffn_up_exps.weight", [2560, 640, 512], 18) for i in range(3)] + ts += [T("blk.0.ffn_down_exps.weight", [640, 2560, 512], 20), T("blk.1.ffn_down_exps.weight", [640, 2560, 512], 20), + T("blk.2.ffn_down_exps.weight", [640, 2560, 512], 42)] + ts += [T("blk.0.attn_q.weight", [2560, 4096], 18), T("blk.1.attn_q.weight", [2560, 4096], 18), # one shape + T("blk.0.attn_k.weight", [2560, 512], 8), # Q8_0: not iq_mmvq + T("token_embd.weight", [2560, 248320], 18), # an embedding + T("blk.3.ffn_gate_exps.weight", [2560, 640, 512], 21)] # no down: skipped + sh = AT.shapes_from_tensors(ts) + self.assertEqual(sh["experts"], [(18, 20), (18, 42)]) + self.assertEqual((sh["n_embd"], sh["n_ff"]), (2560, 640)) + self.assertEqual(sh["mmvq"], [(18, 2560, 4096)]) + a = AT.tuner_args(sh, Path("/x/t.txt"), quick=True) + self.assertEqual(a.count("--expert"), 2) + self.assertIn("18:42", a) + self.assertIn("18:2560:4096", a) + self.assertIn("--rounds", a) + + def test_non_native_down_is_skipped(self): + sh = AT.shapes_from_tensors([T("blk.0.ffn_gate_exps.weight", [2560, 640, 512], 18), + T("blk.0.ffn_down_exps.weight", [640, 2560, 512], 12)]) # Q4_K down + self.assertEqual(sh["experts"], []) + + +class TunerOutput(unittest.TestCase): + def test_parse_results(self): + out = ("noise\nRESULT experts 18 20 2560 640 default_us=41.20 chosen=16/4 chosen_us=36.10 changed=1\n" + "RESULT mmvq 18 2560 4096 default_us=9.00 chosen=4 chosen_us=9.00 changed=0\nSUMMARY changed=2 refused=0\n") + r = AT.parse_results(out) + self.assertEqual(len(r), 2) + self.assertEqual(r[0]["what"], "experts 18 20 2560 640") + self.assertEqual(r[0]["chosen"], "16/4") + self.assertTrue(r[0]["changed"]) + self.assertAlmostEqual(r[0]["chosen_us"], 36.1) + self.assertFalse(r[1]["changed"]) + + def test_table_shapes(self): + with tempfile.TemporaryDirectory() as d: + p = Path(d) / "t.txt" + p.write_text("# made here\nSTRATA_DECODE_TUNING_V1 gfx1100 1 abc\ngu 18 2560 640 16\n\ndown 20 640 2560 4 # x\n") + self.assertEqual(AT.table_shapes(p), 2) + p.write_text("# nothing won\nSTRATA_DECODE_TUNING_V1 gfx1100 1 abc\n") + self.assertEqual(AT.table_shapes(p), 0) + + +class FakeEngine: + def __init__(self, rate, tokens): + self.rate, self.tokens, self.last, self.proc = rate, tokens, {}, None + + def generate(self, ids, max_new, sampling, cancel): + toks = self.tokens(ids) + yield from toks + self.last = {"decode_ms": len(toks) / self.rate * 1000.0} + + +class EndToEnd(unittest.TestCase): + IDS = [[1, 2], [3, 4], [5, 6]] + + def verdict(self, rate_on, tokens_on, tokens_off=None, pairs=2): + n = {"off": 0} + + def start(table): + if table: + return FakeEngine(rate_on, tokens_on) + n["off"] += 1 + return FakeEngine(50.0, tokens_off(n["off"]) if tokens_off else (lambda ids: list(range(20)))) + return AT.end_to_end(start, "/t.txt", self.IDS, pairs=pairs, say=lambda *_: None, close=lambda e: None) + + def test_faster_and_identical_is_kept(self): + v = self.verdict(53.0, lambda ids: list(range(20))) + self.assertTrue(v["keep"]) + self.assertEqual(v["tokens"], "identical tokens") + self.assertAlmostEqual(v["gain"], 0.06, places=3) + + def test_noise_level_slower_is_still_kept(self): + self.assertTrue(self.verdict(49.5, lambda ids: list(range(20)))["keep"]) # -1%: within MAX_SLOWDOWN + + def test_slower_is_dropped(self): + self.assertFalse(self.verdict(45.0, lambda ids: list(range(20)))["keep"]) + + def test_different_tokens_are_dropped(self): + v = self.verdict(60.0, lambda ids: list(range(19)) + [99]) + self.assertFalse(v["keep"]) + self.assertEqual(v["tokens"], "TOKENS DIFFER") + + def test_nondeterministic_default_is_judged_on_speed(self): + v = self.verdict(55.0, lambda ids: [7] * 20, tokens_off=lambda k: (lambda ids: [k] * 20)) + self.assertTrue(v["keep"]) + self.assertIn("not checkable", v["tokens"]) + + def test_quick_cannot_excuse_different_tokens(self): + v = self.verdict(60.0, lambda ids: [7] * 20, pairs=1) + self.assertFalse(v["keep"]) + + +class Config(unittest.TestCase): + def test_save_and_off(self): + with tempfile.TemporaryDirectory() as d: + p = Path(d) / "strata-iq3_xxs.json" + cfg = {"exe": "x", "args": [], "env": {"STRATA_HIPBLASLT_TUNING": "/h.txt"}, "port": 1} + p.write_text(json.dumps(cfg)) + AT.save_config(p, json.loads(p.read_text()), "/t.txt") + w = json.loads(p.read_text()) + self.assertEqual(w["env"], {"STRATA_HIPBLASLT_TUNING": "/h.txt", AT.ENV_KEY: "/t.txt"}) + self.assertTrue((Path(d) / "strata-iq3_xxs.json.bak-autotune").exists()) + self.assertEqual(AT.main([str(p), "--off"]), 0) + self.assertEqual(json.loads(p.read_text())["env"], {"STRATA_HIPBLASLT_TUNING": "/h.txt"}) + # the backup is the ORIGINAL config, not overwritten by later saves + self.assertEqual(json.loads((Path(d) / "strata-iq3_xxs.json.bak-autotune").read_text()), cfg) + + def test_env_dropped_when_empty(self): + with tempfile.TemporaryDirectory() as d: + p = Path(d) / "c.json" + p.write_text(json.dumps({"exe": "x", "args": [], "env": {AT.ENV_KEY: "/t.txt"}})) + AT.save_config(p, json.loads(p.read_text()), None) + self.assertNotIn("env", json.loads(p.read_text())) + + +if __name__ == "__main__": + unittest.main() diff --git a/tools/hip/tune_decode.cpp b/tools/hip/tune_decode.cpp new file mode 100644 index 0000000000..04837440b4 --- /dev/null +++ b/tools/hip/tune_decode.cpp @@ -0,0 +1,574 @@ +// tools/hip/tune_decode.cpp - measure the decode kernels' block shapes on THIS GPU and write a tuning table. +// +// What it tunes (include/strata/kernels/decode_tuning.hpp): rows per thread block of the grouped GPU experts +// (gate/up and down, `native_expert_grouped`) and of `iq_mmvq`. Every candidate computes every row with the same +// warp, lanes and order, so it is bitwise equal to the default - and this tool checks that on the machine before a +// candidate may win (a candidate that differs by one bit is reported and never written). +// +// The workload is the verify pass's: windows of 1, 3 and 5 tokens (the 0.1.27-era MTP and lookup mix; weighted +// 0.2 / 0.3 / 0.5 - override both with --windows / --weights when the running policy drafts wider, and the report +// line names what was measured), `--hit-frac` of each token's 10 experts on the GPU, experts drawn from a pool +// larger than the 7900 XTX's 96 MB Infinity Cache (decode streams ~1 GB of experts per token, so a cache-warm +// benchmark would lie), every candidate measured in every round, interleaved, and a winner kept only when its +// weighted median beats the default by more than --min-gain AND a second interleaved confirmation agrees. +// +// Run it with the Strata server STOPPED (it needs ~1 GB of free VRAM, and a busy GPU makes timings meaningless): +// +// build-hip/tune_decode --expert 18:20 --out tools/hip/gfx1100-decode-tuning.txt +// +// tools/hip/autotune.py finds the model's types and shapes, runs this, then confirms the table end to end. +#include "strata/kernels/decode_tuning.hpp" +#include "strata/kernels/iq_kernels.hpp" + +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace K = strata::kernels; + +namespace { +#define HIP_CHECK(call) \ + do { \ + const hipError_t e_ = (call); \ + if (e_ != hipSuccess) { \ + std::fprintf(stderr, "tune_decode: %s:%d %s: %s\n", __FILE__, __LINE__, #call, hipGetErrorString(e_)); \ + std::exit(2); \ + } \ + } while (0) + +struct Dev { + void* p = nullptr; + size_t n = 0; + explicit Dev(size_t bytes) : n(bytes) { if (bytes) HIP_CHECK(hipMalloc(&p, bytes)); } + ~Dev() { if (p) (void) hipFree(p); } + Dev(const Dev&) = delete; + Dev& operator=(const Dev&) = delete; + template T* as() const { return static_cast(p); } +}; + +struct Options { + std::vector> experts; // (gate/up type, down type) + std::vector> mmvq; // (type, n_in, n_out) + int64_t n_embd = 2560, n_ff = 640; + int k_used = 10; + double hit_frac = 0.6; + int pool_mib = 768; + int rounds = 9, reps = 40; + double min_gain = 0.03; + // the verify-window workload the timings blend: window widths (tokens) and their weights. + // The defaults are the 0.1.27-era MTP proxy {1,3,5} x {0.2,0.3,0.5}; a build running a wider + // speculative policy should measure its own mix (--windows / --weights), and the report line + // names what was used. + int windows[3] = {1, 3, 5}; + double weights[3] = {0.2, 0.3, 0.5}; + std::string out; + bool quiet = false; +}; + +// the parsed options, for the free functions below (the tool is single-run; set once in main) +const Options* g_opt = nullptr; + +int block_elems(int type) { + switch (type) { + case 20: return 32; // IQ4_NL + case 42: return 64; // Q2_0 + default: return 256; + } +} + +// random block bytes, then a small finite fp16 scale where the format starts with one (all but IQ1_M), so timings +// run on ordinary numbers rather than NaN/Inf paths +void fill_blocks(std::vector& host, int type, std::mt19937& rng) { + std::uniform_int_distribution byte(0, 255); + for (uint8_t& b : host) b = (uint8_t) byte(rng); + if (type == 29) return; + const size_t bb = K::iq_row_bytes(type, block_elems(type)); + if (bb == 0) return; + const uint16_t half_small = 0x2C00; // 0.0625 + for (size_t off = 0; off + bb <= host.size(); off += bb) std::memcpy(&host[off], &half_small, 2); +} + +double median(std::vector v) { + if (v.empty()) return 0.0; + std::sort(v.begin(), v.end()); + const size_t m = v.size() / 2; + return v.size() % 2 ? v[m] : 0.5 * (v[m - 1] + v[m]); +} + +double free_mib() { + size_t fr = 0, tot = 0; + HIP_CHECK(hipMemGetInfo(&fr, &tot)); + return (double) fr / (1024.0 * 1024.0); +} + +// ---------------------------------------------------------------- grouped experts +struct Plan { // one verify window's GPU groups (device arrays) + Dev ptr, start, count, dst, tok; + int groups = 0, entries = 0; + Plan(int cap) : ptr(sizeof(unsigned long long) * cap), start(sizeof(int32_t) * (cap + 1)), + count(sizeof(int32_t)), dst(sizeof(int32_t) * cap), tok(sizeof(int32_t) * cap) {} +}; + +struct ExpertBench { + K::NativeExpertLayout L; + int64_t pool = 0; + Dev blobs; + std::vector> plans[3]; // per window, several distinct expert sets + int cap[3] = {0, 0, 0}; + std::unique_ptr xq[3], scratch[3], out[3]; + + ExpertBench(const K::NativeExpertLayout& lay, int pool_mib) + : L(lay), pool(std::max(64, ((int64_t) pool_mib << 20) / (int64_t) lay.bytes)), + blobs((size_t) pool * lay.bytes) {} +}; + +void make_expert_bench(ExpertBench& b, const Options& o, std::mt19937& rng) { + // the blobs: gate | up | down, each part filled as its own format + { + std::vector gu((size_t) b.L.down_off), dn(b.L.bytes - b.L.down_off); + for (int64_t i = 0; i < b.pool; ++i) { + fill_blocks(gu, b.L.gu_type, rng); + fill_blocks(dn, b.L.d_type, rng); + uint8_t* base = b.blobs.as() + (size_t) i * b.L.bytes; + HIP_CHECK(hipMemcpy(base, gu.data(), gu.size(), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(base + b.L.down_off, dn.data(), dn.size(), hipMemcpyHostToDevice)); + } + } + const int kPlans = 16; + for (int w = 0; w < 3; ++w) { + const int T = g_opt->windows[w]; + const int cap = T * o.k_used; + b.cap[w] = cap; + // activations: T tokens of n_embd, quantized by the engine's own q8_1 kernel + std::normal_distribution nd(0.0f, 1.0f); + std::vector xh((size_t) T * (size_t) b.L.n_embd); + for (float& v : xh) v = nd(rng); + Dev xf(xh.size() * sizeof(float)); + HIP_CHECK(hipMemcpy(xf.p, xh.data(), xf.n, hipMemcpyHostToDevice)); + b.xq[w] = std::make_unique((size_t) T * (size_t) (b.L.n_embd / 32) * 36); + K::quantize_q8_1_rows(xf.as(), T, b.L.n_embd, b.xq[w]->p, nullptr); + HIP_CHECK(hipDeviceSynchronize()); + b.scratch[w] = std::make_unique(K::native_expert_scratch_bytes(cap, b.L.n_ff)); + b.out[w] = std::make_unique((size_t) cap * (size_t) b.L.n_embd * sizeof(float)); + // entries: hit_frac of the window's T*K picks; distinct experts, a few shared by two tokens + const int entries = std::max(1, (int) std::lround(o.hit_frac * cap)); + const int groups = std::max(1, std::min(entries, (int) std::lround(entries * (T == 1 ? 1.0 : 0.8)))); + for (int pl = 0; pl < kPlans; ++pl) { + auto p = std::make_unique(cap); + p->groups = groups; + p->entries = entries; + std::vector pick(b.pool); + for (int64_t i = 0; i < b.pool; ++i) pick[(size_t) i] = i; + std::shuffle(pick.begin(), pick.end(), rng); + std::vector ptr(cap, 0); + std::vector start(cap + 1, entries), dst(cap, 0), tok(cap, 0); + for (int g = 0; g < groups; ++g) + ptr[g] = (unsigned long long) (uintptr_t) (b.blobs.as() + (size_t) pick[(size_t) g] * b.L.bytes); + // spread the entries over the groups (the first entries - groups get a second entry) + int e = 0; + for (int g = 0; g < groups; ++g) { + start[g] = e; + const int n = 1 + (g < entries - groups ? 1 : 0); + for (int j = 0; j < n; ++j, ++e) { + dst[e] = e; + tok[e] = (g + j) % T; + } + } + start[groups] = e; + const int32_t cnt = groups; + HIP_CHECK(hipMemcpy(p->ptr.p, ptr.data(), ptr.size() * sizeof(ptr[0]), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(p->start.p, start.data(), start.size() * sizeof(int32_t), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(p->count.p, &cnt, sizeof(cnt), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(p->dst.p, dst.data(), dst.size() * sizeof(int32_t), hipMemcpyHostToDevice)); + HIP_CHECK(hipMemcpy(p->tok.p, tok.data(), tok.size() * sizeof(int32_t), hipMemcpyHostToDevice)); + b.plans[w].push_back(std::move(p)); + } + } +} + +void run_grouped(ExpertBench& b, int w, int plan, int gu, int dn, hipStream_t s) { + const Plan& p = *b.plans[w][(size_t) plan]; + K::native_expert_grouped_rows(b.L, p.ptr.as(), p.start.as(), p.count.as(), + p.dst.as(), p.tok.as(), b.cap[w], b.cap[w], b.xq[w]->p, + b.scratch[w]->p, b.out[w]->as(), gu, dn, s); +} + +// The engine's own default: the real entry point, no table active in this process. The default candidate and the +// bitwise reference are THIS call (its heuristics may pick the multi sub-warp kernels), not an explicit-rows pair - +// so the tuner's numbers describe the launch the engine actually makes without a table. +void run_default(ExpertBench& b, int w, int plan, hipStream_t s) { + const Plan& p = *b.plans[w][(size_t) plan]; + K::native_expert_grouped(b.L, p.ptr.as(), p.start.as(), p.count.as(), + p.dst.as(), p.tok.as(), b.cap[w], b.cap[w], b.xq[w]->p, + b.scratch[w]->p, b.out[w]->as(), s, 0); +} + +std::vector grouped_output(ExpertBench& b, int w, int gu, int dn, hipStream_t s) { + HIP_CHECK(hipMemsetAsync(b.out[w]->p, 0, b.out[w]->n, s)); + run_grouped(b, w, 0, gu, dn, s); + HIP_CHECK(hipStreamSynchronize(s)); + std::vector h(b.out[w]->n); + HIP_CHECK(hipMemcpy(h.data(), b.out[w]->p, h.size(), hipMemcpyDeviceToHost)); + return h; +} + +std::vector default_output(ExpertBench& b, int w, hipStream_t s) { + HIP_CHECK(hipMemsetAsync(b.out[w]->p, 0, b.out[w]->n, s)); + run_default(b, w, 0, s); + HIP_CHECK(hipStreamSynchronize(s)); + std::vector h(b.out[w]->n); + HIP_CHECK(hipMemcpy(h.data(), b.out[w]->p, h.size(), hipMemcpyDeviceToHost)); + return h; +} + +// microseconds per call: `reps` calls cycling through the plans, timed by events +double time_grouped(ExpertBench& b, int w, int gu, int dn, int reps, hipStream_t s, hipEvent_t e0, hipEvent_t e1) { + const int np = (int) b.plans[w].size(); + for (int i = 0; i < 2; ++i) run_grouped(b, w, i % np, gu, dn, s); + HIP_CHECK(hipEventRecord(e0, s)); + for (int i = 0; i < reps; ++i) run_grouped(b, w, i % np, gu, dn, s); + HIP_CHECK(hipEventRecord(e1, s)); + HIP_CHECK(hipEventSynchronize(e1)); + float ms = 0.0f; + HIP_CHECK(hipEventElapsedTime(&ms, e0, e1)); + return 1000.0 * ms / reps; +} + +double time_default(ExpertBench& b, int w, int reps, hipStream_t s, hipEvent_t e0, hipEvent_t e1) { + const int np = (int) b.plans[w].size(); + for (int i = 0; i < 2; ++i) run_default(b, w, i % np, s); + HIP_CHECK(hipEventRecord(e0, s)); + for (int i = 0; i < reps; ++i) run_default(b, w, i % np, s); + HIP_CHECK(hipEventRecord(e1, s)); + HIP_CHECK(hipEventSynchronize(e1)); + float ms = 0.0f; + HIP_CHECK(hipEventElapsedTime(&ms, e0, e1)); + return 1000.0 * ms / reps; +} + +// ---------------------------------------------------------------- iq_mmvq +struct MmvqBench { + int type; + int n_in, n_out; + int64_t copies; + size_t wbytes; + Dev w; + std::unique_ptr xq[3], y[3]; + MmvqBench(int t, int ni, int no, int pool_mib) + : type(t), n_in(ni), n_out(no), copies(0), wbytes(K::iq_row_bytes(t, ni) * (size_t) no), + w(std::max(1, wbytes) * (size_t) std::max(2, ((int64_t) pool_mib << 20) / (int64_t) std::max(1, wbytes))) { + copies = (int64_t) (w.n / wbytes); + } +}; + +void make_mmvq_bench(MmvqBench& b, std::mt19937& rng) { + std::vector h(b.wbytes); + for (int64_t c = 0; c < b.copies; ++c) { + fill_blocks(h, b.type, rng); + HIP_CHECK(hipMemcpy(b.w.as() + (size_t) c * b.wbytes, h.data(), h.size(), hipMemcpyHostToDevice)); + } + std::normal_distribution nd(0.0f, 1.0f); + for (int wi = 0; wi < 3; ++wi) { + const int T = g_opt->windows[wi]; + std::vector xh((size_t) T * (size_t) b.n_in); + for (float& v : xh) v = nd(rng); + Dev xf(xh.size() * sizeof(float)); + HIP_CHECK(hipMemcpy(xf.p, xh.data(), xf.n, hipMemcpyHostToDevice)); + b.xq[wi] = std::make_unique((size_t) T * (size_t) (b.n_in / 32) * 36); + K::quantize_q8_1_rows(xf.as(), T, b.n_in, b.xq[wi]->p, nullptr); + b.y[wi] = std::make_unique((size_t) T * (size_t) b.n_out * sizeof(float)); + } + HIP_CHECK(hipDeviceSynchronize()); +} + +void run_mmvq(MmvqBench& b, int wi, int copy, int rows, hipStream_t s) { + K::iq_mmvq_rows(b.type, b.w.as() + (size_t) copy * b.wbytes, b.xq[wi]->p, b.y[wi]->as(), b.n_in, + b.n_out, g_opt->windows[wi], rows, s); +} + +std::vector mmvq_output(MmvqBench& b, int wi, int rows, hipStream_t s) { + HIP_CHECK(hipMemsetAsync(b.y[wi]->p, 0, b.y[wi]->n, s)); + run_mmvq(b, wi, 0, rows, s); + HIP_CHECK(hipStreamSynchronize(s)); + std::vector h(b.y[wi]->n); + HIP_CHECK(hipMemcpy(h.data(), b.y[wi]->p, h.size(), hipMemcpyDeviceToHost)); + return h; +} + +double time_mmvq(MmvqBench& b, int wi, int rows, int reps, hipStream_t s, hipEvent_t e0, hipEvent_t e1) { + for (int i = 0; i < 2; ++i) run_mmvq(b, wi, (int) (i % b.copies), rows, s); + HIP_CHECK(hipEventRecord(e0, s)); + for (int i = 0; i < reps; ++i) run_mmvq(b, wi, (int) (i % b.copies), rows, s); + HIP_CHECK(hipEventRecord(e1, s)); + HIP_CHECK(hipEventSynchronize(e1)); + float ms = 0.0f; + HIP_CHECK(hipEventElapsedTime(&ms, e0, e1)); + return 1000.0 * ms / reps; +} + +// ---------------------------------------------------------------- the search +// Weighted (by window) median microseconds of each candidate, measured interleaved over `rounds` rounds. +template +std::map measure(const std::vector& cands, int rounds, TimeFn&& t) { + std::map> per[3]; + for (int r = 0; r < rounds; ++r) + for (int w = 0; w < 3; ++w) { + // rotate the order each round so no candidate always runs first (clock ramp, cache state) + for (size_t i = 0; i < cands.size(); ++i) { + const Cand& c = cands[(i + (size_t) r) % cands.size()]; + per[w][c].push_back(t(c, w)); + } + } + std::map score; + for (const Cand& c : cands) { + double s = 0.0; + for (int w = 0; w < 3; ++w) s += g_opt->weights[w] * median(per[w][c]); + score[c] = s; + } + return score; +} + +struct Choice { + bool changed = false; + double default_us = 0.0, best_us = 0.0, confirm_default_us = 0.0, confirm_best_us = 0.0; +}; + +template +Cand choose(const std::vector& cands, const Cand& dflt, const Options& o, TimeFn&& t, Choice& ch) { + const std::map s = measure(cands, o.rounds, t); + Cand best = dflt; + for (const auto& kv : s) + if (kv.second < s.at(best)) best = kv.first; + ch.default_us = s.at(dflt); + ch.best_us = s.at(best); + if (best == dflt || ch.best_us > ch.default_us * (1.0 - o.min_gain)) return dflt; + // confirm: the winner against the default alone, interleaved, as many rounds again + const std::map c = measure(std::vector{dflt, best}, o.rounds, t); + ch.confirm_default_us = c.at(dflt); + ch.confirm_best_us = c.at(best); + if (ch.confirm_best_us > ch.confirm_default_us * (1.0 - o.min_gain)) return dflt; + ch.changed = true; + return best; +} + +bool parse_pair(const char* s, int& a, int& b) { return std::sscanf(s, "%d:%d", &a, &b) == 2; } + +// comma-separated lists for --windows / --weights: exactly three values each +bool parse_list3(const char* s, double* out) { + return std::sscanf(s, "%lf,%lf,%lf", &out[0], &out[1], &out[2]) == 3; +} + +void usage() { + std::fprintf(stderr, + "usage: tune_decode [--expert GU:DOWN]... [--mmvq TYPE:N_IN:N_OUT]... [--n-embd N] [--n-ff N]\n" + " [--k K] [--hit-frac F] [--pool-mib M] [--rounds R] [--reps N] [--min-gain F]\n" + " [--windows T,T,T] [--weights W,W,W] [--out FILE] [--quiet]\n" + " --expert ggml types of the experts' gate/up and down (default 18:20 = IQ3_XXS:IQ4_NL); repeatable\n" + " --mmvq an i-quant dense matrix iq_mmvq runs in decode; repeatable (none by default)\n" + " --hit-frac share of a token's K experts served from VRAM (default 0.6)\n" + " --pool-mib distinct weights the timings cycle through (default 768; keep it well above 96 MB)\n" + " --windows verify-window widths blended in the timings (default 1,3,5 - the 0.1.27-era MTP mix;\n" + " a wider speculative policy should measure its own, e.g. 2,4,6 for --spec 5)\n" + " --weights the windows' weights, normalized (default 0.2,0.3,0.5)\n" + " --out write the table (only shapes whose winner differs from the default are listed)\n"); +} +} // namespace + +int main(int argc, char** argv) { + Options o; + for (int i = 1; i < argc; ++i) { + const std::string a = argv[i]; + auto next = [&]() -> const char* { + if (i + 1 >= argc) { usage(); std::exit(2); } + return argv[++i]; + }; + if (a == "--expert") { + int gu = 0, dn = 0; + if (!parse_pair(next(), gu, dn)) { usage(); return 2; } + o.experts.push_back({gu, dn}); + } else if (a == "--mmvq") { + long long t = 0, ni = 0, no = 0; + if (std::sscanf(next(), "%lld:%lld:%lld", &t, &ni, &no) != 3) { usage(); return 2; } + o.mmvq.push_back({t, ni, no}); + } else if (a == "--n-embd") o.n_embd = std::atoll(next()); + else if (a == "--n-ff") o.n_ff = std::atoll(next()); + else if (a == "--k") o.k_used = std::atoi(next()); + else if (a == "--hit-frac") o.hit_frac = std::atof(next()); + else if (a == "--pool-mib") o.pool_mib = std::atoi(next()); + else if (a == "--rounds") o.rounds = std::max(3, std::atoi(next())); + else if (a == "--reps") o.reps = std::max(5, std::atoi(next())); + else if (a == "--min-gain") o.min_gain = std::atof(next()); + else if (a == "--windows") { + double w[3]; + if (!parse_list3(next(), w)) { usage(); return 2; } + for (int j = 0; j < 3; ++j) o.windows[j] = (int) w[j]; + } + else if (a == "--weights") { + if (!parse_list3(next(), o.weights)) { usage(); return 2; } + } + else if (a == "--out") o.out = next(); + else if (a == "--quiet") o.quiet = true; + else { usage(); return a == "--help" || a == "-h" ? 0 : 2; } + } + if (o.experts.empty()) o.experts.push_back({18, 20}); + if (o.hit_frac <= 0.0 || o.hit_frac > 1.0 || o.k_used < 1 || o.n_embd % 256 || o.n_ff % 32) { + std::fprintf(stderr, "tune_decode: bad shape or --hit-frac\n"); + return 2; + } + for (int j = 0; j < 3; ++j) + if (o.windows[j] < 1 || o.windows[j] > 64 || o.weights[j] <= 0.0) { + std::fprintf(stderr, "tune_decode: bad --windows/--weights (widths 1..64, weights > 0)\n"); + return 2; + } + { + const double sum = o.weights[0] + o.weights[1] + o.weights[2]; + for (int j = 0; j < 3; ++j) o.weights[j] /= sum; + } + g_opt = &o; + const K::DecodeIdentity id = K::decode_identity(); + std::fprintf(stderr, "tune_decode: %s, HIP runtime %lld, toolchain %s\n", id.arch.c_str(), id.runtime, + id.toolchain.c_str()); + std::fprintf(stderr, "tune_decode: workload windows %d/%d/%d, weights %.2f/%.2f/%.2f\n", o.windows[0], + o.windows[1], o.windows[2], o.weights[0], o.weights[1], o.weights[2]); + if (id.arch != "gfx1100") + std::fprintf(stderr, "tune_decode: note - Strata's HIP backend is validated on gfx1100 only; measuring anyway\n"); + const double need = o.pool_mib * (1.0 + (double) o.mmvq.size()) + 256.0; + if (free_mib() < need) { + std::fprintf(stderr, "tune_decode: %.0f MiB of VRAM free, %.0f needed - stop the Strata server (it holds the expert " + "cache) or lower --pool-mib\n", free_mib(), need); + return 3; + } + hipStream_t s; + HIP_CHECK(hipStreamCreateWithFlags(&s, hipStreamNonBlocking)); + hipEvent_t e0, e1; + HIP_CHECK(hipEventCreate(&e0)); + HIP_CHECK(hipEventCreate(&e1)); + std::mt19937 rng(20261004u); + std::vector rows; + std::ostringstream report; + int refused = 0; + + for (const auto& [gu_t, dn_t] : o.experts) { + if (!K::iq_supported(gu_t) || (dn_t != 20 && dn_t != 23 && dn_t != 42)) { + std::fprintf(stderr, "tune_decode: expert types %d:%d are not native expert formats; skipped\n", gu_t, dn_t); + continue; + } + ExpertBench b(K::native_expert_layout(gu_t, dn_t, o.n_embd, o.n_ff), o.pool_mib); + make_expert_bench(b, o, rng); + // every candidate pair, checked bitwise against the engine's true default before it may be timed + const std::vector gc = K::decode_rows_candidates(K::DecodeKernel::GateUp); + const std::vector dc = K::decode_rows_candidates(K::DecodeKernel::Down); + const int d0 = K::kDefaultExpertRows; + std::vector> cands; + std::vector> ref(3); + for (int w = 0; w < 3; ++w) ref[w] = default_output(b, w, s); + for (int g : gc) + for (int d : dc) { + if (g == d0 && d == d0) { cands.push_back({g, d}); continue; } // the default pair IS the default + // path (timed as such below); comparing it bitwise to itself decides nothing + bool same = true; + for (int w = 0; w < 3 && same; ++w) same = grouped_output(b, w, g, d, s) == ref[w]; + if (same) cands.push_back({g, d}); + else { + ++refused; + std::fprintf(stderr, "tune_decode: experts %d:%d gu=%d down=%d is NOT bitwise equal to the default - " + "excluded (please report this build)\n", gu_t, dn_t, g, d); + } + } + Choice ch; + const std::pair best = choose(cands, {d0, d0}, o, + [&](const std::pair& c, int w) { + if (c.first == d0 && c.second == d0) + return time_default(b, w, o.reps, s, e0, e1); + return time_grouped(b, w, c.first, c.second, o.reps, s, e0, e1); + }, ch); + char line[256]; + std::snprintf(line, sizeof(line), + "experts %d:%d (n_embd %lld, n_ff %lld): default chain %.1f us, best %d/%d %.1f us%s\n", gu_t, + dn_t, (long long) o.n_embd, (long long) o.n_ff, ch.default_us, best.first, best.second, + ch.changed ? ch.confirm_best_us : ch.best_us, + ch.changed ? " (confirmed)" : " (kept the default)"); + report << line; + if (!o.quiet) std::fputs(line, stderr); + if (best.first != d0) rows.push_back({K::DecodeKernel::GateUp, gu_t, o.n_embd, o.n_ff, best.first}); + if (best.second != d0) rows.push_back({K::DecodeKernel::Down, dn_t, o.n_ff, o.n_embd, best.second}); + std::printf("RESULT experts %d %d %lld %lld default_us=%.2f chosen=%d/%d chosen_us=%.2f changed=%d\n", gu_t, dn_t, + (long long) o.n_embd, (long long) o.n_ff, ch.default_us, best.first, best.second, + ch.changed ? ch.confirm_best_us : ch.default_us, ch.changed ? 1 : 0); + } + + for (const auto& m : o.mmvq) { + const int t = (int) m[0], ni = (int) m[1], no = (int) m[2]; + if (!K::iq_supported(t) || ni % 256 || no <= 0) { + std::fprintf(stderr, "tune_decode: mmvq %d:%d:%d is not an i-quant shape; skipped\n", t, ni, no); + continue; + } + MmvqBench b(t, ni, no, o.pool_mib); + make_mmvq_bench(b, rng); + const int d0 = K::kDefaultMmvqRows; + std::vector cands; + std::vector> ref(3); + for (int w = 0; w < 3; ++w) ref[w] = mmvq_output(b, w, d0, s); + for (int r : K::decode_rows_candidates(K::DecodeKernel::Mmvq)) { + bool same = true; + for (int w = 0; w < 3 && same; ++w) same = mmvq_output(b, w, r, s) == ref[w]; + if (same) cands.push_back(r); + else { + ++refused; + std::fprintf(stderr, "tune_decode: mmvq %d:%d:%d rows=%d is NOT bitwise equal to the default - excluded\n", + t, ni, no, r); + } + } + Choice ch; + const int best = choose(cands, d0, o, [&](int r, int w) { return time_mmvq(b, w, r, o.reps, s, e0, e1); }, ch); + char line[256]; + std::snprintf(line, sizeof(line), "mmvq %d %dx%d: default %d %.1f us, best %d %.1f us%s\n", t, ni, no, d0, + ch.default_us, best, ch.changed ? ch.confirm_best_us : ch.best_us, + ch.changed ? " (confirmed)" : " (kept the default)"); + report << line; + if (!o.quiet) std::fputs(line, stderr); + if (best != d0) rows.push_back({K::DecodeKernel::Mmvq, t, ni, no, best}); + std::printf("RESULT mmvq %d %d %d default_us=%.2f chosen=%d chosen_us=%.2f changed=%d\n", t, ni, no, ch.default_us, + best, ch.changed ? ch.confirm_best_us : ch.default_us, ch.changed ? 1 : 0); + } + + std::ostringstream table; + table << "# made by tools/hip/tune_decode on this GPU; only shapes whose measured winner beat the default by more\n" + << "# than " << o.min_gain * 100.0 << "% (twice, interleaved) are listed - every listed variant was checked " + << "bitwise equal to the default\n"; + { + std::istringstream rep(report.str()); + std::string l; + while (std::getline(rep, l)) table << "# " << l << "\n"; + } + table << K::DecodeTuningTable::header(id) << "\n"; + for (const K::DecodeTuningRow& r : rows) table << K::DecodeTuningTable::line(r) << "\n"; + if (!o.out.empty()) { + std::ofstream f(o.out); + if (!(f << table.str())) { + std::fprintf(stderr, "tune_decode: cannot write %s\n", o.out.c_str()); + return 1; + } + std::fprintf(stderr, "tune_decode: wrote %s (%zu shapes changed)\n", o.out.c_str(), rows.size()); + } else { + std::fputs(table.str().c_str(), stdout); + } + std::printf("SUMMARY changed=%zu refused=%d\n", rows.size(), refused); + (void) hipEventDestroy(e0); + (void) hipEventDestroy(e1); + (void) hipStreamDestroy(s); + return refused ? 4 : 0; +} diff --git a/tools/test_calibrate.py b/tools/test_calibrate.py index 61adf0ca41..037620a435 100644 --- a/tools/test_calibrate.py +++ b/tools/test_calibrate.py @@ -113,7 +113,7 @@ def test_run_measures_with_the_servers_args(self): (tok / "token_type.json").write_text(json.dumps([1, 1])) seen = [] saved = CAL.measure - CAL.measure = lambda args, ids_list, start_engine, say=print, extra=(): seen.append(args) or {} + CAL.measure = lambda *a, **k: seen.append(a[0]) or {} fake = type("ST", (), {"Tokenizer": lambda *a: type("T", (), {"encode": lambda s, t, **k: [0]})()}) try: with mock.patch.dict(sys.modules, {"strata_tokenizer": fake}): @@ -154,6 +154,108 @@ def test_pick(self): self.assertEqual(CAL.pick({}, "a"), "a") +class ArgsEngine(FakeEngine): + """Speed also depends on the drafting flags it was started with (--spec, --suffix-draft) and records every + per-request tune key it was sent.""" + + def __init__(self, args, speed, info_workers=6, starts=None, sent=None): + super().__init__(args, lambda f, p, w: 0.0, info_workers, starts) + self.full_speed = speed + self.sent = sent if sent is not None else [] + self.spec = int(CAL.arg_value(args, "--spec") or 4) + self.sfx = int(CAL.arg_value(args, "--suffix-draft") or 3) + + def generate(self, ids, max_new, sampling, cancel): + tune = sampling.get("strata_tune") or {} + self.sent.append(dict(tune)) + rate = self.full_speed(tune.get("pcie_frac"), tune.get("spec_min_p", self.info["spec_min_p"]), self.workers, + self.spec, self.sfx) + for _ in range(max_new): + yield 1 + self.last = {"generated": max_new, "decode_ms": max_new / rate * 1000.0} + + +HIP_BASE = ["--pack", "p", "--spec", "4", "--spec-min-p", "0.5", "--max-context", "8192", "--pcie-frac", "0", + "--mmap-experts", "--resident-cpu-experts", "--adapt-every", "0"] + + +class HipAndDraft(unittest.TestCase): + def run_with(self, speed, base=HIP_BASE, backend="hip", draft=False, workers=6): + starts, sent = [], [] + res = CAL.measure(base, [[1, 2, 3]] * 3, lambda a: ArgsEngine(a, speed, workers, starts, sent), + say=lambda *_: None, backend=backend, draft=draft) + return res, starts, sent + + def test_hip_keeps_the_pcie_share(self): + res, starts, sent = self.run_with(lambda f, p, w, s, x: 50.0) + self.assertEqual(res["settings"], {}) + self.assertTrue(all("pcie_frac" not in t for t in sent)) # never swept, never overridden per request + self.assertTrue(all(CAL.arg_value(a, "--pcie-frac") == "0" for a in starts)) + self.assertFalse(res["report"]["pcie_swept"]) + + def test_hip_finds_a_higher_draft_floor(self): + # a CPU-bound hybrid: every extra draft row costs more, so a confident floor wins + res, _, _ = self.run_with(lambda f, p, w, s, x: 60.0 if abs(p - 0.8) < 1e-6 else 50.0) + self.assertEqual(res["settings"], {"--spec-min-p": "0.80"}) + + def test_hip_apply_keeps_pcie_and_the_resident_flags(self): + a = CAL.apply(HIP_BASE, {"--spec-min-p": "0.70"}, backend="hip") + self.assertEqual(CAL.arg_value(a, "--pcie-frac"), "0") + self.assertEqual(CAL.arg_value(a, "--spec-min-p"), "0.70") + self.assertIn("--resident-cpu-experts", a) + b = CAL.apply(a, {}, backend="hip") + self.assertEqual(CAL.arg_value(b, "--pcie-frac"), "0") + self.assertEqual(CAL.arg_value(b, "--spec-min-p"), "0.5") + # NVIDIA's apply still returns the PCIe share to the engine's own choice + self.assertIsNone(CAL.arg_value(CAL.apply(HIP_BASE, {}), "--pcie-frac")) + + def test_draft_finds_window_and_lookup(self): + def speed(f, p, w, spec, sfx): + return 50.0 + (8.0 if spec == 3 else 0.0) + (6.0 if sfx == 6 else 0.0) + res, starts, _ = self.run_with(speed, draft=True) + self.assertEqual(res["settings"].get("--spec"), "3") + self.assertEqual(res["settings"].get("--suffix-draft"), "6") + self.assertNotIn("--pcie-frac", res["settings"]) + self.assertTrue(res["report"]["draft"]) + # the sweep, three windows, three lookup lengths, then the workers on top of the winners + self.assertEqual(len(starts), 1 + len(CAL.SPECS) + len(CAL.SUFFIX_DRAFTS) + len(CAL.worker_candidates(6))) + last = starts[-1] + self.assertEqual(CAL.arg_value(last, "--spec"), "3") + self.assertEqual(CAL.arg_value(last, "--suffix-draft"), "6") + + def test_draft_default_kept_when_flat(self): + res, starts, _ = self.run_with(lambda f, p, w, s, x: 50.0, draft=True) + self.assertEqual(res["settings"], {}) + self.assertTrue(all(CAL.arg_value(a, "--suffix-draft") in (None, "0", "6") for a in starts)) + + def test_draft_apply_resets_old_draft_values(self): + a = CAL.apply(HIP_BASE, {"--spec": "3", "--suffix-draft": "0"}, backend="hip") + self.assertEqual(CAL.arg_value(a, "--spec"), "3") + self.assertEqual(CAL.arg_value(a, "--suffix-draft"), "0") + b = CAL.apply(a, {}, backend="hip", draft=True) # a new draft calibration found the defaults + self.assertEqual(CAL.arg_value(b, "--spec"), "4") + self.assertIsNone(CAL.arg_value(b, "--suffix-draft")) + c = CAL.apply(a, {}, backend="hip") # one without the draft sweep leaves them + self.assertEqual(CAL.arg_value(c, "--spec"), "3") + + def test_load_prompts(self): + with tempfile.TemporaryDirectory() as d: + j = Path(d) / "p.json" + j.write_text(json.dumps(["fix this", "", "and this"])) + self.assertEqual(CAL.load_prompts(j), ["fix this", "and this"]) + t = Path(d) / "p.txt" + t.write_text("first prompt\nline two\n---\nsecond\n") + self.assertEqual(CAL.load_prompts(t), ["first prompt\nline two", "second"]) + bad = Path(d) / "bad.json" + bad.write_text(json.dumps({"a": 1})) + with self.assertRaises(ValueError): + CAL.load_prompts(bad) + empty = Path(d) / "empty.txt" + empty.write_text("\n---\n") + with self.assertRaises(ValueError): + CAL.load_prompts(empty) + + class SetupIntegration(unittest.TestCase): def setUp(self): self.tmp = tempfile.TemporaryDirectory() @@ -192,6 +294,24 @@ def test_calibrate_config_writes_and_remembers(self): other = dict(written, args=CAL.with_arg(written["args"], "--max-context", "131072")) self.assertIsNone(self.S.saved_calibration(other)) + def test_hip_calibration_keeps_pcie_and_remembers_draft(self): + cfg_path = Path(self.tmp.name) / "strata-iq3_xxs.json" + cfg = {"exe": "x", "args": list(HIP_BASE), "model_name": "qwen3.8-flash-next-iq3_xxs", "backend": "hip"} + cfg_path.write_text(json.dumps(cfg)) + seen = {} + + def fake_run(c, say=print, start_engine=None, prompts=None, draft=False, sweep_pcie=None): + seen.update(draft=draft, prompts=prompts) + return {"settings": {"--spec-min-p": "0.70", "--spec": "3"}, "report": {"tok_s": 74.0}} + CAL.run = fake_run + self.assertTrue(self.S.calibrate_config(cfg_path, draft=True)) + self.assertTrue(seen["draft"]) + written = json.loads(cfg_path.read_text()) + self.assertEqual(CAL.arg_value(written["args"], "--pcie-frac"), "0") + self.assertEqual(CAL.arg_value(written["args"], "--spec"), "3") + self.assertEqual(CAL.arg_value(written["args"], "--spec-min-p"), "0.70") + self.assertTrue(self.S.saved_calibration(written)["draft"]) + def test_failed_calibration_keeps_defaults(self): cfg_path = Path(self.tmp.name) / "strata-q2_0.json" cfg_path.write_text(json.dumps({"exe": "x", "args": list(BASE), "model_name": "m"}))