tl;dr: Big sparse MoE models like GLM-5 are surprisingly usable even with significant fractions of the weights on SSD, due to caching dynamics.

I've recently been experimenting with MoE expert weights partially spilling over to SSD. It's much better than I expected. I realized my previous intuition for weights on SSD ("sounds painfully slow") came from analogy to GPU->CPU spillover with dense models, where Amdahl's Law guarantees things will be pretty bad. However, MoE experts on SSD is a very different situation.

Two factors make it work really well. First, you don't have a chunk of weights permanently living on SSD. Rather, you have an mmap'd file paging chunks of the model in/out of RAM, based on the kernel's caching logic, I think basically LRU. If an expert on SSD is not used by a token, then for that token it doesn't matter that you spilled that expert onto SSD. Second, experts' hotness follows a power law distribution, which is very friendly to caching.

Specifically, I recorded activations of GLM-5 from a diverse set of prompts, ~100k tokens generated. I analyzed 128 token chunks within each response, looking at how many experts are used within a given chunk (U), and how many used experts tend to carry over from chunk to chunk - or rather, how many don't (D). The idea: if you have enough RAM to hold at least U experts at once, then every chunk you have to page in no more than D. If you have a shortfall of S (i.e. can only fit U - S experts), then you have to page in D + S. Particularly unfortunate access patterns could make that worse (more likely the less RAM you have, I bet), but for these purposes I think it's ok to simplify. Anyways, for GLM-5 I found U to be 13190 (69.6% of all GLM-5 experts) and D to be 2412. If down/up/gate are all IQ4_XS, that's 19.125MiB per expert, so U is 246 GiB and D is 45GiB. D, and your machine's RAM's shortfall from U, are how much latency you will pay every 128 tokens from SSD reads; beyond that it's just the same latency as if your machine had infinite RAM. So with 256GB RAM and a 7GB/s SSD, the SSD penalty averages 45/(128*7) = 50ms per token.

I wrote a script following this logic to estimate token generation speed for a given hardware config. Can't guarantee I'm not missing something, but at least it reasonably accurately predicted my own machine's performance (just under 3tok/s on GLM-5 with IQ4_XS experts; 3xP40, 288GiB of 2400 DDR4, ~3.5GB/s SSD). It suggests that a 3090 + 128GiB DDR5 (and 7GB/s SSD) should get 3.5tok/s with IQ4_XS down, IQ3_S gate+up, or over 6tok/s if increasing SSD bandwidth to 28GB/s (4 drives in RAID0). Pretty good for a 3.7-bit quant of a ~800B model on 152GB VRAM+RAM!

I think this logic is significantly over-optimistic when RAM is really low, though. I modeled it as, for every expert of the average used expert set size that you can't keep in RAM, you have to load that expert once per 128 tokens. But that's clearly wrong if you imagine 0GB RAM: there you have to load every active expert every token. I think this estimation transitions from wildly optimistic at very low RAM amounts, to basically perfect when you have enough RAM to hold the typical expert set size (S=0, in my terminology). Not sure what shape that transition has, though.

Side notes

Not all layers are equal for this caching logic. The more uniform a layer's activations, the worse for caching. I thought to identify these to prioritize GPU offloading, but it turns out to be pretty consistently the earliest layers with the most uniform activations. So nothing fancy needed; simply offload the first few layers. You will help yourself beyond what would be expected from just upgrading a small chunk of weights from RAM to VRAM bandwidths, by taking pressure off of your SSD.

I believe the SSD friendliness described here should increase as sparsity increases. GLM-5's 8-of-256 is pretty sparse. I think you would not haved wanted to try SSD spillover with Mixtral 8x22B and its 2-of-8. Maybe the lower sparsity of those early open MoE models contributed to all of this not being more common knowledge.

The Smart Expert Reduction feature of ik_llama.cpp could help here. You would add a condition to never skip an already-paged-in expert, and maybe probabilistically page-in/skip based on probability relative to top expert to avoid colder experts being paged out forever.

