diff --git a/src/CMakeLists.txt b/src/CMakeLists.txt index 789221bbe..0baa990aa 100644 --- a/src/CMakeLists.txt +++ b/src/CMakeLists.txt @@ -503,6 +503,7 @@ list(APPEND FLM_ENGINE_LINK_LIBS qwen3vl_flash) list(APPEND FLM_ENGINE_LINK_LIBS qwen3_5vl_npu minicpm_v_4_7_npu + minicpm_ocr_npu minicpm_v_4_7_moe_npu qwen3_5_omni_npu qwen3_6_moe_npu diff --git a/src/common/AutoModel/builtin_backends.cpp b/src/common/AutoModel/builtin_backends.cpp index 3cf755365..8fabbc6f0 100644 --- a/src/common/AutoModel/builtin_backends.cpp +++ b/src/common/AutoModel/builtin_backends.cpp @@ -46,6 +46,7 @@ void register_builtin_backends(BackendRegistry& registry) { RegisterFlm(registry, "qwen3vl-flash"); RegisterFlm(registry, "qwen3.5"); RegisterFlm(registry, "minicpm-v-4.7"); + RegisterFlm(registry, "minicpm-ocr"); RegisterFlm(registry, "minicpm-v-4.7-moe"); RegisterFlm(registry, "qwen3.6-moe"); RegisterFlm(registry, "qwen3.8-mtp"); diff --git a/src/common/AutoModel/modeling_minicpm_ocr.cpp b/src/common/AutoModel/modeling_minicpm_ocr.cpp new file mode 100644 index 000000000..0c79d4faf --- /dev/null +++ b/src/common/AutoModel/modeling_minicpm_ocr.cpp @@ -0,0 +1,472 @@ +/// \file modeling_minicpm_ocr.cpp +/// \brief MiniCPM_OCR class +/// \author FastFlowLM Team +/// \date 2026-10-02 +/// \version 0.9.28 +/// \note This is a source file for the MiniCPM_OCR class + + +#include "AutoModel/modeling_minicpm_ocr.hpp" + +#include +#include + + +/************ MiniCPM_OCR family **************/ +MiniCPM_OCR::MiniCPM_OCR(flm_rt::device* npu_device_inst) : AutoModel(npu_device_inst, "MiniCPM_OCR") {} + +void MiniCPM_OCR::load_model(std::string model_path, json model_info, int default_context_length, bool enable_preemption, const std::string& backend) { + this->_shared_load_backend(model_path, model_info, default_context_length, enable_preemption, backend); + this->setup_tokenizer(model_path); + this->sampler.reset(); + + // The shipped template renders tools in the same XML format as Qwen3.5. + this->enable_tool = true; + + // MiniCPM-V ships no sampling defaults of its own; its text tower is + // Qwen3.5-0.8B's, so start from the Qwen3.5 non-thinking recommendation. + sampler_config config; + config.top_k = 20; + config.top_p = 0.8; + config.min_p = 0.0; + config.temperature = 0.7; + config.rep_penalty = 1.0; + config.freq_penalty = 1.0; + config.pre_penalty = 1.5f; + + this->set_sampler(config); + for (size_t i = 0; i < PROFILER_TYPE_NUM; i++) { + this->profiler_list[i].reset(); + } + this->_prefix_attach(this->lm_engine, minicpm_ocr_state_api(this->lm_engine), 248056); // 248056 = <|image_pad|> +} + +void MiniCPM_OCR::setup_tokenizer(std::string model_path) { + auto tokenizer_config = this->_shared_setup_tokenizer(model_path); +} + +std::string MiniCPM_OCR::apply_chat_template(nlohmann::ordered_json& messages, nlohmann::ordered_json tools, bool continue_final_message) { + minja::chat_template_inputs inputs; + inputs.messages = messages; + inputs.extra_context = this->extra_context; + inputs.extra_context["enable_thinking"] = this->enable_think; + if (!tools.empty() && this->enable_tool) + inputs.tools = tools; + return this->_render_chat_template(inputs, continue_final_message); +} + +/// \brief Report why inference failed, then raise the caller-facing error. +/// +/// Every call site reaches here from a `catch (...)`, which on its own destroys the only +/// description of what went wrong: the user is told "inference failed" and the engine's actual +/// message -- an XRT error, a bad shape, a throw from the vision tower -- is gone. That matters +/// most for the failures you cannot reproduce on demand, where the one occurrence is all the +/// evidence there will ever be. +/// +/// `std::current_exception()` is still live inside a catch block, so rethrowing it here recovers +/// the message without changing any control flow. +void MiniCPM_OCR::fail_inference() { + if (std::current_exception()) { + try { + throw; + } catch (const std::exception& e) { + header_print("ERROR", "inference failed: " << e.what()); + } catch (...) { + header_print("ERROR", "inference failed: non-std exception"); + } + } + + const bool poisoned = this->backend_ && this->backend_->poisoned(); + this->_shared_after_inference_failure(poisoned); + throw ModelRequestError(500, true, poisoned + ? "Inference failed; unload/reload is required because the model is poisoned" + : "Inference failed; the current conversation was cleared"); +} + +bool MiniCPM_OCR::insert(chat_meta_info_t& meta_info, lm_uniform_input_t& input, std::function is_cancelled) { + this->_shared_guard_poisoned(); + constexpr int image_pad = 248056; + // preprocess + this->profiler_list[TKOEN_ENCODE_TIME].start(); + std::string templated_text; + if (input.messages.empty() && input.prompt.empty()) { + header_print("WARNING", "No messages or prompt provided"); + return false; + } + + if (!input.audios.empty()) { + header_print("WARNING", "MiniCPM-V has no audio input, ignoring " << input.audios.size() << " audio input(s)"); + input.audios.clear(); + input.audio_payload_types.clear(); + } + + // is_vlm is set from config.json's vision_model_weight; without it the engine + // has no vision tower, so images are dropped here rather than failing there. + const bool vlm = this->lm_config && this->lm_config->get("is_vlm", false); + + // One layout per picture, in prompt order; the payload holds their VIEWS. + minicpm_ocr_image_payload_t image_payload; + image_payload.num_images = 0; + std::vector layouts; + auto add_image = [&](const std::string& source, bool base64) { + image_data_t chw; + if (!this->load_image(source, base64, chw)) { + header_print("ERROR", "Skipping image that failed to load"); + return false; + } + layouts.push_back(this->preprocess_image(chw, image_payload)); + image_reader_.recycle(chw); + return true; + }; + + if (!input.messages.empty()) { // already a formated messages, usually from REST API + nlohmann::ordered_json messages = nlohmann::ordered_json::array(); + for (auto& message : input.messages) { + message.erase("audios"); + if (!message.contains("images")) { + messages.push_back(message); + continue; + } + if (!vlm) { + header_print("WARNING", "this MiniCPM-V package has no vision tower, ignoring " << message["images"].size() << " image(s)"); + message.erase("images"); + messages.push_back(message); + continue; + } + // Decode/preprocess up front so an invalid image never gets a + // placeholder in the template below. + nlohmann::ordered_json content = nlohmann::ordered_json::array(); + for (const auto& img : message["images"]) { + if (add_image(img.get(), true)) { + content.push_back({ {"type", "image"} }); + } + } + content.push_back({ {"type", "text"}, {"text", message["content"]} }); + messages.push_back({ {"role", message["role"]}, {"content", content} }); + } + templated_text = this->apply_chat_template(messages, input.tools, meta_info.continue_final_message); + } + else if (!input.prompt.empty()) { // a pure text, usually from the cli + nlohmann::ordered_json messages; + nlohmann::ordered_json content = nlohmann::ordered_json::array(); + if (!input.images.empty() && !vlm) { + header_print("WARNING", "this MiniCPM-V package has no vision tower, ignoring " << input.images.size() << " image(s)"); + } + else { + for (const auto& path : input.images) { + if (add_image(path, false)) { + content.push_back({ {"type", "image"} }); + } + } + } + content.push_back({ {"type", "text"}, {"text", input.prompt} }); + messages.push_back({ {"role", "user"}, {"content", content} }); + templated_text = this->apply_chat_template(messages); + } + input.images.clear(); + input.image_payload_types.clear(); + + // Expand each picture's single <|image_pad|> to the processor's placeholder. + std::vector tokens_init = this->tokenizer->encode(templated_text); + std::vector tokens; + { + size_t extra = 0; + for (const auto& l : layouts) extra += (size_t)l.tokens + 16 + 3 * (size_t)l.views; + tokens.reserve(tokens_init.size() + extra); + size_t picture = 0; + for (int id : tokens_init) { + if (id != image_pad) { + tokens.push_back(id); + continue; + } + if (picture >= layouts.size()) { + header_print("ERROR", "the prompt has more <|image_pad|> than images"); + return false; + } + const std::vector ph = this->image_placeholder(layouts[picture], (int)picture); + tokens.insert(tokens.end(), ph.begin(), ph.end()); + picture++; + } + if (picture != layouts.size()) { + header_print("ERROR", "the prompt has " << picture << " <|image_pad|> for " << layouts.size() << " images"); + return false; + } + } + + this->profiler_list[TKOEN_ENCODE_TIME].stop(tokens.size()); + + // ---- prefix slots: resume from the longest saved prefix of this prompt ---- + // The generation prompt ends in "\n" (2 tokens) or, with thinking off, "\n\n\n\n" (4): decode() feeds them, + // the next turn does not render them, so they are left out of what is saved. A continued message has no generation + // prompt: all its tokens are the client's text and stay. + const size_t think_trim = meta_info.continue_final_message ? 0 : (this->enable_think ? 2 : 4); + const bool slot_mode = this->_prefix_slot_mode(meta_info, tokens.size(), think_trim); + size_t match_len = 0; + std::vector full_tokens; // what this prompt leaves in the engine + std::vector image_hashes; // content of each view of its pictures + if (slot_mode) { + minicpm_ocr_npu* slot_eng = dynamic_cast(this->lm_engine); + full_tokens.assign(tokens.begin(), tokens.end() - think_trim); + image_hashes = prefix_image_hashes(image_payload, slot_eng->MINICPM_OCR_PATCH_SIZE, slot_eng->MINICPM_OCR_TEMPORAL_PATCH_SIZE); + match_len = this->_prefix_resume(this->prefix_cache_, full_tokens, image_hashes); + // The single checkpoint keeps only the recurrent state and relies on the attention rows staying as they were: a slot request + // overwrites them, so drop its token history; the path below then has nothing to restore. + checkpoint_his.clear(); + } + else { + this->_prefix_legacy_path(); // the single-checkpoint path below moves the engine on without telling the cache + } + + if (!slot_mode && meta_info.restore_allowed && !checkpoint_his.empty()) { + const int restore_idx = this->lm_engine->restore(); + if (restore_idx >= 0) { + this->total_tokens = restore_idx; + this->token_history = checkpoint_his; // restore the token history to be consistent with the restored KV cache, which is crucial for correct functioning of _shared_insert's prefix-matching logic + } + } + + // The generation prompt ends in "\n" (2 tokens) or, with thinking + // off, "\n\n\n\n" (4 tokens). Leave them out of the prompt so + // the checkpoint lands right after "<|im_start|>assistant\n"; decode() + // feeds them back. + tokens.resize(tokens.size() - think_trim); + + // ---------------------------------------------------------------------- + // Prompt-cache aware image alignment. _shared_insert erases the prefix of + // `tokens` that matches token_history -- only when it matches ALL of + // token_history, else it clears the context and skips nothing. The payload + // must lose exactly the pictures that erased prefix fully covers (all of + // their views) so it lines up with the image rows that are prefilled. + // + // This is _shared_insert's rule on _shared_insert's inputs: token_history + // AFTER the restore above and `tokens` AFTER the think trim. Computing it + // against checkpoint_his instead (Qwen3_5VL's way) disagrees in the CLI, + // where token_history also holds the previous answer: a repeated question + // then "hits" here, its image is dropped, _shared_insert clears and + // re-prefills the whole prompt, and the engine is handed image rows with + // no pixels. + // ---------------------------------------------------------------------- + size_t prefix_skip_count = 0; + { + const size_t idx = this->token_history.size(); + while (prefix_skip_count < idx && prefix_skip_count < tokens.size() && + tokens[prefix_skip_count] == this->token_history[prefix_skip_count]) { + prefix_skip_count++; + } + if (prefix_skip_count != idx) prefix_skip_count = 0; + + if (prefix_skip_count > 0 && !layouts.empty()) { + int skipped_rows = 0; + for (size_t i = 0; i < prefix_skip_count; i++) skipped_rows += (tokens[i] == image_pad); + size_t drop_pictures = 0, drop_views = 0, drop_values = 0; + int consumed = 0; + for (const auto& l : layouts) { + if (consumed + l.tokens > skipped_rows) break; + consumed += l.tokens; + drop_pictures++; + drop_views += (size_t)l.views; + drop_values += l.values; + } + if (drop_pictures > 0) { + image_payload.images.erase(image_payload.images.begin(), image_payload.images.begin() + drop_views); + image_payload._data__processed.erase(image_payload._data__processed.begin(), + image_payload._data__processed.begin() + drop_values); + image_payload.num_images = (unsigned)image_payload.images.size(); + header_print("FLM", "Prompt-cache hit: dropped " << drop_pictures << " cached image(s) from payload"); + } + } + } + + // The payload rides on the first prefill chunk only, so that chunk must hold + // every image GROUP whole -- through the / that close it, + // which the canvas m-rope positions too -- not just up to the last pad row. + int first_len_run = 0; + { + int last = -1; + for (int i = (int)prefix_skip_count; i < (int)tokens.size(); i++) { + if (tokens[i] == image_pad) last = i; + } + if (last >= 0) { + auto structural = [](int t) { + return t == 248078 || t == 248079 || t == 248088 || t == 248089; + }; + while (last + 1 < (int)tokens.size() && structural(tokens[last + 1])) last++; + first_len_run = last + 1 - (int)prefix_skip_count; + } + } + const bool has_images = !image_payload.images.empty(); + + // FLM_MINICPM_OCR_DUMP=: the payload and the prompt exactly as the engine + // gets them, for scoring the preprocessing against the HF processor + // (pixels.bf16 raw, views.txt "h w" per view, tokens.txt one id per line). + // NOTE tokens.txt is the PREFILL set, already shortened by think_trim above, so it + // ends at "<|im_start|>assistant\n" and is 2 (or 4) tokens shorter than what the + // model actually sees -- decode() feeds the "" tokens back. Do not compare it + // against HF's apply_chat_template() output directly and conclude the prompts differ. + if (const char* dump = std::getenv("FLM_MINICPM_OCR_DUMP"); dump && *dump) { + const std::string dir(dump); + std::ofstream(dir + "/pixels.bf16", std::ios::binary) + .write(reinterpret_cast(image_payload._data__processed.data()), + image_payload._data__processed.size() * sizeof(bf16)); + std::ofstream vf(dir + "/views.txt"); + for (const auto& v : image_payload.images) vf << v.grid_h << " " << v.grid_w << "\n"; + std::ofstream tf(dir + "/tokens.txt"); + for (int t : tokens) tf << t << "\n"; + header_print("FLM", "dumped the image payload and prompt to " << dir); + } + + if (slot_mode) { + try { + return this->_prefix_insert(this->prefix_cache_, meta_info, tokens, full_tokens, image_hashes, match_len, "<|im_start|>user", + has_images ? static_cast(&image_payload) : nullptr, first_len_run, is_cancelled, + input.requested_max_new_tokens); + } catch (const ModelRequestError&) { + throw; + } catch (...) { + this->fail_inference(); + } + } + + bool success = false; + try { + success = this->_shared_insert(meta_info, tokens, is_cancelled, has_images ? &image_payload : nullptr, + has_images ? first_len_run : 0, input.requested_max_new_tokens); + } catch (const ModelRequestError&) { + throw; + } catch (...) { + this->fail_inference(); + } + + checkpoint_his = token_history; + this->lm_engine->checkpoint(); + return success; +} + +std::string MiniCPM_OCR::generate(chat_meta_info_t& meta_info, int length_limit, std::ostream& os, std::function is_cancelled) { + this->_shared_guard_poisoned(); + try { + return this->decode(meta_info, length_limit, os, std::move(is_cancelled)); + } catch (const ModelRequestError&) { + throw; + } catch (...) { + this->fail_inference(); + } +} + +std::string MiniCPM_OCR::decode(chat_meta_info_t& meta_info, int length_limit, std::ostream& os, std::function is_cancelled) { + std::string result; + assert(this->last_token != -1); + stop_reason_t reason = EOT_DETECTED; + + this->profiler_list[DECODING_TIME].reset(); + this->profiler_list[TKOEN_DECODE_TIME].reset(); + // Some backends refuse to decode past a limit of their own, below MAX_L. + const uint32_t decode_cap = this->decode_cap(); + + // The tail of the generation prompt that insert() held back. A continued + // message held nothing back: it goes on from the sample insert() made. + const std::vector think_prefix = meta_info.continue_final_message + ? std::vector{} + : this->enable_think + ? std::vector{ think_start_id, 198 } // \n + : std::vector{ think_start_id, 271, think_end_id, 271 }; // \n\n\n\n + if (this->total_tokens + think_prefix.size() >= decode_cap) { + header_print("WARNING", "Max length reached, stopping generation..."); + meta_info.stop_reason = MAX_LENGTH_REACHED; + return result; + } + // The first prefix forward is left out of DECODING_TIME, as Qwen3_5VL does, + // so `flm bench` measures decode speed the same way as for qwen3.5:0.8b. + buffer y; + if (!this->enable_think && !think_prefix.empty()) { + // fed for the KV cache only, it never reaches the output + this->_skip_insert_sample(meta_info); + } + for (size_t i = 0; i < think_prefix.size(); i++) { + if (this->enable_think) this->_force_token(meta_info, think_prefix[i]); + this->token_history.push_back(think_prefix[i]); + if (i > 0) this->profiler_list[DECODING_TIME].start(); + y = this->lm_engine->forward(think_prefix[i]); + if (i > 0) this->profiler_list[DECODING_TIME].stop(1); + this->total_tokens++; + } + if (this->enable_think) { + // Echo the opening tag so the output parser enters reasoning mode. + for (int id : think_prefix) { + std::string token_str = this->tokenizer->run_time_decoder(id); + result += token_str; + os << token_str << std::flush; + } + } + int sampled_token = think_prefix.empty() ? this->last_token : this->_sample(y, meta_info); + + while (true) { + this->profiler_list[TKOEN_DECODE_TIME].start(); + if (this->is_normal_token(sampled_token)){ // filter out special tokens + std::string token_str = this->tokenizer->run_time_decoder(sampled_token); + os << token_str << std::flush; + result += token_str; + } + this->profiler_list[TKOEN_DECODE_TIME].stop(1); + this->token_history.push_back(sampled_token); + meta_info.generated_tokens++; + + if (this->is_eos(sampled_token)){ + // Keep the kv cache aligned with token_history for a following turn. + if (this->forward_on_eos && + (!this->backend_ || this->backend_->forwards_past_eos()) && + this->total_tokens < decode_cap) { + this->lm_engine->forward(sampled_token); + this->total_tokens++; + } else { + // The engine never ran it, so the next prompt has to carry it. + this->token_history.pop_back(); + } + break; + } + if ((length_limit > 0) && (meta_info.generated_tokens >= length_limit)){ + reason = MAX_LENGTH_REACHED; + break; + } + if (this->total_tokens >= decode_cap){ + header_print("WARNING", "Max length reached, stopping generation..."); + reason = MAX_LENGTH_REACHED; + break; + } + if (is_cancelled()) { + reason = CANCEL_DETECTED; + break; + } + + this->profiler_list[DECODING_TIME].start(); + y = this->lm_engine->forward(sampled_token); + this->profiler_list[DECODING_TIME].stop(1); + this->total_tokens++; + + this->profiler_list[SAMPLING_TIME].start(); + sampled_token = this->_sample(y, meta_info); + this->profiler_list[SAMPLING_TIME].stop(1); + } + meta_info.decoding_duration = (uint64_t)(time_utils::cast_to_us(this->profiler_list[DECODING_TIME].get_total_time()).first) * 1e3; + meta_info.stop_reason = reason; + if (reason != EOT_DETECTED) { + // Cut short after sampling its last token: no engine has run that one, + // so the next prompt has to carry it. + this->token_history.pop_back(); + } + + std::cout << std::endl; + if (this->log_raw_output) { + header_print("FLM", "Model RAW Output: \n" + result); + } + return result; +} + +std::string MiniCPM_OCR::generate_with_prompt(chat_meta_info_t& meta_info, lm_uniform_input_t& input, int length_limit, std::ostream& os) { + if (!this->insert(meta_info, input)) { + return ""; + } + // Not _shared_generate: insert() held back the think prefix, which only + // decode() knows to feed. + return this->generate(meta_info, length_limit, os); +} diff --git a/src/common/AutoModel/modeling_minicpm_ocr_image.cpp b/src/common/AutoModel/modeling_minicpm_ocr_image.cpp new file mode 100644 index 000000000..33dc5a46c --- /dev/null +++ b/src/common/AutoModel/modeling_minicpm_ocr_image.cpp @@ -0,0 +1,331 @@ +/// \file modeling_minicpm_ocr_image.cpp +/// \brief MiniCPM_OCR image preprocessing: MiniCPMV4ImageProcessorPil in C++. +/// \author FastFlowLM Team +/// \date 2026-10-05 +/// \note Port of the checkpoint's own image_processing_pil_minicpmv4.py and of +/// the placeholder expansion in processing_minicpmv4.py, rule for rule. +/// Every number here (the 448 scale, the factor 4 in ensure_divide, the +/// slice-grid search) changes how many views and tokens an image produces, +/// and the engine checks that the prompt's image rows match its views +/// exactly -- a wrong rule fails loudly there rather than quietly here. +/// +/// Layout handed to the engine (minicpm_ocr_npu.cpp, _prefill_with_mm): one +/// minicpm_ocr_image_t per VIEW, and `_data__processed` holding every view's +/// patches back to back -- raster order within a view, each patch flattened +/// (c, kh, kw) as 3*14*14 bf16. That is the NaViT `[3, 14, T*14]` tensor the +/// HF processor emits, transposed per patch. + +#include "AutoModel/modeling_minicpm_ocr.hpp" + +#include +#include +#include + +namespace { + +// ---- the processor's arithmetic, as Python computes it ------------------------ +// Python's round() is half-to-even and int() truncates; nearbyint and a cast +// are those two exactly. + +/// `max(round(length / divisor) * divisor, divisor)` +int ensure_divide(double length, int divisor) { + return std::max((int)std::nearbyint(length / divisor) * divisor, divisor); +} + +/// `find_best_resize()`: (height, width) in pixels, multiples of `patch * merge`. +/// The inputs are floats on the refine path. +/// +/// `merge` is NOT a constant, and getting it wrong is silent: every image comes out at a +/// slightly different resolution than the HF processor chose, so the token counts drift and +/// the prompt stops being byte-identical to upstream -- with no error anywhere. +/// +/// image_processing_minicpmv4_6.py: +/// # "4x" merges one 2x2 block, "16x" merges twice (ViT insert merger + downsample MLP) +/// merge_factor = 2 if downsample_mode == "4x" else 4 +/// +/// MiniCPM-V-4.7 is "16x" and merges twice; MiniCPM-OCR is "4x" and merges once, because the +/// windowed vit_merger is skipped entirely. `merge` is therefore sqrt(DOWNSAMPLE_FACTOR): +/// 4 for 16x, 2 for 4x. +void find_best_resize(double height, double width, int scale, int patch, int merge, + bool allow_upscale, int& best_h, int& best_w) { + if (height * width > (double)scale * scale || allow_upscale) { + const double ar = width / height; + const int h = (int)(scale / std::sqrt(ar)); + const int w = (int)(h * ar); + height = h; + width = w; + } + best_w = ensure_divide(width, patch * merge); + best_h = ensure_divide(height, patch * merge); +} + +/// `get_sliced_grid()`: (rows, cols) of slices, or false for no slicing. +bool sliced_grid(int h, int w, int max_slices, int scale, int& rows, int& cols) { + const double log_ratio = std::log((double)w / h); + const double ratio = (double)w * h / ((double)scale * scale); + const int multiple = std::min((int)std::ceil(ratio), max_slices); + if (multiple <= 1) return false; + + rows = cols = 1; + double min_error = std::numeric_limits::infinity(); + for (int num_slices : {multiple - 1, multiple, multiple + 1}) { + if (num_slices == 1 || num_slices > max_slices) continue; + for (int r = 1; r <= num_slices; r++) { + if (num_slices % r) continue; + const int c = num_slices / r; + const double error = std::fabs(log_ratio - std::log((double)c / r)); + if (error < min_error) { rows = r; cols = c; min_error = error; } + else if (error == min_error && r > rows) { rows = r; cols = c; } + } + } + return true; +} + +// ---- Pillow's bicubic resample, bit for bit ------------------------------------ +// +// The HF processor resizes the uint8 image with PIL (Image.resize, BICUBIC; the +// output is exactly PIL's, checked). The shared AVX-512 bicubic +// (imgproc::avx512::resize_bicubic_antialias_rgb_planar_avx512) is close but not +// the same function -- it disagrees with PIL by up to 29 uint8 levels at +// high-contrast pixels (relL2 0.013 on the reference photo), and SigLIP +// amplifies that ~9x: image_embeds relL2 0.116 through HF's own fp32 tower, +// more than every device rounding in the tower put together (0.080). So this +// tower gets Pillow's own algorithm (libImaging/Resample.c, ImagingResample): +// double-precision coefficients normalised per output pixel, quantised to +// 22-bit fixed point, a horizontal pass into a CLAMPED uint8 image, then a +// vertical pass. + +constexpr int PIL_PRECISION_BITS = 32 - 8 - 2; + +double pil_bicubic(double x) { + constexpr double a = -0.5; + if (x < 0.0) x = -x; + if (x < 1.0) return ((a + 2.0) * x - (a + 3.0)) * x * x + 1; + if (x < 2.0) return (((x - 5) * x + 8) * x - 4) * a; + return 0.0; +} + +/// precompute_coeffs() + normalize_coeffs_8bpc(): per output pixel, the first +/// input index, the tap count, and ksize fixed-point weights. +int pil_coeffs(int in_size, int out_size, std::vector& bounds, std::vector& kk) { + const double scale = (double)in_size / out_size; + const double filterscale = std::max(scale, 1.0); + const double support = 2.0 * filterscale; // bicubic support 2 + const int ksize = (int)std::ceil(support) * 2 + 1; + bounds.assign((size_t)out_size * 2, 0); + kk.assign((size_t)out_size * ksize, 0); + std::vector k((size_t)ksize); + for (int xx = 0; xx < out_size; xx++) { + const double center = (xx + 0.5) * scale; + const double ss = 1.0 / filterscale; + int xmin = (int)(center - support + 0.5); + if (xmin < 0) xmin = 0; + int xmax = (int)(center + support + 0.5); + if (xmax > in_size) xmax = in_size; + xmax -= xmin; + double ww = 0.0; + for (int x = 0; x < xmax; x++) { + const double w = pil_bicubic((x + xmin - center + 0.5) * ss); + k[(size_t)x] = w; + ww += w; + } + for (int x = 0; x < xmax; x++) { + double w = ww != 0.0 ? k[(size_t)x] / ww : k[(size_t)x]; + kk[(size_t)xx * ksize + x] = w < 0 ? (int32_t)(-0.5 + w * (1 << PIL_PRECISION_BITS)) + : (int32_t)(0.5 + w * (1 << PIL_PRECISION_BITS)); + } + bounds[(size_t)xx * 2] = xmin; + bounds[(size_t)xx * 2 + 1] = xmax; + } + return ksize; +} + +inline uint8_t pil_clip8(int32_t v) { + v >>= PIL_PRECISION_BITS; // arithmetic: floor, as Pillow's lookup table indexes + return (uint8_t)(v < 0 ? 0 : v > 255 ? 255 : v); +} + +/// Planar uint8 [3][h][w] -> [3][oh][ow], identical to PIL's Image.resize(BICUBIC). +std::vector pil_resize_bicubic(const uint8_t* src, int w, int h, int ow, int oh) { + std::vector bh, bv; + std::vector kh, kv; + const int ksh = pil_coeffs(w, ow, bh, kh); + const int ksv = pil_coeffs(h, oh, bv, kv); + const bool need_h = ow != w, need_v = oh != h; + + // Only the input rows the vertical pass reads, as ImagingResample does. + const int y_first = bv[0]; + const int y_last = bv[(size_t)oh * 2 - 2] + bv[(size_t)oh * 2 - 1]; + + std::vector out((size_t)3 * oh * ow); + for (int c = 0; c < 3; c++) { + const uint8_t* in = src + (size_t)c * h * w; + // horizontal pass, into a clamped uint8 temporary + const int th = need_h ? (y_last - y_first) : h; + const int y0 = need_h ? y_first : 0; + std::vector tmp; + const uint8_t* mid = in; + int mid_w = w; + if (need_h) { + tmp.resize((size_t)th * ow); +#pragma omp parallel for schedule(static) + for (int y = 0; y < th; y++) { + const uint8_t* row = in + (size_t)(y + y0) * w; + for (int xx = 0; xx < ow; xx++) { + const int xmin = bh[(size_t)xx * 2], xmax = bh[(size_t)xx * 2 + 1]; + const int32_t* k = &kh[(size_t)xx * ksh]; + int32_t ss = 1 << (PIL_PRECISION_BITS - 1); + for (int x = 0; x < xmax; x++) ss += (int32_t)row[x + xmin] * k[x]; + tmp[(size_t)y * ow + xx] = pil_clip8(ss); + } + } + mid = tmp.data(); + mid_w = ow; + } + // vertical pass; bounds shifted by the rows the horizontal pass skipped + uint8_t* o = out.data() + (size_t)c * oh * ow; + if (need_v) { + const int shift = need_h ? y_first : 0; +#pragma omp parallel for schedule(static) + for (int yy = 0; yy < oh; yy++) { + const int ymin = bv[(size_t)yy * 2] - shift, ymax = bv[(size_t)yy * 2 + 1]; + const int32_t* k = &kv[(size_t)yy * ksv]; + for (int xx = 0; xx < mid_w; xx++) { + int32_t ss = 1 << (PIL_PRECISION_BITS - 1); + for (int y = 0; y < ymax; y++) ss += (int32_t)mid[(size_t)(y + ymin) * mid_w + xx] * k[y]; + o[(size_t)yy * ow + xx] = pil_clip8(ss); + } + } + } else { + std::copy(mid, mid + (size_t)oh * ow, o); + } + } + return out; +} + +/// Append one view's patches -- the region (y0, x0, vh, vw) of a normalised +/// planar [3][H][W] image -- as raster rows of (c, kh, kw). +void patchify(const float* img, int H, int W, int y0, int x0, int vh, int vw, int p, + std::vector& out) { + const int gh = vh / p, gw = vw / p; + const size_t row = (size_t)3 * p * p; + const size_t base = out.size(); + out.resize(base + (size_t)gh * gw * row); +#pragma omp parallel for schedule(static) + for (int t = 0; t < gh * gw; t++) { + const int pr = t / gw, pc = t % gw; + bf16* dst = out.data() + base + (size_t)t * row; + for (int c = 0; c < 3; c++) + for (int kh = 0; kh < p; kh++) { + const float* src = img + ((size_t)c * H + (y0 + pr * p + kh)) * W + x0 + pc * p; + for (int kw = 0; kw < p; kw++) *dst++ = (bf16)src[kw]; + } + } +} + +} // namespace + +bool MiniCPM_OCR::load_image(const std::string& source, bool base64, image_data_t& chw) { + image_data_t decoded; + const bool ok = base64 ? image_reader_.load_image_base64(source, decoded) + : image_reader_.load_image(source, decoded); + if (!ok) return false; + const bool reordered = image_reader_.reorder_hwc_to_chw(decoded, chw); + image_reader_.recycle(decoded); + return reordered && chw.width > 0 && chw.height > 0; +} + +MiniCPM_OCR::image_layout MiniCPM_OCR::preprocess_image(const image_data_t& chw, + minicpm_ocr_image_payload_t& payload) { + auto* eng = reinterpret_cast(this->lm_engine); + const int patch = (int)eng->MINICPM_OCR_PATCH_SIZE; // 14 + const int max_slices = (int)eng->MINICPM_OCR_MAX_SLICE_NUMS; // 9 + const int divisor = (int)eng->MINICPM_OCR_DOWNSAMPLE_FACTOR; // 4 (16 on 4.7) + // One 2x2 merge per factor-of-4: 2 for "4x", 4 for "16x". See find_best_resize(). + const int merge = (divisor == 16) ? 4 : 2; + constexpr int scale = 448; // scale_resolution: the processor's default; not in any config + const float rescale = eng->MINICPM_OCR_VISION_RESCALE_FACTOR; + const float mean = eng->MINICPM_OCR_VISION_RESCALE_IMAGE_MEAN; + const float stdv = eng->MINICPM_OCR_VISION_RESCALE_IMAGE_STD; + + const int H = chw.height, W = chw.width; + const uint8_t* src = chw.pixels.data(); + + image_layout l; + const bool sliced = sliced_grid(H, W, max_slices, scale, l.rows, l.cols); + if (!sliced) l.rows = l.cols = 0; + const size_t before = payload._data__processed.size(); + + // resize (Pillow's bicubic, exactly), rescale + normalise, return planar fp32 + std::vector norm; + auto resized = [&](int h, int w) { + std::vector r = pil_resize_bicubic(src, W, H, w, h); + norm.resize((size_t)3 * h * w); + imgproc::avx512::rescale_and_normalize_avx512(r.data(), norm.data(), w, h, 3, true, rescale, + true, mean, stdv); + }; + auto push_view = [&](int y0, int x0, int vh, int vw, int img_h, int img_w) { + patchify(norm.data(), img_h, img_w, y0, x0, vh, vw, patch, payload._data__processed); + minicpm_ocr_image_t v{}; + v.height = H; + v.width = W; + v.height_resized = vh; + v.width_resized = vw; + v.grid_h = vh / patch; + v.grid_w = vw / patch; + payload.images.push_back(v); + return v.grid_h * v.grid_w / divisor; + }; + + // the source view: upscaled only when the image is not sliced + int sh, sw; + find_best_resize(H, W, scale, patch, merge, /*allow_upscale=*/!sliced, sh, sw); + resized(sh, sw); + l.src_tokens = push_view(0, 0, sh, sw, sh, sw); + l.views = 1; + l.tokens = l.src_tokens; + + if (sliced) { + // get_refine_size(): the whole image resized so it tiles exactly + const int rw0 = ensure_divide(W, l.cols), rh0 = ensure_divide(H, l.rows); + int bh, bw; + find_best_resize((double)rh0 / l.rows, (double)rw0 / l.cols, scale, patch, merge, true, bh, bw); + const int rh = bh * l.rows, rw = bw * l.cols; + resized(rh, rw); + // divide_to_patches(): row-major tiles + for (int r = 0; r < l.rows; r++) + for (int c = 0; c < l.cols; c++) { + l.slice_tokens = push_view(r * bh, c * bw, bh, bw, rh, rw); + l.tokens += l.slice_tokens; + l.views++; + } + } + payload.num_images = (unsigned)payload.images.size(); + l.values = payload._data__processed.size() - before; + return l; +} + +std::vector MiniCPM_OCR::image_placeholder(const image_layout& l, int index) { + constexpr int image_pad = 248056; + constexpr int image_start = 248078, image_end = 248079; // + constexpr int slice_start = 248088, slice_end = 248089; // + constexpr int image_id_start = 248090, image_id_end = 248091; // + constexpr int newline = 198; + + std::vector t; + t.push_back(image_id_start); + for (int d : this->tokenizer->encode(std::to_string(index))) t.push_back(d); + t.push_back(image_id_end); + t.push_back(image_start); + t.insert(t.end(), (size_t)l.src_tokens, image_pad); + t.push_back(image_end); + for (int r = 0; r < l.rows; r++) { + if (r > 0) t.push_back(newline); // "\n".join over slice rows + for (int c = 0; c < l.cols; c++) { + t.push_back(slice_start); + t.insert(t.end(), (size_t)l.slice_tokens, image_pad); + t.push_back(slice_end); + } + } + return t; +} diff --git a/src/include/AutoModel/all_models.hpp b/src/include/AutoModel/all_models.hpp index ccdce734d..9977d5f99 100644 --- a/src/include/AutoModel/all_models.hpp +++ b/src/include/AutoModel/all_models.hpp @@ -20,6 +20,7 @@ #include "modeling_qwen3vl.hpp" #include "modeling_qwen3_5vl.hpp" #include "modeling_minicpm_v.hpp" +#include "modeling_minicpm_ocr.hpp" #include "modeling_minicpm_v_moe.hpp" #include "modeling_qwen3_5_omni.hpp" #include "modeling_qwen3_6_moe.hpp" @@ -44,6 +45,7 @@ typedef enum { qwen3vl_flash, qwen3_5, minicpm_v, + minicpm_ocr, minicpm_v_moe, qwen3_5_omni, qwen3_6_moe, @@ -80,6 +82,7 @@ inline std::pair> get_auto_model(const s {"qwen3vl-flash", SupportedModelFamily::qwen3vl_flash}, {"qwen3.5", SupportedModelFamily::qwen3_5}, {"minicpm-v-4.7", SupportedModelFamily::minicpm_v}, + {"minicpm-ocr", SupportedModelFamily::minicpm_ocr}, {"minicpm-v-4.7-moe", SupportedModelFamily::minicpm_v_moe}, {"qwen3.5-omni", SupportedModelFamily::qwen3_5_omni}, {"qwen3.6-moe", SupportedModelFamily::qwen3_6_moe}, @@ -172,6 +175,9 @@ inline std::pair> get_auto_model(const s case SupportedModelFamily::minicpm_v: auto_chat_engine = std::make_unique(npu_device_inst); break; + case SupportedModelFamily::minicpm_ocr: + auto_chat_engine = std::make_unique(npu_device_inst); + break; case SupportedModelFamily::minicpm_v_moe: auto_chat_engine = std::make_unique(npu_device_inst); break; diff --git a/src/include/AutoModel/automodel.hpp b/src/include/AutoModel/automodel.hpp index b814bf1bb..297782267 100644 --- a/src/include/AutoModel/automodel.hpp +++ b/src/include/AutoModel/automodel.hpp @@ -34,6 +34,7 @@ #include "models/qwen3vl_flash/flm/aie2p/qwen3vl_flash.hpp" #include "models/qwen3_5vl/flm/aie2p/qwen3_5vl_npu.hpp" #include "models/minicpm_v_4_7/flm/aie2p/minicpm_v_4_7_npu.hpp" +#include "models/minicpm_ocr/flm/aie2p/minicpm_ocr_npu.hpp" #include "models/minicpm_v_4_7_moe/flm/aie2p/minicpm_v_4_7_moe_npu.hpp" #include "models/qwen3_6_moe/flm/aie2p/qwen3_6_moe_npu.hpp" #include "models/qwen3_8mtp/flm/aie2p/qwen3_8mtp_npu.hpp" diff --git a/src/include/AutoModel/modeling_minicpm_ocr.hpp b/src/include/AutoModel/modeling_minicpm_ocr.hpp new file mode 100644 index 000000000..1b4d83ab0 --- /dev/null +++ b/src/include/AutoModel/modeling_minicpm_ocr.hpp @@ -0,0 +1,112 @@ +/// \file modeling_minicpm_ocr.hpp +/// \brief MiniCPM_OCR class +/// \author FastFlowLM Team +/// \date 2026-10-02 +/// \version 0.9.28 +/// \note This is a header file for the MiniCPM_OCR class +/// \note MiniCPM-V-4.7-1B's text tower is config-identical to Qwen3.5-0.8B and +/// shares its tokenizer ids, so the chat flow follows Qwen3_5VL. +/// \note Images (S7b): each picture is sliced the way MiniCPMV4ImageProcessorPil +/// slices it -- a source view plus up to max_slice_nums slices -- and each +/// VIEW becomes one minicpm_ocr_image_t in the payload. The prompt's single +/// <|image_pad|> per picture is expanded to the processor's placeholder +/// (Npad..pad....). See +/// modeling_minicpm_ocr_image.cpp. Audio is dropped with a warning; video is +/// not supported. + +#pragma once +#include "AutoModel/automodel.hpp" +#include "image/image_reader.hpp" +#include "image_process_utils/imageproc.hpp" +#include "image_process_utils/imageprocAVX512.hpp" + +#include + +/************ MiniCPM_OCR **************/ +class MiniCPM_OCR : public AutoModel { +private: + + bool enable_think = false; + bool enable_tool = false; + static constexpr int think_start_id = 248068; + static constexpr int think_end_id = 248069; + + void setup_tokenizer(std::string model_path); + + // ---- images -------------------------------------------------------------- + /// \brief How one picture was sliced: what the prompt placeholder needs. + struct image_layout { + int views = 0; ///< 1 + rows*cols + int rows = 0, cols = 0;///< slice grid, 0 x 0 when unsliced + int src_tokens = 0; ///< LLM rows of the source view + int slice_tokens = 0; ///< LLM rows of each slice + int tokens = 0; ///< all image_token_id rows of the picture + size_t values = 0; ///< bf16 values it adds to _data__processed + }; + ImageReader image_reader_; + /// \brief Decode to uint8 CHW; false on failure. + bool load_image(const std::string& source, bool base64, image_data_t& chw); + /// \brief Slice, resize, normalise and patchify one picture into `payload`. + image_layout preprocess_image(const image_data_t& chw, minicpm_ocr_image_payload_t& payload); + /// \brief The token sequence `<|image_pad|>` expands to for picture `index`. + std::vector image_placeholder(const image_layout& l, int index); + + /// \brief Clear the conversation and throw after a failed inference + [[noreturn]] void fail_inference(); + + std::string decode(chat_meta_info_t& meta_info, int length_limit, std::ostream& os, std::function is_cancelled); + +public: + MiniCPM_OCR(flm_rt::device* npu_device_inst); + + void load_model(std::string model_path, json model_inf, int default_context_length = -1, bool enable_preemption = false, const std::string& backend = "") override; + bool insert(chat_meta_info_t& meta_info, lm_uniform_input_t& input, std::function is_cancelled = [] { return false; }) override; + std::string generate(chat_meta_info_t& meta_info, int length_limit, std::ostream& os, std::function is_cancelled = [] { return false; }) override; + std::string generate_with_prompt(chat_meta_info_t& meta_info, lm_uniform_input_t& input, int length_limit, std::ostream& os = std::cout) override; + std::string apply_chat_template(nlohmann::ordered_json& messages, nlohmann::ordered_json tools = nlohmann::ordered_json::object(), bool continue_final_message = false) override; + + bool thinking_enabled() const override { + return this->enable_think; + } + bool supports_continue_final_message() const override { + return true; + } + bool tools_enabled() const override { + return this->enable_tool; + } + + /// \brief Configure a parameter with type-erased value + /// \param parameter_name the name of the parameter + /// \param value the value to set (can be any type) + /// \return true if the parameter was configured successfully, false otherwise + bool configure_parameter(std::string parameter_name, const std::any& value) override{ + if (parameter_name == "enable_think") { + try { + this->enable_think = std::any_cast(value); + return true; + } catch (const std::bad_any_cast&) { + return false; + } + } + else if (parameter_name == "reasoning_effort") { + std::string reasoning_effort; + try { + reasoning_effort = std::any_cast(value); + if (reasoning_effort == "high" || reasoning_effort == "medium" || reasoning_effort == "low") + this->enable_think = true; + else if (reasoning_effort == "none") + this->enable_think = false; + else + header_print("WARNING", "Reasoning effort must be 'none', 'low', 'medium' or 'high'!"); + return true; + } catch (const std::bad_any_cast&) { + return false; + } + } + else if (parameter_name == "toggle_think") { + this->enable_think = !this->enable_think; + return true; + } + return AutoModel::configure_parameter(parameter_name, value); + } +}; diff --git a/src/include/models/minicpm_ocr/flm/aie2p/minicpm_ocr_npu.hpp b/src/include/models/minicpm_ocr/flm/aie2p/minicpm_ocr_npu.hpp new file mode 100644 index 000000000..a36d70118 --- /dev/null +++ b/src/include/models/minicpm_ocr/flm/aie2p/minicpm_ocr_npu.hpp @@ -0,0 +1,183 @@ +/// \file minicpm_ocr_npu.hpp +/// \brief minicpm_ocr_npu class -- MiniCPM-V-4.7-1B engine (text tower + SigLIP vision tower) +/// \author FastFlowLM Team +/// \date 2026-10-02 +/// \version 0.9.28 +/// +/// \note CONTRACT FILE. This is the engine ABI: S3 implements it, S6 (the AutoModel +/// wrapper) and S7 (vision) compile against it. Do not edit without the brain's +/// agreement -- see /scratch/michyu/minicpm-1b/CONTRACTS.md. +/// +/// MiniCPM-V-4.7-1B's text tower is config-identical to Qwen3.5-0.8B (hidden 1024, +/// 8/2 heads, head_dim 256, FFN 3584, 24 layers, 16x128 GatedDeltaNet heads, conv 4, +/// full_attention_interval 4, rope_theta 1e7, mrope_section [11,11,10]), so the text +/// path reuses the Qwen3.5-0.8B xclbins unchanged. The vision tower is SigLIP and +/// shares nothing with Qwen3.5-VL's ViT. +#pragma once +#include "lm_config.hpp" +#include "npu_utils/npu_utils.hpp" +#include "tensor_utils/q4_npu_eXpress.hpp" +#include "modules/embedding.hpp" +#include "modules/lm_head.hpp" +#include "modules/gemm.hpp" +#include "modules/dequant.hpp" +#include "tensor_2d.hpp" +#include "utils/utils.hpp" +#include "causal_lm.hpp" +#include "flm_state_api.h" +#if USEAVX2 +#include // For AVX intrinsics +#endif + +/// \brief One decoded image, before and after preprocessing. +/// +/// MiniCPM slices a large image into at most `MINICPM_OCR_MAX_SLICE_NUMS` tiles plus one +/// global view; each tile becomes one of these. `grid_h`/`grid_w` are the patch-grid +/// dimensions fed to the vision tower as `target_sizes`, i.e. `height_resized / +/// MINICPM_OCR_PATCH_SIZE` and `width_resized / MINICPM_OCR_PATCH_SIZE`. +typedef struct { + int height; + int width; + int height_resized; // assigned by image preprocessing + int width_resized; + int grid_h; // patch rows == height_resized / patch_size + int grid_w; // patch cols == width_resized / patch_size + + bytes _data; +} minicpm_ocr_image_t; + +/// \brief The multimodal blob handed to prefill() as `void* payload`. +/// +/// `_data__processed` is the concatenated, normalised pixel buffer for every image in +/// `images`, laid out tile-major. The engine turns it into soft tokens and splices them +/// over the rows whose id equals `image_token_id` (248056) -- see ยง7 of +/// ADDING_A_NEW_MODEL.md: those rows must bypass any embedding scale. +typedef struct { + std::vector images; + std::vector _data__processed; + unsigned int num_images; +} minicpm_ocr_image_payload_t; + +class minicpm_ocr_npu : public causal_lm { +public: + /// \brief initialize the minicpm_ocr_npu + /// \param config the configuration + /// \param npu_instance the npu instance + minicpm_ocr_npu(LM_Config config, npu_xclbin_manager *npu_instance, int MAX_L = 4096); + ~minicpm_ocr_npu(); + + /// \brief forward the minicpm_ocr_npu -- decode one token against the cached KV + /// \param ids the token id + /// \return the logits + buffer forward(int ids) override; + + /// \brief prefill the whole prompt + /// \param ids the token ids + /// \param payload nullptr for text-only, else a minicpm_ocr_image_payload_t* + buffer prefill(std::vector& ids, void* payload = nullptr) override; + + /// \brief set the context length + /// \param L the context length + void set_context_length(int L) override; + + /// \brief load the weights + /// \param q4nx the q4nx + void load_weights(Q4NX& q4nx) override; + + /// \brief clear the conversation state + void clear_context() override; + + /// \brief get the k cache + /// \param layer_idx the layer index + /// \param idx the index + /// \return the k cache + buffer get_k_cache(int layer_idx, int idx) override; + + /// \brief get the v cache + /// \param layer_idx the layer index + /// \param idx the index + /// \return the v cache + buffer get_v_cache(int layer_idx, int idx) override; + + /// \brief grow the caches without dropping the conversation + /// \param MAX_L the max length + void update_max_length(uint32_t MAX_L) override; + + /// \brief get the current context length + /// \return the current context length + int get_current_context_length() override; + int checkpoint() override; + int restore() override; + + // ---------------------------------------------------------------------- + // Vision preprocessing parameters. + // + // These live in preprocessor_config.json and config.json's vision_config, NOT in + // the top-level config.json, so they are read through a + // preprocessor -> config -> literal fallback chain. A checkpoint packaged before + // preprocessor_config.json existed still preprocesses correctly. + // ---------------------------------------------------------------------- + unsigned int MINICPM_OCR_PATCH_SIZE; ///< vision_config.patch_size (14) + unsigned int MINICPM_OCR_IMAGE_SIZE; ///< vision_config.image_size (980) + unsigned int MINICPM_OCR_MAX_SLICE_NUMS; ///< preprocessor.max_slice_nums (9) + unsigned int MINICPM_OCR_INSERT_LAYER_ID; ///< config.insert_layer_id (6) + unsigned int MINICPM_OCR_SPATIAL_MERGE_SIZE; ///< 2x2 per merge stage (2) + unsigned int MINICPM_OCR_TEMPORAL_PATCH_SIZE; ///< always 1 -- MiniCPM has no temporal + ///< patching; video frames are independent + ///< images, unlike Qwen3.5-VL's 2. + unsigned int MINICPM_OCR_DOWNSAMPLE_FACTOR; ///< config.downsample_mode "16x" (16) + float MINICPM_OCR_VISION_RESCALE_FACTOR; ///< 1/255 + float MINICPM_OCR_VISION_RESCALE_IMAGE_MEAN; ///< preprocessor.image_mean[0] (0.5) + float MINICPM_OCR_VISION_RESCALE_IMAGE_STD; ///< preprocessor.image_std[0] (0.5) + + /// \brief Read the vision preprocessing parameters. + /// + /// \note Called by Impl's constructor. + /// + /// Source order follows the shipped-config convention (see + /// `Qwen3.5-0.8B-NPU2/config.json`): the packaged `config.json` is a FLATTENED + /// config whose `vision_config` block carries ENGINE-named keys + /// (`MINICPM_OCR_PATCH_SIZE`, ...), not the upstream HuggingFace names. Each field + /// falls back to the upstream HF name, then to the literal from the shipping + /// MiniCPM-V-4.7-1B checkpoint, so a checkpoint packaged either way works and + /// nothing throws. + inline void load_vision_preprocess_parameters(LM_Config& config) { + const nlohmann::json& vc = config.sub("vision_config"); + + MINICPM_OCR_PATCH_SIZE = + cfg_get(vc, "MINICPM_OCR_PATCH_SIZE", cfg_get(vc, "patch_size", 14u)); + MINICPM_OCR_IMAGE_SIZE = + cfg_get(vc, "MINICPM_OCR_IMAGE_SIZE", cfg_get(vc, "image_size", 980u)); + MINICPM_OCR_MAX_SLICE_NUMS = + cfg_get(vc, "MINICPM_OCR_MAX_SLICE_NUMS", + cfg_get(config._json_config, "max_slice_nums", 9u)); + MINICPM_OCR_INSERT_LAYER_ID = + cfg_get(vc, "MINICPM_OCR_INSERT_LAYER_ID", + cfg_get(config._json_config, "insert_layer_id", 6u)); + MINICPM_OCR_SPATIAL_MERGE_SIZE = cfg_get(vc, "MINICPM_OCR_SPATIAL_MERGE_SIZE", 2u); + MINICPM_OCR_TEMPORAL_PATCH_SIZE = cfg_get(vc, "MINICPM_OCR_TEMPORAL_PATCH_SIZE", 1u); + + // downsample_mode is a string ("16x" / "4x") upstream, so the packaged config + // carries the resolved integer instead. + MINICPM_OCR_DOWNSAMPLE_FACTOR = cfg_get(vc, "MINICPM_OCR_DOWNSAMPLE_FACTOR", 0u); + if (MINICPM_OCR_DOWNSAMPLE_FACTOR == 0u) { + const std::string mode = + cfg_get(config._json_config, "downsample_mode", std::string("16x")); + MINICPM_OCR_DOWNSAMPLE_FACTOR = (mode == "4x") ? 4u : 16u; + } + + MINICPM_OCR_VISION_RESCALE_FACTOR = + cfg_get(vc, "MINICPM_OCR_VISION_RESCALE_FACTOR", 1.0f / 255.0f); + MINICPM_OCR_VISION_RESCALE_IMAGE_MEAN = + cfg_get(vc, "MINICPM_OCR_VISION_RESCALE_IMAGE_MEAN", 0.5f); + MINICPM_OCR_VISION_RESCALE_IMAGE_STD = + cfg_get(vc, "MINICPM_OCR_VISION_RESCALE_IMAGE_STD", 0.5f); + } + +private: + struct Impl; + Impl* _impl; +}; + +/// \brief the saved-state table of this engine (flm_state_api.h); nullptr if `engine` is not a minicpm_ocr_npu +extern "C" const flm_state_api* minicpm_ocr_state_api(void* engine); diff --git a/src/inno/flm.iss b/src/inno/flm.iss index 6ca2d71bb..303bcc469 100644 --- a/src/inno/flm.iss +++ b/src/inno/flm.iss @@ -66,6 +66,7 @@ Source: "qwen3vl_npu.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "qwen3vl_flash.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "qwen3_5vl_npu.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "minicpm_v_4_7_npu.dll"; DestDir: "{app}"; Flags: ignoreversion +Source: "minicpm_ocr_npu.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "minicpm_v_4_7_moe_npu.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "qwen3_5_omni_npu.dll"; DestDir: "{app}"; Flags: ignoreversion Source: "qwen3_6_moe_npu.dll"; DestDir: "{app}"; Flags: ignoreversion diff --git a/src/lib/hrx/libminicpm_ocr_npu.so b/src/lib/hrx/libminicpm_ocr_npu.so new file mode 100755 index 000000000..9074e85eb Binary files /dev/null and b/src/lib/hrx/libminicpm_ocr_npu.so differ diff --git a/src/lib/hrx/minicpm_ocr_npu.dll b/src/lib/hrx/minicpm_ocr_npu.dll new file mode 100644 index 000000000..ec43f521c Binary files /dev/null and b/src/lib/hrx/minicpm_ocr_npu.dll differ diff --git a/src/lib/hrx/minicpm_ocr_npu.lib b/src/lib/hrx/minicpm_ocr_npu.lib new file mode 100644 index 000000000..dcc9f31f7 Binary files /dev/null and b/src/lib/hrx/minicpm_ocr_npu.lib differ diff --git a/src/lib/xrt/libminicpm_ocr_npu.so b/src/lib/xrt/libminicpm_ocr_npu.so new file mode 100755 index 000000000..f65855785 Binary files /dev/null and b/src/lib/xrt/libminicpm_ocr_npu.so differ diff --git a/src/lib/xrt/minicpm_ocr_npu.dll b/src/lib/xrt/minicpm_ocr_npu.dll new file mode 100644 index 000000000..e9fc41b64 Binary files /dev/null and b/src/lib/xrt/minicpm_ocr_npu.dll differ diff --git a/src/lib/xrt/minicpm_ocr_npu.lib b/src/lib/xrt/minicpm_ocr_npu.lib new file mode 100644 index 000000000..1c65cda33 Binary files /dev/null and b/src/lib/xrt/minicpm_ocr_npu.lib differ diff --git a/src/model_info.json b/src/model_info.json index a698aa7aa..418dc6572 100644 --- a/src/model_info.json +++ b/src/model_info.json @@ -1206,6 +1206,62 @@ "path": "vision_weight.q4nx" } ], + "minicpm-ocr:1.0": [ + { + "type": "file", + "oid": "1007f10aa61d9b53080cbd51a17fd62ac3cd84ad", + "size": 7344, + "path": "chat_template.jinja" + }, + { + "type": "file", + "oid": "5a2f54b09cd34cbb781bc586cb022fdcbbf54a96", + "size": 2961, + "path": "config.json" + }, + { + "type": "file", + "oid": "e2dd11ef440ac6f4dddd2393ade881329297be00", + "size": 1068886688, + "lfs": { + "oid": "90993e9aa44fccc827962290a675647e46d3dd0eec3095ea83a0eb22d8dfa26d", + "size": 1068886688, + "pointerSize": 135 + }, + "xetHash": "eadb818a46375c81ff43946fa4b1efe91f75dac55ea4bf91c5d6897a3ebaa96a", + "path": "model.q4nx" + }, + { + "type": "file", + "oid": "b79f27c7a333e226a639156752449665ee8ab44d", + "size": 19993565, + "lfs": { + "oid": "6269eb38f25c20c0f2ba676577f2e5603f79eab3be573dfdaa727ca6eeb5522d", + "size": 19993565, + "pointerSize": 133 + }, + "xetHash": "19f1dcc31c14c7081b9f7589cd57ad2eed4b1187cb532439755df5ef9ac9f24d", + "path": "tokenizer.json" + }, + { + "type": "file", + "oid": "703aae0e239e63ba0801c9e0711684ecbe3934ce", + "size": 1538, + "path": "tokenizer_config.json" + }, + { + "type": "file", + "oid": "e572272e0fd5d2e1000c305d1f8d44373b8c46d1", + "size": 925394784, + "lfs": { + "oid": "94bb3108738ff4f4ac4564cc5c7423f68a36bcc72ae517e7a937b1c9b8bca90f", + "size": 925394784, + "pointerSize": 134 + }, + "xetHash": "0aa11f316f8e3d5a04706a2449b688b83a6b802bd6eb4b16d19742e08721c2c7", + "path": "vision_weight.q4nx" + } + ], "minicpm-v-4.7-moe:35b-a3b": [ { "type": "file", diff --git a/src/model_list.json b/src/model_list.json index f645763f8..d7064e860 100644 --- a/src/model_list.json +++ b/src/model_list.json @@ -566,6 +566,42 @@ } }, + "minicpm-ocr": { + "1.0": { + "supported_platforms": ["aie2p"], + "name": "MiniCPM-OCR-1.0-NPU2", + "url": "https://huggingface.co/FastFlowLM/MiniCPM-OCR-1.0-NPU2/resolve/main", + "file_url": "https://huggingface.co/api/models/FastFlowLM/MiniCPM-OCR-1.0-NPU2/tree/main", + "ms_url": "https://modelscope.cn/models/amd/MiniCPM-OCR-1.0-NPU2", + "size": 1300000000, + "flm_min_version": "1.0.12", + "files": [ + "config.json", + "model.q4nx", + "tokenizer.json", + "tokenizer_config.json", + "vision_weight.q4nx", + "chat_template.jinja" + ], + "vlm": true, + "default_context_length": 32768, + "max_prefill_len": 4096, + "details": { + "format": "NPU2", + "family": "minicpm-ocr", + "think": true, + "parameter_size": "1B", + "quantization_level": "Q4_K" + }, + "label": [ + "vision", + "reasoning", + "tool-calling" + ], + "footprint": 2.4 + } + }, + "minicpm-v-4.7-moe": { "35b-a3b": { "supported_platforms": ["aie2p"], diff --git a/src/test/minicpm_ocr_npu/CMakeLists.txt b/src/test/minicpm_ocr_npu/CMakeLists.txt new file mode 100644 index 000000000..6664d646d --- /dev/null +++ b/src/test/minicpm_ocr_npu/CMakeLists.txt @@ -0,0 +1,27 @@ +cmake_minimum_required(VERSION 3.22) +project(minicpm_ocr_npu VERSION 1.0.0 LANGUAGES CXX) + +include(${CMAKE_CURRENT_LIST_DIR}/../CMakeLists.txt) +npu_test_setup() + +add_npu_test( + test_minicpm_ocr_npu + test/minicpm_ocr_npu + USE_AUTOMODEL + USE_TOKENIZER + USE_SAMPLER + SOURCES + "${CMAKE_SOURCE_DIR}/../../common/AutoModel/model_backend.cpp" + "${CMAKE_SOURCE_DIR}/../../common/AutoModel/modeling_minicpm_ocr.cpp" +) + +target_link_libraries(test_minicpm_ocr_npu PUBLIC + minicpm_ocr_npu + xrt_coreutil +) + +# Add test target +add_custom_target(test_minicpm_ocr_npu_target + DEPENDS test_minicpm_ocr_npu + COMMENT "Building test_minicpm_ocr_npu executable" +) diff --git a/src/test/minicpm_ocr_npu/Makefile b/src/test/minicpm_ocr_npu/Makefile new file mode 100644 index 000000000..f9fe963eb --- /dev/null +++ b/src/test/minicpm_ocr_npu/Makefile @@ -0,0 +1,106 @@ +# ============================================================================= +# MiniCPM-V NPU test Makefile +# ============================================================================= +# +# Usage: +# make - Build the test +# make test - Build and run minicpm-ocr:1b (prompt, prompt-cache follow-up, +# fresh prompt with thinking) +# make clean - Remove all built files +# +# ============================================================================= +-include ../common.mk + +LENGTH ?= 256 + +SOURCES += test.cpp +SOURCES += ../../common/AutoModel/automodel.cpp +SOURCES += $(CHAT_SOURCES) +SOURCES += ../../common/AutoModel/model_backend.cpp +SOURCES += ../../common/AutoModel/modeling_minicpm_ocr.cpp +# Image support. modeling_minicpm_ocr_image.cpp holds the slicing, the PIL-exact +# resize and image_placeholder(); image_reader.cpp decodes the file (libjpeg-turbo/lodepng); +# imageprocAVX512.cpp supplies rescale_and_normalize_avx512(), and pulls in +# imageproc.cpp for resize_bicubic_* at link time. +# Same file list as the other VLM harnesses. MiniCPM does NOT call the shared +# resize -- it ships a PIL-exact one in modeling_minicpm_ocr_image.cpp, because the +# shared bicubic differs from PIL by up to 29 levels and SigLIP amplifies that +# more than every device rounding combined +# (notes/tricks/2026-10-05-siglip-amplifies-resize-differences.md in FLM_IRON). +# imageproc.cpp is still linked: imageprocAVX512.cpp references it regardless. +SOURCES += ../../common/AutoModel/modeling_minicpm_ocr_image.cpp +SOURCES += ../../common/image/image_reader.cpp +SOURCES += ../../include/lodepng/lodepng.cpp +SOURCES += ../../common/image_process_utils/imageproc.cpp +SOURCES += ../../common/image_process_utils/imageprocAVX512.cpp +SOURCES += ../../common/tokenizer/tokenizer.cpp +SOURCES += ../../common/modules/sampler.cpp +# automodel.cpp reaches the prompt cache through PrefixSlotCache; without this the link +# fails on PrefixSlotCache::resume/save/attach and prefix_hash_bytes. No other harness +# lists it either, so they share this gap. +SOURCES += ../../common/prefix_slot_cache.cpp + +HEADERS += ../../include/models/minicpm_ocr/flm/aie2p/minicpm_ocr_npu.hpp +HEADERS += ../../include/AutoModel/modeling_minicpm_ocr.hpp + + +ifeq ($(WSL), 0) + +LDFLAGS += -lminicpm_ocr_npu +# image_reader.cpp decodes through libjpeg-turbo (+ vendored lodepng), same as every other VLM harness +LDFLAGS += -ljpeg +# image_reader.cpp also calls WebPGetInfo/WebPDecodeRGBInto unconditionally. +LDFLAGS += -lwebp + +CPP_SOURCES := $(filter %.cpp,$(SOURCES)) +OBJECTS := $(addprefix $(BUILD_DIR)/,$(notdir $(CPP_SOURCES:.cpp=.o))) +DEPS := $(OBJECTS:.o=.d) +VPATH := $(sort $(dir $(CPP_SOURCES))) + +all: $(BUILD_DIR)/test_minicpm_ocr_npu + +$(BUILD_DIR): + mkdir -p $(BUILD_DIR) + +$(BUILD_DIR)/%.o: %.cpp | $(BUILD_DIR) + $(CXX) $(CXX_FLAGS) -c $< -o $@ + +$(BUILD_DIR)/test_minicpm_ocr_npu: $(OBJECTS) + $(CXX) $(CXX_FLAGS) -o $@ $^ $(LDFLAGS) $(DEPENDENCY_LDFLAGS) + +clean: + rm -rf $(BUILD_DIR) + +test: $(BUILD_DIR)/test_minicpm_ocr_npu + cp ../../model_list.json $(BUILD_DIR)/model_list.json + cd $(BUILD_DIR) && ./test_minicpm_ocr_npu --model minicpm-ocr:1b --Length $(LENGTH) + +-include $(DEPS) +.PHONY: all clean test + +else + +# WSL build environment +# Use CMake to invoke the Visual Studio +PWSH := powershell.exe + +all: directories test + +directories: + mkdir -p $(BUILD_DIR) + +$(BUILD_DIR)/test_minicpm_ocr_npu.exe: $(SOURCES) + cd $(BUILD_DIR) && $(PWSH) -Command "cmake ../../../test/minicpm_ocr_npu" + cd $(BUILD_DIR) && $(PWSH) -Command "cmake --build . --config Release --target test_minicpm_ocr_npu" + +clean: + rm -rf $(BUILD_DIR) + + +test: directories $(BUILD_DIR)/test_minicpm_ocr_npu.exe + cp ../../model_list.json $(BUILD_DIR)/model_list.json + cd $(BUILD_DIR) && ${PWSH} -Command "\$$env:PATH = '..\..\..\lib;' + \$$env:PATH; .\test_minicpm_ocr_npu.exe -m minicpm-ocr:1b -l $(LENGTH)" + +.PHONY: all clean test directories + +endif diff --git a/src/test/minicpm_ocr_npu/activate.sh b/src/test/minicpm_ocr_npu/activate.sh new file mode 100755 index 000000000..30c55360e --- /dev/null +++ b/src/test/minicpm_ocr_npu/activate.sh @@ -0,0 +1,119 @@ +#!/usr/bin/env bash +# Set up the environment for the MiniCPM-V-4.7-1B test harness. +# +# source activate.sh # stage whatever engine is already built +# source activate.sh host # rebuild the all-host engine first (no NPU) +# source activate.sh npu # rebuild the all-NPU engine first (S4/S5) +# +# then: +# make test # or: make test LENGTH=64 +# # (the thinking turn generates 4x LENGTH) +# +# Why this is needed: the harness links -lminicpm_ocr_npu out of src/lib/xrt/, but +# that engine is built in the *other* repo (FastFlowLM_IRON) and is deliberately +# not committed here -- it is a build artifact whose host/NPU flavour changes. +# Without staging it you get: +# /usr/bin/ld: cannot find -lminicpm_ocr_npu: No such file or directory +# +# It also exports FLM_MODEL_PATH so the harness finds the model in /scratch +# rather than ~/.config/flm. model_list.json's "model_path" is "models", and +# model_root_path is FLM_MODEL_PATH/models, so point this at the PARENT of the +# model folders. + +# --- settings, override by exporting before sourcing ------------------------- +IRON_REPO="${IRON_REPO:-/scratch/michyu/Projects/FastFlowLM_IRON}" +FLM_MODEL_PATH="${FLM_MODEL_PATH:-/scratch/michyu}" +MODEL_NAME="${MODEL_NAME:-MiniCPM-OCR-1.0-NPU2}" + +# Host build: every stage on the CPU. Correct but slow (~20 tok/s decode). +# The polarity is opt-OUT -- a plain `make` is all-NPU. +HOST_FLAGS="PREFILL_MM=cpu PREFILL_ATTN=cpu PREFILL_GDN=cpu PREFILL_CONV=cpu DECODE=cpu" + +# Tolerate `bash activate.sh` as well as `source activate.sh`; only the latter +# makes the exports stick, so say so rather than silently doing half the job. +_sourced=1 +[ "${BASH_SOURCE[0]}" = "${0}" ] && _sourced=0 + +_here="$(cd "$(dirname "${BASH_SOURCE[0]}")" && pwd)" +_libdir="$(cd "$_here/../../lib/xrt" && pwd)" +_engdir="$IRON_REPO/FLM_DLL/detail/minicpm_ocr_npu" +_built="$IRON_REPO/FLM_DLL/build/lib/libminicpm_ocr_npu.so" + +_fail() { echo "[activate] ERROR: $*" >&2; return 1; } + +_main() { + [ -d "$_engdir" ] || { _fail "engine dir not found: $_engdir (set IRON_REPO)"; return 1; } + + case "${1:-}" in + host) echo "[activate] building all-host engine ..." + make -C "$_engdir" -j"$(nproc)" $HOST_FLAGS >/tmp/minicpm_ocr_activate_build.log 2>&1 \ + || { tail -25 /tmp/minicpm_ocr_activate_build.log; _fail "build failed"; return 1; } ;; + npu) echo "[activate] building all-NPU engine ..." + make -C "$_engdir" -j"$(nproc)" >/tmp/minicpm_ocr_activate_build.log 2>&1 \ + || { tail -25 /tmp/minicpm_ocr_activate_build.log; _fail "build failed"; return 1; } ;; + "") ;; + *) _fail "unknown argument '$1' (expected: host, npu, or nothing)"; return 1 ;; + esac + + [ -f "$_built" ] || { _fail "no engine built yet: $_built + run: source activate.sh host (all-CPU, works without an NPU) + or: source activate.sh npu (needs S4/S5 and the NPU box)"; return 1; } + + # Stale-.so check. build/ is gitignored and therefore per-worktree, so + # several copies of this file exist on disk. AGENTS.md section 8: a stale + # .so makes experiments lie. Compare against the newest engine source. + local newest + newest="$(find "$_engdir" -name '*.cpp' -o -name '*.hpp' | xargs ls -t 2>/dev/null | head -1)" + if [ -n "$newest" ] && [ "$newest" -nt "$_built" ]; then + echo "[activate] WARNING: engine is OLDER than $(basename "$newest")." + echo "[activate] rebuild with: source activate.sh host|npu" + fi + + install -m 755 "$_built" "$_libdir/libminicpm_ocr_npu.so" || { _fail "could not stage into $_libdir"; return 1; } + echo "[activate] staged $(basename "$_built") -> src/lib/xrt/ ($(date -r "$_built" '+%H:%M'))" + + # Report which flavour was staged, so a surprising tok/s is explained before + # it is chased. A -D define leaves no trace in the binary, so test for the + # "layer.xclbin" literal instead: its register_xclbin() call sits inside + # #if MINICPM_OCR_DECODE_ON_NPU, so the host build does not contain it. + # (Do not test lm_head.xclbin -- that string appears in both, from the + # shared detail/lm_head objects linked into every engine.) + if strings "$_libdir/libminicpm_ocr_npu.so" | grep -q 'layer\.xclbin'; then + echo "[activate] flavour: NPU (expect ~44-47 tok/s decode in make test)" + # The harness is built with -DDEV_BUILD, which hardwires LM_Config's + # exec_path to "../../../" relative to its cwd -- i.e. src/ -- so it reads + # src/xclbins// and IGNORES FLM_XCLBIN_PATH. Check that directory. + local xdir="$_here/../../xclbins/$MODEL_NAME" + local x missing="" + for x in layer lm_head mm attn conv GateDeltaNet_prefill; do + [ -f "$xdir/$x.xclbin" ] || missing="$missing $x.xclbin" + done + if [ -n "$missing" ]; then + echo "[activate] WARNING: missing in src/xclbins/$MODEL_NAME/:$missing" + echo "[activate] the text tower reuses Qwen3.5-0.8B's bitstreams unchanged:" + echo "[activate] mkdir -p $xdir && cp $_here/../../xclbins/Qwen3.5-0.8B-NPU2/{layer,lm_head,mm,attn,conv,GateDeltaNet_prefill}.xclbin $xdir/" + fi + else + echo "[activate] flavour: ALL-HOST -- decode will be ~17-20 tok/s, NOT the NPU number." + echo "[activate] for NPU numbers: source activate.sh npu" + fi + + local mdir="$FLM_MODEL_PATH/models/$MODEL_NAME" + [ -d "$mdir" ] || echo "[activate] WARNING: model folder not found: $mdir" + [ -f "$mdir/model.q4nx" ] || echo "[activate] WARNING: no model.q4nx in $mdir" + + export FLM_MODEL_PATH + echo "[activate] FLM_MODEL_PATH=$FLM_MODEL_PATH (resolves $FLM_MODEL_PATH/models/$MODEL_NAME)" + + if [ "$_sourced" = "0" ]; then + echo "[activate] NOTE: run with 'source activate.sh' -- executed directly, FLM_MODEL_PATH will not persist." + else + echo "[activate] ready: make test # add LENGTH=64 for a quicker run" + fi +} + +_main "${1:-}" +_activate_rc=$? +unset -f _main _fail +unset _here _libdir _engdir _built _sourced +return $_activate_rc 2>/dev/null || exit $_activate_rc diff --git a/src/test/minicpm_ocr_npu/test.cpp b/src/test/minicpm_ocr_npu/test.cpp new file mode 100644 index 000000000..8384f143c --- /dev/null +++ b/src/test/minicpm_ocr_npu/test.cpp @@ -0,0 +1,512 @@ +/// \file test.cpp +/// \brief standalone harness for MiniCPM_OCR +/// \note Four turns: a CLI-style prompt, a REST-style follow-up that restores +/// the prompt-cache checkpoint, then clear_context() and a fresh prompt +/// with thinking on, then an image. Each turn reports the token that +/// stopped it. +/// \note --multi-slot instead runs three interleaved conversations through the +/// prefix slot cache and checks each one gets its own state back (run_multi_slot). +#include +#include +#include +#include +#include +#include +#include "base64.hpp" +#include "utils/utils.hpp" +#include "utils/vm_args.hpp" +#include "AutoModel/modeling_minicpm_ocr.hpp" +#include "AutoModel/flm_backend.hpp" +#include "model_list.hpp" + +flm_rt::device npu_device_global; + +/// \note builtin_backends.cpp names every engine and so links every prebuilt +/// library; this harness registers just the one it drives. +namespace flm::backend { +void register_builtin_backends(BackendRegistry& registry) { + registry.register_backend("minicpm-ocr", kFlmBackendId, flm_factory()); +} +} // namespace flm::backend + +/// \brief Name a picture the way a person would: "720p", "1080p", ... +/// \note The label keys off the SHORT side, which is what the p-numbers mean -- +/// 960x720 and 1280x720 are both 720p. It is a label, not a measurement: +/// anything that is not within 5% of a standard height is reported by its +/// pixel count instead of being rounded into the nearest bucket. +static std::string resolution_label(int w, int h) { + const int shortest = std::min(w, h); + static const int kStandard[] = {480, 576, 720, 1080, 1440, 2160, 4320}; + for (int s : kStandard) { + if (std::abs(shortest - s) * 20 <= s) return std::to_string(s) + "p"; + } + std::ostringstream mp; + mp.precision(1); + mp << std::fixed << (w * (double)h / 1e6) << " MP"; + return mp.str(); +} + +/// \brief Decode the image header and say how big it is. +/// \note MiniCPM's TTFT is driven by how many VIEWS the slicer makes, which is a +/// function of both the pixel count and the aspect ratio (get_sliced_grid +/// picks a grid, capped by max_slice_nums), so print both. The view count +/// itself is private to MiniCPM_OCR -- read it off the "vision:" line that +/// FLM_MINICPM_OCR_VISION_PROFILE=1 prints. +static void report_image(const std::string& path) { + ImageReader reader; + image_data_t decoded; + if (!reader.load_image(path, decoded)) { + header_print("WARNING", "could not decode " << path << " to report its size"); + return; + } + const int w = decoded.width, h = decoded.height; + reader.recycle(decoded); + std::ostringstream aspect; + aspect.precision(2); + aspect << std::fixed << (w / (double)h); + header_print("info", "image: " << w << "x" << h << " (" << resolution_label(w, h) + << ", aspect " << aspect.str() << ")"); +} + +static void report_turn(AutoModel& chat, const chat_meta_info_t& meta_info) { + std::vector history = chat.get_history().second; + std::cout << std::endl; + std::cout << "prompt_tokens: " << meta_info.prompt_tokens + << ", cached_prompt_tokens: " << meta_info.cached_prompt_tokens + << ", generated_tokens: " << meta_info.generated_tokens + << ", stop_reason: " << stop_reason_to_string(meta_info.stop_reason) << std::endl; + std::cout << "last token: " << (history.empty() ? -1 : history.back()) + << ", context length: " << chat.get_current_context_length() << std::endl; + std::cout << chat.show_profile() << std::endl; +} + +// --------------------------------------------------------------------------------------------- +// --multi-slot +// --------------------------------------------------------------------------------------------- + +/// \brief MiniCPM_OCR plus the prefix-slot controls the test needs (all protected in AutoModel) +class MiniCPMOCRProbe : public MiniCPM_OCR { +public: + using MiniCPM_OCR::MiniCPM_OCR; + /// \brief turn the prefix slots on (-1 = the default count) or off (0), as --prefix-slots does + void set_slots(int slots) { this->set_prefix_limits(-1, slots); } + bool slots_enabled() const { return this->prefix_cache_.enabled(); } + /// \brief back to a freshly loaded model: no live context, no saved snapshots + void reset_everything() { + this->prefix_cache_.forget_all(); + this->clear_context(); + } + /// \brief what the slot cache did for the last request (valid only if it went through the slots) + prefix_cache_report_t last_report() const { return this->prefix_cache_.report(); } + size_t count_tokens(const std::string& text) { return this->tokenizer->encode(text).size(); } +}; + +/// \brief one conversation of the test: a first user message, then one follow-up question +struct ms_conversation_t { + std::string name; + nlohmann::ordered_json first_user; //!< {"role": "user", "content": ..., maybe "images": [base64]} + std::string question2; +}; + +/// \brief what one turn did +struct ms_turn_t { + std::string reply; + int prompt_tokens = 0; + int cached_tokens = 0; //!< served from a saved state instead of prefilled + double prefill_ms = 0.0; + bool slot_hit = false; //!< resumed from a slot (only when the slots are on) + int slots_used = 0; //!< slots held after the turn (only when the slots are on) +}; + +/// \brief one request, sent the way `flm serve` sends every request +static ms_turn_t ms_run_turn(MiniCPMOCRProbe& chat, const std::string& title, + const nlohmann::ordered_json& messages, int length) { + std::cout << "\n---- " << title << " ----" << std::endl; + chat_meta_info_t m; + // the server sets this on every request: the model may reuse what it already holds. + // Without it neither the slots nor the single checkpoint are ever consulted. + m.restore_allowed = true; + lm_uniform_input_t in; + in.messages = messages; + chat.start_total_timer(); + chat.insert(m, in); + ms_turn_t t; + t.reply = chat.generate(m, length, std::cout); + chat.stop_total_timer(); + t.prompt_tokens = m.prompt_tokens; + t.cached_tokens = m.cached_prompt_tokens; + t.prefill_ms = m.prefill_duration / 1e6; + if (chat.slots_enabled()) { + const prefix_cache_report_t r = chat.last_report(); + t.slot_hit = r.valid && r.hit; + t.slots_used = r.slots_used; + } + std::cout << "\n prompt " << t.prompt_tokens << " tokens, " << t.cached_tokens + << " reused, " << t.prompt_tokens - t.cached_tokens << " prefilled in " + << t.prefill_ms << " ms"; + if (chat.slots_enabled()) { + std::cout << "; slot " << (t.slot_hit ? "HIT" : "miss") << ", " << t.slots_used << " slots held"; + } + std::cout << std::endl; + return t; +} + +/// \brief turn 2 of a conversation: turn 1, its reply, and the follow-up question +static nlohmann::ordered_json ms_turn2_messages(const ms_conversation_t& c, const std::string& reply1) { + nlohmann::ordered_json msgs = nlohmann::ordered_json::array(); + msgs.push_back(c.first_user); + msgs.push_back({ {"role", "assistant"}, {"content", reply1} }); + msgs.push_back({ {"role", "user"}, {"content", c.question2} }); + return msgs; +} + +/// \brief a made-up inventory list of at least `min_tokens` tokens. Long enough (>= 1536) that its +/// prefill runs the gated delta rule on the NPU, while the short follow-up runs it on the +/// host -- so a resume has to hand the recurrent state from one backend to the other. +static std::string ms_long_document(MiniCPMOCRProbe& chat, size_t min_tokens) { + static const char* kCities[] = {"Lisbon", "Oslo", "Lima", "Cairo", "Hanoi", "Quito", "Perth", "Riga"}; + static const char* kGoods[] = {"copper wire", "rice", "glass jars", "wool", "tea", "steel pipes", "paper", "salt"}; + std::string doc = "Warehouse inventory report.\n"; + int item = 1; + while (chat.count_tokens(doc) < min_tokens) { + // ten lines between token counts: the encode is over the whole document each time + for (int k = 0; k < 10; k++, item++) { + const int crates = 100 + (37 * item) % 900; + doc += "Item " + std::to_string(item) + ": the warehouse in " + kCities[item % 8] + " holds " + + std::to_string(crates) + " crates of " + kGoods[(item * 3) % 8] + ".\n"; + } + } + return doc + "\nRead the report. How many items does it list? Answer in one sentence."; +} + +/// \brief --multi-slot: do interleaved conversations each get their own state back from the slots? +/// \note Three conversations, each a first turn plus a short follow-up: +/// A a ~2k-token document (GDN on the NPU for turn 1, on the host for the follow-up) +/// B a short prompt (GDN on the host throughout) +/// C an image and a question about it (skipped with --image "") +/// Phase 1, reference: slots OFF, each conversation alone from a clean model, `rounds` +/// times. The single-checkpoint path rewinds turn 2 to the end of turn 1's prompt and +/// prefills the rest. Repeating it measures whether the engine repeats itself at all. +/// Phase 2, multi-slot: slots ON, all first turns, then all follow-ups (A1 B1 C1 A2 B2 C2). +/// Every follow-up must resume from ITS conversation's snapshot, at the same point the +/// reference resumed (same cached-token count, so the same tokens are prefilled), and +/// every reply must equal the reference's byte for byte. +/// Phase 3, timing only: slots OFF, same interleaved order -- what the slots save. Each +/// follow-up there has to prefill its whole conversation again. +/// Sampling is greedy with the penalties off, so a reply depends on the model state alone. +static int run_multi_slot(MiniCPMOCRProbe& chat, const std::vector& convs, + int length, int rounds) { + chat.set_topk(1); + chat.set_penalty_window(0); // penalties off: resuming resets their history, the legacy path does not + chat.configure_parameter("enable_think", false); + const size_t n = convs.size(); + + // ---- phase 1: reference, slots off, one conversation at a time ---- + std::vector> ref1(n), ref2(n); // [conversation][round] + chat.set_slots(0); + for (int r = 0; r < rounds; r++) { + for (size_t i = 0; i < n; i++) { + chat.reset_everything(); + const std::string tag = "reference round " + std::to_string(r + 1) + ", " + convs[i].name; + ref1[i].push_back(ms_run_turn(chat, tag + " turn 1", nlohmann::ordered_json::array({convs[i].first_user}), length)); + ref2[i].push_back(ms_run_turn(chat, tag + " turn 2", ms_turn2_messages(convs[i], ref1[i].back().reply), length)); + } + } + + // ---- phase 2: slots on, interleaved ---- + chat.set_slots(-1); + if (!chat.slots_enabled()) { + header_print("ERROR", "the prefix slots did not turn on (FLM_DISABLE_PREFIX_SLOTS set, or no state api)"); + return 1; + } + chat.reset_everything(); + std::vector ms1(n), ms2(n); + for (size_t i = 0; i < n; i++) { + ms1[i] = ms_run_turn(chat, "multi-slot " + convs[i].name + " turn 1", nlohmann::ordered_json::array({convs[i].first_user}), length); + } + for (size_t i = 0; i < n; i++) { + // the SAME conversation as the reference's turn 2, so the prompts are byte identical + ms2[i] = ms_run_turn(chat, "multi-slot " + convs[i].name + " turn 2", ms_turn2_messages(convs[i], ref1[i][0].reply), length); + } + + // ---- phase 3: slots off, interleaved, timing only ---- + chat.set_slots(0); + chat.reset_everything(); + std::vector no1(n), no2(n); + for (size_t i = 0; i < n; i++) { + no1[i] = ms_run_turn(chat, "no slots " + convs[i].name + " turn 1", nlohmann::ordered_json::array({convs[i].first_user}), length); + } + for (size_t i = 0; i < n; i++) { + no2[i] = ms_run_turn(chat, "no slots " + convs[i].name + " turn 2", ms_turn2_messages(convs[i], ref1[i][0].reply), length); + } + chat.set_slots(-1); // back to the default for anything that runs after + chat.reset_everything(); + + // ---- verdict ---- + std::cout << "\n" << std::string(100, '=') << "\nmulti-slot summary\n" << std::string(100, '=') << std::endl; + std::cout << "follow-up turn: tokens reused + prefill ms\n"; + std::printf("%-6s | %7s | %-22s | %-22s | %-22s | %s\n", "conv", "prompt", "reference (alone)", + "slots on, interleaved", "slots off, interleaved", "replies vs reference"); + bool all_pass = true, all_repeat = true; + for (size_t i = 0; i < n; i++) { + // does the engine repeat itself? every reference round must agree + bool repeats = true; + for (int r = 1; r < rounds; r++) { + repeats = repeats && ref1[i][r].reply == ref1[i][0].reply && ref2[i][r].reply == ref2[i][0].reply; + } + all_repeat = all_repeat && repeats; + const ms_turn_t& ref = ref2[i][0]; + const bool hit = ms2[i].slot_hit && ms2[i].cached_tokens > 0; + const bool same_split = ms2[i].cached_tokens == ref.cached_tokens; + const bool same1 = ms1[i].reply == ref1[i][0].reply; + const bool same2 = ms2[i].reply == ref.reply; + const bool pass = hit && same_split && same1 && same2; + all_pass = all_pass && pass; + + char a[64], b[64], c[64]; + std::snprintf(a, sizeof(a), "%5d + %7.1f ms", ref.cached_tokens, ref.prefill_ms); + std::snprintf(b, sizeof(b), "%5d + %7.1f ms", ms2[i].cached_tokens, ms2[i].prefill_ms); + std::snprintf(c, sizeof(c), "%5d + %7.1f ms", no2[i].cached_tokens, no2[i].prefill_ms); + std::string verdict = std::string("turn 1 ") + (same1 ? "same" : "DIFFERS") + ", turn 2 " + (same2 ? "same" : "DIFFERS"); + if (!hit) verdict += ", NO SLOT HIT"; + else if (!same_split) verdict += ", resumed at a different point"; + if (!repeats) verdict += " (engine does not repeat itself here)"; + std::printf("%-6s | %7d | %-22s | %-22s | %-22s | %s\n", convs[i].name.c_str(), ref.prompt_tokens, a, b, c, verdict.c_str()); + } + if (all_pass) { + std::cout << "\nPASS: every interleaved follow-up resumed from its own conversation's snapshot, at the same" + " point as the reference, and every reply matches it byte for byte" << std::endl; + return 0; + } + std::cout << (all_repeat + ? "\nFAIL: the engine repeats itself without slots, but the interleaved run with slots differs" + : "\nINCONCLUSIVE: the engine did not repeat itself across reference rounds, so an exact comparison" + " cannot separate the slots from its own spread") << std::endl; + return 1; +} + +/// \brief read a whole file as bytes, base64-encoded (a message's "images" entry) +static std::string ms_file_base64(const std::string& path) { + std::ifstream f(path, std::ios::binary); + std::stringstream buf; + buf << f.rdbuf(); + return base64::to_base64(buf.str()); +} + +int main(int argc, char* argv[]) { + arg_utils::po::options_description desc("Allowed options"); + arg_utils::po::variables_map vm; + desc.add_options()("model,m", arg_utils::po::value()->default_value("minicpm-ocr:1b"), "Model tag"); + desc.add_options()("Length,l", arg_utils::po::value()->default_value(256), "Max generated tokens per turn"); + desc.add_options()("Preemption,p", arg_utils::po::value()->default_value(false), "Preemption"); + // A single timed turn on a prompt from a file, for measuring prefill against length + // (models_simple/test_input/{1k,2k,8k}.txt). Skips the four scripted turns entirely. + desc.add_options()("prompt-file", arg_utils::po::value()->default_value(""), + "Run one turn on this file's contents and exit"); + // With --prompt-file: a second turn, sent the way `flm serve` sends a follow-up (whole + // conversation, restore_allowed). Turn one is sent CLI-style (no restore_allowed), so with + // the prefix slots on (the default) it saves no snapshot, and turn two prefills the whole + // conversation in two calls -- the head up to its last user turn, then the tail. Their + // lengths pick the GDN backend of each, which is what makes it the test for handing the + // recurrent state between the two inside one request. + desc.add_options()("follow-up", arg_utils::po::value()->default_value(""), + "With --prompt-file: file holding a follow-up user turn"); + // With --follow-up: send turn one with restore_allowed as well, so it saves a snapshot and + // turn two resumes from it, prefilling only the reply and the new user turn -- the short + // cached follow-up of a real chat. + desc.add_options()("resume", arg_utils::po::bool_switch()->default_value(false), + "With --follow-up: turn two resumes from turn one's saved state"); + // Default is the repo's own test image, found relative to the executable the + // same way model_list.json is. Pass --image "" to skip turn 4 (and --multi-slot's + // image conversation) outright. + desc.add_options()("image,i", arg_utils::po::value()->default_value(""), + "Image for turn 4 (and for --multi-slot's image conversation); \"\" skips it"); + desc.add_options()("multi-slot", arg_utils::po::bool_switch()->default_value(false), + "Run three interleaved conversations with the prefix slots off and on and compare (run_multi_slot)"); + desc.add_options()("rounds", arg_utils::po::value()->default_value(2), + "With --multi-slot: reference rounds (>= 2 tells whether the engine repeats itself)"); + arg_utils::po::store(arg_utils::po::parse_command_line(argc, argv, desc), vm); + + std::string tag = vm["model"].as(); + int length = vm["Length"].as(); + bool preemption = vm["Preemption"].as(); + std::cout << "Model: " << tag << std::endl; + std::string exe_dir = utils::get_executable_directory(); + std::string model_dir = utils::get_models_directory(); + std::string model_list_path = exe_dir + "/model_list.json"; + model_list model_list(model_list_path, model_dir); + + header_print("info", "Initializing chat model..."); + std::string model_path = model_list.get_model_path(tag); + nlohmann::json model_info = model_list.get_model_info(tag).second; + std::cout << "Model path: " << model_path << std::endl; + + std::unique_ptr chat = std::make_unique(&npu_device_global); + npu_device_global = flm_rt::device(0); + chat->load_model(model_path, model_info, -1, preemption); + header_print("info", "Model loaded"); + chat->set_topk(1); + chat->configure_parameter("enable_think", false); + + // Default is the repo's own test image, found relative to the executable the + // same way model_list.json is. + std::string image = vm["image"].as(); + if (image == "") + image = (std::filesystem::path(exe_dir) / ".." / ".." / ".." / ".." + / ".github" / "script" / "assets" / "test_image1.jpg").lexically_normal().string(); + + if (vm["multi-slot"].as()) { + std::vector convs; + // A: long enough that turn 1 runs the GDN on the NPU (>= 1536 tokens), or --prompt-file + std::string doc; + const std::string doc_file = vm["prompt-file"].as(); + if (!doc_file.empty()) { + std::ifstream df(doc_file); + if (!df) { header_print("ERROR", "cannot open " << doc_file); return 1; } + std::stringstream buf; buf << df.rdbuf(); + doc = buf.str(); + } else { + doc = ms_long_document(*chat, 2000); + } + convs.push_back({"A", { {"role", "user"}, {"content", doc} }, + "Which city holds the crates listed as item 7, and what is in them? Answer in one sentence."}); + // B: short, the GDN on the host throughout + convs.push_back({"B", { {"role", "user"}, {"content", "My name is Ada and I live in Lisbon. In one sentence, what is Lisbon known for?"} }, + "What is my name, and which city did I mention?"}); + // C: a picture; the slot must only be reused for the same picture + if (image.empty()) { + header_print("info", "--multi-slot: image conversation skipped (--image \"\")"); + } else if (!std::filesystem::exists(image)) { + header_print("WARNING", "--multi-slot: image conversation skipped, no such file: " << image); + } else { + nlohmann::ordered_json user = { {"role", "user"}, {"content", "Describe this image."} }; + user["images"] = nlohmann::ordered_json::array({ ms_file_base64(image) }); + convs.push_back({"C", user, "What is the main colour in it? Answer in a few words."}); + } + return run_multi_slot(*chat, convs, length, std::max(1, vm["rounds"].as())); + } + + // --prompt-file: one turn, nothing else, so the reported prefill is this prompt alone. + const std::string prompt_file = vm["prompt-file"].as(); + if (!prompt_file.empty()) { + std::ifstream pf(prompt_file); + if (!pf) { header_print("ERROR", "cannot open " << prompt_file); return 1; } + std::stringstream buf; buf << pf.rdbuf(); + lm_uniform_input_t one; + one.prompt = buf.str(); + std::cout << "Prompt file: " << prompt_file << " (" << one.prompt.size() << " bytes)" << std::endl; + chat_meta_info_t m; + m.restore_allowed = vm["resume"].as(); // see --resume + chat->start_total_timer(); + chat->insert(m, one); + const std::string answer = chat->generate(m, length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, m); + + const std::string follow_file = vm["follow-up"].as(); + if (!follow_file.empty()) { + std::ifstream ff(follow_file); + if (!ff) { header_print("ERROR", "cannot open " << follow_file); return 1; } + std::stringstream fbuf; fbuf << ff.rdbuf(); + lm_uniform_input_t two; + two.messages = nlohmann::ordered_json::array(); + two.messages.push_back({ {"role", "user"}, {"content", one.prompt} }); + two.messages.push_back({ {"role", "assistant"}, {"content", answer} }); + two.messages.push_back({ {"role", "user"}, {"content", fbuf.str()} }); + std::cout << "Follow-up file: " << follow_file << " (" << fbuf.str().size() << " bytes)" << std::endl; + chat_meta_info_t m2; + m2.restore_allowed = true; + chat->start_total_timer(); + chat->insert(m2, two); + chat->generate(m2, length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, m2); + } + return 0; + } + + // Turn 1: a pure prompt, as `flm run` sends it. + lm_uniform_input_t input; + input.prompt = "My name is Ada and I live in Lisbon. In one sentence, what is Lisbon known for?"; + std::cout << "Prompt: " << input.prompt << std::endl << "Response: " << std::endl; + chat_meta_info_t meta_info; + chat->start_total_timer(); + chat->insert(meta_info, input); + std::string answer = chat->generate(meta_info, length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, meta_info); + + // Turn 2: the whole conversation, as `flm serve` sends it on a prompt-cache + // hit. restore() rewinds to the end of turn 1's prompt, so only the + // assistant reply and the new user turn should be prefilled. + lm_uniform_input_t follow_up; + follow_up.messages = nlohmann::ordered_json::array(); + follow_up.messages.push_back({ {"role", "user"}, {"content", input.prompt} }); + follow_up.messages.push_back({ {"role", "assistant"}, {"content", answer} }); + follow_up.messages.push_back({ {"role", "user"}, {"content", "What is my name, and which city did I mention?"} }); + std::cout << "Prompt: " << follow_up.messages.back()["content"].get() << std::endl << "Response: " << std::endl; + chat_meta_info_t meta_info2; + meta_info2.restore_allowed = true; + chat->start_total_timer(); + chat->insert(meta_info2, follow_up); + chat->generate(meta_info2, length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, meta_info2); + + // Turn 3: a fresh conversation with thinking on. + chat->clear_context(); + chat->configure_parameter("enable_think", true); + lm_uniform_input_t fresh; + fresh.prompt = "What is 17 * 23? Answer briefly."; + std::cout << "Prompt: " << fresh.prompt << std::endl << "Response: " << std::endl; + chat_meta_info_t meta_info3; + chat->start_total_timer(); + chat->insert(meta_info3, fresh); + chat->generate(meta_info3, 4 * length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, meta_info3); + + // Turn 4: an image, as `flm run` sends it after /input. The vision tower only + // runs if the package has one -- config.json needs vision_model_weight, which + // is what sets is_vlm. Without it MiniCPM_OCR warns and drops the image, and + // this turn degenerates to a text question about nothing, so say so rather + // than let a 25-token prefill read as success. (`image` is resolved above.) + if (image.empty()) { + header_print("info", "turn 4 (image) skipped: --image \"\""); + } else if (!std::filesystem::exists(image)) { + header_print("WARNING", "turn 4 (image) skipped: no such file: " << image); + } else { + chat->clear_context(); + chat->configure_parameter("enable_think", false); + lm_uniform_input_t with_image; + with_image.prompt = "Describe this image."; + with_image.images.push_back(image); + std::cout << "Image: " << image << std::endl; + report_image(image); + std::cout << "Prompt: " << with_image.prompt << std::endl << "Response: " << std::endl; + chat_meta_info_t meta_info4; + chat->start_total_timer(); + chat->insert(meta_info4, with_image); + chat->generate(meta_info4, length, std::cout); + chat->stop_total_timer(); + report_turn(*chat, meta_info4); + + // The tower turns the repo's 960x720 test image into 5 views / 5040 + // patches / 315 rows, so a real image prefill is ~330 tokens against ~15 + // for the question alone. Anything near the latter means the image never + // reached the tower -- a missing vision_model_weight, or a path the + // reader could not decode -- which is otherwise easy to miss, because the + // model still answers, just from the question alone. + constexpr int kMinImagePromptTokens = 100; + if (meta_info4.prompt_tokens < kMinImagePromptTokens) { + header_print("WARNING", "prompt_tokens=" << meta_info4.prompt_tokens + << " is too low for an image turn; the image was NOT consumed" + " (is_vlm false, or the file failed to decode)"); + } else { + header_print("info", "image consumed: " << meta_info4.prompt_tokens + << " prompt tokens (question alone would be ~15)"); + } + } + + return 0; +} diff --git a/src/wix/flm.wxs b/src/wix/flm.wxs index b3382e11b..2e6bc146d 100644 --- a/src/wix/flm.wxs +++ b/src/wix/flm.wxs @@ -173,6 +173,9 @@ + + + diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/GateDeltaNet_prefill.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/GateDeltaNet_prefill.xclbin new file mode 100644 index 000000000..e28233856 Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/GateDeltaNet_prefill.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/attn.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/attn.xclbin new file mode 100644 index 000000000..6bcd5eceb Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/attn.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/conv.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/conv.xclbin new file mode 100644 index 000000000..7d70327b8 Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/conv.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/layer.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/layer.xclbin new file mode 100644 index 000000000..5aa3c61a0 Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/layer.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/lm_head.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/lm_head.xclbin new file mode 100644 index 000000000..27fde16af Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/lm_head.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/mm.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/mm.xclbin new file mode 100644 index 000000000..b5393ec63 Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/mm.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_attn.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_attn.xclbin new file mode 100644 index 000000000..0c369d79d Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_attn.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_a.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_a.xclbin new file mode 100644 index 000000000..45eae121b Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_a.xclbin differ diff --git a/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_b.xclbin b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_b.xclbin new file mode 100644 index 000000000..8dd2e10b2 Binary files /dev/null and b/src/xclbins/MiniCPM-OCR-1.0-NPU2/vision_mm_b.xclbin differ