Guard large unaligned Qwen expert gathers

docs/issue110-sorted-gather/2048-cold-state.json

@@ -0,0 +1 @@
+{"case":"turn","chunk":2048,"phase":"cold","cached":0,"prompt_len":7374,"state":{"ids":[9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,10Line truncated

docs/issue110-sorted-gather/2048-warm-state.json

@@ -0,0 +1 @@
+{"case":"turn","chunk":2048,"phase":"warm","cached":7377,"prompt_len":10655,"state":{"ids":[9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,1Line truncated

docs/issue110-sorted-gather/4096-cold-state.json

@@ -0,0 +1 @@
+{"case":"turn","chunk":4096,"phase":"cold","cached":0,"prompt_len":7374,"state":{"ids":[9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,10Line truncated

docs/issue110-sorted-gather/4096-warm-state.json

@@ -0,0 +1 @@
+{"case":"turn","chunk":4096,"phase":"warm","cached":7377,"prompt_len":10655,"state":{"ids":[9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,11,1075,374,264,1092,13,9707,1Line truncated

docs/issue110-sorted-gather/README.md

@@ -0,0 +1,85 @@
+# Issue 110: guarded stock expert gathers
+
+Qwen's stock MoE gathers now mirror pinned MTPLX's
+`moe_sorted_gather._call_pad` / `gather_qmm` wrapper. On hardware that actually
+uses NAX, an expert-sorted affine transposed `[rows, 1, K]` call above 32,767
+rows is padded to the next multiple of 64 unless already aligned. Activations
+are zero padded, the final expert ID is repeated, explicit left indices append
+an arange in their own dtype, and the output is sliced back to its real rows.
+Both separate/fused Gate/Up and Down use this entry point. The ordinary runtime
+primitive and pinned 0.32.2 shader remain unchanged.
+
+The hardware test deliberately ignores the MTPLX route-fallback switch, as the
+stock MLX shader does. Unit controls exercise non-NAX and ineligible shapes;
+no physical M1–M4 hardware result is claimed. This implementation is confined
+to Qwen; DS4/GLM have no callers of the new wrapper. It is an oracle correctness
+fix, not a new DS4Server feature.
+
+Old checkpoint manifests lack both producer hardware and prefill chunk sizes.
+Files with at least 3,277 tokens may contain affected rows and are rejected
+before state restoration unless they carry the new guard policy. Short old
+files remain compatible; new files retain the existing container format.
+The admission test verifies rejection leaves healthy live state unchanged.
+
+## Verification
+
+- Original initialized-NaN low-level sentinels still demonstrate the pinned
+  shader defect. The same five token counts (3,276 / 3,277 / 3,404 / 4,095 /
+  4,096) now additionally execute the real guarded graph and require every
+  returned row to equal the exact expected BF16 value, with original shape.
+- Installed layer-0 weights exercise 512 experts, fused Gate/Up K=2560/N=1280
+  and Down K=640/N=2560, group64/4-bit. All output bytes match the pinned guarded
+  oracle at 20,480 / 32,760 / 32,770 / 34,040 / 40,950 / 40,960 rows. Explicit
+  reverse left indices at 34,040 rows also match. Sentinel weights control
+  group32/4-bit; the unchanged runtime-gather suites cover other dispatches.
+- The current optimized oracle's row-map SwiGLU equals its guarded stock chain
+  byte-for-byte at 2,048 / 3,277 / 4,095 / 4,096 tokens, with successful first-use
+  canaries in both narrow/wide regimes and actual kernel engagement recorded.
+  This is an oracle control. The native optimized route belongs to #112.
+- Allocation refusal releases the failed padded graph; a fresh valid retry
+  computes all rows exactly and returns allocator activity to zero.
+- Full-model cold and warm checks use actual server execution defaults, greedy
+  generation, 7,374 cold tokens, then a 3,278-token suffix. Both 2,048 and 4,096
+  chunk sizes are exercised. `MTPLX_GDN_BOUNDARY_TAIL_INTERVAL=0` is explicit:
+  default boundary splitting would divide the final partial span into 256-row
+  pieces and conceal this reproduction. Cold and warm affected body spans are
+  3,277 tokens; cancellation occurs after a settled affected chunk and preserves
+  the committed state for retry. The default 2,048 size remains a control.
+
+All four full-model cases match the oracle byte-for-byte in final logits, token
+IDs and every saved trunk/cache tensor. The required formatting, warnings-denied
+Clippy, signed bundle and ordinary tests pass (309 application + 8 tool tests).
+The official focused upstream suite passes 16 tests, including both block
+equivalence variants with the canary/kernel active. One ordinary-suite run hit
+the unrelated numeric Git-revision YAML fixture flake tracked as #115; the
+unchanged suite passed on rerun.
+
+The oracle is MTPLX `9882703f3105363ddc37eca9f97aa09a1d387112` with MLX exactly
+0.32.2 and the unchanged installed Bare-Speed revision from
+`../issue104-ar/environment.json`. No runtime/assets/default chunk changes or
+performance claims accompany this correctness fix. See `checks.json` and
+adjacent receipts for actual completed results.
+
+## Reproduction
+
+```sh
+local-eval-results/mtplx-reference-env-0.32.2/bin/python \
+  tools/issue110-reference.py /path/to/pinned/mtplx /path/to/model \
+  > /tmp/issue110-oracle.jsonl
+cargo test --all-features --bin ds4-server --no-run
+cargo build --bin test-supervisor
+export DS4_QWEN38_ARTIFACTS=/path/to/model
+export DS4_ISSUE110_ORACLE=/tmp/issue110-oracle.jsonl
+# Use the executable printed by cargo; run GPU processes serially.
+target/debug/test-supervisor 4096 60 120 --command TEST_EXECUTABLE \
+  issue103_sorted_gather_receipts --ignored --nocapture --test-threads=1
+target/debug/test-supervisor 16384 60 120 --command TEST_EXECUTABLE \
+  issue110_actual_projection_rows_match_oracle --ignored --nocapture --test-threads=1
+export DS4_ISSUE110_DIR=/tmp/issue110-product
+export MTPLX_GDN_BOUNDARY_TAIL_INTERVAL=0
+target/debug/test-supervisor 114688 60 120 --command TEST_EXECUTABLE \
+  issue110_large_cold_and_warm_prefill --ignored --nocapture --test-threads=1
+local-eval-results/mtplx-reference-env-0.32.2/bin/python \
+  tools/issue110-model-reference.py /path/to/pinned/mtplx \
+  "$DS4_QWEN38_ARTIFACTS" "$DS4_ISSUE110_DIR" /tmp/issue110-model-oracle
+```

docs/issue110-sorted-gather/allocation.jsonl

@@ -0,0 +1,2 @@
+{"case":"allocation_refusal","limit":1,"partial_peak":4096,"released":true,"retry_exact":true}
+{"case":"allocation_refusal","limit":2,"partial_peak":151552,"released":true,"retry_exact":true}

docs/issue110-sorted-gather/checks.json

@@ -0,0 +1,138 @@
+[
+  {
+    "check": "fmt",
+    "passed": true
+  },
+  {
+    "check": "clippy",
+    "passed": true
+  },
+  {
+    "check": "bundle",
+    "passed": true
+  },
+  {
+    "check": "tests",
+    "passed": true,
+    "results": [
+      "test result: ok. 309 passed; 0 failed; 224 ignored; 0 measured; 0 filtered out; finished in 5.31s",
+      "test result: ok. 4 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out; finished in 0.01s",
+      "test result: ok. 4 passed; 0 failed; 0 ignored; 0 measured; 0 filtered out; finished in 2.36s"
+    ],
+    "prior_run": "Unrelated all-decimal Git SHA YAML fixture failure (9839678), tracked separately as #115; unchanged rerun passed."
+  },
+  {
+    "check": "sentinel",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 531 filtered out; finished in 2.15s"
+    ],
+    "resources": {
+      "elapsed_ms": 3104,
+      "error": null,
+      "peak_physical_bytes": 551748736
+    }
+  },
+  {
+    "check": "projections",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 531 filtered out; finished in 57.88s"
+    ],
+    "resources": {
+      "elapsed_ms": 57987,
+      "error": null,
+      "peak_physical_bytes": 5888003816
+    }
+  },
+  {
+    "check": "allocation",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 0.69s"
+    ],
+    "resources": {
+      "elapsed_ms": 1711,
+      "error": null,
+      "peak_physical_bytes": 224527224
+    }
+  },
+  {
+    "check": "codec",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 0.06s"
+    ],
+    "resources": {
+      "elapsed_ms": 105,
+      "error": null,
+      "peak_physical_bytes": 180296
+    }
+  },
+  {
+    "check": "final-product",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 119.17s"
+    ],
+    "profile": "Server defaults, MTPLX_GDN_BOUNDARY_TAIL_INTERVAL=0; chunks 2048/4096",
+    "resources": {
+      "elapsed_ms": 120155,
+      "error": null,
+      "peak_physical_bytes": 102111559024
+    }
+  },
+  {
+    "check": "model-oracle",
+    "passed": true,
+    "profile": "Server defaults, MTPLX_GDN_BOUNDARY_TAIL_INTERVAL=0; chunks 2048/4096",
+    "resources": {
+      "elapsed_ms": 34036,
+      "error": null,
+      "peak_physical_bytes": 93475651360
+    }
+  },
+  {
+    "check": "graph-control",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 0.72s"
+    ],
+    "resources": {
+      "elapsed_ms": 751,
+      "error": null,
+      "peak_physical_bytes": 155943800
+    }
+  },
+  {
+    "check": "runtime-control",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 7.43s"
+    ],
+    "resources": {
+      "elapsed_ms": 7515,
+      "error": null,
+      "peak_physical_bytes": 313934784
+    }
+  },
+  {
+    "check": "route-switch",
+    "passed": true,
+    "results": [
+      "test result: ok. 1 passed; 0 failed; 0 ignored; 0 measured; 532 filtered out; finished in 2.23s"
+    ],
+    "resources": {
+      "elapsed_ms": 2264,
+      "error": null,
+      "peak_physical_bytes": 528237648
+    }
+  },
+  {
+    "check": "upstream",
+    "passed": true,
+    "results": [
+      "16 passed in 0.74s"
+    ]
+  }
+]

docs/issue110-sorted-gather/model-oracle.jsonl

@@ -0,0 +1,16 @@
+{"case":"token","chunk":2048,"phase":"cold","ids":[374]}
+{"case":"token","chunk":2048,"phase":"cold","ids":[264]}
+{"case":"token","chunk":2048,"phase":"cold","ids":[1092]}
+{"case":"turn","chunk":2048,"phase":"cold","exact":true,"cached":0,"tokens":[374,264,1092],"guard_before":{"padded_calls":0,"padded_rows":0},"guard_after":{"padded_calls":0,"padded_rows":0},"kernel":{"calls":192,"fallbacks":0,"canaries":1,"canary_failures":0,"header_failures":0}}
+{"case":"token","chunk":2048,"phase":"warm","ids":[1075]}
+{"case":"token","chunk":2048,"phase":"warm","ids":[374]}
+{"case":"token","chunk":2048,"phase":"warm","ids":[264]}
+{"case":"turn","chunk":2048,"phase":"warm","exact":true,"cached":7377,"tokens":[1075,374,264],"guard_before":{"padded_calls":0,"padded_rows":0},"guard_after":{"padded_calls":0,"padded_rows":0},"kernel":{"calls":288,"fallbacks":0,"canaries":1,"canary_failures":0,"header_failures":0}}
+{"case":"token","chunk":4096,"phase":"cold","ids":[374]}
+{"case":"token","chunk":4096,"phase":"cold","ids":[264]}
+{"case":"token","chunk":4096,"phase":"cold","ids":[1092]}
+{"case":"turn","chunk":4096,"phase":"cold","exact":true,"cached":0,"tokens":[374,264,1092],"guard_before":{"padded_calls":0,"padded_rows":0},"guard_after":{"padded_calls":48,"padded_rows":2976},"kernel":{"calls":384,"fallbacks":0,"canaries":2,"canary_failures":0,"header_failures":0}}
+{"case":"token","chunk":4096,"phase":"warm","ids":[1075]}
+{"case":"token","chunk":4096,"phase":"warm","ids":[374]}
+{"case":"token","chunk":4096,"phase":"warm","ids":[264]}
+{"case":"turn","chunk":4096,"phase":"warm","exact":true,"cached":7377,"tokens":[1075,374,264],"guard_before":{"padded_calls":48,"padded_rows":2976},"guard_after":{"padded_calls":96,"padded_rows":5952},"kernel":{"calls":432,"fallbacks":0,"canaries":2,"canary_failures":0,"header_failures":0}}

docs/issue110-sorted-gather/product.jsonl

@@ -0,0 +1,5 @@
+{"case":"turn","chunk":2048,"phase":"cold","cached":0,"prompt_len":7374,"generated":[374,264,1092],"padded_calls":0}
+{"case":"turn","chunk":2048,"phase":"warm","cached":7377,"prompt_len":10655,"generated":[1075,374,264],"padded_calls":0}
+{"case":"turn","chunk":4096,"phase":"cold","cached":0,"prompt_len":7374,"generated":[374,264,1092],"padded_calls":96}
+{"case":"cancel","chunk":4096,"saved_state_unchanged":true}
+{"case":"turn","chunk":4096,"phase":"warm","cached":7377,"prompt_len":10655,"generated":[1075,374,264],"padded_calls":96}

docs/issue110-sorted-gather/projection-oracle.jsonl

@@ -0,0 +1,18 @@
+{"case": "projection", "projection": "gate_up", "rows": 20480, "k": 2560, "n": 1280, "left": false, "sha256": "ae97282624851e970b417c37b5a0f3896c6f1c56cbfb0735d76790d81cfc1d1b", "pad": 0, "guard_before": {"padded_calls": 0, "padded_rows": 0}, "guard_after": {"padded_calls": 0, "padded_rows": 0}}
+{"case": "projection", "projection": "down", "rows": 20480, "k": 640, "n": 2560, "left": false, "sha256": "2d1a21b02bedf3f6c34df282d898c18fa9f7d71a67b5245f5967605e5457e13a", "pad": 0, "guard_before": {"padded_calls": 0, "padded_rows": 0}, "guard_after": {"padded_calls": 0, "padded_rows": 0}}
+{"case": "projection", "projection": "gate_up", "rows": 32760, "k": 2560, "n": 1280, "left": false, "sha256": "f3d7fd1a7df2efd8a43d365f1a30aa480efe61b7655641fcf71cb0f0492d7c52", "pad": 0, "guard_before": {"padded_calls": 0, "padded_rows": 0}, "guard_after": {"padded_calls": 0, "padded_rows": 0}}
+{"case": "projection", "projection": "down", "rows": 32760, "k": 640, "n": 2560, "left": false, "sha256": "ee5f7dd6cc53fa680c985a1c56ca752dd7243f6a2663a9a870eccb4f2b8f68e8", "pad": 0, "guard_before": {"padded_calls": 0, "padded_rows": 0}, "guard_after": {"padded_calls": 0, "padded_rows": 0}}
+{"case": "projection", "projection": "gate_up", "rows": 32770, "k": 2560, "n": 1280, "left": false, "sha256": "92dd7f20255bcf2e205a6d8e708ae450b47024ff12839cb393c1160db5af33aa", "pad": 62, "guard_before": {"padded_calls": 0, "padded_rows": 0}, "guard_after": {"padded_calls": 1, "padded_rows": 62}}
+{"case": "projection", "projection": "down", "rows": 32770, "k": 640, "n": 2560, "left": false, "sha256": "66e37f134a0183977ee1a39a3afd70e09160efab3ceee2b68d8caa7c49390e2a", "pad": 62, "guard_before": {"padded_calls": 1, "padded_rows": 62}, "guard_after": {"padded_calls": 2, "padded_rows": 124}}
+{"case": "projection", "projection": "gate_up", "rows": 34040, "k": 2560, "n": 1280, "left": false, "sha256": "ca7dcdcd30ad27c0936420180e42d637009d90372f1e702f79847851831de071", "pad": 8, "guard_before": {"padded_calls": 2, "padded_rows": 124}, "guard_after": {"padded_calls": 3, "padded_rows": 132}}
+{"case": "projection", "projection": "gate_up", "rows": 34040, "k": 2560, "n": 1280, "left": true, "sha256": "d79b05579846810a34b955ef65535adba962d9d02a17bb0ecace423c6ce0791c", "pad": 8, "guard_before": {"padded_calls": 3, "padded_rows": 132}, "guard_after": {"padded_calls": 4, "padded_rows": 140}}
+{"case": "projection", "projection": "down", "rows": 34040, "k": 640, "n": 2560, "left": false, "sha256": "53768e60df9fad3dc06641f5245e6f67ac28e1278dc8fec8e62486858f2de341", "pad": 8, "guard_before": {"padded_calls": 4, "padded_rows": 140}, "guard_after": {"padded_calls": 5, "padded_rows": 148}}
+{"case": "projection", "projection": "down", "rows": 34040, "k": 640, "n": 2560, "left": true, "sha256": "f02f45549d38acd5979cdd768305b39268493449191a0a232514978a4015a658", "pad": 8, "guard_before": {"padded_calls": 5, "padded_rows": 148}, "guard_after": {"padded_calls": 6, "padded_rows": 156}}
+{"case": "projection", "projection": "gate_up", "rows": 40950, "k": 2560, "n": 1280, "left": false, "sha256": "393e33c2576b26e23f6b91bf2ff2ea7e590322a8cfb88a237d7e90cfea74acc8", "pad": 10, "guard_before": {"padded_calls": 6, "padded_rows": 156}, "guard_after": {"padded_calls": 7, "padded_rows": 166}}
+{"case": "projection", "projection": "down", "rows": 40950, "k": 640, "n": 2560, "left": false, "sha256": "e66a8a49018a918e26e3f56a8859ab7894e3d75f0951858e43856fbe6052616e", "pad": 10, "guard_before": {"padded_calls": 7, "padded_rows": 166}, "guard_after": {"padded_calls": 8, "padded_rows": 176}}
+{"case": "projection", "projection": "gate_up", "rows": 40960, "k": 2560, "n": 1280, "left": false, "sha256": "29b3628139c5f7c49d5b5b46134a3141958e9a9bc3360956a770ec3d5b65af5d", "pad": 0, "guard_before": {"padded_calls": 8, "padded_rows": 176}, "guard_after": {"padded_calls": 8, "padded_rows": 176}}
+{"case": "projection", "projection": "down", "rows": 40960, "k": 640, "n": 2560, "left": false, "sha256": "d40d0c9c449d3ac4631a43f775815d1baea299dbfbee39c7585bad4997e61237", "pad": 0, "guard_before": {"padded_calls": 8, "padded_rows": 176}, "guard_after": {"padded_calls": 8, "padded_rows": 176}}
+{"case": "optimized_oracle_control", "tokens": 2048, "rows": 20480, "exact": true, "sha256": "c53aac85f49f4c1b6c11e3da858d2abf4e43de241bf3c0347d2590d0181e2249", "kernel_before": {"calls": 0, "fallbacks": 0, "canaries": 0, "canary_failures": 0, "header_failures": 0}, "kernel_after": {"calls": 1, "fallbacks": 0, "canaries": 1, "canary_failures": 0, "header_failures": 0}}
+{"case": "optimized_oracle_control", "tokens": 3277, "rows": 32770, "exact": true, "sha256": "10f573d0eaf5169475781c4f75eb923fa06d7f88d1684a6d5d2ce8acfdc526a4", "kernel_before": {"calls": 1, "fallbacks": 0, "canaries": 1, "canary_failures": 0, "header_failures": 0}, "kernel_after": {"calls": 2, "fallbacks": 0, "canaries": 2, "canary_failures": 0, "header_failures": 0}}
+{"case": "optimized_oracle_control", "tokens": 4095, "rows": 40950, "exact": true, "sha256": "738212404e537f0608d4e9c9b8d0e5ff2061587197c3391b1220975f93dd8d9f", "kernel_before": {"calls": 2, "fallbacks": 0, "canaries": 2, "canary_failures": 0, "header_failures": 0}, "kernel_after": {"calls": 3, "fallbacks": 0, "canaries": 2, "canary_failures": 0, "header_failures": 0}}
+{"case": "optimized_oracle_control", "tokens": 4096, "rows": 40960, "exact": true, "sha256": "e1a6630ec5cd7e523c98eb352e328560fdd733a640fc7dd69a5a96e14b462d15", "kernel_before": {"calls": 3, "fallbacks": 0, "canaries": 2, "canary_failures": 0, "header_failures": 0}, "kernel_after": {"calls": 4, "fallbacks": 0, "canaries": 2, "canary_failures": 0, "header_failures": 0}}

docs/issue110-sorted-gather/projections.jsonl

@@ -0,0 +1,13 @@
+{"case":"projection","projection":"down","rows":20480,"k":640,"n":2560,"left":false,"exact":true,"sha256":"2d1a21b02bedf3f6c34df282d898c18fa9f7d71a67b5245f5967605e5457e13a","dispatches":"[(\"arangeuint32\", [20480, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_32_bn_64_bk_64_wm_2_wn_2\", [40, 640, 1], [32, 2, 2], false)]"}
+{"case":"projection","projection":"gate_up","rows":32760,"k":2560,"n":1280,"left":false,"exact":true,"sha256":"f3d7fd1a7df2efd8a43d365f1a30aa480efe61b7655641fcf71cb0f0492d7c52","dispatches":"[(\"arangeuint32\", [32760, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_32_bn_64_bk_64_wm_2_wn_2\", [20, 1024, 1], [32, 2, 2], false)]"}
+{"case":"projection","projection":"down","rows":32760,"k":640,"n":2560,"left":false,"exact":true,"sha256":"ee5f7dd6cc53fa680c985a1c56ca752dd7243f6a2663a9a870eccb4f2b8f68e8","dispatches":"[(\"arangeuint32\", [32760, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_32_bn_64_bk_64_wm_2_wn_2\", [40, 1024, 1], [32, 2, 2], false)]"}
+{"case":"projection","projection":"gate_up","rows":32770,"k":2560,"n":1280,"left":false,"exact":true,"sha256":"92dd7f20255bcf2e205a6d8e708ae450b47024ff12839cb393c1160db5af33aa","dispatches":"[(\"sn_copybfloat16bfloat16\", [39680, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [32770, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [62, 1, 1], [32, 1, 1], true), (\"arangeuint32\", [32832, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [83891200, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [158720, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [20, 513, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [41945600, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"down","rows":32770,"k":640,"n":2560,"left":false,"exact":true,"sha256":"66e37f134a0183977ee1a39a3afd70e09160efab3ceee2b68d8caa7c49390e2a","dispatches":"[(\"s_copybfloat16bfloat16\", [39680, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [32770, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [62, 1, 1], [32, 1, 1], true), (\"arangeuint32\", [32832, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [20972800, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [39680, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [40, 513, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [83891200, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"gate_up","rows":34040,"k":2560,"n":1280,"left":false,"exact":true,"sha256":"ca7dcdcd30ad27c0936420180e42d637009d90372f1e702f79847851831de071","dispatches":"[(\"s_copybfloat16bfloat16\", [20480, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [8, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [34048, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [87142400, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [20480, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [20, 532, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [43571200, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"gate_up","rows":34040,"k":2560,"n":1280,"left":true,"exact":true,"sha256":"d79b05579846810a34b955ef65535adba962d9d02a17bb0ecace423c6ce0791c","dispatches":"[(\"arangeint32\", [8, 1, 1], [8, 1, 1], true), (\"gg1_copyint32int32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyint32int32\", [8, 1, 1], [8, 1, 1], true), (\"s_copybfloat16bfloat16\", [20480, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [8, 1, 1], [8, 1, 1], true), (\"v_copyint32uint32\", [34048, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [87142400, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [20480, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmv_fast_bfloat16_t_gs_64_b_4\", [1, 160, 34048], [32, 2, 1], false), (\"g1_copybfloat16bfloat16\", [43571200, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"down","rows":34040,"k":640,"n":2560,"left":false,"exact":true,"sha256":"53768e60df9fad3dc06641f5245e6f67ac28e1278dc8fec8e62486858f2de341","dispatches":"[(\"s_copybfloat16bfloat16\", [5120, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [8, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [34048, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [21785600, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [5120, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [40, 532, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [87142400, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"down","rows":34040,"k":640,"n":2560,"left":true,"exact":true,"sha256":"f02f45549d38acd5979cdd768305b39268493449191a0a232514978a4015a658","dispatches":"[(\"arangeint32\", [8, 1, 1], [8, 1, 1], true), (\"gg1_copyint32int32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyint32int32\", [8, 1, 1], [8, 1, 1], true), (\"s_copybfloat16bfloat16\", [5120, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [8, 1, 1], [8, 1, 1], true), (\"v_copyint32uint32\", [34048, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [21785600, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [5120, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmv_bfloat16_t_gs_64_b_4\", [1, 320, 34048], [32, 2, 1], false), (\"g1_copybfloat16bfloat16\", [87142400, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"gate_up","rows":40950,"k":2560,"n":1280,"left":false,"exact":true,"sha256":"393e33c2576b26e23f6b91bf2ff2ea7e590322a8cfb88a237d7e90cfea74acc8","dispatches":"[(\"s_copybfloat16bfloat16\", [25600, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [40950, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [10, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [104832000, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [25600, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [20, 640, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [52416000, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"down","rows":40950,"k":640,"n":2560,"left":false,"exact":true,"sha256":"e66a8a49018a918e26e3f56a8859ab7894e3d75f0951858e43856fbe6052616e","dispatches":"[(\"s_copybfloat16bfloat16\", [6400, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [40950, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [10, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [26208000, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [6400, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [40, 640, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [104832000, 1, 1], [1024, 1, 1], true)]"}
+{"case":"projection","projection":"gate_up","rows":40960,"k":2560,"n":1280,"left":false,"exact":true,"sha256":"29b3628139c5f7c49d5b5b46134a3141958e9a9bc3360956a770ec3d5b65af5d","dispatches":"[(\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [20, 640, 1], [32, 2, 2], false)]"}
+{"case":"projection","projection":"down","rows":40960,"k":640,"n":2560,"left":false,"exact":true,"sha256":"d40d0c9c449d3ac4631a43f775815d1baea299dbfbee39c7585bad4997e61237","dispatches":"[(\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_64_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [40, 640, 1], [32, 2, 2], false)]"}

docs/issue110-sorted-gather/sentinel.jsonl

@@ -0,0 +1,9 @@
+{"case":"guarded_sentinel","tokens":3276,"rows":32760,"bad_rows":0,"shape":[32760,1,256],"dispatches":"[(\"arangeuint32\", [32760, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 512, 1], [32, 2, 2], false)]"}
+{"case":"sorted_gather","tokens":3277,"routed_rows":32770,"bad_rows":32,"first_bad":0,"last_bad":31,"dispatches":"[(\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 513, 1], [32, 2, 2], false)]"}
+{"case":"guarded_sentinel","tokens":3277,"rows":32770,"bad_rows":0,"shape":[32770,1,256],"dispatches":"[(\"s_copybfloat16bfloat16\", [15872, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [32770, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [62, 1, 1], [32, 1, 1], true), (\"arangeuint32\", [32832, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [8389120, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [15872, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 513, 1], [32, 2, 2], false), (\"g1_copybfloat16bfloat16\", [8389120, 1, 1], [1024, 1, 1], true)]"}
+{"case":"sorted_gather","tokens":3404,"routed_rows":34040,"bad_rows":1280,"first_bad":0,"last_bad":1279,"dispatches":"[(\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 532, 1], [32, 2, 2], false)]"}
+{"case":"guarded_sentinel","tokens":3404,"rows":34040,"bad_rows":0,"shape":[34040,1,256],"dispatches":"[(\"s_copybfloat16bfloat16\", [2048, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [34040, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [8, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [34048, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [8714240, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [2048, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 532, 1], [32, 2, 2], false)]"}
+{"case":"sorted_gather","tokens":4095,"routed_rows":40950,"bad_rows":8192,"first_bad":0,"last_bad":8191,"dispatches":"[(\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 640, 1], [32, 2, 2], false)]"}
+{"case":"guarded_sentinel","tokens":4095,"rows":40950,"bad_rows":0,"shape":[40950,1,256],"dispatches":"[(\"s_copybfloat16bfloat16\", [2560, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [40950, 1, 1], [1024, 1, 1], true), (\"gg1_copyuint32uint32\", [10, 1, 1], [8, 1, 1], true), (\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [10483200, 1, 1], [1024, 1, 1], true), (\"gg1_copybfloat16bfloat16\", [2560, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 640, 1], [32, 2, 2], false)]"}
+{"case":"sorted_gather","tokens":4096,"routed_rows":40960,"bad_rows":0,"first_bad":null,"last_bad":null,"dispatches":"[(\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 640, 1], [32, 2, 2], false)]"}
+{"case":"guarded_sentinel","tokens":4096,"rows":40960,"bad_rows":0,"shape":[40960,1,256],"dispatches":"[(\"arangeuint32\", [40960, 1, 1], [1024, 1, 1], true), (\"affine_gather_qmm_rhs_nax_nt_bfloat16_t_gs_32_b_4_bm_64_bn_64_bk_64_wm_2_wn_2\", [4, 640, 1], [32, 2, 2], false)]"}

src/engine/metal/qwen_mtplx.rs

@@ -39,6 +39,8 @@ mod issue106;
 mod issue108;
 #[cfg(test)]
 mod issue109;
+#[cfg(test)]
+mod issue110;
 mod linear;
 mod load;
 mod mlp;

src/engine/metal/qwen_mtplx/allocator.rs

@@ -230,6 +230,11 @@ impl Allocator {
         }
     }
 
+    #[cfg(test)]
+    pub(super) fn set_resource_limit_for_test(&self, limit: usize) -> usize {
+        std::mem::replace(&mut self.state.lock().unwrap().resource_limit, limit)
+    }
+
     pub(in super::super) fn stats(&self) -> Stats {
         let state = self.state.lock().unwrap();
         Stats {

src/engine/metal/qwen_mtplx/issue103_analysis.rs

@@ -250,13 +250,14 @@ fn issue103_sorted_gather_receipts() {
     super::configure_sources().unwrap();
     let _context = super::super::gpu::Context::open_qwen(0).unwrap();
     assert!(super::sparse_attention::nax_available());
-    let scope = super::encoder::Scope::new().unwrap();
+    let streams = Streams::new(Some(super::allocator::Allocator::new().unwrap()));
     let buffer =
         |count: usize, bits: u16| super::scalar_buffer(&bits.to_le_bytes().repeat(count)).unwrap();
     let w = super::scalar_buffer(&vec![0; 64 * 256 * 256 / 8 * 4]).unwrap();
     let scales = buffer(64 * 256 * 8, 0);
     let biases = buffer(64 * 256 * 8, 0x3f80); // all dequantized weights exactly 1
     for tokens in [3276_u32, 3277, 3404, 4095, 4096] {
+        let scope = super::encoder::Scope::new().unwrap();
         let rows = tokens * 10;
         let x = buffer(rows as usize * 256, 0x3f80);
         let lhs = super::scalar_buffer(&(0..rows).flat_map(u32::to_le_bytes).collect::<Vec<_>>())
@@ -295,6 +296,8 @@ fn issue103_sorted_gather_receipts() {
             "{}",
             serde_json::json!({"case":"sorted_gather","tokens":tokens,"routed_rows":rows,"bad_rows":bad.len(),"first_bad":bad.first(),"last_bad":bad.last(),"dispatches":format!("{dispatches:?}")})
         );
+        drop(scope);
+        super::issue110::sentinel(tokens, &streams);
         if tokens == 3276 || tokens == 4096 {
             assert!(bad.is_empty());
         } else {

src/engine/metal/qwen_mtplx/issue104.rs

@@ -12,7 +12,7 @@ fn fingerprint(execution: &Execution, array: &Array) -> serde_json::Value {
         super::super::qwen_mtplx_tests::hex_digest(&values.iter().flat_map(|v|v.to_le_bytes()).collect::<Vec<_>>())})
 }
 
-fn state(execution: &Execution, saved: &SessionSnapshot) -> serde_json::Value {
+pub(super) fn state(execution: &Execution, saved: &SessionSnapshot) -> serde_json::Value {
     let states = saved
         .trunk
         .iter()

src/engine/metal/qwen_mtplx/issue110.rs

@@ -0,0 +1,344 @@
+//! Isolated stock sorted-gather guard checks; the fused row-map port is #112.
+use super::array::{Array, Dtype};
+use super::stream::{Device, Streams};
+use super::{linear, ops};
+use crate::engine::metal::qwen_mtplx_tests::{hex_digest, pattern};
+
+fn bytes(a: &Array, streams: &Streams) -> Vec<u8> {
+    let gpu = streams.default_stream(Device::Gpu).unwrap();
+    let a = ops::contiguous(a, false, gpu).unwrap();
+    ops::evaluate(streams, std::slice::from_ref(&a), gpu, false).unwrap();
+    let mut data = vec![0; a.layout().nbytes()];
+    a.buffer().read(a.offset(), &mut data).unwrap();
+    data
+}
+
+fn unit_gather(rows: i32, streams: &Streams) -> Array {
+    let gpu = streams.default_stream(Device::Gpu).unwrap();
+    let array = |shape: &[i32], bits: u16| {
+        let count = shape.iter().map(|&n| n as usize).product::<usize>();
+        Array::new(
+            shape,
+            Dtype::BF16,
+            super::scalar_buffer(&bits.to_le_bytes().repeat(count)).unwrap(),
+        )
+        .unwrap()
+    };
+    let x = array(&[rows, 1, 256], 0x3f80);
+    let w = Array::new(
+        &[64, 256, 32],
+        Dtype::U32,
+        super::scalar_buffer(&vec![0; 64 * 256 * 32 * 4]).unwrap(),
+    )
+    .unwrap();
+    let s = array(&[64, 256, 8], 0);
+    let b = array(&[64, 256, 8], 0x3f80);
+    let rhs = Array::new(
+        &[rows],
+        Dtype::U32,
+        super::scalar_buffer(
+            &(0..rows as u32)
+                .flat_map(|r| (r * 64 / rows as u32).to_le_bytes())
+                .collect::<Vec<_>>(),
+        )
+        .unwrap(),
+    )
+    .unwrap();
+    linear::guarded_gather_qmm_bf16(
+        &x,
+        [&w, &s, &b],
+        [None, Some(&rhs)],
+        [4, 32],
+        true,
+        streams,
+        gpu,
+    )
+    .unwrap()
+}
+
+pub(super) fn sentinel(tokens: u32, streams: &Streams) {
+    let rows = (tokens * 10) as i32;
+    let out = unit_gather(rows, streams);
+    assert_eq!(out.layout().shape(), &[rows, 1, 256]);
+    let (data, dispatches) =
+        crate::engine::metal::qwen_mtplx_tests::capture_dispatches(|| bytes(&out, streams));
+    let bad = data
+        .chunks_exact(512)
+        .enumerate()
+        .filter_map(|(i, row)| {
+            row.chunks_exact(2)
+                .any(|b| u16::from_le_bytes(b.try_into().unwrap()) != 0x4380)
+                .then_some(i)
+        })
+        .collect::<Vec<_>>();
+    assert!(bad.is_empty(), "unwritten or incorrect rows: {bad:?}");
+    println!(
+        "{}",
+        serde_json::json!({"case":"guarded_sentinel","tokens":tokens,"rows":rows,"bad_rows":bad.len(),"shape":out.layout().shape(),"dispatches":format!("{dispatches:?}")})
+    );
+}
+
+#[test]
+#[ignore = "installed layer weights and pinned oracle receipts; external memory supervisor required"]
+fn issue110_actual_projection_rows_match_oracle() {
+    super::configure_sources().unwrap();
+    let _context = crate::engine::metal::gpu::Context::open_qwen(0).unwrap();
+    let streams = Streams::new(Some(super::allocator::Allocator::new().unwrap()));
+    let gpu = streams.default_stream(Device::Gpu).unwrap();
+    assert!(super::sparse_attention::nax_available());
+    let root = std::path::PathBuf::from(std::env::var_os("DS4_QWEN38_ARTIFACTS").unwrap());
+    let receipts =
+        std::fs::read_to_string(std::env::var_os("DS4_ISSUE110_ORACLE").unwrap()).unwrap();
+    let weights = super::weights::load_shards(&root, &streams).unwrap();
+    let parts = |name: &str| {
+        ["weight", "scales", "biases"].map(|part| {
+            weights[&format!("language_model.model.layers.0.mlp.switch_mlp.{name}.{part}")].clone()
+        })
+    };
+    let gate = parts("gate_proj");
+    let up = parts("up_proj");
+    let gu = std::array::from_fn::<_, 3, _>(|i| {
+        ops::concatenate(&[gate[i].clone(), up[i].clone()], 1, gpu).unwrap()
+    });
+    let down = parts("down_proj");
+    ops::evaluate(&streams, &[gu.to_vec(), down.to_vec()].concat(), gpu, false).unwrap();
+    for line in receipts.lines() {
+        let row: serde_json::Value = serde_json::from_str(line).unwrap();
+        if row["case"] != "projection" {
+            continue;
+        }
+        let rows = row["rows"].as_i64().unwrap() as i32;
+        let k = row["k"].as_i64().unwrap() as i32;
+        let n = row["n"].as_i64().unwrap() as i32;
+        let layer = if row["projection"] == "gate_up" {
+            &gu
+        } else {
+            &down
+        };
+        let x = Array::new(&[rows, 1, k], Dtype::BF16, pattern((rows * k) as u32, 7)).unwrap();
+        let rhs = Array::new(
+            &[rows],
+            Dtype::U32,
+            super::scalar_buffer(
+                &(0..rows as u32)
+                    .flat_map(|r| (r * 512 / rows as u32).to_le_bytes())
+                    .collect::<Vec<_>>(),
+            )
+            .unwrap(),
+        )
+        .unwrap();
+        let lhs = row["left"].as_bool().unwrap().then(|| {
+            Array::new(
+                &[rows],
+                Dtype::I32,
+                super::scalar_buffer(
+                    &(0..rows)
+                        .rev()
+                        .flat_map(i32::to_le_bytes)
+                        .collect::<Vec<_>>(),
+                )
+                .unwrap(),
+            )
+            .unwrap()
+        });
+        let out = linear::guarded_gather_qmm_bf16(
+            &x,
+            [&layer[0], &layer[1], &layer[2]],
+            [lhs.as_ref(), Some(&rhs)],
+            [4, 64],
+            true,
+            &streams,
+            gpu,
+        )
+        .unwrap();
+        assert_eq!(out.layout().shape(), &[rows, 1, n]);
+        let (data, dispatches) =
+            crate::engine::metal::qwen_mtplx_tests::capture_dispatches(|| bytes(&out, &streams));
+        assert_eq!(hex_digest(&data), row["sha256"].as_str().unwrap(), "{row}");
+        println!(
+            "{}",
+            serde_json::json!({"case":"projection","projection":row["projection"],"rows":rows,"k":k,"n":n,"left":lhs.is_some(),"exact":true,"sha256":hex_digest(&data),"dispatches":format!("{dispatches:?}")})
+        );
+    }
+}
+
+thread_local! {
+    static PADS: std::cell::RefCell<Option<Vec<serde_json::Value>>> = const { std::cell::RefCell::new(None) };
+}
+
+pub(super) fn record(rows: i32, pad: i32, k: i32, n: i32) {
+    PADS.with_borrow_mut(|slot| {
+        if let Some(receipts) = slot {
+            receipts.push(serde_json::json!({"rows":rows,"pad":pad,"k":k,"n":n}));
+        }
+    });
+}
+
+#[test]
+#[ignore = "full installed model and large partial chunks; serial memory supervisor required"]
+fn issue110_large_cold_and_warm_prefill() {
+    assert_eq!(
+        std::env::var("MTPLX_GDN_BOUNDARY_TAIL_INTERVAL").as_deref(),
+        Ok("0")
+    );
+    use super::session_cache::{DecodeOptions, TurnOptions, TurnProgress};
+    let root = std::path::PathBuf::from(std::env::var_os("DS4_QWEN38_ARTIFACTS").unwrap());
+    let directory = std::path::PathBuf::from(std::env::var_os("DS4_ISSUE110_DIR").unwrap());
+    std::fs::create_dir_all(&directory).unwrap();
+    let settings = crate::settings::effective_settings(
+        crate::model::ModelChoice::Qwen38FlashNext,
+        &crate::settings::GenerationPreferences {
+            context_tokens: 16384,
+            max_generated_tokens: 3,
+            temperature: Some(0.),
+            ..Default::default()
+        },
+        &crate::settings::RuntimePreferences::default(),
+        root.parent().unwrap(),
+    )
+    .unwrap();
+    let (mut execution, _) = super::execution::Execution::open(&settings.engine).unwrap();
+    let tokens = |len: usize| {
+        (0..len)
+            .map(|i| [9707, 11, 1075, 374, 264, 1092, 13][i % 7])
+            .collect::<Vec<_>>()
+    };
+    for chunk in [2048, 4096] {
+        let sampler = super::sampling::SamplerConfig {
+            temperature: 0.,
+            ..Default::default()
+        };
+        let options = TurnOptions {
+            prefill: super::prefill::Options {
+                chunk_size: Some(chunk),
+                final_logits_only: true,
+                external_cache_only: false,
+                external_emit_logits: true,
+                return_hidden: false,
+            },
+            decode: DecodeOptions::Ar(super::ar::Options {
+                max_tokens: 3,
+                seed: 1,
+                stop_tokens: &[],
+                sync_eval: false,
+                return_hidden: false,
+                capture_final_state: true,
+            }),
+            sampler,
+            draft_sampler: sampler,
+            lazy_kv: false,
+            store_min_suffix: None,
+            fused_max: 512,
+        };
+        let mut prompt = tokens(7374);
+        let mut saved = None;
+        for phase in ["cold", "warm"] {
+            if phase == "warm" && chunk == 4096 {
+                let original = super::issue104::state(&execution, saved.as_ref().unwrap());
+                let cancelled = execution.session_turn(
+                    &prompt,
+                    saved.as_ref(),
+                    &options,
+                    |event| {
+                        if matches!(
+                            event,
+                            TurnProgress::Prefill(super::prefill::Progress::Settled { .. })
+                        ) {
+                            Err("cancelled after padded chunk".into())
+                        } else {
+                            Ok(())
+                        }
+                    },
+                    |_| Ok(()),
+                    |_| Ok(()),
+                );
+                assert!(matches!(cancelled,Err(e) if e=="cancelled after padded chunk"));
+                assert_eq!(
+                    super::issue104::state(&execution, saved.as_ref().unwrap()),
+                    original
+                );
+                println!(
+                    "{}",
+                    serde_json::json!({"case":"cancel","chunk":chunk,"saved_state_unchanged":true})
+                );
+            }
+            PADS.with_borrow_mut(|slot| *slot = Some(Vec::new()));
+            let out = execution
+                .session_turn(
+                    &prompt,
+                    saved.as_ref(),
+                    &options,
+                    |event| {
+                        if let TurnProgress::Prefill(super::prefill::Progress::Settled {
+                            tokens,
+                            total,
+                        }) = event
+                        {
+                            eprintln!("#110 {chunk}/{phase} prefill {tokens}/{total}");
+                        }
+                        Ok(())
+                    },
+                    |_| Ok(()),
+                    |_| Ok(()),
+                )
+                .unwrap();
+            let pads = PADS.with_borrow_mut(Option::take).unwrap();
+            assert_eq!(pads.is_empty(), chunk == 2048);
+            assert!(out.final_snapshot_error.is_none());
+            let snapshot = out.snapshot.unwrap();
+            let state = super::issue104::state(&execution, &snapshot);
+            let record = serde_json::json!({"case":"turn","chunk":chunk,"phase":phase,"cached":out.cached_tokens,"prompt_len":prompt.len(),"state":state,"pads":pads});
+            std::fs::write(
+                directory.join(format!("{chunk}-{phase}.json")),
+                serde_json::to_vec(&record).unwrap(),
+            )
+            .unwrap();
+            println!(
+                "{}",
+                serde_json::json!({"case":"turn","chunk":chunk,"phase":phase,"cached":out.cached_tokens,"prompt_len":prompt.len(),"generated":&snapshot.token_ids[prompt.len()..],"padded_calls":pads.len()})
+            );
+            prompt = snapshot.token_ids.clone();
+            prompt.extend(tokens(3278));
+            saved = Some(snapshot);
+        }
+    }
+}
+
+#[test]
+#[ignore = "Metal allocation refusal and retry; external supervisor required"]
+fn issue110_guard_allocation_refusal_releases_graph_and_retries() {
+    super::configure_sources().unwrap();
+    let _context = crate::engine::metal::gpu::Context::open_qwen(0).unwrap();
+    for limit in [0, 1, 2] {
+        let allocator = super::allocator::Allocator::new().unwrap();
+        let streams = Streams::new(Some(allocator.clone()));
+        let gpu = streams.default_stream(Device::Gpu).unwrap();
+        let out = unit_gather(34040, &streams);
+        allocator.clear_cache();
+        let old = allocator.set_resource_limit_for_test(limit);
+        let result = ops::evaluate(&streams, std::slice::from_ref(&out), gpu, false);
+        assert!(matches!(result, Err(e) if e.contains("resource limit")));
+        let partial_peak = allocator.stats().peak;
+        assert_eq!(partial_peak > 0, limit > 0);
+        drop(out);
+        drop(streams);
+        allocator.set_resource_limit_for_test(old);
+        allocator.clear_cache();
+        assert_eq!(allocator.stats().active, 0);
+        let streams = Streams::new(Some(allocator.clone()));
+        let out = unit_gather(34040, &streams);
+        let data = bytes(&out, &streams);
+        assert!(
+            data.chunks_exact(2)
+                .all(|v| u16::from_le_bytes(v.try_into().unwrap()) == 0x4380)
+        );
+        drop(out);
+        drop(streams);
+        allocator.clear_cache();
+        assert_eq!(allocator.stats().active, 0);
+        println!(
+            "{}",
+            serde_json::json!({"case":"allocation_refusal","limit":limit,"partial_peak":partial_peak,"released":true,"retry_exact":true})
+        );
+    }
+}

src/engine/metal/qwen_mtplx/linear.rs

@@ -297,6 +297,157 @@ fn ensure_row_contiguous_matrix(input: &Array) -> Result<Array, String> {
     Ok(copy)
 }
 
+/// MTPLX moe_sorted_gather._call_pad for our pinned 0.32.2 affine,
+/// transposed primitive. Hardware truth matters here, not the route fallback
+/// switch: the stock MLX kernel does not read that switch.
+fn sorted_gather_pad(
+    x: &[i32],
+    lhs: Option<&[i32]>,
+    rhs: Option<&[i32]>,
+    sorted: bool,
+    nax: bool,
+) -> i32 {
+    let Some(&[rows]) = rhs else { return 0 };
+    if !sorted
+        || !nax
+        || rows <= 32767
+        || rows % 64 == 0
+        || x.len() < 2
+        || x[x.len() - 2] != 1
+        || x[0] != rows
+        || lhs.is_some_and(|shape| shape != [rows])
+    {
+        return 0;
+    }
+    64 - rows % 64
+}
+
+#[test]
+fn sorted_gather_guard_matches_pinned_call_predicates() {
+    for (rows, pad) in [
+        (1, 0),
+        (32767, 0),
+        (32768, 0),
+        (32769, 63),
+        (32770, 62),
+        (34040, 8),
+        (40950, 10),
+        (40960, 0),
+        (65535, 1),
+        (65537, 63),
+        (81910, 10),
+    ] {
+        assert_eq!(
+            sorted_gather_pad(&[rows, 1, 256], None, Some(&[rows]), true, true),
+            pad
+        );
+        assert_eq!(
+            sorted_gather_pad(&[rows, 1, 256], Some(&[rows]), Some(&[rows]), true, true),
+            pad
+        );
+        assert_eq!(
+            sorted_gather_pad(&[rows, 1, 256], None, Some(&[rows]), false, true),
+            0
+        );
+        assert_eq!(
+            sorted_gather_pad(&[rows, 1, 256], None, Some(&[rows]), true, false),
+            0
+        );
+    }
+    for (x, lhs, rhs) in [
+        (vec![34040, 2, 256], None, Some(vec![34040])),
+        (vec![1, 1, 256], None, Some(vec![34040])),
+        (vec![34040, 1, 256], None, None),
+        (vec![34040, 1, 256], None, Some(vec![1, 34040])),
+        (vec![34040, 1, 256], Some(vec![1]), Some(vec![34040])),
+        (vec![34040, 1, 256], Some(vec![1, 34040]), Some(vec![34040])),
+    ] {
+        assert_eq!(
+            sorted_gather_pad(&x, lhs.as_deref(), rhs.as_deref(), true, true),
+            0
+        );
+    }
+}
+
+/// Guard both stock Gate/Up and Down gathers, including future fused-path
+/// fallbacks. Keep the underlying runtime primitive unchanged.
+pub(super) fn guarded_gather_qmm_bf16(
+    x: &Array,
+    weights: [&Array; 3],
+    indices: [Option<&Array>; 2],
+    quant: [u32; 2],
+    sorted: bool,
+    streams: &Streams,
+    stream: Stream,
+) -> Result<Array, String> {
+    use super::{indexing, ops};
+    let shape = x.layout().shape().to_vec();
+    let pad = sorted_gather_pad(
+        &shape,
+        indices[0].map(|a| a.layout()).as_ref().map(|l| l.shape()),
+        indices[1].map(|a| a.layout()).as_ref().map(|l| l.shape()),
+        sorted,
+        super::sparse_attention::nax_available(),
+    );
+    if pad == 0 {
+        return gather_qmm_bf16(x, weights, indices, quant, sorted, streams, stream);
+    }
+    let rows = shape[0];
+    #[cfg(test)]
+    super::issue110::record(
+        rows,
+        pad,
+        *shape.last().unwrap(),
+        weights[0].layout().dim(-2)?,
+    );
+    let end = rows
+        .checked_add(pad)
+        .ok_or("sorted gather padding overflow")?;
+    let mut tail = shape.clone();
+    tail[0] = pad;
+    let x = ops::concatenate(
+        &[x.clone(), ops::zeros(&tail, x.layout().dtype(), stream)?],
+        0,
+        stream,
+    )?;
+    let rhs = indices[1].unwrap();
+    let last = ops::slice(rhs, &[rows - 1], &[rows], &[1], stream)?;
+    let rhs = ops::concatenate(
+        &[rhs.clone(), ops::broadcast_to(&last, &[pad], stream)?],
+        0,
+        stream,
+    )?;
+    let lhs = indices[0]
+        .map(|lhs| {
+            let tail = ops::astype(
+                &indexing::arange_range_i32(rows, end, stream)?,
+                lhs.layout().dtype(),
+                false,
+                stream,
+            )?;
+            ops::concatenate(&[lhs.clone(), tail], 0, stream)
+        })
+        .transpose()?;
+    let out = gather_qmm_bf16(
+        &x,
+        weights,
+        [lhs.as_ref(), Some(&rhs)],
+        quant,
+        sorted,
+        streams,
+        stream,
+    )?;
+    let mut stop = out.layout().shape().to_vec();
+    stop[0] = rows;
+    ops::slice(
+        &out,
+        &vec![0; stop.len()],
+        &stop,
+        &vec![1; stop.len()],
+        stream,
+    )
+}
+
 /// Installed affine/BF16, transposed GatherQMM from ops.cpp. Explicit index
 /// casts use the default stream, as indices_or_default does in the reference.
 pub(super) fn gather_qmm_bf16(

src/engine/metal/qwen_mtplx/moe.rs

@@ -27,7 +27,7 @@ fn gather(
     streams: &Streams,
     stream: Stream,
 ) -> Result<Array, String> {
-    linear::gather_qmm_bf16(
+    linear::guarded_gather_qmm_bf16(
         input,
         [&layer.weight, &layer.scales, &layer.biases],
         [None, Some(ids)],

src/engine/metal/qwen_mtplx/session_checkpoint.rs

@@ -32,6 +32,8 @@ struct Manifest {
     #[serde(default)]
     verify_gate: bool,
     #[serde(default)]
+    sorted_gather_guard: bool,
+    #[serde(default)]
     images: Vec<crate::engine::qwen_vision::Identity>,
     tokens: Vec<i32>,
     lazy_kv: bool,
@@ -119,6 +121,7 @@ impl LiveSession {
             let offset = file.stream_position().map_err(|e| e.to_string())?;
             let manifest = serde_json::to_vec(&Manifest {
                 verify_gate: true,
+                sorted_gather_guard: true,
                 identity,
                 tag,
                 images: self.current_images.clone(),
@@ -193,6 +196,14 @@ impl LiveSession {
                 "Qwen checkpoint uses obsolete MTP attention-gate math; rebuild the cache".into(),
             );
         }
+        // Old files do not record prefill chunk sizes or producer hardware.
+        // At ten routes/token this length can contain an affected chunk;
+        // shorter sessions cannot. Reject before allocating/restoring state.
+        if !manifest.sorted_gather_guard && manifest.tokens.len() >= 3277 {
+            return Err(
+                "Qwen checkpoint predates the sorted-gather row guard; rebuild the cache".into(),
+            );
+        }
         if !manifest.payload["mtp_history_snapshot"].is_null() != identity.mtp {
             return Err("Qwen checkpoint MTP payload differs".into());
         }
@@ -480,6 +491,40 @@ fn qwen_checkpoint_roundtrip_and_failed_restore_preserve_live_state() {
             .unwrap()
     );
     assert_eq!(fingerprint(&live), ar_original);
+    let bytes = fs::read(&path).unwrap();
+    let offset = u64::from_le_bytes(bytes[8..16].try_into().unwrap()) as usize;
+    let mut manifest: serde_json::Value = serde_json::from_slice(&bytes[offset..]).unwrap();
+    manifest
+        .as_object_mut()
+        .unwrap()
+        .remove("sorted_gather_guard");
+    let write_manifest = |manifest: &serde_json::Value| {
+        let mut old = bytes[..offset].to_vec();
+        old.extend(serde_json::to_vec(manifest).unwrap());
+        fs::write(&broken, old).unwrap();
+    };
+    write_manifest(&manifest);
+    assert!(
+        live.load_checkpoint(&broken, ar_identity, &streams, stream, |_| {}, || false)
+            .unwrap()
+    );
+    assert_eq!(fingerprint(&live), ar_original);
+    manifest["tokens"] = serde_json::json!(vec![1; 3277]);
+    manifest["identity"]["context"] = 4096.into();
+    write_manifest(&manifest);
+    let large_identity = Identity {
+        context: 4096,
+        ..ar_identity
+    };
+    assert!(
+        live.load_checkpoint(&broken, large_identity, &streams, stream, |_| {}, || false)
+            .unwrap_err()
+            .contains("sorted-gather row guard")
+    );
+    assert_eq!(fingerprint(&live), ar_original);
+    println!(
+        "Qwen sorted-gather checkpoint policy: old short state exact, possibly affected long state rejected before restore"
+    );
     println!(
         "Qwen checkpoint: exact roundtrip, cancellation/corruption controls, old MTP rejected and old AR restored"
     );

tools/issue110-model-reference.py

@@ -0,0 +1,53 @@
+"""Full installed-model cold/warm chunk-guard reference, with exact state checks."""
+import argparse,hashlib,json,os,subprocess,sys
+from pathlib import Path
+checkout,root,native,out=map(Path,sys.argv[1:5])
+assert subprocess.check_output(['git','-C',str(checkout),'rev-parse','HEAD'],text=True).strip()=='9882703f3105363ddc37eca9f97aa09a1d387112'
+sys.path.insert(0,str(checkout))
+from mtplx.server.openai import _server_runtime_env_overrides
+if os.environ.get('DS4_ISSUE110_STOCK_CONTROL')!='1':
+    os.environ.update(_server_runtime_env_overrides(argparse.Namespace(model=str(root),generation_mode='ar'),None))
+os.environ.update(MTPLX_COMPILED_GDN='1',MTPLX_QWEN4EXP_COMPILE='1',MTPLX_SUSTAINED_PREFILL='1')
+import mlx.core as mx
+import numpy as np
+from mlx_lm.utils import load_model
+from mtplx.models import qwen4_exp
+from mtplx.runtime import MTPLXRuntime,_load_tokenizer_resilient
+from mtplx.sampling import SamplerConfig
+from mtplx.generation import generate_ar
+from mtplx.cache_state import snapshot_cache
+from mtplx.session_bank import SessionBank
+from mtplx import moe_sorted_gather as msg
+from mtplx.kernels import moe_sorted_gather_nax as kernel
+assert mx.__version__=='0.32.2'
+assert os.environ.get('MTPLX_GDN_BOUNDARY_TAIL_INTERVAL')=='0'
+out.mkdir(parents=True,exist_ok=True)
+model,config=load_model(root,lazy=False,strict=True,get_model_classes=lambda **_:(qwen4_exp.Model,qwen4_exp.ModelArgs))
+model.post_weight_load(root)
+tokenizer=_load_tokenizer_resilient(root,config)
+rt=MTPLXRuntime(model,tokenizer,root,False,None)
+def fingerprint(a):
+    values=np.asarray(a.astype(mx.float32)).astype('<f4')
+    assert np.isfinite(values).all()
+    return dict(shape=list(a.shape),sha256=hashlib.sha256(values.tobytes()).hexdigest())
+for chunk in [2048,4096]:
+    os.environ['MTPLX_PREFILL_CHUNK_SIZE']=str(chunk)
+    bank=SessionBank()
+    for phase in ['cold','warm']:
+        expected=json.loads((native/f'{chunk}-{phase}.json').read_text())
+        prompt=expected['state']['ids'][:expected['prompt_len']]
+        guard_before=msg.stats()
+        result=generate_ar(rt,prompt,max_tokens=3,sampler=SamplerConfig(temperature=0.,top_p=.95,top_k=20),seed=1,
+            stop_token_ids=set(),capture_final_state=True,session_bank=bank,session_id='issue110',store_prefix_snapshot=False,
+            token_callback=lambda ids:print(json.dumps(dict(case='token',chunk=chunk,phase=phase,ids=ids)),flush=True))
+        final=result.final_state
+        assert final.safe_to_commit
+        states=[[None if a is None else fingerprint(a) for a in state] if state is not None else None for state in snapshot_cache(final.final_trunk_cache).states]
+        state=dict(ids=prompt+result.tokens,logits=fingerprint(final.final_logits),states=states)
+        (out/f'{chunk}-{phase}.json').write_text(json.dumps(state))
+        assert result.stats.cached_tokens==expected['cached']
+        assert state==expected['state'],(chunk,phase,'state differs')
+        print(json.dumps(dict(case='turn',chunk=chunk,phase=phase,exact=True,cached=result.stats.cached_tokens,tokens=result.tokens,guard_before=guard_before,guard_after=msg.stats(),kernel=kernel.stats())),flush=True)
+        if phase=='cold':
+            assert bank.put(runtime=rt,token_ids=state['ids'],cache=final.final_trunk_cache,logits=final.final_logits,hidden=None,
+                hidden_variant='post_norm',session_id='issue110',mtp_history_policy='cycle',snapshot_epoch=len(state['ids'])) is not None

tools/issue110-reference.py

@@ -0,0 +1,55 @@
+"""Pinned stock guard, installed projection geometry, and fused-oracle controls."""
+import hashlib,json,subprocess,sys
+from pathlib import Path
+checkout,root=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 mlx.nn as nn
+import numpy as np
+from mtplx import moe_sorted_gather as msg
+from mtplx.kernels import moe_sorted_gather_nax as kernel
+assert mx.__version__=='0.32.2' and msg.guard_active()
+index=json.loads((root/'model.safetensors.index.json').read_text())['weight_map']
+base='language_model.model.layers.0.mlp.switch_mlp.'
+weights=mx.load(str(root/index[base+'gate_proj.weight']))
+def parts(proj): return [weights[base+proj+'.'+part] for part in ['weight','scales','biases']]
+gu=[mx.concatenate(a,axis=1) for a in zip(parts('gate_proj'),parts('up_proj'))]
+down=parts('down_proj')
+mx.eval(*gu,*down)
+def pattern(rows,k):
+    v=((np.arange(rows*k,dtype=np.uint32)*17+7*13)%257).astype(np.int32)-128
+    return mx.array((v.astype(np.float32)/128).reshape(rows,1,k)).astype(mx.bfloat16)
+def digest(a): return hashlib.sha256(np.asarray(a.view(mx.uint16)).tobytes()).hexdigest()
+for rows in [20480,32760,32770,34040,40950,40960]:
+    for name,layer,k,n in [('gate_up',gu,2560,1280),('down',down,640,2560)]:
+        x=pattern(rows,k)
+        ids=mx.array(np.arange(rows,dtype=np.uint32)*512//rows)
+        for left in ([False,True] if rows==34040 else [False]):
+            lhs=mx.arange(rows-1,-1,-1,dtype=mx.int32) if left else None
+            before=msg.stats()
+            y=msg.gather_qmm(x,*layer,lhs_indices=lhs,rhs_indices=ids,group_size=64,bits=4,sorted_indices=True)
+            mx.eval(y)
+            assert y.shape==(rows,1,n)
+            print(json.dumps(dict(case='projection',projection=name,rows=rows,k=k,n=n,left=left,sha256=digest(y),pad=msg.pad_rows(rows),guard_before=before,guard_after=msg.stats())),flush=True)
+            del y
+        del x,ids
+        mx.clear_cache()
+for tokens in [2048,3277,4095,4096]:
+    rows=tokens*10
+    x=pattern(tokens,2560)
+    row_map=mx.arange(rows,dtype=mx.uint32)%tokens
+    ids=mx.array(np.arange(rows,dtype=np.uint32)*512//rows)
+    projection=msg.gather_qmm(x[row_map],*gu,rhs_indices=ids,group_size=64,bits=4,sorted_indices=True)
+    gate,up=mx.split(projection,2,axis=-1)
+    stock=nn.silu(gate)*up
+    before=kernel.stats()
+    fused=msg.swiglu_rows(x,row_map,*gu,ids,group_size=64,bits=4)
+    mx.eval(stock,fused)
+    assert digest(stock)==digest(fused)
+    after=kernel.stats()
+    assert after['calls']>before['calls'],(before,after)
+    assert after['canaries']>0 and after['canary_failures']==0,after
+    print(json.dumps(dict(case='optimized_oracle_control',tokens=tokens,rows=rows,exact=True,sha256=digest(stock),kernel_before=before,kernel_after=after)),flush=True)
+    del x,row_map,ids,projection,gate,up,stock,fused
+    mx.clear_cache()