Code

Here is what I used for this investigation. You can use estimate_speed.py as-is to estimate any hardware's speed on GLM-5 at any quant. You could use the llama.cpp instrumentation patch + analyze_experts.py to gather the necessary --average_active_per_batch and --average_expert_turnover_per_batch parameters to use estimate_speed.py on other models.

The llama.cpp instrumentation for collecting activation data through stdout:

diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c
index 7486acc2b..6c04b70bc 100644
--- a/ggml/src/ggml-cpu/ggml-cpu.c
+++ b/ggml/src/ggml-cpu/ggml-cpu.c
@@ -1607,6 +1607,13 @@ static void ggml_compute_forward_mul_mat_id(

                 MMID_MATRIX_ROW(i02, matrix_row_counts[i02]) = (struct mmid_row_mapping) {id, iid1};
                 matrix_row_counts[i02] += 1;
+
+                if (strstr(src0->name, "up"))
+                {
+                    int layer_id = -1;
+                    sscanf(src0->name, "blk.%d.", &layer_id);
+                    printf("layer %d expert %d\n", layer_id, i02);
+                }
             }
         }
     }
diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu
index 07b10167b..2a56e03b6 100644
--- a/ggml/src/ggml-cuda/mmvq.cu
+++ b/ggml/src/ggml-cuda/mmvq.cu
@@ -1111,6 +1111,21 @@ void ggml_cuda_mul_mat_vec_q(

     const int64_t ids_stride = ids ? ids->nb[1] / ggml_type_size(ids->type) : 0;

+    if (ids_d)
+    {
+        int32_t host_ids[ids->ne[0]];
+        cudaMemcpyAsync(host_ids, ids_d, ids->ne[0] * sizeof(int32_t), cudaMemcpyDeviceToHost, stream);
+        cudaStreamSynchronize(stream);
+        if (strstr(src0->name, "up"))
+        {
+            int layer_id = -1;
+            sscanf(src0->name, "blk.%d.", &layer_id);
+            for (int i = 0; i < ids->ne[0]; ++i)
+                printf("layer %d expert %d\n", layer_id, host_ids[i]);
+            fflush(stdout);
+        }
+    }
+
     mul_mat_vec_q_switch_type(
         src0->data, src0->type, src1_q8_1.get(), ids_d, fusion_local, dst_d, ne00,
         ne01,              ncols_dst,     s01, stride_col_y,     stride_col_dst,
diff --git a/tools/server/server-context.cpp b/tools/server/server-context.cpp
index 6f737d94d..b7e2e6f84 100644
--- a/tools/server/server-context.cpp
+++ b/tools/server/server-context.cpp
@@ -1967,6 +1967,7 @@ private:

             if (all_idle) {
                 SRV_INF("%s", "all slots are idle\n");
+                printf("\n"); fflush(stdout); // otherwise last few printf's from CUDA don't come out

                 return;
             }
@@ -2160,6 +2161,7 @@ private:
                         slot.t_start_generation = 0;

                         slot.state = SLOT_STATE_PROCESSING_PROMPT;
+                        printf("---began prompt processing---\n");

                         SLT_INF(slot, "new prompt, n_ctx_slot = %d, n_keep = %d, task.n_tokens = %d\n",
                                 slot.n_ctx, slot.task->params.n_keep, slot.task->n_tokens());
@@ -2819,6 +2821,7 @@ private:

                     // prompt evaluated for next-token prediction
                     slot.state = SLOT_STATE_GENERATING;
+                    printf("---ended prompt processing---\n");

                     if (slot.can_speculate()) {
                         common_speculative_begin(slot.spec, slot.prompt.tokens.get_text_tokens());

analyze_experts.py

import glob
from collections import Counter

def process_chunk(flat_experts, thresholds):
    """
    Given a flat list of expert activations for a 128-token chunk,
    return the minimum K experts needed to hit each threshold,
    and the set of those top-K experts.
    """
    if not flat_experts:
        return None, None

    counts = Counter(flat_experts)
    total_activations = len(flat_experts)

    # Sort experts by frequency descending
    sorted_experts = counts.most_common()

    k_dict = {}
    set_dict = {}

    cum_sum = 0
    for i, (expert, count) in enumerate(sorted_experts):
        cum_sum += count

        # Check if we've crossed any thresholds we haven't recorded yet
        for t in thresholds:
            if cum_sum >= total_activations * t and t not in k_dict:
                k_dict[t] = i + 1  # +1 because i is 0-indexed
                # Store the set of experts up to this point
                set_dict[t] = {e for e, _ in sorted_experts[:i+1]}

        # Early exit: if we've found all thresholds, stop looping
        if len(k_dict) == len(thresholds):
            break

    return k_dict, set_dict

def main():
    files = sorted(glob.glob('single_response_*.txt'))
    if not files:
        print("No single_response_*.txt files found in the current directory.")
        return

    thresholds = [0.10, 0.25, 0.50, 0.80, 0.95, 0.99, 0.999, 1.0]

    # Accumulators for averages
    sum_k = {t: 0.0 for t in thresholds}
    count_k = {t: 0 for t in thresholds}

    sum_delta = {t: 0.0 for t in thresholds}
    count_delta = {t: 0 for t in thresholds}

    print(f"Processing {len(files)} files...")

    for fpath in files:
        with open(fpath, 'r', encoding='utf-8') as f:
            current_token_experts = []
            all_tokens = []

            for line in f:
                line = line.strip()
                if 'prompt' in line:
                    raise ValueError('you must first remove all ---began/ended prompt processing--- blocks')
                if line.startswith('---token'):
                    if current_token_experts:
                        all_tokens.append(current_token_experts)
                        current_token_experts = []
                elif line:
                    # Line is like "layer 7 expert 188"
                    current_token_experts.append(line)

            # Catch the last token if the file doesn't end with a separator
            if current_token_experts:
                all_tokens.append(current_token_experts)

        # Chunk into exactly 128 tokens, discard remainder
        num_chunks = len(all_tokens) // 128
        chunks = [all_tokens[i*128 : (i+1)*128] for i in range(num_chunks)]

        chunk_results = []
        for chunk in chunks:
            # Flatten the 128 tokens into a single list of all expert activations in the chunk
            flat_experts = [exp for token in chunk for exp in token]
            k_dict, set_dict = process_chunk(flat_experts, thresholds)

            if k_dict:
                chunk_results.append((k_dict, set_dict))

        # Accumulate K averages
        for k_dict, _ in chunk_results:
            for t in thresholds:
                sum_k[t] += k_dict[t]
                count_k[t] += 1

        # Accumulate Delta averages (Turnover)
        # Compare chunk i to chunk i+1 ONLY within the same file
        for i in range(len(chunk_results) - 1):
            _, set_prev = chunk_results[i]
            _, set_curr = chunk_results[i+1]

            for t in thresholds:
                # Experts in current chunk's hot-set that weren't in previous chunk's hot-set
                new_experts = set_curr[t] - set_prev[t]
                sum_delta[t] += len(new_experts)
                count_delta[t] += 1

    # Print Results
    print("\n" + "="*80)
    print(f"{'Threshold':<12} | {'Avg Active Experts (K)':<25} | {'Avg New Experts (Delta)':<25}")
    print("-" * 80)

    for t in thresholds:
        avg_k = sum_k[t] / count_k[t] if count_k[t] > 0 else 0
        avg_d = sum_delta[t] / count_delta[t] if count_delta[t] > 0 else 0

        # Format percentages nicely
        t_str = f"{t*100:>6.1f}%"
        if t == 0.999:
            t_str = " 99.9%"

        print(f"{t_str:<12} | {avg_k:<25.2f} | {avg_d:<25.2f}")

    print("="*80)

if __name__ == '__main__':
    main()

estimate_speed.py

import argparse

def parse_arguments():
    parser = argparse.ArgumentParser(description='Calculate token throughput for MoE model inference')

    # machine specs
    parser.add_argument('--ram_bytes', type=int, default=128*1024*1024*1024,
                        help='RAM size in bytes (default: 128GB)')
    parser.add_argument('--memory_bw_bytes_per_sec', type=float, default=100000000000,
                        help='Memory bandwidth in bytes/sec (default: 100 GB/s for dual-channel 6400MT/s DDR5)')
    parser.add_argument('--ssd_bw_bytes_per_sec_raw', type=float, default=7000000000,
                        help='Raw SSD bandwidth in bytes/sec (default: 7 GB/s)')
    parser.add_argument('--ssds_in_raid', type=int, default=1,
                        help='Number of SSDs in RAID configuration (default: 1)')
    parser.add_argument('--dense_gpu_latency_per_token', type=float, default=0.022,
                        help='GPU latency for dense layers in seconds/token (default: estimating 0.022s for 20GB through a 3090)')

    # model+quant size
    parser.add_argument('--bytes_per_expert', type=float, default=20054016,
                        help='Bytes per expert (default: 20,054,016 for GLM-5 IQ4_XS down+up+gate)')
    parser.add_argument('--num_layers', type=int, default=74,
                        help='Number of layers (default: 74 for GLM-5)')
    parser.add_argument('--experts_per_layer', type=int, default=256,
                        help='Experts per layer (default: 256 for GLM-5)')
    parser.add_argument('--active_per_layer', type=int, default=8,
                        help='Active experts per layer (default: 8 for GLM-5)')

    # expert dynamics
    parser.add_argument('--average_active_per_batch', type=float, default=13190,
                        help='Average active experts per batch (default: 13190, measured for GLM-5, 128 token batch)')
    parser.add_argument('--average_expert_turnover_per_batch', type=float, default=2412,
                        help='Average expert turnover per batch (default: 2412, measured for GLM-5, 128 token batch)')

    # how my evaluation script batches tokens, semi-arbitrary, meant to represent "the model thinking in a given 'mode'"
    parser.add_argument('--batch_size', type=int, default=128, help='Batch size in tokens (default: 128)')
    return parser.parse_args()

def main():
    args = parse_arguments()

    # Calculate SSD bandwidth with RAID
    ssd_bw_bytes_per_sec = args.ssds_in_raid * args.ssd_bw_bytes_per_sec_raw

    total_experts = args.num_layers * args.experts_per_layer

    # Calculate RAM shortfall
    per_batch_ram_shortfal_experts = args.average_active_per_batch - args.ram_bytes / args.bytes_per_expert
    if per_batch_ram_shortfal_experts < 0:
        per_batch_ram_shortfal_experts = 0

    # Calculate paging and SSD latency
    experts_paged_in_per_batch = args.average_expert_turnover_per_batch + per_batch_ram_shortfal_experts
    bytes_paged_in_per_batch = experts_paged_in_per_batch * args.bytes_per_expert
    ssd_latency_per_batch = bytes_paged_in_per_batch / ssd_bw_bytes_per_sec
    ssd_latency_per_token = ssd_latency_per_batch / args.batch_size

    # Calculate memory latency
    expert_bytes_through_memory_per_token = args.active_per_layer * args.num_layers * args.bytes_per_expert
    memory_latency_per_token = expert_bytes_through_memory_per_token / args.memory_bw_bytes_per_sec

    # Calculate total latency and throughput
    latency_per_token = args.dense_gpu_latency_per_token + memory_latency_per_token + ssd_latency_per_token
    tokens_per_second = 1.0 / latency_per_token

    print(f'{tokens_per_second} tok/s')

if __name__ == '__main__':
    main()
Edit

Pub: 11 Apr 2026 19:48 UTC

Edit: 11 Apr 2026 20:43 UTC

Views: 258