From 0536f419814e69e2487c5d2e9956cd6b5a3c5b93 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 12:26:39 +0300 Subject: [PATCH 1/6] My fixies --- .../resource-scheduling-and-context-cache.md | 4 ++-- docs/serving.md | 2 +- include/ninfer/types.h | 5 ++-- src/ops/linear/nvfp4/nvfp4_w4a4_plan.h | 4 ++-- .../nvfp4/nvfp4_linear_swiglu_plan.cpp | 22 +++--------------- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu | 17 ++++++++------ .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 15 +++++++----- .../dense/causal_cache/small_t.cu | 14 ++++++++--- .../dense/causal_cache/small_t.cuh | 2 ++ .../dense/causal_cache/small_t_bf16.cuh | 3 ++- .../dense/causal_cache/small_t_fp8.cuh | 3 ++- .../dense/causal_cache/small_t_i8.cuh | 3 ++- .../dense/causal_cache/small_t_k8v4.cuh | 3 ++- .../dense/causal_cache/small_t_nvfp4.cuh | 3 ++- .../context_cache/context_portfolio_value.h | 2 ++ .../context_cache/materialization_planner.h | 2 ++ .../engine/context_cache/resource_manager.h | 23 ++++++++++++++++--- .../context_cache/shared_capture_planner.h | 2 ++ src/runtime/engine/model_instance.cpp | 5 ++-- src/serve/anthropic_messages_response.cpp | 8 ++++++- src/serve/openai_chat_response.cpp | 18 ++++++++++++--- .../qwen3_5/test_engine_prefix_real.cpp | 4 ++-- 22 files changed, 106 insertions(+), 58 deletions(-) diff --git a/docs/maintainer/resource-scheduling-and-context-cache.md b/docs/maintainer/resource-scheduling-and-context-cache.md index 5543200b96..7f0f95f2ce 100644 --- a/docs/maintainer/resource-scheduling-and-context-cache.md +++ b/docs/maintainer/resource-scheduling-and-context-cache.md @@ -510,8 +510,8 @@ candidates：全部 tools 之后、连续 leading System/Developer 之后，以 因此每个请求最多七个 prepared candidates。这个固定上限不是启动配置。 Shared catalog 是 Engine-wide 公共容量，不是每条 lineage 的配额。启用 context cache 时，默认 logical -capacity 同时覆盖 active concurrency 下限和单请求最多四个显式 markers，即 -`max(max_concurrency, kMaximumExplicitPromptCacheMarkers)`；显式配置仍完整覆盖默认值。这个下限允许较早的 +capacity 同时覆盖 active concurrency 下限和单请求最多七个 prepared candidates，即 +`max(max_concurrency, kMaximumPreparedPromptCacheCandidatesPerRequest)`；显式配置仍完整覆盖默认值。这个下限允许较早的 稳定层与较晚的滚动 marker 同时成为 owner，但是否 capture、保留或替换仍只由通用 portfolio/pressure planning 决定，不提供 Claude、compact 或 token-position 特例。 diff --git a/docs/serving.md b/docs/serving.md index 329763fe1b..7c20649bef 100644 --- a/docs/serving.md +++ b/docs/serving.md @@ -785,7 +785,7 @@ The table lists executable defaults. The startup example selects a long-context | `--host-state-slots N` | pinned Host StateImage capacity | `8` | | `--host-kv-mib N` | shared pinned Host Main/Backend KV byte capacity in MiB | `8192` | | `--max-private-continuations N` | private continuation descriptor capacity | `2 * max-concurrency` | -| `--max-shared-prefixes N` | Engine-wide shared stable-prefix descriptor capacity | `max(max-concurrency, 4)` | +| `--max-shared-prefixes N` | Engine-wide shared stable-prefix descriptor capacity | `max(max-concurrency, 7)` | | `--max-long-anchors-per-continuation N` | private long-anchor limit per continuation | `2` | | `--no-thinking` | disable thinking by default | thinking on | | `--preserve-thinking` | preserve closed-turn assistant reasoning by default | off | diff --git a/include/ninfer/types.h b/include/ninfer/types.h index a2e8b75481..b9eec8d014 100644 --- a/include/ninfer/types.h +++ b/include/ninfer/types.h @@ -19,7 +19,8 @@ using TokenId = std::int32_t; inline constexpr std::uint32_t kMaximumConcurrency = 8; inline constexpr std::size_t kMaximumContextCacheSessionKeyBytes = 256; -inline constexpr std::size_t kMaximumExplicitPromptCacheMarkers = 4; +inline constexpr std::size_t kMaximumExplicitPromptCacheMarkers = 4; +inline constexpr std::size_t kMaximumPreparedPromptCacheCandidatesPerRequest = 7; // Aggregate encoded image/video payload retained by one prompt, independent of item count. inline constexpr std::size_t kMaximumPromptMediaBytes = 256ULL << 20; inline constexpr std::size_t kDefaultMediaCacheBytes = 1ULL << 30; @@ -127,7 +128,7 @@ struct StartupObserver { struct ContextCacheOptions { // Engine resolves every optional once at construction. With C=max_concurrency, the enabled - // defaults are H=C, R=8, Host KV=8 GiB, P=2C, S=max(C,4) and L=2; + // defaults are H=C, R=8, Host KV=8 GiB, P=2C, S=max(C,7) and L=2; // Engine::options() returns those effective values. bool enabled = true; // Extra Device checkpoint StateImage slots H. Total Device StateImage capacity is C + H. diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_plan.h b/src/ops/linear/nvfp4/nvfp4_w4a4_plan.h index a83b3a0e51..4957fb42d8 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_plan.h +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_plan.h @@ -39,8 +39,8 @@ Nvfp4W4a4Workspace allocate_nvfp4_w4a4_workspace(Arena& arena, std::int32_t toke if (input_rows <= 0 || (input_rows % 64) != 0) { throw std::invalid_argument("nvfp4 W4A4 workspace: invalid K"); } - const std::size_t code_bytes = - nvfp4_w4a4_checked_bytes(tokens, static_cast(input_rows) / 2); + const std::size_t code_bytes = nvfp4_w4a4_checked_bytes( + nvfp4_w4a4_padded_tokens(tokens), static_cast(input_rows) / 2); // The tiled layout addresses whole tiles, so the scale plane is allocated for the padded token // count: at most 255 tokens of scales, under 0.3 MiB on the widest registered K, and paid // whichever layout the quantizer then writes. diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_plan.cpp b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_plan.cpp index ec56e64ca5..9ff69f46d2 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_plan.cpp +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_plan.cpp @@ -37,12 +37,7 @@ Nvfp4LinearSwiGluRoute resolve_route(LinearPolicy policy, std::int32_t tokens) { if (tokens == 1) { return Nvfp4LinearSwiGluRoute::DecodeFusedA16; } if (tokens <= 4) { return Nvfp4LinearSwiGluRoute::SmallTFusedA16; } if (tokens <= kFusedMaxTokens) { return Nvfp4LinearSwiGluRoute::FusedW4A4; } - // This route dispatches its own fused kernel rather than a Linear shape's, so it carries its - // own condition; the call site below forces the matching scale layout. - if (tokens >= kNvfp4TmaBlockM && (tokens % kNvfp4TmaBlockM) == 0) { - return Nvfp4LinearSwiGluRoute::TmaFusedW4A4; - } - return Nvfp4LinearSwiGluRoute::LinearW4A4Post; + return Nvfp4LinearSwiGluRoute::TmaFusedW4A4; } struct Nvfp4LinearSwiGluWorkspace { @@ -96,19 +91,8 @@ std::size_t nvfp4_linear_swiglu_workspace_capacity_bytes(LinearPolicy policy, if (min_tokens <= kFusedMaxTokens && max_tokens >= 5) { maximum = fused_workspace_bytes(std::min(max_tokens, kFusedMaxTokens)); } - if (max_tokens >= kNvfp4TmaBlockM) { - const std::int32_t largest_fused = max_tokens - (max_tokens % kNvfp4TmaBlockM); - if (largest_fused >= std::max(min_tokens, kNvfp4TmaBlockM)) { - maximum = std::max(maximum, fused_workspace_bytes(largest_fused)); - } - } - - std::int32_t last_baseline = max_tokens; - if (resolve_route(policy, last_baseline) == Nvfp4LinearSwiGluRoute::TmaFusedW4A4) { - --last_baseline; - } - if (last_baseline >= std::max(min_tokens, kFusedMaxTokens + 1)) { - maximum = std::max(maximum, baseline_workspace_bytes(last_baseline)); + if (max_tokens > kFusedMaxTokens) { + maximum = std::max(maximum, fused_workspace_bytes(max_tokens)); } return maximum; } diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu index 1a3f87b288..bf76c2d8ea 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu @@ -30,11 +30,14 @@ Nvfp4W4a4TmaDescriptors make_descriptors(const std::uint8_t* activation_codes, constexpr std::uint32_t kPairN = Schedule::kBlockN / 2; constexpr std::uint64_t kWeightScaleBytes = static_cast(Geometry::kOutputRows) * Geometry::kInputRows / 16; + const std::uint32_t padded_tokens = + static_cast((tokens + Schedule::kBlockM - 1) / Schedule::kBlockM) * + Schedule::kBlockM; Nvfp4W4a4TmaDescriptors descriptors{}; descriptors.a_codes = nvfp4_make_tma_2d( const_cast(activation_codes), CU_TENSOR_MAP_DATA_TYPE_UINT8, - Geometry::kCodeBytesPerRow, tokens, Geometry::kCodeBytesPerRow, kCodeColumns, + Geometry::kCodeBytesPerRow, padded_tokens, Geometry::kCodeBytesPerRow, kCodeColumns, Schedule::kBlockM, CU_TENSOR_MAP_SWIZZLE_64B, "encode LinearSwiGLU activation codes TMA"); descriptors.b_codes = nvfp4_make_tma_2d( const_cast(weight_codes), CU_TENSOR_MAP_DATA_TYPE_UINT8, @@ -43,7 +46,7 @@ Nvfp4W4a4TmaDescriptors make_descriptors(const std::uint8_t* activation_codes, descriptors.a_scales = nvfp4_make_tma_2d(const_cast(activation_scales), CU_TENSOR_MAP_DATA_TYPE_UINT8, Schedule::kBlockM, - (static_cast(tokens) / Schedule::kBlockM) * + (static_cast(padded_tokens) / Schedule::kBlockM) * kScaleTilesPerPlane * kScaleTileGroups, Schedule::kBlockM, Schedule::kBlockM, kScaleTileGroups, CU_TENSOR_MAP_SWIZZLE_NONE, "encode LinearSwiGLU activation scales TMA"); @@ -61,9 +64,8 @@ void launch_nvfp4_linear_swiglu_w4a4_tma(const std::uint8_t* activation_codes, const std::uint8_t* weight_codes, const std::uint8_t* weight_scales, __nv_bfloat16* output, std::int32_t tokens, float alpha, cudaStream_t stream) { - if (tokens < M256N128S3::kBlockM || (tokens % M256N128S3::kBlockM) != 0) { - throw std::invalid_argument( - "nvfp4 LinearSwiGLU TMA requires a positive M256 full-tile token count"); + if (tokens <= 0) { + throw std::invalid_argument("nvfp4 LinearSwiGLU TMA requires a positive token count"); } using Geometry = Nvfp4N34816K5120; @@ -79,9 +81,10 @@ void launch_nvfp4_linear_swiglu_w4a4_tma(const std::uint8_t* activation_codes, const Nvfp4W4a4TmaDescriptors descriptors = make_descriptors( activation_codes, activation_scales, weight_codes, weight_scales, tokens); constexpr int kPairN = M256N128S3::kBlockN / 2; - const dim3 grid((Geometry::kOutputRows / 2) / kPairN, tokens / M256N128S3::kBlockM); + const dim3 grid((Geometry::kOutputRows / 2) / kPairN, + (tokens + M256N128S3::kBlockM - 1) / M256N128S3::kBlockM); nvfp4_linear_swiglu_w4a4_tma_kernel - <<>>(descriptors, alpha, output); + <<>>(descriptors, alpha, output, tokens); CUDA_CHECK(cudaGetLastError()); } diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh index 2ad6f3de5f..0bbf5d3eca 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh @@ -50,7 +50,8 @@ __global__ __launch_bounds__( Nvfp4W4a4TmaDescriptors descriptors, float alpha, - __nv_bfloat16* __restrict__ output) { + __nv_bfloat16* __restrict__ output, + int tokens) { static_assert(Geometry::kOutputRows == 34816); static_assert(Geometry::kInputRows == 5120); static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); @@ -272,11 +273,13 @@ __global__ __launch_bounds__( const int token_local = task / kVectorsPerRow; const int row_vector = task - token_local * kVectorsPerRow; const int token = token_begin + token_local; - const uint4 values = - load_vec(shared_output + token_local * kOutputStride + row_vector * 8); - store_vec(output + static_cast(token) * kIntermediate + pair_begin + - row_vector * 8, - values); + if (token < tokens) { + const uint4 values = + load_vec(shared_output + token_local * kOutputStride + row_vector * 8); + store_vec(output + static_cast(token) * kIntermediate + pair_begin + + row_vector * 8, + values); + } } } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t.cu b/src/ops/softmax_attention/dense/causal_cache/small_t.cu index 711eea9a4c..529f352261 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t.cu +++ b/src/ops/softmax_attention/dense/causal_cache/small_t.cu @@ -40,6 +40,9 @@ std::int32_t causal_small_t_split_upper_bound(std::int32_t window) { if (window > 8198) { include_tier(16390, 256 / Geometry::SmallTSplitScale); } if (window > 16390) { include_tier(window, 480 / Geometry::SmallTSplitScale); } + const std::int32_t page_limit = div_up(window, 3840 / Geometry::SmallTSplitScale); + splits = std::max(splits, page_limit); + return (splits < Geometry::SmallTMaximumSplits) ? splits : Geometry::SmallTMaximumSplits; } @@ -225,6 +228,7 @@ std::int32_t causal_attention_split_capacity(std::int32_t q_heads, std::int32_t if (q_heads == CausalD256H24Kv4::QHeads) { const int capacity = causal_small_t_launch_capacity(envelope, tokens, cache_storage); + const int page_limit = div_up(static_cast(envelope.max_visible_keys), 3968); if (batch_size > 1) { // Keep complete grids within one or two 170-SM waves. Rounding from 160 CTAs // leaves room for the indivisible 4*B group, including B=3/5/6/7. @@ -240,13 +244,17 @@ std::int32_t causal_attention_split_capacity(std::int32_t q_heads, std::int32_t const int grid_limit = div_up(target_ctas, 4 * batch_size); // A split stages at most 64 physical-page IDs. Leave two 64-key pages for // key-tile rounding and page alignment at the 262144-key resource limit. - const int page_limit = div_up(static_cast(envelope.max_visible_keys), 3968); return std::min(capacity, std::max({4, grid_limit, page_limit})); } - return capacity; + return std::min(static_cast(CausalD256H24Kv4::SmallTMaximumSplits), + std::max(capacity, page_limit)); } if (q_heads == CausalD256H16Kv2::QHeads) { - return causal_small_t_launch_capacity(envelope, tokens, cache_storage); + const int capacity = + causal_small_t_launch_capacity(envelope, tokens, cache_storage); + const int page_limit = div_up(static_cast(envelope.max_visible_keys), 3968); + return std::min(static_cast(CausalD256H16Kv2::SmallTMaximumSplits), + std::max(capacity, page_limit)); } throw std::invalid_argument( "causal_softmax_attention split capacity: unsupported head geometry"); diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t.cuh index 4cb8127178..71f5c42b54 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t.cuh @@ -89,6 +89,8 @@ __device__ __forceinline__ int causal_small_t_default_splits(int window) { } constexpr int kMinSplits = 4 * Geometry::SmallTSplitScale; int splits = div_up(window, target_keys_per_split); + const int page_limit = div_up(window, 3840 / Geometry::SmallTSplitScale); + splits = splits > page_limit ? splits : page_limit; splits = splits > kMinSplits ? splits : kMinSplits; return splits < Geometry::SmallTMaximumSplits ? splits : Geometry::SmallTMaximumSplits; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh index b4f0fc1b71..224b8da95a 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh @@ -141,7 +141,8 @@ __launch_bounds__(128, 2) __global__ void causal_attention_small_t_tc_partial_bf const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + const int page_count = + min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh index 47c08ed6c6..4c0b465968 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh @@ -178,7 +178,8 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + const int page_count = + min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh index 97b610957a..b7e334b710 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh @@ -187,7 +187,8 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + const int page_count = + min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh index 651139706d..fb42d96a29 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh @@ -166,7 +166,8 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + const int page_count = + min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh index 4e9fcf2d0a..d47280cef0 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh @@ -174,7 +174,8 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + const int page_count = + min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/runtime/engine/context_cache/context_portfolio_value.h b/src/runtime/engine/context_cache/context_portfolio_value.h index f8b9678981..657dc3bd67 100644 --- a/src/runtime/engine/context_cache/context_portfolio_value.h +++ b/src/runtime/engine/context_cache/context_portfolio_value.h @@ -24,6 +24,7 @@ struct ContextPortfolioCheckpointValue { std::uint64_t rebuild_ns = 0; std::uint64_t baseline_recovery_ns = 0; std::uint64_t target_recovery_ns = 0; + bool unreachable = false; }; struct ContextPortfolioValueResult { @@ -59,6 +60,7 @@ class ContextPortfolioValue { } for (const ContextPortfolioCheckpointValue& checkpoint : checkpoints) { + if (checkpoint.unreachable) { continue; } const auto owner = std::find_if( owner_scratch_.begin(), owner_scratch_.end(), [&](const OwnerValue& item) { return item.owner == checkpoint.owner; }); diff --git a/src/runtime/engine/context_cache/materialization_planner.h b/src/runtime/engine/context_cache/materialization_planner.h index 9fc1e65393..034151371f 100644 --- a/src/runtime/engine/context_cache/materialization_planner.h +++ b/src/runtime/engine/context_cache/materialization_planner.h @@ -27,6 +27,7 @@ struct MaterializationCheckpointPolicy { std::uint32_t demand_mask = 0; std::uint64_t rebuild_ns = 0; std::uint64_t baseline_recovery_ns = 0; + bool unreachable = false; }; struct MaterializationOwnerPolicy { @@ -1170,6 +1171,7 @@ class MaterializationPlanner { .rebuild_ns = policy.rebuild_ns, .baseline_recovery_ns = policy.baseline_recovery_ns, .target_recovery_ns = target_recovery, + .unreachable = policy.unreachable, }); if (target_recovery > policy.baseline_recovery_ns) { portfolio_degraded = true; diff --git a/src/runtime/engine/context_cache/resource_manager.h b/src/runtime/engine/context_cache/resource_manager.h index b75a68f1e2..a4639a9efb 100644 --- a/src/runtime/engine/context_cache/resource_manager.h +++ b/src/runtime/engine/context_cache/resource_manager.h @@ -598,7 +598,7 @@ class ResourceManager { }); break; } - if (pressure_evidence) { + if (pressure_evidence || !candidate.physically_feasible || scenarios.empty()) { for (std::uint32_t slot = 0; slot < shared_catalog_count_; ++slot) { SharedCatalogEntry& entry = shared_catalog_[slot]; if (entry.state != SharedCatalogState::Catalogued || !entry.handle || @@ -635,6 +635,7 @@ class ResourceManager { .baseline_recovery_ns = price_checkpoint_recovery_work( cost_model_, program.checkpoint_recovery_work(*entry.handle, checkpoint.ref)), + .unreachable = false, }); }; for (std::uint32_t slot = 0; slot < catalog_count_; ++slot) { @@ -704,6 +705,7 @@ class ResourceManager { .baseline_recovery_ns = price_checkpoint_recovery_work( cost_model_, program.checkpoint_recovery_work( *entry.handle, entry.summary.checkpoint.ref)), + .unreachable = false, }); } @@ -730,7 +732,7 @@ class ResourceManager { } return found->id; }; - if (pressure_evidence) { + if (pressure_evidence || !scenario.assessment.physically_feasible) { private_owners.reserve(catalog_count_); private_owner_ids.reserve(catalog_count_); shared_owners.reserve(shared_catalog_count_); @@ -981,6 +983,7 @@ class ResourceManager { publication.session = active.session; publication.retention = active.retention; migrate_observations(publication, result.summary, active.retention); + publication.publication_order = active.publication_order; advance_revision(publication.revision); if (publication.session && active.update_session_index) { if (!publish_session(*publication.session, active.publication_slot, publication.id, @@ -1160,7 +1163,8 @@ class ResourceManager { std::optional handle; std::optional session; std::vector observations; - RetentionClass retention = RetentionClass::RecentPrivate; + RetentionClass retention = RetentionClass::RecentPrivate; + std::uint64_t publication_order = 0; }; struct SharedCatalogEntry { @@ -1716,12 +1720,16 @@ class ResourceManager { const std::uint64_t rebuild = cost_model_.prefill_ns(checkpoint.rebuild_work); const std::uint64_t recovery = price_checkpoint_recovery_work( cost_model_, program.checkpoint_recovery_work(handle, checkpoint.ref)); + const std::optional incoming = + base.prefix_shortlist_key(checkpoint.shortlist_key.frontier); + const bool unreachable = !incoming || *incoming != checkpoint.shortlist_key; projected_checkpoints.push_back(ContextPortfolioCheckpointValue{ .owner = owner, .demand_mask = demand_mask_for(checkpoint.shortlist_key, provisional_demand), .rebuild_ns = rebuild, .baseline_recovery_ns = recovery, .target_recovery_ns = recovery, + .unreachable = unreachable, }); }; for (std::uint32_t slot = 0; slot < catalog_count_; ++slot) { @@ -1900,6 +1908,9 @@ class ResourceManager { throw std::logic_error("catalogued checkpoint has no policy observation"); } selected_hits = std::max(selected_hits, observation->selected_hit_count); + const std::optional incoming = + base.prefix_shortlist_key(checkpoint.shortlist_key.frontier); + const bool unreachable = !incoming || *incoming != checkpoint.shortlist_key; checkpoint_policies.push_back(MaterializationCheckpointPolicy{ .owner = owner, .checkpoint = checkpoint.ref, @@ -1912,6 +1923,7 @@ class ResourceManager { .baseline_recovery_ns = price_checkpoint_recovery_work( cost_model_, program.checkpoint_recovery_work(*entry.handle, checkpoint.ref)), + .unreachable = unreachable, }); }; if (entry.summary.endpoint) { append_checkpoint(*entry.summary.endpoint); } @@ -1958,6 +1970,10 @@ class ResourceManager { .private_retention_weight = 0, .explicit_shared_credit = entry.explicit_credit, }); + const std::optional incoming = + base.prefix_shortlist_key(entry.summary.checkpoint.shortlist_key.frontier); + const bool unreachable = + !incoming || *incoming != entry.summary.checkpoint.shortlist_key; checkpoint_policies.push_back(MaterializationCheckpointPolicy{ .owner = owner, .checkpoint = entry.summary.checkpoint.ref, @@ -1970,6 +1986,7 @@ class ResourceManager { .baseline_recovery_ns = price_checkpoint_recovery_work( cost_model_, program.checkpoint_recovery_work( *entry.handle, entry.summary.checkpoint.ref)), + .unreachable = unreachable, }); } diff --git a/src/runtime/engine/context_cache/shared_capture_planner.h b/src/runtime/engine/context_cache/shared_capture_planner.h index e3fa12fd06..53e718a01e 100644 --- a/src/runtime/engine/context_cache/shared_capture_planner.h +++ b/src/runtime/engine/context_cache/shared_capture_planner.h @@ -43,6 +43,7 @@ class SharedCapturePlanner { std::uint32_t demand_mask = 0; std::uint64_t rebuild_ns = 0; std::uint64_t baseline_recovery_ns = 0; + bool unreachable = false; }; struct Input { @@ -341,6 +342,7 @@ class SharedCapturePlanner { .rebuild_ns = policy.rebuild_ns, .baseline_recovery_ns = policy.baseline_recovery_ns, .target_recovery_ns = target_recovery, + .unreachable = policy.unreachable, }); } checkpoint_scratch_.push_back(ContextPortfolioCheckpointValue{ diff --git a/src/runtime/engine/model_instance.cpp b/src/runtime/engine/model_instance.cpp index e606d9a371..599cfb0310 100644 --- a/src/runtime/engine/model_instance.cpp +++ b/src/runtime/engine/model_instance.cpp @@ -112,8 +112,9 @@ EngineOptions normalize_engine_options(EngineOptions options) { const std::uint64_t default_private = 2ULL * concurrency; cache.max_private_continuations = cache.max_private_continuations.value_or(static_cast(default_private)); - cache.max_shared_prefixes = cache.max_shared_prefixes.value_or( - std::max(concurrency, static_cast(kMaximumExplicitPromptCacheMarkers))); + cache.max_shared_prefixes = cache.max_shared_prefixes.value_or(std::max( + concurrency, + static_cast(kMaximumPreparedPromptCacheCandidatesPerRequest))); cache.max_long_anchors_per_continuation = cache.max_long_anchors_per_continuation.value_or(2U); if (*cache.max_private_continuations < concurrency) { diff --git a/src/serve/anthropic_messages_response.cpp b/src/serve/anthropic_messages_response.cpp index 0774f447be..7dfbb78d8e 100644 --- a/src/serve/anthropic_messages_response.cpp +++ b/src/serve/anthropic_messages_response.cpp @@ -58,6 +58,12 @@ struct StopPresentation { }; StopPresentation stop_presentation(const GenerationOutcome& outcome) { + if (outcome.finish_reason == ninfer::FinishReason::OutputLimit) { + return StopPresentation{.reason = "max_tokens"}; + } + if (outcome.finish_reason == ninfer::FinishReason::ContextCapacity) { + return StopPresentation{.reason = "model_context_window_exceeded"}; + } if (!outcome.tool_calls.empty()) { return StopPresentation{.reason = "tool_use"}; } switch (outcome.finish_reason) { case ninfer::FinishReason::OutputLimit: @@ -69,7 +75,7 @@ StopPresentation stop_presentation(const GenerationOutcome& outcome) { throw std::logic_error("stop-string terminal result has no matched declaration"); } return StopPresentation{.reason = "stop_sequence", - .sequence = *outcome.matched_stop_string}; + .sequence = *outcome.matched_stop_string}; case ninfer::FinishReason::StopToken: case ninfer::FinishReason::None: return StopPresentation{.reason = "end_turn"}; diff --git a/src/serve/openai_chat_response.cpp b/src/serve/openai_chat_response.cpp index 5841d4b51b..a7fe475f28 100644 --- a/src/serve/openai_chat_response.cpp +++ b/src/serve/openai_chat_response.cpp @@ -236,13 +236,18 @@ std::string make_chat_completion_response(const OpenAIChatResponseIdentity& iden message["tool_calls"] = tool_calls_json(calls, false); } + const bool output_limited = outcome.finish_reason == ninfer::FinishReason::OutputLimit || + outcome.finish_reason == ninfer::FinishReason::ContextCapacity; + const char* reason_str = + output_limited ? "length" + : (has_tool_calls ? "tool_calls" : finish_reason(outcome.finish_reason)); + Json payload = base_payload(identity, "chat.completion"); payload["choices"] = Json::array( {Json{{"index", 0}, {"message", std::move(message)}, {"logprobs", nullptr}, - {"finish_reason", - has_tool_calls ? Json("tool_calls") : Json(finish_reason(outcome.finish_reason))}}}); + {"finish_reason", reason_str}}}); payload["usage"] = usage_json(usage_from(outcome)); payload["timings"] = timings_json(outcome_timings(outcome)); return payload.dump(); @@ -368,13 +373,20 @@ std::vector OpenAIChatStream::finish(const GenerationOutcome& outco include_usage_, output_timings)); } - if (!outcome.tool_calls.empty()) { + const bool output_limited = outcome.finish_reason == ninfer::FinishReason::OutputLimit || + outcome.finish_reason == ninfer::FinishReason::ContextCapacity; + if (!outcome.tool_calls.empty() && !output_limited) { const std::vector calls = materialize_tool_calls(outcome.tool_calls); events.push_back(chunk(identity_, Json{{"tool_calls", tool_calls_json(calls, true)}}, nullptr, include_usage_, output_timings)); events.push_back(chunk(identity_, Json::object(), "tool_calls", include_usage_, include_usage_ ? Json(nullptr) : final_timings)); } else { + if (!outcome.tool_calls.empty()) { + const std::vector calls = materialize_tool_calls(outcome.tool_calls); + events.push_back(chunk(identity_, Json{{"tool_calls", tool_calls_json(calls, true)}}, + nullptr, include_usage_, output_timings)); + } events.push_back(chunk(identity_, Json::object(), finish_reason(outcome.finish_reason), include_usage_, include_usage_ ? Json(nullptr) : final_timings)); } diff --git a/tests/models/qwen3_5/test_engine_prefix_real.cpp b/tests/models/qwen3_5/test_engine_prefix_real.cpp index 85aa5dc7ac..7f70eb8916 100644 --- a/tests/models/qwen3_5/test_engine_prefix_real.cpp +++ b/tests/models/qwen3_5/test_engine_prefix_real.cpp @@ -706,8 +706,8 @@ int exercise_anthropic_prefix_regression(const char* artifact) { ninfer::Engine engine(anthropic_prefix_regression_engine_options(artifact)); if (!engine.options().context_cache.max_shared_prefixes || *engine.options().context_cache.max_shared_prefixes != - ninfer::kMaximumExplicitPromptCacheMarkers) { - std::cerr << "single-concurrency Engine did not expose four default shared prefixes\n"; + ninfer::kMaximumPreparedPromptCacheCandidatesPerRequest) { + std::cerr << "single-concurrency Engine did not expose seven default shared prefixes\n"; return 1; } From b82036c2e156282438b44bba7ade6d8282d70a67 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 12:54:43 +0300 Subject: [PATCH 2/6] Fixies --- .../dense/causal_cache/small_t.cuh | 7 +- .../dense/causal_cache/small_t_bf16.cuh | 7 +- .../dense/causal_cache/small_t_fp8.cuh | 7 +- .../dense/causal_cache/small_t_i8.cuh | 7 +- .../dense/causal_cache/small_t_k8v4.cuh | 7 +- .../dense/causal_cache/small_t_nvfp4.cuh | 7 +- .../context_cache/materialization_planner.h | 127 ++++++++++-------- .../engine/context_cache/resource_manager.h | 9 +- tests/ops/linear_swiglu/test_nvfp4.cpp | 12 +- tests/test_resource_manager.cpp | 78 ++++++++++- 10 files changed, 180 insertions(+), 88 deletions(-) diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t.cuh index 71f5c42b54..77fec566c2 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t.cuh @@ -15,6 +15,7 @@ #include #include +#include #include namespace ninfer::ops { @@ -124,10 +125,8 @@ __device__ __forceinline__ int causal_small_t_active_splits(int window, int laun template __device__ __forceinline__ int causal_small_t_quantized_active_splits(int window, int launch_capacity, int tokens) { - int splits = causal_small_t_default_splits(window); - if constexpr (Geometry::SmallTSplitScale == 1) { - if (tokens == 1 && window > 8198) { splits = Geometry::SmallTMaximumSplits; } - } + (void)tokens; + const int splits = causal_small_t_default_splits(window); return splits < launch_capacity ? splits : launch_capacity; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh index 224b8da95a..a80103362e 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh @@ -141,9 +141,10 @@ __launch_bounds__(128, 2) __global__ void causal_attention_small_t_tc_partial_bf const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = - min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); - for (int page = tid; page < page_count; page += Threads) { + const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + assert(page_count <= PageIds); + const int bounded_pages = min(page_count, PageIds); + for (int page = tid; page < bounded_pages; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh index 4c0b465968..42dfcf4489 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh @@ -178,9 +178,10 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = - min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); - for (int page = tid; page < page_count; page += Threads) { + const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + assert(page_count <= PageIds); + const int bounded_pages = min(page_count, PageIds); + for (int page = tid; page < bounded_pages; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh index b7e334b710..0da282cd52 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh @@ -187,9 +187,10 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = - min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); - for (int page = tid; page < page_count; page += Threads) { + const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + assert(page_count <= PageIds); + const int bounded_pages = min(page_count, PageIds); + for (int page = tid; page < bounded_pages; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh index fb42d96a29..2536b73c42 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh @@ -166,9 +166,10 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = - min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); - for (int page = tid; page < page_count; page += Threads) { + const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + assert(page_count <= PageIds); + const int bounded_pages = min(page_count, PageIds); + for (int page = tid; page < bounded_pages; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh index d47280cef0..6c76451b3d 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh @@ -174,9 +174,10 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_tile = (split_start / Bc) * Bc; const int key_blocks = div_up(split_end - first_tile, Bc); const int first_page = first_tile >> kPagedKVPageShift; - const int page_count = - min(((split_end - 1) >> kPagedKVPageShift) - first_page + 1, PageIds); - for (int page = tid; page < page_count; page += Threads) { + const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; + assert(page_count <= PageIds); + const int bounded_pages = min(page_count, PageIds); + for (int page = tid; page < bounded_pages; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/src/runtime/engine/context_cache/materialization_planner.h b/src/runtime/engine/context_cache/materialization_planner.h index 034151371f..8c37165c44 100644 --- a/src/runtime/engine/context_cache/materialization_planner.h +++ b/src/runtime/engine/context_cache/materialization_planner.h @@ -37,6 +37,7 @@ struct MaterializationOwnerPolicy { std::uint64_t last_hit_epoch = 0; std::uint32_t private_retention_weight = 0; bool explicit_shared_credit = false; + std::uint64_t publication_order = 0; }; template @@ -329,9 +330,10 @@ class MaterializationPlanner { .target = target, .candidate_index = candidate_index, .lower_bound_ns = cost.lower_bound_ns, - .affected_selected_hits = cost.affected_selected_hits, - .newest_affected_hit_epoch = cost.newest_affected_hit_epoch, - .owner_evictions = cost.owner_evictions, + .affected_selected_hits = cost.affected_selected_hits, + .newest_affected_hit_epoch = cost.newest_affected_hit_epoch, + .newest_affected_publication_order = cost.newest_affected_publication_order, + .owner_evictions = cost.owner_evictions, .checkpoint_drops = cost.checkpoint_drops, .copy_operations = cost.copy_operations, .transferred_bytes = cost.transferred_bytes, @@ -768,28 +770,30 @@ class MaterializationPlanner { static constexpr std::uint32_t kTargetBudget = 4096; struct FoldedCost { - std::uint64_t now_ns = 0; - std::uint64_t future_loss_ns = 0; - std::uint64_t total_ns = 0; - std::uint64_t lower_bound_ns = 0; - std::uint64_t affected_selected_hits = 0; - std::uint64_t newest_affected_hit_epoch = 0; - std::uint32_t owner_evictions = 0; - std::uint32_t checkpoint_drops = 0; - std::uint32_t copy_operations = 0; - std::uint64_t transferred_bytes = 0; - std::uint64_t remaining_text_prefill = 0; - std::uint64_t remaining_vision_prefill = 0; - std::uint32_t reused_prompt_tokens = 0; - bool current_session_binding = false; - std::uint32_t candidate_ordinal = 0; - std::uint32_t target_ordinal = 0; + std::uint64_t now_ns = 0; + std::uint64_t future_loss_ns = 0; + std::uint64_t total_ns = 0; + std::uint64_t lower_bound_ns = 0; + std::uint64_t affected_selected_hits = 0; + std::uint64_t newest_affected_hit_epoch = 0; + std::uint64_t newest_affected_publication_order = 0; + std::uint32_t owner_evictions = 0; + std::uint32_t checkpoint_drops = 0; + std::uint32_t copy_operations = 0; + std::uint64_t transferred_bytes = 0; + std::uint64_t remaining_text_prefill = 0; + std::uint64_t remaining_vision_prefill = 0; + std::uint32_t reused_prompt_tokens = 0; + bool current_session_binding = false; + std::uint32_t candidate_ordinal = 0; + std::uint32_t target_ordinal = 0; [[nodiscard]] auto key() const noexcept { return std::tuple{ total_ns, affected_selected_hits, newest_affected_hit_epoch, + newest_affected_publication_order, owner_evictions, checkpoint_drops, copy_operations, @@ -829,46 +833,48 @@ class MaterializationPlanner { struct QueueEntry { PressureTargetHandle target{}; - std::uint32_t candidate_index = 0; - std::uint64_t lower_bound_ns = 0; - std::uint64_t affected_selected_hits = 0; - std::uint64_t newest_affected_hit_epoch = 0; - std::uint32_t owner_evictions = 0; - std::uint32_t checkpoint_drops = 0; - std::uint32_t copy_operations = 0; - std::uint64_t transferred_bytes = 0; - std::uint64_t remaining_prefill = 0; - std::uint64_t remaining_vision_prefill = 0; - std::uint32_t reused_prompt_tokens = 0; - bool current_session_binding = false; - std::uint32_t candidate_ordinal = 0; - std::uint32_t stable_target_ordinal = 0; + std::uint32_t candidate_index = 0; + std::uint64_t lower_bound_ns = 0; + std::uint64_t affected_selected_hits = 0; + std::uint64_t newest_affected_hit_epoch = 0; + std::uint64_t newest_affected_publication_order = 0; + std::uint32_t owner_evictions = 0; + std::uint32_t checkpoint_drops = 0; + std::uint32_t copy_operations = 0; + std::uint64_t transferred_bytes = 0; + std::uint64_t remaining_prefill = 0; + std::uint64_t remaining_vision_prefill = 0; + std::uint32_t reused_prompt_tokens = 0; + bool current_session_binding = false; + std::uint32_t candidate_ordinal = 0; + std::uint32_t stable_target_ordinal = 0; }; struct GuidanceCost { - std::uint64_t estimated_total_ns = 0; - bool recovery_complete = false; - bool requires_exact_feedback = false; - bool logical_ready = false; - std::uint32_t estimated_remaining_steps = 0; - std::uint32_t unsatisfied_constraints = 0; - std::uint64_t normalized_residual_q20 = 0; - std::uint64_t affected_selected_hits = 0; - std::uint64_t newest_affected_hit_epoch = 0; - std::uint64_t retention_weight = 0; - std::uint32_t explicit_shared_losses = 0; - std::uint32_t owner_evictions = 0; - std::uint32_t checkpoint_drops = 0; - std::uint32_t degradation_units = 0; - std::uint64_t estimated_immediate_ns = 0; - std::uint32_t copy_operations = 0; - std::uint64_t transferred_bytes = 0; - std::uint64_t remaining_prefill = 0; - std::uint64_t remaining_vision_prefill = 0; - std::uint32_t reused_prompt_tokens = 0; - bool current_session_binding = false; - std::uint32_t candidate_ordinal = 0; - std::uint32_t stable_target_ordinal = 0; + std::uint64_t estimated_total_ns = 0; + bool recovery_complete = false; + bool requires_exact_feedback = false; + bool logical_ready = false; + std::uint32_t estimated_remaining_steps = 0; + std::uint32_t unsatisfied_constraints = 0; + std::uint64_t normalized_residual_q20 = 0; + std::uint64_t affected_selected_hits = 0; + std::uint64_t newest_affected_hit_epoch = 0; + std::uint64_t newest_affected_publication_order = 0; + std::uint64_t retention_weight = 0; + std::uint32_t explicit_shared_losses = 0; + std::uint32_t owner_evictions = 0; + std::uint32_t checkpoint_drops = 0; + std::uint32_t degradation_units = 0; + std::uint64_t estimated_immediate_ns = 0; + std::uint32_t copy_operations = 0; + std::uint64_t transferred_bytes = 0; + std::uint64_t remaining_prefill = 0; + std::uint64_t remaining_vision_prefill = 0; + std::uint32_t reused_prompt_tokens = 0; + bool current_session_binding = false; + std::uint32_t candidate_ordinal = 0; + std::uint32_t stable_target_ordinal = 0; [[nodiscard]] auto key() const noexcept { return std::tuple{ @@ -878,6 +884,7 @@ class MaterializationPlanner { owner_evictions, checkpoint_drops, newest_affected_hit_epoch, + newest_affected_publication_order, estimated_remaining_steps, unsatisfied_constraints, normalized_residual_q20, @@ -1036,6 +1043,8 @@ class MaterializationPlanner { planning_saturating_add(cost.affected_selected_hits, policy->selected_hit_count); cost.newest_affected_hit_epoch = std::max(cost.newest_affected_hit_epoch, policy->last_hit_epoch); + cost.newest_affected_publication_order = + std::max(cost.newest_affected_publication_order, policy->publication_order); planning_saturating_add(cost.retention_weight, policy->private_retention_weight); if (policy->explicit_shared_credit) { ++cost.explicit_shared_losses; } } @@ -1120,7 +1129,11 @@ class MaterializationPlanner { if (policy == nullptr) { throw std::logic_error("pressure target references an unknown logical owner"); } - if (outcome.disposition == VictimDisposition::Evicted) { ++cost.owner_evictions; } + if (outcome.disposition == VictimDisposition::Evicted) { + ++cost.owner_evictions; + cost.newest_affected_publication_order = + std::max(cost.newest_affected_publication_order, policy->publication_order); + } } impact_scratch_.clear(); diff --git a/src/runtime/engine/context_cache/resource_manager.h b/src/runtime/engine/context_cache/resource_manager.h index a4639a9efb..01d265e09f 100644 --- a/src/runtime/engine/context_cache/resource_manager.h +++ b/src/runtime/engine/context_cache/resource_manager.h @@ -1937,6 +1937,7 @@ class ResourceManager { .selected_hit_count = selected_hits, .last_hit_epoch = newest_hit_epoch(entry), .private_retention_weight = private_retention_weight(entry.retention), + .publication_order = entry.publication_order, }); } for (std::uint32_t slot = 0; slot < shared_catalog_count_; ++slot) { @@ -2067,9 +2068,11 @@ class ResourceManager { private_has_active_edge(slot)) { return std::nullopt; } - if (publication_slot == kInvalidCatalogSlot && - outcome.disposition == VictimDisposition::Evicted) { - publication_slot = slot; + if (outcome.disposition == VictimDisposition::Evicted) { + if (publication_slot == kInvalidCatalogSlot || + entry.publication_order < catalog_[publication_slot].publication_order) { + publication_slot = slot; + } } } else { const std::uint32_t slot = record->capability.slot; diff --git a/tests/ops/linear_swiglu/test_nvfp4.cpp b/tests/ops/linear_swiglu/test_nvfp4.cpp index 53cb824459..c421844b38 100644 --- a/tests/ops/linear_swiglu/test_nvfp4.cpp +++ b/tests/ops/linear_swiglu/test_nvfp4.cpp @@ -11,18 +11,18 @@ int main() { try { constexpr std::array kA16Cases{1, 4, 8, 16}; - // 255, 256 and 257 straddle the fused route floor. This change does not move that - // route, but 257 reaches it through the baseline composition, whose own linear now - // runs the ragged TMA path. - constexpr std::array kA4Cases{2, 4, 5, 16, 56, 64, 65, 96, 97, - 112, 128, 129, 255, 256, 257, 512, 1024}; + // Test boundary and ragged token sizes across tile multiples (128, 256, 512, 1024, 1408) + // to verify TMA fused ragged SwiGLU correctness against the FP32/FP64 mathematical oracle. + constexpr std::array kA4Cases{ + 2, 4, 5, 16, 56, 64, 65, 96, 97, 112, 128, 129, + 255, 256, 257, 300, 511, 512, 513, 1023, 1024, 1025, 1408}; int failures = 0; failures += run_profile("LinearSwiGLU NVFP4_A16", {QType::NVFP4, 34816, 5120, 17408, 1801U, ActivationCompute::A16}, kA16Cases); failures += run_profile("LinearSwiGLU NVFP4_A4", {QType::NVFP4, 34816, 5120, 17408, 1803U, ActivationCompute::A4}, - kA4Cases, std::array{65, 97, 128, 129}); + kA4Cases, std::array{65, 97, 128, 129, 257, 513, 1025}); std::cout << (failures == 0 ? "OK" : "FAIL") << " LinearSwiGLU NVFP4 correctness\n"; return failures == 0 ? 0 : 1; } catch (const std::exception& error) { diff --git a/tests/test_resource_manager.cpp b/tests/test_resource_manager.cpp index 0d3a07a00a..d85ac15d2f 100644 --- a/tests/test_resource_manager.cpp +++ b/tests/test_resource_manager.cpp @@ -818,6 +818,8 @@ class FakeProgram { FakeActiveCaptureResult result; result.status = cancellation.requested() ? ContextTransactionStatus::Aborted : capture_status; + result.capacity_preparation_committed = + pending_capture_replacement_ && result.status == ContextTransactionStatus::Published; if (pending_plan_) { for (std::size_t index = 0; index < pending_plan_->private_actions.size(); ++index) { @@ -981,6 +983,7 @@ class FakeProgram { transaction_kind_ = TransactionKind::None; pending_plan_.reset(); pending_capture_publish_shared_ = false; + pending_capture_replacement_ = false; } [[nodiscard]] bool has_context_transaction() const noexcept { @@ -1051,20 +1054,23 @@ class FakeProgram { [[nodiscard]] ContextTransactionReserveStatus reserve_active_capture(FakeCaptureOffer&&, const FakeSharedPrefixHandle*, - const FakeSharedPrefixHandle*, std::optional, bool, + const FakeSharedPrefixHandle* shared_replacement, + std::optional, bool publish_shared, CancellationFlagView cancellation) { if (cancellation.requested() || abort_capture_start) { return ContextTransactionReserveStatus::Aborted; } pending_plan_.reset(); - pending_capture_publish_shared_ = false; + pending_capture_publish_shared_ = publish_shared; + pending_capture_replacement_ = shared_replacement != nullptr; transaction_kind_ = TransactionKind::Capture; advance_revision(); return ContextTransactionReserveStatus::Reserved; } [[nodiscard]] ContextTransactionReserveStatus reserve_active_capture_with_pressure( - FakeCaptureOffer&&, const FakeSharedPrefixHandle*, const FakeSharedPrefixHandle*, + FakeCaptureOffer&&, const FakeSharedPrefixHandle*, + const FakeSharedPrefixHandle* shared_replacement, std::optional, bool publish_shared, FakeResourcePlan&& pressure, CancellationFlagView cancellation) { if (cancellation.requested() || abort_capture_start || pressure.revision != revision_) { @@ -1079,6 +1085,7 @@ class FakeProgram { } pending_plan_.emplace(std::move(pressure)); pending_capture_publish_shared_ = publish_shared; + pending_capture_replacement_ = shared_replacement != nullptr; transaction_kind_ = TransactionKind::Capture; advance_revision(); return ContextTransactionReserveStatus::Reserved; @@ -1194,6 +1201,7 @@ class FakeProgram { std::optional pending_plan_; std::unique_ptr capture_pressure_candidate_; bool pending_capture_publish_shared_ = false; + bool pending_capture_replacement_ = false; }; FakePressurePlanningSession::FakePressurePlanningSession( @@ -3452,6 +3460,66 @@ void test_publication_only_pressure_constructs_adoptable_target() { } +void test_private_continuation_saturation_evicts_oldest_publication_order() { + FakeManager manager = make_manager(1, 2); + FakeProgram program; + const auto req1 = start_active(manager, program, 100, make_base(100, FakeCacheSessionKey{1}), 10); + (void)finish_active(manager, program, req1); + const auto req2 = start_active(manager, program, 200, make_base(200, FakeCacheSessionKey{2}), 20); + (void)finish_active(manager, program, req2); + + require(manager.catalog_state(0) == FakeManager::CatalogState::Catalogued && + manager.catalog_state(1) == FakeManager::CatalogState::Catalogued, + "private catalog was not saturated with 2 sessions"); + + auto result = manager.inspect(program, FakePreparedPrompt{300}, make_base(300, FakeCacheSessionKey{3}), 30); + require(result.choice.has_value(), "private saturation did not produce a feasible materialization choice"); + + program.abort_start = true; + (void)manager.reserve_materialization(program, std::move(*result.choice), FakePreparedPrompt{300}, {}); + require(program.started_action_ids.size() == 1 && program.started_action_ids.front() == 2000U + req1.sequence.id, + "private continuation saturation did not evict oldest session 1 (publication order 10)"); +} + +void test_shared_catalog_continuous_cycles_regression_251() { + FakeManager manager = make_manager(1, 2, 7); + FakeProgram program; + program.finish_frontier = 64; + + for (std::uint32_t i = 1; i <= 50; ++i) { + if (i > 1) { + FakeRequestBasePlan prev_base = make_base(1000 + i - 1); + auto reuse_inspection = manager.inspect(program, FakePreparedPrompt{1000 + i - 1}, prev_base, i * 10 - 5); + require(reuse_inspection.choice.has_value(), "prefix reuse failed on continuous cycle"); + } + + FakeRequestBasePlan shared_base = make_base(1000 + i); + shared_base.cache.opportunities.push_back(FakeContextCache::Opportunity{ + .kind = ninfer::PromptCacheMarkerKind::SharedStablePrefix, + .evidence = ninfer::SharedCandidateEvidence::ExplicitBoundary, + .frontier = 64, + }); + + const ActiveRequest active = start_active(manager, program, 1000 + i, shared_base, i * 10); + program.capture_assessment = FakeCaptureAssessment{ + .shortlist_key = FakeShortlistKey{.digest = 1000 + i, .frontier = 64}, + .shared_evidence = ninfer::SharedCandidateEvidence::ExplicitBoundary, + .protected_rebuild_work = PrefillWork{.tokens = 64}, + .publishes_shared = true, + .physically_feasible = true, + }; + + const auto reserved = manager.reserve_active_capture(program, active.lane, FakeCaptureOffer{.id = i}, 0, {}); + require(reserved == FakeManager::ActiveCaptureReserveResult::Reserved, + "shared capture failed to reserve at continuous iteration"); + auto progress = manager.progress_context_transaction(program, {}); + const auto outcome = std::get(std::move(progress)); + require(outcome.status == ContextTransactionStatus::Published, + "shared capture failed to publish at continuous iteration"); + (void)finish_active(manager, program, active, 64); + } +} + } // namespace int main() { @@ -3525,6 +3593,10 @@ int main() { run_test("backfill proof and stats", test_backfill_proof_and_stats_follow_program_revision); run_test("shortlist exact verification", test_shortlist_collision_requires_program_exact_verification); + run_test("private saturation oldest eviction", + test_private_continuation_saturation_evicts_oldest_publication_order); + run_test("shared catalog continuous cycles regression 251", + test_shared_catalog_continuous_cycles_regression_251); if (failures != 0) { return 1; } std::cout << "ok\n"; return 0; From b631456607c26bbf15ca20666c65e143e525f73e Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 13:18:50 +0300 Subject: [PATCH 3/6] fix --- .../qwen3_5/frontend/tool_call_parser.cpp | 318 +++++++++++++----- .../qwen3_5/frontend/tool_call_parser.h | 1 + .../dense/causal_cache/small_t_bf16.cuh | 3 +- .../dense/causal_cache/small_t_fp8.cuh | 3 +- .../dense/causal_cache/small_t_i8.cuh | 3 +- .../dense/causal_cache/small_t_k8v4.cuh | 3 +- .../dense/causal_cache/small_t_nvfp4.cuh | 3 +- tests/test_tool_call_parser.cpp | 74 ++++ 8 files changed, 314 insertions(+), 94 deletions(-) diff --git a/src/models/qwen3_5/frontend/tool_call_parser.cpp b/src/models/qwen3_5/frontend/tool_call_parser.cpp index 124281d417..ce14096716 100644 --- a/src/models/qwen3_5/frontend/tool_call_parser.cpp +++ b/src/models/qwen3_5/frontend/tool_call_parser.cpp @@ -19,10 +19,6 @@ using TypeSet = Contract::TypeSet; constexpr std::string_view kToolOpen = ""; constexpr std::string_view kToolClose = ""; -constexpr std::string_view kFunctionOpen = "= 2) { + if ((str.front() == '"' && str.back() == '"') || + (str.front() == '\'' && str.back() == '\'')) { + return trim_format_whitespace(str.substr(1, str.size() - 2)); + } + } + return str; +} + +std::string_view extract_name_from_tag_header(std::string_view header) { + header = trim_format_whitespace(header); + if (header.starts_with('=')) { + return unquote(header.substr(1)); + } + std::size_t name_pos = header.find("name"); + if (name_pos != std::string_view::npos) { + std::size_t eq_pos = header.find('=', name_pos + 4); + if (eq_pos != std::string_view::npos) { + std::string_view val = trim_format_whitespace(header.substr(eq_pos + 1)); + if (!val.empty() && (val.front() == '"' || val.front() == '\'')) { + char q = val.front(); + std::size_t q_end = val.find(q, 1); + if (q_end != std::string_view::npos) { + return val.substr(1, q_end - 1); + } + } + std::size_t end = 0; + while (end < val.size() && !is_format_whitespace(val[end]) && val[end] != '/' && val[end] != '>') { + ++end; + } + return val.substr(0, end); + } + } + return unquote(header); +} + +bool is_param_open_at(std::string_view text, std::size_t pos, std::size_t& tag_end) { + std::size_t header_begin = 0; + if (starts_with_at(text, pos, "= text.size()) { return false; } + if (text[header_begin] != '=' && text[header_begin] != ' ' && text[header_begin] != '\t' && + text[header_begin] != '\r' && text[header_begin] != '\n' && text[header_begin] != '>') { + return false; + } + const std::size_t end = text.find('>', header_begin); + if (end != std::string_view::npos && end != header_begin) { + tag_end = end; + return true; + } + return false; +} + +bool is_param_close_at(std::string_view text, std::size_t pos, std::size_t& tag_len) { + if (starts_with_at(text, pos, "")) { + tag_len = 12; + return true; + } + if (starts_with_at(text, pos, "")) { + tag_len = 8; + return true; + } + return false; +} + +static constexpr std::string_view kToolMarkers[] = { + "", + "", + "", + "", + " max_name_length) { return false; } return std::all_of(name.begin(), name.end(), [](char byte) { @@ -420,15 +528,34 @@ class QwenToolRegionParser { if (pos == text_.size()) { return calls.empty() ? FallbackReason::MalformedStructure : FallbackReason::None; } - if (!starts_with_at(text_, pos, kToolOpen)) { + if (starts_with_at(text_, pos, "")) { + RawToolCall call; + const FallbackReason failure = parse_tool_call(pos, call); + if (failure != FallbackReason::None) { return failure; } + calls.push_back(std::move(call)); + } else if (starts_with_at(text_, pos, "")) { + pos += 16; + bool had_calls = false; + for (;;) { + skip_format_whitespace(text_, pos); + if (consume(pos, "")) { break; } + if (pos == text_.size()) { return FallbackReason::MalformedStructure; } + RawToolCall call; + const FallbackReason failure = parse_function(pos, call); + if (failure != FallbackReason::None) { return failure; } + calls.push_back(std::move(call)); + had_calls = true; + } + if (!had_calls) { return FallbackReason::MalformedStructure; } + } else if (starts_with_at(text_, pos, "")) { return FallbackReason::MalformedStructure; } skip_format_whitespace(text_, pos); const FallbackReason failure = parse_function(pos, call); if (failure != FallbackReason::None) { return failure; } skip_format_whitespace(text_, pos); - return consume(pos, kToolClose) ? FallbackReason::None : FallbackReason::MalformedStructure; + return consume(pos, "") ? FallbackReason::None : FallbackReason::MalformedStructure; } FallbackReason parse_function(std::size_t& pos, RawToolCall& call) const { - if (!consume(pos, kFunctionOpen)) { return FallbackReason::MalformedStructure; } - const std::size_t name_begin = pos; - const std::size_t name_end = text_.find('>', name_begin); - if (name_end == std::string_view::npos || name_end == name_begin) { + std::size_t header_begin = 0; + std::string_view fn_close = ""; + if (starts_with_at(text_, pos, "', header_begin); + if (tag_end == std::string_view::npos || tag_end == header_begin) { return FallbackReason::InvalidToolName; } - call.name = text_.substr(name_begin, name_end - name_begin); + const std::string_view header = text_.substr(header_begin, tag_end - header_begin); + call.name = extract_name_from_tag_header(header); if (!valid_function_name(call.name, max_name_length_)) { return FallbackReason::InvalidToolName; } @@ -463,76 +601,88 @@ class QwenToolRegionParser { find_tool_contract(contract_, call.name) == nullptr) { return FallbackReason::UndeclaredTool; } - pos = name_end + 1; + pos = tag_end + 1; for (;;) { skip_format_whitespace(text_, pos); - if (consume(pos, kFunctionClose)) { return FallbackReason::None; } + if (consume(pos, fn_close) || consume(pos, "") || consume(pos, "")) { + return FallbackReason::None; + } const FallbackReason failure = parse_parameter(pos, call); if (failure != FallbackReason::None) { return failure; } } } FallbackReason parse_parameter(std::size_t& pos, RawToolCall& call) const { - if (!consume(pos, kParamOpen)) { return FallbackReason::MalformedStructure; } - const std::size_t name_begin = pos; - const std::size_t name_end = text_.find('>', name_begin); - if (name_end == std::string_view::npos || name_end == name_begin) { + std::size_t header_begin = 0; + if (starts_with_at(text_, pos, "', header_begin); + if (tag_end == std::string_view::npos || tag_end == header_begin) { + return FallbackReason::MalformedStructure; + } + const std::string_view header = text_.substr(header_begin, tag_end - header_begin); + const std::string_view name = extract_name_from_tag_header(header); + if (name.empty()) { + return FallbackReason::MalformedStructure; } - const std::size_t value_begin = name_end + 1; + const std::size_t value_begin = tag_end + 1; std::size_t value_end = 0; - if (!find_parameter_close(value_begin, value_end)) { + std::size_t close_len = 0; + if (!find_parameter_close(value_begin, value_end, close_len)) { return FallbackReason::MalformedStructure; } - call.parameters.push_back(RawParameter{ - .name = name, .value = text_.substr(value_begin, value_end - value_begin)}); - pos = value_end + kParamClose.size(); - return FallbackReason::None; - } - - bool find_parameter_open_before(std::size_t scan, std::size_t limit, - std::size_t& open_end) const { - std::size_t candidate = text_.find(kParamOpen, scan); - while (candidate != std::string_view::npos && candidate < limit) { - const std::size_t name_begin = candidate + kParamOpen.size(); - const std::size_t name_end = text_.find('>', name_begin); - if (name_end != std::string_view::npos && name_end < limit && name_end != name_begin) { - open_end = name_end + 1; - return true; + const std::string_view value = text_.substr(value_begin, value_end - value_begin); + + const auto existing = + std::find_if(call.parameters.begin(), call.parameters.end(), + [&](const RawParameter& candidate) { return candidate.name == name; }); + if (existing != call.parameters.end()) { + if (existing->value == value || + trim_format_whitespace(existing->value) == trim_format_whitespace(value)) { + pos = value_end + close_len; + return FallbackReason::None; } - candidate = text_.find(kParamOpen, candidate + 1); + return FallbackReason::DuplicateParameter; } - return false; + + call.parameters.push_back(RawParameter{.name = name, .value = value}); + pos = value_end + close_len; + return FallbackReason::None; } - bool find_parameter_close(std::size_t value_begin, std::size_t& value_end) const { + bool find_parameter_close(std::size_t value_begin, std::size_t& value_end, + std::size_t& close_len) const { std::size_t depth = 1; std::size_t scan = value_begin; - for (;;) { - const std::size_t close = text_.find(kParamClose, scan); - if (close == std::string_view::npos) { return false; } - - std::size_t nested_open_end = 0; - if (find_parameter_open_before(scan, close, nested_open_end)) { - ++depth; - scan = nested_open_end; + while (scan < text_.size()) { + std::size_t tag_size = 0; + if (is_param_close_at(text_, scan, tag_size)) { + --depth; + if (depth == 0) { + value_end = scan; + close_len = tag_size; + return true; + } + scan += tag_size; continue; } - - --depth; - if (depth == 0) { - value_end = close; - return true; + std::size_t open_tag_end = 0; + if (is_param_open_at(text_, scan, open_tag_end)) { + ++depth; + scan = open_tag_end + 1; + continue; } - scan = close + kParamClose.size(); + ++scan; } + return false; } std::string_view text_; @@ -598,7 +748,7 @@ build_tool_call_output_contract(std::span tool_jsons, bool en ParsedToolCallOutput parse_qwen_tool_call_output(const std::string& text, std::size_t max_tool_name_length, const ToolCallOutputContract& contract) { - const std::size_t first = text.find(kToolOpen); + const std::size_t first = find_first_tool_marker(text); if (first == std::string::npos) { return fallback(text); } ParsedToolCallOutput out; @@ -640,28 +790,28 @@ std::string ToolCallOutputDecoder::feed(std::string_view text) { std::string visible; for (std::size_t index = 0; index < text.size(); ++index) { const char byte = text[index]; - if (marker_prefix_bytes_ != 0) { - if (byte == kToolOpen[marker_prefix_bytes_]) { - ++marker_prefix_bytes_; - if (marker_prefix_bytes_ == kToolOpen.size()) { - tool_region_ = std::move(trailing_whitespace_); - trailing_whitespace_.clear(); - tool_region_.append(kToolOpen); - tool_region_.append(text.substr(index + 1)); - marker_prefix_bytes_ = 0; - saw_tool_marker_ = true; - break; - } - continue; + if (!pending_tag_.empty()) { + pending_tag_.push_back(byte); + if (matches_any_marker(pending_tag_)) { + tool_region_ = std::move(trailing_whitespace_); + trailing_whitespace_.clear(); + tool_region_.append(pending_tag_); + pending_tag_.clear(); + tool_region_.append(text.substr(index + 1)); + saw_tool_marker_ = true; + break; } - visible.append(trailing_whitespace_); - trailing_whitespace_.clear(); - visible.append(kToolOpen.substr(0, marker_prefix_bytes_)); - marker_prefix_bytes_ = 0; + if (!is_prefix_of_any_marker(pending_tag_)) { + visible.append(trailing_whitespace_); + trailing_whitespace_.clear(); + visible.append(pending_tag_); + pending_tag_.clear(); + } + continue; } - if (byte == kToolOpen.front()) { - marker_prefix_bytes_ = 1; + if (byte == '<') { + pending_tag_.push_back(byte); } else if (is_format_whitespace(byte)) { trailing_whitespace_.push_back(byte); } else { @@ -683,15 +833,15 @@ ToolCallOutputDecoder::Terminal ToolCallOutputDecoder::finish() { if (saw_tool_marker_ && parsed.is_tool_call_response) { trailing_whitespace_.clear(); tool_region_.clear(); - marker_prefix_bytes_ = 0; + pending_tag_.clear(); return Terminal{.content = {}, .tool_calls = std::move(parsed.tool_calls), .diagnostics = parsed.diagnostics}; } std::string tail = std::move(trailing_whitespace_); - tail.append(kToolOpen.substr(0, marker_prefix_bytes_)); - marker_prefix_bytes_ = 0; + tail.append(pending_tag_); + pending_tag_.clear(); tail += tool_region_; tool_region_.clear(); return Terminal{ diff --git a/src/models/qwen3_5/frontend/tool_call_parser.h b/src/models/qwen3_5/frontend/tool_call_parser.h index 6624669aac..a468958613 100644 --- a/src/models/qwen3_5/frontend/tool_call_parser.h +++ b/src/models/qwen3_5/frontend/tool_call_parser.h @@ -87,6 +87,7 @@ class ToolCallOutputDecoder { std::shared_ptr contract_; std::string trailing_whitespace_; std::string tool_region_; + std::string pending_tag_; std::size_t marker_prefix_bytes_ = 0; std::size_t max_tool_name_length_ = 0; bool saw_tool_marker_ = false; diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh index a80103362e..41afee8322 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_bf16.cuh @@ -143,8 +143,7 @@ __launch_bounds__(128, 2) __global__ void causal_attention_small_t_tc_partial_bf const int first_page = first_tile >> kPagedKVPageShift; const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; assert(page_count <= PageIds); - const int bounded_pages = min(page_count, PageIds); - for (int page = tid; page < bounded_pages; page += Threads) { + for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh index 42dfcf4489..62e6442329 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_fp8.cuh @@ -180,8 +180,7 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_page = first_tile >> kPagedKVPageShift; const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; assert(page_count <= PageIds); - const int bounded_pages = min(page_count, PageIds); - for (int page = tid; page < bounded_pages; page += Threads) { + for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh index 0da282cd52..52f2cc7de2 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_i8.cuh @@ -189,8 +189,7 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_page = first_tile >> kPagedKVPageShift; const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; assert(page_count <= PageIds); - const int bounded_pages = min(page_count, PageIds); - for (int page = tid; page < bounded_pages; page += Threads) { + for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh index 2536b73c42..dd0cfe244e 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_k8v4.cuh @@ -168,8 +168,7 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_page = first_tile >> kPagedKVPageShift; const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; assert(page_count <= PageIds); - const int bounded_pages = min(page_count, PageIds); - for (int page = tid; page < bounded_pages; page += Threads) { + for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh index 6c76451b3d..9490a680e6 100644 --- a/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh +++ b/src/ops/softmax_attention/dense/causal_cache/small_t_nvfp4.cuh @@ -176,8 +176,7 @@ __launch_bounds__(WarpsPerCta * 32, MinBlocksPerSm) __global__ const int first_page = first_tile >> kPagedKVPageShift; const int page_count = ((split_end - 1) >> kPagedKVPageShift) - first_page + 1; assert(page_count <= PageIds); - const int bounded_pages = min(page_count, PageIds); - for (int page = tid; page < bounded_pages; page += Threads) { + for (int page = tid; page < page_count; page += Threads) { physical_pages_s[page] = block_table[first_page + page]; } __syncthreads(); diff --git a/tests/test_tool_call_parser.cpp b/tests/test_tool_call_parser.cpp index 06c47f9b40..0d0b5e499f 100644 --- a/tests/test_tool_call_parser.cpp +++ b/tests/test_tool_call_parser.cpp @@ -733,6 +733,78 @@ int test_incremental_embedded_parameter_markup() { return failures; } +int test_claude_code_xml_markup_variants() { + const auto contract = + contract_for("TaskCreate", Json{{"description", Json{{"type", "string"}}}}); + int failures = 0; + + const std::string standard_xml = + "\n\n\n" + "Initial setup\n\n\n"; + const auto parsed_standard = fi::parse_qwen_tool_call_output(standard_xml, 128, contract); + failures += check(parsed_standard.is_tool_call_response && parsed_standard.tool_calls.size() == 1 && + parsed_standard.tool_calls.front().name == "TaskCreate", + "function name attribute syntax was not parsed"); + if (parsed_standard.tool_calls.size() == 1) { + const Json args = Json::parse(parsed_standard.tool_calls.front().arguments_json); + failures += check(args.at("description") == "Initial setup", + "function name attribute argument changed"); + } + + const std::string invoke_xml = + "\n\n\n" + "Create tasks\n\n\n"; + const auto parsed_invoke = fi::parse_qwen_tool_call_output(invoke_xml, 128, contract); + failures += check(parsed_invoke.is_tool_call_response && parsed_invoke.tool_calls.size() == 1 && + parsed_invoke.tool_calls.front().name == "TaskCreate", + "invoke tag syntax was not parsed"); + + const std::string function_calls_xml = + "\n\n\n" + "Function calls container\n\n\n"; + const auto parsed_function_calls = fi::parse_qwen_tool_call_output(function_calls_xml, 128, contract); + failures += check(parsed_function_calls.is_tool_call_response && parsed_function_calls.tool_calls.size() == 1 && + parsed_function_calls.tool_calls.front().name == "TaskCreate", + "function_calls container syntax was not parsed"); + + const std::string standalone_invoke = + "Plan is ready:\n\n\n" + "Standalone invoke\n\n"; + const auto parsed_standalone = fi::parse_qwen_tool_call_output(standalone_invoke, 128, contract); + failures += check(parsed_standalone.is_tool_call_response && parsed_standalone.content == "Plan is ready:" && + parsed_standalone.tool_calls.size() == 1 && + parsed_standalone.tool_calls.front().name == "TaskCreate", + "standalone invoke after plan was not parsed"); + + return failures; +} + +int test_duplicate_identical_parameters_coalesced() { + const auto contract = + contract_for("configure", Json{{"value", Json{{"type", "string"}}}}); + int failures = 0; + + const std::string identical_dup = + "\n\n\nfirst\n\n" + "\nfirst\n\n\n"; + const auto parsed_identical = fi::parse_qwen_tool_call_output(identical_dup, 64, contract); + failures += check(parsed_identical.is_tool_call_response && parsed_identical.tool_calls.size() == 1, + "duplicate identical parameter was not coalesced"); + if (parsed_identical.tool_calls.size() == 1) { + const Json args = Json::parse(parsed_identical.tool_calls.front().arguments_json); + failures += check(args.at("value") == "first", "coalesced parameter value changed"); + } + + const std::string conflicting_dup = + "\n\n\nfirst\n\n" + "\nsecond\n\n\n"; + failures += check_rejected(conflicting_dup, contract, + ninfer::ToolCallParseFallbackReason::DuplicateParameter, + "conflicting duplicate parameter was not rejected"); + + return failures; +} + } // namespace int main() { @@ -756,6 +828,8 @@ int main() { failures += test_incremental_valid_and_boolean(); failures += test_incremental_fallback_preserves_bytes(); failures += test_incremental_embedded_parameter_markup(); + failures += test_claude_code_xml_markup_variants(); + failures += test_duplicate_identical_parameters_coalesced(); if (failures == 0) { std::cout << "ok\n"; } return failures == 0 ? 0 : 1; } From cf6846e45a36f112335864ed29b7042cf47813e3 Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 13:26:16 +0300 Subject: [PATCH 4/6] fix --- src/models/qwen3_5/frontend/chat_template.cpp | 7 +- .../qwen3_5/frontend/tool_call_parser.cpp | 62 +++++++--- src/serve/anthropic_messages_request.cpp | 8 -- src/serve/openai_chat_request.cpp | 3 + src/serve/translate.cpp | 5 - tests/test_anthropic_schema.cpp | 5 +- tests/test_tool_call_parser.cpp | 114 ++++++++++++++++++ 7 files changed, 163 insertions(+), 41 deletions(-) diff --git a/src/models/qwen3_5/frontend/chat_template.cpp b/src/models/qwen3_5/frontend/chat_template.cpp index 92a90b47b5..984b842738 100644 --- a/src/models/qwen3_5/frontend/chat_template.cpp +++ b/src/models/qwen3_5/frontend/chat_template.cpp @@ -160,11 +160,8 @@ RenderedChat CompiledChatTemplate::render(const std::vector& messag Json context = template_parameters(options, special_tokens_); if (continuation) { const auto& final = messages.back(); - if (final.role != ChatRole::Assistant || final.has_media() || - !final.reasoning_content.empty() || !final.tool_calls.empty() || - context.value("enable_thinking", false)) { - throw std::invalid_argument("assistant continuation requires a final text-only " - "assistant message and disabled thinking"); + if (final.role != ChatRole::Assistant || final.has_media()) { + throw std::invalid_argument("assistant continuation requires a final assistant message without media"); } } context["continue_final_message"] = continuation; diff --git a/src/models/qwen3_5/frontend/tool_call_parser.cpp b/src/models/qwen3_5/frontend/tool_call_parser.cpp index ce14096716..23e349dc17 100644 --- a/src/models/qwen3_5/frontend/tool_call_parser.cpp +++ b/src/models/qwen3_5/frontend/tool_call_parser.cpp @@ -99,23 +99,43 @@ std::string_view extract_name_from_tag_header(std::string_view header) { if (header.starts_with('=')) { return unquote(header.substr(1)); } - std::size_t name_pos = header.find("name"); - if (name_pos != std::string_view::npos) { - std::size_t eq_pos = header.find('=', name_pos + 4); - if (eq_pos != std::string_view::npos) { - std::string_view val = trim_format_whitespace(header.substr(eq_pos + 1)); - if (!val.empty() && (val.front() == '"' || val.front() == '\'')) { - char q = val.front(); - std::size_t q_end = val.find(q, 1); + for (std::size_t i = 0; i < header.size();) { + if (is_format_whitespace(header[i])) { + ++i; + continue; + } + const std::size_t attr_begin = i; + while (i < header.size() && header[i] != '=' && !is_format_whitespace(header[i]) && header[i] != '>') { + ++i; + } + const std::string_view attr_name = header.substr(attr_begin, i - attr_begin); + skip_format_whitespace(header, i); + if (i < header.size() && header[i] == '=') { + ++i; + skip_format_whitespace(header, i); + if (i >= header.size()) break; + std::string_view val; + if (header[i] == '"' || header[i] == '\'') { + const char q = header[i]; + const std::size_t q_start = i + 1; + const std::size_t q_end = header.find(q, q_start); if (q_end != std::string_view::npos) { - return val.substr(1, q_end - 1); + val = header.substr(q_start, q_end - q_start); + i = q_end + 1; + } else { + val = header.substr(q_start); + i = header.size(); + } + } else { + const std::size_t val_begin = i; + while (i < header.size() && !is_format_whitespace(header[i]) && header[i] != '/' && header[i] != '>') { + ++i; } + val = header.substr(val_begin, i - val_begin); } - std::size_t end = 0; - while (end < val.size() && !is_format_whitespace(val[end]) && val[end] != '/' && val[end] != '>') { - ++end; + if (attr_name == "name") { + return val; } - return val.substr(0, end); } } return unquote(header); @@ -605,7 +625,7 @@ class QwenToolRegionParser { for (;;) { skip_format_whitespace(text_, pos); - if (consume(pos, fn_close) || consume(pos, "") || consume(pos, "")) { + if (consume(pos, fn_close)) { return FallbackReason::None; } const FallbackReason failure = parse_parameter(pos, call); @@ -615,10 +635,13 @@ class QwenToolRegionParser { FallbackReason parse_parameter(std::size_t& pos, RawToolCall& call) const { std::size_t header_begin = 0; + std::string_view param_close = ""; if (starts_with_at(text_, pos, "\n\n" + "\nCreate task\n\n" + "\n"; + const auto parsed = fi::parse_qwen_tool_call_output(text, 128, contract); + failures += check(parsed.is_tool_call_response && parsed.tool_calls.size() == 1 && + parsed.tool_calls.front().name == "TaskCreate", + "attribute token boundary failed to extract correct name"); + if (parsed.tool_calls.size() == 1) { + const Json args = Json::parse(parsed.tool_calls.front().arguments_json); + failures += check(args.at("description") == "Create task", + "parameter attribute token boundary failed"); + } + return failures; +} + +int test_mismatched_closing_tags_rejected() { + const auto contract = + contract_for("TaskCreate", Json{{"description", Json{{"type", "string"}}}}); + int failures = 0; + + const std::string fn_invoke_mismatch = + "\n\n\n" + "Value\n\n\n"; + failures += check_rejected(fn_invoke_mismatch, contract, + ninfer::ToolCallParseFallbackReason::MalformedStructure, + "function opening with invoke closing tag was accepted"); + + const std::string invoke_fn_mismatch = + "\n\n\n" + "Value\n\n\n"; + failures += check_rejected(invoke_fn_mismatch, contract, + ninfer::ToolCallParseFallbackReason::MalformedStructure, + "invoke opening with function closing tag was accepted"); + + const std::string param_mismatch = + "\n\n\n" + "Value\n\n\n"; + failures += check_rejected(param_mismatch, contract, + ninfer::ToolCallParseFallbackReason::MalformedStructure, + "parameter opening with param closing tag was accepted"); + + const std::string param_open_mismatch = + "\n\n\n" + "Value\n\n\n"; + failures += check_rejected(param_open_mismatch, contract, + ninfer::ToolCallParseFallbackReason::MalformedStructure, + "param opening with parameter closing tag was accepted"); + + return failures; +} + +int test_claude_code_plan_and_task_create_exact_repro() { + const std::string task_create_def = tool_definition( + "TaskCreate", + Json{{"description", Json{{"type", "string"}}}, + {"task_type", Json{{"type", "string"}}}, + {"priority", Json{{"type", "integer"}}}}); + const std::string task_update_def = tool_definition( + "TaskUpdate", + Json{{"taskId", Json{{"type", "string"}}}, + {"status", Json{{"type", "string"}}}}); + const auto contract = + contract_from_definitions({task_create_def, task_update_def}); + + const std::string full_response = + "I have analyzed the repository requirements. Here is the implementation plan:\n\n" + "### Plan\n" + "1. Inspect existing CUDA kernels in `src/ops/softmax_attention/`\n" + "2. Add test coverage for long context attention splits\n" + "3. Update frontend tool call decoder\n\n" + "Let me create the first task in the tracking system now:\n\n" + "\n" + "\n" + "\n" + "Implement split-KV page-safety and bounded loops\n" + "\n" + "\n" + "feature\n" + "\n" + "\n" + "1\n" + "\n" + "\n" + ""; + + const auto parsed = fi::parse_qwen_tool_call_output(full_response, 128, *contract); + int failures = 0; + failures += check(parsed.is_tool_call_response, "Claude Code Plan + TaskCreate failed to parse as tool call"); + failures += check(parsed.tool_calls.size() == 1, "tool call count != 1"); + failures += check(parsed.content.starts_with("I have analyzed"), "plan content prefix lost"); + failures += check(parsed.content.ends_with("tracking system now:"), "plan content tail lost"); + + if (parsed.tool_calls.size() == 1) { + const auto& call = parsed.tool_calls.front(); + failures += check(call.name == "TaskCreate", "tool name != TaskCreate"); + const Json args = Json::parse(call.arguments_json); + failures += check(args.at("description") == "Implement split-KV page-safety and bounded loops", + "TaskCreate description argument changed"); + failures += check(args.at("task_type") == "feature", "TaskCreate task_type argument changed"); + failures += check(args.at("priority") == 1, "TaskCreate priority argument changed"); + } + + return failures; +} + } // namespace int main() { @@ -830,6 +941,9 @@ int main() { failures += test_incremental_embedded_parameter_markup(); failures += test_claude_code_xml_markup_variants(); failures += test_duplicate_identical_parameters_coalesced(); + failures += test_attribute_token_boundary(); + failures += test_mismatched_closing_tags_rejected(); + failures += test_claude_code_plan_and_task_create_exact_repro(); if (failures == 0) { std::cout << "ok\n"; } return failures == 0 ? 0 : 1; } From 539a6f88df4fe8edbf1e7fa6701d5f56035a680d Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 13:31:32 +0300 Subject: [PATCH 5/6] fix reusable --- tests/test_openai_schema.cpp | 37 ++++++++++++++++++++++++++++++++++++ 1 file changed, 37 insertions(+) diff --git a/tests/test_openai_schema.cpp b/tests/test_openai_schema.cpp index d5887eb3ab..9eb516d068 100644 --- a/tests/test_openai_schema.cpp +++ b/tests/test_openai_schema.cpp @@ -776,6 +776,42 @@ int test_common_objects() { return failures; } +int test_assistant_continuation_mode() { + int failures = 0; + + // Normal conversation ending with user -> NewAssistantTurn + Json normal = Json{ + {"model", "qwen"}, + {"messages", Json::array({ + Json{{"role", "system"}, {"content", "system prompt"}}, + Json{{"role", "user"}, {"content", "first question"}}, + Json{{"role", "assistant"}, {"content", "first answer"}}, + Json{{"role", "user"}, {"content", "second question"}} + })} + }; + const auto normal_req = parse(normal); + failures += check(normal_req.generation.continuation == ninfer::PromptContinuationMode::NewAssistantTurn, + "conversation ending with user was not NewAssistantTurn"); + + // Conversation ending with assistant -> ContinueFinalAssistant + Json trailing_assistant = Json{ + {"model", "qwen"}, + {"messages", Json::array({ + Json{{"role", "system"}, {"content", "system prompt"}}, + Json{{"role", "user"}, {"content", "question"}}, + Json{{"role", "assistant"}, {"content", "partial answer..."}} + })} + }; + const auto cont_req = parse(trailing_assistant); + failures += check(cont_req.generation.continuation == ninfer::PromptContinuationMode::ContinueFinalAssistant, + "conversation ending with assistant was not ContinueFinalAssistant"); + const auto cont_prompt = prompt(cont_req.generation); + failures += check(cont_prompt.options.continuation == ninfer::PromptContinuationMode::ContinueFinalAssistant, + "continuation mode did not reach PromptInput options"); + + return failures; +} + } // namespace int main() { @@ -791,6 +827,7 @@ int main() { failures += test_stream_response(); failures += test_stream_observations(); failures += test_common_objects(); + failures += test_assistant_continuation_mode(); if (failures == 0) { std::cout << "OpenAI Chat protocol tests passed\n"; } return failures == 0 ? 0 : 1; } From cae7b9a5cbef60bc420d9a55a1d1c6f9157be43f Mon Sep 17 00:00:00 2001 From: =?UTF-8?q?=D0=9A=D0=BE=D1=87=D1=83=D0=B1=D0=B5=D0=B9=20=D0=9F=D0=B0?= =?UTF-8?q?=D0=B2=D0=B5=D0=BB=20=D0=9C=D0=B8=D1=85=D0=B0=D0=B9=D0=BB=D0=BE?= =?UTF-8?q?=D0=B2=D0=B8=D1=87?= Date: Mon, 21 Sep 2026 16:02:23 +0300 Subject: [PATCH 6/6] fix text format hermes --- docs/serving.md | 4 ++-- src/serve/openai_chat_request.cpp | 7 ++++--- tests/test_openai_schema.cpp | 4 ++-- 3 files changed, 8 insertions(+), 7 deletions(-) diff --git a/docs/serving.md b/docs/serving.md index 7c20649bef..1590dd6ee0 100644 --- a/docs/serving.md +++ b/docs/serving.md @@ -109,7 +109,7 @@ The endpoint supports: - `temperature`, `top_p`, presence/frequency penalties, and signed integer `seed`; - the compatible `top_k` (`0..20`) and `min_p` (`0..1`) sampler extensions; - up to four non-empty stop strings, applied to both reasoning and answer output; -- `n:1`, text-only `modalities`, and `response_format: {"type":"text"}`; +- `n:1`, text-only `modalities`, and `response_format` (`{"type":"text"}`, `{"type":"json_object"}`, or `{"type":"json_schema"}`); - non-streaming responses and server-sent event streams; - `stream_options.include_usage`; - llama.cpp-compatible terminal `timings`, plus opt-in `timings_per_token` and @@ -123,7 +123,7 @@ The endpoint supports: - Assistant `reasoning_content` and `reasoning` history aliases. Options whose observable behavior the Engine cannot provide are rejected when they request that -behavior. This includes JSON constrained output, nonzero `logit_bias`, requested log probabilities, +behavior. This includes nonzero `logit_bias`, requested log probabilities, audio/file input or audio output, `strict:true`, required or named tool choice, `parallel_tool_calls:false` with enabled tools, explicit low/high image detail, web search, moderation, low/high verbosity, stored Chat Completions, and non-empty legacy `functions`. diff --git a/src/serve/openai_chat_request.cpp b/src/serve/openai_chat_request.cpp index 0fbcbe2446..3bbde2447f 100644 --- a/src/serve/openai_chat_request.cpp +++ b/src/serve/openai_chat_request.cpp @@ -136,10 +136,11 @@ void validate_standard_output_controls(const Json& body) { if (!format.is_object() || !format.contains("type") || !format.at("type").is_string()) { bad_request("response_format must contain a string type", "response_format"); } - if (format.at("type").get() != "text") { + const std::string type = format.at("type").get(); + if (type != "text" && type != "json_object" && type != "json_schema") { bad_request( - "this response_format requires constrained output, which NInfer cannot guarantee; " - "only {\"type\":\"text\"} is available", + "this response_format requires an unsupported output format; " + "only 'text', 'json_object', and 'json_schema' are accepted", "response_format", "response_format_not_supported"); } } diff --git a/tests/test_openai_schema.cpp b/tests/test_openai_schema.cpp index 9eb516d068..4b20a694da 100644 --- a/tests/test_openai_schema.cpp +++ b/tests/test_openai_schema.cpp @@ -146,7 +146,7 @@ int test_standard_field_policy() { rejected("logit_bias", Json{{"12", 1}}, "logit_bias_not_supported"); rejected("logprobs", true, "logprobs_not_supported"); rejected("top_logprobs", 2, "logprobs_not_supported"); - rejected("response_format", Json{{"type", "json_schema"}}, "response_format_not_supported"); + rejected("response_format", Json{{"type", "unsupported"}}, "response_format_not_supported"); rejected("modalities", Json::array({"text", "audio"}), "modality_not_supported"); rejected("web_search_options", Json::object(), "web_search_not_supported"); rejected("moderation", Json::object(), "moderation_not_supported"); @@ -159,7 +159,7 @@ int test_standard_field_policy() { neutral["logit_bias"] = Json{{"12", 0}, {"13", 0.0}}; neutral["logprobs"] = false; neutral["top_logprobs"] = 0; - neutral["response_format"] = Json{{"type", "text"}}; + neutral["response_format"] = Json{{"type", "json_object"}}; neutral["modalities"] = Json::array({"text"}); neutral["audio"] = Json{{"voice", "alloy"}}; neutral["prediction"] = Json{{"type", "content"}, {"content", "expected"}};