From 892f8a70b95d73311f0135f8eb9716ec2426bc34 Mon Sep 17 00:00:00 2001 From: Georg Bauer Date: Sun, 4 Oct 2026 14:11:27 +0200 Subject: [PATCH] Apply Qwen vision RoPE per row across table boundaries --- docs/issue108-vision-boundary/README.md | 57 +++++++++ docs/issue108-vision-boundary/attention.jsonl | 8 ++ docs/issue108-vision-boundary/checks.json | 108 ++++++++++++++++++ docs/issue108-vision-boundary/native.jsonl | 56 +++++++++ docs/issue108-vision-boundary/oracle.jsonl | 56 +++++++++ src/engine/metal/qwen_mtplx.rs | 2 + .../metal/qwen_mtplx/issue103_analysis.rs | 80 ++++++++++++- src/engine/metal/qwen_mtplx/issue108.rs | 90 +++++++++++++++ src/engine/metal/qwen_mtplx/qsa_attention.rs | 25 ++-- src/engine/metal/qwen_mtplx/rope.rs | 64 +++++++++++ tools/issue108-reference.py | 55 +++++++++ 11 files changed, 580 insertions(+), 21 deletions(-) create mode 100644 docs/issue108-vision-boundary/README.md create mode 100644 docs/issue108-vision-boundary/attention.jsonl create mode 100644 docs/issue108-vision-boundary/checks.json create mode 100644 docs/issue108-vision-boundary/native.jsonl create mode 100644 docs/issue108-vision-boundary/oracle.jsonl create mode 100644 src/engine/metal/qwen_mtplx/issue108.rs create mode 100644 tools/issue108-reference.py diff --git a/docs/issue108-vision-boundary/README.md b/docs/issue108-vision-boundary/README.md new file mode 100644 index 0000000..01ed778 --- /dev/null +++ b/docs/issue108-vision-boundary/README.md @@ -0,0 +1,57 @@ +# Issue 108: row-wise attention positions + +Qwen attention now implements pinned MTPLX's `_vision_chunk_cos_sin` route: +covered rows use the original table slice and M-RoPE graph; uncovered rows use +plain RoPE at `index + delta`; mixed chunks concatenate those two tables. The +YaRN amplitude applies to both portions. With the shipped scaling of one, wholly +inside/outside calls build the same numerical graphs as before. The text-only +branch and indexer remain unchanged. Production changes are confined to two +private Qwen modules, with no DS4/GLM callers or shared numerical changes. + +This fixes a **latent module mismatch**. Review of `vision::prepare`, owner setup, +model dispatch, suffix prefill and MTP history found no new product trigger: +`prepare` builds a table over the complete rendered prompt; prefill stops within +it, and decoding starts after it. The final post-image text positions are equal +across axes at `index + delta`. A normal MTP window crossing that text boundary +is therefore not a reproduction. No state rebase, prompt mutation or stale-table +restore was added. No new DS4Server feature is involved. + +## Verification + +`tools/issue108-reference.py` runs the pinned MTPLX revision +`9882703f3105363ddc37eca9f97aa09a1d387112` in isolated MLX 0.32.2. It uses the +oracle's actual frequency constructor and chunk helper, plus independent NumPy +per-row checks. Its 56 cases cover small and shipped head geometry (dim256, +24 query heads, rotary64), F32/BF16, scaling 1/1.25, wholly inside/outside, +unequal image-axis straddling, equal axes, absent/empty tables, and a full-prompt +post-image text tail. Native cosine/sine and Q/K outputs match every oracle byte. +The text-tail and equal-axis crossings also match plain shifted RoPE exactly. + +The owned `vision_positions` case in `issue103_vision_position_receipts` now +requires corrected cached keys. Actual Attention calls compare one mixed chunk +with the same input split at the table end: outputs, valid K/V, offsets and the +next call's output/state agree exactly. The unrelated sparse-QSA and attention +Gate diagnostic cases retain their current separate ownership (#107/#109). + +Existing whole-attention fixtures exercise the unchanged text/vision controls, +cache hashes and execution routes in both TF32 modes, with lazy and prepared +inputs. The mandatory Rust formatting, Clippy, bundle and complete ordinary test +gates are recorded in `checks.json`; the adjacent JSONL receipts retain the +owned test results. These are correctness checks, not performance claims. + +## Reproduction + +```sh +local-eval-results/mtplx-reference-env-0.32.2/bin/python \ + tools/issue108-reference.py /path/to/pinned/mtplx /tmp/issue108-oracle +cargo test --all-features --bin ds4-server --no-run +cargo build --bin test-supervisor +export DS4_ISSUE108_ORACLE=/tmp/issue108-oracle +# Use the test executable printed by cargo, not cargo itself in the supervisor. +target/debug/test-supervisor 4096 60 120 --command TEST_EXECUTABLE \ + issue108_attention_chunk_positions --ignored --nocapture --test-threads=1 +target/debug/test-supervisor 4096 60 120 --command TEST_EXECUTABLE \ + issue103_vision_position_receipts --ignored --nocapture --test-threads=1 +# Run mtplx_qsa_attention_module_matches_reference the same way with +# MLX_ENABLE_TF32=1 and then MLX_ENABLE_TF32=0, serially. +``` diff --git a/docs/issue108-vision-boundary/attention.jsonl b/docs/issue108-vision-boundary/attention.jsonl new file mode 100644 index 0000000..920bf8c --- /dev/null +++ b/docs/issue108-vision-boundary/attention.jsonl @@ -0,0 +1,8 @@ +{"case":"vision_positions","start":0,"rows":4,"table_end":4,"delta":-1,"different_rows":[],"actual_keys":[0.1982421875,0.396484375,0.59375,0.79296875,0.98828125,1.1875,1.3828125,1.5859375,-0.478515625,0.671875,0.85546875,0.9453125,1.1484375,1.171875,1.1875,1.2578125,-1.265625,0.515625,0.875,0.96484375,0.326171875,1.28125,1.15625,1.1640625,-1.2890625,0.55859375,0.90234375,0.9765625,0.373046875,1.265625,1.125,1.1171875],"expected_keys":[0.1982421875,0.396484375,0.59375,0.79296875,0.98828125,1.1875,1.3828125,1.5859375,-0.478515625,0.671875,0.85546875,0.9453125,1.1484375,1.171875,1.1875,1.2578125,-1.265625,0.515625,0.875,0.96484375,0.326171875,1.28125,1.15625,1.1640625,-1.2890625,0.55859375,0.90234375,0.9765625,0.373046875,1.265625,1.125,1.1171875]} +{"case":"vision_chunk_next_state","start":0,"rows":4,"offset":5,"exact":true} +{"case":"vision_positions","start":4,"rows":2,"table_end":4,"delta":-1,"different_rows":[],"actual_keys":[-0.3359375,0.02783203125,0.55078125,0.7890625,-0.94921875,1.25,1.3984375,1.5859375,0.3125,0.294921875,0.8203125,0.94140625,-1.203125,1.3203125,1.2109375,1.2578125],"expected_keys":[-0.3359375,0.02783203125,0.55078125,0.7890625,-0.94921875,1.25,1.3984375,1.5859375,0.3125,0.294921875,0.8203125,0.94140625,-1.203125,1.3203125,1.2109375,1.2578125]} +{"case":"vision_chunk_next_state","start":4,"rows":2,"offset":7,"exact":true} +{"case":"vision_positions","start":2,"rows":4,"table_end":4,"delta":-1,"different_rows":[],"actual_keys":[-0.98046875,0.02783203125,0.5390625,0.7890625,-0.2314453125,1.25,1.40625,1.5859375,-1.2265625,0.423828125,0.8203125,0.94140625,0.216796875,1.28125,1.2109375,1.2578125,-0.9609375,0.515625,0.88671875,0.96484375,-0.890625,1.28125,1.140625,1.1640625,0.1962890625,0.4296875,0.90234375,0.9765625,-1.328125,1.3203125,1.125,1.125],"expected_keys":[-0.98046875,0.02783203125,0.5390625,0.7890625,-0.2314453125,1.25,1.40625,1.5859375,-1.2265625,0.423828125,0.8203125,0.94140625,0.216796875,1.28125,1.2109375,1.2578125,-0.9609375,0.515625,0.88671875,0.96484375,-0.890625,1.28125,1.140625,1.1640625,0.1962890625,0.4296875,0.90234375,0.9765625,-1.328125,1.3203125,1.125,1.125]} +{"case":"vision_chunk_next_state","start":2,"rows":4,"offset":7,"exact":true} +{"case":"vision_sparse_qsa","prefix_rows":12,"topk_blocks":1,"ratio":4,"position_delta":0,"text_output":[0.10986328125,-0.1923828125,0.08544921875,0.498046875,-0.0791015625,0.16796875,0.31640625,-0.03173828125],"vision_output":[0.10986328125,-0.1923828125,0.08544921875,0.498046875,-0.0791015625,0.16796875,0.31640625,-0.03173828125]} +{"case":"active_attention_gate","gate":-6.84375,"actual":[0.0003185272216796875,-0.00074005126953125,0.00021266937255859375,0.00095367431640625,-0.00021266937255859375,0.0004253387451171875,0.000637054443359375,-0.00010633468627929688],"expected_verify":[0.0003204345703125,-0.0007476806640625,0.000213623046875,0.0009613037109375,-0.000213623046875,0.00042724609375,0.000640869140625,-0.0001068115234375]} diff --git a/docs/issue108-vision-boundary/checks.json b/docs/issue108-vision-boundary/checks.json new file mode 100644 index 0000000..82ddf74 --- /dev/null +++ b/docs/issue108-vision-boundary/checks.json @@ -0,0 +1,108 @@ +[ + { + "check": "fmt", + "passed": true + }, + { + "check": "clippy", + "passed": true + }, + { + "check": "bundle", + "passed": true + }, + { + "check": "tests", + "passed": true, + "results": [ + "test result: ok. 308 passed; 0 failed; 219 ignored; 0 measured; 0 filtered out; finished in 5.51s", + "test result: ok. 4 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out; finished in 0.00s", + "test result: ok. 4 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out; finished in 2.28s" + ] + }, + { + "check": "native", + "passed": true, + "results": [ + "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 526 filtered out; finished in 0.27s" + ], + "resources": { + "event": "test_resource_summary", + "elapsed_ms": 316, + "peak": { + "cpu_time_ns": 172161583, + "resident_bytes": 80969728, + "physical_bytes": 159974240, + "peak_physical_bytes": 160072544, + "disk_read_bytes": 0, + "disk_write_bytes": 0 + }, + "error": null + } + }, + { + "check": "attention", + "passed": true, + "results": [ + "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 526 filtered out; finished in 0.14s" + ], + "resources": { + "event": "test_resource_summary", + "elapsed_ms": 210, + "peak": { + "cpu_time_ns": 86198750, + "resident_bytes": 81559552, + "physical_bytes": 176505672, + "peak_physical_bytes": 176505672, + "disk_read_bytes": 0, + "disk_write_bytes": 0 + }, + "error": null + } + }, + { + "check": "module-1", + "passed": true, + "results": [ + "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 526 filtered out; finished in 14.78s" + ], + "resources": { + "event": "test_resource_summary", + "elapsed_ms": 17201, + "peak": { + "cpu_time_ns": 13461448125, + "resident_bytes": 529022976, + "physical_bytes": 1574700328, + "peak_physical_bytes": 1602028864, + "disk_read_bytes": 8667136, + "disk_write_bytes": 393216 + }, + "error": null + } + }, + { + "check": "module-0", + "passed": true, + "results": [ + "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 526 filtered out; finished in 15.32s" + ], + "resources": { + "event": "test_resource_summary", + "elapsed_ms": 15380, + "peak": { + "cpu_time_ns": 14550850750, + "resident_bytes": 551796736, + "physical_bytes": 1599587768, + "peak_physical_bytes": 1607976376, + "disk_read_bytes": 0, + "disk_write_bytes": 131072 + }, + "error": null + } + }, + { + "check": "upstream", + "passed": true, + "tests": 10 + } +] diff --git a/docs/issue108-vision-boundary/native.jsonl b/docs/issue108-vision-boundary/native.jsonl new file mode 100644 index 0000000..606dc51 --- /dev/null +++ b/docs/issue108-vision-boundary/native.jsonl @@ -0,0 +1,56 @@ +{"case":"attention_chunk","key":"8-inside-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-inside-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-inside-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-inside-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-outside-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-outside-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-outside-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-outside-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-straddle-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-straddle-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-straddle-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-straddle-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-full_text_tail-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-full_text_tail-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-full_text_tail-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-full_text_tail-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-equal-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-equal-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-equal-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-equal-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-none-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-none-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-none-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-none-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-empty-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-empty-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-empty-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"8-empty-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-inside-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-inside-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-inside-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-inside-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-outside-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-outside-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-outside-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-outside-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-straddle-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-straddle-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-straddle-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-straddle-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-full_text_tail-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-full_text_tail-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-full_text_tail-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-full_text_tail-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-equal-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-equal-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-equal-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-equal-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-none-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-none-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-none-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-none-1.25-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-empty-1.0-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-empty-1.0-bf16","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-empty-1.25-f32","exact_q_k_and_tables":true} +{"case":"attention_chunk","key":"256-empty-1.25-bf16","exact_q_k_and_tables":true} diff --git a/docs/issue108-vision-boundary/oracle.jsonl b/docs/issue108-vision-boundary/oracle.jsonl new file mode 100644 index 0000000..d2426ba --- /dev/null +++ b/docs/issue108-vision-boundary/oracle.jsonl @@ -0,0 +1,56 @@ +{"key":"8-inside-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"4c5ad613aa42e9289b941fda7f83ab7e0d4b7a29292ed4962d6c664fa9bf3c4a","sin_sha256":"8832c4a47d3a73f1b3af4a836d42cf0e886520442cd14babb3beba9fba5fe70b","q_sha256":"37d290b59d0162dac59e0aee718210ad450dbdf3b9fc8f5276101fcef50d897c","k_sha256":"26d3123cd9295182186f3e8c056fc05c77e99d1475ab14fe580fd8f35247d857"} +{"key":"8-inside-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"4c5ad613aa42e9289b941fda7f83ab7e0d4b7a29292ed4962d6c664fa9bf3c4a","sin_sha256":"8832c4a47d3a73f1b3af4a836d42cf0e886520442cd14babb3beba9fba5fe70b","q_sha256":"46f6571ff18f449d64c9db397c2b8b59ecfa80d6bcdd40ec2dea5fd43ddf8d29","k_sha256":"90019757e8b38b805b200c9bf860f86371f0e118f714d211ffdd18ccd548f881"} +{"key":"8-inside-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"b33b483634eb4e50faa2050d6efd606c14fb9dae4df4a37efa90c058764dfc58","sin_sha256":"20bf788c0af8c1c96835d4fdfecfa22548f7461c32431ec8229c1bf6b89fd0f0","q_sha256":"0958a24bd3c04f8f1b85de20767fad155bca39baf38194299dd5452fd47d971d","k_sha256":"8415a43c9d0cece7fff556589f0f60a8fe0aecb2c3e48fc11770102d72f77d76"} +{"key":"8-inside-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"b33b483634eb4e50faa2050d6efd606c14fb9dae4df4a37efa90c058764dfc58","sin_sha256":"20bf788c0af8c1c96835d4fdfecfa22548f7461c32431ec8229c1bf6b89fd0f0","q_sha256":"1f9dde5ec35085dd3a3a91f3e1c7eab105e5f0a442a1cbf38f64e3f2100a2b9f","k_sha256":"d5bbac3e7d547cf7422b4120d6ec12b7ccd277ffb52e9b494397add827e51fc1"} +{"key":"8-outside-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"cfdd4c78aecc2e702a87dd914fe858580f91999e2fb1756c10fbc14e6a694557","k_sha256":"e93d00fbfd1f29403de526fc7fdcf0f73b368a817210b58814df77b7d2955246"} +{"key":"8-outside-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"16da54ce9aef87bfef69928dbd1c87800b62a3c6c7bb3d0e2c751351c366e448","k_sha256":"c97ea7805a64b06ae1d2c5f0c2cae128c6a11159cb4d0c22698e86286e92cdf4"} +{"key":"8-outside-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"c693451f235fda763aed834cafddd4219d1d35e86d2ad9a45bbb07908ae04d70","k_sha256":"e02eb2ffee19d0aa2b8a39ecc1b58584dffbcea3d4ccc8269a1928e829bd2f76"} +{"key":"8-outside-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"f24a2cc645f202052db55ace4a09ea7754858b3397cc3eeb497714edae1118e4","k_sha256":"43b0a12d9d97c39f141539be95f877305e85d8e49080abc04109aedff99e7a35"} +{"key":"8-straddle-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"c2795e059020713438fad524723534bd8ed201bc1faa9d23e6af4059cba3ea23","sin_sha256":"9d88ad2d1b6cbe738f47f37f12b169988e1b89c74fb5621b8dc1a46b57cc6457","q_sha256":"a4976c7740385d1392c1a12b585edf2000f502010bc48e31a92b776181763314","k_sha256":"c6a32b83c7b227089b51fd3168c9e6c0473a68e23b3c25dfff5053b4fd9ccffc"} +{"key":"8-straddle-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"c2795e059020713438fad524723534bd8ed201bc1faa9d23e6af4059cba3ea23","sin_sha256":"9d88ad2d1b6cbe738f47f37f12b169988e1b89c74fb5621b8dc1a46b57cc6457","q_sha256":"303190dd3d5514b4952e48353c7d5aa6efcdc3ffa2b985ae6de6796cad84b69f","k_sha256":"3c26adb50d0181f9c4e34904146460ff299aea078d4343737a78d3c57e5944e5"} +{"key":"8-straddle-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"667522d8ce6fb8b16883df34b235c881b3a90a3eb5edcfcae85984a9225d1ece","sin_sha256":"ba57f1dc77bc8cbe0998c5fdcbae4725c598fb39e5758d1138fefbbdf77988d9","q_sha256":"1d8ce44d7fca02aae1640ce4c7e7054a72e600bc750e7da091f2ffd8b2e4a156","k_sha256":"4c68bf15f527f6530941c2b705207703a6b7fee31446f564b124c19a396c583b"} +{"key":"8-straddle-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"667522d8ce6fb8b16883df34b235c881b3a90a3eb5edcfcae85984a9225d1ece","sin_sha256":"ba57f1dc77bc8cbe0998c5fdcbae4725c598fb39e5758d1138fefbbdf77988d9","q_sha256":"68e71d870c8660acd954873dad34180033e6d683da661cdc6fea5e30cbbabf65","k_sha256":"5b347b699915cf4a76aeab754c4b14afcf560ff2cbb18e33b674bbe4fd0a22c3"} +{"key":"8-full_text_tail-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"1d7db40071c56957b00e95c623d8fe33b00844923603b4f22437290844ce5189","sin_sha256":"2e86a7b75f51570f50969a7eaff682cb1c4c740a7df3047ba7d34b881d336131","q_sha256":"bd0241f4af2045f3ccc497eeedace4c31286fac1b30a7976a73d3381740d0853","k_sha256":"4aa491e9bf45098932cfac394540fe176a3cf31cd449e56b670642930c0c0cd8"} +{"key":"8-full_text_tail-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"1d7db40071c56957b00e95c623d8fe33b00844923603b4f22437290844ce5189","sin_sha256":"2e86a7b75f51570f50969a7eaff682cb1c4c740a7df3047ba7d34b881d336131","q_sha256":"5bb0fbb54b84be503ee8d1dfbec37688597d992559a2a8ffc3159433349066cc","k_sha256":"6c6a5ed10c94b751b48cb7d14837f12bf905c8b7e9a790754c71badf4f1d9335"} +{"key":"8-full_text_tail-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"011662c43e69e3be2807d33b029946094bca520f1a043630b6ae2c18a716dfda","sin_sha256":"79215ac3f560b4e8624eaa8d1baf411e7d5f9a639f1c8f90d0fd807418ca9037","q_sha256":"573d6b46366220fdf9ca71d24b1dc91906644501a6b900abc7478e5fa2dccf0c","k_sha256":"e28a2891a9cf01bd8ef3769477ea45d5363d221439b059c33e2f11f985a119f0"} +{"key":"8-full_text_tail-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"011662c43e69e3be2807d33b029946094bca520f1a043630b6ae2c18a716dfda","sin_sha256":"79215ac3f560b4e8624eaa8d1baf411e7d5f9a639f1c8f90d0fd807418ca9037","q_sha256":"3b3c3cfde7cc07dfa518fcbb5f7bf017cf79eb9e86b2b405209411762aee1d8d","k_sha256":"6df327a507adaae71834850a788a5cacf9608ab1b83f03f2260c636dbf84d9b7"} +{"key":"8-equal-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.0,"dtype":"f32","cos_sha256":"77ac83401bc7565ad83fedcfa8dab1df4e52d0a827c8a1b017b1ef5ec93ad5c1","sin_sha256":"b3dd5ba5e17b6ceb6fdeb51065df1d57d7bf35bdbd5dac4925bb75ccd87458a8","q_sha256":"eec5ad30b4abbf15079d9973083c5a70720e6a0c8aa6875200b9582b45fcd864","k_sha256":"0426043a33440a71b3830a89c34b008bcec51ed961fdee51bc7819d46b1fe478"} +{"key":"8-equal-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.0,"dtype":"bf16","cos_sha256":"77ac83401bc7565ad83fedcfa8dab1df4e52d0a827c8a1b017b1ef5ec93ad5c1","sin_sha256":"b3dd5ba5e17b6ceb6fdeb51065df1d57d7bf35bdbd5dac4925bb75ccd87458a8","q_sha256":"5f277c2a557848f452f484388ac5e46500b9942df921b72254de2bf691745fa5","k_sha256":"e200b0ee24778e290b5bb3a0d5b5dd5bff63b6754da56baa6a2c0911b6873fdd"} +{"key":"8-equal-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.25,"dtype":"f32","cos_sha256":"dd4cd7320fa29d59162a9630e078fa84c4661769091f396e75b38a5d798b1ba4","sin_sha256":"ce60e7dc9cd397d885780551f1b12f9f8cd5b6740d6673324793ecd0516b83d2","q_sha256":"aaf430af221f87419869232fe5bfba69aeaebf5dc9b0b30b9e5e80779a21e7d4","k_sha256":"36225b9a070cbc3537ce2e9b32b1b6d035de57325df1da55874c43d0a937fd68"} +{"key":"8-equal-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.25,"dtype":"bf16","cos_sha256":"dd4cd7320fa29d59162a9630e078fa84c4661769091f396e75b38a5d798b1ba4","sin_sha256":"ce60e7dc9cd397d885780551f1b12f9f8cd5b6740d6673324793ecd0516b83d2","q_sha256":"3f5eca7a6aa94afc626d070bdc69f0ae3a197790bcb75c6cac19093dfa54d9ed","k_sha256":"4bf23203d9db6382c09bd3c9abd85a473b1f557825f84983d73a94243a596b70"} +{"key":"8-none-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"cfdd4c78aecc2e702a87dd914fe858580f91999e2fb1756c10fbc14e6a694557","k_sha256":"e93d00fbfd1f29403de526fc7fdcf0f73b368a817210b58814df77b7d2955246"} +{"key":"8-none-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"16da54ce9aef87bfef69928dbd1c87800b62a3c6c7bb3d0e2c751351c366e448","k_sha256":"c97ea7805a64b06ae1d2c5f0c2cae128c6a11159cb4d0c22698e86286e92cdf4"} +{"key":"8-none-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"c693451f235fda763aed834cafddd4219d1d35e86d2ad9a45bbb07908ae04d70","k_sha256":"e02eb2ffee19d0aa2b8a39ecc1b58584dffbcea3d4ccc8269a1928e829bd2f76"} +{"key":"8-none-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"f24a2cc645f202052db55ace4a09ea7754858b3397cc3eeb497714edae1118e4","k_sha256":"43b0a12d9d97c39f141539be95f877305e85d8e49080abc04109aedff99e7a35"} +{"key":"8-empty-1.0-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"cfdd4c78aecc2e702a87dd914fe858580f91999e2fb1756c10fbc14e6a694557","k_sha256":"e93d00fbfd1f29403de526fc7fdcf0f73b368a817210b58814df77b7d2955246"} +{"key":"8-empty-1.0-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"81ca98b954346d7fabada07182b71708d7d44379f6ccf4f7b1d20592e9b2c5b1","sin_sha256":"f12f35d499314f536706031806a03760888a398360797344ed2aac4b0ca45fab","q_sha256":"16da54ce9aef87bfef69928dbd1c87800b62a3c6c7bb3d0e2c751351c366e448","k_sha256":"c97ea7805a64b06ae1d2c5f0c2cae128c6a11159cb4d0c22698e86286e92cdf4"} +{"key":"8-empty-1.25-f32","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"c693451f235fda763aed834cafddd4219d1d35e86d2ad9a45bbb07908ae04d70","k_sha256":"e02eb2ffee19d0aa2b8a39ecc1b58584dffbcea3d4ccc8269a1928e829bd2f76"} +{"key":"8-empty-1.25-bf16","dim":8,"rotary":8,"sections":[2,1,1],"heads":2,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"ffa186135e2e51b2cdf09e44efe7ba62dc7e0354b5744fd041ca98c1014382bb","sin_sha256":"f423178f62b5714a8ea8aca5d4eeaef1654b422b6b6ba3c869e664b16c8e1580","q_sha256":"f24a2cc645f202052db55ace4a09ea7754858b3397cc3eeb497714edae1118e4","k_sha256":"43b0a12d9d97c39f141539be95f877305e85d8e49080abc04109aedff99e7a35"} +{"key":"256-inside-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"6b8626092624533479bc18ba46f6de54cbd643f10f7e64fb93e8fa4118f6b361","sin_sha256":"140df4ec915b4d7cc25a17eba853438c6c6ed1f444831caff3d78ab8c51c9b8e","q_sha256":"c75ebcf3f39d56a35f6d806ef65900bec4f7941fd20d6f8fe7be8b040ae6adc4","k_sha256":"14b94cf0d44c9715f7b6b19638dffe012e255907ddd6fd5fc4170f23f57eb9de"} +{"key":"256-inside-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"6b8626092624533479bc18ba46f6de54cbd643f10f7e64fb93e8fa4118f6b361","sin_sha256":"140df4ec915b4d7cc25a17eba853438c6c6ed1f444831caff3d78ab8c51c9b8e","q_sha256":"41411eedb50ec9eb3905660429d8ccc74ac533eb24f516e2c4aabe6fb56619f1","k_sha256":"21351db5fc62c81b18dedd17079cfb94afce8fc866194991775ce5d809838b67"} +{"key":"256-inside-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"243b58223a86ef69b1528fc8e5ab57ce2d6131c95aa5fe7a11c4c54b1c5c2477","sin_sha256":"b157572375863b5a67db173d729737f5c5cdc0c19df7ae233455f373efca78e9","q_sha256":"c3e554dd73cbd076257399a1c681a73ec7146a42e64f458121bcf551fc85beda","k_sha256":"03fb9884e7adc7c6592a2b5ac8feb2121848543281aabe64596ccd02d41e87a2"} +{"key":"256-inside-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"inside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":0,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"243b58223a86ef69b1528fc8e5ab57ce2d6131c95aa5fe7a11c4c54b1c5c2477","sin_sha256":"b157572375863b5a67db173d729737f5c5cdc0c19df7ae233455f373efca78e9","q_sha256":"3768c85c6483d544e1d1332cb5b1576a5d294c45cb6d108bc9973a60da272da8","k_sha256":"fff583259aeb40f31f18482dc06be5b39bcdb887c39d5d23dd747b2a25d2517d"} +{"key":"256-outside-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"46470cd3fb0fbf04ec5a202def18080387a17e26609dfd2c898b1544922a72e3","k_sha256":"442720910c2f9e81a692fda617e7f291e7958c47767c5681dae2dc152604d2fd"} +{"key":"256-outside-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"66eaaa400cd8c426e69466f385c0da2755a27fad14533e121ac71ddff643dd22","k_sha256":"6327db96552af4550847e68302fc35db26f6ce4c6bfb974c3e86753532461436"} +{"key":"256-outside-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"9f4f70f54dc3c2a96b685867d93eebf893a2771568da279f1af09c4e290fff6f","k_sha256":"aba719fab18624f0bf4e18ca13f1445907b40ed37fc16469bcc12bc7dcf8f2e9"} +{"key":"256-outside-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"outside","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"dd02a4903be48b0a5164dc8aa7a038b4f31542d0502976229993b484b0afb161","k_sha256":"2d14d0767981255265a1c9e3de6d5ca5b9c423796643a7890e0a38defce745fc"} +{"key":"256-straddle-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"6a4ca7c67ef29078baab8a90fe31a9032962beaa63b9407daff32b23338d94d2","sin_sha256":"8ba60c8c231b9bab6af1d5e2cb09076e65b84d86b8b3e82c15a80e80da24ad19","q_sha256":"42fc5eb0f6209819b498c060256f339af41108ff6a2914aa58a582f681fce0de","k_sha256":"5388b4ea514dd361df01b5f6b741b0412f7a5c0cc5c4878b2b184b957e247ba4"} +{"key":"256-straddle-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"6a4ca7c67ef29078baab8a90fe31a9032962beaa63b9407daff32b23338d94d2","sin_sha256":"8ba60c8c231b9bab6af1d5e2cb09076e65b84d86b8b3e82c15a80e80da24ad19","q_sha256":"ad0297ee044b92d621dc63aac720816ca901d42480bd8623e367a87546a8a966","k_sha256":"ec36ced2f229b255cb47b1a048bd874d01e53194ebf13fecc53fe6e6f913e6f3"} +{"key":"256-straddle-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"569b55172bb1106384a5c3642bb90a77ac2997f93fbeefbcd7e43aac01ddece5","sin_sha256":"c3e314142270c58f2d6c769fe4eeb3f7074d7bf69b0fb79612308d440991d7ab","q_sha256":"6ea51e92e8611cc748425fd13b55b1c76409d7ba9ad55093d8e8c0915dcb2b7b","k_sha256":"dc92525bc9de0dbf52f6d0ea4aa9a2ab8e6fe4f5734dedb82737e94fde402f35"} +{"key":"256-straddle-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"straddle","table":[[0,1,2,2],[0,1,3,3],[0,1,4,4]],"start":2,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"569b55172bb1106384a5c3642bb90a77ac2997f93fbeefbcd7e43aac01ddece5","sin_sha256":"c3e314142270c58f2d6c769fe4eeb3f7074d7bf69b0fb79612308d440991d7ab","q_sha256":"cf12a1f9355e3f3fe0851f85eb4eb672e08cfa476b1caa9d5d30836aede2a7f0","k_sha256":"13b434bbdb9ac1da04ece39ff0fbbf32af6a3be4f65bc2c76955bc2eea29918e"} +{"key":"256-full_text_tail-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"02d54605178198f102a284547292f68800035f2ac14a556055fa8e2f69be6d95","sin_sha256":"94252e38d9026e6f504896c9006e9d00233c127eb35c61dd657d9339033c79da","q_sha256":"f95457ad86bec0f74c712d22ca7060b287e8d879649a2a2f6183614e371c272e","k_sha256":"94441d6b3975d66b0b856fa7a961b9e3cc9061e0dcf5ca4e99e3578531bc4eab"} +{"key":"256-full_text_tail-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"02d54605178198f102a284547292f68800035f2ac14a556055fa8e2f69be6d95","sin_sha256":"94252e38d9026e6f504896c9006e9d00233c127eb35c61dd657d9339033c79da","q_sha256":"11d14d87457a3056f2b0002e7ecdcaf731582f7977af57253f63e8e43934e794","k_sha256":"ee38a4c600c702b04372cf296321ebc6c417b99c8f26a92be582723370827bf0"} +{"key":"256-full_text_tail-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"4debf71d26dc4a7f4c90063db81e4cc281794397a6280856630cf69088f0736d","sin_sha256":"181ff89796c3ea57e75507139b06048799088b35a5b1ca734969f61ac5aa13a0","q_sha256":"243387c34260d98b0b7fc7c0402d7a22fa6a2b12dd175bac608989cb3b088e0f","k_sha256":"f01f02933997f089ad7da74e6e9e5330d48d8d5dce05e57221c82fa6616b9d97"} +{"key":"256-full_text_tail-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"full_text_tail","table":[[0,1,2,2,3,4,5,6],[0,1,3,3,3,4,5,6],[0,1,4,4,3,4,5,6]],"start":6,"rows":4,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"4debf71d26dc4a7f4c90063db81e4cc281794397a6280856630cf69088f0736d","sin_sha256":"181ff89796c3ea57e75507139b06048799088b35a5b1ca734969f61ac5aa13a0","q_sha256":"89e21bbff037e9c37b6644bba10c5147ed2f2b2ef581ecfb7a48382adca15dc0","k_sha256":"7696741b386796d165823c076759b8b333e41b38793c5f0930d8748b019644aa"} +{"key":"256-equal-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.0,"dtype":"f32","cos_sha256":"81db85ea7f6e9e5e296a301e27e4559ef47209c07bb8d5e1ef7fbde58cf31f43","sin_sha256":"ab7951404c5316d70cf2c4d23c22bda6688735aadbbbac1154c5561fefac5814","q_sha256":"a5ed90958f9cb48bbaecd95378dd4d5b799e89119a23948a21d6dbf152bfb6fd","k_sha256":"39020d2fe651af5740f3d25b75ab629b4939e756a014a2fd3ea29cc83b707a75"} +{"key":"256-equal-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.0,"dtype":"bf16","cos_sha256":"81db85ea7f6e9e5e296a301e27e4559ef47209c07bb8d5e1ef7fbde58cf31f43","sin_sha256":"ab7951404c5316d70cf2c4d23c22bda6688735aadbbbac1154c5561fefac5814","q_sha256":"a436e9110ac8c4d7caaa5dc8b3b2f26b65a1769c7e0740d1b250ed0b18044f25","k_sha256":"46c7938fc8efb1db7af8c0bb93b023b27c74a9f3c28563c601e10d1f60957be3"} +{"key":"256-equal-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.25,"dtype":"f32","cos_sha256":"aa7a3bff877ac876da92f531df5cdaeee6030357a759617c67773662257afa9c","sin_sha256":"634e86affba6725e9a548e130dc0f78519cc3b5b6d9595d0647299a9b742fcf4","q_sha256":"652d70474ef0a2f5ea94b13af4f06a9c47519dcb2875db09b8f7e2297dc1b537","k_sha256":"8200796c4b50eaeea4cfa542816bf432cd57aca9bd0e681af23efdce0965af7e"} +{"key":"256-equal-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"equal","table":[[0,1,2,3],[0,1,2,3],[0,1,2,3]],"start":2,"rows":4,"delta":0,"scaling":1.25,"dtype":"bf16","cos_sha256":"aa7a3bff877ac876da92f531df5cdaeee6030357a759617c67773662257afa9c","sin_sha256":"634e86affba6725e9a548e130dc0f78519cc3b5b6d9595d0647299a9b742fcf4","q_sha256":"f155c3d76915163001c8cfdf3a450fef536011a9ae9506062005556329bfacca","k_sha256":"f098bae1d55103eb101fe518b6461514939f6d84e0c8aefc21cc337985f4a6a7"} +{"key":"256-none-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"46470cd3fb0fbf04ec5a202def18080387a17e26609dfd2c898b1544922a72e3","k_sha256":"442720910c2f9e81a692fda617e7f291e7958c47767c5681dae2dc152604d2fd"} +{"key":"256-none-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"66eaaa400cd8c426e69466f385c0da2755a27fad14533e121ac71ddff643dd22","k_sha256":"6327db96552af4550847e68302fc35db26f6ce4c6bfb974c3e86753532461436"} +{"key":"256-none-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"9f4f70f54dc3c2a96b685867d93eebf893a2771568da279f1af09c4e290fff6f","k_sha256":"aba719fab18624f0bf4e18ca13f1445907b40ed37fc16469bcc12bc7dcf8f2e9"} +{"key":"256-none-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"none","table":null,"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"dd02a4903be48b0a5164dc8aa7a038b4f31542d0502976229993b484b0afb161","k_sha256":"2d14d0767981255265a1c9e3de6d5ca5b9c423796643a7890e0a38defce745fc"} +{"key":"256-empty-1.0-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"f32","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"46470cd3fb0fbf04ec5a202def18080387a17e26609dfd2c898b1544922a72e3","k_sha256":"442720910c2f9e81a692fda617e7f291e7958c47767c5681dae2dc152604d2fd"} +{"key":"256-empty-1.0-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.0,"dtype":"bf16","cos_sha256":"b74e1412fd4c010fe4680296bce1face53fb27a160c5942d25babe45740f532c","sin_sha256":"690e9805f9d0f5b319dd189bccb902fc94f387d41136060e3bccde090ff6f35e","q_sha256":"66eaaa400cd8c426e69466f385c0da2755a27fad14533e121ac71ddff643dd22","k_sha256":"6327db96552af4550847e68302fc35db26f6ce4c6bfb974c3e86753532461436"} +{"key":"256-empty-1.25-f32","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"f32","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"9f4f70f54dc3c2a96b685867d93eebf893a2771568da279f1af09c4e290fff6f","k_sha256":"aba719fab18624f0bf4e18ca13f1445907b40ed37fc16469bcc12bc7dcf8f2e9"} +{"key":"256-empty-1.25-bf16","dim":256,"rotary":64,"sections":[11,11,10],"heads":24,"mode":"empty","table":[[],[],[]],"start":4,"rows":2,"delta":-1,"scaling":1.25,"dtype":"bf16","cos_sha256":"08deae1f38cc2ff0673dfcee084a24d2f5202a58c65012855d97b60140e18198","sin_sha256":"8410a14d0c8c5ea90587039a782c7fab9bad800e5365cf3736f40de8146c046e","q_sha256":"dd02a4903be48b0a5164dc8aa7a038b4f31542d0502976229993b484b0afb161","k_sha256":"2d14d0767981255265a1c9e3de6d5ca5b9c423796643a7890e0a38defce745fc"} diff --git a/src/engine/metal/qwen_mtplx.rs b/src/engine/metal/qwen_mtplx.rs index 47bc1c4..20a3c08 100644 --- a/src/engine/metal/qwen_mtplx.rs +++ b/src/engine/metal/qwen_mtplx.rs @@ -35,6 +35,8 @@ mod issue104; mod issue105; #[cfg(test)] mod issue106; +#[cfg(test)] +mod issue108; mod linear; mod load; mod mlp; diff --git a/src/engine/metal/qwen_mtplx/issue103_analysis.rs b/src/engine/metal/qwen_mtplx/issue103_analysis.rs index ee64e4f..884d20d 100644 --- a/src/engine/metal/qwen_mtplx/issue103_analysis.rs +++ b/src/engine/metal/qwen_mtplx/issue103_analysis.rs @@ -430,7 +430,7 @@ fn issue103_vision_position_receipts() { gpu, ) .unwrap(); - ops::evaluate(&streams, &[output], gpu, false).unwrap(); + ops::evaluate(&streams, std::slice::from_ref(&output), gpu, false).unwrap(); let projected = ops::reshape( &ops::rms_norm(&x, Some(&norm), 1e-6, gpu).unwrap(), &[1, rows, 1, 8], @@ -490,11 +490,81 @@ fn issue103_vision_position_receipts() { "{}", serde_json::json!({"case":"vision_positions","start":pos,"rows":rows,"table_end":4,"delta":-1,"different_rows":bad,"actual_keys":a,"expected_keys":e}) ); - if pos == 2 { - assert_eq!(bad, vec![0, 1]); - } else { - assert!(bad.is_empty()); + assert!(bad.is_empty(), "each row must use its own vision position"); + // An equivalent split preserves outputs, offsets and the next call. + let mut split = QsaCache::new(4); + if pos > 0 { + let prefix = bf(&[1, 1, pos, 8], &vec![0.; pos as usize * 8]); + split.kv.update_and_fetch(&prefix, &prefix, gpu).unwrap(); } + let boundary = (4 - pos).clamp(0, rows); + let mut pieces = Vec::new(); + for (begin, end) in [(0, boundary), (boundary, rows)] { + if begin == end { + continue; + } + let piece = ops::slice(&x, &[0, begin, 0], &[1, end, 8], &[1; 3], gpu).unwrap(); + pieces.push( + attention + .call( + &piece, + &mut split, + &options, + Some((Some(&table), -1)), + &streams, + gpu, + ) + .unwrap(), + ); + } + let joined = ops::concatenate(&pieces, 1, gpu).unwrap(); + assert_eq!(read(&output), read(&joined)); + assert_eq!(cache.kv.offset, pos + rows); + assert_eq!(split.kv.offset, cache.kv.offset); + let next = bf(&[1, 1, 8], &[0.3, -0.7, 0.2, 0.9, -0.2, 0.4, 0.6, -0.1]); + let next_whole = attention + .call( + &next, + &mut cache, + &options, + Some((Some(&table), -1)), + &streams, + gpu, + ) + .unwrap(); + let next_split = attention + .call( + &next, + &mut split, + &options, + Some((Some(&table), -1)), + &streams, + gpu, + ) + .unwrap(); + assert_eq!(read(&next_whole), read(&next_split)); + assert_eq!(cache.kv.offset, pos + rows + 1); + assert_eq!(split.kv.offset, cache.kv.offset); + for (actual, expected) in [ + (&cache.kv.keys, &split.kv.keys), + (&cache.kv.values, &split.kv.values), + ] { + let valid = |array: &Option| { + ops::slice( + &array.as_ref().unwrap().borrow(), + &[0, 0, 0, 0], + &[1, 1, pos + rows + 1, 8], + &[1; 4], + gpu, + ) + .unwrap() + }; + assert_eq!(read(&valid(actual)), read(&valid(expected))); + } + println!( + "{}", + serde_json::json!({"case":"vision_chunk_next_state","start":pos,"rows":rows,"offset":cache.kv.offset,"exact":true}) + ); } let prefix = bf( &[1, 12, 8], diff --git a/src/engine/metal/qwen_mtplx/issue108.rs b/src/engine/metal/qwen_mtplx/issue108.rs new file mode 100644 index 0000000..3e9b26e --- /dev/null +++ b/src/engine/metal/qwen_mtplx/issue108.rs @@ -0,0 +1,90 @@ +use super::array::{Array, Dtype}; +use super::{ops, rope}; + +#[test] +#[ignore = "Metal and pinned MTPLX chunk fixtures; external supervisor required"] +fn issue108_attention_chunk_positions() { + super::configure_sources().unwrap(); + let _context = crate::engine::metal::gpu::Context::open_qwen(0).unwrap(); + let streams = super::stream::Streams::new(Some(super::allocator::Allocator::new().unwrap())); + let gpu = streams.default_stream(super::stream::Device::Gpu).unwrap(); + let directory = std::path::PathBuf::from(std::env::var_os("DS4_ISSUE108_ORACLE").unwrap()); + let cases: Vec = + serde_json::from_slice(&std::fs::read(directory.join("manifest.json")).unwrap()).unwrap(); + let ints = |shape: &[i32], values: &[i32]| { + Array::new( + shape, + Dtype::I32, + super::scalar_buffer( + &values + .iter() + .flat_map(|n| n.to_le_bytes()) + .collect::>(), + ) + .unwrap(), + ) + .unwrap() + }; + for case in &cases { + let n = |name: &str| case[name].as_i64().unwrap() as i32; + let key = case["key"].as_str().unwrap(); + let sections: Vec = serde_json::from_value(case["sections"].clone()).unwrap(); + let axes = ints(&[n("rotary") / 2], &rope::build_mrope_axes(§ions, true)); + let table: Option>> = serde_json::from_value(case["table"].clone()).unwrap(); + let table = table.map(|t| ints(&[3, t[0].len() as i32], &t.concat())); + let (freq, _) = rope::inv_freq_and_scaling(n("rotary"), 10000000., None, gpu).unwrap(); + let scaling = case["scaling"].as_f64().unwrap(); + let (cos, sin) = rope::vision_chunk_cos_sin( + (table.as_ref(), n("delta")), + n("start"), + n("rows"), + &freq, + &axes, + scaling, + gpu, + ) + .unwrap(); + let check = |part: &str, value: &Array| { + let actual = ops::read_f32_as_f64( + &ops::astype(value, Dtype::F32, false, gpu).unwrap(), + &streams, + gpu, + ) + .unwrap() + .into_iter() + .flat_map(|v| (v as f32).to_le_bytes()) + .collect::>(); + assert_eq!( + actual, + std::fs::read(directory.join(format!("{key}-{part}.f32"))).unwrap(), + "{key} {part}" + ); + }; + check("cos", &cos); + check("sin", &sin); + for (part, heads, seed) in [("q", n("heads"), 3), ("k", 1, 11)] { + let shape = [1, n("rows"), heads, n("dim")]; + let values = (0..shape.iter().product::()) + .map(|i| ((i * 17 + seed) % 97 - 48) as f32 / 32.) + .collect::>(); + let dtype = if case["dtype"] == "bf16" { + Dtype::BF16 + } else { + Dtype::F32 + }; + let raw = ops::astype( + &super::vision::f32_array(&shape, &values).unwrap(), + dtype, + false, + gpu, + ) + .unwrap(); + check(part, &rope::apply_partial(&raw, &cos, &sin, gpu).unwrap()); + } + println!( + "{}", + serde_json::json!({"case":"attention_chunk","key":key,"exact_q_k_and_tables":true}) + ); + } + assert_eq!(cases.len(), 56); +} diff --git a/src/engine/metal/qwen_mtplx/qsa_attention.rs b/src/engine/metal/qwen_mtplx/qsa_attention.rs index 3aaa252..5bce867 100644 --- a/src/engine/metal/qwen_mtplx/qsa_attention.rs +++ b/src/engine/metal/qwen_mtplx/qsa_attention.rs @@ -171,22 +171,15 @@ impl Attention<'_> { rope::cos_sin(&positions, self.inv_freq, self.rope_scaling, stream) }; let (cos, sin) = match (vision, self.mrope_axes) { - (Some((table, delta)), Some(axes)) => { - if let Some(table) = - table.filter(|t| t.layout().shape().get(1).is_some_and(|&n| end <= n)) - { - let table = ops::slice( - table, - &[0, pos], - &[table.layout().shape()[0], end], - &[1, 1], - stream, - )?; - rope::mrope_cos_sin(&table, self.inv_freq, axes, stream)? - } else { - plain_rope(delta)? - } - } + (Some((table, delta)), Some(axes)) => rope::vision_chunk_cos_sin( + (table, delta), + pos, + rows, + self.inv_freq, + axes, + self.rope_scaling, + stream, + )?, _ => plain_rope(0)?, }; let q = ops::transpose( diff --git a/src/engine/metal/qwen_mtplx/rope.rs b/src/engine/metal/qwen_mtplx/rope.rs index a56268a..b748170 100644 --- a/src/engine/metal/qwen_mtplx/rope.rs +++ b/src/engine/metal/qwen_mtplx/rope.rs @@ -246,6 +246,70 @@ pub(super) fn mrope_cos_sin( )) } +/// MTPLX _vision_chunk_cos_sin: preserve the inside/outside graphs and +/// concatenate their tables when one attention chunk contains both kinds of row. +pub(super) fn vision_chunk_cos_sin( + (table, delta): (Option<&Array>, i32), + start: i32, + rows: i32, + inv_freq: &Array, + axes: &Array, + scaling: f64, + stream: Stream, +) -> Result<(Array, Array), String> { + let end = start + .checked_add(rows) + .ok_or("QSA vision position overflow")?; + let table_len = match table { + Some(table) => { + let layout = table.layout(); + let shape = layout.shape(); + if shape.len() != 2 || shape[0] != 3 { + return Err("vision positions require a three-axis table".into()); + } + shape[1] + } + None => 0, + }; + let inside = (end.min(table_len) - start).max(0); + let covered = if inside > 0 { + let positions = ops::slice( + table.unwrap(), + &[0, start], + &[3, start + inside], + &[1, 1], + stream, + )?; + let (mut cos, mut sin) = mrope_cos_sin(&positions, inv_freq, axes, stream)?; + if scaling != 1.0 { + cos = binary(&cos, &scalar(scaling)?, Binary::Multiply, stream)?; + sin = binary(&sin, &scalar(scaling)?, Binary::Multiply, stream)?; + } + if inside == rows { + return Ok((cos, sin)); + } + Some((cos, sin)) + } else { + None + }; + let first = start + .checked_add(inside) + .and_then(|p| p.checked_add(delta)) + .ok_or("QSA vision position overflow")?; + let stop = first + .checked_add(rows - inside) + .ok_or("QSA vision position overflow")?; + let positions = indexing::arange_range_i32(first, stop, stream)?; + let (cos, sin) = cos_sin(&positions, inv_freq, scaling, stream)?; + match covered { + Some((cos_in, sin_in)) => Ok(( + ops::concatenate(&[cos_in, cos], 0, stream)?, + ops::concatenate(&[sin_in, sin], 0, stream)?, + )), + None => Ok((cos, sin)), + } +} + /// MTPLX _vision_position_cos_sin: choose table versus text+delta per row. pub(super) fn vision_position_cos_sin( positions: &Array, diff --git a/tools/issue108-reference.py b/tools/issue108-reference.py new file mode 100644 index 0000000..ce38a26 --- /dev/null +++ b/tools/issue108-reference.py @@ -0,0 +1,55 @@ +"""Pinned attention chunk RoPE oracle, with independent per-row controls.""" +import hashlib +import json +from pathlib import Path +import subprocess +import sys + +checkout, output = map(Path, sys.argv[1:3]) +assert subprocess.check_output(['git', '-C', str(checkout), 'rev-parse', 'HEAD'], text=True).strip() == '9882703f3105363ddc37eca9f97aa09a1d387112' +sys.path.insert(0, str(checkout)) +import mlx.core as mx +import numpy as np +from mtplx.models.qwen4_exp import _vision_chunk_cos_sin, _apply_partial_rope, _build_mrope_axes, _rope_cos_sin, _rope_inv_freq_and_scaling, TextArgs + +output.mkdir(parents=True, exist_ok=True) +manifest = [] +for dim, rotary, sections, heads in [(8, 8, [2,1,1], 2), (256,64,[11,11,10],24)]: + axes = mx.array(_build_mrope_axes(sections, True), dtype=mx.int32) + freq, _ = _rope_inv_freq_and_scaling(TextArgs(head_dim=dim, rope_parameters=dict(rope_type="default", rope_theta=10000000, partial_rotary_factor=rotary/dim))) + short = [[0,1,2,2],[0,1,3,3],[0,1,4,4]] + full = [r + [3,4,5,6] for r in short] + equal = [list(range(4))] * 3 + for mode, table, start, rows, delta in [ + ('inside',short,0,4,-1), ('outside',short,4,2,-1), + ('straddle',short,2,4,-1), ('full_text_tail',full,6,4,-1), + ('equal',equal,2,4,0), ('none',None,4,2,-1), ('empty',[[],[],[]],4,2,-1), + ]: + for scaling in [1.,1.25]: + t = None if table is None else mx.array(table,dtype=mx.int32) + cos, sin = _vision_chunk_cos_sin(t,delta,start,rows,freq,axes,scaling) + # Independent row-position construction; no chunk helper is used. + pos3 = np.stack([np.array(table)[:,p] if table is not None and p