From bc9dca339e3a0baea0769dab065158a95b128b63 Mon Sep 17 00:00:00 2001 From: pelebel Date: Wed, 19 Aug 2026 14:13:23 -0400 Subject: [PATCH 01/20] feat(windows): build and run the engine on Windows 11 (MSVC + CUDA) Build and run ninfer / ninfer-serve with cl (VS2022) + nvcc (CUDA 13.x) + Ninja, resolving FFmpeg and libcurl from a vcpkg x64-windows prefix. - CMake: per-language MSVC flags (/Zc:preprocessor, NOMINMAX, UTF8PROC_STATIC) so CUDA objects keep their cache; new cmake/NInferMediaDeps.cmake locates FFmpeg/curl on Windows and copies their runtime DLLs next to the executables at build time. - Portability shims: Winsock (media_acquire), CreateFileW/MapViewOfFile artifact reader, localtime_s / GetCurrentProcessId / _isatty. - Pass the 128-byte NVFP4 TMA descriptors by device pointer so cl accepts the aligned by-value parameters (linear / linear-swiglu kernels). - Give the qwen3_6 plan move constructors explicit bodies: MSVC does not emit an out-of-line '= default' explicit specialization unless ODR-used. --- CMakeLists.txt | 31 ++++- apps/CMakeLists.txt | 21 +++ cmake/NInferMediaDeps.cmake | 124 ++++++++++++++++++ src/CMakeLists.txt | 12 +- src/artifact/reader.cpp | 81 +++++++++++- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu | 56 +++++++- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh | 10 +- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu | 54 +++++++- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 20 ++- src/product/load_progress/load_progress.cpp | 14 +- src/product/media_acquire/acquire.cpp | 28 +++- src/serve/console_log.cpp | 4 + src/serve/request_log.cpp | 16 ++- src/targets/qwen3_6/impl/runtime/api_impl.h | 14 +- tests/test_request_log.cpp | 16 ++- 15 files changed, 466 insertions(+), 35 deletions(-) create mode 100644 cmake/NInferMediaDeps.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index ca3f6c48e3..dec864d721 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -55,14 +55,37 @@ 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) + +# FFmpeg and libcurl are located cross-platform: pkg-config on POSIX, an +# install prefix / vcpkg tree on Windows. Both are exposed as NInfer::FFmpeg +# and NInfer::Curl (see cmake/NInferMediaDeps.cmake). +list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake") +include(NInferMediaDeps) +ninfer_find_ffmpeg() if(NINFER_BUILD_MEDIA_ACQUIRE) - pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + ninfer_find_curl() endif() + find_package(Threads REQUIRED) +if(MSVC) + # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. The + # flag is forwarded to the cl host compiler nvcc invokes, including for the + # host stubs it generates. (The 128-byte-aligned TMA descriptor parameters are + # passed by device pointer rather than by value precisely so MSVC's stub + # compiler does not reject them.) + add_compile_options($<$:/Zc:preprocessor>) + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=/Zc:preprocessor") + # windows.h defines min/max as macros which clobber std::min/std::max; NOMINMAX + # suppresses them. UTF8PROC_STATIC keeps the vendored utf8proc header from + # marking its symbols __declspec(dllimport) since we compile it statically. Both + # are scoped to C/CXX (no CUDA file needs them) so the CUDA objects keep their + # cache. + add_compile_definitions( + $<$:NOMINMAX> + $<$:UTF8PROC_STATIC>) +endif() + # --- Subdirectories ---------------------------------------------------------- add_subdirectory(src) if(NINFER_BUILD_APPS) diff --git a/apps/CMakeLists.txt b/apps/CMakeLists.txt index fe80287518..8d95296894 100644 --- a/apps/CMakeLists.txt +++ b/apps/CMakeLists.txt @@ -17,3 +17,24 @@ 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_load_progress) + +# On Windows the FFmpeg/libcurl runtime DLLs (from the vcpkg tree or an explicit +# prefix) are not on the default search path, so copy them next to each +# executable to make the build-tree binaries run in place. +if(WIN32) + set(_ninfer_runtime_dll_dirs "") + foreach(_prefix ${NINFER_FFMPEG_PREFIX} ${NINFER_CURL_PREFIX}) + if(_prefix AND EXISTS "${_prefix}/bin") + list(APPEND _ninfer_runtime_dll_dirs "${_prefix}/bin") + endif() + endforeach() + list(REMOVE_DUPLICATES _ninfer_runtime_dll_dirs) + foreach(_app ninfer ninfer-serve) + foreach(_dll_dir IN LISTS _ninfer_runtime_dll_dirs) + add_custom_command(TARGET ${_app} POST_BUILD + COMMAND ${CMAKE_COMMAND} -E copy_directory "${_dll_dir}" + "$" + COMMENT "Copying ${_app} runtime dependencies") + endforeach() + endforeach() +endif() diff --git a/cmake/NInferMediaDeps.cmake b/cmake/NInferMediaDeps.cmake new file mode 100644 index 0000000000..2b0d3bd52c --- /dev/null +++ b/cmake/NInferMediaDeps.cmake @@ -0,0 +1,124 @@ +# NInferMediaDeps.cmake +# +# Locates FFmpeg and libcurl and exposes them as INTERFACE imported targets +# NInfer::FFmpeg and NInfer::Curl, which the ninfer_media_* libraries link. +# +# * POSIX : discovered through pkg-config (unchanged behaviour). +# * Windows : located from a vcpkg install tree or an explicit prefix, because +# pkg-config is not part of the Windows toolchain. The prefix is +# chosen, per package, in this order: NINFER_FFMPEG_ROOT / +# NINFER_CURL_ROOT (per package), then NINFER_MEDIA_ROOT, then +# $ENV{VCPKG_ROOT}/installed/. +# +# On Windows the resolved prefixes are cached as NINFER_FFMPEG_PREFIX and +# NINFER_CURL_PREFIX so the app targets can copy the dependency DLLs next to the +# built executables at build time. + +if(NOT WIN32) + find_package(PkgConfig REQUIRED) +endif() + +function(ninfer_media_resolve_prefix _out) + set(result "") + if(DEFINED NINFER_MEDIA_ROOT) + set(result "${NINFER_MEDIA_ROOT}") + elseif(DEFINED ENV{NINFER_MEDIA_ROOT}) + set(result "$ENV{NINFER_MEDIA_ROOT}") + elseif(DEFINED ENV{VCPKG_ROOT}) + set(triplet "x64-windows") + if(DEFINED ENV{VCPKG_TARGET_TRIPLET}) + set(triplet "$ENV{VCPKG_TARGET_TRIPLET}") + endif() + if(EXISTS "$ENV{VCPKG_ROOT}/installed/${triplet}/include") + set(result "$ENV{VCPKG_ROOT}/installed/${triplet}") + endif() + endif() + set(${_out} "${result}" PARENT_SCOPE) +endfunction() + +function(ninfer_require_prefix _label _root _hint) + if(NOT _root OR NOT EXISTS "${_root}/include") + message(FATAL_ERROR + "${_label} not found. Point the build at an install prefix containing " + "include/ and lib/. ${_hint}") + endif() +endfunction() + +# Bundle an include directory and a link list (imported libs plus, on Windows, +# the system libraries the package needs) into an INTERFACE imported target. +function(ninfer_make_imported_target interface include_dir libs) + add_library(${interface} INTERFACE IMPORTED) + set_target_properties(${interface} PROPERTIES + INTERFACE_INCLUDE_DIRECTORIES "${include_dir}" + INTERFACE_LINK_LIBRARIES "${libs}") +endfunction() + +function(ninfer_find_ffmpeg) + if(TARGET NInfer::FFmpeg) + return() + endif() + if(NOT WIN32) + pkg_check_modules(_ninfer_ffmpeg REQUIRED IMPORTED_TARGET + libavformat>=60 libavcodec>=60 libavutil>=58 libswscale>=7) + ninfer_make_imported_target(NInfer::FFmpeg + "${_ninfer_ffmpeg_INCLUDE_DIRS}" "${_ninfer_ffmpeg_LIBRARIES}") + return() + endif() + + if(DEFINED NINFER_FFMPEG_ROOT) + set(root "${NINFER_FFMPEG_ROOT}") + elseif(DEFINED ENV{NINFER_FFMPEG_ROOT}) + set(root "$ENV{NINFER_FFMPEG_ROOT}") + else() + ninfer_media_resolve_prefix(root) + endif() + ninfer_require_prefix("FFmpeg" "${root}" + "set NINFER_FFMPEG_ROOT or NINFER_MEDIA_ROOT to an FFmpeg prefix, or install it " + "with `vcpkg install ffmpeg --triplet x64-windows` and set VCPKG_ROOT.") + + find_path(NINFER_FFMPEG_INCLUDE_DIR libavcodec/avcodec.h + PATHS "${root}/include" NO_DEFAULT_PATH REQUIRED) + foreach(module IN ITEMS avformat avcodec swscale avutil) + find_library(NINFER_FFMPEG_${module}_LIB NAMES ${module} + PATHS "${root}/lib" NO_DEFAULT_PATH REQUIRED) + endforeach() + + # The FFmpeg import libraries reference the Windows multimedia/network API. + # (The list contains only import libraries that exist in the Windows SDK; e.g. + # there is no Ntmapi.lib, and Nt*/Rtl* symbols, if ever needed, come from ntdll.) + ninfer_make_imported_target(NInfer::FFmpeg "${NINFER_FFMPEG_INCLUDE_DIR}" + "${NINFER_FFMPEG_avformat_LIB};${NINFER_FFMPEG_avcodec_LIB};${NINFER_FFMPEG_swscale_LIB};${NINFER_FFMPEG_avutil_LIB};Strmiids;Ws2_32;Advapi32;Ole32;Oleaut32;User32;Iphlpapi;Userenv") + set(NINFER_FFMPEG_PREFIX "${root}" CACHE INTERNAL "Resolved FFmpeg install prefix") +endfunction() + +function(ninfer_find_curl) + if(TARGET NInfer::Curl) + return() + endif() + if(NOT WIN32) + pkg_check_modules(_ninfer_curl REQUIRED IMPORTED_TARGET libcurl>=7.85) + ninfer_make_imported_target(NInfer::Curl + "${_ninfer_curl_INCLUDE_DIRS}" "${_ninfer_curl_LIBRARIES}") + return() + endif() + + if(DEFINED NINFER_CURL_ROOT) + set(root "${NINFER_CURL_ROOT}") + elseif(DEFINED ENV{NINFER_CURL_ROOT}) + set(root "$ENV{NINFER_CURL_ROOT}") + else() + ninfer_media_resolve_prefix(root) + endif() + ninfer_require_prefix("libcurl" "${root}" + "set NINFER_CURL_ROOT or NINFER_MEDIA_ROOT to a libcurl prefix, or install it " + "with `vcpkg install curl --triplet x64-windows` and set VCPKG_ROOT.") + + find_path(NINFER_CURL_INCLUDE_DIR curl/curl.h + PATHS "${root}/include" NO_DEFAULT_PATH REQUIRED) + find_library(NINFER_CURL_LIB NAMES curl libcurl + PATHS "${root}/lib" NO_DEFAULT_PATH REQUIRED) + + ninfer_make_imported_target(NInfer::Curl "${NINFER_CURL_INCLUDE_DIR}" + "${NINFER_CURL_LIB};Ws2_32;Crypt32;Cryptui;Normaliz;Bcrypt;Advapi32;Userenv") + set(NINFER_CURL_PREFIX "${root}" CACHE INTERNAL "Resolved libcurl install prefix") +endfunction() \ No newline at end of file diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index f5590f3f77..33167d61cc 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -273,14 +273,18 @@ ninfer_internal_includes(ninfer_text) add_library(ninfer_media_decode STATIC media/decode/decode.cpp) ninfer_internal_includes(ninfer_media_decode) -target_link_libraries(ninfer_media_decode PRIVATE PkgConfig::FFMPEG) +target_link_libraries(ninfer_media_decode PRIVATE NInfer::FFmpeg) if(NINFER_BUILD_MEDIA_ACQUIRE) # Product-only path/data/HTTP acquisition. No target package links this library. add_library(ninfer_media_acquire STATIC product/media_acquire/acquire.cpp) ninfer_internal_includes(ninfer_media_acquire) - target_link_libraries(ninfer_media_acquire PRIVATE PkgConfig::LIBCURL) + target_link_libraries(ninfer_media_acquire PRIVATE NInfer::Curl) + if(WIN32) + # Host name resolution / Winsock (getaddrinfo, inet_ntop, WSAStartup). + target_link_libraries(ninfer_media_acquire PRIVATE Ws2_32) + endif() endif() if(NINFER_BUILD_PROMPT_INPUT) @@ -340,4 +344,8 @@ if(NINFER_BUILD_SERVE) target_link_libraries(ninfer_serve PUBLIC ninfer_engine Threads::Threads PRIVATE ninfer_media_acquire CUDA::cudart) + if(WIN32) + # cpp-httplib (plain HTTP) uses the Winsock API. + target_link_libraries(ninfer_serve PRIVATE Ws2_32) + endif() endif() diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 1dc3afd1ea..4f61bd2237 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -15,10 +15,18 @@ #include #include +#ifdef _WIN32 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#define WIN32_LEAN_AND_MEAN +#include +#else #include #include #include #include +#endif namespace ninfer::artifact { namespace { @@ -181,6 +189,55 @@ struct TransparentStringHash { class MappedFile { public: explicit MappedFile(const std::filesystem::path& path) { +#ifdef _WIN32 + const std::wstring wide = path.wstring(); + HANDLE file = ::CreateFileW(wide.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, + OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, nullptr); + if (file == INVALID_HANDLE_VALUE) { + const auto error = ::GetLastError(); + throw std::system_error(error, std::system_category(), "open " + path.string()); + } + + LARGE_INTEGER size_info {}; + + if (!::GetFileSizeEx(file, &size_info)) { + const auto error = ::GetLastError(); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), "GetFileSizeEx " + path.string()); + } + if (size_info.QuadPart < 0 || + size_info.QuadPart > static_cast(std::numeric_limits::max())) { + ::CloseHandle(file); + throw ArtifactError("artifact size does not fit the process address space"); + } + + const auto size = static_cast(size_info.QuadPart); + HANDLE mapping = nullptr; + void* view = nullptr; + if (size != 0) { + mapping = ::CreateFileMappingW(file, nullptr, PAGE_READONLY, 0, 0, nullptr); + if (mapping == nullptr) { + const auto error = ::GetLastError(); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), + "CreateFileMapping " + path.string()); + } + // A zero byte count maps the entire section, which avoids the 32-bit + // dwNumberOfBytesToMap limit for artifacts larger than 4 GiB. + view = ::MapViewOfFile(mapping, FILE_MAP_READ, 0, 0, 0); + if (view == nullptr) { + const auto error = ::GetLastError(); + ::CloseHandle(mapping); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), "MapViewOfFile " + path.string()); + } + } + file_ = file; + mapping_ = mapping; + view_ = view; + data_ = static_cast(view); + 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()); @@ -212,11 +269,18 @@ class MappedFile { fd_ = fd; data_ = static_cast(mapping); size_ = size; +#endif } ~MappedFile() { +#ifdef _WIN32 + if (view_ != nullptr) { ::UnmapViewOfFile(view_); } + if (mapping_ != nullptr) { ::CloseHandle(mapping_); } + if (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; @@ -232,6 +296,14 @@ class MappedFile { reinterpret_cast(destination.data()) % alignment != 0) { throw ArtifactError("direct artifact read is not 4096-byte aligned"); } +#ifdef _WIN32 + // The whole file is already mapped, so a direct read is a copy out of the view. + if (absolute_offset > size_ || destination.size() > size_ - absolute_offset) { + throw ArtifactError("direct artifact read exceeds the mapped file"); + } + std::memcpy(destination.data(), data_ + absolute_offset, destination.size()); + return destination.size(); +#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"); @@ -246,10 +318,17 @@ class MappedFile { throw std::system_error(errno, std::generic_category(), "direct artifact read"); } return static_cast(bytes); +#endif } private: - int fd_ = -1; +#ifdef _WIN32 + HANDLE file_ = INVALID_HANDLE_VALUE; + HANDLE mapping_ = nullptr; + void* view_ = nullptr; +#else + int fd_ = -1; +#endif const std::byte* data_ = nullptr; std::size_t size_ = 0; }; diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu index 19dacec69f..9b7e1e167c 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu @@ -9,6 +9,7 @@ #include #include +#include namespace ninfer::ops::detail { namespace { @@ -16,6 +17,54 @@ namespace { using TmaM256N128 = Nvfp4W4a4TmaSchedule<256, 3, 1>; using TmaM256N128S2 = Nvfp4W4a4TmaSchedule<256, 2, 1>; +// MSVC's cl rejects a by-value 128-byte-aligned TMA descriptor parameter in the +// nvcc-generated kernel host stub (C2719); Clang/GCC accept it. So the kernels +// take a device pointer. One persistent host copy (a stable source for the HtoD +// memcpy node captured into decode CUDA graphs) and one device copy are kept per +// distinct descriptor; a descriptor is constant for a given (buffers, tokens), +// so the device copy stays valid across graph replays. +struct Nvfp4TmaDescKey { + const void* activation_codes = nullptr; + const void* activation_scales = nullptr; + const void* weight_codes = nullptr; + const void* weight_scales = nullptr; + std::int32_t tokens = 0; + + bool operator==(const Nvfp4TmaDescKey& o) const { + return activation_codes == o.activation_codes && + activation_scales == o.activation_scales && weight_codes == o.weight_codes && + weight_scales == o.weight_scales && tokens == o.tokens; + } +}; + +struct Nvfp4TmaDescSlot { + Nvfp4TmaDescKey key; + Nvfp4W4a4TmaDescriptors* host = nullptr; + Nvfp4W4a4TmaDescriptors* device = nullptr; +}; + +Nvfp4W4a4TmaDescriptors* nvfp4_tma_desc_device(const Nvfp4TmaDescKey& key, + const Nvfp4W4a4TmaDescriptors& host_desc, + cudaStream_t stream) { + static std::vector slots; + for (Nvfp4TmaDescSlot& slot : slots) { + if (slot.key == key) { + *slot.host = host_desc; + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + return slot.device; + } + } + Nvfp4TmaDescSlot slot; + slot.key = key; + slot.host = new Nvfp4W4a4TmaDescriptors(host_desc); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&slot.device), sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + slots.push_back(slot); + return slot.device; +} + constexpr std::int32_t kQueryRows = 6144; constexpr std::int32_t kKeyRows = 1024; constexpr std::int32_t kGateRows = 6144; @@ -70,9 +119,14 @@ void launch_tma(const std::uint8_t* activation_codes, const std::uint8_t* activa }(); (void)kConfigured; + const Nvfp4W4a4TmaDescriptors* device_descriptors = nvfp4_tma_desc_device( + Nvfp4TmaDescKey{activation_codes, activation_scales, weight_codes, weight_scales, tokens}, + descriptors, stream); + const dim3 grid(Geometry::kOutputRows / Schedule::kBlockN, tokens / Schedule::kBlockM); nvfp4_w4a4_tma_kernel - <<>>(descriptors, alpha, epilogue, output); + <<>>(device_descriptors, alpha, epilogue, + output); CUDA_CHECK(cudaGetLastError()); } diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh index aa6914eca5..694a3d1fd1 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh @@ -182,7 +182,7 @@ __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 Nvfp4W4a4TmaDescriptors* descriptors, 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); @@ -222,17 +222,17 @@ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4 nvfp4_mbarrier_arrive_expect_tx(&shared.full[stage], kTransactionBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors->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], &descriptors->b_codes, k_tile * Schedule::kCodeRowBytes, row_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors.a_scales, (k_tile / 2) * 16, + nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors->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], &descriptors->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..e90906345e 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 @@ -8,12 +8,61 @@ #include #include #include +#include namespace ninfer::ops::detail { namespace { using M256N128S3 = Nvfp4W4a4TmaSchedule<256, 3, 1>; +// MSVC's cl rejects a by-value 128-byte-aligned TMA descriptor parameter in the +// nvcc-generated kernel host stub (C2719); Clang/GCC accept it. So the kernel +// takes a device pointer. One persistent host copy (a stable source for the HtoD +// memcpy node captured into decode CUDA graphs) and one device copy are kept per +// distinct descriptor; a descriptor is constant for a given (buffers, tokens), +// so the device copy stays valid across graph replays. +struct Nvfp4TmaDescKey { + const void* activation_codes = nullptr; + const void* activation_scales = nullptr; + const void* weight_codes = nullptr; + const void* weight_scales = nullptr; + std::int32_t tokens = 0; + + bool operator==(const Nvfp4TmaDescKey& o) const { + return activation_codes == o.activation_codes && + activation_scales == o.activation_scales && weight_codes == o.weight_codes && + weight_scales == o.weight_scales && tokens == o.tokens; + } +}; + +struct Nvfp4TmaDescSlot { + Nvfp4TmaDescKey key; + Nvfp4W4a4TmaDescriptors* host = nullptr; + Nvfp4W4a4TmaDescriptors* device = nullptr; +}; + +Nvfp4W4a4TmaDescriptors* nvfp4_tma_desc_device(const Nvfp4TmaDescKey& key, + const Nvfp4W4a4TmaDescriptors& host_desc, + cudaStream_t stream) { + static std::vector slots; + for (Nvfp4TmaDescSlot& slot : slots) { + if (slot.key == key) { + *slot.host = host_desc; + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + return slot.device; + } + } + Nvfp4TmaDescSlot slot; + slot.key = key; + slot.host = new Nvfp4W4a4TmaDescriptors(host_desc); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&slot.device), sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + slots.push_back(slot); + return slot.device; +} + template Nvfp4W4a4TmaDescriptors make_descriptors(const std::uint8_t* activation_codes, const std::uint8_t* activation_scales, @@ -70,9 +119,12 @@ 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 Nvfp4W4a4TmaDescriptors* device_descriptors = nvfp4_tma_desc_device( + Nvfp4TmaDescKey{activation_codes, activation_scales, weight_codes, weight_scales, tokens}, + descriptors, stream); const dim3 grid((Geometry::kOutputRows / 2) / kPairN, tokens / M256N128S3::kBlockM); nvfp4_linear_swiglu_w4a4_tma_kernel - <<>>(descriptors, alpha, output); + <<>>(device_descriptors, alpha, output); 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 a7664c7da7..65f0354e0b 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,9 @@ 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) void nvfp4_linear_swiglu_w4a4_tma_kernel( + const Nvfp4W4a4TmaDescriptors* descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { static_assert(Geometry::kOutputRows == 34816); static_assert(Geometry::kInputRows == 5120); static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); @@ -98,16 +96,16 @@ __global__ __launch_bounds__( nvfp4_mbarrier_arrive_expect_tx(&shared.full[stage], kTransactionBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors->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], &descriptors->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, + &descriptors->b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin + kIntermediate, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors.a_scales, (k_tile / 2) * 16, + nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors->a_scales, (k_tile / 2) * 16, token_begin, &shared.full[stage]); const int gate_scale_row = ((pair_begin / 128) * Geometry::kScaleTilesPerRow + @@ -117,9 +115,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], &descriptors->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], &descriptors->b_scales, 0, up_scale_row, &shared.full[stage]); } } diff --git a/src/product/load_progress/load_progress.cpp b/src/product/load_progress/load_progress.cpp index 2617b825b9..40d28bb6cb 100644 --- a/src/product/load_progress/load_progress.cpp +++ b/src/product/load_progress/load_progress.cpp @@ -1,10 +1,15 @@ #include "product/load_progress/load_progress.h" +#ifdef _WIN32 +#include +#else #include +#endif #include #include #include +#include #include #include #include @@ -60,7 +65,14 @@ std::string format_line(std::string_view phase, std::uint64_t done, std::uint64_ } // namespace LoadProgressRendererOptions stderr_load_progress_options() noexcept { - if (::isatty(STDERR_FILENO) == 1) { + const bool is_interactive = []() { +#ifdef _WIN32 + return _isatty(_fileno(stderr)) == 1; +#else + return ::isatty(STDERR_FILENO) == 1; +#endif + }(); + if (is_interactive) { return LoadProgressRendererOptions{ .mode = LoadProgressOutputMode::Interactive, .min_refresh_interval = std::chrono::milliseconds(200), diff --git a/src/product/media_acquire/acquire.cpp b/src/product/media_acquire/acquire.cpp index 1f03ac9c64..bf9d17c015 100644 --- a/src/product/media_acquire/acquire.cpp +++ b/src/product/media_acquire/acquire.cpp @@ -2,9 +2,14 @@ #include +#ifdef _WIN32 +#include +#include +#else #include #include #include +#endif #include #include @@ -93,17 +98,21 @@ bool private_ipv4(std::uint32_t address) { bool private_address(const sockaddr* address) { if (address->sa_family == AF_INET) { - return private_ipv4(reinterpret_cast(address)->sin_addr.s_addr); + std::uint32_t v4 = 0; + std::memcpy(&v4, &reinterpret_cast(address)->sin_addr, sizeof(v4)); + return private_ipv4(v4); } if (address->sa_family != AF_INET6) { return true; } const in6_addr& a = reinterpret_cast(address)->sin6_addr; + std::array raw{}; + std::memcpy(raw.data(), &a, sizeof(raw)); if (IN6_IS_ADDR_UNSPECIFIED(&a) || IN6_IS_ADDR_LOOPBACK(&a) || IN6_IS_ADDR_LINKLOCAL(&a) || - IN6_IS_ADDR_MULTICAST(&a) || (a.s6_addr[0] & 0xfeU) == 0xfcU) { + IN6_IS_ADDR_MULTICAST(&a) || (raw[0] & 0xfeU) == 0xfcU) { return true; } if (IN6_IS_ADDR_V4MAPPED(&a)) { std::uint32_t v4 = 0; - std::memcpy(&v4, &a.s6_addr[12], sizeof(v4)); + std::memcpy(&v4, raw.data() + 12, sizeof(v4)); return private_ipv4(v4); } return false; @@ -203,6 +212,12 @@ std::vector fetch_url(std::string url, const Policy& policy) { if (!policy.allow_remote) { throw std::invalid_argument("remote media URLs are disabled"); } static std::once_flag init; std::call_once(init, [] { +#ifdef _WIN32 + WSADATA wsa{}; + if (WSAStartup(MAKEWORD(2, 2), &wsa) != 0) { + throw std::runtime_error("failed to initialize Winsock"); + } +#endif if (curl_global_init(CURL_GLOBAL_DEFAULT) != CURLE_OK) { throw std::runtime_error("failed to initialize libcurl"); } @@ -287,7 +302,12 @@ 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 (ec || relative.empty() || relative.native().starts_with("..")) { + // A leading ".." component means the path escapes the root. Using the path + // iterator (rather than native().starts_with("..")) stays portable: on + // Windows native() is a wide string and would not match the narrow literal. + const bool escapes_root = + relative.begin() != relative.end() && *relative.begin() == ".."; + if (ec || relative.empty() || escapes_root) { throw std::invalid_argument("media path is outside configured media root"); } } diff --git a/src/serve/console_log.cpp b/src/serve/console_log.cpp index 7c58004793..7056c39d50 100644 --- a/src/serve/console_log.cpp +++ b/src/serve/console_log.cpp @@ -38,7 +38,11 @@ std::string format_console_log_prefix(std::chrono::system_clock::time_point time 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(&local, &wall_seconds); +#else localtime_r(&wall_seconds, &local); +#endif std::ostringstream out; out << '[' << std::put_time(&local, "%Y-%m-%d %H:%M:%S") << '.' << std::setfill('0') diff --git a/src/serve/request_log.cpp b/src/serve/request_log.cpp index b2dd984b69..57b2d393f2 100644 --- a/src/serve/request_log.cpp +++ b/src/serve/request_log.cpp @@ -15,13 +15,27 @@ #include #include +#ifdef _WIN32 +#define NOMINMAX +#define WIN32_LEAN_AND_MEAN +#include +#else #include +#endif namespace ninfer::serve { namespace { using Json = nlohmann::json; +std::uint32_t process_id() { +#ifdef _WIN32 + return GetCurrentProcessId(); +#else + return static_cast(::getpid()); +#endif +} + std::uint64_t unix_time_ms() { const auto now = std::chrono::system_clock::now().time_since_epoch(); return static_cast( @@ -31,7 +45,7 @@ std::uint64_t unix_time_ms() { std::string new_server_instance_id() { const auto now = std::chrono::system_clock::now().time_since_epoch(); const auto micros = std::chrono::duration_cast(now).count(); - return "serve-" + std::to_string(static_cast(::getpid())) + '-' + + return "serve-" + std::to_string(static_cast(process_id())) + '-' + std::to_string(micros); } diff --git a/src/targets/qwen3_6/impl/runtime/api_impl.h b/src/targets/qwen3_6/impl/runtime/api_impl.h index 0eadcde9bc..03797e9bf5 100644 --- a/src/targets/qwen3_6/impl/runtime/api_impl.h +++ b/src/targets/qwen3_6/impl/runtime/api_impl.h @@ -17,8 +17,15 @@ SequencePlan::SequencePlan( std::unique_ptr> impl) noexcept : impl_(std::move(impl)) {} +// MSVC does not emit an out-of-line '= default' explicit specialization unless +// it is ODR-used in this translation unit; another TU (the engine's +// std::optional usage) move-constructs these, so the move constructors must be +// emitted here or they are unresolved at link time on Windows. GCC/Clang emit +// '= default' unconditionally, which is why this only breaks the MSVC build. +// An explicit (still noexcept, still just moving the unique_ptr member) body is +// always emitted. The move-assignment and destructor keep '= default'. template <> -SequencePlan::SequencePlan(SequencePlan&&) noexcept = default; +SequencePlan::SequencePlan(SequencePlan&& other) noexcept : impl_(std::move(other.impl_)) {} template <> SequencePlan& SequencePlan::operator=(SequencePlan&&) noexcept = default; template <> @@ -85,7 +92,8 @@ 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; template <> @@ -102,7 +110,7 @@ RequestPlan::RequestPlan(std::unique_ptr -RequestPlan::RequestPlan(RequestPlan&&) noexcept = default; +RequestPlan::RequestPlan(RequestPlan&& other) noexcept : impl_(std::move(other.impl_)) {} template <> RequestPlan& RequestPlan::operator=(RequestPlan&&) noexcept = default; template <> diff --git a/tests/test_request_log.cpp b/tests/test_request_log.cpp index d557ace416..8e241f7d00 100644 --- a/tests/test_request_log.cpp +++ b/tests/test_request_log.cpp @@ -12,13 +12,27 @@ #include #include +#ifdef _WIN32 +#define NOMINMAX +#define WIN32_LEAN_AND_MEAN +#include +#else #include +#endif namespace { using namespace ninfer::serve; using Json = nlohmann::json; +std::uint32_t process_id() { +#ifdef _WIN32 + return GetCurrentProcessId(); +#else + return static_cast(::getpid()); +#endif +} + int check(bool condition, const char* message) { if (condition) { return 0; } std::cerr << message << '\n'; @@ -347,7 +361,7 @@ int main() { const std::filesystem::path log_path = std::filesystem::temp_directory_path() / - ("ninfer-request-log-test-" + std::to_string(static_cast(::getpid())) + + ("ninfer-request-log-test-" + std::to_string(static_cast(process_id())) + ".jsonl"); std::filesystem::remove(log_path); { From 1c41cf7e316dd1bc2b2004b6165c91ccabd86a48 Mon Sep 17 00:00:00 2001 From: natpate <80576637+natpate@users.noreply.github.com> Date: Sun, 16 Aug 2026 22:25:38 -0500 Subject: [PATCH 02/20] serve: accept stock llama.cpp WebUI dialect (API-compatible with tools/ui) - parse chat_template_kwargs.enable_thinking (top-level or in kwargs; conflict between the two is a 400) - missing/empty 'model' defaults to the loaded artifact - max_tokens <= 0 means server default (WebUI sends -1) - /v1/models entries carry status {value: loaded} - GET /props stub (n_ctx, n_predict, speculative, modalities, chat_template enable_thinking marker) for role/thinking detection - CORS: OPTIONS handler echoes preflight Access-Control-Request-Headers (WebUI sends custom x-conversation-id), keeps Authorization/Content-Type floor Verified end-to-end against the static tools/ui build: model list, thinking toggle, non-stream + SSE streaming with reasoning_content. --- src/serve/http_server.cpp | 172 ++++++++++++++++++++++++++++++++++- src/serve/http_server.h | 1 + src/serve/openai_schema.cpp | 125 ++++++++++++++++++------- src/serve/openai_schema.h | 23 ++++- src/serve/request.h | 2 +- tests/test_openai_schema.cpp | 161 +++++++++++++++++++++++++++++++- 6 files changed, 445 insertions(+), 39 deletions(-) diff --git a/src/serve/http_server.cpp b/src/serve/http_server.cpp index f99e6d9021..27a2bc9ce9 100644 --- a/src/serve/http_server.cpp +++ b/src/serve/http_server.cpp @@ -9,6 +9,7 @@ #include #include +#include #include #include #include @@ -17,6 +18,7 @@ #include #include #include +#include namespace ninfer::serve { namespace { @@ -54,6 +56,7 @@ void write_error(httplib::Response& res, const ApiError& error) { res.set_content(make_error_body(error), "application/json"); } + // Anthropic-shaped error body ({"type":"error","error":{...}}), used by the // /v1/messages endpoints so Claude clients see the error format they expect. void write_messages_error(httplib::Response& res, const ApiError& error) { @@ -220,8 +223,67 @@ void HttpServer::register_routes() { {"Access-Control-Allow-Methods", "GET, POST, DELETE, OPTIONS"}}); // CORS preflight: browsers send OPTIONS with no credentials before the real // request; answer it without auth so the actual GET/POST can carry the key. - server_.Options(R"(.*)", - [](const httplib::Request&, httplib::Response& res) { res.status = 204; }); + // Echo the preflight's requested headers (Access-Control-Request-Headers) + // so clients using custom headers (e.g. llama.cpp webui's x-conversation-id) + // pass the browser's CORS check. The static default above remains the + // floor for direct (non-browser) requests. + server_.Options(R"(.*)", [](const httplib::Request& req, httplib::Response& res) { + res.status = 204; + auto it = req.get_header_value("Access-Control-Request-Headers"); + if (!it.empty()) { + std::string joined; + std::vector seen; + std::string part; + const std::string& raw = it; + for (size_t i = 0; i <= raw.size(); i++) { + const char c = i < raw.size() ? raw[i] : ','; + if (c == ',') { + size_t b = part.find_first_not_of(" \t"); + if (b == std::string::npos) { + part.clear(); + continue; + } + size_t e = part.find_last_not_of(" \t") + 1; + part = part.substr(b, e - b); + if (!part.empty()) { + std::string key = part; + for (auto& ch : key) { + if (ch >= 'A' && ch <= 'Z') { + ch = char(ch - 'A' + 'a'); + } + } + const bool dup = std::any_of(seen.begin(), seen.end(), + [&](const std::string& s) { return s == key; }); + if (!dup) { + if (!joined.empty()) { + joined += ", "; + } + joined += part; + seen.push_back(std::move(key)); + } + } + part.clear(); + } else { + part += c; + } + } + if (!joined.empty()) { + // Build the full header set explicitly: httplib's set_header + // uses emplace and cannot replace the statically seeded + // Access-Control-Allow-Headers default. + httplib::Headers headers; + for (const auto& h : res.headers) { + // The seeded default uses exactly this key spelling. + if (h.first == "Access-Control-Allow-Headers") { + headers.emplace(h.first, joined); + } else { + headers.emplace(h.first, h.second); + } + } + res.headers = std::move(headers); + } + } + }); } server_.set_pre_routing_handler([this](const httplib::Request& req, httplib::Response& res) { @@ -279,6 +341,9 @@ void HttpServer::register_routes() { [this](const httplib::Request& req, httplib::Response& res) { handle_chat_completions(req, res); }); + server_.Get("/props", [this](const httplib::Request& req, httplib::Response& res) { + handle_props(req, res); + }); server_.Post("/v1/responses", [this](const httplib::Request& req, httplib::Response& res) { handle_responses(req, res); }); @@ -315,6 +380,103 @@ void HttpServer::register_routes() { }); } +// llama.cpp webui dialect: /props is the client's server introspection endpoint +// (role detection, context size, default params, thinking-capability probe, api +// key validation). NInfer has no llama.cpp server behind it, so serve a faithful +// stub derived from the process configuration. Only process-level overrides are +// reported as parameter values; everything else stays at neutral zeros so client +// side defaults never swallow a user-set request parameter. +nlohmann::json make_props_stub(const ServeOptions& options, const std::string& model_id) { + (void)model_id; + const auto& ov = options.sampling_overrides; + nlohmann::json params = nlohmann::json::object(); + params["n_predict"] = options.default_max_tokens; + params["seed"] = ov.seed ? static_cast(*ov.seed) : 0; + params["temperature"] = ov.temperature ? static_cast(*ov.temperature) : 0; + params["dynatemp_range"] = 0; + params["dynatemp_exponent"] = 0; + params["top_k"] = ov.top_k ? *ov.top_k : 0; + params["top_p"] = ov.top_p ? static_cast(*ov.top_p) : 0; + params["min_p"] = ov.min_p ? static_cast(*ov.min_p) : 0; + params["top_n_sigma"] = 0; + params["xtc_probability"] = 0; + params["xtc_threshold"] = 0; + params["typ_p"] = 0; + params["repeat_last_n"] = 0; + params["repeat_penalty"] = 0; + params["presence_penalty"] = + ov.presence_penalty ? static_cast(*ov.presence_penalty) : 0; + params["frequency_penalty"] = + ov.frequency_penalty ? static_cast(*ov.frequency_penalty) : 0; + params["dry_multiplier"] = 0; + params["dry_base"] = 0; + params["dry_allowed_length"] = 0; + params["dry_penalty_last_n"] = 0; + params["dry_sequence_breakers"] = nlohmann::json::array(); + params["mirostat"] = 0; + params["mirostat_tau"] = 0; + params["mirostat_eta"] = 0; + params["stop"] = nlohmann::json::array(); + params["max_tokens"] = options.default_max_tokens; + params["n_keep"] = 0; + params["n_discard"] = 0; + params["ignore_eos"] = false; + params["stream"] = false; + params["logit_bias"] = nlohmann::json::array(); + params["n_probs"] = 0; + params["min_keep"] = 0; + params["grammar"] = ""; + params["grammar_lazy"] = false; + params["grammar_triggers"] = nlohmann::json::array(); + params["preserved_tokens"] = nlohmann::json::array(); + params["chat_format"] = ""; + params["reasoning_format"] = ""; + params["reasoning_in_content"] = false; + params["generation_prompt"] = ""; + params["samplers"] = nlohmann::json::array(); + params["backend_sampling"] = false; + params["speculative.n_max"] = 0; + params["speculative.n_min"] = 0; + params["speculative.p_min"] = 0.0; + params["timings_per_token"] = false; + params["post_sampling_probs"] = false; + params["lora"] = nlohmann::json::array(); + + nlohmann::json props = nlohmann::json::object(); + props["default_generation_settings"] = { + {"id", 0}, + {"id_task", 0}, + {"n_ctx", static_cast(options.max_context)}, + {"speculative", options.speculative.backend != SpeculativeBackend::None}, + {"is_processing", false}, + {"params", params}, + {"prompt", ""}, + {"next_token", + {{"has_next_token", false}, + {"has_new_line", false}, + {"n_remain", 0}, + {"n_decoded", 0}, + {"stopping_word", ""}}}, + }; + props["total_slots"] = 1; + props["model_path"] = options.artifact_path; + props["role"] = "model"; + props["modalities"] = {{"vision", options.enable_vision}, {"audio", false}, {"video", false}}; + // Capability marker only: clients that probe the chat template (e.g. the + // webui's thinking-support heuristic) need `enable_thinking` to appear; the + // real template is embedded in the loaded artifact. + props["chat_template"] = + "{# ninfer-serve: capability marker; the real chat template is embedded in " + "the loaded artifact #}\n" + "{%- if enable_thinking is defined %}\n" + " {%- set thinking = enable_thinking %}\n" + "{% endif %}"; + props["bos_token"] = ""; + props["eos_token"] = ""; + props["build_info"] = "ninfer-serve"; + return props; +} + void HttpServer::handle_models(const httplib::Request&, httplib::Response& res) const { res.set_content(make_models_list(public_model_id_, unix_time_now()), "application/json"); } @@ -333,6 +495,10 @@ void HttpServer::handle_model(const httplib::Request& req, httplib::Response& re res.set_content(make_model_object(public_model_id_, unix_time_now()), "application/json"); } +void HttpServer::handle_props(const httplib::Request&, httplib::Response& res) const { + res.set_content(make_props_stub(options_, public_model_id_).dump(), "application/json"); +} + void HttpServer::handle_chat_completions(const httplib::Request& req, httplib::Response& res) { nlohmann::json body; try { @@ -349,7 +515,7 @@ void HttpServer::handle_chat_completions(const httplib::Request& req, httplib::R try { RequestLimits limits; limits.default_max_tokens = options_.default_max_tokens; - request = parse_chat_completion_request(body, limits); + request = parse_chat_completion_request(body, limits, public_model_id_); if (request.model != public_model_id_) { ApiError error; error.status = 404; diff --git a/src/serve/http_server.h b/src/serve/http_server.h index c0557e9d27..1461c40f80 100644 --- a/src/serve/http_server.h +++ b/src/serve/http_server.h @@ -51,6 +51,7 @@ class HttpServer { void handle_response_compact(const httplib::Request& req, httplib::Response& res); void handle_models(const httplib::Request& req, httplib::Response& res) const; void handle_model(const httplib::Request& req, httplib::Response& res) const; + void handle_props(const httplib::Request& req, httplib::Response& res) const; // The process-wide console logger serializes lines from request and reporter threads. void log_line(const std::string& line); diff --git a/src/serve/openai_schema.cpp b/src/serve/openai_schema.cpp index 6975ccf7d8..b8f734f4e1 100644 --- a/src/serve/openai_schema.cpp +++ b/src/serve/openai_schema.cpp @@ -495,7 +495,8 @@ std::optional parse_openai_preserve_thinking(const Json& body) { bad_request("chat_template_kwargs must be an object", "chat_template_kwargs"); } for (auto it = kwargs.begin(); it != kwargs.end(); ++it) { - if (it.key() != "preserve_thinking" && !it.value().is_null()) { + if (it.key() != "preserve_thinking" && it.key() != "enable_thinking" && + !it.value().is_null()) { bad_request("chat_template_kwargs." + it.key() + " is not supported", "chat_template_kwargs", "chat_template_option_not_supported"); } @@ -516,32 +517,83 @@ std::optional parse_openai_preserve_thinking(const Json& body) { return template_value ? template_value : top_level; } -void parse_openai_reasoning_effort(const Json& body, GenerationRequest& out) { - if (!body.contains("reasoning_effort") || body.at("reasoning_effort").is_null()) { return; } - if (!body.at("reasoning_effort").is_string()) { - bad_request("reasoning_effort must be a string or null", "reasoning_effort"); +std::optional parse_chat_enable_thinking(const Json& body) { + // Public top-level extension, or the llama.cpp webui dialect where the switch + // lives under chat_template_kwargs. Nulls stay unset (server default wins). + std::optional top_level; + if (body.contains("enable_thinking") && !body.at("enable_thinking").is_null()) { + if (!body.at("enable_thinking").is_boolean()) { + bad_request("enable_thinking must be a boolean or null", "enable_thinking"); + } + top_level = body.at("enable_thinking").get(); } - const std::string value = body.at("reasoning_effort").get(); - const std::optional effort = parse_requested_reasoning_effort(value); - if (!effort) { - bad_request("reasoning_effort must be one of none, minimal, low, medium, high, xhigh, or " - "max", - "reasoning_effort"); + std::optional kwargs_value; + if (body.contains("chat_template_kwargs") && body.at("chat_template_kwargs").is_object()) { + const Json& kwargs = body.at("chat_template_kwargs"); + if (kwargs.contains("enable_thinking") && !kwargs.at("enable_thinking").is_null()) { + if (!kwargs.at("enable_thinking").is_boolean()) { + bad_request("chat_template_kwargs.enable_thinking must be a boolean or null", + "chat_template_kwargs"); + } + kwargs_value = kwargs.at("enable_thinking").get(); + } } - out.reasoning_effort = *effort; - out.reasoning_effort_param = "reasoning_effort"; + if (top_level && kwargs_value && *top_level != *kwargs_value) { + bad_request("conflicting enable_thinking values", "enable_thinking", + "conflicting_template_option"); + } + return top_level ? top_level : kwargs_value; } -GenerationRequest parse_chat_completion_request(const Json& body, const RequestLimits& limits) { +void parse_openai_chat_thinking(const Json& body, std::optional* enable_thinking, + std::optional* reasoning_effort, + std::string* reasoning_effort_param, + const std::string& conflict_param) { + *enable_thinking = parse_chat_enable_thinking(body); + + std::optional effort; + if (body.contains("reasoning_effort") && !body.at("reasoning_effort").is_null()) { + if (!body.at("reasoning_effort").is_string()) { + bad_request("reasoning_effort must be a string or null", "reasoning_effort"); + } + const std::string value = body.at("reasoning_effort").get(); + const std::optional parsed = parse_requested_reasoning_effort(value); + if (!parsed) { + bad_request("reasoning_effort must be one of none, minimal, low, medium, high, xhigh, " + "or max", + "reasoning_effort"); + } + effort = *parsed; + } + + if (effort) { + *reasoning_effort = *effort; + *reasoning_effort_param = "reasoning_effort"; + if (enable_thinking->has_value() && + *enable_thinking != (*effort != RequestedReasoningEffort::None)) { + bad_request("reasoning effort conflicts with " + conflict_param, + "reasoning_effort", "conflicting_template_option"); + } + } +} + +GenerationRequest parse_chat_completion_request(const Json& body, const RequestLimits& limits, + const std::string& default_model_id) { require_object(body); reject_unsupported_features(body); GenerationRequest out; - if (!body.contains("model") || !body.at("model").is_string() || - body.at("model").get().empty()) { - bad_request("missing required field: model", "model"); + if (body.contains("model") && body.at("model").is_string() && + !body.at("model").get().empty()) { + out.model = body.at("model").get(); + } else { + // Single-model clients (e.g. the llama.cpp webui) omit `model`; serve + // against the loaded artifact instead of rejecting the request. + if (default_model_id.empty()) { + bad_request("missing required field: model", "model"); + } + out.model = default_model_id; } - out.model = body.at("model").get(); parse_tools(body, out); parse_tool_choice(body, out); @@ -553,18 +605,22 @@ GenerationRequest parse_chat_completion_request(const Json& body, const RequestL if (body.contains("stream_options") && body.at("stream_options").is_object()) { out.include_usage = get_bool(body.at("stream_options"), "include_usage", false); } - if (body.contains("enable_thinking") && !body.at("enable_thinking").is_null()) { - out.enable_thinking = get_bool(body, "enable_thinking", false); - } - parse_openai_reasoning_effort(body, out); + parse_openai_chat_thinking(body, &out.enable_thinking, &out.reasoning_effort, + &out.reasoning_effort_param, "enable_thinking"); out.preserve_thinking = parse_openai_preserve_thinking(body); std::optional max_tokens = get_int(body, "max_completion_tokens"); if (!max_tokens) { max_tokens = get_int(body, "max_tokens"); } if (max_tokens) { - if (*max_tokens <= 0) { bad_request("max_tokens must be positive", "max_tokens"); } - out.max_tokens = *max_tokens; - out.max_tokens_set = true; + // Non-positive (llama.cpp `-1` = unlimited) falls back to the server + // default, which the Engine clamps to its effective context capacity. + if (*max_tokens <= 0) { + out.max_tokens = limits.default_max_tokens; + out.max_tokens_set = false; + } else { + out.max_tokens = *max_tokens; + out.max_tokens_set = true; + } } else { out.max_tokens = limits.default_max_tokens; out.max_tokens_set = false; @@ -682,17 +738,24 @@ std::string make_chat_chunk_usage(const std::string& id, const std::string& mode std::string sse_done() { return "data: [DONE]\n\n"; } std::string make_models_list(const std::string& model_id, std::int64_t created) { + // `status` is a llama.cpp webui extension: the client reads status.value to + // decide whether a model is loaded. A single loaded artifact is always loaded. const Json payload = {{"object", "list"}, - {"data", Json::array({Json{{"id", model_id}, - {"object", "model"}, - {"created", created}, - {"owned_by", "ninfer"}}})}}; + {"data", + Json::array({Json{{"id", model_id}, + {"object", "model"}, + {"created", created}, + {"owned_by", "ninfer"}, + {"status", Json{{"value", "loaded"}}}}})}}; return payload.dump(); } std::string make_model_object(const std::string& model_id, std::int64_t created) { - const Json payload = { - {"id", model_id}, {"object", "model"}, {"created", created}, {"owned_by", "ninfer"}}; + const Json payload = {{"id", model_id}, + {"object", "model"}, + {"created", created}, + {"owned_by", "ninfer"}, + {"status", Json{{"value", "loaded"}}}}; return payload.dump(); } diff --git a/src/serve/openai_schema.h b/src/serve/openai_schema.h index 57f430abf2..0554ac2738 100644 --- a/src/serve/openai_schema.h +++ b/src/serve/openai_schema.h @@ -5,6 +5,7 @@ // This layer knows nothing about the engine; it only speaks the OpenAI schema. #include "serve/request.h" +#include "serve/serve_options.h" #include @@ -20,11 +21,31 @@ namespace ninfer::serve { // Parse an already-decoded JSON body into a GenerationRequest. Throws ApiException // on malformed or unsupported requests (n>1, tools, non-text response_format, ...). +// `default_model_id` fills in when the request omits `model` (clients such as the +// llama.cpp webui run a single loaded model and do not send it). GenerationRequest parse_chat_completion_request(const nlohmann::json& body, - const RequestLimits& limits); + const RequestLimits& limits, + const std::string& default_model_id = {}); + +// llama.cpp webui dialect, mapped onto the public thinking controls: +// chat_template_kwargs.enable_thinking and reasoning_effort=low|medium are +// accepted instead of being rejected. `conflict_param` names the source field of +// any clash with `reasoning_effort` ("enable_thinking" or "reasoning_effort"). +void parse_openai_chat_thinking(const nlohmann::json& body, + std::optional* enable_thinking, + std::optional* reasoning_effort, + std::string* reasoning_effort_param, + const std::string& conflict_param); std::optional parse_openai_preserve_thinking(const nlohmann::json& body); +// llama.cpp webui dialect: /props payload derived from the process configuration. +// The webui probes it for role, context size, default params, and the chat +// template; only process-level overrides are reported, everything else is +// neutral so client defaults never swallow a user-set request parameter. +nlohmann::json make_props_stub(const ServeOptions& options, + const std::string& model_id); + // Non-streaming chat completion response body (JSON string). When `reasoning` is // non-empty it is attached as `message.reasoning_content` (the DeepSeek/vLLM-style // convention consumed by Chatbox, Open WebUI, etc.), leaving `content` = answer. diff --git a/src/serve/request.h b/src/serve/request.h index 3088a6b849..5e3f718b32 100644 --- a/src/serve/request.h +++ b/src/serve/request.h @@ -171,7 +171,7 @@ struct GenerationRequest { std::size_t tool_name_max_length = 64; ToolChoice tool_choice; std::vector stop_strings; - int max_tokens = 0; // 0 => use server default + int max_tokens = 0; // 0 => use server default; set via max_tokens_set when client pinned a value bool max_tokens_set = false; bool stream = false; bool include_usage = false; diff --git a/tests/test_openai_schema.cpp b/tests/test_openai_schema.cpp index 6fef3ecf0f..13b95c1049 100644 --- a/tests/test_openai_schema.cpp +++ b/tests/test_openai_schema.cpp @@ -5,6 +5,7 @@ #include "serve/openai_schema.h" #include "serve/request.h" +#include "serve/serve_options.h" #include "serve/translate.h" #include @@ -52,6 +53,9 @@ RequestLimits default_limits() { return limits; } +// Public model id used by the webui-dialect tests (they send it explicitly). +const char* webui_model() { return "qwen3.6-27b"; } + ServeOptions default_server() { return ServeOptions{}; } ninfer::PromptCapabilities effort_capabilities() { @@ -371,9 +375,17 @@ int test_reject_unsupported() { failures += check(text_ok, "text response_format accepted"); Json no_model = {{"messages", Json::array({Json{{"role", "user"}, {"content", "hi"}}})}}; - failures += - check(throws_api([&] { (void)parse_chat_completion_request(no_model, default_limits()); }), - "missing model rejected"); + failures += check(parse_chat_completion_request(no_model, default_limits(), "qwen3.6-27b").model == + "qwen3.6-27b", + "omitted model filled from default model id"); + failures += check( + throws_api([&] { (void)parse_chat_completion_request(no_model, default_limits()); }), + "missing model rejected without a default model id"); + Json empty_model = {{"model", ""}, + {"messages", Json::array({Json{{"role", "user"}, {"content", "hi"}}})}}; + failures += check(parse_chat_completion_request(empty_model, default_limits(), "qwen3.6-27b") + .model == "qwen3.6-27b", + "empty model falls back to default model id"); Json function_role = { {"model", "m"}, {"messages", Json::array({Json{{"role", "function"}, {"content", "x"}}})}}; @@ -675,9 +687,13 @@ int test_models_and_error() { failures += check(list.at("data").at(0).at("object") == "model", "models list entry object"); failures += check(list.at("data").at(0).at("owned_by") == "ninfer", "models list owner"); + failures += check(list.at("data").at(0).at("status").at("value") == "loaded", + "models list entry reports loaded status"); const Json one = Json::parse(make_model_object("qwen3.6-27b", 1)); failures += check(one.at("id") == "qwen3.6-27b" && one.at("object") == "model", "model object"); failures += check(one.at("owned_by") == "ninfer", "model owner"); + failures += check(one.at("status").at("value") == "loaded", + "model object reports loaded status"); ApiError error; error.status = 400; @@ -706,6 +722,143 @@ int test_finish_reason_wire() { } // namespace +int test_llama_webui_dialect() { + const Json base = { + {"model", "qwen3.6-27b"}, + {"messages", Json::array({Json{{"role", "user"}, {"content", "hello"}}})}, + }; + int failures = 0; + + // chat_template_kwargs.enable_thinking (always sent by the webui) + Json et = base; + et["chat_template_kwargs"] = Json{{"enable_thinking", true}}; + const GenerationRequest et_request = + parse_chat_completion_request(et, default_limits(), webui_model()); + failures += check(et_request.enable_thinking == true, "kwargs enable_thinking parsed"); + failures += check(translate(et_request).options.enable_thinking, + "kwargs enable_thinking reached prompt"); + + Json et_false = base; + et_false["chat_template_kwargs"] = Json{{"enable_thinking", false}}; + failures += check( + parse_chat_completion_request(et_false, default_limits(), webui_model()).enable_thinking == + false, + "kwargs enable_thinking=false parsed"); + + Json both = base; + both["enable_thinking"] = true; + both["chat_template_kwargs"] = Json{{"enable_thinking", true}}; + failures += check( + parse_chat_completion_request(both, default_limits(), webui_model()).enable_thinking == + true, + "matching top-level and kwargs enable_thinking accepted"); + Json et_conflict = both; + et_conflict["enable_thinking"] = false; + failures += check(api_code([&] { + (void)parse_chat_completion_request(et_conflict, default_limits(), + webui_model()); + }) == "conflicting_template_option", + "conflicting enable_thinking values rejected"); + Json et_bad = base; + et_bad["chat_template_kwargs"] = Json{{"enable_thinking", "yes"}}; + failures += check( + throws_api([&] { + (void)parse_chat_completion_request(et_bad, default_limits(), webui_model()); + }), + "non-boolean kwargs enable_thinking rejected"); + + // webui reasoning_effort=low|medium accepted; none is a public value + Json low = base; + low["reasoning_effort"] = "low"; + const GenerationRequest low_request = + parse_chat_completion_request(low, default_limits(), webui_model()); + failures += check(low_request.reasoning_effort == RequestedReasoningEffort::Low, + "webui reasoning_effort=low accepted"); + const ninfer::PromptInput low_prompt = + translate(parse_chat_completion_request(low, default_limits(), webui_model())); + failures += check(low_prompt.options.enable_thinking && + low_prompt.options.reasoning_effort == ninfer::ReasoningEffort::Low, + "webui low effort reached prompt"); + + // conflicting enable_thinking=false vs effort low rejected at parse time + Json c = low; + c["chat_template_kwargs"] = Json{{"enable_thinking", false}}; + failures += check(api_code([&] { + (void)parse_chat_completion_request(c, default_limits(), webui_model()); + }) == "conflicting_template_option", + "conflicting enable_thinking and reasoning_effort rejected"); + + // high/max are public protocol values: parse ok, template capability decides + Json high = base; + high["reasoning_effort"] = "high"; + failures += check( + api_code([&] { + (void)resolve_prompt_semantics( + parse_chat_completion_request(high, default_limits(), webui_model()), + default_server(), effort_capabilities()); + }) == "reasoning_effort_not_supported", + "webui high effort rejected by template capability"); + + // max_tokens=-1 (webui "unlimited") falls back to the server default + Json unlimited = base; + unlimited["max_tokens"] = -1; + const GenerationRequest unlimited_request = + parse_chat_completion_request(unlimited, default_limits(), webui_model()); + failures += check(unlimited_request.max_tokens == 512 && !unlimited_request.max_tokens_set, + "max_tokens=-1 falls back to server default"); + + // full webui-shaped body parses cleanly + Json webui = base; + webui["stream"] = true; + webui["return_progress"] = true; + webui["sse_ping_interval"] = 1; + webui["reasoning_format"] = "auto"; + webui["chat_template_kwargs"] = Json{{"enable_thinking", true}}; + webui["reasoning_control"] = true; + webui["thinking_budget_tokens"] = 2048; + webui["max_tokens"] = -1; + webui["temperature"] = 0.7; + webui["timings_per_token"] = true; + failures += check(!throws_api([&] { + (void)parse_chat_completion_request(webui, default_limits(), + webui_model()); + }), + "full webui-shaped body accepted"); + return failures; +} + +int test_props_stub() { + int failures = 0; + ServeOptions options; + options.artifact_path = "models/qwen3_6_27b.ninfer"; + options.max_context = 16384; + options.default_max_tokens = 4096; + options.enable_vision = true; + options.speculative.backend = SpeculativeBackend::Mtp; + options.sampling_overrides.temperature = 1.0F; + options.sampling_overrides.top_k = 20; + + const Json props = make_props_stub(options, "qwen3.6-27b"); + failures += check(props.at("role") == "model", "props role is model"); + failures += check(props.at("modalities").at("vision") == true, "props vision follows --vision"); + failures += check(props.at("modalities").at("audio") == false, "props audio off"); + const Json params = props.at("default_generation_settings").at("params"); + failures += check(params.at("n_ctx") == 16384, "props n_ctx from --max-context"); + failures += check(params.at("n_predict") == 4096, "props n_predict from default max tokens"); + failures += check(params.at("temperature") == 1.0, "props temperature override reported"); + failures += check(params.at("top_k") == 20, "props top_k override reported"); + failures += check(params.at("presence_penalty") == 0, "unreported param stays neutral zero"); + failures += check(params.at("dry_base") == 0, "dry params stay neutral zero"); + failures += check(props.at("default_generation_settings").at("speculative") == true, + "props speculative follows --spec"); + failures += check(props.at("default_generation_settings").at("is_processing") == false, + "props is_processing false"); + failures += check(props.at("chat_template").get().find("enable_thinking") != + std::string::npos, + "props chat_template exposes enable_thinking for capability probes"); + return failures; +} + int main() { int failures = 0; failures += test_parse_string_content(); @@ -724,6 +877,8 @@ int main() { failures += test_chunk_serialization(); failures += test_tool_chunk_serialization(); failures += test_models_and_error(); + failures += test_llama_webui_dialect(); + failures += test_props_stub(); failures += test_finish_reason_wire(); if (failures == 0) { std::cout << "ok\n"; } return failures == 0 ? 0 : 1; From 1a496d7ddb0946a4f6b202c61bac2b94874fdbb6 Mon Sep 17 00:00:00 2001 From: natpate <80576637+natpate@users.noreply.github.com> Date: Sun, 16 Aug 2026 23:33:34 -0500 Subject: [PATCH 03/20] serve: add --webui / --webui-dir to serve the stock llama.cpp WebUI in-process - --webui: at startup, auto-sync the prebuilt webui from the ggml-org/llama-ui Hugging Face bucket (WinHTTP, version-marker check, staged atomic swap) so the UI stays current without a rebuild; standalone mirror script: webui-update.bat - --webui-dir DIR: serve an existing built static dir - static mount at / via httplib set_mount_point; asset MIME pins - SPA fallback in the error handler (404-only, GET/HEAD, dotless, non-API paths) so it can never shadow a real asset or route; /slots, /tools and other probe paths keep honest 404s - API-key auth now gates API paths only, so the UI shell and static assets load freely (llama-server parity) - build wiring: webui_update.cpp in src/CMakeLists.txt + winhttp (MSBuild tree updated on-disk, gitignored) --- apps/serve/main.cpp | 21 +- src/CMakeLists.txt | 5 +- src/serve/http_server.cpp | 87 +++++++- src/serve/http_server.h | 6 + src/serve/serve_options.cpp | 12 ++ src/serve/serve_options.h | 4 + src/serve/webui_update.cpp | 383 ++++++++++++++++++++++++++++++++++++ src/serve/webui_update.h | 25 +++ 8 files changed, 539 insertions(+), 4 deletions(-) create mode 100644 src/serve/webui_update.cpp create mode 100644 src/serve/webui_update.h diff --git a/apps/serve/main.cpp b/apps/serve/main.cpp index 9263e100cd..f78c750210 100644 --- a/apps/serve/main.cpp +++ b/apps/serve/main.cpp @@ -3,11 +3,13 @@ #include "serve/generation_service.h" #include "serve/http_server.h" #include "serve/serve_options.h" +#include "serve/webui_update.h" #include #include #include #include +#include #include #include #include @@ -40,12 +42,29 @@ std::string format_bytes(std::size_t bytes) { int main(int argc, char** argv) { try { - const ninfer::serve::ServeOptions options = ninfer::serve::parse_serve_options(argc, argv); + ninfer::serve::ServeOptions options = ninfer::serve::parse_serve_options(argc, argv); if (options.help_requested) { std::cout << ninfer::serve::serve_usage_text(argv[0]); return 0; } + // Resolve (and, in --webui mode, auto-download) the webui directory before + // the port is taken so a failed download aborts startup cleanly. In + // --webui-dir mode the directory is trusted to already hold a built UI; + // fail early if it does not. + if (options.webui_auto) { + options.webui_dir = + ninfer::serve::ensure_webui_available(ninfer::serve::resolve_webui_dir(options)); + } else if (!options.webui_dir.empty()) { + std::error_code ec; + const bool have_index = + std::filesystem::exists(std::filesystem::path(options.webui_dir) / "index.html", ec); + if (!std::filesystem::is_directory(options.webui_dir, ec) || !have_index) { + throw std::invalid_argument( + "--webui-dir must be a directory containing index.html: " + options.webui_dir); + } + } + using Clock = std::chrono::steady_clock; ninfer::serve::HttpServer server(options); if (!server.bind()) { diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 33167d61cc..05aa0c24eb 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -337,7 +337,8 @@ if(NINFER_BUILD_SERVE) serve/request_log.cpp serve/http_server.cpp serve/responses_http.cpp - serve/serve_options.cpp) + serve/serve_options.cpp + serve/webui_update.cpp) ninfer_internal_includes(ninfer_serve) target_include_directories(ninfer_serve PRIVATE ${PROJECT_SOURCE_DIR}/third_party/cpp-httplib) @@ -347,5 +348,7 @@ if(NINFER_BUILD_SERVE) if(WIN32) # cpp-httplib (plain HTTP) uses the Winsock API. target_link_libraries(ninfer_serve PRIVATE Ws2_32) + # webui_update fetches the stock WebUI bundle over WinHTTP. + target_link_libraries(ninfer_serve PRIVATE winhttp) endif() endif() diff --git a/src/serve/http_server.cpp b/src/serve/http_server.cpp index 27a2bc9ce9..85827771af 100644 --- a/src/serve/http_server.cpp +++ b/src/serve/http_server.cpp @@ -5,6 +5,7 @@ #include "serve/openai_schema.h" #include "serve/request_log.h" #include "serve/translate.h" +#include "serve/webui_update.h" #include @@ -12,6 +13,8 @@ #include #include #include +#include +#include #include #include #include @@ -138,9 +141,71 @@ HttpServer::HttpServer(ServeOptions options) return new httplib::ThreadPool(worker_count, queued_requests); }; server_.set_payload_max_length(options_.max_request_bytes); + if (!options_.webui_dir.empty()) { + mount_webui(options_.webui_dir); + register_webui_mime(); + } register_routes(); } +void HttpServer::mount_webui(const std::string& webui_dir) { + std::ifstream index(webui_dir + "/index.html", std::ios::binary); + if (!index) { + throw std::runtime_error("webui dir has no index.html: " + webui_dir); + } + webui_index_html_ = std::string((std::istreambuf_iterator(index)), + std::istreambuf_iterator()); + webui_serving_ = true; + if (!server_.set_mount_point("/", webui_dir)) { + throw std::runtime_error("cannot mount webui directory: " + webui_dir); + } + log_line("serving webui from " + webui_dir); +} + +void HttpServer::register_webui_mime() { + // The vendored httplib's built-in map covers the stock webui's asset types; + // pin the ones the UI's runtime loading depends on. + server_.set_file_extension_and_mimetype_mapping("js", "text/javascript"); + server_.set_file_extension_and_mimetype_mapping("css", "text/css"); + server_.set_file_extension_and_mimetype_mapping("html", "text/html"); + server_.set_file_extension_and_mimetype_mapping("json", "application/json"); + server_.set_file_extension_and_mimetype_mapping("svg", "image/svg+xml"); + server_.set_file_extension_and_mimetype_mapping("ico", "image/x-icon"); +} + +bool HttpServer::webui_spa_path(const std::string& path) const { + // The SPA fallback must only catch client-side routes: the static file handler + // already served every real asset, and API paths (/v1/...), /props, /health, + // and hashed bundle paths (_app/...) are never SPA routes. A missing _app file + // or an unknown /v1 path must keep falling through to the 404 handler. + if (path.size() < 2 || path[0] != '/') { return false; } + if (path == "/") { return false; } + if (path.rfind("/v1", 0) == 0 && (path.size() == 3 || path[3] == '/')) { return false; } + // llama.cpp server endpoints the webui probes but ninfer does not implement, plus + // ninfer's own status endpoints: keep them at their natural 404/405 so the UI + // degrades exactly as it did against a real llama-server, rather than getting the + // SPA shell back for an API call. + if (path == "/props" || path == "/health" || path == "/slots" || path == "/tools" || + path == "/v1/streams/lookup") { + return false; + } + if (path.rfind("/_app/", 0) == 0) { return false; } + if (path.find_first_of('.') != std::string::npos) { return false; } + return true; +} + +bool HttpServer::is_api_path(const std::string& path) const { + // The endpoints that require an API key. Deliberately dot-free: static assets + // (favicon.ico, app.js) and the UI shell (/, index.html, SPA routes) are never + // API paths, so the UI loads freely even when --api-key is set, exactly like + // llama-server. The webui supplies the key itself on API calls. + if (path.rfind("/v1", 0) == 0 && (path.size() == 3 || path[3] == '/')) { return true; } + if (path == "/props" || path == "/slots" || path == "/tools") { return true; } + if (path.rfind("/v1/streams/lookup", 0) == 0) { return true; } + if (path.rfind("/_app/", 0) == 0) { return false; } // served by the static mount + return false; +} + void HttpServer::log_line(const std::string& line) { write_console_log(ConsoleLogLevel::Info, line); } @@ -214,7 +279,21 @@ void HttpServer::stop_stats_reporter() { void HttpServer::register_routes() { server_.set_error_handler([this](const httplib::Request& request, httplib::Response& response) { - return handle_unrendered_http_error(options_, request, response); + const auto rendered = handle_unrendered_http_error(options_, request, response); + if (rendered == httplib::Server::HandlerResponse::Handled) { + return rendered; + } + // SPA fallback: the file handler already served every real asset and every + // registered API route has already been tried, so an unmatched GET/HEAD is + // a client-side route (e.g. /chat/123). Hand it the SPA shell; leave every + // other error (405 on a real API path, 404 on a missing asset) untouched. + if (webui_serving_ && response.status == 404 && + (request.method == "GET" || request.method == "HEAD") && + webui_spa_path(request.path)) { + response.set_content(webui_index_html_, "text/html"); + return httplib::Server::HandlerResponse::Handled; + } + return httplib::Server::HandlerResponse::Unhandled; }); if (options_.enable_cors) { server_.set_default_headers( @@ -287,7 +366,11 @@ void HttpServer::register_routes() { } server_.set_pre_routing_handler([this](const httplib::Request& req, httplib::Response& res) { - if (options_.api_key.empty() || req.path == "/health" || req.method == "OPTIONS") { + // Only API endpoints require a key. The UI shell and every static asset load + // freely so the webui can prompt for and send the key on API calls (same + // policy as llama-server). /health stays open and OPTIONS is a CORS preflight. + if (options_.api_key.empty() || req.method == "OPTIONS" || !is_api_path(req.path) || + req.path == "/health") { return httplib::Server::HandlerResponse::Unhandled; } // Accept both the OpenAI-style bearer token and the Anthropic-style diff --git a/src/serve/http_server.h b/src/serve/http_server.h index 1461c40f80..1207e03c4a 100644 --- a/src/serve/http_server.h +++ b/src/serve/http_server.h @@ -39,6 +39,10 @@ class HttpServer { private: void register_routes(); + void mount_webui(const std::string& webui_dir); + void register_webui_mime(); + [[nodiscard]] bool webui_spa_path(const std::string& path) const; + [[nodiscard]] bool is_api_path(const std::string& path) const; void handle_chat_completions(const httplib::Request& req, httplib::Response& res); void handle_messages(const httplib::Request& req, httplib::Response& res); void handle_count_tokens(const httplib::Request& req, httplib::Response& res); @@ -66,6 +70,8 @@ class HttpServer { GenerationService* service_ = nullptr; ServeOptions options_; std::string public_model_id_; + bool webui_serving_ = false; // true once a static webui dir is mounted + std::string webui_index_html_; // cached index.html for the SPA fallback ResponseStore response_store_; JsonlRequestLog request_jsonl_; httplib::Server server_; diff --git a/src/serve/serve_options.cpp b/src/serve/serve_options.cpp index c991e2cc85..9e8c46c97a 100644 --- a/src/serve/serve_options.cpp +++ b/src/serve/serve_options.cpp @@ -75,6 +75,7 @@ std::string serve_usage_text(const char* argv0) { "[--default-max-tokens N] " "[--vision] [--no-cuda-graph] [--no-prefix-reuse] " "[--lm-head-draft] [--no-thinking] [--preserve-thinking] [--cors] " + "[--webui | --webui-dir DIR] " "[--temperature F] [--top-p F] [--top-k N] [--min-p F] [--presence-penalty F] " "[--frequency-penalty F] [--seed N] [--greedy]\n" " serves OpenAI Responses/Chat Completions and Anthropic Messages endpoints\n" @@ -98,6 +99,10 @@ std::string serve_usage_text(const char* argv0) { " --preserve-thinking retains closed-turn assistant reasoning in later prompts\n" " sampler defaults come from the loaded model and resolved thinking mode; " "server flags and request fields override individual values.\n" + " --webui auto-downloads the prebuilt llama.cpp webui (ggml-org/llama-ui " + "HF bucket) into the webui dir and serves it at / alongside the API\n" + " --webui-dir DIR serves (and for --webui, downloads into) DIR; " + "defaults to /webui\n" " --greedy forces temperature 0 (exact argmax).\n"; } @@ -235,6 +240,13 @@ ServeOptions parse_serve_options(int argc, char** argv) { options.preserve_thinking = true; } else if (arg == "--cors") { options.enable_cors = true; + } else if (arg == "--webui") { + options.webui_auto = true; + } else if (arg == "--webui-dir") { + options.webui_dir = require_value("--webui-dir"); + if (options.webui_dir.empty()) { + throw std::invalid_argument("--webui-dir must not be empty"); + } } else if (arg == "--temperature") { options.sampling_overrides.temperature = parse_float_in(require_value("--temperature"), "temperature", 0.0f, 2.0f); diff --git a/src/serve/serve_options.h b/src/serve/serve_options.h index b6db8e4fd5..f8d21c7152 100644 --- a/src/serve/serve_options.h +++ b/src/serve/serve_options.h @@ -50,6 +50,10 @@ struct ServeOptions { bool preserve_thinking = false; int default_max_tokens = kDefaultMaxTokens; bool enable_cors = false; // send permissive CORS headers for browser UIs + bool webui_auto = false; // --webui: auto-download the prebuilt llama.cpp + // webui from the ggml-org/llama-ui HF bucket + std::string webui_dir; // --webui-dir: serve this dir; also the download + // location for --webui (default: /webui) // Process-level explicit overrides layered between registered model/mode defaults and request // fields. An omitted seed is replaced per request with a fresh random seed. SamplingOverrides sampling_overrides; diff --git a/src/serve/webui_update.cpp b/src/serve/webui_update.cpp new file mode 100644 index 0000000000..b1c1cb04ea --- /dev/null +++ b/src/serve/webui_update.cpp @@ -0,0 +1,383 @@ +// webui_update.cpp — keeps a local copy of the prebuilt llama.cpp webui current. +// +// Upstream llama.cpp publishes the built webui (tools/ui static output) to the +// Hugging Face bucket ggml-org/llama-ui after each release, under both a +// release-tag folder and a rolling "latest" pointer. Layout: +// +// https://huggingface.co/api/buckets/ggml-org/llama-ui/tree/latest +// -> JSON file list: [{"type":"file","path":"latest/","size":N}, ...] +// https://huggingface.co/buckets/ggml-org/llama-ui/resolve/latest/ +// -> 302 -> CDN byte stream for that file +// +// ensure_webui_available() compares the local marker file against +// latest/_app/version.json and, when the local copy is missing/stale/incomplete, +// downloads the full set into a staging directory and atomically swaps it in. +// Downloads use WinHTTP directly (the vendored httplib is built without TLS). + +#include "serve/webui_update.h" + +#include "serve/console_log.h" + +#include +#include + +#include + +#include +#include +#include +#include +#include +#include + +#pragma comment(lib, "winhttp.lib") + +namespace ninfer::serve { +namespace { + +namespace fs = std::filesystem; + +constexpr const char* kBucketApi = "https://huggingface.co/api/buckets/ggml-org/llama-ui/tree/latest"; +constexpr const char* kBucketBase = "https://huggingface.co/buckets/ggml-org/llama-ui/resolve/latest/"; +constexpr const char* kUserAgent = "ninfer-serve"; +constexpr const char* kMarkerFile = ".ninfer-webui-version"; +constexpr const char* kVersionPath = "_app/version.json"; +constexpr int kMaxAttempts = 3; +constexpr DWORD kReadChunkBytes = 1 << 20; + +struct WebuiFile { + std::string relative_path; // e.g. "index.html", "_app/immutable/..." + uint64_t size = 0; +}; + +// RAII closer for an HINTERNET handle. The vendored httplib has no TLS, so all +// webui downloads go through WinHTTP directly; a guard keeps every handle closed +// on both the success and every throw path. +struct HandleGuard { + HINTERNET handle = nullptr; + explicit HandleGuard(HINTERNET h) : handle(h) {} + ~HandleGuard() { + if (handle != nullptr) { + ::WinHttpCloseHandle(handle); + handle = nullptr; + } + } + HandleGuard(const HandleGuard&) = delete; + HandleGuard& operator=(const HandleGuard&) = delete; +}; + +std::string to_utf8(const fs::path& p) { + const std::wstring w = p.wstring(); + if (w.empty()) { return std::string(); } + const int len = + ::WideCharToMultiByte(CP_UTF8, 0, w.c_str(), static_cast(w.size()), nullptr, 0, nullptr, nullptr); + std::string out(static_cast(len), '\0'); + ::WideCharToMultiByte(CP_UTF8, 0, w.c_str(), static_cast(w.size()), out.data(), len, nullptr, nullptr); + return out; +} + +std::wstring to_wide(const std::string& s) { + if (s.empty()) { return std::wstring(); } + const int len = ::MultiByteToWideChar(CP_UTF8, 0, s.data(), static_cast(s.size()), nullptr, 0); + std::wstring out(static_cast(len), L'\0'); + ::MultiByteToWideChar(CP_UTF8, 0, s.data(), static_cast(s.size()), out.data(), len); + return out; +} + +// Splits "https://host/path..." into (host, path-and-query). +void split_url(const std::string& url, std::string& host, std::string& path) { + const size_t scheme_end = url.find("://"); + const size_t path_start = url.find('/', scheme_end + 3); + host = url.substr(0, path_start == std::string::npos ? url.size() : path_start); + path = path_start == std::string::npos ? "/" : url.substr(path_start); +} + +// One WinHTTP GET to completion. Returns the body. Follows 301/302/303/307/308. +// Throws std::runtime_error on any failure. +std::string http_get(const std::string& url, const std::string& user_agent) { + std::string host, path; + split_url(url, host, path); + + const std::wstring w_agent = to_wide(user_agent); + HINTERNET session = ::WinHttpOpen(w_agent.c_str(), WINHTTP_ACCESS_TYPE_DEFAULT_PROXY, nullptr, nullptr, 0); + if (session == nullptr) { + throw std::runtime_error("WinHttpOpen failed: " + std::to_string(::GetLastError())); + } + HandleGuard session_guard(session); + + std::string body; + for (int redirect = 0; redirect < 8; ++redirect) { + const std::wstring w_host = to_wide(host); + const std::wstring w_path = to_wide(path); + + // A fresh connect handle each pass: a redirect may cross to another host. + HINTERNET connect = ::WinHttpConnect(session, w_host.c_str(), INTERNET_DEFAULT_HTTPS_PORT, 0); + if (connect == nullptr) { + throw std::runtime_error("WinHttpConnect(" + host + ") failed: " + std::to_string(::GetLastError())); + } + HandleGuard connect_guard(connect); + + HINTERNET request = ::WinHttpOpenRequest(connect, L"GET", w_path.c_str(), nullptr, nullptr, nullptr, + WINHTTP_FLAG_SECURE); + if (request == nullptr) { + throw std::runtime_error("WinHttpOpenRequest(" + url + ") failed: " + + std::to_string(::GetLastError())); + } + HandleGuard request_guard(request); + + const std::wstring accept = L"Accept: application/json, text/html, */*"; + ::WinHttpAddRequestHeaders(request, accept.c_str(), static_cast(accept.size()), + WINHTTP_ADDREQ_FLAG_ADD); + + if (!::WinHttpSendRequest(request, 0, 0, nullptr, 0, 0, 0) || + !::WinHttpReceiveResponse(request, nullptr)) { + throw std::runtime_error("request to " + url + " failed: " + std::to_string(::GetLastError())); + } + + DWORD status = 0; + DWORD status_size = sizeof(status); + ::WinHttpQueryHeaders(request, WINHTTP_QUERY_STATUS_CODE | WINHTTP_QUERY_FLAG_NUMBER, L"__WinHttpStatus", + &status, &status_size, nullptr); + if (status >= 300 && status < 400) { + DWORD location_size = 0; + ::WinHttpQueryHeaders(request, WINHTTP_QUERY_LOCATION, nullptr, nullptr, &location_size, nullptr); + if (location_size <= 1) { + throw std::runtime_error("redirect from " + url + " had no Location header"); + } + std::string location(location_size - 1, '\0'); + ::WinHttpQueryHeaders(request, WINHTTP_QUERY_LOCATION, nullptr, &location[0], &location_size, nullptr); + split_url(location, host, path); // reconnect against the new host on the next pass + continue; + } + if (status >= 400) { + throw std::runtime_error("HTTP " + std::to_string(status) + " from " + url); + } + + for (;;) { + DWORD available = 0; + if (!::WinHttpQueryDataAvailable(request, &available) || available == 0) { break; } + std::vector buffer(std::min(available, kReadChunkBytes)); + DWORD read = 0; + if (!::WinHttpReadData(request, buffer.data(), static_cast(buffer.size()), &read)) { + throw std::runtime_error("read from " + url + " failed: " + std::to_string(::GetLastError())); + } + body.append(buffer.data(), read); + if (read < available) { break; } + } + break; + } + return body; +} + +std::string with_retry(const std::string& url, const std::string& what) { + std::string last_error; + for (int attempt = 1; attempt <= kMaxAttempts; ++attempt) { + try { + return http_get(url, kUserAgent); + } catch (const std::exception& e) { + last_error = e.what(); + if (attempt < kMaxAttempts) { + ::Sleep(1000 * attempt); + } + } + } + throw std::runtime_error(what + ": " + last_error); +} + +// File list for the "latest" pointer, with the "latest/" prefix stripped. +std::vector fetch_file_list() { + const std::string body = with_retry(kBucketApi, "listing the webui bucket failed"); + const auto json = nlohmann::json::parse(body); + std::vector files; + for (const auto& entry : json) { + if (entry.value("type", "") != "file") { continue; } + const std::string path = entry.value("path", ""); + if (path.rfind("latest/", 0) != 0) { continue; } // "latest" pointer only + WebuiFile file; + file.relative_path = path.substr(7); + file.size = entry.value("size", 0ULL); + if (file.relative_path.empty()) { continue; } + files.push_back(std::move(file)); + } + if (files.empty()) { + throw std::runtime_error("webui bucket listing contained no files"); + } + return files; +} + +uint64_t file_size_or_zero(const fs::path& p) { + std::error_code ec; + const auto size = fs::file_size(p, ec); + return ec ? 0 : size; +} + +bool directory_exists(const fs::path& p) { + std::error_code ec; + return fs::is_directory(p, ec); +} + +std::string read_file_text(const fs::path& p) { + std::ifstream in(p, std::ios::binary); + if (!in) { return std::string(); } + return std::string((std::istreambuf_iterator(in)), std::istreambuf_iterator()); +} + +// Downloads one file to dest, verifying the final byte count against expected_size +// when it is non-zero. +void download_file(const std::string& relative_path, const fs::path& dest, uint64_t expected_size) { + const std::string url = std::string(kBucketBase) + relative_path; + std::string last_error; + for (int attempt = 1; attempt <= kMaxAttempts; ++attempt) { + try { + const std::string body = http_get(url, kUserAgent); + std::error_code ec; + fs::create_directories(dest.parent_path(), ec); + { + std::ofstream out(dest, std::ios::binary | std::ios::trunc); + if (!out) { throw std::runtime_error("cannot open " + to_utf8(dest)); } + out.write(body.data(), static_cast(body.size())); + out.flush(); + if (!out) { throw std::runtime_error("write failed for " + to_utf8(dest)); } + } + if (expected_size != 0 && file_size_or_zero(dest) != expected_size) { + throw std::runtime_error("size mismatch for " + relative_path); + } + return; + } catch (const std::exception& e) { + last_error = e.what(); + if (attempt < kMaxAttempts) { + ::Sleep(1000 * attempt); + } + } + } + throw std::runtime_error("downloading " + relative_path + " failed: " + last_error); +} + +// The rolling "latest" pointer is only meaningful if it actually points at a +// published release folder; a 404 on version.json means the bucket is empty. +std::string fetch_latest_version() { + const std::string body = with_retry(std::string(kBucketBase) + kVersionPath, + "reading the webui version marker failed"); + const auto json = nlohmann::json::parse(body); + const std::string version = json.value("version", ""); + if (version.empty()) { + throw std::runtime_error("webui bucket has no version.json"); + } + return version; +} + +bool local_copy_is_current(const fs::path& webui_dir, const std::string& version) { + if (!directory_exists(webui_dir)) { return false; } + if (read_file_text(webui_dir / kMarkerFile) != version) { return false; } + return file_size_or_zero(webui_dir / "index.html") != 0; +} + +std::vector staged_stale_dirs(const fs::path& webui_dir) { + // Staging dirs from a previously interrupted run: .ninfer-webui..tmp + std::vector stale; + std::error_code ec; + const fs::path parent = webui_dir.parent_path(); + const std::string prefix = ".ninfer-webui."; + if (!directory_exists(parent)) { return stale; } + for (const auto& entry : fs::directory_iterator(parent, ec)) { + const std::string name = entry.path().filename().string(); + if (name.rfind(prefix, 0) == 0 && name.size() > prefix.size() + 4 && + name.compare(name.size() - 4, 4, ".tmp") == 0) { + stale.push_back(entry.path()); + } + } + return stale; +} + +} // namespace + +std::string resolve_webui_dir(const ServeOptions& options) { + if (!options.webui_dir.empty()) { return options.webui_dir; } + std::error_code ec; + fs::path artifact(options.artifact_path); + fs::path dir = artifact.parent_path(); + if (dir.empty()) { dir = fs::path("."); } + return to_utf8((dir / "webui").lexically_normal()); +} + +std::string ensure_webui_available(const std::string& webui_dir) { + const fs::path target(webui_dir); + + // Sweep staging dirs left behind by a previously interrupted download. + for (const auto& stale : staged_stale_dirs(target)) { + std::error_code ec; + fs::remove_all(stale, ec); + write_console_log(ConsoleLogLevel::Info, "removed stale webui staging directory " + to_utf8(stale)); + } + + const std::string version = fetch_latest_version(); + + if (local_copy_is_current(target, version)) { + write_console_log(ConsoleLogLevel::Info, + "webui up to date (version " + version + ") at " + webui_dir); + return webui_dir; + } + + write_console_log(ConsoleLogLevel::Info, "downloading latest webui (version " + version + + ") from ggml-org/llama-ui..."); + + const std::vector files = fetch_file_list(); + + // Stage in a sibling directory, then swap atomically: the served directory is + // never half-written, and an interrupted run leaves the previous copy intact. + const fs::path staging = + target.parent_path() / (".ninfer-webui." + std::to_string(::GetCurrentProcessId()) + ".tmp"); + { + std::error_code ec; + fs::remove_all(staging, ec); + fs::create_directories(staging, ec); + if (ec) { + throw std::runtime_error("cannot create staging directory " + to_utf8(staging) + ": " + ec.message()); + } + } + + uint64_t total_bytes = 0; + for (const auto& file : files) { + const fs::path dest = staging / file.relative_path; + download_file(file.relative_path, dest, file.size); + total_bytes += file.size; + } + + // version.json is downloaded with the rest; publish its value as the marker. + const std::string staged_version = fetch_latest_version(); + { + std::ofstream out(staging / kMarkerFile, std::ios::binary | std::ios::trunc); + out << staged_version; + } + + // Atomic swap: rename target aside, move staging in, drop the old copy. + const fs::path old = + target.parent_path() / (".ninfer-webui.old." + std::to_string(::GetCurrentProcessId())); + std::error_code ec; + fs::remove_all(old, ec); + if (directory_exists(target)) { + fs::rename(target, old, ec); + if (ec) { + throw std::runtime_error("cannot move aside " + to_utf8(target) + ": " + ec.message()); + } + } + ec.clear(); + fs::rename(staging, target, ec); + if (ec) { + // Roll the served copy back before failing. + std::error_code ec2; + if (directory_exists(old)) { fs::rename(old, target, ec2); } + throw std::runtime_error("cannot install webui at " + to_utf8(target) + ": " + ec.message()); + } + fs::remove_all(old, ec); + + const double mib = static_cast(total_bytes) / (1024.0 * 1024.0); + write_console_log(ConsoleLogLevel::Info, + std::to_string(files.size()) + " webui files (" + + (mib >= 10 ? std::to_string(static_cast(mib)) : + std::to_string(static_cast(mib * 10) / 10)) + + " MiB) installed at " + webui_dir); + return webui_dir; +} + +} // namespace ninfer::serve \ No newline at end of file diff --git a/src/serve/webui_update.h b/src/serve/webui_update.h new file mode 100644 index 0000000000..97e1b50297 --- /dev/null +++ b/src/serve/webui_update.h @@ -0,0 +1,25 @@ +#pragma once + +// Keep a local copy of the prebuilt llama.cpp webui (tools/ui static output) +// current. Upstream publishes it to the Hugging Face bucket ggml-org/llama-ui +// under both release-tag folders and a rolling "latest" pointer. + +#include "serve/serve_options.h" + +#include +#include +#include + +namespace ninfer::serve { + +// Resolves the webui directory for --webui auto mode: --webui-dir if set, +// otherwise /webui. +std::string resolve_webui_dir(const ServeOptions& options); + +// Downloads the current prebuilt llama.cpp webui (ggml-org/llama-ui HF bucket, +// "latest" pointer) into webui_dir when the local copy is missing, stale, or +// incomplete, then returns webui_dir ready to serve. Throws std::runtime_error +// on an unrecoverable download failure. No-op when the local copy is current. +std::string ensure_webui_available(const std::string& webui_dir); + +} // namespace ninfer::serve \ No newline at end of file From dba5fc32ccb9b2db8f104c2dce4047c6bf5fc759 Mon Sep 17 00:00:00 2001 From: natpate <80576637+natpate@users.noreply.github.com> Date: Mon, 17 Aug 2026 01:31:56 -0500 Subject: [PATCH 04/20] fix(webui): strip scheme from host in split_url WinHttpConnect was being handed "https://huggingface.co" instead of "huggingface.co", which cannot DNS-resolve: every --webui download failed with 12005 on any machine. Affected both the initial bucket URL and CDN redirect Location headers. Caught by the portable-zip release test (first clean-dir run of the auto-download path). --- src/serve/webui_update.cpp | 6 ++++-- 1 file changed, 4 insertions(+), 2 deletions(-) diff --git a/src/serve/webui_update.cpp b/src/serve/webui_update.cpp index b1c1cb04ea..21250d0c8b 100644 --- a/src/serve/webui_update.cpp +++ b/src/serve/webui_update.cpp @@ -87,8 +87,10 @@ std::wstring to_wide(const std::string& s) { // Splits "https://host/path..." into (host, path-and-query). void split_url(const std::string& url, std::string& host, std::string& path) { const size_t scheme_end = url.find("://"); - const size_t path_start = url.find('/', scheme_end + 3); - host = url.substr(0, path_start == std::string::npos ? url.size() : path_start); + // Host starts after the scheme so WinHttpConnect never sees "https://host". + const size_t host_start = scheme_end == std::string::npos ? 0 : scheme_end + 3; + const size_t path_start = url.find('/', host_start); + host = url.substr(host_start, path_start == std::string::npos ? std::string::npos : path_start - host_start); path = path_start == std::string::npos ? "/" : url.substr(path_start); } From a836c865fc94902ec75b6aafa9d2b596faff383d Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 11:30:06 -0400 Subject: [PATCH 05/20] fix(win32): compare artifact size against size_t without signed cast The Windows mapping path casted size_t's maximum to LONGLONG before comparing it with the file size. That cast wraps to -1 on a 64-bit build, so every non-empty artifact compared greater and was rejected with "artifact size does not fit the process address space" before the mapping was ever attempted. Compare in unsigned terms instead, matching the POSIX path below. Co-Authored-By: Claude Opus 5 --- src/artifact/reader.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 4f61bd2237..918b015f02 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -5,6 +5,7 @@ #include #include #include +#include #include #include #include @@ -206,7 +207,8 @@ class MappedFile { throw std::system_error(error, std::system_category(), "GetFileSizeEx " + path.string()); } if (size_info.QuadPart < 0 || - size_info.QuadPart > static_cast(std::numeric_limits::max())) { + static_cast(size_info.QuadPart) > + std::numeric_limits::max()) { ::CloseHandle(file); throw ArtifactError("artifact size does not fit the process address space"); } From 98269828197130d94a3da22b6ac23d8b09f39924 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:10:16 -0400 Subject: [PATCH 06/20] fix(win32): mirror pread short-read semantics in the mapped direct read The Windows read_direct rejected any request running past the last byte, but direct I/O rounds every request up to the 4096-byte alignment, so the final read of an artifact overruns EOF by design. The POSIX path returns a short pread() count there and the materializer relies on it, comparing the result against min(request, remaining). Clamp to the mapped size and return the bytes actually copied, which fixes ninfer_artifact_materialization_test on Windows. Co-Authored-By: Claude Opus 5 --- src/artifact/reader.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 918b015f02..12e4d80e6e 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -300,11 +300,16 @@ class MappedFile { } #ifdef _WIN32 // The whole file is already mapped, so a direct read is a copy out of the view. - if (absolute_offset > size_ || destination.size() > size_ - absolute_offset) { + if (absolute_offset > size_) { throw ArtifactError("direct artifact read exceeds the mapped file"); } - std::memcpy(destination.data(), data_ + absolute_offset, destination.size()); - return destination.size(); + // Direct I/O rounds every request up to the alignment, so the final one runs + // past the last byte by design. pread() answers that with a short count and + // callers rely on it, so clamp and report what was actually copied. + const auto offset = static_cast(absolute_offset); + const std::size_t available = std::min(destination.size(), size_ - offset); + std::memcpy(destination.data(), data_ + offset, available); + return available; #else if (absolute_offset > static_cast(std::numeric_limits::max()) || destination.size() > static_cast(std::numeric_limits::max())) { From 5146f4465e596b152f02f2898f5a2b76656c181d Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:10:18 -0400 Subject: [PATCH 07/20] fix(serve): keep the reasoning-effort conflict in prompt resolution The imported WebUI dialect commit also rejected an effort that disagrees with enable_thinking during parsing. resolve_prompt_semantics already raises that exact error with the same code, so the check was a duplicate that moved a semantic decision into protocol-shape validation and broke the existing contract test, which asserts that parsing accepts the body. Parsing threw outside the test's api_code() helper, aborting the binary. Drop the parse-time check and align the imported test with the layer that owns it. The conflicting top-level/kwargs enable_thinking check stays in the parser, where two spellings of one field genuinely disagree. Also read n_ctx from default_generation_settings rather than from params: that is where make_props_stub puts it, matching llama.cpp's /props shape. Co-Authored-By: Claude Opus 5 --- src/serve/openai_schema.cpp | 9 ++++----- tests/test_openai_schema.cpp | 18 +++++++++++++----- 2 files changed, 17 insertions(+), 10 deletions(-) diff --git a/src/serve/openai_schema.cpp b/src/serve/openai_schema.cpp index b8f734f4e1..957fe01868 100644 --- a/src/serve/openai_schema.cpp +++ b/src/serve/openai_schema.cpp @@ -569,11 +569,10 @@ void parse_openai_chat_thinking(const Json& body, std::optional* enable_th if (effort) { *reasoning_effort = *effort; *reasoning_effort_param = "reasoning_effort"; - if (enable_thinking->has_value() && - *enable_thinking != (*effort != RequestedReasoningEffort::None)) { - bad_request("reasoning effort conflicts with " + conflict_param, - "reasoning_effort", "conflicting_template_option"); - } + // An effort that disagrees with enable_thinking is a semantic conflict, not a + // protocol-shape one: resolve_prompt_semantics owns that check so it can weigh + // the request against the loaded template's capabilities. Rejecting it here too + // would only move the same error earlier and break that layering. } } diff --git a/tests/test_openai_schema.cpp b/tests/test_openai_schema.cpp index 13b95c1049..9727022aaa 100644 --- a/tests/test_openai_schema.cpp +++ b/tests/test_openai_schema.cpp @@ -780,11 +780,17 @@ int test_llama_webui_dialect() { low_prompt.options.reasoning_effort == ninfer::ReasoningEffort::Low, "webui low effort reached prompt"); - // conflicting enable_thinking=false vs effort low rejected at parse time + // conflicting enable_thinking=false vs effort low: the body is protocol-valid, so + // parsing accepts it and resolve_prompt_semantics reports the semantic conflict. Json c = low; c["chat_template_kwargs"] = Json{{"enable_thinking", false}}; + const GenerationRequest c_request = + parse_chat_completion_request(c, default_limits(), webui_model()); + failures += check(c_request.enable_thinking == false, + "kwargs enable_thinking parsed alongside reasoning_effort"); failures += check(api_code([&] { - (void)parse_chat_completion_request(c, default_limits(), webui_model()); + (void)resolve_prompt_semantics(c_request, default_server(), + effort_capabilities()); }) == "conflicting_template_option", "conflicting enable_thinking and reasoning_effort rejected"); @@ -834,7 +840,7 @@ int test_props_stub() { options.max_context = 16384; options.default_max_tokens = 4096; options.enable_vision = true; - options.speculative.backend = SpeculativeBackend::Mtp; + options.speculative.backend = ninfer::SpeculativeBackend::Mtp; options.sampling_overrides.temperature = 1.0F; options.sampling_overrides.top_k = 20; @@ -842,8 +848,10 @@ int test_props_stub() { failures += check(props.at("role") == "model", "props role is model"); failures += check(props.at("modalities").at("vision") == true, "props vision follows --vision"); failures += check(props.at("modalities").at("audio") == false, "props audio off"); - const Json params = props.at("default_generation_settings").at("params"); - failures += check(params.at("n_ctx") == 16384, "props n_ctx from --max-context"); + const Json settings = props.at("default_generation_settings"); + const Json params = settings.at("params"); + // llama.cpp reports n_ctx on the settings object itself, not inside params. + failures += check(settings.at("n_ctx") == 16384, "props n_ctx from --max-context"); failures += check(params.at("n_predict") == 4096, "props n_predict from default max tokens"); failures += check(params.at("temperature") == 1.0, "props temperature override reported"); failures += check(params.at("top_k") == 20, "props top_k override reported"); From 07afb8cefec3850885380254e3b71311f80a1068 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:10:27 -0400 Subject: [PATCH 08/20] test(win32): build and run the test suite under MSVC The suite was excluded from the default build, so it had never been compiled on Windows. Four gaps kept it from running: - std::aligned_alloc is absent from the MSVC CRT, and its aligned blocks need _aligned_free rather than free. - std::sqrt is not constexpr before C++26; libstdc++ accepts it as an extension, MSVC does not. These initializers only need const. - The FFmpeg/libcurl DLLs are copied next to the apps but not next to the tests, so every test linking them failed to start. Prepend the prefix bin directories to each test's PATH instead of duplicating the DLLs. - test_frontend reads an official tokenizer from a hard-coded local HF checkout and aborted when absent. Report the standard skip code, and register the test with SKIP_RETURN_CODE like its peers. ctest now runs 83 tests green with 6 skipped for missing local artifacts. Co-Authored-By: Claude Opus 5 --- tests/CMakeLists.txt | 26 ++++++++++++++++ .../test_gated_delta_net_replay_record.cpp | 2 +- tests/ops/test_gdn_replay_fold.cpp | 4 +-- tests/ops/test_vision_attention.cpp | 2 +- tests/targets/qwen3_6/test_frontend.cpp | 30 +++++++++++++++++-- tests/test_gdn_replay_records.cpp | 22 ++++++++++++-- 6 files changed, 77 insertions(+), 9 deletions(-) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 674da12a62..e7a680fc75 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1,10 +1,33 @@ find_package(Python3 REQUIRED COMPONENTS Interpreter) +# On Windows the FFmpeg/libcurl runtime DLLs are not on the default search path. +# The app targets copy them next to their executables; tests instead get those +# directories prepended to PATH, so the suite runs in place without duplicating +# the DLLs into every test output directory. +set(NINFER_TEST_DLL_DIRS "") +if(WIN32) + foreach(_prefix ${NINFER_FFMPEG_PREFIX} ${NINFER_CURL_PREFIX}) + if(_prefix AND EXISTS "${_prefix}/bin") + file(TO_NATIVE_PATH "${_prefix}/bin" _native_bin) + list(APPEND NINFER_TEST_DLL_DIRS "${_native_bin}") + endif() + endforeach() + list(REMOVE_DUPLICATES NINFER_TEST_DLL_DIRS) +endif() + +function(ninfer_test_runtime_path name) + foreach(_dir IN LISTS NINFER_TEST_DLL_DIRS) + set_property(TEST ${name} APPEND PROPERTY + ENVIRONMENT_MODIFICATION "PATH=path_list_prepend:${_dir}") + endforeach() +endfunction() + # This target deliberately receives no src/, CUDA, artifact, kernel, or target # include root. It proves that the installed product headers stand alone. add_executable(ninfer_public_api_test test_public_api.cpp) target_include_directories(ninfer_public_api_test PRIVATE ${PROJECT_SOURCE_DIR}/include) add_test(NAME ninfer_public_api_test COMMAND ninfer_public_api_test) +ninfer_test_runtime_path(ninfer_public_api_test) function(ninfer_add_test name) cmake_parse_arguments(arg "NEEDS_SOURCE_DIR" "" "SOURCES;LIBRARIES" ${ARGN}) @@ -26,6 +49,7 @@ function(ninfer_add_test name) NINFER_PYTHON_EXECUTABLE="${Python3_EXECUTABLE}") endif() add_test(NAME ${name} COMMAND ${name}) + ninfer_test_runtime_path(${name}) endfunction() function(ninfer_add_op_test name) @@ -101,6 +125,8 @@ ninfer_add_test(ninfer_qwen3_6_frontend_test target_include_directories(ninfer_qwen3_6_frontend_test PRIVATE ${PROJECT_SOURCE_DIR}/src/targets/qwen3_6/export ${PROJECT_SOURCE_DIR}/src/targets/qwen3_6/impl) +# Needs a local Qwen3.6-27B HF checkout for the official tokenizer fixtures. +set_tests_properties(ninfer_qwen3_6_frontend_test PROPERTIES SKIP_RETURN_CODE 77) ninfer_add_test(ninfer_qwen3_6_runtime_mechanisms_test SOURCES targets/qwen3_6/test_runtime_mechanisms.cpp LIBRARIES ninfer_engine ninfer_core) diff --git a/tests/ops/test_gated_delta_net_replay_record.cpp b/tests/ops/test_gated_delta_net_replay_record.cpp index f2175dd04a..7e5e7d3708 100644 --- a/tests/ops/test_gated_delta_net_replay_record.cpp +++ b/tests/ops/test_gated_delta_net_replay_record.cpp @@ -115,7 +115,7 @@ int run_case(std::int32_t value_heads, std::int32_t width, std::int32_t batch, Tensor value_record_tensor(value_record.p, DType::BF16, {kStateDim, value_heads, width, batch}); Tensor gate_record_tensor(gate_record.p, DType::FP32, {2, value_heads, width, batch}); - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); ops::gated_delta_net_snapshot(q, k, v, g_tensor, beta_tensor, kScale, true, snapshot_states, valid, initial, bases, snapshot_output, nullptr); ops::gated_delta_net_replay_record(q, k, v, g_tensor, beta_tensor, kScale, record_states, valid, diff --git a/tests/ops/test_gdn_replay_fold.cpp b/tests/ops/test_gdn_replay_fold.cpp index 037f462f19..3dabcaa59f 100644 --- a/tests/ops/test_gdn_replay_fold.cpp +++ b/tests/ops/test_gdn_replay_fold.cpp @@ -318,7 +318,7 @@ int run_case(const FoldProfile profile, std::int32_t width, std::int32_t rows, Tensor output(out.p, DType::BF16, {kStateDim, profile.value_heads, width, 1}); Tensor initial_selector(initial_device.p, DType::I32, {1}); Tensor base_selector(base_device.p, DType::I32, {1}); - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); for (std::int32_t layer = 0; layer < profile.layers; ++layer) { const GdnReplayRecordLayer layer_records = records.layer(layer, rows); @@ -503,7 +503,7 @@ int run_record_fold_rounds() { constexpr std::int32_t kStateSlots = 3; constexpr std::int32_t kInitialSlot = 2; constexpr std::int32_t kSnapshotBase = 0; - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); DevicePackedWeight parent( quantized_weight::make_patterned_weight(QType::W8G32_F16S, kParentRows, kHidden, 1901U)); diff --git a/tests/ops/test_vision_attention.cpp b/tests/ops/test_vision_attention.cpp index 6c6ce630bb..5327533c21 100644 --- a/tests/ops/test_vision_attention.cpp +++ b/tests/ops/test_vision_attention.cpp @@ -41,7 +41,7 @@ std::vector bf16_bits(const std::vector& values) { void vision_attention_oracle(const std::vector& q, const std::vector& k, const std::vector& v, const std::vector& cu_seqlens, std::vector& out) { - constexpr double scale = 1.0 / std::sqrt(72.0); + const double scale = 1.0 / std::sqrt(72.0); out.assign(q.size(), 0.0); for (std::size_t segment = 0; segment + 1 < cu_seqlens.size(); ++segment) { diff --git a/tests/targets/qwen3_6/test_frontend.cpp b/tests/targets/qwen3_6/test_frontend.cpp index 4e29017dae..d9a8ff148a 100644 --- a/tests/targets/qwen3_6/test_frontend.cpp +++ b/tests/targets/qwen3_6/test_frontend.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -88,13 +89,31 @@ const fi::CompiledChatTemplate& reasoning_effort_template() { return value; } +// The official tokenizer fixtures are a local HF checkout, not part of this +// repository, so a checkout without them skips instead of failing. +constexpr const char* kOfficialTokenizerDir = + "/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16"; + +std::string official_tokenizer_file(const char* name) { + return (std::filesystem::path(kOfficialTokenizerDir) / name).string(); +} + +bool official_tokenizer_available() { + for (const char* name : + {"tokenizer.json", "tokenizer_config.json", "generation_config.json"}) { + std::error_code error; + if (!std::filesystem::exists(official_tokenizer_file(name), error)) { return false; } + } + return true; +} + const fi::Tokenizer& official_tokenizer() { static const std::string tokenizer_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/tokenizer.json"); + read_file(official_tokenizer_file("tokenizer.json").c_str()); static const std::string tokenizer_config_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/tokenizer_config.json"); + read_file(official_tokenizer_file("tokenizer_config.json").c_str()); static const std::string generation_config_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/generation_config.json"); + read_file(official_tokenizer_file("generation_config.json").c_str()); static const fi::Tokenizer tokenizer({.tokenizer_json = tokenizer_json, .tokenizer_config_json = tokenizer_config_json, .generation_config_json = generation_config_json}); @@ -1305,6 +1324,11 @@ int test_media_preparation_cancellation() { } // namespace int main() { + if (!official_tokenizer_available()) { + std::cerr << "skipping: official tokenizer fixtures not found under " + << kOfficialTokenizerDir << '\n'; + return 77; + } const FrontendResources owned = resources(); const Frontend frontend = FrontendFactory::create_component(owned); int failures = 0; diff --git a/tests/test_gdn_replay_records.cpp b/tests/test_gdn_replay_records.cpp index 2c777b5fe2..938ffff64b 100644 --- a/tests/test_gdn_replay_records.cpp +++ b/tests/test_gdn_replay_records.cpp @@ -9,14 +9,32 @@ #include #include +#ifdef _WIN32 +#include +#endif + namespace { -using AlignedBacking = std::unique_ptr; +// The MSVC CRT provides no std::aligned_alloc, and blocks from _aligned_malloc +// must be released with _aligned_free rather than free. +void aligned_release(void* data) { +#ifdef _WIN32 + ::_aligned_free(data); +#else + std::free(data); +#endif +} + +using AlignedBacking = std::unique_ptr; AlignedBacking make_backing(std::size_t bytes) { +#ifdef _WIN32 + void* data = ::_aligned_malloc(bytes, 256); +#else void* data = std::aligned_alloc(256, bytes); +#endif if (data == nullptr) { throw std::bad_alloc(); } - return AlignedBacking(data, &std::free); + return AlignedBacking(data, &aligned_release); } int fail(const char* label) { From 50f3ea09f242400a460735da8716302734a84f50 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:54:44 -0400 Subject: [PATCH 09/20] feat(serve): announce a reachable URL and the webui at startup MIME-Version: 1.0 Content-Type: text/plain; charset=UTF-8 Content-Transfer-Encoding: 8bit The startup line echoed the bind address verbatim, so a default --host 0.0.0.0 advertised http://0.0.0.0:8080 — an address browsers refuse with ERR_ADDRESS_INVALID, since a wildcard bind is not a destination. Anyone copying the announced URL landed on an error page while the server was serving normally. Keep reporting the bound address, marking it as covering all interfaces, and add the loopback URL that can actually be opened. Announce the webui entry point when one is mounted, plus the API base. Co-Authored-By: Claude Opus 5 --- apps/serve/main.cpp | 22 ++++++++++++++++++++-- 1 file changed, 20 insertions(+), 2 deletions(-) diff --git a/apps/serve/main.cpp b/apps/serve/main.cpp index f78c750210..42158c1b12 100644 --- a/apps/serve/main.cpp +++ b/apps/serve/main.cpp @@ -119,12 +119,30 @@ int main(int argc, char** argv) { std::signal(SIGINT, handle_signal); std::signal(SIGTERM, handle_signal); + // A wildcard bind address is not reachable as a destination: browsers reject + // http://0.0.0.0/ with ERR_ADDRESS_INVALID. Announce a loopback URL that can + // actually be opened, alongside the address the socket is bound to. + const bool wildcard_bind = options.host == "0.0.0.0" || options.host == "::" || + options.host == "[::]"; + const std::string browse_host = + !wildcard_bind ? options.host : (options.host == "0.0.0.0" ? "127.0.0.1" : "[::1]"); + const std::string browse_url = + "http://" + browse_host + ':' + std::to_string(options.port); + std::ostringstream listening; - listening << "listening on http://" << options.host << ':' << options.port - << " (model id: " << server.public_model_id() + listening << "listening on http://" << options.host << ':' << options.port; + if (wildcard_bind) { listening << " (all interfaces)"; } + listening << " (model id: " << server.public_model_id() << ", auth: " << (options.api_key.empty() ? "disabled" : "bearer") << ')'; ninfer::serve::write_console_log(ninfer::serve::ConsoleLogLevel::Info, listening.str()); + if (!options.webui_dir.empty()) { + ninfer::serve::write_console_log(ninfer::serve::ConsoleLogLevel::Info, + "webui: open " + browse_url + '/'); + } + ninfer::serve::write_console_log(ninfer::serve::ConsoleLogLevel::Info, + "api base: " + browse_url + "/v1"); + const bool ok = server.listen(); g_server.store(nullptr); if (!ok) { From 116de42487831ca740042c4bc1a7b7983d21d4ad Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:57:11 -0400 Subject: [PATCH 10/20] chore: ignore runtime-fetched artifacts and local tool settings The .ninfer model file and the WebUI bundle that --webui downloads both land inside tracked directories, so a plain `git add -A` would otherwise stage roughly 17 GiB of model weights. Also cover the Qwen CLI settings directory and the credential files the GitHub Actions auth action writes into the working tree. Co-Authored-By: Claude Opus 5 --- .gitignore | 11 +++++++++++ 1 file changed, 11 insertions(+) diff --git a/.gitignore b/.gitignore index 70aa014fc1..1729a3cc80 100644 --- a/.gitignore +++ b/.gitignore @@ -46,3 +46,14 @@ profiles/ # Claude Code personal/local settings (may contain API tokens) .claude/settings.local.json + +# Qwen CLI personal/local settings +.qwen/ + +# Credential files written by the GitHub Actions auth action +gha-creds-*.json + +# Model artifacts and the WebUI bundle downloaded by --webui: both are fetched +# into the working tree at runtime and must never be committed. +*.ninfer +model-cards/*/webui/ From 2e5aa8690675e6adfae97da0fc7cefe4c0f8afd1 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:57:11 -0400 Subject: [PATCH 11/20] ci: track the Qwen automation workflows These workflows were sitting untracked in the working tree. They mint a GitHub App token at run time and read every credential from `secrets`, so nothing sensitive is committed here. Co-Authored-By: Claude Opus 5 --- .github/workflows/qwen-dispatch.yml | 204 +++++++++++++++++++ .github/workflows/qwen-invoke.yml | 116 +++++++++++ .github/workflows/qwen-review.yml | 104 ++++++++++ .github/workflows/qwen-scheduled-triage.yml | 208 ++++++++++++++++++++ .github/workflows/qwen-triage.yml | 152 ++++++++++++++ 5 files changed, 784 insertions(+) create mode 100644 .github/workflows/qwen-dispatch.yml create mode 100644 .github/workflows/qwen-invoke.yml create mode 100644 .github/workflows/qwen-review.yml create mode 100644 .github/workflows/qwen-scheduled-triage.yml create mode 100644 .github/workflows/qwen-triage.yml diff --git a/.github/workflows/qwen-dispatch.yml b/.github/workflows/qwen-dispatch.yml new file mode 100644 index 0000000000..e05a98e42d --- /dev/null +++ b/.github/workflows/qwen-dispatch.yml @@ -0,0 +1,204 @@ +name: '🔀 Qwen Code Dispatch' + +on: + pull_request_review_comment: + types: + - 'created' + pull_request_review: + types: + - 'submitted' + pull_request: + types: + - 'opened' + issues: + types: + - 'opened' + - 'reopened' + issue_comment: + types: + - 'created' + +defaults: + run: + shell: 'bash' + +jobs: + debugger: + if: |- + ${{ fromJSON(vars.DEBUG || vars.ACTIONS_STEP_DEBUG || false) }} + runs-on: 'ubuntu-latest' + permissions: + contents: 'read' + steps: + - name: 'Print context for debugging' + env: + DEBUG_event_name: '${{ github.event_name }}' + DEBUG_event__action: '${{ github.event.action }}' + DEBUG_event__comment__author_association: '${{ github.event.comment.author_association }}' + DEBUG_event__issue__author_association: '${{ github.event.issue.author_association }}' + DEBUG_event__pull_request__author_association: '${{ github.event.pull_request.author_association }}' + DEBUG_event__review__author_association: '${{ github.event.review.author_association }}' + DEBUG_event: '${{ toJSON(github.event) }}' + run: |- + env | grep '^DEBUG_' + + dispatch: + # For PRs: only if not from a fork + # For issues: only on open/reopen + # For comments: only if user types @gemini-cli and is OWNER/MEMBER/COLLABORATOR + if: |- + ( + github.event_name == 'pull_request' && + github.event.pull_request.head.repo.fork == false + ) || ( + github.event_name == 'issues' && + contains(fromJSON('["opened", "reopened"]'), github.event.action) + ) || ( + github.event.sender.type == 'User' && + startsWith(github.event.comment.body || github.event.review.body || github.event.issue.body, '@qwen-code') && + contains(fromJSON('["OWNER", "MEMBER", "COLLABORATOR"]'), github.event.comment.author_association || github.event.review.author_association || github.event.issue.author_association) + ) + runs-on: 'ubuntu-latest' + permissions: + contents: 'read' + issues: 'write' + pull-requests: 'write' + outputs: + command: '${{ steps.extract_command.outputs.command }}' + request: '${{ steps.extract_command.outputs.request }}' + additional_context: '${{ steps.extract_command.outputs.additional_context }}' + issue_number: '${{ github.event.pull_request.number || github.event.issue.number }}' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Extract command' + id: 'extract_command' + uses: 'actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea' # ratchet:actions/github-script@v7 + env: + EVENT_TYPE: '${{ github.event_name }}.${{ github.event.action }}' + REQUEST: '${{ github.event.comment.body || github.event.review.body || github.event.issue.body }}' + with: + script: | + const eventType = process.env.EVENT_TYPE; + const request = process.env.REQUEST; + core.setOutput('request', request); + + if (eventType === 'pull_request.opened') { + core.setOutput('command', 'review'); + } else if (['issues.opened', 'issues.reopened'].includes(eventType)) { + core.setOutput('command', 'triage'); + } else if (request.startsWith("@qwen-code /review")) { + core.setOutput('command', 'review'); + const additionalContext = request.replace(/^@qwen-code \/review/, '').trim(); + core.setOutput('additional_context', additionalContext); + } else if (request.startsWith("@qwen-code /triage")) { + core.setOutput('command', 'triage'); + } else if (request.startsWith("@qwen-code")) { + const additionalContext = request.replace(/^@qwen-code/, '').trim(); + core.setOutput('command', 'invoke'); + core.setOutput('additional_context', additionalContext); + } else { + core.setOutput('command', 'fallthrough'); + } + + - name: 'Acknowledge request' + env: + GITHUB_TOKEN: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + ISSUE_NUMBER: '${{ github.event.pull_request.number || github.event.issue.number }}' + MESSAGE: |- + 🤖 Hi @${{ github.actor }}, I've received your request, and I'm working on it now! You can track my progress [in the logs](${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}) for more details. + REPOSITORY: '${{ github.repository }}' + run: |- + gh issue comment "${ISSUE_NUMBER}" \ + --body "${MESSAGE}" \ + --repo "${REPOSITORY}" + + review: + needs: 'dispatch' + if: |- + ${{ needs.dispatch.outputs.command == 'review' }} + uses: './.github/workflows/qwen-review.yml' + permissions: + contents: 'read' + id-token: 'write' + issues: 'write' + pull-requests: 'write' + with: + additional_context: '${{ needs.dispatch.outputs.additional_context }}' + secrets: 'inherit' + + triage: + needs: 'dispatch' + if: |- + ${{ needs.dispatch.outputs.command == 'triage' }} + uses: './.github/workflows/qwen-triage.yml' + permissions: + contents: 'read' + id-token: 'write' + issues: 'write' + pull-requests: 'write' + with: + additional_context: '${{ needs.dispatch.outputs.additional_context }}' + secrets: 'inherit' + + invoke: + needs: 'dispatch' + if: |- + ${{ needs.dispatch.outputs.command == 'invoke' }} + uses: './.github/workflows/qwen-invoke.yml' + permissions: + contents: 'read' + id-token: 'write' + issues: 'write' + pull-requests: 'write' + with: + additional_context: '${{ needs.dispatch.outputs.additional_context }}' + secrets: 'inherit' + + fallthrough: + needs: + - 'dispatch' + - 'review' + - 'triage' + - 'invoke' + if: |- + ${{ always() && !cancelled() && (failure() || needs.dispatch.outputs.command == 'fallthrough') }} + runs-on: 'ubuntu-latest' + permissions: + contents: 'read' + issues: 'write' + pull-requests: 'write' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Send failure comment' + env: + GITHUB_TOKEN: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + ISSUE_NUMBER: '${{ github.event.pull_request.number || github.event.issue.number }}' + MESSAGE: |- + 🤖 I'm sorry @${{ github.actor }}, but I was unable to process your request. Please [see the logs](${{ github.server_url }}/${{ github.repository }}/actions/runs/${{ github.run_id }}) for more details. + REPOSITORY: '${{ github.repository }}' + run: |- + gh issue comment "${ISSUE_NUMBER}" \ + --body "${MESSAGE}" \ + --repo "${REPOSITORY}" diff --git a/.github/workflows/qwen-invoke.yml b/.github/workflows/qwen-invoke.yml new file mode 100644 index 0000000000..0d4bdb33d8 --- /dev/null +++ b/.github/workflows/qwen-invoke.yml @@ -0,0 +1,116 @@ +name: '▶️ Qwen Code Invoke' + +on: + workflow_call: + inputs: + additional_context: + type: 'string' + description: 'Any additional context from the request' + required: false + +concurrency: + group: '${{ github.workflow }}-invoke-${{ github.event_name }}-${{ github.event.pull_request.number || github.event.issue.number }}' + cancel-in-progress: false + +defaults: + run: + shell: 'bash' + +jobs: + invoke: + runs-on: 'ubuntu-latest' + permissions: + contents: 'read' + id-token: 'write' + issues: 'write' + pull-requests: 'write' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Run Qwen Code CLI' + id: 'run_qwen' + uses: 'QwenLM/qwen-code-action@v1' # ratchet:exclude + env: + TITLE: '${{ github.event.pull_request.title || github.event.issue.title }}' + DESCRIPTION: '${{ github.event.pull_request.body || github.event.issue.body }}' + EVENT_NAME: '${{ github.event_name }}' + GITHUB_TOKEN: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + IS_PULL_REQUEST: '${{ !!github.event.pull_request }}' + ISSUE_NUMBER: '${{ github.event.pull_request.number || github.event.issue.number }}' + REPOSITORY: '${{ github.repository }}' + ADDITIONAL_CONTEXT: '${{ inputs.additional_context }}' + with: + openai_api_key: '${{ secrets.QWEN_API_KEY }}' + openai_base_url: '${{ vars.QWEN_BASE_URL }}' + openai_model: '${{ vars.QWEN_MODEL }}' + qwen_cli_version: '${{ vars.QWEN_CLI_VERSION }}' + qwen_debug: '${{ fromJSON(vars.DEBUG || vars.ACTIONS_STEP_DEBUG || false) }}' + upload_artifacts: '${{ vars.UPLOAD_ARTIFACTS }}' + workflow_name: 'qwen-invoke' + settings: |- + { + "model": { + "maxSessionTurns": 25 + }, + "telemetry": { + "enabled": true, + "target": "local", + "outfile": ".qwen/telemetry.log" + }, + "mcpServers": { + "github": { + "command": "docker", + "args": [ + "run", + "-i", + "--rm", + "-e", + "GITHUB_PERSONAL_ACCESS_TOKEN", + "ghcr.io/github/github-mcp-server:v0.18.0" + ], + "includeTools": [ + "add_issue_comment", + "get_issue", + "get_issue_comments", + "list_issues", + "search_issues", + "create_pull_request", + "pull_request_read", + "list_pull_requests", + "search_pull_requests", + "create_branch", + "create_or_update_file", + "delete_file", + "fork_repository", + "get_commit", + "get_file_contents", + "list_commits", + "push_files", + "search_code" + ], + "env": { + "GITHUB_PERSONAL_ACCESS_TOKEN": "${GITHUB_TOKEN}" + } + } + }, + "tools": { + "core": [ + "run_shell_command(cat)", + "run_shell_command(echo)", + "run_shell_command(grep)", + "run_shell_command(head)", + "run_shell_command(tail)" + ] + } + } + prompt: '/qwen-invoke' diff --git a/.github/workflows/qwen-review.yml b/.github/workflows/qwen-review.yml new file mode 100644 index 0000000000..f140e8c008 --- /dev/null +++ b/.github/workflows/qwen-review.yml @@ -0,0 +1,104 @@ +name: '🔎 Qwen Code Review' + +on: + workflow_call: + inputs: + additional_context: + type: 'string' + description: 'Any additional context from the request' + required: false + +concurrency: + group: '${{ github.workflow }}-review-${{ github.event_name }}-${{ github.event.pull_request.number || github.event.issue.number }}' + cancel-in-progress: true + +defaults: + run: + shell: 'bash' + +jobs: + review: + runs-on: 'ubuntu-latest' + timeout-minutes: 7 + permissions: + contents: 'read' + id-token: 'write' + issues: 'write' + pull-requests: 'write' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Checkout repository' + uses: 'actions/checkout@08c6903cd8c0fde910a37f88322edcfb5dd907a8' # ratchet:actions/checkout@v5 + + - name: 'Run Qwen Code pull request review' + uses: 'QwenLM/qwen-code-action@v1' # ratchet:exclude + id: 'qwen_pr_review' + env: + GITHUB_TOKEN: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + ISSUE_TITLE: '${{ github.event.pull_request.title || github.event.issue.title }}' + ISSUE_BODY: '${{ github.event.pull_request.body || github.event.issue.body }}' + PULL_REQUEST_NUMBER: '${{ github.event.pull_request.number || github.event.issue.number }}' + REPOSITORY: '${{ github.repository }}' + ADDITIONAL_CONTEXT: '${{ inputs.additional_context }}' + with: + openai_api_key: '${{ secrets.QWEN_API_KEY }}' + openai_base_url: '${{ vars.QWEN_BASE_URL }}' + openai_model: '${{ vars.QWEN_MODEL }}' + qwen_cli_version: '${{ vars.QWEN_CLI_VERSION }}' + qwen_debug: '${{ fromJSON(vars.DEBUG || vars.ACTIONS_STEP_DEBUG || false) }}' + upload_artifacts: '${{ vars.UPLOAD_ARTIFACTS }}' + workflow_name: 'qwen-review' + settings: |- + { + "model": { + "maxSessionTurns": 25 + }, + "telemetry": { + "enabled": true, + "target": "local", + "outfile": ".qwen/telemetry.log" + }, + "mcpServers": { + "github": { + "command": "docker", + "args": [ + "run", + "-i", + "--rm", + "-e", + "GITHUB_PERSONAL_ACCESS_TOKEN", + "ghcr.io/github/github-mcp-server:v0.18.0" + ], + "includeTools": [ + "add_comment_to_pending_review", + "create_pending_pull_request_review", + "pull_request_read", + "submit_pending_pull_request_review" + ], + "env": { + "GITHUB_PERSONAL_ACCESS_TOKEN": "${GITHUB_TOKEN}" + } + } + }, + "tools": { + "core": [ + "run_shell_command(cat)", + "run_shell_command(echo)", + "run_shell_command(grep)", + "run_shell_command(head)", + "run_shell_command(tail)" + ] + } + } + prompt: '/qwen-review' diff --git a/.github/workflows/qwen-scheduled-triage.yml b/.github/workflows/qwen-scheduled-triage.yml new file mode 100644 index 0000000000..930c63d498 --- /dev/null +++ b/.github/workflows/qwen-scheduled-triage.yml @@ -0,0 +1,208 @@ +name: '📋 Qwen Code Scheduled Issue Triage' + +on: + schedule: + - cron: '0 * * * *' # Runs every hour + pull_request: + branches: + - 'main' + - 'release/**/*' + paths: + - '.github/workflows/qwen-scheduled-triage.yml' + push: + branches: + - 'main' + - 'release/**/*' + paths: + - '.github/workflows/qwen-scheduled-triage.yml' + workflow_dispatch: + +concurrency: + group: '${{ github.workflow }}' + cancel-in-progress: true + +defaults: + run: + shell: 'bash' + +jobs: + triage: + runs-on: 'ubuntu-latest' + timeout-minutes: 7 + permissions: + contents: 'read' + id-token: 'write' + issues: 'read' + pull-requests: 'read' + outputs: + available_labels: '${{ steps.get_labels.outputs.available_labels }}' + triaged_issues: '${{ env.TRIAGED_ISSUES }}' + steps: + - name: 'Get repository labels' + id: 'get_labels' + uses: 'actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea' # ratchet:actions/github-script@v7.0.1 + with: + # NOTE: we intentionally do not use the minted token. The default + # GITHUB_TOKEN provided by the action has enough permissions to read + # the labels. + script: |- + const labels = []; + for await (const response of github.paginate.iterator(github.rest.issues.listLabelsForRepo, { + owner: context.repo.owner, + repo: context.repo.repo, + per_page: 100, // Maximum per page to reduce API calls + })) { + labels.push(...response.data); + } + + if (!labels || labels.length === 0) { + core.setFailed('There are no issue labels in this repository.') + } + + const labelNames = labels.map(label => label.name).sort(); + core.setOutput('available_labels', labelNames.join(',')); + core.info(`Found ${labelNames.length} labels: ${labelNames.join(', ')}`); + return labelNames; + + - name: 'Find untriaged issues' + id: 'find_issues' + env: + GITHUB_REPOSITORY: '${{ github.repository }}' + GITHUB_TOKEN: '${{ secrets.GITHUB_TOKEN || github.token }}' + run: |- + echo '🔍 Finding unlabeled issues and issues marked for triage...' + ISSUES="$(gh issue list \ + --state 'open' \ + --search 'no:label label:"status/needs-triage"' \ + --json number,title,body \ + --limit '100' \ + --repo "${GITHUB_REPOSITORY}" + )" + + echo '📝 Setting output for GitHub Actions...' + echo "issues_to_triage=${ISSUES}" >> "${GITHUB_OUTPUT}" + + ISSUE_COUNT="$(echo "${ISSUES}" | jq 'length')" + echo "✅ Found ${ISSUE_COUNT} issue(s) to triage! 🎯" + + - name: 'Run Qwen Code Issue Analysis' + id: 'qwen_issue_analysis' + if: |- + ${{ steps.find_issues.outputs.issues_to_triage != '[]' }} + uses: 'QwenLM/qwen-code-action@v1' # ratchet:exclude + env: + GITHUB_TOKEN: '' # Do not pass any auth token here since this runs on untrusted inputs + ISSUES_TO_TRIAGE: '${{ steps.find_issues.outputs.issues_to_triage }}' + REPOSITORY: '${{ github.repository }}' + AVAILABLE_LABELS: '${{ steps.get_labels.outputs.available_labels }}' + with: + openai_api_key: '${{ secrets.QWEN_API_KEY }}' + openai_base_url: '${{ vars.QWEN_BASE_URL }}' + openai_model: '${{ vars.QWEN_MODEL }}' + qwen_cli_version: '${{ vars.QWEN_CLI_VERSION }}' + qwen_debug: '${{ fromJSON(vars.DEBUG || vars.ACTIONS_STEP_DEBUG || false) }}' + upload_artifacts: '${{ vars.UPLOAD_ARTIFACTS }}' + workflow_name: 'qwen-scheduled-triage' + settings: |- + { + "model": { + "maxSessionTurns": 25 + }, + "telemetry": { + "enabled": true, + "target": "local", + "outfile": ".qwen/telemetry.log" + }, + "tools": { + "core": [ + "run_shell_command(echo)", + "run_shell_command(jq)", + "run_shell_command(printenv)" + ] + } + } + prompt: '/qwen-scheduled-triage' + + label: + runs-on: 'ubuntu-latest' + needs: + - 'triage' + if: |- + needs.triage.outputs.available_labels != '' && + needs.triage.outputs.available_labels != '[]' && + needs.triage.outputs.triaged_issues != '' && + needs.triage.outputs.triaged_issues != '[]' + permissions: + contents: 'read' + issues: 'write' + pull-requests: 'write' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Apply labels' + env: + AVAILABLE_LABELS: '${{ needs.triage.outputs.available_labels }}' + TRIAGED_ISSUES: '${{ needs.triage.outputs.triaged_issues }}' + uses: 'actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea' # ratchet:actions/github-script@v7.0.1 + with: + # Use the provided token so that the "gemini-cli" is the actor in the + # log for what changed the labels. + github-token: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + script: |- + // Parse the available labels + const availableLabels = (process.env.AVAILABLE_LABELS || '').split(',') + .map((label) => label.trim()) + .sort() + + // Parse out the triaged issues + const triagedIssues = (JSON.parse(process.env.TRIAGED_ISSUES || '{}')) + .sort((a, b) => a.issue_number - b.issue_number) + + core.debug(`Triaged issues: ${JSON.stringify(triagedIssues)}`); + + // Iterate over each label + for (const issue of triagedIssues) { + if (!issue) { + core.debug(`Skipping empty issue: ${JSON.stringify(issue)}`); + continue; + } + + const issueNumber = issue.issue_number; + if (!issueNumber) { + core.debug(`Skipping issue with no data: ${JSON.stringify(issue)}`); + continue; + } + + // Extract and reject invalid labels - we do this just in case + // someone was able to prompt inject malicious labels. + let labelsToSet = (issue.labels_to_set || []) + .map((label) => label.trim()) + .filter((label) => availableLabels.includes(label)) + .sort() + + core.debug(`Identified labels to set: ${JSON.stringify(labelsToSet)}`); + + if (labelsToSet.length === 0) { + core.info(`Skipping issue #${issueNumber} - no labels to set.`) + continue; + } + + core.debug(`Setting labels on issue #${issueNumber} to ${labelsToSet.join(', ')} (${issue.explanation || 'no explanation'})`) + + await github.rest.issues.setLabels({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: issueNumber, + labels: labelsToSet, + }); + } diff --git a/.github/workflows/qwen-triage.yml b/.github/workflows/qwen-triage.yml new file mode 100644 index 0000000000..cfea89d2a7 --- /dev/null +++ b/.github/workflows/qwen-triage.yml @@ -0,0 +1,152 @@ +name: '🔀 Qwen Code Triage' + +on: + workflow_call: + inputs: + additional_context: + type: 'string' + description: 'Any additional context from the request' + required: false + +concurrency: + group: '${{ github.workflow }}-triage-${{ github.event_name }}-${{ github.event.pull_request.number || github.event.issue.number }}' + cancel-in-progress: true + +defaults: + run: + shell: 'bash' + +jobs: + triage: + runs-on: 'ubuntu-latest' + timeout-minutes: 7 + outputs: + available_labels: '${{ steps.get_labels.outputs.available_labels }}' + selected_labels: '${{ env.SELECTED_LABELS }}' + permissions: + contents: 'read' + id-token: 'write' + issues: 'read' + pull-requests: 'read' + steps: + - name: 'Get repository labels' + id: 'get_labels' + uses: 'actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea' # ratchet:actions/github-script@v7.0.1 + with: + # NOTE: we intentionally do not use the given token. The default + # GITHUB_TOKEN provided by the action has enough permissions to read + # the labels. + script: |- + const labels = []; + for await (const response of github.paginate.iterator(github.rest.issues.listLabelsForRepo, { + owner: context.repo.owner, + repo: context.repo.repo, + per_page: 100, // Maximum per page to reduce API calls + })) { + labels.push(...response.data); + } + + if (!labels || labels.length === 0) { + core.setFailed('There are no issue labels in this repository.') + } + + const labelNames = labels.map(label => label.name).sort(); + core.setOutput('available_labels', labelNames.join(',')); + core.info(`Found ${labelNames.length} labels: ${labelNames.join(', ')}`); + return labelNames; + + - name: 'Run Qwen Code issue analysis' + id: 'qwen_analysis' + if: |- + ${{ steps.get_labels.outputs.available_labels != '' }} + uses: 'QwenLM/qwen-code-action@v1' # ratchet:exclude + env: + GITHUB_TOKEN: '' # Do NOT pass any auth tokens here since this runs on untrusted inputs + ISSUE_TITLE: '${{ github.event.issue.title }}' + ISSUE_BODY: '${{ github.event.issue.body }}' + AVAILABLE_LABELS: '${{ steps.get_labels.outputs.available_labels }}' + with: + openai_api_key: '${{ secrets.QWEN_API_KEY }}' + openai_base_url: '${{ vars.QWEN_BASE_URL }}' + openai_model: '${{ vars.QWEN_MODEL }}' + qwen_cli_version: '${{ vars.QWEN_CLI_VERSION }}' + qwen_debug: '${{ fromJSON(vars.DEBUG || vars.ACTIONS_STEP_DEBUG || false) }}' + upload_artifacts: '${{ vars.UPLOAD_ARTIFACTS }}' + workflow_name: 'qwen-triage' + settings: |- + { + "model": { + "maxSessionTurns": 25 + }, + "telemetry": { + "enabled": true, + "target": "local", + "outfile": ".qwen/telemetry.log" + }, + "tools": { + "core": [ + "run_shell_command(echo)" + ] + } + } + prompt: '/qwen-triage' + + label: + runs-on: 'ubuntu-latest' + needs: + - 'triage' + if: |- + ${{ needs.triage.outputs.selected_labels != '' }} + permissions: + contents: 'read' + issues: 'write' + pull-requests: 'write' + steps: + - name: 'Mint identity token' + id: 'mint_identity_token' + if: |- + ${{ vars.APP_ID }} + uses: 'actions/create-github-app-token@a8d616148505b5069dccd32f177bb87d7f39123b' # ratchet:actions/create-github-app-token@v2 + with: + app-id: '${{ vars.APP_ID }}' + private-key: '${{ secrets.APP_PRIVATE_KEY }}' + permission-contents: 'read' + permission-issues: 'write' + permission-pull-requests: 'write' + + - name: 'Apply labels' + env: + ISSUE_NUMBER: '${{ github.event.issue.number }}' + AVAILABLE_LABELS: '${{ needs.triage.outputs.available_labels }}' + SELECTED_LABELS: '${{ needs.triage.outputs.selected_labels }}' + uses: 'actions/github-script@60a0d83039c74a4aee543508d2ffcb1c3799cdea' # ratchet:actions/github-script@v7.0.1 + with: + # Use the provided token so that the "gemini-cli" is the actor in the + # log for what changed the labels. + github-token: '${{ steps.mint_identity_token.outputs.token || secrets.GITHUB_TOKEN || github.token }}' + script: |- + // Parse the available labels + const availableLabels = (process.env.AVAILABLE_LABELS || '').split(',') + .map((label) => label.trim()) + .sort() + + // Parse the label as a CSV, reject invalid ones - we do this just + // in case someone was able to prompt inject malicious labels. + const selectedLabels = (process.env.SELECTED_LABELS || '').split(',') + .map((label) => label.trim()) + .filter((label) => availableLabels.includes(label)) + .sort() + + // Set the labels + const issueNumber = process.env.ISSUE_NUMBER; + if (selectedLabels && selectedLabels.length > 0) { + await github.rest.issues.setLabels({ + owner: context.repo.owner, + repo: context.repo.repo, + issue_number: issueNumber, + labels: selectedLabels, + }); + core.info(`Successfully set labels: ${selectedLabels.join(',')}`); + } else { + core.info(`Failed to determine labels to set. There may not be enough information in the issue or pull request.`) + } From c024bd81f961260a96efa2d62a7a84f69624ce87 Mon Sep 17 00:00:00 2001 From: pelebel Date: Wed, 19 Aug 2026 14:13:23 -0400 Subject: [PATCH 12/20] feat(windows): build and run the engine on Windows 11 (MSVC + CUDA) Build and run ninfer / ninfer-serve with cl (VS2022) + nvcc (CUDA 13.x) + Ninja, resolving FFmpeg and libcurl from a vcpkg x64-windows prefix. - CMake: per-language MSVC flags (/Zc:preprocessor, NOMINMAX, UTF8PROC_STATIC) so CUDA objects keep their cache; new cmake/NInferMediaDeps.cmake locates FFmpeg/curl on Windows and copies their runtime DLLs next to the executables at build time. - Portability shims: Winsock (media_acquire), CreateFileW/MapViewOfFile artifact reader, localtime_s / GetCurrentProcessId / _isatty. - Pass the 128-byte NVFP4 TMA descriptors by device pointer so cl accepts the aligned by-value parameters (linear / linear-swiglu kernels). - Give the qwen3_6 plan move constructors explicit bodies: MSVC does not emit an out-of-line '= default' explicit specialization unless ODR-used. --- CMakeLists.txt | 31 ++++- apps/CMakeLists.txt | 21 +++ cmake/NInferMediaDeps.cmake | 124 ++++++++++++++++++ src/CMakeLists.txt | 12 +- src/artifact/reader.cpp | 81 +++++++++++- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu | 56 +++++++- src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh | 10 +- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu | 54 +++++++- .../nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh | 20 ++- src/product/load_progress/load_progress.cpp | 14 +- src/product/media_acquire/acquire.cpp | 28 +++- src/serve/console_log.cpp | 4 + src/serve/request_log.cpp | 16 ++- src/targets/qwen3_6/impl/runtime/api_impl.h | 14 +- tests/test_request_log.cpp | 16 ++- 15 files changed, 466 insertions(+), 35 deletions(-) create mode 100644 cmake/NInferMediaDeps.cmake diff --git a/CMakeLists.txt b/CMakeLists.txt index ca3f6c48e3..dec864d721 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -55,14 +55,37 @@ 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) + +# FFmpeg and libcurl are located cross-platform: pkg-config on POSIX, an +# install prefix / vcpkg tree on Windows. Both are exposed as NInfer::FFmpeg +# and NInfer::Curl (see cmake/NInferMediaDeps.cmake). +list(APPEND CMAKE_MODULE_PATH "${PROJECT_SOURCE_DIR}/cmake") +include(NInferMediaDeps) +ninfer_find_ffmpeg() if(NINFER_BUILD_MEDIA_ACQUIRE) - pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + ninfer_find_curl() endif() + find_package(Threads REQUIRED) +if(MSVC) + # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. The + # flag is forwarded to the cl host compiler nvcc invokes, including for the + # host stubs it generates. (The 128-byte-aligned TMA descriptor parameters are + # passed by device pointer rather than by value precisely so MSVC's stub + # compiler does not reject them.) + add_compile_options($<$:/Zc:preprocessor>) + set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=/Zc:preprocessor") + # windows.h defines min/max as macros which clobber std::min/std::max; NOMINMAX + # suppresses them. UTF8PROC_STATIC keeps the vendored utf8proc header from + # marking its symbols __declspec(dllimport) since we compile it statically. Both + # are scoped to C/CXX (no CUDA file needs them) so the CUDA objects keep their + # cache. + add_compile_definitions( + $<$:NOMINMAX> + $<$:UTF8PROC_STATIC>) +endif() + # --- Subdirectories ---------------------------------------------------------- add_subdirectory(src) if(NINFER_BUILD_APPS) diff --git a/apps/CMakeLists.txt b/apps/CMakeLists.txt index fe80287518..8d95296894 100644 --- a/apps/CMakeLists.txt +++ b/apps/CMakeLists.txt @@ -17,3 +17,24 @@ 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_load_progress) + +# On Windows the FFmpeg/libcurl runtime DLLs (from the vcpkg tree or an explicit +# prefix) are not on the default search path, so copy them next to each +# executable to make the build-tree binaries run in place. +if(WIN32) + set(_ninfer_runtime_dll_dirs "") + foreach(_prefix ${NINFER_FFMPEG_PREFIX} ${NINFER_CURL_PREFIX}) + if(_prefix AND EXISTS "${_prefix}/bin") + list(APPEND _ninfer_runtime_dll_dirs "${_prefix}/bin") + endif() + endforeach() + list(REMOVE_DUPLICATES _ninfer_runtime_dll_dirs) + foreach(_app ninfer ninfer-serve) + foreach(_dll_dir IN LISTS _ninfer_runtime_dll_dirs) + add_custom_command(TARGET ${_app} POST_BUILD + COMMAND ${CMAKE_COMMAND} -E copy_directory "${_dll_dir}" + "$" + COMMENT "Copying ${_app} runtime dependencies") + endforeach() + endforeach() +endif() diff --git a/cmake/NInferMediaDeps.cmake b/cmake/NInferMediaDeps.cmake new file mode 100644 index 0000000000..2b0d3bd52c --- /dev/null +++ b/cmake/NInferMediaDeps.cmake @@ -0,0 +1,124 @@ +# NInferMediaDeps.cmake +# +# Locates FFmpeg and libcurl and exposes them as INTERFACE imported targets +# NInfer::FFmpeg and NInfer::Curl, which the ninfer_media_* libraries link. +# +# * POSIX : discovered through pkg-config (unchanged behaviour). +# * Windows : located from a vcpkg install tree or an explicit prefix, because +# pkg-config is not part of the Windows toolchain. The prefix is +# chosen, per package, in this order: NINFER_FFMPEG_ROOT / +# NINFER_CURL_ROOT (per package), then NINFER_MEDIA_ROOT, then +# $ENV{VCPKG_ROOT}/installed/. +# +# On Windows the resolved prefixes are cached as NINFER_FFMPEG_PREFIX and +# NINFER_CURL_PREFIX so the app targets can copy the dependency DLLs next to the +# built executables at build time. + +if(NOT WIN32) + find_package(PkgConfig REQUIRED) +endif() + +function(ninfer_media_resolve_prefix _out) + set(result "") + if(DEFINED NINFER_MEDIA_ROOT) + set(result "${NINFER_MEDIA_ROOT}") + elseif(DEFINED ENV{NINFER_MEDIA_ROOT}) + set(result "$ENV{NINFER_MEDIA_ROOT}") + elseif(DEFINED ENV{VCPKG_ROOT}) + set(triplet "x64-windows") + if(DEFINED ENV{VCPKG_TARGET_TRIPLET}) + set(triplet "$ENV{VCPKG_TARGET_TRIPLET}") + endif() + if(EXISTS "$ENV{VCPKG_ROOT}/installed/${triplet}/include") + set(result "$ENV{VCPKG_ROOT}/installed/${triplet}") + endif() + endif() + set(${_out} "${result}" PARENT_SCOPE) +endfunction() + +function(ninfer_require_prefix _label _root _hint) + if(NOT _root OR NOT EXISTS "${_root}/include") + message(FATAL_ERROR + "${_label} not found. Point the build at an install prefix containing " + "include/ and lib/. ${_hint}") + endif() +endfunction() + +# Bundle an include directory and a link list (imported libs plus, on Windows, +# the system libraries the package needs) into an INTERFACE imported target. +function(ninfer_make_imported_target interface include_dir libs) + add_library(${interface} INTERFACE IMPORTED) + set_target_properties(${interface} PROPERTIES + INTERFACE_INCLUDE_DIRECTORIES "${include_dir}" + INTERFACE_LINK_LIBRARIES "${libs}") +endfunction() + +function(ninfer_find_ffmpeg) + if(TARGET NInfer::FFmpeg) + return() + endif() + if(NOT WIN32) + pkg_check_modules(_ninfer_ffmpeg REQUIRED IMPORTED_TARGET + libavformat>=60 libavcodec>=60 libavutil>=58 libswscale>=7) + ninfer_make_imported_target(NInfer::FFmpeg + "${_ninfer_ffmpeg_INCLUDE_DIRS}" "${_ninfer_ffmpeg_LIBRARIES}") + return() + endif() + + if(DEFINED NINFER_FFMPEG_ROOT) + set(root "${NINFER_FFMPEG_ROOT}") + elseif(DEFINED ENV{NINFER_FFMPEG_ROOT}) + set(root "$ENV{NINFER_FFMPEG_ROOT}") + else() + ninfer_media_resolve_prefix(root) + endif() + ninfer_require_prefix("FFmpeg" "${root}" + "set NINFER_FFMPEG_ROOT or NINFER_MEDIA_ROOT to an FFmpeg prefix, or install it " + "with `vcpkg install ffmpeg --triplet x64-windows` and set VCPKG_ROOT.") + + find_path(NINFER_FFMPEG_INCLUDE_DIR libavcodec/avcodec.h + PATHS "${root}/include" NO_DEFAULT_PATH REQUIRED) + foreach(module IN ITEMS avformat avcodec swscale avutil) + find_library(NINFER_FFMPEG_${module}_LIB NAMES ${module} + PATHS "${root}/lib" NO_DEFAULT_PATH REQUIRED) + endforeach() + + # The FFmpeg import libraries reference the Windows multimedia/network API. + # (The list contains only import libraries that exist in the Windows SDK; e.g. + # there is no Ntmapi.lib, and Nt*/Rtl* symbols, if ever needed, come from ntdll.) + ninfer_make_imported_target(NInfer::FFmpeg "${NINFER_FFMPEG_INCLUDE_DIR}" + "${NINFER_FFMPEG_avformat_LIB};${NINFER_FFMPEG_avcodec_LIB};${NINFER_FFMPEG_swscale_LIB};${NINFER_FFMPEG_avutil_LIB};Strmiids;Ws2_32;Advapi32;Ole32;Oleaut32;User32;Iphlpapi;Userenv") + set(NINFER_FFMPEG_PREFIX "${root}" CACHE INTERNAL "Resolved FFmpeg install prefix") +endfunction() + +function(ninfer_find_curl) + if(TARGET NInfer::Curl) + return() + endif() + if(NOT WIN32) + pkg_check_modules(_ninfer_curl REQUIRED IMPORTED_TARGET libcurl>=7.85) + ninfer_make_imported_target(NInfer::Curl + "${_ninfer_curl_INCLUDE_DIRS}" "${_ninfer_curl_LIBRARIES}") + return() + endif() + + if(DEFINED NINFER_CURL_ROOT) + set(root "${NINFER_CURL_ROOT}") + elseif(DEFINED ENV{NINFER_CURL_ROOT}) + set(root "$ENV{NINFER_CURL_ROOT}") + else() + ninfer_media_resolve_prefix(root) + endif() + ninfer_require_prefix("libcurl" "${root}" + "set NINFER_CURL_ROOT or NINFER_MEDIA_ROOT to a libcurl prefix, or install it " + "with `vcpkg install curl --triplet x64-windows` and set VCPKG_ROOT.") + + find_path(NINFER_CURL_INCLUDE_DIR curl/curl.h + PATHS "${root}/include" NO_DEFAULT_PATH REQUIRED) + find_library(NINFER_CURL_LIB NAMES curl libcurl + PATHS "${root}/lib" NO_DEFAULT_PATH REQUIRED) + + ninfer_make_imported_target(NInfer::Curl "${NINFER_CURL_INCLUDE_DIR}" + "${NINFER_CURL_LIB};Ws2_32;Crypt32;Cryptui;Normaliz;Bcrypt;Advapi32;Userenv") + set(NINFER_CURL_PREFIX "${root}" CACHE INTERNAL "Resolved libcurl install prefix") +endfunction() \ No newline at end of file diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index f5590f3f77..33167d61cc 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -273,14 +273,18 @@ ninfer_internal_includes(ninfer_text) add_library(ninfer_media_decode STATIC media/decode/decode.cpp) ninfer_internal_includes(ninfer_media_decode) -target_link_libraries(ninfer_media_decode PRIVATE PkgConfig::FFMPEG) +target_link_libraries(ninfer_media_decode PRIVATE NInfer::FFmpeg) if(NINFER_BUILD_MEDIA_ACQUIRE) # Product-only path/data/HTTP acquisition. No target package links this library. add_library(ninfer_media_acquire STATIC product/media_acquire/acquire.cpp) ninfer_internal_includes(ninfer_media_acquire) - target_link_libraries(ninfer_media_acquire PRIVATE PkgConfig::LIBCURL) + target_link_libraries(ninfer_media_acquire PRIVATE NInfer::Curl) + if(WIN32) + # Host name resolution / Winsock (getaddrinfo, inet_ntop, WSAStartup). + target_link_libraries(ninfer_media_acquire PRIVATE Ws2_32) + endif() endif() if(NINFER_BUILD_PROMPT_INPUT) @@ -340,4 +344,8 @@ if(NINFER_BUILD_SERVE) target_link_libraries(ninfer_serve PUBLIC ninfer_engine Threads::Threads PRIVATE ninfer_media_acquire CUDA::cudart) + if(WIN32) + # cpp-httplib (plain HTTP) uses the Winsock API. + target_link_libraries(ninfer_serve PRIVATE Ws2_32) + endif() endif() diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 1dc3afd1ea..4f61bd2237 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -15,10 +15,18 @@ #include #include +#ifdef _WIN32 +#ifndef NOMINMAX +#define NOMINMAX +#endif +#define WIN32_LEAN_AND_MEAN +#include +#else #include #include #include #include +#endif namespace ninfer::artifact { namespace { @@ -181,6 +189,55 @@ struct TransparentStringHash { class MappedFile { public: explicit MappedFile(const std::filesystem::path& path) { +#ifdef _WIN32 + const std::wstring wide = path.wstring(); + HANDLE file = ::CreateFileW(wide.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, + OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, nullptr); + if (file == INVALID_HANDLE_VALUE) { + const auto error = ::GetLastError(); + throw std::system_error(error, std::system_category(), "open " + path.string()); + } + + LARGE_INTEGER size_info {}; + + if (!::GetFileSizeEx(file, &size_info)) { + const auto error = ::GetLastError(); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), "GetFileSizeEx " + path.string()); + } + if (size_info.QuadPart < 0 || + size_info.QuadPart > static_cast(std::numeric_limits::max())) { + ::CloseHandle(file); + throw ArtifactError("artifact size does not fit the process address space"); + } + + const auto size = static_cast(size_info.QuadPart); + HANDLE mapping = nullptr; + void* view = nullptr; + if (size != 0) { + mapping = ::CreateFileMappingW(file, nullptr, PAGE_READONLY, 0, 0, nullptr); + if (mapping == nullptr) { + const auto error = ::GetLastError(); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), + "CreateFileMapping " + path.string()); + } + // A zero byte count maps the entire section, which avoids the 32-bit + // dwNumberOfBytesToMap limit for artifacts larger than 4 GiB. + view = ::MapViewOfFile(mapping, FILE_MAP_READ, 0, 0, 0); + if (view == nullptr) { + const auto error = ::GetLastError(); + ::CloseHandle(mapping); + ::CloseHandle(file); + throw std::system_error(error, std::system_category(), "MapViewOfFile " + path.string()); + } + } + file_ = file; + mapping_ = mapping; + view_ = view; + data_ = static_cast(view); + 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()); @@ -212,11 +269,18 @@ class MappedFile { fd_ = fd; data_ = static_cast(mapping); size_ = size; +#endif } ~MappedFile() { +#ifdef _WIN32 + if (view_ != nullptr) { ::UnmapViewOfFile(view_); } + if (mapping_ != nullptr) { ::CloseHandle(mapping_); } + if (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; @@ -232,6 +296,14 @@ class MappedFile { reinterpret_cast(destination.data()) % alignment != 0) { throw ArtifactError("direct artifact read is not 4096-byte aligned"); } +#ifdef _WIN32 + // The whole file is already mapped, so a direct read is a copy out of the view. + if (absolute_offset > size_ || destination.size() > size_ - absolute_offset) { + throw ArtifactError("direct artifact read exceeds the mapped file"); + } + std::memcpy(destination.data(), data_ + absolute_offset, destination.size()); + return destination.size(); +#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"); @@ -246,10 +318,17 @@ class MappedFile { throw std::system_error(errno, std::generic_category(), "direct artifact read"); } return static_cast(bytes); +#endif } private: - int fd_ = -1; +#ifdef _WIN32 + HANDLE file_ = INVALID_HANDLE_VALUE; + HANDLE mapping_ = nullptr; + void* view_ = nullptr; +#else + int fd_ = -1; +#endif const std::byte* data_ = nullptr; std::size_t size_ = 0; }; diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu index 19dacec69f..9b7e1e167c 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu @@ -9,6 +9,7 @@ #include #include +#include namespace ninfer::ops::detail { namespace { @@ -16,6 +17,54 @@ namespace { using TmaM256N128 = Nvfp4W4a4TmaSchedule<256, 3, 1>; using TmaM256N128S2 = Nvfp4W4a4TmaSchedule<256, 2, 1>; +// MSVC's cl rejects a by-value 128-byte-aligned TMA descriptor parameter in the +// nvcc-generated kernel host stub (C2719); Clang/GCC accept it. So the kernels +// take a device pointer. One persistent host copy (a stable source for the HtoD +// memcpy node captured into decode CUDA graphs) and one device copy are kept per +// distinct descriptor; a descriptor is constant for a given (buffers, tokens), +// so the device copy stays valid across graph replays. +struct Nvfp4TmaDescKey { + const void* activation_codes = nullptr; + const void* activation_scales = nullptr; + const void* weight_codes = nullptr; + const void* weight_scales = nullptr; + std::int32_t tokens = 0; + + bool operator==(const Nvfp4TmaDescKey& o) const { + return activation_codes == o.activation_codes && + activation_scales == o.activation_scales && weight_codes == o.weight_codes && + weight_scales == o.weight_scales && tokens == o.tokens; + } +}; + +struct Nvfp4TmaDescSlot { + Nvfp4TmaDescKey key; + Nvfp4W4a4TmaDescriptors* host = nullptr; + Nvfp4W4a4TmaDescriptors* device = nullptr; +}; + +Nvfp4W4a4TmaDescriptors* nvfp4_tma_desc_device(const Nvfp4TmaDescKey& key, + const Nvfp4W4a4TmaDescriptors& host_desc, + cudaStream_t stream) { + static std::vector slots; + for (Nvfp4TmaDescSlot& slot : slots) { + if (slot.key == key) { + *slot.host = host_desc; + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + return slot.device; + } + } + Nvfp4TmaDescSlot slot; + slot.key = key; + slot.host = new Nvfp4W4a4TmaDescriptors(host_desc); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&slot.device), sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + slots.push_back(slot); + return slot.device; +} + constexpr std::int32_t kQueryRows = 6144; constexpr std::int32_t kKeyRows = 1024; constexpr std::int32_t kGateRows = 6144; @@ -70,9 +119,14 @@ void launch_tma(const std::uint8_t* activation_codes, const std::uint8_t* activa }(); (void)kConfigured; + const Nvfp4W4a4TmaDescriptors* device_descriptors = nvfp4_tma_desc_device( + Nvfp4TmaDescKey{activation_codes, activation_scales, weight_codes, weight_scales, tokens}, + descriptors, stream); + const dim3 grid(Geometry::kOutputRows / Schedule::kBlockN, tokens / Schedule::kBlockM); nvfp4_w4a4_tma_kernel - <<>>(descriptors, alpha, epilogue, output); + <<>>(device_descriptors, alpha, epilogue, + output); CUDA_CHECK(cudaGetLastError()); } diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh index aa6914eca5..694a3d1fd1 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh @@ -182,7 +182,7 @@ __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 Nvfp4W4a4TmaDescriptors* descriptors, 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); @@ -222,17 +222,17 @@ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4 nvfp4_mbarrier_arrive_expect_tx(&shared.full[stage], kTransactionBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors->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], &descriptors->b_codes, k_tile * Schedule::kCodeRowBytes, row_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors.a_scales, (k_tile / 2) * 16, + nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors->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], &descriptors->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..e90906345e 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 @@ -8,12 +8,61 @@ #include #include #include +#include namespace ninfer::ops::detail { namespace { using M256N128S3 = Nvfp4W4a4TmaSchedule<256, 3, 1>; +// MSVC's cl rejects a by-value 128-byte-aligned TMA descriptor parameter in the +// nvcc-generated kernel host stub (C2719); Clang/GCC accept it. So the kernel +// takes a device pointer. One persistent host copy (a stable source for the HtoD +// memcpy node captured into decode CUDA graphs) and one device copy are kept per +// distinct descriptor; a descriptor is constant for a given (buffers, tokens), +// so the device copy stays valid across graph replays. +struct Nvfp4TmaDescKey { + const void* activation_codes = nullptr; + const void* activation_scales = nullptr; + const void* weight_codes = nullptr; + const void* weight_scales = nullptr; + std::int32_t tokens = 0; + + bool operator==(const Nvfp4TmaDescKey& o) const { + return activation_codes == o.activation_codes && + activation_scales == o.activation_scales && weight_codes == o.weight_codes && + weight_scales == o.weight_scales && tokens == o.tokens; + } +}; + +struct Nvfp4TmaDescSlot { + Nvfp4TmaDescKey key; + Nvfp4W4a4TmaDescriptors* host = nullptr; + Nvfp4W4a4TmaDescriptors* device = nullptr; +}; + +Nvfp4W4a4TmaDescriptors* nvfp4_tma_desc_device(const Nvfp4TmaDescKey& key, + const Nvfp4W4a4TmaDescriptors& host_desc, + cudaStream_t stream) { + static std::vector slots; + for (Nvfp4TmaDescSlot& slot : slots) { + if (slot.key == key) { + *slot.host = host_desc; + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + return slot.device; + } + } + Nvfp4TmaDescSlot slot; + slot.key = key; + slot.host = new Nvfp4W4a4TmaDescriptors(host_desc); + CUDA_CHECK(cudaMalloc(reinterpret_cast(&slot.device), sizeof(Nvfp4W4a4TmaDescriptors))); + CUDA_CHECK(cudaMemcpyAsync(slot.device, slot.host, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); + slots.push_back(slot); + return slot.device; +} + template Nvfp4W4a4TmaDescriptors make_descriptors(const std::uint8_t* activation_codes, const std::uint8_t* activation_scales, @@ -70,9 +119,12 @@ 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 Nvfp4W4a4TmaDescriptors* device_descriptors = nvfp4_tma_desc_device( + Nvfp4TmaDescKey{activation_codes, activation_scales, weight_codes, weight_scales, tokens}, + descriptors, stream); const dim3 grid((Geometry::kOutputRows / 2) / kPairN, tokens / M256N128S3::kBlockM); nvfp4_linear_swiglu_w4a4_tma_kernel - <<>>(descriptors, alpha, output); + <<>>(device_descriptors, alpha, output); 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 a7664c7da7..65f0354e0b 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,9 @@ 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) void nvfp4_linear_swiglu_w4a4_tma_kernel( + const Nvfp4W4a4TmaDescriptors* descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { static_assert(Geometry::kOutputRows == 34816); static_assert(Geometry::kInputRows == 5120); static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); @@ -98,16 +96,16 @@ __global__ __launch_bounds__( nvfp4_mbarrier_arrive_expect_tx(&shared.full[stage], kTransactionBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors->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], &descriptors->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, + &descriptors->b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin + kIntermediate, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors.a_scales, (k_tile / 2) * 16, + nvfp4_tma_load_2d(tensors.a_scale4[stage], &descriptors->a_scales, (k_tile / 2) * 16, token_begin, &shared.full[stage]); const int gate_scale_row = ((pair_begin / 128) * Geometry::kScaleTilesPerRow + @@ -117,9 +115,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], &descriptors->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], &descriptors->b_scales, 0, up_scale_row, &shared.full[stage]); } } diff --git a/src/product/load_progress/load_progress.cpp b/src/product/load_progress/load_progress.cpp index 2617b825b9..40d28bb6cb 100644 --- a/src/product/load_progress/load_progress.cpp +++ b/src/product/load_progress/load_progress.cpp @@ -1,10 +1,15 @@ #include "product/load_progress/load_progress.h" +#ifdef _WIN32 +#include +#else #include +#endif #include #include #include +#include #include #include #include @@ -60,7 +65,14 @@ std::string format_line(std::string_view phase, std::uint64_t done, std::uint64_ } // namespace LoadProgressRendererOptions stderr_load_progress_options() noexcept { - if (::isatty(STDERR_FILENO) == 1) { + const bool is_interactive = []() { +#ifdef _WIN32 + return _isatty(_fileno(stderr)) == 1; +#else + return ::isatty(STDERR_FILENO) == 1; +#endif + }(); + if (is_interactive) { return LoadProgressRendererOptions{ .mode = LoadProgressOutputMode::Interactive, .min_refresh_interval = std::chrono::milliseconds(200), diff --git a/src/product/media_acquire/acquire.cpp b/src/product/media_acquire/acquire.cpp index 1f03ac9c64..bf9d17c015 100644 --- a/src/product/media_acquire/acquire.cpp +++ b/src/product/media_acquire/acquire.cpp @@ -2,9 +2,14 @@ #include +#ifdef _WIN32 +#include +#include +#else #include #include #include +#endif #include #include @@ -93,17 +98,21 @@ bool private_ipv4(std::uint32_t address) { bool private_address(const sockaddr* address) { if (address->sa_family == AF_INET) { - return private_ipv4(reinterpret_cast(address)->sin_addr.s_addr); + std::uint32_t v4 = 0; + std::memcpy(&v4, &reinterpret_cast(address)->sin_addr, sizeof(v4)); + return private_ipv4(v4); } if (address->sa_family != AF_INET6) { return true; } const in6_addr& a = reinterpret_cast(address)->sin6_addr; + std::array raw{}; + std::memcpy(raw.data(), &a, sizeof(raw)); if (IN6_IS_ADDR_UNSPECIFIED(&a) || IN6_IS_ADDR_LOOPBACK(&a) || IN6_IS_ADDR_LINKLOCAL(&a) || - IN6_IS_ADDR_MULTICAST(&a) || (a.s6_addr[0] & 0xfeU) == 0xfcU) { + IN6_IS_ADDR_MULTICAST(&a) || (raw[0] & 0xfeU) == 0xfcU) { return true; } if (IN6_IS_ADDR_V4MAPPED(&a)) { std::uint32_t v4 = 0; - std::memcpy(&v4, &a.s6_addr[12], sizeof(v4)); + std::memcpy(&v4, raw.data() + 12, sizeof(v4)); return private_ipv4(v4); } return false; @@ -203,6 +212,12 @@ std::vector fetch_url(std::string url, const Policy& policy) { if (!policy.allow_remote) { throw std::invalid_argument("remote media URLs are disabled"); } static std::once_flag init; std::call_once(init, [] { +#ifdef _WIN32 + WSADATA wsa{}; + if (WSAStartup(MAKEWORD(2, 2), &wsa) != 0) { + throw std::runtime_error("failed to initialize Winsock"); + } +#endif if (curl_global_init(CURL_GLOBAL_DEFAULT) != CURLE_OK) { throw std::runtime_error("failed to initialize libcurl"); } @@ -287,7 +302,12 @@ 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 (ec || relative.empty() || relative.native().starts_with("..")) { + // A leading ".." component means the path escapes the root. Using the path + // iterator (rather than native().starts_with("..")) stays portable: on + // Windows native() is a wide string and would not match the narrow literal. + const bool escapes_root = + relative.begin() != relative.end() && *relative.begin() == ".."; + if (ec || relative.empty() || escapes_root) { throw std::invalid_argument("media path is outside configured media root"); } } diff --git a/src/serve/console_log.cpp b/src/serve/console_log.cpp index 7c58004793..7056c39d50 100644 --- a/src/serve/console_log.cpp +++ b/src/serve/console_log.cpp @@ -38,7 +38,11 @@ std::string format_console_log_prefix(std::chrono::system_clock::time_point time 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(&local, &wall_seconds); +#else localtime_r(&wall_seconds, &local); +#endif std::ostringstream out; out << '[' << std::put_time(&local, "%Y-%m-%d %H:%M:%S") << '.' << std::setfill('0') diff --git a/src/serve/request_log.cpp b/src/serve/request_log.cpp index b2dd984b69..57b2d393f2 100644 --- a/src/serve/request_log.cpp +++ b/src/serve/request_log.cpp @@ -15,13 +15,27 @@ #include #include +#ifdef _WIN32 +#define NOMINMAX +#define WIN32_LEAN_AND_MEAN +#include +#else #include +#endif namespace ninfer::serve { namespace { using Json = nlohmann::json; +std::uint32_t process_id() { +#ifdef _WIN32 + return GetCurrentProcessId(); +#else + return static_cast(::getpid()); +#endif +} + std::uint64_t unix_time_ms() { const auto now = std::chrono::system_clock::now().time_since_epoch(); return static_cast( @@ -31,7 +45,7 @@ std::uint64_t unix_time_ms() { std::string new_server_instance_id() { const auto now = std::chrono::system_clock::now().time_since_epoch(); const auto micros = std::chrono::duration_cast(now).count(); - return "serve-" + std::to_string(static_cast(::getpid())) + '-' + + return "serve-" + std::to_string(static_cast(process_id())) + '-' + std::to_string(micros); } diff --git a/src/targets/qwen3_6/impl/runtime/api_impl.h b/src/targets/qwen3_6/impl/runtime/api_impl.h index 0eadcde9bc..03797e9bf5 100644 --- a/src/targets/qwen3_6/impl/runtime/api_impl.h +++ b/src/targets/qwen3_6/impl/runtime/api_impl.h @@ -17,8 +17,15 @@ SequencePlan::SequencePlan( std::unique_ptr> impl) noexcept : impl_(std::move(impl)) {} +// MSVC does not emit an out-of-line '= default' explicit specialization unless +// it is ODR-used in this translation unit; another TU (the engine's +// std::optional usage) move-constructs these, so the move constructors must be +// emitted here or they are unresolved at link time on Windows. GCC/Clang emit +// '= default' unconditionally, which is why this only breaks the MSVC build. +// An explicit (still noexcept, still just moving the unique_ptr member) body is +// always emitted. The move-assignment and destructor keep '= default'. template <> -SequencePlan::SequencePlan(SequencePlan&&) noexcept = default; +SequencePlan::SequencePlan(SequencePlan&& other) noexcept : impl_(std::move(other.impl_)) {} template <> SequencePlan& SequencePlan::operator=(SequencePlan&&) noexcept = default; template <> @@ -85,7 +92,8 @@ 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; template <> @@ -102,7 +110,7 @@ RequestPlan::RequestPlan(std::unique_ptr -RequestPlan::RequestPlan(RequestPlan&&) noexcept = default; +RequestPlan::RequestPlan(RequestPlan&& other) noexcept : impl_(std::move(other.impl_)) {} template <> RequestPlan& RequestPlan::operator=(RequestPlan&&) noexcept = default; template <> diff --git a/tests/test_request_log.cpp b/tests/test_request_log.cpp index d557ace416..8e241f7d00 100644 --- a/tests/test_request_log.cpp +++ b/tests/test_request_log.cpp @@ -12,13 +12,27 @@ #include #include +#ifdef _WIN32 +#define NOMINMAX +#define WIN32_LEAN_AND_MEAN +#include +#else #include +#endif namespace { using namespace ninfer::serve; using Json = nlohmann::json; +std::uint32_t process_id() { +#ifdef _WIN32 + return GetCurrentProcessId(); +#else + return static_cast(::getpid()); +#endif +} + int check(bool condition, const char* message) { if (condition) { return 0; } std::cerr << message << '\n'; @@ -347,7 +361,7 @@ int main() { const std::filesystem::path log_path = std::filesystem::temp_directory_path() / - ("ninfer-request-log-test-" + std::to_string(static_cast(::getpid())) + + ("ninfer-request-log-test-" + std::to_string(static_cast(process_id())) + ".jsonl"); std::filesystem::remove(log_path); { From 1fecf28e997c06d054b9166c28f2eb980832e3eb Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 11:30:06 -0400 Subject: [PATCH 13/20] fix(win32): compare artifact size against size_t without signed cast The Windows mapping path casted size_t's maximum to LONGLONG before comparing it with the file size. That cast wraps to -1 on a 64-bit build, so every non-empty artifact compared greater and was rejected with "artifact size does not fit the process address space" before the mapping was ever attempted. Compare in unsigned terms instead, matching the POSIX path below. Co-Authored-By: Claude Opus 5 --- src/artifact/reader.cpp | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 4f61bd2237..918b015f02 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -5,6 +5,7 @@ #include #include #include +#include #include #include #include @@ -206,7 +207,8 @@ class MappedFile { throw std::system_error(error, std::system_category(), "GetFileSizeEx " + path.string()); } if (size_info.QuadPart < 0 || - size_info.QuadPart > static_cast(std::numeric_limits::max())) { + static_cast(size_info.QuadPart) > + std::numeric_limits::max()) { ::CloseHandle(file); throw ArtifactError("artifact size does not fit the process address space"); } From 9b41398049a9e1f1cef8b690da10ec96d2a785f0 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:10:16 -0400 Subject: [PATCH 14/20] fix(win32): mirror pread short-read semantics in the mapped direct read The Windows read_direct rejected any request running past the last byte, but direct I/O rounds every request up to the 4096-byte alignment, so the final read of an artifact overruns EOF by design. The POSIX path returns a short pread() count there and the materializer relies on it, comparing the result against min(request, remaining). Clamp to the mapped size and return the bytes actually copied, which fixes ninfer_artifact_materialization_test on Windows. Co-Authored-By: Claude Opus 5 --- src/artifact/reader.cpp | 11 ++++++++--- 1 file changed, 8 insertions(+), 3 deletions(-) diff --git a/src/artifact/reader.cpp b/src/artifact/reader.cpp index 918b015f02..12e4d80e6e 100644 --- a/src/artifact/reader.cpp +++ b/src/artifact/reader.cpp @@ -300,11 +300,16 @@ class MappedFile { } #ifdef _WIN32 // The whole file is already mapped, so a direct read is a copy out of the view. - if (absolute_offset > size_ || destination.size() > size_ - absolute_offset) { + if (absolute_offset > size_) { throw ArtifactError("direct artifact read exceeds the mapped file"); } - std::memcpy(destination.data(), data_ + absolute_offset, destination.size()); - return destination.size(); + // Direct I/O rounds every request up to the alignment, so the final one runs + // past the last byte by design. pread() answers that with a short count and + // callers rely on it, so clamp and report what was actually copied. + const auto offset = static_cast(absolute_offset); + const std::size_t available = std::min(destination.size(), size_ - offset); + std::memcpy(destination.data(), data_ + offset, available); + return available; #else if (absolute_offset > static_cast(std::numeric_limits::max()) || destination.size() > static_cast(std::numeric_limits::max())) { From 85c606ceb60314cde1e5a1cdb461b8d2a16f6574 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 12:10:27 -0400 Subject: [PATCH 15/20] test(win32): build and run the test suite under MSVC The suite was excluded from the default build, so it had never been compiled on Windows. Four gaps kept it from running: - std::aligned_alloc is absent from the MSVC CRT, and its aligned blocks need _aligned_free rather than free. - std::sqrt is not constexpr before C++26; libstdc++ accepts it as an extension, MSVC does not. These initializers only need const. - The FFmpeg/libcurl DLLs are copied next to the apps but not next to the tests, so every test linking them failed to start. Prepend the prefix bin directories to each test's PATH instead of duplicating the DLLs. - test_frontend reads an official tokenizer from a hard-coded local HF checkout and aborted when absent. Report the standard skip code, and register the test with SKIP_RETURN_CODE like its peers. ctest now runs 83 tests green with 6 skipped for missing local artifacts. Co-Authored-By: Claude Opus 5 --- tests/CMakeLists.txt | 26 ++++++++++++++++ .../test_gated_delta_net_replay_record.cpp | 2 +- tests/ops/test_gdn_replay_fold.cpp | 4 +-- tests/ops/test_vision_attention.cpp | 2 +- tests/targets/qwen3_6/test_frontend.cpp | 30 +++++++++++++++++-- tests/test_gdn_replay_records.cpp | 22 ++++++++++++-- 6 files changed, 77 insertions(+), 9 deletions(-) diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index 674da12a62..e7a680fc75 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -1,10 +1,33 @@ find_package(Python3 REQUIRED COMPONENTS Interpreter) +# On Windows the FFmpeg/libcurl runtime DLLs are not on the default search path. +# The app targets copy them next to their executables; tests instead get those +# directories prepended to PATH, so the suite runs in place without duplicating +# the DLLs into every test output directory. +set(NINFER_TEST_DLL_DIRS "") +if(WIN32) + foreach(_prefix ${NINFER_FFMPEG_PREFIX} ${NINFER_CURL_PREFIX}) + if(_prefix AND EXISTS "${_prefix}/bin") + file(TO_NATIVE_PATH "${_prefix}/bin" _native_bin) + list(APPEND NINFER_TEST_DLL_DIRS "${_native_bin}") + endif() + endforeach() + list(REMOVE_DUPLICATES NINFER_TEST_DLL_DIRS) +endif() + +function(ninfer_test_runtime_path name) + foreach(_dir IN LISTS NINFER_TEST_DLL_DIRS) + set_property(TEST ${name} APPEND PROPERTY + ENVIRONMENT_MODIFICATION "PATH=path_list_prepend:${_dir}") + endforeach() +endfunction() + # This target deliberately receives no src/, CUDA, artifact, kernel, or target # include root. It proves that the installed product headers stand alone. add_executable(ninfer_public_api_test test_public_api.cpp) target_include_directories(ninfer_public_api_test PRIVATE ${PROJECT_SOURCE_DIR}/include) add_test(NAME ninfer_public_api_test COMMAND ninfer_public_api_test) +ninfer_test_runtime_path(ninfer_public_api_test) function(ninfer_add_test name) cmake_parse_arguments(arg "NEEDS_SOURCE_DIR" "" "SOURCES;LIBRARIES" ${ARGN}) @@ -26,6 +49,7 @@ function(ninfer_add_test name) NINFER_PYTHON_EXECUTABLE="${Python3_EXECUTABLE}") endif() add_test(NAME ${name} COMMAND ${name}) + ninfer_test_runtime_path(${name}) endfunction() function(ninfer_add_op_test name) @@ -101,6 +125,8 @@ ninfer_add_test(ninfer_qwen3_6_frontend_test target_include_directories(ninfer_qwen3_6_frontend_test PRIVATE ${PROJECT_SOURCE_DIR}/src/targets/qwen3_6/export ${PROJECT_SOURCE_DIR}/src/targets/qwen3_6/impl) +# Needs a local Qwen3.6-27B HF checkout for the official tokenizer fixtures. +set_tests_properties(ninfer_qwen3_6_frontend_test PROPERTIES SKIP_RETURN_CODE 77) ninfer_add_test(ninfer_qwen3_6_runtime_mechanisms_test SOURCES targets/qwen3_6/test_runtime_mechanisms.cpp LIBRARIES ninfer_engine ninfer_core) diff --git a/tests/ops/test_gated_delta_net_replay_record.cpp b/tests/ops/test_gated_delta_net_replay_record.cpp index f2175dd04a..7e5e7d3708 100644 --- a/tests/ops/test_gated_delta_net_replay_record.cpp +++ b/tests/ops/test_gated_delta_net_replay_record.cpp @@ -115,7 +115,7 @@ int run_case(std::int32_t value_heads, std::int32_t width, std::int32_t batch, Tensor value_record_tensor(value_record.p, DType::BF16, {kStateDim, value_heads, width, batch}); Tensor gate_record_tensor(gate_record.p, DType::FP32, {2, value_heads, width, batch}); - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); ops::gated_delta_net_snapshot(q, k, v, g_tensor, beta_tensor, kScale, true, snapshot_states, valid, initial, bases, snapshot_output, nullptr); ops::gated_delta_net_replay_record(q, k, v, g_tensor, beta_tensor, kScale, record_states, valid, diff --git a/tests/ops/test_gdn_replay_fold.cpp b/tests/ops/test_gdn_replay_fold.cpp index 037f462f19..3dabcaa59f 100644 --- a/tests/ops/test_gdn_replay_fold.cpp +++ b/tests/ops/test_gdn_replay_fold.cpp @@ -318,7 +318,7 @@ int run_case(const FoldProfile profile, std::int32_t width, std::int32_t rows, Tensor output(out.p, DType::BF16, {kStateDim, profile.value_heads, width, 1}); Tensor initial_selector(initial_device.p, DType::I32, {1}); Tensor base_selector(base_device.p, DType::I32, {1}); - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); for (std::int32_t layer = 0; layer < profile.layers; ++layer) { const GdnReplayRecordLayer layer_records = records.layer(layer, rows); @@ -503,7 +503,7 @@ int run_record_fold_rounds() { constexpr std::int32_t kStateSlots = 3; constexpr std::int32_t kInitialSlot = 2; constexpr std::int32_t kSnapshotBase = 0; - constexpr float kScale = 1.0F / std::sqrt(128.0F); + const float kScale = 1.0F / std::sqrt(128.0F); DevicePackedWeight parent( quantized_weight::make_patterned_weight(QType::W8G32_F16S, kParentRows, kHidden, 1901U)); diff --git a/tests/ops/test_vision_attention.cpp b/tests/ops/test_vision_attention.cpp index 6c6ce630bb..5327533c21 100644 --- a/tests/ops/test_vision_attention.cpp +++ b/tests/ops/test_vision_attention.cpp @@ -41,7 +41,7 @@ std::vector bf16_bits(const std::vector& values) { void vision_attention_oracle(const std::vector& q, const std::vector& k, const std::vector& v, const std::vector& cu_seqlens, std::vector& out) { - constexpr double scale = 1.0 / std::sqrt(72.0); + const double scale = 1.0 / std::sqrt(72.0); out.assign(q.size(), 0.0); for (std::size_t segment = 0; segment + 1 < cu_seqlens.size(); ++segment) { diff --git a/tests/targets/qwen3_6/test_frontend.cpp b/tests/targets/qwen3_6/test_frontend.cpp index 4e29017dae..d9a8ff148a 100644 --- a/tests/targets/qwen3_6/test_frontend.cpp +++ b/tests/targets/qwen3_6/test_frontend.cpp @@ -17,6 +17,7 @@ #include #include #include +#include #include #include #include @@ -88,13 +89,31 @@ const fi::CompiledChatTemplate& reasoning_effort_template() { return value; } +// The official tokenizer fixtures are a local HF checkout, not part of this +// repository, so a checkout without them skips instead of failing. +constexpr const char* kOfficialTokenizerDir = + "/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16"; + +std::string official_tokenizer_file(const char* name) { + return (std::filesystem::path(kOfficialTokenizerDir) / name).string(); +} + +bool official_tokenizer_available() { + for (const char* name : + {"tokenizer.json", "tokenizer_config.json", "generation_config.json"}) { + std::error_code error; + if (!std::filesystem::exists(official_tokenizer_file(name), error)) { return false; } + } + return true; +} + const fi::Tokenizer& official_tokenizer() { static const std::string tokenizer_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/tokenizer.json"); + read_file(official_tokenizer_file("tokenizer.json").c_str()); static const std::string tokenizer_config_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/tokenizer_config.json"); + read_file(official_tokenizer_file("tokenizer_config.json").c_str()); static const std::string generation_config_json = - read_file("/home/neroued/models/llm/qwen/Qwen3.6-27B/base-hf-bf16/generation_config.json"); + read_file(official_tokenizer_file("generation_config.json").c_str()); static const fi::Tokenizer tokenizer({.tokenizer_json = tokenizer_json, .tokenizer_config_json = tokenizer_config_json, .generation_config_json = generation_config_json}); @@ -1305,6 +1324,11 @@ int test_media_preparation_cancellation() { } // namespace int main() { + if (!official_tokenizer_available()) { + std::cerr << "skipping: official tokenizer fixtures not found under " + << kOfficialTokenizerDir << '\n'; + return 77; + } const FrontendResources owned = resources(); const Frontend frontend = FrontendFactory::create_component(owned); int failures = 0; diff --git a/tests/test_gdn_replay_records.cpp b/tests/test_gdn_replay_records.cpp index 2c777b5fe2..938ffff64b 100644 --- a/tests/test_gdn_replay_records.cpp +++ b/tests/test_gdn_replay_records.cpp @@ -9,14 +9,32 @@ #include #include +#ifdef _WIN32 +#include +#endif + namespace { -using AlignedBacking = std::unique_ptr; +// The MSVC CRT provides no std::aligned_alloc, and blocks from _aligned_malloc +// must be released with _aligned_free rather than free. +void aligned_release(void* data) { +#ifdef _WIN32 + ::_aligned_free(data); +#else + std::free(data); +#endif +} + +using AlignedBacking = std::unique_ptr; AlignedBacking make_backing(std::size_t bytes) { +#ifdef _WIN32 + void* data = ::_aligned_malloc(bytes, 256); +#else void* data = std::aligned_alloc(256, bytes); +#endif if (data == nullptr) { throw std::bad_alloc(); } - return AlignedBacking(data, &std::free); + return AlignedBacking(data, &aligned_release); } int fail(const char* label) { From 9b1a58958262a64f7ff99d19555f685a3ae1c1b7 Mon Sep 17 00:00:00 2001 From: pelebel Date: Wed, 19 Aug 2026 16:46:39 -0400 Subject: [PATCH 16/20] build: add hosted Windows build workflow (windows-2022, build-only) --- .github/workflows/windows-build.yml | 97 +++++++++++++++++++++++++++++ 1 file changed, 97 insertions(+) create mode 100644 .github/workflows/windows-build.yml diff --git a/.github/workflows/windows-build.yml b/.github/workflows/windows-build.yml new file mode 100644 index 0000000000..0cfcf53c83 --- /dev/null +++ b/.github/workflows/windows-build.yml @@ -0,0 +1,97 @@ +name: Windows build + +on: + push: + branches: + - master + pull_request: + workflow_dispatch: + +# Build-only check on a hosted Windows runner: configure and compile `ninfer` +# and `ninfer-serve` with MSVC (VS2022 `cl`) + `nvcc` (CUDA) + Ninja, resolving +# FFmpeg/libcurl from a vcpkg `x64-windows` prefix. +# +# Hosted runners have no GPU, so the CUDA test suite is NOT executed here — +# running `ctest` (numerical Op suites, real-engine routes) requires a +# self-hosted sm_120a runner. This job guards the Windows *compile* path +# (MSVC + nvcc, the NVFP4 TMA by-value fix, the plan move-ctor emission, and the +# Winsock/reader/datetime shims) and smoke-launches both executables. +# +# `sm_120a` requires CUDA >= 13.1, which is not on the stock windows-2022 image. +# The CUDA step uses a preinstalled toolkit if present, otherwise falls back to +# `choco install cuda`, then hard-guards the version so the job fails loudly +# (with instructions) rather than as a confusing `sm_120a` compile error. +env: + VCPKG_TRIPLET: x64-windows + BUILD_DIR: build + NINJA_VERSION: v1.12.1 + +jobs: + build-windows: + name: Build + smoke (windows-2022, no GPU) + runs-on: windows-2022 + timeout-minutes: 120 + + steps: + - name: Checkout + uses: actions/checkout@v4 + + - name: Install vcpkg dependencies (libcurl, FFmpeg) + # FFmpeg is built from source by vcpkg and is the slow step of the job. + # The resolved prefix (vcpkg/installed/x64-windows) is the NINFER_MEDIA_ROOT. + shell: pwsh + run: | + $vcpkg = Join-Path $env:RUNNER_TEMP 'vcpkg' + git clone --depth 1 https://github.com/microsoft/vcpkg.git $vcpkg + & (Join-Path $vcpkg 'bootstrap-vcpkg.bat') + & (Join-Path $vcpkg 'vcpkg.exe') install "curl:$env:VCPKG_TRIPLET" "ffmpeg:$env:VCPKG_TRIPLET" + if ($LASTEXITCODE -ne 0) { throw 'vcpkg install failed' } + Write-Host "NINFER_MEDIA_ROOT=$vcpkg\installed\$env:VCPKG_TRIPLET" + + - name: Locate Visual Studio 2022 (vcvars64) + shell: pwsh + run: | + $vswhere = "${env:ProgramFiles(x86)}\Microsoft Visual Studio\Installer\vswhere.exe" + if (-not (Test-Path $vswhere)) { throw 'vswhere.exe not found; is VS2022 present?' } + $install = & $vswhere -latest -products * ` + -requires Microsoft.VisualStudio.Component.VC.Tools.x86.x64 ` + -property installationPath + if (-not $install) { throw 'VS2022 C++ workload not found on the runner' } + $vcvars = Join-Path $install 'VC\Auxiliary\Build\vcvars64.bat' + Write-Host "vcvars=$vcvars" + "VCVARS=$vcvars" | Add-Content $env:GITHUB_ENV + + - name: Configure, build, and smoke-launch + shell: pwsh + run: | + # --- Ninja (added to PATH for the cmake steps below) --- + curl.exe -sSL -o ninja.zip "https://github.com/ninja-build/ninja/releases/download/$env:NINJA_VERSION/ninja-win.zip" + New-Item -ItemType Directory -Force tools | Out-Null + Expand-Archive ninja.zip -DestinationPath tools -Force + $env:PATH = "$PWD\tools;$env:PATH" + & ninja --version + + # --- CUDA 13.x (>= 13.1 for sm_120a) --- + $cudaRoot = "C:\Program Files\NVIDIA GPU Computing Toolkit\CUDA" + if (-not (Test-Path $cudaRoot)) { + Write-Host 'CUDA toolkit not present on the runner; installing via choco (best effort).' + choco install -y cuda + } + $cudaDir = Get-ChildItem $cudaRoot -Directory | Sort-Object Name | Select-Object -Last 1 + if (-not $cudaDir) { throw "no CUDA toolkit found under $cudaRoot" } + $env:PATH = "$($cudaDir.FullName)\bin;$env:PATH" + $nvcc = & nvcc --version + $nvcc -join "`n" | Write-Host + $rel = ($nvcc | Select-String 'release (\d+\.\d+)').Matches[0].Groups[1].Value + $num = [int]$rel.Split('.')[0] * 100 + [int]$rel.Split('.')[1] + if ($num -lt 1310) { + throw "CUDA $rel found; sm_120a requires CUDA >= 13.1. Use a runner image with CUDA 13.x, or provision it before this step." + } + + # --- Configure + build + smoke, all in one cmd so the vcvars env is shared --- + $media = Join-Path $env:RUNNER_TEMP 'vcpkg\installed\x64-windows' + $cmake = "cmake -S . -B $env:BUILD_DIR -G Ninja -DCMAKE_BUILD_TYPE=Release -DCMAKE_CUDA_ARCHITECTURES=120a -DNINFER_MEDIA_ROOT=`"$media`" -DBUILD_TESTING=OFF" + $build = "cmake --build $env:BUILD_DIR -j" + $smoke = "$env:BUILD_DIR\apps\ninfer.exe --help && $env:BUILD_DIR\apps\ninfer-serve.exe --help" + cmd /c "call `"$env:VCVARS`" && $cmake && $build && $smoke" + if ($LASTEXITCODE -ne 0) { throw "Windows build failed (exit $LASTEXITCODE)" } \ No newline at end of file From 9028b8fa59c5660d9d1b570a0bc010a2aac06747 Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 15:38:48 -0400 Subject: [PATCH 17/20] build(win32): keep the MSVC guards from mis-flagging clang-cl CMake sets MSVC for any MSVC-ABI compiler, clang-cl included, so both Windows guards fired for a compiler they were not written for. /Zc:preprocessor is a cl flag. clang-cl's preprocessor is already conforming and reports the flag as an unknown argument on every translation unit, so ask for it only when the C++ compiler really is cl. The -Xcompiler copy stays unconditional: nvcc drives cl as its host compiler on Windows whatever CMake uses for C++, and its generated host stubs need it. The op tests disable fast-math and FP contraction so the CPU reference they compare GPU kernels against is not reassociated. clang-cl reports its compiler id as Clang but takes cl-style flags and ignores the GNU spellings silently, which would have dropped that guarantee without a diagnostic; route them through /clang: instead. Neither change affects the cl or POSIX builds: verified by a clean configure and build plus the full ctest suite (89/89) under VS2022 cl. Co-Authored-By: Claude Opus 5 --- CMakeLists.txt | 17 ++++++++++++----- tests/CMakeLists.txt | 13 ++++++++++++- 2 files changed, 24 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index dec864d721..36e56962f5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,13 +68,20 @@ endif() find_package(Threads REQUIRED) +# MSVC is true for any MSVC-ABI compiler, which includes clang-cl. Settings that +# every such compiler needs go here; settings specific to cl are gated further. if(MSVC) - # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. The - # flag is forwarded to the cl host compiler nvcc invokes, including for the - # host stubs it generates. (The 128-byte-aligned TMA descriptor parameters are - # passed by device pointer rather than by value precisely so MSVC's stub + # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. cl needs + # to be told; clang-cl's preprocessor is already conforming and rejects the flag + # as unknown, so ask only when the C++ frontend really is cl. The -Xcompiler + # form is unconditional because nvcc drives cl as its host compiler on Windows + # regardless of which compiler CMake uses for C++, and the host stubs it + # generates need the flag too. (The 128-byte-aligned TMA descriptor parameters + # are passed by device pointer rather than by value precisely so MSVC's stub # compiler does not reject them.) - add_compile_options($<$:/Zc:preprocessor>) + if(CMAKE_CXX_COMPILER_ID STREQUAL "MSVC") + add_compile_options($<$:/Zc:preprocessor>) + endif() set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=/Zc:preprocessor") # windows.h defines min/max as macros which clobber std::min/std::max; NOMINMAX # suppresses them. UTF8PROC_STATIC keeps the vendored utf8proc header from diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index e7a680fc75..ef71eb3151 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -55,7 +55,18 @@ endfunction() function(ninfer_add_op_test name) ninfer_add_test(${name} ${ARGN}) set_tests_properties(${name} PROPERTIES SKIP_RETURN_CODE 77) - if(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang") + # Op tests check GPU kernels against a CPU reference, so the host build must not + # reassociate or fuse the reference arithmetic. clang-cl reports its id as Clang + # but takes cl-style flags, and silently ignores the GNU spelling: route those + # through /clang: instead, or the reference loses the guarantee without a word. + # cl keeps its default, which is what the suite is verified against today. + if(CMAKE_CXX_COMPILER_FRONTEND_VARIANT STREQUAL "MSVC") + if(CMAKE_CXX_COMPILER_ID MATCHES "Clang") + target_compile_options(${name} PRIVATE + $<$:/clang:-fno-fast-math> + $<$:/clang:-ffp-contract=off>) + endif() + elseif(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang") target_compile_options(${name} PRIVATE $<$:-fno-fast-math> $<$:-ffp-contract=off>) From 9eec19589a0b6bea09a8154879807f079643c8cf Mon Sep 17 00:00:00 2001 From: pelebel Date: Fri, 21 Aug 2026 15:38:48 -0400 Subject: [PATCH 18/20] build(win32): keep the MSVC guards from mis-flagging clang-cl CMake sets MSVC for any MSVC-ABI compiler, clang-cl included, so both Windows guards fired for a compiler they were not written for. /Zc:preprocessor is a cl flag. clang-cl's preprocessor is already conforming and reports the flag as an unknown argument on every translation unit, so ask for it only when the C++ compiler really is cl. The -Xcompiler copy stays unconditional: nvcc drives cl as its host compiler on Windows whatever CMake uses for C++, and its generated host stubs need it. The op tests disable fast-math and FP contraction so the CPU reference they compare GPU kernels against is not reassociated. clang-cl reports its compiler id as Clang but takes cl-style flags and ignores the GNU spellings silently, which would have dropped that guarantee without a diagnostic; route them through /clang: instead. Neither change affects the cl or POSIX builds: verified by a clean configure and build plus the full ctest suite (89/89) under VS2022 cl. Co-Authored-By: Claude Opus 5 --- CMakeLists.txt | 17 ++++++++++++----- tests/CMakeLists.txt | 13 ++++++++++++- 2 files changed, 24 insertions(+), 6 deletions(-) diff --git a/CMakeLists.txt b/CMakeLists.txt index dec864d721..36e56962f5 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -68,13 +68,20 @@ endif() find_package(Threads REQUIRED) +# MSVC is true for any MSVC-ABI compiler, which includes clang-cl. Settings that +# every such compiler needs go here; settings specific to cl are gated further. if(MSVC) - # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. The - # flag is forwarded to the cl host compiler nvcc invokes, including for the - # host stubs it generates. (The 128-byte-aligned TMA descriptor parameters are - # passed by device pointer rather than by value precisely so MSVC's stub + # CUDA 13's CCCL (cuda/std) requires the conforming MSVC preprocessor. cl needs + # to be told; clang-cl's preprocessor is already conforming and rejects the flag + # as unknown, so ask only when the C++ frontend really is cl. The -Xcompiler + # form is unconditional because nvcc drives cl as its host compiler on Windows + # regardless of which compiler CMake uses for C++, and the host stubs it + # generates need the flag too. (The 128-byte-aligned TMA descriptor parameters + # are passed by device pointer rather than by value precisely so MSVC's stub # compiler does not reject them.) - add_compile_options($<$:/Zc:preprocessor>) + if(CMAKE_CXX_COMPILER_ID STREQUAL "MSVC") + add_compile_options($<$:/Zc:preprocessor>) + endif() set(CMAKE_CUDA_FLAGS "${CMAKE_CUDA_FLAGS} -Xcompiler=/Zc:preprocessor") # windows.h defines min/max as macros which clobber std::min/std::max; NOMINMAX # suppresses them. UTF8PROC_STATIC keeps the vendored utf8proc header from diff --git a/tests/CMakeLists.txt b/tests/CMakeLists.txt index e7a680fc75..ef71eb3151 100644 --- a/tests/CMakeLists.txt +++ b/tests/CMakeLists.txt @@ -55,7 +55,18 @@ endfunction() function(ninfer_add_op_test name) ninfer_add_test(${name} ${ARGN}) set_tests_properties(${name} PROPERTIES SKIP_RETURN_CODE 77) - if(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang") + # Op tests check GPU kernels against a CPU reference, so the host build must not + # reassociate or fuse the reference arithmetic. clang-cl reports its id as Clang + # but takes cl-style flags, and silently ignores the GNU spelling: route those + # through /clang: instead, or the reference loses the guarantee without a word. + # cl keeps its default, which is what the suite is verified against today. + if(CMAKE_CXX_COMPILER_FRONTEND_VARIANT STREQUAL "MSVC") + if(CMAKE_CXX_COMPILER_ID MATCHES "Clang") + target_compile_options(${name} PRIVATE + $<$:/clang:-fno-fast-math> + $<$:/clang:-ffp-contract=off>) + endif() + elseif(CMAKE_CXX_COMPILER_ID MATCHES "GNU|Clang") target_compile_options(${name} PRIVATE $<$:-fno-fast-math> $<$:-ffp-contract=off>) From 636ea0061948fb561e0f58365bd6d689e7b50eff Mon Sep 17 00:00:00 2001 From: pelebel Date: Sat, 22 Aug 2026 15:50:34 -0400 Subject: [PATCH 19/20] feat(serve): answer the announced API base with an endpoint index --- docs/serving.md | 1 + src/serve/http_server.cpp | 43 +++++++++++++++++++++++++++++++++++ src/serve/http_server.h | 1 + src/serve/openai_schema.h | 4 ++++ tests/test_openai_schema.cpp | 25 ++++++++++++++++++++ tools/smoke/serve_contract.py | 7 ++++++ 6 files changed, 81 insertions(+) diff --git a/docs/serving.md b/docs/serving.md index 3fe656cbf5..5494737c15 100644 --- a/docs/serving.md +++ b/docs/serving.md @@ -43,6 +43,7 @@ cannot be combined with `--vision`. A later request cannot enable a capability o | Method and path | Behavior | |---|---| | `GET /health` | process health | +| `GET /v1` | endpoint index for the announced API base | | `GET /v1/models` | configured OpenAI model alias | | `GET /v1/models/{id}` | lookup of the configured alias | | `POST /v1/chat/completions` | OpenAI-style chat generation | diff --git a/src/serve/http_server.cpp b/src/serve/http_server.cpp index 85827771af..6ad41781a7 100644 --- a/src/serve/http_server.cpp +++ b/src/serve/http_server.cpp @@ -414,6 +414,14 @@ void HttpServer::register_routes() { server_.Get("/health", [](const httplib::Request&, httplib::Response& res) { res.set_content(nlohmann::json{{"status", "ok"}}.dump(), "application/json"); }); + // The startup banner prints the API base as a URL and terminals make it clickable, + // so the bare base must answer rather than 404. Both spellings, since a browser + // address bar happily produces either. + for (const char* base : {"/v1", "/v1/"}) { + server_.Get(base, [this](const httplib::Request& req, httplib::Response& res) { + handle_api_index(req, res); + }); + } server_.Get("/v1/models", [this](const httplib::Request& req, httplib::Response& res) { handle_models(req, res); }); @@ -463,6 +471,37 @@ void HttpServer::register_routes() { }); } +// The /v1 discovery document. An API base is not itself an OpenAI resource, so +// there is no upstream shape to match: report the configured model alias and the +// endpoint table this build serves, which is what someone who opened the URL (or +// pointed a client at it) needs to see. +nlohmann::json make_api_index(const std::string& model_id) { + const auto endpoint = [](const char* method, const char* path, const char* description) { + return nlohmann::json{ + {"method", method}, {"path", path}, {"description", description}}; + }; + nlohmann::json index = nlohmann::json::object(); + index["object"] = "api_base"; + index["service"] = "ninfer-serve"; + index["model"] = model_id; + index["endpoints"] = nlohmann::json::array( + {endpoint("GET", "/health", "process health"), + endpoint("GET", "/v1/models", "configured OpenAI model alias"), + endpoint("GET", "/v1/models/{id}", "lookup of the configured alias"), + endpoint("POST", "/v1/chat/completions", "OpenAI-style chat generation"), + endpoint("POST", "/v1/responses", "OpenAI Responses generation, state, and SSE"), + endpoint("POST", "/v1/responses/input_tokens", + "Responses prompt-token count without generation"), + endpoint("GET", "/v1/responses/{id}", "retrieve a locally stored terminal Response"), + endpoint("DELETE", "/v1/responses/{id}", "delete a locally stored Response"), + endpoint("GET", "/v1/responses/{id}/input_items", + "list that Response's normalized input Items"), + endpoint("POST", "/v1/messages", "Anthropic-style message generation"), + endpoint("POST", "/v1/messages/count_tokens", + "checkpoint-native expanded input-token count")}); + return index; +} + // llama.cpp webui dialect: /props is the client's server introspection endpoint // (role detection, context size, default params, thinking-capability probe, api // key validation). NInfer has no llama.cpp server behind it, so serve a faithful @@ -560,6 +599,10 @@ nlohmann::json make_props_stub(const ServeOptions& options, const std::string& m return props; } +void HttpServer::handle_api_index(const httplib::Request&, httplib::Response& res) const { + res.set_content(make_api_index(public_model_id_).dump(), "application/json"); +} + void HttpServer::handle_models(const httplib::Request&, httplib::Response& res) const { res.set_content(make_models_list(public_model_id_, unix_time_now()), "application/json"); } diff --git a/src/serve/http_server.h b/src/serve/http_server.h index 1207e03c4a..2856df08f3 100644 --- a/src/serve/http_server.h +++ b/src/serve/http_server.h @@ -53,6 +53,7 @@ class HttpServer { void handle_response_input_items(const httplib::Request& req, httplib::Response& res); void handle_response_cancel(const httplib::Request& req, httplib::Response& res); void handle_response_compact(const httplib::Request& req, httplib::Response& res); + void handle_api_index(const httplib::Request& req, httplib::Response& res) const; void handle_models(const httplib::Request& req, httplib::Response& res) const; void handle_model(const httplib::Request& req, httplib::Response& res) const; void handle_props(const httplib::Request& req, httplib::Response& res) const; diff --git a/src/serve/openai_schema.h b/src/serve/openai_schema.h index 0554ac2738..b72b0dab98 100644 --- a/src/serve/openai_schema.h +++ b/src/serve/openai_schema.h @@ -90,6 +90,10 @@ std::string sse_done(); std::string make_models_list(const std::string& model_id, std::int64_t created); std::string make_model_object(const std::string& model_id, std::int64_t created); +// /v1 discovery document: the API base is announced as a URL at startup, so the +// bare path answers with the endpoints this build exposes instead of a 404. +nlohmann::json make_api_index(const std::string& model_id); + // Error object body. std::string make_error_body(const ApiError& error); diff --git a/tests/test_openai_schema.cpp b/tests/test_openai_schema.cpp index 9727022aaa..08f3310fb8 100644 --- a/tests/test_openai_schema.cpp +++ b/tests/test_openai_schema.cpp @@ -867,6 +867,30 @@ int test_props_stub() { return failures; } +int test_api_index() { + int failures = 0; + const Json index = make_api_index("qwen3.6-27b"); + failures += check(index.at("object") == "api_base", "api index object"); + failures += check(index.at("model") == "qwen3.6-27b", "api index reports the model alias"); + const Json endpoints = index.at("endpoints"); + failures += check(endpoints.is_array() && !endpoints.empty(), "api index lists endpoints"); + // Every listed endpoint must be one the server actually registers, otherwise the + // discovery document sends clients to a 404 of its own. + bool has_chat = false; + bool well_formed = true; + for (const Json& entry : endpoints) { + well_formed = well_formed && entry.at("method").is_string() && + entry.at("path").get().rfind('/', 0) == 0 && + entry.at("description").is_string(); + if (entry.at("path") == "/v1/chat/completions" && entry.at("method") == "POST") { + has_chat = true; + } + } + failures += check(well_formed, "api index entries carry method, absolute path, description"); + failures += check(has_chat, "api index advertises chat completions"); + return failures; +} + int main() { int failures = 0; failures += test_parse_string_content(); @@ -887,6 +911,7 @@ int main() { failures += test_models_and_error(); failures += test_llama_webui_dialect(); failures += test_props_stub(); + failures += test_api_index(); failures += test_finish_reason_wire(); if (failures == 0) { std::cout << "ok\n"; } return failures == 0 ? 0 : 1; diff --git a/tools/smoke/serve_contract.py b/tools/smoke/serve_contract.py index c51982a539..db6ec1eb7b 100644 --- a/tools/smoke/serve_contract.py +++ b/tools/smoke/serve_contract.py @@ -324,6 +324,13 @@ def parse_responses_stream(response: Response) -> tuple[str, str, dict[str, Any] def exercise(base_url: str, model: str) -> dict[str, Any]: + # The startup banner prints the API base as an openable URL; it must answer. + index = json_response(base_url, "GET", "/v1") + if index.get("object") != "api_base" or index.get("model") != model: + raise ContractError("API base index has the wrong shape") + if not any(entry.get("path") == "/v1/chat/completions" for entry in index.get("endpoints", [])): + raise ContractError("API base index does not advertise chat completions") + models = json_response(base_url, "GET", "/v1/models") entries = models.get("data") if models.get("object") != "list" or not isinstance(entries, list) or len(entries) != 1: From cd98b520fdb5e57542f4d57884a9698f607e3176 Mon Sep 17 00:00:00 2001 From: pelebel Date: Sat, 22 Aug 2026 15:51:06 -0400 Subject: [PATCH 20/20] docs(readme): document the Windows build --- README.md | 39 +++++++++++++++++++++++++++++++++------ 1 file changed, 33 insertions(+), 6 deletions(-) diff --git a/README.md b/README.md index b0a28379a9..6f51451cd4 100644 --- a/README.md +++ b/README.md @@ -121,14 +121,16 @@ notes. NInfer currently requires: -- 64-bit Linux; +- 64-bit Linux or Windows 11; - NVIDIA GeForce RTX 5090 (`sm_120a`); - NVIDIA driver support for CUDA 13.1 and the CUDA Toolkit 13.1 or newer; -- CMake 3.28 or newer and a C++20-capable host compiler; -- `pkg-config`; +- CMake 3.28 or newer and a C++20-capable host compiler (GCC or Clang on Linux, MSVC from + Visual Studio 2022 on Windows); +- `pkg-config` on Linux; - FFmpeg development libraries: `libavformat >= 60`, `libavcodec >= 60`, - `libavutil >= 58`, and `libswscale >= 7`; -- `libcurl >= 7.85`; + `libavutil >= 58`, and `libswscale >= 7` (on Windows, `vcpkg install ffmpeg --triplet + x64-windows`); +- `libcurl >= 7.85` (on Windows, `vcpkg install curl --triplet x64-windows`); - Ninja, when using the commands below. The build rejects CUDA architectures other than `120a`. There is no install target or packaged @@ -153,11 +155,35 @@ build/apps/ninfer-serve Tests, benchmarks, and maintainer tools are excluded from the default build. +### Windows + +Install the Visual Studio 2022 C++ build tools (the "Desktop development with C++" workload), +the CUDA Toolkit, CMake, Ninja, and [vcpkg](https://github.com/microsoft/vcpkg), then install +the media dependencies: + +```bat +vcpkg install ffmpeg curl --triplet x64-windows +``` + +Configure and build from a prompt where `cl` is on `PATH` (a VS 2022 Developer Command Prompt, +or after calling `vcvars64.bat`), pointing the build at the vcpkg install prefix: + +```bat +cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DNINFER_MEDIA_ROOT=%VCPKG_ROOT%\installed\x64-windows +cmake --build build --parallel +``` + +The default configuration produces `build\apps\ninfer.exe` and `build\apps\ninfer-serve.exe`. +The FFmpeg and libcurl runtime DLLs are copied next to the executables at build time, so both +apps run directly from the build tree. In the remaining sections, use +`build\apps\ninfer.exe` and `build\apps\ninfer-serve.exe` in place of the Linux paths. + ## Docker Build the runtime image on a 64-bit Linux host with an RTX 5090, a CUDA 13.1-compatible NVIDIA driver, Docker, and the [NVIDIA Container Toolkit](https://docs.nvidia.com/datacenter/cloud-native/container-toolkit/latest/install-guide.html). +Windows hosts use the native build above. ```bash docker build --tag ninfer:local . @@ -228,7 +254,8 @@ python3 -m tools.artifact.migrate_v1_to_v2 models/qwen3_6_27b.ninfer Use the same command with `qwen3_6_27b_nvfp4.ninfer` or `qwen3_6_35b_a3b.ninfer` for those artifacts. The migration updates only container metadata; it does not rewrite the weight payload. -Alternatively, download the current version-2 file again from its Hugging Face repository. +Alternatively, download the current version-2 file again from its Hugging Face repository. On +Windows, invoke the same module with `python -m tools.artifact.migrate_v1_to_v2`. Each `.ninfer` file contains the weights and frontend resources needed by NInfer. It is not a Transformers checkpoint, Safetensors distribution, or GGUF file.