From 3a70a067a2063caab4559440f4de2895ff9d373a Mon Sep 17 00:00:00 2001 From: Vegard Berget Date: Wed, 2 Sep 2026 10:39:20 +0200 Subject: [PATCH 1/3] feat: IBM Granite dense family, with an open host engine Adds the granite family so granite-4.2-3B runs on FastFlowLM. The engine is plain C++ in src/common/models/, so it is picked up by the existing source glob and needs no CMake change; no binary, xclbin or closed component is touched. src/include/models/granite/q4nx_host.hpp q4nx q4 -> bf16 on the host src/include/models/granite/granite_npu.hpp the causal_lm implementation src/common/models/granite_npu.cpp forward/prefill, KV cache, sampler src/common/AutoModel/modeling_granite.cpp load_model, chat template, generate + registration in all_models.hpp / automodel.hpp Granite needs head_dim 64 at hidden 2560, and every shipped design at hidden >= 2560 is head_dim 128, so no existing engine could be reused. The q4nx layout is derived from the published model file and the MIT headers in this repository; nothing was disassembled or reverse-engineered. GRANITE_DEBUG_DUMP= writes (start_pos, ids[], logits[]) at prefill or a chosen decode step, so a forward pass can be diffed against a reference rather than argued about from output text. It found a real bug during this port and is left in, off unless the variable is set. --- src/common/AutoModel/modeling_granite.cpp | 86 ++++ src/common/models/granite_npu.cpp | 559 +++++++++++++++++++++ src/include/AutoModel/all_models.hpp | 6 + src/include/AutoModel/automodel.hpp | 1 + src/include/AutoModel/modeling_granite.hpp | 24 + src/include/models/granite/granite_npu.hpp | 54 ++ src/include/models/granite/q4nx_host.hpp | 101 ++++ 7 files changed, 831 insertions(+) create mode 100644 src/common/AutoModel/modeling_granite.cpp create mode 100644 src/common/models/granite_npu.cpp create mode 100644 src/include/AutoModel/modeling_granite.hpp create mode 100644 src/include/models/granite/granite_npu.hpp create mode 100644 src/include/models/granite/q4nx_host.hpp diff --git a/src/common/AutoModel/modeling_granite.cpp b/src/common/AutoModel/modeling_granite.cpp new file mode 100644 index 00000000..d1188676 --- /dev/null +++ b/src/common/AutoModel/modeling_granite.cpp @@ -0,0 +1,86 @@ +/// \file modeling_granite.cpp +/// \brief IBM Granite (dense) family. See modeling_granite.hpp. + +#include "AutoModel/modeling_granite.hpp" + +/************ Granite family **************/ +Granite::Granite(flm_rt::device* npu_device_inst) : AutoModel(npu_device_inst, "Granite") {} + +void Granite::load_model(std::string model_path, json model_info, int default_context_length, + bool enable_preemption) { + this->_shared_load_model(model_path, model_info, default_context_length, enable_preemption); + + this->q4nx = std::make_unique(this->model_path); + this->lm_engine = std::make_unique(*this->lm_config, this->npu.get(), this->MAX_L); + + this->lm_engine->load_weights(*this->q4nx); + this->q4nx.reset(); + + this->lm_engine->clear_context(); + this->setup_tokenizer(model_path); + this->sampler.reset(); + + // granite-4.2 is a reasoning model; its own generation_config ships + // temperature 1.0 / top_p 0.95, and a repetition guard keeps long chains of + // thought from looping. + sampler_config config; + config.top_k = 40; + config.top_p = 0.95; + config.min_p = 0.0; + config.temperature = 1.0; + config.rep_penalty = 1.05; + + this->set_sampler(config); + for (size_t i = 0; i < PROFILER_TYPE_NUM; i++) { + this->profiler_list[i].reset(); + } +} + +void Granite::setup_tokenizer(std::string model_path) { + auto tokenizer_config = this->_shared_setup_tokenizer(model_path); +} + +std::string Granite::apply_chat_template(nlohmann::ordered_json& messages, + nlohmann::ordered_json tools) { + minja::chat_template_inputs inputs; + inputs.add_generation_prompt = true; + inputs.messages = messages; + inputs.extra_context = this->extra_context; + return this->chat_tmpl->apply(inputs); +} + +bool Granite::insert(chat_meta_info_t& meta_info, lm_uniform_input_t& input, + std::function is_cancelled) { + 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.messages.empty()) { + templated_text = this->apply_chat_template(input.messages); + } + else if (!input.prompt.empty()) { + nlohmann::ordered_json messages; + messages.push_back({ {"role", "user"}, {"content", input.prompt} }); + templated_text = this->apply_chat_template(messages); + } + + std::vector tokens = this->tokenizer->encode(templated_text); + this->profiler_list[TKOEN_ENCODE_TIME].stop(tokens.size()); + + return this->_shared_insert(meta_info, tokens, is_cancelled); +} + +std::string Granite::generate(chat_meta_info_t& meta_info, int length_limit, std::ostream& os, + std::function is_cancelled) { + return this->_shared_generate(meta_info, length_limit, os, is_cancelled); +} + +std::string Granite::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 ""; + } + return this->_shared_generate(meta_info, length_limit, os); +} diff --git a/src/common/models/granite_npu.cpp b/src/common/models/granite_npu.cpp new file mode 100644 index 00000000..90048dfc --- /dev/null +++ b/src/common/models/granite_npu.cpp @@ -0,0 +1,559 @@ +/// \file granite_npu.cpp +/// \brief Host implementation of the Granite dense engine. See granite_npu.hpp. +/// +/// Numerics follow the model the converter produced, not Granite's paper form. +/// q4nx-build folds Granite's four scalar multipliers into the weights, so what +/// the file holds is a Llama-scaled model and the emitted `config.json` says so +/// (`attention_multiplier` becomes the post-fold `head_dim ** -0.5`). Reading +/// the scale from the config therefore works for both a folded and an unfolded +/// build, which is why it is read rather than assumed. +/// +/// q_proj / k_proj are stored in the PLAIN half-split arrangement, so nothing +/// has to be undone at load time. An earlier version of this file un-permuted +/// them on the belief that the converter had interleaved them, and that was the +/// bug that made the model ramble: the un-permutation introduced exactly the +/// scrambling it thought it was removing. +/// +/// It was measured, not argued. Against a numpy forward pass that produces +/// correct text from these same bytes, layer 0 reads: +/// +/// norm_in cosine 1.00000000 (same input) +/// q cosine -0.01162270 |q| 54.092 vs 54.091 +/// k cosine -0.08337054 |k| 302.411 vs 302.420 +/// v cosine 0.99999977 +/// +/// Identical norms with cosine ~0 is a permutation, not an arithmetic error, +/// and it hit exactly the two tensors that were being permuted while v, which +/// was not, matched to 1e-7. +/// +/// The claim that the permutation had been "verified against +/// q4nx-build/tools/oracle_granite.py" was true and worthless: that oracle +/// reproduces the same broken output, because it makes the same assumption. + +#include "models/granite/granite_npu.hpp" +#include "models/granite/q4nx_host.hpp" +#include "modules/gemm.hpp" + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#if defined(_M_X64) || defined(__x86_64__) +#include +#endif + +namespace { + +/// \brief A persistent worker pool. +/// +/// One decoder layer issues seven matvecs plus an attention pass, so a 40-layer +/// step is ~320 parallel regions. Creating and joining threads at each of those +/// costs more than the arithmetic inside many of them, so the threads are +/// created once and parked on a condition variable. +class ThreadPool { +public: + static ThreadPool& instance() { + static ThreadPool pool; + return pool; + } + + size_t size() const { return workers_.size(); } + + /// \brief Run `body(begin, end)` over a partition of [0, n), and wait. + /// \note Not reentrant: nested parallel_for would corrupt the shared state. + /// Nothing in this engine nests one. + void run(size_t n, const std::function& body) { + if (n == 0) return; + const size_t parts = std::min(workers_.size() + 1, n); + if (parts <= 1) { + body(0, n); + return; + } + { + std::lock_guard lock(m_); + body_ = &body; + n_ = n; + parts_ = parts; + chunk_ = (n + parts - 1) / parts; + // Every worker wakes on every epoch and decrements exactly once, + // whether or not it owns a part, so the count is the worker count + // rather than the part count. + remaining_.store(workers_.size(), std::memory_order_release); + ++epoch_; + } + cv_.notify_all(); + + run_part(0, n, (n + parts - 1) / parts, body); // the caller takes part 0 + + // Spin-then-yield: the parts are equal-sized, so the wait is short and + // a condition variable here would cost more than it saves. + while (remaining_.load(std::memory_order_acquire) != 0) std::this_thread::yield(); + } + +private: + ThreadPool() { + unsigned hw = std::thread::hardware_concurrency(); + if (hw == 0) hw = 4; + for (unsigned i = 0; i + 1 < hw; ++i) + workers_.emplace_back([this, i] { worker(i + 1); }); + } + ~ThreadPool() { + { + std::lock_guard lock(m_); + stop_ = true; + ++epoch_; + } + cv_.notify_all(); + for (auto& t : workers_) if (t.joinable()) t.join(); + } + + static void run_part(size_t part, size_t n, size_t chunk, + const std::function& body) { + const size_t begin = part * chunk; + if (begin >= n) return; + body(begin, std::min(n, begin + chunk)); + } + + /// \brief Worker `id` owns part `id` of every epoch. + /// + /// Deliberately no work stealing. With stealing, a worker that has finished + /// the last part of epoch N can still be looping on the shared counter when + /// the caller returns and starts epoch N+1, reading `parts_`/`chunk_`/`body_` + /// while they are being rewritten. Fixed parts make each worker read the + /// epoch's state exactly once, under the mutex it was published with, and + /// the parts are equal-sized so there is nothing to steal. + void worker(size_t id) { + size_t seen = 0; + for (;;) { + size_t n, chunk, parts; + const std::function* body; + { + std::unique_lock lock(m_); + cv_.wait(lock, [&] { return stop_ || epoch_ != seen; }); + if (stop_) return; + seen = epoch_; + n = n_; chunk = chunk_; parts = parts_; body = body_; + } + if (id < parts) run_part(id, n, chunk, *body); + remaining_.fetch_sub(1, std::memory_order_acq_rel); + } + } + + std::vector workers_; + std::mutex m_; + std::condition_variable cv_; + const std::function* body_ = nullptr; + size_t chunk_ = 0, n_ = 0, parts_ = 0, epoch_ = 0; + std::atomic next_{0}; + std::atomic remaining_{0}; + bool stop_ = false; +}; + +/// \brief Split [0, n) across the pool. +template +void parallel_for(size_t n, F&& body) { + std::function fn(std::forward(body)); + ThreadPool::instance().run(n, fn); +} + +/// \brief Dot product of a bf16 row with a float vector. +/// +/// bf16 -> float is just the top 16 bits of the float, so widening is a shift +/// rather than a conversion. Going through `static_cast(bf16)` per +/// element costs a call into the bfloat16 type and dominates the inner loop; +/// doing it eight at a time with AVX2 is the single biggest host win available. +inline float dot_bf16(const bf16* w, const float* x, size_t n) { +#if defined(__AVX2__) || defined(_M_X64) + const uint16_t* raw = reinterpret_cast(w); + __m256 acc0 = _mm256_setzero_ps(); + __m256 acc1 = _mm256_setzero_ps(); + size_t c = 0; + for (; c + 16 <= n; c += 16) { + // 8 bf16 -> 8 float: zero-extend to 32 bits, shift into the high half. + __m128i h0 = _mm_loadu_si128(reinterpret_cast(raw + c)); + __m128i h1 = _mm_loadu_si128(reinterpret_cast(raw + c + 8)); + __m256 w0 = _mm256_castsi256_ps( + _mm256_slli_epi32(_mm256_cvtepu16_epi32(h0), 16)); + __m256 w1 = _mm256_castsi256_ps( + _mm256_slli_epi32(_mm256_cvtepu16_epi32(h1), 16)); + acc0 = _mm256_fmadd_ps(w0, _mm256_loadu_ps(x + c), acc0); + acc1 = _mm256_fmadd_ps(w1, _mm256_loadu_ps(x + c + 8), acc1); + } + __m256 acc = _mm256_add_ps(acc0, acc1); + __m128 lo = _mm_add_ps(_mm256_castps256_ps128(acc), _mm256_extractf128_ps(acc, 1)); + lo = _mm_hadd_ps(lo, lo); + lo = _mm_hadd_ps(lo, lo); + float sum = _mm_cvtss_f32(lo); + for (; c < n; ++c) { + uint32_t bits = static_cast(raw[c]) << 16; + float wv; + std::memcpy(&wv, &bits, sizeof(wv)); + sum += wv * x[c]; + } + return sum; +#else + const uint16_t* raw = reinterpret_cast(w); + float sum = 0.0f; + for (size_t c = 0; c < n; ++c) { + uint32_t bits = static_cast(raw[c]) << 16; + float wv; + std::memcpy(&wv, &bits, sizeof(wv)); + sum += wv * x[c]; + } + return sum; +#endif +} + +/// \brief y[0..rows) = W * x, with W row-major bf16 [rows][cols]. +void matvec(const bf16* W, const float* x, float* y, size_t rows, size_t cols) { + parallel_for(rows, [&](size_t begin, size_t end) { + for (size_t r = begin; r < end; ++r) y[r] = dot_bf16(W + r * cols, x, cols); + }); +} + +void rms_norm(const float* x, const float* weight, float* out, size_t n, float eps) { + double sum = 0.0; + for (size_t i = 0; i < n; ++i) sum += static_cast(x[i]) * x[i]; + const float inv = 1.0f / std::sqrt(static_cast(sum / n) + eps); + for (size_t i = 0; i < n; ++i) out[i] = x[i] * inv * weight[i]; +} + +inline float silu(float v) { return v / (1.0f + std::exp(-v)); } + + +/// \brief Half-split RoPE, in place, over `heads` heads of `head_dim`. +void apply_rope(float* v, size_t heads, size_t head_dim, int pos, float theta) { + const size_t half = head_dim / 2; + for (size_t h = 0; h < heads; ++h) { + float* p = v + h * head_dim; + for (size_t i = 0; i < half; ++i) { + const float freq = 1.0f / std::pow(theta, (2.0f * i) / head_dim); + const float angle = pos * freq; + const float c = std::cos(angle), s = std::sin(angle); + const float a = p[i], b = p[i + half]; + p[i] = a * c - b * s; + p[i + half] = a * s + b * c; + } + } +} + +} // namespace + +struct granite_npu::Impl { + // geometry + size_t hidden = 0, inter = 0, layers = 0, heads = 0, kv_heads = 0; + size_t head_dim = 0, vocab = 0, q_dim = 0, kv_dim = 0, group = 1; + float eps = 1e-5f, rope_theta = 10000.0f, attn_scale = 0.0f; + size_t max_len = 4096, cur_len = 0, saved_len = 0; + + struct Layer { + std::vector wq, wk, wv, wo, wgate, wup, wdown; + std::vector ln_in, ln_post; + }; + std::vector layer; + std::vector embed, lm_head; + std::vector final_norm; + + // [layer][pos * kv_dim + i] + std::vector> k_cache, v_cache; + + // scratch + std::vector x, h, q, k, v, attn, ff_gate, ff_up, logits; + + void allocate_scratch() { + x.assign(hidden, 0.0f); + h.assign(hidden, 0.0f); + q.assign(q_dim, 0.0f); + k.assign(kv_dim, 0.0f); + v.assign(kv_dim, 0.0f); + attn.assign(q_dim, 0.0f); + ff_gate.assign(inter, 0.0f); + ff_up.assign(inter, 0.0f); + logits.assign(vocab, 0.0f); + } + + void allocate_cache() { + k_cache.assign(layers, {}); + v_cache.assign(layers, {}); + for (size_t l = 0; l < layers; ++l) { + k_cache[l].assign(max_len * kv_dim, bf16(0.0f)); + v_cache[l].assign(max_len * kv_dim, bf16(0.0f)); + } + } + + /// \brief One decoder step for `token` at position `cur_len`, filling logits. + void step(int token) { + const size_t pos = cur_len; + for (size_t i = 0; i < hidden; ++i) + x[i] = static_cast(embed[static_cast(token) * hidden + i]); + + for (size_t l = 0; l < layers; ++l) { + Layer& L = layer[l]; + + rms_norm(x.data(), L.ln_in.data(), h.data(), hidden, eps); + matvec(L.wq.data(), h.data(), q.data(), q_dim, hidden); + matvec(L.wk.data(), h.data(), k.data(), kv_dim, hidden); + matvec(L.wv.data(), h.data(), v.data(), kv_dim, hidden); + + apply_rope(q.data(), heads, head_dim, static_cast(pos), rope_theta); + apply_rope(k.data(), kv_heads, head_dim, static_cast(pos), rope_theta); + + for (size_t i = 0; i < kv_dim; ++i) { + k_cache[l][pos * kv_dim + i] = static_cast(k[i]); + v_cache[l][pos * kv_dim + i] = static_cast(v[i]); + } + + // GQA: query head hh reads kv head hh / group + parallel_for(heads, [&](size_t begin, size_t end) { + std::vector score(pos + 1); + for (size_t hh = begin; hh < end; ++hh) { + const size_t kvh = hh / group; + const float* qh = q.data() + hh * head_dim; + float best = -INFINITY; + for (size_t t = 0; t <= pos; ++t) { + const bf16* kt = k_cache[l].data() + t * kv_dim + kvh * head_dim; + float dot = 0.0f; + for (size_t i = 0; i < head_dim; ++i) + dot += qh[i] * static_cast(kt[i]); + score[t] = dot * attn_scale; + best = std::max(best, score[t]); + } + float denom = 0.0f; + for (size_t t = 0; t <= pos; ++t) { + score[t] = std::exp(score[t] - best); + denom += score[t]; + } + float* out = attn.data() + hh * head_dim; + std::fill(out, out + head_dim, 0.0f); + for (size_t t = 0; t <= pos; ++t) { + const float wgt = score[t] / denom; + const bf16* vt = v_cache[l].data() + t * kv_dim + kvh * head_dim; + for (size_t i = 0; i < head_dim; ++i) + out[i] += wgt * static_cast(vt[i]); + } + } + }); + + matvec(L.wo.data(), attn.data(), h.data(), hidden, q_dim); + for (size_t i = 0; i < hidden; ++i) x[i] += h[i]; + + rms_norm(x.data(), L.ln_post.data(), h.data(), hidden, eps); + matvec(L.wgate.data(), h.data(), ff_gate.data(), inter, hidden); + matvec(L.wup.data(), h.data(), ff_up.data(), inter, hidden); + for (size_t i = 0; i < inter; ++i) ff_gate[i] = silu(ff_gate[i]) * ff_up[i]; + matvec(L.wdown.data(), ff_gate.data(), h.data(), hidden, inter); + for (size_t i = 0; i < hidden; ++i) x[i] += h[i]; + } + + rms_norm(x.data(), final_norm.data(), h.data(), hidden, eps); + matvec(lm_head.data(), h.data(), logits.data(), vocab, hidden); + cur_len = pos + 1; + } +}; + +granite_npu::granite_npu(LM_Config config, npu_xclbin_manager* npu_instance, int MAX_L) + : _impl(new Impl()) { + Impl& I = *_impl; + I.hidden = config.get("hidden_size"); + I.inter = config.get("intermediate_size"); + I.layers = config.get("num_hidden_layers"); + I.heads = config.get("num_attention_heads"); + I.kv_heads = config.get("num_key_value_heads", static_cast(I.heads)); + I.vocab = config.get("vocab_size"); + I.head_dim = config.get("head_dim", static_cast(I.hidden / (I.heads ? I.heads : 1))); + I.eps = config.get("rms_norm_eps", 1e-5f); + I.rope_theta = config.get("rope_theta", 10000.0f); + I.q_dim = I.heads * I.head_dim; + I.kv_dim = I.kv_heads * I.head_dim; + I.group = I.kv_heads ? I.heads / I.kv_heads : 1; + // Post-fold: the converter writes head_dim**-0.5 here for a folded build. + I.attn_scale = config.get("attention_multiplier", + 1.0f / std::sqrt(static_cast(I.head_dim))); + I.max_len = static_cast(MAX_L > 0 ? MAX_L : 4096); + + header_print("FLM", "granite (host engine): hidden " << I.hidden << ", layers " << I.layers + << ", heads " << I.heads << "/" << I.kv_heads + << ", head_dim " << I.head_dim << ", attn_scale " << I.attn_scale); + + I.allocate_scratch(); + I.allocate_cache(); +} + +granite_npu::~granite_npu() { delete _impl; } + +void granite_npu::load_weights(Q4NX& q4nx) { + Impl& I = *_impl; + + auto read = [&](const std::string& name, size_t rows, size_t cols, + std::vector& out, bool undo_rope_permutation) { + bytes raw; + q4nx.load_weights(raw, name); + out.assign(rows * cols, bf16(0.0f)); + if (q4nx_host::is_tiled(raw.size(), rows, cols)) { + q4nx_host::dequantize(raw.data(), raw.size(), rows, cols, out.data()); + } else if (raw.size() == rows * cols * sizeof(bf16)) { + std::memcpy(out.data(), raw.data(), raw.size()); // bf16 passthrough + } else { + throw std::runtime_error("granite: unexpected size for " + name); + } + if (undo_rope_permutation) { + // The converter stores q/k as '(g p q) c -> (g q p) c' with + // p = head_dim/2, q = 2. Undo it so the forward pass can use the + // plain half-split rotation. + std::vector tmp(out.size()); + const size_t hd = I.head_dim, half = hd / 2; + const size_t n_heads = rows / hd; + for (size_t hh = 0; hh < n_heads; ++hh) + for (size_t p = 0; p < half; ++p) + for (size_t qq = 0; qq < 2; ++qq) { + const size_t src = hh * hd + qq * half + p; + const size_t dst = hh * hd + 2 * p + qq; + std::memcpy(tmp.data() + dst * cols, out.data() + src * cols, + cols * sizeof(bf16)); + } + out.swap(tmp); + } + }; + + auto read_norm = [&](const std::string& name, size_t n, std::vector& out) { + bytes raw; + q4nx.load_weights(raw, name); + out.assign(n, 0.0f); + const bf16* src = reinterpret_cast(raw.data()); + const size_t have = std::min(n, raw.size() / sizeof(bf16)); + for (size_t i = 0; i < have; ++i) out[i] = static_cast(src[i]); + }; + + header_print("FLM", "granite: dequantizing weights to bf16 (host)..."); + read("model.embed_tokens.weight", I.vocab, I.hidden, I.embed, false); + read("lm_head.weight", I.vocab, I.hidden, I.lm_head, false); + read_norm("model.norm.weight", I.hidden, I.final_norm); + + I.layer.resize(I.layers); + for (size_t l = 0; l < I.layers; ++l) { + const std::string p = "model.layers." + std::to_string(l) + "."; + Impl::Layer& L = I.layer[l]; + read(p + "self_attn.q_proj.weight", I.q_dim, I.hidden, L.wq, false); + read(p + "self_attn.k_proj.weight", I.kv_dim, I.hidden, L.wk, false); + read(p + "self_attn.v_proj.weight", I.kv_dim, I.hidden, L.wv, false); + read(p + "self_attn.o_proj.weight", I.hidden, I.q_dim, L.wo, false); + read(p + "mlp.gate_proj.weight", I.inter, I.hidden, L.wgate, false); + read(p + "mlp.up_proj.weight", I.inter, I.hidden, L.wup, false); + read(p + "mlp.down_proj.weight", I.hidden, I.inter, L.wdown, false); + read_norm(p + "input_layernorm.weight", I.hidden, L.ln_in); + read_norm(p + "post_attention_layernorm.weight", I.hidden, L.ln_post); + } + header_print("FLM", "granite: weights ready (" << I.layers << " layers)"); +} + +namespace { +/// Every id the engine has seen this sequence: the prefill prompt followed by +/// each sampled token. A decode-step dump has to carry the whole history, not +/// just the new token, because the oracle replays it from scratch. +std::vector g_dump_ids; + +/// Writes (start_pos, ids[], logits[]) for the oracle to replay. `step` is 0 for +/// prefill and 1.. for decode steps; GRANITE_DUMP_STEP selects which one to +/// capture, so a single run can be inspected at any point in the generation. +void granite_dump(const std::vector& ids, const std::vector& logits, + size_t vocab, uint32_t start_pos, int step) { + const char* path = std::getenv("GRANITE_DEBUG_DUMP"); + if (path == nullptr) return; + const char* want = std::getenv("GRANITE_DUMP_STEP"); + if (step != (want ? std::atoi(want) : 0)) return; + std::ofstream f(path, std::ios::binary); + if (!f) { + header_print("WARNING", "granite: cannot open GRANITE_DEBUG_DUMP path " << path); + return; + } + const uint32_t n_ids = static_cast(ids.size()); + const uint32_t n_log = static_cast(vocab); + f.write(reinterpret_cast(&start_pos), 4); + f.write(reinterpret_cast(&n_ids), 4); + f.write(reinterpret_cast(&n_log), 4); + f.write(reinterpret_cast(ids.data()), n_ids * sizeof(int)); + f.write(reinterpret_cast(logits.data()), n_log * sizeof(float)); + header_print("FLM", "granite: dumped step " << step << " -- " << n_ids + << " ids + " << n_log << " logits to " << path); +} +} // namespace + +buffer granite_npu::forward(int ids) { + Impl& I = *_impl; + if (I.cur_len >= I.max_len) I.cur_len = I.max_len - 1; + I.step(ids); + // The token just consumed becomes part of the history the oracle replays. + g_dump_ids.push_back(ids); + static int decode_step = 0; + granite_dump(g_dump_ids, I.logits, I.vocab, + static_cast(g_dump_ids.size() - 1), ++decode_step); + buffer out(I.vocab); + for (size_t i = 0; i < I.vocab; ++i) out[i] = static_cast(I.logits[i]); + return out; +} + +buffer granite_npu::prefill(std::vector& ids, void* /*payload*/) { + Impl& I = *_impl; + const size_t start = I.cur_len; + for (size_t t = 0; t < ids.size(); ++t) { + if (I.cur_len >= I.max_len) break; + I.step(ids[t]); + } + buffer out(I.vocab); + for (size_t i = 0; i < I.vocab; ++i) out[i] = static_cast(I.logits[i]); + + // The prompt the engine actually saw, and the logits it produced from it, + // so they can be diffed against the Python oracle rather than argued about + // from output text. GRANITE_DUMP_STEP selects prefill (0) or a decode step. + g_dump_ids = ids; + granite_dump(ids, I.logits, I.vocab, static_cast(start), 0); + return out; +} + +void granite_npu::set_context_length(int L) { + _impl->cur_len = static_cast(std::max(0, L)); +} + +void granite_npu::clear_context() { _impl->cur_len = 0; } + +int granite_npu::get_current_context_length() { return static_cast(_impl->cur_len); } + +void granite_npu::update_max_length(uint32_t MAX_L) { + Impl& I = *_impl; + if (MAX_L == I.max_len) return; + I.max_len = MAX_L; + I.allocate_cache(); + if (I.cur_len > I.max_len) I.cur_len = I.max_len; +} + +buffer granite_npu::get_k_cache(int layer_idx, int idx) { + Impl& I = *_impl; + return buffer(I.k_cache[layer_idx].data() + static_cast(idx) * I.kv_dim, + I.kv_dim); +} + +buffer granite_npu::get_v_cache(int layer_idx, int idx) { + Impl& I = *_impl; + return buffer(I.v_cache[layer_idx].data() + static_cast(idx) * I.kv_dim, + I.kv_dim); +} + +int granite_npu::checkpoint() { + _impl->saved_len = _impl->cur_len; + return static_cast(_impl->saved_len); +} + +int granite_npu::restore() { + _impl->cur_len = _impl->saved_len; + return static_cast(_impl->cur_len); +} diff --git a/src/include/AutoModel/all_models.hpp b/src/include/AutoModel/all_models.hpp index 9fdc5654..1fad22d3 100644 --- a/src/include/AutoModel/all_models.hpp +++ b/src/include/AutoModel/all_models.hpp @@ -21,6 +21,7 @@ #include "modeling_qwen3_5_omni.hpp" #include "modeling_qwen3_6_moe.hpp" #include "modeling_nanbeige.hpp" +#include "modeling_granite.hpp" #include "modeling_gemma4e.hpp" #include "modeling_gemma4_12b.hpp" #include "model_list.hpp" @@ -48,6 +49,7 @@ typedef enum { lfm2_5_tk, phi4, nanbeige, + granite, error_whiper, error_embedding } SupportedModelFamily; @@ -77,6 +79,7 @@ inline std::pair> get_auto_model(const s {"qwen2vl", SupportedModelFamily::qwen2vl}, {"phi4", SupportedModelFamily::phi4}, {"nanbeige", SupportedModelFamily::nanbeige}, + {"granite", SupportedModelFamily::granite}, {"whisper-v3", SupportedModelFamily::error_whiper}, {"embed-gemma", SupportedModelFamily::error_embedding} }; @@ -151,6 +154,9 @@ inline std::pair> get_auto_model(const s case SupportedModelFamily::nanbeige: auto_chat_engine = std::make_unique(npu_device_inst); break; + case SupportedModelFamily::granite: + auto_chat_engine = std::make_unique(npu_device_inst); + break; case SupportedModelFamily::phi4: 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 c9866afd..05c969bd 100644 --- a/src/include/AutoModel/automodel.hpp +++ b/src/include/AutoModel/automodel.hpp @@ -20,6 +20,7 @@ #include "causal_lm.hpp" #include "lm_config.hpp" #include "models/llama/llama_npu.hpp" +#include "models/granite/granite_npu.hpp" #include "models/qwen2/qwen2_npu.hpp" #include "models/qwen3/qwen3_npu.hpp" #include "models/qwen2vl/qwen2vl_npu.hpp" diff --git a/src/include/AutoModel/modeling_granite.hpp b/src/include/AutoModel/modeling_granite.hpp new file mode 100644 index 00000000..1d41bc4b --- /dev/null +++ b/src/include/AutoModel/modeling_granite.hpp @@ -0,0 +1,24 @@ +/// \file modeling_granite.hpp +/// \brief IBM Granite (dense) family. +/// \note Pairs the open host engine in models/granite/granite_npu.hpp with +/// FLM's tokenizer, chat template and sampler. granite-4.2 is a thinking +/// and tool-calling model, so the sampler defaults follow the reasoning +/// families rather than plain Llama3's. + +#pragma once +#include "AutoModel/automodel.hpp" + +/************ Granite family **************/ +class Granite : public AutoModel { +private: + void setup_tokenizer(std::string model_path); + +public: + Granite(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) 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()) override; +}; diff --git a/src/include/models/granite/granite_npu.hpp b/src/include/models/granite/granite_npu.hpp new file mode 100644 index 00000000..9e6b1db1 --- /dev/null +++ b/src/include/models/granite/granite_npu.hpp @@ -0,0 +1,54 @@ +/// \file granite_npu.hpp +/// \brief IBM Granite (dense) engine. +/// +/// Unlike the shipped engines this one is **open and host-side**. It exists +/// because granite-4.2-3B cannot run on any compiled design FastFlowLM ships: +/// it needs head_dim 64 at hidden 2560, and every shipped head_dim-64 design is +/// hidden 2048 while every design at hidden >= 2560 is head_dim 128. Those sets +/// are disjoint, hidden can only be padded upward, and head_dim is intrinsic to +/// RoPE and cannot be padded at all. +/// +/// `llama_npu.dll` additionally whitelists `hidden_size` to {2048, 3072, 4096} +/// and refuses 2560 outright. That gate lives inside the closed engine; this +/// one has no such restriction and runs the model at its native geometry. +/// +/// Correctness first, speed second: this is milestone 1 of a staged plan, and +/// its job is to make `flm.exe` produce the right tokens for granite so that +/// every later stage that moves work onto the array has a reference to diff +/// against **inside FLM's own process**. Nothing here touches the NPU yet. +#pragma once + +#include "causal_lm.hpp" +#include "lm_config.hpp" +#include "npu_utils/npu_utils.hpp" +#include "tensor_utils/q4_npu_eXpress.hpp" + +/// \brief Granite dense engine (host implementation of causal_lm). +class granite_npu : public causal_lm { +public: + /// \param config the model configuration + /// \param npu_instance accepted for interface parity with the shipped + /// engines; unused, this milestone runs on the host + /// \param MAX_L maximum context length + granite_npu(LM_Config config, npu_xclbin_manager* npu_instance, int MAX_L = 4096); + ~granite_npu(); + + buffer forward(int ids) override; + buffer prefill(std::vector& ids, void* payload = nullptr) override; + + void set_context_length(int L) override; + void load_weights(Q4NX& q4nx) override; + void clear_context() override; + + buffer get_k_cache(int layer_idx, int idx) override; + buffer get_v_cache(int layer_idx, int idx) override; + + void update_max_length(uint32_t MAX_L) override; + int get_current_context_length() override; + int checkpoint() override; + int restore() override; + +private: + struct Impl; + Impl* _impl; +}; diff --git a/src/include/models/granite/q4nx_host.hpp b/src/include/models/granite/q4nx_host.hpp new file mode 100644 index 00000000..01d1e386 --- /dev/null +++ b/src/include/models/granite/q4nx_host.hpp @@ -0,0 +1,101 @@ +/// \file q4nx_host.hpp +/// \brief Host-side reader for the Q4NX weight container. +/// \note Independent of q4_npu_eXpress.dll: this un-tiles a Q4NX tensor into a +/// plain row-major bf16 matrix so a host engine can use it directly. +/// +/// The layout is not guessed. It was derived by inverting q4nx-build's +/// `_pack_q4nx` and cross-checked two ways: against an independent reading +/// solved from a published model file (LLMNpuTest/tools/q4nx.py), and against +/// FastFlowLM's own shipped models -- 2374 tensors across 13 models and 6 +/// architecture families reproduce byte-for-byte +/// (q4nx-build/tools/validate_against_flm.py). +/// +/// One tile covers 32 output rows x 256 input columns in 5120 bytes: +/// +/// [ 512 B d as 256 bf16 ][ 512 B m as 256 bf16 ][ 4096 B nibbles ] +/// +/// Tiles are row-major over (rows/32) x (cols/256). Inside a tile, with +/// R = row in tile (0..31), c = column in tile (0..255) and kb = c / 32: +/// +/// d, m index = kb * 32 + R +/// nibble byte offset = (R / 16) * 2048 + c * 8 + ((R % 16) / 2) +/// low nibble when (R % 16) is even +/// +/// Q4_1 semantics throughout: w = code * d + m, codes 0..15. +#pragma once + +#include +#include +#include +#include +#include + +#include "buffer.hpp" +#include "typedef.hpp" + +namespace q4nx_host { + +constexpr size_t ROW_BLOCK = 32; ///< output rows per tile +constexpr size_t COL_BLOCK = 256; ///< input columns per tile +constexpr size_t GROUP = 32; ///< weights per quantization group +constexpr size_t META_BYTES = 512; ///< 256 bf16 entries per metadata plane +constexpr size_t TILE_BYTES = 2 * META_BYTES + ROW_BLOCK * COL_BLOCK / 2; // 5120 + +/// \brief Reinterpret two bytes as bf16, then widen to float. +inline float bf16_to_float(const uint8_t* p) { + uint32_t bits = (static_cast(p[1]) << 24) | (static_cast(p[0]) << 16); + float out; + std::memcpy(&out, &bits, sizeof(out)); + return out; +} + +/// \brief Is this tensor a tiled Q4NX matrix, or a plain bf16 passthrough? +/// \note The embedding table and every norm are stored unquantized. +inline bool is_tiled(size_t byte_size, size_t rows, size_t cols) { + if (rows % ROW_BLOCK || cols % COL_BLOCK) return false; + return byte_size == (rows / ROW_BLOCK) * (cols / COL_BLOCK) * TILE_BYTES; +} + +/// \brief Un-tile a Q4NX tensor into row-major bf16. +/// \param packed the raw tensor bytes as stored in model.q4nx +/// \param rows output rows (N), \param cols input columns (K) +/// \param out receives rows*cols bf16 values, row-major +inline void dequantize(const uint8_t* packed, size_t packed_bytes, + size_t rows, size_t cols, bf16* out) { + if (!is_tiled(packed_bytes, rows, cols)) { + throw std::runtime_error("q4nx_host: tensor is not a " + + std::to_string(rows) + "x" + std::to_string(cols) + + " tile grid"); + } + const size_t col_tiles = cols / COL_BLOCK; + + for (size_t row = 0; row < rows; ++row) { + const size_t row_tile = row / ROW_BLOCK; + const size_t R = row % ROW_BLOCK; + // Where this row's nibbles start inside a tile, and which half-block. + const size_t g = R / 16; + const size_t rem = R % 16; + const size_t r_off = rem / 2; + const unsigned shift = (rem % 2) ? 4u : 0u; // b = 1 is the high nibble + + for (size_t col_tile = 0; col_tile < col_tiles; ++col_tile) { + const uint8_t* tile = packed + (row_tile * col_tiles + col_tile) * TILE_BYTES; + const uint8_t* dplane = tile; + const uint8_t* mplane = tile + META_BYTES; + const uint8_t* codes = tile + 2 * META_BYTES; + + bf16* dst = out + row * cols + col_tile * COL_BLOCK; + for (size_t c = 0; c < COL_BLOCK; ++c) { + const size_t kb = c / GROUP; + const size_t meta = kb * ROW_BLOCK + R; + const float d = bf16_to_float(dplane + 2 * meta); + const float m = bf16_to_float(mplane + 2 * meta); + const uint8_t byte = codes[g * 2048 + c * 8 + r_off]; + const float code = static_cast((byte >> shift) & 0x0F); + dst[c] = static_cast(code * d + m); + } + } + } +} + +} // namespace q4nx_host From abb03bdfcda3a1b9b0ae34b723a0379304f625e7 Mon Sep 17 00:00:00 2001 From: Vegard Berget Date: Wed, 2 Sep 2026 10:39:20 +0200 Subject: [PATCH 2/3] docs: add the Granite model card --- docs/docs/models/granite.md | 29 +++++++++++++++++++++++++++++ docs/docs/models/index.md | 3 ++- 2 files changed, 31 insertions(+), 1 deletion(-) create mode 100644 docs/docs/models/granite.md diff --git a/docs/docs/models/granite.md b/docs/docs/models/granite.md new file mode 100644 index 00000000..024c380c --- /dev/null +++ b/docs/docs/models/granite.md @@ -0,0 +1,29 @@ +--- +layout: docs +title: Granite +nav_order: 14 +parent: Models +--- + +## 🧩 Model Card: [ibm-granite/granite-4.2-3b](https://huggingface.co/ibm-granite/granite-4.2-3b) + +- **Type:** Text-to-Text +- **Think:** Yes +- **Tool Calling Support:** Yes +- **Base Model:** [ibm-granite/granite-4.2-3b](https://huggingface.co/ibm-granite/granite-4.2-3b) +- **Quantization:** Q4_1 +- **Max Context Length:** 128k tokens +- **Default Context Length:** 8k tokens ([change default](https://fastflowlm.com/docs/instructions/cli/#-change-default-context-length-max)) +- **[Set Context Length at Launch](https://fastflowlm.com/docs/instructions/cli/#-set-context-length-at-launch)** + +▶️ Run with FastFlowLM in PowerShell: + +```shell +flm run granite:3b +``` + +Granite 4.2 is a reasoning model: the chat template opens a `` block in +the generation prompt, so the model emits a reasoning trace, closes it with +``, and then answers. + +--- diff --git a/docs/docs/models/index.md b/docs/docs/models/index.md index ea9613c7..fedb3b5e 100644 --- a/docs/docs/models/index.md +++ b/docs/docs/models/index.md @@ -27,4 +27,5 @@ has_children: true - [Nanbeige](nanbeige/) - [Whisper](whisper/) - [EmbeddingGemma](embeddinggemma/) -- [SmolVLA](smolvla/) \ No newline at end of file +- [SmolVLA](smolvla/) +- [Granite](granite/) \ No newline at end of file From 28e22070d9bece9301ae8e1a20ba246c784b003d Mon Sep 17 00:00:00 2001 From: Vegard Berget Date: Wed, 2 Sep 2026 20:49:57 +0200 Subject: [PATCH 3/3] feat: register granite:3b in model_list.json and model_info.json Weights: https://huggingface.co/vegahyo/Granite-4.2-3B-NPU2 -- a q4nx conversion of IBM's Apache-2.0 granite-4.2-3b, redistributed under the same licence with attribution. Happy to move it under the FastFlowLM org if you would rather host it. model_info.json is the hub's own tree listing for that repo, fetched from the file_url the entry records, so the two cannot disagree. Two things worth a look: * `chat_template.jinja` is in the `files` list. This model has no chat_template in its tokenizer_config, so without the file automodel.cpp reports "No template file found and no chat_template in tokenizer_config.json" and the model will not start. Five shipped entries already list it for the same reason. * `flm_min_version` is 1.0.4 on the assumption that this lands in the next release; adjust if it lands elsewhere. On an older binary the entry is correctly flagged rather than silently unusable. Verified against a local build: `flm list` shows the entry, every file passes the hash check against the published repo, and `flm serve granite:3b` answers correctly. --- src/model_info.json | 50 +++++++++++++++++++++++++++++++++++++++++++++ src/model_list.json | 30 +++++++++++++++++++++++++++ 2 files changed, 80 insertions(+) diff --git a/src/model_info.json b/src/model_info.json index c47e0fb1..b6b66afb 100644 --- a/src/model_info.json +++ b/src/model_info.json @@ -1,4 +1,54 @@ { + "granite:3b": [ + { + "type": "file", + "oid": "2ace5b8c3964855369daf18f16907dfdba648708", + "size": 1566, + "path": ".gitattributes" + }, + { + "type": "file", + "oid": "cd7c1a742f6729a928449bf587514cdfa21873b8", + "size": 3380, + "path": "README.md" + }, + { + "type": "file", + "oid": "f56b1268cb3addc51811411951ff1dc7fd1beacf", + "size": 9197, + "path": "chat_template.jinja" + }, + { + "type": "file", + "oid": "5ed1abc98df2487bfaa76956eb3a4312de3ee24f", + "size": 1043, + "path": "config.json" + }, + { + "type": "file", + "oid": "459034424fdeb26089a5cf634b35ab4ccf3943e9", + "size": 2640901312, + "lfs": { + "oid": "4900c9c2993a4285d7fe4f49e12d59eee3b2ef775d249d4dd27b27af7a201dea", + "size": 2640901312, + "pointerSize": 135 + }, + "xetHash": "bf1151be3bd02f9847d6a1436fe361c881f594033a0b5647af19803759140771", + "path": "model.q4nx" + }, + { + "type": "file", + "oid": "d7e1714703eb97dcef3435aa50eb1de1cf241d62", + "size": 7153421, + "path": "tokenizer.json" + }, + { + "type": "file", + "oid": "23b026c635949bf5941e4884448dccf8d465e283", + "size": 18533, + "path": "tokenizer_config.json" + } + ], "nanbeige4.1:3b": [ { "type": "file", diff --git a/src/model_list.json b/src/model_list.json index 2142a7c2..b3675bca 100644 --- a/src/model_list.json +++ b/src/model_list.json @@ -1,6 +1,36 @@ { "model_path": "models", "models": { + "granite": { + "3b": { + "name": "Granite-4.2-3B-NPU2", + "url": "https://huggingface.co/vegahyo/Granite-4.2-3B-NPU2", + "file_url": "https://huggingface.co/api/models/vegahyo/Granite-4.2-3B-NPU2/tree/main", + "ms_url": "", + "modified_at": "2026-09-02T00:00:00Z", + "size": 2648083506, + "flm_min_version": "1.0.5", + "default_context_length": 8192, + "max_prefill_len": 4096, + "files": [ + "config.json", + "model.q4nx", + "tokenizer.json", + "tokenizer_config.json", + "chat_template.jinja" + ], + "details": { + "family": "granite", + "think": true, + "parameter_size": "3B", + "quantization_level": "Q4_1" + }, + "label": [ + "reasoning" + ], + "footprint": 2.6 + } + }, "nanbeige4.1": { "3b": { "name": "Nanbeige4.1-3B-NPU2",