Repository navigation
[None][feat] Share Kimi K3 FP4 modeling and add NVFP4 support on SM103 - #20532
Closed
VALLIS-NERIA wants to merge 345 commits into
Closed
VALLIS-NERIA wants to merge 345 commits into
VALLIS-NERIA wants to merge 345 commits into
Conversation
… split when it raises As MnnvlWorkspace.create (G1): the workspaces' create() splits an MPI communicator for each call, and on the two failures every rank agrees on (a rank cannot allocate; the allocation failed on a rank) it raised and kept the split. Every rank now frees it on both paths; under Ray the communicator is c10d's TP group and is only dropped. The sandwich, latent exchange and MoE front matrices also assert that every rank freed the communicators its refused creates split. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…nnvl_allreduce_attn_res (G1) Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…n, k3_drafter_attn_qknorm) A CuTe DSL kernel for the drafter's block attention, on SM 100. Each of up to 8 requests brings a draft block of up to 8 tokens. Each block attends densely to its request's paged context and to its own K / V, which are read from qkv and not written to the cache. GQA uses 6 query heads per KV head, head dim 64, HND pages of 64. - trtllm::k3_drafter_attn takes normalized and roped q / k. - trtllm::k3_drafter_attn_qknorm applies the per-head q / k RMSNorm and the NeoX RoPE in the kernel, with fused_qk_norm_rope's arithmetic. The kernel compiles on its first call for its head counts, page stride and whether there is more than one request. That first call must be outside CUDA-graph capture. Lengths, page tables and qkv are read on the device, so captured calls replay with them rewritten. test_k3_drafter_attn.py checks every split the engine schedules, with mixed context lengths: - against an fp64 reference; - against the DFlash TRTLLM path the kernel replaces (append_paged_kv_cache and the trtllm-gen context kernel, non-causal); - no NaN, and the pool left untouched; - reruns bit-identical, and requests isolated; - CUDA-graph replays with rewritten inputs. It runs in l0_b200.yml. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…er_attn_qknorm; sm_100 cells for fused_qk_norm_rope - attention/k3_drafter_attn and attention/k3_drafter_attn_qknorm: contracts, wrappers (argument for argument with the op schemas) and GPU tests on sm_100. The tests run against fp64 references written in the test: - both head layouts Kimi K3 runs (6 / 1, 24 / 4); - eight request x token splits; - context lengths 0 to 2041; - page tables as strided or dense rows, padded page strides; - int32 and int64 positions; - CUDA-graph replay with rewritten inputs; - the out-of-contract calls that must raise. The tests skip off SM 10.0. - attention/fused_qk_norm_rope gains the drafter's cell (head_dim 64, 6 query heads per KV head, 1-64 tokens, eps 1e-5, base 10000, NeoX), which skips off SM 10.0, and an sm_100 receipt for the whole test file. The sm_103 receipt stays, noted as pre-dating that cell. - index.yaml and README record the three sm_100 entries and the rule for reused entries. - The three test files run in l0_b200.yml. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…s ranks only Every Kimi K3 state's create() made its communicator by splitting the session communicator, a collective over every rank of the session: with pipeline or context parallelism every rank of the job had to call create() together. Both create helpers now take it from _get_mnnvl_tp_group_comm, as MnnvlWorkspace.create does: a communicator of the TP group's ranks alone (MPI_Comm_create_group; the TP ProcessGroup under Ray), freed on the agreed failures. The latent exchange matrix checks the helper with the job split into two TP groups, one group making its communicator while the other does not call; the contracts say which ranks call create(). Also: the MoE front's compile cache is keyed by device, and the k3_moe contract states that the op picks its head_flags build from the head tensors it is given, which the entry and K3MoeLayer tie to the state. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…he TP group's ranks only As the workspace's create() now does: every rank of the TP group calls it, ranks outside the group take no part, and on a refusal each rank frees the communicator it made for the call. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…post-attention steps The attention modules expose what a fused post-attention step needs: - will_run_decode_branch(attn_metadata, step), on K3DecodeKDA, K3DecodeMLA and KimiMLARuntime: whether forward runs the step on the decode branch. KDA: every step decode_step classifies, outside a breakable CUDA graph. MLA: a decode step whose paged cache the decode kernels read. MLA's KV writes allow one run per step, so a caller asks before the call. - reduce_output=False, on K3DecodeKDA and KimiMLARuntime: o_proj's TP partial, without the all-reduce. - project_output=False, on K3DecodeKDA, K3DecodeMLA and KimiMLARuntime: the input of o_proj, which is KDA's post-o_norm core or MLA's gated attention output. The flags raise ValueError where the decode branch does not run, except KimiMLARuntime's reduce_output, which holds on every step. The default call path is unchanged. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…dules, copied Targets share no files, so the tp16_moetp16ep1 target ships a copy of the tp16_moetp4ep4 target's modeling.py, weights.py and decode_gemv.py. The copy differs only inside blocks marked "# >>> route B: <why>" ... "# <<< route B": * the module docstrings: experts split 16 x 1, no speculative decoding; * the construction checks: topology (16, 16, 1, 16, 1, no attention DP), the expert split set explicitly (left unset, Kimi K3 runs the experts expert-parallel over the 16 ranks), and no speculative decoding config (the message names the 4 x 4 split for DSpark, DFlash and SA); * the registered name and class; * the speculative path: no SA / DFlash / DSpark admission, no drafter LM head hand-off, no per-token KDA verify states (kda_token_states) and no KDA verify kernels (k3_kda_attn, k3_kda_verify, trtllm::kda_mtp_decode), with REQUIRED_TRTLLM_OPS and UNCERTIFIED_GENERIC_CALLS to match. decode_gemv.py is byte-identical. The causal LM stays the stock one-engine shell the built-in model builds on; without a speculative config it builds no drafter or worker, and the text model's hidden-state taps never fire. test_modeling_v2_kimi_k3_drift.py reads both targets' files, imports neither, and fails on any difference outside the blocks; a control checks that it flags changes outside blocks only. test_modeling_v2_kimi_k3_construction.py runs both targets' construction checks on the host. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…tateful entries and comm/mnnvl_allreduce_attn_res Brings P2 NVIDIA#19828 and P3 NVIDIA#19836 with it. The catalog index and README rows are the union of both sides: 13 entries certified on sm_100. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…its modeling_v2 catalog entries Brings U1 NVIDIA#19816, U3 NVIDIA#19814 and P7 NVIDIA#19827 with it. The catalog index and README rows are the union of both sides. The receipt count is the number of contracts with an sm_100 key in this tree: 22 (19 Kimi K3 entries, cublas_mm, flashinfer_rmsnorm and fused_qk_norm_rope). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… W4A16_MXFP4 KimiK3Config surfaces the checkpoint's compressed-tensors declaration (text_config.quantization_config, mxfp4-pack-quantized), and the model config reads it as W4A16_MXFP4 with no per-layer declarations. The routing criterion and the target's construction assert expected no quantization, so TRTLLM_MODELING_V2=require refused the checkpoint. The routing tests now derive the quantization through KimiK3Config and ModelConfig.load_hf_quant_config from the checkpoint's declaration, and an unquantized checkpoint of the same shape does not match. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…FP4 checkpoint routed as W4A16_MXFP4) The tp16_moetp4ep4 target's text model is trimmed to its settings, and the routing criterion and construction assert read the MXFP4 checkpoint as the model config does, W4A16_MXFP4. routing.py's criterion serves both keys. The routing tests' conflict: the tp16_moetp16ep1 cases build their model configs with the checkpoint's quantization, as the tp16_moetp4ep4 ones now do (_k3_model_config). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… the W4A16_MXFP4 fix into the copy The tp16_moetp4ep4 target's text model is trimmed to its settings (f01a46b), and its construction assert reads the MXFP4 checkpoint as W4A16_MXFP4 (03df8ee). The copy here takes the same changes. Two of them name the layout: the routed-expert split's docstring (16 x 1, a new route B block) and the module docstring's list of asserted settings (inside its block). The construction test's model config carries the checkpoint's quantization, W4A16_MXFP4. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ry path K3DecodeKDA honors reduce_output=False and project_output=False off the decode branch too (unclassified steps and breakable CUDA graphs): the core is the built-in forward's (maybe_bcg_kda_core_inplace inside a breakable graph, else _forward_impl), and the partial is o_proj without the all-reduce. The fused post-attention step runs on steps of at most 16 tokens, classified or not. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… TP-size refusal K3LatentExchange.create now refuses a TP size other than 4, 8 or 16 before any collective step, so a group of one raises ValueError, eagerly and under CUDA-graph capture alike; the test asserts that. The capture refusal at a supported size is a collective agreement, which the 4-rank matrix checks. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…f a router-dump replay The rtb2 cases replayed a recorded router dump named by K3_ROUTING_DUMP and skipped without it, so CI never ran them (256 cases). The skewed cases draw the same kind of routing in the test, seeded: per draw a skewed expert popularity (log-normal), each token's 16 distinct experts a weighted draw without replacement, and the draw's busiest EP group mapped onto this rank, at every M of 1..64. Nothing in the file reads the environment any more. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…the CPU stage test_modeling_v2_kimi_k3_decode_step.py is cpu_only. l0_b300 collects unittest/_torch/modeling_v2 with -m "not cpu_only", and no CPU list named it, so no CI stage ran it. l0_cpu now lists the file. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
The receipt status in index.yaml and the package README still said 12, the count before the Kimi K3 KDA and MLA decode entries were added. 18 contracts have an sm_100 receipt: the 16 Kimi K3 entries (k3_*, attn_res_* and ssm/kda_decode) plus cublas_mm and flashinfer_rmsnorm. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
index.yaml and the package README still said the catalog is fully certified on sm_103. The six Kimi K3 entries this PR adds (k3_kda_attn, k3_kda_decode_attn, k3_kda_verify, kda_decode, k3_mla_qkv, k3_mla_attn_vb_out) carry sm_100 receipts only, and their tests skip on any other architecture. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…s (l0_cpu decode-step entry, sm_100 count) The trim and the W4A16_MXFP4 routing fix came in with u5b. The receipt banner and README keep this tree's count, 22 on sm_100, with the follow-up's wording of the Kimi K3 paragraph (ssm/kda_decode named). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…eceipts The nine entries get their sm_100 receipts: the seven collective ones at 4 ranks on one GB200 tray (each matrix's collected test file in its CI form, pytest starting mpirun, and the same rank bodies started by srun), moe/k3_moe and moe/k3_route_quant with their single-GPU test files' counts. comm/allgather, which the Kimi K3 LM head calls, gets an sm_100 key; its sm_103 receipt stays, taken on the same, unchanged test file. The catalog's counts follow, and the MNNVL all-reduce contract states the evidence for its one-shot kernel's change (main's test bodies, bit-identical outputs, per- call time against the previous kernel). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… their modeling_v2 catalog entries Brings G2 NVIDIA#19833's head through G4's merge f39b747 (already in NVIDIA#19839). The catalog index and README rows are the union of both sides; the receipt count is the number of contracts with an sm_100 key in this tree: 32 (28 Kimi K3 entries, cublas_mm, flashinfer_rmsnorm, allgather and fused_qk_norm_rope). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…xperts of one or two decode tokens Two weight-stream CuTe DSL kernels for decode steps of one or two tokens when every expert is on the rank (moe TP16 x EP1: all 896 experts, a 192-wide slice of each intermediate zero-padded to 256): FC1 (MXFP4 x MXFP8), SiTU, the MXFP8 intermediate, FC2 and the routing-weighted combine in trtllm::k3_moe's order, so the result has k3_moe's bits except where FC1's two partial sums round an intermediate value differently. - trtllm::k3_moe_m1 (one or two tokens) and trtllm::k3_moe_m2 (two tokens, an intermediate of at most 256, per-group FC1 -> FC2 hand-offs) are torch ops over a caller-owned workspace; mutates_args names the intermediate rows, the counts, the epochs, the exchange's multicast words and out. - K3MoeM1State / K3MoeM2State own the workspace: every call re-arms the count set the next call uses and advances the CTAs' epochs. create() compiles the plain build and the push builds it is given, before capture. K3MoeM1Layer / K3MoeM2Layer hold a layer's weights. - The push build stores the partial into every rank's K3LatentExchange for trtllm::k3_latent_reduce instead of returning it. Tests: both kernels against trtllm::k3_moe on the same buffers and an fp32 reference over the dequantized experts (TP16, and TP4 x EP4 for m1 at one token), run to run, layers interleaved on one state, the epochs across the int32 wrap, and the two-token builds' refusals. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
trtllm::k3_moe takes a K3LatentExchange's words (this rank's and the multicast mapping) and flags, and a slot. With them it is the push build, the fused all-reduce's push-only mode: the routed partial goes into that slot of every rank's exchange instead of a returned tensor, for trtllm::k3_latent_reduce, and the call returns an empty tensor. mutates_args gains the exchange's words. K3MoeLayer.push makes the call after trtllm::k3_route_quant or trtllm::k3_moe_front. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…Reduce test_k3_moe_push.py, one process per GPU over the run's TP group (4 on one GB200 tray), every rank with its own TP16 experts and the same tokens: - push: trtllm::k3_moe_m1 at 1 and 2 tokens and trtllm::k3_moe_m2 at 2; - k3_moe_push: trtllm::k3_moe's push build at 1, 3 and 8 tokens, after trtllm::k3_route_quant and after trtllm::k3_moe_front. Per token count and routing set, the pushed partials' reduce equals MNNVLAllReduce's one-shot of the plain partials bit for bit, on every rank and run to run, and leaves the exchange empty with its call count advanced; the same into a 16-slot exchange filled 4 slots per rank (the order of TP16's receive side); and all of it again across the int32 wrap of the call count. Listed in l0_gb200_multi_gpus. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…r caller-owned states Two stateful entries for the routed experts of one or two decode tokens at moe TP16 x EP1, each with its contract (a State section for the caller-owned K3MoeM1State / K3MoeM2State: contents, creator, sharing, call order, what a later call reads, re-arm) and its wrappers (the plain call and the push form into a K3LatentExchange), in index.yaml. Tests (one GPU, l0_b200): single calls against the fp32 reference, and the call sequences the contracts certify: 12 steps of 3 layers on one state, capture and replay with rewritten inputs, two states interleaved, the epochs across the int32 wrap, and create() refusing capture. test_k3_moe_push.py gains the push form's call sequences at 4 ranks, with a negative control (one rank pushing two token sets in the other order). Both entries carry sm_100 receipts with their test counts, and the README's sm_100 count follows (25 entries, 22 of them Kimi K3). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
test_k3_tcgen05_fences.py reads the two weight-stream kernels' sources too: every tcgen05 operation after a wait follows fence::after_thread_sync, and every arrive or barrier after a TMEM load wait follows fence::before_thread_sync. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
k3_moe_push: the M <= 8 partial stored into a slot of every rank's K3LatentExchange for comm/k3_latent_reduce instead of returned, through trtllm::k3_moe's push build. The contract gains the push form's semantics, signature and arguments, the exchange's call order and its preconditions. test_k3_moe_push.py's k3_moe_push check pushes into the 16-slot exchange through the entry (K3MoeLayer.push still fills the run's exchange). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…1 / k3_moe_m2), k3_moe's push build and their catalog entries The receipt count is the number of contracts with an sm_100 key in this tree: 34 (30 Kimi K3 entries, cublas_mm, flashinfer_rmsnorm, allgather and fused_qk_norm_rope). Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…eration request's KDA replay caches On a disaggregated generation server a request's SSM slot is allocated at gen-init, before its conv and SSM states are transferred, and its replay caches are seeded from the transferred conv state only when its first generation step is prepared. A request whose transfer has completed is evictable before that step. Its suspend committed the slot's replay caches, which still held another request's window, over the transferred conv state, and the resume seeded the replay caches from it. The manager now tracks the requests whose slot is allocated at gen-init and whose replay caches are not seeded yet. A suspend of such a request keeps its transferred V2 states as they are; its resume seeds the replay caches from them. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…nto k3-cons/b-runtime Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… k3-cons/c-kernels Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
… k3-cons/d-model Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…he single decode/generic dispatch Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…ng from the managed states Review follow-ups on the V2 manager's KDA draft records: - The constructor flag kda_token_states is now kda_draft_records (with keeps_kda_draft_records, _kda_draft_record_bytes_per_slot and _live_floor_draft_record_slots), so it is not confused with the kda_state_tok tensor that holds the records. The model contract that asks for them follows (`kda_draft_records` on the backbone). - _managed_state_bytes_per_slot gives the bytes of the states the V2 pools manage (SSM, conv and PLE); _mamba_state_bytes_per_slot is that plus the draft records, as before. - _build_cache_config keeps the device quota as total_gpu_quota and names the draft-record reservation and the remaining pool quota separately. - _quota_filling_request_capacity uses ceil_div. No behavior change. Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…s) into k3-cons/c-kernels Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…s) into k3-cons/d-model Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
…untime Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com> # Conflicts: # tensorrt_llm/_torch/_experimental/modeling_v2/README.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/attention/load_paged_kv_cache_for_mla.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/attention/mla_rope_append_paged_kv_assign_q.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/attention/mla_rope_generation.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/attention/thop_attention.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/gemm/cublas_mm.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/index.yaml # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/moe/mxe4m3_mxe2m1_block_scale_moe_runner.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/moe/noaux_tc_op.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/norm/flashinfer_rmsnorm.md # tensorrt_llm/_torch/_experimental/modeling_v2/catalog/quantization/mxfp8_quantize.md # tensorrt_llm/_torch/_experimental/modeling_v2/models/deepseek_v3/r1_0528_nvfp4__sm_103__dep4/ops/attention/kimi_k3_noaux_tc_mxfp8_quant.md # tensorrt_llm/_torch/_experimental/modeling_v2/models/deepseek_v3/r1_0528_nvfp4__sm_103__dep4/ops/attention/kimi_k3_noaux_tc_mxfp8_quant.py # tests/unittest/_torch/modeling_v2/attention/test_modeling_v2_thop_attention.py # tests/unittest/_torch/modeling_v2/models/deepseek_v3/r1_0528_nvfp4__sm_103__dep4/moe/test_modeling_v2_kimi_k3_noaux_tc_mxfp8_quant.py # tests/unittest/_torch/modeling_v2/moe/test_modeling_v2_noaux_tc_op.py # tests/unittest/_torch/modeling_v2/quantization/test_modeling_v2_mxfp8_quantize.py
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Vasanth Sabavat <vsabavat@nvidia.com>
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
Signed-off-by: Xiwen Yu <13230610+VALLIS-NERIA@users.noreply.github.com>
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
This draft extends the Kimi K3
modeling_v2MXFP4 implementation with a shared FP4 modeling layer and an NVFP4 deployment path for SM103/B300.Dependency and review scope: this branch builds on the MXFP4/MNNVL/collective-decode/DSpark work in #19991, which is still open. Because this draft targets
NVIDIA:main, the overall diff includes that inherited foundation. The inherited snapshot isae7c7a9fd7; the NVFP4 and shared-model changes are in46d2b3d958and5c870f584c.Incremental diff from the inherited snapshot (61 files, +4,541 / -9,651).
Test Coverage
Validation below was run from
5c870f584cb05f0fb37c12d74b35d7b1d3ce93cc.FAST_BUILD=OFFand compiler caching disabled.modeling_v2=require, CUDA Graph enabled, normal asynchronous CUDA execution. Actual target class and all 16 SM103 workers were verified.All three accuracy results exceed the repository GSM8K gate derived from its 96.5% reference (threshold approximately 93.2975%). Question IDs, few-shot prompt hashes, and target hashes were checked across all configurations. Each batch size used its own native response cache to resume across short compute allocations. No CUDA errors occurred during these complete batch=1/2/4 inference runs.
Runner outcome: the original full-generation wrapper jobs exited 1 after the official evaluator had saved all outputs. Batch=2/4 hit a driver assertion that counted the two per-filter records as two questions; batch=1 hit NumPy-boolean JSON serialization in the summary writer. The unchanged original per-question records were independently audited on CPU: all 1,319 unique questions appear exactly once per filter, and exact-match counts reproduce the official scores above. The test helper also needed to drain in-flight requests before ending an intermediate allocation window. These corrections were confined to the external validation harness.
Known Issues / Draft Status
max_batch_size=8still encounters an illegal memory access during CUDA Graph warmup. This remains unresolved and blocks readiness for merge.PR Checklist