feat(platform): native Windows (MSVC + CUDA) build for ninfer-serve - #84
Open
devan-carlin wants to merge 9 commits into
Open
feat(platform): native Windows (MSVC + CUDA) build for ninfer-serve#84devan-carlin wants to merge 9 commits into
devan-carlin wants to merge 9 commits into
Conversation
Port the engine to build and run natively on Windows (no WSL2), serving the same .ninfer artifacts with byte-identical output and equal-or-better throughput. Platform code: - artifact/reader.cpp: MappedFile Windows branch (CreateFileW/MapViewOfFile/SetFilePointerEx+ReadFile); fix LARGE_INTEGER aggregate-init truncating file offsets >= 4 GiB (must set .QuadPart) - request_log.cpp: getpid -> GetCurrentProcessId; load_progress.cpp: isatty -> GetConsoleMode; acquire.cpp: Winsock branch - CMake: NINFER_BUILD_MEDIA option (OFF on Windows) + decode/acquire stubs; MSVC C++20 friction fixes NVFP4 TMA fix (Windows-only crash at T>=1024 prefill): MSVC cannot pass the 128-aligned CUtensorMap by value (C2719), so the kernels take a pointer. The TMA unit reads tensor maps through a separate tensormap proxy: kernel-side (generic-proxy) writes to a descriptor staged in local memory are invisible to it without fence.proxy.tensormap, which surfaced as 'Illegal instruction' at the first cp.async.bulk.tensor. Both TMA kernels now read the descriptor directly from the host-written global buffer (cudaMalloc'd, 256-byte-aligned, cudaMemcpyAsync H2D) - no in-kernel staging, no fence. Verified: 1024-token repro + 24k needle (ZEBRA-42-QUARTZ-7719) pass; WSL<->Windows byte-identical parity (seed 42, content + reasoning); compute-sanitizer clean (zero memory errors); prefill 6378 tok/s, decode 172.5 tok/s (WSL baseline 115-125).
![chatgpt-codex-connector[bot]](https://avatars.githubusercontent.com/in/1144995?s=60&v=4){kind=small}
There was a problem hiding this comment.
💡 Codex Review
Here are some automated review suggestions for this pull request.
Reviewed commit: ac146e6fbe
ℹ️ About Codex in GitHub
Codex has been enabled to automatically review pull requests in this repo. Reviews are triggered when you
- Open a pull request for review
- Mark a draft as ready
- Comment "@codex review".
If Codex has suggestions, it will comment; otherwise it will react with 👍.
When you sign up for Codex through ChatGPT, Codex can also answer questions or update the PR, like "@codex address that feedback".
| // multi-stream caller must supply per-stream buffers. | ||
| inline const std::uint64_t* nvfp4_stage_tma_descriptor(const Nvfp4W4a4TmaDescriptors& descriptors, | ||
| cudaStream_t stream) { | ||
| static std::uint64_t* d_descriptor = [] { |
There was a problem hiding this comment.
When two Programs or streams execute an NVFP4 TMA prefill concurrently, both asynchronous copies target this single process-wide buffer, so one launch can overwrite the tensor maps while the other kernel is still consuming them. Because descriptors contain activation, weight, and output addresses, this can cause illegal accesses or silently corrupt logits; make the buffer Program- or stream-owned rather than static.
AGENTS.md reference: AGENTS.md:L124-L133
Useful? React with 👍 / 👎.
| # (NINFER_BUILD_MEDIA=OFF) compiles API-compatible stubs instead and rejects | ||
| # vision requests at runtime. | ||
| if(WIN32) | ||
| set(NINFER_MEDIA_DEFAULT OFF) |
There was a problem hiding this comment.
On the default Windows build this disables media, but serve_options.cpp still accepts --vision and GenerationService treats that flag as proof that media is available. The server therefore starts with Vision allocations enabled, then every media request reaches acquire_stub.cpp, throws an unclassified std::runtime_error, and is returned as HTTP 500. Either retain the supported Vision route on Windows or wire this build capability into startup validation so --vision is rejected before serving.
AGENTS.md reference: AGENTS.md:L107-L113
Useful? React with 👍 / 👎.
Enable vision (media acquire/decode) on the native Windows build: - CMakeLists.txt: on Windows, discover FFMPEG and libcurl via find_path/find_library against third_party roots (BtbN shared DLLs and a source-built libcurl with SCHANNEL TLS) instead of pkg-config, exposing them as PkgConfig::FFMPEG / PkgConfig::LIBCURL imported targets. Non-Windows pkg-config path is unchanged. - src/CMakeLists.txt: link ws2_32 for ninfer_media_acquire on Windows. - artifact/reader.cpp: ReadFile returns 0 bytes at EOF with GetLastError() == 0; aligned read spans can extend past file content (e.g. vision tensors at the end of the artifact). Break on a zero-byte read and let the caller's short-read check decide, instead of throwing a confusing 'direct artifact read: success' error. - media_acquire/acquire.cpp: compare against the wide literal L'..' on Windows (path::native() is wstring) instead of the narrow literal.
igorls
added a commit
to igorls/ninfer
that referenced
this pull request
Aug 25, 2026
… (issue #6, phase 1) Both picks patch-identical to upstream Neroued#84. Native ninfer-serve.exe builds (MSVC 19.44 + CUDA 13.3, sm_120a); Linux container lane verified green by coordinator after merge-gating on the cross-platform changes. WSL2-tax A/B remains per RUNBOOK.md. Claude-Session: https://claude.ai/code/session_01Wv1ehCcaeL86hBzw74iqgr
Integrates 58 upstream dev commits (context-cost engine, FP8/INT8 KV caches, paged KV containers, request-log rewrite, TTFT fixtures, adaptive-graph work) into the Windows-native port. Conflict resolution: - src/targets/qwen3_6/impl/runtime/api_impl.h: dev renamed RequestPlan -> AdmissionPlan. Kept the fork's explicit MSVC-style move ctor/assign/dtor specializations (MSVC mishandles = default on these template specializations), applied to the new name. Windows-compatibility review (held for follow-up, not yet fixed): - src/runtime/contract/types.h: make_prefill_work uses unsigned __int128 (MSVC has no __int128). Core header pulled into the engine/serve build -> will break the MSVC build. - src/runtime/engine/context_cost.cpp: unsigned __int128 (U128) plus <unistd.h>/getpid() -> MSVC blockers. - bench/context_cost/qwen3_6_context_fixture.cpp: __int128 (bench target only, NINFER_BUILD_BENCHMARKS=OFF, not in serve build).
Upstream dev's context-cost engine uses unsigned __int128 and getpid(), neither of which is available on MSVC. This breaks the Windows build of ninfer-serve. - src/runtime/contract/types.h: make_prefill_work now computes prefix*suffix + suffix*(suffix+1)/2 with pure 64-bit saturating arithmetic (each term saturates, then the sum saturates — equivalent to the original min(true_sum, max) clamp). - src/runtime/engine/context_cost.cpp: added a portable 128-bit unsigned multiply (32-bit limb decomposition) used by saturating_product and q32_product_ns; guarded <unistd.h> with a _WIN32 branch and replaced getpid() with GetCurrentProcessId(). Both rewrites were verified against a reference implementation (500k random 64-bit pairs plus boundary cases). Verified: full ninfer-serve build on MSVC 19.44 + CUDA 13.3.
Tool messages now use parse_content_parts() to handle both plain strings and arrays of content parts (text + image_url), matching the behavior of user/assistant messages. This fixes a 400 error when Copilot Chat sends a screenshot as a tool result with image content parts: tool messages must contain string content The OpenAI spec allows content to be either a string or an array of content parts; only the tool role branch was rejecting the array form.
# Conflicts: # src/serve/openai_schema.cpp # src/targets/qwen3_6/impl/runtime/api_impl.h
…server) before upstream merge
This was referenced Aug 31, 2026
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
1 participant
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.
Summary
Port
ninfer-serveto build and run natively on Windows (MSVC + CUDA, no WSL2), serving the same.ninferartifacts with byte-identical output and equal-or-better throughput versus the WSL2 baseline.Changes
Platform code:
artifact/reader.cpp:MappedFileWindows branch (CreateFileW/MapViewOfFile/SetFilePointerEx+ReadFile); fixes aLARGE_INTEGERaggregate-init that truncated file offsets >= 4 GiB (must set.QuadPart)request_log.cpp:getpid->GetCurrentProcessId;load_progress.cpp:isatty->GetConsoleMode;acquire.cpp: Winsock branchNINFER_BUILD_MEDIAoption (OFF on Windows) + decode/acquire stubs; MSVC C++20 friction fixesNVFP4 TMA fix (Windows-only crash at T>=1024 prefill):
MSVC cannot pass the 128-aligned
CUtensorMapby value (C2719), so the kernels take a pointer. The TMA unit reads tensor maps through a separate tensormap proxy: kernel-side (generic-proxy) writes to a descriptor staged in local memory are invisible to it withoutfence.proxy.tensormap, which surfaced asIllegal instructionat the firstcp.async.bulk.tensor. Both TMA kernels now read the descriptor directly from a host-written global buffer (cudaMalloc'd, 256-byte-aligned,cudaMemcpyAsyncH2D) — no in-kernel staging, no fence.Verification
compute-sanitizerclean (zero memory errors)Environment
AI disclosure
Per CONTRIBUTING.md: this implementation was produced with AI assistance (GitHub Copilot, model: 256K Context - llama.cpp - RTX 5090). The contributor reviewed the complete diff and performed the verification above.