diff --git a/common/CMakeLists.txt b/common/CMakeLists.txt index 799d227519f9..0af436b54b69 100644 --- a/common/CMakeLists.txt +++ b/common/CMakeLists.txt @@ -41,7 +41,7 @@ configure_file(${TEMPLATE_FILE} ${OUTPUT_FILE}) set(TARGET llama-common-base) add_library(${TARGET} STATIC ${OUTPUT_FILE}) -target_include_directories(${TARGET} PUBLIC .) +target_include_directories(${TARGET} PUBLIC . ../src) if (BUILD_SHARED_LIBS) set_target_properties(${TARGET} PROPERTIES POSITION_INDEPENDENT_CODE ON) diff --git a/common/arg.cpp b/common/arg.cpp index 4cb853c7a4b0..d07e7c592976 100644 --- a/common/arg.cpp +++ b/common/arg.cpp @@ -6,6 +6,7 @@ #include "download.h" #include "json-schema-to-grammar.h" #include "llama.h" +#include "llama-expert-preload.h" #include "log.h" #include "sampling.h" #include "speculative.h" @@ -35,6 +36,7 @@ #include #include #include +#include #include // for hardware_concurrency #include @@ -884,6 +886,22 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context params.cors_origins = "localhost"; } + // manual hot store slots need all MoE weights in the CPU (host pointers); + // auto-activate -cmoe unless the user already did (or wants autofit slots) + if (params.expert_hot_s > 0) { + bool has_cmoe = false; + for (const auto & o : params.tensor_buft_overrides) { + if (o.pattern != nullptr && strcmp(o.pattern, LLM_FFN_EXPS_REGEX) == 0) { + has_cmoe = true; + break; + } + } + if (!has_cmoe) { + params.tensor_buft_overrides.push_back(llm_ffn_exps_cpu_override()); + LOG_WRN("manually selecting --expert-hot-s slots activates --cmoe (all MoE weights kept in the CPU)\n"); + } + } + // pad tensor_buft_overrides for llama_params_fit: const size_t ntbo = llama_max_tensor_buft_overrides(); while (params.tensor_buft_overrides.size() < ntbo) { @@ -2679,6 +2697,88 @@ common_params_context common_params_parser_init(common_params & params, llama_ex } } ).set_env("LLAMA_ARG_N_CPU_MOE")); + add_opt(common_arg( + {"--expert-heat-decay"}, "F", + "expert heatmap decay rate per update (default: 0.999)", + [](common_params & params, const std::string & value) { + params.expert_heat_decay = std::stof(value); + } + ).set_env("LLAMA_ARG_EXPERT_HEAT_DECAY")); + add_opt(common_arg( + {"--expert-heat-log-period"}, "N", + "print the expert heatmap at generation end (default: 0, 0 = off)", + [](common_params & params, int value) { + params.expert_heat_log_period = value; + } + ).set_env("LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD")); + add_opt(common_arg( + {"--expert-sync-period"}, "N", + "expert hot store re-sync cadence in tokens (default: 1)", + [](common_params & params, int value) { + params.expert_sync_period = value; + } + ).set_env("LLAMA_ARG_EXPERT_SYNC_PERIOD")); + add_opt(common_arg( + {"--expert-hyst"}, "F", + "expert hot store hysteresis ratio (default: 1.3, 0 = off)", + [](common_params & params, const std::string & value) { + params.expert_hyst = std::stof(value); + } + ).set_env("LLAMA_ARG_EXPERT_HYST")); + add_opt(common_arg( + {"--expert-dwell"}, "N", + "expert hot store minimum dwell updates before swap (default: 0 = off)", + [](common_params & params, int value) { + params.expert_dwell = value; + } + ).set_env("LLAMA_ARG_EXPERT_DWELL")); + add_opt(common_arg( + {"-ehs", "--expert-hot-s"}, "N", + "-1 = autofit slots from free VRAM, 0 = disabled, N = manual top-N slots", + [](common_params & params, int value) { + params.expert_hot_s = value; + llama_expert_preload::set_slots(value); + } + ).set_env("LLAMA_ARG_EXPERT_HOT_S")); + add_opt(common_arg( + {"--expert-pin"}, "N", + "fraction (percent) of cold experts to keep pinned in RAM via madvise, " + "0 = off, -1 = auto (hot store sets 40, else 0)", + [](common_params & params, int value) { + params.expert_pin_pct = value; + } + ).set_env("LLAMA_ARG_EXPERT_PIN")); + add_opt(common_arg( + {"--expert-no-evict"}, + {}, + "never evict experts from the hot store (fill-only, no move-back)", + [](common_params &, bool value) { + llama_expert_preload::set_no_evict(value); + } + )); + add_opt(common_arg( + {"--expert-move-mode"}, "N", + "expert store mode: 0 = auto, 1 = copy (keep RAM copy), 2 = move " + "(free RAM after verified transfer)", + [](common_params & params, int value) { + params.expert_move_mode = value; + } + ).set_env("LLAMA_ARG_EXPERT_MOVE_MODE")); + add_opt(common_arg( + {"--expert-sidecar"}, + {}, + "load the expert heatmap sidecar (.tier) at start, save it at exit", + [](common_params & params, bool value) { + params.expert_sidecar = value; + } + ).set_env("LLAMA_ARG_EXPERT_SIDECAR")); + add_opt(common_arg( + {"--expert-gpu"}, "N", + "put the expert store on this GPU index (default: -1 = all GPUs)", + [](common_params & params, int value) { + params.expert_gpu = value; + } + ).set_env("LLAMA_ARG_EXPERT_GPU")); GGML_ASSERT(params.n_gpu_layers < 0); // string_format would need to be extended for a default >= 0 add_opt(common_arg( {"-ngl", "--gpu-layers", "--n-gpu-layers"}, "N", diff --git a/common/common.cpp b/common/common.cpp index ffe3e7761bfa..41f12dd08094 100644 --- a/common/common.cpp +++ b/common/common.cpp @@ -1239,12 +1239,48 @@ common_init_result::common_init_result(common_params & params, bool model_only) if (params.fit_params) { COM_TRC("%s", "fitting params to device memory ...\n"); COM_TRC("%s", "(for bugs during this step try to reproduce them with -fit off, or provide --verbose logs if the bug only occurs with -fit on)\n"); - common_fit_params(params.model.path.c_str(), &mparams, &cparams, + int n_expert_hot_s = params.expert_hot_s; + int * p_expert_hot_s = params.expert_hot_s == -1 ? &n_expert_hot_s : nullptr; + const common_params_fit_status fit_status = common_fit_params(params.model.path.c_str(), &mparams, &cparams, params.tensor_split, params.tensor_buft_overrides.data(), params.fit_params_target.data(), params.fit_params_min_ctx, - params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR); + params.verbosity >= LOG_LEVEL_DEBUG ? GGML_LOG_LEVEL_DEBUG : GGML_LOG_LEVEL_ERROR, + p_expert_hot_s); + if (params.expert_hot_s == -1) { + // -1 = autofit slots from what the fit leaves on GPU; send all experts + // to CPU so the hot store copy reads host pointers (<=> -cmoe). + params.expert_hot_s = n_expert_hot_s > 0 ? n_expert_hot_s : 0; + cparams.expert_hot_s = params.expert_hot_s; + if (params.expert_hot_s > 0) { + for (auto & o : params.tensor_buft_overrides) { + if (o.pattern == nullptr) { + o.buft = ggml_backend_cpu_buffer_type(); + o.pattern = LLM_FFN_EXPS_REGEX; + break; + } + } + } else if (fit_status == COMMON_PARAMS_FIT_STATUS_FAILURE) { + LOG_WRN("%s: --expert-hot-s -1 autofit aborted (explicit -ngl/-ncmoe or fit error); expert cache is OFF\n", + __func__); + } else { + LOG_WRN("%s: --expert-hot-s -1 autofit found no free VRAM for expert slots; expert cache is OFF\n", + __func__); + } + } + } else if (params.expert_hot_s == -1) { + // autofit only runs inside --fit; without it -1 is meaningless + params.expert_hot_s = 0; + cparams.expert_hot_s = params.expert_hot_s; + LOG_WRN("%s: --expert-hot-s -1 requires --fit (disabled by -fit off or explicit -ngl/-ncmoe); expert cache is OFF\n", + __func__); + } + + // --expert-pin -1 auto: 40 with the hot store, 0 without + if (params.expert_pin_pct == -1) { + params.expert_pin_pct = params.expert_hot_s != 0 ? 40 : 0; + cparams.expert_pin_pct = params.expert_pin_pct; } llama_model * model = llama_model_load_from_file(params.model.path.c_str(), mparams); @@ -1667,6 +1703,18 @@ struct llama_context_params common_context_params_to_llama(const common_params & cparams.type_k = params.cache_type_k; cparams.type_v = params.cache_type_v; + cparams.expert_heat_decay = params.expert_heat_decay; + cparams.expert_heat_log_period = params.expert_heat_log_period; + cparams.expert_hot_s = params.expert_hot_s; + cparams.expert_sync_period = params.expert_sync_period; + cparams.expert_hyst = params.expert_hyst; + cparams.expert_dwell = params.expert_dwell; + cparams.expert_pin_pct = params.expert_pin_pct; + cparams.expert_move_mode = params.expert_move_mode; + cparams.expert_sidecar = params.expert_sidecar; + cparams.expert_gpu = params.expert_gpu; + cparams.model_path = params.model.path.c_str(); + return cparams; } diff --git a/common/common.h b/common/common.h index 4811345f9864..6d4cc47cb61f 100644 --- a/common/common.h +++ b/common/common.h @@ -523,6 +523,17 @@ struct common_params { int32_t verbosity = 3; // LOG_LEVEL_INFO int32_t control_vector_layer_start = -1; // layer range for control vector int32_t control_vector_layer_end = -1; // layer range for control vector + + float expert_heat_decay = 0.999f; // multiplicative decay per update + int expert_heat_log_period = 0; // print heatmap at generation end (0 = off) + int expert_hot_s = 0; // top-S expert slots (0 = disabled) + int expert_sync_period = 1; // hot store re-sync cadence in tokens + float expert_hyst = 1.3f; // hysteresis ratio: only swap when cold >= hyst x hot + int expert_dwell = 0; // minimum updates a resident slot must keep before a swap + int expert_pin_pct = -1; // percent of cold experts to keep pinned (madvise); -1 = auto + int expert_move_mode = 0; // expert store mode: 0 = auto, 1 = copy, 2 = move + bool expert_sidecar = false; // load/save the expert heatmap sidecar (.tier) + int expert_gpu = -1; // expert store GPU index (-1 = all GPUs) bool offline = false; int32_t ppl_stride = 0; // stride for perplexity calculations. If left at 0, the pre-existing approach will be used. diff --git a/common/fit.cpp b/common/fit.cpp index dd1f3ef76619..2390a102fec6 100644 --- a/common/fit.cpp +++ b/common/fit.cpp @@ -178,7 +178,7 @@ common_device_memory_data_vec common_get_device_memory_data( static void common_params_fit_impl( const char * path_model, struct llama_model_params * mparams, struct llama_context_params * cparams, float * tensor_split, struct llama_model_tensor_buft_override * tensor_buft_overrides, - size_t * margins_s, uint32_t n_ctx_min, enum ggml_log_level log_level) { + size_t * margins_s, uint32_t n_ctx_min, enum ggml_log_level log_level, int * n_expert_hot_s) { if (mparams->split_mode == LLAMA_SPLIT_MODE_TENSOR) { throw common_params_fit_exception("llama_params_fit is not implemented for SPLIT_MODE_TENSOR, abort"); } @@ -228,6 +228,8 @@ static void common_params_fit_impl( int64_t sum_projected_free = 0; int64_t sum_projected_used = 0; int64_t sum_projected_model = 0; + int64_t total_moe_bytes = 0; // MoE expert tensor bytes (for slot autofit) + int64_t dense_model_gpu = 0; // dense-only model bytes on GPU (for slot autofit) std::vector projected_free_per_device; projected_free_per_device.reserve(nd); @@ -541,6 +543,8 @@ static void common_params_fit_impl( for (size_t id = 0; id < nd; id++) { global_surplus_cpu_moe += dmds_cpu_moe[id].free; global_surplus_cpu_moe -= int64_t(dmds_cpu_moe[id].mb.total()) + margins[id]; + total_moe_bytes += int64_t(dmds_full[id].mb.model) - int64_t(dmds_cpu_moe[id].mb.model); + dense_model_gpu += int64_t(dmds_cpu_moe[id].mb.model); } if (global_surplus_cpu_moe > 0) { @@ -641,6 +645,10 @@ static void common_params_fit_impl( } if (hp_nex == 0 || global_surplus_cpu_moe <= 0) { set_ngl_tensor_split_tbo(ngl_per_device, overflow_bufts, *mparams); + if (n_expert_hot_s) { + // all MoE stays on CPU (no surplus), so no GPU hot slots fit + *n_expert_hot_s = 0; + } return; } @@ -786,6 +794,29 @@ static void common_params_fit_impl( } set_ngl_tensor_split_tbo(ngl_per_device, overflow_bufts, *mparams); + + // step 5: autofit the expert hot store slots when --expert-hot-s -1 is set. + // the fit above left some MoE bytes on GPU (final_gpu_model - dense_model_gpu); + // s = experts-per-layer that fit, minus one plane for the sentinel slot. + if (n_expert_hot_s && total_moe_bytes > 0) { + const dmds_t dmds_final = common_get_device_memory_data_impl( + path_model, mparams, cparams, devs, hp_ngl, hp_nct, hp_nex, log_level); + int64_t final_gpu_model = 0; + bool is_vulkan = false; + for (size_t id = 0; id < nd; id++) { + final_gpu_model += dmds_final[id].mb.model; + if (dev_names[id].find("Vulkan") != std::string::npos) { + is_vulkan = true; + } + } + // Vulkan reserves per-layer descriptor/pool memory the fit does not + // account for; subtract an 8 MiB per offloaded layer estimate so S + // does not overshoot and OOM at graph capture. + const int64_t vulkan_padding = is_vulkan ? int64_t(hp_ngl) * 8 * MiB : 0; + const int64_t moe_on_gpu = final_gpu_model - dense_model_gpu - vulkan_padding; + const int64_t s = moe_on_gpu > 0 ? int64_t(hp_nex) * moe_on_gpu / total_moe_bytes : 0; + *n_expert_hot_s = s > 1 ? (int) (s - 1) : 0; + } } enum common_params_fit_status common_fit_params( @@ -796,11 +827,12 @@ enum common_params_fit_status common_fit_params( llama_model_tensor_buft_override * tensor_buft_overrides, size_t * margins, uint32_t n_ctx_min, - ggml_log_level log_level) { + ggml_log_level log_level, + int * n_expert_hot_s) { const int64_t t0_us = llama_time_us(); common_params_fit_status status = COMMON_PARAMS_FIT_STATUS_SUCCESS; try { - common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, log_level); + common_params_fit_impl(path_model, mparams, cparams, tensor_split, tensor_buft_overrides, margins, n_ctx_min, log_level, n_expert_hot_s); LOG_TRC("%s: successfully fit params to free device memory\n", __func__); } catch (const common_params_fit_exception & e) { LOG_WRN("%s: failed to fit params to free device memory: %s\n", __func__, e.what()); diff --git a/common/fit.h b/common/fit.h index 208fc30694e0..95dd153a5d7c 100644 --- a/common/fit.h +++ b/common/fit.h @@ -24,7 +24,8 @@ common_params_fit_status common_fit_params( llama_model_tensor_buft_override * tensor_buft_overrides, // writable buffer for overrides, needs at least llama_max_tensor_buft_overrides elements size_t * margins, // margins of memory to leave per device in bytes uint32_t n_ctx_min, // minimum context size to set when trying to reduce memory use - ggml_log_level log_level); // minimum log level to print during fitting, lower levels go to debug log + ggml_log_level log_level, // minimum log level to print during fitting, lower levels go to debug log + int * n_expert_hot_s = nullptr);// out: fitted expert hot store slots, untouched if not applicable // print estimated memory to stdout void common_fit_print( diff --git a/counter.md b/counter.md new file mode 100644 index 000000000000..f668a91730b4 --- /dev/null +++ b/counter.md @@ -0,0 +1,42 @@ +# counter + +Times the user's input produced a materially better decision than my default. +Update only when the user asks. + +## Count: 17 + +## Examples (2026-08-06 session) + +1. **Rotation + cooldown=0 test** - I concluded the corruption was AMD-specific; the user's test ("removing the cooldown should instantly corrupt the rtx?") proved it is Vulkan-wide and duplicate-id driven. + +2. **"It's not the model"** - I attributed a failure to model randomness; the user's 200-run knowledge + the CUDA IQ2 test proved it is the tier/Vulkan. + +3. **Sentinel + mask must stay** - I claimed copy-on-read eliminates them; the user asked "is the sentinel not still necessary?" and was right (alignment + Vulkan safety). + +4. **-no-cnv invalidates corruption tests** - EOS "failures" were ambiguous without conversation mode; I had judged corruption from run counts. + +5. **--fit-target 64 was missing** - the correct fit flag changed the measured config. + +6. **-ehs -1 autofit** - the valid config revealed the tier is ~61 tok/s (faster than my invalid S=96 numbers). + +7. **Native+lazy+madvise instead of the custom pool** - "use llama's native rampool and send a release... load them in vram from disk" replaced two committed pool phases with a simpler, better design. + +8. **"RAM allocation is not actually instant"** - caught that the pool's thousands of per-slice mallocs are slow vs one native allocation. + +9. **Hash is of the memory bytes, not the output** - "I said generate a hash that can only be generated from the memory" - avoided a float-tolerance mess. + +10. **"Try 1024"** - shrinking the hash sample from 16KB to 1KB recovered ~4 tok/s. + +11. **Copy-at-init doubles VRAM** - "will we lose the ability to use the 3gb for actual slots?" - caught the transient double-buffer that would OOM an 8GB card. + +12. **Uniform first-S startup** - the user chose it over my heatmap-seeded idea. + +13. **32+8 memory-fits constraint** - "we crash and oom if the model cant fit in the ram + gpu" - shaped the startup as memory distribution, not just warming. + +14. **"Are you sure there is no other way?"** - led to lifting the gate and discovering the real n_tokens>1 blocker was a mask shape assert, not the predicted kernel crash. + +15. **Kernel speed loss unacceptable** - pushed to the count+rank kernel fix (v3 reference) over batch-split. + +16. **Deferred release** - madvise after verification, less CPU overhead. + +17. **Streaming to GPU instead of second disk read** - "move the layers into the gpu in 128mb chunks" - better than my re-read fallback. diff --git a/ggml/include/ggml-rpc.h b/ggml/include/ggml-rpc.h index 276aea00ea1b..4b4f3ba07c84 100644 --- a/ggml/include/ggml-rpc.h +++ b/ggml/include/ggml-rpc.h @@ -8,10 +8,10 @@ extern "C" { #define RPC_PROTO_MAJOR_VERSION 5 #define RPC_PROTO_MINOR_VERSION 0 -#define RPC_PROTO_PATCH_VERSION 0 +#define RPC_PROTO_PATCH_VERSION 2 #ifdef __cplusplus -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT has changed - update RPC_PROTO_PATCH_VERSION"); #endif #define GGML_RPC_MAX_SERVERS 16 diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 5cb49d0ee482..59f636b87f83 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -590,6 +590,9 @@ extern "C" { GGML_OP_GLU, + GGML_OP_MUL_MAT_ID_COLD, + GGML_OP_MOE_COLD, + GGML_OP_COUNT, }; @@ -1448,6 +1451,42 @@ extern "C" { struct ggml_tensor * b, struct ggml_tensor * ids); + // mul_mat_id restricted to cold experts only: computes only rows whose + // expert is marked 1 in cold_mask (i32 [n_expert], 1 = cold); hot slots + // are zeroed in the result + // counts (optional) accumulates per-expert routed hits, index [n_expert] = total; + // ptrs (optional) is reserved for the RAM pool and is not used by this op + GGML_API struct ggml_tensor * ggml_mul_mat_id_cold( + struct ggml_context * ctx, + struct ggml_tensor * as, + struct ggml_tensor * b, + struct ggml_tensor * ids, + struct ggml_tensor * cold_mask, + struct ggml_tensor * counts, + struct ggml_tensor * ptrs); + + // fused cold-expert MoE for one layer: down(act(gate(x)) * up(x)) computed + // on the CPU for cold experts only (cold_mask[i] == 1); hot slots zeroed. + // act is 0 = silu (separate gate/up), 1 = gelu (fused gate_up tensor). + // counts (optional) accumulates per-expert routed hits, index [n_expert] = total. + // x must be [n_embd, 1, n_tokens]; result is [down->ne[1], ids->ne[0], n_tokens] + GGML_API struct ggml_tensor * ggml_moe_cold( + struct ggml_context * ctx, + struct ggml_tensor * gate, + struct ggml_tensor * up, + struct ggml_tensor * down, + struct ggml_tensor * x, + struct ggml_tensor * ids, + struct ggml_tensor * cold_mask, + struct ggml_tensor * counts, + int32_t act); + + // optional per-expert slice source for the MoE cold op; base lib so the + // dlopen'd ggml-cpu module and the llama lib both resolve it + typedef const uint8_t * (*ggml_mmid_cold_slice_fn)(const struct ggml_tensor * src0, int expert); + GGML_API void ggml_mmid_cold_set_slice_fn(ggml_mmid_cold_slice_fn fn); + GGML_API const uint8_t * ggml_mmid_cold_get_slice(const struct ggml_tensor * src0, int expert); + // A: m columns, n rows, // B: p columns, n rows, // result is m columns, p rows diff --git a/ggml/src/ggml-cpu/CMakeLists.txt b/ggml/src/ggml-cpu/CMakeLists.txt index 836bae4d05a7..f22fe6af4bda 100644 --- a/ggml/src/ggml-cpu/CMakeLists.txt +++ b/ggml/src/ggml-cpu/CMakeLists.txt @@ -29,6 +29,10 @@ function(ggml_add_cpu_backend_variant_impl tag_name) list (APPEND GGML_CPU_SOURCES ggml-cpu/ggml-cpu.c ggml-cpu/ggml-cpu.cpp + ggml-cpu/ggml-cpu-mul-mat-id-cold.c + ggml-cpu/ggml-cpu-mul-mat-id-cold.h + ggml-cpu/ggml-cpu-moe-cold.c + ggml-cpu/ggml-cpu-moe-cold.h ggml-cpu/repack.cpp ggml-cpu/repack.h ggml-cpu/hbm.cpp diff --git a/ggml/src/ggml-cpu/ggml-cpu-moe-cold.c b/ggml/src/ggml-cpu/ggml-cpu-moe-cold.c new file mode 100644 index 000000000000..4af91baf976b --- /dev/null +++ b/ggml/src/ggml-cpu/ggml-cpu-moe-cold.c @@ -0,0 +1,320 @@ +#define _CRT_SECURE_NO_DEPRECATE // Disables "unsafe" warnings on Windows +#define _USE_MATH_DEFINES // For M_PI on MSVC + +#include "ggml.h" +#include "ggml-impl.h" +#include "ggml-cpu-impl.h" +#include "ggml-cpu.h" +#include "ops.h" +#include "ggml-cpu-moe-cold.h" +#include "ggml-cpu-mul-mat-id-cold.h" + +#if defined(_MSC_VER) || defined(__MINGW32__) +#include // using malloc.h with MSC/MINGW +#elif !defined(__FreeBSD__) && !defined(__NetBSD__) && !defined(__OpenBSD__) +#include +#endif + +#include +#include +#include + +#if defined(_WIN32) + +#define WIN32_LEAN_AND_MEAN +#ifndef NOMINMAX + #define NOMINMAX +#endif +#include + +#endif // _WIN32 + +#if defined(_MSC_VER) && !defined(__clang__) + +typedef volatile LONG atomic_int; + +typedef enum { + memory_order_relaxed, + memory_order_consume, + memory_order_acquire, + memory_order_release, + memory_order_acq_rel, + memory_order_seq_cst +} memory_order; + +static LONG atomic_fetch_add_explicit(atomic_int * ptr, LONG inc, memory_order mo) { + return InterlockedExchangeAdd(ptr, inc); +} + +#else // clang +#include +#endif + +// __builtin_prefetch is a GCC/Clang builtin; MSVC has no equivalent, so +// compile it to a no-op there +#if defined(_MSC_VER) && !defined(__clang__) +#define PREFETCH(p) ((void) 0) +#else +#define PREFETCH(p) __builtin_prefetch(p, 0, 3) +#endif + +// resolve a cold expert's weight slice through the registered hook, else the +// tensor's own data (mmap or the model buffer) +static const char * moe_cold_slice(const struct ggml_tensor * w, int64_t e, const char * fallback) { + const uint8_t * s = ggml_mmid_cold_get_slice(w, (int) e); + return s ? (const char *) s : fallback; +} + +// fused cold-expert MoE: computes down(act(gate(x)) * up(x)) only for slots +// whose expert is cold (cold_mask[e] == 1); hot slots are zeroed. +// src = {gate, up, down, x, ids, cold_mask}; op param 0 = act (0 silu, 1 gelu) +void ggml_compute_forward_moe_cold( + const struct ggml_compute_params * params, + struct ggml_tensor * dst) { + + const struct ggml_tensor * w_gate = dst->src[0]; + const struct ggml_tensor * w_up = dst->src[1]; + const struct ggml_tensor * w_down = dst->src[2]; + const struct ggml_tensor * x = dst->src[3]; + const struct ggml_tensor * ids = dst->src[4]; + const struct ggml_tensor * mask = dst->src[5]; + // counts (optional) accumulates per-expert routed hits, index [n_as] = total + int32_t * counts = dst->src[6] ? (int32_t *) dst->src[6]->data : NULL; + const int32_t act = ggml_get_op_params_i32(dst, 0); + + const int32_t * cold_mask = (const int32_t *) mask->data; + + const int ith = params->ith; + const int nth = params->nth; + + // gate/up share a type; down may use a different quant + const enum ggml_type type_g = w_gate->type; + const enum ggml_type type_d = w_down->type; + GGML_ASSERT(w_up->type == type_g); + GGML_ASSERT(x->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(x->ne[1] == 1 && x->nb[0] == sizeof(float)); + + const int64_t ne_embd = x->ne[0]; + const int64_t n_ff = (w_gate == w_up) ? (w_gate->ne[1] / 2) : w_gate->ne[1]; + const int64_t n_tokens = x->ne[2]; + const int64_t n_out = w_down->ne[1]; + const int n_ids = ids->ne[0]; + const int n_as = w_gate->ne[2]; + + ggml_vec_dot_t const vec_dot_g = ggml_get_type_traits_cpu(type_g)->vec_dot; + enum ggml_type const vdt_g = ggml_get_type_traits_cpu(type_g)->vec_dot_type; + ggml_from_float_t const from_fx = ggml_get_type_traits_cpu(vdt_g)->from_float; + ggml_vec_dot_t const vec_dot_d = ggml_get_type_traits_cpu(type_d)->vec_dot; + enum ggml_type const vdt_d = ggml_get_type_traits_cpu(type_d)->vec_dot_type; + ggml_from_float_t const from_fa = ggml_get_type_traits_cpu(vdt_d)->from_float; + + const size_t q_embd = ggml_row_size(vdt_g, ne_embd); + const size_t q_ff = ggml_row_size(vdt_d, n_ff); + + void * wdata_cur = params->wdata; + + char * xq = (char *) incr_ptr_aligned(&wdata_cur, n_tokens*q_embd, sizeof(int64_t)); + + int64_t * matrix_row_counts = + (int64_t *) incr_ptr_aligned(&wdata_cur, n_as*sizeof(int64_t), sizeof(int64_t)); + + struct mmid_row_mapping * matrix_rows = + (struct mmid_row_mapping *) incr_ptr_aligned(&wdata_cur, n_as*(int64_t)n_ids*n_tokens*sizeof(struct mmid_row_mapping), sizeof(int64_t)); + + int64_t * col0 = + (int64_t *) incr_ptr_aligned(&wdata_cur, n_as*sizeof(int64_t), sizeof(int64_t)); + + char (*atomic_current_chunk)[CACHE_LINE_SIZE] = + (char (*)[CACHE_LINE_SIZE]) incr_ptr_aligned(&wdata_cur, CACHE_LINE_SIZE*n_as, CACHE_LINE_SIZE); + + // intermediates for the cold slots, indexed by global cold column + const int64_t maxc = (int64_t) n_ids*n_tokens; + float * gate_out = (float *) incr_ptr_aligned(&wdata_cur, 2*n_ff*maxc*sizeof(float), CACHE_LINE_SIZE); + float * up_out = gate_out + n_ff*maxc; + char * act_q = (char *) incr_ptr_aligned(&wdata_cur, q_ff*maxc, CACHE_LINE_SIZE); + + GGML_ASSERT(params->wsize >= (size_t)((char *) wdata_cur - (char *) params->wdata)); + + // quantize x once, shared by all experts + for (int64_t t = ith; t < n_tokens; t++) { + from_fx((const float *) ((const char *) x->data + t*x->nb[2]), xq + t*q_embd, ne_embd); + } + + if (ith == 0) { + memset(dst->data, 0, ggml_nbytes(dst)); + memset(matrix_row_counts, 0, n_as*sizeof(int64_t)); + + for (int64_t t = 0; t < n_tokens; t++) { + for (int id = 0; id < n_ids; id++) { + const int32_t e = *(const int32_t *) ((const char *) ids->data + t*ids->nb[1] + id*ids->nb[0]); + GGML_ASSERT(e >= 0 && e < n_as); + if (counts) { + counts[e]++; + counts[n_as]++; + } + if (cold_mask[e] == 0) { + continue; + } + matrix_rows[e*(int64_t)n_ids*n_tokens + matrix_row_counts[e]] = (struct mmid_row_mapping) {id, (int32_t) t}; + matrix_row_counts[e]++; + } + } + + int64_t coff = 0; + for (int e = 0; e < n_as; e++) { + col0[e] = coff; + coff += matrix_row_counts[e]; + } + } + + // reset phase-A chunk counters + for (int e = ith; e < n_as; e += nth) { + atomic_int * ctr = (atomic_int *)(atomic_current_chunk + e); + *ctr = nth; + } + + ggml_barrier(params->threadpool); + + // phase A: gate/up dots for all cold slots into gate_out/up_out + for (int cur_a = 0; cur_a < n_as; ++cur_a) { + const int64_t cne1 = matrix_row_counts[cur_a]; + if (cne1 == 0) { + continue; + } + const char * wg = moe_cold_slice(w_gate, cur_a, (const char *) w_gate->data + cur_a*w_gate->nb[2]); + const char * wu = (w_gate == w_up) ? (wg + n_ff*w_gate->nb[1]) + : moe_cold_slice(w_up, cur_a, (const char *) w_up->data + cur_a*w_up->nb[2]); + + const int64_t nr0 = n_ff; + const int64_t nr1 = cne1; + + int chunk_size = 16; + if (nr1 == 1) { + chunk_size = 64; + } + const bool disable_chunking = ggml_is_numa(); + int64_t nchunk0 = (nr0 + chunk_size - 1)/chunk_size; + int64_t nchunk1 = (nr1 + chunk_size - 1)/chunk_size; + if (nchunk0*nchunk1 < nth*4 || disable_chunking) { + nchunk0 = nr0 > nr1 ? nth : 1; + nchunk1 = nr0 > nr1 ? 1 : nth; + } + const int64_t dr0 = (nr0 + nchunk0 - 1)/nchunk0; + const int64_t dr1 = (nr1 + nchunk1 - 1)/nchunk1; + + int current_chunk = ith; + atomic_int * ctr = (atomic_int *)(atomic_current_chunk + cur_a); + + while (current_chunk < nchunk0*nchunk1) { + const int64_t ith0 = current_chunk % nchunk0; + const int64_t ith1 = current_chunk / nchunk0; + const int64_t ir0_start = dr0*ith0, ir0_end = MIN(ir0_start + dr0, nr0); + const int64_t ir1_start = dr1*ith1, ir1_end = MIN(ir1_start + dr1, nr1); + + for (int64_t c = ir1_start; c < ir1_end; c++) { + const struct mmid_row_mapping rm = matrix_rows[cur_a*(int64_t)n_ids*n_tokens + c]; + const char * xcol = xq + (int64_t) rm.i2*q_embd; + float * gout = gate_out + (col0[cur_a] + c)*n_ff; + float * uout = up_out + (col0[cur_a] + c)*n_ff; + for (int64_t i = ir0_start; i < ir0_end; i++) { + PREFETCH(wg + (i + 1)*w_gate->nb[1]); + PREFETCH(wu + (i + 1)*w_up->nb[1]); + vec_dot_g(ne_embd, &gout[i], 0, wg + i*w_gate->nb[1], 0, xcol, 0, 1); + vec_dot_g(ne_embd, &uout[i], 0, wu + i*w_up->nb[1], 0, xcol, 0, 1); + } + } + + if (nth >= nchunk0*nchunk1) { + break; + } + current_chunk = atomic_fetch_add_explicit(ctr, 1, memory_order_relaxed); + } + } + + ggml_barrier(params->threadpool); + + // phase B (multi-threaded): activation (SwiGLU/GELU) + quantize the intermediate + const int64_t total_cols = n_as > 0 ? (col0[n_as - 1] + matrix_row_counts[n_as - 1]) : 0; + + for (int64_t c = ith; c < total_cols; c += nth) { + float * go = gate_out + c*n_ff; + const float * uo = up_out + c*n_ff; + if (act == 1) { + for (int64_t i = 0; i < n_ff; i++) { + const float g = go[i]; + const float gelu_g = 0.5f * g * (1.0f + tanhf(0.7978845608028654f * g * (1.0f + 0.044715f * g * g))); + go[i] = gelu_g * uo[i]; + } + } else { + for (int64_t i = 0; i < n_ff; i++) { + const float g = go[i]; + const float silu_g = g / (1.0f + expf(-g)); + go[i] = silu_g * uo[i]; + } + } + from_fa(go, act_q + c*q_ff, n_ff); + } + + if (ith == 0) { + // reset phase-C chunk counters + for (int e = 0; e < n_as; e++) { + atomic_int * ctr = (atomic_int *)(atomic_current_chunk + e); + *ctr = nth; + } + } + + ggml_barrier(params->threadpool); + + // phase C: down dots for all cold slots, scattered into dst + for (int cur_a = 0; cur_a < n_as; ++cur_a) { + const int64_t cne1 = matrix_row_counts[cur_a]; + if (cne1 == 0) { + continue; + } + const char * wd = moe_cold_slice(w_down, cur_a, (const char *) w_down->data + cur_a*w_down->nb[2]); + + const int64_t nr0 = n_out; + const int64_t nr1 = cne1; + + int chunk_size = 16; + if (nr1 == 1) { + chunk_size = 64; + } + const bool disable_chunking = ggml_is_numa(); + int64_t nchunk0 = (nr0 + chunk_size - 1)/chunk_size; + int64_t nchunk1 = (nr1 + chunk_size - 1)/chunk_size; + if (nchunk0*nchunk1 < nth*4 || disable_chunking) { + nchunk0 = nr0 > nr1 ? nth : 1; + nchunk1 = nr0 > nr1 ? 1 : nth; + } + const int64_t dr0 = (nr0 + nchunk0 - 1)/nchunk0; + const int64_t dr1 = (nr1 + nchunk1 - 1)/nchunk1; + + int current_chunk = ith; + atomic_int * ctr = (atomic_int *)(atomic_current_chunk + cur_a); + + while (current_chunk < nchunk0*nchunk1) { + const int64_t ith0 = current_chunk % nchunk0; + const int64_t ith1 = current_chunk / nchunk0; + const int64_t ir0_start = dr0*ith0, ir0_end = MIN(ir0_start + dr0, nr0); + const int64_t ir1_start = dr1*ith1, ir1_end = MIN(ir1_start + dr1, nr1); + + for (int64_t c = ir1_start; c < ir1_end; c++) { + const struct mmid_row_mapping rm = matrix_rows[cur_a*(int64_t)n_ids*n_tokens + c]; + const char * acol = act_q + (col0[cur_a] + c)*q_ff; + float * dst_col = (float *) ((char *) dst->data + rm.i1*dst->nb[1] + (int64_t) rm.i2*dst->nb[2]); + for (int64_t j = ir0_start; j < ir0_end; j++) { + float res = 0.0f; + vec_dot_d(n_ff, &res, 0, wd + j*w_down->nb[1], 0, acol, 0, 1); + dst_col[j] += res; + } + } + + if (nth >= nchunk0*nchunk1) { + break; + } + current_chunk = atomic_fetch_add_explicit(ctr, 1, memory_order_relaxed); + } + } +} diff --git a/ggml/src/ggml-cpu/ggml-cpu-moe-cold.h b/ggml/src/ggml-cpu/ggml-cpu-moe-cold.h new file mode 100644 index 000000000000..43bc58817883 --- /dev/null +++ b/ggml/src/ggml-cpu/ggml-cpu-moe-cold.h @@ -0,0 +1,18 @@ +#pragma once + +// Fused cold-expert MoE kernel: down(act(gate(x)) * up(x)) for cold experts +// only, in a single CPU op. see ggml_moe_cold() in ggml.c for the op contract. + +#include "ggml-cpu-impl.h" + +#ifdef __cplusplus +extern "C" { +#endif + +void ggml_compute_forward_moe_cold( + const struct ggml_compute_params * params, + struct ggml_tensor * dst); + +#ifdef __cplusplus +} +#endif diff --git a/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.c b/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.c new file mode 100644 index 000000000000..fee27d9d6e8e --- /dev/null +++ b/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.c @@ -0,0 +1,229 @@ +#define _CRT_SECURE_NO_DEPRECATE // Disables "unsafe" warnings on Windows +#define _USE_MATH_DEFINES // For M_PI on MSVC + +#include "ggml.h" +#include "ggml-impl.h" +#include "ggml-cpu-impl.h" +#include "ggml-cpu.h" +#include "ops.h" +#include "ggml-cpu-mul-mat-id-cold.h" + +#if defined(_MSC_VER) || defined(__MINGW32__) +#include // using malloc.h with MSC/MINGW +#elif !defined(__FreeBSD__) && !defined(__NetBSD__) && !defined(__OpenBSD__) +#include +#endif + +#include +#include +#include +#include +#include +#include + +#if defined(_WIN32) + +#define WIN32_LEAN_AND_MEAN +#ifndef NOMINMAX + #define NOMINMAX +#endif +#include + +#endif // _WIN32 + +#if defined(_MSC_VER) && !defined(__clang__) + +typedef volatile LONG atomic_int; + +typedef enum { + memory_order_relaxed, + memory_order_consume, + memory_order_acquire, + memory_order_release, + memory_order_acq_rel, + memory_order_seq_cst +} memory_order; + +static LONG atomic_fetch_add_explicit(atomic_int * ptr, LONG inc, memory_order mo) { + return InterlockedExchangeAdd(ptr, inc); +} + +#else // clang +#include +#endif + +void ggml_compute_forward_mul_mat_id_cold( + const struct ggml_compute_params * params, + struct ggml_tensor * dst) { + + const struct ggml_tensor * src0 = dst->src[0]; + const struct ggml_tensor * src1 = dst->src[1]; + const struct ggml_tensor * ids = dst->src[2]; + const struct ggml_tensor * mask = dst->src[3]; + // counts (optional) accumulates per-expert routed hits; ptrs (optional) is + // reserved for the RAM pool and is not used by this op + int32_t * counts = dst->src[4] ? (int32_t *) dst->src[4]->data : NULL; + + const int32_t * cold_mask = (const int32_t *) mask->data; + + GGML_TENSOR_BINARY_OP_LOCALS + + const int ith = params->ith; + const int nth = params->nth; + + const enum ggml_type type = src0->type; + + const bool src1_cont = ggml_is_contiguous(src1); + + const struct ggml_type_traits_cpu * ttype = ggml_get_type_traits_cpu(type); + enum ggml_type const vec_dot_type = ttype->vec_dot_type; + ggml_from_float_t const from_float = ggml_get_type_traits_cpu(vec_dot_type)->from_float; + + GGML_ASSERT(nb00 == ggml_type_size(type)); + GGML_ASSERT(nb10 == ggml_type_size(src1->type)); + + GGML_ASSERT(nb0 == sizeof(float)); + GGML_ASSERT(nb0 <= nb1); + GGML_ASSERT(nb1 <= nb2); + GGML_ASSERT(nb2 <= nb3); + + const int n_ids = ids->ne[0]; + const int n_as = ne02; + + void * wdata_cur = params->wdata; + + if (src1->type != vec_dot_type) { + incr_ptr_aligned(&wdata_cur, ggml_row_size(vec_dot_type, ggml_nelements(src1)), sizeof(int64_t)); + } + + int64_t * matrix_row_counts = + incr_ptr_aligned(&wdata_cur, n_as*sizeof(int64_t), sizeof(int64_t)); + + struct mmid_row_mapping * matrix_rows = + incr_ptr_aligned(&wdata_cur, n_as*ids->ne[0]*ids->ne[1]*sizeof(struct mmid_row_mapping), sizeof(int64_t)); + + char (*atomic_current_chunk)[CACHE_LINE_SIZE] = + incr_ptr_aligned(&wdata_cur, CACHE_LINE_SIZE * n_as, CACHE_LINE_SIZE); + + GGML_ASSERT(params->wsize >= (size_t)((char *) wdata_cur - (char *) params->wdata)); + + if (src1->type != vec_dot_type) { + char * wdata = params->wdata; + + const size_t nbw0 = ggml_type_size(vec_dot_type); + const size_t nbw1 = ggml_row_size(vec_dot_type, ne10); + const size_t nbw2 = nbw1*ne11; + const size_t nbw3 = nbw2*ne12; + + assert(params->wsize >= ne13*nbw3); + GGML_ASSERT(src1->type == GGML_TYPE_F32); + + for (int64_t i13 = 0; i13 < ne13; ++i13) { + for (int64_t i12 = 0; i12 < ne12; ++i12) { + for (int64_t i11 = 0; i11 < ne11; ++i11) { + size_t bs = ggml_blck_size(vec_dot_type); + int64_t ne10_block_start = (ith * ne10/bs) / nth; + int64_t ne10_block_end = ((ith + 1) * ne10/bs) / nth; + from_float((float *)((char *) src1->data + i13*nb13 + i12*nb12 + i11*nb11 + ne10_block_start*bs*nb10), + (void *) (wdata + i13*nbw3 + i12*nbw2 + i11*nbw1 + ne10_block_start*nbw0), + (ne10_block_end - ne10_block_start) * bs); + } + } + } + } + + if (ith == 0) { + memset(dst->data, 0, ggml_nbytes(dst)); + memset(matrix_row_counts, 0, n_as*sizeof(int64_t)); + + for (int64_t iid1 = 0; iid1 < ids->ne[1]; ++iid1) { + for (int id = 0; id < n_ids; ++id) { + const int32_t i02 = *(const int32_t *) ((const char *) ids->data + iid1*ids->nb[1] + id*ids->nb[0]); + + assert(i02 >= 0 && i02 < n_as); + + if (counts) { + counts[i02]++; + counts[n_as]++; + } + + if (cold_mask[i02] == 0) { + continue; + } + + MMID_MATRIX_ROW(i02, matrix_row_counts[i02]) = (struct mmid_row_mapping) {id, iid1}; + matrix_row_counts[i02] += 1; + } + } + } + + for (int cur_a = ith; cur_a < n_as; cur_a += nth) { + atomic_int * current_chunk_ctr = (atomic_int *)(atomic_current_chunk + cur_a); + *current_chunk_ctr = nth; + } + + ggml_barrier(params->threadpool); + + for (int cur_a = 0; cur_a < n_as; ++cur_a) { + const int64_t cne1 = matrix_row_counts[cur_a]; + if (cne1 == 0) { + continue; + } + + // prefer the registered per-expert slice source (expert preload), else + // read from the tensor's own data + const char * src0_cur; + const uint8_t * s = ggml_mmid_cold_get_slice(src0, (int) cur_a); + src0_cur = s ? (const char *) s : (const char *) src0->data + cur_a * nb02; + const void * wdata = (src1->type == vec_dot_type) ? src1->data : params->wdata; + const size_t row_size = ggml_row_size(vec_dot_type, ne10); + + const int64_t nr0 = ne01; + const int64_t nr1 = cne1; + + int chunk_size = 16; + if (nr0 == 1 || nr1 == 1) { + chunk_size = 64; + } + + const bool disable_chunking = ggml_is_numa(); + + int64_t nchunk0 = (nr0 + chunk_size - 1) / chunk_size; + int64_t nchunk1 = (nr1 + chunk_size - 1) / chunk_size; + + if (nchunk0 * nchunk1 < nth * 4 || disable_chunking) { + nchunk0 = nr0 > nr1 ? nth : 1; + nchunk1 = nr0 > nr1 ? 1 : nth; + } + + const int64_t dr0 = (nr0 + nchunk0 - 1) / nchunk0; + const int64_t dr1 = (nr1 + nchunk1 - 1) / nchunk1; + + int current_chunk = ith; + + atomic_int * current_chunk_ctr = (atomic_int *)(atomic_current_chunk + cur_a); + + while (current_chunk < nchunk0 * nchunk1) { + const int64_t ith0 = current_chunk % nchunk0; + const int64_t ith1 = current_chunk / nchunk0; + + const int64_t ir0_start = dr0 * ith0; + const int64_t ir0_end = MIN(ir0_start + dr0, nr0); + + const int64_t ir1_start = dr1 * ith1; + const int64_t ir1_end = MIN(ir1_start + dr1, nr1); + + ggml_compute_forward_mul_mat_id_one_chunk( + dst, src0, src1, ids, cur_a, + ir0_start, ir0_end, ir1_start, ir1_end, + src0_cur, matrix_rows, row_size, src1_cont, wdata + ); + + if (nth >= nchunk0 * nchunk1) { + break; + } + + current_chunk = atomic_fetch_add_explicit(current_chunk_ctr, 1, memory_order_relaxed); + } + } +} diff --git a/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.h b/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.h new file mode 100644 index 000000000000..3a5aa03d5036 --- /dev/null +++ b/ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.h @@ -0,0 +1,46 @@ +#pragma once + +// Shared helpers between ggml-cpu.c (stock GGML_OP_MUL_MAT_ID) and the +// dedicated MUL_MAT_ID_COLD kernel in ggml-cpu-mul-mat-id-cold.cpp. + +#include "ggml-cpu-impl.h" + +#include +#include + +#define MMID_MATRIX_ROW(row_id, i1) matrix_rows[(row_id)*ids->ne[0]*ids->ne[1] + (i1)] + +struct mmid_row_mapping { + int32_t i1; + int32_t i2; +}; + +#ifdef __cplusplus +extern "C" { +#endif + +void * incr_ptr_aligned(void ** p, size_t size, size_t align); + +void ggml_compute_forward_mul_mat_id_one_chunk( + struct ggml_tensor * dst, + const struct ggml_tensor * src0, + const struct ggml_tensor * src1, + const struct ggml_tensor * ids, + const int64_t cur_a, + const int64_t ir0_start, + const int64_t ir0_end, + const int64_t ir1_start, + const int64_t ir1_end, + const char * src0_cur, + const struct mmid_row_mapping * matrix_rows, + const size_t row_size, + const bool src1_cont, + const void * wdata); + +void ggml_compute_forward_mul_mat_id_cold( + const struct ggml_compute_params * params, + struct ggml_tensor * dst); + +#ifdef __cplusplus +} +#endif diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 491316f74912..87f09a5f9cf7 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -5,6 +5,8 @@ #include "ggml-backend.h" #include "traits.h" #include "ggml-cpu-impl.h" +#include "ggml-cpu-mul-mat-id-cold.h" +#include "ggml-cpu-moe-cold.h" #include "ggml-impl.h" #include "quants.h" #include "ggml-threading.h" @@ -1453,14 +1455,7 @@ UseGgmlGemm2:; // ggml_compute_forward_mul_mat_id -#define MMID_MATRIX_ROW(row_id, i1) matrix_rows[(row_id)*ids->ne[0]*ids->ne[1] + (i1)] - -struct mmid_row_mapping { - int32_t i1; - int32_t i2; -}; - -static void ggml_compute_forward_mul_mat_id_one_chunk( +void ggml_compute_forward_mul_mat_id_one_chunk( struct ggml_tensor * dst, const struct ggml_tensor * src0, const struct ggml_tensor * src1, @@ -1523,7 +1518,7 @@ static void ggml_compute_forward_mul_mat_id_one_chunk( } } -static void * incr_ptr_aligned(void ** p, size_t size, size_t align) { +void * incr_ptr_aligned(void ** p, size_t size, size_t align) { void * ptr = *p; ptr = (void *) GGML_PAD((uintptr_t) ptr, align); @@ -1841,6 +1836,14 @@ static void ggml_compute_forward(struct ggml_compute_params * params, struct ggm { ggml_compute_forward_mul_mat_id(params, tensor); } break; + case GGML_OP_MUL_MAT_ID_COLD: + { + ggml_compute_forward_mul_mat_id_cold(params, tensor); + } break; + case GGML_OP_MOE_COLD: + { + ggml_compute_forward_moe_cold(params, tensor); + } break; case GGML_OP_OUT_PROD: { ggml_compute_forward_out_prod(params, tensor); @@ -2217,6 +2220,7 @@ static void set_numa_thread_affinity(int thread_n) { UNUSED(thread_n); } static void clear_numa_thread_affinity(void) {} #endif + static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { int n_tasks = 0; @@ -2329,6 +2333,8 @@ static int ggml_get_n_tasks(struct ggml_tensor * node, int n_threads) { case GGML_OP_CONCAT: case GGML_OP_MUL_MAT: case GGML_OP_MUL_MAT_ID: + case GGML_OP_MUL_MAT_ID_COLD: + case GGML_OP_MOE_COLD: case GGML_OP_OUT_PROD: { n_tasks = n_threads; @@ -2854,6 +2860,7 @@ struct ggml_cplan ggml_graph_plan( } } break; case GGML_OP_MUL_MAT_ID: + case GGML_OP_MUL_MAT_ID_COLD: { cur = 0; const struct ggml_tensor * src0 = node->src[0]; @@ -2872,6 +2879,30 @@ struct ggml_cplan ggml_graph_plan( // atomic_current_chunk cur += CACHE_LINE_SIZE*n_as + CACHE_LINE_SIZE; } break; + case GGML_OP_MOE_COLD: + { + cur = 0; + const struct ggml_tensor * w_gate = node->src[0]; + const struct ggml_tensor * w_down = node->src[2]; + const struct ggml_tensor * x = node->src[3]; + const struct ggml_tensor * ids = node->src[4]; + const enum ggml_type vdt_g = type_traits_cpu[w_gate->type].vec_dot_type; + const enum ggml_type vdt_d = type_traits_cpu[w_down->type].vec_dot_type; + const int n_as = w_gate->ne[2]; + const int64_t maxc = ids->ne[0]*ids->ne[1]; + // quantized x + cur += ggml_row_size(vdt_g, ggml_nelements(x)) + sizeof(int64_t); + // matrix_row_counts + col0 + cur += 2*n_as*sizeof(int64_t) + 2*sizeof(int64_t); + // matrix_rows + cur += n_as*maxc*sizeof(struct mmid_row_mapping) + sizeof(int64_t); + // atomic chunk counters + cur += CACHE_LINE_SIZE*n_as + CACHE_LINE_SIZE; + // gate_out + up_out + cur += 2*w_gate->ne[1]*maxc*sizeof(float) + CACHE_LINE_SIZE; + // act_q + cur += ggml_row_size(vdt_d, w_gate->ne[1])*maxc + CACHE_LINE_SIZE; + } break; case GGML_OP_OUT_PROD: { if (ggml_is_quantized(node->src[0]->type) || diff --git a/ggml/src/ggml-cuda/mmf.cu b/ggml/src/ggml-cuda/mmf.cu index 646a5899c803..f60be42e8569 100644 --- a/ggml/src/ggml-cuda/mmf.cu +++ b/ggml/src/ggml-cuda/mmf.cu @@ -85,7 +85,7 @@ void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * sr GGML_ASSERT(sis1 > 0); ggml_cuda_launch_mm_ids_helper(ids_d, ids_src_compact_dev.get(), ids_dst_compact_dev.get(), expert_bounds_dev.get(), - static_cast(n_experts), static_cast(n_tokens), static_cast(n_expert_used), static_cast(ne11), si1, sis1, /*write_inverse =*/ false, ctx.stream()); + static_cast(n_experts), static_cast(n_tokens), static_cast(n_expert_used), static_cast(ne11), si1, sis1, false, ctx.stream(), false); CUDA_CHECK(cudaGetLastError()); ids_info.ids_src_compact = ids_src_compact_dev.get(); diff --git a/ggml/src/ggml-cuda/mmid.cu b/ggml/src/ggml-cuda/mmid.cu index f80442fbe4e8..e2dbca2ba0ff 100644 --- a/ggml/src/ggml-cuda/mmid.cu +++ b/ggml/src/ggml-cuda/mmid.cu @@ -39,23 +39,37 @@ static __global__ void mm_ids_helper( int it_compact = 0; // Running index for the compact slice of this expert. if constexpr (n_expert_used_template == 0) { - // Generic implementation: + // Generic implementation. With expert tiering, several slots of one + // token can map to the same (sentinel) expert, so count hits per + // token and give each hitting lane its own store slot via its rank. for (int it = 0; it < n_tokens; ++it) { - int iex_used = -1; // The index at which the expert is used, if any. + int cnt = 0; for (int iex = threadIdx.x; iex < n_expert_used; iex += warp_size) { const int expert_used = ids[it*si1 + iex]; nex_prev += expert_used < expert; - if (expert_used == expert) { - iex_used = iex; - } + cnt += expert_used == expert; } - if (iex_used != -1) { - store[it_compact] = mm_ids_helper_store(it, iex_used); + // inclusive warp scan of cnt; n_hit broadcast from the last lane + int rank = cnt; +#pragma unroll + for (int offset = 1; offset < warp_size; offset <<= 1) { + const int tmp = __shfl_up_sync(0xFFFFFFFF, rank, offset, warp_size); + if (threadIdx.x >= static_cast(offset)) { + rank += tmp; + } } - - if (warp_reduce_any(iex_used != -1)) { - it_compact++; + const int n_hit = __shfl_sync(0xFFFFFFFF, rank, warp_size - 1, warp_size); + rank -= cnt; // exclusive scan: hits held by lower lanes + + if (n_hit > 0) { + for (int iex = threadIdx.x; iex < n_expert_used; iex += warp_size) { + if (ids[it*si1 + iex] == expert) { + store[it_compact + rank] = mm_ids_helper_store(it, iex); + rank++; + } + } + it_compact += n_hit; } } } else { @@ -68,28 +82,43 @@ static __global__ void mm_ids_helper( const int iex = threadIdx.x % neu_padded; // The index at which the expert is used, if any. const int expert_used = (neu_padded == n_expert_used || iex < n_expert_used) && it < n_tokens ? ids[it*si1 + iex] : INT_MAX; - const int iex_used = expert_used == expert ? iex : -1; + const bool hit = expert_used == expert; nex_prev += expert_used < expert; - // Whether the threads at this token position have used the expert: - const int it_compact_add_self = warp_reduce_any(iex_used != -1); + // With expert tiering, several slots of one token can map to the same + // (sentinel) expert, so count hits per token instead of "any" and give + // each hitting lane its own store slot via its rank within the token. + + // Inclusive scan of hit within the token's subgroup of lanes: + int scan = hit ? 1 : 0; +#pragma unroll + for (int offset = 1; offset < neu_padded; offset <<= 1) { + const int tmp = __shfl_up_sync(0xFFFFFFFF, scan, offset, warp_size); + if ((threadIdx.x % neu_padded) >= static_cast(offset)) { + scan += tmp; + } + } + + // Hits of this token; rank of this lane among the token's hitting lanes: + const int n_hit = __shfl_sync(0xFFFFFFFF, scan, threadIdx.x | (neu_padded - 1), warp_size); + const int rank = scan - 1; // Do a scan over threads at lower token positions in warp to get the correct index for writing data: int it_compact_add_lower = 0; #pragma unroll for (int offset = neu_padded; offset < warp_size; offset += neu_padded) { - const int tmp = __shfl_up_sync(0xFFFFFFFF, it_compact_add_self, offset, warp_size); + const int tmp = __shfl_up_sync(0xFFFFFFFF, n_hit, offset, warp_size); if (threadIdx.x >= static_cast(offset)) { it_compact_add_lower += tmp; } } - if (iex_used != -1) { - store[it_compact + it_compact_add_lower] = mm_ids_helper_store(it, iex_used); + if (hit) { + store[it_compact + it_compact_add_lower + rank] = mm_ids_helper_store(it, iex); } // The thread with the highest index in the warp always has the sum over the whole warp, use it to increment all threads: - it_compact += __shfl_sync(0xFFFFFFFF, it_compact_add_lower + it_compact_add_self, warp_size - 1, warp_size); + it_compact += __shfl_sync(0xFFFFFFFF, it_compact_add_lower + n_hit, warp_size - 1, warp_size); } } nex_prev = warp_reduce_sum(nex_prev); @@ -123,7 +152,7 @@ static __global__ void mm_ids_helper( template static void launch_mm_ids_helper( const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds, - const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) { + const int n_experts, const int n_tokens, const int n_expert_used_var, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream, bool tiered_hot) { GGML_ASSERT(n_tokens < (1 << 22) && "too few bits in mm_ids_helper_store"); GGML_ASSERT(n_expert_used_var < (1 << 10) && "too few bits in mm_ids_helper_store"); @@ -134,7 +163,9 @@ static void launch_mm_ids_helper( const dim3 num_blocks(n_experts, 1, 1); const dim3 block_size(warp_size, 1, 1); - const size_t nbytes_shared = n_tokens*sizeof(mm_ids_helper_store); + // with expert tiering (.hot stores), ids can map several slots of one token + // to the same (sentinel) expert; only then size for the true worst case + const size_t nbytes_shared = (tiered_hot ? n_tokens*n_expert_used_var : n_tokens)*sizeof(mm_ids_helper_store); GGML_ASSERT(nbytes_shared <= smpbo); mm_ids_helper<<>> (ids, ids_src1, ids_dst, expert_bounds, n_tokens, n_expert_used_var, nchannels_y, si1, sis1, write_inverse); @@ -142,28 +173,28 @@ static void launch_mm_ids_helper( void ggml_cuda_launch_mm_ids_helper( const int32_t * __restrict__ ids, int32_t * __restrict__ ids_src1, int32_t * __restrict__ ids_dst, int32_t * __restrict__ expert_bounds, - const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream) { + const int n_experts, const int n_tokens, const int n_expert_used, const int nchannels_y, const int si1, const int sis1, const bool write_inverse, cudaStream_t stream, bool tiered_hot) { switch (n_expert_used) { case 2: - launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper< 2>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; case 4: - launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper< 4>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; case 6: - launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper< 6>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; case 8: - launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper< 8>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; case 16: - launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper<16>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; case 32: - launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper<32>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; default: - launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream); + launch_mm_ids_helper< 0>(ids, ids_src1, ids_dst, expert_bounds, n_experts, n_tokens, n_expert_used, nchannels_y, si1, sis1, write_inverse, stream, tiered_hot); break; } } diff --git a/ggml/src/ggml-cuda/mmid.cuh b/ggml/src/ggml-cuda/mmid.cuh index 74c2db43385e..4ba2cfe07e3f 100644 --- a/ggml/src/ggml-cuda/mmid.cuh +++ b/ggml/src/ggml-cuda/mmid.cuh @@ -2,4 +2,4 @@ void ggml_cuda_launch_mm_ids_helper( const int32_t * ids, int32_t * ids_src1, int32_t * ids_dst, int32_t * expert_bounds, - int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, bool write_inverse, cudaStream_t stream); + int n_experts, int n_tokens, int n_expert_used, int nchannels_y, int si1, int sis1, bool write_inverse, cudaStream_t stream, bool tiered_hot); diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 707437ea3e52..b89dc17a2042 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -197,8 +197,10 @@ void ggml_cuda_mul_mat_q( const int si1 = ids->nb[1] / ggml_element_size(ids); const int sis1 = nb12 / nb11; + const size_t src0_name_len = strlen(src0->name); + const bool tiered_hot_ids = src0_name_len >= 4 && strcmp(src0->name + src0_name_len - 4, ".hot") == 0; ggml_cuda_launch_mm_ids_helper((const int32_t *) ids->data, ids_src1.get(), ids_dst.get(), expert_bounds.get(), - ne02, ne12, n_expert_used, ne11, si1, sis1, /*write_inverse =*/ dedup_bcast, stream); + ne02, ne12, n_expert_used, ne11, si1, sis1, dedup_bcast, stream, tiered_hot_ids); CUDA_CHECK(cudaGetLastError()); } @@ -244,6 +246,11 @@ void ggml_cuda_mul_mat_q( ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int)); const int64_t s13 = ne12*s12; + // Tiered hot stores (name suffix .hot) map duplicate ids to a sentinel slot, so one + // expert can hold up to ne_get_rows columns; stock ids keep the tighter n_tokens bound. + const size_t src0_name_len = strlen(src0->name); + const bool is_tiered_hot = src0_name_len >= 4 && strcmp(src0->name + src0_name_len - 4, ".hot") == 0; + // Note that ne02 is used instead of ne12 because the number of y channels determines the z dimension of the CUDA grid. const mmq_args args = { src0_d, src0->type, (const int *) src1_q8_1.get(), ids_dst.get(), expert_bounds.get(), dst_d, @@ -251,7 +258,7 @@ void ggml_cuda_mul_mat_q( ne00, ne01, ne_get_rows, s01, ne_get_rows, s1, ne02, ne02, s02, s12, s2, ne03, ne13, s03, s13, s3, - ne12}; + is_tiered_hot ? ne_get_rows : ne12}; ggml_cuda_mul_mat_q_switch_type(ctx, args, stream); } diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index da7f3a5f2e30..fca71c73c768 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -1098,9 +1098,11 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "OPT_STEP_SGD", "GLU", + "MUL_MAT_ID_COLD", + "MOE_COLD", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1213,9 +1215,11 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "sgd(x)", "glu(x)", + "mul_mat_id_cold(x,x,x,x)", + "moe_cold(x,x,x,x,x,x)", }; -static_assert(GGML_OP_COUNT == 101, "GGML_OP_COUNT != 101"); +static_assert(GGML_OP_COUNT == 103, "GGML_OP_COUNT != 103"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -3352,6 +3356,107 @@ struct ggml_tensor * ggml_mul_mat_id( return result; } +// ggml_mul_mat_id_cold + +struct ggml_tensor * ggml_mul_mat_id_cold( + struct ggml_context * ctx, + struct ggml_tensor * as, + struct ggml_tensor * b, + struct ggml_tensor * ids, + struct ggml_tensor * cold_mask, + struct ggml_tensor * counts, + struct ggml_tensor * ptrs) { + GGML_ASSERT(!ggml_is_transposed(as)); + GGML_ASSERT(ids->type == GGML_TYPE_I32); + // cold_mask treated as integer for 0/non-zero check + GGML_ASSERT(cold_mask->ne[0] == as->ne[2]); + GGML_ASSERT(as->ne[3] == 1); + GGML_ASSERT(b->ne[3] == 1); + GGML_ASSERT(ids->ne[2] == 1 && ids->ne[3] == 1); + GGML_ASSERT(ids->ne[1] == b->ne[2]); + GGML_ASSERT(as->ne[0] == b->ne[0]); + GGML_ASSERT(ids->ne[0] % b->ne[1] == 0); + if (counts) { + GGML_ASSERT(counts->type == GGML_TYPE_I32); + GGML_ASSERT(counts->ne[0] >= as->ne[2] + 1); + } + if (ptrs) { + GGML_ASSERT(ptrs->type == GGML_TYPE_I64); + GGML_ASSERT(ptrs->ne[0] == as->ne[2]); + } + + const int64_t ne[4] = { as->ne[1], ids->ne[0], b->ne[2], 1 }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + + result->op = GGML_OP_MUL_MAT_ID_COLD; + result->src[0] = as; + result->src[1] = b; + result->src[2] = ids; + result->src[3] = cold_mask; + result->src[4] = counts; + result->src[5] = ptrs; + + return result; +} + +// ggml_moe_cold + +struct ggml_tensor * ggml_moe_cold( + struct ggml_context * ctx, + struct ggml_tensor * gate, + struct ggml_tensor * up, + struct ggml_tensor * down, + struct ggml_tensor * x, + struct ggml_tensor * ids, + struct ggml_tensor * cold_mask, + struct ggml_tensor * counts, + int32_t act) { + GGML_ASSERT(!ggml_is_transposed(gate) && !ggml_is_transposed(up) && !ggml_is_transposed(down)); + GGML_ASSERT(ids->type == GGML_TYPE_I32); + GGML_ASSERT(act == 0 || act == 1); + + GGML_ASSERT(gate->ne[3] == 1 && up->ne[3] == 1 && down->ne[3] == 1); + GGML_ASSERT(x->ne[3] == 1 && x->ne[1] == 1); + GGML_ASSERT(ids->ne[2] == 1 && ids->ne[3] == 1); + GGML_ASSERT(ids->ne[1] == x->ne[2]); + GGML_ASSERT(gate->ne[0] == x->ne[0] && up->ne[0] == x->ne[0]); + GGML_ASSERT(gate->ne[1] == up->ne[1]); + GGML_ASSERT(down->ne[0] == (gate == up ? gate->ne[1] / 2 : gate->ne[1])); + GGML_ASSERT(gate->ne[2] == up->ne[2] && up->ne[2] == down->ne[2]); + GGML_ASSERT(ids->ne[0] % x->ne[1] == 0); + GGML_ASSERT(cold_mask->type == GGML_TYPE_I32); + GGML_ASSERT(cold_mask->ne[0] == gate->ne[2]); + if (counts) { + GGML_ASSERT(counts->type == GGML_TYPE_I32); + GGML_ASSERT(counts->ne[0] >= gate->ne[2] + 1); + } + + const int64_t ne[4] = { down->ne[1], ids->ne[0], x->ne[2], 1 }; + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, 4, ne); + + result->op = GGML_OP_MOE_COLD; + result->src[0] = gate; + result->src[1] = up; + result->src[2] = down; + result->src[3] = x; + result->src[4] = ids; + result->src[5] = cold_mask; + result->src[6] = counts; + ggml_set_op_params_i32(result, 0, act); + + return result; +} + +static ggml_mmid_cold_slice_fn g_mmid_cold_slice_fn = NULL; + +GGML_API void ggml_mmid_cold_set_slice_fn(ggml_mmid_cold_slice_fn fn) { + g_mmid_cold_slice_fn = fn; +} + +GGML_API const uint8_t * ggml_mmid_cold_get_slice(const struct ggml_tensor * src0, int expert) { + return g_mmid_cold_slice_fn ? g_mmid_cold_slice_fn(src0, expert) : NULL; +} + // ggml_out_prod static inline bool ggml_can_out_prod(const struct ggml_tensor * t0, const struct ggml_tensor * t1) { diff --git a/include/llama.h b/include/llama.h index a14498925f14..67f6cf6a1b43 100644 --- a/include/llama.h +++ b/include/llama.h @@ -403,6 +403,20 @@ extern "C" { struct llama_sampler_seq_config * samplers; size_t n_samplers; + // model file path; used for the expert heatmap sidecar (.tier) + const char * model_path; + + float expert_heat_decay; // expert heatmap decay per update + int expert_heat_log_period; // expert heatmap log interval + int expert_hot_s; // number of top-S expert slots for GPU hot store + int expert_sync_period; // hot store re-sync cadence in tokens + float expert_hyst; // hysteresis ratio for slot swaps + int expert_dwell; // min updates a resident slot keeps before swap + int expert_pin_pct; // percent of cold experts to keep pinned (madvise); -1 = auto + bool expert_sidecar; // load/save the expert heatmap sidecar (.tier) + int expert_move_mode; // expert store mode: 0 = auto, 1 = copy, 2 = move + int expert_gpu; // put the expert store on this GPU index (-1 = all GPUs) + // a source/target/parent context // can be utilized in various ways, for example by sharing results or llama_memory between 2 contexts struct llama_context * ctx_other; @@ -569,6 +583,9 @@ extern "C" { LLAMA_API llama_memory_t llama_get_memory (const struct llama_context * ctx); LLAMA_API enum llama_pooling_type llama_pooling_type(const struct llama_context * ctx); // TODO: rename to llama_get_pooling_type + // print the expert heatmap (no-op when the tier is not active) + LLAMA_API void llama_print_expert_heatmap(const struct llama_context * ctx); + LLAMA_API const struct llama_vocab * llama_model_get_vocab(const struct llama_model * model); LLAMA_API enum llama_rope_type llama_model_rope_type(const struct llama_model * model); diff --git a/manuallog3.md b/manuallog3.md new file mode 100644 index 000000000000..a9ec39a17962 --- /dev/null +++ b/manuallog3.md @@ -0,0 +1,755 @@ +# manuallog: Project6 - Expert-Granular MoE Tiering (GPU Hot Store) + +> Session log + design/state reference for the agent working in +> `/run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6/repo2`. +> Read AGENTS.md first for the authoritative rule set and current build/test +> commands. This file is the long-form reference: state, architecture, why past +> decisions failed, and the launch/benchmark etiquette. ASCII only. +> Last updated: 2026-08-05 (section 14 CLEANED: corruption root cause = 3-node +> cold path; fix = port fused MOE_COLD op. Section 14.7 is the current truth). + +## 0. One-line status (2026-08-05) + +Single-stream MoE decode accelerated ~2.5x via an expert-granular GPU hot +store (S=96 slots) + a dedicated `MUL_MAT_ID_COLD` CPU op. Verified at ~40 +tok/s on Qwen3.6-35B-A3B IQ2_M (vs ~16 stock -cmoe) on CUDA. Multi-slot +server batches freeze the hot store (committed). **VULKAN SUPPORT DROPPED** +(2026-08-05): after a full investigation session, all tier implementations +(repo2, v2, v3.bak, v5) corrupt on Vulkan - there is no known-good Vulkan +tier to mirror. Working tree is clean at commit 1a850ec17. +MULTI-SLOT (2C gate, n_tokens>1) IS DROPPED FOR NOW: the mmid count+rank +CUDA port remains reverted, and -np>1 server batches fall back to stock. +CUDA single-stream is the supported tier path. See section 12 for eliminated +theories and the one remaining (unconfirmed) Vulkan lead. + +## 1. Project context & constraints + +- **Repo**: `Project6/repo2` (a fork of llama.cpp, merge-request targeted). + Keep the upstream diff minimal and reviewable. Read `repo2/AGENTS.md` too. +- **Goal (SCOPE - do not overextend)**: MANUAL expert-granular GPU hot store + for MoE. User sets `--expert-hot-s N` (0 = disabled). We: + 1. Find expert weight tensors per layer, compute bytes/slot. + 2. Allocate N GPU slot buffers at context init (VRAM committed at init). + 3. After the heatmap warms, copy the top-N hottest expert weight slices + from CPU to the GPU hot store (one-shot first fill). + 4. Periodically re-sync the GPU slots to mirror the current top-N + (plain top-N mirror on a token cadence, NO hysteresis gate / NO dwell). + 5. Hook the MoE ffn graph so hot experts compute from the GPU hot store, + cold experts from a dedicated CPU op, summed per routed expert. +- **Strict constraint - Rule 6 new-file isolation**: all new logic lives in + new .cpp/.h/.c files. Edits to upstream files stay minimal: single-line + hooks, struct-member or CLI-flag additions, or (where unavoidable) dropping + `static` to share a helper. The cold kernel is the only exception allowed + and it now lives in its OWN file (see Architecture). + +### Out of scope / NOT in scope +- Hysteresis ratio gate + dwell (Trick 5/6): DEFERRED. (v3.bak/v5 DO use + hysteresis + dwell internally, but repo2 deliberately keeps it simple.) +- Cross-session warm start (Trick 21): repo2 fresh heatmap each session. + (v3.bak/v5 have a `.tier` sidecar warm seed - see section 6.J.) +- Prompt harvesting / MOE_COUNT op / TMAX gate (Trick 11, 24). +- Perf tuning tricks (9/10/15) beyond what shipped. +- mmid.cu count+rank CUDA port: ATTEMPTED and REVERTED (see 6.H) - it made + multi-slot output fully corrupt. The 2C gate stays. + +If the user asks for any of the above, treat it as a new phase with its own +scope discussion. oldtricks.md is a REFERENCE for in-scope tricks, not a +feature list. + +## 2. Environment & paths (ABSOLUTE) + +- Project root: `/run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6` +- Active source: `Project6/repo2` +- Build dir (CUDA): `Project6/build` (`-DGGML_CUDA=ON -DGGML_CUDA_FA=ON -DCMAKE_BUILD_TYPE=Release`) +- Build dir (Vulkan): `Project6/build_vulkan` (`-DGGML_VULKAN=ON -DGGML_CUDA=OFF`) +- Reference / pre-fix tree: `Project5/repo2` +- Colibri original: `Project1/folder2/llama.cpp` +- **wackMall family (CRITICAL REFERENCES for the Vulkan question)**: + - `Project3/llama-wackMall_v3.bak` - "the one that does not complicate + things". Tier is auto-enabled (no CLI flag), auto-fit S, native 2D LUT, + mask on CPU buffer, discards w_s on tiered path. WORKS on Vulkan. + - `Project3/llama-wackMall_v5` - v3 + more features (RAM pool, sidecar, + autofit ALWAYS ON, hysteresis env vars). WORKS on Vulkan, faster than v3 + (IQ2 22.35 vs 18.65 tok/s on this machine). + - Both have `build_vulkan/` dirs created 2026-08-04 for this investigation. +- Models: `/run/media/miltos/Boost drive/Models/` + - Qwen3.6-35B-A3B IQ2_M (`Qwen3.6-35B-A3B-abliterated-MAX.i1-IQ2_M.gguf`): + 40 MoE layers (0-39, ALL MoE), n_expert=256, n_expert_used=8. Has per-expert + scale tensors (`*.scale` in GGUF) - relevant to the scale bug. + - Qwen3.6-35B-A3B Q5_K_P (`.../BoostHome/Models/Qwen3.6-35B-A3B-Uncensored-HauhauCS-Aggressive-Q5_K_P.gguf`): + bigger quant, per-expert scale too. Has a `.tier` sidecar (v5 warm seed). + - Non-MoE guard: Bonsai-27B-Q1_0 (3.8GB) - heatmap+hotstore stay inert. + +## 3. Architecture (current, accurate) + +### 3.1 Files we touched (in repo2) +``` +src/llama-expert-heatmap.{h,cpp} - per-layer expert usage tracking + log +src/llama-expert-hotstore.{h,cpp} - GPU hot store: alloc, copy_top_s, resync, LUTs +src/llama-expert-tier.{h,cpp} - in-graph dual-path build hook +ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.{c,h} - the MUL_MAT_ID_COLD kernel + shared mmid helpers +common/arg.cpp - --expert-hot-s / -las / --expert-heat-* flags +common/common.{h,cpp} - param plumbing +include/llama.h - context_params fields +src/llama-context.cpp - hotstore/heatmap init + post-compute hook +src/llama-graph.{cpp,h} - moe_sel_experts capture + tier dispatch hook +ggml/include/ggml.h - GGML_OP_MUL_MAT_ID_COLD enum + ctor decl +ggml/src/ggml.c - op name/symbol + constructor + OP_COUNT 102 +ggml/src/ggml-cpu/ggml-cpu.c - 3 switch hooks (forward/n_tasks/work_size) +ggml/src/ggml-cpu/CMakeLists.txt - new source added +``` + +### 3.2 Data flow +1. At context construction (gated: `hparams.n_expert > 0 && !warmup && + hot_s != 0`), build heatmap + hotstore. Hotstore picks the first non-CPU + backend's default buffer type for the GPU hot store. +2. After each ubatch's graph compute (process_ubatch), if heatmap is set, + `synchronize()` then `update_from_graph(res->moe_sel_experts)` reads the + selected-experts tensors back and bumps per-(layer,expert) heat. +3. First fill: `copy_top_s` copies the top-N hottest expert weight slices + CPU->GPU, builds hot_lut/cold_mask per layer, registers each expert + weight tensor with the tier hook. +4. Cadence re-sync: `maybe_resync` mirrors the current top-N into the slots. + NEW (committed 5c5020d80): when the ubatch has n_tokens>1 (multi-slot + server batch), `maybe_resync` is called with multi_slot=true and SKIPS + swapping - the hot store is frozen/static for the batch. +5. At graph build (`build_lora_mm_id`, called by build_moe_ffn), if `w` is + registered with the tier, `llama_expert_tier_build` constructs the dual path: + - hot: remap real ids -> slot ids via hot_lut (sentinel S for cold), + `ggml_mul_mat_id(dst_hot, cur, ids_hot)` on GPU. + - cold: `ggml_mul_mat_id_cold(w, cur, ids, cold_mask)` on CPU - skips + hot experts (cold_mask==0), computes only cold-selected rows. + - result = `add(hot, cold)`. + Returns nullptr -> caller falls back to stock `ggml_mul_mat_id`. + NOTE: the per-expert w_s scale was REMOVED from this path (see 3.5). + +### 3.3 The MUL_MAT_ID_COLD op +- Declared in `ggml.h` after `GGML_OP_GLU`; constructor `ggml_mul_mat_id_cold` + takes `(ctx, as, b, ids, cold_mask)` - 5 sources. +- `cold_mask` is f32[n_experts], 1.0f=cold/compute, 0.0f=hot/skip; the kernel + reads it as int32 zero-check. +- CPU-only. CUDA never sees it. The hot path uses stock `ggml_mul_mat_id`. +- Lives in `ggml/src/ggml-cpu/ggml-cpu-mul-mat-id-cold.c`. Uses the SAME + chunking helper as stock mul_mat_id (`ggml_compute_forward_mul_mat_id_one_chunk`), + exposed via `ggml-cpu-mul-mat-id-cold.h`. Reads type info via the public + `ggml_get_type_traits_cpu()` accessor. + +### 3.4 Working-tree changes (REVERTED 2026-08-05 - Vulkan attempt dropped) +Two candidate Vulkan fixes (Option A: removed the in-graph `w_s` scale block; +LUT native 2D) were applied and tested but were INSUFFICIENT to fix Vulkan +corruption, and the Vulkan path is now dropped entirely. They have been +REVERTED; the working tree is clean at 1a850ec17. See section 12. + +### 3.5 v3.bak reference architecture (what "works") +- Tier auto-enabled (no CLI flag; unconditional init), auto-fit S from free VRAM. +- `s.lut` native 2D `ggml_new_tensor_2d(g_ctx_gpu, I32, 1, n_expert)` on GPU. +- `s.mask` `ggml_new_tensor_1d(g_ctx_cpu, I32, n_expert)` on CPU (i32, 1=cold). + IMPORTANT: mask is on a CPU buffer - read directly by the CPU cold op. +- `s.w_hot` 3D on g_ctx_gpu (no_alloc ctx + alloc_ctx_tensors_from_buft, WEIGHTS usage). +- `build_mul_mat_id`: same remap -> get_rows(lut) -> ids_hot -> mul_mat_id(w_hot) + -> mul_mat_id_cold(w, ids, mask, ptrs) -> add. Discards w_s. +- Cold op takes an extra `s.ptrs` (i64 host weight addresses for RAM pool). +- Gate: `ids->ne[1] > g_tmax` (g_tmax default 16 in v3, 1 in v5). +- v5 adds: RAM pool (SR slots), sidecar warm seed, LLAMA_EXPERT_DECAY (default + 1.0) / LLAMA_EXPERT_HYSTERESIS (default 1.5) env vars, dwell>=32, autofit + ALWAYS ON, LLAMA_EXPERT_S/HOT/ADAPT env vars. + +## 4. Build & test (authoritative; AGENTS.md is the source of truth) + +### Configure from clean (CUDA) +```sh +cd /run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6 +mkdir -p build && cd build +cmake ../repo2 -DGGML_CUDA=ON -DGGML_CUDA_FA=ON -DCMAKE_BUILD_TYPE=Release +``` +### Configure from clean (Vulkan) +```sh +mkdir -p build_vulkan && cd build_vulkan +cmake ../repo2 -DGGML_VULKAN=ON -DGGML_CUDA=OFF -DCMAKE_BUILD_TYPE=Release +``` +### Build (incremental) +```sh +cd /run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6/build +cmake --build . -j$(nproc) --target llama-completion llama-server +# Vulkan: +cd /run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6/build_vulkan +cmake --build . -j$(nproc) --target llama-completion +``` +### Device selection (Vulkan - IMPORTANT, differs from vulkaninfo) +`./bin/llama-completion --list-devices` is authoritative. In llama.cpp's +enumeration order on this machine: +- Vulkan0 = NVIDIA GeForce RTX 3070 Laptop GPU (the FAST one) +- Vulkan1 = AMD Radeon RX 570 (RADV POLARIS10) +This is OPPOSITE to `vulkaninfo` (which shows AMD as GPU0). Use +`GGML_VK_VISIBLE_DEVICES=0` for the RTX, `=1` for the AMD. A default run +picks Vulkan0 = RTX. +### PASS/FAIL test (CUDA single-stream, the headline case) +```sh +cd /run/media/miltos/Linuxaddon/AI_Workspace/Opencode/Project6/build +ulimit -c 0 +./bin/llama-completion \ + -m "/run/media/miltos/Boost drive/Models/Qwen3.6-35B-A3B-abliterated-MAX.i1-IQ2_M.gguf" \ + -ngl 99 -cmoe -c 8000 -ctk q8_0 -ctv q8_0 -b 256 -ub 256 -t 6 \ + --kv-offload --flash-attn on -fit off --temp 0 -no-cnv \ + --expert-hot-s 96 \ + -p "Write a comprehensive technical guide to setting up a home Linux server" \ + -n 256 60 ms/token (stock ~16 tok/s), +crash, or incoherent output. +### NOTE ON --no-mmap (REMOVED PERMANENTLY, 2026-08-04) +Do NOT pass `--no-mmap` to test commands anymore. The user said it "hurts us" +permanently. Reason observed: with --no-mmap each server reads the FULL model +into RAM; two servers = ~23GB, maxed 31GB RAM + 17GB swap (looked like a leak, +was not - killing the servers freed all of it instantly). Drop it. + +## 5. Launch & benchmark etiquette (READ BEFORE ANY TEST) + +1. **Check for concurrent builds/tests first** (`ps -eo pid,pmem,comm --sort=-pmem | head`). + RAM/CPU pressure causes OOMs and spurious timeouts. Wait for it. +2. **Minimum -n 128** for generation tests. -n 16 masks bugs and is noisy. + For the wackMall/v5 caching tier use -n 256+ (it has not warmed up until then). +3. **Non-interactive**: pass -no-cnv and redirect stdin from /dev/null + (`< /dev/null`) so the test exits on its own. +4. **-st exits cleanly** on llama-cli (single-turn). Keep `timeout 60` as a + hard safety net for llama-completion (it has no -st; it exits after -n). +5. **Do NOT change user-specified flags** without asking. -ngl 99 -cmoe etc. +6. **Never use decay 0.9 for "diagnostic" tests** unless explicitly told. +7. **Use LLAMA_LOG, not LLAMA_LOG_INFO**, for any line you must see. +8. **Use the grep tool narrow**: scope `path` to a dir or a file. +9. **Quote model paths** with spaces. +10. **Do NOT pass -ngl to the CPU-only build**. +11. **Per backend build dirs**: build/ = CUDA, build_vulkan/ = Vulkan. Do not + confuse binaries. +12. **Launching a server without softlocking the harness**: do NOT use + `pkill -f llama-server` (kills your own shell). Use `pkill -x llama-server`. + To launch detached WITHOUT the harness killing it on timeout: use + `setsid --fork ./bin/llama-server ... log 2>&1 < /dev/null` + then poll /health in a SEPARATE command. The earlier nohup+disown+same-cmd + poll softlocks the harness (the timeout then SIGTERMs the process group). +13. **Commit style**: user's name (Miltos22), lowercase, no period, no + Co-authored-by/Assisted-by tags. Draft the message WITH the user before + committing. Local commits only, no push, no PR. +14. **No git mutations** (commit/push/amend/rebase) unless explicitly asked. + The user explicitly asked for the -las commit amend twice. +15. **push needs a token**: `git push https://miltos22:TOKEN@github.com/...` + (username/password auth is rejected by GitHub). The remote is the user's + fork `miltos22/llama.cpp-wackMall-merge-request.git`. Only push when asked. + +## 6. Debugging history - WHY past decisions didn't work + +### A. The `ggml_backend_tensor_get` signature trap +5 args vs real 4: `ggml_backend_tensor_get(tensor, data, offset, size)`. + +### B. OOM on CPU-only inference +35B fully on CPU (-ngl 0) triggers kernel OOM killer during PP graph alloc. +Fix: reduce -c. Watch RAM before builds/tests. + +### C. LLAMA_LOG_INFO is invisible (CRITICAL QUIRK) +LLAMA_LOG_INFO -> LOG_LEVEL_TRACE(4) filtered by INFO(3). Use LLAMA_LOG (0). + +### D. "Hysteresis" naming confusion +Heatmap accumulation rate vs swap-policy ratio gate (Trick 6). repo2 uses pure +additive counting + decay only; no ratio gate / dwell (deferred). + +### E. CLI float lambda +`common_arg` lambdas only accept `int` and `std::string`. Use std::stof. + +### F. Fabricated test data - NEVER LIE ABOUT OUTPUT +NEVER invent log lines or test output. AGENTS.md Rule 1: evidence or it did +not happen. + +### G. GPU tensor readback bug (the big one) +Async CUDA + recycled scratch buffers -> garbage expert ids. Fix: +`synchronize()` before readback AND `ggml_set_output(selected_experts)`. + +### H. The 4-step cold chain REGRESSION -> MUL_MAT_ID_COLD +Original chain: 5-node remap x 4 x 3 matrices x 40 layers = ~3,500 extra +graph nodes + dummy expert-0 rows. 1.3 tok/s. Ported GGML_OP_MUL_MAT_ID_COLD +(colibri Trick 7 Path A): single skip-hot CPU op. 40 tok/s. +Then moved the kernel to its own file (`d1cdf1b35`, pushed via sync-fork). + +### I. The cold kernel from_float NULL segfault +`type_traits_cpu[type].from_float` is NULL for IQ2_M. Fix: fetch from +`type_traits_cpu[vec_dot_type].from_float` (matches stock mul_mat_id). + +### J. The "self-contained new-file kernel" rewrite attempt +The working version REUSES stock chunking helpers; a fresh self-contained +rewrite segfaulted (bug I). Do not re-attempt without a new decision. + +### K. **mmid.cu count+rank port ATTEMPTED AND REVERTED (2026-08-04)** +Goal: lift the 2C gate (n_tokens==1) so the tier works at n_tokens>1 +(multi-slot server). Ported the colibri count+rank fix into BOTH the generic +and templated paths of repo2's mmid.cu (kept write_inverse), sized shared +memory for the worst case (n_tokens*n_expert_used), and REMOVED the gate +`if (cur->ne[2] > 1) return nullptr;` in tier.cpp. +RESULT: **fully corrupt multi-slot output** - worse than before (before = slow +but correct stock fallback; after = garbage). User: "the slots thing does not +work for multiple ones, it is actually worse than before." +REVERTED both files to committed state (git checkout). The 2C gate stays. +The user decided to instead FREEZE the hot store during multi-slot batches +(see 3.2, committed 5c5020d80). The mmid count+rank fix may be revisited only +as a separate deliberate effort; the corrupt result suggests our sentinel +duplicate handling interacts badly with the CUDA MMQ path beyond what colibri's +fix covers (or the port had a subtle bug). + +### L. Freeze-exchange analysis (and why it was chosen) +Earlier analysis said freezing does NOT fix multi-slot speed (the gate is the +cause, not resync churn). But the mmid port made things CORRUPT, so the user +chose freeze + keep the gate: multi-slot falls back to stock (correct, slow) +and the hot store stops swapping during multi-slot batches. Committed as +`5c5020d80`. Tested: multi-slot -np 8 coherent (stock speed ~2.45 t/s/slot), +single-stream 38-41 tok/s no regression. This is the current committed state. + +### M. Emoji-loop degenerate output is NOT the tier +--ignore-eos + temp-0 + this Qwen3 model drifts into repetition loops. +Isolated: happens with tier ON and OFF; only --ignore-eos triggers it. +Do not blame the tier; check --ignore-eos first. + +### N. The earlier "random server crash" was a self-inflicted kill +`pkill -f llama-server` killed the launching shell. Use `pkill -x`. + +### O. Plant_static and log_hit_rate are diagnostic, not core +Diagnostic helpers, not part of the core feature path. + +### P. **Vulkan device ordering trap (2026-08-04)** +`vulkaninfo` shows AMD as GPU0, NVIDIA as GPU1. But llama.cpp's OWN +`--list-devices` shows Vulkan0 = NVIDIA RTX 3070, Vulkan1 = AMD RX 570. +A test with GGML_VK_VISIBLE_DEVICES=1 silently ran on the AMD (9.94 tok/s, +"it still used the 570"). Correct: GGML_VK_VISIBLE_DEVICES=0 = RTX. + +### Q. **RAM was NOT leaking - it was --no-mmap (2026-08-04)** +Two --no-mmap servers held 2x 11.6GB model in RAM -> 30Gi used, 17Gi swap, +user thought it was a leak. Killing them freed 23GB instantly. Also: removing +`-c` from a run let context default huge and eat RAM alongside the v5 RAM pool +(18Gi). Always keep `-c` on multi-GB model tests. + +### R. **repo2 tier corrupts on Vulkan; v3.bak/v5 are coherent (2026-08-04)** +This is the OPEN issue. Full detail in section 9. + +### S. **v5 "autofit always on" gotcha (2026-08-04)** +v3.bak/v5 enable the tier UNCONDITIONALLY (no CLI flag needed) and autofit S +from free VRAM. So EVERY v3/v5 run is tiered - there is no "stock" v3/v5 run +to compare against without disabling the tier. My early "v5 stock-ish" 23.52 +tok/s number was tiered too (run-to-run variance / tuning churn). + +### T. **v5 cold first-run is slow (2026-08-04)** +v5 Q5 Vulkan first run 2.96 tok/s (init churn: RAM pool, sidecar build). +Warm run (2nd, sidecar present) 11.72 tok/s. First-run numbers are meaningless +for v5; use warm runs. Also the `.tier` sidecar file is a warm seed - deleting +it forces cold start. User asked to delete sidecar + LLAMA_EXPERT_DECAY=0.999 +LLAMA_EXPERT_HYSTERESIS=1.3 + no -cmoe/-ngl/-t for a "clean" v5 test: result +was coherent, 1.61 tok/s with RAM pool 18Gi (slow but works). + +## 7. CLI flags (repo2) + +| Flag | Type | Default | Env | Purpose | +|------|------|---------|-----|---------| +| `-las`, `--expert-hot-s N` | int | 0 | `LLAMA_ARG_EXPERT_HOT_S` | top-S expert slots for GPU hot store (0=disabled). -las added this session (pushed) | +| `--expert-heat-decay F` | float | 0.99 | `LLAMA_ARG_EXPERT_HEAT_DECAY` | multiplicative decay per update | +| `--expert-heat-log-period N` | int | 100 | `LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD` | log + re-sync cadence in updates | + +There is NO `LLAMA_EXPERT_S` env fallback - removed earlier, CLI only. +(v3.bak/v5 DO use LLAMA_EXPERT_S / LLAMA_EXPERT_DECAY / LLAMA_EXPERT_HYSTERESIS +/ LLAMA_EXPERT_ADAPT / LLAMA_EXPERT_HOT - but those are THEIR flags.) + +## 8. Commits & roadmap + +### Commits on this branch (origin/master has them up to 6e75753ab) +- `fcaac3d75` expert heatmap: decay-tracked usage counters per layer +- `610802380` expert heatmap: add top-S ranking, log top-8 per layer +- `21537c16f` expert heatmap: add --expert-hot-s flag for GPU hot store slot count +- `b27d2f59e` expert heatmap: fix GPU tensor readback via ggml_set_output and synchronize +- `413e65dde` expert heatmap: move readback logic into heatmap module +- `3fd577828` Merge remote-tracking branch 'origin/master' +- `23ff80ec3` expert hotstore: add per-layer expert slot sizing +- `289c41ef5` expert hotstore: allocate GPU hot store buffers for S slots +- `284215c04` expert hotstore: reduced cross contamination when expert args off +- `8e955700f` expert heatmap: count updates in tokens not layers +- `08e657d41` expert heatmap: fix decay and log trigger +- `1aefbbe58` expert hotstore: copy top-S experts to GPU after first ubatch +- `cee1740bb` expert hotstore: re-sync hot slots on cadence (stable slots) +- `a3c442540` expert hotstore: add sentinel slot for zero-contribution routing +- `904090276` expert hotstore: per-layer LUTs and masks for in-graph routing +- `8d1a68e5b` Merge remote-tracking branch 'origin/master' +- `3c9616b3a` expert tier: hook GPU hot store into graph with MUL_MAT_ID_COLD cold op +- `d1cdf1b35` ggml-cpu: move MUL_MAT_ID_COLD kernel to its own file +- `8687ee2b6` common: added -las short flag for expert hot store slots (alias of --expert-hot-s) [PUSHED] +- `6e75753ab` Merge branch 'ggml-org:master' into master [on origin/master - created by GitHub "Sync fork" button, upstream pulled in; local fast-forwarded after tar backup] +- `5c5020d80` expert hotstore: freeze swapping during multi-slot batches [LOCAL, NOT pushed, 1 ahead] +- `1a850ec17` llama : gate expert hot store to CUDA only, with force override [LOCAL, NOT pushed, 2 ahead] - CUDA-only guard + LLAMA_EXPERT_HOT_FORCE + +### Working tree (CLEAN as of 2026-08-05) +- Clean at `1a850ec17`. The Option A (scale removed) + LUT 2D Vulkan-fix + changes were reverted, and the CUDA-only guard commit is in. Nothing + uncommitted. + +### DEFERRED / NOT DONE +- Step 3f: auto-S via native fit (`--expert-hot-s -1`). NOT DONE in repo2. +- Hysteresis ratio gate + dwell (Trick 5/6): deferred in repo2. +- mmid.cu count+rank CUDA port: ATTEMPTED, CORRUPT, REVERTED. Do not re-approach without a new decision. +- Per-backend gate (option B): let Vulkan/Metal enable tier at n_tokens>1 - irrelevant now (gate is by design after the mmid failure). +- Fused MoE cold (GELU/gate_up): not ported. + +## 9. KNOWN LIMITATIONS + THE OPEN VULKAN ISSUE + +### 9.1 The 2C gate (n_tokens==1) - MULTI-SLOT DROPPED FOR NOW +Tier only engages at single-token decode. Multi-slot server batches fall back +to stock. This is now BY DESIGN after the mmid port failed (see 6.K/6.L). +DECISION (2026-08-05): multi-slot (n_tokens>1) tiered support is DROPPED FOR +NOW. The mmid.cu count+rank port remains reverted; do not re-approach without +an explicit new decision. -np>1 server batches use stock CPU MoE speed. +Single-stream CUDA is the supported tier path. + +### 9.2 ~~OPEN: repo2 tier corrupts on Vulkan~~ RESOLVED BY DROPPING VULKAN (2026-08-05) +Historical record only. The Vulkan tier path is DROPPED and the CUDA-only +guard (1a850ec17) prevents accidental use. Facts below are preserved for +reference; see section 12 for the eliminated theories and the sole remaining +lead. FACTS (all verified 2026-08-04): +- repo2 tier + Vulkan (RTX) + IQ2_M or Q5_K_P => `` loop / garbage. +- repo2 tier + CUDA + same models => fully coherent. +- v3.bak tier + Vulkan + same models => coherent (IQ2 18.65, Q5 5.29 tok/s). +- v5 tier + Vulkan + same models => coherent (IQ2 22.35, Q5 warm 11.72). +- repo2 stock (tier disabled) + Vulkan => coherent. +- repo2 tier + Vulkan + `-ngl 0` => still `` loop (rules out GPU->CPU + cur sync for the cold op). +- repo2 tier + Vulkan + S=1 => coherent (but ends early, not a full test). +- v3.bak tier + Vulkan + `-ngl 0` => coherent (the decisive A/B: same model, + same flags, same Vulkan backend, v3 works / repo2 loops). + +ROOT CAUSE STATUS: two candidate fixes applied but NOT sufficient: +1. Scale removed (Option A) - fixed binary garbage, still loops. +2. LUT native 2D - still loops. +Both were structurally necessary to match v3.bak but the corruption persists. + +NEXT LEAD (2026-08-04, highest priority to investigate): +**cold_mask buffer placement.** v3.bak allocates the cold `mask` on a CPU +buffer (`ggml_new_tensor_1d(g_ctx_cpu, I32, n_expert)`). repo2 allocates +`cold_mask` in the SAME ggml context as dst_hot/hot_lut, which is allocated +with the GPU buft - so `cold_mask` lives on the GPU buffer. The CPU cold op +reads `mask->data` DIRECTLY (`const int32_t * cold_mask = (const int32_t *) +mask->data;`). On Vulkan, a GPU-buffer tensor's `->data` is NOT host-readable +=> garbage mask => wrong skip/compute decisions => degenerate output. On CUDA +it apparently works (host-mapped or scheduler copies). THE FIX would be to +allocate cold_mask on a CPU/host buffer (like v3.bak), or read the mask via a +host-accessible path. VERIFY THIS FIRST in the next session. +Secondary candidates if the mask fix is not sufficient: +- The `ggml_mul_mat_id` hot call at n_tokens==1 uses the Vulkan VEC-id path + (`mul_mat_vec_id`); upstream `ggml_vk_get_dequantize_mul_mat_vec_id` omits + GGML_TYPE_IQ2_M from its supported list (has IQ2_XXS/XS/S but not IQ2_M). + v3.bak runs IQ2_M on the same backend, so likely not the blocker, but worth + confirming repo2 does not take a different dispatch. +- `ggml_mul_mat_id_cold` signature differs: v3 takes `(mask, ptrs)`, repo2 + takes `(cold_mask)` only. The ptrs arg is a RAM-pool optimization, likely + irrelevant to correctness. +- The `dst_hot` buffer content after `ggml_backend_tensor_set` copies on + Vulkan - verify by dumping dst_hot vs source on both backends. + +### 9.3 Hotstore segfault without -cmoe +repo2 tier requires the MoE expert weights to stay on CPU (needs `-cmoe` or +`-ncmoe`). Running `--expert-hot-s 96` WITHOUT -cmoe segfaulted on Vulkan +(2026-08-04): the hotstore copy reads `w->data` as a host pointer, but +without -cmoe the experts are GPU-offloaded. This is a latent bug - the tier +should reject or handle non-CPU experts. Not the Vulkan corruption (CUDA needs +-cmoe too and works). + +### 9.4 IQ2_M missing from Vulkan mul_mat_vec_id supported list +Upstream `ggml_vk_get_dequantize_mul_mat_vec_id` lists IQ2_XXS/XS/S but not +IQ2_M. Not the primary suspect (v3.bak runs IQ2_M cleanly on the same backend). + +### 9.5 Server -np>1 = stock speed +Multi-slot decode batches (n_tokens>1) fall back to stock CPU MoE via the +gate - ~20 tok/s aggregate regardless of S. Single-stream is where the +40 tok/s lives. This is by design after the mmid failure. + +### 9.6 Hot store competes for VRAM +8GB VRAM + 11GB IQ2_M model + S=96 hotstore barely coexist. Q5_K_P needs S<=48 +(96 slots would be 9424 MiB, exceeds 8GB). Larger GPU = more headroom. + +### 9.7 Stale comments / minor +- `ggml.h` cold_mask comment calls it "i32" but the tensor is f32 (read as int32). +- `ggml-rpc.h:14` static_assert 101 vs ggml.c 102 (harmless, RPC off). +- `llama-expert-tier.h` header comment describes the OLD dual-mul_mat_id design + (cold_lut/hot_mask) - stale, code is correct. Low priority. +- v3.bak/v5 ggml-vulkan.cpp matches upstream - their Vulkan success is NOT a + backend patch; it is their tier code. + +## 10. Reference docs + +- `oldtricks.md` - 25 tricks from colibri/wackMall. In scope: Trick 13 + (WEIGHTS usage tag), Trick 18 (non-owning ggml contexts / no_alloc batch), + Trick 23 (create-then-allocate), Trick 7 Path A (MUL_MAT_ID_COLD, ported). + Others OUT OF SCOPE unless re-discussed. +- wackMall references: `Project3/llama-wackMall_v3.bak` (simple, works on + Vulkan) and `Project3/llama-wackMall_v5` (feature-rich, faster, works on + Vulkan). Their tier files are the reference for the Vulkan fix. +- Colibri mmid.cu count+rank diff (the "revisit later" reference): the change + to `mm_ids_helper` replacing `iex_used` (one hit) with `cnt` + inclusive + warp scan `rank` + `n_hit`, so every duplicate sentinel hit gets its own + compact row. `Project1/folder2/llama.cpp/ggml/src/ggml-cuda/mmid.cu`. + NOTE: attempted in repo2 2026-08-04 and produced corrupt output - see 6.K. +- `Project6/repo2/AGENTS.md` - llama.cpp contributor rules. + +## 11. Vulkan benchmark table (2026-08-04, all RTX 3070 unless noted) + +| Build | Model | Flags | Result | tok/s | +|-------|-------|-------|--------|-------| +| repo2 CUDA | IQ2 | -cmoe -ngl99 S96 | coherent | 26.2-41 | +| repo2 Vulkan | IQ2 | -cmoe -ngl99 S96 | GARBAGE-> loop | 14-16 | +| repo2 Vulkan | Q5 | -cmoe -ngl99 S48 | GARBAGE | 5.24 | +| repo2 Vulkan | IQ2 | -cmoe -ngl0 S96 | loop | 4.66 | +| repo2 Vulkan | IQ2 | -cmoe -ngl0 S1 | coherent (short) | 5.9 | +| v3.bak Vulkan | IQ2 | -cmoe -ngl99 autoS | coherent | 16.4-18.7 | +| v3.bak Vulkan | Q5 | -cmoe -ngl99 autoS=51 | coherent | 5.29 | +| v3.bak Vulkan | IQ2 | -ngl0 autoS | coherent | 4.35 | +| v5 Vulkan | IQ2 | -cmoe -ngl99 autoS | coherent | 13.96-22.35 | +| v5 Vulkan | Q5 | -cmoe -ngl99 autoS=51 warm | coherent | 11.72 | +| v5 Vulkan | Q5 | cold first run | coherent | 2.96 | +| v5 Vulkan | Q5 | decay.999 hyst1.3 no-cmoe c8000 n256 | coherent (slow, RAM pool 18Gi) | 3.92 | +| v5 Vulkan | IQ2 | AMD RX570 (=1) | coherent | 9.94 | + +Takeaways: v3/v5 coherent on Vulkan; repo2 not. v5 faster than v3 (IQ2 22 vs 19, +Q5 11.7 vs 5.3). v3/v5 autofit S; repo2 manual S. The v3/v5 `.tier` sidecar is +a warm seed (seed coverage 84.6% warm vs 0% cold). + +## 12. VULKAN DECISION - DROPPED (2026-08-05) + +The Vulkan tier attempt is DROPPED after a full session of investigation. +CUDA tier (40 tok/s, coherent) remains the supported path. This section +records the eliminated theories and remaining leads so future sessions do not +re-litigate. + +### ELIMINATED THEORIES (do not re-test without new evidence) + +1. **IQ2_M unsupported type theory - DEAD.** The model file is named + "i1-IQ2_M" but the ACTUAL expert tensor types are `iq2_s` (gate/up, 82MiB) + and `iq3_s` (down, 110MiB), verified from loader logs. GGML_TYPE_IQ2_M does + NOT exist in this tree's ggml.h enum (lines 400-430). IQ2_S/IQ3_S ARE in + both the Vulkan dmmv getter and supports_op lists. So the hot node runs the + vec-id path fine (no abort, exit 0). Type was never the problem. +2. **supports_op probe as a discriminator - DEAD.** CUDA's MUL_MAT_ID + supports_op also lacks whatever (its list ends at IQ2_XXS/XS/S, IQ3_*, IQ4_*, + BF16) yet CUDA works. A supports_op probe would disable the working CUDA + path too. +3. **Mask buffer placement / dtype - DEAD.** cold_mask being F32-on-GPU (repo2) + vs I32-on-CPU (v3) is irrelevant: the scheduler copies the mask to CPU + correctly (seen as CPU#leaf_1012#0 in dumps) and the F32 zero-check is + equivalent to I32. +4. **tensor_set / buffer_clear asyncness - DEAD.** Both ggml_vk_buffer_write_2d + and ggml_vk_buffer_memset are synchronous (host-visible memcpy or + transfer+fence-wait). No race. +5. **IQ2_M missing from Vulkan vec-id pipeline - DEAD.** Irrelevant since the + type is actually IQ2_S/IQ3_S, and v3/v5 also lack IQ2_M in that getter yet + (see #6) none of them are actually coherent on Vulkan anyway. +6. **"v3/v5 work on Vulkan" - FALSE PREMISE, DEAD.** Directly tested 2026-08-05: + v2 (Project1/llama-wackMall_v2, its OWN vulkan build), v3.bak, and v5 ALL + corrupt on Vulkan with the tiered IQ2 config (same `` loop) including + with `-ngl 99 -cmoe`. The manuallog's earlier "coherent on Vulkan" rows for + v3/v5 do not reproduce under these exact commands and were likely measured + under different conditions (Q5 model, or runs that terminated early). + Conclusion: there is NO known-good Vulkan tier implementation to mirror. +7. **Per-matrix cold+ADD on Vulkan vs fused cold - DEAD.** v2/v3 share + byte-identical build_mul_mat_id + begin/end_moe_cold + ggml_moe_cold fused + cold, and v2 corrupts while the "fix" claim for v3 is void (both corrupt). + The v2-vs-v3 code diff is exclusively the prefetch/predict/poolB feature, + which is inactive by default; replacing v2's 3 differing files with v3's + changed nothing (still corrupt). So the graph structure is NOT the cause. +8. **Scheduler "1.wgt" weights-rule difference - DEAD.** The rule is + identical across repo2/v3/v5. All assign the hot node to Vulkan. + +### REMAINING LEAD (unconfirmed, low priority since Vulkan is dropped) + +The tiered hot node (`mul_mat_id(dst_hot[S+1 planes], cur, ids_hot)` with +slot indices 0..S incl. the zeroed sentinel plane) runs on Vulkan's vec-id +path and corrupts on ALL implementations (repo2/v2/v3/v5). Stock Vulkan (tier +off, MoE on CPU) is coherent. So the bug is in how Vulkan's `mul_mat_id` +handles the slot-indexed S+1-plane hot tensor with the sentinel plane, NOT in +repo2's tier code. Candidate mechanisms, unverified: + - coopmat2 path on Ampere (RTX 3070) mishandling the slot tensor + - vec-id shader plane indexing with a zeroed sentinel plane + - row_ids/result-tile shared memory sizing for S+1 planes +If Vulkan support is ever re-attempted, start here with a minimal standalone +repro (mul_mat_id on a 3-plane slot tensor vs the full tensor) before touching +the tier. + +### ACTIONS TAKEN +- Reverted uncommitted Vulkan fix attempts (Option A scale removal + LUT 2D) + in tier.cpp/hotstore.cpp -> working tree back to clean committed state + (5c5020d80). CUDA path untouched and verified. +- Added CUDA-only guard in llama-context.cpp (2026-08-05): the hotstore now + only allocates into a backend whose device name starts with "CUDA". On Vulkan + (bugged) it logs a WARN and skips -> tier off -> coherent stock output. + On CPU the existing non-CPU check already skips. LLAMA_EXPERT_HOT_FORCE=1 + overrides the guard (re-enables the tier on any non-CUDA GPU for testing/ + emergency only). Verified: Vulkan skips+coherent, CUDA enables+coherent + (re-sync swapped 2471), Vulkan+HOT_FORCE engages tier. COMMITTED as + `1a850ec17`. HEAD is now 1a850ec17 (2 ahead of origin, nothing uncommitted). + +## 13. Next actions (for the next session) + +State as of 2026-08-05: ALL WORK COMMITTED. Working tree clean, 5 commits +ahead of origin. Feature set complete for this pass: +- Auto-S via fit (`--expert-hot-s -1`) - DONE (f201ef52a) +- Hysteresis gate (--expert-hyst, default 1.3) - DONE (ca624bc06) +- decay default 0.999 - DONE (ca624bc06) +- --expert-dwell exists (default 0 = off) - DONE but OFF by default + +Committed this session (oldest->newest): +- `1a850ec17` llama : gate expert hot store to CUDA only, with force override +- `b18c4266c` common: expert hot store manual slots activate --cmoe +- `f201ef52a` common: autofit expert hot store slots via --expert-hot-s -1 +- `ca624bc06` expert hot store: hysteresis gate for slot swaps + +Verified on Qwen3.6-35B-A3B IQ2_M (255-run decode): stock autofit 33.25 +tok/s, autofit tier (S=133) 41 tok/s, hyst gate on 34.6-35.1 tok/s +(hyst=1.3 dwell=0: 35.1, 2600 swaps vs 2985 no-gate). Dwell default 0 +(user decision: not worth the speed cost). Dwell aging counts real tokens +and initial fill is eligible (else first sync defers -> speed crash). + +## 14. CPU OVERHEAD + CORRUPTION INVESTIGATION (2026-08-05) - CLEANED REFERENCE + +This section is the authoritative record. It is CHRONOLOGICAL and marks +superseded conclusions explicitly. The final truth (14.7) is the current +state. All measurements on CUDA build, Qwen3.6-35B-A3B IQ2_M, -c 8000, +--temp 0, --fit-target as noted, unless stated. Corruption judgements are +DIRECT READS of full output (see AGENTS.md rule 10), never grep counts. + +### 14.1 Phase A - the "5-10% CPU overhead" (t6) is CLOSED as a non-issue + +The user-reported +5-10% CPU vs stock was investigated thoroughly and +proved to be a 256-token short-run artifact of the tier's initial +convergence burst, not a real overhead. + +Key measurements (n=256 burst vs n=1024/4096 real): +- -t 6, n=256: ours 178.6 ms CPU/token vs stock 160.6 = +11%. Real per + token, but ONLY during the first ~100-token convergence burst. +- -t 6, n=4096 (VALID, fresh hotset): dynamic 60.87 tok/s / 98.1 ms + CPU/token vs static 58.45 / 101.3. Dynamic is FASTER and uses LESS CPU + than a frozen set. Both crush stock (34 tok/s, 160.6 ms/token). +- Verdict: dynamic tier at real lengths = ~1.8x faster and ~39% less CPU + per token than stock. CLOSED. Benchmark at -n 1024+ only. + +Superseded intermediate claims (for history only): 14.9 "ROOT CAUSE = +heatmap bookkeeping" was WRONG (conflated readback+resync+fill quality; +controlled test showed readback only ~1.8 ms/token); 14.11 "swap moves +~20ms/token" was the 256-burst artifact; 14.7 PASSIVE / 14.8 GGML_OPENMP=OFF +were dead ends (spin is load-bearing; native pool identical cost). + +### 14.2 Phase B - the -t 12 collapse is REAL and structural + +The real remaining issue is -t 12 (ours 11-19 tok/s vs stock 25). Isolated: +- Not heatmap (static also collapses), not swaps (0 in static), not + convergence (n=1024), not cold-gpu. It is the tier graph at 12 threads. +- Energy (perf power/energy-pkg, t12, n=1024): stock 3604 J / 3.48 J/tok + vs ours 5435 J / 5.25 J/tok. Ours does 51% MORE real work, not spin. +- Root cause: repo2 emits 3 separate MUL_MAT_ID_COLD CPU nodes per MoE + layer (gate, up, down) + 3 mul + 1 add = ~120 CPU nodes/token, each with + a full ggml_barrier (ggml-cpu.c:3116). v3/v5 use ONE fused GGML_OP_MOE_COLD + node per layer computing the whole chain under ONE barrier (~40 + nodes/token, 3x fewer barriers). + +### 14.3 Phase C - corruption discovered, mechanism LOCALIZED + +While tuning repo2's swap cadence for the t12 fix, output CORRUPTION was +found (mid-word splits: "Here 's", "summar izing", "Back ups", "/ usr", +"x 8 6 _ 6 4"). This VOIDED the tuning speed gains (they were measured on +corrupt output). + +Localization (direct read of interleaved stdout): a resync marker lands +BETWEEN the halves of a generated word: + ...using a **De=== re-sync swapped X ===bian-based** system... +The model emitted "De", a resync fired, then "bian" - the expert's weights +changed between two tokens of one word. Cadence sensitivity: cadence-1 = +pervasive, cadence-2 = intermittent, cadence-20 = rare (1/~700 words), +cadence-100 = unobserved. Resyncs cluster in the initial convergence burst. + +GRAPHLOG experiment (ggml_cuda_graph_compute instrumentation): the corrupt +token is produced in DIRECT mode (no CUDA graph active yet) exactly at a +resync marker. This RULED OUT the CUDA-graph-replay theory. + +### 14.4 Phase D - candidate fixes (ALL SUPERSEDED by 14.7) + +Multiple fixes reduced corruption frequency but NONE reached zero: +1. host-mask (cold_mask on host/CPU buffer) - closed a cross-stream race + on the mask write but corruption persisted. +2. post-resync ggml_backend_sched_synchronize - orders LUT copies but not + the logical hot/cold flip. +3. active-expert guard (don't evict experts routed this token) - helped + only partially (protects THIS token, not the word's prefix). +4. PP-priming (fill at first decode) - reasonable perf idea, not a fix. +5. rate-limit bursts (suppress swaps N tokens after a big resync) - capped + the burst but residual corruption remained. +6. word-boundary gating (Enhancement B: only resync when the just-sampled + token is ENTIRELY whitespace/punct, no cap) - reduced to rare, NOT zero. +7. atomic-pair LUT+mask write ordering - helped, not zero. +8. CUDA set_tensor device-wide sync (option 1) - helped, not zero. + +None achieved the required ZERO. All are superseded by 14.7. + +### 14.5 Phase E - static store is 100% clean (the control) + +LLAMA_EXPERT_STATIC_FILE (planted fixed set, heatmap+resync OFF, 0 swaps): +t6, n=1024, 30.36 tok/s, 755 words DIRECT READ = COMPLETELY CLEAN. Zero +mid-word splits. PROVES: zero swaps = zero corruption. The swap itself is +the sole corruption source. (Autofit S is nondeterministic and can OOM +graph capture; pin S manually when testing.) + +### 14.6 Phase F - v3 aggressive-repin: fused op is STRUCTURALLY immune + +To discriminate "v3 clean because swaps are RARE" (frequency) vs "v3 clean +because fused op is EXACT" (structural): forced v3's repin gate to be +aggressive (dwell 0, ratio 0.3 instead of 32/1.5). Repins = 10,280 +(massive). t6, n=1024, cold start, 57.56 tok/s. DIRECT READ of 656 words: +COMPLETELY CLEAN. "iptables", "systemd-resolved", "/etc/netplan/ +01-netcfg.yaml", "SSH: 22" all correct. ZERO artifacts at ~10 swaps/token. + +CONCLUSION: frequency theory WRONG. The fused single-node MOE_COLD op is +structurally immune to the mid-word corruption regardless of swap rate. +repo2's 3-node separate path is the corruption source. + +### 14.7 FINAL TRUTH (current state) - port the fused op + +The fix is to PORT v3/v5's fused GGML_OP_MOE_COLD op into repo2: +- One graph node computes the entire cold MoE chain (gate->silu/gelu-> + up->down) under ONE threadpool barrier, with a fixed intermediate + quantization contract (act_q) that matches what the hot path produces, + so hot<->cold flips are numerically identical (no divergence to corrupt). +- Repo2's 3-node separate MUL_MAT_ID_COLD path is the REGRESSION. +- Porting it fixes BOTH the -t 12 collapse (3x fewer barriers) AND the + corruption (structural immunity). Genuine cadence-1/2 becomes possible. +- v3 tree reverted to original gate after the experiment (git diff clean). + +FILES for the port (see AGENTS.md "Files you may touch" + approved +llama-graph.cpp): +- ggml/include/ggml.h: GGML_OP_MOE_COLD enum + ggml_moe_cold() decl +- ggml/src/ggml.c: op name, constructor, shape inference +- ggml/src/ggml-cpu/ggml-cpu.c: ggml_compute_forward_moe_cold kernel + + dispatch registration (~250 lines, the core) +- src/llama-expert-tier.h/.cpp: begin_moe_cold / end_moe_cold + g_hot_only +- src/llama-graph.cpp (APPROVED): cold_ok gate + begin/end calls in + build_moe_ffn (v5 lines 1974/2111) +- src/llama-context.cpp: wire the fused update path + +## 15. Current state and next actions (for the next session) + +STATE (2026-08-05): working tree has UNCOMMITTED experimental changes from +the corruption investigation (host-mask, word-boundary gating, rate-limit, +PP-priming, atomic-pair, active-guard, device-sync in ggml-cuda.cu, +graphlog instrumentation, diagnostic env hooks). These are all SUPERSEDED +by the fused-op port decision (14.7) and should be DISCARDED before +porting (git checkout -- the dirty files) to keep the port diff clean. + +COMMITTED (unpushed, keep): +- `5425deb6b` sentinel autofit S-1 fix +- `919168e1e` MSVC Interlocked fallback in cold op + +OPEN THREADS: +1. ACTIVE: port the fused GGML_OP_MOE_COLD op (14.7). Clean base first. +2. Decide whether to push the 2 committed commits (needs token). +3. v3 tree is untouched (reverted); v3/v5 remain reference trees for the + fused op implementation. +4. SYCL backend port of the tier exists (user's agent); the CUDA-only + guard (strncmp "CUDA") is more restrictive than the backend-agnostic + code requires - consider an allowlist if shipping multi-backend. diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 24f05cc91673..af8abf5a8f4d 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -17,6 +17,11 @@ add_library(llama llama-chat.cpp llama-context.cpp llama-cparams.cpp + llama-expert-heatmap.cpp + llama-expert-hotstore.cpp + llama-expert-pin.cpp + llama-expert-preload.cpp + llama-expert-tier.cpp llama-grammar.cpp llama-graph.cpp llama-hparams.cpp diff --git a/src/llama-context.cpp b/src/llama-context.cpp index 19cca7df1e9d..5b70e1165442 100644 --- a/src/llama-context.cpp +++ b/src/llama-context.cpp @@ -1,5 +1,8 @@ #include "llama-context.h" +#include "llama-expert-pin.h" +#include "llama-expert-preload.h" + #include "ggml.h" #include "llama-arch.h" #include "llama-graph.h" @@ -23,8 +26,7 @@ // llama_context // -static llm_graph_type ctx_type_to_graph_type(llama_context_type ctx_type) { - switch (ctx_type) { +static llm_graph_type ctx_type_to_graph_type(llama_context_type ctx_type) { switch (ctx_type) { case LLAMA_CONTEXT_TYPE_DEFAULT: return LLM_GRAPH_TYPE_DEFAULT; case LLAMA_CONTEXT_TYPE_MTP : return LLM_GRAPH_TYPE_DECODER_MTP; } @@ -463,6 +465,59 @@ llama_context::llama_context( } } + // heatmap exists for the tier, the heat log, or standalone pinning + // (LLAMA_EXPERT_PIN without the hot store) + llama_expert_pin::set_pct(params.expert_pin_pct); + if (hparams.n_expert > 0 && !cparams.warmup && + (params.expert_heat_log_period != 0 || params.expert_hot_s != 0 || + llama_expert_pin::active())) { + expert_heatmap = std::make_unique( + hparams.n_layer(), hparams.n_expert, + params.expert_heat_decay, + params.expert_heat_log_period, + params.expert_hot_s); + } + + // expert sidecar: restore the heatmap from .tier if present; the + // store starts pre-warmed and keeps adapting from there + expert_sidecar_enabled = params.expert_sidecar; + expert_sidecar_path = std::string(params.model_path ? params.model_path : "") + ".tier"; + if (expert_heatmap && params.expert_sidecar) { + expert_heatmap->load(expert_sidecar_path.c_str()); + } + + if (hparams.n_expert > 0 && !cparams.warmup && params.expert_hot_s != 0) { + const int sync_period = params.expert_sync_period; + expert_hotstore = std::make_unique( + &model, hparams.n_layer(), hparams.n_expert, + params.expert_hot_s, sync_period, + params.expert_hyst, params.expert_dwell, params.expert_move_mode); + // enable the GPU hot store on any GPU backend (CUDA, Vulkan, ROCm, + // SYCL, Metal, ...). + bool cache_enabled = false; + std::vector gpu_bufts; + for (auto & backend : backends) { + ggml_backend_dev_t dev = ggml_backend_get_device(backend.get()); + const enum ggml_backend_dev_type type = ggml_backend_dev_type(dev); + if (type == GGML_BACKEND_DEVICE_TYPE_CPU || type == GGML_BACKEND_DEVICE_TYPE_ACCEL) { + continue; + } + if (params.expert_gpu >= 0 && (int) gpu_bufts.size() != params.expert_gpu) { + continue; // store pinned to a specific GPU: skip the others + } + gpu_bufts.push_back(ggml_backend_get_default_buffer_type(backend.get())); + } + if (!gpu_bufts.empty()) { + cache_enabled = expert_hotstore->allocate(gpu_bufts, model.tensor_split(), (int) gpu_bufts.size()); + } + // launch hint: cache did not engage, usually no GPU accelerator + if (!cache_enabled) { + LLAMA_LOG_WARN("%s: expert cache is OFF: %d slots requested but no GPU backend in use\n", + __func__, params.expert_hot_s); + } + expert_hotstore->log(); + } + // Initialize the full vocabulary token ids for backend samplers. { const int n_vocab = model.vocab.n_tokens(); @@ -475,6 +530,10 @@ llama_context::llama_context( } llama_context::~llama_context() { + if (expert_heatmap && expert_sidecar_enabled && !cparams.warmup) { + const std::string sc = expert_sidecar_path; + expert_heatmap->save(sc.c_str()); + } // wait for any pending asynchronous copies into the output buffers before they are freed synchronize(); @@ -1378,6 +1437,17 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll //LLAMA_LOG_INFO("graph set inputs time: %.3f ms\n", (ggml_time_us() - t_start_us)/1000.0); } + // the store fill is deferred to the first token: copy the startup batch + // from the gguf file (sequentially, hash-verified) so the graph never + // computes against an unverified or load-time-corrupted store. + if (expert_heatmap && expert_hotstore && !expert_hotstore->is_filled) { + expert_hotstore->copy_top_s(*expert_heatmap); + } + + if (expert_hotstore) { + expert_hotstore->reset_counts(); // zero the cold-op tallies for this token + } + const auto status = graph_compute(res->get_gf(), ubatch.n_tokens > 1); if (status != GGML_STATUS_SUCCESS) { LLAMA_LOG_ERROR("%s: failed to compute graph, compute status: %d\n", __func__, status); @@ -1385,6 +1455,46 @@ llm_graph_result * llama_context::process_ubatch(const llama_ubatch & ubatch, ll return nullptr; } + // cold-op counts feed the heatmap directly (no D2H readback, no sync). + // decode only: prefill routing is uniform and would dilute the ranking. + if (expert_heatmap && expert_hotstore && expert_hotstore->is_filled && ubatch.n_tokens == 1) { + expert_hotstore->read_counts(*expert_heatmap, ubatch.n_tokens); + } + if (expert_heatmap && expert_hotstore && expert_hotstore->is_filled) { + expert_hotstore->maybe_resync(*expert_heatmap, ubatch.n_tokens > 1); + if (ubatch.n_tokens == 1 && getenv("LLAMA_EXPERT_HITRATE")) { + expert_hotstore->log_hit_rate(res->moe_sel_experts); + } + } + + // standalone pin mode (no hot store): feed the heatmap from the graph + // readback. decode only (see above). + if (expert_heatmap && !expert_hotstore && llama_expert_pin::active() && ubatch.n_tokens == 1) { + expert_heatmap->tick(1); + synchronize(); + for (const auto & [il, tensor] : res->moe_sel_experts) { + if (!tensor || !tensor->data) { + continue; + } + const int n_ids = (int) tensor->ne[0]; + const int n_tokens = (int) tensor->ne[1]; + std::vector ids((size_t) n_ids * n_tokens); + ggml_backend_tensor_get(tensor, ids.data(), 0, ids.size() * sizeof(int32_t)); + expert_heatmap->update_ids(il, ids.data(), n_ids, n_tokens); + } + } + + // mmap page hints for the expert tier (hot store active: GPU-aware sets; + // standalone: everything is cold, warm the top fraction) + if (expert_heatmap && llama_expert_pin::active()) { + llama_expert_hotstore * hs = expert_hotstore.get(); + llama_expert_pin::maybe_run(&model, expert_heatmap.get(), + [](void * ud, int il, int e) -> bool { + return ud != nullptr && static_cast(ud)->is_resident(il, e); + }, + hs); + } + ret = GGML_STATUS_SUCCESS; return res; @@ -2366,6 +2476,19 @@ uint32_t llama_context::graph_max_nodes(uint32_t n_tokens) const { return res; } +void llama_context::print_expert_heatmap() const { + // --expert-heat-log-period: print at generation end (0 = off) + if (expert_heatmap && expert_heatmap->log_period != 0) { + expert_heatmap->log(); + } +} + +void llama_print_expert_heatmap(const struct llama_context * ctx) { + if (ctx) { + ctx->print_expert_heatmap(); + } +} + llm_graph_result * llama_context::get_gf_res_reserve() const { return static_cast(gf_res_reserve.get()); } @@ -3517,6 +3640,17 @@ llama_context_params llama_context_default_params() { /*.kv_unified =*/ false, /*.sampler =*/ nullptr, /*.n_sampler =*/ 0, + /*.model_path =*/ nullptr, + /*.expert_heat_decay =*/ 0.999f, + /*.expert_heat_log_period =*/ 0, + /*.expert_hot_s =*/ 0, + /*.expert_sync_period =*/ 1, + /*.expert_hyst =*/ 1.3f, + /*.expert_dwell =*/ 0, + /*.expert_pin_pct =*/ -1, + /*.expert_move_mode =*/ 0, + /*.expert_sidecar =*/ false, + /*.expert_gpu =*/ -1, /*.ctx_other =*/ nullptr, }; diff --git a/src/llama-context.h b/src/llama-context.h index bf91daa8b562..94ae88018d82 100644 --- a/src/llama-context.h +++ b/src/llama-context.h @@ -7,6 +7,8 @@ #include "llama-adapter.h" #include "llama-impl.h" #include "llama-memory.h" +#include "llama-expert-heatmap.h" +#include "llama-expert-hotstore.h" #include "ggml-cpp.h" #include "ggml-opt.h" @@ -241,6 +243,9 @@ struct llama_context { public: uint32_t graph_max_nodes(uint32_t n_tokens) const; + // print the expert heatmap (no-op when not active) + void print_expert_heatmap() const; + // can reuse the llm_graph_result instance of the context (for example to update a memory module) llm_graph_result * get_gf_res_reserve() const; @@ -367,6 +372,13 @@ struct llama_context { llm_graph_result_ptr gf_res_prev; llm_graph_result_ptr gf_res_reserve; + std::unique_ptr expert_heatmap; + std::unique_ptr expert_hotstore; + + // expert sidecar: .tier path and flag, kept for the destructor save + bool expert_sidecar_enabled = false; + std::string expert_sidecar_path; + // host buffer for the model output (logits and embeddings) ggml_backend_buffer_ptr buf_output; diff --git a/src/llama-expert-heatmap.cpp b/src/llama-expert-heatmap.cpp new file mode 100644 index 000000000000..899f5083f56a --- /dev/null +++ b/src/llama-expert-heatmap.cpp @@ -0,0 +1,184 @@ +#include "llama-expert-heatmap.h" +#include "llama-impl.h" + +#include +#include +#include +#include +#include + +llama_expert_heatmap::llama_expert_heatmap( + int n_layers, int n_experts, + float decay_rate, int log_period, int hot_s) : + n_layers(n_layers), + n_experts(n_experts), + hot_s(hot_s), + decay_rate(decay_rate), + log_period(log_period), + tokens_total(0), + generated_tokens_count(0), + heat(n_layers * n_experts, 0.0f), + last_reuse(n_layers * n_experts, 0) { +} + +void llama_expert_heatmap::update_counts(const std::vector & per_layer, int n_tokens) { + tick(n_tokens); + for (int il = 0; il < n_layers && il < (int) per_layer.size(); il++) { + const int32_t * cnt = per_layer[il]; + if (!cnt) { + continue; + } + float * layer_heat = &heat[il * n_experts]; + int64_t * layer_reuse = &last_reuse[il * n_experts]; + for (int e = 0; e < n_experts; e++) { + if (cnt[e] > 0) { + layer_heat[e] += (float) cnt[e]; + layer_reuse[e] = tokens_total; + } + } + } +} + +void llama_expert_heatmap::tick(int n_tokens) { + // batched decay: apply decay_rate^2 every two updates instead of + // decay_rate every update (identical long-run weighting, half the work) + ++update_counter; + if (update_counter % 2 == 0) { + const float rate = decay_rate * decay_rate; + for (int i = 0; i < n_layers * n_experts; i++) { + heat[i] *= rate; + } + } + if (n_tokens == 1) { + generated_tokens_count++; + } + tokens_total += n_tokens; +} + +void llama_expert_heatmap::update_ids(int layer_idx, const int32_t * expert_ids, int n_ids, int n_tokens) { + if (layer_idx < 0 || layer_idx >= n_layers || !expert_ids) { + return; + } + float * layer_heat = &heat[layer_idx * n_experts]; + int64_t * layer_reuse = &last_reuse[layer_idx * n_experts]; + for (int t = 0; t < n_tokens; t++) { + for (int e = 0; e < n_ids; e++) { + const int32_t id = expert_ids[t * n_ids + e]; + if (id >= 0 && id < n_experts) { + layer_heat[id] += 1.0f; + layer_reuse[id] = tokens_total; + } + } + } +} + +void llama_expert_heatmap::log() const { + LLAMA_LOG("expert_heatmap: tokens %" PRId64 "\n", tokens_total); + + for (int l = 0; l < n_layers; l++) { + const float * layer_heat = heat.data() + l * n_experts; + int active_count = 0; + float max_heat = 0.0f; + int max_id = -1; + + for (int e = 0; e < n_experts; e++) { + if (layer_heat[e] > 0.01f) { + active_count++; + } + if (layer_heat[e] > max_heat) { + max_heat = layer_heat[e]; + max_id = e; + } + } + + if (active_count > 0) { + LLAMA_LOG(" layer %3d: %d warm experts, max heat=%.2f (expert %d)", + l, active_count, max_heat, max_id); + + auto top = get_top_s(l, 8); + LLAMA_LOG(" top-8="); + for (size_t i = 0; i < top.size(); i++) { + LLAMA_LOG("%s%d", i > 0 ? "," : "{", top[i]); + } + LLAMA_LOG("}\n"); + } + } +} + +float llama_expert_heatmap::get_score(int layer_idx, int expert_id) const { + if (layer_idx < 0 || layer_idx >= n_layers || expert_id < 0 || expert_id >= n_experts) { + return 0.0f; + } + float s = heat[layer_idx * n_experts + expert_id]; + if (generated_tokens_count <= 3) { + s *= 5.0f; // start-up boost: force the store toward the true hot set + } + return s; +} + +std::vector llama_expert_heatmap::get_top_s(int layer_idx, int s) const { + std::vector result; + if (layer_idx < 0 || layer_idx >= n_layers || s <= 0) { + return result; + } + + const float * layer_heat = heat.data() + layer_idx * n_experts; + + std::vector indices(n_experts); + for (int i = 0; i < n_experts; i++) { + indices[i] = i; + } + + int k = std::min(s, n_experts); + std::partial_sort(indices.begin(), indices.begin() + k, indices.end(), + [layer_heat](int a, int b) { + return layer_heat[a] > layer_heat[b]; + }); + + result.assign(indices.begin(), indices.begin() + k); + return result; +} + +bool llama_expert_heatmap::save(const char * path) const { + FILE * f = fopen(path, "wb"); + if (!f) { + return false; + } + const char magic[4] = {'H', 'M', 'S', 'D'}; + const int32_t nl = n_layers, ne = n_experts, uc = update_counter; + const int64_t tt = tokens_total, gc = generated_tokens_count; + bool ok = fwrite(magic, 1, 4, f) == 4 && + fwrite(&nl, sizeof(nl), 1, f) == 1 && + fwrite(&ne, sizeof(ne), 1, f) == 1 && + fwrite(&tt, sizeof(tt), 1, f) == 1 && + fwrite(&gc, sizeof(gc), 1, f) == 1 && + fwrite(&uc, sizeof(uc), 1, f) == 1 && + fwrite(heat.data(), sizeof(float), heat.size(), f) == heat.size(); + fclose(f); + return ok; +} + +bool llama_expert_heatmap::load(const char * path) { + FILE * f = fopen(path, "rb"); + if (!f) { + return false; + } + char magic[4] = {0}; + int32_t nl = 0, ne = 0, uc = 0; + int64_t tt = 0, gc = 0; + const bool ok = fread(magic, 1, 4, f) == 4 && memcmp(magic, "HMSD", 4) == 0 && + fread(&nl, sizeof(nl), 1, f) == 1 && nl == n_layers && + fread(&ne, sizeof(ne), 1, f) == 1 && ne == n_experts && + fread(&tt, sizeof(tt), 1, f) == 1 && + fread(&gc, sizeof(gc), 1, f) == 1 && + fread(&uc, sizeof(uc), 1, f) == 1 && + fread(heat.data(), sizeof(float), heat.size(), f) == heat.size(); + fclose(f); + if (!ok) { + return false; + } + tokens_total = tt; + generated_tokens_count = gc; + update_counter = uc; + return true; +} diff --git a/src/llama-expert-heatmap.h b/src/llama-expert-heatmap.h new file mode 100644 index 000000000000..c5937dd7c228 --- /dev/null +++ b/src/llama-expert-heatmap.h @@ -0,0 +1,44 @@ +#pragma once + +#include +#include + +struct ggml_tensor; + +struct llama_expert_heatmap { + int n_layers; + int n_experts; + int hot_s; + float decay_rate; + int log_period; + int64_t tokens_total; // real tokens seen (not multiplied by layers) + int64_t generated_tokens_count; // decode tokens seen; drives first-fill deferral + int update_counter = 1; // batched decay every 2 updates; start at 1 for early stability + std::vector heat; + std::vector last_reuse; // last token each expert was routed (dwell gate) + + llama_expert_heatmap(int n_layers, int n_experts, + float decay_rate = 0.99f, + int log_period = 100, + int hot_s = 0); + + // cold-op counts path: per-layer tallies of the selected experts (host + // memory). advances tokens_total; the first-3-token boost lives in get_score. + void update_counts(const std::vector & per_layer, int n_tokens); + + // standalone path (no cold op): per-expert increment from graph readback + void update_ids(int layer_idx, const int32_t * expert_ids, int n_ids, int n_tokens); + + // once per ubatch: decay + token counters + void tick(int n_tokens); + void log() const; + + float get_score(int layer_idx, int expert_id) const; + std::vector get_top_s(int layer_idx, int s) const; + + // persist the heat to a sidecar file (binary) and restore it on a later + // run. load() returns false when the file is missing or incompatible + // (different n_layers/n_experts). + bool save(const char * path) const; + bool load(const char * path); +}; diff --git a/src/llama-expert-hotstore.cpp b/src/llama-expert-hotstore.cpp new file mode 100644 index 000000000000..7ab47c465ed9 --- /dev/null +++ b/src/llama-expert-hotstore.cpp @@ -0,0 +1,1205 @@ +#include "llama-expert-hotstore.h" +#include "llama-expert-heatmap.h" +#include "llama-expert-preload.h" +#include "llama-expert-tier.h" +#include "llama-impl.h" +#include "llama-model.h" + +#include "ggml.h" +#include "ggml-backend.h" +#include "ggml-cpu.h" + +#include +#include +#include +#include +#include + +#if !defined(_WIN32) +#include +#endif + +#ifdef _WIN32 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#include +#endif + +// FNV-1a of the first min(n, 1024) bytes - matches the launch hash sampling +static uint64_t hash_slice_at(const uint8_t * p, size_t n) { + const size_t m = std::min(n, (size_t) 1024); + uint64_t h = 0xcbf29ce484222325ULL; + for (size_t j = 0; j < m; j++) { + h ^= p[j]; + h *= 1099511628211ULL; + } + return h; +} + +static bool verify_gpu_copy(ggml_tensor * dst, size_t slot_off, const ggml_tensor * src_tensor, int expert, const uint8_t * src_data, size_t len) { + // strict: hash the same first 1024 bytes of the GPU copy and compare to + // the authoritative data - the launch hash (no-mmap) or the source slice + // (mmap, where the model tensor is the file-backed ground truth). + const size_t n = std::min(len, (size_t) 1024); + std::vector buf(n); + ggml_backend_tensor_get(dst, buf.data(), slot_off, n); + uint64_t h = 0xcbf29ce484222325ULL; + for (size_t j = 0; j < n; j++) { + h ^= buf[j]; + h *= 1099511628211ULL; + } + uint64_t expected = llama_expert_preload::expected_hash(src_tensor, expert); + if (expected == 0 && src_data) { + expected = 0xcbf29ce484222325ULL; + for (size_t j = 0; j < n; j++) { + expected ^= src_data[j]; + expected *= 1099511628211ULL; + } + } + if (expected == 0) { + return false; // no ground truth - do not trust the copy + } + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "hotstore: verify expert %d expected=%016llx gpu=%016llx %s\n", + expert, (unsigned long long) expected, (unsigned long long) h, + h == expected ? "MATCH" : "MISMATCH"); + } + return h == expected; +} + +// matches the weight tensor of an expert tensor, e.g.: +// blk.0.ffn_gate_exps.weight +// blk.3.ffn_down_chexps.weight +// follows the same convention as LLM_FFN_EXPS_REGEX in common.h +static const std::regex g_re_exps_weight("blk\\.(\\d+)\\.ffn_(up|down|gate|gate_up)_(ch|)exps\\.weight"); + +llama_expert_hotstore::llama_expert_hotstore( + const llama_model * model, int n_layers, int n_experts, int hot_s, int sync_period, + float hyst, int dwell, int mode) : + n_layers(n_layers), + n_experts(n_experts), + hot_s(hot_s), + bytes_per_slot(n_layers, 0), + sync_period(sync_period), + hyst(hyst), + dwell(dwell), + mode(mode) { + if (n_layers <= 0) { + return; + } + if (this->hot_s > this->n_experts) { + LLAMA_LOG_WARN("%s: clamping expert hot store S=%d to n_experts=%d\n", __func__, this->hot_s, this->n_experts); + this->hot_s = this->n_experts; + } + + for (const auto & [name, tensor] : llama_internal_get_tensor_map(model)) { + std::smatch m; + if (std::regex_search(name, m, g_re_exps_weight)) { + const int il = std::stoi(m[1].str()); + if (il >= 0 && il < n_layers && tensor->ne[2] > 0) { + // a slot holds nbytes/n_experts of this tensor + bytes_per_slot[il] += ggml_nbytes(tensor) / (size_t) tensor->ne[2]; + entries.push_back({il, tensor, {}}); + } + } + } + + // entries is fixed from here on; build a per-layer index of stable + // pointers so copy/resync do not iterate the whole entries vector. + entries_by_layer.assign(n_layers, {}); + for (auto & e : entries) { + entries_by_layer[e.layer_idx].push_back(&e); + } + + // resolve mode: 1 = copy, 2 = move, 0 = auto. copy keeps the RAM copy of + // promoted experts (needs the full exps set resident); move frees it + // (RAM-tight). mmap exps are file-backed page cache (kernel-reclaimable), + // so copy is always right there; the RAM check only matters on no-mmap. + if (mode == 1) { + copy_mode = true; + } else if (mode == 2) { + copy_mode = false; + } else if (!copy_mode) { + const bool mmap_mode = !entries.empty() && llama_expert_preload::index_of(entries[0].src) < 0; + if (mmap_mode) { + copy_mode = true; + fprintf(stderr, "hotstore: copy mode (mmap: exps are file-backed page cache)\n"); + } else { + int64_t exps_bytes = 0; + for (int il = 0; il < n_layers; il++) { + exps_bytes += bytes_per_slot[il]; + } + exps_bytes *= n_experts; // full exps set: per-slot bytes x experts per layer + // available RAM: MemAvailable includes reclaimable page cache; + // _SC_AVPHYS_PAGES only counts truly-free pages and underreports + int64_t free_ram = 0; +#if defined(_WIN32) + MEMORYSTATUSEX ms; + memset(&ms, 0, sizeof(ms)); + ms.dwLength = sizeof(ms); + if (GlobalMemoryStatusEx(&ms)) { + free_ram = (int64_t) ms.ullAvailPhys; + } +#else +#if defined(__linux__) + FILE * f = fopen("/proc/meminfo", "r"); + if (f) { + char line[256]; + while (fgets(line, sizeof(line), f)) { + if (sscanf(line, "MemAvailable: %" PRId64 " kB", &free_ram) == 1) { + free_ram *= 1024; + break; + } + } + fclose(f); + } +#endif + if (free_ram <= 0) { + // fallback: total pages (Linux AVPHYS is the available count; + // other POSIX lack it, so use PHYS as a coarse bound) + const long page = sysconf(_SC_PAGESIZE); +#if defined(_SC_AVPHYS_PAGES) + free_ram = (int64_t) sysconf(_SC_AVPHYS_PAGES) * page; +#else + free_ram = (int64_t) sysconf(_SC_PHYS_PAGES) * page; +#endif + } +#endif + if (exps_bytes > 0 && free_ram > exps_bytes * 2) { + copy_mode = true; + fprintf(stderr, "hotstore: copy mode: exps (%d MiB) fit in RAM (%d MiB free)\n", + (int) (exps_bytes / (1024*1024)), (int) (free_ram / (1024*1024))); + } else if (exps_bytes > 0) { + fprintf(stderr, "hotstore: move mode: exps (%d MiB) need more RAM than free (%d MiB)\n", + (int) (exps_bytes / (1024*1024)), (int) (free_ram / (1024*1024))); + } + } + } + + if (this->hot_s > 0) { + slot_to_expert.assign(n_layers, std::vector(this->hot_s, -1)); + dwell_count.assign(n_layers, std::vector(this->hot_s, 0)); + gpu_routed.assign(n_layers, std::vector(this->n_experts, 0)); + pending_in.assign(n_layers, {}); + pending_out.assign(n_layers, {}); + // staging: one expert slot per exps tensor, per layer + cpu_staging_off.assign(n_layers, 0); + size_t off = 0; + for (int il = 0; il < n_layers; il++) { + cpu_staging_off[il] = off; + for (entry * e : entries_by_layer[il]) { + off += ggml_nbytes(e->src) / (size_t) e->src->ne[2]; + } + } + cpu_staging.resize(off, 0); + } +} + +bool llama_expert_hotstore::allocate( + const std::vector & bufts, + const float * tensor_split, int n_split) { + if (hot_s <= 0 || entries.empty()) { + return false; + } + if (hot_s > n_experts) { + throw std::runtime_error(format("%s: hot store S=%d exceeds n_experts=%d", + __func__, hot_s, n_experts)); + } + if (n_split <= 0 || (int) bufts.size() < n_split) { + n_split = (int) bufts.size() > 0 ? (int) bufts.size() : 1; + } + + n_devices = n_split; + slot_start.assign(n_devices, 0); + slot_end.assign(n_devices, 0); + + // per-device slot ranges from the tensor_split fractions (-ts); even split + // when the fractions are all zero + { + float total = 0.0f; + for (int g = 0; g < n_devices; g++) { + total += tensor_split ? tensor_split[g] : 0.0f; + } + if (total <= 0.0f) { + total = (float) n_devices; + } + int acc = 0; + for (int g = 0; g < n_devices; g++) { + slot_start[g] = acc; + const float frac = tensor_split && tensor_split[g] > 0.0f ? tensor_split[g] : 1.0f; + slot_end[g] = acc + (int) ((float) hot_s * frac / total); + acc = slot_end[g]; + } + slot_end[n_devices - 1] = hot_s; // last device absorbs the remainder + } + + // per-device no_alloc contexts holding that device's dst + hot_lut tensors + ctx_dev.resize(n_devices); + buf_dev.resize(n_devices); + luts.assign(n_layers, layer_lut{}); + for (int g = 0; g < n_devices; g++) { + const int local_slots = slot_end[g] - slot_start[g]; + + ggml_init_params p = { + /*.mem_size =*/ ggml_tensor_overhead() * (entries.size() + 2 * n_layers), + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ctx_dev[g] = ggml_context_ptr(ggml_init(p)); + if (!ctx_dev[g]) { + LLAMA_LOG_ERROR("%s: hot store: failed to create device %d context\n", __func__, g); + return false; + } + + // one hot tensor per expert weight tensor: local_slots slot planes + // plus a zeroed sentinel plane (index local_slots) + for (auto & e : entries) { + e.dst.resize(n_devices); + e.dst[g] = ggml_new_tensor_3d(ctx_dev[g].get(), e.src->type, e.src->ne[0], e.src->ne[1], local_slots + 1); + ggml_set_name(e.dst[g], (std::string(e.src->name) + ".hot").c_str()); + } + for (int il = 0; il < n_layers; il++) { + luts[il].hot_lut.resize(n_devices); + luts[il].hot_lut[g] = ggml_new_tensor_2d(ctx_dev[g].get(), GGML_TYPE_I32, 1, n_experts); + luts[il].mask_lut.resize(n_devices); + luts[il].mask_lut[g] = ggml_new_tensor_2d(ctx_dev[g].get(), GGML_TYPE_F32, 1, local_slots + 1); + } + + // adopt the loader-streamed store buffer if present (the startup batch + // is already resident and the sentinel planes are zeroed), else + // allocate a fresh buffer + ggml_backend_buffer_t pre = llama_expert_preload::take_buffer(); + if (pre) { + char * base = (char *) ggml_backend_buffer_get_base(pre); + for (auto & e : entries) { + const llama_expert_preload::entry * pe = nullptr; + for (size_t i = 0; i < llama_expert_preload::num_entries(); i++) { + const auto * cand = llama_expert_preload::entry_at(i); + if (cand->src == e.src) { + pe = cand; + break; + } + } + if (!pe) { + throw std::runtime_error(format("%s: preload entry missing for %s", __func__, e.src->name)); + } + ggml_backend_tensor_alloc(pre, e.dst[g], base + pe->gpu_offset); + if (getenv("LLAMA_EXPERT_DEBUG")) { + static int once = 0; + if (!once) { + once = 1; + std::vector plane(1024); + for (int ex = 0; ex < (int) e.dst[g]->ne[2]; ex++) { + ggml_backend_tensor_get(e.dst[g], plane.data(), (size_t) ex * pe->plane_bytes, 1024); + uint64_t h = 0xcbf29ce484222325ULL; + for (int i = 0; i < 1024; i++) { + h ^= plane[i]; + h *= 0x100000001b3ULL; + } + const uint64_t exp = ex < hot_s ? llama_expert_preload::expected_hash(e.src, ex) : 0; + fprintf(stderr, "hotstore: slot %d fnv=%016llx expected=%016llx %s\n", + ex, (unsigned long long) h, (unsigned long long) exp, + h == exp ? "MATCH" : "MISMATCH"); + } + } + } + } + size_t lut_off = llama_expert_preload::align_up256(llama_expert_preload::entries_size()); + for (int il = 0; il < n_layers; il++) { + lut_off = llama_expert_preload::align_up256(lut_off); + ggml_backend_tensor_alloc(pre, luts[il].hot_lut[g], base + lut_off); + lut_off += llama_expert_preload::align_up256((size_t) n_experts * sizeof(int32_t)); + lut_off = llama_expert_preload::align_up256(lut_off); + ggml_backend_tensor_alloc(pre, luts[il].mask_lut[g], base + lut_off); + lut_off += llama_expert_preload::align_up256((size_t) (local_slots + 1) * sizeof(float)); + } + buf_dev[g] = ggml_backend_buffer_ptr(pre); + ggml_backend_buffer_set_usage(buf_dev[g].get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + preloaded = true; + } else { + // check the buffer would fit before committing any VRAM + const size_t need = ggml_backend_alloc_ctx_tensors_from_buft_size(ctx_dev[g].get(), bufts[g]); + if (need == 0) { + LLAMA_LOG_ERROR("%s: hot store: zero-sized buffer on device %d, disabled\n", __func__, g); + return false; + } + size_t free_mem = 0, total_mem = 0; + ggml_backend_dev_t dev = ggml_backend_buft_get_device(bufts[g]); + if (dev) { + ggml_backend_dev_memory(dev, &free_mem, &total_mem); + } + if (dev && free_mem < need) { + throw std::runtime_error(format("%s: not enough memory to allocate the GPU hot store of %d slots (%zu MiB needed, %zu MiB free on %s)", + __func__, hot_s, need / (1024 * 1024), free_mem / (1024 * 1024), + ggml_backend_dev_name(dev))); + } + ggml_backend_buffer_t b = ggml_backend_alloc_ctx_tensors_from_buft(ctx_dev[g].get(), bufts[g]); + if (b == nullptr) { + throw std::runtime_error(format("%s: unable to allocate hot store buffer of %d slots (%zu MiB)", + __func__, hot_s, need / (1024 * 1024))); + } + buf_dev[g] = ggml_backend_buffer_ptr(b); + ggml_backend_buffer_set_usage(buf_dev[g].get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + ggml_backend_buffer_clear(buf_dev[g].get(), 0); + } + + // sentinel mask: 1.0 for real slots, 0.0 for the sentinel plane, so + // sentinel-routed hot rows are zeroed after the GPU mul_mat_id. + std::vector mask_h(local_slots + 1, 1.0f); + mask_h[local_slots] = 0.0f; + for (int il = 0; il < n_layers; il++) { + ggml_backend_tensor_set(luts[il].mask_lut[g], mask_h.data(), 0, + (local_slots + 1) * sizeof(float)); + if (getenv("LLAMA_EXPERT_DEBUG") && il == 0 && g == 0) { + static int once = 0; + if (!once) { + once = 1; + std::vector rb(local_slots + 1); + ggml_backend_tensor_get(luts[il].mask_lut[g], rb.data(), 0, + (local_slots + 1) * sizeof(float)); + fprintf(stderr, "hotstore: mask_lut[0..3]=%.1f %.1f %.1f %.1f [95]=%.1f [96]=%.1f\n", + rb[0], rb[1], rb[2], rb[3], rb[95], rb[96]); + } + } + } + } + + // CPU context for the cold_mask tensors + ggml_init_params params_cpu = { + /*.mem_size =*/ ggml_tensor_overhead() * (2 * n_layers) + 1024 * 1024, + /*.mem_buffer =*/ nullptr, + /*.no_alloc =*/ true, + }; + ctx_cpu = ggml_context_ptr(ggml_init(params_cpu)); + for (int il = 0; il < n_layers; il++) { + luts[il].cold_mask = ggml_new_tensor_1d(ctx_cpu.get(), GGML_TYPE_I32, n_experts); + luts[il].counts = ggml_new_tensor_1d(ctx_cpu.get(), GGML_TYPE_I32, n_experts + 1); + } + ggml_backend_buffer_type_t cpu_buft = ggml_backend_cpu_buffer_type(); + ggml_backend_buffer_t b_cpu = ggml_backend_alloc_ctx_tensors_from_buft(ctx_cpu.get(), cpu_buft); + if (b_cpu) { + buf_cpu = ggml_backend_buffer_ptr(b_cpu); + ggml_backend_buffer_set_usage(buf_cpu.get(), GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + } + + // the preloaded store carries the startup batch already; publish the + // matching LUTs now so the tier is correct from the very first token + // (copy_top_s later re-runs them, idempotently). + if (preloaded) { + for (int il = 0; il < n_layers; il++) { + for (int p = 0; p < hot_s; p++) { + slot_to_expert[il][p] = p; + } + } + update_luts(); + } + + + // register each expert weight tensor with the tier hook so build_lora_mm_id + // can find its per-device GPU hot tensors and per-device LUTs. + for (const auto & e : entries) { + const auto & L = luts[e.layer_idx]; + llama_expert_tier_register(e.src, e.dst, L.hot_lut, L.mask_lut, L.cold_mask, L.counts); + } + + return true; +} + +llama_expert_hotstore::~llama_expert_hotstore() { + llama_expert_tier_clear(); + llama_expert_preload::clear(); +} + +// device owning a global slot index, or -1 (slot ranges are contiguous) +static int slot_device(const std::vector & slot_start, const std::vector & slot_end, int p) { + for (int g = 0; g < (int) slot_start.size(); g++) { + if (p >= slot_start[g] && p < slot_end[g]) { + return g; + } + } + return -1; +} + +bool llama_expert_hotstore::copy_top_s(const llama_expert_heatmap & heatmap) { + if (is_filled || hot_s <= 0 || entries.empty() || buf_dev.empty()) { + return false; + } + + for (int il = 0; il < n_layers; il++) { + auto & ste = slot_to_expert[il]; + auto & dc = dwell_count[il]; + // startup batch: the first S experts of each layer go to the GPU + for (int p = 0; p < hot_s; p++) { + ste[p] = p; + dc[p] = dwell; // initial fill is eligible to be corrected next sync + } + + for (entry * e : entries_by_layer[il]) { + const size_t slot = ggml_nbytes(e->src) / (size_t) e->src->ne[2]; + if (preloaded) { + // the startup slices were streamed into the store at load + // (write_entry), so the GPU store is already sound. + continue; + } + const char * src = e->src->data ? (const char *) ggml_get_data(e->src) : nullptr; + if (!src) { + continue; + } + for (int p = 0; p < hot_s; p++) { + const int ex = ste[p]; + if (ex < 0) { + continue; + } + const int g = slot_device(slot_start, slot_end, p); + if (g < 0) { + continue; + } + ggml_backend_tensor_set(e->dst[g], src + (size_t) ex * slot, (size_t) (p - slot_start[g]) * slot, slot); + // page hints are handled by llama_expert_pin (mmap only); + // DONTNEED here would refault from disk on eviction + } + } + // startup batch is trusted on the GPU (filled from the file, hash-checked); + // the GPU output counts for it from the first token + for (int p = 0; p < hot_s; p++) { + if (ste[p] >= 0) { + gpu_routed[il][ste[p]] = 1; + } + } + } + + last_sync_tokens = heatmap.tokens_total; + is_filled = true; + update_luts(); + fprintf(stderr, "hotstore: startup batch moved to GPU\n"); + return true; +} + +// LLAMA_EXPERT_FULL_SYNC: direct swap (1 per layer per token). copy the +// new expert to a free or gate-cleared slot, verify the GPU copy, route +// immediately. no handshake pacing queues - the verify is the safety net. +bool llama_expert_hotstore::resync_full_mirror(const llama_expert_heatmap & heatmap, int budget) { + if (!is_filled || hot_s <= 0 || buf_dev.empty()) { + return false; + } + const int64_t elapsed = heatmap.tokens_total - last_sync_tokens; + int swapped = 0; + for (int il = 0; il < n_layers; il++) { + auto & ste = slot_to_expert[il]; + auto & dc = dwell_count[il]; + auto & rout = gpu_routed[il]; + std::vector resident_set(n_experts, 0); + for (int p = 0; p < hot_s; p++) { + if (ste[p] >= 0) { + resident_set[ste[p]] = 1; + } + } + const std::vector top = heatmap.get_top_s(il, hot_s); + + auto find_slot = [&](int e_cold) -> int { + for (int p = 0; p < hot_s; p++) { + if (ste[p] < 0) { + return p; + } + } + const float s_cold = heatmap.get_score(il, e_cold); + int p_worst = -1; + float worst_score = 1e9f; + for (int p = 0; p < hot_s; p++) { + if (ste[p] < 0) { + continue; + } + if (dc[p] < dwell) { + continue; + } + if (s_cold >= hyst * heatmap.get_score(il, ste[p])) { + const float s_inc = heatmap.get_score(il, ste[p]); + if (s_inc < worst_score) { + worst_score = s_inc; + p_worst = p; + } + } + } + return p_worst; + }; + + int swapped_in_layer = 0; + for (int e_cold : top) { + if (swapped_in_layer >= budget) { + break; + } + if (resident_set[e_cold]) { + continue; + } + const int p = find_slot(e_cold); + if (p < 0) { + break; + } + const int g = slot_device(slot_start, slot_end, p); + if (g < 0) { + continue; + } + const int e_out = ste[p]; + // evicted expert: copy its slices back out (copy-on-read), route it to the CPU + if (e_out >= 0) { + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx < 0) { + continue; + } + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off = (size_t) (p - slot_start[g]) * slot; + std::vector out(slot); + ggml_backend_tensor_get(ent->dst[g], out.data(), off, slot); + llama_expert_preload::set_cpu_slice(pidx, e_out, out.data()); + } + rout[e_out] = 0; + } + // new expert: copy from the host, route immediately (verify dropped) + bool ok = true; + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off = (size_t) (p - slot_start[g]) * slot; + const char * model_src = ent->src->data ? (const char *) ggml_get_data(ent->src) : nullptr; + const uint8_t * src = pidx >= 0 + ? llama_expert_preload::cpu_slice(pidx, e_cold) + : ((const uint8_t *) model_src) + (size_t) e_cold * slot; + if (!src) { + ok = false; + continue; + } + ggml_backend_tensor_set(ent->dst[g], src, off, slot); + } + if (ok) { + ste[p] = e_cold; + rout[e_cold] = 1; + dc[p] = -elapsed; // fresh dwell + swapped++; + swapped_in_layer++; + } + } + } + + last_sync_tokens = heatmap.tokens_total; + if (swapped > 0) { + update_luts(); + } + return swapped > 0; +} + +// byte offset of an exps tensor's single slot within its layer's staging region +static size_t staging_slot_off(const std::vector & layer, const llama_expert_hotstore::entry * ent) { + size_t off = 0; + for (const auto * e : layer) { + if (e == ent) { + break; + } + off += ggml_nbytes(e->src) / (size_t) e->src->ne[2]; + } + return off; +} + +bool llama_expert_hotstore::resync_top_s(const llama_expert_heatmap & heatmap) { + if (!is_filled || hot_s <= 0 || buf_dev.empty()) { + return false; + } + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "hotstore: resync tok=%lld\n", (long long) heatmap.tokens_total); + } + + // tokens elapsed since the previous sync, used to age dwell counters + const int64_t elapsed = heatmap.tokens_total - last_sync_tokens; + int changed = 0; + std::vector dirty(n_layers, 0); // layers whose LUTs need rebuilding + for (int il = 0; il < n_layers; il++) { + auto & ste = slot_to_expert[il]; + auto & dc = dwell_count[il]; + auto & rout = gpu_routed[il]; + auto & pin = pending_in[il]; + auto & pout = pending_out[il]; + auto & stg = cpu_staging; // global staging + const size_t stg_off = cpu_staging_off[il]; + + // ---- pending move-outs: verify the staging copy (all per token), + // then hand the expert to the CPU and free the slot ------------ + for (auto it = pout.begin(); it != pout.end();) { + if (!it->verified) { + if (getenv("LLAMA_EXPERT_NO_VERIFY")) { + it->verified = true; // safety dropped: trust the staging copy + } else { + bool ok = true; + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx < 0) { + continue; // mmap: file-backed, nothing to verify + } + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const uint64_t exp = llama_expert_preload::expected_hash(ent->src, it->expert); + const uint64_t got = hash_slice_at(&stg[stg_off + staging_slot_off(entries_by_layer[il], ent)], slot); + if (exp == 0 || exp != got) { + ok = false; + break; + } + } + if (ok) { + it->verified = true; + } else if (++it->failures >= 5) { + // corruption recovery: the GPU -> staging copy kept failing, + // so re-read the expert straight from the gguf file into the + // CPU slice (ground truth); only then route it cold + bool recovered = true; + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx < 0) { + continue; // mmap: data stays in the file + } + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + std::vector buf(slot); + if (!llama_expert_preload::read_expert(pidx, it->expert, buf.data(), slot)) { + recovered = false; + continue; + } + if (llama_expert_preload::expected_hash(ent->src, it->expert) != hash_slice_at(buf.data(), slot)) { + recovered = false; + } + llama_expert_preload::set_cpu_slice(pidx, it->expert, buf.data()); + } + if (recovered) { + rout[it->expert] = 0; + ste[it->p] = -1; + dirty[il] = 1; + changed++; + it = pout.erase(it); + continue; + } + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "hotstore: move-out ABORT expert %d after %d fails\n", it->expert, it->failures); + } + it = pout.erase(it); + continue; + } + } + } + if (it->verified && --it->countdown <= 0) { + // hand the expert to the CPU and free the slot + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx < 0) { + continue; // mmap: data stays in the file + } + llama_expert_preload::set_cpu_slice(pidx, it->expert, &stg[stg_off + staging_slot_off(entries_by_layer[il], ent)]); + } + rout[it->expert] = 0; // the CPU output counts now + ste[it->p] = -1; // the GPU slot is freed + dirty[il] = 1; + changed++; + it = pout.erase(it); + } else { + ++it; + } + } + + // ---- pending move-ins: verify ALL copies every token; a confirmed + // one is routed to the GPU on the following token --------------- + for (auto it = pin.begin(); it != pin.end();) { + if (it->verified) { + ++it; + continue; + } + if (getenv("LLAMA_EXPERT_NO_VERIFY")) { + it->verified = true; // safety dropped: trust the GPU copy + ++it; + continue; + } + const int g = slot_device(slot_start, slot_end, it->p); + bool ok = g >= 0; + if (g >= 0) { + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off = (size_t) (it->p - slot_start[g]) * slot; + const char * model_src = ent->src->data ? (const char *) ggml_get_data(ent->src) : nullptr; + const uint8_t * src = pidx >= 0 + ? llama_expert_preload::cpu_slice(pidx, it->expert) + : ((const uint8_t *) model_src) + (size_t) it->expert * slot; + if (!src || !verify_gpu_copy(ent->dst[g], off, ent->src, it->expert, src, slot)) { + ok = false; + } + } + } + if (ok) { + it->verified = true; + it->countdown = 1; // route to the GPU on the next token + ++it; + } else if (++it->failures >= 5) { + // corruption recovery: the CPU slice kept failing to verify, so + // re-read the expert straight from the gguf file (ground truth) + // and re-copy it; keep the move-in only if the disk copy verifies + bool recovered = true; + const int g2 = it->p >= 0 ? slot_device(slot_start, slot_end, it->p) : -1; + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + std::vector buf(slot); + if (pidx < 0 || g2 < 0 || !llama_expert_preload::read_expert(pidx, it->expert, buf.data(), slot)) { + recovered = false; + continue; + } + const size_t off = (size_t) (it->p - slot_start[g2]) * slot; + ggml_backend_tensor_set(ent->dst[g2], buf.data(), off, slot); + if (!verify_gpu_copy(ent->dst[g2], off, ent->src, it->expert, buf.data(), slot)) { + recovered = false; + } + } + if (recovered) { + it->verified = true; + it->countdown = 1; // route to the GPU on the next token + ++it; + } else { + ste[it->p] = -1; + it = pin.erase(it); + } + } else { + ++it; + } + } + for (auto it = pin.begin(); it != pin.end();) { + if (it->verified && --it->countdown <= 0) { + rout[it->expert] = 1; // the GPU output counts now + if (!copy_mode) { + // move: the expert now lives on the GPU, free its RAM copy + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx >= 0) { + llama_expert_preload::free_cpu_slice(pidx, it->expert); + } + } + } + dirty[il] = 1; + changed++; + it = pin.erase(it); + } else { + ++it; + } + } + + // ---- boundary gate: extract the actual bottom-GPU and top-CPU heats + // from the routed sets (the store lags the ideal ordering) and + // gate BOTH the eviction and the move-in on them. + float lowest_gpu = INFINITY, highest_cpu = -INFINITY; + for (int e = 0; e < n_experts; e++) { + const float s = heatmap.get_score(il, e); + if (rout[e]) { + lowest_gpu = std::min(lowest_gpu, s); + } else { + highest_cpu = std::max(highest_cpu, s); + } + } + if (highest_cpu > lowest_gpu) { + // ---- eviction selection (1/token): a GPU expert leaves only when + // a CPU expert beats it by the hysteresis margin and it has + // dwelled. gated on the eviction queue capacity. + int evict_candidate = -1; + if (!llama_expert_preload::get_no_evict() && !getenv("LLAMA_EXPERT_NO_EVICT") && (int) pout.size() < max_concurrent_moves) { + // never evict an expert still in the top-S: a swap must move a + // genuinely cold expert out, not shuffle two GPU-bound experts + const std::vector top = heatmap.get_top_s(il, hot_s); + std::vector in_top(n_experts, 0); + for (int te : top) { + in_top[te] = 1; + } + std::vector routed; + for (int e = 0; e < n_experts; e++) { + if (rout[e]) { + routed.push_back(e); + } + } + std::sort(routed.begin(), routed.end(), [&](int a, int b) { + return heatmap.get_score(il, a) < heatmap.get_score(il, b); + }); + for (int e : routed) { + if (in_top[e]) { + break; // monotonic: warmer residents are in the top-S too + } + int p = -1; + for (int pp = 0; pp < hot_s; pp++) { + if (ste[pp] == e) { + p = pp; + break; + } + } + if (p < 0 || dc[p] < dwell) { + continue; // not an applicant yet + } + if (highest_cpu >= hyst * heatmap.get_score(il, e)) { + evict_candidate = e; + } + break; // monotonic: hotter experts fail too + } + } + if (evict_candidate >= 0 && (int) pout.size() < max_concurrent_moves) { + // queue: at most 2 move-outs per layer at a time + const int p = [&]() { for (int pp = 0; pp < hot_s; pp++) if (ste[pp] == evict_candidate) return pp; return -1; }(); + const int dg = p >= 0 ? slot_device(slot_start, slot_end, p) : -1; + if (p >= 0 && dg >= 0) { + // copy the evicted expert's slices into the staging buffer + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + if (pidx < 0) { + continue; // mmap: file-backed, no staging needed + } + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off = (size_t) (p - slot_start[dg]) * slot; + ggml_backend_tensor_get(ent->dst[dg], &stg[stg_off + staging_slot_off(entries_by_layer[il], ent)], off, slot); + } + pout.push_back({evict_candidate, p, false, 0, 0}); + dirty[il] = 1; + changed++; // the LUTs do not change yet, but the transition started + } + } + } // boundary gate + + // ---- promote/demote (D2D): copy mode keeps RAM copies, so moves go + // through the host pool; D2D would double the bandwidth for nothing + for (int g = 0; g + 1 < n_devices && !copy_mode; g++) { + float worst_bound = INFINITY; + int worst_slot = -1; + for (int p = slot_start[g]; p < slot_end[g]; p++) { + if (ste[p] < 0) { + continue; + } + const float s = heatmap.get_score(il, ste[p]); + if (s < worst_bound) { + worst_bound = s; + worst_slot = p; + } + } + if (worst_slot < 0 || worst_bound <= 0.0f) { + continue; // dead resident (no heat yet); eviction will clear it + } + if (dc[worst_slot] < dwell + 2) { + continue; // not aged enough since its last change + } + int best_slot = -1; + float best_score = -INFINITY; + for (int p = slot_start[g+1]; p < slot_end[g+1]; p++) { + if (ste[p] < 0) { + continue; + } + const float s = heatmap.get_score(il, ste[p]); + if (s > best_score) { + best_score = s; + best_slot = p; + } + } + if (best_slot < 0 || best_score < hyst * worst_bound) { + continue; + } + const int e_demote = ste[worst_slot]; + const int e_promote = ste[best_slot]; + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + const size_t eslot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off_w = (size_t) (worst_slot - slot_start[g]) * eslot; + const size_t off_b = (size_t) (best_slot - slot_start[g+1]) * eslot; + const uint8_t * src_demote = pidx >= 0 + ? llama_expert_preload::cpu_slice(pidx, e_demote) + : nullptr; + const uint8_t * src_promote = pidx >= 0 + ? llama_expert_preload::cpu_slice(pidx, e_promote) + : nullptr; + if (src_demote && src_promote) { + // no-mmap: both experts are already host-resident (the + // cold pool); write straight into the store, no GPU reads + ggml_backend_tensor_set(ent->dst[g+1], src_demote, off_b, eslot); + ggml_backend_tensor_set(ent->dst[g], src_promote, off_w, eslot); + } else { + std::vector s1(eslot), s2(eslot); + ggml_backend_tensor_get(ent->dst[g], s1.data(), off_w, eslot); // demote + ggml_backend_tensor_get(ent->dst[g+1], s2.data(), off_b, eslot); // promote + ggml_backend_tensor_set(ent->dst[g+1], s1.data(), off_b, eslot); + ggml_backend_tensor_set(ent->dst[g], s2.data(), off_w, eslot); + } + } + ste[worst_slot] = e_promote; + ste[best_slot] = e_demote; + // dwell is in resyncs; the counter ages by elapsed tokens per resync + dc[worst_slot] = -elapsed * (dwell + 2); + dc[best_slot] = -elapsed * (dwell + 2); + dirty[il] = 1; + changed++; + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "hotstore: d2d swap promote=%d demote=%d (dev %d -> %d) promote_s=%.2f worst=%.2f hyst=%.2f\n", + e_promote, e_demote, g+1, g, best_score, worst_bound, hyst * worst_bound); + } + break; // one D2D swap per layer per tick + } + + // ---- move-in (1/token, gated): fill a free slot with the hottest CPU + // expert that should be on the GPU. check the free slot first (the + // common blocker), then the move-in queue (rarer), then evaluate. + int free_slot = -1; + for (int pp = 0; pp < hot_s; pp++) { + if (ste[pp] < 0) { + free_slot = pp; + break; + } + } + if (free_slot >= 0) { + if ((int) pin.size() < max_concurrent_moves) { // the queue is rarer than a missing slot + const std::vector top = heatmap.get_top_s(il, hot_s); + for (int e_cold : top) { + if (e_cold < 0 || e_cold >= n_experts || rout[e_cold]) { + continue; + } + const int dg = slot_device(slot_start, slot_end, free_slot); + if (dg < 0) { + break; + } + // copy the new expert in; keep it CPU-counted until verified + for (entry * ent : entries_by_layer[il]) { + const int pidx = llama_expert_preload::index_of(ent->src); + const size_t slot = ggml_nbytes(ent->src) / (size_t) ent->src->ne[2]; + const size_t off = (size_t) (free_slot - slot_start[dg]) * slot; + const char * model_src = ent->src->data ? (const char *) ggml_get_data(ent->src) : nullptr; + const uint8_t * src = pidx >= 0 + ? llama_expert_preload::cpu_slice(pidx, e_cold) + : ((const uint8_t *) model_src) + (size_t) e_cold * slot; + if (!src) { + continue; + } + ggml_backend_tensor_set(ent->dst[dg], src, off, slot); + } + ste[free_slot] = e_cold; + pin.push_back({e_cold, free_slot, false, 0, 0}); + dc[free_slot] = -elapsed; + dirty[il] = 1; + changed++; + break; // one move-in per layer per token + } + } + } + + for (int p = 0; p < hot_s; p++) { + if (ste[p] >= 0) { + dc[p] += (int) std::max(elapsed, 0); + } + } + } + + last_sync_tokens = heatmap.tokens_total; + if (changed > 0) { + update_luts(dirty); + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "hotstore: re-sync changed %d slots\n", changed); + } + } + return changed > 0; +} + +bool llama_expert_hotstore::maybe_resync(const llama_expert_heatmap & heatmap, bool multi_slot) { + // n_tokens>1 (multi-slot) freezes the hot store: no swapping during the batch + if (multi_slot || heatmap.tokens_total <= 0) { + return false; + } + // adaptive cadence: floor 10 (5 in copy mode: host-pool moves are half + // the bandwidth), stretching to 32 as the hit rate climbs toward 6x target + float ratio = hit_rate_valid ? hit_rate / target_hit_rate() : 0.0f; + const int floor = copy_mode ? 5 : 10; + const int eff_period = std::max(floor, std::min(32, floor + (int) (27.0f * (ratio - 1.0f) / 5.0f))); + if (heatmap.tokens_total / eff_period > last_sync_tokens / eff_period) { + if (getenv("LLAMA_EXPERT_FULL_SYNC")) { + // test: mirror the whole store to the top-S on each sync instead of + // the incremental handshake (pair with --expert-sync-period N) + return resync_full_mirror(heatmap, std::max(1, hot_s)); + } + return resync_top_s(heatmap); + } + return false; +} + +int llama_expert_hotstore::slot_of(int layer_idx, int expert_id) const { + if (layer_idx < 0 || layer_idx >= n_layers || hot_s <= 0) { + return -1; + } + const auto & ste = slot_to_expert[layer_idx]; + for (int p = 0; p < hot_s; p++) { + if (ste[p] == expert_id) { + return p; + } + } + return -1; +} + +void llama_expert_hotstore::update_luts(const std::vector & dirty) { + if (hot_s <= 0 || luts.empty() || buf_dev.empty()) { + return; + } + if (getenv("LLAMA_EXPERT_DEBUG")) { + static int once = 0; + if (!once) { + once = 1; + const char * base = (const char *) ggml_backend_buffer_get_base(buf_dev[0].get()); + fprintf(stderr, "hotstore: update_luts hot_lut[0] data_off=%td buf_base=%p\n", + (const char *) luts[0].hot_lut[0]->data - base, (const void *) base); + } + } + + std::vector cold_mask_h(n_experts); + + for (int il = 0; il < n_layers; il++) { + if (!dirty.empty() && !dirty[il]) { + continue; // this layer's store did not change; LUT already current + } + const auto & ste = slot_to_expert[il]; + const auto & rout = gpu_routed[il]; + + // per-device LUTs: an expert whose slot is on device g AND whose output + // currently comes from the GPU maps to the LOCAL slot index there; + // everything else (cold / pending) maps to the sentinel slot. + for (int g = 0; g < n_devices; g++) { + const int local_slots = slot_end[g] - slot_start[g]; + std::vector hot_lut_h(n_experts, local_slots); + for (int p = slot_start[g]; p < slot_end[g]; p++) { + const int e = ste[p]; + if (e >= 0 && e < n_experts && rout[e]) { + hot_lut_h[e] = p - slot_start[g]; + } + } + ggml_backend_tensor_set(luts[il].hot_lut[g], hot_lut_h.data(), 0, + n_experts * sizeof(int32_t)); + } + + // defaults: everyone cold + for (int e = 0; e < n_experts; e++) { + cold_mask_h[e] = 1; + } + // GPU-counted experts override + for (int e = 0; e < n_experts; e++) { + if (rout[e]) { + cold_mask_h[e] = 0; + } + } + ggml_backend_tensor_set(luts[il].cold_mask, cold_mask_h.data(), 0, + n_experts * sizeof(int32_t)); + if (getenv("LLAMA_EXPERT_DEBUG") && il == 0) { + static int once = 0; + if (!once) { + once = 1; + std::vector rb_hot(n_experts); + ggml_backend_tensor_get(luts[0].hot_lut[0], rb_hot.data(), 0, n_experts * sizeof(int32_t)); + fprintf(stderr, "hotstore: hot_lut[0..7]=%d %d %d %d %d %d %d %d cold_mask[0..3]=%d %d %d %d [96]=%d [97]=%d\n", + rb_hot[0], rb_hot[1], rb_hot[2], rb_hot[3], rb_hot[4], rb_hot[5], rb_hot[6], rb_hot[7], + ((const int32_t *) luts[0].cold_mask->data)[0], + ((const int32_t *) luts[0].cold_mask->data)[1], + ((const int32_t *) luts[0].cold_mask->data)[2], + ((const int32_t *) luts[0].cold_mask->data)[3], + ((const int32_t *) luts[0].cold_mask->data)[96], + ((const int32_t *) luts[0].cold_mask->data)[97]); + } + } + } +} + +void llama_expert_hotstore::log_hit_rate(const std::vector> & moe_sel) { + if (moe_sel.empty() || !is_filled) { + return; + } + size_t hits = 0, total = 0; + for (const auto & kv : moe_sel) { + const int il = kv.first; + const ggml_tensor * t = kv.second; + if (!t || !t->data || t->type != GGML_TYPE_I32) { + continue; + } + const size_t n = ggml_nelements(t); + std::vector ids(n); + ggml_backend_tensor_get(t, ids.data(), 0, n * sizeof(int32_t)); + for (size_t i = 0; i < n; i++) { + const int32_t id = ids[i]; + if (id >= 0 && id < n_experts) { + total++; + if (slot_of(il, id) >= 0) { + hits++; + } + } + } + } + if (total > 0) { + fprintf(stderr, "hotstore: hit rate %zu/%zu = %.1f%%\n", hits, total, 100.0f * (float) hits / (float) total); + } +} + +void llama_expert_hotstore::reset_counts() { + for (int il = 0; il < n_layers; il++) { + ggml_tensor * c = luts[il].counts; + if (c && c->data) { + memset(c->data, 0, (size_t) (n_experts + 1) * sizeof(int32_t)); + } + } +} + +void llama_expert_hotstore::read_counts(llama_expert_heatmap & heatmap, int n_tokens) { + if (luts.size() != (size_t) n_layers) { + return; + } + std::vector per_layer((size_t) n_layers, nullptr); + int64_t total = 0, hot = 0; + for (int il = 0; il < n_layers; il++) { + ggml_tensor * c = luts[il].counts; + const int32_t * cnt = c && c->data ? (const int32_t *) c->data : nullptr; + per_layer[il] = cnt; + if (!cnt) { + continue; + } + const auto & rout = gpu_routed[il]; + for (int e = 0; e < n_experts; e++) { + if (cnt[e] > 0) { + total += cnt[e]; + if (rout[e]) { + hot += cnt[e]; + } + } + } + } + if (total > 0) { + const float inst = (float) hot / (float) total; + hit_rate = hit_rate_valid ? 0.9f * hit_rate + 0.1f * inst : inst; + hit_rate_valid = true; + } + heatmap.update_counts(per_layer, n_tokens); +} + +float llama_expert_hotstore::target_hit_rate() const { + if (n_experts <= 0) { + return 0.5f; + } + const float frac = (float) hot_s / (float) n_experts; + const float t = 0.8f * std::pow(frac, 0.6f); + return std::max(0.2f, std::min(0.9f, t)); +} + +void llama_expert_hotstore::log() const { + fprintf(stderr, "hotstore: sizing (S=%d)\n", hot_s); + const bool debug = getenv("LLAMA_EXPERT_DEBUG") != nullptr; + size_t total = 0; for (int il = 0; il < n_layers; il++) { + total += bytes_per_slot[il]; + if (debug) { + fprintf(stderr, " layer %3d: bytes/slot = %zu\n", il, bytes_per_slot[il]); + } + } + fprintf(stderr, " total bytes/slot across all layers = %zu (%zu MiB)\n", + total, total / (1024 * 1024)); + if (!buf_dev.empty()) { + fprintf(stderr, " GPU hot store allocated: %s, %zu bytes (%zu MiB) for %d+1 slots across %d device(s) (%d expert + 1 sentinel per device)\n", + ggml_backend_buffer_name(buf_dev[0].get()), + ggml_backend_buffer_get_size(buf_dev[0].get()), + ggml_backend_buffer_get_size(buf_dev[0].get()) / (1024 * 1024), + hot_s, n_devices, hot_s); + } else if (hot_s > 0) { + fprintf(stderr, " hot store DISABLED (%d slots requested)\n", hot_s); + } +} + diff --git a/src/llama-expert-hotstore.h b/src/llama-expert-hotstore.h new file mode 100644 index 000000000000..83d140371a8b --- /dev/null +++ b/src/llama-expert-hotstore.h @@ -0,0 +1,194 @@ +#pragma once + +#include +#include +#include +#include +#include + +#include "ggml-cpp.h" + +struct llama_model; +struct llama_expert_heatmap; + +// stores per-layer sizing for the Mixture of Experts GPU hot store. +// one "slot" holds a single expert's weights for one layer. +struct llama_expert_hotstore { + int n_layers; + int n_experts; + int hot_s; + + // bytes of a single expert slot per layer, summed over that layer's + // expert weight tensors (gate/up/down, incl. chexps variants) + std::vector bytes_per_slot; + + // one hot tensor per expert weight tensor per device, shape {ne0, ne1, + // local_slots_g + 1}; the last plane is the zeroed sentinel slot + struct entry { + int layer_idx; + ggml_tensor* src; // model tensor holding all n_experts slices + std::vector dst; // per-device hot tensors + }; + std::vector entries; + + // per-layer index into entries (built once in ctor, entries stable after) + std::vector> entries_by_layer; + + // slot_to_expert[il][p] = expert id held in slot p of layer il, or -1 if empty. + // stable across re-syncs: an expert that stays hot keeps its slot. + std::vector> slot_to_expert; + + // per-layer LUT and mask for in-graph routing. + // hot_lut[g][e] = LOCAL slot index if e is hot on device g, else the + // device's local sentinel slot (zero contribution). + // cold_mask[e] = 1 if e is cold, else 0 (read as int zero-check by + // mul_mat_id_cold). + struct layer_lut { + std::vector hot_lut; // per-device i32[n_experts] + std::vector mask_lut; // per-device f32[local_slots+1], 0 at sentinel + ggml_tensor * cold_mask = nullptr; // i32[n_experts] + ggml_tensor * counts = nullptr; // i32[n_experts+1], tallied by the cold op + }; + std::vector luts; // size n_layers + + // per-device hot store: each device owns a contiguous slot range and its + // own no_alloc context + GPU buffer (dst tensors and hot_luts inside). + int n_devices = 1; + std::vector slot_start; // per-device slot range start (inclusive) + std::vector slot_end; // per-device slot range end (exclusive) + std::vector ctx_dev; + std::vector buf_dev; + + // CPU context and buffer for host-side tensors (like cold_mask) + ggml_context_ptr ctx_cpu; + ggml_backend_buffer_ptr buf_cpu; + + // true once the first copy of the top-S experts landed (once per session) + bool is_filled = false; // true when the store buffer was streamed by the loader (startup batch + // already resident), so the fill does not copy again + bool preloaded = false; + + // true when expert e of layer il is GPU-counted (output comes from the store) + bool is_resident(int il, int e) const { + return il >= 0 && il < (int) gpu_routed.size() && + e >= 0 && e < (int) gpu_routed[il].size() && gpu_routed[il][e] != 0; + } + + // re-sync cadence in tokens; 0 disables periodic re-sync + int sync_period = 0; + // tokens_total at the last sync (fill or re-sync) for boundary-cross check + int64_t last_sync_tokens = 0; + + // adaptive cadence: smoothed hit rate + the target the store must reach + // before the resync slows. target = 0.8 * (hot_s/n_experts)^0.6, clamped to + // [0.2, 0.9]: small models (store ~ full) aim ~80%, 1/5-fit models ~30%. + // hit_rate >= 0.95 disables the resync entirely. + float hit_rate = 0.0f; + bool hit_rate_valid = false; + + float target_hit_rate() const; + + // start-up full sync: after the 3-token heat boost, mirror the store to the + // top-S over the next `full_sync_remaining` tokens (budget hot_s/4 per token) + bool full_sync_done = false; + int full_sync_remaining = 0; + + // hysteresis gate: a resident slot is only swapped when a cold + // expert scores >= hyst * the incumbent AND the slot has dwelled long enough + float hyst = 0.0f; // 0 = gate off (swap freely) + int dwell = 0; // minimum syncs a resident must keep; 0 = off + bool copy_mode = false; // resolved: keep the RAM copy of promoted experts + int mode = 0; // user mode: 0 = auto, 1 = copy, 2 = move + // max concurrent transfers in flight per layer, per direction (eviction + // queue + promotion queue). higher = faster store adaptation, at the cost + // of more slots temporarily in transition (not counted). + int max_concurrent_moves = 3; + // dwell_count[il][p] = syncs since slot p last changed (0 = fresh/empty) + std::vector> dwell_count; + + // swap lifecycle handshake. slot_to_expert holds the physical slot + // content; gpu_routed holds the experts whose output actually comes from + // the GPU. a moved-in expert stays CPU-counted (gpu_routed=0) while it is + // pending verification; all pending copies are verified every token, and a + // confirmed one is routed to the GPU only on the following token. a + // moved-out expert is copied back on read, becomes CPU-counted, and runs + // one full generation on the CPU before its GPU copy is considered gone. + struct pending_move_in { + int expert; + int p; // slot holding the copy + bool verified; + int countdown; // 1 = route to the GPU on the next token + int failures; // consecutive failed verifies (anti-stall) + }; + struct pending_move_out { + int expert; + int p; // slot to free once the CPU takes over + bool verified; // staging copy matches the launch hash + int countdown; // 1 = CPU output counts on the next token + int failures; // consecutive failed verifies (anti-stall) + }; + std::vector> gpu_routed; // [il][e] + std::vector> pending_in; // [il] + std::vector> pending_out;// [il] + + // per-layer staging buffer: one expert slot per exps tensor of the layer, + // holding the in-flight evicted expert's slices while the CPU copy is + // verified against the launch hash (the GPU output stays valid meanwhile). + std::vector cpu_staging; + std::vector cpu_staging_off; // [il] offset into cpu_staging + +llama_expert_hotstore(const llama_model * model, int n_layers, + int n_experts, int hot_s, int sync_period = 0, + float hyst = 0.0f, int dwell = 0, int mode = 0); + + ~llama_expert_hotstore(); + + // allocate the GPU hot store for `hot_s` slots, split across the given + // device buffer types by tensor_split (fractions, one per device). returns + // false (and leaves the store disabled) on failure or shortage of VRAM. + bool allocate(const std::vector & bufts, + const float * tensor_split, int n_split); + + // copy the top-S expert slices for every layer into the GPU hot store, + // using the given heatmap for the ranking. one-shot (guarded by is_filled). + // copy the top-S expert slices for every layer into the GPU hot store, + // using the given heatmap for the ranking. one-shot (guarded by is_filled). + // returns true if a fill happened (caller should synchronize the GPU). + bool copy_top_s(const llama_expert_heatmap & heatmap); + + // re-sync the hot store to the current heatmap ranking, swapping only + // the experts that changed (stable slots; unchanged experts not re-copied). + // returns true if any slot changed (caller should synchronize the GPU). + bool resync_top_s(const llama_expert_heatmap & heatmap); + + + // LLAMA_EXPERT_FULL_SYNC: mirror the store to the top-S every token + // (direct swaps, hash-verified, copy-on-read). no pacing queues. + bool resync_full_mirror(const llama_expert_heatmap & heatmap, int budget = 1); + + // cadence-gated wrapper: re-sync only if tokens_total crossed sync_period; + // multi_slot freezes the hot store (static slots, no swapping). returns + // true if a re-sync ran and swapped slots. + bool maybe_resync(const llama_expert_heatmap & heatmap, bool multi_slot); + + // returns the GPU slot index holding expert_id in layer il, or -1 if none + int slot_of(int layer_idx, int expert_id) const; + + // zero the cold-op per-expert counts (call before the graph compute) + void reset_counts(); + + // feed the cold-op counts into the heatmap (host memory, no D2H readback). + // call after the graph compute; n_tokens advances the heatmap clock. + void read_counts(llama_expert_heatmap & heatmap, int n_tokens); + + // diagnostic: count how many router-selected expert ids hit a hot slot. + // reads the selected_experts tensors (call after synchronize). + void log_hit_rate(const std::vector> & moe_sel); + + // rebuild hot_lut/cold_mask from slot_to_expert for every layer + // and copy them into the tensors. called from copy_top_s (initial + // fill) and resync_top_s (swaps). empty dirty = all layers. + void update_luts(const std::vector & dirty = {}); + + void log() const; +}; diff --git a/src/llama-expert-pin.cpp b/src/llama-expert-pin.cpp new file mode 100644 index 000000000000..84528fde5422 --- /dev/null +++ b/src/llama-expert-pin.cpp @@ -0,0 +1,245 @@ +#include "llama-expert-pin.h" + +#include "llama-expert-heatmap.h" +#include "llama-impl.h" +#include "llama-model.h" + +#include "ggml.h" +#include "ggml-backend.h" + +#include +#include +#include +#include +#include +#include + +#if !defined(_WIN32) +#include +#include +#else +#ifndef NOMINMAX +#define NOMINMAX +#endif +#include +#endif + +namespace llama_expert_pin { + +static void madvise_range(const void * p, size_t len, bool keep) { +#if !defined(_WIN32) + static const long page = sysconf(_SC_PAGESIZE); + const uintptr_t a = (uintptr_t) p & ~(uintptr_t) (page - 1); + const uintptr_t b = ((uintptr_t) p + len + page - 1) & ~(uintptr_t) (page - 1); + if (b > a) { + madvise((void *) a, b - a, keep ? MADV_WILLNEED : MADV_DONTNEED); + } +#else + (void) p; (void) len; (void) keep; +#endif +} + +// only hint file-backed (mmap) host tensors; DONTNEED on anonymous buffers +// would zero them. no data pointer = no pages to hint. +static bool hintable(const ggml_tensor * w) { + if (!w || !w->data) { + return false; + } + const ggml_backend_buffer_t buf = w->view_src ? w->view_src->buffer : w->buffer; + return buf != nullptr && ggml_backend_buffer_is_host(buf); +} + +static const config g_config = [] { + config c; + if (const char * e = getenv("LLAMA_EXPERT_PIN_PERIOD")) { + c.period = std::max(1, atoi(e)); + } + if (const char * e = getenv("LLAMA_EXPERT_PIN_START")) { + c.start_tokens = std::max(0, atoi(e)); + } + if (const char * e = getenv("LLAMA_EXPERT_PIN_DONTNEED_GPU")) { + c.dontneed_gpu = (float) atof(e); + } + if (const char * e = getenv("LLAMA_EXPERT_PIN_WILLNEED_GPU")) { + c.willneed_gpu = (float) atof(e); + } + if (const char * e = getenv("LLAMA_EXPERT_PIN_WILLNEED_COLD")) { + c.willneed_cold = (float) atof(e); + } + return c; +}(); + +const config & get_config() { + return g_config; +} + +static int g_pct = -1; + +void set_pct(int pct) { + g_pct = pct; +} + +int get_pct() { + return g_pct; +} + +bool active() { + return getenv("LLAMA_EXPERT_PIN") != nullptr || g_pct > 0; +} + +// exps tensors per layer, e.g. blk.0.ffn_gate_exps.weight +static const std::regex g_re_exps("blk\\.(\\d+)\\.ffn_(up|down|gate|gate_up)_(ch|)exps\\.weight"); + +static void hint_expert(const ggml_tensor * w, int expert, bool keep) { + if (!hintable(w) || expert < 0 || expert >= (int) w->ne[2]) { + return; + } + const size_t plane = ggml_nbytes(w) / (size_t) w->ne[2]; + madvise_range((const char *) w->data + (size_t) expert * plane, plane, keep); +} + +static void hint_experts(const llama_model * model, + const llama_expert_heatmap & heatmap, + bool (*is_gpu_resident)(void * ud, int il, int e), + void * ud) { + const int n_layers = heatmap.n_layers; + const int n_experts = heatmap.n_experts; + + // RAM-pressure gate: only evict cold pages when the system is genuinely + // tight (MemAvailable under 10% of total). + bool ram_tight = false; + int64_t mem_total = 0, mem_avail = 0; +#if defined(__linux__) + FILE * f = fopen("/proc/meminfo", "r"); + if (f) { + char line[256]; + while (fgets(line, sizeof(line), f)) { + if (sscanf(line, "MemTotal: %" PRId64 " kB", &mem_total) == 1) { + mem_total *= 1024; + } else if (sscanf(line, "MemAvailable: %" PRId64 " kB", &mem_avail) == 1) { + mem_avail *= 1024; + } + } + fclose(f); + } +#endif + ram_tight = mem_total > 0 && mem_avail * 10 < mem_total; + + // group the exps tensors by layer + std::vector> per_layer(n_layers); + for (const auto & [name, tensor] : llama_internal_get_tensor_map(model)) { + std::smatch m; + if (std::regex_search(name, m, g_re_exps)) { + const int il = std::stoi(m[1].str()); + if (il >= 0 && il < n_layers) { + per_layer[il].push_back(tensor); + } + } + } + + std::vector idx(n_experts); + std::vector score(n_experts); + + for (int il = 0; il < n_layers; il++) { + if (per_layer[il].empty()) { + continue; + } + for (int e = 0; e < n_experts; e++) { + score[e] = heatmap.get_score(il, e); + idx[e] = e; + } + + // partition by residency, sort each by score + std::vector gpu, cold; + gpu.reserve(n_experts); + cold.reserve(n_experts); + for (int e = 0; e < n_experts; e++) { + if (is_gpu_resident && is_gpu_resident(ud, il, e)) { + gpu.push_back(e); + } else { + cold.push_back(e); + } + } + auto by_score = [&](int a, int b) { return score[a] > score[b]; }; + + // GPU residents: drop the top fraction, warm the bottom fraction + if (!gpu.empty()) { + std::sort(gpu.begin(), gpu.end(), by_score); + const int n_drop = std::min((int) gpu.size(), + (int) (g_config.dontneed_gpu * (float) gpu.size())); + const int n_warm = std::min((int) gpu.size(), + (int) (g_config.willneed_gpu * (float) gpu.size())); + for (int i = 0; i < n_drop; i++) { + for (const ggml_tensor * w : per_layer[il]) { + hint_expert(w, gpu[i], false); + } + } + for (int i = (int) gpu.size() - n_warm; i < (int) gpu.size(); i++) { + for (const ggml_tensor * w : per_layer[il]) { + hint_expert(w, gpu[i], true); + } + } + } + + // cold: warm the top fraction (most likely promoted next) + if (!cold.empty()) { + std::sort(cold.begin(), cold.end(), by_score); + float cold_pct = g_config.willneed_cold; + if (const char * e = getenv("LLAMA_EXPERT_PIN_WILLNEED_COLD")) { + cold_pct = (float) atof(e); + } else if (g_pct > 0) { + cold_pct = (float) g_pct / 100.0f; + } + const int n_warm = std::min((int) cold.size(), + (int) (cold_pct * (float) cold.size())); + for (int i = 0; i < n_warm; i++) { + for (const ggml_tensor * w : per_layer[il]) { + hint_expert(w, cold[i], true); + } + } + } + + // RAM pressure: evict the bottom 10% of total heat that is resident + // in the mmap pool and has not been reused for a full dwell cycle. + if (ram_tight) { + std::vector all(n_experts); + for (int e = 0; e < n_experts; e++) { + all[e] = e; + } + std::sort(all.begin(), all.end(), by_score); + const int n_evict = std::max(1, (int) (0.10f * (float) n_experts)); + for (int i = 0; i < n_evict; i++) { + const int e = all[i]; + if (is_gpu_resident && is_gpu_resident(ud, il, e)) { + continue; // on GPU: its mmap pages are already dropped + } + if (heatmap.tokens_total - heatmap.last_reuse[il * n_experts + e] < 8) { + continue; // reused within the dwell cycle: keep it + } + for (const ggml_tensor * w : per_layer[il]) { + hint_expert(w, e, false); + } + } + } + } +} + +void maybe_run(const llama_model * model, + const llama_expert_heatmap * heatmap, + bool (*is_gpu_resident)(void * ud, int il, int e), + void * ud) { + if (!model || !heatmap) { + return; + } + static int64_t last_run = -1; + if (heatmap->tokens_total < g_config.start_tokens) { + return; + } + if (last_run >= 0 && heatmap->tokens_total - last_run < g_config.period) { + return; + } + last_run = heatmap->tokens_total; + hint_experts(model, *heatmap, is_gpu_resident, ud); +} + +} diff --git a/src/llama-expert-pin.h b/src/llama-expert-pin.h new file mode 100644 index 000000000000..760f37ce5742 --- /dev/null +++ b/src/llama-expert-pin.h @@ -0,0 +1,37 @@ +#pragma once + +#include + +struct llama_expert_heatmap; +struct llama_model; + +// mmap page hints for the expert tier. periodic madvise pass: drop the hottest +// GPU experts' pages, warm the ones most likely needed next. +namespace llama_expert_pin { + + // dials, all env-overridable with these defaults + struct config { + int period = 24; // tokens between passes + int start_tokens = 128; // first pass at this many tokens + float dontneed_gpu = 0.20f; // top fraction of GPU experts to drop + float willneed_gpu = 0.20f; // bottom fraction of GPU experts to warm + float willneed_cold = 0.30f; // top fraction of cold experts to warm + }; + + const config & get_config(); + + // resolved pin fraction from --expert-pin (percent, -1 = unset) + void set_pct(int pct); + int get_pct(); + + // true when pinning is enabled (LLAMA_EXPERT_PIN set or pct > 0) + bool active(); + + // periodic madvise pass. is_gpu_resident marks store residents, + // nullptr = standalone (all cold) + void maybe_run(const llama_model * model, + const llama_expert_heatmap * heatmap, + bool (*is_gpu_resident)(void * ud, int il, int e), + void * ud); + +} diff --git a/src/llama-expert-preload.cpp b/src/llama-expert-preload.cpp new file mode 100644 index 000000000000..5074b0199772 --- /dev/null +++ b/src/llama-expert-preload.cpp @@ -0,0 +1,384 @@ +#include "llama-expert-preload.h" + +#include "ggml.h" +#include "ggml-backend.h" +#include "ggml-cpu.h" + +#include +#include +#include +#include +#include + +#ifdef _WIN32 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#include +#include +#include +#include +// MSVC POSIX layer: map open/close to _open/_close and emulate pread +#define open _open +#define close _close +#define O_RDONLY _O_RDONLY +static long long pread(int fd, void * buf, size_t n, long long off) { + if (_lseeki64(fd, off, SEEK_SET) < 0) { + return -1; + } + return _read(fd, buf, (unsigned int) n); +} +#else +#include +#include +#include +#endif + +namespace llama_expert_preload { + +namespace { + // matches an expert weight tensor, e.g. blk.0.ffn_gate_exps.weight + const std::regex g_re_exps("blk\\.(\\d+)\\.ffn_(up|down|gate|gate_up)_(ch|)exps\\.weight"); + + int g_slots = 0; + + struct ggml_backend_buffer * g_gpu_buf = nullptr; + struct ggml_context * g_ctx = nullptr; // temp context for the store write tensors + uint8_t * g_cpu_buf = nullptr; // host buffer for the cold experts + size_t g_cpu_size = 0; + std::vector g_entries; + std::unordered_map g_src_idx; // src -> entry index + std::vector> g_addrs; // [entry][expert] -> host address + std::vector> g_hashes; // [entry][expert] -> gguf FNV-1a(1024B) + size_t g_gpu_cursor = 0; // next free offset in the store + size_t g_cpu_cursor = 0; // next free offset in the host buffer + + static uint64_t fnv1a(const uint8_t * p, size_t n) { + uint64_t h = 0xcbf29ce484222325ULL; + for (size_t i = 0; i < n; i++) { + h ^= p[i]; + h *= 0x100000001b3ULL; + } + return h; + } + + std::string g_path; + int g_fd = -1; + + // C callback for the ggml cold op: resolve a cold expert's address from the + // table (set at write time, so it always matches where the slice landed) + const uint8_t * preload_slice_cb(const struct ggml_tensor * src0, int expert) { + if (!g_cpu_buf || !src0 || expert < 0) { + return nullptr; + } + const int i = index_of(src0); + if (i < 0 || expert >= (int) g_addrs[i].size()) { + return nullptr; + } + return (const uint8_t *) (uintptr_t) g_addrs[i][expert]; + } +} + +LLAMA_API void set_model_path(const char * path) { + g_path = path ? path : ""; + if (g_fd >= 0) { + close(g_fd); + g_fd = -1; + } + g_fd = open(g_path.c_str(), O_RDONLY); +} +LLAMA_API void set_slots(int s) { + g_slots = s; +} + +int get_slots() { + return g_slots; +} + +bool g_no_evict = false; + +LLAMA_API void set_no_evict(bool no_evict) { + g_no_evict = no_evict; +} + +bool get_no_evict() { + return g_no_evict; +} + +bool tier_will_engage() { + if (g_slots <= 0) { + return false; + } + for (size_t i = 0; i < ggml_backend_dev_count(); i++) { + const ggml_backend_dev_t dev = ggml_backend_dev_get(i); + if (dev && ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_GPU) { + return true; // any GPU backend (CUDA, Vulkan, ROCm, SYCL, Metal, ...) + } + } + return false; +} + +size_t align_up256(size_t x) { + const size_t align = 256; + return (x + align - 1) & ~(align - 1); +} + +size_t store_total_bytes(size_t entries_bytes, int n_layers, int n_experts, int slots) { + size_t off = align_up256(entries_bytes); + for (int il = 0; il < n_layers; il++) { + off = align_up256(off); + off += align_up256((size_t) n_experts * sizeof(int32_t)); + off = align_up256(off); + off += align_up256((size_t) (slots + 1) * sizeof(float)); + } + return off; +} + +bool is_exps(const char * name, int & layer_idx) { + if (!name) { + return false; + } + std::cmatch m; + if (!std::regex_search(name, m, g_re_exps)) { + return false; + } + layer_idx = std::stoi(m[1].str()); + return true; +} + +void begin(struct ggml_backend_buffer_type * buft, size_t gpu_bytes, size_t cpu_bytes, int max_entries) { + if (g_gpu_buf || g_cpu_buf) { + return; // load_all_data can run more than once (fit estimate + real load) + } + g_gpu_buf = ggml_backend_buft_alloc_buffer(buft, gpu_bytes); + // zero the whole store up front: the sentinel plane and any unwritten gaps + // would otherwise carry per-launch garbage into the graph output (Vulkan) + ggml_backend_buffer_clear(g_gpu_buf, 0); + g_ctx = ggml_init({ ggml_tensor_overhead() * (max_entries + 8), nullptr, true }); + g_cpu_buf = (uint8_t *) malloc(cpu_bytes); + g_cpu_size = cpu_bytes; + g_entries.clear(); + g_src_idx.clear(); + g_addrs.clear(); + g_hashes.clear(); + g_gpu_cursor = 0; + g_cpu_cursor = 0; +} + +size_t register_tensor(const ggml_tensor * src, size_t plane_bytes, int n_experts, int startup, + size_t file_off, int fd) { + for (size_t i = 0; i < g_entries.size(); i++) { + if (g_entries[i].src == src) { + g_entries[i].fd = fd; // refresh: the file may be reopened between passes + return i; // already registered (second load pass) + } + } + const size_t gpu_off = g_gpu_cursor; + const size_t cpu_off = g_cpu_cursor; + g_entries.push_back({src, plane_bytes, gpu_off, cpu_off, file_off, fd, n_experts, startup}); + g_src_idx[src] = g_entries.size() - 1; + g_addrs.emplace_back((size_t) n_experts, 0); + g_hashes.emplace_back((size_t) n_experts, 0); + g_gpu_cursor += (size_t) (startup + 1) * plane_bytes; + g_cpu_cursor += (size_t) n_experts * plane_bytes; // all experts, cold committed at load + return g_entries.size() - 1; +} + +bool write_entry(size_t idx, const uint8_t * data, size_t nbytes) { + if (idx >= g_entries.size() || !data) { + return false; + } + const entry & e = g_entries[idx]; + if (nbytes < (size_t) e.startup * e.plane_bytes) { + return false; + } + // stream the startup slices into the store buffer at load (fast startup, + // no token-1 file reads). the store adopts this buffer and re-allocates + // its dst tensors at the same gpu_offset; the scratch tensor below dies + // with g_ctx at take_buffer(). the cold op's slice callback ignores the + // startup experts (they are GPU-counted), so no host copy is needed. + if (g_gpu_buf && g_ctx) { + ggml_tensor * t = ggml_new_tensor_1d(g_ctx, GGML_TYPE_I8, e.plane_bytes); + uint8_t * base = (uint8_t *) ggml_backend_buffer_get_base(g_gpu_buf); + for (int ex = 0; ex < e.startup; ex++) { + uint8_t * dst = base + e.gpu_offset + (size_t) ex * e.plane_bytes; + if (ex == 0) { + ggml_backend_tensor_alloc(g_gpu_buf, t, dst); + } else { + t->data = dst; // manual view into the adopted store buffer + } + ggml_backend_tensor_set(t, data + (size_t) ex * e.plane_bytes, 0, e.plane_bytes); + } + } + const size_t chunk = e.plane_bytes < 1024 ? e.plane_bytes : 1024; + for (int ex = 0; ex < e.startup; ex++) { + g_hashes[idx][ex] = fnv1a(data + (size_t) ex * e.plane_bytes, chunk); + } + return true; +} + +bool read_expert(size_t idx, int expert, void * out, size_t n) { + if (idx >= g_entries.size() || expert < 0 || expert >= g_entries[idx].n_experts) { + return false; + } + const entry & e = g_entries[idx]; + if (n > e.plane_bytes) { + return false; + } + const long long got = pread(g_fd >= 0 ? g_fd : e.fd, out, n, + (long long) (e.file_off + (size_t) expert * e.plane_bytes)); + return got == (long long) n; +} + +bool write_cold(size_t idx, const uint8_t * data, size_t nbytes) { + if (idx >= g_entries.size() || !g_cpu_buf || !data) { + fprintf(stderr, "write_cold: FAIL idx=%zu cpu_buf=%p data=%p nbytes=%zu\n", + idx, (void *) g_cpu_buf, (const void *) data, nbytes); + return false; + } + const entry & e = g_entries[idx]; + const size_t dst = e.cpu_offset + (size_t) e.startup * e.plane_bytes; + if (dst + nbytes > g_cpu_size) { + fprintf(stderr, "write_cold: OUT OF BOUNDS\n"); + return false; + } + std::memcpy(g_cpu_buf + dst, data, nbytes); + const size_t chunk = e.plane_bytes < 1024 ? e.plane_bytes : 1024; + for (int ex = e.startup; ex < e.n_experts; ex++) { + g_addrs[idx][ex] = (uint64_t) (uintptr_t) (g_cpu_buf + e.cpu_offset + (size_t) ex * e.plane_bytes); + g_hashes[idx][ex] = fnv1a(data + (size_t) (ex - e.startup) * e.plane_bytes, chunk); + } + if (getenv("LLAMA_EXPERT_DEBUG") && idx == 0) { + fprintf(stderr, "write_cold: idx0 %s first bytes: %02x %02x %02x %02x\n", + e.src->name, g_cpu_buf[dst], g_cpu_buf[dst+1], g_cpu_buf[dst+2], g_cpu_buf[dst+3]); + } + ggml_mmid_cold_set_slice_fn(preload_slice_cb); + return true; +} + +struct ggml_backend_buffer * take_buffer() { + struct ggml_backend_buffer * b = g_gpu_buf; + g_gpu_buf = nullptr; + if (g_ctx) { + ggml_free(g_ctx); + g_ctx = nullptr; + } + return b; +} + +size_t num_entries() { + return g_entries.size(); +} + +const entry * entry_at(size_t idx) { + return idx < g_entries.size() ? &g_entries[idx] : nullptr; +} + +int index_of(const ggml_tensor * src) { + if (!src) { + return -1; + } + for (size_t i = 0; i < g_entries.size(); i++) { + if (g_entries[i].src == src) { + return (int) i; + } + } + // the hotstore's entry tensors can be distinct objects with the same name + // (created in a different context); fall back to matching by name + for (size_t i = 0; i < g_entries.size(); i++) { + if (g_entries[i].src && g_entries[i].src->name[0] && + strcmp(g_entries[i].src->name, src->name) == 0) { + return (int) i; + } + } + return -1; +} + +uint64_t expected_hash(const ggml_tensor * src, int expert) { + const int idx = index_of(src); + if (idx < 0 || expert < 0 || expert >= (int) g_hashes[idx].size()) { + return 0; + } + return g_hashes[idx][expert]; +} + +size_t entries_size() { + return g_gpu_cursor; +} + +const uint8_t * cpu_slice(size_t idx, int expert) { + if (idx >= g_entries.size() || !g_cpu_buf) { + return nullptr; + } + const entry & e = g_entries[idx]; + if (expert < 0 || expert >= e.n_experts) { + return nullptr; + } + return g_cpu_buf + e.cpu_offset + (size_t) expert * e.plane_bytes; +} + +static void release_pages(void * ptr, size_t len) { +#ifdef _WIN32 + SYSTEM_INFO si; + GetSystemInfo(&si); + const size_t page = si.dwPageSize; +#else + const long page = sysconf(_SC_PAGESIZE); +#endif + const uintptr_t base = (uintptr_t) ptr; + const uintptr_t start = (base + (uintptr_t) page - 1) & ~((uintptr_t) page - 1); + const uintptr_t end = (base + len) & ~((uintptr_t) page - 1); + if (start < end) { +#ifdef _WIN32 + VirtualFree((LPVOID) start, end - start, MEM_RESET); +#else + madvise((void *) start, end - start, MADV_DONTNEED); +#endif + } +} + +void free_cpu_slice(size_t idx, int expert) { + if (idx >= g_entries.size() || !g_cpu_buf) { + return; + } + const entry & e = g_entries[idx]; + if (expert < 0 || expert >= e.n_experts) { + return; + } + release_pages(g_cpu_buf + e.cpu_offset + (size_t) expert * e.plane_bytes, e.plane_bytes); +} + +void set_cpu_slice(size_t idx, int expert, const uint8_t * data) { + if (idx >= g_entries.size() || !g_cpu_buf || !data) { + return; + } + const entry & e = g_entries[idx]; + if (expert < 0 || expert >= e.n_experts) { + return; + } + std::memcpy(g_cpu_buf + e.cpu_offset + (size_t) expert * e.plane_bytes, data, e.plane_bytes); + g_addrs[idx][expert] = (uint64_t) (uintptr_t) (g_cpu_buf + e.cpu_offset + (size_t) expert * e.plane_bytes); +} + +void clear() { + if (g_gpu_buf) { + ggml_backend_buffer_free(g_gpu_buf); + g_gpu_buf = nullptr; + } + if (g_ctx) { + ggml_free(g_ctx); + g_ctx = nullptr; + } + if (g_cpu_buf) { + free(g_cpu_buf); + g_cpu_buf = nullptr; + } + g_cpu_size = 0; + g_entries.clear(); + g_src_idx.clear(); + g_addrs.clear(); + g_gpu_cursor = 0; + g_cpu_cursor = 0; +} + +} // namespace llama_expert_preload diff --git a/src/llama-expert-preload.h b/src/llama-expert-preload.h new file mode 100644 index 000000000000..c16c07e27be3 --- /dev/null +++ b/src/llama-expert-preload.h @@ -0,0 +1,84 @@ +#pragma once + +#include "llama.h" + +#include +#include + +struct ggml_tensor; +struct ggml_backend_buffer; +struct ggml_backend_buffer_type; + +// pre-load handoff between the model loader and the expert hot store. +// +// when the tier is active (--expert-hot-s set), the model loader does not +// allocate the expert weight tensors at all; instead it streams their data +// into two buffers owned here: a GPU store holding the first S experts of +// every layer, and a host buffer holding the remaining (cold) experts. this +// keeps the experts entirely outside the model's own buffer plan, so a model +// larger than RAM can load when the GPU store supplies the missing capacity. +// the hotstore adopts the GPU store at context init; the cold op reads the +// host buffer through a slice hook. +namespace llama_expert_preload { + + LLAMA_API void set_slots(int s); + int get_slots(); + LLAMA_API void set_model_path(const char * path); // for the debug disk hash + LLAMA_API void set_no_evict(bool no_evict); // --expert-no-evict + bool get_no_evict(); + + // true when the tier will actually build a hot store: slots requested AND + // (forced OR a CUDA device is present). mirrors the gate in llama-context.cpp. + bool tier_will_engage(); + + // backend-agnostic 256-byte alignment for the store's LUT region: some + // backends (Vulkan) require minStorageBufferOffsetAlignment for get_rows + // sources. returns the total store bytes (entries + padded LUTs). + size_t store_total_bytes(size_t entries_bytes, int n_layers, int n_experts, int slots); + size_t align_up256(size_t x); + + // true if `name` matches an exps weight tensor; sets layer_idx on match + bool is_exps(const char * name, int & layer_idx); + + struct entry { + const ggml_tensor * src; + size_t plane_bytes; // bytes of one expert slice in this tensor + size_t gpu_offset; // offset of this entry's planes in the store buffer + size_t cpu_offset; // offset of this entry's cold experts in the host buffer + size_t file_off; // absolute gguf data offset of this tensor + int fd; // model file descriptor (for the debug disk read) + int n_experts; + int startup; // first S experts (live on the GPU) + }; + + // loader side. begin() allocates the two buffers (single call, after the + // loader computed the sizes); register_tensor() adds an entry; write_entry() + // writes the startup slices to the store; write_cold() writes the rest to + // the host buffer. + void begin(struct ggml_backend_buffer_type * buft, size_t gpu_bytes, size_t cpu_bytes, int max_entries); + size_t register_tensor(const ggml_tensor * src, size_t plane_bytes, int n_experts, int startup, + size_t file_off, int fd); + bool write_entry(size_t idx, const uint8_t * data, size_t nbytes); + bool read_expert(size_t idx, int expert, void * out, size_t n); // from the gguf file + bool write_cold(size_t idx, const uint8_t * data, size_t nbytes); + + // hotstore + cold-op side. + struct ggml_backend_buffer * take_buffer(); + size_t num_entries(); + const entry * entry_at(size_t idx); + int index_of(const ggml_tensor * src); + size_t entries_size(); // bytes used by the entry regions in the store + + // cold expert slice access (nullptr when the slice is on the GPU) + const uint8_t * cpu_slice(size_t idx, int expert); + void free_cpu_slice(size_t idx, int expert); + void set_cpu_slice(size_t idx, int expert, const uint8_t * data); + + // debug: ground-truth FNV-1a of the first 1024 bytes of an expert slice, + // taken straight from the GGUF at load. compare against the hash of the + // slice actually read to verify a copy/routing is correct. + uint64_t expected_hash(const ggml_tensor * src, int expert); + + void clear(); + +} // namespace llama_expert_preload diff --git a/src/llama-expert-tier.cpp b/src/llama-expert-tier.cpp new file mode 100644 index 000000000000..0a0ee60dd1e3 --- /dev/null +++ b/src/llama-expert-tier.cpp @@ -0,0 +1,162 @@ +#include "llama-expert-tier.h" + +#include +#include + +namespace { + struct tier_entry { + std::vector dst_hot; // per-device hot tensors + std::vector hot_lut; // per-device LUTs + std::vector mask_lut; // per-device sentinel masks + ggml_tensor * cold_mask; + ggml_tensor * counts; // i32[n_experts+1], tallied by the cold op + }; + + std::mutex g_mtx; + std::unordered_map g_table; + + // fused cold path: while a layer's experts are being built, the tier build + // returns hot-only results; the fused op (built by end_fused) covers cold. + bool g_fused_active = false; + + // fused path only kicks in for batches up to this many tokens (gated on + // ids->ne[1]); larger batches fall back to the per-op cold path. + static int g_tmax = 16; +} + +void llama_expert_tier_register(ggml_tensor * src, + const std::vector & dst_hot, + const std::vector & hot_lut, + const std::vector & mask_lut, + ggml_tensor * cold_mask, + ggml_tensor * counts) { + std::lock_guard lk(g_mtx); + g_table[src] = {dst_hot, hot_lut, mask_lut, cold_mask, counts}; +} + +void llama_expert_tier_clear() { + std::lock_guard lk(g_mtx); + g_table.clear(); +} + +bool llama_expert_tier_has(ggml_tensor * w) { + std::lock_guard lk(g_mtx); + return g_table.find(w) != g_table.end(); +} + +// Remap real expert ids through a LUT to slot indices, returning a 2d +// [n_expert_used, n_tokens] i32 tensor usable as ids for ggml_mul_mat_id. +// The lut is a [1, n_experts] table, so ggml_get_rows picks one scalar per +// id. ggml_cont guards against argsort views that may not be contiguous. +static ggml_tensor * remap_ids(ggml_context * ctx, + ggml_tensor * lut, + ggml_tensor * selected, + int n_experts, + int n_expert_used, + int n_tokens) { + (void)n_experts; + ggml_tensor * flat_ids = ggml_reshape_1d(ctx, + ggml_cont(ctx, selected), n_expert_used * n_tokens); + ggml_tensor * r = ggml_get_rows(ctx, lut, flat_ids); + return ggml_reshape_2d(ctx, r, n_expert_used, n_tokens); +} + +ggml_tensor * llama_expert_tier_build(ggml_context * ctx, + ggml_tensor * w, + ggml_tensor * cur, + ggml_tensor * ids, + ggml_tensor * w_s) { + // the count+rank mmid helper (see mmid.cu) handles duplicate expert ids per + // token, so the tier is safe for any batch size. + + tier_entry ent; + { + std::lock_guard lk(g_mtx); + auto it = g_table.find(w); + if (it == g_table.end()) { + return nullptr; + } + ent = it->second; + } + + const int n_experts = (int) w->ne[2]; + const int n_expert_used = (int) ids->ne[0]; + const int n_tokens = (int) cur->ne[2]; + + // hot: for each device, remap real expert ids to that device's LOCAL slot + // indices; experts whose slot lives on another device (or are cold) map to + // this device's zeroed sentinel slot and contribute nothing. Sum the + // per-device results (the scheduler inserts any cross-device copies). + ggml_tensor * hot = nullptr; + for (size_t g = 0; g < ent.dst_hot.size(); g++) { + ggml_tensor * ids_hot = remap_ids(ctx, ent.hot_lut[g], ids, n_experts, n_expert_used, n_tokens); + ggml_tensor * h = ggml_mul_mat_id(ctx, ent.dst_hot[g], cur, ids_hot); + hot = hot ? ggml_add(ctx, hot, h) : h; + } + + // fused cold path active for this layer: the CPU cold op is deferred to + // end_fused, so return the hot contribution only + if (g_fused_active) { + (void)w_s; + return hot; + } + + // cold: dedicated CPU op that computes only the cold-selected experts, + // skipping hot ones via the integer zero-check on cold_mask. the same op + // tallies the selected ids into ent.counts (host memory) for the heatmap. + ggml_tensor * cold = ggml_mul_mat_id_cold(ctx, w, cur, ids, ent.cold_mask, ent.counts, nullptr); + + (void)w_s; // per-expert quant scale is intentionally discarded on the tiered path + + return ggml_add(ctx, hot, cold); +} + +bool llama_expert_tier_begin_fused(ggml_tensor * gate_w, + ggml_tensor * up_w, + ggml_tensor * down_w, + ggml_tensor * ids) { + g_fused_active = false; + if (!gate_w || !up_w || !down_w) { + return false; + } + if (ids->ne[1] > (int64_t) g_tmax) { + return false; // batch too large for the fused path + } + const int n_experts = (int) down_w->ne[2]; + std::lock_guard lk(g_mtx); + if (g_table.find(gate_w) == g_table.end() || + g_table.find(up_w) == g_table.end() || + g_table.find(down_w) == g_table.end()) { + return false; // some tensors of this layer are not tiered + } + if (up_w != gate_w && (int) up_w->ne[2] != n_experts) { + return false; + } + g_fused_active = true; + return true; +} + +ggml_tensor * llama_expert_tier_end_fused(ggml_context * ctx, + ggml_tensor * gate_w, + ggml_tensor * up_w, + ggml_tensor * down_w, + ggml_tensor * x, + ggml_tensor * ids, + int32_t act) { + if (!g_fused_active) { + return nullptr; + } + g_fused_active = false; + + tier_entry ent; + { + std::lock_guard lk(g_mtx); + auto it = g_table.find(down_w); + if (it == g_table.end()) { + return nullptr; + } + ent = it->second; + } + + return ggml_moe_cold(ctx, gate_w, up_w, down_w, x, ids, ent.cold_mask, ent.counts, act); +} \ No newline at end of file diff --git a/src/llama-expert-tier.h b/src/llama-expert-tier.h new file mode 100644 index 000000000000..ec0161b0d7a0 --- /dev/null +++ b/src/llama-expert-tier.h @@ -0,0 +1,77 @@ +#pragma once + +#include + +#include "ggml.h" + +// Expert tier hook: drop-in replacement for ggml_mul_mat_id on expert weight +// tensors that have a registered GPU hot store. Pure stock ggml ops, no +// custom kernels. +// +// A registered expert tensor w is split between a GPU hot store (the top-S +// experts, held in dst_hot with hot_s+1 slot planes) and the CPU cold store +// (the remaining experts, still inside w). build_lora_mm_id calls back into +// llama_expert_tier_build, which computes: +// - hot: expert ids remapped through hot_lut to slot indices, then a +// mul_mat_id on dst_hot. Cold experts land on the zeroed sentinel +// plane (index hot_s) and therefore contribute zero on the GPU. +// - cold: mul_mat_id_cold on w, which skips hot experts entirely +// (cold_mask[e] == 0) and computes only the cold-selected rows. +// - result = hot + cold. +// The result has the same shape as a stock mul_mat_id output and feeds +// straight back into the caller's downstream ops. +// +// The per-expert quant scale w_s is discarded on the tiered path. It is an +// intentional approximation: applying it would add get_rows/mul nodes per +// layer, and the scale factors are close to 1. + +// register one expert weight tensor -> its per-device GPU hot tensors and +// per-device LUTs. called by llama_expert_hotstore::allocate() after creating +// dst_hot/hot_lut and the cold_mask tensors. multiple entries per layer share +// the same luts[i]. +void llama_expert_tier_register(ggml_tensor * src, + const std::vector & dst_hot, + const std::vector & hot_lut, + const std::vector & mask_lut, + ggml_tensor * cold_mask, + ggml_tensor * counts); + +// drop the entire table (called by hotstore destructor) +void llama_expert_tier_clear(); + +// cheap check: is `w` registered? (used so callers can short-circuit lora) +bool llama_expert_tier_has(ggml_tensor * w); + +// drop-in hook called from build_lora_mm_id. Returns the combined +// hot+cold output tensor when `w` is registered; returns nullptr to let the +// caller fall back to stock ggml_mul_mat_id. +// ctx : graph context (ctx0 of the calling llm_graph_context) +// w : expert weight tensor, ne = [in, out, n_experts], 3d, ne[3]==1 +// cur : activation, ne = [in, 1, n_tokens], 3d +// ids : selected_experts, ne = [n_expert_used, n_tokens], 2d i32, REAL ids +// w_s : per-expert quant scale, ne = [n_experts], f32; ignored by the +// tiered path (see above) +ggml_tensor * llama_expert_tier_build(ggml_context * ctx, + ggml_tensor * w, + ggml_tensor * cur, + ggml_tensor * ids, + ggml_tensor * w_s); + +// fused cold path (GGML_OP_MOE_COLD), one call pair per MoE layer: +// begin_fused is called before the layer's expert matmuls and makes +// llama_expert_tier_build return hot-only results; end_fused (after the down +// matmul) returns the cold contribution tensor to add to the down result, or +// nullptr when the fused path is not active. gate/up/down must all be +// registered. act = 0 (silu, separate gate/up) or 1 (gelu, fused gate_up). +bool llama_expert_tier_begin_fused(ggml_tensor * gate_w, + ggml_tensor * up_w, + ggml_tensor * down_w, + ggml_tensor * ids); + +ggml_tensor * llama_expert_tier_end_fused(ggml_context * ctx, + ggml_tensor * gate_w, + ggml_tensor * up_w, + ggml_tensor * down_w, + ggml_tensor * x, + ggml_tensor * ids, + int32_t act); \ No newline at end of file diff --git a/src/llama-graph.cpp b/src/llama-graph.cpp index 2be3b75fb982..e8feea3ea3d0 100644 --- a/src/llama-graph.cpp +++ b/src/llama-graph.cpp @@ -1,5 +1,6 @@ #include "llama-graph.h" +#include "llama-expert-tier.h" #include "llama-impl.h" #include "llama-model.h" #include "llama-batch.h" @@ -1311,6 +1312,7 @@ void llm_graph_result::reset() { inputs.clear(); fused_nodes.clear(); + moe_sel_experts.clear(); buf_compute_meta.resize(ggml_tensor_overhead()*max_nodes + ggml_graph_overhead_custom(max_nodes, false)); @@ -1519,6 +1521,14 @@ ggml_tensor * llm_graph_context::build_lora_mm_id( ggml_tensor * cur, // ggml_tensor * b ggml_tensor * ids, ggml_tensor * w_s) const { + // expert tier hook: if `w` has a registered GPU hot store, build the + // dual-path (hot GPU slots + cold CPU experts) result and return it. + // Skip when loras are active so build_lora_mm_id's lora loop still runs. + if (loras->empty()) { + if (auto * r = llama_expert_tier_build(ctx0, w, cur, ids, w_s)) { + return r; + } + } ggml_tensor * res = ggml_mul_mat_id(ctx0, w, cur, ids); if (w_s) { @@ -2029,9 +2039,12 @@ ggml_tensor * llm_graph_context::build_moe_ffn( if (selected_experts == nullptr) { selected_experts = ggml_argsort_top_k(ctx0, selection_probs, n_expert_used); // [n_expert_used, n_tokens] cb(selected_experts->src[0], "ffn_moe_argsort", il); + ggml_set_output(selected_experts); } cb(selected_experts, "ffn_moe_topk", il); + res->moe_sel_experts.emplace_back(il, selected_experts); + if (arch == LLM_ARCH_GROVEMOE && n_expert != hparams.n_expert) { // TODO: Use scalar div instead when/if implemented ggml_tensor * f_sel = ggml_cast(ctx0, selected_experts, GGML_TYPE_F32); @@ -2084,6 +2097,25 @@ ggml_tensor * llm_graph_context::build_moe_ffn( cb(cur, "ffn_moe_weighted", il); } + // expert tiering: fused cold path. when active, build_lora_mm_id returns + // hot-only results and the fused op (built after the down matmul so the + // CPU cold work overlaps the GPU hot work) covers the cold experts. + ggml_tensor * gw = gate_exps ? gate_exps : gate_up_exps; + ggml_tensor * uw = up_exps ? up_exps : gate_up_exps; + constexpr float clamp_eps = 1e-6f; + const bool swiglu_clamped = (type_op == LLM_FFN_SILU) && gate_exps && il >= 0 && + hparams.swiglu_clamp_exp[il] > clamp_eps; + const bool cold_ok = !weight_before_ffn && gw && uw && down_exps && + !up_exps_b && !gate_exps_b && !down_exps_b && !gate_up_exps_b && + !up_exps_s && !gate_exps_s && + (type_op == LLM_FFN_SILU || type_op == LLM_FFN_GELU) && + arch != LLM_ARCH_STEP35 && + loras->empty() && !swiglu_clamped; + const int32_t act = type_op == LLM_FFN_GELU ? 1 : 0; + const bool moe_cold = cold_ok && + llama_expert_tier_begin_fused(gw, uw, down_exps, selected_experts); + ggml_tensor * x_in = cur; + ggml_tensor * up = nullptr; ggml_tensor * experts = nullptr; @@ -2217,6 +2249,20 @@ ggml_tensor * llm_graph_context::build_moe_ffn( cb(experts, "ffn_moe_down_scaled", il); } + if (moe_cold) { + ggml_tensor * cold = llama_expert_tier_end_fused(ctx0, gw, uw, down_exps, x_in, selected_experts, act); + if (cold) { + if (down_exps_s) { + ggml_tensor * s = ggml_reshape_3d(ctx0, down_exps_s, 1, down_exps_s->ne[0], 1); + s = ggml_repeat_4d(ctx0, s, 1, down_exps_s->ne[0], cold->ne[2], 1); + s = ggml_get_rows(ctx0, s, selected_experts); + cold = ggml_mul(ctx0, cold, s); + } + experts = ggml_add(ctx0, experts, cold); + cb(experts, "ffn_moe_down_cold", il); + } + } + if (down_exps_b) { experts = ggml_add_id(ctx0, experts, down_exps_b, selected_experts); cb(experts, "ffn_moe_down_biased", il); diff --git a/src/llama-graph.h b/src/llama-graph.h index 32d8d395aa45..a3a0f5addac9 100644 --- a/src/llama-graph.h +++ b/src/llama-graph.h @@ -909,6 +909,8 @@ class llm_graph_result { std::map t_sampled; std::map t_sampled_probs; + std::vector> moe_sel_experts; + std::vector inputs; std::vector fused_nodes; diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 71bc9f7ef0aa..1ed0fcdfc169 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -1,4 +1,5 @@ #include "llama-model-loader.h" +#include "llama-expert-preload.h" #include "ggml-alloc.h" #include "ggml.h" @@ -530,6 +531,7 @@ llama_model_loader::llama_model_loader( const llama_model_kv_override * param_overrides_p, const llama_model_tensor_buft_override * param_tensor_buft_overrides_p) : metadata(meta), set_tensor_data(set_tensor_data), set_tensor_data_ud(set_tensor_data_ud) { + llama_expert_preload::set_model_path(fname.c_str()); int trace = 0; if (getenv("LLAMA_TRACE")) { trace = atoi(getenv("LLAMA_TRACE")); @@ -1310,6 +1312,28 @@ struct ggml_tensor * llama_model_loader::create_tensor( struct ggml_tensor * tensor = ggml_dup_tensor(ctx, &t_meta); ggml_set_name(tensor, ggml_get_name(&t_meta)); + // expert tier + no-mmap: the exps data lives in our buffers (GPU store for + // the hot experts, host buffer for the cold). give the tensor a valid ghost + // buffer so the model allocation skips it and the graph stays satisfied; + // the tier and the cold-op hook never read the tensor's own data. + if (!use_mmap && llama_expert_preload::tier_will_engage()) { + int il = -1; + if (llama_expert_preload::is_exps(tn.str().c_str(), il)) { + static ggml_backend_buffer_t ghost = nullptr; + static size_t ghost_size = 0; + const size_t nbytes = ggml_nbytes(&t_meta); + if (!ghost || nbytes > ghost_size) { + if (ghost) { + ggml_backend_buffer_free(ghost); + } + ghost = ggml_backend_buft_alloc_buffer(ggml_backend_cpu_buffer_type(), nbytes); + ggml_backend_buffer_set_usage(ghost, GGML_BACKEND_BUFFER_USAGE_WEIGHTS); + ghost_size = nbytes; + } + ggml_backend_tensor_alloc(ghost, tensor, ggml_backend_buffer_get_base(ghost)); + } + } + if (duplicated) { size_data += ggml_nbytes(&t_meta); } else { @@ -1525,6 +1549,49 @@ bool llama_model_loader::load_all_data( ggml_backend_name(upload_backend)); } + // tier + no-mmap + manual -ehs: stream the startup batch (first S experts + // per layer) into a GPU store buffer so those slices never commit RAM. + if (!use_mmap && llama_expert_preload::tier_will_engage() && !bufs.empty()) { + size_t gpu_total = 0; + size_t cpu_total = 0; + int n_entries = 0; + int n_layers = 0; + int n_experts = 0; + for (const auto & [name, w] : weights_map) { + int il = -1; + if (llama_expert_preload::is_exps(name.c_str(), il)) { + const size_t plane = ggml_nbytes(w.tensor) / (size_t) w.tensor->ne[2]; + gpu_total += (size_t) (llama_expert_preload::get_slots() + 1) * plane; + cpu_total += (size_t) w.tensor->ne[2] * plane; + n_entries++; + n_layers = std::max(n_layers, il + 1); + n_experts = (int) w.tensor->ne[2]; + } + } + if (n_entries > 0 && gpu_total > 0) { + // room for the per-layer hot_lut (i32[n_experts]) and mask_lut + // (f32[slots+1]) tensors after the slots; 256-aligned so backends + // with minStorageBufferOffsetAlignment can bind them for get_rows + gpu_total = llama_expert_preload::store_total_bytes( + gpu_total, n_layers, n_experts, llama_expert_preload::get_slots()); + // the store should live in VRAM when possible, else fall back to a + // host buffer (slower but functional) + ggml_backend_buffer_type_t buft = nullptr; + for (const auto & [idx, b] : bufs) { + ggml_backend_buffer_type_t bt = ggml_backend_buffer_get_type(b); + ggml_backend_dev_t d = ggml_backend_buft_get_device(bt); + if (d && ggml_backend_dev_type(d) == GGML_BACKEND_DEVICE_TYPE_GPU) { + buft = bt; + break; + } + } + if (!buft) { + buft = ggml_backend_cpu_buffer_type(); + } + llama_expert_preload::begin(buft, gpu_total, cpu_total, n_entries); + } + } + for (struct ggml_tensor * cur = ggml_get_first_tensor(ctx); cur != NULL; cur = ggml_get_next_tensor(ctx, cur)) { const auto * weight = get_weight(ggml_get_name(cur)); if (weight == nullptr) { @@ -1572,12 +1639,40 @@ bool llama_model_loader::load_all_data( const auto & file = files.at(weight->idx); if (ggml_backend_buffer_is_host(cur->buffer)) { - file->seek(weight->offs, SEEK_SET); - file->read_raw(cur->data, n_size); - if (check_tensors) { - validation_result.emplace_back(std::async(std::launch::async, [cur, n_size] { - return std::make_pair(cur, ggml_validate_row_data(cur->type, cur->data, n_size)); - })); + int il = -1; + if (llama_expert_preload::tier_will_engage() && llama_expert_preload::is_exps(ggml_get_name(cur), il)) { + // stream the startup batch: first S slices to the GPU store, + // the rest (cold) into our host buffer (the model tensor's + // data is only a placeholder; the cold-op hook reads ours) + const size_t plane = n_size / (size_t) cur->ne[2]; + const int S = llama_expert_preload::get_slots(); + if (getenv("LLAMA_EXPERT_DEBUG")) { + fprintf(stderr, "loader: exps %s off=%lld n_size=%zu plane=%zu\n", + ggml_get_name(cur), (long long) weight->offs, n_size, plane); + } + std::vector buf(n_size); + file->seek(weight->offs, SEEK_SET); + file->read_raw(buf.data(), n_size); + const size_t idx = llama_expert_preload::register_tensor(cur, plane, (int) cur->ne[2], S, + weight->offs, file->file_id()); + std::vector store_data((size_t) (S + 1) * plane, 0); + std::memcpy(store_data.data(), buf.data(), (size_t) S * plane); + llama_expert_preload::write_entry(idx, store_data.data(), store_data.size()); + llama_expert_preload::write_cold(idx, buf.data() + (size_t) S * plane, + (size_t) (cur->ne[2] - S) * plane); + if (check_tensors) { + validation_result.emplace_back(std::async(std::launch::async, [cur, n_size] { + return std::make_pair(cur, ggml_validate_row_data(cur->type, cur->data, n_size)); + })); + } + } else { + file->seek(weight->offs, SEEK_SET); + file->read_raw(cur->data, n_size); + if (check_tensors) { + validation_result.emplace_back(std::async(std::launch::async, [cur, n_size] { + return std::make_pair(cur, ggml_validate_row_data(cur->type, cur->data, n_size)); + })); + } } } else { // If upload_backend is valid load the tensor in chunks to pinned memory and upload the buffers asynchronously to the GPU. diff --git a/tools/cli/README.md b/tools/cli/README.md index 640d4fee80e8..57a65d8d1d46 100644 --- a/tools/cli/README.md +++ b/tools/cli/README.md @@ -64,6 +64,13 @@ | `-ot, --override-tensor =,...` | override tensor buffer type
(env: LLAMA_ARG_OVERRIDE_TENSOR) | | `-cmoe, --cpu-moe` | keep all Mixture of Experts (MoE) weights in the CPU
(env: LLAMA_ARG_CPU_MOE) | | `-ncmoe, --n-cpu-moe N` | keep the Mixture of Experts (MoE) weights of the first N layers in the CPU
(env: LLAMA_ARG_N_CPU_MOE) | +| `-ehs, --expert-hot-s N` | number of hot experts cached on the GPU (-1 = autofit, 0 = disabled)
(env: LLAMA_ARG_EXPERT_HOT_S) | +| `--ecf, --expert-cache-force` | enable the expert cache on non-CUDA backends (testing/emergency only) | +| `--expert-heat-decay N` | expert heatmap decay rate per update (default: 0.999)
(env: LLAMA_ARG_EXPERT_HEAT_DECAY) | +| `--expert-heat-log-period N` | expert heatmap log interval in updates (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD) | +| `--expert-sync-period N` | expert hot store re-sync cadence in tokens (default: 1)
(env: LLAMA_ARG_EXPERT_SYNC_PERIOD) | +| `--expert-hyst N` | expert hot store hysteresis ratio (default: 1.3, 0 = off)
(env: LLAMA_ARG_EXPERT_HYST) | +| `--expert-dwell N` | expert hot store minimum dwell updates before swap (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_DWELL) | | `-ngl, --gpu-layers, --n-gpu-layers N` | max. number of layers to store in VRAM, either an exact number, 'auto', or 'all' (default: auto)
(env: LLAMA_ARG_N_GPU_LAYERS) | | `-sm, --split-mode {none,layer,row,tensor}` | how to split the model across multiple GPUs, one of:
- none: use one GPU only
- layer (default): split layers and KV across GPUs (pipelined)
- row: split weight across GPUs by rows (parallelized)
- tensor: split weights and KV across GPUs (parallelized, EXPERIMENTAL)
(env: LLAMA_ARG_SPLIT_MODE) | | `-ts, --tensor-split N0,N1,N2,...` | fraction of the model to offload to each GPU, comma-separated list of proportions, e.g. 3,1
(env: LLAMA_ARG_TENSOR_SPLIT) | @@ -98,6 +105,22 @@ | `--spec-draft-type-k, -ctkd, --cache-type-k-draft TYPE` | KV cache data type for K for the draft model
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_SPEC_DRAFT_CACHE_TYPE_K) | | `--spec-draft-type-v, -ctvd, --cache-type-v-draft TYPE` | KV cache data type for V for the draft model
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_SPEC_DRAFT_CACHE_TYPE_V) | +### Expert tier params + +| Argument | Explanation | +| -------- | ----------- | +| `-ehs, --expert-hot-s N` | expert hot store slots: -1 = autofit from free VRAM, 0 = disabled, N = manual top-N slots
(env: LLAMA_ARG_EXPERT_HOT_S) | +| `--expert-pin N` | fraction (percent) of cold experts to keep pinned in RAM via madvise, 0 = off, -1 = auto
(env: LLAMA_ARG_EXPERT_PIN) | +| `--expert-move-mode N` | expert store mode: 0 = auto, 1 = copy (keep RAM copy), 2 = move (free RAM after verified transfer)
(env: LLAMA_ARG_EXPERT_MOVE_MODE) | +| `--expert-sidecar` | load the expert heatmap sidecar (`.tier`) at start, save it at exit
(env: LLAMA_ARG_EXPERT_SIDECAR) | +| `--expert-gpu N` | put the expert store on this GPU index (default: -1 = all GPUs)
(env: LLAMA_ARG_EXPERT_GPU) | +| `--expert-sync-period N` | expert hot store re-sync cadence in tokens (default: 1)
(env: LLAMA_ARG_EXPERT_SYNC_PERIOD) | +| `--expert-hyst F` | expert hot store hysteresis ratio (default: 1.3, 0 = off)
(env: LLAMA_ARG_EXPERT_HYST) | +| `--expert-dwell N` | expert hot store minimum dwell updates before swap (default: 0 = off)
(env: LLAMA_ARG_EXPERT_DWELL) | +| `--expert-heat-decay F` | expert heatmap decay rate per update (default: 0.999)
(env: LLAMA_ARG_EXPERT_HEAT_DECAY) | +| `--expert-heat-log-period N` | print the expert heatmap at generation end (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD) | +| `--expert-no-evict` | never evict experts from the hot store (fill-only, no move-back) | + ### Sampling params diff --git a/tools/completion/README.md b/tools/completion/README.md index e0923ea3005e..7628a82f0776 100644 --- a/tools/completion/README.md +++ b/tools/completion/README.md @@ -147,6 +147,13 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1 | `-ot, --override-tensor =,...` | override tensor buffer type
(env: LLAMA_ARG_OVERRIDE_TENSOR) | | `-cmoe, --cpu-moe` | keep all Mixture of Experts (MoE) weights in the CPU
(env: LLAMA_ARG_CPU_MOE) | | `-ncmoe, --n-cpu-moe N` | keep the Mixture of Experts (MoE) weights of the first N layers in the CPU
(env: LLAMA_ARG_N_CPU_MOE) | +| `-ehs, --expert-hot-s N` | number of hot experts cached on the GPU (-1 = autofit, 0 = disabled)
(env: LLAMA_ARG_EXPERT_HOT_S) | +| `--ecf, --expert-cache-force` | enable the expert cache on non-CUDA backends (testing/emergency only) | +| `--expert-heat-decay N` | expert heatmap decay rate per update (default: 0.999)
(env: LLAMA_ARG_EXPERT_HEAT_DECAY) | +| `--expert-heat-log-period N` | expert heatmap log interval in updates (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD) | +| `--expert-sync-period N` | expert hot store re-sync cadence in tokens (default: 1)
(env: LLAMA_ARG_EXPERT_SYNC_PERIOD) | +| `--expert-hyst N` | expert hot store hysteresis ratio (default: 1.3, 0 = off)
(env: LLAMA_ARG_EXPERT_HYST) | +| `--expert-dwell N` | expert hot store minimum dwell updates before swap (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_DWELL) | | `-ngl, --gpu-layers, --n-gpu-layers N` | max. number of layers to store in VRAM, either an exact number, 'auto', or 'all' (default: auto)
(env: LLAMA_ARG_N_GPU_LAYERS) | | `-sm, --split-mode {none,layer,row,tensor}` | how to split the model across multiple GPUs, one of:
- none: use one GPU only
- layer (default): split layers and KV across GPUs (pipelined)
- row: split weight across GPUs by rows (parallelized)
- tensor: split weights and KV across GPUs (parallelized, EXPERIMENTAL)
(env: LLAMA_ARG_SPLIT_MODE) | | `-ts, --tensor-split N0,N1,N2,...` | fraction of the model to offload to each GPU, comma-separated list of proportions, e.g. 3,1
(env: LLAMA_ARG_TENSOR_SPLIT) | diff --git a/tools/server/README.md b/tools/server/README.md index 64f0b03269d3..908fc3003d5b 100644 --- a/tools/server/README.md +++ b/tools/server/README.md @@ -115,6 +115,22 @@ For the full list of features, please refer to [server's changelog](https://gith | `--spec-draft-type-k, -ctkd, --cache-type-k-draft TYPE` | KV cache data type for K for the draft model
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_SPEC_DRAFT_CACHE_TYPE_K) | | `--spec-draft-type-v, -ctvd, --cache-type-v-draft TYPE` | KV cache data type for V for the draft model
allowed values: f32, f16, bf16, q8_0, q4_0, q4_1, iq4_nl, q5_0, q5_1
(default: f16)
(env: LLAMA_ARG_SPEC_DRAFT_CACHE_TYPE_V) | +### Expert tier params + +| Argument | Explanation | +| -------- | ----------- | +| `-ehs, --expert-hot-s N` | expert hot store slots: -1 = autofit from free VRAM, 0 = disabled, N = manual top-N slots
(env: LLAMA_ARG_EXPERT_HOT_S) | +| `--expert-pin N` | fraction (percent) of cold experts to keep pinned in RAM via madvise, 0 = off, -1 = auto
(env: LLAMA_ARG_EXPERT_PIN) | +| `--expert-move-mode N` | expert store mode: 0 = auto, 1 = copy (keep RAM copy), 2 = move (free RAM after verified transfer)
(env: LLAMA_ARG_EXPERT_MOVE_MODE) | +| `--expert-sidecar` | load the expert heatmap sidecar (`.tier`) at start, save it at exit
(env: LLAMA_ARG_EXPERT_SIDECAR) | +| `--expert-gpu N` | put the expert store on this GPU index (default: -1 = all GPUs)
(env: LLAMA_ARG_EXPERT_GPU) | +| `--expert-sync-period N` | expert hot store re-sync cadence in tokens (default: 1)
(env: LLAMA_ARG_EXPERT_SYNC_PERIOD) | +| `--expert-hyst F` | expert hot store hysteresis ratio (default: 1.3, 0 = off)
(env: LLAMA_ARG_EXPERT_HYST) | +| `--expert-dwell N` | expert hot store minimum dwell updates before swap (default: 0 = off)
(env: LLAMA_ARG_EXPERT_DWELL) | +| `--expert-heat-decay F` | expert heatmap decay rate per update (default: 0.999)
(env: LLAMA_ARG_EXPERT_HEAT_DECAY) | +| `--expert-heat-log-period N` | print the expert heatmap at generation end (default: 0, 0 = off)
(env: LLAMA_ARG_EXPERT_HEAT_LOG_PERIOD) | +| `--expert-no-evict` | never evict experts from the hot store (fill-only, no move-back) | + ### Sampling params diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp index 38d2e5c7a05a..78482be9e927 100644 --- a/tools/server/server-context.cpp +++ b/tools/server/server-context.cpp @@ -3170,6 +3170,7 @@ struct server_context_impl { SLT_WRN(slot, "%s", "empty prompt - releasing slot\n"); slot.print_timings(); + llama_print_expert_heatmap(slot.ctx_tgt); send_final_response(slot); slot.release(); @@ -3853,6 +3854,7 @@ struct server_context_impl { if (!process_token(result, slot)) { // release slot because of stop condition slot.print_timings(); + llama_print_expert_heatmap(slot.ctx_tgt); send_final_response(slot); metrics.on_prediction(slot); slot.release();