diff --git a/ds4_cuda.cu b/ds4_cuda.cu index aaa4df113..169665f66 100644 --- a/ds4_cuda.cu +++ b/ds4_cuda.cu @@ -149,6 +149,66 @@ typedef struct { static cuda_stream_selected_cache g_stream_selected_cache; +/* Resident SSD-streaming expert cache: keeps hot routed experts in VRAM so a + * decode step turns most selected-expert loads into device-to-device copies + * instead of host uploads. Caches are kept per (gate,down) byte-size class: + * mixed-quant models (e.g. q2-q4-imatrix GGUFs) interleave layers whose + * routed experts have different byte sizes, and a single cache keyed on the + * last-seen size would be released and rebuilt on every layer-size + * transition, never accumulating hits. */ +enum { + DS4_CUDA_STREAM_EXPERT_DEFAULT = 8u * 64u, + DS4_CUDA_STREAM_EXPERT_MAX = 61u * 384u +}; +#define DS4_CUDA_STREAM_EXPERT_CLASSES 4 + +struct cuda_stream_expert_cache_slot { + int valid; + const void *model_map; + uint64_t model_size; + uint32_t layer; + uint32_t n_total_expert; + uint32_t expert; + uint64_t gate_offset; + uint64_t up_offset; + uint64_t down_offset; + uint64_t gate_expert_bytes; + uint64_t down_expert_bytes; + uint64_t age; +}; + +struct cuda_stream_expert_cache { + int valid; + uint32_t capacity; + uint32_t count; + uint64_t tick; + uint64_t gate_expert_bytes; + uint64_t down_expert_bytes; + char *gate_ptr; + char *up_ptr; + char *down_ptr; + uint64_t gate_capacity; + uint64_t up_capacity; + uint64_t down_capacity; + std::vector slots; +}; + +static cuda_stream_expert_cache g_stream_expert_caches[DS4_CUDA_STREAM_EXPERT_CLASSES]; +static uint32_t g_stream_expert_budget_override; +static uint32_t g_stream_expert_runtime_caps[DS4_CUDA_STREAM_EXPERT_CLASSES]; +static uint32_t g_stream_expert_memory_cap_notices[DS4_CUDA_STREAM_EXPERT_CLASSES]; +static uint64_t g_stream_expert_class_gate_bytes[DS4_CUDA_STREAM_EXPERT_CLASSES]; +static uint64_t g_stream_expert_class_down_bytes[DS4_CUDA_STREAM_EXPERT_CLASSES]; +/* Requested cap the current cache allocation of each class settled at. + * When the VRAM-capped capacity ends up below the requested budget, the + * cheap "capacity >= target" fast path in prepare() can never hit, and + * every per-layer selected load would fall through to cudaMemGetInfo + * (~1ms per call — at 61 layers that is most of the decode step). + * Remembering the settled request lets steady-state loads return + * immediately. Cleared on release and re-evaluated when the requested + * budget changes. */ +static uint32_t g_stream_expert_settled_caps[DS4_CUDA_STREAM_EXPERT_CLASSES]; + static void cuda_stream_selected_cache_invalidate(void) { g_stream_selected_cache.valid = 0; } @@ -412,6 +472,11 @@ static void *g_stream_selected_stage[4]; static cudaEvent_t g_stream_selected_stage_event[4]; static uint64_t g_stream_selected_stage_bytes; static cudaStream_t g_stream_selected_upload_stream; +/* Staging-buffer cursor shared across grouped async copies: buffer reuse is + * fenced by the per-buffer events, so back-to-back async copies (e.g. the + * three tensors of one streamed expert) pipeline the host read of one + * tensor with the device upload of the previous one. */ +static uint64_t g_stream_selected_stage_cursor; static int cuda_ok(cudaError_t err, const char *what); static const char *cuda_model_range_ptr_from_fd( @@ -1354,6 +1419,15 @@ static void cuda_model_load_progress_note(uint64_t cached_bytes) { } } +static void cuda_model_load_progress_finish(void) { + if (!g_model_load_progress_started) return; + if (g_model_load_progress_tty) { + fputc('\n', stderr); + fflush(stderr); + } + g_model_load_progress_started = 0; +} + static int cuda_model_prefetch_range(const void *model_map, uint64_t model_size, uint64_t map_offset, uint64_t map_size) { if (!model_map || map_size == 0 || map_offset > model_size || map_size > model_size - map_offset) return 0; if (getenv("DS4_CUDA_NO_MODEL_PREFETCH") != NULL || @@ -1650,22 +1724,53 @@ static int cuda_stream_selected_stage_pool_alloc(uint64_t bytes) { } } g_stream_selected_stage_bytes = bytes; + g_stream_selected_stage_cursor = 0; return 1; } -static int cuda_model_copy_to_device_streamed( +static int cuda_model_copy_to_device_streamed_flush(const char *what) { + if (!g_stream_selected_upload_stream) return 1; + const cudaError_t err = + cudaStreamSynchronize(g_stream_selected_upload_stream); + if (err != cudaSuccess) { + fprintf(stderr, + "ds4: CUDA streaming selected upload sync failed for %s: %s\n", + what ? what : "expert", cudaGetErrorString(err)); + (void)cudaGetLastError(); + return 0; + } + return 1; +} + +static int cuda_model_copy_to_device_streamed_ex( char *dst, const void *model_map, uint64_t model_size, uint64_t offset, uint64_t bytes, - const char *what) { + const char *what, + int flush) { if (!dst || !model_map || offset > model_size || bytes > model_size - offset) { return 0; } if (bytes == 0) return 1; - if (g_model_fd < 0 || + /* Upload straight from the (pageable) model map when there is no + * O_DIRECT read path: for page-cache/tmpfs-backed maps the driver's + * pageable upload runs ~3x faster than the pread bounce (measured + * ~2.3ms -> ~0.8ms per 13.5 MiB expert, decode 23.5 -> 30 tok/s). + * Real direct-I/O SSD streaming keeps the staged pread path. + * DS4_CUDA_STREAM_MISS_PLAIN_COPY=1/0 forces either way. */ + static int plain_copy = -1; + if (plain_copy < 0) { + const char *env = getenv("DS4_CUDA_STREAM_MISS_PLAIN_COPY"); + if (env && env[0]) { + plain_copy = strcmp(env, "0") != 0; + } else { + plain_copy = g_model_direct_fd < 0 ? 1 : 0; + } + } + if (g_model_fd < 0 || plain_copy || (g_model_fd_host_base != NULL && model_map != g_model_fd_host_base)) { return cuda_ok(cudaMemcpy(dst, (const char *)model_map + offset, @@ -1674,18 +1779,32 @@ static int cuda_model_copy_to_device_streamed( what ? what : "stream selected expert copy"); } - const uint64_t chunk = cuda_model_copy_chunk_bytes(); + /* Small chunks let the host read of chunk i+1 overlap the device upload + * of chunk i through the 4-buffer ring; the 64 MiB model-load chunk + * would serialize each ~4.5 MiB expert tensor into one read + upload. */ + static uint64_t stream_chunk; + if (stream_chunk == 0) { + uint64_t mb = 2; + const char *env = getenv("DS4_CUDA_STREAM_COPY_CHUNK_MB"); + if (env && env[0]) { + char *end = NULL; + unsigned long long v = strtoull(env, &end, 10); + if (end != env && v > 0) mb = (uint64_t)v; + } + stream_chunk = mb * 1048576ull; + } + uint64_t chunk = cuda_model_copy_chunk_bytes(); + if (chunk > stream_chunk) chunk = stream_chunk; const uint64_t stage_bytes = chunk + (g_model_direct_align > 1 ? g_model_direct_align : 1); if (!cuda_stream_selected_stage_pool_alloc(stage_bytes)) return 0; uint64_t copied = 0; - uint64_t chunk_idx = 0; while (copied < bytes) { const uint64_t n = bytes - copied < chunk ? bytes - copied : chunk; - const uint64_t bi = chunk_idx % 4u; + const uint64_t bi = g_stream_selected_stage_cursor % 4u; cudaError_t err; - if (chunk_idx >= 4u) { + if (g_stream_selected_stage_cursor >= 4u) { err = cudaEventSynchronize(g_stream_selected_stage_event[bi]); if (err != cudaSuccess) { fprintf(stderr, @@ -1729,19 +1848,22 @@ static int cuda_model_copy_to_device_streamed( cuda_model_discard_source_pages(model_map, model_size, offset + copied, n); copied += n; - chunk_idx++; + g_stream_selected_stage_cursor++; } - const cudaError_t err = - cudaStreamSynchronize(g_stream_selected_upload_stream); - if (err != cudaSuccess) { - fprintf(stderr, - "ds4: CUDA streaming selected upload sync failed for %s: %s\n", - what ? what : "expert", cudaGetErrorString(err)); - (void)cudaGetLastError(); - return 0; - } - return 1; + if (!flush) return 1; + return cuda_model_copy_to_device_streamed_flush(what); +} + +static int cuda_model_copy_to_device_streamed( + char *dst, + const void *model_map, + uint64_t model_size, + uint64_t offset, + uint64_t bytes, + const char *what) { + return cuda_model_copy_to_device_streamed_ex(dst, model_map, model_size, + offset, bytes, what, 1); } static uint64_t cuda_model_cache_limit_bytes(void) { @@ -22950,10 +23072,690 @@ static int cuda_stream_selected_ranges_valid( down_bytes <= table->model_size - table->down_offset; } +static void cuda_stream_expert_cache_release_class(int class_idx) { + cuda_stream_expert_cache *cache = &g_stream_expert_caches[class_idx]; + if (cache->gate_ptr) { + (void)cudaFree(cache->gate_ptr); + } + if (cache->up_ptr) { + (void)cudaFree(cache->up_ptr); + } + if (cache->down_ptr) { + (void)cudaFree(cache->down_ptr); + } + std::vector().swap(cache->slots); + g_stream_expert_settled_caps[class_idx] = 0; + cache->valid = 0; + cache->capacity = 0; + cache->count = 0; + cache->tick = 0; + cache->gate_expert_bytes = 0; + cache->down_expert_bytes = 0; + cache->gate_ptr = NULL; + cache->up_ptr = NULL; + cache->down_ptr = NULL; + cache->gate_capacity = 0; + cache->up_capacity = 0; + cache->down_capacity = 0; +} + +static void cuda_stream_expert_cache_release_all(void) { + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + cuda_stream_expert_cache_release_class(i); + g_stream_expert_class_gate_bytes[i] = 0; + g_stream_expert_class_down_bytes[i] = 0; + g_stream_expert_runtime_caps[i] = 0; + g_stream_expert_memory_cap_notices[i] = 0; + } +} + +static void cuda_stream_expert_cache_invalidate(void) { + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + cuda_stream_expert_cache *cache = &g_stream_expert_caches[i]; + for (cuda_stream_expert_cache_slot &slot : cache->slots) { + slot.valid = 0; + } + cache->valid = 0; + cache->count = 0; + cache->tick = 0; + } +} + +static uint32_t cuda_stream_expert_cache_requested_budget(void) { + uint32_t cap = g_stream_expert_budget_override != 0 ? + g_stream_expert_budget_override : DS4_CUDA_STREAM_EXPERT_DEFAULT; + const char *env = getenv("DS4_CUDA_STREAMING_EXPERT_CACHE_N"); + if (env && env[0]) { + char *end = NULL; + errno = 0; + unsigned long v = strtoul(env, &end, 10); + while (end && (*end == ' ' || *end == '\t')) end++; + if (end != env && errno == 0 && end && *end == '\0') { + cap = v > DS4_CUDA_STREAM_EXPERT_MAX ? + DS4_CUDA_STREAM_EXPERT_MAX : (uint32_t)v; + } + } + if (cap > DS4_CUDA_STREAM_EXPERT_MAX) { + cap = DS4_CUDA_STREAM_EXPERT_MAX; + } + return cap; +} + +static uint32_t cuda_stream_expert_cache_configured_budget_class(int class_idx) { + uint32_t cap = cuda_stream_expert_cache_requested_budget(); + if (g_stream_expert_runtime_caps[class_idx] != 0 && + cap > g_stream_expert_runtime_caps[class_idx]) { + cap = g_stream_expert_runtime_caps[class_idx]; + } + return cap; +} + +/* Class-less view (reporting/shared queries): apply the most restrictive + * runtime cap across the active size classes. */ +static uint32_t cuda_stream_expert_cache_configured_budget(void) { + uint32_t cap = cuda_stream_expert_cache_requested_budget(); + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + if ((g_stream_expert_class_gate_bytes[i] != 0 || + g_stream_expert_class_down_bytes[i] != 0) && + g_stream_expert_runtime_caps[i] != 0 && + cap > g_stream_expert_runtime_caps[i]) { + cap = g_stream_expert_runtime_caps[i]; + } + } + return cap; +} + +/* The shared graph code sizes its expert working-set estimates on the cache + * budget; only report a budget there when the operator explicitly opted in, + * so default runs keep the upstream working-set behavior. */ +static int cuda_stream_expert_cache_budget_visible_to_shared(void) { + if (!g_ssd_streaming_mode) return 0; + if (g_stream_expert_budget_override != 0) return 1; + const char *env = getenv("DS4_CUDA_STREAMING_EXPERT_CACHE_N"); + if (env && env[0]) return 1; + env = getenv("DS4_CUDA_ENABLE_STREAMING_EXPERT_HOTLIST"); + if (!env || !env[0]) { + env = getenv("DS4_CUDA_STREAMING_EXPERT_HOTLIST"); + } + return env && env[0] && strcmp(env, "0") != 0; +} + +static uint64_t cuda_stream_expert_cache_reserve_bytes(void) { + uint64_t gb = 16; + const char *env = getenv("DS4_CUDA_STREAMING_EXPERT_CACHE_RESERVE_GB"); + if (env && env[0]) { + char *end = NULL; + errno = 0; + unsigned long long v = strtoull(env, &end, 10); + while (end && (*end == ' ' || *end == '\t')) end++; + if (end != env && errno == 0 && end && *end == '\0') { + gb = (uint64_t)v; + } + } + if (gb > UINT64_MAX / 1073741824ull) return UINT64_MAX; + return gb * 1073741824ull; +} + +static uint32_t cuda_stream_expert_cache_live_budget( + int class_idx, + uint32_t requested, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes, + uint64_t reclaim_bytes, + int report) { + if (requested == 0 || + gate_expert_bytes == 0 || + down_expert_bytes == 0 || + gate_expert_bytes > (UINT64_MAX - down_expert_bytes) / 2ull) { + return 0; + } + const uint64_t per_expert_bytes = + gate_expert_bytes * 2ull + down_expert_bytes; + if (per_expert_bytes == 0) return 0; + + size_t free_b = 0; + size_t total_b = 0; + cudaError_t err = cudaMemGetInfo(&free_b, &total_b); + if (err != cudaSuccess) { + fprintf(stderr, + "ds4: CUDA streaming expert cache memory query failed: %s; " + "using direct selected loads\n", + cudaGetErrorString(err)); + (void)cudaGetLastError(); + return 0; + } + + uint64_t free_bytes = (uint64_t)free_b; + if (reclaim_bytes > UINT64_MAX - free_bytes) { + free_bytes = UINT64_MAX; + } else { + free_bytes += reclaim_bytes; + } + uint64_t reserve = cuda_stream_expert_cache_reserve_bytes(); + const uint64_t total_bytes = (uint64_t)total_b; + if (total_bytes != 0 && reserve > total_bytes / 2ull) { + reserve = total_bytes / 2ull; + } + if (free_bytes <= reserve) { + if (report && g_stream_expert_memory_cap_notices[class_idx] != requested) { + cuda_model_load_progress_finish(); + fprintf(stderr, + "ds4: CUDA streaming expert cache disabled: available %.2f GiB <= reserve %.2f GiB\n", + (double)free_bytes / 1073741824.0, + (double)reserve / 1073741824.0); + g_stream_expert_memory_cap_notices[class_idx] = requested; + } + return 0; + } + + uint64_t usable = free_bytes - reserve; + uint64_t max_slots64 = usable / per_expert_bytes; + if (max_slots64 > UINT32_MAX) max_slots64 = UINT32_MAX; + uint32_t capped = requested; + if ((uint64_t)capped > max_slots64) capped = (uint32_t)max_slots64; + if (report && capped != requested && + g_stream_expert_memory_cap_notices[class_idx] != capped) { + cuda_model_load_progress_finish(); + fprintf(stderr, + "ds4: CUDA streaming expert cache capped from %u to %u experts " + "(available %.2f GiB, reserve %.2f GiB, %.2f MiB/expert)\n", + requested, + capped, + (double)free_bytes / 1073741824.0, + (double)reserve / 1073741824.0, + (double)per_expert_bytes / 1048576.0); + g_stream_expert_memory_cap_notices[class_idx] = capped; + } + return capped; +} + +static uint64_t cuda_stream_expert_cache_expert_bytes( + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes) { + if (gate_expert_bytes == 0 || + down_expert_bytes == 0 || + gate_expert_bytes > (UINT64_MAX - down_expert_bytes) / 2ull) { + return 0; + } + return gate_expert_bytes * 2ull + down_expert_bytes; +} + +/* Return the size-class index for (gate,down), assigning a free slot on + * first appearance. If distinct classes ever exceed + * DS4_CUDA_STREAM_EXPERT_CLASSES (never observed in official GGUFs), class 0 + * is recycled: that size falls back to rebuild-on-transition without + * touching the other classes. */ +static int cuda_stream_expert_cache_class_index( + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes) { + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + if (g_stream_expert_class_gate_bytes[i] == gate_expert_bytes && + g_stream_expert_class_down_bytes[i] == down_expert_bytes) { + return i; + } + } + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + if (g_stream_expert_class_gate_bytes[i] == 0 && + g_stream_expert_class_down_bytes[i] == 0) { + g_stream_expert_class_gate_bytes[i] = gate_expert_bytes; + g_stream_expert_class_down_bytes[i] = down_expert_bytes; + g_stream_expert_runtime_caps[i] = 0; + g_stream_expert_memory_cap_notices[i] = 0; + return i; + } + } + cuda_stream_expert_cache_release_class(0); + g_stream_expert_class_gate_bytes[0] = gate_expert_bytes; + g_stream_expert_class_down_bytes[0] = down_expert_bytes; + g_stream_expert_runtime_caps[0] = 0; + g_stream_expert_memory_cap_notices[0] = 0; + return 0; +} + +static uint32_t cuda_stream_expert_cache_shrunken_cap(uint32_t cap) { + if (cap == 0) return 0; + const uint32_t release = (cap + 9u) / 10u; + return cap > release ? cap - release : 0; +} + +static void cuda_stream_expert_cache_note_oom_cap( + int class_idx, + uint32_t failed_cap, + uint32_t new_cap, + uint64_t expert_bytes, + const char *errstr) { + if (g_stream_expert_runtime_caps[class_idx] != 0 && + g_stream_expert_runtime_caps[class_idx] <= new_cap) { + return; + } + g_stream_expert_runtime_caps[class_idx] = new_cap; + const uint32_t released = + failed_cap > new_cap ? failed_cap - new_cap : 0; + cuda_model_load_progress_finish(); + fprintf(stderr, + "ds4: CUDA streaming expert cache allocation failed at %u experts " + "/ %.2f GiB%s%s\n", + failed_cap, + expert_bytes != 0 ? + (double)((uint64_t)failed_cap * expert_bytes) / 1073741824.0 : + 0.0, + errstr && errstr[0] ? ": " : "", + errstr && errstr[0] ? errstr : ""); + if (new_cap != 0) { + fprintf(stderr, + "ds4: shrinking resident cache margin by %u experts / %.2f GiB; " + "runtime cache cap now %u experts\n", + released, + expert_bytes != 0 ? + (double)((uint64_t)released * expert_bytes) / 1073741824.0 : + 0.0, + new_cap); + } else { + fprintf(stderr, + "ds4: disabling resident expert cache after OOM; using direct selected loads\n"); + } +} + +static int cuda_stream_expert_cache_try_alloc( + uint32_t cap, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes, + char **gate_ptr, + char **up_ptr, + char **down_ptr, + const char **errstr) { + *gate_ptr = NULL; + *up_ptr = NULL; + *down_ptr = NULL; + if (errstr) *errstr = NULL; + if (cap == 0 || + (uint64_t)cap > UINT64_MAX / gate_expert_bytes || + (uint64_t)cap > UINT64_MAX / down_expert_bytes) { + return 0; + } + const uint64_t gate_bytes = (uint64_t)cap * gate_expert_bytes; + const uint64_t down_bytes = (uint64_t)cap * down_expert_bytes; + + void *gate = NULL; + void *up = NULL; + void *down = NULL; + cudaError_t err = cudaMalloc(&gate, (size_t)gate_bytes); + if (err != cudaSuccess) goto fail; + err = cudaMalloc(&up, (size_t)gate_bytes); + if (err != cudaSuccess) goto fail; + err = cudaMalloc(&down, (size_t)down_bytes); + if (err != cudaSuccess) goto fail; + + *gate_ptr = (char *)gate; + *up_ptr = (char *)up; + *down_ptr = (char *)down; + return 1; + +fail: + if (errstr) *errstr = cudaGetErrorString(err); + (void)cudaGetLastError(); + if (gate) (void)cudaFree(gate); + if (up) (void)cudaFree(up); + if (down) (void)cudaFree(down); + return 0; +} + +static cuda_stream_expert_cache *cuda_stream_expert_cache_prepare( + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes, + uint32_t target_cap) { + /* The resident cache assumes the single-GPU streaming placement (all + * buffers live on device 0, like the selected staging cache). */ + if (g_n_gpus != 1) return NULL; + const uint64_t expert_bytes = + cuda_stream_expert_cache_expert_bytes(gate_expert_bytes, + down_expert_bytes); + if (expert_bytes == 0) return NULL; + const int class_idx = + cuda_stream_expert_cache_class_index(gate_expert_bytes, + down_expert_bytes); + cuda_stream_expert_cache *cache = &g_stream_expert_caches[class_idx]; + + const uint32_t requested_cap = + cuda_stream_expert_cache_configured_budget_class(class_idx); + if (requested_cap == 0) return NULL; + if (target_cap == 0 || target_cap > requested_cap) target_cap = requested_cap; + if (target_cap == 0) return NULL; + const uint32_t entry_target = target_cap; + const int same_dims = + cache->valid && + cache->gate_expert_bytes == gate_expert_bytes && + cache->down_expert_bytes == down_expert_bytes; + if (!same_dims && cache->valid) { + cuda_stream_expert_cache_release_class(class_idx); + } + if (same_dims && + cache->capacity != 0 && + cache->slots.size() == cache->capacity && + (cache->capacity >= target_cap || + g_stream_expert_settled_caps[class_idx] == entry_target)) { + return cache; + } + + uint64_t reclaim_bytes = 0; + if (same_dims && + cache->capacity != 0 && + (uint64_t)cache->capacity <= UINT64_MAX / expert_bytes) { + reclaim_bytes = (uint64_t)cache->capacity * expert_bytes; + } + uint32_t cap = + cuda_stream_expert_cache_live_budget(class_idx, + target_cap, + gate_expert_bytes, + down_expert_bytes, + reclaim_bytes, + reclaim_bytes == 0); + if (cap == 0) return NULL; + if (same_dims && + cache->capacity != 0 && + cache->capacity >= cap && + cache->slots.size() == cache->capacity) { + g_stream_expert_settled_caps[class_idx] = entry_target; + return cache; + } + + cuda_stream_expert_cache_release_class(class_idx); + while (cap != 0) { + if ((uint64_t)cap > UINT64_MAX / gate_expert_bytes || + (uint64_t)cap > UINT64_MAX / down_expert_bytes) { + fprintf(stderr, "ds4: CUDA streaming expert cache size overflow\n"); + return NULL; + } + + char *gate_ptr = NULL; + char *up_ptr = NULL; + char *down_ptr = NULL; + const char *alloc_error = NULL; + if (!cuda_stream_expert_cache_try_alloc(cap, + gate_expert_bytes, + down_expert_bytes, + &gate_ptr, + &up_ptr, + &down_ptr, + &alloc_error)) { + const uint32_t new_cap = + cuda_stream_expert_cache_shrunken_cap(cap); + cuda_stream_expert_cache_note_oom_cap(class_idx, + cap, + new_cap, + expert_bytes, + alloc_error); + cap = new_cap; + if (cap != 0) { + cap = cuda_stream_expert_cache_live_budget(class_idx, + cap, + gate_expert_bytes, + down_expert_bytes, + 0, + 1); + } + continue; + } + + try { + cache->slots.resize(cap); + } catch (...) { + fprintf(stderr, "ds4: CUDA streaming expert cache metadata allocation failed\n"); + (void)cudaFree(gate_ptr); + (void)cudaFree(up_ptr); + (void)cudaFree(down_ptr); + cuda_stream_expert_cache_release_class(class_idx); + return NULL; + } + + cache->valid = 1; + cache->capacity = cap; + cache->count = 0; + cache->tick = 0; + cache->gate_expert_bytes = gate_expert_bytes; + cache->down_expert_bytes = down_expert_bytes; + cache->gate_ptr = gate_ptr; + cache->up_ptr = up_ptr; + cache->down_ptr = down_ptr; + cache->gate_capacity = + (uint64_t)cap * gate_expert_bytes; + cache->up_capacity = + (uint64_t)cap * gate_expert_bytes; + cache->down_capacity = + (uint64_t)cap * down_expert_bytes; + g_stream_expert_settled_caps[class_idx] = entry_target; + return cache; + } + return NULL; +} + +static int cuda_stream_expert_cache_find( + cuda_stream_expert_cache *cache, + const void *model_map, + uint64_t model_size, + uint32_t layer, + uint32_t n_total_expert, + uint32_t expert, + uint64_t gate_offset, + uint64_t up_offset, + uint64_t down_offset, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes) { + if (!cache || !cache->valid) return -1; + for (uint32_t i = 0; i < cache->capacity; i++) { + const cuda_stream_expert_cache_slot &slot = cache->slots[i]; + if (slot.valid && + slot.model_map == model_map && + slot.model_size == model_size && + slot.layer == layer && + slot.n_total_expert == n_total_expert && + slot.expert == expert && + slot.gate_offset == gate_offset && + slot.up_offset == up_offset && + slot.down_offset == down_offset && + slot.gate_expert_bytes == gate_expert_bytes && + slot.down_expert_bytes == down_expert_bytes) { + return (int)i; + } + } + return -1; +} + +static uint32_t cuda_stream_expert_cache_lru_slot( + cuda_stream_expert_cache *cache) { + for (uint32_t i = 0; i < cache->capacity; i++) { + if (!cache->slots[i].valid) return i; + } + uint32_t slot = 0; + uint64_t best_age = cache->slots[0].age; + for (uint32_t i = 1; i < cache->capacity; i++) { + if (cache->slots[i].age < best_age) { + best_age = cache->slots[i].age; + slot = i; + } + } + return slot; +} + +static int cuda_stream_expert_cache_copy_to_compact( + cuda_stream_expert_cache *cache, + uint32_t cache_slot, + uint32_t compact_slot, + char *compact_gate, + char *compact_up, + char *compact_down) { + const uint64_t gate_src = (uint64_t)cache_slot * cache->gate_expert_bytes; + const uint64_t down_src = (uint64_t)cache_slot * cache->down_expert_bytes; + const uint64_t gate_dst = (uint64_t)compact_slot * cache->gate_expert_bytes; + const uint64_t down_dst = (uint64_t)compact_slot * cache->down_expert_bytes; + return cuda_ok(cudaMemcpy(compact_gate + gate_dst, + cache->gate_ptr + gate_src, + (size_t)cache->gate_expert_bytes, + cudaMemcpyDeviceToDevice), + "streaming selected gate cache copy") && + cuda_ok(cudaMemcpy(compact_up + gate_dst, + cache->up_ptr + gate_src, + (size_t)cache->gate_expert_bytes, + cudaMemcpyDeviceToDevice), + "streaming selected up cache copy") && + cuda_ok(cudaMemcpy(compact_down + down_dst, + cache->down_ptr + down_src, + (size_t)cache->down_expert_bytes, + cudaMemcpyDeviceToDevice), + "streaming selected down cache copy"); +} + +static int cuda_stream_expert_cache_load_slot( + cuda_stream_expert_cache *cache, + const void *model_map, + uint64_t model_size, + uint32_t slot, + uint32_t layer, + uint32_t n_total_expert, + uint32_t expert, + uint64_t gate_offset, + uint64_t up_offset, + uint64_t down_offset, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes) { + const uint64_t gate_src = + gate_offset + (uint64_t)expert * gate_expert_bytes; + const uint64_t up_src = + up_offset + (uint64_t)expert * gate_expert_bytes; + const uint64_t down_src = + down_offset + (uint64_t)expert * down_expert_bytes; + const uint64_t gate_dst = (uint64_t)slot * gate_expert_bytes; + const uint64_t down_dst = (uint64_t)slot * down_expert_bytes; + /* Group the three tensor uploads of the expert: async copies pipeline + * the host read of one tensor with the device upload of the previous + * one; a single flush fences the group. */ + if (!cuda_model_copy_to_device_streamed_ex(cache->gate_ptr + gate_dst, + model_map, + model_size, + gate_src, + gate_expert_bytes, + "cached moe_gate", 0) || + !cuda_model_copy_to_device_streamed_ex(cache->up_ptr + gate_dst, + model_map, + model_size, + up_src, + gate_expert_bytes, + "cached moe_up", 0) || + !cuda_model_copy_to_device_streamed_ex(cache->down_ptr + down_dst, + model_map, + model_size, + down_src, + down_expert_bytes, + "cached moe_down", 0)) { + (void)cuda_model_copy_to_device_streamed_flush("cached expert"); + return 0; + } + if (!cuda_model_copy_to_device_streamed_flush("cached expert")) return 0; + cuda_stream_expert_cache_slot &entry = cache->slots[slot]; + entry.valid = 1; + entry.model_map = model_map; + entry.model_size = model_size; + entry.layer = layer; + entry.n_total_expert = n_total_expert; + entry.expert = expert; + entry.gate_offset = gate_offset; + entry.up_offset = up_offset; + entry.down_offset = down_offset; + entry.gate_expert_bytes = gate_expert_bytes; + entry.down_expert_bytes = down_expert_bytes; + entry.age = ++cache->tick; + return 1; +} + +static int cuda_stream_expert_cache_seed_one( + cuda_stream_expert_cache *cache, + const void *model_map, + uint64_t model_size, + uint32_t layer, + uint32_t n_total_expert, + uint32_t expert, + uint64_t gate_offset, + uint64_t up_offset, + uint64_t down_offset, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes) { + int cache_slot = cuda_stream_expert_cache_find(cache, + model_map, + model_size, + layer, + n_total_expert, + expert, + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes); + if (cache_slot >= 0) { + cache->slots[(uint32_t)cache_slot].age = ++cache->tick; + return 1; + } + + const uint32_t load_slot = cuda_stream_expert_cache_lru_slot(cache); + const int append = !cache->slots[load_slot].valid; + if (!cuda_stream_expert_cache_load_slot(cache, + model_map, + model_size, + load_slot, + layer, + n_total_expert, + expert, + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes)) { + return 0; + } + if (append && cache->count < cache->capacity) cache->count++; + return 1; +} + +static int cuda_stream_layer_expert_ranges_valid( + uint64_t model_size, + uint32_t n_total_expert, + uint64_t gate_offset, + uint64_t up_offset, + uint64_t down_offset, + uint64_t gate_expert_bytes, + uint64_t down_expert_bytes, + const char *what) { + if (n_total_expert == 0 || + gate_expert_bytes == 0 || + down_expert_bytes == 0 || + (uint64_t)n_total_expert > UINT64_MAX / gate_expert_bytes || + (uint64_t)n_total_expert > UINT64_MAX / down_expert_bytes) { + fprintf(stderr, + "ds4: CUDA streaming %s expert size overflow\n", + what ? what : "selected"); + return 0; + } + const uint64_t full_gate_bytes = + (uint64_t)n_total_expert * gate_expert_bytes; + const uint64_t full_down_bytes = + (uint64_t)n_total_expert * down_expert_bytes; + if (gate_offset > model_size || up_offset > model_size || + down_offset > model_size || + full_gate_bytes > model_size - gate_offset || + full_gate_bytes > model_size - up_offset || + full_down_bytes > model_size - down_offset) { + fprintf(stderr, + "ds4: CUDA streaming %s expert range outside model map\n", + what ? what : "selected"); + return 0; + } + return 1; +} + static int cuda_stream_selected_cache_begin_load( const ds4_gpu_stream_expert_table *table, const int32_t *selected_ids, - uint32_t slot_count) { + uint32_t slot_count, + int allow_cache_evict) { cuda_stream_selected_cache_invalidate(); if (!g_ssd_streaming_mode) return 1; if (!cuda_stream_selected_ranges_valid(table) || !selected_ids || @@ -23009,53 +23811,186 @@ static int cuda_stream_selected_cache_begin_load( g_stream_selected_cache.slot_selected_ptr)) { cuda_stream_selected_cache_release(); } - if (ds4_gpu_set_current_device(logical_tier) != 0 || - !cuda_stream_selected_ensure_bytes( + if (ds4_gpu_set_current_device(logical_tier) != 0) { + cuda_stream_selected_cache_invalidate(); + return 0; + } + /* Try to allocate the staging buffers with the resident expert cache + * still in place: an unconditional release here would throw away the + * warm decode working set on every request. Only when VRAM is really + * short (very long prompts) sacrifice the cache and retry once. */ + int selected_ok = + cuda_stream_selected_ensure_bytes( &g_stream_selected_cache.gate_ptr, &g_stream_selected_cache.gate_capacity, - gate_bytes, "gate experts") || - !cuda_stream_selected_ensure_bytes( + gate_bytes, "gate experts") && + cuda_stream_selected_ensure_bytes( &g_stream_selected_cache.up_ptr, &g_stream_selected_cache.up_capacity, - gate_bytes, "up experts") || - !cuda_stream_selected_ensure_bytes( + gate_bytes, "up experts") && + cuda_stream_selected_ensure_bytes( &g_stream_selected_cache.down_ptr, &g_stream_selected_cache.down_capacity, - down_bytes, "down experts") || - !cuda_stream_selected_ensure_i32(slot_count)) { + down_bytes, "down experts") && + cuda_stream_selected_ensure_i32(slot_count); + if (!selected_ok) { + cuda_stream_expert_cache_release_all(); + selected_ok = + cuda_stream_selected_ensure_bytes( + &g_stream_selected_cache.gate_ptr, + &g_stream_selected_cache.gate_capacity, + gate_bytes, "gate experts") && + cuda_stream_selected_ensure_bytes( + &g_stream_selected_cache.up_ptr, + &g_stream_selected_cache.up_capacity, + gate_bytes, "up experts") && + cuda_stream_selected_ensure_bytes( + &g_stream_selected_cache.down_ptr, + &g_stream_selected_cache.down_capacity, + down_bytes, "down experts") && + cuda_stream_selected_ensure_i32(slot_count); + } + if (!selected_ok) { cuda_stream_selected_cache_invalidate(); return 0; } + const uint32_t configured_cache_budget = + cuda_stream_expert_cache_configured_budget_class( + cuda_stream_expert_cache_class_index(table->gate_expert_bytes, + table->down_expert_bytes)); + cuda_stream_expert_cache *expert_cache = configured_cache_budget != 0 ? + cuda_stream_expert_cache_prepare(table->gate_expert_bytes, + table->down_expert_bytes, + configured_cache_budget) : + NULL; + int expert_cache_disabled = expert_cache == NULL; + const uint32_t cache_count_before = + expert_cache && expert_cache->valid ? expert_cache->count : 0; + uint32_t cache_hits = 0; + uint32_t cache_misses = 0; + uint32_t direct_loads = 0; + for (uint32_t i = 0; i < compact_ids.size(); i++) { const uint64_t expert = (uint32_t)compact_ids[i]; - const uint64_t gate_src = - table->gate_offset + expert * table->gate_expert_bytes; - const uint64_t up_src = - table->up_offset + expert * table->gate_expert_bytes; - const uint64_t down_src = - table->down_offset + expert * table->down_expert_bytes; const uint64_t gate_dst = (uint64_t)i * table->gate_expert_bytes; const uint64_t down_dst = (uint64_t)i * table->down_expert_bytes; - if (!cuda_model_copy_to_device_streamed( - g_stream_selected_cache.gate_ptr + gate_dst, - table->model_map, table->model_size, - gate_src, table->gate_expert_bytes, - "stream gate expert copy") || - !cuda_model_copy_to_device_streamed( - g_stream_selected_cache.up_ptr + gate_dst, - table->model_map, table->model_size, - up_src, table->gate_expert_bytes, - "stream up expert copy") || - !cuda_model_copy_to_device_streamed( - g_stream_selected_cache.down_ptr + down_dst, - table->model_map, table->model_size, - down_src, table->down_expert_bytes, - "stream down expert copy")) { - cuda_stream_selected_cache_invalidate(); - return 0; + int copied_from_global_cache = 0; + + if (!expert_cache_disabled) { + int cache_slot = + cuda_stream_expert_cache_find(expert_cache, + table->model_map, + table->model_size, + table->layer, + table->n_total_expert, + (uint32_t)expert, + table->gate_offset, + table->up_offset, + table->down_offset, + table->gate_expert_bytes, + table->down_expert_bytes); + if (cache_slot >= 0) { + cache_hits++; + expert_cache->slots[(uint32_t)cache_slot].age = + ++expert_cache->tick; + } else { + cache_misses++; + /* Prefill batch loads (allow_cache_evict=0) may append while + * there is free capacity but must NOT evict valid slots: a + * long prompt would cycle the whole LRU cache and destroy + * the decode working set. */ + if (allow_cache_evict || + expert_cache->count < expert_cache->capacity) { + const uint32_t load_slot = + cuda_stream_expert_cache_lru_slot(expert_cache); + const int append = !expert_cache->slots[load_slot].valid; + if (cuda_stream_expert_cache_load_slot( + expert_cache, + table->model_map, + table->model_size, + load_slot, + table->layer, + table->n_total_expert, + (uint32_t)expert, + table->gate_offset, + table->up_offset, + table->down_offset, + table->gate_expert_bytes, + table->down_expert_bytes)) { + if (append && + expert_cache->count < expert_cache->capacity) { + expert_cache->count++; + } + cache_slot = (int)load_slot; + } else { + cuda_stream_expert_cache_invalidate(); + expert_cache_disabled = 1; + cache_slot = -1; + } + } + } + + if (cache_slot >= 0) { + copied_from_global_cache = + cuda_stream_expert_cache_copy_to_compact( + expert_cache, + (uint32_t)cache_slot, + i, + g_stream_selected_cache.gate_ptr, + g_stream_selected_cache.up_ptr, + g_stream_selected_cache.down_ptr); + if (!copied_from_global_cache) { + cuda_stream_expert_cache_invalidate(); + expert_cache_disabled = 1; + } + } + } + + if (!copied_from_global_cache) { + const uint64_t gate_src = + table->gate_offset + expert * table->gate_expert_bytes; + const uint64_t up_src = + table->up_offset + expert * table->gate_expert_bytes; + const uint64_t down_src = + table->down_offset + expert * table->down_expert_bytes; + direct_loads++; + if (!cuda_model_copy_to_device_streamed( + g_stream_selected_cache.gate_ptr + gate_dst, + table->model_map, table->model_size, + gate_src, table->gate_expert_bytes, + "stream gate expert copy") || + !cuda_model_copy_to_device_streamed( + g_stream_selected_cache.up_ptr + gate_dst, + table->model_map, table->model_size, + up_src, table->gate_expert_bytes, + "stream up expert copy") || + !cuda_model_copy_to_device_streamed( + g_stream_selected_cache.down_ptr + down_dst, + table->model_map, table->model_size, + down_src, table->down_expert_bytes, + "stream down expert copy")) { + cuda_stream_selected_cache_invalidate(); + return 0; + } } } + + if (getenv("DS4_CUDA_STREAMING_EXPERT_CACHE_VERBOSE")) { + cuda_model_load_progress_finish(); + fprintf(stderr, + "ds4: CUDA streaming selected layer=%u slots=%u compact=%u " + "global_budget=%u before=%u after=%u hits=%u misses=%u direct=%u\n", + table->layer, + slot_count, + (uint32_t)compact_ids.size(), + expert_cache && expert_cache->valid ? expert_cache->capacity : 0, + cache_count_before, + expert_cache && expert_cache->valid ? expert_cache->count : 0, + cache_hits, + cache_misses, + direct_loads); + } if (!cuda_ok(cudaMemcpy(g_stream_selected_cache.slot_selected_ptr, slot_ids.data(), (size_t)slot_count * sizeof(int32_t), @@ -26713,7 +27648,7 @@ extern "C" int ds4_gpu_glm_stream_expert_cache_begin_selected_load_tensor( "GLM streaming selected-id read")) { return 0; } - return cuda_stream_selected_cache_begin_load(table, ids.data(), n_selected); + return cuda_stream_selected_cache_begin_load(table, ids.data(), n_selected, 1); } __global__ static void glm_value_project_q8_0_batch_heads_kernel( @@ -27254,15 +28189,21 @@ extern "C" int ds4_gpu_stream_expert_cache_begin_selected_load( const int32_t *selected_ids, uint32_t n_selected) { return cuda_stream_selected_cache_begin_load(table, selected_ids, - n_selected); + n_selected, 1); } extern "C" uint32_t ds4_gpu_stream_expert_cache_budget_for_expert_size( uint64_t gate_expert_bytes, uint64_t down_expert_bytes) { - (void)gate_expert_bytes; - (void)down_expert_bytes; - return 0; + if (!cuda_stream_expert_cache_budget_visible_to_shared() || + cuda_stream_expert_cache_expert_bytes(gate_expert_bytes, + down_expert_bytes) == 0) { + return 0; + } + const int class_idx = + cuda_stream_expert_cache_class_index(gate_expert_bytes, + down_expert_bytes); + return cuda_stream_expert_cache_configured_budget_class(class_idx); } extern "C" int ds4_gpu_tensor_copy_f32_to_f16(ds4_gpu_tensor *dst, uint64_t dst_offset, @@ -27373,12 +28314,22 @@ extern "C" void ds4_gpu_set_glm_model(bool enabled) { extern "C" void ds4_gpu_set_ssd_streaming(bool enabled) { g_ssd_streaming_mode = enabled ? 1 : 0; + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + g_stream_expert_runtime_caps[i] = 0; + g_stream_expert_memory_cap_notices[i] = 0; + } cuda_stream_selected_cache_invalidate(); - if (!g_ssd_streaming_mode) cuda_stream_selected_cache_release(); + if (!g_ssd_streaming_mode) { + cuda_stream_selected_cache_release(); + cuda_stream_expert_cache_release_all(); + } } extern "C" void ds4_gpu_set_streaming_expert_cache_budget(uint32_t experts) { - (void)experts; + g_stream_expert_budget_override = experts; + cuda_stream_selected_cache_invalidate(); + /* release_all also clears the size classes, runtime caps and notices. */ + cuda_stream_expert_cache_release_all(); } extern "C" void ds4_gpu_set_streaming_expert_cache_expert_bytes(uint64_t bytes) { @@ -27386,18 +28337,23 @@ extern "C" void ds4_gpu_set_streaming_expert_cache_expert_bytes(uint64_t bytes) } extern "C" uint32_t ds4_gpu_stream_expert_cache_configured_count(void) { - return 0; + if (!cuda_stream_expert_cache_budget_visible_to_shared()) return 0; + return cuda_stream_expert_cache_configured_budget(); } extern "C" uint32_t ds4_gpu_stream_expert_cache_current_count(void) { - return g_stream_selected_cache.valid ? - g_stream_selected_cache.compact_count : 0; + uint32_t count = 0; + for (int i = 0; i < DS4_CUDA_STREAM_EXPERT_CLASSES; i++) { + count += g_stream_expert_caches[i].count; + } + return count; } extern "C" void ds4_gpu_stream_expert_cache_reset_route_hotness(void) { } extern "C" void ds4_gpu_stream_expert_cache_release_resident(void) { + cuda_stream_expert_cache_release_all(); cuda_stream_selected_cache_release(); } @@ -27405,7 +28361,59 @@ extern "C" int ds4_gpu_stream_expert_cache_seed_selected( const ds4_gpu_stream_expert_table *table, const int32_t *selected_ids, uint32_t n_selected) { - (void)table; (void)selected_ids; (void)n_selected; + if (!g_ssd_streaming_mode) return 1; + if (!table) return 0; + const void *model_map = table->model_map; + const uint64_t model_size = table->model_size; + const uint32_t layer = table->layer; + const uint32_t n_total_expert = table->n_total_expert; + const uint64_t gate_offset = table->gate_offset; + const uint64_t up_offset = table->up_offset; + const uint64_t down_offset = table->down_offset; + const uint64_t gate_expert_bytes = table->gate_expert_bytes; + const uint64_t down_expert_bytes = table->down_expert_bytes; + if (!model_map || !selected_ids || n_selected == 0 || + n_selected > n_total_expert || + !cuda_stream_layer_expert_ranges_valid(model_size, + n_total_expert, + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes, + "seed selected")) { + return 0; + } + + cuda_stream_expert_cache *cache = + cuda_stream_expert_cache_prepare(gate_expert_bytes, + down_expert_bytes, + n_selected); + if (!cache) return 1; + for (uint32_t i = 0; i < n_selected; i++) { + if (selected_ids[i] < 0 || (uint32_t)selected_ids[i] >= n_total_expert) { + fprintf(stderr, + "ds4: CUDA streaming seed selected expert id %d is outside 0..%u at layer %u\n", + selected_ids[i], + n_total_expert, + layer); + return 0; + } + if (!cuda_stream_expert_cache_seed_one(cache, + model_map, + model_size, + layer, + n_total_expert, + (uint32_t)selected_ids[i], + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes)) { + cuda_stream_expert_cache_invalidate(); + return 1; + } + } return 1; } @@ -27418,8 +28426,11 @@ extern "C" int ds4_gpu_stream_expert_cache_prepare_selected_batch( (uint64_t)n_tokens * n_selected > UINT32_MAX) { return 0; } + /* Prefill batches read the resident expert cache too (hits become + * device-to-device copies instead of host uploads) but never evict: + * see allow_cache_evict in the load loop. */ return cuda_stream_selected_cache_begin_load( - table, selected_ids, n_tokens * n_selected); + table, selected_ids, n_tokens * n_selected, 0); } extern "C" int ds4_gpu_stream_expert_cache_seed_experts( @@ -27427,7 +28438,99 @@ extern "C" int ds4_gpu_stream_expert_cache_seed_experts( const int32_t *expert_ids, const uint32_t *expert_priorities, uint32_t n_experts) { - (void)table; (void)expert_ids; (void)expert_priorities; (void)n_experts; + if (!g_ssd_streaming_mode) return 1; + if (!table) return 0; + const void *model_map = table->model_map; + const uint64_t model_size = table->model_size; + const uint32_t layer = table->layer; + const uint32_t n_total_expert = table->n_total_expert; + const uint64_t gate_offset = table->gate_offset; + const uint64_t up_offset = table->up_offset; + const uint64_t down_offset = table->down_offset; + const uint64_t gate_expert_bytes = table->gate_expert_bytes; + const uint64_t down_expert_bytes = table->down_expert_bytes; + if (!model_map || !expert_ids || n_experts == 0 || + !cuda_stream_layer_expert_ranges_valid(model_size, + n_total_expert, + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes, + "seed hotlist")) { + return 0; + } + + cuda_stream_expert_cache *cache = + cuda_stream_expert_cache_prepare(gate_expert_bytes, + down_expert_bytes, + n_experts); + if (!cache || cache->capacity == 0) return 1; + + const uint32_t layer_seed_cap = + n_experts < cache->capacity ? n_experts : cache->capacity; + std::vector chosen; + try { + chosen.reserve(layer_seed_cap); + } catch (...) { + return 1; + } + + for (uint32_t i = 0; i < n_experts; i++) { + const int32_t expert = expert_ids[i]; + if (expert < 0 || (uint32_t)expert >= n_total_expert) { + fprintf(stderr, + "ds4: CUDA streaming hotlist seed expert id %d is outside 0..%u at layer %u\n", + expert, + n_total_expert, + layer); + return 0; + } + const uint32_t priority = + expert_priorities ? expert_priorities[i] : (n_experts - i); + uint32_t pos = 0; + while (pos < chosen.size()) { + const uint32_t other = chosen[pos]; + const uint32_t other_priority = + expert_priorities ? expert_priorities[other] : + (n_experts - other); + if (priority > other_priority) break; + pos++; + } + if (chosen.size() < layer_seed_cap) { + chosen.insert(chosen.begin() + pos, i); + } else if (pos < chosen.size()) { + chosen.insert(chosen.begin() + pos, i); + chosen.pop_back(); + } + } + + const uint32_t n = (uint32_t)chosen.size(); + for (uint32_t ri = 0; ri < n; ri++) { + const uint32_t i = chosen[n - 1u - ri]; + if (!cuda_stream_expert_cache_seed_one(cache, + model_map, + model_size, + layer, + n_total_expert, + (uint32_t)expert_ids[i], + gate_offset, + up_offset, + down_offset, + gate_expert_bytes, + down_expert_bytes)) { + cuda_stream_expert_cache_invalidate(); + return 1; + } + } + if (getenv("DS4_CUDA_STREAMING_EXPERT_CACHE_VERBOSE")) { + fprintf(stderr, + "ds4: CUDA streaming hotlist seeded layer=%u requested=%u cached=%u cap=%u\n", + layer, + n_experts, + n, + cache->capacity); + } return 1; }