1use crate::attention::{self, QwenAttnCfg};
10use crate::inference;
11use crate::kv_cache::KvCache;
12use crate::linear_core::{
13 GdnCfg, GdnWeights, ShortConvCfg, ShortConvWeights, VmfPhaseCfg, VmfPhaseWeights, gdn_forward,
14 gdn_pair, short_conv_forward, short_conv_forward_batch, short_conv_pair, vmf_phase_forward,
15 vmf_phase_pair,
16};
17use crate::pool::Pool;
18use crate::qtensor::QTensor;
19use crate::sampler::{self, SamplerConfig, SamplerScratch, SplitMix64};
20use crate::tokenizer::Tokenizer;
21use cortiq_core::mask::TaskMask;
22use cortiq_core::types::NormStyle;
23
24pub static GLOBAL_USE_GPU: std::sync::atomic::AtomicBool =
25 std::sync::atomic::AtomicBool::new(false);
26
27struct ForwardScratch {
31 n1: Vec<f32>,
32 n2: Vec<f32>,
33 p1: Vec<f32>,
34 p2: Vec<f32>,
35}
36
37impl ForwardScratch {
38 fn new(hidden: usize) -> Self {
39 Self {
40 n1: vec![0.0; hidden],
41 n2: vec![0.0; hidden],
42 p1: vec![0.0; hidden],
43 p2: vec![0.0; hidden],
44 }
45 }
46}
47
48pub struct Pipeline {
50 gpu_plan: Option<std::sync::Arc<Vec<(usize, usize, usize)>>>,
55 pub tokenizer: std::sync::Arc<Tokenizer>,
58 pub kv_cache: KvCache,
59 pub sampler_config: SamplerConfig,
60 pub weights: PipelineWeights,
61 pub hidden_size: usize,
62 pub intermediate_size: usize,
63 pub num_heads: usize,
64 pub num_kv_heads: usize,
65 pub head_dim: usize,
66 pub num_layers: usize,
68 pub physical_layers: usize,
70 pub loop_final_norm: bool,
72 pub vocab_size: usize,
73 pub rms_eps: f64,
74 pub rope_base: f32,
75 pub norm_style: NormStyle,
76 pub rotary_dim: usize,
78 pub attention_heads_per_layer: Option<Vec<usize>>,
80 pub vmf_cfg: Option<VmfPhaseCfg>,
82 pub gdn_cfg: Option<GdnCfg>,
84 pub logit_multiplier: Option<f32>,
86 pub cancel: std::sync::Arc<std::sync::atomic::AtomicBool>,
91 pub kv_history: Vec<u32>,
96 pub kda_cfg: Option<crate::linear_core::KdaCfg>,
98 pub g3n: Option<Box<(crate::g3n::G3nGlobals, Vec<crate::g3n::G3nLayer>)>>,
101 pub dsv4: Option<
105 Box<(
106 crate::dsv4::Dsv4Globals,
107 Vec<crate::dsv4::Dsv4Layer>,
108 crate::dsv4::Dsv4Cfg,
109 crate::dsv4::Dsv4State,
110 )>,
111 >,
112 pub dsv4_mtp: Vec<crate::dsv4::Dsv4Mtp>,
116 pub dspark: Option<crate::dsv4::DsparkState>,
118 pub dspark_pending: Vec<(usize, Vec<u32>, bool, usize)>,
121 pub dspark_hist: Vec<usize>,
123 pub dspark_real: Vec<u32>,
127 pub dspark_trunk_picks: Vec<Vec<(usize, Vec<usize>)>>,
131 pub dspark_exp: Vec<(usize, usize, usize, usize)>,
133 pub dspark_draft_ns: u128,
137 pub short_conv_cfg: Option<ShortConvCfg>,
140 pub mtp: Option<MtpModule>,
142 pub speculative: bool,
144 rng: SplitMix64,
145 sampler_scratch: SamplerScratch,
146 pub(crate) inv_freq: std::sync::Arc<Vec<f32>>,
150 ws: ForwardScratch,
154 pool: Option<std::sync::Arc<Pool>>,
156 pub(crate) model: Option<std::sync::Arc<cortiq_core::CmfModel>>,
160 pub(crate) dyn_force_f32: bool,
162 pub(crate) dyn_skill_layers: Vec<Option<Vec<usize>>>,
167 pub(crate) dyn_active: Option<usize>,
173 pub(crate) dyn_blend_loaded: bool,
177 pub(crate) dyn_phi_layer: Option<usize>,
180 dyn_phi_ema: Vec<f32>,
182 dyn_phi_seen: usize,
183 pub dyn_router: Option<crate::swarm::DynRouter>,
186 o1_cfg: Option<crate::nystrom::O1Cfg>,
189 o1_epoch: u64,
192 o1_flags: Vec<bool>,
194 trace: bool,
197 calib_temp: f32,
200 #[cfg_attr(not(target_os = "macos"), allow(dead_code))]
202 graph_kv_id: u64,
203 #[cfg_attr(not(target_os = "macos"), allow(dead_code))]
206 graph_want_logits: bool,
207 graph_logits: Option<Vec<f32>>,
210 pub embed_multiplier: f32,
212 pub attn_scale: f32,
215 pub swa: Option<(usize, usize)>,
218 pub sliding_layers: Option<Vec<bool>>,
221 pub inv_freq_local: Option<std::sync::Arc<Vec<f32>>>,
224 pub rotary_dim_local: Option<usize>,
225 pub rope_scale: f32,
226 pub rope_scale_local: f32,
227 pub global_attn: Option<(usize, usize)>,
230 pub inv_freq_global: Option<std::sync::Arc<Vec<f32>>>,
233 pub attn_v_norm: bool,
235 pub final_softcap: Option<f32>,
237 pub attn_softcap: f32,
239 confidence_on: bool,
243}
244
245#[cfg(target_os = "macos")]
246impl Drop for Pipeline {
247 fn drop(&mut self) {
248 crate::gpu::kv_mirror_drop(self.graph_kv_id);
249 }
250}
251
252pub struct PipelineWeights {
257 pub embed_tokens: QTensor,
259 pub layers: Vec<LayerWeights>,
261 pub lm_head: QTensor,
263 pub final_norm: Vec<f32>,
265}
266
267pub struct LayerWeights {
269 pub input_norm: Vec<f32>,
270 pub post_norm: Vec<f32>,
273 pub attn_out_norm: Option<Vec<f32>>,
276 pub layer_scale: Option<f32>,
278 pub ffn_out_norm: Option<Vec<f32>>,
281 pub ffn: FfnKind,
282 pub attn: AttnKind,
283}
284
285#[derive(Clone, Copy, PartialEq, Debug, Default)]
288pub enum Act {
289 #[default]
290 Silu,
291 GeluTanh,
292 Situ {
295 beta: f32,
296 linear_beta: f32,
297 },
298}
299
300impl Act {
301 pub fn from_arch(name: &str) -> Self {
302 if name == "gelu_tanh" {
303 Self::GeluTanh
304 } else {
305 Self::Silu
306 }
307 }
308
309 pub fn from_arch_full(arch: &cortiq_core::ModelArch) -> Self {
311 match arch.hidden_act.as_str() {
312 "situ" => Self::Situ {
313 beta: arch.activation_situ_beta.unwrap_or(1.0) as f32,
314 linear_beta: arch.activation_situ_linear_beta.unwrap_or(0.0) as f32,
315 },
316 other => Self::from_arch(other),
317 }
318 }
319
320 #[inline]
321 pub fn apply(self, x: f32) -> f32 {
322 match self {
323 Self::Silu => inference::silu(x),
324 Self::GeluTanh => inference::gelu_tanh(x),
325 Self::Situ { beta, .. } => beta * (x / beta).tanh() * (1.0 / (1.0 + (-x).exp())),
326 }
327 }
328
329 #[inline]
332 pub fn combine(self, g: f32, u: f32) -> f32 {
333 match self {
334 Self::Situ { linear_beta, .. } if linear_beta > 0.0 => {
335 self.apply(g) * (linear_beta * (u / linear_beta).tanh())
336 }
337 _ => self.apply(g) * u,
338 }
339 }
340}
341
342pub struct DenseFfn {
344 pub gate_proj: QTensor,
345 pub up_proj: QTensor,
346 pub down_proj: QTensor,
347 pub act: Act,
349}
350
351pub enum FfnKind {
354 Dense(DenseFfn),
355 Moe(MoeFfn),
359 DenseMoe(Box<DenseMoeFfn>),
366}
367
368pub struct DenseMoeFfn {
370 pub dense: DenseFfn,
371 pub moe: MoeFfn,
372 pub post_norm_1: Vec<f32>,
374 pub pre_norm_2: Vec<f32>,
377 pub post_norm_2: Vec<f32>,
379}
380
381pub struct MoeFfn {
382 pub router: QTensor,
384 pub experts: Vec<DenseFfn>,
385 pub top_k: usize,
386 pub norm_topk_prob: bool,
387 pub router_sigmoid: bool,
390 pub expert_bias: Option<Vec<f32>>,
394 pub routed_scaling: f32,
397 pub route_tau: Option<f32>,
403 pub shared: Option<(DenseFfn, Option<QTensor>)>,
406 pub stats: std::cell::RefCell<Vec<u64>>,
410 pub act_sq: std::cell::RefCell<Vec<f64>>,
417 pub act_rows: std::cell::RefCell<Vec<f32>>,
423 pub mask: Option<Vec<bool>>,
428 pub per_expert_scale: Option<Vec<f32>>,
431 pub router_input_norm: bool,
435}
436
437pub enum AttnKind {
440 Full {
442 wq: QTensor,
443 wk: QTensor,
444 wv: QTensor,
445 wo: QTensor,
446 q_norm: Option<Vec<f32>>,
447 k_norm: Option<Vec<f32>>,
448 output_gate: bool,
449 softplus_gate: Option<(QTensor, bool)>,
453 bias: Option<(Vec<f32>, Vec<f32>, Vec<f32>)>,
455 },
456 Linear(VmfPhaseWeights),
458 LinearGdn(GdnWeights),
460 ShortConv(ShortConvWeights),
463 Mla(Box<MlaWeights>),
471 Kda(Box<crate::linear_core::KdaWeights>),
475}
476
477pub struct MlaWeights {
479 pub q_proj: QTensor,
483 pub q_a: Option<QTensor>,
486 pub q_a_norm: Option<Vec<f32>>,
487 pub kv_a: QTensor,
489 pub kv_a_norm: Vec<f32>,
491 pub kv_b: QTensor,
493 pub o_proj: QTensor,
495 pub nh: usize,
496 pub qk_rope: usize,
497 pub qk_nope: usize,
498 pub v_dim: usize,
499 pub lora: usize,
500 pub scale: f32,
502 pub nope: bool,
504}
505
506pub struct MtpModule {
511 pub enorm: Vec<f32>,
512 pub hnorm: Vec<f32>,
513 pub eh_proj: QTensor,
515 pub layer: LayerWeights,
516 pub final_norm: Vec<f32>,
517 pub kv: crate::kv_cache::LayerKvCache,
518}
519
520pub struct GenerateResult {
522 pub text: String,
523 pub token_ids: Vec<u32>,
524 pub prompt_tokens: usize,
525 pub tokens_generated: usize,
526 pub finish_reason: String,
527 pub mtp_drafted: usize,
529 pub mtp_accepted: usize,
530 pub token_confidence: Vec<f32>,
535 pub traces: Vec<TokenTrace>,
538}
539
540#[derive(Clone, Debug)]
545pub struct TokenTrace {
546 pub t: usize,
548 pub token_id: u32,
550 pub confidence: f32,
552 pub active_skill: Option<String>,
554 pub recon: Option<f32>,
558 pub switched: bool,
561}
562
563fn top1_prob_t(logits: &[f32], id: u32, temp: f32) -> f32 {
568 let t = if temp > 1e-3 { temp } else { 1.0 };
569 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
570 let sum: f32 = logits.iter().map(|&v| ((v - max) / t).exp()).sum();
571 if sum > 0.0 {
572 (((logits[id as usize] - max) / t).exp()) / sum
573 } else {
574 0.0
575 }
576}
577
578fn prefill_batched() -> bool {
581 std::env::var("CMF_PREFILL")
582 .map(|v| v != "seq")
583 .unwrap_or(true)
584}
585
586#[derive(Clone, Copy)]
590enum PrefillIn<'a> {
591 Ids(&'a [u32]),
592 Hidden(&'a [f32]),
593}
594
595impl Pipeline {
602 fn can_prefill_batched(&self) -> bool {
603 prefill_batched() && !self.weights.layers.is_empty()
604 }
605}
606
607pub fn prefill_chunk() -> usize {
614 if let Some(n) = std::env::var("CMF_PREFILL_CHUNK")
615 .ok()
616 .and_then(|v| v.parse::<usize>().ok())
617 {
618 return n.max(1);
619 }
620 if cfg!(target_os = "macos") {
621 512
622 } else if cfg!(target_arch = "aarch64") {
623 256
626 } else {
627 48
628 }
629}
630
631pub type TokenCallback = Box<dyn FnMut(&str) -> bool + Send>;
633
634impl Pipeline {
635 #[inline]
639 pub fn phys_layer(&self, virtual_idx: usize) -> usize {
640 virtual_idx % self.physical_layers
641 }
642
643 #[inline]
646 pub fn is_loop_end(&self, virtual_idx: usize) -> bool {
647 self.loop_final_norm && (virtual_idx + 1) % self.physical_layers == 0
648 }
649
650 #[allow(clippy::too_many_arguments)]
652
653 #[cfg(target_os = "macos")]
672 fn graph_prefill_preferred(&self) -> bool {
673 if !crate::gpu::enabled_here()
674 || !crate::gpu::q1_force()
675 || std::env::var("CMF_GPU_BLOCK")
676 .map(|v| v == "0")
677 .unwrap_or(false)
678 {
679 return false;
680 }
681 self.weights
682 .layers
683 .iter()
684 .any(|lw| matches!(&lw.attn, AttnKind::LinearGdn(w) if w.in_proj_qkv.is_q1()))
685 }
686
687 #[cfg(not(target_os = "macos"))]
688 fn graph_prefill_preferred(&self) -> bool {
689 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Prefill);
697 if !graph_on || !crate::gpu::enabled_here() {
698 return false;
699 }
700 if self.o1_active() {
709 return false;
710 }
711 self.weights
712 .layers
713 .iter()
714 .any(|lw| matches!(&lw.attn, AttnKind::LinearGdn(_)))
715 }
716
717 #[cfg(target_os = "macos")]
718 fn q1_graph_gpu(
719 &mut self,
720 start: usize,
721 upto: Option<usize>,
722 position: usize,
723 h: &mut [f32],
724 ) -> usize {
725 use crate::gpu::{AttnGpuLayer, GdnGpuCfg, GdnGpuLayer, GraphDims, MetalFfn, TokenGraph};
726 if self.attn_softcap > 0.0 || !crate::gpu::enabled_here()
728 || !crate::gpu::q1_force()
729 || std::env::var("CMF_GPU_BLOCK")
730 .map(|v| v == "0")
731 .unwrap_or(false)
732 {
733 if std::env::var("CMF_GRAPH_DBG").is_ok() {
734 eprintln!(
735 "block-graph: front gate (softcap={} enabled_here={} q1_force={})",
736 self.attn_softcap > 0.0,
737 crate::gpu::enabled_here(),
738 crate::gpu::q1_force(),
739 );
740 }
741 return start;
742 }
743 if self.swa.is_some()
748 || self.global_attn.is_some()
749 || self.attention_heads_per_layer.is_some()
750 || self.attn_v_norm
751 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
752 || self.weights.layers.iter().any(|lw| {
753 lw.attn_out_norm.is_some()
754 || lw.ffn_out_norm.is_some()
755 || lw.layer_scale.is_some()
756 || matches!(&lw.ffn, FfnKind::Dense(d) if d.act != Act::Silu)
757 })
758 {
759 if std::env::var("CMF_GRAPH_DBG").is_ok() {
760 eprintln!(
761 "block-graph: arch ineligible (swa={} gattn={} hpl={} vnorm={} scale_delta={:.2e})",
762 self.swa.is_some(),
763 self.global_attn.is_some(),
764 self.attention_heads_per_layer.is_some(),
765 self.attn_v_norm,
766 (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs(),
767 );
768 }
769 return start;
770 }
771 let limit = upto
774 .map(|u| u + 1)
775 .unwrap_or(self.num_layers)
776 .min(self.num_layers);
777
778 enum Item<'a> {
779 Gdn {
780 run: Vec<GdnGpuLayer<'a>>,
781 first: usize,
782 },
783 Attn {
784 l: AttnGpuLayer<'a>,
785 li: usize,
786 q_norm: Option<&'a [f32]>,
787 k_norm: Option<&'a [f32]>,
788 output_gate: bool,
789 bias: Option<(&'a [f32], &'a [f32], &'a [f32])>,
790 full_gpu: bool,
793 },
794 }
795
796 let attend_mode = std::env::var("CMF_GPU_ATTEND").unwrap_or_else(|_| "auto".into());
803 let attend_contract = attend_mode != "0"
804 && attend_mode != "off"
805 && self.head_dim % 4 == 0
806 && self.head_dim <= 256
807 && self.rotary_dim >= 2
808 && self.rotary_dim <= self.head_dim
809 && (self.rotary_dim / 2) % 32 == 0
810 && self.num_kv_heads > 0
811 && self.num_heads % self.num_kv_heads == 0;
812
813 let mut plan: Vec<Item> = Vec::new();
814 let mut model_ref: Option<std::sync::Arc<cortiq_core::CmfModel>> = None;
815 let block_diag = std::env::var("CMF_GRAPH_DBG").is_ok();
817 let mut scan = start;
818 while scan < limit {
819 let lw = &self.weights.layers[self.phys_layer(scan)];
820 let ffn = match &lw.ffn {
821 FfnKind::Dense(d) => {
822 let (Some(g), Some(u), Some(dn)) = (
823 d.gate_proj.q1_parts(),
824 d.up_proj.q1_parts(),
825 d.down_proj.q1_parts(),
826 ) else {
827 if block_diag {
828 eprintln!(
829 "block-graph: L{scan} FFN trio not graph-mappable — run ends"
830 );
831 }
832 break;
833 };
834 MetalFfn::Dense {
835 gate: g,
836 up: u,
837 down: dn,
838 }
839 }
840 FfnKind::Moe(m) => {
841 let Some(moe) = metal_moe_graph_parts(m, self.hidden_size) else {
842 if block_diag {
843 eprintln!(
844 "block-graph: L{scan} MoE outside the graph contract — run ends"
845 );
846 }
847 break;
848 };
849 if let QTensor::Mapped { model, .. } = &m.experts[0].gate_proj {
850 model_ref.get_or_insert_with(|| model.clone());
851 }
852 MetalFfn::Moe(moe)
853 }
854 _ => {
855 if block_diag {
856 eprintln!("block-graph: L{scan} non-graph FFN — run ends");
857 }
858 break;
859 }
860 };
861 match &lw.attn {
862 AttnKind::LinearGdn(w) if self.gdn_cfg.is_some() => {
863 let parts = (
864 w.in_proj_qkv.q1_parts(),
865 w.in_proj_z.q1_parts(),
866 w.in_proj_a.f32_parts(),
867 w.in_proj_b.f32_parts(),
868 w.out_proj.q1_parts(),
869 );
870 let (Some(qkv), Some(z), Some(a), Some(b), Some(out)) = parts else {
871 if block_diag {
872 eprintln!(
873 "block-graph: L{scan} GDN parts refused (qkv={} z={} a_f32={} b_f32={} out={})",
874 w.in_proj_qkv.q1_parts().is_some(),
875 w.in_proj_z.q1_parts().is_some(),
876 w.in_proj_a.f32_parts().is_some(),
877 w.in_proj_b.f32_parts().is_some(),
878 w.out_proj.q1_parts().is_some(),
879 );
880 }
881 break;
882 };
883 if let QTensor::Mapped { model, .. } = &w.in_proj_qkv {
884 model_ref.get_or_insert_with(|| model.clone());
885 }
886 let gl = GdnGpuLayer {
887 attn_norm: &lw.input_norm,
888 post_norm: &lw.post_norm,
889 qkv,
890 z,
891 a,
892 b,
893 out,
894 ffn,
895 conv1d: &w.conv1d,
896 a_log: &w.a_log,
897 dt_bias: &w.dt_bias,
898 gnorm: &w.norm,
899 };
900 match plan.last_mut() {
901 Some(Item::Gdn { run, .. }) => run.push(gl),
902 _ => plan.push(Item::Gdn {
903 run: vec![gl],
904 first: scan,
905 }),
906 }
907 }
908 AttnKind::Full {
909 wq,
910 wk,
911 wv,
912 wo,
913 q_norm,
914 k_norm,
915 output_gate,
916 softplus_gate: None,
917 bias,
918 } if !self.kv_cache.layers[scan].o1_sealed() => {
919 let parts = (wq.q1_parts(), wk.q1_parts(), wv.q1_parts(), wo.q1_parts());
920 let (Some(pq), Some(pk), Some(pv), Some(po)) = parts else {
921 break;
922 };
923 if let QTensor::Mapped { model, .. } = wq {
924 model_ref.get_or_insert_with(|| model.clone());
925 }
926 let cache = &self.kv_cache.layers[scan];
927 let full_gpu = attend_contract
928 && cache.mode == crate::kv_cache::KvMode::F32
929 && cache.o1.is_none()
930 && bias.is_none()
931 && pq.1 == self.num_heads * self.head_dim * (1 + *output_gate as usize)
932 && pk.1 == self.num_kv_heads * self.head_dim
933 && pv.1 == self.num_kv_heads * self.head_dim
934 && po.2 == self.num_heads * self.head_dim;
935 plan.push(Item::Attn {
936 l: AttnGpuLayer {
937 attn_norm: &lw.input_norm,
938 post_norm: &lw.post_norm,
939 wq: pq,
940 wk: pk,
941 wv: pv,
942 wo: po,
943 ffn,
944 },
945 li: scan,
946 q_norm: q_norm.as_deref(),
947 k_norm: k_norm.as_deref(),
948 output_gate: *output_gate,
949 bias: bias
950 .as_ref()
951 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
952 full_gpu,
953 });
954 }
955 _ => break,
956 }
957 scan += 1;
958 }
959 let Some(model) = model_ref else {
960 if std::env::var("CMF_GRAPH_DBG").is_ok() {
961 eprintln!("q1-graph: no model ref (start {start}, scanned to {scan})");
962 }
963 return start;
964 };
965 if plan.is_empty() {
966 if std::env::var("CMF_GRAPH_DBG").is_ok() {
967 eprintln!("q1-graph: empty plan at layer {start}");
968 }
969 return start;
970 }
971 let has_moe = plan.iter().any(|it| match it {
972 Item::Gdn { run, .. } => run.iter().any(|l| matches!(l.ffn, MetalFfn::Moe(_))),
973 Item::Attn { l, .. } => matches!(l.ffn, MetalFfn::Moe(_)),
974 });
975 let dev_attend = attend_contract
976 && (self.head_dim <= 128
977 || has_moe
978 || attend_mode == "force"
979 || attend_mode == "256");
980 if !dev_attend {
981 for it in &mut plan {
982 if let Item::Attn { full_gpu, .. } = it {
983 *full_gpu = false;
984 }
985 }
986 }
987 if std::env::var("CMF_GRAPH_DBG").is_ok() {
988 use std::sync::atomic::{AtomicBool, Ordering};
989 static SAID: AtomicBool = AtomicBool::new(false);
990 if !SAID.swap(true, Ordering::Relaxed) {
991 let fg = plan
992 .iter()
993 .filter(|it| matches!(it, Item::Attn { full_gpu: true, .. }))
994 .count();
995 let att = plan
996 .iter()
997 .filter(|it| matches!(it, Item::Attn { .. }))
998 .count();
999 eprintln!(
1000 "q1-graph: plan of {} items from layer {start} to {scan} | dev_attend={dev_attend} full_gpu {fg}/{att} | hd={} rd={} nkv={} nh={}",
1001 plan.len(),
1002 self.head_dim,
1003 self.rotary_dim,
1004 self.num_kv_heads,
1005 self.num_heads,
1006 );
1007 }
1008 }
1009 let dims = GraphDims {
1010 hidden: self.hidden_size,
1011 eps: self.rms_eps as f32,
1012 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1013 };
1014 let Some(mut graph) = TokenGraph::new(&model, dims, h) else {
1015 return start;
1016 };
1017 let gcfg = self.gdn_cfg.map(|cfg| GdnGpuCfg {
1018 nv: cfg.num_v_heads,
1019 nk: cfg.num_k_heads,
1020 dk: cfg.key_head_dim,
1021 dv: cfg.value_head_dim,
1022 kk: cfg.conv_kernel,
1023 hidden: self.hidden_size,
1024 inter: self.intermediate_size,
1025 c_dim: cfg.conv_dim(),
1026 eps: cfg.rms_eps as f32,
1027 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1028 });
1029 let mut valid = 0usize;
1033 let mut end = start;
1034 for item in &plan {
1035 let ok = match item {
1036 Item::Gdn { run, .. } => gcfg
1037 .as_ref()
1038 .map(|gc| run.iter().all(|l| graph.gdn_ok(l, gc)))
1039 .unwrap_or(false),
1040 Item::Attn { l, .. } => graph.attn_ok(l),
1041 };
1042 if !ok {
1043 if block_diag {
1044 eprintln!(
1045 "block-graph: plan item {} ({}) failed graph preflight",
1046 valid,
1047 match item {
1048 Item::Gdn { run, first } =>
1049 format!("GDN run L{first}+{}", run.len()),
1050 Item::Attn { li, .. } => format!("Attn L{li}"),
1051 }
1052 );
1053 }
1054 break;
1055 }
1056 valid += 1;
1057 end += match item {
1058 Item::Gdn { run, .. } => run.len(),
1059 Item::Attn { .. } => 1,
1060 };
1061 }
1062 plan.truncate(valid);
1063 if plan.is_empty() {
1064 return start;
1065 }
1066
1067 let inv_freq = self.inv_freq.clone();
1068 let pool = self.pool.clone();
1069 let (nh, nkv, hd, hs, rd, eps) = (
1070 self.num_heads,
1071 self.num_kv_heads,
1072 self.head_dim,
1073 self.hidden_size,
1074 self.rotary_dim,
1075 self.rms_eps,
1076 );
1077 let norm_style = self.norm_style;
1078 let gemma = norm_style == cortiq_core::NormStyle::Gemma;
1079 let want = self.gdn_cfg.map(|c| c.state_len()).unwrap_or(0);
1080 let kv_id = self.graph_kv_id;
1081 let mut pending: Vec<(usize, usize)> = Vec::new();
1084 let mut dev_attn: Vec<usize> = Vec::new();
1087 for item in &plan {
1088 if self.loop_final_norm {
1090 let item_start = match item {
1091 Item::Gdn { first, .. } => *first,
1092 Item::Attn { li, .. } => *li,
1093 };
1094 if item_start > start && self.is_loop_end(item_start - 1) {
1095 graph.encode_loop_norm(&self.weights.final_norm);
1096 }
1097 }
1098 match item {
1099 Item::Gdn { run, first } => {
1100 for l in &mut self.kv_cache.layers[*first..*first + run.len()] {
1101 if l.linear_state.len() != want {
1102 l.linear_state = vec![0f32; want];
1103 }
1104 }
1105 let ro: Vec<&[f32]> = self.kv_cache.layers[*first..*first + run.len()]
1106 .iter()
1107 .map(|l| l.linear_state.as_slice())
1108 .collect();
1109 if !graph.encode_gdn_run(run, &ro, gcfg.as_ref().unwrap()) {
1110 tracing::error!("q1 graph: GDN run refused after validation");
1112 return start;
1113 }
1114 graph.commit();
1117 pending.push((*first, run.len()));
1118 }
1119 Item::Attn {
1120 l,
1121 li,
1122 q_norm,
1123 k_norm,
1124 output_gate,
1125 bias,
1126 full_gpu,
1127 } => {
1128 if *full_gpu {
1130 let cache = &self.kv_cache.layers[*li];
1131 let cpu_k: Vec<&[f32]> = (0..nkv).map(|g| cache.head_keys(g)).collect();
1132 let cpu_v: Vec<&[f32]> = (0..nkv).map(|g| cache.head_values(g)).collect();
1133 let cpu_stored = cpu_k[0].len() / hd;
1134 let p = crate::gpu::AttnDeviceParams {
1135 kv_id,
1136 layer: *li,
1137 nh,
1138 nkv,
1139 hd,
1140 rd,
1141 position,
1142 eps: eps as f32,
1143 gemma,
1144 output_gate: *output_gate,
1145 q_norm: *q_norm,
1146 k_norm: *k_norm,
1147 inv_freq: &inv_freq,
1148 cpu_k,
1149 cpu_v,
1150 cpu_stored,
1151 };
1152 if graph.attn_device_ok(l, &p) && graph.encode_attn_device(l, &p) {
1153 graph.commit();
1154 dev_attn.push(*li);
1155 continue;
1156 }
1157 }
1159 graph.encode_attn_prefix(l);
1160 graph.sync();
1161 if !pending.is_empty() {
1162 let idxs: Vec<usize> =
1163 pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1164 let mut outs: Vec<&mut [f32]> = self
1165 .kv_cache
1166 .layers
1167 .iter_mut()
1168 .enumerate()
1169 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1170 .map(|(_, s)| s.linear_state.as_mut_slice())
1171 .collect();
1172 graph.read_states(&mut outs);
1173 }
1174 let mut q_raw = attention::take_buf(l.wq.1);
1175 let mut k = attention::take_buf(l.wk.1);
1176 let mut v = attention::take_buf(l.wv.1);
1177 graph.read_qkv(&mut q_raw, &mut k, &mut v);
1178 let cfg = QwenAttnCfg {
1179 num_heads: nh,
1180 num_kv_heads: nkv,
1181 head_dim: hd,
1182 hidden_size: hs,
1183 position,
1184 inv_freq: &inv_freq,
1185 rotary_dim: rd,
1186 scale: self.attn_scale,
1187 softcap: self.attn_softcap,
1188 window: None,
1189 v_norm: false,
1190 q_norm: *q_norm,
1191 k_norm: *k_norm,
1192 output_gate: *output_gate,
1193 softplus_gate: None,
1194 rope_scale: 1.0,
1195 bias: *bias,
1196 rms_eps: eps,
1197 norm_style,
1198 pool: pool.as_deref(),
1199 };
1200 let mut ao = attention::qwen_attention_core(
1201 q_raw,
1202 k,
1203 v,
1204 &mut self.kv_cache.layers[*li],
1205 &cfg,
1206 );
1207 graph.encode_attn_suffix(l, &ao);
1208 graph.commit();
1211 attention::recycle_buf(&mut ao);
1212 }
1213 }
1214 }
1215 let mut lm_rows = None;
1220 if self.graph_want_logits
1221 && upto.is_none()
1222 && end == self.num_layers
1223 && std::env::var("CMF_GPU_LMHEAD")
1224 .map(|v| v != "0")
1225 .unwrap_or(true)
1226 {
1227 if let Some(lm) = self.weights.lm_head.q1_parts() {
1228 if graph.lm_head_ok(lm) {
1229 graph.encode_lm_head(&self.weights.final_norm, lm);
1230 lm_rows = Some(lm.1);
1231 }
1232 }
1233 }
1234 graph.sync();
1235 if !pending.is_empty() {
1236 let idxs: Vec<usize> = pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1237 let mut outs: Vec<&mut [f32]> = self
1238 .kv_cache
1239 .layers
1240 .iter_mut()
1241 .enumerate()
1242 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1243 .map(|(_, s)| s.linear_state.as_mut_slice())
1244 .collect();
1245 graph.read_states(&mut outs);
1246 }
1247 if let Some(rows) = lm_rows {
1248 let mut lg = attention::take_buf(rows.min(self.vocab_size));
1249 graph.read_logits(&mut lg);
1250 lg.resize(self.vocab_size, 0.0);
1251 if let Some(c) = self.final_softcap {
1252 for l in lg.iter_mut() {
1253 *l = c * (*l / c).tanh();
1254 }
1255 }
1256 self.graph_logits = Some(lg);
1257 }
1258 graph.finish(h);
1259 for li in dev_attn {
1263 let mut krow = attention::take_buf(nkv * hd);
1264 let mut vrow = attention::take_buf(nkv * hd);
1265 if crate::gpu::kv_mirror_read_last(kv_id, li, nkv, hd, &mut krow, &mut vrow) {
1266 let cache = &mut self.kv_cache.layers[li];
1267 cache.append(&krow, &vrow, &[]);
1268 let n = cache.seq_len;
1269 let mut imp = attention::take_buf(n);
1270 crate::gpu::kv_mirror_take_imp(kv_id, li, &mut imp);
1271 cache.accumulate_imp(&imp);
1272 attention::recycle_buf(&mut imp);
1273 }
1274 attention::recycle_buf(&mut krow);
1275 attention::recycle_buf(&mut vrow);
1276 }
1277 end
1278 }
1279
1280 pub fn new(
1281 tokenizer: Tokenizer,
1282 weights: PipelineWeights,
1283 hidden_size: usize,
1284 intermediate_size: usize,
1285 num_heads: usize,
1286 num_kv_heads: usize,
1287 head_dim: usize,
1288 num_layers: usize,
1289 physical_layers: usize,
1290 loop_final_norm: bool,
1291 vocab_size: usize,
1292 rms_eps: f64,
1293 rope_base: f32,
1294 norm_style: NormStyle,
1295 max_seq_len: usize,
1296 sampler_config: SamplerConfig,
1297 ) -> Self {
1298 let rng = match sampler_config.seed {
1299 Some(s) => SplitMix64::new(s),
1300 None => SplitMix64::from_entropy(),
1301 };
1302 let inv_freq = std::sync::Arc::new(attention::rope_inv_freq(head_dim, rope_base));
1303 let pool = Pool::from_env();
1304 if let Some(p) = &pool {
1305 tracing::info!("worker pool: {} threads", p.n_workers());
1306 }
1307 Self {
1308 gpu_plan: None,
1309 tokenizer: std::sync::Arc::new(tokenizer),
1310 kv_cache: KvCache::new(num_layers, num_kv_heads, head_dim, max_seq_len),
1311 sampler_config,
1312 weights,
1313 hidden_size,
1314 intermediate_size,
1315 num_heads,
1316 num_kv_heads,
1317 head_dim,
1318 num_layers,
1319 physical_layers,
1320 loop_final_norm,
1321 vocab_size,
1322 rms_eps,
1323 rope_base,
1324 norm_style,
1325 rotary_dim: head_dim,
1326 attention_heads_per_layer: None,
1327 vmf_cfg: None,
1328 gdn_cfg: None,
1329 kda_cfg: None,
1330 g3n: None,
1331 dsv4: None,
1332 dsv4_mtp: Vec::new(),
1333 dspark: None,
1334 dspark_pending: Vec::new(),
1335 dspark_hist: Vec::new(),
1336 dspark_real: Vec::new(),
1337 dspark_trunk_picks: Vec::new(),
1338 dspark_exp: Vec::new(),
1339 dspark_draft_ns: 0,
1340 logit_multiplier: None,
1341 cancel: std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)),
1342 kv_history: Vec::new(),
1343 short_conv_cfg: None,
1344 mtp: None,
1345 speculative: std::env::var("CMF_MTP").map(|v| v != "0").unwrap_or(true),
1346 rng,
1347 sampler_scratch: SamplerScratch::default(),
1348 inv_freq,
1349 ws: ForwardScratch::new(hidden_size),
1350 pool,
1351 model: None,
1352 dyn_force_f32: false,
1353 dyn_skill_layers: Vec::new(),
1354 dyn_active: None,
1355 dyn_blend_loaded: false,
1356 dyn_phi_layer: None,
1357 dyn_phi_ema: Vec::new(),
1358 dyn_phi_seen: 0,
1359 dyn_router: None,
1360 o1_cfg: None,
1361 o1_epoch: 0,
1362 o1_flags: Vec::new(),
1363 trace: false,
1364 calib_temp: 1.0,
1365 confidence_on: true,
1366 embed_multiplier: 1.0,
1367 attn_scale: 1.0 / (head_dim as f32).sqrt(),
1368 swa: None,
1369 sliding_layers: None,
1370 inv_freq_local: None,
1371 rotary_dim_local: None,
1372 rope_scale: 1.0,
1373 rope_scale_local: 1.0,
1374 global_attn: None,
1375 inv_freq_global: None,
1376 attn_v_norm: false,
1377 final_softcap: None,
1378 attn_softcap: 0.0,
1379 graph_want_logits: false,
1380 graph_logits: None,
1381 graph_kv_id: {
1382 static NEXT: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(1);
1383 NEXT.fetch_add(1, std::sync::atomic::Ordering::Relaxed)
1384 },
1385 }
1386 }
1387
1388 pub fn set_o1(&mut self, cfg: Option<crate::nystrom::O1Cfg>) {
1395 self.o1_flags = match &cfg {
1396 Some(c) => {
1397 let mut flags = c.layer_flags(self.num_layers);
1398 for (li, f) in flags.iter_mut().enumerate() {
1399 if *f
1400 && !matches!(
1401 self.weights.layers[self.phys_layer(li)].attn,
1402 AttnKind::Full { .. }
1403 )
1404 {
1405 *f = false;
1406 }
1407 }
1408 flags
1409 }
1410 None => Vec::new(),
1411 };
1412 if let Some(c) = &cfg {
1413 let n = self.o1_flags.iter().filter(|&&f| f).count();
1414 tracing::info!(
1415 "o1 nystrom attention: {n}/{} layer(s), m={} w={} sink={} rect={:?}",
1416 self.num_layers,
1417 c.m,
1418 c.w,
1419 c.sink,
1420 c.rect
1421 );
1422 }
1423 self.o1_cfg = cfg;
1424 }
1425
1426 pub fn o1_active(&self) -> bool {
1428 self.o1_cfg.is_some() && self.o1_flags.iter().any(|&f| f)
1429 }
1430
1431 pub fn o1_begin(&mut self) {
1436 if let Some(c) = &self.o1_cfg {
1437 let (m, w, sink, rect) = (c.m, c.w, c.sink, c.rect);
1438 for (li, &f) in self.o1_flags.iter().enumerate() {
1439 if f {
1440 self.kv_cache.layers[li].o1_begin(m, w, sink, rect);
1441 }
1442 }
1443 }
1444 }
1445
1446 pub fn o1_seal(&mut self) {
1450 self.o1_epoch = self.o1_epoch.wrapping_add(1);
1451 if self.o1_cfg.is_none() {
1452 return;
1453 }
1454 for li in 0..self.num_layers {
1455 if self.o1_flags.get(li).copied().unwrap_or(false) {
1456 self.kv_cache.layers[li].o1_seal(self.num_heads);
1457 }
1458 }
1459 }
1460
1461 pub fn set_trace(&mut self, on: bool) {
1463 self.trace = on;
1464 }
1465
1466 pub fn set_sampler_config(&mut self, config: SamplerConfig) {
1469 self.rng = match config.seed {
1470 Some(seed) => SplitMix64::new(seed),
1471 None => SplitMix64::from_entropy(),
1472 };
1473 self.sampler_config = config;
1474 }
1475
1476 pub fn set_confidence(&mut self, on: bool) {
1481 self.confidence_on = on;
1482 }
1483
1484 pub fn set_calib_temp(&mut self, t: f32) {
1487 self.calib_temp = if t > 1e-3 { t } else { 1.0 };
1488 }
1489
1490 pub fn calib_temp(&self) -> f32 {
1492 self.calib_temp
1493 }
1494
1495 pub fn set_rotary(&mut self, rotary_dim: usize, base: f32) {
1498 self.rotary_dim = rotary_dim.min(self.head_dim);
1499 self.inv_freq = std::sync::Arc::new(attention::rope_inv_freq(self.rotary_dim, base));
1500 }
1501
1502 fn attn_cfg(&self, position: usize) -> QwenAttnCfg<'_> {
1503 QwenAttnCfg {
1504 num_heads: self.num_heads,
1505 num_kv_heads: self.num_kv_heads,
1506 head_dim: self.head_dim,
1507 hidden_size: self.hidden_size,
1508 position,
1509 inv_freq: &self.inv_freq,
1510 rotary_dim: self.rotary_dim,
1511 scale: self.attn_scale,
1512 softcap: self.attn_softcap,
1513 window: None,
1514 v_norm: false,
1515 q_norm: None,
1516 k_norm: None,
1517 output_gate: false,
1518 softplus_gate: None,
1519 rope_scale: self.rope_scale,
1520 bias: None,
1521 rms_eps: self.rms_eps,
1522 norm_style: self.norm_style,
1523 pool: self.pool.as_deref(),
1524 }
1525 }
1526
1527 pub fn generate(
1529 &mut self,
1530 prompt: &str,
1531 max_tokens: usize,
1532 task_mask: Option<&TaskMask>,
1533 on_token: Option<TokenCallback>,
1534 ) -> Result<GenerateResult, String> {
1535 let input_ids = self.tokenizer.with_bos(self.tokenizer.encode(prompt));
1536 self.generate_from_ids(&input_ids, max_tokens, task_mask, on_token)
1537 }
1538
1539 pub fn generate_from_ids(
1547 &mut self,
1548 input_ids: &[u32],
1549 max_tokens: usize,
1550 task_mask: Option<&TaskMask>,
1551 mut on_token: Option<TokenCallback>,
1552 ) -> Result<GenerateResult, String> {
1553 if std::env::var("CMF_TRACE_H").is_ok() {
1554 eprintln!("input_ids: {input_ids:?}");
1555 }
1556 if input_ids.is_empty() {
1557 return Err("empty prompt: nothing to generate from".to_string());
1558 }
1559
1560 let reuse_from = {
1568 let on = !std::env::var("CMF_KV_REUSE").is_ok_and(|v| v == "0");
1569 let h = &self.kv_history;
1570 if on
1571 && task_mask.is_none()
1572 && self.mtp.is_none()
1573 && self.o1_cfg.is_none()
1574 && !h.is_empty()
1575 && h.len() < input_ids.len()
1576 && input_ids[..h.len()] == h[..]
1577 {
1578 h.len()
1579 } else {
1580 0
1581 }
1582 };
1583 if reuse_from == 0 {
1584 self.kv_cache.clear();
1586 self.kv_history.clear();
1587 crate::gpu::graph_kv_reset(self.graph_kv_id);
1588 } else if std::env::var("CMF_PREFILL_PROF").is_ok() {
1589 eprintln!(
1590 "kv-reuse: {} of {} prompt positions already cached",
1591 reuse_from,
1592 input_ids.len()
1593 );
1594 }
1595 crate::gpu::graph_race_begin_generation();
1596 self.o1_begin();
1597
1598 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
1604 let graph_spec = self.speculative
1620 && graph_on
1621 && self.mtp.is_some()
1622 && task_mask.is_none()
1623 && !self.o1_active()
1624 && self.sampler_config.temperature < 1e-6
1625 && self.sampler_config.repetition_penalty == 1.0
1626 && std::env::var("CMF_GRAPH_SPEC").is_ok_and(|v| v != "0");
1627 let pair_pays = self.gdn_cfg.is_none()
1634 || std::env::var("CMF_MTP").as_deref() == Ok("1");
1635 let spec_active = self.speculative
1636 && self.mtp.is_some()
1637 && task_mask.is_none()
1638 && !self.o1_active()
1639 && ((!graph_on && pair_pays) || graph_spec)
1640 && self.sampler_config.temperature < 1e-6;
1641 let mut mtp = if spec_active { self.mtp.take() } else { None };
1644 if std::env::var("CMF_MTP_CHAIN_PROBE").is_ok() {
1645 eprintln!(
1646 "mtp-probe gate: spec_active={spec_active} mtp={} speculative={} graph_on={graph_on} temp_ok={}",
1647 mtp.is_some(),
1648 self.speculative,
1649 self.sampler_config.temperature < 1e-6,
1650 );
1651 }
1652 if let Some(m) = &mut mtp {
1653 m.kv.clear();
1654 }
1655 let mut router = if mtp.is_none() {
1659 self.dyn_router.take()
1660 } else {
1661 None
1662 };
1663 if let Some(r) = &mut router {
1664 r.reset(); self.dyn_phi_seen = 0; let _ = self.set_active_skill(None);
1667 }
1668
1669 let mut all_ids = input_ids.to_vec();
1670 let mut generated = 0usize;
1671 let mut finish_reason = "max_tokens".to_string();
1672 let mut drafted = 0usize;
1673 let mut accepted = 0usize;
1674 let mut confidence: Vec<f32> = Vec::new();
1675 let trace_on = self.trace;
1676 let calib_temp = self.calib_temp;
1677 let mut traces: Vec<TokenTrace> = Vec::new();
1678
1679 let mut hidden = vec![0.0f32; self.hidden_size];
1685 let mut pos = reuse_from;
1686 let fuse_lm = mtp.is_none()
1695 && router.is_none()
1696 && std::env::var("CMF_GPU_LMHEAD").as_deref() != Ok("0");
1697 self.graph_logits = None;
1698 self.graph_want_logits = false;
1699 let _tpf = std::time::Instant::now();
1700 let batch_k = std::env::var("CMF_BATCH_K")
1701 .ok()
1702 .and_then(|v| v.parse::<usize>().ok())
1703 .unwrap_or(0);
1704 while self.dsv4.is_some()
1715 && mtp.is_none()
1716 && pos < input_ids.len()
1717 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1718 {
1719 let end = (pos + prefill_chunk()).min(input_ids.len());
1720 let ids: Vec<u32> = input_ids[pos..end].to_vec();
1721 let mut lg = Vec::new();
1722 if let Some(b) = &mut self.dsv4 {
1723 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
1724 crate::dsv4::forward_chunk(
1725 g,
1726 layers,
1727 &cfg,
1728 st,
1729 &ids,
1730 pos,
1731 &self.inv_freq,
1732 self.pool.as_deref(),
1733 &mut lg,
1734 end == input_ids.len(),
1735 );
1736 }
1737 if end == input_ids.len() {
1738 self.graph_logits = Some(lg);
1739 }
1740 pos = end;
1741 hidden = vec![0.0; self.hidden_size];
1742 }
1743 let dyn_prefill = router.is_some();
1748 let graph_prefill = self.graph_prefill_preferred();
1754 if task_mask.is_none()
1755 && !dyn_prefill
1756 && !graph_prefill
1757 && self.can_prefill_batched()
1758 && self.g3n.is_none()
1759 && input_ids.len() > 2
1760 {
1761 let chunk = prefill_chunk();
1767 let hs = self.hidden_size;
1768 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
1769 let end = (pos + chunk).min(input_ids.len());
1770 let hb = self.prefill_batch(&input_ids[pos..end], pos);
1771 if let Some(m) = &mut mtp {
1772 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1773 .ok()
1774 .and_then(|v| v.parse().ok())
1775 .unwrap_or(0);
1776 for p in pos..end {
1777 if p + 1 < input_ids.len() {
1778 if probe >= 1 && p + 2 < input_ids.len() {
1779 let (d1, mut hx) = self.mtp_step_h(
1783 m,
1784 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1785 input_ids[p + 1],
1786 p,
1787 );
1788 let mut ok = d1 == input_ids[p + 2];
1789 Self::chain_probe_note(0, ok);
1790 let mut d_prev = d1;
1791 let mut extra = 0usize;
1792 for j in 1..probe {
1793 if p + 2 + j >= input_ids.len() {
1794 break;
1795 }
1796 let (dj, hj) =
1797 self.mtp_step_h(m, &hx, d_prev, p + 1 + j);
1798 extra += 1;
1799 ok = ok && dj == input_ids[p + 2 + j];
1800 Self::chain_probe_note(j, ok);
1801 d_prev = dj;
1802 hx = hj;
1803 }
1804 m.kv.truncate_last(extra);
1805 } else {
1806 let _ = self.mtp_step(
1807 m,
1808 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1809 input_ids[p + 1],
1810 p,
1811 );
1812 }
1813 }
1814 }
1815 }
1816 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
1817 pos = end;
1818 }
1819 }
1820 let pair_off = std::env::var("CMF_PAIR").is_ok_and(|v| v == "0");
1821 if task_mask.is_none()
1822 && !dyn_prefill
1823 && !graph_prefill
1824 && !pair_off
1825 && self.pair_supported()
1826 {
1827 while pos + 1 < input_ids.len()
1828 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1829 {
1830 let e1 = self.embed_single(input_ids[pos]);
1831 let e2 = self.embed_single(input_ids[pos + 1]);
1832 let (h1, h2) = self.forward_pair(&e1, &e2, pos);
1833 self.commit_linear_scratch();
1835 if let Some(m) = &mut mtp {
1836 let _ = self.mtp_step(m, &h1, input_ids[pos + 1], pos);
1837 if pos + 2 < input_ids.len() {
1838 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1839 .ok()
1840 .and_then(|v| v.parse().ok())
1841 .unwrap_or(0);
1842 if probe >= 1 && pos + 3 < input_ids.len() {
1843 let (d1, mut hx) =
1847 self.mtp_step_h(m, &h2, input_ids[pos + 2], pos + 1);
1848 let mut ok = d1 == input_ids[pos + 3];
1849 Self::chain_probe_note(0, ok);
1850 let mut d_prev = d1;
1851 let mut extra = 0usize;
1852 for j in 1..probe {
1853 if pos + 3 + j >= input_ids.len() {
1854 break;
1855 }
1856 let (dj, hj) =
1857 self.mtp_step_h(m, &hx, d_prev, pos + 2 + j);
1858 extra += 1;
1859 ok = ok && dj == input_ids[pos + 3 + j];
1860 Self::chain_probe_note(j, ok);
1861 d_prev = dj;
1862 hx = hj;
1863 }
1864 m.kv.truncate_last(extra);
1865 } else {
1866 let _ = self.mtp_step(m, &h2, input_ids[pos + 2], pos + 1);
1867 }
1868 }
1869 }
1870 hidden = h2;
1871 pos += 2;
1872 }
1873 }
1874 if batch_k > 0
1883 && graph_prefill
1884 && task_mask.is_none()
1885 && !self.o1_active()
1886 && mtp.is_none()
1887 && !dyn_prefill
1888 && pos + 1 < input_ids.len()
1889 {
1890 let hs = self.hidden_size;
1891 let chunk = batch_k;
1892 while pos < input_ids.len() {
1893 let end = (pos + chunk).min(input_ids.len());
1894 let bk = end - pos;
1895 let mut hiddens = vec![0f32; bk * hs];
1896 for (j, &id) in input_ids[pos..end].iter().enumerate() {
1897 hiddens[j * hs..(j + 1) * hs].copy_from_slice(&self.embed_single(id));
1898 }
1899 let positions: Vec<usize> = (pos..end).collect();
1900 let t_chunk = std::time::Instant::now();
1901 let ok_b = self.try_batch_graph_wgpu(&mut hiddens, &positions, bk, None);
1902 if std::env::var("CMF_GRAPH_PROF").is_ok() {
1903 let ms = t_chunk.elapsed().as_secs_f64() * 1000.0;
1904 eprintln!(
1905 "batch-chunk: k={bk} ok={ok_b} {ms:.1} ms ({:.1} tok/s)",
1906 bk as f64 / (ms / 1000.0)
1907 );
1908 }
1909 {
1910 use std::sync::atomic::{AtomicBool, Ordering};
1911 static SAID: AtomicBool = AtomicBool::new(false);
1912 if !SAID.swap(true, Ordering::Relaxed) {
1913 if ok_b {
1914 tracing::info!("batched prefill: ACTIVE (k={bk})");
1915 } else {
1916 tracing::warn!("batched prefill declined — per-position graph");
1917 }
1918 }
1919 }
1920 if ok_b {
1921 hidden.copy_from_slice(&hiddens[(bk - 1) * hs..]);
1922 pos = end;
1923 } else {
1924 break; }
1926 }
1927 }
1928 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
1929 self.graph_want_logits = fuse_lm && pos + 1 == input_ids.len();
1930 hidden = self.forward_layers(&self.embed_single(input_ids[pos]), pos, task_mask);
1931 if let Some(m) = &mut mtp {
1932 if pos + 1 < input_ids.len() {
1933 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1939 .ok()
1940 .and_then(|v| v.parse().ok())
1941 .unwrap_or(0);
1942 if probe >= 1 && pos + 2 < input_ids.len() {
1943 let (d1, mut hx) = self.mtp_step_h(m, &hidden, input_ids[pos + 1], pos);
1944 let mut ok = d1 == input_ids[pos + 2];
1945 Self::chain_probe_note(0, ok);
1946 let mut d_prev = d1;
1947 let mut extra = 0usize;
1948 for j in 1..probe {
1949 if pos + 2 + j >= input_ids.len() {
1950 break;
1951 }
1952 let (dj, hj) =
1953 self.mtp_step_h(m, &hx, d_prev, pos + 1 + j);
1954 extra += 1;
1955 ok = ok && dj == input_ids[pos + 2 + j];
1956 Self::chain_probe_note(j, ok);
1957 d_prev = dj;
1958 hx = hj;
1959 }
1960 m.kv.truncate_last(extra);
1963 } else {
1964 let _ = self.mtp_step(m, &hidden, input_ids[pos + 1], pos);
1965 }
1966 }
1967 }
1968 pos += 1;
1969 }
1970 if std::env::var("CMF_PREFILL_PROF").is_ok() {
1971 eprintln!(
1972 "prefill: {} tokens in {:.1} ms (batch_k={batch_k})",
1973 input_ids.len(),
1974 _tpf.elapsed().as_secs_f64() * 1000.0
1975 );
1976 }
1977 if self
1980 .cancel
1981 .swap(false, std::sync::atomic::Ordering::Relaxed)
1982 {
1983 self.kv_history.clear();
1984 if let Some(m) = mtp {
1985 self.mtp = Some(m);
1986 }
1987 return Ok(GenerateResult {
1988 text: String::new(),
1989 token_ids: Vec::new(),
1990 prompt_tokens: input_ids.len(),
1991 tokens_generated: 0,
1992 finish_reason: "cancelled".to_string(),
1993 mtp_drafted: 0,
1994 mtp_accepted: 0,
1995 token_confidence: Vec::new(),
1996 traces: Vec::new(),
1997 });
1998 }
1999
2000 self.o1_seal();
2003
2004 macro_rules! commit {
2006 ($id:expr) => {{
2007 all_ids.push($id);
2008 generated += 1;
2009 if self.tokenizer.is_eos($id) {
2010 finish_reason = "stop".to_string();
2011 false
2012 } else {
2013 let token_text = self.tokenizer.decode_token($id);
2014 let mut go = true;
2015 if let Some(ref mut cb) = on_token {
2016 if !cb(&token_text) {
2017 finish_reason = "cancelled".to_string();
2018 go = false;
2019 }
2020 }
2021 go
2022 }
2023 }};
2024 }
2025
2026 let mut next_pos = input_ids.len();
2028 'decode: while generated < max_tokens {
2029 if self
2030 .cancel
2031 .swap(false, std::sync::atomic::Ordering::Relaxed)
2032 {
2033 finish_reason = "cancelled".to_string();
2034 break 'decode;
2035 }
2036 let mut logits = match self.graph_logits.take() {
2037 Some(lg) => lg,
2038 None => {
2039 inference::rms_norm_into(
2040 &hidden,
2041 &self.weights.final_norm,
2042 self.rms_eps,
2043 self.norm_style,
2044 &mut self.ws.n1,
2045 );
2046 self.lm_head_forward(&self.ws.n1)
2047 }
2048 };
2049 let t_next = sampler::sample_with_scratch(
2050 &logits,
2051 &self.sampler_config,
2052 &all_ids,
2053 &mut self.rng,
2054 &mut self.sampler_scratch,
2055 );
2056 if self.confidence_on {
2057 confidence.push(top1_prob_t(&logits, t_next, calib_temp));
2058 }
2059 attention::recycle_buf(&mut logits);
2060 if trace_on {
2061 let skill = router.as_ref().and_then(|r| r.active_id());
2065 traces.push(TokenTrace {
2066 t: generated,
2067 token_id: t_next,
2068 confidence: confidence.last().copied().unwrap_or(0.0),
2069 active_skill: skill,
2070 recon: None,
2071 switched: false,
2072 });
2073 }
2074 if !commit!(t_next) {
2075 break 'decode;
2076 }
2077 if generated >= max_tokens {
2078 break 'decode;
2079 }
2080
2081 if self.kv_cache.needs_eviction() {
2082 static SAID: std::sync::Once = std::sync::Once::new();
2088 SAID.call_once(|| {
2089 tracing::warn!(
2090 "KV cache full at {} positions — evicting half; quality \
2091 will degrade. Raise CMF_MAX_SEQ.",
2092 self.kv_cache.max_seq_len,
2093 );
2094 });
2095 let keep = (self.kv_cache.max_seq_len / 2).max(1);
2096 self.kv_cache.evict(keep);
2097 }
2098
2099 match &mut mtp {
2100 #[cfg(feature = "gpu")]
2102 Some(m) if graph_spec && generated + 1 < max_tokens && next_pos > 0 => {
2103 if let Some((extra, n_pos, new_h)) = self.graph_spec_step(
2104 m,
2105 &hidden,
2106 t_next,
2107 next_pos,
2108 &mut drafted,
2109 &mut accepted,
2110 ) {
2111 next_pos = n_pos;
2112 hidden = new_h;
2113 let mut stopped = false;
2114 for &id in &extra {
2115 if self.confidence_on {
2116 confidence.push(0.0);
2117 }
2118 if !commit!(id) {
2119 stopped = true;
2120 break;
2121 }
2122 }
2123 if stopped {
2124 break 'decode;
2125 }
2126 continue 'decode;
2127 }
2128 hidden = self.forward_layers(&self.embed_single(t_next), next_pos, task_mask);
2130 next_pos += 1;
2131 continue 'decode;
2132 }
2133 Some(m) if !graph_spec && generated + 1 < max_tokens => {
2135 let draft = self.mtp_step(m, &hidden, t_next, next_pos - 1);
2136 drafted += 1;
2137 let emb1 = self.embed_single(t_next);
2138 let emb2 = self.embed_single(draft);
2139 let (h1, h2) = self.forward_pair(&emb1, &emb2, next_pos);
2140
2141 inference::rms_norm_into(
2142 &h1,
2143 &self.weights.final_norm,
2144 self.rms_eps,
2145 self.norm_style,
2146 &mut self.ws.n1,
2147 );
2148 let mut logits1 = self.lm_head_forward(&self.ws.n1);
2149 let t_after = sampler::sample_with_scratch(
2150 &logits1,
2151 &self.sampler_config,
2152 &all_ids,
2153 &mut self.rng,
2154 &mut self.sampler_scratch,
2155 );
2156 if self.confidence_on {
2157 confidence.push(top1_prob_t(&logits1, t_after, calib_temp));
2158 }
2159 attention::recycle_buf(&mut logits1);
2160 if trace_on {
2161 traces.push(TokenTrace {
2164 t: generated,
2165 token_id: t_after,
2166 confidence: confidence.last().copied().unwrap_or(0.0),
2167 active_skill: None,
2168 recon: None,
2169 switched: false,
2170 });
2171 }
2172 let stop = !commit!(t_after);
2173
2174 if t_after == draft {
2175 accepted += 1;
2176 self.commit_linear_scratch();
2177 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2178 hidden = h2;
2179 next_pos += 2;
2180 } else {
2181 for layer in &mut self.kv_cache.layers {
2183 layer.truncate_last(1);
2184 }
2185 if !stop {
2186 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2187 hidden = self.forward_layers(
2188 &self.embed_single(t_after),
2189 next_pos + 1,
2190 None,
2191 );
2192 }
2193 next_pos += 2;
2194 }
2195 if stop {
2196 break 'decode;
2197 }
2198 }
2199 _ => {
2201 #[cfg(feature = "gpu")]
2206 if Self::dsv4_spec_on() && self.dsv4.is_some() {
2207 static SAID: std::sync::Once = std::sync::Once::new();
2208 SAID.call_once(|| {
2209 eprintln!(
2210 "dsv4-spec гейт: mtp={} mask={} router={} trace={} temp={} rep={} ",
2211 !self.dsv4_mtp.is_empty(),
2212 task_mask.is_none(),
2213 router.is_none(),
2214 !trace_on,
2215 self.sampler_config.temperature < 1e-6,
2216 self.sampler_config.repetition_penalty == 1.0,
2217 );
2218 });
2219 }
2220 #[cfg(feature = "gpu")]
2221 if Self::dsv4_spec_on()
2222 && self.dsv4.is_some()
2223 && !self.dsv4_mtp.is_empty()
2224 && task_mask.is_none()
2225 && router.is_none()
2226 && !trace_on
2227 && self.sampler_config.temperature < 1e-6
2228 && self.sampler_config.repetition_penalty == 1.0
2229 && generated + 1 < max_tokens
2230 && all_ids.len() >= 2
2231 {
2232 let tip_token = all_ids[all_ids.len() - 2];
2233 if let Some((extra, n_pos)) = self.dsv4_spec_step(
2234 tip_token,
2235 t_next,
2236 next_pos,
2237 &mut drafted,
2238 &mut accepted,
2239 ) {
2240 next_pos = n_pos;
2241 let mut stopped = false;
2242 for &id in &extra {
2243 if self.confidence_on {
2244 confidence.push(0.0);
2245 }
2246 if !commit!(id) {
2247 stopped = true;
2248 break;
2249 }
2250 }
2251 if stopped {
2252 break 'decode;
2253 }
2254 continue 'decode;
2255 }
2256 }
2257 self.graph_want_logits = fuse_lm;
2258 let mut t_fwd = t_next;
2264 let pure_greedy = self.sampler_config.temperature < 1e-6
2265 && self.sampler_config.repetition_penalty == 1.0
2266 && self.sampler_config.suppress_tokens.is_empty();
2267 let burst_k = std::env::var("CMF_MULTISTEP")
2272 .ok()
2273 .and_then(|v| v.parse::<usize>().ok())
2274 .unwrap_or(0);
2275 if pure_greedy
2276 && burst_k >= 1
2277 && fuse_lm
2278 && task_mask.is_none()
2279 && router.is_none()
2280 && !trace_on
2281 && !self.confidence_on
2282 {
2283 let mut stopped = false;
2284 loop {
2285 let room = max_tokens.saturating_sub(generated);
2286 if room <= 2 {
2287 break;
2288 }
2289 let k = burst_k.min(room - 1);
2290 if k < 1 {
2291 break;
2292 }
2293 let Some(ids) = self.try_multi_burst(t_fwd, next_pos, k) else {
2294 break;
2295 };
2296 next_pos += k;
2297 for &id in &ids {
2298 if !commit!(id) {
2299 stopped = true;
2300 break;
2301 }
2302 }
2303 if stopped {
2304 break;
2305 }
2306 t_fwd = *ids.last().unwrap();
2307 }
2308 if stopped {
2309 break 'decode;
2310 }
2311 }
2312 hidden = self.forward_layers(&self.embed_single(t_fwd), next_pos, task_mask);
2313 next_pos += 1;
2314 if let Some(r) = &mut router {
2317 let phi = self.dyn_phi_ema.clone();
2318 let decision = r.step(&phi, generated);
2319 if let Some(new_active) = decision {
2320 let _ = self.set_active_skill(new_active);
2321 }
2322 if trace_on {
2325 if let Some(last) = traces.last_mut() {
2326 let e = r.last_best_e();
2327 last.recon = e.is_finite().then_some(e);
2328 last.switched = decision.is_some();
2329 }
2330 }
2331 }
2332 }
2333 }
2334 }
2335
2336 self.graph_want_logits = false;
2337 self.graph_logits = None;
2338 if router.is_some() {
2340 let _ = self.set_active_skill(None);
2341 }
2342 self.dyn_router = router.or(self.dyn_router.take());
2343 self.mtp = mtp.or(self.mtp.take());
2344
2345 let output_ids = &all_ids[input_ids.len()..];
2346 let forwarded = input_ids.len() + output_ids.len().saturating_sub(1);
2350 self.kv_history = all_ids[..forwarded.min(all_ids.len())].to_vec();
2351 confidence.truncate(output_ids.len()); traces.truncate(output_ids.len());
2353 Ok(GenerateResult {
2354 text: self.tokenizer.decode(output_ids),
2355 token_ids: output_ids.to_vec(),
2356 prompt_tokens: input_ids.len(),
2357 tokens_generated: generated,
2358 finish_reason,
2359 mtp_drafted: drafted,
2360 mtp_accepted: accepted,
2361 token_confidence: confidence,
2362 traces,
2363 })
2364 }
2365
2366 fn mtp_step(
2370 &mut self,
2371 m: &mut MtpModule,
2372 hidden: &[f32],
2373 next_token: u32,
2374 position: usize,
2375 ) -> u32 {
2376 self.mtp_step_h(m, hidden, next_token, position).0
2377 }
2378
2379 fn chain_probe_note(depth: usize, prefix_ok: bool) {
2383 use std::sync::Mutex;
2384 static T: Mutex<Vec<(u64, u64)>> = Mutex::new(Vec::new());
2385 let mut t = T.lock().unwrap();
2386 if t.len() <= depth {
2387 t.resize(depth + 1, (0, 0));
2388 }
2389 t[depth].0 += 1;
2390 t[depth].1 += prefix_ok as u64;
2391 if depth == 0 && t[0].0 % 128 == 0 {
2392 let line: Vec<String> = t
2393 .iter()
2394 .enumerate()
2395 .map(|(d, (n, k))| format!("d{}={:.0}%({n})", d + 1, 100.0 * *k as f64 / (*n).max(1) as f64))
2396 .collect();
2397 eprintln!("mtp-chain: {}", line.join(" "));
2398 }
2399 }
2400
2401 fn mtp_step_h(
2405 &mut self,
2406 m: &mut MtpModule,
2407 hidden: &[f32],
2408 next_token: u32,
2409 position: usize,
2410 ) -> (u32, Vec<f32>) {
2411 let e = self.embed_single(next_token);
2415 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2416 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2417 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2418 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2419 let mut x = vec![0.0f32; self.hidden_size];
2420 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2421
2422 let lw = &m.layer;
2424 inference::rms_norm_into(
2425 &x,
2426 &lw.input_norm,
2427 self.rms_eps,
2428 self.norm_style,
2429 &mut self.ws.n1,
2430 );
2431 let attn = match &lw.attn {
2432 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2434 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2435 AttnKind::Full {
2436 wq,
2437 wk,
2438 wv,
2439 wo,
2440 q_norm,
2441 k_norm,
2442 output_gate,
2443 softplus_gate,
2444 bias,
2445 } => {
2446 let mut cfg = self.attn_cfg(position);
2447 cfg.q_norm = q_norm.as_deref();
2448 cfg.k_norm = k_norm.as_deref();
2449 cfg.output_gate = *output_gate;
2450 cfg.softplus_gate = softplus_gate
2451 .as_ref()
2452 .map(|(gate, per_head)| (gate, *per_head));
2453 cfg.bias = bias
2454 .as_ref()
2455 .map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2456 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2457 }
2458 AttnKind::Linear(_) | AttnKind::LinearGdn(_) | AttnKind::ShortConv(_) => {
2459 unreachable!("MTP block is full attention")
2460 }
2461 };
2462 for (i, &a) in attn.iter().enumerate() {
2463 x[i] += a;
2464 }
2465 inference::rms_norm_into(
2466 &x,
2467 &lw.post_norm,
2468 self.rms_eps,
2469 self.norm_style,
2470 &mut self.ws.p1,
2471 );
2472 let ffn = ffn_forward(&lw.ffn, &self.ws.p1, self.pool.as_deref(), None);
2473 for (i, &f) in ffn.iter().enumerate() {
2474 x[i] += f;
2475 }
2476
2477 inference::rms_norm_into(
2478 &x,
2479 &m.final_norm,
2480 self.rms_eps,
2481 self.norm_style,
2482 &mut self.ws.n1,
2483 );
2484 let mut lg = self.lm_head_forward(&self.ws.n1);
2485 let draft = sampler::argmax(&lg);
2486 attention::recycle_buf(&mut lg);
2487 (draft, x)
2488 }
2489
2490 fn mtp_warm(&mut self, m: &mut MtpModule, hidden: &[f32], next_token: u32, position: usize) {
2494 let e = self.embed_single(next_token);
2495 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2496 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2497 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2498 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2499 let mut x = vec![0.0f32; self.hidden_size];
2500 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2501 inference::rms_norm_into(&x, &m.layer.input_norm, self.rms_eps, self.norm_style, &mut self.ws.n1);
2502 let attn = match &m.layer.attn {
2503 AttnKind::Full { wq, wk, wv, wo, q_norm, k_norm, output_gate, softplus_gate, bias } => {
2504 let mut cfg = self.attn_cfg(position);
2505 cfg.q_norm = q_norm.as_deref();
2506 cfg.k_norm = k_norm.as_deref();
2507 cfg.output_gate = *output_gate;
2508 cfg.softplus_gate = softplus_gate.as_ref().map(|(g, p)| (g, *p));
2509 cfg.bias = bias.as_ref().map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2510 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2511 }
2512 _ => return,
2513 };
2514 let _ = attn;
2515 }
2516
2517 #[cfg(feature = "gpu")]
2524 #[allow(clippy::too_many_arguments)]
2525 fn graph_spec_step(
2526 &mut self,
2527 m: &mut MtpModule,
2528 hidden: &[f32],
2529 t_next: u32,
2530 next_pos: usize,
2531 drafted: &mut usize,
2532 accepted: &mut usize,
2533 ) -> Option<(Vec<u32>, usize, Vec<f32>)> {
2534 let k_spec: usize = std::env::var("CMF_GRAPH_SPEC_K")
2540 .ok()
2541 .and_then(|v| v.parse().ok())
2542 .filter(|&v| (1..=8).contains(&v))
2543 .unwrap_or(3);
2544 if next_pos == 0 {
2545 return None;
2546 }
2547 let t_round = std::time::Instant::now();
2548 let subs = || crate::gpu_wgpu::SUBMITS.load(std::sync::atomic::Ordering::Relaxed);
2564 let sub0 = subs();
2565 let mut drafts = Vec::with_capacity(k_spec);
2570 let (d1, mut hx) = self.mtp_step_h(m, hidden, t_next, next_pos - 1);
2571 drafts.push(d1);
2572 for j in 1..k_spec {
2573 let (dj, hj) = self.mtp_step_h(m, &hx, drafts[j - 1], next_pos - 1 + j);
2574 drafts.push(dj);
2575 hx = hj;
2576 }
2577 *drafted += k_spec;
2578 let t_draft = t_round.elapsed();
2579 let sub_draft = subs();
2580 let b = k_spec + 1;
2583 let mut hiddens = vec![0.0f32; b * self.hidden_size];
2584 for (i, &t) in std::iter::once(&t_next).chain(drafts.iter()).enumerate() {
2585 let e = self.embed_single(t);
2586 hiddens[i * self.hidden_size..(i + 1) * self.hidden_size].copy_from_slice(&e);
2587 }
2588 let positions: Vec<usize> = (next_pos..next_pos + b).collect();
2589 let (lm_gw, lm_rows) = {
2590 let (_, i, kind, rs) = self.weights.lm_head.graph_weight()?;
2591 (
2592 crate::gpu::GraphW { idx: i, kind, row_scale: rs, data: &[] },
2593 self.weights.lm_head.rows(),
2594 )
2595 };
2596 let mut logits = Vec::new();
2597 let final_norm = self.weights.final_norm.clone();
2598 let ok = self.try_batch_graph_wgpu(
2599 &mut hiddens,
2600 &positions,
2601 b,
2602 Some(crate::gpu::SpecTail {
2603 lm: lm_gw,
2604 lm_rows,
2605 final_norm: &final_norm,
2606 logits_out: &mut logits,
2607 }),
2608 );
2609 if !ok {
2610 m.kv.truncate_last(k_spec);
2613 return None;
2614 }
2615 let t_verify = t_round.elapsed();
2616 let sub_verify = subs();
2617 let ids: Vec<u32> = (0..b)
2619 .map(|i| sampler::argmax(&logits[i * lm_rows..(i + 1) * lm_rows]))
2620 .collect();
2621 let mut a = 0usize;
2622 while a < k_spec && ids[a] == drafts[a] {
2623 a += 1;
2624 }
2625 if a + 1 < b {
2627 crate::gpu::gdn_spec_restore(self.graph_kv_id, a);
2628 }
2629 *accepted += a;
2630 m.kv.truncate_last(k_spec.saturating_sub(1));
2641 let warm_off = std::env::var("CMF_SPEC_WARM").is_ok_and(|v| v == "0");
2642 if !warm_off {
2643 for j in 0..a {
2644 let row = &hiddens[j * self.hidden_size..(j + 1) * self.hidden_size];
2645 let row = row.to_vec();
2646 self.mtp_warm(m, &row, ids[j], next_pos + j);
2647 }
2648 }
2649 let mut row = logits[a * lm_rows..(a + 1) * lm_rows].to_vec();
2651 row.resize(self.vocab_size, 0.0);
2652 if let Some(c) = self.final_softcap {
2653 for l in row.iter_mut() {
2654 *l = c * (*l / c).tanh();
2655 }
2656 }
2657 self.graph_logits = Some(row);
2658 let new_hidden = hiddens[a * self.hidden_size..(a + 1) * self.hidden_size].to_vec();
2659 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
2665 let end = subs();
2666 eprintln!(
2667 "spec-round: draft {:.1} ms/{} sub | verify {:.1} ms/{} sub | \
2668 commit {:.1} ms/{} sub (accepted {a} of {k_spec})",
2669 t_draft.as_secs_f64() * 1e3,
2670 sub_draft - sub0,
2671 (t_verify - t_draft).as_secs_f64() * 1e3,
2672 sub_verify - sub_draft,
2673 (t_round.elapsed() - t_verify).as_secs_f64() * 1e3,
2674 end - sub_verify,
2675 );
2676 }
2677 Some((drafts[..a].to_vec(), next_pos + a + 1, new_hidden))
2678 }
2679
2680 pub fn measure_pair_fusion(&mut self, iters: usize) -> (f64, f64) {
2689 if !self.pair_supported() {
2690 return (0.0, 0.0);
2691 }
2692 let emb1 = self.embed_single(1);
2693 let emb2 = self.embed_single(2);
2694 let pos = self.kv_cache.seq_len();
2695
2696 let t0 = std::time::Instant::now();
2697 for _ in 0..iters {
2698 let _ = self.forward_layers(&emb1, pos, None);
2699 let _ = self.forward_layers(&emb2, pos + 1, None);
2700 for l in &mut self.kv_cache.layers {
2701 l.truncate_last(2);
2702 }
2703 }
2704 let singles_ms = t0.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2705
2706 let t1 = std::time::Instant::now();
2707 for _ in 0..iters {
2708 let _ = self.forward_pair(&emb1, &emb2, pos);
2709 for l in &mut self.kv_cache.layers {
2710 l.truncate_last(2);
2711 }
2712 }
2713 let pair_ms = t1.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2714 (singles_ms, pair_ms)
2715 }
2716
2717 fn pair_supported(&self) -> bool {
2725 !self.weights.layers.is_empty()
2732 && self.g3n.is_none()
2733 && !self
2734 .weights
2735 .layers
2736 .iter()
2737 .any(|lw| matches!(&lw.attn, AttnKind::Mla(_) | AttnKind::Kda(_)))
2738 }
2739
2740 fn forward_pair(
2741 &mut self,
2742 emb1: &[f32],
2743 emb2: &[f32],
2744 position: usize,
2745 ) -> (Vec<f32>, Vec<f32>) {
2746 let mut h1 = emb1.to_vec();
2747 let mut h2 = emb2.to_vec();
2748 let (_nkv, _hd, hs, _rd, eps) = (
2749 self.num_kv_heads,
2750 self.head_dim,
2751 self.hidden_size,
2752 self.rotary_dim,
2753 self.rms_eps,
2754 );
2755 let pool = self.pool.clone();
2756
2757 for li in 0..self.num_layers {
2758 let lw = &self.weights.layers[self.phys_layer(li)];
2759 inference::rms_norm_into(
2762 &h1,
2763 &lw.input_norm,
2764 self.rms_eps,
2765 self.norm_style,
2766 &mut self.ws.n1,
2767 );
2768 inference::rms_norm_into(
2769 &h2,
2770 &lw.input_norm,
2771 self.rms_eps,
2772 self.norm_style,
2773 &mut self.ws.n2,
2774 );
2775
2776 let (a1, a2) = match &lw.attn {
2777 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2778 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2779 AttnKind::Linear(w) => {
2780 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
2781 let layer = &mut self.kv_cache.layers[li];
2782 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2783 vmf_phase_pair(
2784 &self.ws.n1,
2785 &self.ws.n2,
2786 w,
2787 &cfg,
2788 state,
2789 scratch,
2790 self.pool.as_deref(),
2791 )
2792 }
2793 AttnKind::LinearGdn(w) => {
2794 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
2795 let layer = &mut self.kv_cache.layers[li];
2796 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2797 gdn_pair(
2798 &self.ws.n1,
2799 &self.ws.n2,
2800 w,
2801 &cfg,
2802 state,
2803 scratch,
2804 self.pool.as_deref(),
2805 )
2806 }
2807 AttnKind::ShortConv(w) => {
2808 let cfg = self
2809 .short_conv_cfg
2810 .expect("short-conv layer without short_conv_cfg");
2811 let layer = &mut self.kv_cache.layers[li];
2812 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2813 short_conv_pair(
2814 &self.ws.n1,
2815 &self.ws.n2,
2816 w,
2817 &cfg,
2818 state,
2819 scratch,
2820 self.pool.as_deref(),
2821 )
2822 }
2823 AttnKind::Full {
2824 wq,
2825 wk,
2826 wv,
2827 wo,
2828 q_norm,
2829 k_norm,
2830 output_gate,
2831 softplus_gate,
2832 bias,
2833 } => {
2834 let inv_freq_l = self.layer_inv_freq(li);
2835 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
2836 let cfg = QwenAttnCfg {
2837 num_heads: self.layer_num_heads(li),
2838 num_kv_heads: nkv_l,
2839 head_dim: hd_l,
2840 hidden_size: hs,
2841 position,
2842 inv_freq: &inv_freq_l,
2843 rotary_dim: rd_l,
2844 scale: self.attn_scale,
2845 softcap: self.attn_softcap,
2846 window: self.layer_window(li),
2847 v_norm: self.attn_v_norm,
2848 q_norm: q_norm.as_deref(),
2849 k_norm: k_norm.as_deref(),
2850 output_gate: *output_gate,
2851 softplus_gate: softplus_gate
2852 .as_ref()
2853 .map(|(gate, per_head)| (gate, *per_head)),
2854 rope_scale: self.layer_rope_scale(li),
2855 bias: bias
2856 .as_ref()
2857 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
2858 rms_eps: eps,
2859 norm_style: self.norm_style,
2860 pool: pool.as_deref(),
2861 };
2862 attention::qwen_attention_pair(
2863 &self.ws.n1,
2864 &self.ws.n2,
2865 wq,
2866 wk,
2867 wv,
2868 wo,
2869 &mut self.kv_cache.layers[li],
2870 &cfg,
2871 )
2872 }
2873 };
2874 let (a1, a2) = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
2875 Some(w) => (
2876 inference::rms_norm(&a1, w, self.rms_eps, self.norm_style),
2877 inference::rms_norm(&a2, w, self.rms_eps, self.norm_style),
2878 ),
2879 None => (a1, a2),
2880 };
2881 for i in 0..self.hidden_size {
2882 h1[i] += a1[i];
2883 h2[i] += a2[i];
2884 }
2885 let (mut a1, mut a2) = (a1, a2);
2886 attention::recycle_buf(&mut a1);
2887 attention::recycle_buf(&mut a2);
2888
2889 let lw = &self.weights.layers[self.phys_layer(li)];
2890 inference::rms_norm_into(
2891 &h1,
2892 &lw.post_norm,
2893 self.rms_eps,
2894 self.norm_style,
2895 &mut self.ws.p1,
2896 );
2897 inference::rms_norm_into(
2898 &h2,
2899 &lw.post_norm,
2900 self.rms_eps,
2901 self.norm_style,
2902 &mut self.ws.p2,
2903 );
2904 let (f1, f2) = match &lw.ffn {
2905 FfnKind::DenseMoe(dm) => (
2908 dense_moe_ffn(
2909 dm,
2910 &self.ws.p1,
2911 &h1,
2912 self.rms_eps,
2913 self.norm_style,
2914 self.pool.as_deref(),
2915 ),
2916 dense_moe_ffn(
2917 dm,
2918 &self.ws.p2,
2919 &h2,
2920 self.rms_eps,
2921 self.norm_style,
2922 self.pool.as_deref(),
2923 ),
2924 ),
2925 _ => ffn_forward_pair(
2926 &lw.ffn,
2927 &self.ws.p1,
2928 &self.ws.p2,
2929 self.pool.as_deref(),
2930 None,
2931 ),
2932 };
2933 let (f1, f2) = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
2934 Some(w) => (
2935 inference::rms_norm(&f1, w, self.rms_eps, self.norm_style),
2936 inference::rms_norm(&f2, w, self.rms_eps, self.norm_style),
2937 ),
2938 None => (f1, f2),
2939 };
2940 for i in 0..self.hidden_size {
2941 h1[i] += f1[i];
2942 h2[i] += f2[i];
2943 }
2944 let (mut f1, mut f2) = (f1, f2);
2945 attention::recycle_buf(&mut f1);
2946 attention::recycle_buf(&mut f2);
2947 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
2948 for i in 0..self.hidden_size {
2949 h1[i] *= sc;
2950 h2[i] *= sc;
2951 }
2952 }
2953 if self.is_loop_end(li) && li + 1 < self.num_layers {
2955 h1 = inference::rms_norm(
2956 &h1,
2957 &self.weights.final_norm,
2958 self.rms_eps,
2959 self.norm_style,
2960 );
2961 h2 = inference::rms_norm(
2962 &h2,
2963 &self.weights.final_norm,
2964 self.rms_eps,
2965 self.norm_style,
2966 );
2967 }
2968 }
2969 (h1, h2)
2970 }
2971
2972 fn commit_linear_scratch(&mut self) {
2974 for layer in &mut self.kv_cache.layers {
2975 if !layer.linear_scratch.is_empty() {
2976 std::mem::swap(&mut layer.linear_state, &mut layer.linear_scratch);
2977 layer.linear_scratch.clear();
2978 }
2979 }
2980 }
2981
2982 pub fn forward_ids(
2985 &mut self,
2986 ids: &[u32],
2987 task_mask: Option<&TaskMask>,
2988 ) -> Result<Vec<f32>, String> {
2989 if ids.is_empty() {
2990 return Err("empty id sequence".to_string());
2991 }
2992 self.kv_cache.clear();
2993 self.kv_history.clear();
2994 self.o1_begin();
2995 let mut hidden = vec![0.0f32; self.hidden_size];
2996 let mut pos = 0usize;
2997 if self.can_prefill_batched() && !self.graph_prefill_preferred() && ids.len() > 2 {
3005 let chunk = prefill_chunk();
3009 let hs = self.hidden_size;
3010 while pos < ids.len() {
3011 let end = (pos + chunk).min(ids.len());
3012 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3013 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
3014 pos = end;
3015 }
3016 }
3017 if task_mask.is_none()
3026 && !self.graph_prefill_preferred()
3027 && !std::env::var("CMF_PAIR").is_ok_and(|v| v == "0")
3028 && self.pair_supported()
3029 {
3030 while pos + 1 < ids.len() {
3031 let e1 = self.embed_single(ids[pos]);
3032 let e2 = self.embed_single(ids[pos + 1]);
3033 let (_, h2) = self.forward_pair(&e1, &e2, pos);
3034 self.commit_linear_scratch();
3035 hidden = h2;
3036 pos += 2;
3037 }
3038 }
3039 while pos < ids.len() {
3040 hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3041 pos += 1;
3042 }
3043 self.o1_seal();
3047 let normed = inference::rms_norm(
3048 &hidden,
3049 &self.weights.final_norm,
3050 self.rms_eps,
3051 self.norm_style,
3052 );
3053 Ok(self.lm_head_forward(&normed))
3054 }
3055
3056 pub fn ppl_ids(&mut self, ids: &[u32]) -> f64 {
3063 let (nll, cnt) = self.nll_ids_from(ids, 0);
3064 (nll / cnt.max(1) as f64).exp()
3065 }
3066
3067 pub fn probe_ffn_mass(&mut self, ids: &[u32]) -> Vec<Vec<f64>> {
3072 self.kv_cache.clear();
3073 self.kv_history.clear();
3074 FFN_PROBE.with(|p| {
3075 *p.borrow_mut() = Some(vec![vec![0f64; self.intermediate_size]; self.num_layers]);
3076 });
3077 crate::gpu::cpu_scope(|| {
3078 for (pos, &id) in ids.iter().enumerate() {
3079 let emb = self.embed_single(id);
3080 let _ = self.forward_layers(&emb, pos, None);
3081 }
3082 });
3083 self.kv_cache.clear();
3084 self.kv_history.clear();
3085 FFN_PROBE
3086 .with(|p| p.borrow_mut().take())
3087 .unwrap_or_default()
3088 }
3089
3090 pub fn ppl_ids_masked(&mut self, ids: &[u32], mask: &TaskMask) -> f64 {
3094 self.kv_cache.clear();
3095 self.kv_history.clear();
3096 let mut nll = 0f64;
3097 let mut cnt = 0usize;
3098 let mut hidden = vec![0f32; self.hidden_size];
3099 for (pos, &id) in ids.iter().enumerate() {
3100 if pos > 0 {
3101 inference::rms_norm_into(
3102 &hidden,
3103 &self.weights.final_norm,
3104 self.rms_eps,
3105 self.norm_style,
3106 &mut self.ws.n1,
3107 );
3108 let mut logits = self.lm_head_forward(&self.ws.n1);
3109 let max = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
3110 let sum: f64 = logits.iter().map(|&v| ((v - max) as f64).exp()).sum();
3111 let p = ((logits[id as usize] - max) as f64).exp() / sum.max(1e-300);
3112 nll -= p.max(1e-300).ln();
3113 cnt += 1;
3114 attention::recycle_buf(&mut logits);
3115 }
3116 let emb = self.embed_single(id);
3117 hidden = self.forward_layers(&emb, pos, Some(mask));
3118 }
3119 self.kv_cache.clear();
3120 self.kv_history.clear();
3121 (nll / cnt.max(1) as f64).exp()
3122 }
3123
3124 pub fn nll_ids_masked(
3143 &mut self,
3144 ids: &[u32],
3145 start: usize,
3146 task_mask: Option<&TaskMask>,
3147 ) -> (f64, usize) {
3148 self.nll_ids_inner(ids, start, task_mask)
3149 }
3150
3151 pub fn nll_ids_from(&mut self, ids: &[u32], start: usize) -> (f64, usize) {
3152 self.nll_ids_inner(ids, start, None)
3153 }
3154
3155 fn nll_ids_inner(
3156 &mut self,
3157 ids: &[u32],
3158 start: usize,
3159 task_mask: Option<&TaskMask>,
3160 ) -> (f64, usize) {
3161 self.kv_cache.clear();
3162 self.kv_history.clear();
3163 let mut nll = 0f64;
3164 let mut cnt = 0usize;
3165 if self.can_prefill_batched() {
3166 const CHUNK: usize = 128;
3172 const LM_SUB: usize = 32;
3173 let n = ids.len().saturating_sub(1);
3174 let hs = self.hidden_size;
3175 let rows = self.weights.lm_head.rows();
3176 let mut pos = 0usize;
3177 while pos < n {
3178 let end = (pos + CHUNK).min(n);
3179 let bsz = end - pos;
3180 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3181 let mut k0 = 0usize;
3182 while k0 < bsz {
3183 let k1 = (k0 + LM_SUB).min(bsz);
3184 let sb = k1 - k0;
3185 if pos + k1 <= start {
3188 k0 = k1;
3189 continue;
3190 }
3191 let mut normed = vec![0.0f32; sb * hs];
3192 for k in 0..sb {
3193 let r = inference::rms_norm(
3194 &hb[(k0 + k) * hs..(k0 + k + 1) * hs],
3195 &self.weights.final_norm,
3196 self.rms_eps,
3197 self.norm_style,
3198 );
3199 normed[k * hs..(k + 1) * hs].copy_from_slice(&r);
3200 }
3201 let mut logits = vec![0.0f32; sb * rows];
3202 self.weights
3203 .lm_head
3204 .matmat(&normed, sb, &mut logits, self.pool.as_deref());
3205 for k in 0..sb {
3206 if pos + k0 + k < start {
3207 continue;
3208 }
3209 let lg = &mut logits[k * rows..k * rows + self.vocab_size.min(rows)];
3210 if let Some(mu) = self.logit_multiplier {
3211 for v in lg.iter_mut() {
3212 *v *= mu;
3213 }
3214 }
3215 if let Some(c) = self.final_softcap {
3219 for v in lg.iter_mut() {
3220 *v = c * (*v / c).tanh();
3221 }
3222 }
3223 let lg = &logits[k * rows..k * rows + self.vocab_size.min(rows)];
3224 let target = ids[pos + k0 + k + 1] as usize;
3225 let max = lg.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3226 let lse: f64 = lg
3227 .iter()
3228 .map(|&v| ((v - max) as f64).exp())
3229 .sum::<f64>()
3230 .ln()
3231 + max as f64;
3232 nll += lse - lg[target] as f64;
3233 cnt += 1;
3234 if std::env::var("CMF_PPL_TRACE").is_ok() {
3235 let top = lg
3236 .iter()
3237 .enumerate()
3238 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3239 .map(|(i, _)| i)
3240 .unwrap_or(0);
3241 eprintln!(
3242 "BTRACE pos {} target {} nll {:.4} top {} lg_t {:.3} lg_top {:.3}",
3243 pos + k0 + k,
3244 target,
3245 lse - lg[target] as f64,
3246 top,
3247 lg[target],
3248 lg[top]
3249 );
3250 }
3251 }
3252 k0 = k1;
3253 }
3254 pos = end;
3255 }
3256 self.kv_cache.clear();
3257 self.kv_history.clear();
3258 return (nll, cnt);
3259 }
3260 for pos in 0..ids.len().saturating_sub(1) {
3261 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3262 let out_of_band = self.graph_logits.take();
3270 if pos < start {
3271 continue;
3272 }
3273 let logits = match out_of_band {
3274 Some(lg) => lg,
3275 None => {
3276 let normed = inference::rms_norm(
3277 &hidden,
3278 &self.weights.final_norm,
3279 self.rms_eps,
3280 self.norm_style,
3281 );
3282 self.lm_head_forward(&normed)
3286 }
3287 };
3288 let target = ids[pos + 1] as usize;
3289 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3290 let lse: f64 = logits
3291 .iter()
3292 .map(|&v| ((v - max) as f64).exp())
3293 .sum::<f64>()
3294 .ln()
3295 + max as f64;
3296 let tok_nll = lse - logits[target] as f64;
3297 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3298 let top = logits
3299 .iter()
3300 .enumerate()
3301 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3302 .map(|(i, _)| i)
3303 .unwrap_or(0);
3304 eprintln!(
3305 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3306 logits[target], logits[top]
3307 );
3308 }
3309 nll += tok_nll;
3310 cnt += 1;
3311 }
3312 self.kv_cache.clear();
3313 self.kv_history.clear();
3314 (nll, cnt)
3315 }
3316
3317 pub fn nll_ids_o1(&mut self, ids: &[u32], prefill: usize) -> (f64, usize) {
3333 self.kv_cache.clear();
3334 self.kv_history.clear();
3335 self.o1_begin();
3336 let n = ids.len().saturating_sub(1);
3337 let p = prefill.min(n);
3338 let mut pos = 0usize;
3340 if self.can_prefill_batched() {
3341 const CHUNK: usize = 128;
3342 while pos < p {
3343 let end = (pos + CHUNK).min(p);
3344 let _ = self.prefill_batch(&ids[pos..end], pos);
3345 pos = end;
3346 }
3347 } else {
3348 while pos < p {
3349 let _ = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3350 pos += 1;
3351 }
3352 }
3353 self.o1_seal();
3354
3355 let mut nll = 0f64;
3356 let mut cnt = 0usize;
3357 for pos in p..n {
3358 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3359 let normed = inference::rms_norm(
3360 &hidden,
3361 &self.weights.final_norm,
3362 self.rms_eps,
3363 self.norm_style,
3364 );
3365 let logits = self.lm_head_forward(&normed);
3369 let target = ids[pos + 1] as usize;
3370 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3371 let lse: f64 = logits
3372 .iter()
3373 .map(|&v| ((v - max) as f64).exp())
3374 .sum::<f64>()
3375 .ln()
3376 + max as f64;
3377 let tok_nll = lse - logits[target] as f64;
3378 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3379 let top = logits
3380 .iter()
3381 .enumerate()
3382 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3383 .map(|(i, _)| i)
3384 .unwrap_or(0);
3385 eprintln!(
3386 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3387 logits[target], logits[top]
3388 );
3389 }
3390 nll += tok_nll;
3391 cnt += 1;
3392 }
3393 self.kv_cache.clear();
3394 self.kv_history.clear();
3395 (nll, cnt)
3396 }
3397
3398 pub fn calib_ids(&mut self, ids: &[u32], temps: &[f32]) -> (Vec<bool>, Vec<Vec<f32>>) {
3406 self.kv_cache.clear();
3407 self.kv_history.clear();
3408 let n = ids.len().saturating_sub(1);
3409 let mut correct = Vec::with_capacity(n);
3410 let mut pmax = Vec::with_capacity(n);
3411 for pos in 0..n {
3412 let emb = self.embed_single(ids[pos]);
3413 let hidden = self.forward_layers(&emb, pos, None);
3414 let normed = inference::rms_norm(
3415 &hidden,
3416 &self.weights.final_norm,
3417 self.rms_eps,
3418 self.norm_style,
3419 );
3420 let logits = self.lm_head_forward(&normed);
3424 let target = ids[pos + 1] as usize;
3425 let (mut amax, mut mval) = (0usize, f32::NEG_INFINITY);
3426 for (i, &v) in logits.iter().enumerate() {
3427 if v > mval {
3428 mval = v;
3429 amax = i;
3430 }
3431 }
3432 correct.push(amax == target);
3433 let row: Vec<f32> = temps
3434 .iter()
3435 .map(|&t| {
3436 let tt = t.max(1e-3);
3437 let s: f32 = logits.iter().map(|&v| ((v - mval) / tt).exp()).sum();
3438 1.0 / s.max(1e-12) })
3440 .collect();
3441 pmax.push(row);
3442 }
3443 self.kv_cache.clear();
3444 self.kv_history.clear();
3445 (correct, pmax)
3446 }
3447
3448 pub fn ppl_ids_dynamic(&mut self, ids: &[u32]) -> (f64, usize) {
3455 let mut router = match self.dyn_router.take() {
3456 Some(r) => r,
3457 None => return (self.ppl_ids(ids), 0),
3458 };
3459 router.reset();
3460 self.dyn_phi_seen = 0;
3461 let _ = self.set_active_skill(None);
3462
3463 self.kv_cache.clear();
3464
3465 self.kv_history.clear();
3466 let mut nll = 0f64;
3467 let mut cnt = 0usize;
3468 for pos in 0..ids.len().saturating_sub(1) {
3469 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3470 let normed = inference::rms_norm(
3471 &hidden,
3472 &self.weights.final_norm,
3473 self.rms_eps,
3474 self.norm_style,
3475 );
3476 let logits = self.lm_head_forward(&normed);
3480 let target = ids[pos + 1] as usize;
3481 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3482 let lse: f64 = logits
3483 .iter()
3484 .map(|&v| ((v - max) as f64).exp())
3485 .sum::<f64>()
3486 .ln()
3487 + max as f64;
3488 let tok_nll = lse - logits[target] as f64;
3489 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3490 let top = logits
3491 .iter()
3492 .enumerate()
3493 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3494 .map(|(i, _)| i)
3495 .unwrap_or(0);
3496 eprintln!(
3497 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3498 logits[target], logits[top]
3499 );
3500 }
3501 nll += tok_nll;
3502 cnt += 1;
3503 let phi = self.dyn_phi_ema.clone();
3505 if let Some(new_active) = router.step(&phi, pos) {
3506 let _ = self.set_active_skill(new_active);
3507 }
3508 }
3509 let switches = router.switches.len();
3510 let _ = self.set_active_skill(None);
3511 self.dyn_router = Some(router);
3512 self.kv_cache.clear();
3513 self.kv_history.clear();
3514 ((nll / cnt.max(1) as f64).exp(), switches)
3515 }
3516
3517 pub fn probe_phi(&mut self, ids: &[u32], layer: usize) -> Vec<f32> {
3519 self.kv_cache.clear();
3520 self.kv_history.clear();
3521 let mut acc = vec![0f32; self.hidden_size];
3522 for (pos, &id) in ids.iter().enumerate() {
3523 let h = self.forward_layers_upto(&self.embed_single(id), pos, None, Some(layer));
3524 for (a, v) in acc.iter_mut().zip(&h) {
3525 *a += v;
3526 }
3527 }
3528 let n = ids.len().max(1) as f32;
3529 for a in acc.iter_mut() {
3530 *a /= n;
3531 }
3532 self.kv_cache.clear();
3533 self.kv_history.clear();
3534 acc
3535 }
3536
3537 fn prefill_batch(&mut self, ids: &[u32], start_pos: usize) -> Vec<f32> {
3543 self.prefill_batch_masked(ids, start_pos, None)
3544 }
3545
3546 fn prefill_batch_masked(
3552 &mut self,
3553 ids: &[u32],
3554 start_pos: usize,
3555 task_mask: Option<&TaskMask>,
3556 ) -> Vec<f32> {
3557 self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, usize::MAX)
3558 }
3559
3560 fn prefill_batch_span(
3566 &mut self,
3567 input: PrefillIn<'_>,
3568 start_pos: usize,
3569 task_mask: Option<&TaskMask>,
3570 from: usize,
3571 upto_excl: usize,
3572 ) -> Vec<f32> {
3573 let hs = self.hidden_size;
3574 let b = match input {
3575 PrefillIn::Ids(ids) => ids.len(),
3576 PrefillIn::Hidden(hb) => hb.len() / hs,
3577 };
3578 let upto_excl = upto_excl.min(self.num_layers);
3579 let mut h: Vec<f32>;
3583 let mut h_ready;
3584 match input {
3585 PrefillIn::Ids(_) => {
3586 h = vec![0.0; b * hs];
3587 h_ready = false;
3588 }
3589 PrefillIn::Hidden(hb) => {
3590 h = hb.to_vec();
3591 h_ready = true;
3592 }
3593 }
3594 let fill_h = |h: &mut Vec<f32>, me: &Self| {
3595 if let PrefillIn::Ids(ids) = input {
3596 for (bi, &id) in ids.iter().enumerate() {
3597 let e = me.embed_single(id);
3598 h[bi * hs..(bi + 1) * hs].copy_from_slice(&e);
3599 }
3600 }
3601 };
3602 let (_nkv, _hd, _rd, eps) = (
3603 self.num_kv_heads,
3604 self.head_dim,
3605 self.rotary_dim,
3606 self.rms_eps,
3607 );
3608 let pool = self.pool.clone();
3609 let norm_style = self.norm_style;
3610
3611 #[cfg(target_os = "macos")]
3612 let mut chunk_skip_until = 0usize;
3613 for li in from..upto_excl {
3614 crate::gpu::set_layer(li as i64); #[cfg(target_os = "macos")]
3621 if task_mask.is_none() {
3622 if li < chunk_skip_until {
3623 continue;
3624 }
3625 if !h_ready && li == 0 && self.weights.embed_tokens.q8_row_parts().is_none() {
3631 fill_h(&mut h, self);
3632 h_ready = true;
3633 }
3634 let ids_for_embed = match input {
3635 PrefillIn::Ids(ids) => (!h_ready && li == 0).then_some(ids),
3636 PrefillIn::Hidden(_) => None,
3637 };
3638 let end = self.chunk_run_gpu(li, &mut h, b, start_pos, ids_for_embed, upto_excl);
3639 if end > li {
3640 h_ready = true;
3641 chunk_skip_until = end;
3642 if self.is_loop_end(end - 1) && end < self.num_layers {
3645 for bi in 0..b {
3646 let normed = inference::rms_norm(
3647 &h[bi * hs..(bi + 1) * hs],
3648 &self.weights.final_norm,
3649 eps,
3650 norm_style,
3651 );
3652 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
3653 }
3654 }
3655 continue;
3656 }
3657 }
3658 if !h_ready {
3659 fill_h(&mut h, self);
3660 h_ready = true;
3661 }
3662 let lw = &self.weights.layers[self.phys_layer(li)];
3663 match &lw.attn {
3665 AttnKind::Kda(w) => {
3666 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
3668 let mut normed = vec![0.0f32; b * hs];
3669 for bi in 0..b {
3670 inference::rms_norm_into(
3671 &h[bi * hs..(bi + 1) * hs],
3672 &lw.input_norm,
3673 eps,
3674 norm_style,
3675 &mut normed[bi * hs..(bi + 1) * hs],
3676 );
3677 }
3678 let attn = crate::linear_core::kda_forward_batch(
3679 &normed,
3680 b,
3681 w,
3682 &cfg,
3683 &mut self.kv_cache.layers[li].linear_state,
3684 pool.as_deref(),
3685 );
3686 for (dst, &a) in h.iter_mut().zip(&attn) {
3687 *dst += a;
3688 }
3689 }
3690 AttnKind::LinearGdn(w) => {
3691 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
3693 let mut normed = vec![0.0f32; b * hs];
3694 for bi in 0..b {
3695 let r = inference::rms_norm(
3696 &h[bi * hs..(bi + 1) * hs],
3697 &lw.input_norm,
3698 eps,
3699 norm_style,
3700 );
3701 normed[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3702 }
3703 let attn = crate::linear_core::gdn_forward_batch(
3704 &normed,
3705 b,
3706 w,
3707 &cfg,
3708 &mut self.kv_cache.layers[li].linear_state,
3709 pool.as_deref(),
3710 );
3711 for (dst, &a) in h.iter_mut().zip(&attn) {
3712 *dst += a;
3713 }
3714 }
3715 AttnKind::ShortConv(w) => {
3716 let cfg = self
3719 .short_conv_cfg
3720 .expect("short-conv layer without short_conv_cfg");
3721 let mut normed = vec![0.0f32; b * hs];
3722 for bi in 0..b {
3723 inference::rms_norm_into(
3724 &h[bi * hs..(bi + 1) * hs],
3725 &lw.input_norm,
3726 eps,
3727 norm_style,
3728 &mut normed[bi * hs..(bi + 1) * hs],
3729 );
3730 }
3731 let attn = short_conv_forward_batch(
3732 &normed,
3733 b,
3734 w,
3735 &cfg,
3736 &mut self.kv_cache.layers[li].linear_state,
3737 pool.as_deref(),
3738 );
3739 for (dst, &a) in h.iter_mut().zip(&attn) {
3740 *dst += a;
3741 }
3742 }
3743 AttnKind::Mla(w) => {
3744 let inv_freq_l = self.layer_inv_freq(li);
3747 let rs = self.layer_rope_scale(li);
3748 let mut normed = vec![0.0f32; hs];
3749 for bi in 0..b {
3750 inference::rms_norm_into(
3751 &h[bi * hs..(bi + 1) * hs],
3752 &lw.input_norm,
3753 eps,
3754 norm_style,
3755 &mut normed,
3756 );
3757 let ao = mla_attention(
3758 w,
3759 &normed,
3760 &mut self.kv_cache.layers[li],
3761 start_pos + bi,
3762 &inv_freq_l,
3763 rs,
3764 eps,
3765 pool.as_deref(),
3766 );
3767 for (dst, &a) in h[bi * hs..(bi + 1) * hs].iter_mut().zip(&ao) {
3768 *dst += a;
3769 }
3770 }
3771 }
3772 AttnKind::Full {
3773 wq,
3774 wk,
3775 wv,
3776 wo,
3777 q_norm,
3778 k_norm,
3779 output_gate,
3780 softplus_gate,
3781 bias,
3782 } => {
3783 let mut normed = vec![0.0f32; b * hs];
3787 for bi in 0..b {
3788 inference::rms_norm_into(
3789 &h[bi * hs..(bi + 1) * hs],
3790 &lw.input_norm,
3791 eps,
3792 norm_style,
3793 &mut normed[bi * hs..(bi + 1) * hs],
3794 );
3795 }
3796 let inv_freq_l = self.layer_inv_freq(li);
3797 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
3798 let cfg = QwenAttnCfg {
3799 num_heads: self.layer_num_heads(li),
3800 num_kv_heads: nkv_l,
3801 head_dim: hd_l,
3802 hidden_size: hs,
3803 position: start_pos,
3804 inv_freq: &inv_freq_l,
3805 rotary_dim: rd_l,
3806 scale: self.attn_scale,
3807 softcap: self.attn_softcap,
3808 window: self.layer_window(li),
3809 v_norm: self.attn_v_norm,
3810 q_norm: q_norm.as_deref(),
3811 k_norm: k_norm.as_deref(),
3812 output_gate: *output_gate,
3813 softplus_gate: softplus_gate
3814 .as_ref()
3815 .map(|(gate, per_head)| (gate, *per_head)),
3816 rope_scale: self.layer_rope_scale(li),
3817 bias: bias
3818 .as_ref()
3819 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
3820 rms_eps: eps,
3821 norm_style,
3822 pool: pool.as_deref(),
3823 };
3824 let mut attn = attention::qwen_attention_batch(
3825 &normed,
3826 b,
3827 wq,
3828 wk,
3829 wv,
3830 wo,
3831 &mut self.kv_cache.layers[li],
3832 &cfg,
3833 );
3834 if let Some(w) = &lw.attn_out_norm {
3835 for bi in 0..b {
3836 inference::rms_norm_into(
3837 &attn[bi * hs..(bi + 1) * hs],
3838 w,
3839 eps,
3840 norm_style,
3841 &mut normed[bi * hs..(bi + 1) * hs],
3842 );
3843 }
3844 attn.copy_from_slice(&normed);
3845 }
3846 for (dst, &a) in h.iter_mut().zip(&attn) {
3847 *dst += a;
3848 }
3849 }
3850 AttnKind::Linear(w) => {
3851 for bi in 0..b {
3852 let normed = inference::rms_norm(
3853 &h[bi * hs..(bi + 1) * hs],
3854 &lw.input_norm,
3855 eps,
3856 norm_style,
3857 );
3858 vmf_phase_forward(
3859 &normed,
3860 w,
3861 &self.vmf_cfg.expect("linear layer without vmf_cfg"),
3862 &mut self.kv_cache.layers[li].linear_state,
3863 pool.as_deref(),
3864 )
3865 .iter()
3866 .enumerate()
3867 .for_each(|(i, &a)| h[bi * hs + i] += a);
3868 }
3869 }
3870 }
3871
3872 let lw = &self.weights.layers[self.phys_layer(li)];
3874 let mut post = vec![0.0f32; b * hs];
3875 for bi in 0..b {
3876 let r =
3877 inference::rms_norm(&h[bi * hs..(bi + 1) * hs], &lw.post_norm, eps, norm_style);
3878 post[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3879 }
3880 let mask_row = task_mask
3883 .filter(|m| m.ffn_active_count(li) < self.intermediate_size)
3884 .and_then(|m| m.ffn_masks.get(li))
3885 .map(|v| v.as_slice());
3886 let mut ffn = match &lw.ffn {
3887 FfnKind::Dense(d) => dense_ffn_batch(d, &post, b, pool.as_deref(), mask_row),
3888 FfnKind::Moe(m) => moe_ffn_batch(m, &post, b, hs, pool.as_deref(), None),
3889 FfnKind::DenseMoe(dm) => {
3892 let mut out = vec![0.0f32; b * hs];
3893 for bi in 0..b {
3894 let r = dense_moe_ffn(
3895 dm,
3896 &post[bi * hs..(bi + 1) * hs],
3897 &h[bi * hs..(bi + 1) * hs],
3898 eps,
3899 norm_style,
3900 pool.as_deref(),
3901 );
3902 out[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3903 }
3904 out
3905 }
3906 };
3907 if let Some(w) = &lw.ffn_out_norm {
3908 for bi in 0..b {
3909 inference::rms_norm_into(
3910 &ffn[bi * hs..(bi + 1) * hs],
3911 w,
3912 eps,
3913 norm_style,
3914 &mut post[bi * hs..(bi + 1) * hs],
3915 );
3916 }
3917 ffn.copy_from_slice(&post);
3918 }
3919 for (dst, &f) in h.iter_mut().zip(&ffn) {
3920 *dst += f;
3921 }
3922 if let Some(sc) = lw.layer_scale {
3923 for v in h.iter_mut() {
3924 *v *= sc;
3925 }
3926 }
3927 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
3928 if let Some(t) = tp.parse::<usize>().ok() {
3929 if t >= start_pos && t < start_pos + b {
3930 let bi = t - start_pos;
3931 let row = &h[bi * hs..(bi + 1) * hs];
3932 let n: f32 = row.iter().map(|x| x * x).sum::<f32>().sqrt();
3933 eprintln!(
3934 "BATCH pos {t} after layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
3935 row[0], row[1]
3936 );
3937 }
3938 }
3939 }
3940 if std::env::var("CMF_DEBUG_LAYERS").is_ok() {
3944 let row = &h[(b - 1) * hs..b * hs];
3945 let rms =
3946 (row.iter().map(|&v| (v as f64) * (v as f64)).sum::<f64>() / hs as f64).sqrt();
3947 let mx = row.iter().fold(0f32, |m, &v| m.max(v.abs()));
3948 eprintln!(
3949 "layer {li:>3} {:>10} ffn={:<5} rms={rms:>12.4} max={mx:>12.4}",
3950 match &self.weights.layers[self.phys_layer(li)].attn {
3951 AttnKind::LinearGdn(_) => "gdn",
3952 AttnKind::Linear(_) => "vmf",
3953 AttnKind::ShortConv(_) => "conv",
3954 _ => "attn",
3955 },
3956 match &lw.ffn {
3957 FfnKind::Moe(_) => "moe",
3958 FfnKind::Dense(_) => "dense",
3959 FfnKind::DenseMoe(_) => "dense+moe",
3960 },
3961 );
3962 }
3963 if self.is_loop_end(li) && li + 1 < self.num_layers {
3965 for bi in 0..b {
3966 let normed = inference::rms_norm(
3967 &h[bi * hs..(bi + 1) * hs],
3968 &self.weights.final_norm,
3969 eps,
3970 norm_style,
3971 );
3972 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
3973 }
3974 }
3975 if std::env::var("CMF_TRACE_H").is_ok() {
3976 let n = h[..hs].iter().map(|v| v.abs()).sum::<f32>() / hs as f32;
3977 let mx = h[..hs].iter().fold(0.0f32, |a, &v| a.max(v.abs()));
3978 eprintln!(
3979 "layer {li}: mean|h|={n:.4} max|h|={mx:.2} scale={:?}",
3980 lw.layer_scale
3981 );
3982 }
3983 }
3984 crate::gpu::set_layer(-1); h
3986 }
3987
3988 fn embed_single(&self, id: u32) -> Vec<f32> {
3990 let mut out = vec![0.0f32; self.hidden_size];
3991 if (id as usize) < self.weights.embed_tokens.rows() {
3992 self.weights.embed_tokens.row_f32(id as usize, &mut out);
3993 }
3994 if self.embed_multiplier != 1.0 {
3995 for v in out.iter_mut() {
3996 *v *= self.embed_multiplier;
3997 }
3998 }
3999 if self.dsv4.is_some() {
4003 let mut v = vec![0.0f32; self.hidden_size.max(1)];
4004 v[0] = id as f32;
4005 return v;
4006 }
4007 if let Some(b) = &self.g3n {
4010 return b.0.extend_embedding(id, &out, self.pool.as_deref());
4011 }
4012 out
4013 }
4014
4015 #[cfg(target_os = "macos")]
4021 fn chunk_run_gpu(
4022 &mut self,
4023 li0: usize,
4024 h: &mut [f32],
4025 b: usize,
4026 pos0: usize,
4027 embed_ids: Option<&[u32]>,
4028 cap: usize,
4029 ) -> usize {
4030 if !crate::gpu::enabled_here()
4034 || std::env::var("CMF_GPU_CHUNK")
4035 .map(|v| v == "0")
4036 .unwrap_or(false)
4037 || b < 32
4038 || self.swa.is_some()
4039 || self.global_attn.is_some()
4040 || self.attn_v_norm
4041 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
4042 {
4043 return li0;
4044 }
4045 let Some(model) = self.model.clone() else {
4046 return li0;
4047 };
4048 let inv_freq = self.inv_freq.clone();
4049 let (nh, nkv, hd, hs) = (
4050 self.num_heads,
4051 self.num_kv_heads,
4052 self.head_dim,
4053 self.hidden_size,
4054 );
4055 let loop_end = if self.loop_final_norm {
4059 ((li0 / self.physical_layers) + 1) * self.physical_layers
4060 } else {
4061 self.num_layers
4062 };
4063 let mut layers: Vec<crate::gpu_metal::ChunkLayer> = Vec::new();
4064 let mut stored_at: Vec<usize> = Vec::new();
4065 for li in li0..self.num_layers.min(loop_end).min(cap) {
4066 let lw = &self.weights.layers[self.phys_layer(li)];
4067 if lw.attn_out_norm.is_some() || lw.ffn_out_norm.is_some() || lw.layer_scale.is_some() {
4068 break;
4069 }
4070 let AttnKind::Full {
4071 wq,
4072 wk,
4073 wv,
4074 wo,
4075 q_norm,
4076 k_norm,
4077 output_gate: false,
4078 softplus_gate: None,
4079 bias,
4080 } = &lw.attn
4081 else {
4082 break;
4083 };
4084 let FfnKind::Dense(d) = &lw.ffn else { break };
4085 if d.act != Act::Silu {
4086 break;
4087 }
4088 fn cw(t: &QTensor) -> Option<(usize, usize, usize, &[f32])> {
4093 t.q8_row_parts()
4094 .or_else(|| t.q4t_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4095 .or_else(|| t.q4tp_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4096 }
4097 let parts = (
4098 cw(wq),
4099 cw(wk),
4100 cw(wv),
4101 cw(wo),
4102 cw(&d.gate_proj),
4103 cw(&d.up_proj),
4104 cw(&d.down_proj),
4105 );
4106 let (Some(pq), Some(pk), Some(pv), Some(po), Some(pg), Some(pu), Some(pd)) = parts
4107 else {
4108 break;
4109 };
4110 let layer = &self.kv_cache.layers[li];
4111 if layer.mode != crate::kv_cache::KvMode::F32 || layer.o1.is_some() {
4112 break;
4113 }
4114 stored_at.push(layer.head_len(0));
4115 layers.push(crate::gpu_metal::ChunkLayer {
4116 model: &model,
4117 kv_id: self.graph_kv_id,
4118 layer: li,
4119 wq: pq,
4120 wk: pk,
4121 wv: pv,
4122 wo: po,
4123 gate: pg,
4124 up: pu,
4125 down: pd,
4126 input_norm: &lw.input_norm,
4127 post_norm: &lw.post_norm,
4128 bias: bias
4129 .as_ref()
4130 .map(|(a, bb, cc)| (a.as_slice(), bb.as_slice(), cc.as_slice())),
4131 q_norm: q_norm.as_deref(),
4132 k_norm: k_norm.as_deref(),
4133 inv_freq: &inv_freq,
4134 rd: self.rotary_dim,
4135 nh,
4136 nkv,
4137 hd,
4138 hs,
4139 inter: d.gate_proj.rows(),
4140 gemma: matches!(self.norm_style, cortiq_core::NormStyle::Gemma),
4141 eps: self.rms_eps as f32,
4142 });
4143 }
4144 if layers.is_empty() {
4145 return li0;
4146 }
4147 let row = nkv * hd;
4148 let mut store: Vec<(Vec<f32>, Vec<f32>, Vec<f32>)> = stored_at
4149 .iter()
4150 .map(|&st| (vec![0f32; b * row], vec![0f32; b * row], vec![0f32; st + b]))
4151 .collect();
4152 let mut io: Vec<crate::gpu_metal::ChunkIo> = Vec::with_capacity(layers.len());
4153 for (i, (ok, ov, oi)) in store.iter_mut().enumerate() {
4154 let li = layers[i].layer;
4155 let layer = &self.kv_cache.layers[li];
4156 io.push(crate::gpu_metal::ChunkIo {
4157 cpu_stored: stored_at[i],
4158 cpu_k: (0..nkv).map(|g| layer.head_keys(g)).collect(),
4159 cpu_v: (0..nkv).map(|g| layer.head_values(g)).collect(),
4160 out_k: ok,
4161 out_v: ov,
4162 imp: oi,
4163 });
4164 }
4165 let n_run = layers.len();
4166 let last = layers.last().map(|l| l.layer + 1).unwrap_or(li0);
4167 let ep = embed_ids.and_then(|ids| {
4170 self.weights
4171 .embed_tokens
4172 .q8_row_parts()
4173 .map(|(idx, rows, _c, rs)| crate::gpu_metal::ChunkEmbed {
4174 idx,
4175 rows,
4176 row_scale: rs,
4177 ids,
4178 mult: self.embed_multiplier,
4179 })
4180 });
4181 if embed_ids.is_some() && ep.is_none() {
4182 return li0;
4183 }
4184 if !crate::gpu_metal::chunk_run_gpu(&layers, &mut io, h, b, pos0, ep.as_ref()) {
4185 return li0;
4186 }
4187 drop(io);
4188 drop(layers);
4189 for (i, (ok, ov, oi)) in store.iter().enumerate().take(n_run) {
4192 let li = li0 + i;
4193 let layer = &mut self.kv_cache.layers[li];
4194 for bi in 0..b {
4195 layer.append(
4196 &ok[bi * row..(bi + 1) * row],
4197 &ov[bi * row..(bi + 1) * row],
4198 &[],
4199 );
4200 }
4201 layer.accumulate_imp(oi);
4202 }
4203 last
4204 }
4205
4206 fn layer_is_local(&self, li: usize) -> bool {
4209 if let Some(layers) = &self.sliding_layers {
4210 return layers.get(li).copied().unwrap_or(false);
4211 }
4212 match self.swa {
4213 Some((_, pattern)) => (li + 1) % pattern.max(1) != 0,
4214 None => false,
4215 }
4216 }
4217
4218 fn layer_inv_freq(&self, li: usize) -> std::sync::Arc<Vec<f32>> {
4221 if self.layer_is_local(li) {
4222 if let Some(f) = &self.inv_freq_local {
4223 return f.clone();
4224 }
4225 } else if let Some(f) = &self.inv_freq_global {
4226 return f.clone();
4227 }
4228 self.inv_freq.clone()
4229 }
4230
4231 fn layer_window(&self, li: usize) -> Option<usize> {
4233 self.swa
4234 .and_then(|(w, _)| self.layer_is_local(li).then_some(w))
4235 }
4236
4237 fn layer_num_heads(&self, li: usize) -> usize {
4238 self.attention_heads_per_layer
4239 .as_ref()
4240 .and_then(|v| v.get(li).copied())
4241 .unwrap_or(self.num_heads)
4242 }
4243
4244 fn layer_rope_scale(&self, li: usize) -> f32 {
4245 if self.layer_is_local(li) {
4246 self.rope_scale_local
4247 } else {
4248 self.rope_scale
4249 }
4250 }
4251
4252 fn layer_geom(&self, li: usize) -> (usize, usize, usize) {
4255 if !self.layer_is_local(li) {
4256 if let Some((ghd, gkv)) = self.global_attn {
4257 return (gkv, ghd, ghd);
4258 }
4259 }
4260 (
4261 self.num_kv_heads,
4262 self.head_dim,
4263 if self.layer_is_local(li) {
4264 self.rotary_dim_local.unwrap_or(self.rotary_dim)
4265 } else {
4266 self.rotary_dim
4267 },
4268 )
4269 }
4270
4271 fn forward_layers(
4273 &mut self,
4274 hidden: &[f32],
4275 position: usize,
4276 task_mask: Option<&TaskMask>,
4277 ) -> Vec<f32> {
4278 self.forward_layers_upto(hidden, position, task_mask, None)
4279 }
4280
4281 pub fn embed_id(&self, id: u32) -> Vec<f32> {
4289 self.embed_single(id)
4290 }
4291
4292 pub fn split_supported(&self) -> Result<(), String> {
4296 if self.dsv4.is_some() {
4297 return Err(
4298 "network split: DeepSeek-V4 runs its own fused stack (not splittable yet)".into(),
4299 );
4300 }
4301 if self.g3n.is_some() {
4302 return Err(
4303 "network split: Gemma-3n runs its own AltUp stack (not splittable yet)".into(),
4304 );
4305 }
4306 Ok(())
4307 }
4308
4309 pub fn forward_span(
4314 &mut self,
4315 hidden: &[f32],
4316 position: usize,
4317 from: usize,
4318 upto: usize,
4319 task_mask: Option<&TaskMask>,
4320 ) -> Result<Vec<f32>, String> {
4321 self.split_supported()?;
4322 if from > upto || upto >= self.num_layers {
4323 return Err(format!(
4324 "forward_span: layer range {from}..={upto} outside 0..{}",
4325 self.num_layers
4326 ));
4327 }
4328 if hidden.len() != self.hidden_size {
4329 return Err(format!(
4330 "forward_span: hidden len {} ≠ hidden_size {}",
4331 hidden.len(),
4332 self.hidden_size
4333 ));
4334 }
4335 Ok(self.forward_layers_span(hidden, position, task_mask, from, Some(upto)))
4336 }
4337
4338 pub fn logits_from_hidden(&mut self, hidden: &[f32]) -> Vec<f32> {
4341 let normed = inference::rms_norm(
4342 hidden,
4343 &self.weights.final_norm,
4344 self.rms_eps,
4345 self.norm_style,
4346 );
4347 self.lm_head_forward(&normed)
4348 }
4349
4350 pub fn sample_next(&mut self, logits: &[f32], past_tokens: &[u32]) -> u32 {
4352 sampler::sample_with_scratch(
4353 logits,
4354 &self.sampler_config,
4355 past_tokens,
4356 &mut self.rng,
4357 &mut self.sampler_scratch,
4358 )
4359 }
4360
4361 pub fn reset_session(&mut self) {
4363 self.kv_cache.clear();
4364 self.kv_history.clear();
4365 crate::gpu::graph_kv_reset(self.graph_kv_id);
4366 }
4367
4368 pub fn prefill_span_ids(
4374 &mut self,
4375 ids: &[u32],
4376 start_pos: usize,
4377 upto: usize,
4378 task_mask: Option<&TaskMask>,
4379 ) -> Result<Vec<f32>, String> {
4380 self.split_supported()?;
4381 if upto >= self.num_layers {
4382 return Err(format!(
4383 "prefill_span_ids: upto {upto} outside 0..{}",
4384 self.num_layers
4385 ));
4386 }
4387 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4391 Ok(self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, upto + 1))
4392 } else {
4393 let hs = self.hidden_size;
4394 let mut out = Vec::with_capacity(ids.len() * hs);
4395 for (i, &id) in ids.iter().enumerate() {
4396 let emb = self.embed_id(id);
4397 out.extend_from_slice(&self.forward_span(
4398 &emb,
4399 start_pos + i,
4400 0,
4401 upto,
4402 task_mask,
4403 )?);
4404 }
4405 Ok(out)
4406 }
4407 }
4408
4409 pub fn prefill_span_hidden(
4412 &mut self,
4413 hidden: &[f32],
4414 start_pos: usize,
4415 from: usize,
4416 upto: usize,
4417 task_mask: Option<&TaskMask>,
4418 ) -> Result<Vec<f32>, String> {
4419 self.split_supported()?;
4420 let hs = self.hidden_size;
4421 if hidden.is_empty() || hidden.len() % hs != 0 {
4422 return Err(format!(
4423 "prefill_span_hidden: {} floats is not a multiple of hidden {hs}",
4424 hidden.len()
4425 ));
4426 }
4427 if from > upto || upto >= self.num_layers {
4428 return Err(format!(
4429 "prefill_span_hidden: layer range {from}..={upto} outside 0..{}",
4430 self.num_layers
4431 ));
4432 }
4433 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4434 Ok(self.prefill_batch_span(
4435 PrefillIn::Hidden(hidden),
4436 start_pos,
4437 task_mask,
4438 from,
4439 upto + 1,
4440 ))
4441 } else {
4442 let b = hidden.len() / hs;
4443 let mut out = Vec::with_capacity(hidden.len());
4444 for i in 0..b {
4445 let h = self.forward_span(
4446 &hidden[i * hs..(i + 1) * hs],
4447 start_pos + i,
4448 from,
4449 upto,
4450 task_mask,
4451 )?;
4452 out.extend_from_slice(&h);
4453 }
4454 Ok(out)
4455 }
4456 }
4457
4458 fn try_token_graph_wgpu(
4462 &self,
4463 hidden: &[f32],
4464 position: usize,
4465 logits_out: &mut Vec<f32>,
4466 layers_run: &mut usize,
4467 ) -> Option<Vec<f32>> {
4468 self.try_token_graph_wgpu_steps(
4469 hidden,
4470 position,
4471 logits_out,
4472 1,
4473 None,
4474 Some(layers_run),
4475 0,
4476 self.num_layers,
4477 )
4478 }
4479
4480 fn try_token_graph_wgpu_span(
4484 &self,
4485 hidden: &[f32],
4486 position: usize,
4487 logits_out: &mut Vec<f32>,
4488 from: usize,
4489 upto_excl: usize,
4490 layers_run: &mut usize,
4491 ) -> Option<Vec<f32>> {
4492 self.try_token_graph_wgpu_steps(
4493 hidden,
4494 position,
4495 logits_out,
4496 1,
4497 None,
4498 Some(layers_run),
4499 from,
4500 upto_excl,
4501 )
4502 }
4503
4504 fn try_multi_burst(&self, t_next: u32, position: usize, k: usize) -> Option<Vec<u32>> {
4508 if self.o1_active() || self.attn_softcap > 0.0 {
4509 return None;
4510 }
4511 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
4512 if !graph_on || crate::gpu::graph_unsupported() {
4513 return None;
4520 }
4521 let emb = self.embed_single(t_next);
4522 let mut lg = Vec::new();
4523 let mut ids = Vec::new();
4524 self.try_token_graph_wgpu_steps(
4525 &emb,
4526 position,
4527 &mut lg,
4528 k,
4529 Some(&mut ids),
4530 None,
4531 0,
4532 self.num_layers,
4533 )?;
4534 (ids.len() == k).then_some(ids)
4535 }
4536
4537 fn try_token_graph_wgpu_steps(
4541 &self,
4542 hidden: &[f32],
4543 position: usize,
4544 logits_out: &mut Vec<f32>,
4545 steps: usize,
4546 ids_out: Option<&mut Vec<u32>>,
4547 layers_run: Option<&mut usize>,
4548 from: usize,
4549 upto_excl: usize,
4550 ) -> Option<Vec<f32>> {
4551 let o1_gpu = std::env::var("CMF_O1_GPU").as_deref() == Ok("1");
4554 if (self.o1_active() && !o1_gpu) || self.attn_softcap > 0.0 {
4555 return None;
4559 }
4560 let o1_views: Vec<Option<Vec<crate::nystrom::O1DeviceView<'_>>>> = (from..upto_excl)
4565 .map(|li| {
4566 if !o1_gpu {
4567 return None;
4568 }
4569 self.kv_cache.layers[self.phys_layer(li)].o1_views()
4570 })
4571 .collect();
4572 if self.o1_active() && o1_gpu {
4573 let want: usize = (from..upto_excl)
4576 .filter(|li| !matches!(self.kv_cache.layers[self.phys_layer(*li)].o1, None))
4577 .count();
4578 let have = o1_views.iter().filter(|v| v.is_some()).count();
4579 if want == 0 || have != want {
4580 use std::sync::atomic::{AtomicUsize, Ordering};
4590 static LAST: AtomicUsize = AtomicUsize::new(usize::MAX);
4591 let code = have * 1000 + want;
4592 if LAST.swap(code, Ordering::Relaxed) != code {
4593 tracing::warn!(
4594 "o1 graph: {have} of {want} layers sealed — per-op until all seal"
4595 );
4596 }
4597 return None;
4598 }
4599 }
4600 let nh = self.num_heads;
4601 let (nkv, hd, rd) = self.layer_geom(0);
4602 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
4603 let mut layers = Vec::with_capacity(upto_excl - from);
4604 let mut model = None;
4605 let dbg = std::env::var("CMF_GRAPH_DEBUG").is_ok();
4606 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
4607 if let Some((_, i, kind, rs)) = t.graph_weight() {
4608 return Some(crate::gpu::GraphW {
4609 idx: i,
4610 kind,
4611 row_scale: rs,
4612 data: &[],
4613 });
4614 }
4615 t.as_f32().map(|d| crate::gpu::GraphW {
4617 idx: 0,
4618 kind: 4,
4619 row_scale: &[],
4620 data: d,
4621 })
4622 }
4623 for li in from..upto_excl {
4624 let lw = &self.weights.layers[self.phys_layer(li)];
4625 if dbg {
4626 let ak = match &lw.attn {
4627 AttnKind::Mla(_) => "Mla".into(),
4628 AttnKind::Full {
4629 output_gate, bias, ..
4630 } => format!("Full gate={output_gate} bias={}", bias.is_some()),
4631 AttnKind::LinearGdn(_) => "LinearGdn".into(),
4632 AttnKind::Kda(_) => "Kda".into(),
4633 AttnKind::Linear(_) => "Linear".into(),
4634 AttnKind::ShortConv(_) => "ShortConv".into(),
4635 };
4636 let fk = match &lw.ffn {
4637 FfnKind::Dense(_) => "Dense",
4638 FfnKind::Moe(_) => "Moe",
4639 FfnKind::DenseMoe(_) => "DenseMoe",
4640 };
4641 eprintln!("graph L{li}: attn={ak} ffn={fk}");
4642 }
4643 let gffn = match &lw.ffn {
4644 FfnKind::DenseMoe(_) => return None, FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
4646 gate: gw(&d.gate_proj)?,
4647 up: gw(&d.up_proj)?,
4648 down: gw(&d.down_proj)?,
4649 },
4650 FfnKind::Moe(m) => {
4651 if m.router_sigmoid
4656 || m.expert_bias.is_some()
4657 || m.route_tau.is_some()
4658 || m.mask.is_some()
4659 {
4660 return None;
4661 }
4662 let (se, sg) = m.shared.as_ref()?;
4663 let sgate = gw(sg.as_ref()?)?;
4664 let router = gw(&m.router)?;
4665 let inter = m.experts.first()?.gate_proj.rows();
4666 let mut experts = Vec::with_capacity(m.experts.len() + 1);
4667 let mut q4tp: Option<bool> = None;
4670 let mut gu_q2: Option<bool> = None;
4673 for e in m.experts.iter().chain(std::iter::once(se)) {
4674 if !matches!(e.act, Act::Silu)
4675 || e.gate_proj.rows() != inter
4676 || e.up_proj.rows() != inter
4677 {
4678 return None;
4679 }
4680 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
4681 Some((mm, gi)) => (
4682 mm,
4683 gi,
4684 e.up_proj.mapped_q4t()?.1,
4685 e.down_proj.mapped_q4t()?.1,
4686 false,
4687 false,
4688 ),
4689 None => match e.gate_proj.mapped_q2tp() {
4690 Some((mm, gi)) => (
4691 mm,
4692 gi,
4693 e.up_proj.mapped_q2tp()?.1,
4694 e.down_proj.mapped_q4tp()?.1,
4695 true,
4696 true,
4697 ),
4698 None => {
4699 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
4700 (
4701 mm,
4702 gi,
4703 e.up_proj.mapped_q4tp()?.1,
4704 e.down_proj.mapped_q4tp()?.1,
4705 true,
4706 false,
4707 )
4708 }
4709 },
4710 };
4711 if *q4tp.get_or_insert(is_p) != is_p || *gu_q2.get_or_insert(is_q2) != is_q2
4712 {
4713 tracing::warn!(
4719 "MoE layer mixes expert layouts (q4tp={is_p}, q2tp gate/up={is_q2}) — every expert of a layer, INCLUDING the shared one, must share a layout. The whole-token graph declines this layer."
4720 );
4721 return None;
4722 }
4723 model.get_or_insert_with(|| mm.clone());
4724 experts.push((gi, ui, di));
4725 }
4726 crate::gpu::GraphFfn::Moe {
4727 router,
4728 shared_gate: sgate,
4729 experts,
4730 n_exp: m.experts.len(),
4731 top_k: std::env::var("CMF_TOPK_PROBE")
4737 .ok()
4738 .and_then(|v| v.parse::<usize>().ok())
4739 .filter(|k| *k > 0 && *k <= m.top_k)
4740 .unwrap_or(m.top_k),
4741 inter,
4742 norm_topk: m.norm_topk_prob,
4743 q4tp: q4tp?,
4744 gu_q2: gu_q2.unwrap_or(false),
4745 }
4746 }
4747 };
4748 let attn = match &lw.attn {
4749 AttnKind::Full {
4750 wq,
4751 wk,
4752 wv,
4753 wo,
4754 q_norm,
4755 k_norm,
4756 output_gate,
4757 softplus_gate,
4758 bias,
4759 } => {
4760 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
4761 return None;
4762 }
4763 let (m, _, _, _) = wq.graph_weight()?;
4764 model = Some(m.clone());
4765 crate::gpu::GraphAttn::Full {
4766 wq: gw(wq)?,
4767 wk: gw(wk)?,
4768 wv: gw(wv)?,
4769 wo: gw(wo)?,
4770 q_norm: q_norm.as_deref(),
4771 k_norm: k_norm.as_deref(),
4772 bias: bias
4773 .as_ref()
4774 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
4775 output_gate: *output_gate,
4776 cpu_k: self.kv_cache.layers[li].k_heads(),
4777 cpu_v: self.kv_cache.layers[li].v_heads(),
4778 }
4779 }
4780 AttnKind::LinearGdn(w) => {
4781 let cfg = self.gdn_cfg?;
4782 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
4783 model = Some(m.clone());
4784 crate::gpu::GraphAttn::Gdn {
4785 qkv: gw(&w.in_proj_qkv)?,
4786 z: gw(&w.in_proj_z)?,
4787 a: gw(&w.in_proj_a)?,
4788 b: gw(&w.in_proj_b)?,
4789 out: gw(&w.out_proj)?,
4790 conv1d: &w.conv1d,
4791 a_log: &w.a_log,
4792 dt_bias: &w.dt_bias,
4793 norm: &w.norm,
4794 nv: cfg.num_v_heads,
4795 nk: cfg.num_k_heads,
4796 dk: cfg.key_head_dim,
4797 dv: cfg.value_head_dim,
4798 kk: cfg.conv_kernel,
4799 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
4800 }
4801 }
4802 _ => return None,
4803 };
4804 layers.push(crate::gpu::GraphLayer {
4805 input_norm: &lw.input_norm,
4806 attn,
4807 post_norm: &lw.post_norm,
4808 ffn: gffn,
4809 });
4810 }
4811 let model = model?;
4812 let lm_gw = if upto_excl == self.num_layers
4818 && self.graph_want_logits
4819 && std::env::var("CMF_GPU_LMHEAD")
4820 .map(|v| v != "0")
4821 .unwrap_or(true)
4822 {
4823 self.weights.lm_head.graph_weight().map(|(_, i, kind, rs)| {
4824 (
4825 crate::gpu::GraphW {
4826 idx: i,
4827 kind,
4828 row_scale: rs,
4829 data: &[],
4830 },
4831 self.weights.lm_head.rows(),
4832 )
4833 })
4834 } else {
4835 None
4836 };
4837 let lm = lm_gw.as_ref().map(|(gw, rows)| (gw, *rows));
4838 let emb_gw = if steps > 1 {
4840 self.weights
4841 .embed_tokens
4842 .graph_weight()
4843 .map(|(_, i, kind, rs)| {
4844 (
4845 crate::gpu::GraphW {
4846 idx: i,
4847 kind,
4848 row_scale: rs,
4849 data: &[],
4850 },
4851 self.weights.embed_tokens.rows(),
4852 self.embed_multiplier as f32,
4853 )
4854 })
4855 } else {
4856 None
4857 };
4858
4859 let loop_norm_at: Vec<usize> = if self.loop_final_norm {
4865 (from..upto_excl.min(self.num_layers - 1))
4866 .filter(|&li| (li + 1) % self.physical_layers == 0)
4867 .map(|li| li - from)
4868 .collect()
4869 } else {
4870 Vec::new()
4871 };
4872 let mut h = hidden.to_vec();
4873 crate::gpu::forward_token_graph(
4874 &model,
4875 self.graph_kv_id,
4876 &layers,
4877 &o1_views,
4878 self.o1_epoch,
4879 &self.inv_freq,
4880 &mut h,
4881 nh,
4882 nkv,
4883 hd,
4884 rd,
4885 self.hidden_size,
4886 self.intermediate_size,
4887 position,
4888 self.kv_cache.max_seq_len,
4889 gemma,
4890 self.rms_eps as f32,
4891 lm,
4892 &self.weights.final_norm,
4893 logits_out,
4894 &loop_norm_at,
4895 steps,
4896 emb_gw.as_ref().map(|(gw, rows, m)| (gw, *rows, *m)),
4897 ids_out,
4898 layers_run,
4899 from,
4900 )
4901 .then_some(h)
4902 }
4903
4904 fn try_batch_graph_wgpu(
4909 &self,
4910 hiddens: &mut [f32],
4911 positions: &[usize],
4912 k: usize,
4913 spec: Option<crate::gpu::SpecTail<'_>>,
4914 ) -> bool {
4915 let _tb = std::time::Instant::now();
4916 if self.attn_softcap > 0.0 {
4917 return false; }
4919 if self.o1_active() {
4920 return false;
4921 }
4922 let nh = self.num_heads;
4923 let (nkv, hd, rd) = self.layer_geom(0);
4924 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
4925 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
4926 if let Some((_, i, kind, rs)) = t.graph_weight() {
4927 return Some(crate::gpu::GraphW {
4928 idx: i,
4929 kind,
4930 row_scale: rs,
4931 data: &[],
4932 });
4933 }
4934 t.as_f32().map(|d| crate::gpu::GraphW {
4935 idx: 0,
4936 kind: 4,
4937 row_scale: &[],
4938 data: d,
4939 })
4940 }
4941 let built: Option<(
4942 Vec<crate::gpu::GraphLayer<'_>>,
4943 std::sync::Arc<cortiq_core::CmfModel>,
4944 )> = (|| {
4945 let mut layers = Vec::with_capacity(self.num_layers);
4946 let mut model = None;
4947 for li in 0..self.num_layers {
4948 let lw = &self.weights.layers[self.phys_layer(li)];
4949 let gffn = match &lw.ffn {
4956 FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
4957 gate: gw(&d.gate_proj)?,
4958 up: gw(&d.up_proj)?,
4959 down: gw(&d.down_proj)?,
4960 },
4961 FfnKind::Moe(m) => {
4962 if m.router_sigmoid
4963 || m.expert_bias.is_some()
4964 || m.route_tau.is_some()
4965 || m.mask.is_some()
4966 {
4967 return None;
4968 }
4969 let (se, sg) = m.shared.as_ref()?;
4970 let sgate = gw(sg.as_ref()?)?;
4971 let router = gw(&m.router)?;
4972 let inter = m.experts.first()?.gate_proj.rows();
4973 let mut experts = Vec::with_capacity(m.experts.len() + 1);
4974 let mut q4tp: Option<bool> = None;
4975 let mut gu_q2: Option<bool> = None;
4976 for e in m.experts.iter().chain(std::iter::once(se)) {
4977 if !matches!(e.act, Act::Silu)
4978 || e.gate_proj.rows() != inter
4979 || e.up_proj.rows() != inter
4980 {
4981 return None;
4982 }
4983 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
4987 Some((mm, gi)) => (
4988 mm,
4989 gi,
4990 e.up_proj.mapped_q4t()?.1,
4991 e.down_proj.mapped_q4t()?.1,
4992 false,
4993 false,
4994 ),
4995 None => match e.gate_proj.mapped_q2tp() {
4996 Some((mm, gi)) => (
4997 mm,
4998 gi,
4999 e.up_proj.mapped_q2tp()?.1,
5000 e.down_proj.mapped_q4tp()?.1,
5001 true,
5002 true,
5003 ),
5004 None => {
5005 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
5006 (
5007 mm,
5008 gi,
5009 e.up_proj.mapped_q4tp()?.1,
5010 e.down_proj.mapped_q4tp()?.1,
5011 true,
5012 false,
5013 )
5014 }
5015 },
5016 };
5017 if *q4tp.get_or_insert(is_p) != is_p
5018 || *gu_q2.get_or_insert(is_q2) != is_q2
5019 {
5020 return None;
5021 }
5022 model.get_or_insert_with(|| mm.clone());
5023 experts.push((gi, ui, di));
5024 }
5025 crate::gpu::GraphFfn::Moe {
5026 router,
5027 shared_gate: sgate,
5028 experts,
5029 n_exp: m.experts.len(),
5030 top_k: m.top_k,
5031 inter,
5032 norm_topk: m.norm_topk_prob,
5033 q4tp: q4tp?,
5034 gu_q2: gu_q2.unwrap_or(false),
5035 }
5036 }
5037 _ => return None,
5038 };
5039 let attn = match &lw.attn {
5040 AttnKind::Full {
5041 wq,
5042 wk,
5043 wv,
5044 wo,
5045 q_norm,
5046 k_norm,
5047 output_gate,
5048 softplus_gate,
5049 bias,
5050 } => {
5051 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
5052 return None;
5053 }
5054 let (m, _, _, _) = wq.graph_weight()?;
5055 model = Some(m.clone());
5056 crate::gpu::GraphAttn::Full {
5057 wq: gw(wq)?,
5058 wk: gw(wk)?,
5059 wv: gw(wv)?,
5060 wo: gw(wo)?,
5061 q_norm: q_norm.as_deref(),
5062 k_norm: k_norm.as_deref(),
5063 bias: bias
5064 .as_ref()
5065 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
5066 output_gate: *output_gate,
5067 cpu_k: self.kv_cache.layers[li].k_heads(),
5068 cpu_v: self.kv_cache.layers[li].v_heads(),
5069 }
5070 }
5071 AttnKind::LinearGdn(w) => {
5072 let cfg = self.gdn_cfg?;
5073 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
5074 model = Some(m.clone());
5075 crate::gpu::GraphAttn::Gdn {
5076 qkv: gw(&w.in_proj_qkv)?,
5077 z: gw(&w.in_proj_z)?,
5078 a: gw(&w.in_proj_a)?,
5079 b: gw(&w.in_proj_b)?,
5080 out: gw(&w.out_proj)?,
5081 conv1d: &w.conv1d,
5082 a_log: &w.a_log,
5083 dt_bias: &w.dt_bias,
5084 norm: &w.norm,
5085 nv: cfg.num_v_heads,
5086 nk: cfg.num_k_heads,
5087 dk: cfg.key_head_dim,
5088 dv: cfg.value_head_dim,
5089 kk: cfg.conv_kernel,
5090 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
5091 }
5092 }
5093 _ => return None,
5094 };
5095 layers.push(crate::gpu::GraphLayer {
5096 input_norm: &lw.input_norm,
5097 attn,
5098 post_norm: &lw.post_norm,
5099 ffn: gffn,
5100 });
5101 }
5102 Some((layers, model?))
5103 })();
5104 let Some((layers, model)) = built else {
5105 {
5106 use std::sync::atomic::{AtomicBool, Ordering};
5107 static SAID: AtomicBool = AtomicBool::new(false);
5108 if !SAID.swap(true, Ordering::Relaxed) {
5109 tracing::warn!("batch graph: BUILDER refused (layer weights/kinds)");
5110 }
5111 }
5112 return false;
5113 };
5114 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
5115 eprintln!("batch-build: {:.1} ms", _tb.elapsed().as_secs_f64() * 1e3);
5116 }
5117 crate::gpu::forward_batch_graph(
5118 &model,
5119 self.graph_kv_id,
5120 &layers,
5121 &self.inv_freq,
5122 hiddens,
5123 nh,
5124 nkv,
5125 hd,
5126 rd,
5127 self.hidden_size,
5128 self.intermediate_size,
5129 positions,
5130 self.kv_cache.max_seq_len,
5131 gemma,
5132 self.rms_eps as f32,
5133 k,
5134 spec,
5135 )
5136 }
5137
5138 fn draft_probe() -> bool {
5142 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5143 *ON.get_or_init(|| std::env::var("CMF_DSV4_DRAFT_PROBE").is_ok_and(|v| v != "0"))
5144}
5145
5146 #[cfg(feature = "gpu")]
5158 fn dsv4_spec_on() -> bool {
5159 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5160 *ON.get_or_init(|| std::env::var("CMF_DSV4_SPEC").map(|v| v != "0").unwrap_or(true))
5161 }
5162
5163 #[cfg(feature = "gpu")]
5170 fn dsv4_spec_step(
5171 &mut self,
5172 tip_token: u32,
5173 t_next: u32,
5174 next_pos: usize,
5175 drafted: &mut usize,
5176 accepted_ctr: &mut usize,
5177 ) -> Option<(Vec<u32>, usize)> {
5178 let t_all = std::time::Instant::now();
5179 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5180 thread_local! {
5181 static LAST: std::cell::Cell<Option<std::time::Instant>> =
5182 const { std::cell::Cell::new(None) };
5183 }
5184 LAST.with(|l| {
5185 if let Some(prev) = l.get() {
5186 eprintln!("между раундами {:.1} мс", prev.elapsed().as_secs_f64() * 1e3);
5187 }
5188 l.set(Some(std::time::Instant::now()));
5189 });
5190 }
5191 if std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5192 eprintln!("spec_step: вход pos={next_pos}");
5193 }
5194 let n_layers = self.dsv4.as_ref().map(|b| b.1.len())?;
5195 let cfg = self.dsv4.as_ref().map(|b| b.2)?;
5196 if self.dspark.is_none() {
5198 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5199 if t.is_empty() {
5200 return None;
5201 }
5202 crate::dsv4::dspark_arm(&t, cfg.dim);
5203 self.dspark = Some(crate::dsv4::DsparkState::new(
5204 self.dsv4_mtp.len(),
5205 &cfg,
5206 t.len(),
5207 ));
5208 }
5209 let targets = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5210 let pack = crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg);
5211 if pack.is_none() && std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5212 eprintln!("spec_step: пак не построился (targets {targets:?})");
5213 }
5214 let pack = pack?;
5215 let block = crate::dsv4::dspark_block();
5216 let b_box = self.dsv4.as_mut()?;
5217 let (g, layers, st) = (&b_box.0, &b_box.1, &mut b_box.3);
5218 let ds = self.dspark.as_mut()?;
5219 let dbg = std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok();
5222 if !crate::dsv4::dspark_take(&mut ds.main_hidden) && !ds.have_hidden {
5223 if dbg {
5224 eprintln!("spec_step: нет захвата");
5225 }
5226 return None;
5227 }
5228 ds.have_hidden = true;
5229 let tip_pos = next_pos.checked_sub(1)?;
5230 let draft_started = std::time::Instant::now();
5231 let mut conf = Vec::new();
5232 let props = crate::dsv4::dspark_draft_gpu(
5233 g,
5234 &self.dsv4_mtp,
5235 &cfg,
5236 ds,
5237 pack,
5238 st.kv_id,
5239 tip_token,
5240 tip_pos,
5241 self.pool.as_deref(),
5242 &mut conf,
5243 );
5244 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5245 *drafted += block;
5246 if props.is_empty() || props[0] != t_next {
5247 if dbg {
5248 eprintln!(
5249 "spec_step: черновик {} (props0={:?} t_next={t_next})",
5250 if props.is_empty() { "пуст" } else { "мимо" },
5251 props.first()
5252 );
5253 }
5254 return None;
5255 }
5256 let mut k_verify = crate::dsv4::dspark_verify_k().min(props.len());
5257 let conf_min = {
5263 static M: std::sync::OnceLock<f32> = std::sync::OnceLock::new();
5264 *M.get_or_init(|| {
5265 std::env::var("CMF_DSPARK_CONF_MIN")
5266 .ok()
5267 .and_then(|v| v.parse().ok())
5268 .unwrap_or(0.0)
5269 })
5270 };
5271 if conf_min > 0.0 && conf.len() >= props.len() {
5272 let mut keep = 1usize;
5273 while keep < k_verify && conf.get(keep).copied().unwrap_or(0.0) >= conf_min {
5274 keep += 1;
5275 }
5276 k_verify = k_verify.min(keep.max(2));
5277 }
5278 if k_verify < 2 {
5279 return None;
5280 }
5281 let mut fed = Vec::with_capacity(k_verify);
5282 fed.push(t_next);
5283 fed.extend_from_slice(&props[1..k_verify]);
5284 let mut argmax = Vec::new();
5285 let mut logits_all = Vec::new();
5286 let mut walked = Vec::new();
5287 let txn = crate::dsv4::dsv4_verify_chunk(
5288 g,
5289 layers,
5290 &cfg,
5291 st,
5292 &fed,
5293 next_pos,
5294 &self.inv_freq,
5295 self.pool.as_deref(),
5296 &targets,
5297 &mut argmax,
5298 &mut logits_all,
5299 &mut walked,
5300 );
5301 if txn.is_none() && dbg {
5302 eprintln!("spec_step: verify отказал");
5303 }
5304 let txn = txn?;
5305 let b = fed.len();
5306 let mut accepted = 1usize;
5307 while accepted < b && fed[accepted] == argmax[accepted - 1] {
5308 accepted += 1;
5309 }
5310 if std::env::var("CMF_DSV4_SPEC_FORCE_REJECT").is_ok_and(|v| v != "0") {
5315 accepted = 1;
5316 }
5317 if std::env::var("CMF_DSV4_SPEC_TRACE").is_ok() {
5318 eprintln!(
5319 "spec@{next_pos}: fed={fed:?} argmax={argmax:?} accepted={accepted}"
5320 );
5321 }
5322 let t_fin = std::time::Instant::now();
5323 if !crate::dsv4::dsv4_spec_finish(
5324 g,
5325 layers,
5326 &cfg,
5327 st,
5328 txn,
5329 accepted,
5330 &fed,
5331 &self.inv_freq,
5332 self.pool.as_deref(),
5333 ) {
5334 tracing::warn!("dsv4: спекулятивный откат не удался — состояние подозрительно");
5335 return None;
5336 }
5337 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5338 eprintln!("finish(k={accepted}): {:.1} мс", t_fin.elapsed().as_secs_f64() * 1e3);
5339 }
5340 *accepted_ctr += accepted - 1;
5341 let (hc, dim) = (cfg.hc_mult, cfg.dim);
5346 let dev_caps: Vec<usize> = targets
5353 .iter()
5354 .copied()
5355 .filter(|&t| {
5356 st.dev_set.get(t).copied().unwrap_or(false)
5357 && !st.partial_set.get(t).copied().unwrap_or(false)
5358 })
5359 .collect();
5360 let mut caps_all = vec![0.0f32; dev_caps.len() * b * hc * dim];
5361 if !crate::gpu_wgpu::dsv4_spec_cap_read_all(b, dev_caps.len(), hc * dim, &mut caps_all) {
5362 return None;
5363 }
5364 for t in 0..accepted {
5365 let tip = t + 1 == accepted;
5366 for (slot, &tl) in targets.iter().enumerate() {
5367 if let Some(di) = dev_caps.iter().position(|&d| d == tl) {
5368 let lo = (di * b + t) * hc * dim;
5369 crate::dsv4::dspark_capture(
5370 &caps_all[lo..lo + hc * dim],
5371 &cfg,
5372 slot,
5373 &mut ds.main_hidden,
5374 );
5375 } else if tip
5376 && crate::dsv4::dspark_peek_slot(slot, dim, {
5377 let lo = slot * dim;
5378 &mut ds.main_hidden[lo..lo + dim]
5379 })
5380 {
5381 } else {
5386 crate::dsv4::dspark_capture(
5390 &walked[t * hc * dim..(t + 1) * hc * dim],
5391 &cfg,
5392 slot,
5393 &mut ds.main_hidden,
5394 );
5395 }
5396 }
5397 crate::dsv4::dspark_ring_append(g, &self.dsv4_mtp, &cfg, ds, next_pos + t, self.pool.as_deref());
5398 }
5399 let row = logits_all[(accepted - 1) * cfg.vocab..accepted * cfg.vocab].to_vec();
5400 self.graph_logits = Some(row);
5401 if std::env::var("CMF_DSV4_TRUNK_PICK_DUMP").is_ok() {
5406 crate::dsv4::trunk_freq_note(&crate::dsv4::pick_tally_take());
5407 crate::dsv4::pick_tally_arm();
5408 }
5409 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5410 eprintln!("spec_step total {:.1} мс (k={accepted})", t_all.elapsed().as_secs_f64() * 1e3);
5411 }
5412 Some((fed[1..accepted].to_vec(), next_pos + accepted))
5413 }
5414
5415 fn dspark_probe(&mut self, position: usize, token_id: u32) {
5416 if self.dsv4_mtp.is_empty() || !Self::draft_probe() {
5417 return;
5418 }
5419 let trunk_now = crate::dsv4::pick_tally_take();
5421 crate::dsv4::trunk_freq_note(&trunk_now);
5422 if !trunk_now.is_empty() {
5423 self.dspark_trunk_picks.push(trunk_now);
5424 let keep = crate::dsv4::dspark_block();
5425 if self.dspark_trunk_picks.len() > keep {
5426 self.dspark_trunk_picks.remove(0);
5427 }
5428 }
5429 for p in std::mem::take(&mut self.dspark_pending) {
5432 let Some(i) = position.checked_sub(p.0 + 1) else {
5433 continue;
5434 };
5435 let mut p = p;
5436 if i < p.1.len() {
5437 if p.2 && p.1[i] == token_id {
5438 p.3 = i + 1;
5439 } else {
5440 p.2 = false;
5441 }
5442 if i + 1 < p.1.len() {
5443 self.dspark_pending.push(p);
5444 continue;
5445 }
5446 }
5447 self.dspark_hist.push(p.3);
5448 self.dspark_real.push(token_id);
5449 }
5450 let Some(b) = &mut self.dsv4 else { return };
5451 let (g, layers, cfg) = (&b.0, &b.1, b.2);
5452 let n_layers = layers.len();
5453 if self.dspark.is_none() {
5454 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5455 if t.is_empty() {
5456 return;
5457 }
5458 eprintln!("DSpark: захват со слоёв {t:?}, блок {}", crate::dsv4::dspark_block());
5459 crate::dsv4::dspark_arm(&t, cfg.dim);
5460 self.dspark = Some(crate::dsv4::DsparkState::new(
5461 self.dsv4_mtp.len(),
5462 &cfg,
5463 t.len(),
5464 ));
5465 }
5466 let ds = self.dspark.as_mut().unwrap();
5467 if !crate::dsv4::dspark_take(&mut ds.main_hidden) {
5468 return; }
5470 let mut conf = Vec::new();
5471 crate::dsv4::pick_tally_arm();
5472 let draft_started = std::time::Instant::now();
5477 #[cfg(feature = "gpu")]
5478 let gpu_draft = crate::dsv4::dspark_gpu_on();
5479 #[cfg(not(feature = "gpu"))]
5480 let gpu_draft = false;
5481 let props = if gpu_draft {
5482 #[cfg(feature = "gpu")]
5483 {
5484 let kv_id = b.3.kv_id;
5485 match crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg) {
5486 Some(pk) => crate::dsv4::dspark_draft_gpu(
5487 g,
5488 &self.dsv4_mtp,
5489 &cfg,
5490 ds,
5491 pk,
5492 kv_id,
5493 token_id,
5494 position,
5495 self.pool.as_deref(),
5496 &mut conf,
5497 ),
5498 None => Vec::new(),
5499 }
5500 }
5501 #[cfg(not(feature = "gpu"))]
5502 Vec::new()
5503 } else {
5504 crate::gpu::cpu_scope(|| {
5505 crate::dsv4::dspark_draft(
5506 g,
5507 &self.dsv4_mtp,
5508 &cfg,
5509 ds,
5510 token_id,
5511 position,
5512 self.pool.as_deref(),
5513 &mut conf,
5514 )
5515 })
5516 };
5517 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5518 let draft_picks = crate::dsv4::pick_tally_take();
5519 crate::dsv4::dspark_freq_note(&draft_picks);
5520 crate::dsv4::pick_tally_arm();
5523 if !props.is_empty() {
5524 let (tu, tt) = {
5528 let flat: Vec<(usize, Vec<usize>)> = self
5529 .dspark_trunk_picks
5530 .iter()
5531 .flat_map(|v| v.iter().cloned())
5532 .collect();
5533 let mut per: std::collections::HashMap<usize, Vec<usize>> =
5535 std::collections::HashMap::new();
5536 for (li, picks) in flat {
5537 per.entry(li).or_default().extend(picks);
5538 }
5539 let n = per.len().max(1);
5540 let mut u = 0usize;
5541 let mut t = 0usize;
5542 for (_, v) in per {
5543 t += v.len();
5544 u += v.iter().collect::<std::collections::HashSet<_>>().len();
5545 }
5546 (u / n, t / n)
5547 };
5548 let (du, dt) = crate::dsv4::tally_unique(&draft_picks);
5549 self.dspark_exp.push((tu, tt, du, dt));
5550 self.dspark_pending.push((position, props, true, 0));
5551 }
5552 if self.dspark_hist.len() >= 8 && self.dspark_hist.len() % 8 == 0 {
5553 let n = self.dspark_hist.len() as f32;
5554 let mean: f32 = self.dspark_hist.iter().sum::<usize>() as f32 / n;
5555 let block = crate::dsv4::dspark_block();
5556 let mut at = vec![0usize; block + 1];
5557 for &k in &self.dspark_hist {
5558 at[k] += 1;
5559 }
5560 let mut surv = Vec::with_capacity(block);
5562 for i in 1..=block {
5563 let k = at[i..].iter().sum::<usize>() as f32 / n;
5564 surv.push(format!("{k:.2}"));
5565 }
5566 let distinct = self
5567 .dspark_real
5568 .iter()
5569 .collect::<std::collections::HashSet<_>>()
5570 .len();
5571 let (tu, tt, du, dt) = self.dspark_exp.iter().fold((0, 0, 0, 0), |a, b| {
5572 (a.0 + b.0, a.1 + b.1, a.2 + b.2, a.3 + b.3)
5573 });
5574 let m = self.dspark_exp.len().max(1);
5575 eprintln!(
5576 "DSpark: черновиков {}, принято в среднем {mean:.2} из {block} \
5577 (токенов за проход {:.2}), распределение {at:?}, выживание [{}]",
5578 self.dspark_hist.len(),
5579 mean + 1.0,
5580 surv.join(" ")
5581 );
5582 eprintln!(
5583 "DSpark: разных токенов {distinct} из {} (вырожденность), \
5584 эксперты ствол {}/{} на слой за {block} токенов, \
5585 черновик {}/{} за блок, draft {:.2} мс/блок",
5586 self.dspark_real.len(),
5587 tu / m,
5588 tt / m,
5589 du / m,
5590 dt / m,
5591 self.dspark_draft_ns as f64 / self.dspark_exp.len().max(1) as f64 / 1e6
5592 );
5593 }
5594 }
5595
5596 fn forward_layers_upto(
5597 &mut self,
5598 hidden: &[f32],
5599 position: usize,
5600 task_mask: Option<&TaskMask>,
5601 upto: Option<usize>,
5602 ) -> Vec<f32> {
5603 if let Some(plan) = self.gpu_plan.clone() {
5609 if upto.is_none() && plan.len() > 1 {
5610 let mut h = hidden.to_vec();
5611 for &(dev, from, upto_incl) in plan.iter() {
5612 h = crate::gpu::with_device(dev, || {
5613 self.forward_layers_span(&h, position, task_mask, from, Some(upto_incl))
5614 });
5615 }
5616 return h;
5617 }
5618 }
5619 self.forward_layers_span(hidden, position, task_mask, 0, upto)
5620 }
5621
5622 pub fn set_gpu_plan(&mut self, devices: Option<&[usize]>) -> Result<(), String> {
5627 self.set_gpu_plan_at(devices, None)
5628 }
5629
5630 pub fn set_gpu_plan_at(
5634 &mut self,
5635 devices: Option<&[usize]>,
5636 at: Option<usize>,
5637 ) -> Result<(), String> {
5638 let Some(devs) = devices.filter(|d| d.len() > 1) else {
5639 self.gpu_plan = None;
5640 return Ok(());
5641 };
5642 self.split_supported()?;
5643 let n = self.num_layers;
5644 if devs.len() > n {
5645 return Err(format!("{} devices for {n} layers", devs.len()));
5646 }
5647 if let Some(k) = at {
5648 if k == 0 || k >= n {
5649 return Err(format!("split at {k}: the model has {n} layers"));
5650 }
5651 if devs.len() == 2 {
5652 self.gpu_plan = Some(std::sync::Arc::new(vec![
5653 (devs[0], 0, k - 1),
5654 (devs[1], k, n - 1),
5655 ]));
5656 return Ok(());
5657 }
5658 return Err(format!(
5659 "an explicit split point takes exactly 2 devices, got {}",
5660 devs.len()
5661 ));
5662 }
5663 let per = n.div_ceil(devs.len());
5664 let mut plan = Vec::with_capacity(devs.len());
5665 let mut from = 0usize;
5666 for &d in devs {
5667 if from >= n {
5668 break;
5669 }
5670 let upto = (from + per - 1).min(n - 1);
5671 plan.push((d, from, upto));
5672 from = upto + 1;
5673 }
5674 self.gpu_plan = Some(std::sync::Arc::new(plan));
5675 Ok(())
5676 }
5677
5678 pub fn gpu_plan(&self) -> Option<Vec<(usize, usize, usize)>> {
5680 self.gpu_plan.as_ref().map(|p| p.as_ref().clone())
5681 }
5682
5683 fn forward_layers_span(
5689 &mut self,
5690 hidden: &[f32],
5691 position: usize,
5692 task_mask: Option<&TaskMask>,
5693 from: usize,
5694 upto: Option<usize>,
5695 ) -> Vec<f32> {
5696 debug_assert!(from == 0 || (self.dsv4.is_none() && self.g3n.is_none()));
5697 if let Some(b) = &mut self.dsv4 {
5703 let _ = (task_mask, upto);
5704 let token_id = hidden.first().copied().unwrap_or(0.0) as u32;
5705 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
5706 st.pos = position;
5707 let mut logits = Vec::new();
5708 crate::dsv4::forward_token(
5709 g,
5710 layers,
5711 &cfg,
5712 st,
5713 token_id,
5714 &self.inv_freq,
5715 self.pool.as_deref(),
5716 &mut logits,
5717 );
5718 self.graph_logits = Some(logits);
5719 self.dspark_probe(position, token_id);
5720 return vec![0.0; self.hidden_size];
5723 }
5724 if let Some(b) = &self.g3n {
5727 let _ = (task_mask, upto);
5728 return crate::g3n::g3n_forward(
5729 &b.0,
5730 &b.1,
5731 hidden,
5732 position,
5733 &mut self.kv_cache.layers,
5734 self.num_heads,
5735 self.num_kv_heads,
5736 self.head_dim,
5737 self.pool.as_deref(),
5738 );
5739 }
5740 let mut h = hidden.to_vec();
5741 let (nh, _nkv, _hd, hs, _rd, eps) = (
5744 self.num_heads,
5745 self.num_kv_heads,
5746 self.head_dim,
5747 self.hidden_size,
5748 self.rotary_dim,
5749 self.rms_eps,
5750 );
5751 let pool = self.pool.clone();
5752 let graph_env = std::env::var("CMF_GPU_WGPU_GRAPH").ok();
5764 let graph_on = match graph_env.as_deref() {
5765 Some("0") => false,
5766 Some("prefill") => false, Some(_) => true,
5768 None => crate::gpu::wgpu_graph_default(),
5774 };
5775 let graph_trusted =
5776 graph_env.is_some() || crate::gpu::wgpu_graph_default() || self.gdn_cfg.is_some();
5777 let race_eligible = graph_on
5778 && upto.is_none()
5779 && task_mask.is_none()
5780 && from == 0
5781 && !crate::gpu::graph_unsupported();
5782 let mut tail_start = 0usize;
5783 if race_eligible && crate::gpu::graph_race_use_graph(graph_trusted) {
5784 let t_graph = std::time::Instant::now();
5785 let mut lg = Vec::new();
5786 let mut gl = 0usize;
5787 let built = self.try_token_graph_wgpu(hidden, position, &mut lg, &mut gl);
5788 if built.is_none() && !self.o1_active() && self.attn_softcap == 0.0 {
5793 crate::gpu::graph_mark_unsupported();
5794 }
5795 graph_note(built.is_some());
5796 if let Some(hh) = built {
5797 let dur = t_graph.elapsed();
5798 if std::env::var("CMF_GRAPH_PROF").is_ok() {
5799 eprintln!("graph-call: {:.2} ms total", dur.as_secs_f64() * 1000.0);
5800 }
5801 if gl > 0 && gl < self.num_layers {
5802 h = hh;
5808 tail_start = gl;
5809 } else if graph_trusted || !crate::gpu::graph_race_first_token_hopeless(dur) {
5810 if !graph_trusted {
5811 crate::gpu::graph_race_record(true, dur);
5812 }
5813 if !lg.is_empty() {
5814 lg.resize(self.vocab_size, 0.0);
5817 if let Some(c) = self.final_softcap {
5818 for l in lg.iter_mut() {
5819 *l = c * (*l / c).tanh();
5820 }
5821 }
5822 self.graph_logits = Some(lg);
5823 }
5824 return hh;
5825 }
5826 }
5832 }
5833 let span = from > 0 || upto.is_some();
5857 if span && graph_on && task_mask.is_none() && graph_trusted {
5858 let upto_excl = upto.map_or(self.num_layers, |u| u + 1);
5859 let mut lg = Vec::new();
5860 let mut gl = 0usize;
5861 let span_res =
5862 self.try_token_graph_wgpu_span(hidden, position, &mut lg, from, upto_excl, &mut gl);
5863 graph_note(span_res.is_some() && gl == upto_excl - from);
5864 if std::env::var("CMF_GPU_DEBUG").is_ok() {
5865 static SEEN: std::sync::atomic::AtomicU32 =
5869 std::sync::atomic::AtomicU32::new(0);
5870 if SEEN.fetch_add(1, std::sync::atomic::Ordering::Relaxed) < 4 {
5871 eprintln!(
5872 "span graph: covered {gl} of {} layers [{from}..{upto_excl}) res={}",
5873 upto_excl - from,
5874 span_res.is_some()
5875 );
5876 }
5877 }
5878 if let Some(hh) = span_res {
5879 if gl == upto_excl - from {
5880 if !lg.is_empty() {
5881 lg.resize(self.vocab_size, 0.0);
5882 if let Some(c) = self.final_softcap {
5883 for l in lg.iter_mut() {
5884 *l = c * (*l / c).tanh();
5885 }
5886 }
5887 self.graph_logits = Some(lg);
5888 }
5889 crate::gpu::set_layer(-1);
5890 return hh;
5891 }
5892 h = hh;
5894 tail_start = from + gl;
5895 }
5896 }
5897 let t_race_cpu = (race_eligible && !graph_trusted).then(std::time::Instant::now);
5898
5899 #[cfg(target_os = "macos")]
5900 let mut gpu_skip_until = 0usize;
5901 for li in tail_start.max(from)..self.num_layers {
5902 crate::gpu::set_layer(li as i64); if let Some(u) = upto {
5904 if li > u {
5905 break;
5906 }
5907 }
5908 if let Some(mask) = task_mask {
5909 if !mask.layer_alive(li) {
5910 continue; }
5912 }
5913 #[cfg(target_os = "macos")]
5917 {
5918 if li < gpu_skip_until {
5919 continue;
5920 }
5921 if task_mask.is_none() {
5922 let end = self.q1_graph_gpu(li, upto, position, &mut h);
5923 if end > li {
5924 gpu_skip_until = end;
5925 if self.is_loop_end(end - 1) && end < self.num_layers {
5928 h = inference::rms_norm(
5929 &h,
5930 &self.weights.final_norm,
5931 self.rms_eps,
5932 self.norm_style,
5933 );
5934 }
5935 continue;
5936 }
5937 }
5938 }
5939
5940 let lw = &self.weights.layers[self.phys_layer(li)];
5941 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
5942 if tp.parse::<usize>().ok() == Some(position) {
5943 let n: f32 = h.iter().map(|x| x * x).sum::<f32>().sqrt();
5944 eprintln!(
5945 "TRACE pos {position} layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
5946 h[0], h[1]
5947 );
5948 }
5949 }
5950 inference::rms_norm_into(
5953 &h,
5954 &lw.input_norm,
5955 self.rms_eps,
5956 self.norm_style,
5957 &mut self.ws.n1,
5958 );
5959
5960 let attn_out = match &lw.attn {
5961 AttnKind::Mla(w) => {
5962 let inv_freq_l = self.layer_inv_freq(li);
5963 let rs = self.layer_rope_scale(li);
5964 let eps = self.rms_eps;
5965 let pool = self.pool.clone();
5966 mla_attention(
5967 w,
5968 &self.ws.n1,
5969 &mut self.kv_cache.layers[li],
5970 position,
5971 &inv_freq_l,
5972 rs,
5973 eps,
5974 pool.as_deref(),
5975 )
5976 }
5977 AttnKind::Linear(w) => {
5978 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
5979 vmf_phase_forward(
5980 &self.ws.n1,
5981 w,
5982 &cfg,
5983 &mut self.kv_cache.layers[li].linear_state,
5984 self.pool.as_deref(),
5985 )
5986 }
5987 AttnKind::Kda(w) => {
5988 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
5989 crate::linear_core::kda_forward(
5990 &self.ws.n1,
5991 w,
5992 &cfg,
5993 &mut self.kv_cache.layers[li].linear_state,
5994 self.pool.as_deref(),
5995 )
5996 }
5997 AttnKind::LinearGdn(w) => {
5998 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
5999 gdn_forward(
6000 &self.ws.n1,
6001 w,
6002 &cfg,
6003 &mut self.kv_cache.layers[li].linear_state,
6004 self.pool.as_deref(),
6005 )
6006 }
6007 AttnKind::ShortConv(w) => {
6008 let cfg = self
6009 .short_conv_cfg
6010 .expect("short-conv layer without short_conv_cfg");
6011 short_conv_forward(
6012 &self.ws.n1,
6013 w,
6014 &cfg,
6015 &mut self.kv_cache.layers[li].linear_state,
6016 self.pool.as_deref(),
6017 )
6018 }
6019 AttnKind::Full {
6020 wq,
6021 wk,
6022 wv,
6023 wo,
6024 q_norm,
6025 k_norm,
6026 output_gate,
6027 softplus_gate,
6028 bias,
6029 } if self.kv_cache.layers[li].o1_sealed() => {
6030 let inv_freq_l = self.layer_inv_freq(li);
6033 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6034 let cfg = QwenAttnCfg {
6035 num_heads: self.layer_num_heads(li),
6036 num_kv_heads: nkv_l,
6037 head_dim: hd_l,
6038 hidden_size: hs,
6039 position,
6040 inv_freq: &inv_freq_l,
6041 rotary_dim: rd_l,
6042 scale: self.attn_scale,
6043 softcap: self.attn_softcap,
6044 window: None,
6045 v_norm: self.attn_v_norm,
6046 q_norm: q_norm.as_deref(),
6047 k_norm: k_norm.as_deref(),
6048 output_gate: *output_gate,
6049 softplus_gate: softplus_gate
6050 .as_ref()
6051 .map(|(gate, per_head)| (gate, *per_head)),
6052 rope_scale: self.layer_rope_scale(li),
6053 bias: bias
6054 .as_ref()
6055 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6056 rms_eps: eps,
6057 norm_style: self.norm_style,
6058 pool: pool.as_deref(),
6059 };
6060 attention::qwen_attention_nystrom(
6061 &self.ws.n1,
6062 wq,
6063 wk,
6064 wv,
6065 wo,
6066 &mut self.kv_cache.layers[li],
6067 &cfg,
6068 )
6069 }
6070 AttnKind::Full {
6071 wq,
6072 wk,
6073 wv,
6074 wo,
6075 q_norm,
6076 k_norm,
6077 output_gate,
6078 softplus_gate,
6079 bias,
6080 } => 'attn: {
6081 if graph_on
6084 && !*output_gate
6085 && softplus_gate.is_none()
6086 && self.attention_heads_per_layer.is_none()
6087 && bias.is_none()
6088 && task_mask.is_none()
6089 {
6090 let inv_freq_l = self.layer_inv_freq(li);
6091 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6092 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
6093 if let (Some((gm, qi)), Some((_, ki)), Some((_, vi)), Some((_, oi))) = (
6094 wq.mapped_q1(),
6095 wk.mapped_q1(),
6096 wv.mapped_q1(),
6097 wo.mapped_q1(),
6098 ) {
6099 let gm = gm.clone();
6100 let mut out = vec![0f32; hs];
6101 let cache = &self.kv_cache.layers[li];
6102 if crate::gpu::attn_dropin(
6103 &gm,
6104 self.graph_kv_id,
6105 li,
6106 &self.ws.n1,
6107 qi,
6108 ki,
6109 vi,
6110 oi,
6111 q_norm.as_deref(),
6112 k_norm.as_deref(),
6113 &inv_freq_l,
6114 nh,
6115 nkv_l,
6116 hd_l,
6117 rd_l,
6118 hs,
6119 position,
6120 self.kv_cache.max_seq_len,
6121 gemma,
6122 eps as f32,
6123 cache.k_heads(),
6124 cache.v_heads(),
6125 &mut out,
6126 ) {
6127 break 'attn out;
6128 }
6129 }
6130 }
6131 let masked = task_mask
6132 .map(|m| m.head_flags(li, self.num_heads).iter().any(|&a| !a))
6133 .unwrap_or(false);
6134 let f32_view = (wq.as_f32(), wk.as_f32(), wv.as_f32(), wo.as_f32());
6135 match (masked, f32_view) {
6136 (true, (Some(q), Some(k), Some(v), Some(o))) => {
6139 let active_heads = task_mask.unwrap().head_flags(li, self.num_heads);
6140 attention::multi_head_attention(
6141 &self.ws.n1,
6142 q,
6143 k,
6144 v,
6145 o,
6146 &mut self.kv_cache.layers[li],
6147 self.num_heads,
6148 self.num_kv_heads,
6149 self.head_dim,
6150 self.hidden_size,
6151 position,
6152 &active_heads,
6153 &self.inv_freq,
6154 )
6155 }
6156 (masked, _) => {
6157 if masked {
6158 tracing::warn!(
6159 "layer {li}: head mask on quantized weights not \
6160 supported yet — executing dense"
6161 );
6162 }
6163 let inv_freq_l = self.layer_inv_freq(li);
6164 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6165 let cfg = QwenAttnCfg {
6166 num_heads: self.layer_num_heads(li),
6167 num_kv_heads: nkv_l,
6168 head_dim: hd_l,
6169 hidden_size: hs,
6170 position,
6171 inv_freq: &inv_freq_l,
6172 rotary_dim: rd_l,
6173 scale: self.attn_scale,
6174 softcap: self.attn_softcap,
6175 window: self.layer_window(li),
6176 v_norm: self.attn_v_norm,
6177 q_norm: q_norm.as_deref(),
6178 k_norm: k_norm.as_deref(),
6179 output_gate: *output_gate,
6180 softplus_gate: softplus_gate
6181 .as_ref()
6182 .map(|(gate, per_head)| (gate, *per_head)),
6183 rope_scale: self.layer_rope_scale(li),
6184 bias: bias
6185 .as_ref()
6186 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6187 rms_eps: eps,
6188 norm_style: self.norm_style,
6189 pool: pool.as_deref(),
6190 };
6191 attention::qwen_attention(
6192 &self.ws.n1,
6193 wq,
6194 wk,
6195 wv,
6196 wo,
6197 &mut self.kv_cache.layers[li],
6198 &cfg,
6199 )
6200 }
6201 }
6202 }
6203 };
6204 let attn_out = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
6207 Some(w) => inference::rms_norm(&attn_out, w, self.rms_eps, self.norm_style),
6208 None => attn_out,
6209 };
6210 let lw = &self.weights.layers[self.phys_layer(li)];
6211 inference::add_rmsnorm_fused_into(
6212 &mut h,
6213 &attn_out,
6214 &lw.post_norm,
6215 self.rms_eps,
6216 self.norm_style,
6217 &mut self.ws.p1,
6218 );
6219 let mut attn_out = attn_out;
6220 attention::recycle_buf(&mut attn_out);
6221 let post_normed = &self.ws.p1;
6222
6223 let ffn_masked = task_mask
6224 .map(|m| m.ffn_active_count(li) < self.intermediate_size)
6225 .unwrap_or(false);
6226 let ffn_out = match (ffn_masked, &lw.ffn) {
6238 (true, FfnKind::Dense(d)) => {
6239 let tm = task_mask.unwrap();
6240 let alive = tm.ffn_active_count(li);
6241 let deep = alive * 2 <= self.intermediate_size;
6242 if deep && d.down_proj.sparse_col_ok() {
6243 let active = tm.ffn_active_indices(li);
6244 sparse_ffn_quant(
6245 d,
6246 post_normed,
6247 &active,
6248 self.hidden_size,
6249 self.pool.as_deref(),
6250 )
6251 } else if deep
6252 && let (Some(g), Some(u), Some(dn)) = (
6253 d.gate_proj.as_f32(),
6254 d.up_proj.as_f32(),
6255 d.down_proj.as_f32(),
6256 )
6257 {
6258 let active = tm.ffn_active_indices(li);
6259 inference::sparse_ffn_forward(
6260 post_normed,
6261 g,
6262 u,
6263 dn,
6264 self.hidden_size,
6265 self.intermediate_size,
6266 &active,
6267 self.pool.as_deref(),
6268 )
6269 } else {
6270 let row = tm.ffn_masks.get(li).map(|v| v.as_slice());
6271 dense_ffn_batch(d, post_normed, 1, self.pool.as_deref(), row)
6272 }
6273 }
6274 (true, FfnKind::Moe(m)) => {
6275 let allowed = task_mask.and_then(|tm| tm.expert_flags(li, m.experts.len()));
6279 ffn_forward(
6280 &lw.ffn,
6281 post_normed,
6282 self.pool.as_deref(),
6283 allowed.as_deref(),
6284 )
6285 }
6286 (true, FfnKind::DenseMoe(dm)) => dense_moe_ffn(
6287 dm,
6288 post_normed,
6289 &h,
6290 self.rms_eps,
6291 self.norm_style,
6292 self.pool.as_deref(),
6293 ),
6294 (false, _) => match &lw.ffn {
6295 FfnKind::DenseMoe(dm) => dense_moe_ffn(
6296 dm,
6297 post_normed,
6298 &h,
6299 self.rms_eps,
6300 self.norm_style,
6301 self.pool.as_deref(),
6302 ),
6303 _ => {
6304 let allowed = match (&lw.ffn, task_mask) {
6305 (FfnKind::Moe(m), Some(tm)) => tm.expert_flags(li, m.experts.len()),
6306 _ => None,
6307 };
6308 ffn_forward(
6309 &lw.ffn,
6310 post_normed,
6311 self.pool.as_deref(),
6312 allowed.as_deref(),
6313 )
6314 }
6315 },
6316 };
6317 let ffn_out = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
6318 Some(w) => inference::rms_norm(&ffn_out, w, self.rms_eps, self.norm_style),
6319 None => ffn_out,
6320 };
6321 for (i, &f) in ffn_out.iter().enumerate() {
6322 h[i] += f;
6323 }
6324 let mut ffn_out = ffn_out;
6325 attention::recycle_buf(&mut ffn_out);
6326
6327 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
6329 for v in h.iter_mut() {
6330 *v *= sc;
6331 }
6332 }
6333
6334 if self.is_loop_end(li) && li + 1 < self.num_layers {
6337 h = inference::rms_norm(
6338 &h,
6339 &self.weights.final_norm,
6340 self.rms_eps,
6341 self.norm_style,
6342 );
6343 }
6344
6345 if self.dyn_phi_layer == Some(li) {
6349 self.update_dyn_phi(&h);
6350 }
6351 }
6352 crate::gpu::set_layer(-1); if let Some(t) = t_race_cpu {
6354 crate::gpu::graph_race_record(false, t.elapsed());
6355 }
6356
6357 h
6358 }
6359
6360 fn update_dyn_phi(&mut self, h: &[f32]) {
6363 const A: f32 = 0.2;
6364 if self.dyn_phi_ema.len() != h.len() {
6365 self.dyn_phi_ema = vec![0.0; h.len()];
6366 self.dyn_phi_seen = 0;
6367 }
6368 if self.dyn_phi_seen == 0 {
6369 self.dyn_phi_ema.copy_from_slice(h);
6370 } else {
6371 for (e, &v) in self.dyn_phi_ema.iter_mut().zip(h) {
6372 *e = (1.0 - A) * *e + A * v;
6373 }
6374 }
6375 self.dyn_phi_seen += 1;
6376 }
6377
6378 pub fn dyn_phi(&self) -> &[f32] {
6380 &self.dyn_phi_ema
6381 }
6382
6383 pub fn set_dyn_phi_layer(&mut self, layer: Option<usize>) {
6385 self.dyn_phi_layer = layer;
6386 self.dyn_phi_ema.clear();
6387 self.dyn_phi_seen = 0;
6388 }
6389
6390 pub fn dynamic_skills(&self) -> Vec<(usize, String, usize)> {
6392 let Some(model) = &self.model else {
6393 return Vec::new();
6394 };
6395 model
6396 .header
6397 .skills
6398 .iter()
6399 .enumerate()
6400 .filter_map(|(i, sk)| {
6401 let ok = matches!(self.dyn_skill_layers.get(i), Some(Some(_)));
6402 let sel = sk.selection.as_ref()?;
6403 (ok).then(|| (i, sk.id.clone(), sel.phi_layer))
6404 })
6405 .collect()
6406 }
6407
6408 pub fn active_skill(&self) -> Option<usize> {
6410 self.dyn_active
6411 }
6412
6413 pub fn enable_dynamic_routing(&mut self) -> usize {
6418 use crate::swarm::{DynRouter, RoutableSkill};
6419 let Some(model) = self.model.clone() else {
6420 return 0;
6421 };
6422 if self.dyn_blend_loaded {
6425 tracing::warn!("dynamic routing unavailable on a blend-loaded pipeline");
6426 return 0;
6427 }
6428 if let Some(a) = self.dyn_active {
6432 if !matches!(self.dyn_skill_layers.get(a), Some(Some(_))) {
6433 tracing::warn!("loaded skill is not FFN-eligible — dynamic routing unavailable");
6434 return 0;
6435 }
6436 }
6437 let hidden = self.hidden_size;
6438 let mut skills = Vec::new();
6439 for (idx, id, _phi) in self.dynamic_skills() {
6440 if let Some(sel) = model.header.skills[idx].selection.as_ref() {
6441 if let Some(rs) = RoutableSkill::from_descriptor(idx, id, sel, hidden) {
6442 skills.push(rs);
6443 }
6444 }
6445 }
6446 if skills.is_empty() {
6447 return 0;
6448 }
6449 let phi = skills[0].phi_layer;
6451 if skills.iter().any(|s| s.phi_layer != phi) {
6452 tracing::warn!("routable skills disagree on phi_layer; using {phi}");
6453 }
6454 let n = skills.len();
6455 self.set_dyn_phi_layer(Some(phi));
6456 self.dyn_router = Some(DynRouter::new(skills));
6457 n
6458 }
6459
6460 pub fn route_switches(&self) -> Vec<(usize, Option<String>, Option<String>)> {
6462 self.dyn_router
6463 .as_ref()
6464 .map(|r| r.switches.clone())
6465 .unwrap_or_default()
6466 }
6467
6468 fn lm_head_forward(&self, hidden: &[f32]) -> Vec<f32> {
6471 let rows = self.weights.lm_head.rows();
6472 let mut logits = attention::take_buf(rows.min(self.vocab_size));
6473 self.weights
6474 .lm_head
6475 .matvec(hidden, &mut logits, self.pool.as_deref());
6476 logits.resize(self.vocab_size, 0.0);
6477 if let Some(m) = self.logit_multiplier {
6478 for l in logits.iter_mut() {
6479 *l *= m;
6480 }
6481 }
6482 if let Some(c) = self.final_softcap {
6483 for l in logits.iter_mut() {
6484 *l = c * (*l / c).tanh();
6485 }
6486 }
6487 logits
6488 }
6489
6490 pub fn prefill_next_logits(&mut self, ids: &[u32], task_mask: Option<&TaskMask>) -> Vec<f32> {
6495 self.kv_cache.clear();
6496 self.kv_history.clear();
6497 let mut hidden = vec![0.0f32; self.hidden_size];
6498 for (pos, &id) in ids.iter().enumerate() {
6499 let emb = self.embed_single(id);
6500 hidden = self.forward_layers(&emb, pos, task_mask);
6501 }
6502 inference::rms_norm_into(
6503 &hidden,
6504 &self.weights.final_norm,
6505 self.rms_eps,
6506 self.norm_style,
6507 &mut self.ws.n1,
6508 );
6509 self.lm_head_forward(&self.ws.n1)
6510 }
6511}
6512
6513pub fn create_test_pipeline(
6515 hidden_size: usize,
6516 intermediate_size: usize,
6517 num_heads: usize,
6518 num_kv_heads: usize,
6519 head_dim: usize,
6520 num_layers: usize,
6521 vocab_size: usize,
6522) -> Pipeline {
6523 let synth = |n: usize, salt: usize| -> Vec<f32> {
6526 (0..n)
6527 .map(|i| (((i * 31 + salt * 17 + 7) % 97) as f32 / 97.0 - 0.5) * 0.2)
6528 .collect()
6529 };
6530 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
6531 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
6532 };
6533 let layer_weights: Vec<LayerWeights> = (0..num_layers)
6534 .map(|li| LayerWeights {
6535 input_norm: vec![1.0; hidden_size],
6536 post_norm: vec![1.0; hidden_size],
6537 attn_out_norm: None,
6538 ffn_out_norm: None,
6539 layer_scale: None,
6540 ffn: FfnKind::Dense(DenseFfn {
6541 gate_proj: qt(intermediate_size, hidden_size, li * 10 + 5),
6542 up_proj: qt(intermediate_size, hidden_size, li * 10 + 6),
6543 down_proj: qt(hidden_size, intermediate_size, li * 10 + 7),
6544 act: Act::Silu,
6545 }),
6546 attn: AttnKind::Full {
6547 bias: None,
6548 wq: qt(num_heads * head_dim, hidden_size, li * 10 + 1),
6549 wk: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 2),
6550 wv: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 3),
6551 wo: qt(hidden_size, num_heads * head_dim, li * 10 + 4),
6552 q_norm: None,
6553 k_norm: None,
6554 output_gate: false,
6555 softplus_gate: None,
6556 },
6557 })
6558 .collect();
6559
6560 Pipeline::new(
6561 Tokenizer::byte_level(),
6562 PipelineWeights {
6563 embed_tokens: qt(vocab_size, hidden_size, 100),
6564 layers: layer_weights,
6565 lm_head: qt(vocab_size, hidden_size, 200),
6566 final_norm: vec![1.0; hidden_size],
6567 },
6568 hidden_size,
6569 intermediate_size,
6570 num_heads,
6571 num_kv_heads,
6572 head_dim,
6573 num_layers,
6574 num_layers, false, vocab_size,
6577 1e-6,
6578 10_000.0,
6579 NormStyle::Qwen,
6580 4096,
6581 SamplerConfig {
6582 seed: Some(42),
6583 ..Default::default()
6584 },
6585 )
6586}
6587
6588#[inline]
6593fn mask_bit(row: &[u8], j: usize) -> bool {
6594 (row.get(j >> 3).copied().unwrap_or(0) >> (j & 7)) & 1 != 0
6595}
6596
6597fn zero_masked_cols(g: &mut [f32], rows: usize, inter: usize, row: &[u8]) {
6603 for r in 0..rows {
6604 let base = r * inter;
6605 for (bi, &byte) in row.iter().enumerate() {
6606 if byte == 0xFF {
6607 continue;
6608 }
6609 let j0 = bi * 8;
6610 for bit in 0..8 {
6611 let j = j0 + bit;
6612 if j < inter && byte & (1 << bit) == 0 {
6613 g[base + j] = 0.0;
6614 }
6615 }
6616 }
6617 }
6618}
6619
6620fn dense_ffn_batch(
6621 d: &DenseFfn,
6622 xs: &[f32],
6623 b: usize,
6624 pool: Option<&Pool>,
6625 mask_row: Option<&[u8]>,
6626) -> Vec<f32> {
6627 let inter = d.gate_proj.rows();
6628 let hidden = d.down_proj.rows();
6629 if mask_row.is_none()
6637 && d.act == Act::Silu
6638 && b >= 32
6639 && crate::gpu::enabled_here()
6640 && !crate::gpu::mm_killed()
6641 {
6642 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6643 d.gate_proj.mapped_q4t(),
6644 d.up_proj.mapped_q4t(),
6645 d.down_proj.mapped_q4t(),
6646 ) {
6647 let mut out = vec![0.0f32; b * hidden];
6648 if crate::gpu::q4t_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6649 return out;
6650 }
6651 }
6652 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6657 d.gate_proj.mapped_q4tp(),
6658 d.up_proj.mapped_q4tp(),
6659 d.down_proj.mapped_q4tp(),
6660 ) {
6661 let mut out = vec![0.0f32; b * hidden];
6662 if crate::gpu::q4tp_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6663 return out;
6664 }
6665 }
6666 }
6667 let mut g = vec![0.0f32; b * inter];
6668 d.gate_proj.matmat(xs, b, &mut g, pool);
6669 let mut u = vec![0.0f32; b * inter];
6670 d.up_proj.matmat(xs, b, &mut u, pool);
6671 for i in 0..b * inter {
6672 g[i] = d.act.combine(g[i], u[i]);
6673 }
6674 if let Some(row) = mask_row {
6675 zero_masked_cols(&mut g, b, inter, row);
6676 }
6677 let mut out = vec![0.0f32; b * hidden];
6678 d.down_proj.matmat(&g, b, &mut out, pool);
6679 out
6680}
6681
6682fn accumulate_act(m: &MoeFfn, xs: &[f32], b: usize) {
6687 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6688 static DUMP: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6689 let on = *ON.get_or_init(|| std::env::var("CMF_RMS_TRACE").is_ok());
6690 let dump = *DUMP.get_or_init(|| std::env::var("CMF_ACT_DUMP").is_ok());
6691 if (!on && !dump) || b == 0 {
6692 return;
6693 }
6694 let hidden = xs.len() / b;
6695 if on {
6696 let mut acc = m.act_sq.borrow_mut();
6697 if acc.len() < hidden {
6698 acc.resize(hidden, 0.0);
6699 }
6700 for t in 0..b {
6701 let row = &xs[t * hidden..(t + 1) * hidden];
6702 for (a, &v) in acc.iter_mut().zip(row) {
6703 *a += (v as f64) * (v as f64);
6704 }
6705 }
6706 }
6707 if dump {
6708 let cap: usize = std::env::var("CMF_ACT_DUMP_ROWS")
6711 .ok()
6712 .and_then(|v| v.parse().ok())
6713 .unwrap_or(4096);
6714 let mut rows = m.act_rows.borrow_mut();
6715 if rows.len() < cap * hidden {
6716 let take = b.min((cap * hidden - rows.len()) / hidden.max(1));
6717 rows.extend_from_slice(&xs[..take * hidden]);
6718 }
6719 }
6720}
6721
6722#[derive(Clone, Copy)]
6725struct SendVecs(*mut Vec<f32>);
6726unsafe impl Send for SendVecs {}
6727unsafe impl Sync for SendVecs {}
6728impl SendVecs {
6729 #[inline]
6730 fn at(self, i: usize) -> *mut Vec<f32> {
6731 unsafe { self.0.add(i) }
6732 }
6733}
6734
6735fn moe_ffn_batch(
6736 m: &MoeFfn,
6737 xs: &[f32],
6738 b: usize,
6739 hidden: usize,
6740 pool: Option<&Pool>,
6741 allowed: Option<&[bool]>,
6742) -> Vec<f32> {
6743 accumulate_act(m, xs, b);
6744 let ne = m.experts.len();
6745 let mut logits = vec![0.0f32; b * ne];
6746 m.router.matmat(xs, b, &mut logits, pool);
6747
6748 let mut assign: Vec<Vec<(usize, f32)>> = vec![Vec::new(); ne];
6751 {
6752 let mut st = m.stats.borrow_mut();
6753 if st.len() < ne {
6754 st.resize(ne, 0);
6755 }
6756 for bi in 0..b {
6757 let (idx, p, wsum) = moe_route(&logits[bi * ne..(bi + 1) * ne], m, allowed);
6758 for &e in &idx {
6759 st[e] += 1;
6760 assign[e].push((bi, p[e] / wsum));
6761 }
6762 }
6763 }
6764
6765 let mut out = vec![0.0f32; b * hidden];
6766 let cols = m.experts[0].gate_proj.cols();
6767 let run_expert = |d: &DenseFfn, list: &[(usize, f32)], out: &mut [f32]| {
6768 let sb = list.len();
6769 let mut sub = vec![0.0f32; sb * cols];
6770 for (k, &(bi, _)) in list.iter().enumerate() {
6771 sub[k * cols..(k + 1) * cols].copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6772 }
6773 let eo = dense_ffn_batch(d, &sub, sb, pool, None);
6774 for (k, &(bi, w)) in list.iter().enumerate() {
6775 for i in 0..hidden {
6776 out[bi * hidden + i] += w * eo[k * hidden + i];
6777 }
6778 }
6779 };
6780 let active: Vec<usize> = (0..ne).filter(|&e| !assign[e].is_empty()).collect();
6786 if pool.is_some() && active.len() >= 8 {
6787 let mut panels: Vec<Vec<f32>> = vec![Vec::new(); active.len()];
6788 {
6789 let panel_ptr = SendVecs(panels.as_mut_ptr());
6790 let experts = &m.experts;
6793 let (active_r, assign_r) = (&active, &assign);
6794 let run = |start: usize, end: usize| {
6795 for ai in start..end {
6796 let e = active_r[ai];
6797 let list = &assign_r[e];
6798 let sb = list.len();
6799 let mut sub = vec![0.0f32; sb * cols];
6800 for (k, &(bi, _)) in list.iter().enumerate() {
6801 sub[k * cols..(k + 1) * cols]
6802 .copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6803 }
6804 unsafe {
6806 *panel_ptr.at(ai) =
6807 dense_ffn_batch(&experts[e], &sub, sb, None, None);
6808 }
6809 }
6810 };
6811 match pool {
6812 Some(p) => p.run_rows(active.len(), &run),
6813 None => run(0, active.len()),
6814 }
6815 }
6816 for (ai, &e) in active.iter().enumerate() {
6817 for (k, &(bi, w)) in assign[e].iter().enumerate() {
6818 let eo = &panels[ai][k * hidden..(k + 1) * hidden];
6819 for i in 0..hidden {
6820 out[bi * hidden + i] += w * eo[i];
6821 }
6822 }
6823 }
6824 } else {
6825 for &e in &active {
6826 run_expert(&m.experts[e], &assign[e], &mut out);
6827 }
6828 }
6829 if let Some((se, gate)) = &m.shared {
6830 let all: Vec<(usize, f32)> = if let Some(gate) = gate {
6831 let mut gl = vec![0.0f32; b];
6832 gate.matmat(xs, b, &mut gl, pool);
6833 (0..b)
6834 .map(|bi| (bi, 1.0 / (1.0 + (-gl[bi]).exp())))
6835 .collect()
6836 } else {
6837 (0..b).map(|bi| (bi, 1.0)).collect()
6838 };
6839 run_expert(se, &all, &mut out);
6840 }
6841 out
6842}
6843
6844thread_local! {
6845 static FFN_SCRATCH: std::cell::RefCell<[Vec<f32>; 4]> =
6849 const { std::cell::RefCell::new([Vec::new(), Vec::new(), Vec::new(), Vec::new()]) };
6850}
6851
6852fn dense_ffn(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6854 if crate::gpu::enabled_here()
6865 && (d.gate_proj.rows() >= crate::gpu::min_rows() || d.gate_proj.is_q1())
6866 {
6867 let arm = if d.gate_proj.is_q1() && crate::gpu::q1_force() {
6868 crate::gpu::ProbeArm::Gpu
6869 } else {
6870 crate::gpu::probe_arm(crate::gpu::OpClass::Ffn)
6871 };
6872 match arm {
6873 crate::gpu::ProbeArm::Gpu => {
6874 let t0 = std::time::Instant::now();
6875 if let Some(out) = dense_ffn_gpu(d, x, pool) {
6876 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
6877 return out;
6878 }
6879 }
6880 crate::gpu::ProbeArm::CpuTimed => {
6881 let t0 = std::time::Instant::now();
6882 let out = crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6883 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
6884 return out;
6885 }
6886 crate::gpu::ProbeArm::Cpu => {
6887 return crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6888 }
6889 }
6890 }
6891 dense_ffn_cpu(d, x, pool)
6892}
6893
6894fn dense_ffn_cpu(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6896 let inter = d.gate_proj.rows();
6897 FFN_SCRATCH.with(|s| {
6898 let mut s = s.borrow_mut();
6899 let [g, u, ..] = &mut *s;
6900 g.resize(inter, 0.0);
6901 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
6904 } else {
6906 u.resize(inter, 0.0);
6907 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
6909 for i in 0..inter {
6910 g[i] = d.act.combine(g[i], u[i]);
6911 }
6912 }
6913 FFN_PROBE.with(|pr| {
6916 if let Some(acc) = pr.borrow_mut().as_mut() {
6917 let li = crate::gpu::cur_layer();
6918 if li >= 0 {
6919 if let Some(row) = acc.get_mut(li as usize) {
6920 for (a, &v) in row.iter_mut().zip(g.iter()) {
6921 *a += (v as f64).abs();
6922 }
6923 }
6924 }
6925 }
6926 });
6927 let mut out = attention::take_buf(d.down_proj.rows());
6928 d.down_proj.matvec(g, &mut out, pool);
6929 out
6930 })
6931}
6932
6933thread_local! {
6934 static FFN_PROBE: std::cell::RefCell<Option<Vec<Vec<f64>>>> =
6937 const { std::cell::RefCell::new(None) };
6938}
6939
6940fn dense_ffn_masked(
6945 d: &DenseFfn,
6946 x: &[f32],
6947 pool: Option<&Pool>,
6948 mask_row: &[u8],
6949) -> Vec<f32> {
6950 let inter = d.gate_proj.rows();
6951 FFN_SCRATCH.with(|s| {
6952 let mut s = s.borrow_mut();
6953 let [g, u, ..] = &mut *s;
6954 g.resize(inter, 0.0);
6955 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
6956 } else {
6958 u.resize(inter, 0.0);
6959 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
6960 for i in 0..inter {
6961 g[i] = d.act.combine(g[i], u[i]);
6962 }
6963 }
6964 zero_masked_cols(g, 1, inter, mask_row);
6965 let mut out = attention::take_buf(d.down_proj.rows());
6966 d.down_proj.matvec(g, &mut out, pool);
6967 out
6968 })
6969}
6970
6971fn dense_ffn_gpu(d: &DenseFfn, x: &[f32], _pool: Option<&Pool>) -> Option<Vec<f32>> {
6977 if d.act != Act::Silu {
6979 return None;
6980 }
6981 if d.gate_proj.rows() < crate::gpu::min_rows() && !d.gate_proj.is_q1() {
6984 return None;
6985 }
6986 let mut jobs: Vec<crate::gpu::MoeJob> = Vec::with_capacity(1);
6987 let mut model_ref = None;
6988 moe_push_job(d, x, 1.0, &mut jobs, &mut model_ref)?;
6989 let model = model_ref?;
6990 let hidden = jobs[0].down.1;
6991 let mut out = attention::take_buf(hidden);
6992 if crate::gpu::moe_block(&model, &jobs, &mut out) {
6993 Some(out)
6994 } else {
6995 let mut out = out;
6996 attention::recycle_buf(&mut out);
6997 None
6998 }
6999}
7000
7001#[allow(clippy::type_complexity)]
7006#[allow(clippy::type_complexity)]
7007pub(crate) fn moe_parts(
7008 t: &QTensor,
7009) -> Option<(
7010 &std::sync::Arc<cortiq_core::CmfModel>,
7011 usize,
7012 usize,
7013 usize,
7014 &[f32],
7015 &[f32],
7016 bool,
7017 bool,
7018 bool,
7019)> {
7020 match t {
7021 QTensor::Mapped {
7022 model,
7023 idx,
7024 dtype: dt @ (cortiq_core::TensorDtype::Q8_2f | cortiq_core::TensorDtype::Q8Row),
7025 rows,
7026 cols,
7027 row_scale,
7028 col_field,
7029 ..
7030 } if (*dt == cortiq_core::TensorDtype::Q8Row) || !col_field.is_empty() => Some((
7031 model, *idx, *rows, *cols, row_scale, col_field, false, false, false,
7032 )),
7033 QTensor::Mapped {
7035 model,
7036 idx,
7037 dtype: cortiq_core::TensorDtype::Q1,
7038 rows,
7039 cols,
7040 ..
7041 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], true, false, false)),
7042 QTensor::Mapped {
7044 model,
7045 idx,
7046 dtype: cortiq_core::TensorDtype::Q4Tiled,
7047 rows,
7048 cols,
7049 ..
7050 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7051 QTensor::Mapped {
7053 model,
7054 idx,
7055 dtype: cortiq_core::TensorDtype::Q4TiledP,
7056 rows,
7057 cols,
7058 ..
7059 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7060 QTensor::Mapped {
7064 model,
7065 idx,
7066 dtype: cortiq_core::TensorDtype::Q2TiledP,
7067 rows,
7068 cols,
7069 ..
7070 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, true)),
7071 _ => None,
7072 }
7073}
7074
7075#[cfg(target_os = "macos")]
7081fn metal_moe_graph_parts(m: &MoeFfn, hidden: usize) -> Option<crate::gpu::GpuMoe<'_>> {
7082 if m.router_sigmoid
7083 || m.router_input_norm
7084 || m.expert_bias.is_some()
7085 || m.route_tau.is_some()
7086 || m.mask.is_some()
7087 || m.per_expert_scale.is_some()
7088 || m.experts.is_empty()
7089 || m.top_k == 0
7090 {
7091 return None;
7092 }
7093 let (sh, sg) = match &m.shared {
7096 Some((sh, Some(sg))) => (sh, sg),
7097 _ => return None,
7098 };
7099 let (rf, rr, rc) = m.router.f32_parts()?;
7100 if rr != m.experts.len() || rc != hidden {
7101 return None;
7102 }
7103 let (sf, sr, sc) = sg.f32_parts()?;
7104 if sr * sc != hidden {
7105 return None;
7106 }
7107 let inter = m.experts[0].gate_proj.rows();
7108 let gu_q2 = m.experts[0].gate_proj.mapped_q2tp().is_some();
7111 let trio = |e: &DenseFfn| -> Option<(usize, usize, usize)> {
7112 if e.act != Act::Silu
7113 || e.gate_proj.rows() != inter
7114 || e.gate_proj.cols() != hidden
7115 || e.up_proj.rows() != inter
7116 || e.up_proj.cols() != hidden
7117 || e.down_proj.rows() != hidden
7118 || e.down_proj.cols() != inter
7119 {
7120 return None;
7121 }
7122 let pick = |t: &QTensor| -> Option<usize> {
7123 if gu_q2 {
7124 t.mapped_q2tp().map(|(_, i)| i)
7125 } else {
7126 t.mapped_q4tp().map(|(_, i)| i)
7127 }
7128 };
7129 Some((
7130 pick(&e.gate_proj)?,
7131 pick(&e.up_proj)?,
7132 e.down_proj.mapped_q4tp().map(|(_, i)| i)?,
7133 ))
7134 };
7135 let experts = m
7136 .experts
7137 .iter()
7138 .map(trio)
7139 .collect::<Option<Vec<_>>>()?;
7140 let shared = trio(sh)?;
7141 Some(crate::gpu::GpuMoe {
7142 router: rf,
7143 sgate: sf,
7144 experts,
7145 shared,
7146 n_exp: m.experts.len(),
7147 top_k: m.top_k,
7148 inter,
7149 norm_topk: m.norm_topk_prob,
7150 route_scale: m.routed_scaling,
7151 gu_q2,
7152 })
7153}
7154
7155pub(crate) fn moe_push_job_parts<'a>(
7159 gate: &'a QTensor,
7160 up: &'a QTensor,
7161 down: &'a QTensor,
7162 x: &[f32],
7163 w: f32,
7164 swiglu_limit: f32,
7165 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7166 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7167) -> Option<()> {
7168 use crate::qtensor::prescale;
7169 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(gate)?;
7170 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(up)?;
7171 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(down)?;
7172 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7173 return None; }
7175 if gq2 && (dq2 || !dq4 || down.mapped_q4tp().is_none()) {
7178 return None;
7179 }
7180 if !gq2 && dq2 {
7181 return None;
7182 }
7183 model_ref.get_or_insert_with(|| gm.clone());
7184 let dt = |cf: &[f32]| {
7185 if cf.is_empty() {
7186 cortiq_core::TensorDtype::Q8Row
7187 } else {
7188 cortiq_core::TensorDtype::Q8_2f
7189 }
7190 };
7191 jobs.push(crate::gpu::MoeJob {
7192 gate: (gi, gr, gc, grs),
7193 up: (ui, ur, uc, urs),
7194 down: (di, dr, dc, drs),
7195 xs_gate: prescale(x, gcf, dt(gcf)).into_owned(),
7196 xs_up: prescale(x, ucf, dt(ucf)).into_owned(),
7197 down_col: dcf,
7198 w,
7199 q1: gq1,
7200 q4t: gq4 && !gq2 && gate.mapped_q4tp().is_none(),
7201 q4tp: gq4 && (gq2 || gate.mapped_q4tp().is_some()),
7202 gu_q2: gq2,
7203 swiglu_limit,
7204 });
7205 Some(())
7206}
7207
7208fn moe_push_job<'a>(
7210 d: &'a DenseFfn,
7211 x: &[f32],
7212 w: f32,
7213 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7214 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7215) -> Option<()> {
7216 use crate::qtensor::prescale;
7217 if d.act != Act::Silu {
7218 return None; }
7220 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(&d.gate_proj)?;
7221 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(&d.up_proj)?;
7222 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(&d.down_proj)?;
7223 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7224 return None; }
7226 if gq2 && (dq2 || !dq4 || d.down_proj.mapped_q4tp().is_none()) {
7227 return None;
7228 }
7229 if !gq2 && dq2 {
7230 return None;
7231 }
7232 model_ref.get_or_insert_with(|| gm.clone());
7233 let gdt = if gcf.is_empty() {
7234 cortiq_core::TensorDtype::Q8Row
7235 } else {
7236 cortiq_core::TensorDtype::Q8_2f
7237 };
7238 let udt = if ucf.is_empty() {
7239 cortiq_core::TensorDtype::Q8Row
7240 } else {
7241 cortiq_core::TensorDtype::Q8_2f
7242 };
7243 jobs.push(crate::gpu::MoeJob {
7244 gate: (gi, gr, gc, grs),
7245 up: (ui, ur, uc, urs),
7246 down: (di, dr, dc, drs),
7247 xs_gate: prescale(x, gcf, gdt).into_owned(),
7248 xs_up: prescale(x, ucf, udt).into_owned(),
7249 down_col: dcf,
7250 w,
7251 q1: gq1,
7252 q4t: gq4 && !gq2 && d.gate_proj.mapped_q4tp().is_none(),
7253 q4tp: gq4 && (gq2 || d.gate_proj.mapped_q4tp().is_some()),
7254 gu_q2: gq2,
7255 swiglu_limit: 0.0,
7256 });
7257 Some(())
7258}
7259
7260fn sparse_ffn_quant(
7267 d: &DenseFfn,
7268 x: &[f32],
7269 active: &[u16],
7270 hidden: usize,
7271 pool: Option<&Pool>,
7272) -> Vec<f32> {
7273 let n = active.len();
7274 let inter = d.gate_proj.rows();
7275 let mut act = vec![0.0f32; n];
7276 let need_scratch = !(d.gate_proj.sparse_col_ok() && d.up_proj.sparse_col_ok());
7279 let compute = |ai: usize| -> f32 {
7280 let idx = active[ai] as usize;
7281 if idx >= inter {
7282 return 0.0; }
7284 let mut s = if need_scratch {
7285 vec![0.0f32; hidden]
7286 } else {
7287 Vec::new()
7288 };
7289 let gate = d.gate_proj.row_dot(idx, x, &mut s);
7290 let up = d.up_proj.row_dot(idx, x, &mut s);
7291 d.act.combine(gate, up)
7292 };
7293 match pool {
7294 Some(p) if n >= 256 => {
7295 let ptr = SendMut(act.as_mut_ptr());
7296 p.run(&|widx, nw| {
7297 let chunk = n.div_ceil(nw);
7298 let (s, e) = (widx * chunk, ((widx + 1) * chunk).min(n));
7299 for ai in s..e {
7300 unsafe { *ptr.at(ai) = compute(ai) };
7301 }
7302 });
7303 }
7304 _ => {
7305 for (ai, a) in act.iter_mut().enumerate() {
7306 *a = compute(ai);
7307 }
7308 }
7309 }
7310 let mut out = vec![0.0f32; hidden];
7312 for (ai, &idx) in active.iter().enumerate() {
7313 let w = act[ai];
7314 if w.abs() >= 1e-12 && (idx as usize) < inter {
7315 d.down_proj.add_col_scaled(idx as usize, w, &mut out);
7316 }
7317 }
7318 out
7319}
7320
7321#[doc(hidden)]
7323pub fn sparse_ffn_quant_for_test(
7324 d: &DenseFfn,
7325 x: &[f32],
7326 active: &[u16],
7327 hidden: usize,
7328) -> Vec<f32> {
7329 sparse_ffn_quant(d, x, active, hidden, None)
7330}
7331
7332fn dequant_dense_f32(d: &DenseFfn) -> (Vec<f32>, Vec<f32>, Vec<f32>) {
7336 let deq = |t: &QTensor| -> Vec<f32> {
7337 let (rows, cols) = (t.rows(), t.cols());
7338 let mut out = vec![0.0f32; rows * cols];
7339 for r in 0..rows {
7340 t.row_f32(r, &mut out[r * cols..(r + 1) * cols]);
7341 }
7342 out
7343 };
7344 (deq(&d.gate_proj), deq(&d.up_proj), deq(&d.down_proj))
7345}
7346
7347struct SendMut(*mut f32);
7349unsafe impl Send for SendMut {}
7350unsafe impl Sync for SendMut {}
7351impl SendMut {
7352 #[inline]
7353 #[allow(clippy::mut_from_ref)]
7356 unsafe fn at(&self, i: usize) -> &mut f32 {
7357 unsafe { &mut *self.0.add(i) }
7358 }
7359}
7360
7361fn moe_route(logits: &[f32], m: &MoeFfn, allowed: Option<&[bool]>) -> (Vec<usize>, Vec<f32>, f32) {
7371 let ne = logits.len();
7372 let p: Vec<f32> = if m.router_sigmoid {
7373 logits.iter().map(|&l| 1.0 / (1.0 + (-l).exp())).collect()
7374 } else {
7375 let mx = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
7376 let mut e: Vec<f32> = logits.iter().map(|&l| (l - mx).exp()).collect();
7377 let s: f32 = e.iter().sum();
7378 for v in &mut e {
7379 *v /= s;
7380 }
7381 e
7382 };
7383 let admit = |e: usize| {
7389 m.mask.as_ref().is_none_or(|mk| mk[e])
7390 && allowed.is_none_or(|a| a.get(e).copied().unwrap_or(false))
7391 };
7392 let mut idx: Vec<usize> = (0..ne).filter(|&e| admit(e)).collect();
7393 match &m.expert_bias {
7395 Some(b) => idx.sort_unstable_by(|&x, &y| {
7396 (p[y] + b[y])
7397 .partial_cmp(&(p[x] + b[x]))
7398 .unwrap()
7399 .then(x.cmp(&y))
7400 }),
7401 None => idx.sort_unstable_by(|&x, &y| p[y].partial_cmp(&p[x]).unwrap().then(x.cmp(&y))),
7402 }
7403 idx.truncate(m.top_k);
7404 if let Some(tau) = m.route_tau {
7408 let total: f32 = idx.iter().map(|&e| p[e]).sum();
7409 if total > 0.0 {
7410 let mut acc = 0.0f32;
7411 let mut keep = idx.len();
7412 for (i, &e) in idx.iter().enumerate() {
7413 acc += p[e];
7414 if acc >= tau * total {
7415 keep = i + 1;
7416 break;
7417 }
7418 }
7419 idx.truncate(keep);
7420 }
7421 }
7422 let wsum: f32 = if m.norm_topk_prob {
7423 let s: f32 = idx.iter().map(|&e| p[e]).sum();
7424 (if m.router_sigmoid { s + 1e-6 } else { s }) / m.routed_scaling
7427 } else {
7428 1.0 / m.routed_scaling
7429 };
7430 (idx, p, wsum)
7431}
7432
7433fn moe_ffn(m: &MoeFfn, x: &[f32], pool: Option<&Pool>, allowed: Option<&[bool]>) -> Vec<f32> {
7436 accumulate_act(m, x, 1);
7437 let ne = m.experts.len();
7438 let mut logits = vec![0.0f32; ne];
7439 m.router.matvec(x, &mut logits, pool);
7440 let (idx, p, wsum) = moe_route(&logits, m, allowed);
7441 {
7442 let mut st = m.stats.borrow_mut();
7443 if st.len() < ne {
7444 st.resize(ne, 0);
7445 }
7446 for &e in &idx {
7447 st[e] += 1;
7448 }
7449 }
7450 if crate::gpu::enabled_here() {
7455 match crate::gpu::probe_arm(crate::gpu::OpClass::Ffn) {
7456 crate::gpu::ProbeArm::Gpu => {
7457 let t0 = std::time::Instant::now();
7458 if let Some(out) = moe_ffn_gpu(m, x, &idx, &p, wsum, pool) {
7459 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
7460 return out;
7461 }
7462 }
7463 crate::gpu::ProbeArm::CpuTimed => {
7464 let t0 = std::time::Instant::now();
7465 let out = crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7466 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
7467 return out;
7468 }
7469 crate::gpu::ProbeArm::Cpu => {
7470 return crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7471 }
7472 }
7473 }
7474 moe_ffn_cpu(m, x, &idx, &p, wsum, pool)
7475}
7476
7477fn graph_note(built: bool) {
7481 use std::sync::atomic::{AtomicBool, Ordering};
7482 if built {
7483 GRAPH_TOK_OK.fetch_add(1, Ordering::Relaxed);
7484 } else {
7485 GRAPH_TOK_MISS.fetch_add(1, Ordering::Relaxed);
7486 }
7487 static SAID: AtomicBool = AtomicBool::new(false);
7488 if !SAID.swap(true, Ordering::Relaxed) {
7489 if built {
7490 tracing::info!("wgpu whole-token graph: ACTIVE");
7491 } else {
7492 tracing::warn!("wgpu whole-token graph refused — per-op path");
7493 }
7494 }
7495}
7496
7497pub static GRAPH_TOK_OK: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7501pub static GRAPH_TOK_MISS: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7502
7503fn moe_batch_enabled() -> bool {
7506 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
7507 *ON.get_or_init(|| std::env::var("CMF_MOE_BATCH").as_deref() != Ok("0"))
7508}
7509
7510fn moe_ffn_cpu_batched(
7516 m: &MoeFfn,
7517 x: &[f32],
7518 idx: &[usize],
7519 p: &[f32],
7520 wsum: f32,
7521 pool: Option<&Pool>,
7522) -> Option<Vec<f32>> {
7523 if idx.is_empty() || !moe_batch_enabled() {
7524 return None;
7525 }
7526 if FFN_PROBE.with(|pr| pr.borrow().is_some()) {
7530 return None;
7531 }
7532 let n = idx.len() + usize::from(m.shared.is_some());
7533 let mut pairs = Vec::with_capacity(n);
7534 let mut downs = Vec::with_capacity(n);
7535 let mut ws = Vec::with_capacity(n);
7536 for &e in idx {
7537 let d = &m.experts[e];
7538 if d.act != Act::Silu {
7539 return None;
7540 }
7541 pairs.push((&d.gate_proj, &d.up_proj));
7542 downs.push(&d.down_proj);
7543 ws.push(p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]));
7544 }
7545 if let Some((se, gate)) = &m.shared {
7548 if se.act != Act::Silu {
7549 return None;
7550 }
7551 let g = gate.as_ref().map_or(1.0, |gate| {
7552 let mut gl = [0.0f32; 1];
7553 gate.matvec(x, &mut gl, pool);
7554 1.0 / (1.0 + (-gl[0]).exp())
7555 });
7556 pairs.push((&se.gate_proj, &se.up_proj));
7557 downs.push(&se.down_proj);
7558 ws.push(g);
7559 }
7560 let inter = pairs[0].0.rows();
7561 let mut gs: Vec<Vec<f32>> = (0..pairs.len()).map(|_| vec![0f32; inter]).collect();
7562 if !QTensor::moe_gate_up_many(&pairs, x, &mut gs, pool) {
7563 return None;
7564 }
7565 let mut out = attention::take_buf(x.len());
7566 if !QTensor::moe_down_many(&downs, &gs, &ws, &mut out, pool) {
7567 attention::recycle_buf(&mut out);
7568 return None;
7569 }
7570 Some(out)
7571}
7572
7573fn moe_ffn_cpu(
7575 m: &MoeFfn,
7576 x: &[f32],
7577 idx: &[usize],
7578 p: &[f32],
7579 wsum: f32,
7580 pool: Option<&Pool>,
7581) -> Vec<f32> {
7582 if let Some(out) = moe_ffn_cpu_batched(m, x, idx, p, wsum, pool) {
7583 return out;
7584 }
7585 let mut out = attention::take_buf(x.len());
7586 for &e in idx {
7587 let mut eo = dense_ffn(&m.experts[e], x, pool);
7588 let w = p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]);
7589 for i in 0..out.len() {
7590 out[i] += w * eo[i];
7591 }
7592 attention::recycle_buf(&mut eo);
7593 }
7594 if let Some((se, gate)) = &m.shared {
7595 let mut so = dense_ffn(se, x, pool);
7596 let g = gate.as_ref().map_or(1.0, |gate| {
7597 let mut gl = [0.0f32; 1];
7598 gate.matvec(x, &mut gl, pool);
7599 1.0 / (1.0 + (-gl[0]).exp())
7600 });
7601 for i in 0..out.len() {
7602 out[i] += g * so[i];
7603 }
7604 attention::recycle_buf(&mut so);
7605 }
7606 out
7607}
7608
7609#[allow(clippy::too_many_arguments)]
7617fn mla_attention(
7618 w: &MlaWeights,
7619 normed: &[f32],
7620 cache: &mut crate::kv_cache::LayerKvCache,
7621 position: usize,
7622 inv_freq: &[f32],
7623 rope_scale: f32,
7624 eps: f64,
7625 pool: Option<&Pool>,
7626) -> Vec<f32> {
7627 let (nh, dr, dn, dv, lora) = (w.nh, w.qk_rope, w.qk_nope, w.v_dim, w.lora);
7628 let hd = dr + dn;
7629 let mut q = vec![0.0f32; nh * hd];
7630 match (&w.q_a, &w.q_a_norm) {
7631 (Some(qa), Some(qn)) => {
7632 let mut t = vec![0.0f32; qa.rows()];
7633 qa.matvec(normed, &mut t, pool);
7634 let tn = inference::rms_norm(&t, qn, eps, NormStyle::Qwen);
7635 w.q_proj.matvec(&tn, &mut q, pool);
7636 }
7637 _ => w.q_proj.matvec(normed, &mut q, pool),
7638 }
7639 let mut ca = vec![0.0f32; lora + dr];
7640 w.kv_a.matvec(normed, &mut ca, pool);
7641 let (c_lat, k_rope) = ca.split_at_mut(lora);
7642 let latn = inference::rms_norm(c_lat, &w.kv_a_norm, eps, NormStyle::Qwen);
7643 let mut kvb = vec![0.0f32; nh * (dn + dv)];
7644 w.kv_b.matvec(&latn, &mut kvb, pool);
7645 if !w.nope {
7646 attention::rope_rotate_scaled(k_rope, position, inv_freq, rope_scale);
7647 }
7648 for h in 0..nh {
7649 if !w.nope {
7650 attention::rope_rotate_scaled(
7651 &mut q[h * hd..h * hd + dr],
7652 position,
7653 inv_freq,
7654 rope_scale,
7655 );
7656 }
7657 }
7658 let mut k = vec![0.0f32; nh * hd];
7659 let mut v = vec![0.0f32; nh * hd];
7660 for h in 0..nh {
7661 k[h * hd..h * hd + dr].copy_from_slice(k_rope);
7662 k[h * hd + dr..(h + 1) * hd].copy_from_slice(&kvb[h * (dn + dv)..h * (dn + dv) + dn]);
7663 v[h * hd..h * hd + dv].copy_from_slice(&kvb[h * (dn + dv) + dn..(h + 1) * (dn + dv)]);
7664 }
7665 cache.append(&k, &v, &vec![true; nh]);
7666 let (ao, mut imp) = attention::attend_all_heads(&q, cache, nh, 1, hd, w.scale, None, 0.0);
7667 attention::recycle_buf(&mut imp);
7668 let mut ov = vec![0.0f32; nh * dv];
7669 for h in 0..nh {
7670 ov[h * dv..(h + 1) * dv].copy_from_slice(&ao[h * hd..h * hd + dv]);
7671 }
7672 let mut out = vec![0.0f32; w.o_proj.rows()];
7673 w.o_proj.matvec(&ov, &mut out, pool);
7674 out
7675}
7676
7677fn dense_moe_ffn(
7684 dm: &DenseMoeFfn,
7685 x_normed: &[f32],
7686 h_raw: &[f32],
7687 eps: f64,
7688 norm_style: NormStyle,
7689 pool: Option<&Pool>,
7690) -> Vec<f32> {
7691 let mut d = dense_ffn(&dm.dense, x_normed, pool);
7692 d = inference::rms_norm(&d, &dm.post_norm_1, eps, norm_style);
7693 let m = &dm.moe;
7694 let ne = m.experts.len();
7695 let mut logits = vec![0.0f32; ne];
7696 if m.router_input_norm {
7697 let ss: f32 = h_raw.iter().map(|v| v * v).sum::<f32>() / h_raw.len() as f32;
7698 let inv = 1.0 / (ss + eps as f32).sqrt();
7699 let xr: Vec<f32> = h_raw.iter().map(|v| v * inv).collect();
7700 m.router.matvec(&xr, &mut logits, pool);
7701 } else {
7702 m.router.matvec(h_raw, &mut logits, pool);
7703 }
7704 let (idx, p, wsum) = moe_route(&logits, m, None);
7705 {
7706 let mut st = m.stats.borrow_mut();
7707 if st.len() < ne {
7708 st.resize(ne, 0);
7709 }
7710 for &e in &idx {
7711 st[e] += 1;
7712 }
7713 }
7714 let x2 = inference::rms_norm(h_raw, &dm.pre_norm_2, eps, norm_style);
7715 let mo = moe_ffn_cpu(m, &x2, &idx, &p, wsum, pool);
7716 let mo = inference::rms_norm(&mo, &dm.post_norm_2, eps, norm_style);
7717 for (di, mi) in d.iter_mut().zip(&mo) {
7718 *di += mi;
7719 }
7720 d
7721}
7722
7723fn moe_gpu_refused(why: &'static str) {
7730 use std::sync::atomic::{AtomicBool, Ordering};
7731 static SAID: AtomicBool = AtomicBool::new(false);
7732 if !SAID.swap(true, Ordering::Relaxed) {
7733 tracing::warn!("MoE GPU block refused ({why}) — experts run on the CPU");
7734 }
7735}
7736
7737fn moe_ffn_gpu(
7738 m: &MoeFfn,
7739 x: &[f32],
7740 idx: &[usize],
7741 p: &[f32],
7742 wsum: f32,
7743 pool: Option<&Pool>,
7744) -> Option<Vec<f32>> {
7745 use crate::gpu::MoeJob;
7746
7747 let mut jobs: Vec<MoeJob> = Vec::with_capacity(idx.len() + 1);
7748 let mut model_ref = None;
7749 for &e in idx {
7750 if moe_push_job(&m.experts[e], x, p[e] / wsum, &mut jobs, &mut model_ref).is_none() {
7751 moe_gpu_refused("push_job(expert)");
7752 return None;
7753 }
7754 }
7755 if let Some((se, gate)) = &m.shared {
7756 let g = gate.as_ref().map_or(1.0, |gate| {
7757 let mut gl = [0.0f32; 1];
7758 gate.matvec(x, &mut gl, pool);
7759 1.0 / (1.0 + (-gl[0]).exp())
7760 });
7761 if moe_push_job(se, x, g, &mut jobs, &mut model_ref).is_none() {
7762 moe_gpu_refused("push_job(shared)");
7763 return None;
7764 }
7765 }
7766 let Some(model) = model_ref else {
7767 moe_gpu_refused("no model_ref");
7768 return None;
7769 };
7770 let hidden = jobs[0].down.1;
7771 let mut out = vec![0.0f32; hidden];
7772 if crate::gpu::moe_block(&model, &jobs, &mut out) {
7773 Some(out)
7774 } else {
7775 moe_gpu_refused("gpu::moe_block");
7776 None
7777 }
7778}
7779
7780fn ffn_forward(
7782 ffn: &FfnKind,
7783 x: &[f32],
7784 pool: Option<&Pool>,
7785 experts_allowed: Option<&[bool]>,
7786) -> Vec<f32> {
7787 match ffn {
7788 FfnKind::Dense(d) => dense_ffn(d, x, pool),
7789 FfnKind::Moe(m) => moe_ffn(m, x, pool, experts_allowed),
7790 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7794 }
7795}
7796
7797fn ffn_forward_pair(
7801 ffn: &FfnKind,
7802 x1: &[f32],
7803 x2: &[f32],
7804 pool: Option<&Pool>,
7805 experts_allowed: Option<&[bool]>,
7806) -> (Vec<f32>, Vec<f32>) {
7807 let d = match ffn {
7808 FfnKind::Dense(d) => d,
7809 FfnKind::Moe(m) => {
7810 return (
7811 moe_ffn(m, x1, pool, experts_allowed),
7812 moe_ffn(m, x2, pool, experts_allowed),
7813 );
7814 }
7815 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7816 };
7817 let inter = d.gate_proj.rows();
7818 FFN_SCRATCH.with(|s| {
7819 let mut s = s.borrow_mut();
7820 let [g1, g2, u1, u2] = &mut *s;
7821 g1.resize(inter, 0.0);
7822 g2.resize(inter, 0.0);
7823 u1.resize(inter, 0.0);
7824 u2.resize(inter, 0.0);
7825 QTensor::matvec2_many(
7828 [&d.gate_proj, &d.up_proj],
7829 x1,
7830 x2,
7831 [g1.as_mut_slice(), u1.as_mut_slice()],
7832 [g2.as_mut_slice(), u2.as_mut_slice()],
7833 pool,
7834 );
7835 for i in 0..inter {
7836 g1[i] = d.act.combine(g1[i], u1[i]);
7837 g2[i] = d.act.combine(g2[i], u2[i]);
7838 }
7839 let mut o1 = attention::take_buf(d.down_proj.rows());
7840 let mut o2 = attention::take_buf(d.down_proj.rows());
7841 d.down_proj.matvec2(g1, g2, &mut o1, &mut o2, pool);
7842 (o1, o2)
7843 })
7844}
7845
7846#[cfg(test)]
7847mod tests {
7848
7849 #[test]
7850 fn cancel_flag_stops_generation() {
7851 let mut p = create_test_pipeline(16, 32, 2, 2, 8, 2, 32);
7852 p.cancel.store(true, std::sync::atomic::Ordering::Relaxed);
7855 let r = p.generate_from_ids(&[1, 2, 3], 8, None, None).unwrap();
7856 assert_eq!(r.finish_reason, "cancelled");
7857 assert!(
7858 r.token_ids.is_empty(),
7859 "no tokens after cancel: {:?}",
7860 r.token_ids
7861 );
7862 let r2 = p.generate_from_ids(&[1, 2, 3], 4, None, None).unwrap();
7864 assert_ne!(r2.finish_reason, "cancelled");
7865 }
7866 use super::*;
7867
7868 #[test]
7874 fn sparse_ffn_quant_equals_dense_with_inactive_zeroed() {
7875 let (hidden, inter) = (16usize, 40usize);
7876 let synth = |n: usize, salt: usize| -> Vec<f32> {
7877 (0..n)
7878 .map(|i| (((i * 37 + salt * 11 + 3) % 101) as f32 / 101.0 - 0.5) * 0.4)
7879 .collect()
7880 };
7881 let d = DenseFfn {
7882 gate_proj: QTensor::from_f32(synth(inter * hidden, 1), inter, hidden),
7883 up_proj: QTensor::from_f32(synth(inter * hidden, 2), inter, hidden),
7884 down_proj: QTensor::from_f32(synth(hidden * inter, 3), hidden, inter),
7885 act: Act::Silu,
7886 };
7887 let x = synth(hidden, 9);
7888 let active: Vec<u16> = (0..inter as u16).filter(|i| i % 3 == 0).collect();
7890
7891 let sparse = sparse_ffn_quant(&d, &x, &active, hidden, None);
7892
7893 let mut g = vec![0.0f32; inter];
7895 d.gate_proj.matvec(&x, &mut g, None);
7896 let mut u = vec![0.0f32; inter];
7897 d.up_proj.matvec(&x, &mut u, None);
7898 let act_set: std::collections::HashSet<u16> = active.iter().copied().collect();
7899 for i in 0..inter {
7900 g[i] = if act_set.contains(&(i as u16)) {
7901 inference::silu(g[i]) * u[i]
7902 } else {
7903 0.0
7904 };
7905 }
7906 let mut reference = vec![0.0f32; hidden];
7907 d.down_proj.matvec(&g, &mut reference, None);
7908
7909 let max_d = sparse
7910 .iter()
7911 .zip(&reference)
7912 .map(|(a, b)| (a - b).abs())
7913 .fold(0.0f32, f32::max);
7914 assert!(max_d < 1e-5, "sparse != dense-zeroed: max|Δ| = {max_d}");
7915 }
7916
7917 fn attach_test_mtp(p: &mut Pipeline) {
7919 let (h, inter, heads, kv, hd) = (
7920 p.hidden_size,
7921 p.intermediate_size,
7922 p.num_heads,
7923 p.num_kv_heads,
7924 p.head_dim,
7925 );
7926 let synth = |n: usize, salt: usize| -> Vec<f32> {
7927 (0..n)
7928 .map(|i| (((i * 29 + salt * 23 + 5) % 101) as f32 / 101.0 - 0.5) * 0.2)
7929 .collect()
7930 };
7931 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
7932 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
7933 };
7934 p.mtp = Some(MtpModule {
7935 enorm: vec![1.0; h],
7936 hnorm: vec![1.0; h],
7937 eh_proj: qt(h, 2 * h, 301),
7938 layer: LayerWeights {
7939 input_norm: vec![1.0; h],
7940 post_norm: vec![1.0; h],
7941 attn_out_norm: None,
7942 ffn_out_norm: None,
7943 layer_scale: None,
7944 ffn: FfnKind::Dense(DenseFfn {
7945 gate_proj: qt(inter, h, 315),
7946 up_proj: qt(inter, h, 316),
7947 down_proj: qt(h, inter, 317),
7948 act: Act::Silu,
7949 }),
7950 attn: AttnKind::Full {
7951 bias: None,
7952 wq: qt(heads * hd, h, 311),
7953 wk: qt(kv * hd, h, 312),
7954 wv: qt(kv * hd, h, 313),
7955 wo: qt(h, heads * hd, 314),
7956 q_norm: None,
7957 k_norm: None,
7958 output_gate: false,
7959 softplus_gate: None,
7960 },
7961 },
7962 final_norm: vec![1.0; h],
7963 kv: crate::kv_cache::LayerKvCache::new(kv, hd),
7964 });
7965 }
7966
7967 #[test]
7968 fn speculative_equals_vanilla_greedy() {
7969 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
7973 let run = |spec: bool| {
7974 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
7975 p.sampler_config.temperature = 0.0;
7976 attach_test_mtp(&mut p);
7977 p.speculative = spec;
7978 let r = p.generate("abcdef", 12, None, None).unwrap();
7979 (r.token_ids, r.mtp_drafted, r.mtp_accepted)
7980 };
7981 let (vanilla, d0, _) = run(false);
7982 let (spec, d1, a1) = run(true);
7983 assert_eq!(d0, 0, "vanilla path must not draft");
7984 assert!(d1 > 0, "speculative path must draft");
7985 assert_eq!(
7986 vanilla, spec,
7987 "speculative must reproduce the exact greedy sequence (accepted {a1}/{d1})"
7988 );
7989 }
7990
7991 #[test]
7992 fn speculative_accepts_constant_oracle() {
7993 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
7995 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
7996 p.sampler_config.temperature = 0.0;
7997 p.sampler_config.repetition_penalty = 1.0;
7998 p.weights.lm_head = QTensor::from_f32(vec![0.01; 64 * 8], 64, 8);
8001 attach_test_mtp(&mut p);
8002 p.speculative = true;
8003 let r = p.generate("abcd", 10, None, None).unwrap();
8004 assert!(r.mtp_drafted > 0);
8005 assert_eq!(
8006 r.mtp_accepted, r.mtp_drafted,
8007 "constant logits → every draft accepted"
8008 );
8009 assert!(r.token_ids.windows(2).all(|w| w[0] == w[1]));
8012 }
8013
8014 #[test]
8015 fn empty_prompt_is_an_error_not_a_panic() {
8016 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
8017 let r = p.generate("", 4, None, None);
8018 assert!(r.is_err(), "empty prompt must be a clean error");
8019 }
8020
8021 #[test]
8022 fn every_token_enters_kv_exactly_once() {
8023 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8024 p.sampler_config.temperature = 0.0;
8026 let r = p.generate("abc", 2, None, None).unwrap();
8027 assert_eq!(r.prompt_tokens, 3);
8028 assert_eq!(
8032 p.kv_cache.seq_len(),
8033 3 + r.tokens_generated - 1,
8034 "each token must be cached exactly once (v1 cached the last prompt token twice)"
8035 );
8036 }
8037
8038 #[test]
8039 fn generation_is_reproducible_with_seed() {
8040 let run = || {
8041 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8042 p.generate("hello", 8, None, None).unwrap().token_ids
8043 };
8044 assert_eq!(run(), run());
8045 }
8046
8047 #[test]
8048 fn resetting_sampler_restarts_the_seeded_stream() {
8049 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8050 let config = SamplerConfig {
8051 seed: Some(1234),
8052 ..SamplerConfig::default()
8053 };
8054 p.set_sampler_config(config.clone());
8055 let first = p.generate("hello", 8, None, None).unwrap().token_ids;
8056 p.set_sampler_config(config);
8057 let second = p.generate("hello", 8, None, None).unwrap().token_ids;
8058 assert_eq!(first, second);
8059 }
8060
8061 #[test]
8062 fn eviction_bounds_the_cache() {
8063 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
8064 p.kv_cache.max_seq_len = 6;
8065 p.sampler_config.temperature = 0.0;
8066 let _ = p.generate("abcd", 12, None, None).unwrap();
8067 assert!(
8068 p.kv_cache.seq_len() <= 6 + 1,
8069 "cache must stay bounded by max_seq_len (got {})",
8070 p.kv_cache.seq_len()
8071 );
8072 }
8073
8074 #[test]
8075 fn confidence_matches_tokens_and_is_a_probability() {
8076 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8077 p.sampler_config.temperature = 0.0;
8078 p.sampler_config.repetition_penalty = 1.0;
8079 let r = p.generate("abcd", 10, None, None).unwrap();
8080 assert_eq!(
8081 r.token_confidence.len(),
8082 r.token_ids.len(),
8083 "one confidence per emitted token"
8084 );
8085 for &c in &r.token_confidence {
8086 assert!((0.0..=1.0).contains(&c), "confidence out of [0,1]: {c}");
8087 }
8088 let logits = [1.0f32, 3.0, 0.5, 3.0];
8090 let p0 = top1_prob_t(&logits, 1, 1.0);
8091 let p1 = top1_prob_t(&logits, 3, 1.0);
8092 assert!((p0 - p1).abs() < 1e-6, "equal logits → equal prob");
8093 assert!(p0 > 0.0 && p0 < 1.0);
8094 let sharp = top1_prob_t(&logits, 1, 1.0);
8096 let soft = top1_prob_t(&logits, 1, 2.0);
8097 assert!(soft < sharp, "higher temperature lowers peak confidence");
8098 }
8099
8100 #[test]
8101 fn trace_is_opt_in_and_parallels_the_output() {
8102 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8104 p.sampler_config.temperature = 0.0;
8105 p.sampler_config.repetition_penalty = 1.0;
8106 let r = p.generate("abcd", 10, None, None).unwrap();
8107 assert!(r.traces.is_empty(), "trace must be empty unless enabled");
8108
8109 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8111 p.sampler_config.temperature = 0.0;
8112 p.sampler_config.repetition_penalty = 1.0;
8113 p.set_trace(true);
8114 let r = p.generate("abcd", 10, None, None).unwrap();
8115 assert_eq!(r.traces.len(), r.token_ids.len(), "one trace row per token");
8116 for (i, tr) in r.traces.iter().enumerate() {
8117 assert_eq!(tr.t, i, "trace index is sequential");
8118 assert_eq!(tr.token_id, r.token_ids[i], "trace token_id matches output");
8119 assert_eq!(
8120 tr.confidence, r.token_confidence[i],
8121 "trace confidence matches the confidence channel"
8122 );
8123 assert!(tr.active_skill.is_none() && tr.recon.is_none() && !tr.switched);
8125 }
8126 }
8127
8128 #[test]
8129 fn explain_prefill_logits_match_greedy_first_token() {
8130 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8134 p.sampler_config.temperature = 0.0;
8135 p.sampler_config.repetition_penalty = 1.0;
8136 let ids = p.tokenizer.encode("abcd");
8137 let logits = p.prefill_next_logits(&ids, None);
8138 let argmax = logits
8139 .iter()
8140 .enumerate()
8141 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
8142 .unwrap()
8143 .0 as u32;
8144 let r = p.generate("abcd", 1, None, None).unwrap();
8145 assert_eq!(
8146 argmax, r.token_ids[0],
8147 "explain preview must match greedy emit"
8148 );
8149 }
8150
8151 #[test]
8152 fn laguna_shared_expert_is_unconditionally_added() {
8153 let matrix = |values: Vec<f32>| QTensor::from_f32(values, 2, 2);
8154 let identity = || matrix(vec![1.0, 0.0, 0.0, 1.0]);
8155 let zero_dense = || DenseFfn {
8156 gate_proj: matrix(vec![0.0; 4]),
8157 up_proj: matrix(vec![0.0; 4]),
8158 down_proj: matrix(vec![0.0; 4]),
8159 act: Act::Silu,
8160 };
8161 let shared = DenseFfn {
8162 gate_proj: identity(),
8163 up_proj: identity(),
8164 down_proj: identity(),
8165 act: Act::Silu,
8166 };
8167 let x = [1.0, 2.0];
8168 let expected = dense_ffn(&shared, &x, None);
8169 let moe = MoeFfn {
8170 router: QTensor::from_f32(vec![0.0, 0.0], 1, 2),
8171 experts: vec![zero_dense()],
8172 top_k: 1,
8173 norm_topk_prob: true,
8174 router_sigmoid: true,
8175 expert_bias: None,
8176 routed_scaling: 1.0,
8177 route_tau: None,
8178 shared: Some((shared, None)),
8179 stats: std::cell::RefCell::new(Vec::new()),
8180 act_sq: std::cell::RefCell::new(Vec::new()),
8181 act_rows: std::cell::RefCell::new(Vec::new()),
8182 mask: None,
8183 per_expert_scale: None,
8184 router_input_norm: false,
8185 };
8186 let actual = moe_ffn_cpu(&moe, &x, &[0], &[0.0], 1.0, None);
8187 for (actual, expected) in actual.iter().zip(expected) {
8188 assert!((actual - expected).abs() < 1e-6);
8189 }
8190 }
8191}