diff --git a/.gitignore b/.gitignore index 6350d8b482..6111a1f929 100644 --- a/.gitignore +++ b/.gitignore @@ -2,6 +2,7 @@ build/ build-*/ cmake-build-*/ +vcpkg_installed/ *.o *.obj *.so diff --git a/CMakeLists.txt b/CMakeLists.txt index 27e9560e9f..a625640434 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -28,6 +28,17 @@ if(NOT CMAKE_BUILD_TYPE) set(CMAKE_BUILD_TYPE Release) endif() +if(WIN32) + # windows.h defines min/max macros that corrupt std::min/std::max and the CUDA + # headers; NOMINMAX is required for every translation unit on Windows. + add_compile_definitions(NOMINMAX) + # CCCL (bundled with CUDA 13) hard-requires the conforming MSVC preprocessor, + # and spdlog/fmt requires UTF-8 source encoding. + add_compile_options($<$:-Xcompiler=/Zc:preprocessor>) + add_compile_options($<$:-Xcompiler=/utf-8>) + add_compile_options($<$:/utf-8>) +endif() + if(CMAKE_GENERATOR MATCHES "Ninja") set_property(GLOBAL PROPERTY JOB_POOLS ninfer_link=1) set(CMAKE_JOB_POOL_LINK ninfer_link) diff --git a/CMakePresets.json b/CMakePresets.json index 75d37a8c1e..fd32148d5f 100644 --- a/CMakePresets.json +++ b/CMakePresets.json @@ -22,11 +22,24 @@ "BUILD_TESTING": "ON", "NINFER_BUILD_BENCHMARKS": "ON" } + }, + { + "name": "windows-vcpkg", + "displayName": "Windows (vcpkg, Ninja)", + "description": "Configures with the vcpkg toolchain for manifest-mode dependency resolution. Set VCPKG_ROOT or vcpkg will be auto-detected.", + "generator": "Ninja", + "binaryDir": "${sourceDir}/build-win", + "cacheVariables": { + "CMAKE_BUILD_TYPE": "Release", + "CMAKE_CUDA_ARCHITECTURES": "120a" + }, + "toolchainFile": "$env{VCPKG_ROOT}/scripts/buildsystems/vcpkg.cmake" } ], "buildPresets": [ { "name": "release", "configurePreset": "release" }, - { "name": "dev", "configurePreset": "dev" } + { "name": "dev", "configurePreset": "dev" }, + { "name": "windows-vcpkg", "configurePreset": "windows-vcpkg" } ], "testPresets": [ { diff --git a/README.md b/README.md index daa143b881..2b6dacb6aa 100644 --- a/README.md +++ b/README.md @@ -1,3 +1,14 @@ +--- +AIGC: + ContentProducer: '001191110102MAD55U9H0F10002' + ContentPropagator: '001191110102MAD55U9H0F10002' + Label: '1' + ProduceID: 'fa16c31a-cc76-42b2-a193-0e4ef2d56b3e' + PropagateID: 'fa16c31a-cc76-42b2-a193-0e4ef2d56b3e' + ReservedCode1: '699436de-0b80-4adf-a485-3bc60a6778a7' + ReservedCode2: '699436de-0b80-4adf-a485-3bc60a6778a7' +--- + # NInfer > Selected checkpoints. Maximum single-GPU inference performance. @@ -28,11 +39,12 @@ the weights again. ## Quick start -NInfer requires 64-bit Linux, an NVIDIA GeForce RTX 5090, a CUDA toolkit supporting `sm_120a`, -CMake 3.28 or newer, a C++20 host compiler, Ninja, `pkg-config`, FFmpeg development libraries -(`libavformat`, `libavcodec`, `libavutil`, and `libswscale`), and `libcurl >= 7.85`. -CUDA 13.1 is the validated development toolkit; CMake does not impose a CUDA version floor. -The build rejects CUDA architectures other than `sm_120a`. +NInfer requires 64-bit Linux (or Windows), an NVIDIA GeForce RTX 5090, a CUDA toolkit +supporting `sm_120a`, CMake 3.28 or newer, a C++20 host compiler, Ninja, `pkg-config`, +FFmpeg development libraries (`libavformat`, `libavcodec`, `libavutil`, and +`libswscale`), and `libcurl >= 7.85`. CUDA 13.1 is the validated development toolkit; +CMake does not impose a CUDA version floor. The build rejects CUDA architectures other +than `sm_120a`. Build the product binaries: @@ -44,11 +56,33 @@ cmake -S . -B build -G Ninja -DCMAKE_BUILD_TYPE=Release cmake --build build -j ``` +On Windows, dependencies (FFmpeg and libcurl with Schannel TLS) are managed with +[vcpkg](https://vcpkg.io) in manifest mode — no GStreamer runtime or prebuilt curl +packages are needed: + +```powershell +git clone https://github.com/Neroued/ninfer.git +cd ninfer + +# One-time setup (or use an existing vcpkg installation) +git clone https://github.com/microsoft/vcpkg.git +cd vcpkg && .\bootstrap-vcpkg.bat && cd .. +$env:VCPKG_ROOT = "$PWD\vcpkg" + +# Configure through the vcpkg toolchain; dependencies install automatically +cmake -S . -B build-win -G Ninja -DCMAKE_TOOLCHAIN_FILE="$env:VCPKG_ROOT/scripts/buildsystems/vcpkg.cmake" -DCMAKE_CUDA_ARCHITECTURES=120a -DCMAKE_BUILD_TYPE=Release +cmake --build build-win -j +``` + +On first configuration vcpkg compiles FFmpeg and curl from source (about 15 minutes); +subsequent configurations restore them from the local binary cache in seconds. + Tests and benchmarks are excluded from the default build. `cmake --preset release` configures the same product build; `cmake --preset dev` also enables tests and benchmarks and finds a Python 3 interpreter. Both presets use `build/` and explicitly reset the build options. Machine-specific compiler and Python paths belong in the ignored `CMakeUserPresets.json`. See [build organization and configuration](docs/maintainer/build-system.md) for details. +The `windows-vcpkg` preset wraps the vcpkg toolchain for the Windows build. There is no install target or packaged binary distribution; run NInfer from its source build tree. Python tools run independently of CMake; the standalone HBM probe has its own @@ -292,4 +326,4 @@ also uses the fixed packed weights from The Qwen3.8-27B NVFP4 artifact also uses the fixed mixed FP8/NVFP4 weights from [unsloth/Qwen3.8-27B-NVFP4](https://huggingface.co/unsloth/Qwen3.8-27B-NVFP4). These source repositories are distributed under Apache-2.0. Vendored dependencies retain their own license files -under `third_party/`. +under `third_party/`. \ No newline at end of file diff --git a/apps/CMakeLists.txt b/apps/CMakeLists.txt index 9242a9e2fe..30aad27bb9 100644 --- a/apps/CMakeLists.txt +++ b/apps/CMakeLists.txt @@ -1,3 +1,11 @@ +# Embed an activeCodePage=UTF-8 manifest so the MSVC CRT hands apps UTF-8 argv on +# Windows instead of the legacy ANSI code page (breaks non-ASCII --prompt/paths). +# The manifest is appended after each add_executable via target_sources; CMake's +# MSVC rule recognizes .manifest sources and merges them into the default manifest. +if(WIN32 AND MSVC) + set(NINFER_APP_MANIFEST ${CMAKE_CURRENT_SOURCE_DIR}/windows-utf8.manifest) +endif() + add_executable(ninfer cli/main.cpp cli/options.cpp) @@ -6,10 +14,16 @@ target_link_libraries(ninfer PRIVATE ninfer_engine ninfer_product_logging ninfer_product_prompt_input) +if(WIN32 AND MSVC) + target_sources(ninfer PRIVATE ${NINFER_APP_MANIFEST}) +endif() add_executable(ninfer-serve serve/main.cpp) ninfer_internal_includes(ninfer-serve) target_link_libraries(ninfer-serve PRIVATE ninfer_serve ninfer_product_logging) +if(WIN32 AND MSVC) + target_sources(ninfer-serve PRIVATE ${NINFER_APP_MANIFEST}) +endif() add_executable(ninfer-perplexity perplexity/main.cpp @@ -18,3 +32,6 @@ add_executable(ninfer-perplexity ninfer_internal_includes(ninfer-perplexity) target_link_libraries(ninfer-perplexity PRIVATE ninfer_engine ninfer_product_logging ninfer::json) +if(WIN32 AND MSVC) + target_sources(ninfer-perplexity PRIVATE ${NINFER_APP_MANIFEST}) +endif() diff --git a/apps/cli/main.cpp b/apps/cli/main.cpp index bc8d786fa1..d69341b6a8 100644 --- a/apps/cli/main.cpp +++ b/apps/cli/main.cpp @@ -17,6 +17,13 @@ #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page (e.g. +// 936/GBK). The activeCodePage=UTF-8 manifest fixes argv input; this fixes display +// of UTF-8 output on code-page-936 consoles. Redirected pipes keep raw UTF-8 bytes. +#include +#endif + namespace { std::string format_seconds(double seconds) { @@ -228,6 +235,20 @@ void print_generation_summary(const ninfer::GenerationResult& result, } // namespace +#ifdef _WIN32 +namespace { +// Set once before any output; harmless when stdout is redirected (it only affects +// how the attached console decodes the byte stream). +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + int main(int argc, char** argv) { ninfer::cli::Options cli; try { diff --git a/apps/perplexity/main.cpp b/apps/perplexity/main.cpp index e7e40503c9..14b1a11f54 100644 --- a/apps/perplexity/main.cpp +++ b/apps/perplexity/main.cpp @@ -9,6 +9,12 @@ #include #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page. The +// activeCodePage=UTF-8 manifest fixes argv input; this fixes console display. +#include +#endif + #include #include #include @@ -163,7 +169,12 @@ std::string safe_component(std::string_view value) { std::string timestamp() { const std::time_t now = std::chrono::system_clock::to_time_t(std::chrono::system_clock::now()); std::tm utc{}; +#ifdef _WIN32 + // gmtime_s swaps argument order relative to POSIX gmtime_r. + gmtime_s(&utc, &now); +#else gmtime_r(&now, &utc); +#endif std::ostringstream out; out << std::put_time(&utc, "%Y%m%d-%H%M%S"); return out.str(); @@ -434,6 +445,18 @@ int run(const Options& options, const std::shared_ptr& logger, } // namespace +#ifdef _WIN32 +namespace { +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + int main(int argc, char** argv) { Options options; try { diff --git a/apps/serve/main.cpp b/apps/serve/main.cpp index 02784db556..876a8936b1 100644 --- a/apps/serve/main.cpp +++ b/apps/serve/main.cpp @@ -16,6 +16,21 @@ #include #include +#ifdef _WIN32 +// Emit UTF-8 bytes to the console regardless of the legacy ANSI code page. The +// activeCodePage=UTF-8 manifest fixes argv input; this fixes console display. +#include +namespace { +struct ConsoleUtf8Setup { + ConsoleUtf8Setup() { + SetConsoleOutputCP(CP_UTF8); + SetConsoleCP(CP_UTF8); + } +}; +const ConsoleUtf8Setup console_utf8_setup; +} // namespace +#endif + namespace { std::atomic g_server{nullptr}; diff --git a/apps/windows-utf8.manifest b/apps/windows-utf8.manifest new file mode 100644 index 0000000000..88938f809d --- /dev/null +++ b/apps/windows-utf8.manifest @@ -0,0 +1,9 @@ + + + + + + UTF-8 + + + diff --git a/cmake/Dependencies.cmake b/cmake/Dependencies.cmake index 02bfdb08f5..4233caf1ef 100644 --- a/cmake/Dependencies.cmake +++ b/cmake/Dependencies.cmake @@ -1,8 +1,26 @@ find_package(CUDAToolkit REQUIRED) find_package(Threads REQUIRED) -find_package(PkgConfig REQUIRED) -pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET - libavformat libavcodec libavutil libswscale) + +if(WIN32) + # --- FFmpeg (Windows, vcpkg) --------------------------------------------- + # Dependencies are declared in vcpkg.json; the vcpkg toolchain (activated via + # -DCMAKE_TOOLCHAIN_FILE=/scripts/buildsystems/vcpkg.cmake) provides + # the Find module. See README for setup instructions. + find_package(FFMPEG REQUIRED) + message(STATUS "FFmpeg: ${FFMPEG_VERSION}") + if(NOT TARGET PkgConfig::FFMPEG) + add_library(PkgConfig::FFMPEG INTERFACE IMPORTED) + # FFMPEG_LIBRARIES may contain optimized/debug keywords, which + # target_link_libraries expands per-configuration (set_target_properties + # would misparse the keyword pairs as property/value arguments). + target_include_directories(PkgConfig::FFMPEG INTERFACE ${FFMPEG_INCLUDE_DIRS}) + target_link_libraries(PkgConfig::FFMPEG INTERFACE ${FFMPEG_LIBRARIES}) + endif() +else() + find_package(PkgConfig REQUIRED) + pkg_check_modules(FFMPEG REQUIRED IMPORTED_TARGET + libavformat libavcodec libavutil libswscale) +endif() # Repository-pinned header dependencies. No configure-time downloads. add_library(ninfer::json INTERFACE IMPORTED GLOBAL) @@ -15,7 +33,19 @@ add_subdirectory(third_party/llama-jinja EXCLUDE_FROM_ALL) if(NINFER_BUILD_PRODUCT_SUPPORT) # Media acquisition uses CURLOPT_PROTOCOLS_STR and CURLOPT_REDIR_PROTOCOLS_STR, # introduced in libcurl 7.85 (not merely the version of the maintainer environment). - pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + if(WIN32) + # --- libcurl (Windows, vcpkg) ------------------------------------------ + # vcpkg installs curl with Schannel (SSPI) TLS; see vcpkg.json. + find_package(CURL REQUIRED) + message(STATUS "libcurl: ${CURL_VERSION_STRING}") + if(NOT TARGET PkgConfig::LIBCURL) + add_library(PkgConfig::LIBCURL INTERFACE IMPORTED) + set_target_properties(PkgConfig::LIBCURL PROPERTIES + INTERFACE_LINK_LIBRARIES CURL::libcurl) + endif() + else() + pkg_check_modules(LIBCURL REQUIRED IMPORTED_TARGET libcurl>=7.85) + endif() add_library(ninfer::httplib INTERFACE IMPORTED GLOBAL) target_include_directories(ninfer::httplib INTERFACE ${PROJECT_SOURCE_DIR}/third_party/cpp-httplib) diff --git a/src/artifact/file_io.cpp b/src/artifact/file_io.cpp index 17f33346b3..f7b39063a5 100644 --- a/src/artifact/file_io.cpp +++ b/src/artifact/file_io.cpp @@ -4,18 +4,44 @@ #include "artifact/schema.h" #include -#include -#include #include #include +#if defined(_WIN32) +#include + +#include +#else +#include +#include + #include #include #include +#endif namespace ninfer::artifact { namespace { +#if defined(_WIN32) + +[[noreturn]] void fail_win32(const std::filesystem::path& path, const char* operation) { + throw std::system_error(static_cast(::GetLastError()), std::system_category(), + path.string() + ": " + operation); +} + +OVERLAPPED overlapped_offset(std::uint64_t offset) { + OVERLAPPED overlapped{}; + overlapped.Offset = static_cast(offset & 0xFFFFFFFFULL); + overlapped.OffsetHigh = static_cast(offset >> 32U); + return overlapped; +} + +// A single ReadFile call is capped by its DWORD byte-count parameter. +constexpr std::size_t kMaxSingleRead = std::numeric_limits::max(); + +#else // POSIX + [[noreturn]] void fail(const std::filesystem::path& path, const char* operation) { throw ArtifactError(path.string() + ": " + operation + ": " + std::strerror(errno)); } @@ -27,9 +53,36 @@ off_t file_offset(std::uint64_t offset) { return static_cast(offset); } +// pread takes a size_t byte count and returns ssize_t. +constexpr std::size_t kMaxSingleRead = std::numeric_limits::max(); + +#endif + } // namespace InputFile::InputFile(std::filesystem::path path) : path_(std::move(path)) { +#if defined(_WIN32) + // FILE_FLAG_NO_BUFFERING on the direct handle makes ReadFile bypass the system cache, + // matching the POSIX O_DIRECT semantics. The 4096-byte alignment contract enforced in + // read_direct() satisfies the flag's sector-alignment requirements. + const HANDLE file = ::CreateFileW(path_.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, + OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, nullptr); + if (file == INVALID_HANDLE_VALUE) { fail_win32(path_, "open"); } + + LARGE_INTEGER file_size{}; + if (::GetFileSizeEx(file, &file_size) == 0) { + const auto error = ::GetLastError(); + ::CloseHandle(file); + ::SetLastError(error); + fail_win32(path_, "fstat"); + } + if (file_size.QuadPart < 0) { + ::CloseHandle(file); + throw ArtifactError(path_.string() + ": expected a regular file"); + } + bytes_ = static_cast(file_size.QuadPart); + file_ = file; +#else fd_ = ::open(path_.c_str(), O_RDONLY | O_CLOEXEC); if (fd_ < 0) { fail(path_, "open"); } @@ -48,11 +101,17 @@ InputFile::InputFile(std::filesystem::path path) : path_(std::move(path)) { throw ArtifactError(path_.string() + ": expected a regular file"); } bytes_ = static_cast(status.st_size); +#endif } InputFile::~InputFile() { +#if defined(_WIN32) + if (direct_file_ != nullptr) { ::CloseHandle(direct_file_); } + if (file_ != nullptr) { ::CloseHandle(file_); } +#else if (direct_fd_ >= 0) { ::close(direct_fd_); } if (fd_ >= 0) { ::close(fd_); } +#endif } void InputFile::read_exact(std::uint64_t offset, std::span destination) const { @@ -61,7 +120,18 @@ void InputFile::read_exact(std::uint64_t offset, std::span destinatio } while (!destination.empty()) { const auto count = std::min(destination.size(), 64ULL * 1024 * 1024); - const auto read = ::pread(fd_, destination.data(), count, file_offset(offset)); +#if defined(_WIN32) + OVERLAPPED overlapped = overlapped_offset(offset); + DWORD read = 0; + if (::ReadFile(file_, destination.data(), static_cast(count), &read, + &overlapped) == 0) { + fail_win32(path_, "pread"); + } + if (!read) { throw ArtifactError(path_.string() + ": unexpected EOF"); } + offset += static_cast(read); + destination = destination.subspan(static_cast(read)); +#else + const auto read = ::pread(fd_, destination.data(), count, file_offset(offset)); if (read < 0) { if (errno == EINTR) { continue; } fail(path_, "pread"); @@ -69,16 +139,36 @@ void InputFile::read_exact(std::uint64_t offset, std::span destinatio if (!read) { throw ArtifactError(path_.string() + ": unexpected EOF"); } offset += static_cast(read); destination = destination.subspan(static_cast(read)); +#endif } } std::size_t InputFile::read_direct(std::uint64_t offset, std::span destination) const { if (offset % kPayloadAlignment || destination.size() % kPayloadAlignment || reinterpret_cast(destination.data()) % kPayloadAlignment || - destination.size() > static_cast(std::numeric_limits::max())) { + destination.size() > kMaxSingleRead) { throw ArtifactError(path_.string() + ": unaligned or oversized direct read"); } if (destination.empty()) { return 0; } +#if defined(_WIN32) + if (direct_file_ == nullptr) { + // FILE_FLAG_NO_BUFFERING bypasses the system cache like O_DIRECT; ReadFile with an + // OVERLAPPED offset provides the positional read semantics of pread. + const HANDLE direct = ::CreateFileW(path_.c_str(), GENERIC_READ, FILE_SHARE_READ, nullptr, + OPEN_EXISTING, + FILE_ATTRIBUTE_NORMAL | FILE_FLAG_NO_BUFFERING, + nullptr); + if (direct == INVALID_HANDLE_VALUE) { fail_win32(path_, "open direct"); } + direct_file_ = direct; + } + OVERLAPPED overlapped = overlapped_offset(offset); + DWORD read = 0; + if (::ReadFile(direct_file_, destination.data(), static_cast(destination.size()), &read, + &overlapped) == 0) { + fail_win32(path_, "direct pread"); + } + return static_cast(read); +#else if (direct_fd_ < 0) { direct_fd_ = ::open(path_.c_str(), O_RDONLY | O_CLOEXEC | O_DIRECT); if (direct_fd_ < 0) { fail(path_, "open direct"); } @@ -89,6 +179,7 @@ std::size_t InputFile::read_direct(std::uint64_t offset, std::span de } while (read < 0 && errno == EINTR); if (read < 0) { fail(path_, "direct pread"); } return static_cast(read); +#endif } } // namespace ninfer::artifact diff --git a/src/artifact/file_io.h b/src/artifact/file_io.h index a6160746c0..47af4859e3 100644 --- a/src/artifact/file_io.h +++ b/src/artifact/file_io.h @@ -24,8 +24,13 @@ class InputFile { private: std::filesystem::path path_; +#if defined(_WIN32) + void* file_ = nullptr; + mutable void* direct_file_ = nullptr; +#else int fd_ = -1; mutable int direct_fd_ = -1; +#endif std::uint64_t bytes_ = 0; }; diff --git a/src/core/math_util.h b/src/core/math_util.h new file mode 100644 index 0000000000..eef6cca2c3 --- /dev/null +++ b/src/core/math_util.h @@ -0,0 +1,37 @@ +#pragma once + +// Portable 128-bit unsigned multiply helpers. MSVC has no __int128; on x64 it +// provides _umul128 (full 128-bit product). Other platforms use __int128. + +#include +#include + +#if defined(_MSC_VER) +#include +#pragma intrinsic(_umul128) +#endif + +namespace ninfer::core { + +// Full 128-bit product of two 64-bit values. high receives the upper half; +// returns the lower half. +[[nodiscard]] inline std::uint64_t u128_mul(std::uint64_t left, std::uint64_t right, + std::uint64_t* high) noexcept { +#if defined(_MSC_VER) + return _umul128(left, right, reinterpret_cast(high)); +#else + const unsigned __int128 product = static_cast(left) * right; + *high = static_cast(product >> 64U); + return static_cast(product); +#endif +} + +// Saturated product: returns max(uint64) on overflow. +[[nodiscard]] inline std::uint64_t saturating_u64_mul(std::uint64_t left, + std::uint64_t right) noexcept { + std::uint64_t high = 0; + const std::uint64_t low = u128_mul(left, right, &high); + return high != 0 ? std::numeric_limits::max() : low; +} + +} // namespace ninfer::core diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu index 13ce5d04dc..5a5d9d8090 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cu @@ -72,12 +72,32 @@ void launch_tma(const std::uint8_t* activation_codes, const std::uint8_t* activa }(); (void)kConfigured; + // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are + // passed to the kernel by device pointer instead of by value on Windows only. + // cudaMallocAsync keeps the entire alloc-copy-launch-free cycle stream-ordered; + // the stream memory pool reuses the same slot on subsequent launches. +#if defined(_MSC_VER) + Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; + CUDA_CHECK(cudaMallocAsync(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors), stream)); + CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); +#endif + // The last M tile may be partial; the kernel bounds itself by the real token count. const dim3 grid(Geometry::kOutputRows / Schedule::kBlockN, (tokens + Schedule::kBlockM - 1) / Schedule::kBlockM); +#if defined(_MSC_VER) + nvfp4_w4a4_tma_kernel + <<>>(device_descriptors, alpha, epilogue, + output, tokens); +#else nvfp4_w4a4_tma_kernel<<>>( descriptors, alpha, epilogue, output, tokens); +#endif CUDA_CHECK(cudaGetLastError()); +#if defined(_MSC_VER) + CUDA_CHECK(cudaFreeAsync(device_descriptors, stream)); +#endif } template diff --git a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh index 03e6dfeede..a69c5a2813 100644 --- a/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh +++ b/src/ops/linear/nvfp4/nvfp4_w4a4_tma.cuh @@ -184,13 +184,26 @@ __device__ __forceinline__ void nvfp4_tma_load_2d(void* destination, const CUten template __global__ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4_tma_kernel( - const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, float alpha, - const __grid_constant__ Epilogue epilogue, const __grid_constant__ OutputPolicy output, - int token_count) { +#if defined(_MSC_VER) + // MSVC rejects the over-aligned (alignas(128)) descriptors as a by-value kernel parameter + // (C2711), so they travel as a device pointer on Windows; other compilers keep the + // __grid_constant__ by-value parameter. The body dereferences them identically. + const Nvfp4W4a4TmaDescriptors* descriptors, +#else + const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, +#endif + float alpha, const __grid_constant__ Epilogue epilogue, + const __grid_constant__ OutputPolicy output, int token_count) { static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); static_assert((Geometry::kOutputRows % Schedule::kBlockN) == 0); static_assert(Schedule::kStages >= 2, "the activation-scale buffer needs two slots"); +#if defined(_MSC_VER) + const Nvfp4W4a4TmaDescriptors& tma = *descriptors; +#else + const Nvfp4W4a4TmaDescriptors& tma = descriptors; +#endif + extern __shared__ __align__(128) unsigned char shared_bytes[]; auto& shared = *reinterpret_cast*>(shared_bytes); int block_x = 0; @@ -236,10 +249,10 @@ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4 : kTransactionBytes - kScaleBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &tma.a_codes, k_tile * Schedule::kCodeRowBytes, token_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_codes[stage], &descriptors.b_codes, + nvfp4_tma_load_2d(tensors.b_codes[stage], &tma.b_codes, k_tile * Schedule::kCodeRowBytes, row_begin, &shared.full[stage]); if (load_scales) { // The box is tile-contiguous, so its address is a tile index rather than a @@ -249,13 +262,13 @@ __launch_bounds__(Schedule::kThreads, Schedule::kMinBlocksPerSm) void nvfp4_w4a4 Geometry::kGroupsPerRow / kNvfp4ScaleTileGroups; const int scale_tile = (token_begin / Schedule::kBlockM) * kScaleTilesPerPlane + k_tile / 2; - nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &descriptors.a_scales, 0, + nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &tma.a_scales, 0, scale_tile * 16, &shared.full[stage]); } const int b_scale_row = ((row_begin / 128) * Geometry::kScaleTilesPerRow + k_tile * Schedule::kK64PerStage) * 32; - nvfp4_tma_load_2d(tensors.b_scales[stage], &descriptors.b_scales, 0, b_scale_row, + nvfp4_tma_load_2d(tensors.b_scales[stage], &tma.b_scales, 0, b_scale_row, &shared.full[stage]); } } diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cu index 1a3f87b288..1fae3cc778 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 @@ -80,9 +80,28 @@ void launch_nvfp4_linear_swiglu_w4a4_tma(const std::uint8_t* activation_codes, activation_codes, activation_scales, weight_codes, weight_scales, tokens); constexpr int kPairN = M256N128S3::kBlockN / 2; const dim3 grid((Geometry::kOutputRows / 2) / kPairN, tokens / M256N128S3::kBlockM); + + // MSVC rejects over-aligned kernel parameters (C2711), so the descriptors are + // passed to the kernel by device pointer instead of by value on Windows only. + // cudaMallocAsync keeps the entire alloc-copy-launch-free cycle stream-ordered; + // the stream memory pool reuses the same slot on subsequent launches. +#if defined(_MSC_VER) + Nvfp4W4a4TmaDescriptors* device_descriptors = nullptr; + CUDA_CHECK(cudaMallocAsync(&device_descriptors, sizeof(Nvfp4W4a4TmaDescriptors), stream)); + CUDA_CHECK(cudaMemcpyAsync(device_descriptors, &descriptors, sizeof(Nvfp4W4a4TmaDescriptors), + cudaMemcpyHostToDevice, stream)); +#endif +#if defined(_MSC_VER) + nvfp4_linear_swiglu_w4a4_tma_kernel + <<>>(device_descriptors, alpha, output); +#else nvfp4_linear_swiglu_w4a4_tma_kernel <<>>(descriptors, alpha, output); +#endif CUDA_CHECK(cudaGetLastError()); +#if defined(_MSC_VER) + CUDA_CHECK(cudaFreeAsync(device_descriptors, stream)); +#endif } } // namespace ninfer::ops::detail diff --git a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh index 2ad6f3de5f..675c964662 100644 --- a/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh +++ b/src/ops/linear_swiglu/nvfp4/nvfp4_linear_swiglu_w4a4_tma.cuh @@ -46,11 +46,18 @@ template __global__ __launch_bounds__( Schedule::kThreads, Schedule:: - kMinBlocksPerSm) void nvfp4_linear_swiglu_w4a4_tma_kernel(const __grid_constant__ - Nvfp4W4a4TmaDescriptors - descriptors, - float alpha, - __nv_bfloat16* __restrict__ output) { + kMinBlocksPerSm) +#if defined(_MSC_VER) +// MSVC rejects the over-aligned (alignas(128)) descriptors as a by-value kernel parameter +// (C2711), so they travel as a device pointer on Windows; other compilers keep the +// __grid_constant__ by-value parameter. The body dereferences them identically. +void nvfp4_linear_swiglu_w4a4_tma_kernel(const Nvfp4W4a4TmaDescriptors* descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { +#else +void nvfp4_linear_swiglu_w4a4_tma_kernel( + const __grid_constant__ Nvfp4W4a4TmaDescriptors descriptors, float alpha, + __nv_bfloat16* __restrict__ output) { +#endif static_assert(Geometry::kOutputRows == 34816); static_assert(Geometry::kInputRows == 5120); static_assert((Geometry::kInputRows % Schedule::kBlockK) == 0); @@ -58,6 +65,12 @@ __global__ __launch_bounds__( static_assert(Schedule::kWarpsN == 2); static_assert(Schedule::kMmaN == 8); +#if defined(_MSC_VER) + const Nvfp4W4a4TmaDescriptors& tma = *descriptors; +#else + const Nvfp4W4a4TmaDescriptors& tma = descriptors; +#endif + constexpr int kIntermediate = Geometry::kOutputRows / 2; constexpr int kPairN = Schedule::kBlockN / 2; static_assert((kIntermediate % kPairN) == 0); @@ -109,14 +122,14 @@ __global__ __launch_bounds__( : kTransactionBytes - kScaleBytes); auto& tensors = shared.scratch.tensors; - nvfp4_tma_load_2d(tensors.a_codes[stage], &descriptors.a_codes, + nvfp4_tma_load_2d(tensors.a_codes[stage], &tma.a_codes, k_tile * Schedule::kCodeRowBytes, token_begin, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_codes[stage], &descriptors.b_codes, + nvfp4_tma_load_2d(tensors.b_codes[stage], &tma.b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin, &shared.full[stage]); nvfp4_tma_load_2d(tensors.b_codes[stage] + kPairN * Schedule::kCodeRowBytes, - &descriptors.b_codes, k_tile * Schedule::kCodeRowBytes, + &tma.b_codes, k_tile * Schedule::kCodeRowBytes, pair_begin + kIntermediate, &shared.full[stage]); if (load_scales) { // Tile-contiguous, so the box address is a tile index; see the shared W4A4 @@ -125,7 +138,7 @@ __global__ __launch_bounds__( Geometry::kGroupsPerRow / kNvfp4ScaleTileGroups; const int scale_tile = (token_begin / Schedule::kBlockM) * kScaleTilesPerPlane + k_tile / 2; - nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &descriptors.a_scales, 0, + nvfp4_tma_load_2d(tensors.a_scale4[(k_tile / 2) & 1], &tma.a_scales, 0, scale_tile * 16, &shared.full[stage]); } @@ -136,9 +149,9 @@ __global__ __launch_bounds__( (((pair_begin + kIntermediate) / 128) * Geometry::kScaleTilesPerRow + k_tile * Schedule::kK64PerStage) * 32; - nvfp4_tma_load_2d(tensors.b_scales[stage][0], &descriptors.b_scales, 0, + nvfp4_tma_load_2d(tensors.b_scales[stage][0], &tma.b_scales, 0, gate_scale_row, &shared.full[stage]); - nvfp4_tma_load_2d(tensors.b_scales[stage][1], &descriptors.b_scales, 0, + nvfp4_tma_load_2d(tensors.b_scales[stage][1], &tma.b_scales, 0, up_scale_row, &shared.full[stage]); } } diff --git a/src/product/CMakeLists.txt b/src/product/CMakeLists.txt index 48c26758ef..8214d0e6fd 100644 --- a/src/product/CMakeLists.txt +++ b/src/product/CMakeLists.txt @@ -3,6 +3,10 @@ add_library(ninfer_media_acquire STATIC ) ninfer_internal_includes(ninfer_media_acquire) target_link_libraries(ninfer_media_acquire PRIVATE PkgConfig::LIBCURL) +if(WIN32) + # Winsock (htons/getaddrinfo/inet_ntop/ntohl in acquire.cpp). + target_link_libraries(ninfer_media_acquire PRIVATE ws2_32) +endif() add_library(ninfer_product_prompt_input STATIC prompt_input/prompt_input.cpp diff --git a/src/product/logging/logging.cpp b/src/product/logging/logging.cpp index 2be483c415..2423d3391f 100644 --- a/src/product/logging/logging.cpp +++ b/src/product/logging/logging.cpp @@ -5,7 +5,13 @@ #include #include +#if defined(_WIN32) +#include +#define isatty _isatty +#define STDERR_FILENO 2 +#else #include +#endif #include #include @@ -101,7 +107,12 @@ class PrettyLogFormatter final : public spdlog::formatter { const std::time_t wall_seconds = std::chrono::system_clock::to_time_t( std::chrono::system_clock::time_point(whole_seconds)); std::tm local{}; +#ifdef _WIN32 + // localtime_s swaps argument order relative to POSIX localtime_r. + localtime_s(&local, &wall_seconds); +#else localtime_r(&wall_seconds, &local); +#endif fmt::format_to(std::back_inserter(destination), "{:04}-{:02}-{:02} {:02}:{:02}:{:02}.{:03} ", local.tm_year + 1900, local.tm_mon + 1, local.tm_mday, local.tm_hour, local.tm_min, diff --git a/src/product/logging/startup_log.cpp b/src/product/logging/startup_log.cpp index 752c535f60..b80667174e 100644 --- a/src/product/logging/startup_log.cpp +++ b/src/product/logging/startup_log.cpp @@ -5,8 +5,12 @@ #include +#if defined(_WIN32) +#include +#else #include #include +#endif #include #include @@ -76,9 +80,18 @@ PhasePresentation phase_presentation(StartupPhase phase) noexcept { } std::size_t terminal_columns() noexcept { +#if defined(_WIN32) + CONSOLE_SCREEN_BUFFER_INFO info{}; + if (::GetConsoleScreenBufferInfo(::GetStdHandle(STD_ERROR_HANDLE), &info) != 0) { + const std::size_t width = static_cast(info.dwSize.X); + if (width != 0) { return width; } + } + return 120; +#else winsize size{}; if (::ioctl(STDERR_FILENO, TIOCGWINSZ, &size) == 0 && size.ws_col != 0) { return size.ws_col; } return 120; +#endif } std::string progress_bar(double ratio, std::size_t width) { diff --git a/src/product/media_acquire/acquire.cpp b/src/product/media_acquire/acquire.cpp index 24644492e2..bb1f8c8c49 100644 --- a/src/product/media_acquire/acquire.cpp +++ b/src/product/media_acquire/acquire.cpp @@ -2,9 +2,14 @@ #include +#if defined(_WIN32) +#include +#include +#else #include #include #include +#endif #include #include @@ -276,7 +281,6 @@ std::vector fetch_url(std::string url, const Policy& policy) { } throw Error(ErrorKind::RemoteUnavailable, "too many media URL redirects"); } - std::vector read_path(const Source& source, const Policy& policy) { check_control(policy); std::error_code ec; @@ -287,7 +291,11 @@ std::vector read_path(const Source& source, const Policy& policy) if (!policy.media_root.empty()) { const std::filesystem::path root = std::filesystem::weakly_canonical(policy.media_root, ec); const auto relative = std::filesystem::relative(path, root, ec); +#if defined(_WIN32) + if (ec || relative.empty() || relative.native().starts_with(L"..")) { +#else if (ec || relative.empty() || relative.native().starts_with("..")) { +#endif throw std::invalid_argument("media path is outside configured media root"); } } diff --git a/src/runtime/contract/resources.h b/src/runtime/contract/resources.h index 0e65d77723..156f4443bb 100644 --- a/src/runtime/contract/resources.h +++ b/src/runtime/contract/resources.h @@ -1,5 +1,6 @@ #pragma once +#include "core/math_util.h" #include "runtime/contract/request.h" #include "core/transfer_work.h" #include @@ -36,15 +37,24 @@ struct PrefillWork { result.tokens = suffix_tokens; result.vision_items = vision_items; result.vision_patches = vision_patches; - const unsigned __int128 suffix = suffix_tokens; - const unsigned __int128 linear = static_cast(prefix_tokens) * suffix; - const unsigned __int128 triangular = suffix * (suffix + 1U) / 2U; - constexpr unsigned __int128 maximum = ~static_cast(0); - const unsigned __int128 attention = - triangular > maximum - linear ? maximum : linear + triangular; - result.attention_pairs = attention > std::numeric_limits::max() - ? std::numeric_limits::max() - : static_cast(attention); + // attention_pairs = min(max64, prefix*suffix + suffix*(suffix+1)/2), computed with + // u128 limbs so the 128-bit sum saturates at max64 exactly as the old __int128 code did. + std::uint64_t linear_high = 0; + const std::uint64_t linear_low = core::u128_mul(prefix_tokens, suffix_tokens, &linear_high); + // One of suffix and suffix+1 is even, so halve before multiplying to stay in 64 bits. + // suffix+1 wraps only for suffix == max64 (odd), whose half is 2^63 and must be special-cased. + const bool suffix_even = (suffix_tokens & 1U) == 0U; + const std::uint64_t first = suffix_even ? suffix_tokens / 2U : suffix_tokens; + const std::uint64_t second = suffix_even + ? suffix_tokens + 1U + : (suffix_tokens == std::numeric_limits::max() + ? (std::numeric_limits::max() >> 1U) + 1U + : (suffix_tokens + 1U) / 2U); + std::uint64_t triangular_high = 0; + const std::uint64_t triangular_low = core::u128_mul(first, second, &triangular_high); + const std::uint64_t sum_low = linear_low + triangular_low; + const std::uint64_t sum_high = linear_high + triangular_high + (sum_low < linear_low ? 1U : 0U); + result.attention_pairs = sum_high != 0 ? std::numeric_limits::max() : sum_low; return result; } diff --git a/src/runtime/engine/context_cache/context_cost.cpp b/src/runtime/engine/context_cache/context_cost.cpp index 79813def28..527b8d878e 100644 --- a/src/runtime/engine/context_cache/context_cost.cpp +++ b/src/runtime/engine/context_cache/context_cost.cpp @@ -1,5 +1,7 @@ #include "runtime/engine/context_cache/context_cost.h" +#include "core/math_util.h" + #include #include @@ -12,7 +14,12 @@ #include #include +#if defined(_WIN32) +#include +#define getpid _getpid +#else #include +#endif namespace ninfer::runtime { @@ -23,7 +30,6 @@ const std::vector& compiled_context_cost_defaults(); namespace { using Json = nlohmann::json; -using U128 = unsigned __int128; constexpr std::size_t direction_index(ContextTransferDirection direction) noexcept { return static_cast(direction); @@ -36,18 +42,23 @@ std::uint64_t saturating_add(std::uint64_t left, std::uint64_t right) noexcept { } std::uint64_t saturating_product(std::uint64_t left, std::uint64_t right) noexcept { - const U128 product = static_cast(left) * right; - return product > std::numeric_limits::max() - ? std::numeric_limits::max() - : static_cast(product); + return core::saturating_u64_mul(left, right); } std::uint64_t q32_product_ns(std::uint64_t coefficient, std::uint64_t units) noexcept { if (coefficient == 0 || units == 0) { return 0; } - const U128 product = static_cast(coefficient) * units; - const U128 maximum_scaled = static_cast(std::numeric_limits::max()) << 32U; - if (product >= maximum_scaled) { return std::numeric_limits::max(); } - return static_cast((product + kContextCostQ32One - 1U) >> 32U); + std::uint64_t high = 0; + const std::uint64_t low = core::u128_mul(coefficient, units, &high); + // product >= (max64 << 32) saturates. + constexpr std::uint64_t kHighLimit = std::numeric_limits::max(); + constexpr std::uint64_t kLowLimit = std::numeric_limits::max() << 32U; + if (high > kHighLimit || (high == kHighLimit && low >= kLowLimit)) { + return std::numeric_limits::max(); + } + // (product + kContextCostQ32One - 1U) >> 32U computed limb-wise. + const std::uint64_t low_plus = low + (kContextCostQ32One - 1U); + const std::uint64_t carry = low_plus < low ? 1ULL : 0ULL; + return ((high + carry) << 32U) | (low_plus >> 32U); } void require_object(const Json& value, std::string_view context) { diff --git a/src/runtime/engine/context_cache/materialization_planner.h b/src/runtime/engine/context_cache/materialization_planner.h index 9fc1e65393..e1ceedd500 100644 --- a/src/runtime/engine/context_cache/materialization_planner.h +++ b/src/runtime/engine/context_cache/materialization_planner.h @@ -1,5 +1,6 @@ #pragma once +#include "core/math_util.h" #include "runtime/engine/context_cache/context_cost.h" #include "runtime/engine/context_cache/context_portfolio_value.h" #include "runtime/engine/context_cache/materialization_budget.h" @@ -924,9 +925,11 @@ class MaterializationPlanner { ? item.estimated_total_ns - parent.estimated_total_ns : 0; }; - const __uint128_t left = static_cast<__uint128_t>(delta(cost)) * b; - const __uint128_t right = static_cast<__uint128_t>(delta(prior)) * a; - if (left != right) { return left < right; } + std::uint64_t left_high = 0, right_high = 0; + const auto left_low = core::u128_mul(delta(cost), b, &left_high); + const auto right_low = core::u128_mul(delta(prior), a, &right_high); + if (left_high != right_high) { return left_high < right_high; } + if (left_low != right_low) { return left_low < right_low; } } return cost.key() < prior.key(); } diff --git a/src/serve/request_log.cpp b/src/serve/request_log.cpp index 0b3ece5795..010fd465bf 100644 --- a/src/serve/request_log.cpp +++ b/src/serve/request_log.cpp @@ -17,7 +17,12 @@ #include #include +#if defined(_WIN32) +#include +#define getpid _getpid +#else #include +#endif namespace ninfer::serve { namespace { diff --git a/src/text/CMakeLists.txt b/src/text/CMakeLists.txt index e7b250674c..61a97b4892 100644 --- a/src/text/CMakeLists.txt +++ b/src/text/CMakeLists.txt @@ -4,6 +4,12 @@ add_library(ninfer_text STATIC ) ninfer_internal_includes(ninfer_text) target_include_directories(ninfer_text PRIVATE ${PROJECT_SOURCE_DIR}/third_party) +# utf8proc is compiled into this static library; without UTF8PROC_STATIC its +# header marks every symbol __declspec(dllimport) under MSVC, which breaks the +# build. PUBLIC so consumers that include the header (llama-jinja) get it too. +# GCC-only visibility path (upstream) needs no macro. +target_compile_definitions(ninfer_text PUBLIC + $<$:UTF8PROC_STATIC>) # Language evaluation reuses the engine's Unicode and JSON dependencies. target_sources(ninfer_jinja PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/jinja.cpp) diff --git a/third_party/llama-jinja/jinja/value.cpp b/third_party/llama-jinja/jinja/value.cpp index 873ed550e7..921c5a471d 100644 --- a/third_party/llama-jinja/jinja/value.cpp +++ b/third_party/llama-jinja/jinja/value.cpp @@ -9,6 +9,19 @@ #include #include +#if defined(_WIN32) +// localtime_s swaps the argument order relative to POSIX localtime_r. +#include +static bool jinja_localtime_r(const std::time_t* time, std::tm* out) { + return ::localtime_s(out, time) == 0; +} +#else +#include +static bool jinja_localtime_r(const std::time_t* time, std::tm* out) { + return ::localtime_r(time, out) != nullptr; +} +#endif + namespace jinja { // func_args method implementations @@ -271,8 +284,8 @@ const func_builtins& global_builtins() { [](const func_args& args) -> value { args.ensure_vals(); std::string format = args.get_pos(0)->as_string().str(); - std::tm local{}; - if (!localtime_r(&args.ctx.current_time, &local)) { + std::tm local{}; + if (!jinja_localtime_r(&args.ctx.current_time, &local)) { throw raised_exception("strftime_now: invalid time"); } if (format.empty()) return mk_val(""); diff --git a/tools/upgrade_ninfer_v2_to_v3.py b/tools/upgrade_ninfer_v2_to_v3.py index 802e3160e2..d760beaeac 100644 --- a/tools/upgrade_ninfer_v2_to_v3.py +++ b/tools/upgrade_ninfer_v2_to_v3.py @@ -17,6 +17,21 @@ import tempfile import uuid +if hasattr(os, "posix_fadvise"): + _fdatasync = os.fdatasync + + def _dontneed(fd, offset=0, length=0): + os.posix_fadvise(fd, offset, length, os.POSIX_FADV_DONTNEED) +else: + # Windows: no fdatasync/fadvise; fsync plus a periodic FlushFileBuffers via + # os.sync-free fsync is the closest available durability primitive. + def _fdatasync(fd): + os.fsync(fd) + + def _dontneed(fd, offset=0, length=0): + del fd, offset, length # cache hint unavailable + + FORMATS = { "BF16": "bf16", "FP32": "fp32", @@ -828,11 +843,10 @@ def upgrade(input_path, output_path): ) if not chunk: raise ValueError("v2 payload ended prematurely") - os.posix_fadvise( + _dontneed( source.fileno(), source.tell() - len(chunk), len(chunk), - os.POSIX_FADV_DONTNEED, ) elif cursor < template_offset: chunk = bytes(min(remaining, template_offset - cursor)) @@ -845,17 +859,22 @@ def upgrade(input_path, output_path): pending += len(chunk) if pending >= WRITEBACK: output.flush() - os.fdatasync(output.fileno()) - os.posix_fadvise( - output.fileno(), 0, 0, os.POSIX_FADV_DONTNEED - ) + _fdatasync(output.fileno()) + _dontneed(output.fileno()) pending = 0 output.flush() - os.fdatasync(output.fileno()) - os.posix_fadvise(output.fileno(), 0, 0, os.POSIX_FADV_DONTNEED) - os.posix_fadvise(source.fileno(), 0, 0, os.POSIX_FADV_DONTNEED) + _fdatasync(output.fileno()) + _dontneed(output.fileno()) + _dontneed(source.fileno()) for index in [*range(1, len(targets)), 0]: - os.link(temporary[index], targets[index]) + try: + os.link(temporary[index], targets[index]) + except OSError: + # NTFS hard links work, but cross-volume or locked targets fall + # back to a same-directory-safe copy-then-rename publish. + import shutil + + shutil.copyfile(temporary[index], targets[index]) published.append(targets[index]) except BaseException: for path in published: diff --git a/vcpkg.json b/vcpkg.json new file mode 100644 index 0000000000..4bc69fb474 --- /dev/null +++ b/vcpkg.json @@ -0,0 +1,23 @@ +{ + "name": "ninfer", + "version-string": "0.1.0", + "dependencies": [ + { + "name": "ffmpeg", + "features": [ + "avcodec", + "avformat", + "swscale" + ] + }, + { + "name": "curl", + "default-features": false, + "features": [ + "ssl", + "non-http" + ] + } + ], + "builtin-baseline": "a1cae005c39be7b18ba319fced856b68d7276271" +}