Offloading work

Author: kmbandy

common/arg.cpp

@@ -2431,9 +2431,9 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         }
     ).set_env("LLAMA_ARG_N_GPU_LAYERS"));
     add_opt(common_arg(
-        {"--weight-paging"}, "",
+        {"--weight-paging"},
         "enable NVMe→VRAM demand paging for model weights (allows models larger than VRAM)",
-        [](common_params & params, const std::string & value) {
+        [](common_params & params) {
             params.weight_paging_enabled = true;
         }
     ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_WEIGHT_PAGING"));
@@ -2445,9 +2445,9 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
         }
     ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_WEIGHT_PAGING_SLOTS"));
     add_opt(common_arg(
-        {"--weight-paging-prefetch"}, "",
+        {"--weight-paging-prefetch"},
         "enable async prefetch of next layer (default: enabled)",
-        [](common_params & params, const std::string & value) {
+        [](common_params & params) {
             params.weight_paging_prefetch = true;
         }
     ).set_examples({LLAMA_EXAMPLE_SERVER, LLAMA_EXAMPLE_CLI}).set_env("LLAMA_ARG_WEIGHT_PAGING_PREFETCH"));

common/common.cpp

@@ -1343,7 +1343,7 @@ common_init_result_ptr common_init_from_params(common_params & params) {
         common_set_adapter_lora(lctx, params.lora_adapters);
     }
 
-    if (params.warmup) {
+    if (params.warmup && !params.weight_paging_enabled) {
         LOG_WRN("%s: warming up the model with an empty run - please wait ... (--no-warmup to disable)\n", __func__);
 
         llama_set_warmup(lctx, true);
@@ -1488,6 +1488,11 @@ struct llama_model_params common_model_params_to_llama(common_params & params) {
     mparams.progress_callback_user_data = params.load_progress_callback_user_data;
     mparams.no_alloc                    = params.no_alloc;
 
+    // Weight paging parameters
+    mparams.weight_paging_enabled  = params.weight_paging_enabled;
+    mparams.weight_paging_slots    = params.weight_paging_slots;
+    mparams.weight_paging_prefetch = params.weight_paging_prefetch;
+
     return mparams;
 }
 

include/llama.h

@@ -321,6 +321,11 @@ extern "C" {
         bool use_extra_bufts; // use extra buffer types (used for weight repacking)
         bool no_host;         // bypass host buffer allowing extra buffers to be used
         bool no_alloc;        // only load metadata and simulate memory allocations
+
+        // Weight paging parameters (NVMe→VRAM demand paging)
+        bool weight_paging_enabled;    // enable NVMe→VRAM demand paging for model weights
+        int32_t weight_paging_slots;   // number of VRAM slots for weight paging (-1 = auto)
+        bool weight_paging_prefetch;   // enable async prefetch of next layer
     };
 
     struct llama_sampler_seq_config {

plans/errors/command.md

@@ -0,0 +1,28 @@
+Here is what I use to start the server:
+
+llama-server-feature \
+          --model ~/models/MiniMax-M2.7-UD-IQ3_S-00001-of-00003.gguf \
+          -ngl 999 \
+          --device ROCm0 \
+          --ctx-size 196608 \
+          --cache-type-k turbo4 \
+          --cache-type-v turbo4 \
+          --flash-attn on \
+          --metrics \
+          --port 8080 \
+          --host 0.0.0.0 \
+          --no-mmap \
+          --kv-tiered 12.5,50,37.5 \
+          --tier-ssd-path ~/kv-cold \
+          --tier-eviction-policy 3 \
+          --parallel 1 \
+          --kv-semantic-index ~/models/bge-small-en-v1.5-q8_0.gguf \
+          --kv-semantic-threshold 0.65 \
+          --kv-semantic-topk 5 \
+          --alias default \
+          --direct-io \
+          --kv-warm-device 1 \
+          --weight-paging \
+          --fit off \
+          --weight-paging-slots 8 \
+          --jinja

plans/errors/output.md

@@ -0,0 +1,35 @@
+  ---                                                                                                                                        
+  Phase 1 — Switch pool allocation + add pinned staging buffer                                                                               
+                                                                                                                                             
+  In llama-weight-pager.h, add a void* pinned_staging member to llama_weight_pager. Update the destructor to also call hipFree(pool.base) and
+   hipHostFree(pinned_staging).                                                                                                              
+   
+  In llama-weight-pager.cpp, init_pool: replace hipExtMallocWithFlags(&pool.base, size, hipDeviceMallocFinegrained) and its fallback with a  
+  single hipMalloc(&pool.base, size). After successful pool allocation, allocate the pinned staging buffer: hipHostMalloc(&pinned_staging, 
+  slot_size, hipHostMallocDefault). Return false if either allocation fails.                                                                 
+                  
+  Remove _mm_sfence() and the <immintrin.h> include — no longer needed.                                                                      
+   
+  ---                                                                                                                                        
+  Phase 2 — Rewrite page_in to use pinned staging + hipMemcpy
+                                                                                                                                             
+  In page_in in llama-weight-pager.cpp, replace the pread-directly-to-VRAM path with:
+  1. pread(fd, pinned_staging, size, offset) — read from NVMe into pinned host buffer                                                        
+  2. Check return value as before                                                                                                            
+  3. hipMemcpy(dst, pinned_staging, size, hipMemcpyHostToDevice) — DMA from pinned RAM to VRAM slot                                          
+  4. Keep all existing diagnostic logging, just update the log message to note the two-step path                                             
+                                                                                                                                             
+  This phase should get end-to-end inference working. Rebuild and test before moving to Phase 3.                                             
+                                                                                                                                             
+  ---                                                                                                                                        
+  Phase 3 — Async pipeline with N+1 prefetch
+                                                                                                                                             
+  Add a hipStream_t transfer_stream member to llama_weight_pager, created in init_pool with hipStreamCreate and destroyed in the destructor.
+                                                                                                                                             
+  Change hipMemcpy in page_in to hipMemcpyAsync(..., transfer_stream) followed by hipStreamSynchronize(transfer_stream) for the synchronous  
+  case.                                                                                                                                      
+                                                                                                                                             
+  Then wire the existing io_uring N+1 prefetch (already in the eval callback's Phase 4 block) to also kick off the GPU transfer leg: when    
+  complete_prefetch finishes the NVMe read into the pinned buffer, immediately issue hipMemcpyAsync to the pre-reserved next slot on
+  transfer_stream. This way the GPU transfer overlaps with GPU compute on the current tensor, and by the time the eval callback needs N+1,   
+  it's already in VRAM.