From 3f62db060a95c898f55708b9af31baa3e7662dbd Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:22 +0800 Subject: [PATCH 1/9] build: support the MSVC toolchain with vcpkg-managed dependencies - NOMINMAX for every Windows translation unit (windows.h min/max macros corrupt std::min/std::max and the CUDA headers). - /utf-8 source encoding for MSVC C++ and CUDA; the conforming MSVC preprocessor (-Xcompiler=/Zc:preprocessor) is required by CCCL on CUDA 13. - Windows dependencies are declared in vcpkg.json (baseline-locked ffmpeg + curl with Schannel TLS) and resolved with find_package through the vcpkg toolchain file; CMakePresets.json wires a one-command setup (cmake --preset windows-vcpkg). pkg-config lookup stays untouched on Linux. - FFMPEG_LIBRARIES may carry optimized/debug keywords, so they flow into the interface target via target_link_libraries rather than INTERFACE_LINK_LIBRARIES. - utf8proc compiles statically on Windows (UTF8PROC_STATIC). - README: Windows build instructions (vcpkg bootstrap, manifest install, preset build). vcpkg.json baseline: a1cae005c39be7b18ba319fced856b68d7276271 --- .gitignore | 1 + CMakeLists.txt | 52 +++++++++++++++++++++++++++++++++++++++++----- CMakePresets.json | 17 +++++++++++++++ README.md | 25 ++++++++++++++++++++-- src/CMakeLists.txt | 4 ++++ vcpkg.json | 23 ++++++++++++++++++++ 6 files changed, 115 insertions(+), 7 deletions(-) create mode 100644 CMakePresets.json create mode 100644 vcpkg.json diff --git a/.gitignore b/.gitignore index 6350d8b482..6111a1f929 100644 --- a/.gitignore +++ b/.gitignore @@ -2,6 +2,7 @@ build/ build-*/ cmake-build-*/ +vcpkg_installed/ *.o *.obj *.so diff --git a/CMakeLists.txt b/CMakeLists.txt index 98ac0bbd65..9265b035a8 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -28,6 +28,17 @@ if(NOT CMAKE_BUILD_TYPE) set(CMAKE_BUILD_TYPE Release) endif() +if(WIN32) + # windows.h defines min/max macros that corrupt std::min/std::max and the CUDA + # headers; NOMINMAX is required for every translation unit on Windows. + add_compile_definitions(NOMINMAX) + # CCCL (bundled with CUDA 13) hard-requires the conforming MSVC preprocessor, + # and spdlog/fmt requires UTF-8 source encoding. + add_compile_options($<$:-Xcompiler=/Zc:preprocessor>) + add_compile_options($<$:-Xcompiler=/utf-8>) + add_compile_options($<$:/utf-8>) +endif() + if(CMAKE_GENERATOR MATCHES "Ninja") set_property(GLOBAL PROPERTY JOB_POOLS cuda_link=1) set(CMAKE_JOB_POOL_LINK cuda_link) @@ -55,13 +66,44 @@ if(NINFER_BUILD_APPS OR BUILD_TESTING) endif() find_package(CUDAToolkit REQUIRED) -find_package(PkgConfig REQUIRED) -pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET - libavformat>=60 libavcodec>=60 libavutil>=58 libswscale>=7) +find_package(Threads REQUIRED) + +if(WIN32) + # --- FFmpeg (Windows, vcpkg) --------------------------------------------- + # Dependencies are declared in vcpkg.json; the vcpkg toolchain (activated via + # -DCMAKE_TOOLCHAIN_FILE=/scripts/buildsystems/vcpkg.cmake) provides + # the Find module. See README for setup instructions. + find_package(FFMPEG REQUIRED) + message(STATUS "FFmpeg: ${FFMPEG_VERSION}") + if(NOT TARGET PkgConfig::FFMPEG) + add_library(PkgConfig::FFMPEG INTERFACE IMPORTED) + set_target_properties(PkgConfig::FFMPEG PROPERTIES + INTERFACE_INCLUDE_DIRECTORIES "${FFMPEG_INCLUDE_DIRS}") + # FFMPEG_LIBRARIES from select_library_configurations may contain + # optimized/debug keywords; target_link_libraries handles them correctly. + target_link_libraries(PkgConfig::FFMPEG INTERFACE ${FFMPEG_LIBRARIES}) + endif() +else() + find_package(PkgConfig REQUIRED) + pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET + libavformat>=60 libavcodec>=60 libavutil>=58 libswscale>=7) +endif() + if(NINFER_BUILD_MEDIA_ACQUIRE) - pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + if(WIN32) + # --- libcurl (Windows, vcpkg) ------------------------------------------ + # vcpkg installs curl with Schannel (SSPI) TLS; see vcpkg.json. + find_package(CURL REQUIRED) + message(STATUS "libcurl: ${CURL_VERSION_STRING}") + if(NOT TARGET PkgConfig::LIBCURL) + add_library(PkgConfig::LIBCURL INTERFACE IMPORTED) + set_target_properties(PkgConfig::LIBCURL PROPERTIES + INTERFACE_LINK_LIBRARIES CURL::libcurl) + endif() + else() + pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + endif() endif() -find_package(Threads REQUIRED) if(NINFER_BUILD_APPS OR BUILD_TESTING) # Product operational logging is an offline, repository-pinned dependency. Core-only builds do diff --git a/CMakePresets.json b/CMakePresets.json new file mode 100644 index 0000000000..b645c4aa2d --- /dev/null +++ b/CMakePresets.json @@ -0,0 +1,17 @@ +{ + "version": 6, + "configurePresets": [ + { + "name": "windows-vcpkg", + "displayName": "Windows (vcpkg, Ninja)", + "description": "Configures with the vcpkg toolchain for manifest-mode dependency resolution. Set VCPKG_ROOT or vcpkg will be auto-detected.", + "generator": "Ninja", + "binaryDir": "${sourceDir}/build-win", + "cacheVariables": { + "CMAKE_BUILD_TYPE": "Release", + "CMAKE_CUDA_ARCHITECTURES": "120a" + }, + "toolchainFile": "$env{VCPKG_ROOT}/scripts/buildsystems/vcpkg.cmake" + } + ] +} diff --git a/README.md b/README.md index a29ba87670..2f9cdcdcc6 100644 --- a/README.md +++ b/README.md @@ -22,7 +22,7 @@ tokenizer, chat template, and media frontend resources required by its registere ## Quick start -NInfer requires 64-bit Linux, an NVIDIA GeForce RTX 5090, CUDA Toolkit 13.1 or newer, CMake 3.28 or +NInfer requires 64-bit Linux (or Windows), an NVIDIA GeForce RTX 5090, CUDA Toolkit 13.1 or newer, CMake 3.28 or newer, a C++20 host compiler, Ninja, `pkg-config`, FFmpeg development libraries (`libavformat >= 60`, `libavcodec >= 60`, `libavutil >= 58`, and `libswscale >= 7`), and `libcurl >= 7.85`. The build rejects CUDA architectures other than `sm_120a`. @@ -37,6 +37,27 @@ cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release cmake --build build -j ``` +On Windows, dependencies (FFmpeg and libcurl with Schannel TLS) are managed with +[vcpkg](https://vcpkg.io) in manifest mode — no GStreamer runtime or prebuilt curl +packages are needed: + +```powershell +git clone https://github.com/Neroued/ninfer.git +cd ninfer + +# One-time setup (or use an existing vcpkg installation) +git clone https://github.com/microsoft/vcpkg.git +cd vcpkg && .\bootstrap-vcpkg.bat && cd .. +$env:VCPKG_ROOT = "$PWD\vcpkg" + +# Configure through the vcpkg toolchain; dependencies install automatically +cmake -S . -B build-win -G Ninja -DCMAKE_TOOLCHAIN_FILE="$env:VCPKG_ROOT/scripts/buildsystems/vcpkg.cmake" -DCMAKE_CUDA_ARCHITECTURES=120a -DCMAKE_BUILD_TYPE=Release +cmake --build build-win -j +``` + +On first configuration vcpkg compiles FFmpeg and curl from source (about 15 minutes); +subsequent configurations restore them from the local binary cache in seconds. + Tests, benchmarks, and maintainer tools are excluded from the default build. There is no install target or packaged binary distribution; run NInfer from its source build tree. @@ -277,4 +298,4 @@ also uses the fixed packed weights from The Qwen3.8-27B NVFP4 artifact also uses the fixed mixed FP8/NVFP4 weights from [unsloth/Qwen3.8-27B-NVFP4](https://huggingface.co/unsloth/Qwen3.8-27B-NVFP4). These source repositories are distributed under Apache-2.0. Vendored dependencies retain their own license files -under `third_party/`. +under `third_party/`. \ No newline at end of file diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 3e99d8c7e4..df52949e0b 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -314,6 +314,10 @@ add_library(ninfer_text STATIC text/unicode.cpp ${PROJECT_SOURCE_DIR}/third_party/utf8proc/utf8proc.c) ninfer_internal_includes(ninfer_text) +if(WIN32) + # utf8proc.h defaults to __declspec(dllimport); force static linkage. + target_compile_definitions(ninfer_text PRIVATE UTF8PROC_STATIC) +endif() add_library(ninfer_media_decode STATIC media/decode/decode.cpp) diff --git a/vcpkg.json b/vcpkg.json new file mode 100644 index 0000000000..4bc69fb474 --- /dev/null +++ b/vcpkg.json @@ -0,0 +1,23 @@ +{ + "name": "ninfer", + "version-string": "0.1.0", + "dependencies": [ + { + "name": "ffmpeg", + "features": [ + "avcodec", + "avformat", + "swscale" + ] + }, + { + "name": "curl", + "default-features": false, + "features": [ + "ssl", + "non-http" + ] + } + ], + "builtin-baseline": "a1cae005c39be7b18ba319fced856b68d7276271" +} From 163e65e2dc7070ee61b098ce2d6e3b91b3cad417 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:25 +0800 Subject: [PATCH 2/9] feat(apps): embed a UTF-8 code-page manifest and enable UTF-8 consoles - activeCodePage=UTF-8 manifest (merged through target_sources on MSVC) makes the CRT hand the CLI apps UTF-8 argv instead of the legacy ANSI code page. - SetConsoleCP/SetConsoleOutputCP(CP_UTF8) fix console display of UTF-8 output on ANSI-code-page (e.g. 936/GBK) consoles; redirected pipes keep raw UTF-8. - perplexity timestamp: gmtime_s argument order vs POSIX gmtime_r. --- apps/CMakeLists.txt | 17 +++++++++++++++++ apps/cli/main.cpp | 21 +++++++++++++++++++++ apps/perplexity/main.cpp | 23 +++++++++++++++++++++++ apps/serve/main.cpp | 15 +++++++++++++++ apps/windows-utf8.manifest | 9 +++++++++ 5 files changed, 85 insertions(+) create mode 100644 apps/windows-utf8.manifest diff --git a/apps/CMakeLists.txt b/apps/CMakeLists.txt index 04ab34abf7..1c153ab8dc 100644 --- a/apps/CMakeLists.txt +++ b/apps/CMakeLists.txt @@ -1,3 +1,11 @@ +# Embed an activeCodePage=UTF-8 manifest so the MSVC CRT hands apps UTF-8 argv on +# Windows instead of the legacy ANSI code page (breaks non-ASCII --prompt/paths). +# The manifest is appended after each add_executable via target_sources; CMake's +# MSVC rule recognizes .manifest sources and merges them into the default manifest. +if(WIN32 AND MSVC) + set(NINFER_APP_MANIFEST ${CMAKE_CURRENT_SOURCE_DIR}/windows-utf8.manifest) +endif() + add_executable(ninfer cli/main.cpp cli/options.cpp) @@ -9,6 +17,9 @@ target_link_libraries(ninfer PRIVATE ninfer_engine ninfer_product_logging ninfer_product_prompt_input) +if(WIN32 AND MSVC) + target_sources(ninfer PRIVATE ${NINFER_APP_MANIFEST}) +endif() add_executable(ninfer-serve serve/main.cpp) target_include_directories(ninfer-serve PRIVATE @@ -17,6 +28,9 @@ target_include_directories(ninfer-serve PRIVATE ${PROJECT_SOURCE_DIR}/third_party ${PROJECT_SOURCE_DIR}/third_party/cpp-httplib) target_link_libraries(ninfer-serve PRIVATE ninfer_serve ninfer_product_logging) +if(WIN32 AND MSVC) + target_sources(ninfer-serve PRIVATE ${NINFER_APP_MANIFEST}) +endif() add_executable(ninfer-perplexity perplexity/main.cpp @@ -28,3 +42,6 @@ target_include_directories(ninfer-perplexity PRIVATE ${PROJECT_SOURCE_DIR}/third_party ${CMAKE_CURRENT_SOURCE_DIR}/perplexity) target_link_libraries(ninfer-perplexity PRIVATE ninfer_engine ninfer_product_logging) +if(WIN32 AND MSVC) + target_sources(ninfer-perplexity PRIVATE ${NINFER_APP_MANIFEST}) +endif() diff --git a/apps/cli/main.cpp b/apps/cli/main.cpp index c074b6f5f5..c294f03565 100644 --- a/apps/cli/main.cpp +++ b/apps/cli/main.cpp @@ -17,6 +17,13 @@ #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page (e.g. +// 936/GBK). The activeCodePage=UTF-8 manifest fixes argv input; this fixes display +// of UTF-8 output on code-page-936 consoles. Redirected pipes keep raw UTF-8 bytes. +#include +#endif + namespace { std::string format_seconds(double seconds) { @@ -228,6 +235,20 @@ void print_generation_summary(const ninfer::GenerationResult& result, } // namespace +#ifdef _WIN32 +namespace { +// Set once before any output; harmless when stdout is redirected (it only affects +// how the attached console decodes the byte stream). +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + int main(int argc, char** argv) { ninfer::cli::Options cli; try { diff --git a/apps/perplexity/main.cpp b/apps/perplexity/main.cpp index 2a2b70ece7..1555a0e67f 100644 --- a/apps/perplexity/main.cpp +++ b/apps/perplexity/main.cpp @@ -9,6 +9,12 @@ #include #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page. The +// activeCodePage=UTF-8 manifest fixes argv input; this fixes console display. +#include +#endif + #include #include #include @@ -163,7 +169,12 @@ std::string safe_component(std::string_view value) { std::string timestamp() { const std::time_t now = std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); std::tm utc{}; +#ifdef _WIN32 + // gmtime_s swaps argument order relative to POSIX gmtime_r. + gmtime_s(&utc, &now); +#else gmtime_r(&now, &utc); +#endif std::ostringstream out; out << std::put_time(&utc, "%Y%m%d-%H%M%S"); return out.str(); @@ -433,6 +444,18 @@ int run(const Options& options, const std::shared_ptr& logger, } // namespace +#ifdef _WIN32 +namespace { +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + int main(int argc, char** argv) { Options options; try { diff --git a/apps/serve/main.cpp b/apps/serve/main.cpp index 02784db556..876a8936b1 100644 --- a/apps/serve/main.cpp +++ b/apps/serve/main.cpp @@ -16,6 +16,21 @@ #include #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page. The +// activeCodePage=UTF-8 manifest fixes argv input; this fixes console display. +#include +namespace { +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + namespace { std::atomic g_server{nullptr}; diff --git a/apps/windows-utf8.manifest b/apps/windows-utf8.manifest new file mode 100644 index 0000000000..88938f809d --- /dev/null +++ b/apps/windows-utf8.manifest @@ -0,0 +1,9 @@ + + + + + + UTF-8 + + + From c114d6a651ad091c5f0051e49d15d96d601f575c Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:30 +0800 Subject: [PATCH 3/9] fix(media_acquire): use the Windows socket stack and wide-char path checks - winsock2/ws2tcpip replace arpa/inet.h, netdb.h, and sys/socket.h under _WIN32; the code shares the existing getaddrinfo/inet_ntop flow. - The media-root escape check compares native wide-char paths on Windows (L"..") so the traversal guard keeps working with the same semantics. - Link ws2_32 into ninfer_media_acquire on Windows. --- src/CMakeLists.txt | 3 +++ src/product/media_acquire/acquire.cpp | 10 +++++++++- 2 files changed, 12 insertions(+), 1 deletion(-) diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index df52949e0b..321e559d58 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -330,6 +330,9 @@ if(NINFER_BUILD_MEDIA_ACQUIRE) product/media_acquire/acquire.cpp) ninfer_internal_includes(ninfer_media_acquire) target_link_libraries(ninfer_media_acquire PRIVATE PkgConfig::LIBCURL) + if(WIN32) + target_link_libraries(ninfer_media_acquire PRIVATE ws2_32) + endif() endif() if(NINFER_BUILD_PROMPT_INPUT) diff --git a/src/product/media_acquire/acquire.cpp b/src/product/media_acquire/acquire.cpp index 24644492e2..bb1f8c8c49 100644 --- a/src/product/media_acquire/acquire.cpp +++ b/src/product/media_acquire/acquire.cpp @@ -2,9 +2,14 @@ #include +#if defined(_WIN32) +#include +#include +#else #include #include #include +#endif #include #include @@ -276,7 +281,6 @@ std::vector fetch_url(std::string url, const Policy& policy) { } throw Error(ErrorKind::RemoteUnavailable, "too many media URL redirects"); } - std::vector read_path(const Source& source, const Policy& policy) { check_control(policy); std::error_code ec; @@ -287,7 +291,11 @@ std::vector read_path(const Source& source, const Policy& policy) if (!policy.media_root.empty()) { const std::filesystem::path root = std::filesystem::weakly_canonical(policy.media_root, ec); const auto relative = std::filesystem::relative(path, root, ec); +#if defined(_WIN32) + if (ec || relative.empty() || relative.native().starts_with(L"..")) { +#else if (ec || relative.empty() || relative.native().starts_with("..")) { +#endif throw std::invalid_argument("media path is outside configured media root"); } } From 6c4e29ec57a18e9f6c90902d02e7259e42c76be2 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:33 +0800 Subject: [PATCH 4/9] feat(artifact): add a Win32 mapped-file and overlapped direct-read path CreateFileW/CreateFileMappingW/MapViewOfFile replace POSIX O_DIRECT mmap, and OVERLAPPED ReadFile performs the aligned direct reads. Bounds checks mirror the POSIX branch (LONGLONG offsets, DWORD transfer sizes). --- src/artifact/reader.cpp | 80 ++++++++++++++++++++++++++++++++++++++++- 1 file changed, 79 insertions(+), 1 deletion(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 56bce38c56..57de9ba123 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -15,10 +15,14 @@ #include #include +#if defined(_WIN32) +#include +#else #include #include #include #include +#endif namespace ninfer::artifact { namespace { @@ -177,6 +181,52 @@ struct TransparentStringHash { class MappedFile { public: explicit MappedFile(const std::filesystem::path& path) { +#if defined(_WIN32) + const HANDLE file = ::CreateFileW(path.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, + OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, nullptr); + if (file == INVALID_HANDLE_VALUE) { + throw std::system_error(static_cast(::GetLastError()), std::system_category(), + "open " + path.string()); + } + + LARGE_INTEGER file_size{}; + if (::GetFileSizeEx(file, &file_size) == 0) { + const int error = static_cast(::GetLastError()); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), + "fstat " + path.string()); + } + if (file_size.QuadPart < 0 || + static_cast(file_size.QuadPart) > + std::numeric_limits::max()) { + ::CloseHandle(file); + throw ArtifactError("artifact size does not fit the process address space"); + } + + const auto size = static_cast(file_size.QuadPart); + void* mapping = nullptr; + if (size != 0) { + const HANDLE mapping_handle = + ::CreateFileMappingW(file, nullptr, PAGE_READONLY, 0, 0, nullptr); + if (mapping_handle == nullptr) { + const int error = static_cast(::GetLastError()); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), + "mmap " + path.string()); + } + mapping = ::MapViewOfFile(mapping_handle, FILE_MAP_READ, 0, 0, 0); + ::CloseHandle(mapping_handle); // The view keeps the file object alive. + if (mapping == nullptr) { + const int error = static_cast(::GetLastError()); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), + "mmap " + path.string()); + } + } + file_ = file; + data_ = static_cast(mapping); + size_ = size; +#else const int fd = ::open(path.c_str(), O_RDONLY | O_CLOEXEC | O_DIRECT); if (fd < 0) { throw std::system_error(errno, std::generic_category(), "open " + path.string()); @@ -208,11 +258,17 @@ class MappedFile { fd_ = fd; data_ = static_cast(mapping); size_ = size; +#endif } ~MappedFile() { +#if defined(_WIN32) + if (data_ != nullptr) { ::UnmapViewOfFile(const_cast(data_)); } + if (file_ != nullptr && file_ != INVALID_HANDLE_VALUE) { ::CloseHandle(file_); } +#else if (data_ != nullptr) { ::munmap(const_cast(data_), size_); } if (fd_ >= 0) { ::close(fd_); } +#endif } MappedFile(const MappedFile&) = delete; @@ -228,6 +284,23 @@ class MappedFile { reinterpret_cast(destination.data()) % alignment != 0) { throw ArtifactError("direct artifact read is not 4096-byte aligned"); } +#if defined(_WIN32) + if (absolute_offset > static_cast(std::numeric_limits::max()) || + destination.size() > static_cast(std::numeric_limits::max())) { + throw ArtifactError("direct artifact read exceeds platform I/O limits"); + } + + OVERLAPPED overlapped{}; + overlapped.Offset = static_cast(absolute_offset & 0xFFFFFFFFULL); + overlapped.OffsetHigh = static_cast(absolute_offset >> 32U); + DWORD bytes = 0; + if (::ReadFile(file_, destination.data(), static_cast(destination.size()), &bytes, + &overlapped) == 0) { + throw std::system_error(static_cast(::GetLastError()), std::system_category(), + "direct artifact read"); + } + return static_cast(bytes); +#else if (absolute_offset > static_cast(std::numeric_limits::max()) || destination.size() > static_cast(std::numeric_limits::max())) { throw ArtifactError("direct artifact read exceeds platform I/O limits"); @@ -242,10 +315,15 @@ class MappedFile { throw std::system_error(errno, std::generic_category(), "direct artifact read"); } return static_cast(bytes); +#endif } private: - int fd_ = -1; +#if defined(_WIN32) + HANDLE file_ = nullptr; +#else + int fd_ = -1; +#endif const std::byte* data_ = nullptr; std::size_t size_ = 0; }; From 9860d4324296951ef455f01fdba2c044aa0190ca Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:38 +0800 Subject: [PATCH 5/9] refactor(core): portable 128-bit multiply for context cost MSVC has no __int128; add u128_mul (_umul128 on x64) and saturating_u64_mul, then use them in PrefillWork, the materialization-search tie-break, and context-cost Q32 accounting. PrefillWork keeps the exact saturating semantics of the old 128-bit code: attention_pairs = min(max64, prefix*suffix + suffix*(suffix+1)/2) is summed limb-wise, and the triangular term halves a factor before multiplying so it never saturates early at 2^63. q32_product_ns reproduces the (product + 2^32 - 1) >> 32 rounding limb-wise; the saturating threshold max64 << 32 is compared exactly. --- src/core/math_util.h | 37 ++++++++++++++++++++ src/runtime/contract/types.h | 28 ++++++++++----- src/runtime/engine/context_cost.cpp | 29 ++++++++++----- src/runtime/engine/materialization_planner.h | 9 +++-- 4 files changed, 82 insertions(+), 21 deletions(-) create mode 100644 src/core/math_util.h diff --git a/src/core/math_util.h b/src/core/math_util.h new file mode 100644 index 0000000000..eef6cca2c3 --- /dev/null +++ b/src/core/math_util.h @@ -0,0 +1,37 @@ +#pragma once + +// Portable 128-bit unsigned multiply helpers. MSVC has no __int128; on x64 it +// provides _umul128 (full 128-bit product). Other platforms use __int128. + +#include +#include + +#if defined(_MSC_VER) +#include +#pragma intrinsic(_umul128) +#endif + +namespace ninfer::core { + +// Full 128-bit product of two 64-bit values. high receives the upper half; +// returns the lower half. +[[nodiscard]] inline std::uint64_t u128_mul(std::uint64_t left, std::uint64_t right, + std::uint64_t* high) noexcept { +#if defined(_MSC_VER) + return _umul128(left, right, reinterpret_cast(high)); +#else + const unsigned __int128 product = static_cast(left) * right; + *high = static_cast(product >> 64U); + return static_cast(product); +#endif +} + +// Saturated product: returns max(uint64) on overflow. +[[nodiscard]] inline std::uint64_t saturating_u64_mul(std::uint64_t left, + std::uint64_t right) noexcept { + std::uint64_t high = 0; + const std::uint64_t low = u128_mul(left, right, &high); + return high != 0 ? std::numeric_limits::max() : low; +} + +} // namespace ninfer::core diff --git a/src/runtime/contract/types.h b/src/runtime/contract/types.h index b2754cce47..0d37b01788 100644 --- a/src/runtime/contract/types.h +++ b/src/runtime/contract/types.h @@ -1,5 +1,6 @@ #pragma once +#include "core/math_util.h" #include "core/nvtx.h" #include "core/transfer_work.h" #include "ninfer/types.h" @@ -244,15 +245,24 @@ struct PrefillWork { result.tokens = suffix_tokens; result.vision_items = vision_items; result.vision_patches = vision_patches; - const unsigned __int128 suffix = suffix_tokens; - const unsigned __int128 linear = static_cast(prefix_tokens) * suffix; - const unsigned __int128 triangular = suffix * (suffix + 1U) / 2U; - constexpr unsigned __int128 maximum = ~static_cast(0); - const unsigned __int128 attention = - triangular > maximum - linear ? maximum : linear + triangular; - result.attention_pairs = attention > std::numeric_limits::max() - ? std::numeric_limits::max() - : static_cast(attention); + // attention_pairs = min(max64, prefix*suffix + suffix*(suffix+1)/2), computed with + // u128 limbs so the 128-bit sum saturates at max64 exactly as the old __int128 code did. + std::uint64_t linear_high = 0; + const std::uint64_t linear_low = core::u128_mul(prefix_tokens, suffix_tokens, &linear_high); + // One of suffix and suffix+1 is even, so halve before multiplying to stay in 64 bits. + // suffix+1 wraps only for suffix == max64 (odd), whose half is 2^63 and must be special-cased. + const bool suffix_even = (suffix_tokens & 1U) == 0U; + const std::uint64_t first = suffix_even ? suffix_tokens / 2U : suffix_tokens; + const std::uint64_t second = suffix_even + ? suffix_tokens + 1U + : (suffix_tokens == std::numeric_limits::max() + ? (std::numeric_limits::max() >> 1U) + 1U + : (suffix_tokens + 1U) / 2U); + std::uint64_t triangular_high = 0; + const std::uint64_t triangular_low = core::u128_mul(first, second, &triangular_high); + const std::uint64_t sum_low = linear_low + triangular_low; + const std::uint64_t sum_high = linear_high + triangular_high + (sum_low < linear_low ? 1U : 0U); + result.attention_pairs = sum_high != 0 ? std::numeric_limits::max() : sum_low; return result; } diff --git a/src/runtime/engine/context_cost.cpp b/src/runtime/engine/context_cost.cpp index fe0c885303..f27bccfa26 100644 --- a/src/runtime/engine/context_cost.cpp +++ b/src/runtime/engine/context_cost.cpp @@ -1,5 +1,7 @@ #include "runtime/engine/context_cost.h" +#include "core/math_util.h" + #include #include @@ -12,7 +14,12 @@ #include #include +#if defined(_WIN32) +#include +#define getpid _getpid +#else #include +#endif namespace ninfer::runtime { @@ -23,7 +30,6 @@ const std::vector& compiled_context_cost_defaults(); namespace { using Json = nlohmann::json; -using U128 = unsigned __int128; constexpr std::size_t direction_index(ContextTransferDirection direction) noexcept { return static_cast(direction); @@ -36,18 +42,23 @@ std::uint64_t saturating_add(std::uint64_t left, std::uint64_t right) noexcept { } std::uint64_t saturating_product(std::uint64_t left, std::uint64_t right) noexcept { - const U128 product = static_cast(left) * right; - return product > std::numeric_limits::max() - ? std::numeric_limits::max() - : static_cast(product); + return core::saturating_u64_mul(left, right); } std::uint64_t q32_product_ns(std::uint64_t coefficient, std::uint64_t units) noexcept { if (coefficient == 0 || units == 0) { return 0; } - const U128 product = static_cast(coefficient) * units; - const U128 maximum_scaled = static_cast(std::numeric_limits::max()) << 32U; - if (product >= maximum_scaled) { return std::numeric_limits::max(); } - return static_cast((product + kContextCostQ32One - 1U) >> 32U); + std::uint64_t high = 0; + const std::uint64_t low = core::u128_mul(coefficient, units, &high); + // product >= (max64 << 32) saturates. + constexpr std::uint64_t kHighLimit = std::numeric_limits::max(); + constexpr std::uint64_t kLowLimit = std::numeric_limits::max() << 32U; + if (high > kHighLimit || (high == kHighLimit && low >= kLowLimit)) { + return std::numeric_limits::max(); + } + // (product + kContextCostQ32One - 1U) >> 32U computed limb-wise. + const std::uint64_t low_plus = low + (kContextCostQ32One - 1U); + const std::uint64_t carry = low_plus < low ? 1ULL : 0ULL; + return ((high + carry) << 32U) | (low_plus >> 32U); } void require_object(const Json& value, std::string_view context) { diff --git a/src/runtime/engine/materialization_planner.h b/src/runtime/engine/materialization_planner.h index 64c47ee271..ac3a8b6eae 100644 --- a/src/runtime/engine/materialization_planner.h +++ b/src/runtime/engine/materialization_planner.h @@ -1,5 +1,6 @@ #pragma once +#include "core/math_util.h" #include "runtime/engine/context_cost.h" #include "runtime/engine/context_portfolio_value.h" #include "runtime/engine/materialization_budget.h" @@ -924,9 +925,11 @@ class MaterializationPlanner { ? item.estimated_total_ns - parent.estimated_total_ns : 0; }; - const __uint128_t left = static_cast<__uint128_t>(delta(cost)) * b; - const __uint128_t right = static_cast<__uint128_t>(delta(prior)) * a; - if (left != right) { return left < right; } + std::uint64_t left_high = 0, right_high = 0; + const auto left_low = core::u128_mul(delta(cost), b, &left_high); + const auto right_low = core::u128_mul(delta(prior), a, &right_high); + if (left_high != right_high) { return left_high < right_high; } + if (left_low != right_low) { return left_low < right_low; } } return cost.key() < prior.key(); } From 280334a5e13abd9ab1ae097a1ab603d3db65d3d1 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:43 +0800 Subject: [PATCH 6/9] fix(ops): pass nvfp4 TMA descriptors by device pointer on MSVC MSVC rejects the alignas(128) descriptor struct as a by-value kernel parameter (C2711), so on _MSC_VER builds the launchers copy the descriptors to a per-launch device buffer (cudaMalloc + cudaMemcpyAsync on the compute stream, cudaFreeAsync after the launch) and the kernels dereference a pointer. Other platforms keep the zero-copy __grid_constant__ parameter so the hot decode path is unchanged. --- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu | 18 ++++++++++ src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh | 26 ++++++++++---- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu | 17 +++++++++ .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 35 +++++++++++++------ 4 files changed, 79 insertions(+), 17 deletions(-) diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu index 19dacec69f..4c43c5c993 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu @@ -70,10 +70,28 @@ void launch_tma(const std::uint8_t* activation_codes, const std::uint8_t* activa }(); (void)kConfigured; + // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are + // passed to the kernel by device pointer instead of by value on Windows only. +#if defined(_MSC_VER) + Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; + CUDA_CHECK(cudaMalloc(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); +#endif + const dim3 grid(Geometry::kOutputRows / Schedule::kBlockN, tokens / Schedule::kBlockM); +#if defined(_MSC_VER) + nvfp4_w4a4_tma_kernel + <<>>(device_descriptors, alpha, epilogue, + output); +#else nvfp4_w4a4_tma_kernel <<>>(descriptors, alpha, epilogue, output); +#endif CUDA_CHECK(cudaGetLastError()); +#if defined(_MSC_VER) + CUDA_CHECK(cudaFreeAsync(device_descriptors, stream)); +#endif } template diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh index 6c0d1e3e28..fe6e1f14f7 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh @@ -163,12 +163,26 @@ __device__ __forceinline__ void nvfp4_tma_load_2d(void* destination, const CUten template __global__ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4_tma_kernel( - const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, float alpha, - const __grid_constant__ Epilogue epilogue, const __grid_constant__ OutputPolicy output) { +#if defined(_MSC_VER) + // MSVC rejects the over-aligned (alignas(128)) descriptors as a by-value kernel parameter + // (C2711), so they travel as a device pointer on Windows; other compilers keep the + // __grid_constant__ by-value parameter. The body dereferences them identically. + const Nvfp4W4a4TmaDescriptors* descriptors, +#else + const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, +#endif + float alpha, const __grid_constant__ Epilogue epilogue, + const __grid_constant__ OutputPolicy output) { static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); static_assert((Geometry::kOutputRows % Schedule::kBlockN) == 0); static_assert(Schedule::kStages >= 2, "the activation-scale buffer needs two slots"); +#if defined(_MSC_VER) + const Nvfp4W4a4TmaDescriptors& tma = *descriptors; +#else + const Nvfp4W4a4TmaDescriptors& tma = descriptors; +#endif + extern __shared__ __align__(128) unsigned char shared_bytes[]; auto& shared = *reinterpret_cast*>(shared_bytes); int block_x = 0; @@ -214,19 +228,19 @@ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4 : kTransactionBytes - kScaleBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &tma.a_codes, k_tile * Schedule::kCodeRowBytes, token_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_codes[stage], &descriptors.b_codes, + nvfp4_tma_load_2d(tensors.b_codes[stage], &tma.b_codes, k_tile * Schedule::kCodeRowBytes, row_begin, &shared.full[stage]); if (load_scales) { - nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &descriptors.a_scales, + nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &tma.a_scales, (k_tile / 2) * 16, token_begin, &shared.full[stage]); } const int b_scale_row = ((row_begin / 128) * Geometry::kScaleTilesPerRow + k_tile * Schedule::kK64PerStage) * 32; - nvfp4_tma_load_2d(tensors.b_scales[stage], &descriptors.b_scales, 0, b_scale_row, + nvfp4_tma_load_2d(tensors.b_scales[stage], &tma.b_scales, 0, b_scale_row, &shared.full[stage]); } } 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 0127a8d1d5..3df1e2082e 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 @@ -71,9 +71,26 @@ void launch_nvfp4_linear_swiglu_w4a4_tma(const std::uint8_t* activation_codes, 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); + + // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are + // passed to the kernel by device pointer instead of by value on Windows only. +#if defined(_MSC_VER) + Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; + CUDA_CHECK(cudaMalloc(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); +#endif +#if defined(_MSC_VER) + nvfp4_linear_swiglu_w4a4_tma_kernel + <<>>(device_descriptors, alpha, output); +#else nvfp4_linear_swiglu_w4a4_tma_kernel <<>>(descriptors, alpha, output); +#endif CUDA_CHECK(cudaGetLastError()); +#if defined(_MSC_VER) + CUDA_CHECK(cudaFreeAsync(device_descriptors, stream)); +#endif } } // namespace ninfer::ops::detail 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 7f85a36c85..5dc6f86936 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 @@ -46,11 +46,18 @@ template __global__ __launch_bounds__( Schedule::kThreads, Schedule:: - kMinBlocksPerSm) void nvfp4_linear_swiglu_w4a4_tma_kernel(const __grid_constant__ - Nvfp4W4a4TmaDescriptors - descriptors, - float alpha, - __nv_bfloat16* __restrict__ output) { + kMinBlocksPerSm) +#if defined(_MSC_VER) +// MSVC rejects the over-aligned (alignas(128)) descriptors as a by-value kernel parameter +// (C2711), so they travel as a device pointer on Windows; other compilers keep the +// __grid_constant__ by-value parameter. The body dereferences them identically. +void nvfp4_linear_swiglu_w4a4_tma_kernel(const Nvfp4W4a4TmaDescriptors* descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { +#else +void nvfp4_linear_swiglu_w4a4_tma_kernel( + const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { +#endif static_assert(Geometry::kOutputRows == 34816); static_assert(Geometry::kInputRows == 5120); static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); @@ -58,6 +65,12 @@ __global__ __launch_bounds__( static_assert(Schedule::kWarpsN == 2); static_assert(Schedule::kMmaN == 8); +#if defined(_MSC_VER) + const Nvfp4W4a4TmaDescriptors& tma = *descriptors; +#else + const Nvfp4W4a4TmaDescriptors& tma = descriptors; +#endif + constexpr int kIntermediate = Geometry::kOutputRows / 2; constexpr int kPairN = Schedule::kBlockN / 2; static_assert((kIntermediate % kPairN) == 0); @@ -109,17 +122,17 @@ __global__ __launch_bounds__( : kTransactionBytes - kScaleBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &tma.a_codes, k_tile * Schedule::kCodeRowBytes, token_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_codes[stage], &descriptors.b_codes, + nvfp4_tma_load_2d(tensors.b_codes[stage], &tma.b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin, &shared.full[stage]); nvfp4_tma_load_2d(tensors.b_codes[stage] + kPairN * Schedule::kCodeRowBytes, - &descriptors.b_codes, k_tile * Schedule::kCodeRowBytes, + &tma.b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin + kIntermediate, &shared.full[stage]); if (load_scales) { - nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &descriptors.a_scales, + nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &tma.a_scales, (k_tile / 2) * 16, token_begin, &shared.full[stage]); } @@ -130,9 +143,9 @@ __global__ __launch_bounds__( (((pair_begin + kIntermediate) / 128) * Geometry::kScaleTilesPerRow + k_tile * Schedule::kK64PerStage) * 32; - nvfp4_tma_load_2d(tensors.b_scales[stage][0], &descriptors.b_scales, 0, + nvfp4_tma_load_2d(tensors.b_scales[stage][0], &tma.b_scales, 0, gate_scale_row, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_scales[stage][1], &descriptors.b_scales, 0, + nvfp4_tma_load_2d(tensors.b_scales[stage][1], &tma.b_scales, 0, up_scale_row, &shared.full[stage]); } } From e7dee33dc25f6272c1e1a00393908b32d8fbc956 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:47 +0800 Subject: [PATCH 7/9] fix(targets): user-provided qwen3_6 runtime special members for MSVC linking MSVC does not emit out-of-class defaulted explicit specializations unless they are odr-used in the defining translation unit, which left undefined references that GCC weak symbols had masked. Provide user-defined move constructors, move assignments, and destructors instead, keeping the implicit noexcept of the old defaulted declarations explicit. --- src/targets/qwen3_6/impl/runtime/api_impl.h | 63 +++++++++++++++------ src/targets/qwen3_6/impl/runtime/program.h | 23 +++++--- 2 files changed, 61 insertions(+), 25 deletions(-) diff --git a/src/targets/qwen3_6/impl/runtime/api_impl.h b/src/targets/qwen3_6/impl/runtime/api_impl.h index 4c3f7e2a29..8ab6b2862d 100644 --- a/src/targets/qwen3_6/impl/runtime/api_impl.h +++ b/src/targets/qwen3_6/impl/runtime/api_impl.h @@ -17,12 +17,21 @@ SequencePlan::SequencePlan( std::unique_ptr> impl) noexcept : impl_(std::move(impl)) {} +// NOTE: MSVC does not emit out-of-class defaulted explicit specializations unless they are +// odr-used in the defining translation unit (GCC emits them unconditionally as weak symbols, +// which masked this on Linux). User-provided bodies guarantee the symbols exist for the linker. template <> -SequencePlan::SequencePlan(SequencePlan&&) noexcept = default; +SequencePlan::SequencePlan(SequencePlan&& other) noexcept + : impl_(std::move(other.impl_)) {} + template <> -SequencePlan& SequencePlan::operator=(SequencePlan&&) noexcept = default; +SequencePlan& SequencePlan::operator=(SequencePlan&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} + template <> -SequencePlan::~SequencePlan() = default; +SequencePlan::~SequencePlan() noexcept {} template <> std::uint32_t SequencePlan::capacity() const noexcept { @@ -55,11 +64,17 @@ SequencePlanner::SequencePlanner( : impl_(std::move(impl)) {} template <> -SequencePlanner::SequencePlanner(SequencePlanner&&) noexcept = default; +SequencePlanner::SequencePlanner(SequencePlanner&& other) noexcept + : impl_(std::move(other.impl_)) {} + template <> -SequencePlanner& SequencePlanner::operator=(SequencePlanner&&) noexcept = default; +SequencePlanner& SequencePlanner::operator=(SequencePlanner&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} + template <> -SequencePlanner::~SequencePlanner() = default; +SequencePlanner::~SequencePlanner() noexcept {} template <> const runtime::SequenceCapacityCurve& SequencePlanner::capacity_curve() const noexcept { @@ -80,11 +95,17 @@ RequestBasePlan::RequestBasePlan( : impl_(std::move(impl)) {} template <> -RequestBasePlan::RequestBasePlan(RequestBasePlan&&) noexcept = default; +RequestBasePlan::RequestBasePlan(RequestBasePlan&& other) noexcept + : impl_(std::move(other.impl_)) {} + template <> -RequestBasePlan& RequestBasePlan::operator=(RequestBasePlan&&) noexcept = default; +RequestBasePlan& RequestBasePlan::operator=(RequestBasePlan&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} + template <> -RequestBasePlan::~RequestBasePlan() = default; +RequestBasePlan::~RequestBasePlan() noexcept {} template <> const runtime::RequestPlanSummary& RequestBasePlan::summary() const noexcept { @@ -130,26 +151,34 @@ PressurePlanningSession::PressurePlanningSession( : impl_(std::move(impl)) {} template <> -PressurePlanningSession::PressurePlanningSession(PressurePlanningSession&&) noexcept = - default; +PressurePlanningSession::PressurePlanningSession(PressurePlanningSession&& other) noexcept + : impl_(std::move(other.impl_)) {} template <> PressurePlanningSession& -PressurePlanningSession::operator=(PressurePlanningSession&&) noexcept = default; +PressurePlanningSession::operator=(PressurePlanningSession&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} template <> -PressurePlanningSession::~PressurePlanningSession() = default; +PressurePlanningSession::~PressurePlanningSession() noexcept {} template <> CapturePressurePlanningSession::CapturePressurePlanningSession( - CapturePressurePlanningSession&&) noexcept = default; + CapturePressurePlanningSession&& other) noexcept + : candidate_(std::move(other.candidate_)), session_(std::move(other.session_)) {} template <> CapturePressurePlanningSession& CapturePressurePlanningSession::operator=( - CapturePressurePlanningSession&&) noexcept = default; + CapturePressurePlanningSession&& other) noexcept { + candidate_ = std::move(other.candidate_); + session_ = std::move(other.session_); + return *this; +} template <> -CapturePressurePlanningSession::~CapturePressurePlanningSession() = default; +CapturePressurePlanningSession::~CapturePressurePlanningSession() noexcept {} template <> PressureTargetHandle @@ -309,7 +338,7 @@ Program::Program(std::unique_ptr> impl) no : impl_(std::move(impl)) {} template <> -Program::~Program() noexcept = default; +Program::~Program() noexcept {} template <> RequestBasePlan diff --git a/src/targets/qwen3_6/impl/runtime/program.h b/src/targets/qwen3_6/impl/runtime/program.h index daf47b3861..a3231cbfce 100644 --- a/src/targets/qwen3_6/impl/runtime/program.h +++ b/src/targets/qwen3_6/impl/runtime/program.h @@ -275,15 +275,19 @@ AdmissionCandidate::AdmissionCandidate( : impl_(std::move(impl)) {} template <> -AdmissionCandidate::AdmissionCandidate(AdmissionCandidate&&) noexcept = - default; +AdmissionCandidate::AdmissionCandidate( + AdmissionCandidate&& other) noexcept + : impl_(std::move(other.impl_)) {} template <> AdmissionCandidate& -AdmissionCandidate::operator=(AdmissionCandidate&&) noexcept = default; +AdmissionCandidate::operator=(AdmissionCandidate&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} template <> -AdmissionCandidate::~AdmissionCandidate() = default; +AdmissionCandidate::~AdmissionCandidate() noexcept {} template <> CapturePressureCandidate::CapturePressureCandidate( @@ -292,15 +296,18 @@ CapturePressureCandidate::CapturePressureCandidate( template <> CapturePressureCandidate::CapturePressureCandidate( - CapturePressureCandidate&&) noexcept = default; + CapturePressureCandidate&& other) noexcept + : impl_(std::move(other.impl_)) {} template <> CapturePressureCandidate& -CapturePressureCandidate::operator=(CapturePressureCandidate&&) noexcept = - default; +CapturePressureCandidate::operator=(CapturePressureCandidate&& other) noexcept { + impl_ = std::move(other.impl_); + return *this; +} template <> -CapturePressureCandidate::~CapturePressureCandidate() = default; +CapturePressureCandidate::~CapturePressureCandidate() noexcept {} template <> const runtime::RequestPlanSummary& From 91fd83ca9ffe381f0aa6e8e7295a9263bb1059b1 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sat, 12 Sep 2026 23:59:47 +0800 Subject: [PATCH 8/9] fix(product/serve): Windows console TTY, localtime, pid, and terminal width - logging: _isatty on the redirected/console stderr instead of POSIX isatty; localtime_s argument order vs localtime_r. - startup log: CONSOLE_SCREEN_BUFFER_INFO width when the progress bar targets a console, ioctl TIOCGWINSZ otherwise. - request log and context cost: _getpid on Windows (see refactor(core) for the context-cost side). --- src/product/logging/logging.cpp | 11 +++++++++++ src/product/logging/startup_log.cpp | 13 +++++++++++++ src/serve/request_log.cpp | 5 +++++ 3 files changed, 29 insertions(+) diff --git a/src/product/logging/logging.cpp b/src/product/logging/logging.cpp index 2be483c415..2423d3391f 100644 --- a/src/product/logging/logging.cpp +++ b/src/product/logging/logging.cpp @@ -5,7 +5,13 @@ #include #include +#if defined(_WIN32) +#include +#define isatty _isatty +#define STDERR_FILENO 2 +#else #include +#endif #include #include @@ -101,7 +107,12 @@ class PrettyLogFormatter final : public spdlog::formatter { const std::time_t wall_seconds = std::chrono::system_clock::to_time_t( std::chrono::system_clock::time_point(whole_seconds)); std::tm local{}; +#ifdef _WIN32 + // localtime_s swaps argument order relative to POSIX localtime_r. + localtime_s(&local, &wall_seconds); +#else localtime_r(&wall_seconds, &local); +#endif fmt::format_to(std::back_inserter(destination), "{:04}-{:02}-{:02} {:02}:{:02}:{:02}.{:03} ", local.tm_year + 1900, local.tm_mon + 1, local.tm_mday, local.tm_hour, local.tm_min, diff --git a/src/product/logging/startup_log.cpp b/src/product/logging/startup_log.cpp index f8faf09288..0ce6308623 100644 --- a/src/product/logging/startup_log.cpp +++ b/src/product/logging/startup_log.cpp @@ -5,8 +5,12 @@ #include +#if defined(_WIN32) +#include +#else #include #include +#endif #include #include @@ -76,9 +80,18 @@ PhasePresentation phase_presentation(StartupPhase phase) noexcept { } std::size_t terminal_columns() noexcept { +#if defined(_WIN32) + CONSOLE_SCREEN_BUFFER_INFO info{}; + if (::GetConsoleScreenBufferInfo(::GetStdHandle(STD_ERROR_HANDLE), &info) != 0) { + const std::size_t width = static_cast(info.dwSize.X); + if (width != 0) { return width; } + } + return 120; +#else winsize size{}; if (::ioctl(STDERR_FILENO, TIOCGWINSZ, &size) == 0 && size.ws_col != 0) { return size.ws_col; } return 120; +#endif } std::string progress_bar(double ratio, std::size_t width) { diff --git a/src/serve/request_log.cpp b/src/serve/request_log.cpp index ca246bb6bc..c379fa59d3 100644 --- a/src/serve/request_log.cpp +++ b/src/serve/request_log.cpp @@ -17,7 +17,12 @@ #include #include +#if defined(_WIN32) +#include +#define getpid _getpid +#else #include +#endif namespace ninfer::serve { namespace { From 153b6a26db1bdd34007090fe111b841537432c80 Mon Sep 17 00:00:00 2001 From: troubadour-hell <17946453+troubadour-hell@users.noreply.github.com> Date: Sun, 13 Sep 2026 00:32:50 +0800 Subject: [PATCH 9/9] perf(ops,artifact): stream-ordered TMA descriptor allocation and unbuffered Windows reads Address Codex review feedback on PR #233: - Replace synchronous cudaMalloc with stream-ordered cudaMallocAsync for TMA descriptors on MSVC. The stream memory pool reuses the same slot on subsequent launches, eliminating per-launch allocation overhead in the prefill hot path. (nvfp4_w4a4_tma.cu, nvfp4_linear_swiglu_w4a4_tma.cu) - Add FILE_FLAG_NO_BUFFERING to the Windows artifact file handle so ReadFile bypasses the system cache, matching POSIX O_DIRECT semantics. The existing 4096-byte alignment contract satisfies the flag's sector-alignment requirements. (reader.cpp) --- src/artifact/reader.cpp | 7 ++++++- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu | 4 +++- .../linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu | 4 +++- 3 files changed, 12 insertions(+), 3 deletions(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 57de9ba123..ffd997a9c8 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -182,8 +182,13 @@ class MappedFile { public: explicit MappedFile(const std::filesystem::path& path) { #if defined(_WIN32) + // FILE_FLAG_NO_BUFFERING makes ReadFile bypass the system cache, matching the + // POSIX O_DIRECT semantics. It has no effect on MapViewOfFile, which uses the + // OS paging system independently. The 4096-byte alignment contract enforced in + // read_direct() satisfies the flag's sector-alignment requirements. const HANDLE file = ::CreateFileW(path.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, - OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, nullptr); + OPEN_EXISTING, + FILE_ATTRIBUTE_NORMAL | FILE_FLAG_NO_BUFFERING, nullptr); if (file == INVALID_HANDLE_VALUE) { throw std::system_error(static_cast(::GetLastError()), std::system_category(), "open " + path.string()); diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu index 4c43c5c993..7d2d72ed7c 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu @@ -72,9 +72,11 @@ void launch_tma(const std::uint8_t* activation_codes, const std::uint8_t* activa // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are // passed to the kernel by device pointer instead of by value on Windows only. + // cudaMallocAsync keeps the entire alloc-copy-launch-free cycle stream-ordered; + // the stream memory pool reuses the same slot on subsequent launches. #if defined(_MSC_VER) Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; - CUDA_CHECK(cudaMalloc(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMallocAsync(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors), stream)); CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), cudaMemcpyHostToDevice, stream)); #endif 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 3df1e2082e..998d74b4c3 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 @@ -74,9 +74,11 @@ void launch_nvfp4_linear_swiglu_w4a4_tma(const std::uint8_t* activation_codes, // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are // passed to the kernel by device pointer instead of by value on Windows only. + // cudaMallocAsync keeps the entire alloc-copy-launch-free cycle stream-ordered; + // the stream memory pool reuses the same slot on subsequent launches. #if defined(_MSC_VER) Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; - CUDA_CHECK(cudaMalloc(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMallocAsync(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors), stream)); CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), cudaMemcpyHostToDevice, stream)); #endif