diff --git a/README.md b/README.md index 841eb5f2b..066362a8e 100644 --- a/README.md +++ b/README.md @@ -100,6 +100,10 @@ Use this table to choose a profile. The linked guides contain the complete build | Exact heterogeneous AMD | RX 7900 XT + Strix Halo + DeepSeek V4 | Use the qualified true top-6 profile in [`serve_ds4_dual_rocm_128k.sh`](server/scripts/serve_ds4_dual_rocm_128k.sh). | | Lucebox heterogeneous profile | R9700 + Strix Halo + DeepSeek V4 | Use the opt-in in-process expert-parallel profile. Top-4 routing and sparse prefill remain explicit approximations. [DS4 guide](server/docs/DS4.md#in-process-heterogeneous-expert-parallel) | +The DS4 guide also documents the Strix long-context sparse-verifier profile and +Qwen3-0.6B PFlash integration. PFlash is lossy prompt compression; keep it off +for exact-retrieval and matched true-context benchmarks. + ## Client Harnesses [`harness/`](harness/) runs Lucebox through popular coding clients and checks server compatibility. diff --git a/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu index e7286de55..5f138d6a9 100644 --- a/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu +++ b/server/deps/llama.cpp/ggml/src/ggml-cuda/ds4-indexer.cu @@ -351,12 +351,14 @@ static __global__ void ds4_indexer_score_decode_wmma_kernel( } } -// The speculative verifier scores exactly four query tokens. The general -// WMMA kernel places those four tokens in a 16-row tile and executes twelve -// zero rows for every head. Pack four consecutive heads into the tile instead: -// row = 4*head_in_group + token. The post-WMMA loop still accumulates heads in -// their original order, preserving the established F32 numerical topology. -static __global__ void ds4_indexer_score_wmma_q4_kernel( +// The speculative verifier scores only a few query tokens. The general WMMA +// kernel places them in a 16-row tile and executes the unused rows for every +// head. Pack consecutive heads into the tile instead: +// row = N_TOKENS*head_in_group + token. The post-WMMA loop still accumulates +// heads in their original order, preserving the established F32 numerical +// topology. This is the HIP equivalent of the Vulkan small-CM dispatch. +template +static __global__ void ds4_indexer_score_wmma_small_kernel( float * scores, const float * q, const float * weights, @@ -375,7 +377,14 @@ static __global__ void ds4_indexer_score_wmma_q4_kernel( __shared__ float c_sh[8 * 16 * 16]; __shared__ float weight_sh[16]; - float acc[2] = {0.0f, 0.0f}; + static_assert(N_TOKENS >= 2 && N_TOKENS <= 5, + "small-CM kernel is specialized for verifier widths 2..5"); + constexpr int HEADS_PER_TILE = 16 / N_TOKENS; + constexpr int USED_ROWS = HEADS_PER_TILE * N_TOKENS; + constexpr int ACC_SLOTS = (N_TOKENS + 1) / 2; + float acc[ACC_SLOTS]; +#pragma unroll + for (int slot = 0; slot < ACC_SLOTS; ++slot) acc[slot] = 0.0f; for (int i = tid; i < 128 * 128; i += 256) { const int c = i >> 7; @@ -387,21 +396,30 @@ static __global__ void ds4_indexer_score_wmma_q4_kernel( } __syncthreads(); - for (int head_base = 0; head_base < n_head; head_base += 4) { + for (int head_base = 0; head_base < n_head; + head_base += HEADS_PER_TILE) { for (int pair = tid; pair < 16 * 64; pair += 256) { const int row = pair >> 6; const int d = (pair & 63) * 2; - const int token = row & 3; - const int head = head_base + (row >> 2); - const float2 q_value = *reinterpret_cast( - q + ((size_t) token * n_head + head) * 128 + d); - *reinterpret_cast(a_sh + row * 128 + d) = - __floats2half2_rn(q_value.x, q_value.y); + half2 value = __float2half2_rn(0.0f); + if (row < USED_ROWS) { + const int token = row % N_TOKENS; + const int head = head_base + row / N_TOKENS; + if (head < n_head) { + const float2 q_value = + *reinterpret_cast( + q + ((size_t) token * n_head + head) * 128 + d); + value = __floats2half2_rn(q_value.x, q_value.y); + } + } + *reinterpret_cast(a_sh + row * 128 + d) = value; } if (tid < 16) { - const int token = tid & 3; - const int head = head_base + (tid >> 2); - weight_sh[tid] = weights[(size_t) token * n_head + head]; + const int token = tid % N_TOKENS; + const int head = head_base + tid / N_TOKENS; + weight_sh[tid] = tid < USED_ROWS && head < n_head + ? weights[(size_t) token * n_head + head] + : 0.0f; } __syncthreads(); @@ -431,16 +449,17 @@ static __global__ void ds4_indexer_score_wmma_q4_kernel( __syncthreads(); int slot = 0; - for (int output = tid; output < 4 * 128; + for (int output = tid; output < N_TOKENS * 128; output += 256, ++slot) { const int token = output >> 7; const int local_comp = output & 127; const int comp_tile = local_comp >> 4; const int comp_col = local_comp & 15; #pragma unroll - for (int head_in_group = 0; head_in_group < 4; + for (int head_in_group = 0; + head_in_group < HEADS_PER_TILE; ++head_in_group) { - const int row = 4 * head_in_group + token; + const int row = N_TOKENS * head_in_group + token; const float dot = c_sh[ comp_tile * 16 * 16 + row * 16 + comp_col]; acc[slot] += fmaxf(dot, 0.0f) * weight_sh[row]; @@ -450,7 +469,7 @@ static __global__ void ds4_indexer_score_wmma_q4_kernel( } int slot = 0; - for (int output = tid; output < 4 * 128; + for (int output = tid; output < N_TOKENS * 128; output += 256, ++slot) { const int token = output >> 7; const int comp = tile_c + (output & 127); @@ -553,6 +572,17 @@ void ggml_cuda_op_ds4_indexer_score( warp_size == 32 && (!GGML_CUDA_CC_IS_NVIDIA(device_info.cc) || device_info.cc >= GGML_CUDA_CC_VOLTA); + const char * packed_small_name = "GGML_DS4_INDEXER_PACK_SMALL"; + const char * packed_small_env = std::getenv(packed_small_name); + if (!packed_small_env) { + // Backward-compatible alias for the original q=4-only prototype. + packed_small_name = "GGML_DS4_INDEXER_PACK_Q4"; + packed_small_env = std::getenv(packed_small_name); + } + const bool use_packed_small = packed_small_env + ? ds4_env_flag_enabled(packed_small_name) + : GGML_CUDA_CC_IS_RDNA3_5(device_info.cc) || + GGML_CUDA_CC_IS_RDNA4(device_info.cc); #if DS4_INDEXER_WMMA_AVAILABLE if (wmma_capable && n_tokens == 1) { const dim3 grid((unsigned) ((n_comp + 127) / 128), 1, 1); @@ -564,10 +594,39 @@ void ggml_cuda_op_ds4_indexer_score( visibility_mask ? static_cast(visibility_mask->data) : nullptr, n_comp, kv_start, n_head, ratio); - } else if (wmma_capable && n_tokens == 4 && n_head % 4 == 0 && - ds4_env_flag_enabled("GGML_DS4_INDEXER_PACK_Q4")) { + } else if (wmma_capable && n_tokens == 2 && use_packed_small) { + const dim3 grid((unsigned) ((n_comp + 127) / 128), 1, 1); + ds4_indexer_score_wmma_small_kernel<2><<>>( + static_cast(dst->data), + static_cast(q->data), + static_cast(weights->data), + static_cast(comp->data), + visibility_mask + ? static_cast(visibility_mask->data) : nullptr, + n_comp, kv_start, n_head, ratio); + } else if (wmma_capable && n_tokens == 3 && use_packed_small) { + const dim3 grid((unsigned) ((n_comp + 127) / 128), 1, 1); + ds4_indexer_score_wmma_small_kernel<3><<>>( + static_cast(dst->data), + static_cast(q->data), + static_cast(weights->data), + static_cast(comp->data), + visibility_mask + ? static_cast(visibility_mask->data) : nullptr, + n_comp, kv_start, n_head, ratio); + } else if (wmma_capable && n_tokens == 4 && use_packed_small) { + const dim3 grid((unsigned) ((n_comp + 127) / 128), 1, 1); + ds4_indexer_score_wmma_small_kernel<4><<>>( + static_cast(dst->data), + static_cast(q->data), + static_cast(weights->data), + static_cast(comp->data), + visibility_mask + ? static_cast(visibility_mask->data) : nullptr, + n_comp, kv_start, n_head, ratio); + } else if (wmma_capable && n_tokens == 5 && use_packed_small) { const dim3 grid((unsigned) ((n_comp + 127) / 128), 1, 1); - ds4_indexer_score_wmma_q4_kernel<<>>( + ds4_indexer_score_wmma_small_kernel<5><<>>( static_cast(dst->data), static_cast(q->data), static_cast(weights->data), diff --git a/server/docs/DS4.md b/server/docs/DS4.md index 11be8699d..3abd1155c 100644 --- a/server/docs/DS4.md +++ b/server/docs/DS4.md @@ -455,6 +455,7 @@ Run the converted drafter against a DeepSeek4 target with: ```bash export DFLASH_DS4_SPEC=1 export DFLASH_DS4_FUSED_VERIFY=1 +export DFLASH_DS4_SPARSE_DECODE_FLASH=1 export DFLASH_DS4_DRAFT=/path/to/dflash-draft.gguf export DFLASH_DS4_SPEC_Q=4 @@ -465,12 +466,43 @@ export DFLASH_DS4_SPEC_Q=4 ``` `--ds4-fused-verify-f16-kv` feeds the persistent F16 MLA cache directly to -batched explicit verifier attention instead of converting the full cache to -F32 on every speculative step. Key-side accumulation remains F32 through 512 -attention rows to preserve the short-context quality baseline. The option is -currently qualified only for a single HIP target and remains off by default. -It changes verifier floating-point inputs and can change generated tokens, so -re-run workload quality checks before enabling it for another checkpoint. +batched explicit or sparse verifier attention instead of converting the full +cache to F32 on every speculative step. With +`DFLASH_DS4_SPARSE_DECODE_FLASH=1`, the verifier keeps explicit attention for +short histories and switches each layer to the model's sparse top-k attention +once it removes at least half of the compressed rows. Key-side accumulation +remains F32 through 512 attention rows to preserve the short-context quality +baseline. The option is currently qualified only for a single HIP target and +remains off by default. It changes verifier floating-point inputs and can +change generated tokens, so re-run workload quality checks before enabling it +for another checkpoint. + +On RDNA3.5 and RDNA4, speculative widths 2–5 use the packed small-CM rocWMMA +indexer by default. It is bit-identical to the generic indexer in the GPU unit +test and can be disabled with `GGML_DS4_INDEXER_PACK_SMALL=0` for diagnosis. +The legacy `GGML_DS4_INDEXER_PACK_Q4` variable remains an alias. + +### PFlash prompt compression + +DeepSeek4 supports the shared in-process Qwen3-0.6B PFlash scorer: + +```bash +./server/build-hip/dflash_server /path/to/deepseek4-target.gguf \ + --target-device hip:0 \ + --prefill-compression auto \ + --prefill-drafter /path/to/Qwen3-0.6B-BF16.gguf \ + --prefill-skip-park +``` + +The HTTP path converts target tokens to text, scores Qwen tokens, decodes the +kept Qwen spans, and tokenizes that text for DeepSeek4. This cross-tokenizer +round trip is required; drafter token IDs are never passed directly to the +target. Omit `--prefill-skip-park` when the target, DSpark drafter, and PFlash +drafter do not fit together. + +PFlash reduces TTFT and the effective context used during generation, but it +is lossy prompt compression. Disable it for matched true-context throughput +or exact-retrieval comparisons. `DFLASH_DS4_FUSED_VERIFY=1` is the opt-in throughput profile. Its persistent whole-model GPU graph uses stable padded reduction shapes, so near-tied greedy diff --git a/server/src/common/model_capabilities.h b/server/src/common/model_capabilities.h index f62087141..752b0dc21 100644 --- a/server/src/common/model_capabilities.h +++ b/server/src/common/model_capabilities.h @@ -76,7 +76,7 @@ inline constexpr ArchCapabilities kArchCapabilities[] = { {"laguna", true, false, false, true, kMono, kMono, kMono, kNever, kNever, kNever, kNever}, {"qwen3", false, false, true, false, kNever, kNever, kNever, kNever, kNever, kNever, kNever}, {"gemma4", true, false, false, false, kMono, kNever, kNever, kNever, kBoth, kNever, kNever}, - {"deepseek4", true, false, false, false, kNever, kNever, kNever, kNever, kNever, kNever, kNever}, + {"deepseek4", true, false, true, false, kNever, kNever, kNever, kNever, kNever, kNever, kNever}, }; inline constexpr std::size_t kArchCount = diff --git a/server/src/deepseek4/deepseek4_backend.cpp b/server/src/deepseek4/deepseek4_backend.cpp index c8750bb28..f2ed3ac08 100644 --- a/server/src/deepseek4/deepseek4_backend.cpp +++ b/server/src/deepseek4/deepseek4_backend.cpp @@ -4,7 +4,9 @@ #include "deepseek4_backend.h" #include "deepseek4_budget_hook.h" #include "deepseek4_internal.h" +#include "dflash27b.h" #include "common/dynamic_backend.h" +#include "common/io_utils.h" #include "common/peer_access.h" #include "common/platform_env.h" #include "common/sampler.h" @@ -27,6 +29,7 @@ #include #include #include +#include namespace dflash::common { @@ -1683,6 +1686,11 @@ bool DeepSeek4Backend::park(ParkTarget target) { std::printf("[deepseek4] DSpark drafter parked (VRAM released)\n"); std::fflush(stdout); } + if (want_draft && pflash_drafter_loaded_) { + release_pflash_drafter(); + std::printf("[deepseek4] PFlash drafter parked (VRAM released)\n"); + std::fflush(stdout); + } if (!want_target_model || parked_) return true; maybe_save_routing_stats(); @@ -2599,17 +2607,150 @@ GenerateResult DeepSeek4Backend::restore_and_generate_impl( return result; } +ModelBackend::CompressResult DeepSeek4Backend::compress( + const CompressRequest & req) { + const auto results = compress_batch({req}); + return results.empty() ? CompressResult{} : results.front(); +} + +std::vector DeepSeek4Backend::compress_batch( + const std::vector & requests) { + std::vector results(requests.size()); + if (requests.empty()) return results; + + const CompressRequest * load_request = nullptr; + for (const CompressRequest & request : requests) { + if (request.input_ids.empty() || request.drafter_path.empty() || + request.keep_ratio <= 0.0f || request.keep_ratio > 1.0f) { + continue; + } + if (load_request == nullptr) { + load_request = &request; + } else if (request.drafter_path != load_request->drafter_path || + request.drafter_gpu != load_request->drafter_gpu || + request.skip_park != load_request->skip_park || + request.residency_action != load_request->residency_action) { + // A residency window can host only one drafter configuration. + // Process heterogeneous batches as independent windows instead of + // calling the virtual base fallback, which would recurse through + // DeepSeek4Backend::compress(). + std::vector independent(requests.size()); + for (size_t index = 0; index < requests.size(); ++index) { + const auto one = compress_batch({requests[index]}); + if (!one.empty()) independent[index] = one.front(); + } + return independent; + } + } + if (load_request == nullptr) return results; + + const bool was_parked = parked_; + if (!load_request->skip_park && !parked_ && + !park(ParkTarget::TargetModel)) { + return results; + } + if (backend_) ggml_backend_synchronize(backend_); + if (spec_backend_) ggml_backend_synchronize(spec_backend_); + + if (pflash_drafter_loaded_ && + (pflash_drafter_path_ != load_request->drafter_path || + pflash_drafter_gpu_ != load_request->drafter_gpu)) { + release_pflash_drafter(); + } + if (!pflash_drafter_loaded_) { + if (!load_drafter(load_request->drafter_path, 999, + load_request->drafter_gpu, + pflash_drafter_ctx_)) { + std::fprintf(stderr, "[deepseek4-pflash] load failed: %s\n", + dflash27b_last_error()); + if (!load_request->skip_park && !was_parked) { + unpark(ParkTarget::TargetModel); + } + return results; + } + pflash_drafter_loaded_ = true; + pflash_drafter_path_ = load_request->drafter_path; + pflash_drafter_gpu_ = load_request->drafter_gpu; + } + + for (size_t index = 0; index < requests.size(); ++index) { + const CompressRequest & request = requests[index]; + if (request.input_ids.empty() || request.drafter_path.empty() || + request.keep_ratio <= 0.0f || request.keep_ratio > 1.0f) { + continue; + } + CompressResult & result = results[index]; + result.compressed_ids = drafter_score_and_compress( + pflash_drafter_ctx_, request.input_ids, request.keep_ratio); + result.ok = !result.compressed_ids.empty(); + } + + if (load_request->residency_action == + DraftResidencyAction::ReleaseAfterUse) { + release_pflash_drafter(); + } + if (!load_request->skip_park && !was_parked && + !unpark(ParkTarget::TargetModel)) { + std::fill(results.begin(), results.end(), CompressResult{}); + } + return results; +} + bool DeepSeek4Backend::handle_compress(const std::string & line, - const DaemonIO & io) { - (void)line; (void)io; - std::fprintf(stderr, "[deepseek4] compress not yet supported\n"); - return false; + const DaemonIO & io) { + std::istringstream iss(line.size() > 9 ? line.substr(9) : std::string{}); + std::string prompt_path; + std::string drafter_path; + int keep_x1000 = 0; + if (!(iss >> prompt_path >> keep_x1000)) { + std::fprintf(stderr, "[deepseek4-pflash] bad compress arguments\n"); + io.emit(-1); + return false; + } + + std::getline(iss >> std::ws, drafter_path); + bool skip_park = false; + const std::string suffix = " nopark"; + if (drafter_path.size() > suffix.size() && + drafter_path.compare(drafter_path.size() - suffix.size(), + suffix.size(), suffix) == 0) { + skip_park = true; + drafter_path.resize(drafter_path.size() - suffix.size()); + } + + CompressRequest req; + req.input_ids = read_int32_file(prompt_path); + req.keep_ratio = (float) keep_x1000 / 1000.0f; + req.drafter_path = std::move(drafter_path); + req.skip_park = skip_park; + CompressResult result = compress(req); + if (!result.ok) { + std::fprintf(stderr, "[deepseek4-pflash] compression failed\n"); + io.emit(-1); + return false; + } + + std::printf("[deepseek4-pflash] %zu -> %zu tokens\n", + req.input_ids.size(), result.compressed_ids.size()); + std::fflush(stdout); + for (int32_t token : result.compressed_ids) io.emit(token); + io.emit(-1); + return true; +} + +void DeepSeek4Backend::release_pflash_drafter() { + if (!pflash_drafter_loaded_) return; + dflash::common::free_drafter(pflash_drafter_ctx_); + pflash_drafter_loaded_ = false; + pflash_drafter_path_.clear(); + pflash_drafter_gpu_ = -1; } void DeepSeek4Backend::free_drafter() { // Keep the configured path so request-scoped residency and an explicit // later `unpark draft` can restore the DSpark model. release_spec_drafter(/*mark_parked=*/true); + release_pflash_drafter(); } void DeepSeek4Backend::maybe_save_routing_stats() { diff --git a/server/src/deepseek4/deepseek4_backend.h b/server/src/deepseek4/deepseek4_backend.h index 583132042..84f0745c6 100644 --- a/server/src/deepseek4/deepseek4_backend.h +++ b/server/src/deepseek4/deepseek4_backend.h @@ -15,6 +15,7 @@ #include "../common/moe_hybrid_stream.h" #include "deepseek4_internal.h" #include "deepseek4_dspark.h" +#include "qwen3/qwen3_drafter.h" #include "ggml.h" #include "ggml-backend.h" @@ -70,6 +71,9 @@ class DeepSeek4Backend : public ModelBackend { const GenerateRequest & req, const DaemonIO & io) override; + CompressResult compress(const CompressRequest & req) override; + std::vector compress_batch( + const std::vector & requests) override; bool handle_compress(const std::string & line, const DaemonIO & io) override; void free_drafter() override; @@ -114,12 +118,17 @@ class DeepSeek4Backend : public ModelBackend { ggml_backend_t spec_backend_ = nullptr; std::unique_ptr spec_drafter_; std::vector spec_feat_window_; + DrafterContext pflash_drafter_ctx_; + bool pflash_drafter_loaded_ = false; + std::string pflash_drafter_path_; + int pflash_drafter_gpu_ = -1; // Once a long prompt selects the fragmentation-safe prefill shape, retain // it for later requests so the HIP arenas never switch back under load. int hybrid_prefill_chunk_cap_ = 0; bool load_spec_drafter(); void release_spec_drafter(bool mark_parked); + void release_pflash_drafter(); void keep_spec_feature_tail(std::vector & features, size_t max_rows) const; // True when a wide prefill path returns per-token DSpark features and the diff --git a/server/src/deepseek4/deepseek4_fused_verify.inc b/server/src/deepseek4/deepseek4_fused_verify.inc index ee270cc7d..e732fa6bc 100644 --- a/server/src/deepseek4/deepseek4_fused_verify.inc +++ b/server/src/deepseek4/deepseek4_fused_verify.inc @@ -534,13 +534,30 @@ static bool ds4_build_fused_verify_graph( std::vector i32b; std::vector i32ab; std::vector i64ab; + std::vector f32ab; + // Sparse attention pays a fixed indexer/compaction cost. Keep the + // established explicit kernel while history is small, then switch per + // layer once sparse work can remove at least half of the compressed + // rows. This also avoids a transient identity binding at the exact + // top-k boundary, preserving fused-graph replay. + const bool sparse_attention = + ds4_env_flag("DFLASH_DS4_SPARSE_DECODE_FLASH") && ratio > 0 && + padded > 2 * w.n_indexer_top_k; + const DeepSeek4AttentionImpl attention_impl = sparse_attention + ? DeepSeek4AttentionImpl::SparseFlash + : DeepSeek4AttentionImpl::Explicit; ggml_tensor * normed = build_rms_norm(ctx, attn_in, L.attn_norm, w.rms_eps); ggml_tensor * attn_out = build_mla_attention(ctx, gf, normed, w, L, lc, il, lane_kv_start, lane_q, &ain, - i32b, i32ab, i64ab); + i32b, i32ab, i64ab, + &f32ab, attention_impl); if (!attn_out) return false; - if (!i32b.empty() || !i32ab.empty() || !i64ab.empty()) { - std::fprintf(stderr, "[ds4-fused-verify] layer %d dynamic bindings; cannot fuse\n", il); + if (!i32b.empty() || !i32ab.empty() || !i64ab.empty() || + !f32ab.empty()) { + std::fprintf(stderr, + "[ds4-fused-verify] layer %d dynamic bindings " + "(%zu/%zu/%zu/%zu); cannot fuse\n", + il, i32b.size(), i32ab.size(), i64ab.size(), f32ab.size()); return false; } diff --git a/server/src/deepseek4/deepseek4_graph.cpp b/server/src/deepseek4/deepseek4_graph.cpp index 8986328b1..87f7f0d45 100644 --- a/server/src/deepseek4/deepseek4_graph.cpp +++ b/server/src/deepseek4/deepseek4_graph.cpp @@ -1878,11 +1878,14 @@ static ggml_tensor * build_mla_attention( : n_index_comp_live; ggml_tensor * index_visibility_mask = nullptr; if (masked_kv && n_index_comp > 0) { - index_visibility_mask = ggml_view_2d( + // The fused verifier stores one full attention-mask column per + // proposed token. Extract the compressed section for every token, + // then compact the strided view for the indexer scorer. + index_visibility_mask = ggml_cont(ctx, ggml_view_2d( ctx, cached_inputs->attn_row_mask, - n_index_comp, 1, - (size_t) n_index_comp * sizeof(float), - (size_t) w.n_swa * sizeof(float)); + n_index_comp, n_tokens, + cached_inputs->attn_row_mask->nb[1], + (size_t) w.n_swa * sizeof(float))); } indexer_topk = build_indexer_topk( ctx, qr, cur, w, L, index_comp_kv_source, @@ -1954,18 +1957,21 @@ static ggml_tensor * build_mla_attention( } else { kv_attn = raw_kv_view(0, n_raw); } - const bool fused_explicit_f16_kv = w.fused_verify_f16_kv && + const bool fused_verify_f16_kv = w.fused_verify_f16_kv && masked_kv && n_tokens > 1 && - attention_impl == DeepSeek4AttentionImpl::Explicit && kv_attn->type == GGML_TYPE_F32 && raw_kv_source->type == GGML_TYPE_F16 && (!comp_kv_source || comp_kv_source->type == GGML_TYPE_F16) && (!old_rows_scratch_f16 || old_rows_scratch_f16->type == GGML_TYPE_F16); - if (fused_explicit_f16_kv) { + const bool fused_explicit_f16_kv = fused_verify_f16_kv && + attention_impl == DeepSeek4AttentionImpl::Explicit; + const bool fused_sparse_f16_kv = fused_verify_f16_kv && + attention_impl == DeepSeek4AttentionImpl::SparseFlash; + if (fused_explicit_f16_kv || fused_sparse_f16_kv) { // DS4's persistent MLA caches are already F16. Feed those tensors - // directly to the established explicit attention matmuls instead of - // casting the entire long-context cache to F32 on every verifier step. + // directly to the attention implementation instead of casting the + // entire long-context cache to F32 on every verifier step. // Current writes are consumed through their set_rows results, while // preserved overwritten rows retain the same cached F16 values. kv_attn = ggml_view_2d( @@ -1981,12 +1987,15 @@ static ggml_tensor * build_mla_attention( ctx, kv_attn, old_rows_scratch_f16, 1); } static std::atomic explicit_f16_kv_logged{false}; - if (!explicit_f16_kv_logged.exchange(true)) { + static std::atomic sparse_f16_kv_logged{false}; + std::atomic & logged = fused_sparse_f16_kv + ? sparse_f16_kv_logged : explicit_f16_kv_logged; + if (!logged.exchange(true)) { std::fprintf(stderr, - "[deepseek4] fused explicit F16 K/V active: tokens=%d " + "[deepseek4] fused %s F16 K/V active: tokens=%d " "compressed=%d\n", + fused_sparse_f16_kv ? "sparse" : "explicit", n_tokens, n_comp_attn); - explicit_f16_kv_logged = true; } } else { if (n_comp_attn > 0 && comp_kv_source) { @@ -2240,7 +2249,12 @@ static ggml_tensor * build_mla_attention( // The DS4 D=512 kernel consumes Q strides directly, avoiding a full // [D,H,T] -> [D,T,H] materialization for every layer. ggml_tensor * q_fa = ggml_permute(ctx, q, 0, 2, 1, 3); - ggml_tensor * kv_fa = ds4_cast_if_needed(ctx, kv_attn, GGML_TYPE_F32); + // The DS4 D=512 kernel has native F16 K/V specializations. Keep + // fused verifier caches in their persistent representation and + // avoid a full long-context F16 -> F32 conversion every step. + ggml_tensor * kv_fa = fused_sparse_f16_kv + ? kv_attn + : ds4_cast_if_needed(ctx, kv_attn, GGML_TYPE_F32); ggml_tensor * k_fa = ggml_reshape_3d(ctx, kv_fa, head_dim, n_attn, 1); ggml_tensor * v_fa = k_fa; ggml_tensor * mask_fa = score_mask diff --git a/server/test/test_feature_gate.cpp b/server/test/test_feature_gate.cpp index 404d7c4ee..156e2518c 100644 --- a/server/test/test_feature_gate.cpp +++ b/server/test/test_feature_gate.cpp @@ -659,6 +659,7 @@ void test_model_capability_tables() { CHECK(!arch_has_expert_offload("qwen35")); // deepseek4 is mixture-of-experts but has no hot/cold offload path. CHECK(!arch_has_expert_offload("deepseek4")); + CHECK(arch_supports_pflash_compression("deepseek4")); // Every capability predicate must be false for an architecture the // factory cannot build, so no rule can admit an unbuildable model. diff --git a/server/tests/test_deepseek4_unit.cpp b/server/tests/test_deepseek4_unit.cpp index 0d3d8d361..29daf8a78 100644 --- a/server/tests/test_deepseek4_unit.cpp +++ b/server/tests/test_deepseek4_unit.cpp @@ -1726,12 +1726,18 @@ static void test_dspark_park_all_releases_drafter() { backend.spec_draft_path_ = "/tmp/ds4-dspark-fixture.gguf"; backend.spec_drafter_ = std::make_unique(); backend.spec_enabled_ = true; + backend.pflash_drafter_loaded_ = true; + backend.pflash_drafter_path_ = "/tmp/ds4-pflash-fixture.gguf"; + backend.pflash_drafter_gpu_ = 1; TEST_ASSERT(backend.park(ParkTarget::All)); TEST_ASSERT(backend.parked_); TEST_ASSERT(backend.spec_drafter_ == nullptr); TEST_ASSERT(!backend.spec_enabled_); TEST_ASSERT(backend.spec_drafter_parked_); + TEST_ASSERT(!backend.pflash_drafter_loaded_); + TEST_ASSERT(backend.pflash_drafter_path_.empty()); + TEST_ASSERT(backend.pflash_drafter_gpu_ == -1); backend.free_drafter(); TEST_ASSERT(backend.spec_drafter_ == nullptr); @@ -1741,6 +1747,25 @@ static void test_dspark_park_all_releases_drafter() { std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); } +static void test_pflash_rejects_invalid_requests() { + std::fprintf(stderr, " test_pflash_rejects_invalid_requests ..."); + DeepSeek4BackendConfig cfg; + DeepSeek4Backend backend(cfg); + + ModelBackend::CompressRequest empty; + TEST_ASSERT(!backend.compress(empty).ok); + + ModelBackend::CompressRequest invalid_ratio; + invalid_ratio.input_ids = {1, 2, 3}; + invalid_ratio.keep_ratio = 0.0f; + invalid_ratio.drafter_path = "/nonexistent/drafter.gguf"; + const auto results = backend.compress_batch({empty, invalid_ratio}); + TEST_ASSERT(results.size() == 2); + TEST_ASSERT(!results[0].ok); + TEST_ASSERT(!results[1].ok); + std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); +} + static void test_dspark_raw_ring_rollback_after_wrap(ggml_backend_t backend) { std::fprintf(stderr, " test_dspark_raw_ring_rollback_after_wrap ..."); @@ -2831,28 +2856,16 @@ static void test_ds4_flash_attention_parallel_index_scan_gpu() { std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); } -static void test_ds4_indexer_score_packed_q4_gpu() { - std::fprintf(stderr, " test_ds4_indexer_score_packed_q4_gpu ..."); -#if !defined(GGML_USE_HIP) - std::fprintf(stderr, " skipped (HIP-only candidate)\n"); - return; -#endif - ggml_backend_t backend = ggml_backend_cuda_init(0); - if (!backend) { - std::fprintf(stderr, " skipped (no GPU backend)\n"); - return; - } - +static void run_ds4_indexer_score_packed_small_case( + ggml_backend_t backend, int n_tokens) { constexpr int dim = 128; constexpr int n_heads = 64; - constexpr int n_tokens = 4; constexpr int n_comp = 4160; constexpr int kv_start = 16384; constexpr int ratio = 4; ggml_context * ctx = make_test_context(4u << 20); TEST_ASSERT_MSG(ctx != nullptr, "ggml_init failed"); if (!ctx) { - ggml_backend_free(backend); std::fprintf(stderr, " FAIL\n"); return; } @@ -2867,14 +2880,14 @@ static void test_ds4_indexer_score_packed_q4_gpu() { ctx, q, weights, comp, kv_start, ratio); ggml_set_output(scores); TEST_ASSERT_MSG(ggml_backend_supports_op(backend, scores), - "GPU rejected packed-q4 indexer fixture"); + "GPU rejected packed-small indexer fixture"); ggml_cgraph * graph = ggml_new_graph_custom(ctx, 16, false); ggml_build_forward_expand(graph, scores); ggml_gallocr_t alloc = ggml_gallocr_new( ggml_backend_get_default_buffer_type(backend)); const bool allocated = ggml_gallocr_alloc_graph(alloc, graph); - TEST_ASSERT_MSG(allocated, "packed-q4 indexer graph allocation failed"); + TEST_ASSERT_MSG(allocated, "packed-small indexer graph allocation failed"); if (allocated) { std::vector q_data((size_t) dim * n_heads * n_tokens); std::vector weight_data((size_t) n_heads * n_tokens); @@ -2896,37 +2909,35 @@ static void test_ds4_indexer_score_packed_q4_gpu() { ggml_backend_tensor_set(comp, comp_data.data(), 0, comp_data.size() * sizeof(ggml_fp16_t)); - const char * previous = std::getenv("GGML_DS4_INDEXER_PACK_Q4"); - const std::string previous_value = previous ? previous : ""; std::vector reference((size_t) n_comp * n_tokens); std::vector candidate(reference.size()); ScopedCudaGraphOverrides eager( /*disable_graphs=*/true, /*mmvq_max_ncols=*/0, /*skip_property_check=*/false); - unsetenv("GGML_DS4_INDEXER_PACK_Q4"); + setenv("GGML_DS4_INDEXER_PACK_SMALL", "0", 1); TEST_ASSERT_MSG( ggml_backend_graph_compute(backend, graph) == GGML_STATUS_SUCCESS, - "reference q4 indexer score failed"); + "reference small-CM indexer score failed"); ggml_backend_tensor_get(scores, reference.data(), 0, reference.size() * sizeof(float)); - setenv("GGML_DS4_INDEXER_PACK_Q4", "1", 1); + setenv("GGML_DS4_INDEXER_PACK_SMALL", "1", 1); TEST_ASSERT_MSG( ggml_backend_graph_compute(backend, graph) == GGML_STATUS_SUCCESS, - "packed q4 indexer score failed"); + "packed small-CM indexer score failed"); ggml_backend_tensor_get(scores, candidate.data(), 0, candidate.size() * sizeof(float)); TEST_ASSERT_MSG( std::memcmp(reference.data(), candidate.data(), reference.size() * sizeof(float)) == 0, - "packed q4 indexer changed score bits"); + "packed small-CM indexer changed score bits"); auto measure_us = [&](bool packed) { if (packed) { - setenv("GGML_DS4_INDEXER_PACK_Q4", "1", 1); + setenv("GGML_DS4_INDEXER_PACK_SMALL", "1", 1); } else { - unsetenv("GGML_DS4_INDEXER_PACK_Q4"); + setenv("GGML_DS4_INDEXER_PACK_SMALL", "0", 1); } constexpr int warmups = 3; constexpr int iterations = 30; @@ -2945,18 +2956,29 @@ static void test_ds4_indexer_score_packed_q4_gpu() { }; const double reference_us = measure_us(false); const double packed_us = measure_us(true); - std::fprintf(stderr, " reference=%.1fus packed=%.1fus", + std::fprintf(stderr, " q%d=%.1f->%.1fus", n_tokens, reference_us, packed_us); - - if (previous) { - setenv("GGML_DS4_INDEXER_PACK_Q4", previous_value.c_str(), 1); - } else { - unsetenv("GGML_DS4_INDEXER_PACK_Q4"); - } } ggml_gallocr_free(alloc); ggml_free(ctx); +} + +static void test_ds4_indexer_score_packed_small_gpu() { + std::fprintf(stderr, " test_ds4_indexer_score_packed_small_gpu ..."); +#if !defined(GGML_USE_HIP) + std::fprintf(stderr, " skipped (HIP-only candidate)\n"); + return; +#endif + ggml_backend_t backend = ggml_backend_cuda_init(0); + if (!backend) { + std::fprintf(stderr, " skipped (no GPU backend)\n"); + return; + } + ScopedEnvVar packed_small_guard("GGML_DS4_INDEXER_PACK_SMALL"); + for (int n_tokens = 2; n_tokens <= 5; ++n_tokens) { + run_ds4_indexer_score_packed_small_case(backend, n_tokens); + } ggml_backend_free(backend); std::fprintf(stderr, g_failures ? " done\n" : " ok\n"); } @@ -4153,6 +4175,7 @@ int main() { test_safe_compressor_batch_tokens(); test_hybrid_prefill_chunk_tokens(); test_dspark_park_all_releases_drafter(); + test_pflash_rejects_invalid_requests(); test_dspark_raw_ring_rollback_after_wrap(backend); test_snapshot_save_restore(); test_monolithic_snapshot_preserves_decode_state(); @@ -4168,7 +4191,7 @@ int main() { #if defined(GGML_USE_CUDA) || defined(GGML_USE_HIP) test_ds4_flash_attention_keep_cap_gpu(); test_ds4_flash_attention_parallel_index_scan_gpu(); - test_ds4_indexer_score_packed_q4_gpu(); + test_ds4_indexer_score_packed_small_gpu(); test_ds4_topk_block_radix_gpu(); test_ds4_flash_attention_inverse_rope_fallback_gpu(); test_hc_post_strided_split_gpu();