From ac3ae20529aa22fecacdae0bb53bfda57539db79 Mon Sep 17 00:00:00 2001 From: Graffioh Date: Sat, 11 Jul 2026 09:12:26 +0200 Subject: [PATCH 1/4] fix(deepseek4): hand off full HC state between local shards Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce --- server/src/deepseek4/deepseek4_graph.cpp | 27 ++++++++++++++++--- .../deepseek4_layer_split_adapter.cpp | 22 +++++++-------- .../deepseek4/deepseek4_layer_split_adapter.h | 4 +++ 3 files changed, 36 insertions(+), 17 deletions(-) diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 19adc8b49..c62b7adfb 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -5798,9 +5798,21 @@ bool deepseek4_step_layer_range( // Initialize HC state. // First shard (layer_begin=0): embed is token embeddings [n_embd × n_tokens], // replicate into n_hc streams. - // Later shards: embed is full HC state [hc_dim × n_tokens] from previous shard. - if (hc_state.size() != (size_t)hc_dim * (size_t)n_tokens) { - hc_state.resize((size_t)hc_dim * (size_t)n_tokens); + // Later shards: embed is full HC state [hc_dim × n_tokens] from previous + // shard — either hc_state's own buffer (local in-process handoff) or a + // separate buffer (IPC daemon). Detect the alias before any resize, which + // would invalidate embed. + const size_t hc_state_elems = (size_t)hc_dim * (size_t)n_tokens; + const bool boundary_in_place = embed != nullptr && embed == hc_state.data(); + if (hc_state.size() != hc_state_elems) { + if (boundary_in_place) { + std::fprintf(stderr, + "[deepseek4] HC boundary state size mismatch for layer range [%d,%d): " + "have %zu want %zu\n", + layer_begin, layer_end, hc_state.size(), hc_state_elems); + return false; + } + hc_state.resize(hc_state_elems); } if (layer_begin == 0) { // First shard: replicate embedding into all HC streams @@ -5812,7 +5824,14 @@ bool deepseek4_step_layer_range( } } else { // Later shard: embed contains full HC state from previous shard - memcpy(hc_state.data(), embed, sizeof(float) * (size_t)hc_dim * (size_t)n_tokens); + if (!embed) { + std::fprintf(stderr, "[deepseek4] missing HC boundary state for layer range [%d,%d)\n", + layer_begin, layer_end); + return false; + } + if (!boundary_in_place) { + memcpy(hc_state.data(), embed, sizeof(float) * hc_state_elems); + } } // Keep all reusable layer-range resources under one model owner so park, diff --git a/server/src/deepseek4/deepseek4_layer_split_adapter.cpp b/server/src/deepseek4/deepseek4_layer_split_adapter.cpp index 87b998bf8..afbfe0c9b 100644 --- a/server/src/deepseek4/deepseek4_layer_split_adapter.cpp +++ b/server/src/deepseek4/deepseek4_layer_split_adapter.cpp @@ -124,6 +124,13 @@ size_t DeepSeek4LayerSplitAdapter::hc_state_elements( return (size_t) weights.n_hc * (size_t) weights.n_embd; } +const float * DeepSeek4LayerSplitAdapter::local_shard_input( + size_t shard_index, + const std::vector & token_embeddings, + const std::vector & hc_state) { + return shard_index == 0 ? token_embeddings.data() : hc_state.data(); +} + int DeepSeek4LayerSplitAdapter::compute_auto_split_layers() const { // Check env override first int override_layers = env_int("DFLASH_DS4_CUDA_LAYERS", -1); @@ -443,7 +450,6 @@ bool DeepSeek4LayerSplitAdapter::run_forward( const auto forward_t0 = SplitClock::now(); const int n_tokens = (int)tokens.size(); const int n_embd = shards_[0].weights.n_embd; - const int n_hc = shards_[0].weights.n_hc; // Embed tokens on first shard auto & first_shard = shards_[0]; @@ -456,17 +462,6 @@ bool DeepSeek4LayerSplitAdapter::run_forward( DeepSeek4StepTelemetry tel_acc; if (timing) tel_acc.embed_us = split_elapsed_us(embed_t0, SplitClock::now()); - // Initialize HC state from embedding for new sequences - if (base_pos == 0 && cur_pos_ == 0) { - for (int t = 0; t < n_tokens; ++t) { - for (int h = 0; h < n_hc; ++h) { - std::memcpy(hc_state_.data() + (size_t)h * n_embd, - embed.data() + (size_t)t * n_embd, - (size_t)n_embd * sizeof(float)); - } - } - } - // If using mixed target split (remote Halo shard), delegate to that path if (use_mixed_target_split()) { return run_mixed_forward(tokens, base_pos, last_tok, logits_out); @@ -479,10 +474,11 @@ bool DeepSeek4LayerSplitAdapter::run_forward( const bool is_last = (si == shards_.size() - 1); std::vector * shard_logits = is_last ? logits_out : nullptr; DeepSeek4StepTelemetry step_tel; + const float * shard_input = local_shard_input(si, embed, hc_state_); if (!deepseek4_step_layer_range( shard.backend, shard.weights, shard.cache, - hc_state_, embed.data(), n_tokens, base_pos, + hc_state_, shard_input, n_tokens, base_pos, shard.layer_begin, shard.layer_end, shard_logits, tokens.data(), timing ? &step_tel : nullptr)) { std::fprintf(stderr, "[deepseek4-split] forward failed on shard %zu\n", si); diff --git a/server/src/deepseek4/deepseek4_layer_split_adapter.h b/server/src/deepseek4/deepseek4_layer_split_adapter.h index 1b4edc538..2a52d1866 100644 --- a/server/src/deepseek4/deepseek4_layer_split_adapter.h +++ b/server/src/deepseek4/deepseek4_layer_split_adapter.h @@ -84,6 +84,10 @@ class DeepSeek4LayerSplitAdapter : public LayerSplitAdapter { int compute_auto_split_layers() const; static int estimate_cuda_layers_from_free_bytes(size_t free_bytes); static size_t hc_state_elements(const DeepSeek4Weights & weights); + static const float * local_shard_input( + size_t shard_index, + const std::vector & token_embeddings, + const std::vector & hc_state); DeepSeek4LayerSplitAdapterConfig cfg_; std::vector shards_; From d48edf2ce89e0ddee48b804c1666130a0314773a Mon Sep 17 00:00:00 2001 From: Graffioh Date: Sat, 11 Jul 2026 09:13:57 +0200 Subject: [PATCH 2/4] fix(deepseek4): scope HC runtime caches per shard Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce --- server/src/deepseek4/deepseek4_graph.cpp | 233 ++++++++++++++-------- server/src/deepseek4/deepseek4_internal.h | 8 + 2 files changed, 161 insertions(+), 80 deletions(-) diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index c62b7adfb..3b24b9b9a 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -3707,18 +3707,20 @@ static bool load_tensor_to_f16_cpu(std::vector & dst, ggml_tensor * t) return true; } -static void load_hc_weights_cpu(HcWeightsCpu & dst, ggml_tensor * fn, - ggml_tensor * scale, ggml_tensor * base) { - if (!fn || !scale || !base || dst.loaded) return; +static bool load_hc_weights_cpu(HcWeightsCpu & dst, ggml_tensor * fn, + ggml_tensor * scale, ggml_tensor * base) { + if (dst.loaded) return true; + if (!fn || !scale || !base) return false; if (!load_tensor_to_f16_cpu(dst.fn_data, fn) || !load_tensor_to_f32_cpu(dst.scale_data, scale) || !load_tensor_to_f32_cpu(dst.base_data, base)) { dst.fn_data.clear(); dst.scale_data.clear(); dst.base_data.clear(); - return; + return false; } dst.loaded = true; + return true; } static void reset_hc_weights_cpu(HcWeightsCpu & w) { @@ -4352,9 +4354,15 @@ struct DeepSeek4FusedDecodeCache { } }; -struct DeepSeek4LayerRangeRuntime { +struct DeepSeek4LayerRangeCache { + ~DeepSeek4LayerRangeCache() { reset(); } + + const DeepSeek4Weights * owner_weights = nullptr; const ggml_context * owner_ctx = nullptr; - int loaded_n_layer = 0; + ggml_backend_t backend = nullptr; + int layer_begin = -1; + int layer_end = -1; + bool owns_output = false; std::vector hc_layer_weights; HcWeightsCpu hc_output_weights; std::vector hash_routing_tables; @@ -4370,7 +4378,20 @@ struct DeepSeek4LayerRangeRuntime { Ds4DecodeSharedInputs decode_shared_inputs; DeepSeek4LayerRangeScratch scratch; - void destroy() { + bool matches(const DeepSeek4Weights & w, + ggml_backend_t candidate_backend, + int candidate_begin, + int candidate_end, + bool candidate_owns_output) const { + return owner_weights == &w && + owner_ctx == w.ctx && + backend == candidate_backend && + layer_begin == candidate_begin && + layer_end == candidate_end && + owns_output == candidate_owns_output; + } + + void reset() { for (auto & alloc : cached_attn_allocs) { alloc.free(); } @@ -4403,13 +4424,15 @@ struct DeepSeek4LayerRangeRuntime { hash_routing_tables.clear(); hash_routing_tables.shrink_to_fit(); scratch.clear(); - loaded_n_layer = 0; + owner_weights = nullptr; owner_ctx = nullptr; + backend = nullptr; + layer_begin = -1; + layer_end = -1; + owns_output = false; } }; -static thread_local DeepSeek4LayerRangeRuntime ds4_layer_range_runtime; - static ggml_tensor * ds4_fused_hc_base_f32(ggml_context * ctx, ggml_tensor * base) { if (!base) return nullptr; ggml_tensor * b = base; @@ -5197,9 +5220,6 @@ void deepseek4_release_runtime_graphs(const DeepSeek4Weights & w) { ds4_layer_major_meta_arena.shrink_to_fit(); ds4_layer_major_meta_owner = nullptr; } - if (ds4_layer_range_runtime.owner_ctx == owner) { - ds4_layer_range_runtime.destroy(); - } if (ds4_hybrid_runtime.owner_ctx == owner) { ds4_hybrid_runtime.destroy(); } @@ -5693,6 +5713,95 @@ static int ds4_try_layer_major_prefill( return out_logits.empty() ? -1 : 1; } +static bool ds4_hc_layer_weights_ready(const HcWeightsCpu & weights, + int n_embd, + int n_hc) { + const size_t hc_dim = (size_t)n_embd * (size_t)n_hc; + const size_t mix_dim = (size_t)(2 * n_hc + n_hc * n_hc); + return weights.loaded && + weights.fn_data.size() >= hc_dim * mix_dim && + weights.scale_data.size() >= 3 && + weights.base_data.size() >= mix_dim; +} + +static bool ds4_hc_output_weights_ready(const HcWeightsCpu & weights, + int n_embd, + int n_hc) { + const size_t hc_dim = (size_t)n_embd * (size_t)n_hc; + return weights.loaded && + weights.fn_data.size() >= hc_dim * (size_t)n_hc && + !weights.scale_data.empty() && + weights.base_data.size() >= (size_t)n_hc; +} + +static bool initialize_layer_range_cache( + DeepSeek4LayerRangeCache & runtime, + ggml_backend_t backend, + const DeepSeek4Weights & w, + int layer_begin, + int layer_end, + bool owns_output) { + runtime.reset(); + if (layer_begin < 0 || layer_end < layer_begin || layer_end > w.n_layer) { + std::fprintf(stderr, + "[deepseek4] invalid HC cache layer range [%d,%d) for %d layers\n", + layer_begin, layer_end, w.n_layer); + return false; + } + + runtime.hc_layer_weights.resize((size_t)w.n_layer); + runtime.hash_routing_tables.assign((size_t)w.n_layer, {}); + runtime.cached_attn_allocs.assign((size_t)w.n_layer, {}); + runtime.cached_decode_attn_hc_pre_graphs.assign((size_t)w.n_layer, {}); + runtime.cached_decode_ffn_hc_pre_graphs.assign((size_t)w.n_layer, {}); + runtime.cached_decode_attn_graphs.assign((size_t)w.n_layer, {}); + runtime.cached_decode_ffn_graphs.assign((size_t)w.n_layer, {}); + + for (int il = layer_begin; il < layer_end; ++il) { + const DeepSeek4Layer & layer = w.layers[(size_t)il]; + HcLayerWeightsCpu & cached = runtime.hc_layer_weights[(size_t)il]; + const bool attn_loaded = load_hc_weights_cpu( + cached.attn, layer.hc_attn_fn, layer.hc_attn_scale, layer.hc_attn_base); + const bool ffn_loaded = load_hc_weights_cpu( + cached.ffn, layer.hc_ffn_fn, layer.hc_ffn_scale, layer.hc_ffn_base); + if (!attn_loaded || !ds4_hc_layer_weights_ready(cached.attn, w.n_embd, w.n_hc)) { + std::fprintf(stderr, + "[deepseek4] missing or invalid HC attention weights for layer %d\n", il); + runtime.reset(); + return false; + } + if (!ffn_loaded || !ds4_hc_layer_weights_ready(cached.ffn, w.n_embd, w.n_hc)) { + std::fprintf(stderr, + "[deepseek4] missing or invalid HC FFN weights for layer %d\n", il); + runtime.reset(); + return false; + } + if (il < w.n_hash_layer && layer.ffn_gate_tid2eid) { + load_hash_routing_cpu(runtime.hash_routing_tables[(size_t)il], + layer.ffn_gate_tid2eid); + } + } + + if (owns_output) { + if (!load_hc_weights_cpu(runtime.hc_output_weights, + w.output_hc_fn, + w.output_hc_scale, + w.output_hc_base) || + !ds4_hc_output_weights_ready(runtime.hc_output_weights, w.n_embd, w.n_hc)) { + std::fprintf(stderr, "[deepseek4] missing or invalid HC output weights\n"); + runtime.reset(); + return false; + } + } + + runtime.owner_weights = &w; + runtime.owner_ctx = w.ctx; + runtime.backend = backend; + runtime.layer_begin = layer_begin; + runtime.layer_end = layer_end; + runtime.owns_output = owns_output; + return true; +} bool deepseek4_step_layer_range( ggml_backend_t backend, @@ -5803,9 +5912,9 @@ bool deepseek4_step_layer_range( // separate buffer (IPC daemon). Detect the alias before any resize, which // would invalidate embed. const size_t hc_state_elems = (size_t)hc_dim * (size_t)n_tokens; - const bool boundary_in_place = embed != nullptr && embed == hc_state.data(); + const bool embed_points_to_hc_state = embed != nullptr && embed == hc_state.data(); if (hc_state.size() != hc_state_elems) { - if (boundary_in_place) { + if (embed_points_to_hc_state) { std::fprintf(stderr, "[deepseek4] HC boundary state size mismatch for layer range [%d,%d): " "have %zu want %zu\n", @@ -5829,78 +5938,40 @@ bool deepseek4_step_layer_range( layer_begin, layer_end); return false; } - if (!boundary_in_place) { + if (!embed_points_to_hc_state) { memcpy(hc_state.data(), embed, sizeof(float) * hc_state_elems); } } - // Keep all reusable layer-range resources under one model owner so park, - // unload, and reload have a deterministic teardown boundary. - DeepSeek4LayerRangeRuntime & runtime = ds4_layer_range_runtime; - auto & hc_layer_weights_range = runtime.hc_layer_weights; - auto & hc_output_weights_range = runtime.hc_output_weights; - auto & hash_routing_tables_range = runtime.hash_routing_tables; - auto & cached_attn_allocs = runtime.cached_attn_allocs; - auto & cached_decode_attn_hc_pre_graphs = - runtime.cached_decode_attn_hc_pre_graphs; - auto & cached_decode_ffn_hc_pre_graphs = - runtime.cached_decode_ffn_hc_pre_graphs; - auto & cached_decode_hc_post_graph = runtime.cached_decode_hc_post_graph; - auto & cached_decode_attn_graphs = runtime.cached_decode_attn_graphs; - auto & cached_decode_ffn_graphs = runtime.cached_decode_ffn_graphs; - auto & cached_decode_output_graph = runtime.cached_decode_output_graph; - auto & cached_dynamic_output_alloc = runtime.cached_dynamic_output_alloc; - auto & fused_decode_graph_cache = runtime.fused_decode_graph_cache; - auto & decode_shared_inputs = runtime.decode_shared_inputs; - int & hc_loaded_n_layer = runtime.loaded_n_layer; - const ggml_context * & hc_loaded_ctx = runtime.owner_ctx; - if (hc_loaded_n_layer != w.n_layer || hc_loaded_ctx != w.ctx) { - reset_hc_layer_weights_cpu(hc_layer_weights_range); - reset_hc_weights_cpu(hc_output_weights_range); - hc_layer_weights_range.resize((size_t)w.n_layer); - hash_routing_tables_range.assign((size_t)w.n_layer, {}); - for (auto & alloc : cached_attn_allocs) { - alloc.free(); - } - cached_attn_allocs.assign((size_t)w.n_layer, {}); - for (auto & g : cached_decode_attn_hc_pre_graphs) { - g.free(); - } - cached_decode_attn_hc_pre_graphs.assign((size_t)w.n_layer, {}); - for (auto & g : cached_decode_ffn_hc_pre_graphs) { - g.free(); - } - cached_decode_ffn_hc_pre_graphs.assign((size_t)w.n_layer, {}); - cached_decode_hc_post_graph.free(); - for (auto & per_layer : cached_decode_attn_graphs) { - for (auto & g : per_layer) { - g.free(); - } - } - cached_decode_attn_graphs.assign((size_t)w.n_layer, {}); - for (auto & g : cached_decode_ffn_graphs) { - g.free(); - } - cached_decode_ffn_graphs.assign((size_t)w.n_layer, {}); - cached_decode_output_graph.free(); - cached_dynamic_output_alloc.free(); - fused_decode_graph_cache.destroy(); - decode_shared_inputs.free(); - for (int il = 0; il < w.n_layer; il++) { - const DeepSeek4Layer & L = w.layers[(size_t)il]; - load_hc_weights_cpu(hc_layer_weights_range[il].attn, L.hc_attn_fn, L.hc_attn_scale, L.hc_attn_base); - load_hc_weights_cpu(hc_layer_weights_range[il].ffn, L.hc_ffn_fn, L.hc_ffn_scale, L.hc_ffn_base); - if (il < w.n_hash_layer && L.ffn_gate_tid2eid) { - load_hash_routing_cpu(hash_routing_tables_range[(size_t)il], L.ffn_gate_tid2eid); - } - } - load_hc_weights_cpu(hc_output_weights_range, w.output_hc_fn, w.output_hc_scale, w.output_hc_base); - hc_loaded_n_layer = w.n_layer; - hc_loaded_ctx = w.ctx; + // Keep HC and graph runtime state per DeepSeek4Cache. This isolates model + // and shard instances that may execute interleaved in the same process and + // ties the cached resources to the owning cache's lifetime. + if (!cache.layer_range_cache) { + cache.layer_range_cache = new DeepSeek4LayerRangeCache(); + } + DeepSeek4LayerRangeCache & layer_range_cache = *cache.layer_range_cache; + if (!layer_range_cache.matches(w, backend, layer_begin, layer_end, is_last_shard) && + !initialize_layer_range_cache( + layer_range_cache, backend, w, layer_begin, layer_end, is_last_shard)) { + return false; } + auto & hc_layer_weights_range = layer_range_cache.hc_layer_weights; + auto & hc_output_weights_range = layer_range_cache.hc_output_weights; + auto & hash_routing_tables_range = layer_range_cache.hash_routing_tables; + auto & cached_attn_allocs = layer_range_cache.cached_attn_allocs; + auto & cached_decode_attn_hc_pre_graphs = layer_range_cache.cached_decode_attn_hc_pre_graphs; + auto & cached_decode_ffn_hc_pre_graphs = layer_range_cache.cached_decode_ffn_hc_pre_graphs; + auto & cached_decode_hc_post_graph = layer_range_cache.cached_decode_hc_post_graph; + auto & cached_decode_attn_graphs = layer_range_cache.cached_decode_attn_graphs; + auto & cached_decode_ffn_graphs = layer_range_cache.cached_decode_ffn_graphs; + auto & cached_decode_output_graph = layer_range_cache.cached_decode_output_graph; + auto & cached_dynamic_output_alloc = layer_range_cache.cached_dynamic_output_alloc; + auto & fused_decode_graph_cache = layer_range_cache.fused_decode_graph_cache; + auto & decode_shared_inputs = layer_range_cache.decode_shared_inputs; + // Per-layer execution with CPU-side HC - DeepSeek4LayerRangeScratch & scratch = runtime.scratch; + DeepSeek4LayerRangeScratch & scratch = layer_range_cache.scratch; const int n_expert_used = ds4_effective_expert_count(w); scratch.ensure(w.ctx, n_tokens, n_embd, n_hc, n_expert_used); @@ -6708,6 +6779,8 @@ bool create_deepseek4_cache(ggml_backend_t backend, } void free_deepseek4_cache(DeepSeek4Cache & c) { + delete c.layer_range_cache; + c.layer_range_cache = nullptr; if (c.ctx) { ggml_free(c.ctx); c.ctx = nullptr; } if (c.buf) { ggml_backend_buffer_free(c.buf); c.buf = nullptr; } c.layers.clear(); diff --git a/server/src/deepseek4/deepseek4_internal.h b/server/src/deepseek4/deepseek4_internal.h index d38532af6..07ed34f94 100644 --- a/server/src/deepseek4/deepseek4_internal.h +++ b/server/src/deepseek4/deepseek4_internal.h @@ -265,6 +265,11 @@ struct DeepSeek4LayerCache { DeepSeek4CompressorState indexer_compressor; }; +// Per-shard runtime state for deepseek4_step_layer_range (host-side HC weight +// cache + cached decode graphs). Defined in deepseek4_graph.cpp; owned by the +// DeepSeek4Cache below and released by free_deepseek4_cache(). +struct DeepSeek4LayerRangeCache; + struct DeepSeek4Cache { int cur_pos = 0; int max_ctx = 0; @@ -276,6 +281,9 @@ struct DeepSeek4Cache { // HC residual streams: [n_hc * n_embd] persistent state ggml_tensor * hc_state = nullptr; // [n_hc * n_embd] + // Lazily created on the first deepseek4_step_layer_range call. + DeepSeek4LayerRangeCache * layer_range_cache = nullptr; + ggml_context * ctx = nullptr; ggml_backend_buffer_t buf = nullptr; }; From 685de03ad73228e6d05d80e519a1d859be7cf665 Mon Sep 17 00:00:00 2001 From: Graffioh Date: Mon, 13 Jul 2026 16:54:41 +0200 Subject: [PATCH 3/4] fix(ds4): scope HC scratch by GPU device Also size the shared scratch for the requested shapes: the bounded buffers (mix/post/comb) are allocated at their kMaxHc maxima and the working buffer tracks n_embd capacity, instead of staying at whatever dims the first call happened to use. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce --- server/src/deepseek4/deepseek4_backend.cpp | 8 +- server/src/deepseek4/deepseek4_dspark.h | 2 + .../src/deepseek4/deepseek4_dspark_spec.cpp | 15 +- server/src/deepseek4/deepseek4_graph.cpp | 29 +- server/src/deepseek4/deepseek4_hc_cuda.cu | 189 +++++++++---- server/src/deepseek4/deepseek4_hc_cuda.h | 4 + server/src/deepseek4/deepseek4_internal.h | 2 + .../deepseek4_layer_split_adapter.cpp | 4 +- .../deepseek4_target_shard_ipc_daemon.cpp | 2 +- server/tests/test_deepseek4_unit.cpp | 252 ++++++++++++++++++ 10 files changed, 434 insertions(+), 73 deletions(-) diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp index 9588d2e69..90e8284fc 100644 --- a/server/src/deepseek4/deepseek4_backend.cpp +++ b/server/src/deepseek4/deepseek4_backend.cpp @@ -831,14 +831,14 @@ int DeepSeek4Backend::do_prefill(const std::vector & tokens, } bool ok = false; if (moe_hybrid_) { - ok = deepseek4_step(backend_, w_, cache_, embed.data(), n_tok, pos, + ok = deepseek4_step(backend_, cfg_.device.gpu, w_, cache_, embed.data(), n_tok, pos, logits, moe_hybrid_.get(), tokens.data() + i, &stream_engine_, timing ? &step_tel : nullptr, routing_stats_.get(), hp); } else { std::vector hc_state; ok = deepseek4_step_layer_range( - backend_, w_, cache_, hc_state, embed.data(), n_tok, pos, + backend_, cfg_.device.gpu, w_, cache_, hc_state, embed.data(), n_tok, pos, 0, w_.n_layer, &logits, tokens.data() + i, timing ? &step_tel : nullptr, cfg_.prefill_mode != PrefillAttentionMode::Sparse, hp); @@ -919,7 +919,7 @@ bool DeepSeek4Backend::do_decode(int committed, int n_gen, if (timing) step_tel.embed_us = elapsed_us(embed_t0, Clock::now()); const int pos = std::max(0, committed + generated - 1); - if (!deepseek4_step(backend_, w_, cache_, embed.data(), 1, + if (!deepseek4_step(backend_, cfg_.device.gpu, w_, cache_, embed.data(), 1, pos, logits, moe_hybrid_.get(), &tok_to_eval, moe_hybrid_ ? &stream_engine_ : nullptr, @@ -1023,7 +1023,7 @@ GenerateResult DeepSeek4Backend::generate_impl(const GenerateRequest & req, std::vector spec_toks; spec_ran = true; if (!run_deepseek4_dspark_spec_decode( - backend_, w_, cache_, *spec_drafter_, committed, seed, + backend_, cfg_.device.gpu, w_, cache_, *spec_drafter_, committed, seed, req.n_gen - 1, win_len > 0 ? spec_feat_window_.data() : nullptr, win_len, spec_toks, &accept_rate, diff --git a/server/src/deepseek4/deepseek4_dspark.h b/server/src/deepseek4/deepseek4_dspark.h index adc08167c..0d3a07981 100644 --- a/server/src/deepseek4/deepseek4_dspark.h +++ b/server/src/deepseek4/deepseek4_dspark.h @@ -118,6 +118,7 @@ bool deepseek4_dspark_draft_forward(ggml_backend_t backend, // main_hidden feed. // Advances/updates the target cache exactly like a decode of these tokens. bool deepseek4_dspark_verify_forward(ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, const std::vector & capture_layer_ids, @@ -165,6 +166,7 @@ void deepseek4_spec_rollback_apply(const DeepSeek4SpecRollback & rollback, struct GenerateRequest; // fwd (from common/…); the loop only needs n_gen + committed bool run_deepseek4_dspark_spec_decode( ggml_backend_t backend, + int device, const DeepSeek4Weights & target_w, DeepSeek4Cache & target_cache, const DSparkDrafter & drafter, diff --git a/server/src/deepseek4/deepseek4_dspark_spec.cpp b/server/src/deepseek4/deepseek4_dspark_spec.cpp index d22bf0ac9..ddd388a38 100644 --- a/server/src/deepseek4/deepseek4_dspark_spec.cpp +++ b/server/src/deepseek4/deepseek4_dspark_spec.cpp @@ -43,9 +43,9 @@ namespace dflash::common { class DeepSeek4DFlashTarget : public DFlashTarget { public: DeepSeek4DFlashTarget(const DeepSeek4Weights & w, DeepSeek4Cache & cache, - ggml_backend_t backend, ggml_backend_t snap_backend, + ggml_backend_t backend, int device, ggml_backend_t snap_backend, std::vector capture_ids, int mask_tok) - : w_(w), cache_(cache), backend_(backend), snap_backend_(snap_backend), + : w_(w), cache_(cache), backend_(backend), device_(device), snap_backend_(snap_backend), capture_ids_(std::move(capture_ids)), mask_tok_(mask_tok) {} ~DeepSeek4DFlashTarget() override { clear_snapshot(); } @@ -80,7 +80,7 @@ class DeepSeek4DFlashTarget : public DFlashTarget { std::vector am1; std::vector feat1; std::vector logits1; - if (!deepseek4_dspark_verify_forward(backend_, w_, cache_, capture_ids_, + if (!deepseek4_dspark_verify_forward(backend_, device_, w_, cache_, capture_ids_, embed_buf_.data() + (size_t) t * w_.n_embd, tokens.data() + t, 1, base_pos + t, am1, keep_logits_ ? &logits1 : nullptr, @@ -103,7 +103,7 @@ class DeepSeek4DFlashTarget : public DFlashTarget { std::vector am; // n==1 must take the dynamic (non-reuse) path: the reused decode graph // skips the capture/all-logits hooks (backend HC), which this needs. - if (!deepseek4_dspark_verify_forward(backend_, w_, cache_, capture_ids_, + if (!deepseek4_dspark_verify_forward(backend_, device_, w_, cache_, capture_ids_, embed_buf_.data(), tokens.data(), n, base_pos, am, keep_logits_ ? &verify_logits_ : nullptr, verify_features_, telemetry_, @@ -195,6 +195,7 @@ class DeepSeek4DFlashTarget : public DFlashTarget { const DeepSeek4Weights & w_; DeepSeek4Cache & cache_; ggml_backend_t backend_; + int device_; ggml_backend_t snap_backend_; std::vector capture_ids_; int mask_tok_; @@ -359,6 +360,7 @@ void deepseek4_spec_rollback_apply(const DeepSeek4SpecRollback & rollback, // touches the fused single-token 23 tok/s path, with the Ds4VerifyHooks that // add per-layer mean-over-HC capture and full per-position logits. bool deepseek4_dspark_verify_forward(ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, const std::vector & capture_layer_ids, @@ -378,7 +380,7 @@ bool deepseek4_dspark_verify_forward(ggml_backend_t backend, hooks.capture_layer_ids = &capture_layer_ids; hooks.capture_out = &capture_out; hooks.all_logits_out = &all_logits; - if (!deepseek4_step_layer_range(backend, w, cache, hc_state, embed, n_tokens, kv_start, + if (!deepseek4_step_layer_range(backend, device, w, cache, hc_state, embed, n_tokens, kv_start, 0, w.n_layer, &last_logits, token_ids, telemetry, allow_graph_reuse, &hooks)) { @@ -404,6 +406,7 @@ bool deepseek4_dspark_verify_forward(ggml_backend_t backend, bool run_deepseek4_dspark_spec_decode( ggml_backend_t backend, + int device, const DeepSeek4Weights & target_w, DeepSeek4Cache & target_cache, const DSparkDrafter & drafter, @@ -475,7 +478,7 @@ bool run_deepseek4_dspark_spec_decode( ggml_backend_t snap_backend = ggml_backend_cpu_init(); if (!snap_backend) { std::fprintf(stderr, "[ds4-spec] no CPU snapshot backend\n"); return false; } - DeepSeek4DFlashTarget target(target_w, target_cache, backend, snap_backend, + DeepSeek4DFlashTarget target(target_w, target_cache, backend, device, snap_backend, drafter.capture_layer_ids, drafter.mask_token_id); DraftWeights dw = make_dspark_shim(drafter); DeepSeek4SpecRollback rollback; diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 3b24b9b9a..b14af4474 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -4253,6 +4253,7 @@ static bool deepseek4_step_hybrid( bool deepseek4_step( ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, const float * embed, @@ -4266,6 +4267,12 @@ bool deepseek4_step( MoeHybridRoutingStats * routing_stats, Ds4VerifyHooks * verify_hooks) { if (w.moe_hybrid && moe_hybrid != nullptr) { + if (!deepseek4_cuda_hc_set_device(device)) { + std::fprintf(stderr, + "[deepseek4] failed to select HC device %d for hybrid step\n", + device); + return false; + } return deepseek4_step_hybrid(backend, w, cache, *moe_hybrid, embed, n_tokens, kv_start, out_logits, token_ids, stream_engine, telemetry, routing_stats); @@ -4273,7 +4280,7 @@ bool deepseek4_step( std::vector hc_state; return deepseek4_step_layer_range( - backend, w, cache, hc_state, embed, n_tokens, kv_start, + backend, device, w, cache, hc_state, embed, n_tokens, kv_start, 0, w.n_layer, &out_logits, token_ids, telemetry, /*allow_decode_graph_reuse=*/verify_hooks == nullptr, verify_hooks); } @@ -4360,6 +4367,7 @@ struct DeepSeek4LayerRangeCache { const DeepSeek4Weights * owner_weights = nullptr; const ggml_context * owner_ctx = nullptr; ggml_backend_t backend = nullptr; + int device = -1; int layer_begin = -1; int layer_end = -1; bool owns_output = false; @@ -4380,12 +4388,14 @@ struct DeepSeek4LayerRangeCache { bool matches(const DeepSeek4Weights & w, ggml_backend_t candidate_backend, + int candidate_device, int candidate_begin, int candidate_end, bool candidate_owns_output) const { return owner_weights == &w && owner_ctx == w.ctx && backend == candidate_backend && + device == candidate_device && layer_begin == candidate_begin && layer_end == candidate_end && owns_output == candidate_owns_output; @@ -4427,6 +4437,7 @@ struct DeepSeek4LayerRangeCache { owner_weights = nullptr; owner_ctx = nullptr; backend = nullptr; + device = -1; layer_begin = -1; layer_end = -1; owns_output = false; @@ -5737,6 +5748,7 @@ static bool ds4_hc_output_weights_ready(const HcWeightsCpu & weights, static bool initialize_layer_range_cache( DeepSeek4LayerRangeCache & runtime, ggml_backend_t backend, + int device, const DeepSeek4Weights & w, int layer_begin, int layer_end, @@ -5797,6 +5809,7 @@ static bool initialize_layer_range_cache( runtime.owner_weights = &w; runtime.owner_ctx = w.ctx; runtime.backend = backend; + runtime.device = device; runtime.layer_begin = layer_begin; runtime.layer_end = layer_end; runtime.owns_output = owns_output; @@ -5805,6 +5818,7 @@ static bool initialize_layer_range_cache( bool deepseek4_step_layer_range( ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, std::vector & hc_state, @@ -5820,6 +5834,13 @@ bool deepseek4_step_layer_range( Ds4VerifyHooks * verify_hooks) { const auto step_t0 = Ds4TimingClock::now(); + if (!deepseek4_cuda_hc_set_device(device)) { + std::fprintf(stderr, + "[deepseek4] failed to select HC device %d for layer range [%d,%d)\n", + device, layer_begin, layer_end); + return false; + } + // ── Partial layer-range forward with HC ───────────────────────────── const int n_embd = w.n_embd; const int n_hc = w.n_hc; @@ -5868,7 +5889,7 @@ bool deepseek4_step_layer_range( chunk_hooks_ptr = &chunk_hooks; } if (!deepseek4_step_layer_range( - backend, w, cache, chunk_hc, + backend, device, w, cache, chunk_hc, embed + (size_t) off * input_width, chunk, kv_start + off, layer_begin, layer_end, out_logits ? &chunk_out : nullptr, @@ -5950,9 +5971,9 @@ bool deepseek4_step_layer_range( cache.layer_range_cache = new DeepSeek4LayerRangeCache(); } DeepSeek4LayerRangeCache & layer_range_cache = *cache.layer_range_cache; - if (!layer_range_cache.matches(w, backend, layer_begin, layer_end, is_last_shard) && + if (!layer_range_cache.matches(w, backend, device, layer_begin, layer_end, is_last_shard) && !initialize_layer_range_cache( - layer_range_cache, backend, w, layer_begin, layer_end, is_last_shard)) { + layer_range_cache, backend, device, w, layer_begin, layer_end, is_last_shard)) { return false; } diff --git a/server/src/deepseek4/deepseek4_hc_cuda.cu b/server/src/deepseek4/deepseek4_hc_cuda.cu index 8f784d1f2..1462092eb 100644 --- a/server/src/deepseek4/deepseek4_hc_cuda.cu +++ b/server/src/deepseek4/deepseek4_hc_cuda.cu @@ -7,6 +7,7 @@ #include #include #include +#include #include #include @@ -20,6 +21,7 @@ constexpr int kMaxHc = 8; constexpr int kMaxMixDim = 2 * kMaxHc + kMaxHc * kMaxHc; struct HcCudaScratch { + int owner_device = -1; float * d_state = nullptr; float * d_sums = nullptr; float * d_mix = nullptr; @@ -29,8 +31,14 @@ struct HcCudaScratch { float * d_post = nullptr; float * d_comb = nullptr; size_t state_cap = 0; + size_t working_cap = 0; + + explicit HcCudaScratch(int device = -1) : owner_device(device) {} ~HcCudaScratch() { + if (owner_device >= 0) { + (void) cudaSetDevice(owner_device); + } if (d_state) cudaFree(d_state); if (d_sums) cudaFree(d_sums); if (d_mix) cudaFree(d_mix); @@ -41,14 +49,25 @@ struct HcCudaScratch { if (d_comb) cudaFree(d_comb); } - bool ensure(size_t hc_dim, size_t n_embd, size_t n_hc) { + bool ensure(size_t hc_dim, size_t n_embd) { + // Mix/scale/base/post/comb are bounded by kMaxHc (every entry point + // validates n_hc <= kMaxHc), so allocate them at their maxima once. + // Working and state scale with n_embd, so they track a capacity. if (!d_sums && cudaMalloc(&d_sums, sizeof(float) * kSums) != cudaSuccess) return false; - if (!d_mix && cudaMalloc(&d_mix, sizeof(float) * kMixDim) != cudaSuccess) return false; + if (!d_mix && cudaMalloc(&d_mix, sizeof(float) * kMaxMixDim) != cudaSuccess) return false; if (!d_scale && cudaMalloc(&d_scale, sizeof(float) * kMaxMixDim) != cudaSuccess) return false; if (!d_base && cudaMalloc(&d_base, sizeof(float) * kMaxMixDim) != cudaSuccess) return false; - if (!d_working && cudaMalloc(&d_working, sizeof(float) * n_embd) != cudaSuccess) return false; - if (!d_post && cudaMalloc(&d_post, sizeof(float) * n_hc) != cudaSuccess) return false; - if (!d_comb && cudaMalloc(&d_comb, sizeof(float) * n_hc * n_hc) != cudaSuccess) return false; + if (!d_post && cudaMalloc(&d_post, sizeof(float) * kMaxHc) != cudaSuccess) return false; + if (!d_comb && cudaMalloc(&d_comb, sizeof(float) * kMaxHc * kMaxHc) != cudaSuccess) return false; + if (working_cap < n_embd) { + if (d_working) { + cudaFree(d_working); + d_working = nullptr; + working_cap = 0; + } + if (cudaMalloc(&d_working, sizeof(float) * n_embd) != cudaSuccess) return false; + working_cap = n_embd; + } if (state_cap < hc_dim) { if (d_state) { cudaFree(d_state); @@ -62,8 +81,38 @@ struct HcCudaScratch { } }; -std::mutex g_mu; -HcCudaScratch g_scratch; +struct HcScratchSlot { + explicit HcScratchSlot(int device = -1) : scratch(device) {} + + std::mutex mutex; + HcCudaScratch scratch; +}; + +void hc_log_cuda_error(const char * label, cudaError_t err); + +// CUDA and HIP allocations belong to the device that was current when they +// were created. Keep one independently locked slot per logical device so local +// shards cannot reuse another device's scratch pointers. +std::mutex g_scratch_slots_mutex; +std::vector> g_scratch_slots; + +HcScratchSlot * current_scratch_slot() { + int device = -1; + const cudaError_t err = cudaGetDevice(&device); + if (err != cudaSuccess || device < 0) { + hc_log_cuda_error("get current device", err); + return nullptr; + } + + std::lock_guard lock(g_scratch_slots_mutex); + if (g_scratch_slots.size() <= (size_t) device) { + g_scratch_slots.resize((size_t) device + 1); + } + if (!g_scratch_slots[(size_t) device]) { + g_scratch_slots[(size_t) device] = std::make_unique(device); + } + return g_scratch_slots[(size_t) device].get(); +} void hc_log_cuda_error(const char * label, cudaError_t err) { if (err != cudaSuccess) { @@ -246,7 +295,8 @@ __global__ void hc_finish_kernel(const float * hc_state, } } -bool hc_pre_device_locked(const void * hc_state_device, +bool hc_pre_device_locked(HcCudaScratch & scratch, + const void * hc_state_device, const void * fn_device, const void * scale_device, const void * base_device, @@ -267,8 +317,8 @@ bool hc_pre_device_locked(const void * hc_state_device, return false; }; - if (hc_state_device != g_scratch.d_state) { - cudaError_t err = cudaMemcpy(g_scratch.d_state, hc_state_device, + if (hc_state_device != scratch.d_state) { + cudaError_t err = cudaMemcpy(scratch.d_state, hc_state_device, sizeof(float) * (size_t) hc_dim, cudaMemcpyDeviceToDevice); if (err != cudaSuccess) { @@ -277,28 +327,28 @@ bool hc_pre_device_locked(const void * hc_state_device, } hc_mix_norm_kernel<<>>( - g_scratch.d_state, + scratch.d_state, static_cast(fn_device), hc_dim, mix_dim, eps, - g_scratch.d_mix); + scratch.d_mix); cudaError_t err = cudaGetLastError(); if (err != cudaSuccess) { return fail("mix kernel", err); } hc_finish_kernel<<<1, kThreads>>>( - g_scratch.d_state, - g_scratch.d_mix, + scratch.d_state, + scratch.d_mix, static_cast(scale_device), static_cast(base_device), n_embd, n_hc, sinkhorn_iters, - g_scratch.d_working, - g_scratch.d_post, - g_scratch.d_comb); + scratch.d_working, + scratch.d_post, + scratch.d_comb); err = cudaGetLastError(); if (err != cudaSuccess) { return fail("finish kernel", err); @@ -310,24 +360,24 @@ bool hc_pre_device_locked(const void * hc_state_device, } #endif - if (working_device != g_scratch.d_working) { - err = cudaMemcpy(working_device, g_scratch.d_working, + if (working_device != scratch.d_working) { + err = cudaMemcpy(working_device, scratch.d_working, sizeof(float) * (size_t) n_embd, cudaMemcpyDeviceToDevice); if (err != cudaSuccess) { return fail("copy working d2d", err); } } - if (post_device != g_scratch.d_post) { - err = cudaMemcpy(post_device, g_scratch.d_post, + if (post_device != scratch.d_post) { + err = cudaMemcpy(post_device, scratch.d_post, sizeof(float) * (size_t) n_hc, cudaMemcpyDeviceToDevice); if (err != cudaSuccess) { return fail("copy post d2d", err); } } - if (comb_device != g_scratch.d_comb) { - err = cudaMemcpy(comb_device, g_scratch.d_comb, + if (comb_device != scratch.d_comb) { + err = cudaMemcpy(comb_device, scratch.d_comb, sizeof(float) * (size_t) n_hc * (size_t) n_hc, cudaMemcpyDeviceToDevice); if (err != cudaSuccess) { @@ -344,6 +394,18 @@ bool hc_pre_device_locked(const void * hc_state_device, } // namespace +bool deepseek4_cuda_hc_set_device(int device) { + if (device < 0) { + return false; + } + const cudaError_t err = cudaSetDevice(device); + if (err != cudaSuccess) { + hc_log_cuda_error("set device", err); + return false; + } + return true; +} + bool deepseek4_cuda_hc_pre_mix(const float * hc_state_host, const void * fn_device, int n_embd, @@ -354,31 +416,34 @@ bool deepseek4_cuda_hc_pre_mix(const float * hc_state_host, return false; } const int hc_dim = n_embd * n_hc; - std::lock_guard lock(g_mu); - if (!g_scratch.ensure((size_t)hc_dim, (size_t)n_embd, (size_t)n_hc)) { + HcScratchSlot * slot = current_scratch_slot(); + if (!slot) return false; + std::lock_guard lock(slot->mutex); + HcCudaScratch & scratch = slot->scratch; + if (!scratch.ensure((size_t)hc_dim, (size_t)n_embd)) { return false; } - if (cudaMemcpy(g_scratch.d_state, hc_state_host, sizeof(float) * (size_t)hc_dim, + if (cudaMemcpy(scratch.d_state, hc_state_host, sizeof(float) * (size_t)hc_dim, cudaMemcpyHostToDevice) != cudaSuccess) { return false; } - hc_sumsq_kernel<<>>(g_scratch.d_state, hc_dim, g_scratch.d_sums); + hc_sumsq_kernel<<>>(scratch.d_state, hc_dim, scratch.d_sums); if (cudaGetLastError() != cudaSuccess) return false; std::vector sums(kSums); - if (cudaMemcpy(sums.data(), g_scratch.d_sums, sizeof(float) * sums.size(), + if (cudaMemcpy(sums.data(), scratch.d_sums, sizeof(float) * sums.size(), cudaMemcpyDeviceToHost) != cudaSuccess) { return false; } float ss = 0.0f; for (float v : sums) ss += v; const float inv_rms = 1.0f / std::sqrt(ss / (float)hc_dim + eps); - hc_mix_kernel<<>>(g_scratch.d_state, + hc_mix_kernel<<>>(scratch.d_state, static_cast(fn_device), hc_dim, inv_rms, - g_scratch.d_mix); + scratch.d_mix); if (cudaGetLastError() != cudaSuccess) return false; - if (cudaMemcpy(mix_host, g_scratch.d_mix, sizeof(float) * kMixDim, + if (cudaMemcpy(mix_host, scratch.d_mix, sizeof(float) * kMixDim, cudaMemcpyDeviceToHost) != cudaSuccess) { return false; } @@ -407,47 +472,51 @@ bool deepseek4_cuda_hc_pre(const float * hc_state_host, return false; } - std::lock_guard lock(g_mu); - if (!g_scratch.ensure((size_t) hc_dim, (size_t) n_embd, (size_t) n_hc)) { + HcScratchSlot * slot = current_scratch_slot(); + if (!slot) return false; + std::lock_guard lock(slot->mutex); + HcCudaScratch & scratch = slot->scratch; + if (!scratch.ensure((size_t) hc_dim, (size_t) n_embd)) { return false; } - if (cudaMemcpy(g_scratch.d_state, hc_state_host, sizeof(float) * (size_t) hc_dim, + if (cudaMemcpy(scratch.d_state, hc_state_host, sizeof(float) * (size_t) hc_dim, cudaMemcpyHostToDevice) != cudaSuccess) { return false; } - if (cudaMemcpy(g_scratch.d_scale, scale_host, sizeof(float) * (size_t) mix_dim, + if (cudaMemcpy(scratch.d_scale, scale_host, sizeof(float) * (size_t) mix_dim, cudaMemcpyHostToDevice) != cudaSuccess) { return false; } - if (cudaMemcpy(g_scratch.d_base, base_host, sizeof(float) * (size_t) mix_dim, + if (cudaMemcpy(scratch.d_base, base_host, sizeof(float) * (size_t) mix_dim, cudaMemcpyHostToDevice) != cudaSuccess) { return false; } if (!hc_pre_device_locked( - g_scratch.d_state, + scratch, + scratch.d_state, fn_device, - g_scratch.d_scale, - g_scratch.d_base, + scratch.d_scale, + scratch.d_base, n_embd, n_hc, sinkhorn_iters, eps, - g_scratch.d_working, - g_scratch.d_post, - g_scratch.d_comb, + scratch.d_working, + scratch.d_post, + scratch.d_comb, false)) { return false; } - if (cudaMemcpy(working_host, g_scratch.d_working, sizeof(float) * (size_t) n_embd, + if (cudaMemcpy(working_host, scratch.d_working, sizeof(float) * (size_t) n_embd, cudaMemcpyDeviceToHost) != cudaSuccess) { return false; } - if (cudaMemcpy(post_host, g_scratch.d_post, sizeof(float) * (size_t) n_hc, + if (cudaMemcpy(post_host, scratch.d_post, sizeof(float) * (size_t) n_hc, cudaMemcpyDeviceToHost) != cudaSuccess) { return false; } - if (cudaMemcpy(comb_host, g_scratch.d_comb, sizeof(float) * (size_t) n_hc * (size_t) n_hc, + if (cudaMemcpy(comb_host, scratch.d_comb, sizeof(float) * (size_t) n_hc * (size_t) n_hc, cudaMemcpyDeviceToHost) != cudaSuccess) { return false; } @@ -477,11 +546,15 @@ bool deepseek4_cuda_hc_pre_device(const void * hc_state_device, return false; } - std::lock_guard lock(g_mu); - if (!g_scratch.ensure((size_t) hc_dim, (size_t) n_embd, (size_t) n_hc)) { + HcScratchSlot * slot = current_scratch_slot(); + if (!slot) return false; + std::lock_guard lock(slot->mutex); + HcCudaScratch & scratch = slot->scratch; + if (!scratch.ensure((size_t) hc_dim, (size_t) n_embd)) { return false; } - return hc_pre_device_locked(hc_state_device, + return hc_pre_device_locked(scratch, + hc_state_device, fn_device, scale_device, base_device, @@ -517,30 +590,34 @@ bool deepseek4_cuda_hc_pre_device_params(const void * hc_state_device, return false; } - std::lock_guard lock(g_mu); - if (!g_scratch.ensure((size_t) hc_dim, (size_t) n_embd, (size_t) n_hc)) { + HcScratchSlot * slot = current_scratch_slot(); + if (!slot) return false; + std::lock_guard lock(slot->mutex); + HcCudaScratch & scratch = slot->scratch; + if (!scratch.ensure((size_t) hc_dim, (size_t) n_embd)) { std::fprintf(stderr, "[deepseek4-hc-direct] ensure failed\n"); return false; } - if (cudaMemcpy(g_scratch.d_state, hc_state_device, sizeof(float) * (size_t) hc_dim, + if (cudaMemcpy(scratch.d_state, hc_state_device, sizeof(float) * (size_t) hc_dim, cudaMemcpyDeviceToDevice) != cudaSuccess) { hc_log_cuda_error("copy state d2d", cudaGetLastError()); return false; } - if (cudaMemcpy(g_scratch.d_scale, scale_host, sizeof(float) * (size_t) mix_dim, + if (cudaMemcpy(scratch.d_scale, scale_host, sizeof(float) * (size_t) mix_dim, cudaMemcpyHostToDevice) != cudaSuccess) { hc_log_cuda_error("copy scale", cudaGetLastError()); return false; } - if (cudaMemcpy(g_scratch.d_base, base_host, sizeof(float) * (size_t) mix_dim, + if (cudaMemcpy(scratch.d_base, base_host, sizeof(float) * (size_t) mix_dim, cudaMemcpyHostToDevice) != cudaSuccess) { hc_log_cuda_error("copy base", cudaGetLastError()); return false; } - return hc_pre_device_locked(g_scratch.d_state, + return hc_pre_device_locked(scratch, + scratch.d_state, fn_device, - g_scratch.d_scale, - g_scratch.d_base, + scratch.d_scale, + scratch.d_base, n_embd, n_hc, sinkhorn_iters, diff --git a/server/src/deepseek4/deepseek4_hc_cuda.h b/server/src/deepseek4/deepseek4_hc_cuda.h index d2268ac12..08766e1a4 100644 --- a/server/src/deepseek4/deepseek4_hc_cuda.h +++ b/server/src/deepseek4/deepseek4_hc_cuda.h @@ -4,6 +4,10 @@ namespace dflash::common { +// CUDA and HIP select devices per host thread. Call this at worker/shard entry +// before invoking the direct HC helpers. +bool deepseek4_cuda_hc_set_device(int device); + bool deepseek4_cuda_hc_pre_mix(const float * hc_state_host, const void * fn_device, int n_embd, diff --git a/server/src/deepseek4/deepseek4_internal.h b/server/src/deepseek4/deepseek4_internal.h index 07ed34f94..3d44f0e07 100644 --- a/server/src/deepseek4/deepseek4_internal.h +++ b/server/src/deepseek4/deepseek4_internal.h @@ -357,6 +357,7 @@ struct Ds4VerifyHooks; bool deepseek4_step( ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, const float * embed, @@ -382,6 +383,7 @@ struct Ds4VerifyHooks { bool deepseek4_step_layer_range( ggml_backend_t backend, + int device, const DeepSeek4Weights & w, DeepSeek4Cache & cache, std::vector & hc_state, diff --git a/server/src/deepseek4/deepseek4_layer_split_adapter.cpp b/server/src/deepseek4/deepseek4_layer_split_adapter.cpp index afbfe0c9b..f98feafc9 100644 --- a/server/src/deepseek4/deepseek4_layer_split_adapter.cpp +++ b/server/src/deepseek4/deepseek4_layer_split_adapter.cpp @@ -477,7 +477,7 @@ bool DeepSeek4LayerSplitAdapter::run_forward( const float * shard_input = local_shard_input(si, embed, hc_state_); if (!deepseek4_step_layer_range( - shard.backend, shard.weights, shard.cache, + shard.backend, shard.gpu, shard.weights, shard.cache, hc_state_, shard_input, n_tokens, base_pos, shard.layer_begin, shard.layer_end, shard_logits, tokens.data(), timing ? &step_tel : nullptr)) { @@ -527,7 +527,7 @@ bool DeepSeek4LayerSplitAdapter::run_mixed_forward( // Run only the local shard's layer range, getting hidden state output std::vector hidden_out; - if (!deepseek4_step_layer_range(local_shard.backend, local_shard.weights, + if (!deepseek4_step_layer_range(local_shard.backend, local_shard.gpu, local_shard.weights, local_shard.cache, hc_state_, embed.data(), n_tokens, base_pos, local_shard.layer_begin, local_shard.layer_end, diff --git a/server/src/deepseek4/deepseek4_target_shard_ipc_daemon.cpp b/server/src/deepseek4/deepseek4_target_shard_ipc_daemon.cpp index d6876db0b..d0d3b6dc2 100644 --- a/server/src/deepseek4/deepseek4_target_shard_ipc_daemon.cpp +++ b/server/src/deepseek4/deepseek4_target_shard_ipc_daemon.cpp @@ -238,7 +238,7 @@ int run_deepseek4_target_shard_ipc_daemon( "[deepseek4-target-shard] forward shard %zu: n_tokens=%d base_pos=%d layers=[%d,%d)\n", i, n_tokens, req.base_pos, shard.layer_begin, shard.layer_end); const bool ok = deepseek4_step_layer_range( - shard.backend, shard.weights, shard.cache, hc_state, + shard.backend, shard.gpu, shard.weights, shard.cache, hc_state, shard_input, n_tokens, req.base_pos, shard.layer_begin, shard.layer_end, is_last ? &logits : nullptr, diff --git a/server/tests/test_deepseek4_unit.cpp b/server/tests/test_deepseek4_unit.cpp index 7464722aa..4fc432600 100644 --- a/server/tests/test_deepseek4_unit.cpp +++ b/server/tests/test_deepseek4_unit.cpp @@ -28,6 +28,7 @@ #include #include #include +#include #include #include @@ -2854,6 +2855,253 @@ static void test_hc_pre_kernel_gpu() { ggml_free(ctx); ggml_backend_free(backend); } + +// The stale-boundary guard in deepseek4_step_layer_range must fire before any +// backend, cache, or weight access, so stub weights and a null backend are +// sufficient to pin it. +static void test_layer_range_rejects_stale_hc_boundary() { + std::fprintf(stderr, " test_layer_range_rejects_stale_hc_boundary ..."); + DeepSeek4Weights w; + w.n_layer = 4; + w.n_embd = 8; + w.n_hc = 4; + DeepSeek4Cache cache; + + const int n_tokens = 1; + const size_t hc_dim = (size_t) w.n_embd * (size_t) w.n_hc; + + // Later-shard call where embed aliases hc_state but the carried state has + // a stale size: must fail without resizing, because a resize would free + // the buffer embed points into. + std::vector hc_state(hc_dim / 2, 0.0f); + const float * stale_alias = hc_state.data(); + bool ok = deepseek4_step_layer_range( + /*backend=*/nullptr, /*device=*/0, w, cache, hc_state, + stale_alias, n_tokens, /*kv_start=*/0, + /*layer_begin=*/2, /*layer_end=*/4, /*out_logits=*/nullptr); + TEST_ASSERT_MSG(!ok, "stale aliased HC boundary state must be rejected"); + TEST_ASSERT_MSG(hc_state.size() == hc_dim / 2, + "rejected call must not resize the aliased HC state"); + TEST_ASSERT_MSG(cache.layer_range_cache == nullptr, + "guard must fire before the layer-range cache is created"); + + // Later-shard call with no boundary state at all. + std::vector missing_state; + ok = deepseek4_step_layer_range( + nullptr, 0, w, cache, missing_state, + /*embed=*/nullptr, n_tokens, 0, 2, 4, nullptr); + TEST_ASSERT_MSG(!ok, "missing HC boundary state must be rejected"); + TEST_ASSERT_MSG(cache.layer_range_cache == nullptr, + "guard must fire before the layer-range cache is created"); + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} + +#if defined(GGML_USE_CUDA) || defined(GGML_USE_HIP) +static void test_hc_scratch_per_device() { + std::fprintf(stderr, " test_hc_scratch_per_device ..."); + if (ggml_backend_cuda_get_device_count() < 2) { + std::fprintf(stderr, " skipped (requires two GPU devices)\n"); + return; + } + + constexpr int n_embd = 32; + constexpr int n_hc = 4; + constexpr int sinkhorn_iters = 6; + constexpr float hc_eps = 1.0e-6f; + constexpr int mix_dim = 2 * n_hc + n_hc * n_hc; + constexpr int hc_dim = n_embd * n_hc; + + std::vector hc_state((size_t) hc_dim); + std::iota(hc_state.begin(), hc_state.end(), -16.0f); + for (float & value : hc_state) value *= 0.01f; + std::vector fn_f16((size_t) mix_dim * (size_t) hc_dim, + ggml_fp32_to_fp16(0.0f)); + std::vector scale((size_t) mix_dim, 0.0f); + scale[0] = 1.0f; + scale[1] = 1.0f; + scale[2] = 1.0f; + std::vector base((size_t) mix_dim, 0.0f); + + struct Result { + bool ok = false; + std::vector working; + std::vector post; + std::vector comb; + + Result(int embd, int hc) + : working((size_t) embd), + post((size_t) hc), + comb((size_t) hc * (size_t) hc) {} + }; + + auto run_on_worker = [&](int device, Result & result) { + std::thread worker([&, device]() { + if (!deepseek4_cuda_hc_set_device(device)) { + return; + } + ggml_backend_t dev_backend = ggml_backend_cuda_init(device); + if (!dev_backend) { + return; + } + ggml_context * ctx = make_test_context(1u << 16); + ggml_tensor * fn_t = + ctx ? ggml_new_tensor_2d(ctx, GGML_TYPE_F16, hc_dim, mix_dim) : nullptr; + ggml_backend_buffer_t buf = + fn_t ? ggml_backend_alloc_ctx_tensors(ctx, dev_backend) : nullptr; + if (buf) { + ggml_backend_tensor_set(fn_t, fn_f16.data(), 0, + fn_f16.size() * sizeof(ggml_fp16_t)); + result.ok = deepseek4_cuda_hc_pre( + hc_state.data(), fn_t->data, scale.data(), base.data(), + n_embd, n_hc, sinkhorn_iters, hc_eps, + result.working.data(), result.post.data(), result.comb.data()); + ggml_backend_buffer_free(buf); + } + if (ctx) ggml_free(ctx); + ggml_backend_free(dev_backend); + }); + worker.join(); + }; + + Result first_device(n_embd, n_hc); + Result second_device(n_embd, n_hc); + Result first_device_again(n_embd, n_hc); + run_on_worker(0, first_device); + run_on_worker(1, second_device); + run_on_worker(0, first_device_again); + + TEST_ASSERT_MSG(first_device.ok, "HC direct call failed on GPU device 0"); + TEST_ASSERT_MSG(second_device.ok, "HC direct call failed on GPU device 1"); + TEST_ASSERT_MSG(first_device_again.ok, "HC direct call failed after returning to GPU device 0"); + if (first_device.ok && second_device.ok && first_device_again.ok) { + for (int i = 0; i < n_embd; ++i) { + TEST_ASSERT(nearly_equal(first_device.working[(size_t) i], + second_device.working[(size_t) i])); + TEST_ASSERT(nearly_equal(first_device.working[(size_t) i], + first_device_again.working[(size_t) i])); + } + for (int i = 0; i < n_hc; ++i) { + TEST_ASSERT(nearly_equal(first_device.post[(size_t) i], + second_device.post[(size_t) i])); + } + for (int i = 0; i < n_hc * n_hc; ++i) { + TEST_ASSERT(nearly_equal(first_device.comb[(size_t) i], + second_device.comb[(size_t) i])); + } + } + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} + +static void test_hc_set_device_contract() { + std::fprintf(stderr, " test_hc_set_device_contract ..."); + // Device selection is per host thread; run on a worker so the rejected + // selections cannot leave sticky CUDA error state on the main thread. + std::thread worker([]() { + TEST_ASSERT_MSG(!deepseek4_cuda_hc_set_device(-1), + "negative device must be rejected"); + const int device_count = ggml_backend_cuda_get_device_count(); + TEST_ASSERT_MSG(!deepseek4_cuda_hc_set_device(device_count), + "out-of-range device must be rejected"); + if (device_count > 0) { + TEST_ASSERT_MSG(deepseek4_cuda_hc_set_device(0), + "selecting device 0 failed"); + } + }); + worker.join(); + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} + +static void test_hc_scratch_shape_capacity() { + std::fprintf(stderr, " test_hc_scratch_shape_capacity ..."); + if (!deepseek4_cuda_hc_set_device(0)) { + std::fprintf(stderr, " skipped (no GPU device)\n"); + return; + } + ggml_backend_t backend = ggml_backend_cuda_init(0); + if (!backend) { + std::fprintf(stderr, " skipped (no GPU backend)\n"); + return; + } + + constexpr int sinkhorn_iters = 6; + constexpr float hc_eps = 1.0e-6f; + // Grow both shape axes past the first call, then shrink back: every run + // must stay correct while the per-device scratch reallocates and reuses. + const int shapes[][2] = {{32, 2}, {192, 8}, {32, 2}}; + + std::mt19937 rng(321); + std::uniform_real_distribution dist(-0.2f, 0.2f); + + for (const auto & shape : shapes) { + const int n_embd = shape[0]; + const int n_hc = shape[1]; + const int hc_dim = n_embd * n_hc; + const int mix_dim = 2 * n_hc + n_hc * n_hc; + + std::vector hc_state((size_t) hc_dim); + std::vector fn_f16((size_t) mix_dim * (size_t) hc_dim); + std::vector scale((size_t) mix_dim, 0.0f); + std::vector base((size_t) mix_dim); + for (float & v : hc_state) v = dist(rng); + for (ggml_fp16_t & v : fn_f16) v = ggml_fp32_to_fp16(dist(rng)); + scale[0] = 0.9f; + scale[1] = 1.05f; + scale[2] = 1.0f; + for (float & v : base) v = 0.1f * dist(rng); + + std::vector ref_working; + std::vector ref_post; + std::vector ref_comb; + test_reference_hc_pre(hc_state, fn_f16, scale, base, n_embd, n_hc, + sinkhorn_iters, hc_eps, + ref_working, ref_post, ref_comb); + + ggml_context * ctx = make_test_context(1u << 16); + TEST_ASSERT_MSG(ctx != nullptr, "ggml_init failed"); + if (!ctx) break; + ggml_tensor * fn_t = ggml_new_tensor_2d(ctx, GGML_TYPE_F16, hc_dim, mix_dim); + ggml_backend_buffer_t buf = ggml_backend_alloc_ctx_tensors(ctx, backend); + TEST_ASSERT_MSG(buf != nullptr, "ggml backend buffer alloc failed"); + if (!buf) { + ggml_free(ctx); + break; + } + ggml_backend_tensor_set(fn_t, fn_f16.data(), 0, + fn_f16.size() * sizeof(ggml_fp16_t)); + + std::vector working((size_t) n_embd); + std::vector post((size_t) n_hc); + std::vector comb((size_t) n_hc * (size_t) n_hc); + const bool ok = deepseek4_cuda_hc_pre( + hc_state.data(), fn_t->data, scale.data(), base.data(), + n_embd, n_hc, sinkhorn_iters, hc_eps, + working.data(), post.data(), comb.data()); + TEST_ASSERT_MSG(ok, "HC pre failed after scratch shape change"); + if (ok) { + for (int i = 0; i < n_embd; ++i) { + TEST_ASSERT_MSG( + nearly_equal(working[(size_t) i], ref_working[(size_t) i], 2.0e-4f, 2.0e-4f), + "working mismatch after scratch shape change"); + } + for (int i = 0; i < n_hc; ++i) { + TEST_ASSERT_MSG( + nearly_equal(post[(size_t) i], ref_post[(size_t) i], 2.0e-4f, 2.0e-4f), + "post mismatch after scratch shape change"); + } + for (int i = 0; i < n_hc * n_hc; ++i) { + TEST_ASSERT_MSG( + nearly_equal(comb[(size_t) i], ref_comb[(size_t) i], 2.0e-4f, 2.0e-4f), + "comb mismatch after scratch shape change"); + } + } + ggml_backend_buffer_free(buf); + ggml_free(ctx); + } + + ggml_backend_free(backend); + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} +#endif #endif int main() { @@ -2906,6 +3154,10 @@ int main() { test_ds4_flash_attention_inverse_rope_fallback_gpu(); test_hc_post_strided_split_gpu(); test_hc_pre_kernel_gpu(); + test_layer_range_rejects_stale_hc_boundary(); + test_hc_scratch_per_device(); + test_hc_set_device_contract(); + test_hc_scratch_shape_capacity(); #endif ggml_backend_free(backend); From b47e409429013b1d667c2673aa73dbbca239f021 Mon Sep 17 00:00:00 2001 From: Graffioh Date: Sun, 19 Jul 2026 11:12:52 +0200 Subject: [PATCH 4/4] feat(ds4): opt-in GPU F16 HC fn mirrors for ROCmFPX ROCmFPX-quantized HC fn tensors cannot feed the __half HC pre kernels, so HC pre always fell back to CPU on ROCmFPX builds. Behind DFLASH_DS4_ROCMFPX_HC_GPU (default off), upload the existing host F16 copy of each in-range layer's attn/ffn fn weights once per shard cache and let the GPU HC pre path consume it. Mirrors are owned by the shard's DeepSeek4LayerRangeCache, keyed to the owning device, and freed on that device. A failed upload logs an explicit CPU-HC fallback and degrades gracefully. The output HC projection is not mirrored: it runs once per step, not per layer. Co-Authored-By: Claude Fable 5 Claude-Session: https://claude.ai/code/session_01Mx99uuhcriPweAFggPwfce --- server/src/deepseek4/deepseek4_graph.cpp | 60 ++++++++++++++++++++++- server/src/deepseek4/deepseek4_hc_cuda.cu | 40 +++++++++++++++ server/src/deepseek4/deepseek4_hc_cuda.h | 10 ++++ 3 files changed, 108 insertions(+), 2 deletions(-) diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index b14af4474..4e0ef0ac0 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -54,6 +54,11 @@ static bool ds4_env_flag(const char * name) { return value && value[0] && std::strcmp(value, "0") != 0; } +static bool ds4_rocmfpx_hc_gpu_enabled() { + static const bool enabled = ds4_env_flag("DFLASH_DS4_ROCMFPX_HC_GPU"); + return enabled; +} + static int ds4_effective_expert_count(const DeepSeek4Weights & w) { const int requested = w.routed_expert_top_k; if (requested > 0 && requested < w.n_expert_used) { @@ -3016,6 +3021,9 @@ struct HcWeightsCpu { std::vector fn_data; // [hc_dim * mix_dim] F16 std::vector scale_data; // [3] std::vector base_data; // [2*n_hc + n_hc*n_hc] + void * fn_f16_device = nullptr; // Persistent F16 mirror for quantized HC fn tensors. + size_t fn_f16_device_bytes = 0; + int fn_f16_device_id = -1; bool loaded = false; }; @@ -3723,7 +3731,19 @@ static bool load_hc_weights_cpu(HcWeightsCpu & dst, ggml_tensor * fn, return true; } +static void release_hc_fn_device(HcWeightsCpu & w) { +#if defined(DFLASH27B_BACKEND_CUDA) || defined(DFLASH27B_BACKEND_HIP) || defined(GGML_USE_HIP) + if (w.fn_f16_device) { + deepseek4_cuda_hc_free(w.fn_f16_device_id, w.fn_f16_device); + } +#endif + w.fn_f16_device = nullptr; + w.fn_f16_device_bytes = 0; + w.fn_f16_device_id = -1; +} + static void reset_hc_weights_cpu(HcWeightsCpu & w) { + release_hc_fn_device(w); w.fn_data.clear(); w.scale_data.clear(); w.base_data.clear(); @@ -3755,10 +3775,42 @@ struct DeepSeek4HybridRuntime { static thread_local DeepSeek4HybridRuntime ds4_hybrid_runtime; -static const void * hc_fn_device_ptr(const HcWeightsCpu &, ggml_tensor * fn) { +static bool ensure_hc_fn_device(HcWeightsCpu & w, ggml_tensor * fn, int device) { + if (!fn || !fn->data) return false; + if (fn->type == GGML_TYPE_F16) return true; + if (!ds4_rocmfpx_hc_gpu_enabled()) return false; + if (!w.loaded || w.fn_data.empty()) return false; + +#if defined(DFLASH27B_BACKEND_CUDA) || defined(DFLASH27B_BACKEND_HIP) || defined(GGML_USE_HIP) + const size_t bytes = w.fn_data.size() * sizeof(uint16_t); + if (w.fn_f16_device && w.fn_f16_device_bytes == bytes && + w.fn_f16_device_id == device) { + return true; + } + release_hc_fn_device(w); + if (!deepseek4_cuda_hc_upload_f16(device, w.fn_data.data(), bytes, &w.fn_f16_device)) { + std::fprintf(stderr, + "[deepseek4] HC fn F16 mirror upload failed on device %d; " + "falling back to CPU HC\n", + device); + w.fn_f16_device = nullptr; + w.fn_f16_device_bytes = 0; + w.fn_f16_device_id = -1; + return false; + } + w.fn_f16_device_bytes = bytes; + w.fn_f16_device_id = device; + return true; +#else + return false; +#endif +} + +static const void * hc_fn_device_ptr(const HcWeightsCpu & w, ggml_tensor * fn) { if (!fn) return nullptr; if (fn->type == GGML_TYPE_F16) return fn->data; - return nullptr; + if (!ds4_rocmfpx_hc_gpu_enabled()) return nullptr; + return w.fn_f16_device; } static bool load_hash_routing_cpu(HashRoutingTableCpu & dst, ggml_tensor * table) { @@ -5788,6 +5840,10 @@ static bool initialize_layer_range_cache( runtime.reset(); return false; } + if (ds4_backend_is_gpu(backend) && !ds4_env_flag("DFLASH_DS4_HC_CPU")) { + ensure_hc_fn_device(cached.attn, layer.hc_attn_fn, device); + ensure_hc_fn_device(cached.ffn, layer.hc_ffn_fn, device); + } if (il < w.n_hash_layer && layer.ffn_gate_tid2eid) { load_hash_routing_cpu(runtime.hash_routing_tables[(size_t)il], layer.ffn_gate_tid2eid); diff --git a/server/src/deepseek4/deepseek4_hc_cuda.cu b/server/src/deepseek4/deepseek4_hc_cuda.cu index 1462092eb..895b56941 100644 --- a/server/src/deepseek4/deepseek4_hc_cuda.cu +++ b/server/src/deepseek4/deepseek4_hc_cuda.cu @@ -628,4 +628,44 @@ bool deepseek4_cuda_hc_pre_device_params(const void * hc_state_device, true); } +bool deepseek4_cuda_hc_upload_f16(int device, + const void * host_f16, + size_t bytes, + void ** device_out) { + if (!host_f16 || bytes == 0 || !device_out) { + return false; + } + if (!deepseek4_cuda_hc_set_device(device)) { + return false; + } + void * ptr = nullptr; + cudaError_t err = cudaMalloc(&ptr, bytes); + if (err != cudaSuccess) { + hc_log_cuda_error("hc f16 mirror malloc", err); + return false; + } + err = cudaMemcpy(ptr, host_f16, bytes, cudaMemcpyHostToDevice); + if (err != cudaSuccess) { + hc_log_cuda_error("hc f16 mirror upload", err); + cudaFree(ptr); + return false; + } + *device_out = ptr; + return true; +} + +void deepseek4_cuda_hc_free(int device, void * device_ptr) { + if (!device_ptr) { + return; + } + // Best effort: unified addressing lets cudaFree release memory from any + // current device, so never leak the mirror just because the owning device + // could not be selected (e.g. during shutdown). + (void) deepseek4_cuda_hc_set_device(device); + const cudaError_t err = cudaFree(device_ptr); + if (err != cudaSuccess) { + hc_log_cuda_error("hc f16 mirror free", err); + } +} + } // namespace dflash::common diff --git a/server/src/deepseek4/deepseek4_hc_cuda.h b/server/src/deepseek4/deepseek4_hc_cuda.h index 08766e1a4..6f2718bc6 100644 --- a/server/src/deepseek4/deepseek4_hc_cuda.h +++ b/server/src/deepseek4/deepseek4_hc_cuda.h @@ -1,5 +1,6 @@ #pragma once +#include #include namespace dflash::common { @@ -51,4 +52,13 @@ bool deepseek4_cuda_hc_pre_device_params(const void * hc_state_device, void * post_device, void * comb_device); +// Persistent F16 device mirrors for quantized HC fn weights: upload the host +// F16 copy to the given device, and free it on that device later. +bool deepseek4_cuda_hc_upload_f16(int device, + const void * host_f16, + size_t bytes, + void ** device_out); + +void deepseek4_cuda_hc_free(int device, void * device_ptr); + } // namespace dflash::common