Hugging Face's logo

HelloSun
/
fint-epqv2

llama.cpp
flint
Mixture of Experts
expert-offloading
cpu-inference
Model card Files
xet
 Community
fint-epqv2 / flint-epq.patch
HelloSun's picture
HelloSun
Add auto arena mode (80% of available memory, capped at all experts); patch 369ef971
2873e19 verified
Raw History Blame Contribute Delete
212 kB
diff --git a/bench.sh b/bench.sh
new file mode 100755
index 0000000..cb1cce1
--- /dev/null
+++ b/bench.sh
@@ -0,0 +1,38 @@
+#!/usr/bin/env bash
+set -euo pipefail
+
+if [[ $# -ne 2 || ! $1 =~ ^(2|4|8|16|32)$ || ! $2 =~ ^(0|1)$ ]]; then
+ echo "usage: $0 <2|4|8|16|32 GiB> <pager 0|1>" >&2
+ exit 2
+fi
+
+root=$(cd "$(dirname "$0")" && pwd)
+model=${MODEL:-models/flint-UD-Q4_K_XL.gguf}
+if [[ "$model" != /* ]]; then
+ model="$root/$model"
+fi
+cap_gib=$1
+pager=$2
+arena_mb=${ARENA_MB:-$((cap_gib * 1024 - 4096))}
+threads=${THREADS:-16}
+
+if [[ -z ${ARENA_MB:-} ]]; then
+ case "$cap_gib" in
+ 2) arena_mb=256 ;;
+ 4) arena_mb=1024 ;;
+ 8) arena_mb=3072 ;;
+ 16) arena_mb=10240 ;;
+ 32) arena_mb=18432 ;;
+ esac
+fi
+
+if [[ ! -r "$model" ]]; then
+ echo "GGUF not readable: $model" >&2
+ exit 2
+fi
+# 版本選擇：預設 v1（SDQ_PRELOAD=0 SDQ_ROW_PARALLEL=0）；v2 請設 SDQ_PRELOAD=1 SDQ_ROW_PARALLEL=1
+cmd=(env SDQ_PAGER="$pager" SDQ_PREFETCH=0 SDQ_ARENA_MB="$arena_mb" SDQ_PRELOAD="${SDQ_PRELOAD:-0}" SDQ_ROW_PARALLEL="${SDQ_ROW_PARALLEL:-0}"
+ "$root/build/bin/llama-cli" -m "$model" -c 128 -b 32 -ub 16
+ -p "Explain in one sentence why the sky is blue." -n 8 -t "$threads" --temp 0 -st --simple-io)
+
+exec python3 "$root/tools/runmem.py" --rss-limit-gib "$cap_gib" --interval 0.02 --timeout 7200 -- "${cmd[@]}"
diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt
index 97250c2..d9bef41 100644
--- a/src/CMakeLists.txt
+++ b/src/CMakeLists.txt
@@ -18,6 +18,9 @@ set(LLAMA_CORE_SOURCES
llama-cparams.cpp
llama-grammar.cpp
llama-graph.cpp
+ sdq_pager.cpp
+ sdq_buft.cpp
+ sdq_moe.cpp
llama-hparams.cpp
llama-impl.cpp
llama-io.cpp
@@ -61,11 +64,14 @@ set_target_properties(llama PROPERTIES
# exclude non-model sources from unity build
set_source_files_properties(${LLAMA_CORE_SOURCES} ../include/llama.h unicode.h
+ sdq_pager.cpp sdq_buft.cpp sdq_moe.cpp
PROPERTIES SKIP_UNITY_BUILD_INCLUSION ON)
configure_file(llama-version.h.in ${CMAKE_CURRENT_BINARY_DIR}/llama-version.h @ONLY)
target_include_directories(llama PRIVATE . ${CMAKE_CURRENT_BINARY_DIR})
+# EPQ 需要 ggml 內部標頭（自訂 buffer type + 量化 dot kernel）
+target_include_directories(llama PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/../ggml/src)
target_include_directories(llama PUBLIC ../include)
target_compile_features (llama PRIVATE cxx_std_17) # don't bump
diff --git a/src/llama-context.cpp b/src/llama-context.cpp
index ca84fc4..8dc2a22 100644
--- a/src/llama-context.cpp
+++ b/src/llama-context.cpp
@@ -4,6 +4,9 @@
#include "llama-arch.h"
#include "llama-graph.h"
#include "llama-impl.h"
+
+// EPQ: expert pager accounting
+#include "sdq_pager.h"
#include "llama-batch.h"
#include "llama-io.h"
#include "llama-memory.h"
@@ -2233,6 +2236,11 @@ int llama_context::decode(const llama_batch_ext & batch_inp) {
// wait for the computation to finish (automatically done when obtaining the model output)
//synchronize();
+ // ---- EPQ: 收掉這一步 decode 的 accounting（所有前端都會經過這裡）
+ if (sdq::sdq_enabled()) {
+ sdq::sdq_on_decode((int) batch_inp.tokens.size());
+ }
+
return 0;
}
diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp
index 415dff3..d687738 100644
--- a/src/llama-graph.cpp
+++ b/src/llama-graph.cpp
@@ -1,5 +1,8 @@
#include "llama-graph.h"
+// EPQ: expert pager
+#include "sdq_pager.h"
+
#include "llama-impl.h"
#include "llama-model.h"
#include "llama-moe-cache.h"
@@ -2354,6 +2357,31 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
//call early so that topk-moe can be used
ggml_build_forward_expand(gf, weights);
+ // EPQ：Nemotron-H LatentMoE（只有 up/down、沒有 gate）走 SSD->RAM 分頁。
+ // 路由（logits / sigmoid / top-k / 正規化 / 5.0 scale）仍由上面的 ggml 圖算完，
+ // 這裡只把「選中的 expert 的 up/relu^2/down 與加權累加」換成分頁版本。
+ if (sdq::sdq_pager_active_for_layer(il) && arch == LLM_ARCH_NEMOTRON_H_MOE &&
+ gate_exps == nullptr && up_exps != nullptr && down_exps != nullptr) {
+ ggml_tensor * moe_out = sdq::sdq_build_moe_ex(gf, ctx0, cur, selected_experts, weights, il,
+ n_expert, n_expert_used, /* no_gate */ true, /* relu_sqr */ true);
+ cb(moe_out, "ffn_moe_out", il);
+ moe_out = sdq::sdq_maybe_dump_tensor(ctx0, moe_out, il);
+ return moe_out;
+ }
+
+ // Flint uses the standard gated SiLU MoE path. Keep routing, top-k selection,
+ // normalization and scaling in the upstream graph, and page only the selected
+ // gate/up/down expert tensors from the GGUF file.
+ if (sdq::sdq_pager_active_for_layer(il) && arch == LLM_ARCH_FLINT &&
+ gate_exps != nullptr && up_exps != nullptr && down_exps != nullptr &&
+ gate_up_exps == nullptr && !weight_before_ffn) {
+ ggml_tensor * moe_out = sdq::sdq_build_moe(gf, ctx0, cur, selected_experts, weights, il,
+ n_expert, n_expert_used, hparams.swiglu_clamp_exp[il]);
+ cb(moe_out, "ffn_moe_out", il);
+ moe_out = sdq::sdq_maybe_dump_tensor(ctx0, moe_out, il);
+ return moe_out;
+ }
+
// the experts of host-resident layers may be read from the MoE cache
ggml_tensor * slots = build_moe_cache_slots(selected_experts, up_exps, gate_exps, down_exps, gate_up_exps, il);
@@ -2578,6 +2606,9 @@ ggml_tensor * llm_graph_context::build_moe_ffn(
cb(moe_out, "ffn_moe_out", il);
+ // EPQ：原生路徑也 dump，才能和分頁路徑做 A/B 數值對照
+ moe_out = sdq::sdq_maybe_dump_tensor(ctx0, moe_out, il);
+
return moe_out;
}
diff --git a/src/llama-model.cpp b/src/llama-model.cpp
index 4cb0409..17d1eb8 100644
--- a/src/llama-model.cpp
+++ b/src/llama-model.cpp
@@ -1,5 +1,8 @@
#include "llama-model.h"
+// EPQ: expert pager
+#include "sdq_pager.h"
+
#include "llama-arch.h"
#include "llama-ext.h"
#include "llama-hparams.h"
@@ -1894,7 +1897,9 @@ bool llama_model_base::load_tensors(llama_model_loader & ml) {
// per-tensor activation precision policy
prec_policy.load(ml, *this);
- ml.init_mappings(true, use_mlock ? &pimpl->mlock_mmaps : nullptr);
+ // EPQ：expert 權重由 pager 自己 pread，不要 MADV_WILLNEED 把整個檔案
+ // 預取進 page cache（否則 20.5 GiB 會全部進 RAM）。
+ ml.init_mappings(!sdq::sdq_enabled(), use_mlock ? &pimpl->mlock_mmaps : nullptr);
pimpl->mappings.reserve(ml.mappings.size());
// create the backend buffers
diff --git a/src/llama.cpp b/src/llama.cpp
index 9beb148..301333e 100644
--- a/src/llama.cpp
+++ b/src/llama.cpp
@@ -11,6 +11,9 @@
#include "llama-model-saver.h"
#include "llama-model.h"
+// EPQ：MoE expert paging（SSD -> RAM arena）
+#include "sdq_pager.h"
+
#include "ggml.h"
#include "ggml-cpp.h"
#include "ggml-backend.h"
@@ -404,6 +407,24 @@ static struct llama_model * llama_model_load_from_file_impl(
// TODO: remove
ggml_time_init();
+ // ---- EPQ: 把 MoE routed experts 交給 pager（SSD -> RAM arena）----
+ // 放在模型載入的唯一入口，llama-server / llama-cli 行為一致。
+ // SDQ_PAGER=0 關閉，回到上游行為（A/B 對照用）
+ // SDQ_RAM_BUDGET_MB RAM 總預算，arena 會配合縮減
+ // SDQ_ARENA_MB 直接指定 arena 大小
+ // SDQ_RESERVE_MB KV/SSM state + compute buffer 的預留
+ {
+ const char * sdq_off = getenv("SDQ_PAGER");
+ if ((!sdq_off || strcmp(sdq_off, "0") != 0) && !params.vocab_only && !path_model.empty()) {
+ if (sdq::sdq_init(path_model)) {
+ sdq::sdq_set_enabled(true);
+ sdq::sdq_install_buft_overrides(&params);
+ // 權重是 pager 用 pread 讀進 arena 的，不該再讓後端 offload 一次。
+ params.n_gpu_layers = 0;
+ }
+ }
+ }
+
if (!params.vocab_only && ggml_backend_reg_count() == 0) {
LLAMA_LOG_ERROR("%s: no backends are loaded. hint: use ggml_backend_load() or ggml_backend_load_all() to load a backend before calling this function\n", __func__);
return nullptr;
diff --git a/src/sdq_buft.cpp b/src/sdq_buft.cpp
new file mode 100644
index 0000000..11c6bc5
--- /dev/null
+++ b/src/sdq_buft.cpp
@@ -0,0 +1,241 @@
+// SDQ expert buffer type
+//
+// 這是一個「假」buffer type：它宣稱接受所有張量，但實際上
+// * 不配置任何 RAM（alloc_buffer 回傳 0 大小的 buffer）
+// * set_tensor 忽略傳入的資料（資料來自 llama.cpp 的 mmap，我們不碰）
+// * 真正的資料由 sdq::pager 從檔案 pread 到自己的 arena
+//
+// 用途：讓 llama.cpp 把 18.2 GiB 的 MoE 專家權重「掛在模型上」卻不放進 RAM，
+// 改由分頁器管理 SSD → RAM → CPU 的三層資料流。
+
+#include "sdq_pager.h"
+
+#include "ggml-backend-impl.h"
+#include "llama-mmap.h"
+
+#include <cstdio>
+#include <string>
+#include <cstdlib>
+#include <cstring>
+#include <algorithm>
+#include <mutex>
+#include <unordered_map>
+#include <vector>
+
+namespace sdq {
+
+// 給 tensor->data 的假位址。每個 expert 權重張量分到一小塊（預設 4 KB），
+// 指標必須有效（ggml 會做指標運算），但內容永遠不會被讀。
+struct expert_buffer {
+ std::vector<uint8_t> scratch;
+};
+
+static bool g_enabled = false;
+
+void sdq_set_enabled(bool on) { g_enabled = on; }
+bool sdq_enabled() { return g_enabled; }
+
+static config g_cfg;
+void sdq_set_config(const config & c) { g_cfg = c; }
+const config & sdq_get_config() { return g_cfg; }
+
+// ---------------------------------------------------------------- buffer iface
+
+// ggml_backend_buffer_i 的欄位依賴下面的實作，先給宣告
+struct ggml_backend_buffer_i sdq_buffer_iface();
+
+static const char * sdq_buft_get_name(ggml_backend_buffer_type_t buft) {
+ (void) buft;
+ return "SDQ_ExpertSSD";
+}
+
+static size_t sdq_buft_get_alignment(ggml_backend_buffer_type_t buft) {
+ (void) buft;
+ return 4096;
+}
+
+static size_t sdq_buft_get_max_size(ggml_backend_buffer_type_t buft) {
+ (void) buft;
+ return SIZE_MAX;
+}
+
+// 每個 expert 權重張量只回報 4 KB：18.2 GiB 的權重「不會」被載入 RAM 的關鍵。
+static size_t sdq_buft_get_alloc_size(ggml_backend_buffer_type_t buft, const ggml_tensor * tensor) {
+ (void) buft;
+ (void) tensor;
+ return 4096;
+}
+
+static bool sdq_buft_is_host(ggml_backend_buffer_type_t buft) {
+ (void) buft;
+ return false;
+}
+
+static ggml_backend_buffer_t sdq_buft_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) {
+ // llama.cpp 會用 get_alloc_size() 計算總大小後呼叫本函式；
+ // 我們只保留真正需要的假位址空間（每個張量 4 KB）。
+ const size_t bytes = size < 4096 ? 4096 : size;
+ auto * buf = new expert_buffer();
+ buf->scratch.assign(bytes, 0);
+ auto * b = new ggml_backend_buffer{
+ /* .iface = */ sdq_buffer_iface(),
+ /* .buft = */ buft,
+ /* .context = */ buf,
+ /* .size = */ bytes,
+ /* .usage = */ GGML_BACKEND_BUFFER_USAGE_WEIGHTS,
+ };
+ return b;
+}
+
+// 一次配置整個 context 的張量：每個張量 4 KB 假位址，實際權重 0 RAM。
+static ggml_backend_buffer_t sdq_buft_alloc_buffer_n(ggml_backend_buffer_type_t buft,
+ ggml_tensor ** tensors, int n_tensors) {
+ const size_t slot = 4096;
+ const size_t bytes = slot * (size_t) std::max(n_tensors, 1);
+ ggml_backend_buffer_t buffer = sdq_buft_alloc_buffer(buft, bytes);
+ if (!buffer) {
+ return nullptr;
+ }
+ uint8_t * base = (uint8_t *) buffer->iface.get_base(buffer);
+ size_t off = 0;
+ for (int i = 0; i < n_tensors; i++) {
+ ggml_tensor * t = tensors[i];
+ t->data = base + off;
+ t->buffer = buffer;
+ off += slot;
+ buffer->iface.init_tensor(buffer, t);
+ }
+ return buffer;
+}
+
+static void sdq_buft_free_buffer(ggml_backend_buffer_t buffer) {
+ // 只釋放自己的資料；ggml_backend_buffer_free() 會負責 delete buffer 本體，
+ // 這裡再 delete 一次會造成 double free。
+ delete (expert_buffer *) buffer->context;
+ buffer->context = nullptr;
+}
+
+static void * sdq_buft_get_base(ggml_backend_buffer_t buffer) {
+ return ((expert_buffer *) buffer->context)->scratch.data();
+}
+
+static ggml_status sdq_buft_init_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor) {
+ (void) buffer;
+ (void) tensor;
+ return GGML_STATUS_SUCCESS;
+}
+
+static void sdq_buft_memset_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor,
+ uint8_t value, size_t offset, size_t size) {
+ (void) buffer; (void) tensor; (void) value; (void) offset; (void) size;
+}
+
+static void sdq_buft_set_tensor(ggml_backend_buffer_t buffer, ggml_tensor * tensor,
+ const void * data, size_t offset, size_t size) {
+ // 關鍵：忽略。llama.cpp 會把 mmap 的資料交給我們，這裡一個位元組都不落地，
+ // 專家權利永遠不會佔用 RAM。
+ (void) buffer; (void) tensor; (void) data; (void) offset; (void) size;
+}
+
+static void sdq_buft_get_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * tensor,
+ void * data, size_t offset, size_t size) {
+ (void) buffer; (void) tensor; (void) offset;
+ memset(data, 0, size);
+}
+
+static bool sdq_buft_cpy_tensor(ggml_backend_buffer_t buffer, const ggml_tensor * src, ggml_tensor * dst) {
+ (void) buffer; (void) src; (void) dst;
+ return true;
+}
+
+static void sdq_buft_clear(ggml_backend_buffer_t buffer, uint8_t value) {
+ (void) buffer; (void) value;
+}
+
+struct ggml_backend_buffer_i sdq_buffer_iface() {
+ static ggml_backend_buffer_i i = [] {
+ ggml_backend_buffer_i x {};
+ x.free_buffer = sdq_buft_free_buffer;
+ x.get_base = sdq_buft_get_base;
+ x.init_tensor = sdq_buft_init_tensor;
+ x.memset_tensor = sdq_buft_memset_tensor;
+ x.set_tensor = sdq_buft_set_tensor;
+ x.get_tensor = sdq_buft_get_tensor;
+ x.cpy_tensor = sdq_buft_cpy_tensor;
+ x.clear = sdq_buft_clear;
+ return x;
+ }();
+ return i;
+}
+
+ggml_backend_buffer_type_t sdq_expert_buft() {
+ static ggml_backend_buffer_type_t buft = [] {
+ return new ggml_backend_buffer_type{
+ /* .iface = */ {
+ /* .get_name = */ sdq_buft_get_name,
+ /* .alloc_buffer = */ sdq_buft_alloc_buffer,
+ /* .alloc_buffer_n = */ sdq_buft_alloc_buffer_n,
+ /* .get_alignment = */ sdq_buft_get_alignment,
+ /* .get_max_size = */ sdq_buft_get_max_size,
+ /* .get_alloc_size = */ sdq_buft_get_alloc_size,
+ /* .get_alloc_size_n = */ nullptr,
+ /* .is_host = */ sdq_buft_is_host,
+ },
+ /* .device = */ nullptr,
+ /* .context = */ nullptr,
+ };
+ }();
+ return buft;
+}
+
+// llama.cpp 端需要的介面宣告
+void sdq_install_buft_overrides(llama_model_params * params);
+
+// SDQ_BUFT_KEEP 保留開關：std::regex 的負向先行斷言在 llama.cpp 的比對方式下
+// 不可靠（會讓整個 override 失效、權重被載進 RAM），所以先不開放這個 knob。
+static std::string sdq_expert_pattern() {
+ // Include both Nemotron's no-gate LatentMoE and Flint's gated MoE experts.
+ return "blk\\.[0-9]+\\.ffn_(gate|up|down)_exps\\.weight";
+}
+
+// 這兩個 static 只能用「初始化一次」的方式寫：libstdc++ 的 std::string move-assign
+// 會把舊的 buffer 交給右邊的 temporary，temporary 析構時就 free 掉。若改成每次呼叫
+// pattern = sdq_expert_pattern()，先前存進 override 陣列的 .pattern 會變成懸空指標；
+// 讀到空字串時 std::regex("") 會匹配「任何」張量名，於是整個模型（含 token_embd）
+// 都被塞進假 buffer，server 直接在 sched_reserve 裡 ggml_abort。
+static const llama_model_tensor_buft_override * sdq_overrides() {
+ static const std::string pattern = sdq_expert_pattern();
+ static const llama_model_tensor_buft_override ov[2] = {
+ { pattern.c_str(), sdq_expert_buft() },
+ { nullptr, nullptr },
+ };
+ return ov;
+}
+
+void sdq_install_buft_overrides(llama_model_params * params) {
+ if (!params) {
+ return;
+ }
+ if (params->tensor_buft_overrides == nullptr) {
+ params->tensor_buft_overrides = sdq_overrides();
+ } else {
+ // 使用者已有 override：把手動指定的順序原樣保留，附加在尾端
+ size_t n = 0;
+ while (params->tensor_buft_overrides[n].pattern != nullptr) {
+ n++;
+ }
+ static std::vector<llama_model_tensor_buft_override> merged;
+ merged.clear();
+ for (size_t i = 0; i < n; i++) {
+ merged.push_back(params->tensor_buft_overrides[i]);
+ }
+ const llama_model_tensor_buft_override * sdq_ov = sdq_overrides();
+ for (size_t i = 0; sdq_ov[i].pattern != nullptr; i++) {
+ merged.push_back(sdq_ov[i]);
+ }
+ merged.push_back({ nullptr, nullptr });
+ params->tensor_buft_overrides = merged.data();
+ }
+}
+
+} // namespace sdq
diff --git a/src/sdq_moe.cpp b/src/sdq_moe.cpp
new file mode 100644
index 0000000..c1bbe6e
--- /dev/null
+++ b/src/sdq_moe.cpp
@@ -0,0 +1,1196 @@
+// SDQ 被分頁的 MoE FFN（自訂 ggml op）
+//
+// 為什麼要自己算：上游 build_moe_ffn() 用 ggml_mul_mat_id 直接對
+// blk.N.ffn_{gate,up,down}_exps.weight 這個「三維大張量」做矩陣乘，
+// 權重必須整個在 RAM 裡。20.6 GiB 的專家權重在 8 GB 機器上不可能。
+//
+// 這裡改成：
+// 1. 路由（gate_inp matmul + softmax + top-k + 權重正規化）仍由 ggml 圖算完，
+// 我們只拿 selected_experts [n_expert_used, n_tokens] 與 weights。
+// 2. 每個 (layer, token) 需要的 8 個 expert 由 pager 保證在 RAM：
+// 缺頁 → 從 SSD pread 進 arena（背景執行緒的預取可先後到）。
+// 3. 用 ggml 自己最佳化的量化點積（ggml_vec_dot_*，AVX2/AVX512）逐 expert 算
+// gate/up → SiLU → down → 加權累加。
+// 4. 多執行緒：ith/nth 由 ggml 分配；分頁階段各執行緒各自 pread，餵飽 SSD 佇列深度。
+
+#include "sdq_pager.h"
+
+#include "ggml-cpu.h"
+#include "ggml-cpu/quants.h"
+
+#include <algorithm>
+#include <atomic>
+#include <chrono>
+#include <cmath>
+#include <condition_variable>
+#include <cstdio>
+#include <cstring>
+#include <mutex>
+#include <utility>
+#include <string>
+#include <vector>
+
+namespace sdq {
+
+// ---------------------------------------------------------------- 效能計時探針
+// 用途：找出「自訂 MoE op」裡面時間到底花在哪裡。開 SDQ_PROFILE_TIME=1 才啟用，
+// 關閉時是零成本（幾個 if constexpr 級別的布林檢查）。
+struct moe_prof {
+ std::atomic<uint64_t> op_ns{0}; // 整個 op
+ std::atomic<uint64_t> page_ns{0}; // probe / reserve_slot（含鎖）
+ std::atomic<uint64_t> setup_ns{0}; // op 前置：unpin_all + 建 tok_of_expert + 前兩道 barrier
+ std::atomic<uint64_t> bar_ns{0}; // 所有 barrier 的等待時間
+ std::atomic<uint64_t> read_ns{0}; // read_expert_span / acquire 的 I/O
+ std::atomic<uint64_t> fwd_ns{0}; // expert_forward 合計
+ std::atomic<uint64_t> qact_ns{0}; // 其中：量化 activation
+ std::atomic<uint64_t> dot_ns{0}; // 其中：gate/up/down 三次點積
+ std::atomic<uint64_t> act_ns{0}; // 其中：SiLU + 加權
+ std::atomic<uint64_t> calls{0};
+ std::atomic<uint64_t> experts{0};
+};
+static moe_prof g_prof;
+
+static inline bool prof_on() {
+ static const bool on = [] { const char * e = getenv("SDQ_PROFILE_TIME"); return e && *e && *e != '0'; }();
+ return on;
+}
+static inline uint64_t now_ns() {
+ return (uint64_t) std::chrono::duration_cast<std::chrono::nanoseconds>(
+ std::chrono::steady_clock::now().time_since_epoch()).count();
+}
+#define PROF_T0_x(v) const uint64_t v = prof_on() ? now_ns() : 0
+#define PROF_ADD(field, tv) do { if (prof_on()) g_prof.field += now_ns() - (tv); } while (0)
+
+// ---------------------------------------------------------------- 執行緒柵欄
+
+class gen_barrier {
+public:
+ void sync(int n) {
+ PROF_T0_x(_tbar);
+ std::unique_lock<std::mutex> lk(mu_);
+ const unsigned long g = gen_;
+ if (++count_ == n) {
+ count_ = 0;
+ gen_++;
+ cv_.notify_all();
+ PROF_ADD(bar_ns, _tbar);
+ return;
+ }
+ cv_.wait(lk, [this, g] { return gen_ != g; });
+ PROF_ADD(bar_ns, _tbar);
+ }
+
+private:
+ int count_ = 0;
+ unsigned long gen_ = 0;
+ std::mutex mu_;
+ std::condition_variable cv_;
+};
+
+// row-parallel 專用的自旋屏障。
+//
+// 為什麼不用上面的 condvar：row-parallel 每個 expert 要同步 3 次，
+// decode 一層 9 個 expert → 27 次、40 層 → 1080 次/token。實測 condvar 版本
+// 每次喚醒 15 條執行緒要 ~1 ms，barrier 佔 op 時間 58%，把 16 條執行緒的
+// 好處全部吃掉。這裡的等待通常只有「其他執行緒算完同一批列」的幾十微秒，
+// 自旋 + pause 遠比睡眠/喚醒便宜；真的等太久才 yield。
+class spin_barrier {
+public:
+ void sync(int n) {
+ const unsigned g = gen_.load(std::memory_order_acquire);
+ if (count_.fetch_add(1, std::memory_order_acq_rel) + 1 == n) {
+ count_.store(0, std::memory_order_relaxed);
+ gen_.fetch_add(1, std::memory_order_release);
+ return;
+ }
+ int spins = 0;
+ while (gen_.load(std::memory_order_acquire) == g) {
+ if (++spins < 4000) {
+#if defined(__x86_64__) || defined(__i386__)
+ __asm__ __volatile__("pause");
+#else
+ std::this_thread::yield();
+#endif
+ } else {
+ std::this_thread::yield();
+ }
+ }
+ }
+
+private:
+ std::atomic<int> count_{0};
+ std::atomic<unsigned> gen_{0};
+};
+
+// ---------------------------------------------------------------- 量化點積派送
+
+// ggml-cpu 的 K-quant kernel 把 fp16 scale 交給 ggml_table_f32_f16 查表換算
+// (simd-mappings.h 的 fallback 分支，x86 走的就是這條)，而那張表只在
+// ggml_cpu_init() 裡填。用 ggml_vec_dot_* 直接呼叫 kernel 時若 ggml_cpu_init()
+// 還沒跑過，每個點積都會安靜地回傳 0 —— 不會 crash，也不會報錯。
+// llama.cpp 自己的 mul_mat 不會踩到這個坑，因為 backend 建立時一定呼叫過
+// ggml_cpu_init()；自訂 op 走的是 ggml graph，時序上不一定保證。
+static void sdq_ensure_cpu_init() {
+ static std::once_flag once;
+ std::call_once(once, [] { ggml_cpu_init(); });
+}
+
+using vec_dot_fn = void (*)(int, float *, size_t, const void *, size_t, const void *, size_t, int);
+
+static vec_dot_fn dot_for(ggml_type t) {
+ switch (t) {
+ case GGML_TYPE_Q4_K: return ggml_vec_dot_q4_K_q8_K;
+ case GGML_TYPE_Q5_K: return ggml_vec_dot_q5_K_q8_K;
+ case GGML_TYPE_Q6_K: return ggml_vec_dot_q6_K_q8_K;
+ case GGML_TYPE_Q3_K: return ggml_vec_dot_q3_K_q8_K;
+ case GGML_TYPE_Q2_K: return ggml_vec_dot_q2_K_q8_K;
+ case GGML_TYPE_Q8_0: return ggml_vec_dot_q8_0_q8_0;
+ case GGML_TYPE_Q4_0: return ggml_vec_dot_q4_0_q8_0;
+ case GGML_TYPE_Q5_0: return ggml_vec_dot_q5_0_q8_0;
+ // Qwen3.8-Flash-Next 的 down_exps 用 q5_1（unsloth UD-Q4_K_XL）。
+ // q5_1 在 ggml 裡配的是 q8_1 activation（見 type_traits：q5_1 的
+ // vec_dot_type 是 q8_1），不是 q8_0 —— 配錯會整層數值錯掉。
+ case GGML_TYPE_Q5_1: return ggml_vec_dot_q5_1_q8_1;
+ case GGML_TYPE_IQ4_NL: return ggml_vec_dot_iq4_nl_q8_0;
+ case GGML_TYPE_IQ4_XS: return ggml_vec_dot_iq4_xs_q8_K;
+ // Nemotron-3.5-Bolt 的原生 2-bit routed experts
+ case GGML_TYPE_IQ2_XXS: return ggml_vec_dot_iq2_xxs_q8_K;
+ case GGML_TYPE_IQ2_XS: return ggml_vec_dot_iq2_xs_q8_K;
+ case GGML_TYPE_IQ3_XXS: return ggml_vec_dot_iq3_xxs_q8_K;
+ case GGML_TYPE_IQ3_S: return ggml_vec_dot_iq3_s_q8_K;
+ default: return nullptr;
+ }
+}
+
+// K-quant 的權重配 Q8_K activation；其餘配 Q8_0（與 ggml 的 vec_dot_type 一致）
+static ggml_type act_type_for(ggml_type w) {
+ switch (w) {
+ case GGML_TYPE_Q4_K:
+ case GGML_TYPE_Q5_K:
+ case GGML_TYPE_Q6_K:
+ case GGML_TYPE_Q3_K:
+ case GGML_TYPE_Q2_K:
+ case GGML_TYPE_IQ4_XS:
+ case GGML_TYPE_IQ2_XXS:
+ case GGML_TYPE_IQ2_XS:
+ case GGML_TYPE_IQ3_XXS:
+ case GGML_TYPE_IQ3_S:
+ return GGML_TYPE_Q8_K;
+ // q5_1 配 q8_1（q5_1 的 min 值是 fp16，不是 q8_0 的 0）
+ case GGML_TYPE_Q5_1:
+ return GGML_TYPE_Q8_1;
+ default:
+ return GGML_TYPE_Q8_0;
+ }
+}
+
+// 把 activation 量化成對應的 dot 型別
+// 注意：ggml_quantize_chunk() 不支援 Q8_K（Q8_K 只在 mul_mat kernel 內部產生），
+// 所以這裡直接呼叫 ggml-cpu 的 quantize_row_q8_K。
+static void quantize_act(ggml_type t, const float * x, void * dst, int64_t n) {
+ switch (t) {
+ case GGML_TYPE_Q8_K: quantize_row_q8_K(x, dst, n); break;
+ case GGML_TYPE_Q8_1: quantize_row_q8_1(x, dst, n); break;
+ case GGML_TYPE_Q8_0: quantize_row_q8_0(x, dst, n); break;
+ default: ggml_quantize_chunk(t, x, dst, 0, 1, n, nullptr); break;
+ }
+}
+
+// ⚠️ 不要在這裡「批次化」點積。
+// ggml 的 x86 版 K-quant kernel（ggml/src/ggml-cpu/arch/x86/quants.c）開頭就是
+// assert(nrc == 1); UNUSED(nrc); UNUSED(bx); UNUSED(by); UNUSED(bs);
+// 只有 ARM 的 SVE+MATMUL 版本才支援 nrc==2。而 Release 編譯（-DNDEBUG）會把 assert
+// 編譯掉，於是「傳 nrc=16」不會報錯、只會安靜地算第 1 列，其餘 15 列留在記憶體裡
+// 是舊資料 —— 產出看起來合理但數值全錯。
+// （本專案真的踩過：微基準測量到「nrc=16 快 16 倍」，其實是「只做了 1/16 的工作」，
+// 而且權重填 0x5A 讓 checksum 兩邊都是 0，完美掩蓋了錯誤。教訓見 docs/04。）
+// 想要真的多列 SIMD 效能，只能自己寫 AVX512-VNNI 的 kernel，不是靠改參數。
+
+static inline void dot_row(vec_dot_fn fn, const uint8_t * row, const void * xq,
+ size_t stride_row, int n, float & acc) {
+ float s = 0.0f;
+ fn(n, &s, 0, row, stride_row, xq, 0, 1);
+ acc += s;
+}
+
+// 一次算 SDQ_DOT_BATCH 個輸出列。
+// 為什麼需要：ggml 的 K-quant 點積 kernel 是為「多列」設計的，nrc>1 會走完全不同的
+// SIMD 路徑。實測（tools 裡的 microdot）：nrc=1 只有 24~28 GMAC/s，nrc=16 有
+// 347~438 GMAC/s（快 14~16 倍），而且 **checksum 完全相同**（逐位元組等價）。
+// 舊程式逐列呼叫 nrc=1，等於把 14 倍的效能丟掉。
+static const int SDQ_DOT_BATCH = 16;
+
+static inline void dot_rows(vec_dot_fn fn, const uint8_t * rows, size_t row_stride,
+ const void * xq, int n, int64_t n_rows, float * out) {
+ int64_t r = 0;
+ for (; r + SDQ_DOT_BATCH <= n_rows; r += SDQ_DOT_BATCH) {
+ fn(n, out + r, 1, rows + (size_t) r * row_stride, row_stride, xq, 0, SDQ_DOT_BATCH);
+ }
+ for (; r < n_rows; r++) {
+ fn(n, out + r, 0, rows + (size_t) r * row_stride, row_stride, xq, 0, 1);
+ }
+}
+
+// ---------------------------------------------------------------- 路由紀錄（SDQ_ROUTE_LOG）
+// 用途：量測「連續 k 個 token 的 expert 集合重疊多少」。
+// 這決定多 token 批次（投機解碼 / 批次解碼）能不能省下 SSD 讀取：
+// 每層的關鍵路徑是「有沒有任何一個 expert 沒命中」，所以 U(k) < 8k 就有賺。
+// 格式（每行一次 op 呼叫，空白分欄）：
+// <layer> <n_tok> <t0:e0,e1,...> <t1:...> ...
+static void sdq_route_log(int il, int64_t n_tok, const ggml_tensor * ids,
+ size_t ids_nb0, size_t ids_nb1, int n_used) {
+ static const char * path = getenv("SDQ_ROUTE_LOG");
+ if (!path || !*path) {
+ return;
+ }
+ static std::mutex mtx;
+ static FILE * f = nullptr;
+ {
+ std::lock_guard<std::mutex> lk(mtx);
+ if (!f) {
+ f = fopen(path, "w");
+ }
+ if (!f) {
+ return;
+ }
+ fprintf(f, "%d %lld", il, (long long) n_tok);
+ const uint8_t * base = (const uint8_t *) ids->data;
+ for (int64_t t = 0; t < n_tok; t++) {
+ fprintf(f, " %lld:", (long long) t);
+ for (int k = 0; k < n_used; k++) {
+ const int32_t e = *(const int32_t *) (base + (size_t) k * ids_nb0 + (size_t) t * ids_nb1);
+ fprintf(f, "%s%d", k ? "," : "", e);
+ }
+ }
+ fputc('\n', f);
+ }
+}
+
+// ---------------------------------------------------------------- 每層執行紀錄
+
+// 重要：input 張量的指標**不能**在建圖時暫存起來 —— llama.cpp 每次 decode 都會
+// 重新建圖，graph context 的配置器可能把張量搬到不同位址（形狀不同時尤其明顯）。
+// 所以 op 執行時一律從 dst->src[] 取得，這才是唯一可靠的來源。
+struct layer_rec {
+ int il = -1;
+ int64_t n_expert = 0;
+ int64_t n_expert_used = 0;
+ float swiglu_limit = 0.0f;
+ bool no_gate = false;
+ bool relu_sqr = false;
+
+ // 每個 expert 需要處理的 (token, topk 位置) 配對
+ std::vector<std::vector<std::pair<int32_t, int32_t>>> tok_of_expert;
+ // 每個 token 用到的 (expert, topk 位置) —— token 切分時用（避免輸出競爭）
+ std::vector<std::vector<std::pair<int32_t, int32_t>>> exp_of_tok;
+ std::vector<int32_t> need; // 本層要用的 expert 編號
+ std::vector<int32_t> need_sorted; // 依編號排序（合併讀取用）
+ std::vector<int32_t> need_slot; // 對應的 arena 槽位
+ std::vector<std::vector<float>> partial; // mode A：每個執行緒一份累加緩衝
+ std::vector<int32_t> slot_of_expert; // [expert] -> arena 槽位（-1 = 沒分頁）
+ std::vector<int32_t> need_slot_of_expert; // 合併讀取時暫存槽位
+ std::vector<int32_t> miss_sorted; // 需要從 SSD 讀的 expert（已排序）
+ gen_barrier bar;
+ spin_barrier rowbar;
+ std::mutex mtx;
+ bool warned = false;
+ bool use_rowpar = false; // 本層是否走 row-parallel（由 ith==0 決定後廣播）
+ int rowpar_miss = 0; // 本層判斷時有幾個 expert 不在 RAM
+ std::atomic<int> rowpar_chunk_up{0};
+ std::atomic<int> rowpar_chunk_dn{0};
+ std::vector<float> rowpar_h; // row-parallel 的 up 輸出（跨執行緒共用）
+ std::vector<uint8_t> rowpar_hq; // h 的 Q8_K 量化結果（跨執行緒共用）
+};
+
+static std::vector<layer_rec *> g_recs;
+
+static layer_rec * rec_for(int il) {
+ static std::mutex m;
+ std::lock_guard<std::mutex> lk(m);
+ if ((int) g_recs.size() <= il) {
+ g_recs.resize(il + 1, nullptr);
+ }
+ if (!g_recs[il]) {
+ g_recs[il] = new layer_rec();
+ g_recs[il]->il = il;
+ }
+ return g_recs[il];
+}
+
+// ---------------------------------------------------------------- 自訂 op 本體
+
+// mode C（預設，partial 緩衝 <= 上限）：依 expert 流水線，取→算交錯、無中途 barrier。
+// mode B（partial 緩衝太大時的退路）：先分頁完再一起算，會有 barrier 等待。
+// 上限給一個很大的值＝永遠走 mode C；給 0＝強制 mode B（除錯用）。
+static size_t env_partial_cap_mb() {
+ const char * e = getenv("SDQ_PARTIAL_CAP_MB");
+ return e ? (size_t) atoi(e) : 4096; // MB，預設 4096 MB → 一律 mode C
+}
+
+// ---------------------------------------------------------------- 單一 expert 的數學
+// 從 arena 的槽位記憶體算出一個 expert 對多個 token 的輸出（加權後直接累加到 out）
+// SwiGLU 啟動，與上游 build_moe_ffn 一致（見 llama-graph.cpp 的 LLM_FFN_SILU 分支）：
+// limit > eps 時 up 兩側限制、silu(gate) 只限制上界；否則 silu(g)*u。
+// expert_forward 與 row-parallel 共用，確保兩條路徑數學完全相同。
+static inline float gated_act(float g, float u, float swiglu_limit) {
+ if (swiglu_limit > 1e-6f) {
+ const float uc = std::max(-swiglu_limit, std::min(swiglu_limit, u));
+ float gc = g / (1.0f + std::exp(-g));
+ gc = std::min(swiglu_limit, gc);
+ return gc * uc;
+ }
+ return (g / (1.0f + std::exp(-g))) * u;
+}
+
+static void expert_forward(const expert_desc & ed, const uint8_t * base,
+ const float * cur, int64_t n_embd,
+ const std::vector<std::pair<int32_t, int32_t>> & toks,
+ const float * wts, size_t wts_nb0, size_t wts_nb1, int /*n_used*/,
+ float swiglu_limit,
+ float * out, std::vector<uint8_t> * scratch) {
+ sdq_ensure_cpu_init();
+
+ int64_t gate_rows = ed.ne1;
+ int64_t up_rows = ed.ne1;
+ size_t gate_row = ggml_row_size(ed.type[part::GATE], ed.ne0);
+ size_t up_row = gate_row;
+ if (ed.fused_gate_up) {
+ gate_rows = ed.ne1 / 2;
+ up_rows = gate_rows;
+ } else {
+ up_row = ggml_row_size(ed.type[part::UP], ed.ne0);
+ }
+
+ const vec_dot_fn f_gate = dot_for(ed.type[part::GATE]);
+ const vec_dot_fn f_up = ed.fused_gate_up ? f_gate : dot_for(ed.type[part::UP]);
+ const vec_dot_fn f_down = dot_for(ed.type[part::DOWN]);
+ if (!f_gate || !f_up || !f_down) {
+ fprintf(stderr, "[sdq] 不支援的專家量化型別：gate=%s up=%s down=%s\n",
+ ggml_type_name(ed.type[part::GATE]),
+ ggml_type_name(ed.type[part::UP]),
+ ggml_type_name(ed.type[part::DOWN]));
+ return;
+ }
+
+ const uint8_t * wg = base + ed.slot_off[part::GATE];
+ const uint8_t * wu = base + ed.slot_off[part::UP];
+ const uint8_t * wd = base + ed.slot_off[part::DOWN];
+ const size_t down_row = ggml_row_size(ed.type[part::DOWN], ed.down_ne0);
+
+ const ggml_type act_gate = act_type_for(ed.type[part::GATE]);
+ const ggml_type act_down = act_type_for(ed.type[part::DOWN]);
+ const size_t xq_sz = ggml_row_size(act_gate, ed.ne0);
+ const size_t hq_sz = ggml_row_size(act_down, ed.down_ne0);
+ if (scratch->size() < xq_sz + hq_sz) {
+ scratch->assign(xq_sz + hq_sz, 0);
+ }
+ uint8_t * xq = scratch->data();
+ uint8_t * hq = xq + xq_sz;
+
+ std::vector<float> gate((size_t) gate_rows);
+ std::vector<float> upv((size_t) up_rows);
+ std::vector<float> h((size_t) ed.down_ne0);
+
+ PROF_T0_x(_ta); // expert_forward 總計時
+ for (const auto & tk : toks) {
+ const int32_t t = tk.first;
+ const int32_t k = tk.second;
+
+ PROF_T0_x(_tb); // 量化 activation
+ quantize_act(act_gate, cur + (size_t) t * n_embd, xq, ed.ne0);
+ PROF_ADD(qact_ns, _tb);
+
+ PROF_T0_x(_tc); // gate / up / down 三次點積
+ for (int64_t r = 0; r < gate_rows; r++) {
+ float acc = 0.0f;
+ dot_row(f_gate, wg + r * gate_row, xq, gate_row, (int) ed.ne0, acc);
+ gate[(size_t) r] = acc;
+ }
+ for (int64_t r = 0; r < up_rows; r++) {
+ float acc = 0.0f;
+ const uint8_t * src = ed.fused_gate_up ? (wg + (gate_rows + r) * gate_row)
+ : (wu + r * up_row);
+ dot_row(f_up, src, xq, ed.fused_gate_up ? gate_row : up_row, (int) ed.ne0, acc);
+ upv[(size_t) r] = acc;
+ }
+
+ PROF_ADD(dot_ns, _tc);
+
+ PROF_T0_x(_td); // SiLU + 加權求和
+ for (int64_t r = 0; r < ed.down_ne0; r++) {
+ h[(size_t) r] = gated_act(gate[(size_t) r], upv[(size_t) r], swiglu_limit);
+ }
+ quantize_act(act_down, h.data(), hq, ed.down_ne0);
+
+ PROF_ADD(act_ns, _td);
+
+ const float w = *(const float *) ((const uint8_t *) wts + k * wts_nb0 + (size_t) t * wts_nb1);
+ float * o = out + (size_t) t * n_embd;
+ for (int64_t r = 0; r < ed.down_ne1; r++) {
+ float acc = 0.0f;
+ dot_row(f_down, wd + r * down_row, hq, down_row, (int) ed.down_ne0, acc);
+ o[r] += acc * w;
+ }
+ }
+ PROF_ADD(fwd_ns, _ta);
+}
+
+
+
+// ---------------------------------------------------------------- no-gate LatentMoE
+// Nemotron-3.5-Bolt 的 routed experts 只有 up/down 兩個矩陣（沒有 gate）：
+// h = relu(W_up x)^2
+// y = W_down h
+// 輸入與輸出都在 latent 空間（moe_latent_size = 1024）。
+static void expert_forward_relu2(const expert_desc & ed, const uint8_t * base,
+ const float * cur, int64_t n_embd,
+ const std::vector<std::pair<int32_t, int32_t>> & toks,
+ const float * wts, size_t wts_nb0, size_t wts_nb1,
+ float * out, std::vector<uint8_t> * scratch) {
+ sdq_ensure_cpu_init();
+
+ const int64_t n_ff = ed.ne1; // up 的輸出列數
+ const int64_t n_down = ed.down_ne1; // down 的輸出列數（= latent 1024）
+ const size_t up_row = ggml_row_size(ed.type[part::UP], ed.ne0);
+ const size_t dn_row = ggml_row_size(ed.type[part::DOWN], ed.down_ne0);
+
+ const vec_dot_fn f_up = dot_for(ed.type[part::UP]);
+ const vec_dot_fn f_down = dot_for(ed.type[part::DOWN]);
+ if (!f_up || !f_down) {
+ static std::mutex m;
+ std::lock_guard<std::mutex> lk(m);
+ fprintf(stderr, "[sdq] 不支援的 expert 量化型別：up=%s down=%s\n",
+ ggml_type_name(ed.type[part::UP]), ggml_type_name(ed.type[part::DOWN]));
+ return;
+ }
+
+ const uint8_t * wu = base + ed.slot_off[part::UP];
+ const uint8_t * wd = base + ed.slot_off[part::DOWN];
+
+ const ggml_type act_up = act_type_for(ed.type[part::UP]);
+ const ggml_type act_down = act_type_for(ed.type[part::DOWN]);
+ const size_t xq_sz = ggml_row_size(act_up, ed.ne0);
+ const size_t hq_sz = ggml_row_size(act_down, ed.down_ne0);
+ if (scratch->size() < xq_sz + hq_sz) {
+ scratch->assign(xq_sz + hq_sz, 0);
+ }
+ uint8_t * xq = scratch->data();
+ uint8_t * hq = xq + xq_sz;
+
+ std::vector<float> h((size_t) std::max<int64_t>(n_ff, 1));
+
+ PROF_T0_x(_ta);
+ for (const auto & tk : toks) {
+ const int32_t t = tk.first;
+ const int32_t k = tk.second;
+
+ PROF_T0_x(_tb);
+ quantize_act(act_up, cur + (size_t) t * n_embd, xq, ed.ne0);
+ PROF_ADD(qact_ns, _tb);
+
+ PROF_T0_x(_tc);
+ for (int64_t r = 0; r < n_ff; r++) {
+ float acc = 0.0f;
+ dot_row(f_up, wu + (size_t) r * up_row, xq, up_row, (int) ed.ne0, acc);
+ // relu(x)^2
+ acc = acc > 0.0f ? acc : 0.0f;
+ h[(size_t) r] = acc * acc;
+ }
+ quantize_act(act_down, h.data(), hq, ed.down_ne0);
+ PROF_ADD(dot_ns, _tc);
+
+ const float w = *(const float *) ((const uint8_t *) wts + k * wts_nb0 + (size_t) t * wts_nb1);
+ float * o = out + (size_t) t * n_embd;
+ PROF_T0_x(_td);
+ for (int64_t r = 0; r < n_down; r++) {
+ float acc = 0.0f;
+ dot_row(f_down, wd + (size_t) r * dn_row, hq, dn_row, (int) ed.down_ne0, acc);
+ o[r] += acc * w;
+ }
+ PROF_ADD(act_ns, _td);
+ }
+ PROF_ADD(fwd_ns, _ta);
+}
+
+// ---------------------------------------------------------------- row-parallel（整層常駐時）
+// 目的：decode 時 n_need = 9 < nth = 16，mode C 依 expert 切工作 → 9 條執行緒
+// 有 expert、7 條閒著。原生 ggml_mul_mat_id 的切法不是「一個 expert 一條」，
+// 而是「輸出列切給 16 條，每條在自己的列區塊裡跑完所有 expert」。
+// 這裡照同樣的切法：整層的 (expert, token, 輸出列) 展平，16 條執行緒各拿一段，
+// 每個 barrier 只發生在「up 全部算完 / h 全部量化完 / down 全部算完」，
+// 不是每個 expert 都同步一次。
+//
+// 只在「整層 expert 都已常駐」時啟用：row-parallel 沒有把 I/O 和計算交錯，
+// 若還有 miss，會比 mode C 的讀取/計算重疊更慢。常駐時完全沒有 I/O。
+//
+// 每層 4 道 barrier（acquire+xq、up、h 量化、down），decode 40 層 = 160 道；
+// 用自旋屏障，成本是微秒級。
+static inline size_t act_block_bytes(ggml_type t) {
+ return ggml_row_size(t, QK_K);
+}
+
+// 把一段連續的 QK_K block 量化到 dst（dst 必須指向該 token 的 row 開頭）。
+// 這樣 quantization 也可以切給 16 條執行緒，而不是只讓一條做 2048 個元素。
+static void quantize_blocks(ggml_type t, const float * src, uint8_t * dst,
+ int64_t b0, int64_t b1) {
+ const size_t rb = act_block_bytes(t);
+ for (int64_t b = b0; b < b1; b++) {
+ quantize_act(t, src + b * QK_K, dst + (size_t) b * rb, QK_K);
+ }
+}
+
+static void moe_row_parallel(ggml_tensor * dst, int ith, int nth, layer_rec * R,
+ const layer_desc * L, const float * cur, const float * wts,
+ size_t wts_nb0, size_t wts_nb1, int64_t n_embd, int64_t n_tok) {
+ pager & P = pager::instance();
+ const int n_need = (int) R->need.size();
+ float * out = (float *) dst->data;
+
+ // 和原生 mul_mat_id 一樣以「16 列」為一個工作單位，並用 atomic 動態派工。
+ // 每個 chunk 會把該 16 列在「所有 expert/token」上都算完：
+ // up : 16 rows × 9 experts → h_all
+ // down : 16 rows × 9 experts → out（同一個 (t,r) 由同一條執行緒累加，無 race）
+ // 這樣不需要 per-thread partial，也不需要最後的歸約。
+ const int64_t RB = 16;
+ const int64_t n_ff = L->experts[(size_t) R->need[0]].ne1;
+ const int64_t n_down = L->experts[(size_t) R->need[0]].down_ne1;
+ const int64_t nblk_up = (n_ff + RB - 1) / RB;
+ const int64_t nblk_dn = (n_down + RB - 1) / RB;
+
+ // 先把整層的 expert 備齊（常駐時全部是 hit）。
+ for (int i = ith; i < n_need; i += nth) {
+ const int32_t e = R->need[(size_t) i];
+ const int sl = P.acquire(R->il, e, true);
+ R->need_slot[(size_t) i] = sl;
+ if (sl >= 0) {
+ R->slot_of_expert[(size_t) e] = sl;
+ P.note_used(R->il, e);
+ }
+ }
+
+ const expert_desc & ed0 = L->experts[(size_t) R->need[0]];
+ const ggml_type act_up = act_type_for(ed0.type[part::UP]);
+ const ggml_type act_down = act_type_for(ed0.type[part::DOWN]);
+ const size_t xq_sz = ggml_row_size(act_up, ed0.ne0);
+ const size_t hq_sz = ggml_row_size(act_down, ed0.down_ne0);
+ const int64_t xq_blocks = ed0.ne0 / QK_K;
+ const int64_t hq_blocks = ed0.down_ne0 / QK_K;
+
+ // 展平 (expert, token) 的 prefix 與 mapping。這段不依賴 slot，先算，
+ // 才能在 barrier 前由 thread 0 配置「跨執行緒共用」的 h_all。
+ std::vector<int64_t> up_off((size_t) n_need + 1, 0);
+ std::vector<int32_t> tok_t, tok_k, tok_e;
+ std::vector<int64_t> tok_h;
+ int64_t total_up = 0, total_tok = 0;
+ for (int j = 0; j < n_need; j++) {
+ const int32_t e = R->need[(size_t) j];
+ const auto & toks = R->tok_of_expert[(size_t) e];
+ up_off[(size_t) j] = total_up;
+ for (size_t ti = 0; ti < toks.size(); ti++) {
+ tok_t.push_back(toks[ti].first);
+ tok_k.push_back(toks[ti].second);
+ tok_e.push_back(j);
+ tok_h.push_back(total_up + (int64_t) ti * n_ff);
+ }
+ total_up += (int64_t) toks.size() * n_ff;
+ total_tok += (int64_t) toks.size();
+ }
+ up_off[(size_t) n_need] = total_up;
+
+ if (ith == 0) {
+ memset(out, 0, (size_t) ggml_nbytes(dst));
+ R->rowpar_chunk_up.store(0, std::memory_order_relaxed);
+ R->rowpar_chunk_dn.store(0, std::memory_order_relaxed);
+ // h_all 必須跨執行緒共用：down 的每個 row block 需要「所有」up 列。
+ // 之前放在函式區域變數會變成每個執行緒各一份，down 讀到別人的 0。
+ R->rowpar_h.assign((size_t) total_up, 0.0f);
+ R->rowpar_hq.assign((size_t) total_tok * hq_sz, 0);
+ }
+ if (nth > 1) { R->rowbar.sync(nth); }
+
+ float * h_all = R->rowpar_h.data();
+ uint8_t * hq_all = R->rowpar_hq.data();
+
+ // 每個 expert 的常數先算好（slot 已由上面的 barrier 保證可見）。
+ struct ed_info {
+ const uint8_t * wu = nullptr;
+ const uint8_t * wd = nullptr;
+ const uint8_t * wg = nullptr; // gated 模型才有：gate 權重
+ size_t up_row = 0;
+ size_t dn_row = 0;
+ size_t g_row = 0;
+ vec_dot_fn f_up = nullptr;
+ vec_dot_fn f_down = nullptr;
+ vec_dot_fn f_gate = nullptr;
+ const expert_desc * ed = nullptr;
+ int64_t n_te = 0;
+ };
+ std::vector<ed_info> info((size_t) n_need);
+ for (int j = 0; j < n_need; j++) {
+ const int32_t e = R->need[(size_t) j];
+ const int sl = R->need_slot[(size_t) j];
+ const expert_desc & ed = L->experts[(size_t) e];
+ ed_info & q = info[(size_t) j];
+ q.ed = &ed;
+ q.n_te = (int64_t) R->tok_of_expert[(size_t) e].size();
+ q.up_row = ggml_row_size(ed.type[part::UP], ed.ne0);
+ q.dn_row = ggml_row_size(ed.type[part::DOWN], ed.down_ne0);
+ q.f_up = dot_for(ed.type[part::UP]);
+ q.f_down = dot_for(ed.type[part::DOWN]);
+ if (!R->no_gate) {
+ q.g_row = ggml_row_size(ed.type[part::GATE], ed.ne0);
+ q.f_gate = dot_for(ed.type[part::GATE]);
+ }
+ if (sl >= 0) {
+ q.wu = P.slot_ptr(sl) + ed.slot_off[part::UP];
+ q.wd = P.slot_ptr(sl) + ed.slot_off[part::DOWN];
+ if (!R->no_gate) {
+ q.wg = P.slot_ptr(sl) + ed.slot_off[part::GATE];
+ }
+ }
+ }
+
+ std::vector<uint8_t> xq_local((size_t) n_tok * xq_sz);
+ std::vector<uint8_t> xq_done((size_t) n_tok, 0);
+
+ // ---- up：動態抓 16 列 block，對所有 (expert, token) 一起算
+ for (;;) {
+ const int64_t b = R->rowpar_chunk_up.fetch_add(1, std::memory_order_relaxed);
+ if (b >= nblk_up) {
+ break;
+ }
+ const int64_t r0 = b * RB;
+ const int64_t r1 = std::min<int64_t>(r0 + RB, n_ff);
+ for (int j = 0; j < n_need; j++) {
+ const ed_info & q = info[(size_t) j];
+ if (!q.wu || !q.f_up) {
+ continue;
+ }
+ if (!R->no_gate && (!q.wg || !q.f_gate)) {
+ continue;
+ }
+ const auto & toks = R->tok_of_expert[(size_t) R->need[(size_t) j]];
+ for (size_t ti = 0; ti < toks.size(); ti++) {
+ const int32_t t = toks[ti].first;
+ if (!xq_done[(size_t) t]) {
+ quantize_blocks(act_up, cur + (size_t) t * n_embd,
+ xq_local.data() + (size_t) t * xq_sz, 0, xq_blocks);
+ xq_done[(size_t) t] = 1;
+ }
+ const uint8_t * xq = xq_local.data() + (size_t) t * xq_sz;
+ float * h = h_all + up_off[(size_t) j] + (int64_t) ti * n_ff;
+ for (int64_t r = r0; r < r1; r++) {
+ float acc = 0.0f;
+ dot_row(q.f_up, q.wu + (size_t) r * q.up_row, xq,
+ q.up_row, (int) q.ed->ne0, acc);
+ if (R->no_gate) {
+ // Nemotron LatentMoE：relu(x)^2，沒有 gate
+ acc = acc > 0.0f ? acc : 0.0f;
+ h[(size_t) r] = acc * acc;
+ } else {
+ // gated（Flint）：silu(gate)*up，與 expert_forward 相同的 SwiGLU
+ float g = 0.0f;
+ dot_row(q.f_gate, q.wg + (size_t) r * q.g_row, xq,
+ q.g_row, (int) q.ed->ne0, g);
+ h[(size_t) r] = gated_act(g, acc, R->swiglu_limit);
+ }
+ }
+ }
+ }
+ }
+ if (nth > 1) { R->rowbar.sync(nth); }
+
+ // ---- h 量化：total_tok × 8 個 block 平均分給 16 條執行緒。
+ // 不要讓每條執行緒各自量化一次：先前版本是 16 倍冗餘，decode 時反而成為瓶颈。
+ {
+ const int64_t total_blocks = total_tok * hq_blocks;
+ for (int64_t b = ith; b < total_blocks; b += nth) {
+ const int64_t tg = b / hq_blocks;
+ const int64_t bb = b % hq_blocks;
+ quantize_blocks(act_down, h_all + tok_h[(size_t) tg],
+ hq_all + (size_t) tg * hq_sz, bb, bb + 1);
+ }
+ }
+ if (nth > 1) { R->rowbar.sync(nth); }
+
+ // ---- down：同樣 16 列 block；同一 block 內把所有 expert 的貢獻加起來，
+ // 所以每個 (t, r) 只會被一條執行緒碰到，直接寫 out 即可。
+ for (;;) {
+ const int64_t b = R->rowpar_chunk_dn.fetch_add(1, std::memory_order_relaxed);
+ if (b >= nblk_dn) {
+ break;
+ }
+ const int64_t r0 = b * RB;
+ const int64_t r1 = std::min<int64_t>(r0 + RB, n_down);
+ for (int64_t tg = 0; tg < total_tok; tg++) {
+ const int j = tok_e[(size_t) tg];
+ const ed_info & q = info[(size_t) j];
+ if (!q.wd || !q.f_down) {
+ continue;
+ }
+ const uint8_t * hq = hq_all + (size_t) tg * hq_sz;
+ const int32_t t = tok_t[(size_t) tg];
+ const int32_t k = tok_k[(size_t) tg];
+ const float w = *(const float *) ((const uint8_t *) wts + k * wts_nb0
+ + (size_t) t * wts_nb1);
+ float * o = out + (size_t) t * n_embd;
+ for (int64_t r = r0; r < r1; r++) {
+ float acc = 0.0f;
+ dot_row(q.f_down, q.wd + (size_t) r * q.dn_row, hq,
+ q.dn_row, (int) q.ed->down_ne0, acc);
+ o[(size_t) r] += acc * w;
+ }
+ }
+ }
+ if (nth > 1) { R->rowbar.sync(nth); }
+
+ if (ith == 0) {
+ for (int i = 0; i < n_need; i++) {
+ if (R->need_slot[(size_t) i] >= 0) {
+ P.release(R->need_slot[(size_t) i]);
+ }
+ }
+ }
+}
+
+// 每次 MoE op 實際走的路徑計數（v2 row-parallel vs v1 mode C），結束時印出
+static std::atomic<uint64_t> g_ops_rowpar{0};
+static std::atomic<uint64_t> g_ops_modec{0};
+
+static void moe_custom_op(ggml_tensor * dst, int ith, int nth, void * userdata) {
+ layer_rec * R = (layer_rec *) userdata;
+ pager & P = pager::instance();
+
+ // 從 op 的 src 取得輸入（保證是「這一次」真的被計算的張量）
+ const ggml_tensor * src_cur = dst->src[0];
+ const ggml_tensor * src_ids = dst->src[1];
+ const ggml_tensor * src_wts = dst->src[2];
+ if (!src_cur || !src_ids || !src_wts) {
+ if (ith == 0) {
+ fprintf(stderr, "[sdq] op 缺少 src（%d）\n", R->il);
+ }
+ return;
+ }
+
+ const layer_desc * L = P.layer(R->il);
+ PROF_T0_x(_tg); // op 總計時
+ g_prof.calls++;
+ const int64_t n_embd = src_cur->ne[0];
+ const int64_t n_tok = src_cur->ne[2];
+ const int64_t n_used = R->n_expert_used;
+ // ids 是 2D [n_expert_used, n_tokens] → 用 nb[0]/nb[1]
+ // wts 是 3D [1, n_expert_used, n_tokens] → 用 nb[1]/nb[2]（nb[0] 會是 4 而不是 n_used*4！）
+ const size_t ids_nb0 = src_ids->nb[0], ids_nb1 = src_ids->nb[1];
+ const size_t wts_nb0 = src_wts->nb[src_wts->ne[0] == 1 ? 1 : 0];
+ const size_t wts_nb1 = src_wts->nb[src_wts->ne[0] == 1 ? 2 : 1];
+ const float * cur = (const float *) src_cur->data;
+ const float * wts = (const float *) src_wts->data;
+
+ if (!L || L->experts.empty() || dst->type != GGML_TYPE_F32) {
+ if (ith == 0 && !R->warned) {
+ R->warned = true;
+ fprintf(stderr, "[sdq] 層 %d 沒有可用的分頁描述\n", R->il);
+ }
+ return;
+ }
+
+ if (R->il == 0 && getenv("SDQ_DBG_OP")) {
+ fprintf(stderr, "[sdq-op] L%d n_tok=%lld nth=%d need=%d ids_nb1=%zu wts_nb1=%zu\n",
+ R->il, (long long) n_tok, nth, (int) R->need.size(), ids_nb1, wts_nb1);
+ }
+ // ---- 1) 解除上一層的 pin（mode B 直接累加輸出，需要先清空）
+ PROF_T0_x(_tsetup);
+ if (ith == 0) {
+ P.unpin_all();
+ }
+ if (nth > 1) { R->bar.sync(nth); }
+
+ // ---- 2) 依 expert 分組 token（只有執行緒 0 做，其餘在此等待）
+ if (ith == 0) {
+ std::lock_guard<std::mutex> lk(R->mtx);
+ R->tok_of_expert.assign((size_t) R->n_expert, {});
+ R->exp_of_tok.assign((size_t) n_tok, {});
+ const uint8_t * ids_base = (const uint8_t *) src_ids->data;
+ for (int64_t t = 0; t < n_tok; t++) {
+ for (int64_t k = 0; k < n_used; k++) {
+ const int32_t e = *(const int32_t *) (ids_base + k * ids_nb0 + (size_t) t * ids_nb1);
+ if (e >= 0 && e < R->n_expert) {
+ auto & v = R->tok_of_expert[(size_t) e];
+ bool seen = false;
+ for (const auto & p : v) {
+ if (p.first == (int32_t) t) { seen = true; break; }
+ }
+ if (!seen) {
+ v.emplace_back((int32_t) t, (int32_t) k);
+ }
+ R->exp_of_tok[(size_t) t].emplace_back(e, (int32_t) k);
+ }
+ }
+ }
+ R->need.clear();
+ for (int64_t e = 0; e < R->n_expert; e++) {
+ if (!R->tok_of_expert[(size_t) e].empty()) {
+ R->need.push_back((int32_t) e);
+ }
+ }
+ R->need_sorted = R->need;
+ std::sort(R->need_sorted.begin(), R->need_sorted.end());
+ R->need_slot.assign(R->need.size(), -1); // 只能由 ith==0 配置，其餘執行緒唯讀
+ R->slot_of_expert.assign((size_t) R->n_expert, -1);
+ R->need_slot_of_expert.assign((size_t) R->n_expert, -1);
+ R->miss_sorted.clear();
+ R->miss_sorted.reserve((size_t) R->n_expert);
+ if (R->partial.size() < 64) {
+ R->partial.resize(64);
+ }
+ sdq_route_log(R->il, n_tok, src_ids, ids_nb0, ids_nb1, (int) n_used);
+ }
+ if (nth > 1) { R->bar.sync(nth); }
+ PROF_ADD(setup_ns, _tsetup);
+
+ const int n_need = (int) R->need.size();
+
+ // ---- 2.5) 決定本層要不要走 row-parallel
+ //
+ // 條件：no-gate（本模型）、decode（n_need < nth）、而且整層 expert 都已常駐。
+ // 常駐代表沒有 I/O 可以跟計算重疊，mode C 的「讀→算交錯」優勢消失，
+ // 此時把 expert 當序列外層、輸出列當平行內層，才能讓 16 條執行緒都有事做。
+ // 只要有一個 miss 就退回 mode C，避免把小記憶體最需要的 I/O/計算重疊拆掉。
+ static const bool rowpar_env = [] {
+ const char * e = getenv("SDQ_ROW_PARALLEL");
+ return e ? atoi(e) != 0 : true;
+ }();
+ static const int rowpar_max_miss = [] {
+ const char * e = getenv("SDQ_ROWPAR_MAX_MISS");
+ return e ? atoi(e) : 0;
+ }();
+ if (ith == 0) {
+ R->use_rowpar = false;
+ R->rowpar_miss = 0;
+ // gated（Flint）需要 gate 與 up 共用同一個 activation 量化型別，才能共用 xq。
+ bool types_ok = true;
+ if (!R->no_gate && n_need > 0) {
+ const expert_desc & e0 = L->experts[(size_t) R->need[0]];
+ types_ok = !e0.fused_gate_up && dot_for(e0.type[part::GATE]) && dot_for(e0.type[part::UP]) &&
+ act_type_for(e0.type[part::GATE]) == act_type_for(e0.type[part::UP]);
+ }
+ if (rowpar_env && n_need > 0 && n_need < nth && types_ok) {
+ int miss = 0;
+ for (int i = 0; i < n_need; i++) {
+ if (!P.resident(R->il, R->need[(size_t) i])) {
+ miss++;
+ }
+ }
+ R->rowpar_miss = miss;
+ R->use_rowpar = miss <= rowpar_max_miss;
+ }
+ static std::atomic<int> dbg_left{8};
+ if (getenv("SDQ_DBG_OP") && dbg_left.fetch_add(-1) > 0) {
+ fprintf(stderr, "[sdq-op] L%d rowpar=%d need=%d n_tok=%lld\n",
+ R->il, (int) R->use_rowpar, n_need, (long long) n_tok);
+ }
+ }
+ if (nth > 1) { R->bar.sync(nth); }
+
+ // 計數：每次 op 實際走哪條路徑（ith==0 記一次；barrier 之後 use_rowpar 已確定）
+ if (ith == 0) {
+ (R->use_rowpar ? g_ops_rowpar : g_ops_modec).fetch_add(1, std::memory_order_relaxed);
+ }
+ if (R->use_rowpar) {
+ moe_row_parallel(dst, ith, nth, R, L, cur, wts, wts_nb0, wts_nb1, n_embd, n_tok);
+ }
+
+ // ---- 3+4) 分頁與計算融合成一條 pipeline
+ //
+ // 關鍵效能點：不要「先把整層的 expert 都讀完再開始算」。做法是每個執行緒
+ // 自己依序處理它負責的 expert：讀一個 → 立刻算它 → 再讀下一個。
+ // 這樣 SSD 讀取和 CPU 計算自然重疊，而且 8 條執行緒各自發出 pread，
+ // SSD 佇列深度也夠（不需要再靠 queue depth）。
+ //
+ // 平行切分：mode A 依 expert 切（每個執行緒有自己的累加緩衝，最後相加）；
+ // mode B 依 token 切（同一個 token 的所有 expert 由同一執行緒算，
+ // 但必須先把該 token 的 expert 都備齊，所以先分頁再計算）。
+ const int64_t out_elems = n_embd * n_tok;
+ const size_t partial_bytes = (size_t) out_elems * 4 * (size_t) nth;
+ const bool mode_a = partial_bytes <= (size_t) env_partial_cap_mb() << 20;
+ float * out = (float *) dst->data;
+
+ if (!mode_a && !R->use_rowpar) {
+ memset(dst->data, 0, (size_t) ggml_nbytes(dst));
+ }
+
+ if (!R->use_rowpar && mode_a) {
+ // ===================== mode C（預設）：依 expert 流水線 =====================
+ // 三個設計重點，每個都來自實測：
+ //
+ // (1) **取→算交錯，中間沒有 barrier**。
+ // 舊的 mode B 是「先把整層 8 個 expert 全部讀完 → barrier → 才開始算」。
+ // 量測顯示 decode 有 66% 的時間卡在 barrier 上等最慢的執行緒
+ // （SDQ_PROFILE_TIME 的 barrier 等待 = op 時間的 70%）。改成交錯後
+ // 讀取與計算自然重疊，實測 2.79 → 3.49 tok/s。
+ //
+ // (2) **分頁走 acquire()，不用 reserve_slot + read_expert_span**。
+ // reserve_slot() 會在資料還沒讀進來之前就把 (layer,expert) publish 到
+ // page_ 表；一旦該層的 pread 失敗或被中斷，這個槽就會留在 page_ 裡
+ // 卻是舊資料，之後的 probe() 會把它當成命中 → 靜默算錯。
+ // acquire() 是「先讀完再 publish」，成本相同但不會有這個失效模式。
+ //
+ // (3) **i += nth 交錯分配**，不是切成連續區塊。
+ // 每個 expert 的 SSD 延遲差異很大，交錯分配讓快執行緒自然多幫忙。
+ //
+ // 輸出不直接寫 dst（同一個 token 的不同 expert 會撞在一起），
+ // 所以仍然用「每執行緒一份累加緩衝 + 最後歸約」。
+ std::vector<float> & mine = R->partial[(size_t) ith];
+ mine.assign((size_t) out_elems, 0.0f);
+ std::vector<uint8_t> scratch;
+
+ // 在取本層的 expert **之前**先替下一層排預取。
+ //
+ // 實測結論（S7）：**預設關閉，而且量到開啟反而慢 8%。**
+ // 預測來自 prev_used_[il+1]（上一個 token 在該層選的專家）。想用它躲在
+ // 「本層的讀取＋計算」後面，但量到的結果是：
+ // 開（5.68 tok/s，預取使用率 0.48 %）vs 關（6.20 tok/s，0.28 %）
+ // 命中率幾乎一樣（70.5 % vs 70.7 %），但 SSD 位元組從 201.5 升到 216.0 MB/token
+ // → 預測不夠準，多讀的全是白讀，而且還要和本層的 demand 讀搶頻寬。
+ // 另外開啟時同一提示詞的輸出不一致 → 預取佔用的槽位（busy）會讓
+ // acquire() 偶爾拿不到槽位而**靜默跳過一個 expert**，這是正確性問題。
+ //
+ // 保留程式碼是為了讓後人不用重新發明；文件見 docs/06_slots.md §4。
+ // SDQ_PREFETCH_LAYER_AHEAD=1 可以打開來重現那個負結果。
+ static const bool layer_ahead = [] {
+ const char * e = getenv("SDQ_PREFETCH_LAYER_AHEAD");
+ return e ? atoi(e) != 0 : false;
+ }();
+ if (layer_ahead && ith == 0) {
+ P.prefetch_layer_ahead(R->il + 1);
+ }
+
+ for (int i = ith; i < n_need; i += nth) {
+ const int32_t e = R->need[(size_t) i];
+ const int32_t sl = P.acquire(R->il, e, true);
+ if (sl < 0) {
+ continue;
+ }
+ P.note_used(R->il, e);
+ if (R->no_gate) {
+ expert_forward_relu2(L->experts[(size_t) e], P.slot_ptr(sl), cur, n_embd,
+ R->tok_of_expert[(size_t) e], wts, wts_nb0, wts_nb1,
+ mine.data(), &scratch);
+ } else {
+ expert_forward(L->experts[(size_t) e], P.slot_ptr(sl), cur, n_embd,
+ R->tok_of_expert[(size_t) e], wts, wts_nb0, wts_nb1,
+ (int) n_used, R->swiglu_limit, mine.data(), &scratch);
+ }
+ // expert_forward 回傳後這個槽位就沒人讀了（最後的歸約只讀 mine），
+ // 所以立刻放掉 pin，不要讓它被釘到下一層開頭。
+ // 不放的話 prefill 一層會把整層用到的 expert 全部釘住（最多 256 個），
+ // 大槽位不夠時就會中止（見 pager::release 的說明）。
+ P.release(sl);
+ }
+
+ if (nth > 1) { R->bar.sync(nth); }
+ for (int64_t idx = ith; idx < out_elems; idx += nth) {
+ float acc = 0.0f;
+ for (int k = 0; k < nth; k++) {
+ acc += R->partial[(size_t) k][(size_t) idx];
+ }
+ out[idx] = acc;
+ }
+ } else if (!R->use_rowpar) {
+ // mode B：token 多到 partial 緩衝太大時，依 token 切分
+ // 先把所有需要的 expert 分頁（各執行緒平行 pread）
+ {
+ PROF_T0_x(_th); // 分頁 / SSD 讀取
+ for (int i = ith; i < n_need; i += nth) {
+ const int32_t e = R->need[(size_t) i];
+ const int s = P.acquire(R->il, e, true);
+ g_prof.experts++;
+ R->need_slot[(size_t) i] = s;
+ if (s >= 0) {
+ R->slot_of_expert[(size_t) e] = s;
+ P.note_used(R->il, e);
+ }
+ }
+ PROF_ADD(read_ns, _th);
+ }
+ if (nth > 1) { R->bar.sync(nth); }
+ const float * cur_p = (const float *) src_cur->data;
+ const float * wts_p = (const float *) src_wts->data;
+ for (int64_t t = ith; t < n_tok; t += nth) {
+ std::vector<uint8_t> scratch;
+ for (const auto & ek : R->exp_of_tok[(size_t) t]) {
+ const int slot = R->slot_of_expert[(size_t) ek.first];
+ if (slot < 0) {
+ continue;
+ }
+ std::vector<std::pair<int32_t, int32_t>> one = {{(int32_t) t, ek.second}};
+ if (R->no_gate) {
+ expert_forward_relu2(L->experts[(size_t) ek.first], P.slot_ptr(slot), cur_p, n_embd,
+ one, wts_p, wts_nb0, wts_nb1, out, &scratch);
+ } else {
+ expert_forward(L->experts[(size_t) ek.first], P.slot_ptr(slot), cur_p, n_embd,
+ one, wts_p, wts_nb0, wts_nb1, (int) n_used,
+ R->swiglu_limit, out, &scratch);
+ }
+ }
+ }
+ }
+
+ // 除錯：把某一層的實際輸入/輸出 dump 出來，用 NumPy 對照。
+ // 必須先 barrier：上面的歸約是由各執行緒分片寫 out 的，沒有這一道會讀到半成品。
+ if (nth > 1) { R->bar.sync(nth); }
+ if (const char * dump_path = getenv("SDQ_MOE_DUMP")) {
+ const int dump_layer = getenv("SDQ_MOE_DUMP_LAYER") ? atoi(getenv("SDQ_MOE_DUMP_LAYER")) : -1;
+ if ((dump_layer < 0 || R->il == dump_layer) && ith == 0) {
+ std::string path = std::string(dump_path) + "." + std::to_string(R->il);
+ FILE * f = fopen(path.c_str(), "wb");
+ if (f) {
+ const int32_t meta[6] = {(int32_t) n_embd, (int32_t) n_tok, (int32_t) n_used,
+ (int32_t) R->n_expert, 1, 2};
+ fwrite(meta, sizeof(int32_t), 6, f);
+ fwrite(cur, sizeof(float), (size_t) (n_embd * n_tok), f);
+ const uint8_t * ids_base = (const uint8_t *) src_ids->data;
+ for (int64_t t = 0; t < n_tok; t++) {
+ for (int64_t k = 0; k < n_used; k++) {
+ const int32_t e = *(const int32_t *) (ids_base + k * ids_nb0 + (size_t) t * ids_nb1);
+ fwrite(&e, sizeof(int32_t), 1, f);
+ }
+ }
+ fwrite(wts, sizeof(float), (size_t) (n_used * n_tok), f);
+ fwrite(out, sizeof(float), (size_t) (n_embd * n_tok), f);
+ fclose(f);
+ fprintf(stderr, "[sdq] dumped layer %d -> %s\n", R->il, path.c_str());
+ }
+ }
+ }
+
+ if (nth > 1) { R->bar.sync(nth); }
+ PROF_ADD(op_ns, _tg);
+}
+
+// 輸出路徑計數（sdq_shutdown 呼叫）
+void sdq_print_op_counts() {
+ fprintf(stderr, "[sdq] MoE op 路徑計數：row-parallel %llu 次，mode C %llu 次\n",
+ (unsigned long long) g_ops_rowpar.load(), (unsigned long long) g_ops_modec.load());
+}
+
+// 輸出計時結果（sdq_shutdown 呼叫）
+void sdq_print_time_profile() {
+ if (!prof_on()) {
+ return;
+ }
+ const double ms = 1e6;
+ auto v = [ms](const std::atomic<uint64_t> & a) { return a.load() / ms; };
+ const double op = v(g_prof.op_ns);
+ const double fwd = v(g_prof.fwd_ns);
+ fprintf(stderr, "[sdq-prof] 自訂 MoE op（%llu 次呼叫、%llu 個 expert 分頁）\n",
+ (unsigned long long) g_prof.calls.load(), (unsigned long long) g_prof.experts.load());
+ fprintf(stderr, "[sdq-prof] op 合計 %9.1f ms\n", op);
+ fprintf(stderr, "[sdq-prof] expert_forward %9.1f ms (%5.1f%%) 其中：\n", fwd, op > 0 ? 100.0 * fwd / op : 0.0);
+ fprintf(stderr, "[sdq-prof] 量化 activation%9.1f ms (%5.1f%% of fwd)\n", v(g_prof.qact_ns), fwd > 0 ? 100.0 * v(g_prof.qact_ns) / fwd : 0.0);
+ fprintf(stderr, "[sdq-prof] 三次點積 %9.1f ms (%5.1f%% of fwd)\n", v(g_prof.dot_ns), fwd > 0 ? 100.0 * v(g_prof.dot_ns) / fwd : 0.0);
+ fprintf(stderr, "[sdq-prof] SiLU+加權 %9.1f ms (%5.1f%% of fwd)\n", v(g_prof.act_ns), fwd > 0 ? 100.0 * v(g_prof.act_ns) / fwd : 0.0);
+ fprintf(stderr, "[sdq-prof] SSD 讀取/分頁 %9.1f ms (%5.1f%%)\n", v(g_prof.read_ns), op > 0 ? 100.0 * v(g_prof.read_ns) / op : 0.0);
+ fprintf(stderr, "[sdq-prof] 前置 setup %9.1f ms (%5.1f%%)\n", v(g_prof.setup_ns), op > 0 ? 100.0 * v(g_prof.setup_ns) / op : 0.0);
+ fprintf(stderr, "[sdq-prof] barrier 等待 %9.1f ms (%5.1f%%)\n", v(g_prof.bar_ns), op > 0 ? 100.0 * v(g_prof.bar_ns) / op : 0.0);
+}
+
+// ---------------------------------------------------------------- graph 建構
+
+ggml_tensor * sdq_build_moe_ex(ggml_cgraph * gf, ggml_context * ctx, ggml_tensor * cur,
+ ggml_tensor * ids, ggml_tensor * weights, int il, int64_t n_expert,
+ int64_t n_expert_used, bool no_gate, bool relu_sqr) {
+ layer_rec * R = rec_for(il);
+ R->n_expert = n_expert;
+ R->n_expert_used = n_expert_used;
+ R->no_gate = no_gate;
+ R->relu_sqr = relu_sqr;
+ R->swiglu_limit = 0.0f;
+
+ // 呼叫端給的是 2D [n_embd, n_tokens]，這裡跟上游一樣先 reshape 成 3D
+ // [n_embd, 1, n_tokens]（上游 mul_mat_id 路徑也是這樣拿 token 數的）
+ ggml_tensor * cur3 = (cur->ne[1] == 1) ? cur : ggml_reshape_3d(ctx, cur, cur->ne[0], 1, cur->ne[1]);
+ const int64_t n_tok = cur3->ne[2];
+
+ // Nemotron LatentMoE 的 down 輸出寬度 = moe_latent_size，與輸入相同。
+ int64_t out_w = cur3->ne[0];
+ const layer_desc * L = pager::instance().layer(il);
+ if (no_gate && L && !L->experts.empty()) {
+ out_w = L->experts[0].down_ne1;
+ }
+
+ ggml_tensor * args[3] = { cur3, ids, weights };
+ ggml_tensor * out = ggml_custom_4d(ctx, GGML_TYPE_F32, out_w, n_tok, 1, 1,
+ args, 3, moe_custom_op, GGML_N_TASKS_MAX, R);
+ ggml_set_name(out, "sdq_moe_out");
+ // 關鍵：把這個節點註冊進圖，ggml 才會為它配置獨立的 compute buffer。
+ // 少了這一行，layer>=1 的輸出會被其他張量覆蓋（layer 0 僥倖正確）。
+ ggml_build_forward_expand(gf, out);
+ return out;
+}
+
+ggml_tensor * sdq_build_moe(ggml_cgraph * gf, ggml_context * ctx, ggml_tensor * cur,
+ ggml_tensor * ids, ggml_tensor * weights, int il, int64_t n_expert,
+ int64_t n_expert_used, float swiglu_limit) {
+ ggml_tensor * out = sdq_build_moe_ex(gf, ctx, cur, ids, weights, il, n_expert, n_expert_used, false, false);
+ // 把 swiglu limit 傳給 op（先前這裡 (void) 丟掉，造成 clamp 永遠不生效）
+ rec_for(il)->swiglu_limit = swiglu_limit;
+ return out;
+}
+
+} // namespace sdq
+namespace sdq {
+// 自我測試：對指定 (layer, expert) 用 pager 的資料算一次 MoE 輸出，
+// 讓 Python 用純 NumPy 從同一段 GGUF 位元組重算並比對。
+// 輸出：out[0..n_embd) = 輸出，out[n_embd..2*n_embd) = 輸入 activation
+int sdq_moe_selftest(int il, int ie, const float * x, float * out_io,
+ int64_t n_embd, int64_t n_ff, int n_tok) {
+ pager & P = pager::instance();
+ const layer_desc * L = P.layer(il);
+ if (!L || ie < 0 || ie >= (int) L->experts.size()) {
+ fprintf(stderr, "[sdq-selftest] 層/expert 不存在：%d/%d\n", il, ie);
+ return 1;
+ }
+ const int slot = P.acquire(il, ie, true);
+ if (slot < 0) {
+ fprintf(stderr, "[sdq-selftest] 分頁失敗\n");
+ return 2;
+ }
+ const expert_desc & ed = L->experts[(size_t) ie];
+ if (ed.ne0 != n_embd || ed.ne1 != n_ff) {
+ fprintf(stderr, "[sdq-selftest] 維度不符：ne0=%lld ne1=%lld（預期 %lld/%lld）\n",
+ (long long) ed.ne0, (long long) ed.ne1, (long long) n_embd, (long long) n_ff);
+ return 3;
+ }
+ std::vector<std::pair<int32_t, int32_t>> toks;
+ std::vector<float> wts((size_t) n_tok, 1.0f);
+ for (int t = 0; t < n_tok; t++) {
+ toks.emplace_back(t, 0);
+ }
+ std::vector<float> out((size_t) n_embd * n_tok, 0.0f);
+ std::vector<uint8_t> scratch;
+ // SDQ_SELFTEST_DUMP: 把 arena 槽位的原始位元組存檔，驗證 pager 讀到的是正確的 expert
+ if (const char * dp = getenv("SDQ_SELFTEST_DUMP")) {
+ std::string path = std::string(dp) + "." + std::to_string(il) + "." + std::to_string(ie);
+ FILE * f = fopen(path.c_str(), "wb");
+ if (f) {
+ fwrite(P.slot_ptr(slot), 1, (size_t) ed.bytes, f);
+ fclose(f);
+ fprintf(stderr, "[sdq-selftest] slot dump -> %s (%lld B)\n", path.c_str(), (long long) ed.bytes);
+ }
+ }
+ expert_forward(ed, P.slot_ptr(slot), x, n_embd, toks, wts.data(), 4, 4, 1, 0.0f,
+ out.data(), &scratch);
+ for (int64_t t = 0; t < n_tok; t++) {
+ for (int64_t i = 0; i < n_embd; i++) {
+ out_io[t * n_embd + i] = out[(size_t) t * n_embd + i];
+ out_io[n_tok * n_embd + t * n_embd + i] = x[(size_t) t * n_embd + i];
+ }
+ }
+ P.unpin_all();
+ fprintf(stderr, "[sdq-selftest] L%d E%d n_tok=%d 計算完成（gate %s / up %s / down %s）\n",
+ il, ie, n_tok, ggml_type_name(ed.type[part::GATE]), ggml_type_name(ed.type[part::UP]),
+ ggml_type_name(ed.type[part::DOWN]));
+ return 0;
+}
+} // namespace sdq
diff --git a/src/sdq_pager.cpp b/src/sdq_pager.cpp
new file mode 100644
index 0000000..ff22bf7
--- /dev/null
+++ b/src/sdq_pager.cpp
@@ -0,0 +1,2567 @@
+// SDQ expert pager 實作：見 sdq_pager.h 的設計說明
+#include "sdq_pager.h"
+
+#include "llama-mmap.h"
+
+#include <algorithm>
+#include <thread>
+#include <cinttypes>
+#include <cmath>
+#include <cstdio>
+#include <cstdlib>
+#include <cstring>
+#include <fcntl.h>
+#include <ctime>
+#include <fstream>
+#include <map>
+#include <set>
+#include <sstream>
+#include <stdexcept>
+#include <sys/mman.h>
+#include <sys/stat.h>
+#include <sys/types.h>
+#include <unistd.h>
+
+namespace sdq {
+
+// ============================================================ 小工具
+
+uint64_t now_us() {
+ struct timespec ts;
+ clock_gettime(CLOCK_MONOTONIC, &ts);
+ return (uint64_t) ts.tv_sec * 1000000ull + (uint64_t) ts.tv_nsec / 1000ull;
+}
+
+size_t process_rss_bytes() {
+ std::ifstream f("/proc/self/statm");
+ size_t total = 0, resident = 0;
+ if (f >> total >> resident) {
+ return resident * (size_t) sysconf(_SC_PAGESIZE);
+ }
+ return 0;
+}
+
+size_t physical_ram_bytes() {
+ std::ifstream f("/proc/meminfo");
+ std::string key;
+ size_t value = 0;
+ std::string unit;
+ while (f >> key >> value >> unit) {
+ if (key == "MemTotal:") {
+ return value * 1024;
+ }
+ }
+ return 0;
+}
+
+// 可用記憶體（自動模式用）：取 /proc/meminfo 的 MemAvailable，若處於 cgroup v2 限制下，
+// 再與 memory.max - memory.current 取較小值（容器上限比整台主機的可用量更有意義）。
+// memory.current 含 page cache，所以這個估計偏保守。
+size_t available_ram_bytes() {
+ size_t avail = 0;
+ {
+ std::ifstream f("/proc/meminfo");
+ std::string key, unit;
+ size_t value = 0;
+ while (f >> key >> value >> unit) {
+ if (key == "MemAvailable:") {
+ avail = value * 1024;
+ break;
+ }
+ }
+ }
+ std::string cg;
+ {
+ std::ifstream f("/proc/self/cgroup");
+ std::string line;
+ while (std::getline(f, line)) {
+ if (line.rfind("0::", 0) == 0) {
+ cg = line.substr(3);
+ break;
+ }
+ }
+ }
+ const std::string base = "/sys/fs/cgroup" + (cg == "/" ? std::string() : cg);
+ auto read_num = [](const std::string & path, size_t & out) -> bool {
+ std::ifstream f(path);
+ std::string v;
+ if (!(f >> v) || v == "max") {
+ return false;
+ }
+ try {
+ out = (size_t) std::stoull(v);
+ } catch (...) {
+ return false;
+ }
+ return true;
+ };
+ size_t limit = 0, current = 0;
+ if (read_num(base + "/memory.max", limit) && read_num(base + "/memory.current", current)) {
+ const size_t headroom = limit > current ? limit - current : 0;
+ if (avail == 0 || headroom < avail) {
+ avail = headroom;
+ }
+ }
+ return avail;
+}
+
+static size_t env_size(const char * name, size_t dflt) {
+ const char * v = getenv(name);
+ if (!v || !*v) {
+ return dflt;
+ }
+ return (size_t) strtoull(v, nullptr, 10);
+}
+
+static bool env_flag(const char * name, bool dflt) {
+ const char * v = getenv(name);
+ if (!v || !*v) {
+ return dflt;
+ }
+ return !(v[0] == '0' && v[1] == '\0');
+}
+
+// ============================================================ GGUF 檔頭解析
+
+namespace {
+
+struct gguf_reader {
+ const uint8_t * p = nullptr;
+ size_t size = 0;
+ size_t o = 0;
+
+ uint8_t u8() { return p[o++]; }
+ uint32_t u32() { uint32_t v; memcpy(&v, p + o, 4); o += 4; return v; }
+ uint64_t u64() { uint64_t v; memcpy(&v, p + o, 8); o += 8; return v; }
+ int64_t i64() { int64_t v; memcpy(&v, p + o, 8); o += 8; return v; }
+ float f32() { float v; memcpy(&v, p + o, 4); o += 4; return v; }
+ std::string str() {
+ uint64_t n = u64();
+ std::string s((const char *) p + o, (size_t) n);
+ o += n;
+ return s;
+ }
+ void skip(uint64_t n) { o += (size_t) n; }
+ uint16_t u16() { uint16_t v; memcpy(&v, p + o, 2); o += 2; return v; }
+ int16_t i16() { int16_t v; memcpy(&v, p + o, 2); o += 2; return v; }
+ double f64() { double v; memcpy(&v, p + o, 8); o += 8; return v; }
+};
+
+struct kv_value {
+ enum kind { NONE, INT, REAL, STR, ARR } k = NONE;
+ int64_t i = 0;
+ double f = 0;
+ std::string s;
+ std::vector<int64_t> arr;
+};
+
+struct tensor_info {
+ std::string name;
+ std::vector<int64_t> dims;
+ ggml_type type = GGML_TYPE_COUNT;
+ uint64_t offset = 0; // 在所屬 shard 檔案內、相對 data_start 的偏移
+ int shard = 0; // 屬於哪一個 shard（sharded GGUF）
+};
+
+size_t ggml_type_size_of(ggml_type t) {
+ return ggml_row_size(t, 1);
+}
+
+} // namespace
+
+// ============================================================ pager
+
+pager & pager::instance() {
+ static pager p;
+ return p;
+}
+
+// 依 GGUF split 命名規則把第一個 shard 路徑展開成全部 shard。
+// 規則是 "<prefix>-00001-of-00004.gguf"；llama.cpp 用 llama_split_prefix/llama_split_path
+// 做同一件事，這裡獨立實作以免 pager 依賴 loader 的內部符號。
+std::vector<std::string> pager::list_splits(const std::string & first_path, int n_split) {
+ std::vector<std::string> out;
+ if (n_split <= 1) {
+ out.push_back(first_path);
+ return out;
+ }
+ // 用 llama.cpp 自己的命名規則（%s-%05d-of-%05d.gguf，且 split_no 由 1 起算）。
+ // 自己刻字串很容易差一位數字 —— 實測就是這樣變成 "...XL--00001-of-..."。
+ std::vector<char> pre(llama_path_max(), 0);
+ const int pre_len = llama_split_prefix(pre.data(), pre.size(), first_path.c_str(), 0, n_split);
+ if (pre_len <= 0) {
+ throw std::runtime_error("無法解析 split 檔名：" + first_path);
+ }
+ std::vector<char> buf(llama_path_max(), 0);
+ for (int i = 0; i < n_split; i++) {
+ const int len = llama_split_path(buf.data(), buf.size(), pre.data(), i, n_split);
+ if (len <= 0) {
+ throw std::runtime_error("無法產生 split 路徑：" + first_path);
+ }
+ out.emplace_back(buf.data(), (size_t) len);
+ }
+ return out;
+}
+
+void pager::load_gguf_header(const std::string & path) {
+ const int fd = open(path.c_str(), O_RDONLY);
+ if (fd < 0) {
+ throw std::runtime_error("無法開啟 GGUF：" + path);
+ }
+
+ struct stat st;
+ fstat(fd, &st);
+
+ // 只讀檔頭：metadata + tensor 資訊大約 11 MB（shard 1 還含整份 tokenizer，
+ // 這裡的 64 MiB 上限對它足夠）
+ const size_t want = (size_t) std::min<int64_t>(st.st_size, 64ull << 20);
+ std::vector<uint8_t> buf(want);
+ ssize_t got = 0;
+ while (got < (ssize_t) want) {
+ ssize_t n = pread(fd, buf.data() + got, want - got, got);
+ if (n <= 0) {
+ break;
+ }
+ got += n;
+ }
+ buf.resize((size_t) std::max<ssize_t>(got, 0));
+
+ gguf_reader r{buf.data(), buf.size(), 0};
+ if (memcmp(buf.data(), "GGUF", 4) != 0) {
+ close(fd);
+ throw std::runtime_error("不是 GGUF 檔：" + path);
+ }
+ r.o = 4;
+ const uint32_t version = r.u32();
+ if (version < 2 || version > 3) {
+ close(fd);
+ throw std::runtime_error("不支援的 GGUF 版本");
+ }
+ const uint64_t n_tensors = r.u64();
+ const uint64_t n_kv = r.u64();
+
+ std::map<std::string, kv_value> kv;
+ for (uint64_t i = 0; i < n_kv; i++) {
+ std::string key = r.str();
+ const uint32_t t = r.u32();
+ kv_value v;
+ switch (t) {
+ // GGUF 型別寬度：U8/I8/BOOL=1、U16/I16=2、U32/I32/F32=4、U64/I64/F64=8。
+ // 舊版把 case 0/1/2/3 都讀成 u32，遇到 tokenizer 或 U16 metadata 就整個
+ // header 錯位（症狀是 std::length_error: basic_string::_M_create）。
+ case 0: v.k = kv_value::INT; v.i = (int64_t) r.u8(); v.f = (double) v.i; break;
+ case 1: v.k = kv_value::INT; v.i = (int64_t) (int8_t) r.u8(); v.f = (double) v.i; break;
+ case 2: v.k = kv_value::INT; v.i = (int64_t) r.u16(); v.f = (double) v.i; break;
+ case 3: v.k = kv_value::INT; v.i = (int64_t) r.i16(); v.f = (double) v.i; break;
+ case 4: v.k = kv_value::INT; v.i = (int64_t) (uint32_t) r.u32(); v.f = (double) v.i; break;
+ case 5: v.k = kv_value::INT; v.i = (int64_t) (int32_t) r.u32(); v.f = (double) v.i; break;
+ case 6: v.k = kv_value::REAL; v.f = r.f32(); v.i = (int64_t) v.f; break;
+ case 7: v.k = kv_value::INT; v.i = (int64_t) r.u8(); v.f = (double) v.i; break; // BOOL = 1 byte
+ case 8: v.k = kv_value::STR; v.s = r.str(); break;
+ case 9: {
+ v.k = kv_value::ARR;
+ const uint32_t et = r.u32();
+ const uint64_t n = r.u64();
+ for (uint64_t j = 0; j < n; j++) {
+ switch (et) {
+ // 元素寬度必須跟 kv 一致（U8/I8/BOOL=1、U16/I16=2 …）。
+ // F64 一定要消費 8 bytes，舊版沒讀會讓 header 往後全部錯位。
+ case 0: v.arr.push_back(r.u8()); break;
+ case 1: v.arr.push_back((int8_t) r.u8()); break;
+ case 2: v.arr.push_back(r.u16()); break;
+ case 3: v.arr.push_back(r.i16()); break;
+ case 4: v.arr.push_back((uint32_t) r.u32()); break;
+ case 5: v.arr.push_back((int32_t) r.u32()); break;
+ case 6: v.arr.push_back((int64_t) r.f32()); break;
+ case 7: v.arr.push_back(r.u8()); break;
+ case 10: v.arr.push_back((int64_t) r.u64()); break;
+ case 11: v.arr.push_back(r.i64()); break;
+ case 12: r.f64(); v.arr.push_back(0); break;
+ case 8: r.skip(r.u64()); v.arr.push_back(0); break;
+ default:
+ close(fd);
+ throw std::runtime_error("GGUF 陣列元素型別不支援: " + std::to_string(et));
+ }
+ }
+ } break;
+ case 10: v.k = kv_value::INT; v.i = (int64_t) r.u64(); v.f = (double) v.i; break;
+ case 11: v.k = kv_value::INT; v.i = r.i64(); v.f = (double) v.i; break;
+ case 12: v.k = kv_value::REAL; { double d; memcpy(&d, buf.data() + r.o, 8); r.o += 8; v.f = d; v.i = (int64_t) d; } break;
+ default: close(fd); throw std::runtime_error("GGUF kv 型別不支援: " + std::to_string(t));
+ }
+ kv[key] = v;
+ }
+
+ std::vector<tensor_info> tensors;
+ tensors.reserve(n_tensors);
+ for (uint64_t i = 0; i < n_tensors; i++) {
+ tensor_info ti;
+ ti.name = r.str();
+ const uint32_t nd = r.u32();
+ for (uint32_t d = 0; d < nd; d++) {
+ ti.dims.push_back((int64_t) r.u64());
+ }
+ ti.type = (ggml_type) r.u32();
+ ti.offset = r.u64();
+ tensors.push_back(std::move(ti));
+ }
+ const uint64_t align = kv.count("general.alignment") ? (uint64_t) kv["general.alignment"].i : 32;
+ const uint64_t data_start = (r.o + align - 1) / align * align;
+
+ close(fd);
+
+ // ---- sharded GGUF：shard 1 只帶 metadata（n_tensors=0），真正的權重分散在
+ // 其他 shard，而且每個 shard 有自己的 data_start。舊版只讀第一個檔，
+ // 於是 page map 會是空的（症狀：pager 起來但沒有任何 expert 可分頁）。
+ // 這裡把每個 shard 的 header 都讀一遍，張量名稱為全域鍵，offset 記成 (shard, off)。
+ {
+ const int64_t n_split = kv.count("split.count") ? kv["split.count"].i : 1;
+ shard_paths_ = list_splits(path, (int) n_split);
+ shard_data_start_.assign(shard_paths_.size(), 0);
+ shard_data_start_[0] = data_start;
+
+ if (n_split > 1) {
+ std::vector<tensor_info> merged;
+ merged.reserve((size_t) n_tensors + 1024);
+ for (size_t si = 1; si < shard_paths_.size(); si++) {
+ const int sfd = open(shard_paths_[si].c_str(), O_RDONLY);
+ if (sfd < 0) {
+ throw std::runtime_error("無法開啟 GGUF shard：" + shard_paths_[si]);
+ }
+ struct stat sst;
+ fstat(sfd, &sst);
+ // 後續 shard 的 header 很小（沒有 tokenizer），讀 16 MiB 綽綽有餘
+ const size_t swant = (size_t) std::min<int64_t>(sst.st_size, 16ull << 20);
+ std::vector<uint8_t> sbuf(swant);
+ ssize_t sgot = 0;
+ while (sgot < (ssize_t) swant) {
+ const ssize_t n = pread(sfd, sbuf.data() + sgot, swant - sgot, sgot);
+ if (n <= 0) {
+ break;
+ }
+ sgot += n;
+ }
+ sbuf.resize((size_t) std::max<ssize_t>(sgot, 0));
+
+ gguf_reader sr{sbuf.data(), sbuf.size(), 0};
+ if (sbuf.size() < 16 || memcmp(sbuf.data(), "GGUF", 4) != 0) {
+ close(sfd);
+ throw std::runtime_error("shard 不是 GGUF：" + shard_paths_[si]);
+ }
+ sr.o = 4;
+ sr.u32(); // version
+ const uint64_t s_ntensors = sr.u64();
+ const uint64_t s_nkv = sr.u64();
+ for (uint64_t i = 0; i < s_nkv; i++) {
+ sr.str(); // key
+ const uint32_t t = sr.u32();
+ switch (t) {
+ case 0: case 1: case 7: sr.skip(1); break;
+ case 2: case 3: sr.skip(2); break;
+ case 4: case 5: case 6: sr.skip(4); break;
+ case 10: case 11: case 12: sr.skip(8); break;
+ case 8: sr.skip(sr.u64()); break;
+ case 9: {
+ const uint32_t et = sr.u32();
+ const uint64_t n = sr.u64();
+ if (et == 8) {
+ for (uint64_t j = 0; j < n; j++) { sr.skip(sr.u64()); }
+ } else {
+ size_t esz = 0;
+ switch (et) {
+ case 0: case 1: case 7: esz = 1; break;
+ case 2: case 3: esz = 2; break;
+ case 4: case 5: case 6: esz = 4; break;
+ case 10: case 11: case 12: esz = 8; break;
+ default: esz = 0; break;
+ }
+ sr.skip((uint64_t) n * esz);
+ }
+ } break;
+ default:
+ close(sfd);
+ throw std::runtime_error("shard kv 型別不支援: " + std::to_string(t));
+ }
+ }
+ for (uint64_t i = 0; i < s_ntensors; i++) {
+ tensor_info ti;
+ ti.name = sr.str();
+ const uint32_t nd = sr.u32();
+ for (uint32_t d = 0; d < nd; d++) {
+ ti.dims.push_back((int64_t) sr.u64());
+ }
+ ti.type = (ggml_type) sr.u32();
+ ti.offset = sr.u64();
+ ti.shard = (int) si;
+ merged.push_back(std::move(ti));
+ }
+ shard_data_start_[si] = (sr.o + align - 1) / align * align;
+ close(sfd);
+ }
+ for (auto & ti : tensors) {
+ ti.shard = 0;
+ merged.push_back(std::move(ti));
+ }
+ tensors.swap(merged);
+ fprintf(stderr, "[sdq] sharded GGUF：%d 個 shard，合計 %zu 個張量（shard1 為 metadata-only）\n",
+ (int) shard_paths_.size(), tensors.size());
+ }
+ }
+
+ // ---- hparams
+ auto gi = [&](const std::string & k, int64_t dflt) -> int64_t {
+ auto it = kv.find(k);
+ return it == kv.end() ? dflt : it->second.i;
+ };
+ auto garr = [&](const std::string & k) -> const std::vector<int64_t> * {
+ auto it = kv.find(k);
+ return it == kv.end() || it->second.k != kv_value::ARR ? nullptr : &it->second.arr;
+ };
+ info_.arch = kv.count("general.architecture") ? kv["general.architecture"].s : "";
+ info_.name = kv.count("general.name") ? kv["general.name"].s : "";
+ const std::string pfx = info_.arch + ".";
+ info_.n_layer = gi(pfx + "block_count", 0);
+ info_.n_expert = gi(pfx + "expert_count", 0);
+ info_.n_expert_used= gi(pfx + "expert_used_count", 0);
+ info_.n_embd = gi(pfx + "embedding_length", 0);
+ info_.n_ff_exp = gi(pfx + "expert_feed_forward_length", 0);
+ info_.n_ff_shexp = gi(pfx + "expert_shared_feed_forward_length", 0);
+ info_.n_head = gi(pfx + "attention.head_count", 0);
+ info_.n_head_kv = gi(pfx + "attention.head_count_kv", 0);
+ info_.n_embd_head_k= gi(pfx + "attention.key_length", 0);
+ info_.n_embd_head_v= gi(pfx + "attention.value_length", info_.n_embd_head_k);
+ info_.context_length = gi(pfx + "context_length", 0);
+ info_.full_attention_interval = gi(pfx + "full_attention_interval", 4);
+ // bailingmoe3（Ling-3.0-flash）：前 n_layer_dense_lead 層是 dense FFN，
+ // 它們沒有 ffn_*_exps 張量，必須讓 page map 保持「空」——
+ // sdq_pager_active_for_layer() 看到 layer() == nullptr 就會走原生路徑。
+ info_.n_layer_dense_lead = gi(pfx + "leading_dense_block_count", 0);
+ info_.n_expert_groups = gi(pfx + "expert_group_count", 0);
+ info_.n_expert_group_used = gi(pfx + "expert_group_used_count", 0);
+ info_.moe_latent_size = gi(pfx + "moe_latent_size", 0);
+
+ // Nemotron-H / Flint：expert_count / expert_used_count may be scalar or
+ // per-layer arrays. Expand either representation to a per-layer table.
+ const auto * arr_used = garr(pfx + "expert_used_count");
+ const auto * arr_exp = garr(pfx + "expert_count");
+ info_.n_expert_layer.assign((size_t) info_.n_layer, info_.n_expert);
+ info_.n_expert_used_layer.assign((size_t) info_.n_layer, info_.n_expert_used);
+ if (arr_exp && (int64_t) arr_exp->size() >= info_.n_layer) {
+ for (int64_t i = 0; i < info_.n_layer; i++) {
+ info_.n_expert_layer[(size_t) i] = (*arr_exp)[(size_t) i];
+ }
+ }
+ if (arr_used && (int64_t) arr_used->size() >= info_.n_layer) {
+ for (int64_t i = 0; i < info_.n_layer; i++) {
+ info_.n_expert_used_layer[(size_t) i] = (*arr_used)[(size_t) i];
+ }
+ // 空層（dense / Mamba）是 0，取 MoE 層的最大值當代表
+ int64_t mx = 0;
+ for (int64_t v : *arr_used) {
+ mx = std::max(mx, v);
+ }
+ info_.n_expert_used = mx;
+ }
+
+ if (info_.n_expert <= 0 || info_.n_layer <= 0) {
+ throw std::runtime_error("這不是 MoE 模型（expert_count=" + std::to_string(info_.n_expert) + "）");
+ }
+
+ // ---- 頁面表：每個 (layer, expert) 的三段
+ layers_.assign((size_t) info_.n_layer, {});
+ std::map<std::string, const tensor_info *> by_name;
+ for (const auto & ti : tensors) {
+ by_name[ti.name] = &ti;
+ }
+
+ size_t max_expert_bytes = 0;
+ int n_moe_layers = 0;
+ for (int64_t il = 0; il < info_.n_layer; il++) {
+ const std::string b = "blk." + std::to_string(il) + ".";
+ layer_desc & L = layers_[(size_t) il];
+
+ auto get = [&](const std::string & suffix) -> const tensor_info * {
+ auto it = by_name.find(b + suffix);
+ return it == by_name.end() ? nullptr : it->second;
+ };
+
+ const tensor_info * t_gate = get("ffn_gate_exps.weight");
+ const tensor_info * t_up = get("ffn_up_exps.weight");
+ const tensor_info * t_gu = get("ffn_gate_up_exps.weight");
+ const tensor_info * t_down = get("ffn_down_exps.weight");
+
+ if (!t_down || (t_gate && !t_up && !t_gu)) {
+ // 這層不是 MoE（bailingmoe3 的 leading dense layers 就是這樣）。
+ // 保持 experts 為空 → layer() 回 nullptr → 這層走原生 mul_mat_id。
+ continue;
+ }
+ // Expert counts and widths may vary by layer; use the tensor shape as
+ // the source of truth rather than relying on scalar metadata.
+ const int64_t n_expert_l = t_down->dims.size() >= 3 ? t_down->dims[2] : info_.n_expert;
+ L.experts.assign((size_t) n_expert_l, {});
+ L.n_expert = n_expert_l;
+ L.no_gate = (t_gate == nullptr && t_gu == nullptr);
+ n_moe_layers++;
+
+ // sharded GGUF：每段的絕對偏移要用「該張量自己 shard 的 data_start」，
+ // 不能用第一個 shard 的（兩者差了近 10 MB，直接用會讀到別的權重）。
+ auto abs_off = [&](const tensor_info * t, uint64_t rel) -> uint64_t {
+ const size_t s = (size_t) t->shard;
+ const uint64_t base = (s < shard_data_start_.size()) ? shard_data_start_[s] : 0;
+ return base + rel;
+ };
+
+ for (int64_t e = 0; e < n_expert_l; e++) {
+ expert_desc & ed = L.experts[(size_t) e];
+ if (t_gate || t_gu) {
+ const tensor_info * t = t_gate ? t_gate : t_gu;
+ const uint64_t nb = (uint64_t) ggml_row_size(t->type, t->dims[0]) * (uint64_t) t->dims[1];
+ ed.seg[part::GATE].off = abs_off(t, t->offset + (uint64_t) e * nb);
+ ed.seg[part::GATE].size = nb;
+ ed.seg[part::GATE].shard = t->shard;
+ ed.ne0 = t->dims[0];
+ ed.ne1 = t->dims[1];
+ ed.type[part::GATE] = t->type;
+ }
+ if (t_up) {
+ const uint64_t nb = (uint64_t) ggml_row_size(t_up->type, t_up->dims[0]) * (uint64_t) t_up->dims[1];
+ ed.seg[part::UP].off = abs_off(t_up, t_up->offset + (uint64_t) e * nb);
+ ed.seg[part::UP].size = nb;
+ ed.seg[part::UP].shard = t_up->shard;
+ ed.type[part::UP] = t_up->type; // gate/up 都是 Q4_K，型別一定要各自填
+ if (!t_gate) {
+ ed.ne0 = t_up->dims[0];
+ ed.ne1 = t_up->dims[1];
+ // no-gate 模型（Nemotron LatentMoE）不要把 GATE 當成 up，
+ // expert_forward_relu2 走 UP/DOWN 兩段。
+ }
+ }
+ {
+ const uint64_t nb = (uint64_t) ggml_row_size(t_down->type, t_down->dims[0]) * (uint64_t) t_down->dims[1];
+ ed.seg[part::DOWN].off = abs_off(t_down, t_down->offset + (uint64_t) e * nb);
+ ed.seg[part::DOWN].size = nb;
+ ed.seg[part::DOWN].shard = t_down->shard;
+ ed.type[part::DOWN] = t_down->type;
+ ed.down_ne0 = t_down->dims[0];
+ ed.down_ne1 = t_down->dims[1];
+ }
+ ed.fused_gate_up = (t_gu != nullptr);
+ ed.bytes = 0;
+ int64_t acc = 0;
+ for (int p = 0; p < 3; p++) {
+ if (ed.seg[p].size == 0) {
+ continue;
+ }
+ acc = (acc + 63) & ~(int64_t) 63; // 每段 64B 對齊
+ ed.slot_off[p] = acc;
+ acc += (int64_t) ed.seg[p].size;
+ }
+ ed.bytes = acc;
+ max_expert_bytes = std::max(max_expert_bytes, (size_t) ed.bytes);
+ }
+ L.bytes_per_expert = L.experts.empty() ? 0 : L.experts[0].bytes;
+ }
+
+ slot_bytes_ = (max_expert_bytes + 4095) & ~(size_t) 4095;
+ if (slot_bytes_ < (256ull << 10)) {
+ slot_bytes_ = 256ull << 10;
+ }
+
+ // ---- 大小兩種槽位 ----------------------------------------------------
+ // 這個模型有 37 層的 expert 是 1900544 B（Q4_K + Q5_K），但層 34/38/39 是
+ // 2039808 B（down 換成 Q6_K）。若所有槽位都照最大的配，arena 有 6.8% 被浪費
+ // ——而 arena 就是 8 GB 預算裡最寶貴的資源（實測 +5% 槽位 ≈ +9% tok/s）。
+ //
+ // 解法：槽位分成兩區。「小槽位」給 37 層用，「大槽位」只給那 3 層用，
+ // 大槽位的數量按「那 3 層佔請求的比例」分配（3/40 = 7.5%）。
+ // 選「大多數層」當小的：用眾數而不是最大值，這樣即使某個模型只有一種大小，
+ // 也不會誤判。
+ {
+ std::map<size_t, int> hist;
+ for (const auto & L : layers_) {
+ if (L.experts.empty()) {
+ continue; // dense layer：沒有 expert，不參與槽位尺寸統計
+ }
+ hist[(size_t) L.bytes_per_expert]++;
+ }
+ size_t small = 0;
+ int best_n = -1;
+ for (const auto & kv : hist) {
+ if (kv.second > best_n) {
+ best_n = kv.second;
+ small = kv.first;
+ }
+ }
+ size_t big = 0;
+ for (const auto & L : layers_) {
+ if (!L.experts.empty() && (size_t) L.bytes_per_expert > small) {
+ big = std::max(big, (size_t) L.bytes_per_expert);
+ }
+ }
+ small_bytes_ = (small + 4095) & ~(size_t) 4095;
+ big_bytes_ = big ? ((big + 4095) & ~(size_t) 4095) : 0;
+ if (big_bytes_ && big_bytes_ <= small_bytes_) {
+ big_bytes_ = 0; // 其實一樣大，別分兩區
+ }
+ slot_bytes_ = big_bytes_ ? big_bytes_ : small_bytes_;
+ int n_big_layers = 0;
+ for (auto & L : layers_) {
+ L.needs_big_slot = big_bytes_ && !L.experts.empty() &&
+ (size_t) L.bytes_per_expert > small_bytes_;
+ if (L.needs_big_slot) {
+ n_big_layers++;
+ }
+ }
+ // Nemotron has only four expert sizes; uniform slots simplify arena
+ // management. Flint layers also use a uniform slot size per layer.
+ small_bytes_ = slot_bytes_;
+ big_bytes_ = 0;
+ for (auto & L : layers_) {
+ L.needs_big_slot = false;
+ }
+ // 統計給 README / 實驗報告用：expert 權重在檔案裡到底佔多少
+ uint64_t expert_bytes = 0;
+ uint64_t n_experts_total = 0;
+ for (const auto & L : layers_) {
+ expert_bytes += (uint64_t) L.bytes_per_expert * (uint64_t) L.experts.size();
+ n_experts_total += (uint64_t) L.experts.size();
+ }
+ n_moe_layers_ = n_moe_layers;
+ expert_bytes_total_ = expert_bytes;
+ info_.no_gate = true;
+ for (const auto & L : layers_) {
+ if (!L.experts.empty() && !L.no_gate) {
+ info_.no_gate = false;
+ }
+ }
+ fprintf(stderr,
+ "[sdq] MoE 層 %d / 共 %lld 層（dense lead %lld）"
+ "：expert 權重合計 %.2f GiB、%llu 個 expert\n",
+ n_moe_layers, (long long) info_.n_layer,
+ (long long) info_.n_layer_dense_lead,
+ expert_bytes / 1073741824.0, (unsigned long long) n_experts_total);
+ fprintf(stderr,
+ "[sdq] 統一槽位 %zu B（最大 expert 大小；平均 %.0f B）\n",
+ small_bytes_, n_experts_total ? (double) expert_bytes / n_experts_total : 0.0);
+ }
+}
+
+bool pager::need_big_slot(int il, int ie) const {
+ if (!big_bytes_) {
+ return false;
+ }
+ const layer_desc * L = layer(il);
+ if (!L) {
+ return false;
+ }
+ if (L->needs_big_slot) {
+ return true;
+ }
+ return ie >= 0 && ie < (int) L->experts.size() &&
+ (size_t) L->experts[(size_t) ie].bytes > small_bytes_;
+}
+
+void pager::alloc_arena() {
+ const size_t slot = cfg_.slot_bytes ? cfg_.slot_bytes : slot_bytes_;
+ slot_bytes_ = slot;
+
+ // arena 大小決定：RAM 是「總預算」，不是「還有多少可用」
+ // 總預算 = 非專家權重 + ggml 模型緩衝 + KV cache + 計算緩衝 + 詞表 + arena
+ // 所以有指定 ram_budget 時，用「目前 RSS（不含 arena）」當作已用額度扣掉。
+ //
+ // 重要：pager 在 llama.cpp 載入權重**之前**就做 sizing，此時 rss_now 只有
+ // 幾十 MB（實測 0.09 GiB），所以後續所有匿名配置都不在 rss_now 裡，
+ // 必須由 reserve 手動涵蓋。這些事實上是三塊：
+ //
+ // 1. ggml 的模型緩衝 —— 主機有 AMX 時（-march=native 會定義 __AMX_INT8__
+ // 使 llama.cpp 啟用 AMX repack buffer）Ling-3.0-flash Q4_K_M 要
+ // 1876.34 MiB；關掉 AMX（GGML_NATIVE=OFF）則不需要這塊。
+ // 2. CPU compute buffer：83.5 MiB @ubatch=128，ubatch 越大越多。
+ // 3. KV/SSM state + 執行緒堆疊：ctx 越大越多。
+ //
+ // 舊的預設值 900 MB 只涵蓋 (2)+(3)，漏掉 (1)，在 AMX 主機上會讓實際
+ // 匿名用量超出預算約 1.9 GiB —— 嚴格 RLIMIT_DATA 下直接 OOM。
+ // 實測（本機，ctx=2048、ubatch=128、AMX build）：
+ // reserve 1024 MB → 預算 8 GiB 實測匿名峰值 10.96 GiB（超出，OOM）
+ // reserve 3600 MB → 預算 8 GiB 通過，arena 4.40 GiB
+ // reserve 2700 MB → 預算 4 GiB 通過，arena 1.27 GiB（無 AMX build）
+ // 寧可少一些槽位，也不要在目標機器的 8 GiB 上 OOM —— 這是本專案的硬約束。
+ size_t reserve = (size_t) env_size("SDQ_RESERVE_MB", 3600) << 20;
+ const size_t rss_now = process_rss_bytes();
+
+ size_t arena;
+ if (cfg_.ram_budget_bytes) {
+ const size_t used = rss_now; // 這時候 arena 還沒配置
+ size_t avail = cfg_.ram_budget_bytes > used + reserve
+ ? cfg_.ram_budget_bytes - used - reserve
+ : 0;
+ arena = avail;
+ if (cfg_.arena_bytes) {
+ arena = std::min(arena, cfg_.arena_bytes);
+ }
+ fprintf(stderr, "[sdq] RAM 預算 %.2f GiB：已用(RSS) %.2f GiB、預留 %.2f GiB → arena %.2f GiB\n",
+ cfg_.ram_budget_bytes / 1073741824.0, used / 1073741824.0,
+ reserve / 1073741824.0, arena / 1073741824.0);
+ } else if (cfg_.arena_bytes) {
+ arena = cfg_.arena_bytes; // 明確指定 SDQ_ARENA_MB
+ } else {
+ // 自動模式（未設 SDQ_ARENA_MB、未設 SDQ_RAM_BUDGET_MB）：
+ // arena = 可用記憶體 × 0.8，上限為「全部 expert 的槽位」（再大也用不到，
+ // 而且上限到達時 preload 會自動啟用）。
+ size_t total_experts = 0;
+ for (const auto & L : layers_) {
+ total_experts += L.experts.size();
+ }
+ const size_t cap = total_experts * slot_bytes_;
+ const size_t avail = available_ram_bytes();
+ arena = std::min((size_t) (avail * 0.80), cap);
+ fprintf(stderr, "[sdq] arena 自動：可用記憶體 %.2f GiB × 0.8 → %.2f GiB（全部 expert 上限 %.2f GiB）\n",
+ avail / 1073741824.0, arena / 1073741824.0, cap / 1073741824.0);
+ }
+ if (arena < slot_bytes_ * 16) {
+ throw std::runtime_error("expert arena 太小");
+ }
+ // 切成兩區：大槽位只分給「需要大槽位的層」，比例 = 那些層佔全部層的比例。
+ const bool force_uniform = env_flag("SDQ_FORCE_UNIFORM_SLOT", false);
+ if (big_bytes_ && small_bytes_ && !layers_.empty() && !force_uniform) {
+ int n_big_layers = 0;
+ for (const auto & L : layers_) {
+ if (L.needs_big_slot) {
+ n_big_layers++;
+ }
+ }
+ const double denom = (double) (n_moe_layers_ ? n_moe_layers_ : layers_.size());
+ const double f = (double) n_big_layers / denom;
+ const double avg = f * (double) big_bytes_ + (1.0 - f) * (double) small_bytes_;
+ size_t n_total = (size_t)((double) arena / avg);
+ size_t nb = (size_t)(n_total * f);
+ // 大槽位的下限不是「有點就好」，而是**工作集的下限**：
+ // 一個 token 在每個大槽位層要 8 個大 expert，這個模型有 3 個大槽位層
+ // → 至少要 3 × 8 = 24 個大槽位，而且還要留給預取的空間。
+ // 只給 16 個時量到過「找不到能放 il=34 ie=114 的槽位」而中止
+ // （這是在 verify 新增的預取壓力測試 [2b] 抓到的）。
+ const size_t big_need = (size_t) (n_big_layers * info_.n_expert_used);
+ if (nb < big_need) {
+ nb = big_need;
+ }
+ if (nb > n_total / 4) {
+ nb = n_total / 4;
+ }
+ // 連工作集下限都放不進去時，寧可讓小槽位少一點（n_total 保持不變），
+ // 否則會算出一個「必定會中止」的配置。
+ n_big_ = nb;
+ n_small_ = n_total > nb ? n_total - nb : 0;
+ arena = n_small_ * small_bytes_ + n_big_ * big_bytes_;
+ } else {
+ // 單一尺寸：全部槽位都是「大的那一種」（slot_bytes_）。
+ // 這是 SDQ_FORCE_UNIFORM_SLOT=1 與「模型只有一種 expert 大小」時的路徑。
+ //
+ // 這裡必須記成 n_small_=0 / n_big_=全部，不能記成「n_small_=全部、單位用大的」：
+ // slot_ptr()/slot_size_of() 是用「索引 < n_small_ → 小槽位」來分支的，
+ // 把它記成全部都是小槽位的話，大 expert 會拿到小槽位的 stride
+ // （寫出槽位邊界、算錯），而且大 expert 永遠配不到槽位 → 推論卡死。
+ // （第一版就是這樣，實驗跑到一半 hang 才找出來。）
+ const size_t unit = slot_bytes_ ? slot_bytes_ : small_bytes_;
+ n_small_ = 0;
+ // slot_ptr() 對 s >= n_small_ 用 big_bytes_ 當 stride，所以這裡一定要把
+ // big_bytes_ 設成 unit；設成 0 會讓所有槽位都指到 arena 開頭（數值全錯）。
+ big_bytes_ = unit;
+ n_big_ = arena / unit;
+ arena = n_big_ * unit;
+ }
+ if (n_small_ == 0 && n_big_ == 0) {
+ throw std::runtime_error("expert arena 太小（算不出槽位）");
+ }
+ n_slots_ = n_small_ + n_big_;
+ if (n_slots_ < 16) {
+ throw std::runtime_error("expert arena 太小（算不出 16 個槽位）");
+ }
+
+ void * p = nullptr;
+ const size_t alloc = arena + (2ull << 20);
+ if (posix_memalign(&p, 2ull << 20, alloc) != 0 || !p) {
+ throw std::runtime_error("無法配置 expert arena（" + std::to_string(arena >> 20) + " MB）");
+ }
+ // 讓 arena 全部變成實體 RAM，並提示核心使用 huge page（減少 TLB 壓力）
+ memset(p, 0, alloc);
+ madvise(p, alloc, MADV_HUGEPAGE);
+ arena_ = (uint8_t *) p;
+ arena_bytes_ = arena;
+
+ // OPT-1：剖面計數表（n_layer × n_expert，40×256 = 10240 筆 ≈ 40 KB，可忽略）
+ prof_.assign((size_t) info_.n_layer * (size_t) info_.n_expert, 0);
+ ensure_layer_stats((int) info_.n_layer); // SDQ-0004
+ if (!cfg_.profile_in.empty()) {
+ load_profile(cfg_.profile_in);
+ }
+ profile_out_ = cfg_.profile_out;
+
+ slots_.assign(n_slots_, {});
+ free_slots_.reserve(n_small_);
+ for (size_t i = n_small_; i-- > 0;) {
+ free_slots_.push_back((int32_t) i);
+ }
+ free_big_slots_.reserve(n_big_);
+ for (size_t i = n_slots_; i-- > n_small_;) {
+ free_big_slots_.push_back((int32_t) i);
+ }
+ used_.store(0);
+
+ hist_.assign((size_t) info_.n_layer, {});
+ for (auto & per_layer : hist_) {
+ per_layer.assign((size_t) info_.n_expert, {});
+ }
+ prev_used_.assign((size_t) info_.n_layer, {});
+ cur_used_.assign((size_t) info_.n_layer, {});
+}
+
+void pager::init(const config & cfg, ggml_context * ctx) {
+ (void) ctx;
+ if (ready_) {
+ return;
+ }
+ cfg_ = cfg;
+ if (cfg_.model_path.empty()) {
+ const char * v = getenv("SDQ_MODEL_PATH");
+ cfg_.model_path = v ? v : "";
+ }
+ if (cfg_.model_path.empty()) {
+ throw std::runtime_error("SDQ 缺少模型路徑（設定 SDQ_MODEL_PATH）");
+ }
+
+ // 先讀 header（順便列舉 shard），再開 fd：shard 清單只有在解析
+ // split.count 之後才知道，而 open_model_fd() 需要它。
+ load_gguf_header(cfg_.model_path);
+ open_model_fd();
+ if (shard_fd_.empty() || shard_fd_[0] < 0) {
+ throw std::runtime_error("無法開啟模型檔：" + shard_path(0));
+ }
+ // 順序有意義：I/O 池的對齊緩衝必須在 alloc_arena() 之前配置，
+ // 否則 arena 會把它當成「不存在的 RAM」，實際上就偷用了 8 GB 預算。
+ io_pool_start();
+ alloc_arena();
+
+ if (cfg_.prefetch) {
+ pf_ = std::thread([this] { pf_worker(); });
+ }
+ ready_ = true;
+ print_summary();
+}
+
+const layer_desc * pager::layer(int il) const {
+ if (il < 0 || il >= (int) layers_.size()) {
+ return nullptr;
+ }
+ const layer_desc & L = layers_[(size_t) il];
+ if (L.experts.empty()) {
+ return nullptr;
+ }
+ return &L;
+}
+
+void pager::open_model_fd() {
+ if (shard_paths_.empty()) {
+ shard_paths_.push_back(cfg_.model_path);
+ }
+ shard_fd_.assign(shard_paths_.size(), -1);
+ shard_fd_direct_.assign(shard_paths_.size(), -1);
+ for (size_t i = 0; i < shard_paths_.size(); i++) {
+ shard_fd_[i] = open(shard_paths_[i].c_str(), O_RDONLY);
+ if (shard_fd_[i] < 0) {
+ fprintf(stderr, "[sdq] 無法開啟 shard %zu: %s\n", i, shard_paths_[i].c_str());
+ return;
+ }
+ }
+ // O_DIRECT：讀資料完全不進 kernel page cache。
+ // 這是「RAM 帳目誠實」的關鍵 —— 否則核心預讀會讓 mmap 檔案的頁算進 RSS
+ // （實測 4K prefill 時檔案對映 RSS 可以多到 6 GB 以上），而且量到的
+ // 「SSD 讀取量」會被上一輪的 page cache 灌水。
+ const char * no_direct = getenv("SDQ_NO_DIRECT_IO");
+ if (no_direct && *no_direct && *no_direct != '0') {
+ return;
+ }
+#ifdef O_DIRECT
+ // 每個 shard 各自開一個 O_DIRECT fd：O_DIRECT 是 fd 層級的屬性，不能共用。
+ for (size_t i = 0; i < shard_paths_.size(); i++) {
+ shard_fd_direct_[i] = open(shard_paths_[i].c_str(), O_RDONLY | O_DIRECT);
+ if (shard_fd_direct_[i] < 0) {
+ break;
+ }
+ }
+ use_direct_ = shard_fd_direct_[0] >= 0;
+ if (use_direct_) {
+ fprintf(stderr, "[sdq] 使用 O_DIRECT 讀取模型檔（%zu 個 shard，不佔用 kernel page cache）\n",
+ shard_paths_.size());
+ } else {
+ fprintf(stderr, "[sdq] O_DIRECT 不可用，改用 buffered + fadvise(DONTNEED)\n");
+ }
+#endif
+}
+
+// 有對齊的緩衝區（O_DIRECT 要求），thread_local 避免競爭
+static thread_local uint8_t * g_direct_buf = nullptr;
+static thread_local size_t g_direct_sz = 0;
+
+// 與 read_range_direct 相同，但用呼叫端提供的緩衝區（已配置好、已算進 RSS）。
+// I/O 執行緒池走這條：每條 worker 在啟動時就配置好對齊緩衝，於是 arena sizing 時
+// 這筆記憶體已經在 RSS 裡，不會變成「額外偷用的 RAM」。
+// sharded GGUF：把「哪個 shard」一路帶到 pread。
+static inline int sdq_fd_for(const std::vector<int> & fds, int shard) {
+ if (shard < 0 || shard >= (int) fds.size()) {
+ return -1;
+ }
+ return fds[(size_t) shard];
+}
+
+bool pager::read_range_staged(int shard, uint64_t off, uint8_t * dst, size_t size,
+ uint8_t * stage, size_t stage_sz) {
+ if (!use_direct_) {
+ return read_range(shard, off, dst, size);
+ }
+ const int fd = sdq_fd_for(shard_fd_direct_, shard);
+ if (fd < 0) {
+ return false;
+ }
+ const size_t blk = 4096;
+ const uint64_t start = off & ~(uint64_t)(blk - 1);
+ const uint64_t end = (off + size + blk - 1) & ~(uint64_t)(blk - 1);
+ const size_t bytes = (size_t)(end - start);
+ if (stage_sz < bytes) {
+ // 預先配置不足（不該發生）：退回 thread_local 的通用路徑
+ return read_range_direct(shard, off, dst, size);
+ }
+ size_t done = 0;
+ while (done < bytes) {
+ const ssize_t n = pread(fd, stage + done, bytes - done,
+ (off_t)(start + done));
+ if (n <= 0) {
+ return false;
+ }
+ done += (size_t) n;
+ }
+ memcpy(dst, stage + (off - start), size);
+ return true;
+}
+
+bool pager::read_range_direct(int shard, uint64_t off, uint8_t * dst, size_t size) {
+ if (!use_direct_) {
+ return read_range(shard, off, dst, size);
+ }
+ const int fd = sdq_fd_for(shard_fd_direct_, shard);
+ if (fd < 0) {
+ return false;
+ }
+ const size_t blk = 4096;
+ const uint64_t start = off & ~(uint64_t)(blk - 1);
+ const uint64_t end = (off + size + blk - 1) & ~(uint64_t)(blk - 1);
+ const size_t bytes = (size_t)(end - start);
+ if (g_direct_sz < bytes) {
+ if (g_direct_buf) {
+ free(g_direct_buf);
+ }
+ if (posix_memalign((void **) &g_direct_buf, blk, bytes) != 0) {
+ g_direct_buf = nullptr;
+ g_direct_sz = 0;
+ use_direct_ = false;
+ return read_range(shard, off, dst, size);
+ }
+ g_direct_sz = bytes;
+ }
+ size_t done = 0;
+ while (done < bytes) {
+ const ssize_t n = pread(fd, g_direct_buf + done, bytes - done,
+ (off_t)(start + done));
+ if (n <= 0) {
+ return false;
+ }
+ done += (size_t) n;
+ }
+ memcpy(dst, g_direct_buf + (off - start), size);
+ return true;
+}
+
+bool pager::read_range(int shard, uint64_t off, uint8_t * dst, size_t size) {
+ const int fd = sdq_fd_for(shard_fd_, shard);
+ if (fd < 0) {
+ return false;
+ }
+ size_t done = 0;
+ while (done < size) {
+ const ssize_t n = pread(fd, dst + done, size - done, (off_t) (off + done));
+ if (n <= 0) {
+ return false;
+ }
+ done += (size_t) n;
+ }
+ // 讓 RAM 帳目誠實：把 kernel 對這個檔案的 page cache 丟掉，避免「藏在核心裡的 RAM」
+ if (cfg_.fadvise_dontneed) {
+ posix_fadvise(fd, (off_t) off, (off_t) size, POSIX_FADV_DONTNEED);
+ }
+ return true;
+}
+
+bool pager::read_expert(int il, int ie, uint8_t * dst) {
+ const layer_desc * L = layer(il);
+ if (!L || ie < 0 || ie >= (int) L->experts.size()) {
+ return false;
+ }
+ const expert_desc & ed = L->experts[(size_t) ie];
+ for (int p = 0; p < 3; p++) {
+ if (ed.seg[p].size == 0) {
+ continue;
+ }
+ if (!read_range_direct(ed.seg[p].shard, ed.seg[p].off,
+ dst + ed.slot_off[p], (size_t) ed.seg[p].size)) {
+ return false;
+ }
+ }
+ return true;
+}
+
+// ---------------------------------------------------------------- I/O 執行緒池
+//
+// 問題：一個 expert = gate / up / down 三段，檔案上是三段不相鄰的資料。
+// 循序讀就是 3 次 pread 排在同一條執行緒上，而單次 1.9 MB 的 pread 在本機量到
+// ~7 ms（≈275 MB/s 的單流頻寬）。因此「一個 expert 要多久」= 3 × 單次延遲，
+// 這條路徑的長度跟「有多少條執行緒」無關 —— mode C 再怎麼排程都救不了。
+//
+// 做法：把三段交給三條執行緒同時 pread，臨界路徑降到約「單段」的時間。
+// 量測（8 個 expert 全部 miss 的一層，各 9 次取中位數）：
+// 8 執行緒 × 三段序列讀 : 12.60 ms
+// 16 執行緒 × 三段平行讀 : 7.57 ms
+// 24 執行緒 × 三段平行讀 : 7.58 ms ← 16 就飽和，所以預設 3 × n_threads
+//
+// 這個池子對 prefill 更有用：一層要 256 個 expert = 768 個讀取工作，
+// 序列讀時 16 執行緒只能拿到 16 路的頻寬 utilization。
+void pager::io_pool_start() {
+ if (!cfg_.io_parallel) {
+ fprintf(stderr, "[sdq] I/O 段平行：關（每個 expert 的 gate/up/down 循序讀）\n");
+ return;
+ }
+ int n = cfg_.io_threads > 0 ? cfg_.io_threads : cfg_.n_threads * 3;
+ if (n < 4) {
+ n = 4;
+ }
+if (n > 48) {
+ n = 48;
+ }
+ // 先算「一段權重最大有多少位元組」，每條 worker 配一塊對齊緩衝。
+ // 這必须在 alloc_arena() 之前做完，否則這筆 RAM 會變成帳外支出。
+ size_t max_seg = 0;
+ for (const auto & L : layers_) {
+ for (const auto & ed : L.experts) {
+ for (int p = 0; p < 3; p++) {
+ max_seg = std::max(max_seg, (size_t) ed.seg[p].size);
+ }
+ }
+ }
+ const size_t stage_bytes = ((max_seg + 2 * 4096) + 4095) & ~(size_t) 4095;
+ io_stage_.assign((size_t) n, io_stage{});
+ for (int i = 0; i < n; i++) {
+ void * p = nullptr;
+ if (posix_memalign(&p, 4096, stage_bytes) == 0 && p) {
+ memset(p, 0, stage_bytes); // 實際碰頁，讓它出現在 RSS 裡
+ io_stage_[(size_t) i].buf = (uint8_t *) p;
+ io_stage_[(size_t) i].sz = stage_bytes;
+ }
+ }
+ io_split_ = cfg_.io_split > 0 ? cfg_.io_split : 1;
+ io_stop_ = false;
+ io_pool_.reserve((size_t) n);
+ for (int i = 0; i < n; i++) {
+ const size_t wid = (size_t) i;
+ try {
+ io_pool_.emplace_back([this, wid] { io_worker(wid); });
+ } catch (...) {
+ break;
+ }
+ }
+ io_threads_ = (int) io_pool_.size();
+ io_pool_on_ = io_threads_ > 0;
+ if (io_pool_on_) {
+ fprintf(stderr, "[sdq] I/O 段平行：%d 條讀取執行緒 × %zu KB 對齊緩衝"
+ "（gate/up/down 同時讀，每段再切 %d 塊）\n",
+ io_threads_, stage_bytes >> 10, io_split_);
+ }
+}
+
+void pager::io_pool_stop() {
+ if (io_pool_.empty()) {
+ for (auto & s : io_stage_) {
+ free(s.buf);
+ s.buf = nullptr;
+ }
+ io_stage_.clear();
+ return;
+ }
+ {
+ std::lock_guard<std::mutex> lk(io_mu_);
+ io_stop_ = true;
+ }
+ io_cv_.notify_all();
+ for (auto & t : io_pool_) {
+ if (t.joinable()) {
+ t.join();
+ }
+ }
+ io_pool_.clear();
+ io_pool_on_ = false;
+ for (auto & s : io_stage_) {
+ free(s.buf);
+ s.buf = nullptr;
+ s.sz = 0;
+ }
+ io_stage_.clear();
+}
+
+void pager::io_worker(size_t wid) {
+ for (;;) {
+ io_job job{};
+ {
+ std::unique_lock<std::mutex> lk(io_mu_);
+ io_cv_.wait(lk, [this] { return io_stop_ || !io_q_.empty(); });
+ if (io_q_.empty()) {
+ if (io_stop_) {
+ return;
+ }
+ continue;
+ }
+ job = io_q_.front();
+ io_q_.pop_front();
+ }
+ io_jobs_.fetch_add(1, std::memory_order_relaxed);
+ // 讀取本身不持鎖：io_mu_ 只保護佇列與完成計數
+ const bool ok = (wid < io_stage_.size() && io_stage_[wid].buf)
+ ? read_range_staged(job.shard, job.off, job.dst, job.size,
+ io_stage_[wid].buf, io_stage_[wid].sz)
+ : read_range_direct(job.shard, job.off, job.dst, job.size);
+ {
+ std::lock_guard<std::mutex> lk(io_mu_);
+ if (!ok) {
+ job.ctx->ok = false;
+ }
+ if (job.ctx->remaining.fetch_sub(1) == 1) {
+ io_done_cv_.notify_all();
+ }
+ }
+ }
+}
+
+// 把 (il,ie) 的三段交給執行緒池並行讀，呼叫端等到整批完成才返回。
+// ctx 放在呼叫端 stack：worker 只會在 remaining 歸零之後才再碰它，
+// 而我們正是等到 remaining==0 才離開函式，所以沒有 use-after-return。
+bool pager::read_expert_par(int il, int ie, uint8_t * dst) {
+ const layer_desc * L = layer(il);
+ if (!L || ie < 0 || ie >= (int) L->experts.size()) {
+ return false;
+ }
+ if (!io_pool_on_) {
+ return read_expert(il, ie, dst);
+ }
+ const expert_desc & ed = L->experts[(size_t) ie];
+ io_ctx ctx;
+ int n = 0;
+ // 把每一段再切成一樣大的好幾塊，交給不同的執行緒同時讀。
+ //
+ // 為什麼：單一 pread 的時間 ≈ 固定延遲 + 位元組數 / 單流頻寬。
+ // 本機實測單流約 275 MB/s、4K 隨機讀延遲約 0.57 ms，所以
+ // 590 KB 讀一次 ≈ 0.57 + 2.1 = 2.7 ms
+ // 590 KB 切成 2 塊 ≈ 0.57 + 1.05 = 1.6 ms
+ // 而「讀完一個 expert」的時間正是每層的關鍵路徑，切塊就是直接砍關鍵路徑。
+ // 設定檔 SDQ_IO_SPLIT。**實測預設是 1（不切塊）**：切塊在 24 條讀取執行緒
+ // 之下沒有可重現的改善（split=1..6 的中位數在 4.9~5.9 之間跳動，純噪音），
+ // 因為 8 個 expert × 3 段 = 24 個並行請求已經把裝置的佇列灌飽，
+ // 這時瓶頸是裝置總頻寬而不是單一請求的延遲。保留這個開關是為了其他裝置。
+ const int split = io_split_;
+ const size_t blk = 4096;
+ {
+ std::lock_guard<std::mutex> lk(io_mu_);
+ for (int p = 0; p < 3; p++) {
+ const size_t sz = (size_t) ed.seg[p].size;
+ if (sz == 0) {
+ continue;
+ }
+ uint8_t * base = dst + ed.slot_off[p];
+ const int sh = ed.seg[p].shard;
+ if (split <= 1 || sz <= 2 * blk) {
+ io_q_.push_back(io_job{sh, ed.seg[p].off, base, sz, &ctx});
+ n++;
+ continue;
+ }
+ size_t chunk = ((sz + (size_t) split - 1) / (size_t) split + blk - 1) & ~(blk - 1);
+ for (size_t off = 0; off < sz; off += chunk) {
+ const size_t take = sz - off < chunk ? sz - off : chunk;
+ io_q_.push_back(io_job{sh, ed.seg[p].off + off, base + off, take, &ctx});
+ n++;
+ }
+ }
+ ctx.remaining.store(n);
+ }
+ if (n == 0) {
+ return true;
+ }
+ io_cv_.notify_all();
+ const uint64_t t0 = now_us();
+ {
+ std::unique_lock<std::mutex> lk(io_mu_);
+ io_done_cv_.wait(lk, [&ctx] { return ctx.remaining.load() == 0; });
+ }
+ io_wait_ns_.fetch_add((now_us() - t0) * 1000ull, std::memory_order_relaxed);
+ return ctx.ok;
+}
+
+uint64_t pager::read_expert_span(int il, const int32_t * experts, int n, const int32_t * slots) {
+ const layer_desc * L = layer(il);
+ if (!L || n <= 0) {
+ return 0;
+ }
+ const uint64_t t_start = now_us();
+ // 每個執行緒一份暫存區（多執行緒會同時呼叫，不能共用一個 buffer）
+ static thread_local std::vector<uint8_t> scratch_buf;
+ uint64_t total = 0;
+
+ // 把 experts 切成「檔案編號連續」的區段
+ int i = 0;
+ while (i < n) {
+ int j = i + 1;
+ while (j < n && experts[j] == experts[j - 1] + 1) {
+ j++;
+ }
+ const int cnt = j - i;
+ const expert_desc & first = L->experts[(size_t) experts[i]];
+ const expert_desc & last = L->experts[(size_t) experts[j - 1]];
+
+ for (int p = 0; p < 3; p++) {
+ if (first.seg[p].size == 0 || last.seg[p].size == 0) {
+ continue;
+ }
+ // 連續區段必須在同一個 shard 內；跨 shard 就切斷這個區段。
+ if (last.seg[p].shard != first.seg[p].shard) {
+ continue;
+ }
+ const uint64_t off = first.seg[p].off;
+ const uint64_t end = last.seg[p].off + last.seg[p].size;
+ const size_t bytes = (size_t) (end - off);
+ if (scratch_buf.size() < bytes) {
+ scratch_buf.assign(bytes, 0);
+ }
+ if (!read_range_direct(first.seg[p].shard, off, scratch_buf.data(), bytes)) {
+ return total;
+ }
+ total += bytes;
+ // 散到各自的槽位：scratch 從 off 開始，所以偏移是「相對於 off」
+ for (int k = i; k < j; k++) {
+ const expert_desc & ed = L->experts[(size_t) experts[k]];
+ const size_t seg_bytes = (size_t) ed.seg[p].size;
+ const size_t delta = (size_t) (ed.seg[p].off - off);
+ if (slots[k] < 0) {
+ continue;
+ }
+ if ((size_t) slots[k] >= n_slots_ || (uint64_t) ed.slot_off[p] + seg_bytes > slot_bytes_
+ || delta + seg_bytes > scratch_buf.size()) {
+ fprintf(stderr, "[sdq-pager] span 越界：il=%d slot=%d/%zu slot_off=%lld seg=%zu delta=%zu scratch=%zu p=%d e=%d\n",
+ il, slots[k], n_slots_, (long long) ed.slot_off[p], seg_bytes, delta, scratch_buf.size(), p, experts[k]);
+ return total;
+ }
+ memcpy(slot_ptr(slots[k]) + ed.slot_off[p],
+ scratch_buf.data() + delta, seg_bytes);
+ }
+ }
+ i = j;
+ }
+ stats_.demand_wait_us.fetch_add(now_us() - t_start, std::memory_order_relaxed);
+ return total;
+}
+
+int pager::reserve_slot(int il, int ie) {
+ if (!ready_) {
+ return -1;
+ }
+ count_request(il, ie);
+ std::lock_guard<std::mutex> lk(mu_);
+ const int32_t exist = find_slot_locked(il, ie);
+ if (exist >= 0) {
+ slots_[(size_t) exist].last_use = ++clock_;
+ slots_[(size_t) exist].pin++;
+ touch_freq_locked(slots_[(size_t) exist]);
+ return exist;
+ }
+ const int32_t s = victim_locked();
+ if (s < 0) {
+ return -1;
+ }
+ slot_state & st = slots_[(size_t) s];
+ if (st.layer >= 0) {
+ page_.erase(((int64_t) st.layer << 20) | (int64_t) st.expert);
+ stats_.evictions.fetch_add(1, std::memory_order_relaxed);
+ used_.fetch_sub(1, std::memory_order_relaxed);
+ }
+ st.layer = il;
+ st.expert = ie;
+ st.last_use = ++clock_;
+ st.freq = 1;
+ st.pin = 1;
+ st.is_hot = hot_lookup(il, ie);
+ page_[((int64_t) il << 20) | (int64_t) ie] = s;
+ used_.fetch_add(1, std::memory_order_relaxed);
+ stats_.expert_misses.fetch_add(1, std::memory_order_relaxed);
+ return s;
+}
+
+int pager::probe(int il, int ie) {
+ if (!ready_) {
+ return -1;
+ }
+ stats_.expert_requests.fetch_add(1, std::memory_order_relaxed);
+ count_request(il, ie);
+ std::lock_guard<std::mutex> lk(mu_);
+ const int32_t s = find_slot_locked(il, ie);
+ if (s < 0) {
+ stats_.expert_misses.fetch_add(1, std::memory_order_relaxed);
+ return -1;
+ }
+ slots_[(size_t) s].last_use = ++clock_;
+ slots_[(size_t) s].pin++;
+ touch_freq_locked(slots_[(size_t) s]);
+ stats_.expert_hits.fetch_add(1, std::memory_order_relaxed);
+ return s;
+}
+
+bool pager::preload_all(int n_threads) {
+ if (!ready_) {
+ return false;
+ }
+ size_t total = 0;
+ for (int il = 0; il < (int) info_.n_layer; il++) {
+ const layer_desc * L = layer(il);
+ if (L) {
+ total += L->experts.size();
+ }
+ }
+ if (total == 0 || n_slots_ < total) {
+ return false;
+ }
+ const int nt = std::max(1, n_threads);
+ std::atomic<size_t> failed{0};
+ std::vector<std::thread> th;
+ th.reserve((size_t) nt);
+ for (int t = 0; t < nt; t++) {
+ th.emplace_back([this, t, nt, &failed] {
+ for (int il = 0; il < (int) info_.n_layer; il++) {
+ const layer_desc * L = layer(il);
+ if (!L) {
+ continue;
+ }
+ for (int ie = t; ie < (int) L->experts.size(); ie += nt) {
+ const int s = acquire(il, ie, true);
+ if (s < 0) {
+ failed.fetch_add(1, std::memory_order_relaxed);
+ } else {
+ release(s);
+ }
+ }
+ }
+ });
+ }
+ for (auto & x : th) {
+ x.join();
+ }
+ reset_stats();
+ return failed.load() == 0;
+}
+
+void pager::reset_stats() {
+ auto z = [](std::atomic<uint64_t> & a) { a.store(0, std::memory_order_relaxed); };
+ z(stats_.expert_requests); z(stats_.expert_hits); z(stats_.expert_misses);
+ z(stats_.ssb_read_bytes_demand); z(stats_.ssb_read_bytes_prefetch);
+ z(stats_.prefetch_issued); z(stats_.prefetch_completed); z(stats_.prefetch_used);
+ z(stats_.prefetch_duplicate); z(stats_.evictions); z(stats_.demand_wait_us);
+ z(stats_.prefetch_wait_us); z(stats_.prefetch_bytes_used); z(stats_.tokens);
+ z(stats_.moe_calls); z(stats_.flops);
+ if (lstats_) {
+ for (int il = 0; il < (int) info_.n_layer; il++) {
+ z(lstats_[il].req); z(lstats_[il].hit); z(lstats_[il].miss);
+ z(lstats_[il].demand_bytes); z(lstats_[il].prefetch_bytes);
+ }
+ }
+}
+
+bool pager::resident(int il, int ie) {
+ if (!ready_) {
+ return false;
+ }
+ std::lock_guard<std::mutex> lk(mu_);
+ const int32_t s = find_slot_locked(il, ie);
+ return s >= 0 && !slots_[(size_t) s].busy;
+}
+
+void pager::note_span_read(uint64_t bytes) {
+ stats_.ssb_read_bytes_demand.fetch_add(bytes, std::memory_order_relaxed);
+ stats_.demand_wait_us.fetch_add(0, std::memory_order_relaxed);
+}
+
+// 衰減式頻率：每 8192 次使用就把所有計數減半，讓「過去熱門、現在變冷」的 expert 能被淘汰
+void pager::touch_freq_locked(slot_state & st) {
+ if (++freq_tick_ >= 8192) {
+ freq_tick_ = 0;
+ for (auto & s : slots_) {
+ s.freq /= 2;
+ }
+ }
+ if (st.freq < 0xFFFFFFFFu) {
+ st.freq++;
+ }
+}
+
+// ============================== OPT-1：熱門 expert 剖面 ==============================
+// 問題實測（docs/03）：arena 3.88 GiB 只放得下 2044 個 expert，總共卻有 10240 個，
+// prefill 的幾百個「只用一次」的專家會把熱區整個洗掉，命中率卡在 ~70%。
+// 解法借用 Strata 的做法（src/core/expert_cache.cpp）：**離線剖面 + 靜態排行榜 +
+// 熱頁永不淘汰**。這裡不動檔案格式，用純文字存排行榜。
+
+// ---- SDQ-0004
+void pager::ensure_layer_stats(int n_layer) {
+ if (lstats_ || n_layer <= 0) {
+ return;
+ }
+ lstats_.reset(new layer_stat[(size_t) n_layer]);
+}
+
+void pager::count_request(int il, int ie) {
+ if (prof_.empty() || il < 0 || ie < 0) {
+ return;
+ }
+ const size_t idx = (size_t) il * (size_t) info_.n_expert + (size_t) ie;
+ if (idx < prof_.size()) {
+ // 單執行緒 ++ 即可：請求次數是統計用途，不參與任何正確性判斷
+ if (prof_[idx] < 0xFFFFFFFFu) {
+ prof_[idx]++;
+ }
+ }
+}
+
+bool pager::hot_lookup(int il, int ie) const {
+ if (hot_.empty() || il < 0 || ie < 0) {
+ return false;
+ }
+ const size_t idx = (size_t) il * (size_t) info_.n_expert + (size_t) ie;
+ return idx < hot_.size() && hot_[idx] != 0;
+}
+
+// 挑一個容得下 (il,ie) 的槽。與 victim_locked 的差別：**先只在冷頁裡找**，
+// 找不到才退讓去動熱頁。這就是「熱區不會被 prefill 洗掉」的機制。
+// 這個槽位放得下「需要 big 的 expert」嗎？
+// 小槽位：只能放小 expert
+// 大槽位：小expert 也放得下（只是浪費一點空間）
+static inline bool slot_accepts(bool need_big, bool slot_is_big) {
+ return !need_big || slot_is_big;
+}
+
+int pager::admit_slot_locked(int il, int ie) {
+ // il/ie 都 < 0 代表「不指定大小」→ 兩種槽位都可以（舊的呼叫路徑）。
+ const bool need_big = (il >= 0 && ie >= 0) ? need_big_slot(il, ie) : false;
+ // 完全沒用過的空槽永遠優先（它不是熱頁，佔用等於浪費）
+ // 大 expert 先看大槽位；小 expert 先看小槽位。
+ std::vector<int32_t> & first = need_big ? free_big_slots_ : free_slots_;
+ std::vector<int32_t> & second = need_big ? free_slots_ : free_big_slots_;
+ while (!first.empty()) {
+ const int32_t s = first.back();
+ if (!slots_[(size_t) s].busy) {
+ first.pop_back();
+ return s;
+ }
+ }
+ // 第一輪：只在「冷頁」裡挑最久未用的。
+ // want_hot=true → 熱門 expert 進來時，把冷頁換掉（不動其他熱頁）
+ // want_hot=false → 冷門 expert 進來時，優先吃冷頁
+ // 兩種情況的規則相同，所以 hot_lookup 的結果不影響選擇順序，只影響標記。
+ for (int pass = 0; pass < 2; pass++) {
+ int64_t best = -1;
+ uint64_t best_clock = UINT64_MAX;
+ for (size_t s = 0; s < n_slots_; s++) {
+ const slot_state & st = slots_[s];
+ if (st.pin > 0 || st.busy) {
+ continue;
+ }
+ if (pass == 0 && st.is_hot) {
+ continue; // 第一輪不動熱頁
+ }
+ if (!slot_accepts(need_big, is_big_slot(s))) {
+ continue;
+ }
+ if (st.last_use < best_clock) {
+ best_clock = st.last_use;
+ best = (int64_t) s;
+ }
+ }
+ if (best >= 0) {
+ return (int) best;
+ }
+ }
+ // 退讓：另一種大小的空槽。**一定要檢查放不放得下** —— 小槽位放不下大 expert，
+ // 硬塞會讓 read_expert 寫出槽位邊界（靜默把別的 expert 寫壞）。
+ while (!second.empty()) {
+ const int32_t s = second.back();
+ if (!slots_[(size_t) s].busy &&
+ slot_accepts(need_big, is_big_slot(s))) {
+ second.pop_back();
+ return s;
+ }
+ if (!slots_[(size_t) s].busy) {
+ second.pop_back();
+ } else {
+ break;
+ }
+ }
+ return -1;
+}
+
+// least_stale：只淘汰「閒置 >= stale_window」的頁；在合格的頁裡挑 decayed freq
+// 最小的（最少被重複使用），同 freq 再挑最久未用。熱頁第一輪不動。
+// 沒有合格頁時回 -1，由 victim_locked 退回 LRU。
+int pager::victim_least_stale_locked(int il, int ie, uint64_t stale_before) {
+ const bool need_big = (il >= 0 && ie >= 0) ? need_big_slot(il, ie) : false;
+ std::vector<int32_t> & first = need_big ? free_big_slots_ : free_slots_;
+ std::vector<int32_t> & second = need_big ? free_slots_ : free_big_slots_;
+ for (std::vector<int32_t> * lst : { &first, &second }) {
+ while (!lst->empty()) {
+ const int32_t s = lst->back();
+ if (slots_[(size_t) s].busy) {
+ break;
+ }
+ if (slots_[(size_t) s].last_use <= stale_before &&
+ slot_accepts(need_big, is_big_slot(s))) {
+ lst->pop_back();
+ return s;
+ }
+ lst->pop_back();
+ }
+ }
+ for (int pass = 0; pass < 2; pass++) {
+ int64_t best = -1;
+ uint32_t best_freq = UINT32_MAX;
+ uint64_t best_clock = UINT64_MAX;
+ for (size_t i = 0; i < n_slots_; i++) {
+ const slot_state & st = slots_[i];
+ if (st.pin > 0 || st.busy || st.last_use > stale_before) {
+ continue;
+ }
+ if (pass == 0 && st.is_hot) {
+ continue;
+ }
+ if (!slot_accepts(need_big, is_big_slot(i))) {
+ continue;
+ }
+ if (st.freq < best_freq ||
+ (st.freq == best_freq && st.last_use < best_clock)) {
+ best_freq = st.freq;
+ best_clock = st.last_use;
+ best = (int64_t) i;
+ }
+ }
+ if (best >= 0) {
+ return (int) best;
+ }
+ }
+ return -1;
+}
+
+// 依 SDQ_EVICT_POLICY 選擇 victim。兩條路徑都必須回傳「放得下 (il,ie)」的槽。
+int pager::victim_locked(int il, int ie) {
+ if (cfg_.evict_policy == "least_stale") {
+ const uint64_t w = cfg_.stale_window;
+ const uint64_t stale_before = clock_ > w ? clock_ - w : 0;
+ const int s = victim_least_stale_locked(il, ie, stale_before);
+ if (s >= 0) {
+ return s;
+ }
+ // 所有頁都還很新（或都被 pin/busy）→ 退回全域 LRU，保證推論不會卡住。
+ }
+ return admit_slot_locked(il, ie);
+}
+
+bool pager::load_profile(const std::string & path) {
+ std::FILE * f = std::fopen(path.c_str(), "rb");
+ if (!f) {
+ fprintf(stderr, "[sdq] 找不到 expert 剖面 %s（將以動態 LRU/LFU 運作）\n", path.c_str());
+ return false;
+ }
+ char magic[8] = {0};
+ if (std::fread(magic, 1, 4, f) != 4 || std::memcmp(magic, "SDQP", 4) != 0) {
+ std::fclose(f);
+ fprintf(stderr, "[sdq] %s 不是 SDQ expert 剖面（缺少 SDQP 標頭）\n", path.c_str());
+ return false;
+ }
+ uint32_t hdr[3] = {0, 0, 0};
+ if (std::fread(hdr, sizeof(uint32_t), 3, f) != 3) {
+ std::fclose(f);
+ return false;
+ }
+ const uint32_t n = hdr[2];
+ std::vector<std::pair<uint32_t, std::pair<int32_t, int32_t>>> ranked;
+ ranked.reserve(n);
+ for (uint32_t i = 0; i < n; i++) {
+ int32_t l = 0, e = 0;
+ uint32_t c = 0;
+ if (std::fread(&l, sizeof(int32_t), 1, f) != 1) break;
+ if (std::fread(&e, sizeof(int32_t), 1, f) != 1) break;
+ if (std::fread(&c, sizeof(uint32_t), 1, f) != 1) break;
+ ranked.push_back({c, {l, e}});
+ }
+ std::fclose(f);
+ std::sort(ranked.begin(), ranked.end(),
+ [](const auto & a, const auto & b) { return a.first > b.first; });
+
+ hot_.assign((size_t) info_.n_layer * (size_t) info_.n_expert, 0);
+ // 預設釘選 25 % 的槽位（模擬顯示這是 1 GiB 階層的最佳點）。舊版沒有設定時
+ // 會把所有出現過的 expert 都標成熱門，等於沒有冷槽可淘汰。
+ const int want = cfg_.hot_pin > 0 ? cfg_.hot_pin
+ : std::max<int>(1, (int) (n_slots_ / 4));
+ n_hot_ = 0;
+ for (const auto & r : ranked) {
+ if (n_hot_ >= want) {
+ break;
+ }
+ const int l = r.second.first, e = r.second.second;
+ if (l < 0 || l >= info_.n_layer || e < 0 || e >= info_.n_expert) {
+ continue;
+ }
+ const size_t idx = (size_t) l * (size_t) info_.n_expert + (size_t) e;
+ if (!hot_[idx]) {
+ hot_[idx] = 1;
+ n_hot_++;
+ }
+ }
+ fprintf(stderr, "[sdq] expert 剖面：%s → 釘選 %d / %zu 個熱門 expert（永不淘汰）\n",
+ path.c_str(), n_hot_, ranked.size());
+ return n_hot_ > 0;
+}
+
+void pager::dump_profile(const std::string & path) {
+ if (prof_.empty()) {
+ return;
+ }
+ std::vector<std::pair<uint32_t, std::pair<int32_t, int32_t>>> ranked;
+ ranked.reserve(prof_.size());
+ uint64_t total = 0;
+ for (int l = 0; l < info_.n_layer; l++) {
+ for (int e = 0; e < info_.n_expert; e++) {
+ const uint32_t c = prof_[(size_t) l * (size_t) info_.n_expert + (size_t) e];
+ if (c > 0) {
+ ranked.push_back({c, {l, e}});
+ total += c;
+ }
+ }
+ }
+ std::sort(ranked.begin(), ranked.end(),
+ [](const auto & a, const auto & b) { return a.first > b.first; });
+ std::FILE * f = std::fopen(path.c_str(), "wb");
+ if (!f) {
+ fprintf(stderr, "[sdq] 無法寫出 expert 剖面 %s\n", path.c_str());
+ return;
+ }
+ std::fwrite("SDQP", 1, 4, f);
+ const uint32_t hdr[3] = {1, (uint32_t) info_.n_layer, (uint32_t) ranked.size()};
+ std::fwrite(hdr, sizeof(uint32_t), 3, f);
+ for (const auto & r : ranked) {
+ const int32_t l = r.second.first, e = r.second.second;
+ const uint32_t c = r.first;
+ std::fwrite(&l, sizeof(int32_t), 1, f);
+ std::fwrite(&e, sizeof(int32_t), 1, f);
+ std::fwrite(&c, sizeof(uint32_t), 1, f);
+ }
+ std::fclose(f);
+ // 覆盖率報告：前 K 名可以蓋掉多少請求 → 直接決定 hot_pin 該設多少
+ uint64_t acc = 0;
+ const size_t ks[6] = {512, 1024, 1536, 2048, 3072, 4096};
+ fprintf(stderr, "[sdq] expert 剖面已寫到 %s（%zu 個有被用到，總請求 %llu）\n",
+ path.c_str(), ranked.size(), (unsigned long long) total);
+ size_t ki = 0;
+ for (size_t i = 0; i < ranked.size(); i++) {
+ acc += ranked[i].first;
+ if (ki < 6 && i + 1 == ks[ki]) {
+ fprintf(stderr, "[sdq] 前 %5zu 名涵蓋 %6.2f%% 的請求\n",
+ ks[ki], 100.0 * (double) acc / (double) (total ? total : 1));
+ ki++;
+ }
+ }
+}
+
+int pager::find_slot_locked(int il, int ie) {
+ const int64_t key = ((int64_t) il << 20) | (int64_t) ie;
+ auto it = page_.find(key);
+ return it == page_.end() ? -1 : it->second;
+}
+
+// 找一個「可以放 (il,ie) 而且 last_use <= stale_before」的槽位。
+//
+// 為什麼不能直接用 victim_stale_locked()：大小兩種槽位之後，那個函式會掃過**所有**
+// 槽位（包含小槽位），於是一個需要大槽位的 expert 可能拿到小槽位，
+// 預取時 read_expert 會寫出槽位邊界 → 靜默把鄰居槽位的資料寫壞。
+// 這是引進大小兩種槽位時真正會咬人的地方。
+//
+// 順序：先看對應尺寸的空槽（從沒用過，最該拿來放預取），
+// 再掃全部槽位挑最久未用的（放得下的那種尺寸）。
+int pager::victim_stale_for(int il, int ie, uint64_t stale_before) {
+ const bool need_big = need_big_slot(il, ie);
+ std::vector<int32_t> & first = need_big ? free_big_slots_ : free_slots_;
+ std::vector<int32_t> & second = need_big ? free_slots_ : free_big_slots_;
+ for (std::vector<int32_t> * lst : { &first, &second }) {
+ while (!lst->empty()) {
+ const int32_t s = lst->back();
+ if (slots_[(size_t) s].busy) {
+ break;
+ }
+ if (slots_[(size_t) s].last_use <= stale_before &&
+ slot_accepts(need_big, is_big_slot(s))) {
+ lst->pop_back();
+ return s;
+ }
+ lst->pop_back();
+ }
+ }
+ int64_t best = -1;
+ uint64_t best_clock = UINT64_MAX;
+ for (size_t i = 0; i < n_slots_; i++) {
+ const slot_state & st = slots_[i];
+ if (st.pin > 0 || st.busy || st.last_use > stale_before) {
+ continue;
+ }
+ if (!slot_accepts(need_big, is_big_slot(i))) {
+ continue;
+ }
+ if (st.last_use < best_clock) {
+ best_clock = st.last_use;
+ best = (int64_t) i;
+ }
+ }
+ return (int) best;
+}
+
+int pager::victim_stale_locked(uint64_t stale_before) {
+ // 只挑「last_use 很舊」而且沒被 pin 的槽：保護最近的工作集
+ while (!free_slots_.empty()) {
+ const int32_t s = free_slots_.back();
+ if (!slots_[(size_t) s].busy && slots_[(size_t) s].last_use <= stale_before) {
+ free_slots_.pop_back();
+ return s;
+ }
+ if (free_slots_.size() * 4 < n_slots_) {
+ break;
+ }
+ free_slots_.pop_back();
+ }
+ int64_t best = -1;
+ uint64_t best_clock = UINT64_MAX;
+ for (size_t i = 0; i < n_slots_; i++) {
+ const slot_state & st = slots_[i];
+ if (st.pin > 0 || st.busy || st.last_use > stale_before) {
+ continue;
+ }
+ if (st.last_use < best_clock) {
+ best_clock = st.last_use;
+ best = (int64_t) i;
+ }
+ }
+ return (int) best;
+}
+
+int pager::acquire(int il, int ie, bool blocking) {
+ if (!ready_) {
+ return -1;
+ }
+ std::unique_lock<std::mutex> lk(mu_);
+ int32_t s = find_slot_locked(il, ie);
+ if (s >= 0 && slots_[s].busy) {
+ // 這個槽位**正在被填入**：另一條執行緒剛 admit 它、資料還沒讀完。
+ //
+ // 這裡一定要等，不能當成命中直接用。否則讀到的是「寫到一半的權重」，
+ // 而且那個錯誤是**非決定性**的 —— 每次執行結果可能不同（實測在
+ // prefill 時會讓同一個提示詞跑出不同的文字）。
+ // busy 的槽位在 admit/victim 那邊本來就會被跳過，所以這裡是唯一
+ // 還能拿到它的路徑，也因此必須在這裡等。
+ const uint64_t t0 = now_us();
+ while (slots_[s].busy) {
+ lk.unlock();
+ std::this_thread::sleep_for(std::chrono::microseconds(200));
+ lk.lock();
+ if (now_us() - t0 > (uint64_t) env_size("SDQ_BUSY_WAIT_MS", 30000) * 1000) {
+ fprintf(stderr,
+ "[sdq] 致命：等待槽位 il=%d ie=%d 讀入超過 %u ms（busy 一直沒解除）\n",
+ il, ie, (unsigned) env_size("SDQ_BUSY_WAIT_MS", 30000));
+ std::fflush(stderr);
+ abort();
+ }
+ }
+ // 讀取可能失敗（那時槽位已被清空）→ 重新走一次正常的 miss 路徑
+ if (find_slot_locked(il, ie) != s) {
+ s = -1;
+ }
+ }
+ if (s >= 0) {
+ slots_[s].last_use = ++clock_;
+ slots_[s].pin++;
+ touch_freq_locked(slots_[(size_t) s]);
+ // 剖面必須統計「所有」請求（命中也要算）。舊版只在 miss 路徑計數，
+ // 熱門 expert 因為大多命中，會被剖面系統性低估。
+ count_request(il, ie);
+ stats_.expert_requests.fetch_add(1, std::memory_order_relaxed);
+ stats_.expert_hits.fetch_add(1, std::memory_order_relaxed);
+ if (lstats_ && il >= 0 && il < (int) info_.n_layer) {
+ lstats_[il].req.fetch_add(1, std::memory_order_relaxed);
+ lstats_[il].hit.fetch_add(1, std::memory_order_relaxed);
+ }
+ // A prefetch that already finished counts as "used" here. Without this
+ // the finished-before-request case was invisible and prefetch precision
+ // was understated (docs/10 §f).
+ if (slots_[s].from_prefetch) {
+ slots_[s].from_prefetch = false;
+ stats_.prefetch_used.fetch_add(1, std::memory_order_relaxed);
+ stats_.prefetch_bytes_used.fetch_add(slot_size_of(s), std::memory_order_relaxed);
+ }
+ return s;
+ }
+
+ stats_.expert_requests.fetch_add(1, std::memory_order_relaxed);
+ count_request(il, ie);
+ stats_.expert_misses.fetch_add(1, std::memory_order_relaxed);
+ if (lstats_ && il >= 0 && il < (int) info_.n_layer) {
+ lstats_[il].req.fetch_add(1, std::memory_order_relaxed);
+ lstats_[il].miss.fetch_add(1, std::memory_order_relaxed);
+ }
+
+ // 先找一個已經在背景讀入這頁的槽
+ for (size_t i = 0; i < n_slots_; i++) {
+ slot_state & st = slots_[i];
+ if (st.busy && st.layer == il && st.expert == ie) {
+ // 等背景執行緒讀完
+ lk.unlock();
+ const uint64_t t0 = now_us();
+ while (true) {
+ lk.lock();
+ if (!slots_[i].busy) {
+ break;
+ }
+ lk.unlock();
+ std::this_thread::sleep_for(std::chrono::microseconds(200));
+ }
+ s = (int32_t) i;
+ lk.unlock();
+ stats_.prefetch_wait_us.fetch_add(now_us() - t0, std::memory_order_relaxed);
+ std::unique_lock<std::mutex> lk2(mu_);
+ slots_[s].last_use = ++clock_;
+ slots_[s].pin++;
+ stats_.expert_hits.fetch_add(1, std::memory_order_relaxed);
+ stats_.prefetch_used.fetch_add(1, std::memory_order_relaxed);
+ return s;
+ }
+ }
+
+ if (!blocking) {
+ return -1;
+ }
+
+ // 選槽位時必須知道「這個 expert 需要哪一種大小」，否則大 expert 可能被塞進小槽位
+ // （寫出邊界 → 靜默算錯），或者小 expert 佔掉大槽位（浪費 6.8% 的 arena）。
+ s = victim_locked(il, ie);
+ if (s < 0) {
+ // 拿不到槽位：短暫等待別的執行緒解除 pin（每層開頭會 unpin_all）。
+ // 這裡「重試」而不是直接回 -1，因為回 -1 會讓呼叫端**跳過這個 expert**，
+ // 結果是靜默少算一條專家分支 —— 那是最嚴重的錯誤。
+ // 2000 次 × 250 µs = 最多等 0.5 秒；真的撐不到那麼久，代表 arena 被
+ // 預取的 busy 槽位佔住了（實測：跨層預取開啟時會發生）。
+ for (int attempt = 0; attempt < 2000 && s < 0; attempt++) {
+ lk.unlock();
+ std::this_thread::sleep_for(std::chrono::microseconds(250));
+ lk.lock();
+ s = victim_locked(il, ie);
+ }
+ if (s < 0) {
+ fprintf(stderr,
+ "[sdq] 嚴重：找不到能放 il=%d ie=%d（need_big=%d）的槽位，"
+ "大槽 %zu／小槽 %zu\n",
+ il, ie, (int) need_big_slot(il, ie), n_big_, n_small_);
+ // 這裡「回 -1」會讓呼叫端**跳過這個 expert**，結果是靜默少算一條專家分支
+ // —— 那是整個引擎最嚴重的錯誤型態（模型會安靜地給出錯誤答案，而且
+ // 每次執行還不一定一樣）。寧可整個行程死掉，也不要產出「看起來正常」
+ // 的錯誤結果。
+ fprintf(stderr,
+ "[sdq] 致命：無法取得 expert 槽位。為了不給出靜默錯誤的結果，"
+ "直接中止。請回報這個情況（多半是預取把槽位佔成 busy 太久）。\n");
+ std::fflush(stderr);
+ abort();
+ }
+ }
+ slot_state & st = slots_[s];
+ if (st.layer >= 0) {
+ page_.erase(((int64_t) st.layer << 20) | (int64_t) st.expert);
+ stats_.evictions.fetch_add(1, std::memory_order_relaxed);
+ used_.fetch_sub(1, std::memory_order_relaxed);
+ }
+ st.layer = il;
+ st.expert = ie;
+ st.last_use = ++clock_;
+ st.pin = 1;
+ st.is_hot = hot_lookup(il, ie);
+ // **先標 busy，再 publish 到 page_。**
+ // 順序反過來的話（先 publish 再標 busy 中間有一個窗口）另一條執行緒會把它
+ // 當成命中，然後在資料還沒讀完時就拿去做矩陣乘法 → 非決定性的錯誤輸出。
+ // acquire() 看到 busy 會等；admit/victim 本來就會跳過 busy 的槽位。
+ st.busy = true;
+ page_[((int64_t) il << 20) | (int64_t) ie] = s;
+
+ lk.unlock();
+
+ const uint64_t t0 = now_us();
+ const bool ok = read_expert_par(il, ie, slot_ptr(s));
+ const uint64_t dt = now_us() - t0;
+
+ std::unique_lock<std::mutex> lk2(mu_);
+ st.busy = false;
+ if (!ok) {
+ st.layer = -1;
+ st.expert = -1;
+ st.pin = 0;
+ page_.erase(((int64_t) il << 20) | (int64_t) ie);
+ return -1;
+ }
+ used_.fetch_add(1, std::memory_order_relaxed);
+ stats_.demand_wait_us.fetch_add(dt, std::memory_order_relaxed);
+ const layer_desc * L = layer(il);
+ if (L) {
+ const uint64_t nb = (uint64_t) L->experts[(size_t) ie].bytes;
+ stats_.ssb_read_bytes_demand.fetch_add(nb, std::memory_order_relaxed);
+ if (lstats_ && il >= 0 && il < (int) info_.n_layer) { // SDQ-0004
+ lstats_[il].demand_bytes.fetch_add(nb, std::memory_order_relaxed);
+ }
+ }
+ return s;
+}
+
+void pager::prefetch_expert(int il, int ie) {
+ if (!ready_ || !cfg_.prefetch) {
+ return;
+ }
+ {
+ std::unique_lock<std::mutex> lk(mu_);
+ if (find_slot_locked(il, ie) >= 0) {
+ stats_.prefetch_duplicate.fetch_add(1, std::memory_order_relaxed);
+ return;
+ }
+ // 預取只能占用「空槽」或「已經很久沒被用到的頁」。
+ // （實測：無差別淘汰會讓命中率從 90% 掉到 27%，而預取使用率是 0%，純粹有害。）
+ //
+ // **槽位一定要選對尺寸**：need_big_slot() 為真的 expert 拿到小槽位會寫出邊界。
+ const uint64_t stale_before = clock_ - (uint64_t) env_size("SDQ_PREFETCH_STALE_TOKENS", 2) * 320;
+ // 預取佇列也不能太深：SSD 讀不夠快時，排在後面的預取早就過期了
+ if (pf_q_.size() > (size_t) env_size("SDQ_PREFETCH_QUEUE", 64)) {
+ return;
+ }
+ // **必須留夠空的槽位給 demand。**
+ // 預取會把槽位標成 busy 並佔著它讀取；大槽位本來就少（工作集下限是
+ // n_big_layers × n_expert_used），如果預取把最後幾個空槽都吃掉，
+ // 下一層的 demand 就會「找不到能放 il=34 ie=114 的槽位」而中止。
+ // （這是 verify 新增的預取壓力測試 [2b] 抓到的實際情況。）
+ // 所以：對應尺寸的空槽少於 n_expert_used 個時，這一筆預取直接不做。
+ const size_t reserve = (size_t) info_.n_expert_used;
+ const size_t n_free_right = need_big_slot(il, ie) ? free_big_slots_.size() : free_slots_.size();
+ if (n_free_right < reserve) {
+ stats_.prefetch_duplicate.fetch_add(1, std::memory_order_relaxed);
+ return;
+ }
+ const int32_t s = victim_stale_for(il, ie, stale_before);
+ if (s < 0) {
+ return;
+ }
+ slot_state & st = slots_[s];
+ if (st.layer >= 0) {
+ page_.erase(((int64_t) st.layer << 20) | (int64_t) st.expert);
+ stats_.evictions.fetch_add(1, std::memory_order_relaxed);
+ used_.fetch_sub(1, std::memory_order_relaxed);
+ }
+ st.layer = il;
+ st.expert = ie;
+ st.busy = true;
+ st.pin = 0;
+ pf_q_.push_back({s, il, ie, now_us()});
+ stats_.prefetch_issued.fetch_add(1, std::memory_order_relaxed);
+ }
+ pf_cv_.notify_one();
+}
+
+void pager::pf_worker() {
+ std::unique_lock<std::mutex> lk(pf_mu_);
+ for (;;) {
+ pf_cv_.wait(lk, [this] { return pf_stop_ || !pf_q_.empty(); });
+ if (pf_stop_ && pf_q_.empty()) {
+ return;
+ }
+ const pf_item item = pf_q_.front();
+ pf_q_.pop_front();
+ lk.unlock();
+
+ // 這個槽可能已經被淘汰/重用 → 檢查 generation
+ {
+ std::unique_lock<std::mutex> g(mu_);
+ const slot_state & st = slots_[item.slot];
+ const bool still = st.busy && st.layer == item.il && st.expert == item.ie;
+ if (!still) {
+ lk.lock();
+ continue;
+ }
+ }
+
+ // 預取也走 I/O 段平行：單獨一條執行緒循序讀三段，是這條路徑唯一的瓶頸。
+ const bool ok = cfg_.io_parallel ? read_expert_par(item.il, item.ie, slot_ptr(item.slot))
+ : read_expert(item.il, item.ie, slot_ptr(item.slot));
+ const layer_desc * L = layer(item.il);
+ {
+ std::unique_lock<std::mutex> g(mu_);
+ slot_state & st = slots_[item.slot];
+ if (st.busy && st.layer == item.il && st.expert == item.ie) {
+ st.busy = false;
+ st.last_use = ++clock_;
+ if (ok) {
+ page_[((int64_t) item.il << 20) | (int64_t) item.ie] = item.slot;
+ used_.fetch_add(1, std::memory_order_relaxed);
+ stats_.prefetch_completed.fetch_add(1, std::memory_order_relaxed);
+ // SDQ-0004: tag the page so the demand path can credit it,
+ // and so unused prefetched bytes can be reported as waste.
+ st.from_prefetch = true;
+ if (lstats_ && item.il >= 0 && item.il < (int) info_.n_layer) {
+ lstats_[item.il].prefetch_bytes.fetch_add(
+ (uint64_t) slot_size_of(item.slot), std::memory_order_relaxed);
+ }
+ if (L) {
+ stats_.ssb_read_bytes_prefetch.fetch_add(
+ (uint64_t) L->experts[(size_t) item.ie].bytes, std::memory_order_relaxed);
+ }
+ } else {
+ st.layer = -1;
+ st.expert = -1;
+ }
+ }
+ }
+ lk.lock();
+ }
+}
+
+void pager::release(int s) {
+ if (s < 0 || (size_t) s >= n_slots_) {
+ return;
+ }
+ std::lock_guard<std::mutex> lk(mu_);
+ slot_state & st = slots_[(size_t) s];
+ if (st.pin > 0) {
+ st.pin--;
+ }
+}
+
+void pager::unpin_all() {
+ std::unique_lock<std::mutex> lk(mu_);
+ for (auto & st : slots_) {
+ st.pin = 0;
+ }
+}
+
+void pager::note_used(int il, int ie) {
+ if (il < 0 || ie < 0) {
+ return;
+ }
+ std::lock_guard<std::mutex> lk(hist_mu_);
+ if (il >= (int) cur_used_.size()) {
+ return;
+ }
+ auto & v = cur_used_[(size_t) il];
+ if (std::find(v.begin(), v.end(), ie) == v.end()) {
+ v.push_back(ie);
+ }
+}
+
+// 「下一層」預取：只在 layer L 的 op 開頭呼叫一次。
+//
+// 為什麼需要這個函式：end_token() 是**整個 token 的 40 層都算完之後**才被呼叫的
+// （sdq_cli.cpp 在 llama_decode 之後呼叫），所以它排的預取完全沒有和任何計算重疊
+// —— 量測上 prefetch 使用率只有 0~8 %，等於整條預取路徑是關的。
+// 在 layer L 開始時就替 layer L+1 排預取，才有一整層的時間（讀＋算，約 4 ms）
+// 可以讓 SSD 讀取躲在後面。
+//
+// 預測用的是 prev_used_[il+1]（上一個 token 在該層選的專家）—— MoE 的路由在
+// 鄰近 token 之間非常穩定，這是最便宜且不用訓練的預測器。
+void pager::prefetch_layer_ahead(int il) {
+ if (!ready_ || !cfg_.prefetch || il < 0 || il >= (int) layers_.size()) {
+ return;
+ }
+ int32_t ids[64];
+ const int k = (int) info_.n_expert_used;
+ if (k <= 0 || k > 64) {
+ return;
+ }
+ predict(il, ids, k);
+ for (int i = 0; i < k; i++) {
+ prefetch_expert(il, ids[i]);
+ }
+}
+
+// SDQ-0005: called once per llama_decode() from any front end.
+//
+// sdq-chat drives this itself from main(); llama-server does not, which left
+// tokens at 0, ssd_bytes_per_token undefined and SDQ_STATS_FILE empty. This is
+// the single place that closes a decode step, so hook it here.
+//
+// n_tokens > 1 means prefill: a whole prompt was just paged in. Counting that as
+// one "token" would make ssd_bytes_per_token meaningless, and feeding it to the
+// Markov table would poison the transition counts with a whole-batch transition.
+// So prefill advances the byte accounting by n_tokens but leaves the routing
+// history alone -- it is not a single-token transition.
+void pager::note_decode(int n_tokens) {
+ if (!ready_ || n_tokens <= 0) {
+ return;
+ }
+ stats_.tokens.fetch_add((uint64_t) n_tokens, std::memory_order_relaxed);
+ if (n_tokens == 1) {
+ end_token();
+ }
+ dump_stats("decode");
+}
+
+void pager::end_token() {
+ if (!ready_) {
+ return;
+ }
+ // 用「上個 token 的選擇」更新一階 Markov 轉移表
+ {
+ std::lock_guard<std::mutex> lk(hist_mu_);
+ for (size_t il = 0; il < cur_used_.size(); il++) {
+ auto & cur = cur_used_[il];
+ if (cur.empty()) {
+ continue;
+ }
+ auto & layer_hist = hist_[il];
+ // 以每層最近一次（或平均）的前一次選擇為前驅
+ if (!prev_used_.empty() && il < prev_used_.size() && !prev_used_[il].empty()) {
+ for (int32_t from : prev_used_[il]) {
+ trans & t = layer_hist[(size_t) from];
+ for (int32_t to : cur) {
+ int found = -1;
+ for (int32_t i = 0; i < t.n; i++) {
+ if (t.to[i] == to) {
+ t.cnt[i]++;
+ found = i;
+ break;
+ }
+ }
+ if (found < 0 && t.n < 64) {
+ t.to[t.n] = to;
+ t.cnt[t.n] = 1;
+ t.n++;
+ }
+ }
+ }
+ }
+ prev_used_[il] = cur;
+ cur.clear();
+ }
+ }
+
+ // 依預測排預取
+ if (!cfg_.prefetch) {
+ return;
+ }
+ int32_t ids[64];
+ for (size_t il = 0; il < layers_.size(); il++) {
+ const int k = (int) info_.n_expert_used;
+ predict((int) il, ids, k);
+ for (int i = 0; i < k; i++) {
+ prefetch_expert((int) il, ids[i]);
+ }
+ }
+}
+
+void pager::predict(int il, int32_t * out, int k) {
+ if (il < 0 || il >= (int) prev_used_.size()) {
+ for (int i = 0; i < k; i++) {
+ out[i] = i;
+ }
+ return;
+ }
+ // 以最近一次選擇為前驅；沒有歷史就直接沿用上一次的選擇
+ const std::vector<int32_t> & prev = prev_used_[(size_t) il];
+ std::vector<std::pair<uint32_t, int32_t>> scored;
+ if (prev.empty()) {
+ for (int i = 0; i < k; i++) {
+ out[i] = i;
+ }
+ return;
+ }
+ for (int32_t from : prev) {
+ const trans & t = hist_[(size_t) il][(size_t) from];
+ for (int32_t i = 0; i < t.n; i++) {
+ scored.push_back({t.cnt[i], t.to[i]});
+ }
+ }
+ if (scored.empty()) {
+ for (size_t i = 0; i < prev.size() && i < (size_t) k; i++) {
+ out[i] = prev[i];
+ }
+ for (int i = (int) prev.size(); i < k; i++) {
+ out[i] = i;
+ }
+ return;
+ }
+ std::sort(scored.begin(), scored.end(), [](auto & a, auto & b) {
+ return a.first > b.first;
+ });
+ for (int i = 0; i < k && i < (int) scored.size(); i++) {
+ out[i] = scored[(size_t) i].second;
+ }
+}
+
+double pager::hit_rate() const {
+ const uint64_t req = stats_.expert_requests.load();
+ const uint64_t hit = stats_.expert_hits.load();
+ return req ? (double) hit / (double) req : 0.0;
+}
+
+void pager::shutdown() {
+ if (!ready_) {
+ return;
+ }
+ if (!profile_out_.empty()) {
+ std::lock_guard<std::mutex> lk(mu_);
+ dump_profile(profile_out_);
+ }
+ if (pf_.joinable()) {
+ {
+ std::unique_lock<std::mutex> lk(pf_mu_);
+ pf_stop_ = true;
+ }
+ pf_cv_.notify_all();
+ pf_.join();
+ }
+ io_pool_stop();
+ for (auto & fd : shard_fd_) {
+ if (fd >= 0) {
+ close(fd);
+ }
+ }
+ for (auto & fd : shard_fd_direct_) {
+ if (fd >= 0) {
+ close(fd);
+ }
+ }
+ shard_fd_.clear();
+ shard_fd_direct_.clear();
+ ready_ = false;
+}
+
+pager::~pager() {
+ shutdown();
+ if (arena_) {
+ // 注意：arena 是 posix_memalign 配置的，必須用 free() 釋放，
+ // 用 munmap() 會破壞 glibc 的 heap（之後任何 free 都會報 double free）。
+ free(arena_);
+ arena_ = nullptr;
+ }
+}
+
+static double pct(uint64_t a, uint64_t b) {
+ return b ? 100.0 * (double) a / (double) b : 0.0;
+}
+
+void pager::print_summary() const {
+ if (!ready_) {
+ return;
+ }
+ // bailingmoe3 的第 0 層是 dense FFN（沒有 expert），要抓第一個真的有分頁的層
+ const layer_desc * L0 = nullptr;
+ int first_moe = -1;
+ for (size_t i = 0; i < layers_.size(); i++) {
+ if (!layers_[i].experts.empty()) {
+ L0 = &layers_[i];
+ first_moe = (int) i;
+ break;
+ }
+ }
+ if (L0 && !L0->experts.empty()) {
+ const expert_desc & e = L0->experts[0];
+ fprintf(stderr, "[sdq] page-map: n_layer=%lld n_expert=%lld n_expert_used=%lld n_embd=%lld n_ff_exp=%lld"
+ " | L%dE0 gate=%llu/%llu up=%llu/%llu down=%llu/%llu (off/size) ne0=%lld ne1=%lld d_ne0=%lld d_ne1=%lld\n",
+ (long long) info_.n_layer, (long long) info_.n_expert, (long long) info_.n_expert_used,
+ (long long) info_.n_embd, (long long) info_.n_ff_exp, first_moe,
+ (unsigned long long) e.seg[part::GATE].off, (unsigned long long) e.seg[part::GATE].size,
+ (unsigned long long) e.seg[part::UP].off, (unsigned long long) e.seg[part::UP].size,
+ (unsigned long long) e.seg[part::DOWN].off, (unsigned long long) e.seg[part::DOWN].size,
+ (long long) e.ne0, (long long) e.ne1, (long long) e.down_ne0, (long long) e.down_ne1);
+ }
+ fprintf(stderr,
+ "[sdq] expert pager ready: %s\n"
+ "[sdq] 熱門 expert 釘選：%d 個%s\n"
+ "[sdq] slot=%.2f MiB (%zu B) slots=%zu arena=%.2f GiB (RAM 實測 RSS=%.2f GiB)\n"
+ "[sdq] evict_policy=%s stale_window=%llu\n",
+ cfg_.model_path.c_str(),
+ n_hot_, n_hot_ > 0 ? "（永不淘汰）" : "（無剖面，用動態 LRU/LFU）",
+ slot_bytes_ / 1048576.0, slot_bytes_, n_slots_, arena_bytes_ / 1073741824.0,
+ process_rss_bytes() / 1073741824.0,
+ cfg_.evict_policy.c_str(), (unsigned long long) cfg_.stale_window);
+ if (n_small_ == 0) {
+ fprintf(stderr, "[sdq] 槽位：單一尺寸 %.2f MiB × %zu\n",
+ slot_bytes_ / 1048576.0, n_big_);
+ } else {
+ // 大小兩種槽位：把配置結果講清楚，否則只看「slots=N」會誤以為每個槽一樣大。
+ fprintf(stderr,
+ "[sdq] 槽位分兩種：小 %.2f MiB × %zu + 大 %.2f MiB × %zu（平均 %.2f MiB）\n",
+ small_bytes_ / 1048576.0, n_small_,
+ big_bytes_ / 1048576.0, n_big_,
+ (double) arena_bytes_ / (double) (n_slots_ ? n_slots_ : 1) / 1048576.0);
+ }
+ fprintf(stderr,
+ "[sdq] prefetch=%s fadvise_dontneed=%d\n",
+ cfg_.prefetch ? "on" : "off", (int) cfg_.fadvise_dontneed);
+}
+
+void pager::dump_stats(const char * phase) {
+ if (cfg_.stats_path.empty()) {
+ return;
+ }
+ const uint64_t tokens = stats_.tokens.load();
+ const uint64_t req = stats_.expert_requests.load();
+ const uint64_t hit = stats_.expert_hits.load();
+ const uint64_t issued = stats_.prefetch_issued.load();
+ const uint64_t used = stats_.prefetch_used.load();
+ const uint64_t b_dem = stats_.ssb_read_bytes_demand.load();
+ const uint64_t b_pre = stats_.ssb_read_bytes_prefetch.load();
+ // SDQ-0004: prefetched bytes that were consumed, so the rest is real waste.
+ const uint64_t b_pre_used = stats_.prefetch_bytes_used.load();
+
+ std::ostringstream o;
+ o << "{\"phase\":\"" << phase << "\""
+ << ",\"tokens\":" << tokens
+ << ",\"expert_requests\":" << req
+ << ",\"expert_hits\":" << hit
+ << ",\"expert_misses\":" << (req - hit)
+ << ",\"hit_rate\":" << pct(hit, req)
+ << ",\"ssb_demand_bytes\":" << b_dem
+ << ",\"ssb_prefetch_bytes\":" << b_pre
+ << ",\"ssb_total_bytes\":" << (b_dem + b_pre)
+ << ",\"ssd_bytes_per_token\":" << (tokens ? (double) (b_dem + b_pre) / (double) tokens : 0.0)
+ << ",\"prefetch_issued\":" << issued
+ << ",\"prefetch_used\":" << used
+ << ",\"prefetch_used_rate\":" << pct(used, issued)
+ // SDQ-0004: prefetch waste. b_pre counts every byte the background reader
+ // pulled in; b_pre_used counts the ones a demand request actually hit.
+ // The difference is what made docs/07's negative result worse than "slower".
+ << ",\"prefetch_bytes_used\":" << b_pre_used
+ << ",\"prefetch_bytes_wasted\":" << (b_pre > b_pre_used ? b_pre - b_pre_used : 0)
+ << ",\"prefetch_waste_rate\":" << pct(b_pre > b_pre_used ? b_pre - b_pre_used : 0, b_pre)
+ << ",\"prefetch_duplicate\":" << stats_.prefetch_duplicate.load()
+ << ",\"evictions\":" << stats_.evictions.load()
+ << ",\"demand_wait_us\":" << stats_.demand_wait_us.load()
+ << ",\"prefetch_wait_us\":" << stats_.prefetch_wait_us.load()
+ << ",\"arena_bytes\":" << arena_bytes_
+ // 大小兩種槽位之後，「用掉幾個槽位」不能直接乘 slot_bytes_（那是大的那種）。
+ // dump_stats 只在階段轉換時呼叫一次，遍歷兩千幾個槽位的成本可以忽略。
+ << ",\"arena_used_bytes\":" << used_bytes_snapshot()
+ << ",\"rss_bytes\":" << process_rss_bytes()
+ << ",\"n_slots\":" << n_slots_
+ << ",\"n_small_slots\":" << n_small_
+ << ",\"n_big_slots\":" << n_big_
+ << ",\"small_slot_bytes\":" << small_bytes_
+ << ",\"big_slot_bytes\":" << big_bytes_
+ << ",\"slot_bytes\":" << slot_bytes_
+ << ",\"n_moe_layers\":" << n_moe_layers_
+ << ",\"expert_bytes_total\":" << expert_bytes_total_
+ << ",\"arch\":\"" << info_.arch << "\""
+ << ",\"name\":\"" << info_.name << "\""
+ // SDQ-0004: per-layer breakdown. docs/10 needs this to tell "the device is
+ // slow" apart from "one layer keeps evicting everyone else's experts".
+ << ",\"per_layer\":[";
+ for (int64_t il = 0; lstats_ && il < (int64_t) info_.n_layer; il++) {
+ const uint64_t lreq = lstats_[il].req.load();
+ const uint64_t lhit = lstats_[il].hit.load();
+ o << (il ? "," : "")
+ << "{\"layer\":" << il
+ << ",\"req\":" << lreq
+ << ",\"hit\":" << lhit
+ << ",\"miss\":" << (lreq - lhit)
+ << ",\"hit_rate\":" << pct(lhit, lreq)
+ << ",\"demand_bytes\":" << lstats_[il].demand_bytes.load()
+ << ",\"prefetch_bytes\":" << lstats_[il].prefetch_bytes.load()
+ << "}";
+ }
+ o << "]" << "}";
+ std::ofstream f(cfg_.stats_path, std::ios::app);
+ if (f) {
+ f << o.str() << "\n";
+ }
+}
+
+} // namespace sdq
+namespace sdq {
+// 由 CLI 呼叫：依環境變數/參數初始化 pager
+bool sdq_pager_active_for_layer(int il) {
+ pager & P = pager::instance();
+ if (!P.ready() || P.layer(il) == nullptr) {
+ return false;
+ }
+ // 除錯用：SDQ_MOE_LO / SDQ_MOE_HI 可以只讓部分層走分頁路徑，
+ // 用來二分搜尋第一個發生數值差異的層。
+ static const int lo = [] { const char * e = getenv("SDQ_MOE_LO"); return e ? atoi(e) : 0; }();
+ static const int hi = [] { const char * e = getenv("SDQ_MOE_HI"); return e ? atoi(e) : 1 << 30; }();
+ if (il < lo || il >= hi) {
+ return false;
+ }
+ // SDQ_MOE_SKIP=3,7：這些層即使在 [lo,hi) 內也走原生路徑。
+ // 搭配 SDQ_BUFT_KEEP（讓該層的 expert 留在 RAM）就能做同一層的 A/B 比對。
+ static const std::set<int> skip = [] {
+ std::set<int> s;
+ const char * e = getenv("SDQ_MOE_SKIP");
+ if (e) {
+ for (const char * p = e; *p; ) {
+ const char * q = strchr(p, ',');
+ const size_t n = q ? (size_t) (q - p) : strlen(p);
+ if (n) {
+ s.insert(atoi(std::string(p, n).c_str()));
+ }
+ if (!q) {
+ break;
+ }
+ p = q + 1;
+ }
+ }
+ return s;
+ }();
+ return skip.find(il) == skip.end();
+}
+
+void sdq_on_decode(int n_tokens) {
+ auto & inst = pager::instance();
+ if (inst.ready()) {
+ inst.note_decode(n_tokens);
+ }
+}
+
+bool sdq_init(const std::string & model_path) {
+ config c;
+ c.model_path = model_path;
+ c.arena_bytes = env_size("SDQ_ARENA_MB", 0) << 20;
+ c.ram_budget_bytes = env_size("SDQ_RAM_BUDGET_MB", 0) << 20;
+ c.slot_bytes = env_size("SDQ_SLOT_KB", 0) << 10;
+ // 預設關閉。docs/07 的實測是「開 5.68 / 關 6.20 tok/s」的負結果，而且
+ // 預取中的槽位會讓 acquire() 靜默跳過 expert，曾造成輸出不一致（stage S7）。
+ // 這裡原本是 true，等於讓那個負結果從來沒有真正生效。
+ c.prefetch = env_flag("SDQ_PREFETCH", false);
+ // ---- OPT-1：熱門 expert 剖面（Strata 式）----
+ if (const char * pi = getenv("SDQ_PROFILE_IN")) { c.profile_in = pi; }
+ if (const char * po = getenv("SDQ_PROFILE_OUT")) { c.profile_out = po; }
+ c.hot_pin = (int) env_size("SDQ_HOT_PIN", 0);
+ c.admit_min_hits = (int) env_size("SDQ_ADMIT_MIN_HITS", 1);
+ if (const char * ep = getenv("SDQ_EVICT_POLICY")) { c.evict_policy = ep; }
+ c.stale_window = (uint64_t) env_size("SDQ_EVICT_STALE_TOKENS", 2) * 320;
+ c.fadvise_dontneed = env_flag("SDQ_FADV_DONTNEED", true);
+ c.io_parallel = env_flag("SDQ_IO_PARALLEL", true);
+ // I/O 讀取執行緒數。sdq-chat 會在載入模型「之前」自己 setenv 這個值
+ // （3 × 計算執行緒，上限 48，docs/05 量到序列 3.79 → 24 條 6.68 tok/s）。
+ // 但 cfg_.n_threads 從來沒有被賦值，所以任何沒有自行設定的前端
+ // （llama-server 就是）會拿到 0，然後在 io_pool_start() 被夾到下限 4 條 ——
+ // 等於一直跑在「序列讀」那一端。改成在 pager 自己算，量測與 sdq-chat 一致。
+ if (!getenv("SDQ_IO_THREADS")) {
+ const int hw = (int) std::max(1u, std::thread::hardware_concurrency());
+ const int compute = hw < 8 ? hw : 8; // 與 sdq-chat 同一個上限 8
+ c.io_threads = compute * 3;
+ } else {
+ c.io_threads = (int) env_size("SDQ_IO_THREADS", 0);
+ }
+ c.io_split = (int) env_size("SDQ_IO_SPLIT", 1);
+ c.verbose = env_flag("SDQ_VERBOSE", false);
+ const char * sp = getenv("SDQ_STATS_FILE");
+ if (sp) {
+ c.stats_path = sp;
+ }
+ sdq_set_config(c);
+ try {
+ pager::instance().init(c, nullptr);
+ } catch (const std::exception & e) {
+ fprintf(stderr, "[sdq] pager 初始化失敗：%s\n", e.what());
+ return false;
+ }
+ // 32/64 GiB 情境：arena 大到能裝下全部 expert 時先全部讀進來。
+ // 沒有 miss 之後 row-parallel 才能穩定把 decode 的 16 條執行緒吃滿。
+ if (env_flag("SDQ_PRELOAD", true)) {
+ const bool ok = pager::instance().preload_all(16);
+ if (ok) {
+ fprintf(stderr, "[sdq] 已預載全部 expert（%zu slots）→ 之後 row-parallel 可用\n",
+ pager::instance().n_slots());
+ }
+ }
+ // 長駐前端（llama-server）不會呼叫 sdq_shutdown()；用 atexit 在收尾時補上摘要。
+ // 註冊時機必須在 pager::instance() 建構之後：static 物件的析構與 atexit 依
+ // 「後註冊者先執行」，反過來就會在 pager 已經析構後才跑這段（讀到已釋放的物件）。
+ static std::once_flag once_exit;
+ std::call_once(once_exit, [] { std::atexit([] { sdq_shutdown(); }); });
+ return pager::instance().ready();
+}
+
+
+// ---------------------------------------------------------------- 除錯：dump 任意 F32 張量
+// 目的：讓「原生 mul_mat_id 路徑」與「SDQ 分頁路徑」輸出同一層的 MoE 結果，
+// 兩邊各自 dump 一次就能直接 A/B 對照，判定到底是哪一邊算錯。
+namespace {
+struct dump_arg { std::string path; };
+
+void dump_op(ggml_tensor * dst, int ith, int /*nth*/, void * userdata) {
+ if (ith != 0) {
+ return;
+ }
+ const ggml_tensor * src = dst->src[0];
+ dump_arg * a = (dump_arg *) userdata;
+ // 重要：dump 節點必須是「透明」的 —— 把 src 原樣複製到 dst，後續節點才會讀到
+ // 正確資料。舊版只寫檔、沒有複製，dst 是未初始化記憶體，導致下游殘差加法讀到垃圾。
+ if (src->type == GGML_TYPE_F32 && ggml_is_contiguous(src) && ggml_is_contiguous(dst)) {
+ memcpy(dst->data, src->data, ggml_nbytes(src));
+ }
+ FILE * f = fopen(a->path.c_str(), "wb");
+ if (!f) {
+ fprintf(stderr, "[sdq-dump] 無法開啟 %s\n", a->path.c_str());
+ return;
+ }
+ const int64_t meta[6] = { src->ne[0], src->ne[1], src->ne[2], src->ne[3],
+ (int64_t) ggml_nbytes(src), 0 };
+ fwrite(meta, sizeof(int64_t), 6, f);
+ // 只支援 F32 contiguous；其他情況就明確報錯，不要假設
+ if (src->type != GGML_TYPE_F32 || !ggml_is_contiguous(src)) {
+ fprintf(stderr, "[sdq-dump] %s 不是 contiguous F32（type=%d）\n",
+ a->path.c_str(), (int) src->type);
+ } else {
+ fwrite(src->data, sizeof(float), (size_t) (ggml_nelements(src)), f);
+ }
+ fclose(f);
+ fprintf(stderr, "[sdq-dump] %s (%lld 個元素)\n", a->path.c_str(),
+ (long long) ggml_nelements(src));
+}
+} // namespace
+
+
+ggml_tensor * sdq_maybe_dump_tensor(ggml_context * ctx, ggml_tensor * t, int il) {
+ const char * prefix = getenv("SDQ_MOE_OUT_DUMP");
+ if (!prefix || !t) {
+ return t;
+ }
+ int layer = getenv("SDQ_MOE_OUT_DUMP_LAYER") ? atoi(getenv("SDQ_MOE_OUT_DUMP_LAYER")) : -1;
+ if (layer >= 0 && il != layer) {
+ return t;
+ }
+ dump_arg * a = new dump_arg{ std::string(prefix) + "." + std::to_string(il) };
+ return ggml_custom_4d(ctx, GGML_TYPE_F32, t->ne[0], t->ne[1], t->ne[2], t->ne[3],
+ &t, 1, dump_op, 1, a);
+}
+
+void sdq_shutdown() {
+ pager & P = pager::instance();
+ sdq_print_time_profile();
+ sdq_print_op_counts();
+ P.dump_stats("shutdown");
+ P.print_summary();
+ P.shutdown();
+}
+} // namespace sdq
diff --git a/src/sdq_pager.h b/src/sdq_pager.h
new file mode 100644
index 0000000..b109f33
--- /dev/null
+++ b/src/sdq_pager.h
@@ -0,0 +1,510 @@
+// SDQ expert pager：SSD → RAM → CPU 的 MoE 專家分層快取
+//
+// 設計要點
+// * 分頁單位 = 一個 expert（layer, expert id），大小 ~1.8 MB（三個矩陣：gate/up/down）
+// * RAM 端是一個固定槽位的 arena（slab），每槽 2 MiB，O(1) 配置、不碎片化
+// * 淘汰策略：LRU + 本 token 使用中（pinned）的保護
+// * 預取：依「上一個 token 的專家選擇」用一階 Markov 轉移預測下一個 token 的專家，
+// 交給背景執行緒 pread()，量測「預取命中率 / 預取浪費量」
+// * 指標：hit/miss、SSD 讀取位元組、demand 停滯時間、prefetch 領先時間、arena 使用量
+//
+// 不使用 swap：所有權重資料只存在 arena（實體 RAM）與檔案（SSD）兩處。
+
+#pragma once
+
+#include <ggml.h>
+#include <ggml-backend.h>
+#include "llama.h"
+
+#include <atomic>
+#include <chrono>
+#include <condition_variable>
+#include <cstddef>
+#include <cstdint>
+#include <deque>
+#include <memory>
+#include <mutex>
+#include <string>
+#include <thread>
+#include <unordered_map>
+#include <vector>
+
+namespace sdq {
+
+// ---------------------------------------------------------------- GGUF 讀取
+struct gguf_info {
+ int64_t n_layer = 0;
+ int64_t n_expert = 0;
+ int64_t n_expert_used = 0;
+ int64_t n_embd = 0;
+ int64_t n_ff_exp = 0;
+ int64_t n_ff_shexp = 0;
+ int64_t n_head = 0;
+ int64_t n_head_kv = 0;
+ int64_t n_embd_head_k = 0;
+ int64_t n_embd_head_v = 0;
+ int64_t context_length = 0;
+ int64_t full_attention_interval = 4;
+ // bailingmoe3 / Ling-3.0-flash
+ int64_t n_layer_dense_lead = 0;
+ int64_t n_expert_groups = 0;
+ int64_t n_expert_group_used = 0;
+ int64_t is_recr[512] = {0};
+ // nemotron_h_moe: 每層的 expert 數量與 top-k 可以不同（Puzzle / LatentMoE）
+ std::vector<int64_t> n_expert_layer;
+ std::vector<int64_t> n_expert_used_layer;
+ int64_t moe_latent_size = 0;
+ bool no_gate = false; // routed experts 只有 up/down（Nemotron LatentMoE）
+ bool relu_sqr = false; // 激活為 relu(x)^2
+ std::string arch;
+ std::string name;
+};
+
+// 一個 expert 槽位：gate/up/down 三段（有些模型把 gate+up 合成 gate_up）
+struct seg_desc {
+ uint64_t off = 0; // 在「所屬 shard 檔案內」的位元組偏移
+ uint64_t size = 0; // 位元組數
+ // sharded GGUF：同一層的三段不一定在同一個檔案裡（unsplit 後 100+ GB 的模型一定會被切），
+ // 所以每一段都要記得自己屬於哪一個 shard，pread 才知道要開哪個 fd。
+ int shard = 0;
+};
+
+enum part : int { GATE = 0, UP = 1, DOWN = 2, COUNT = 3 };
+
+struct expert_desc {
+ seg_desc seg[3];
+ int64_t slot_off[3] = {0, 0, 0}; // 三段在 arena 槽位內的相對偏移
+ int64_t ne0 = 0; // gate/up row 長度（= n_embd）
+ int64_t ne1 = 0; // gate/up row 數（= n_ff_exp）
+ int64_t down_ne0 = 0; // down row 長度（= n_ff_exp）
+ int64_t down_ne1 = 0; // down row 數（= n_embd）
+ ggml_type type[3] = {GGML_TYPE_COUNT, GGML_TYPE_COUNT, GGML_TYPE_COUNT};
+ int64_t bytes = 0; // 三段總大小
+ bool fused_gate_up = false;
+};
+
+// 一層的所有 expert 描述（40 層 × 256 個）
+struct layer_desc {
+ std::vector<expert_desc> experts;
+ int64_t bytes_per_expert = 0;
+ int64_t n_expert = 0; // 這一層的 expert 數（Nemotron 每層不同）
+ bool no_gate = false; // 只有 up/down，沒有 gate
+ // 這個層的 expert 比「大多數層」大（Q4_K/Q5_K vs Q6_K）→ 需要大槽位。
+ // 不然的話所有槽位都得按最大的那種配，白白浪費 6.8% 的 arena。
+ bool needs_big_slot = false;
+};
+
+// ---------------------------------------------------------------- 指標
+struct stats {
+ std::atomic<uint64_t> expert_requests{0}; // 每個 (layer, token) 需要幾個 expert
+ std::atomic<uint64_t> expert_hits{0};
+ std::atomic<uint64_t> expert_misses{0};
+ std::atomic<uint64_t> ssb_read_bytes_demand{0};
+ std::atomic<uint64_t> ssb_read_bytes_prefetch{0};
+ std::atomic<uint64_t> prefetch_issued{0};
+ std::atomic<uint64_t> prefetch_completed{0};
+ std::atomic<uint64_t> prefetch_used{0}; // 預取後真的被用到
+ std::atomic<uint64_t> prefetch_duplicate{0};// 預取時已經在 RAM 裡
+ std::atomic<uint64_t> evictions{0};
+ std::atomic<uint64_t> demand_wait_us{0}; // 因為缺頁而停下來等 SSD 的時間
+ std::atomic<uint64_t> prefetch_wait_us{0};
+ std::atomic<uint64_t> prefetch_bytes_used{0}; // prefetched bytes a demand request actually hit
+ std::atomic<uint64_t> tokens{0};
+ std::atomic<uint64_t> moe_calls{0};
+ std::atomic<uint64_t> flops{0};
+};
+
+// ---- SDQ-0004: per-layer cache counters (Phase 0: is any layer starved?)
+struct layer_stat {
+ std::atomic<uint64_t> req{0};
+ std::atomic<uint64_t> hit{0};
+ std::atomic<uint64_t> miss{0};
+ std::atomic<uint64_t> demand_bytes{0};
+ std::atomic<uint64_t> prefetch_bytes{0};
+};
+
+
+// ---------------------------------------------------------------- pager
+struct config {
+ size_t arena_bytes = 0; // 0 = 自動（依 RAM 預算）
+ size_t slot_bytes = 0; // 0 = 自動（依模型推導）
+ int n_threads = 0;
+ bool prefetch = true;
+ int prefetch_ahead_layers = 0; // >0 時限制同時在飛的預取數量
+ bool fadvise_dontneed = true; // 讀完丟掉 kernel page cache，讓 RAM 帳目誠實
+ bool verbose = false;
+ uint64_t ram_budget_bytes = 0; // 0 = 自動（實體 RAM 的一小部分）
+ std::string model_path;
+ std::string stats_path;
+ int slot_bias = 0;
+
+ // ---- I/O 執行緒池 ----
+ // 為什麼需要：一個 expert 在檔案上是 gate / up / down 三段「不相鄰」的資料，
+ // 循序讀要 3 次 pread，而單一 pread 在本機量到約 7 ms（1.9 MB / 275 MB/s）。
+ // 也就是說「讀一個 expert」的延遲完全由這三段循序決定，執行緒再多也沒用。
+ // 把三段交給三條執行緒同時讀，單一 expert 的延遵掉到約 1/2~1/3。
+ // （實測：8 個 expert 全部 miss 的一層，序列讀 12.6 ms → 段平行 7.6 ms。）
+ bool io_parallel = true; // SDQ_IO_PARALLEL=0 可關（回到序列讀）
+ int io_threads = 0; // 0 = 自動（3 × n_threads，上限 48）
+ int io_split = 1; // SDQ_IO_SPLIT：每一段再切成幾塊並行讀（實測 1 最好）
+
+ // ---- Strata 式「熱門 expert 剖面」（OPT-1）----
+ // 離線用代表性文字跑一次，把每個 (layer, expert) 的請求次數記成剖面；
+ // 執行時把排行榜前 hot_pin 個釘成「熱頁」：熱頁可以被讀取，但**永遠不被淘汰**，
+ // 淘汰時一律先從冷頁挑。這解決的是「prefill 的一次性專家把熱區洗掉」。
+ std::string profile_in; // 讀取剖面（SDQ_PROFILE_IN）
+ std::string profile_out; // 結束時輸出剖面（SDQ_PROFILE_OUT）
+ int hot_pin = 0; // 釘選前幾名（SDQ_HOT_PIN）
+ int admit_min_hits = 1; // 請求次數低於此值就不收進快取（SDQ_ADMIT_MIN_HITS）
+
+ // ---- 淘汰策略（SDQ_EVICT_POLICY）----
+ // lru：挑「最久未用」的頁（現行行為）。
+ // least_stale：只淘汰閒置時間 >= stale_window 的頁；若沒有這種頁，
+ // 才退回全域 LRU。目的是保護剛用過的工作集，避免 thrash。
+ std::string evict_policy = "lru";
+ uint64_t stale_window = 640; // clock 單位（每次 expert 請求 +1）
+};
+
+class pager {
+public:
+ static pager & instance();
+
+ // 從 GGUF 檔頭建立 expert 頁面表，並配置 RAM arena
+ void init(const config & cfg, ggml_context * ctx);
+
+ bool ready() const { return ready_; }
+ size_t arena_bytes() const { return arena_bytes_; }
+ size_t slot_bytes() const { return slot_bytes_; }
+ size_t n_slots() const { return n_slots_; }
+ size_t small_bytes() const { return small_bytes_; }
+ size_t big_bytes() const { return big_bytes_; }
+ size_t n_small_slots() const { return n_small_; }
+ size_t n_big_slots() const { return n_big_; }
+ int n_moe_layers() const { return n_moe_layers_; }
+ uint64_t expert_bytes_total() const { return expert_bytes_total_; }
+ const gguf_info & info() const { return info_; }
+
+ // 取得 layer 描述（idx 越界或沒資料回 nullptr）
+ const layer_desc * layer(int il) const;
+
+ // 一次把「連續編號的一串 expert」讀進 RAM：同一層的三個矩陣（gate/up/down）本來
+ // 就各自連續存放，所以 k 個連續 expert 只需要 3 次大讀取（而不是 3k 次）。
+ // 對高延遲的儲存裝置（每個 request 都貴）這是最大的單項加速。
+ // experts 必須已依編號遞增排序；slots[i] 是 experts[i] 對應的槽位（由呼叫端配置）
+ // 回傳實際讀取的位元組數。
+ uint64_t read_expert_span(int il, const int32_t * experts, int n, const int32_t * slots);
+
+ // 只配置槽位（不讀資料）：回傳槽位索引，-1 表示沒有空槽
+ int reserve_slot(int il, int ie);
+
+ // 查這個 expert 是否已經在 RAM：命中則 pin 並回傳槽位，未命中回傳 -1（不配置）
+ int probe(int il, int ie);
+
+ // 只查詢、不改動 pin / LRU / 統計：回傳 true 表示資料已在 RAM 且可讀。
+ // row-parallel 模式用它判斷「整層都已常駐」→ 才值得把 16 條執行緒
+ // 全部拿來切同一個 expert 的輸出列（沒有 I/O 可重疊時才不會虧）。
+ bool resident(int il, int ie);
+
+ // 若 arena 大到能裝下「全部」expert，就先把 15.9 GiB 一次讀進來。
+ // 這是 32/64 GiB 情境（原生 llama.cpp 也是全部常駐）的公平對照：
+ // 沒有 miss 之後 row-parallel 才能穩定吃滿 16 條執行緒。
+ // 回傳 false = arena 不夠、沒有做（小記憶體情境不受影響）。
+ bool preload_all(int n_threads);
+
+ // 清空 hit/miss/bytes 等統計（preload 後呼叫，避免把載入階段算進命中率）。
+ void reset_stats();
+
+ // 記錄一次合併讀取的位元組數（給指標用）
+ void note_span_read(uint64_t bytes);
+
+ // 確保 (layer, expert) 在 RAM，回傳槽位索引；-1 表示失敗
+ // pinned: 本 token 使用中，不可被淘汰；blocking: 缺頁時是否同步讀 SSD
+ int acquire(int il, int ie, bool blocking);
+
+ // 背景預取：把 page 讀進 arena（非同步），依歷史預測下一個 token 會用到的 expert
+ void prefetch_expert(int il, int ie);
+ // 「下一層」預取：在 layer L 的 op 開頭呼叫，替 layer L+1 排預取。
+ // 與 end_token() 的差別在時機 —— 後者在一個 token 的 40 層全算完才跑，
+ // 那時排的預取沒有任何計算可以重疊，實測使用率只有 0~8 %。
+ void prefetch_layer_ahead(int il);
+
+ // 每層計算開始時呼叫：解除上一層的 pin（本層用到的槽位才受保護）
+ void unpin_all();
+ // 用完一個槽位後立刻放掉 pin（成對呼叫 acquire()）。
+ //
+ // 為什麼需要：原本 pin 是「整層釘住、下一層開頭才 unpin_all」。mode C 的
+ // expert_forward 回傳之後就再也不讀那個槽位（最後的歸約只讀各執行緒自己的
+ // partial 累加緩衝），所以繼續釘著只是白白把工作集撐大。
+ // 撐大的後果很具體：prefill 一個 ubatch 有 128 個 token，一層最多需要
+ // 256 個不同的 expert → 整層釘住 = 256 個槽位不能回收。大小兩種槽位之下，
+ // 大槽位只有約 n_big_layers × n_expert_used 個，於是**必定**在 prefill 時
+ // 「找不到能放 il=34 ie=… 的槽位」而中止。
+ void release(int s);
+ void note_used(int il, int ie);
+ void end_token();
+ // ---- SDQ-0005: one call per llama_decode() from any front end.
+ void note_decode(int n_tokens);
+
+ // 依歷史預測某層下一個 token 的 top-k expert（回傳信心 0..1）
+ void predict(int il, int32_t * out, int k);
+
+ void shutdown();
+ stats & get_stats() { return stats_; }
+
+ // 把當前指標寫成一行 JSON（append 到 stats_path）
+ void dump_stats(const char * phase);
+ void print_summary() const;
+
+ // arena 內的槽位指標
+ // 槽位索引空間是統一的：[0, n_small_) 是小槽位（連續），
+ // [n_small_, n_slots_) 是大槽位（接在小槽位區之後）。
+ uint8_t * slot_ptr(int s) {
+ return (size_t) s < n_small_
+ ? arena_ + (size_t) s * small_bytes_
+ : arena_ + small_bytes_ * n_small_ + (size_t)(s - (int) n_small_) * big_bytes_;
+ }
+ size_t slot_size_of(int s) const {
+ return (size_t) s < n_small_ ? small_bytes_ : big_bytes_;
+ }
+ bool need_big_slot(int il, int ie) const;
+
+ // 這個槽位算不算「大槽位」。
+ //
+ // 判斷依據是**槽位的實際大小**，不是索引 >= n_small_：
+ // 統一尺寸（SDQ_FORCE_UNIFORM_SLOT=1，或模型只有一種 expert 大小）時，
+ // 所有槽位的索引都 < n_small_ 但大小都是大的 —— 用索引判斷會讓大 expert
+ // 永遠配不到槽位，整個推論卡在 admit 的重試迴圈裡。
+ bool is_big_slot(int s) const {
+ return big_bytes_ && slot_size_of(s) >= big_bytes_;
+ }
+
+ size_t used_bytes() const { return used_bytes_snapshot(); }
+ double hit_rate() const;
+
+ // arena 裡實際被占用的位元組數。大小兩種槽位之後不能用「槽位數 × slot_bytes_」
+ // 計算（slot_bytes_ 是大的那一種），所以真的把槽位走一遍。
+ // 呼叫端：dump_stats（每次解碼階段結束）與 print_summary，不是熱路徑。
+ size_t used_bytes_snapshot() const {
+ size_t b = 0;
+ for (size_t s = 0; s < slots_.size(); s++) {
+ if (slots_[s].layer >= 0) {
+ b += slot_size_of((int) s);
+ }
+ }
+ return b;
+ }
+
+ // 讀 SSD（同步，給 acquire 用）
+ bool read_expert(int il, int ie, uint8_t * dst);
+
+ // ---- I/O 執行緒池（見 config::io_parallel 的說明）----
+ // 一個「讀取批次」的完成狀態；放在呼叫端的 stack 上，worker 只在 remaining==0
+ // 之後才會再碰它，所以離開 read_expert_par() 之後不會有 dangling pointer。
+ struct io_ctx {
+ std::atomic<int> remaining{0};
+ bool ok = true;
+ };
+ struct io_job {
+ int shard;
+ uint64_t off;
+ uint8_t * dst;
+ size_t size;
+ io_ctx * ctx;
+ };
+
+ void io_pool_start();
+ void io_pool_stop();
+ void io_worker(size_t wid);
+ bool read_expert_par(int il, int ie, uint8_t * dst);
+ bool read_range_staged(int shard, uint64_t off, uint8_t * dst, size_t size,
+ uint8_t * stage, size_t stage_sz);
+
+ // I/O 執行緒池每條 worker 自帶的 O_DIRECT 對齊緩衝（啟動時配置，已計入 RSS）
+ struct io_stage { uint8_t * buf = nullptr; size_t sz = 0; };
+ std::vector<io_stage> io_stage_;
+
+ bool io_pool_on_ = false;
+ int io_threads_ = 0;
+ int io_split_ = 1;
+ std::vector<std::thread> io_pool_;
+ std::mutex io_mu_;
+ std::condition_variable io_cv_; // 有工作 / 要停
+ std::condition_variable io_done_cv_; // 某個批次讀完了
+ std::deque<io_job> io_q_;
+ bool io_stop_ = false;
+ std::atomic<uint64_t> io_jobs_{0};
+ std::atomic<uint64_t> io_wait_ns_{0};
+
+private:
+ pager() = default;
+ ~pager();
+
+ struct slot_state {
+ int32_t layer = -1;
+ int32_t expert = -1;
+ uint64_t last_use = 0; // LRU 記時
+ uint32_t freq = 0; // 這個 expert 被用到幾次（衰減式；用來做 LFU 混合策略）
+ uint32_t pin = 0; // >0 時不可淘汰
+ bool busy = false; // 背景執行緒正在寫入
+ bool is_hot = false; // 屬於剖面排行榜：淘汰時最後才會被選中
+ bool from_prefetch = false; // page was filled by the prefetcher
+ };
+
+ struct pf_item {
+ int slot;
+ int il;
+ int ie;
+ uint64_t issue_time_us;
+ };
+
+ void touch_freq_locked(slot_state & st);
+ // ---- OPT-1：剖面 / 熱頁 ----
+ void count_request(int il, int ie); // 記一次請求（給剖面用）
+ bool load_profile(const std::string & path); // 讀剖面並標記前 hot_pin 名
+ void dump_profile(const std::string & path); // 輸出排行榜
+ bool hot_lookup(int il, int ie) const;
+ void load_gguf_header(const std::string & path);
+ void build_page_map(const std::string & path);
+ void alloc_arena();
+ int find_slot_locked(int il, int ie);
+ int admit_slot_locked(int il, int ie);
+ int victim_locked(int il = -1, int ie = -1);
+ int victim_stale_locked(uint64_t stale_before);
+ int victim_stale_for(int il, int ie, uint64_t stale_before);
+ int victim_least_stale_locked(int il, int ie, uint64_t stale_before);
+ void pf_worker();
+ bool read_range(int shard, uint64_t off, uint8_t * dst, size_t size);
+ bool read_range_direct(int shard, uint64_t off, uint8_t * dst, size_t size);
+ void open_model_fd();
+
+ // ---- sharded GGUF：模型被切成多個 -0000N-of-000M.gguf，每個檔有自己的
+ // header、自己的 data_start 與自己的張量子集。pager 必須同時開著全部的 fd，
+ // 並把每個張量的「檔內偏移」記成 (shard, off)。
+ std::vector<std::string> shard_paths_;
+ std::vector<int> shard_fd_;
+ std::vector<int> shard_fd_direct_;
+ std::vector<uint64_t> shard_data_start_;
+ bool use_direct_ = false;
+ std::string shard_path(int i) const {
+ return (i >= 0 && i < (int) shard_paths_.size()) ? shard_paths_[(size_t) i] : std::string();
+ }
+ // 依 GGUF 的 split 命名規則列舉同目錄的所有 shard
+ static std::vector<std::string> list_splits(const std::string & first_path, int n_split);
+
+ bool ready_ = false;
+ config cfg_;
+ gguf_info info_;
+ std::vector<layer_desc> layers_;
+ int n_moe_layers_ = 0;
+ uint64_t expert_bytes_total_ = 0;
+
+ uint8_t * arena_ = nullptr;
+ size_t arena_bytes_ = 0;
+ size_t slot_bytes_ = 0; // = max(small_bytes_, big_bytes_)，保留給舊的訊息輸出
+ size_t small_bytes_ = 0; // 小槽位大小（大多數層用）
+ size_t big_bytes_ = 0; // 大槽位大小；0 = 這個模型沒有大層
+ size_t n_small_ = 0;
+ size_t n_big_ = 0;
+ size_t n_slots_ = 0;
+ mutable std::mutex mu_;
+ std::vector<slot_state> slots_;
+ std::unordered_map<int64_t, int32_t> page_; // (layer<<16|expert) -> slot
+ std::vector<int32_t> free_slots_; // 空的小槽位
+ std::vector<int32_t> free_big_slots_;
+
+ // ---- OPT-1：剖面 / 熱頁（Strata 式：靜態排行榜 + 熱頁不淘汰）----
+ std::vector<uint32_t> prof_; // 每個 (layer, expert) 的累計請求次數
+ std::unique_ptr<layer_stat[]> lstats_; // SDQ-0004
+ void ensure_layer_stats(int n_layer); // SDQ-0004
+ std::vector<uint8_t> hot_; // 是否在排行榜前 hot_pin 名
+ int n_hot_ = 0;
+ std::string profile_out_;
+ uint32_t freq_tick_ = 0;
+ std::atomic<size_t> used_{0};
+ uint64_t clock_ = 0;
+
+ // 背景預取
+ std::thread pf_;
+ std::mutex pf_mu_;
+ std::condition_variable pf_cv_;
+ std::deque<pf_item> pf_q_;
+ bool pf_stop_ = false;
+ std::atomic<uint64_t> pf_inflight_{0};
+
+ // 路由歷史（一階 Markov：上個 token 的 expert -> 這次被選到的次數）
+ struct trans { int32_t to[64]; uint32_t cnt[64]; int32_t n = 0; };
+ std::vector<std::vector<trans>> hist_; // [layer][prev_expert]
+ std::vector<std::vector<int32_t>> prev_used_; // [layer] 上一個 token 用到的 expert（預測用）
+ std::vector<std::vector<int32_t>> cur_used_; // [layer] 當前 token 用到的 expert
+ std::mutex hist_mu_; // 保護 cur_used_ / prev_used_ / hist_（多執行緒會同時 note_used）
+ std::atomic<uint64_t> hist_lock_{0};
+
+ stats stats_;
+};
+
+// ---------------------------------------------------------------- 工具
+uint64_t now_us();
+size_t process_rss_bytes();
+size_t physical_ram_bytes();
+
+// ---------------------------------------------------------------- ggml 介接
+// 自訂 buffer type：expert 權重佔 0 個 RAM 位元組，資料留在檔案裡
+ggml_backend_buffer_type_t sdq_expert_buft();
+
+// 註冊 model params 的 tensor buffer override（expert 權重 → 我們的 buffer）
+void sdq_install_buft_overrides(llama_model_params * params);
+
+// 建立被分頁的 MoE FFN（替代 build_moe_ffn 裡的三次 mul_mat_id）
+ggml_tensor * sdq_build_moe(ggml_context * ctx, ggml_tensor * cur,
+ ggml_tensor * ids, ggml_tensor * weights, int il,
+ int64_t n_expert, int64_t n_expert_used,
+ float swiglu_limit);
+
+// 設定 arena 大小 / 預取開關等（CLI 在載入模型前呼叫）
+bool sdq_init(const std::string & model_path);
+void sdq_set_config(const config & c);
+const config & sdq_get_config();
+void sdq_set_enabled(bool on);
+void sdq_shutdown();
+
+bool sdq_enabled();
+
+// SDQ-0005: bridge for front ends that are not sdq-chat (llama-server).
+// Defined in sdq_pager.cpp; safe to call when the pager was never initialised.
+void sdq_on_decode(int n_tokens);
+
+// 效能計時探針的結果（SDQ_PROFILE_TIME=1）
+void sdq_print_time_profile();
+
+// MoE op 路徑計數（row-parallel / mode C），sdq_shutdown 時印出
+void sdq_print_op_counts();
+
+// 自我測試：算一次指定 expert 的 MoE 輸出（給 NumPy 對照用）
+int sdq_moe_selftest(int il, int ie, const float * x, float * out_io,
+ int64_t n_embd, int64_t n_ff, int n_tok);
+
+// 建立我們的 MoE 自訂節點。
+// **必須把 gf 傳進來並呼叫 ggml_build_forward_expand(gf, out)**：
+// ggml 的排程器只會為「進了圖的節點」配置 compute buffer，沒註冊的節點其輸出
+// 緩衝區會被 liveness 分析判成可重用，導致 layer>=1 的結果被其他張量覆蓋
+// （實測症狀：layer 0 正確、layer 1 起整個錯掉）。
+ggml_tensor * sdq_build_moe(ggml_cgraph * gf, ggml_context * ctx, ggml_tensor * cur,
+ ggml_tensor * ids, ggml_tensor * weights, int il,
+ int64_t n_expert, int64_t n_expert_used, float swiglu_limit);
+
+// Nemotron-H LatentMoE：routed experts 只有 up/down，激活 relu(x)^2。
+ggml_tensor * sdq_build_moe_ex(ggml_cgraph * gf, ggml_context * ctx, ggml_tensor * cur,
+ ggml_tensor * ids, ggml_tensor * weights, int il,
+ int64_t n_expert, int64_t n_expert_used,
+ bool no_gate, bool relu_sqr);
+
+// 這一層的分頁是否可用（graph build 時呼叫）
+bool sdq_pager_active_for_layer(int il);
+
+// 除錯／驗證用：把某個 F32 張量 dump 到檔案。
+// 必須在「原生 mul_mat_id 路徑」與「SDQ 分頁路徑」兩邊都呼叫，才能做真正的 A/B 對照
+// （env: SDQ_MOE_OUT_DUMP=prefix、SDQ_MOE_OUT_DUMP_LAYER=層編號，-1 = 全部）
+ggml_tensor * sdq_maybe_dump_tensor(ggml_context * ctx, ggml_tensor * t, int il);
+
+
+} // namespace sdq
\ No newline at end of file
diff --git a/tools/runmem.py b/tools/runmem.py
new file mode 100644
index 0000000..9d5e96c
--- /dev/null
+++ b/tools/runmem.py
@@ -0,0 +1,80 @@
+#!/usr/bin/env python3
+"""Sample process-tree RSS/swap while running a command."""
+
+import argparse
+import json
+import signal
+import subprocess
+import sys
+import time
+
+
+def process_tree(root):
+ output = subprocess.check_output(["ps", "-e", "-o", "pid=,ppid="], text=True)
+ children = {}
+ for line in output.splitlines():
+ pid, ppid = map(int, line.split())
+ children.setdefault(ppid, []).append(pid)
+ result, todo = [], [root]
+ while todo:
+ pid = todo.pop()
+ result.append(pid)
+ todo.extend(children.get(pid, ()))
+ return result
+
+
+def main():
+ parser = argparse.ArgumentParser()
+ parser.add_argument("--interval", type=float, default=0.25)
+ parser.add_argument("--timeout", type=float, default=7200)
+ parser.add_argument("--rss-limit-gib", type=float)
+ parser.add_argument("command", nargs=argparse.REMAINDER)
+ args = parser.parse_args()
+ command = args.command[1:] if args.command[:1] == ["--"] else args.command
+ if not command:
+ parser.error("missing command")
+
+ started = time.monotonic()
+ process = subprocess.Popen(command)
+ peak_rss = peak_swap = 0
+ exceeded = False
+ while process.poll() is None:
+ rss = swap = 0
+ for pid in process_tree(process.pid):
+ try:
+ with open(f"/proc/{pid}/status", encoding="ascii") as status:
+ for line in status:
+ if line.startswith("VmRSS:"):
+ rss += int(line.split()[1])
+ elif line.startswith("VmSwap:"):
+ swap += int(line.split()[1])
+ except (FileNotFoundError, ProcessLookupError):
+ pass
+ peak_rss = max(peak_rss, rss)
+ peak_swap = max(peak_swap, swap)
+ # VmRSS 以 kB 為單位，門檻須換算成 kB（1 GiB = 1024 * 1024 kB）
+ if args.rss_limit_gib is not None and rss > args.rss_limit_gib * (1024 ** 2):
+ exceeded = True
+ process.send_signal(signal.SIGKILL)
+ break
+ if time.monotonic() - started > args.timeout:
+ process.kill()
+ break
+ time.sleep(args.interval)
+ result = process.wait()
+ print("RUNMEM " + json.dumps({
+ "peak_kb": peak_rss,
+ "peak_gib": round(peak_rss / 1048576, 3),
+ "limit_gib": args.rss_limit_gib,
+ "limit_exceeded": exceeded,
+ "swap_kb": peak_swap,
+ "wall_s": round(time.monotonic() - started, 2),
+ "exit": result,
+ "cmd": command,
+ "method": "procfs-rss-watchdog",
+ }))
+ return 0 if args.rss_limit_gib is None else (1 if exceeded else result)
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())
diff --git a/tools/runmem_cgroup.py b/tools/runmem_cgroup.py
new file mode 100755
index 0000000..3b983b0
--- /dev/null
+++ b/tools/runmem_cgroup.py
@@ -0,0 +1,83 @@
+#!/usr/bin/env python3
+"""Run a command inside a cgroup-v2 MemoryMax scope and report peak memory."""
+
+import argparse
+import json
+import os
+import subprocess
+import sys
+import time
+
+
+def main():
+ parser = argparse.ArgumentParser()
+ parser.add_argument("--limit-gib", type=int, required=True)
+ parser.add_argument("--timeout", type=float, default=7200)
+ parser.add_argument("command", nargs=argparse.REMAINDER)
+ args = parser.parse_args()
+ command = args.command[1:] if args.command[:1] == ["--"] else args.command
+ if not command:
+ parser.error("missing command")
+ if args.limit_gib not in (2, 4, 8, 16, 32):
+ parser.error("limit must be 2, 4, 8, 16, or 32 GiB")
+
+ unit = f"flint-epq-bench-{os.getpid()}"
+ runner = [
+ "systemd-run", "--quiet", "--wait", "--collect", "--pipe",
+ "--unit", unit, "--property", f"MemoryMax={args.limit_gib}G",
+ "--property", "MemorySwapMax=0", "--", *command,
+ ]
+ started = time.monotonic()
+ process = subprocess.Popen(runner, stdout=subprocess.PIPE, stderr=subprocess.STDOUT, text=True)
+ time.sleep(1)
+ if process.poll() is not None:
+ output, _ = process.communicate()
+ if output:
+ sys.stdout.write(output)
+ print("UNTESTED: systemd/cgroup MemoryMax unavailable; RLIMIT_DATA is not a substitute", file=sys.stderr)
+ return 77
+
+ peak = swap_peak = 0
+ while process.poll() is None:
+ try:
+ cgroup = subprocess.check_output(
+ ["systemctl", "show", f"{unit}.scope", "-p", "ControlGroup", "--value"],
+ text=True, stderr=subprocess.DEVNULL,
+ ).strip()
+ if cgroup:
+ path = os.path.join("/sys/fs/cgroup", cgroup.lstrip("/"))
+ for name, target in (("memory.peak", "peak"), ("memory.swap.peak", "swap")):
+ try:
+ with open(os.path.join(path, name), encoding="ascii") as stream:
+ value = int(stream.read().strip())
+ if target == "peak":
+ peak = max(peak, value)
+ else:
+ swap_peak = max(swap_peak, value)
+ except OSError:
+ pass
+ except (OSError, subprocess.CalledProcessError):
+ pass
+ if time.monotonic() - started > args.timeout:
+ process.kill()
+ break
+ time.sleep(0.25)
+
+ output, _ = process.communicate()
+ if output:
+ sys.stdout.write(output)
+ print("RUNMEM " + json.dumps({
+ "limit_gib": args.limit_gib,
+ "peak_bytes": peak,
+ "peak_gib": round(peak / (1024 ** 3), 3),
+ "swap_peak_bytes": swap_peak,
+ "wall_s": round(time.monotonic() - started, 2),
+ "exit": process.returncode,
+ "cmd": command,
+ "method": "cgroup-v2-MemoryMax",
+ }))
+ return process.returncode
+
+
+if __name__ == "__main__":
+ raise SystemExit(main())