Repository navigation
Feature/qwen iree kv export - #457
Merged
Merged
Conversation
…models (Qwen) Part of #411 (Qwen export harness / compiled IREE path). iree_kv_jni.c and IreeKvSpec handled exactly one model shape (FunctionGemma: nKvHeads == 1, sliding+global layer split, two RoPE bases, three independent parameter archives). Generalizes both axes a real Qwen cartridge needs without touching FunctionGemma's behaviour: - nKvHeads may now be any value with nHeads % nKvHeads == 0 (grouped-query attention -- Qwen2.5-0.5B nKvHeads=2, Qwen3-0.6B nKvHeads=8). The [1, nKV, rows, headDim] cache layout is nKV-major, so a row-range view is a single zero-copy subspan only when nKV == 1; sub_rows now gathers the nKV > 1 case with nKV device-to-device copies into a fresh buffer. This only runs after a prefill/chunk call (padding removal), never per decode step. The attention mask stays the untouched rank-4 [1, nHeads, C, past+C] for every model -- verified against SKaiNET's AttentionOperationsConverter, which reshapes Q into the GQA group form and broadcasts a rank-4 mask onto it internally via dims=[0,1,3,4], so the native side never needs a grouped mask shape. - globalLayerPeriod == 1 (every layer "global": no sliding/global split at all, the qwen-kv-v1 shape) means the traced graph carries no sliding-side RoPE/mask inputs; nativeChunk/nativeStep now gate building and counting those tensors on hasSliding = period > 1 instead of unconditionally computing and discarding them. - nativeCreate opens one shared IREE session (graph_alias, ref-counted via iree_runtime_session_retain) when the same vmfb+irpa path is passed for all three graphs -- the qwen-kv-v1 contract merges prefill/chunk/with-past into one module sharing one parameter archive; three independent sessions against the same ~1-1.2 GB bf16 archive would not fit a 32-bit process. Three distinct paths (FunctionGemma) keep the original one-session-per-graph behaviour. - locate_embedding now mmaps only the page-aligned range covering the embedding table instead of the whole archive file -- immaterial for FunctionGemma's per-graph ~0.5 GB archives, real for the one shared ~1-1.2 GB Qwen archive. - jint_field/jfloat_field clear the pending JNI exception GetFieldID leaves on a miss instead of leaking it into the next JNI call. IreeKvSpec needs no new field: nKvHeads > 1 and globalLayerPeriod == 1 with the existing fields already express both models. Added qwen25_05bInstruct()/qwen3_06b() factories (architecture facts read off the two GGUFs, tvv workspace NLU-QWEN-TRACKING.md Q0.2) and documented the two contracts (functiongemma-kv-v1 / qwen-kv-v1) sharing this class. Verified: cross-compiled for armeabi-v7a (--vulkan) against skainet/iree-android:3.11.0 -- clean build, 278/278 objects, no warnings on iree_kv_jni.c, libskainet_iree_kv.so linked and copied into jniLibs. IreeKvSpecTest (6 cases: both Qwen factories, nHeads%nKvHeads, fromManifest round-trip incl. a qwen-kv-v1-shaped manifest, chunk override, missing-field fallback) plus the pre-existing IreeTaskTopologyTest (3 cases) both green: ./gradlew :llm-runtime:iree-android:testDebugUnitTest -- 9/9 pass. Not yet built: arm64-v8a .so (armeabi-v7a is the box's actual ABI; arm64 stays on the old FunctionGemma-only binary until a Qwen arm64 target is needed). Device bring-up (does the GQA row-gather and the shared session actually work end to end on Vulkan valhall4) is SKaiNET-transformers#411's Q1.7, blocked on the export harness producing a real vmfb+irpa.
…h (Qwen) Part of #411 (Qwen export harness / compiled IREE path). GemmaModel carries the only existing KV-cache with-past forwards (forwardPrefillAt / forwardPrefillWithPast / forwardWithPast), hand-written for Gemma's specific shape: sandwich post-norms, PLE, sliding+global layer split with two RoPE bases, and expandKV before SDPA. Qwen needs the same three-graph contract over a plain decoderTransformerNetwork module (Llama-shaped: one RMSNorm/MHA/Residual/RMSNorm/SwiGLU/Residual block, one RoPE base, every layer causal full attention) instead of copying and stripping down GemmaModel's version. DecoderKvModel resolves its block/embedding/norm/lm-head references from any module decoderTransformerNetwork() builds (EmbeddingAdapter + HybridTransformerBlock x N + RMSNormalization + VoidDense at the top level; each block's own 2 RMSNorm + MHA + SwiGLU by filterIsInstance, mirroring HybridTransformerBlock's own reference resolution) and mirrors GemmaModel's with-past math with two real differences: - GQA-native SDPA, no expandKV: K/V stay at their native nKVHeads and go straight into ops.scaledDotProductAttention. This isn't a simplification for Qwen's sake -- it's already how MultiHeadAttention.attentionImpl itself calls SDPA (the "Grouped-query attention is native to SDPA (SKEEP-005 phase 2)" comment at its own call site), and how the StableHLO AttentionOperationsConverter lowers it (Q reshaped to [b,nKV,nRep,Sq,hd], a rank-4 mask broadcast onto it via dims=[0,1,3,4]) -- expandKV first would only add dead narrow+concat nodes to both the eager path and the export trace. - Adds the Q/K/V/O projection bias when MultiHeadAttention.bias is set. Qwen2.5 (arch qwen2) GGUFs carry real attention biases; GemmaModel never needed this since Gemma has none. One RoPE base / one mask per call throughout (no sliding-vs-global branching -- decoderTransformerNetwork never sets MultiHeadAttention.slidingWindow), no PLE, no final-logit softcapping, no embedding scale (Gemma-only: embed_scale = hidden_size**0.5). Verified: DecoderKvModelTest builds a tiny GQA network (dim=8, nHeads=4, nKVHeads=2, 2 layers, real random weights via decoderTransformerNetwork + WeightMapper -- no dependency on the qwen module, which sits one layer above llm-core) and checks forwardPrefillAt + forwardWithPast reproduce OptimizedLLMRuntime's ordinary eager forward EXACTLY (within 1e-4), token by token, across 4 prompt tokens plus 2 further decode steps, including that the returned K/V cache stays at nKVHeads (never expands) and grows by exactly one row per step. Caught one real bug along the way -- not in DecoderKvModel but in the test itself: sharing one model instance between the eager runtime and DecoderKvModel corrupts the eager side, because forwardPrefillAt's stripCaches() nulls every block's mha.kvCache (needed for its own stateless prefill, same as GemmaModel's stripCaches) -- fixed by giving each path its own model instance built from the same weights. ./gradlew :llm-core:jvmTest -- 21/21 pass (2 new). The chunk path (forwardPrefillWithPast) needs a host-built RoPE-range table and causal mask that only the export harness (#411's next step) and the native iree_kv_jni.c runtime build today; it gets its correctness gate from the export dump test and vmfb parity test against a real compiled graph instead of a unit test here.
…commonMain) Part of #411. Pure-Kotlin, no I/O -- the arg/output order and manifest.json emission the export harness (next commit) and iree_kv_jni.c's native runtime must agree on, following the FunctionGemmaContract pattern with two real adaptations: - QwenKvArch derives the architecture facts (nLayers, headDim, nKvHeads, nHeads, GQA, QK-norm, attention bias, RoPE base) instead of hardcoding one checkpoint the way FunctionGemmaSpec does -- one QwenKvArch instance describes whichever of the two Qwen models (Qwen2.5-0.5B-Instruct, arch qwen2, GQA 14/2, bias, no QK-norm; Qwen3-0.6B, arch qwen3, GQA 16/8, QK-norm, no bias) the harness loaded. Factories for both, matching IreeKvSpec's own qwen25_05bInstruct()/qwen3_06b() numbers in llm-runtime/iree-android. - Every arg/output list is written out explicitly against iree_kv_jni.c's actual ins[]/ outs[] assembly (nativePrefill/nativeChunk/nativeStep, as generalized in the prior commit) rather than copied from FunctionGemmaContract's -- that contract predates the host-gather addendum and omits the embedding tensor those graphs actually take (see its own docstring); this one lists `emb` explicitly so the harness's trace and the native runtime's calls check against the same source of truth. No sliding/global split (every layer "global", QwenKvArch has one RoPE base, not two) means the arg lists are shorter than FunctionGemma's: cos/sin/mask introduced once before/after the first layer, not per layer-type. manifestJson always reports slidingWindow=0, globalLayerPeriod=1 -- what makes IreeKvSession build no sliding-side tensors at all -- and carries every key IreeKvSpec.fromManifest reads, plus qkNorm/attnBias/eos/eot/ropeBase/embeddingKey and the merged-module contract's parameters.shared archive name. Verified: QwenKvContractTest (commonTest, no checkpoint) checks every arg/output list's length and order against the exact native-side formula for both models (e.g. prefillWithPastArgs must be [tokens, emb, cos, sin, l0.k, l0.v, mask, l1.k, l1.v, ..., select], length 6+2*nLayers, matching nativeChunk's nIn with hasSliding=false), that the mask appears exactly once, that manifestJson carries every IreeKvSpec.fromManifest key and reflects each model's qkNorm/attnBias correctly, and that QwenKvArch rejects nHeads % nKvHeads != 0. ./gradlew :llm-inference:qwen:jvmTest -- 8/8 new, whole-module suite green.
…ort (jvmMain) Closes the QwenExportHarness half of #411. Three graphs, following FunctionGemmaExportHarness's proven structure over DecoderKvModel instead of a hand-written model class: - exportPrefillAt (qwen_prefill_at): the catalog-prefix prefill, LM head on one position via a one-hot select row. - exportPrefillWithPast (qwen_prefill_with_past): a fixed chunk against the incoming cache. - exportWithPast (qwen_with_past): one autoregressive step, true-dynamic past (Dim.DYNAMIC). QwenKvArch is derived from the loaded GGUF's metadata and tensor presence (hasQkNorm/ hasAttnBias, the same structural scan QwenNetworkLoader.applyWeightsToNetwork already uses) instead of hardcoded per model, so one harness serves both Qwen2.5-0.5B-Instruct and Qwen3-0.6B. writeSafetensors/writeHandle/rewriteGlobalsToBf16 are copied verbatim from FunctionGemmaExportHarness (also duplicated in SmolLm2ExportHarness) -- established pattern in this codebase, each export module owns its own copy of this generic machinery. Scope, both decided deliberately and verified by tracing real checkpoints rather than assumed: 1. Per-graph archives, not the merged single-module/shared-archive design qwen-kv-v1 calls for. FunctionGemma's per-trace parameter numbering means merging needs the traces combined into one module with stable, shared keys first -- real further work, tracked as a follow-up. Until it lands, IreeKvSession.nativeCreate's non-shared path (three independent sessions, already built and tested in the prior commit) is what a Qwen cartridge uses, at the same three-archive memory cost FunctionGemma already pays. 2. The harness's raw MLIR is pre-host-gather (an in-graph stablehlo.gather against the embedding table, not a separate `emb` function argument) -- exactly like FunctionGemmaExportHarness's own output. QwenKvContract's arg lists (which DO list `emb`, matching iree_kv_jni.c's native runtime) describe the contract AFTER the blueprint plugin's host-gather rewrite, a later pipeline stage this harness doesn't perform, same as FunctionGemma's. Verified by tracing: qwen_with_past for Qwen3-0.6B is exactly `3 + 2*nLayers` args (token, cos, sin, per-layer K/V), qwen_prefill_at is `tokens, select`, qwen_prefill_with_past is `3 + 2*nLayers + mask + select`. Known gap found by tracing Qwen2.5-0.5B-Instruct (attnBias=true) for real: the Q/K/V/O projection bias tensors do not externalize as util.global parameters the way weight matrices do -- they leak into the function signature as genuine runtime arguments instead (verified: one extra tensor<{dim}xf32> arg per layer, never backed by a util.global.load, absent from the safetensors). Qwen3-0.6B (attnBias=false) is unaffected. This is a SKaiNET core tracer/converter gap (ModuleParameter.BiasParameter isn't recognized as an externalizable constant the way WeightParameter is) -- no existing harness in this codebase has ever traced a bias-bearing attention layer before this one, since Gemma/ SmolLM2/Gemma3n have none. Not fixed here; tracked in NLU-QWEN-TRACKING.md. Qwen2.5-0.5B export is blocked on it, Qwen3-0.6B is not. Verified end to end against the real checkpoints (not synthetic fixtures): - Qwen2.5-0.5B-Instruct qwen_with_past: traced, compiled the graph, safetensors written (1.26 GB bf16, 291 externals) -- surfaced the bias gap above. - Qwen3-0.6B: all three graphs traced successfully (7-19s each on this machine), real safetensors written (291-310 externals per graph, ~1.2 GB bf16 each, tied embeddings detected and logged), MLIR inspected directly (arg/result counts match the raw-trace formula exactly, dynamic KV-cache dims present as `x?x` where expected, no leaked args). - QwenExportDumpTest (new, env-gated on QWEN3_GGUF, real checkpoint): runs exportAll end to end and asserts the above -- all green. QwenKvContractTest, QwenDslPipelineTest, QwenConfigParserTest, QwenDslQuantizedTest all still green. ./gradlew :llm-inference:qwen:jvmTest -PqwenTestMaxHeap=24g -- 18/18 pass, 0 skipped (QWEN3_GGUF pointed at the real checkpoint). build.gradle.kts: new jvmMain source set (compile.hlo/compile.dag/backend.cpu, mirroring :llm-inference:functiongemma's shape), exportQwen Gradle task (QWEN_GGUF/QWEN_OUT_DIR/ QWEN_GRAPH/QWEN_SEQ/QWEN_CHUNK/QWEN_DTYPE env contract), bumped test/JavaExec heap to match FunctionGemma's real-checkpoint numbers scaled for Qwen's larger vocab (151936 vs 262144, but two model sizes instead of one -- 24g direct memory is comfortably safe on this machine, verified by actually running the export at that setting).
…template Closes #409. QwenChatTemplate is faithful to Qwen3's official chat_template but wrong for Qwen2.5 in two ways the host harness found (tvv/nlu-llm-harness, P0.3): no default persona when the caller supplies no system message (Qwen2.5's official template still opens with "You are Qwen, created by Alibaba Cloud. You are a helpful assistant."; Qwen3's dropped it), and QwenChatTemplate's enableThinking=false mode pre-fills an empty <think></think> block to suppress Qwen3's native thinking -- exactly the wrong move for Qwen2.5, which predates the thinking-mode convention entirely and never saw a <think> token, empty or not, in training. Measured impact on the host: 0/8 and 1/8 golden-8 under generic ChatML / QwenChatTemplate vs 4-5/8 with the correct template. Reuses QwenChatTemplate's proven multi-turn logic unchanged (tool listing one JSON object per line, consecutive tool results merged into one user turn, assistant tool calls replayed from the structured ChatMessage.toolCalls) rather than the narrower single-turn version the harness's own Qwen25ToolTemplate.kt got away with -- that harness only ever scores one utterance -> one call, so it never needed multi-turn correctness. parseThinkingBlocks/stripThinking are no-ops: there is nothing to find or strip. Verified against the REAL official template, not memory or the harness's port: fetched Qwen/Qwen2.5-0.5B-Instruct's tokenizer_config.json from huggingface.co directly and rendered its chat_template with a real Jinja2 engine for four cases (no system message, caller system message, no tools, and a full multi-turn tool-call-then-response exchange). This caught a real discrepancy along the way -- the GGUF's embedded copy of the same template has a llama.cpp-conversion brace-escaping artifact (`{{"name": ...}}` doubled braces) that isn't in the authoritative HF source (`{"name": ...}`, single braces, confirmed by diffing both); the harness's own port already had this right, this class matches it. The one deliberate, documented departure from literal HF byte output: tool-call JSON uses kotlinx.serialization's default compact separators (`{"k":"v"}`, verified against the library source: prettyPrint=false, no spaces written) instead of Python json.dumps's spaced default (`{"k": "v"}`) -- matching QwenChatTemplate's own established, already-shipped convention; semantically identical, never affects what a model reads. Qwen25ChatTemplateTest pins all four rendered golden strings byte-for-byte (plus no-generation-prompt and no-thinking-mode checks, 6 cases total). ./gradlew :llm-agent:jvmTest -- 6/6 new, whole-module suite green.
Summarizes the four prior commits (native iree-android generalization, DecoderKvModel, QwenKvArch/QwenKvContract/QwenExportHarness, Qwen25ChatTemplate) plus a real finding from compiling the traced Qwen3-0.6B qwen_with_past graph against skainet/iree-compiler:3.11.0 today: llvm-cpu (host and arm32, the box's CPU flavour) compiles cleanly; vulkan-spirv valhall4 crashes the compiler outright (internal stack dump, no diagnostic) on the GQA-native attention lowering. CPU is viable today; Vulkan -- the box's primary, faster target -- is not, until an IREE fix lands or a DecoderKvModel expandKV fallback is added for that target specifically. Recorded here and in NLU-QWEN-TRACKING.md (OQ6).
…heses tested and refuted Follow-up to dd286f4. Actually reproduced the compile with --mlir-disable-threading --mlir-pass-pipeline-crash-reproducer and tested both candidate causes the first entry guessed at: stripping the in-graph argmax tail (still crashes, now on an ordinary FFN matmul) and an isolated/chained matmul of the flagged shape (compiles fine standalone, 28x). Neither reproduces the crash alone -- it needs the full graph's context. Corrects the record rather than leaving a wrong root-cause guess in the changelog.
…the API dumps The commits that added DecoderKvModel (llm-core) and Qwen25ChatTemplate (llm-agent) on this branch never ran apiDump, so jvmApiCheck failed for both modules. Additions only, no existing declaration changes. Tests: ./gradlew :llm-core:apiCheck :llm-agent:apiCheck green.
iree-compile 3.11.0 crashed compiling the traced Qwen3-0.6B graphs for vulkan-spirv valhall4 (SPIRVInitialVectorLoweringPass: arith.constant vector<512xf32> not legalized) while llvm-cpu compiled them. Bisected with real exports, fp32 and bf16, graph by graph. bf16 is not a cause. Three independent causes, each fixed by a post-emit rewrite in QwenExportHarness: 1. argMax over the 151936 vocab. The SPIR-V backend lowers the fused max/compare/iota/select/min reduction only when its extent is a multiple of 2048: 151936 and 152064 fail; 153600, 155648, 163840, 196608 and 262144 compile. FunctionGemma's 262144 vocab is 128 x 2048, which is why it never hit this. rewriteArgMaxPadded pads the logits row with -3e38 (finite, no NaN from the EQ compare) to QwenKvContract.paddedArgMaxExtent(vocab) = 153600 before the argMax. Real logits and the returned id are unchanged; a padded column can never be the maximum. 2. The in-graph token-embedding gather. On the decode step the row lookup fuses into the first residual add's vecmat dispatch and fails there; on the fixed-length prefill the gather fails on a vector.step vector<1024xindex>. rewriteHostGather is the blueprint plugin's host-gather rewrite applied to every function of the module. Opt-in via hostGather / QWEN_HOST_GATHER=1 so the raw export stays the plugin's HostGatherTask input. 3. The chunk graph's mask under grouped-query attention. SKaiNET 0.56.0's SDPA converter emits broadcast_in_dim %mask, dims = [0, 1, 3, 4] onto the grouped scores [1, nKV, nRep, C, ?]. With the dynamic key length that op is invalid on every backend, so qwen_prefill_with_past compiled nowhere, not even llvm-cpu. IREE 3.11 legalizes neither dynamic_reshape of a per-head mask nor dynamic_broadcast_in_dim without expansion hints; it does legalize a head-shared [1, 1, C, ?] mask broadcast with known_expanding_dimensions/known_nonexpanding_dimensions set, which is what rewriteGqaMaskBroadcast emits. The chunk mask rows never depended on the head, so the contract now exports it head-shared (QwenKvContract.MASK_HEADS = 1, manifest "maskHeads": 1) and iree_kv_jni.c builds it with maskHeads rows instead of nHeads. IreeKvSpec gains an optional maskHeads (0 = nHeads, the default, FunctionGemma unchanged; the Qwen factories set 1; fromManifest reads it); the 11-argument constructor stays through @jvmoverloads. The converter itself should emit this; the rewrite stands in until a SKaiNET release does. Result on the real Qwen3-0.6B export (bf16, host-gather, seq 1024, chunk 32): qwen_prefill_at, qwen_prefill_with_past and qwen_with_past all compile for vulkan-spirv valhall4 and llvm-cpu arm32 with skainet/iree-compiler:3.11.0. The qwen API dump now records the export harness, contract and CLI, which earlier commits on this branch had not dumped. Tests: QwenExportRewriteTest 10/10 (new), QwenKvContractTest 8/8, IreeKvSpecTest 7/7, whole :llm-inference:qwen:jvmTest green with QWEN3_GGUF set (QwenExportDumpTest against the real checkpoint), :llm-inference:qwen:apiCheck green. The armeabi-v7a libskainet_iree_kv.so is not rebuilt in this commit.
Rebuilds armeabi-v7a and arm64-v8a (--vulkan) from iree_kv_jni.c as of fbdb9e9, so the native session builds the one-head chunk mask the qwen-kv-v1 chunk graph now takes. The previous binaries predate maskHeads and would pass a [1, nHeads, C, past+C] mask to a graph expecting [1, 1, C, past+C]. FunctionGemma is unaffected: maskHeads defaults to nHeads. Built with skainet/iree-android:3.11.0 (IREE 3.11.0 runtime, NDK r27c, API 29). On an Apple-silicon host that image must be built for linux/amd64: the NDK ships x86-64 host tools, and an arm64 image fails at configure time. A build cache left by an arm64 image also has to be cleared, since its host flatcc tool is aarch64. Verified: both libraries build with no warnings in iree_kv_jni.c and contain the maskHeads field lookup.
…k fix DecoderKvModel's class doc and ChunkContext doc, and iree_kv_jni.c's header, still said the attention converter broadcasts a per-head [b, H, Sq, Sk] mask onto the grouped scores with dims = [0, 1, 3, 4] and that the mask is always [1, nHeads, C, past+C]. That broadcast was the bug reported as SKaiNET#1302 (fixed there in PR #1303). The docs now say the mask is [1, maskH, C, past+C]: nHeads for functiongemma-kv-v1, which expands K/V so the mask equals the scores shape, and 1 for qwen-kv-v1, the only form IREE 3.11 lowers under GQA with a dynamic key length. Comments only; the native libraries are unchanged.
…e/qwen-iree-kv-export # Conflicts: # llm-runtime/iree-android/src/test/kotlin/sk/ainet/transformers/iree/android/IreeKvSpecTest.kt
…self rewriteGqaMaskBroadcast turned SKaiNET 0.56.0's invalid static broadcast_in_dim of the chunk mask into a hinted dynamic_broadcast_in_dim. SKaiNET 0.57.0 fixes that in the attention converter (SKaiNET#1302), so the rewrite's pattern no longer occurs and it is removed with its tests. The head-shared mask (QwenKvContract.MASK_HEADS = 1, maskHeads in the manifest and the native session) stays: a per-head mask with a dynamic key length needs dynamic_reshape, which IREE 3.11 does not lower. The argMax padding and host-gather rewrites stay; they work around IREE limits, not an engine bug. Docs in QwenExportHarness, QwenKvContract, DecoderKvModel and the CHANGELOG now say so. Verified on the real Qwen3-0.6B export (bf16, QWEN_HOST_GATHER=1, seq 1024, chunk 32) on engine 0.57.0: the chunk graph's mask reaches every layer as the engine-emitted "stablehlo.dynamic_broadcast_in_dim"(... known_expanding_dimensions = [1] ...) (28 of 28 layers), and qwen_prefill_at, qwen_prefill_with_past and qwen_with_past compile for vulkan-spirv valhall4 and llvm-cpu arm32 with skainet/iree-compiler:3.11.0, byte-for-byte the same vmfb sizes as with the rewrite. Tests: :llm-inference:qwen:jvmTest 24/24 with QWEN3_GGUF set (incl. QwenExportDumpTest), :llm-core:jvmTest 133/133, apiCheck green.
Replaces references to a private evaluation harness and its tracking notes in the Qwen export harness, the export CLI, Qwen25ChatTemplate and the CHANGELOG with the facts they stood for (an 8-utterance zero-shot tool-calling set; a 1024-token prefill that fits a ~1000-token tool-catalog prefix; device deployment and the native runtime). Qwen25ChatTemplateTest now uses a generic weather-forecast tool instead of a domain-specific one. Verified: the four golden strings in Qwen25ChatTemplateTest were re-rendered from Qwen/Qwen2.5-0.5B-Instruct's official chat_template with Jinja2 (compact tojson separators, the test's documented adjustment) for the new tool and messages, and match byte-for-byte; :llm-agent:jvmTest --tests Qwen25ChatTemplateTest 6/6.
…kens) QwenVmfbParityTest exports the three qwen-kv-v1 graphs exactly as they ship (bf16 weights, host-gather emb argument, padded argMax, head-shared chunk mask), converts each archive to .irpa, compiles for llvm-cpu host with skainet/iree-compiler:3.11.0 and drives them through iree-run-module the way iree_kv_jni.c does: qwen_prefill_at over the first 3 prompt tokens (padded to 16, K/V sliced to the real rows), one qwen_prefill_with_past call over the remaining 3 tokens (padded to the 32-token chunk, split-half RoPE tables and the [1,1,C,past+C] mask built with the native session's formulas), then qwen_with_past steps. Embedding rows come from the exported archive, bf16-widened like the device reads them. The oracle is a new fixture for the checkpoint that actually ships, qwen3-06b/golden-greedy-06b.txt: mainline llama.cpp (0.3.0 build 10621), prompt "The first ten prime numbers are", 32 greedy tokens, smallest top-1/top-2 gap 2.35 nats, so a mismatch is a defect rather than cross-implementation noise. The fixture records the build, the checkpoint digest and the exact commands. Result: all 32 tokens equal the oracle, including the chunk call's first token. The run writes an evidence record; this commit keeps it as llm-inference/qwen/validation/L3-qwen3-06b-host-vmfb.json. The graphs return the argmax token only, so logit drift is not recorded yet; that needs a logits-returning export variant. Gated: skips without QWEN3_06B_GGUF or without docker and the toolchain image. Tests: QwenVmfbParityTest 1/1 against the real checkpoint.
Merged
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
1 participant
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
No description provided.