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 self.weights
701 .layers
702 .iter()
703 .any(|lw| matches!(&lw.attn, AttnKind::LinearGdn(_)))
704 }
705
706 #[cfg(target_os = "macos")]
707 fn q1_graph_gpu(
708 &mut self,
709 start: usize,
710 upto: Option<usize>,
711 position: usize,
712 h: &mut [f32],
713 ) -> usize {
714 use crate::gpu::{AttnGpuLayer, GdnGpuCfg, GdnGpuLayer, GraphDims, MetalFfn, TokenGraph};
715 if self.attn_softcap > 0.0 || !crate::gpu::enabled_here()
717 || !crate::gpu::q1_force()
718 || std::env::var("CMF_GPU_BLOCK")
719 .map(|v| v == "0")
720 .unwrap_or(false)
721 {
722 if std::env::var("CMF_GRAPH_DBG").is_ok() {
723 eprintln!(
724 "block-graph: front gate (softcap={} enabled_here={} q1_force={})",
725 self.attn_softcap > 0.0,
726 crate::gpu::enabled_here(),
727 crate::gpu::q1_force(),
728 );
729 }
730 return start;
731 }
732 if self.swa.is_some()
737 || self.global_attn.is_some()
738 || self.attention_heads_per_layer.is_some()
739 || self.attn_v_norm
740 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
741 || self.weights.layers.iter().any(|lw| {
742 lw.attn_out_norm.is_some()
743 || lw.ffn_out_norm.is_some()
744 || lw.layer_scale.is_some()
745 || matches!(&lw.ffn, FfnKind::Dense(d) if d.act != Act::Silu)
746 })
747 {
748 if std::env::var("CMF_GRAPH_DBG").is_ok() {
749 eprintln!(
750 "block-graph: arch ineligible (swa={} gattn={} hpl={} vnorm={} scale_delta={:.2e})",
751 self.swa.is_some(),
752 self.global_attn.is_some(),
753 self.attention_heads_per_layer.is_some(),
754 self.attn_v_norm,
755 (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs(),
756 );
757 }
758 return start;
759 }
760 let limit = upto
763 .map(|u| u + 1)
764 .unwrap_or(self.num_layers)
765 .min(self.num_layers);
766
767 enum Item<'a> {
768 Gdn {
769 run: Vec<GdnGpuLayer<'a>>,
770 first: usize,
771 },
772 Attn {
773 l: AttnGpuLayer<'a>,
774 li: usize,
775 q_norm: Option<&'a [f32]>,
776 k_norm: Option<&'a [f32]>,
777 output_gate: bool,
778 bias: Option<(&'a [f32], &'a [f32], &'a [f32])>,
779 full_gpu: bool,
782 },
783 }
784
785 let attend_mode = std::env::var("CMF_GPU_ATTEND").unwrap_or_else(|_| "auto".into());
792 let attend_contract = attend_mode != "0"
793 && attend_mode != "off"
794 && self.head_dim % 4 == 0
795 && self.head_dim <= 256
796 && self.rotary_dim >= 2
797 && self.rotary_dim <= self.head_dim
798 && (self.rotary_dim / 2) % 32 == 0
799 && self.num_kv_heads > 0
800 && self.num_heads % self.num_kv_heads == 0;
801
802 let mut plan: Vec<Item> = Vec::new();
803 let mut model_ref: Option<std::sync::Arc<cortiq_core::CmfModel>> = None;
804 let block_diag = std::env::var("CMF_GRAPH_DBG").is_ok();
806 let mut scan = start;
807 while scan < limit {
808 let lw = &self.weights.layers[self.phys_layer(scan)];
809 let ffn = match &lw.ffn {
810 FfnKind::Dense(d) => {
811 let (Some(g), Some(u), Some(dn)) = (
812 d.gate_proj.q1_parts(),
813 d.up_proj.q1_parts(),
814 d.down_proj.q1_parts(),
815 ) else {
816 if block_diag {
817 eprintln!(
818 "block-graph: L{scan} FFN trio not graph-mappable — run ends"
819 );
820 }
821 break;
822 };
823 MetalFfn::Dense {
824 gate: g,
825 up: u,
826 down: dn,
827 }
828 }
829 FfnKind::Moe(m) => {
830 let Some(moe) = metal_moe_graph_parts(m, self.hidden_size) else {
831 if block_diag {
832 eprintln!(
833 "block-graph: L{scan} MoE outside the graph contract — run ends"
834 );
835 }
836 break;
837 };
838 if let QTensor::Mapped { model, .. } = &m.experts[0].gate_proj {
839 model_ref.get_or_insert_with(|| model.clone());
840 }
841 MetalFfn::Moe(moe)
842 }
843 _ => {
844 if block_diag {
845 eprintln!("block-graph: L{scan} non-graph FFN — run ends");
846 }
847 break;
848 }
849 };
850 match &lw.attn {
851 AttnKind::LinearGdn(w) if self.gdn_cfg.is_some() => {
852 let parts = (
853 w.in_proj_qkv.q1_parts(),
854 w.in_proj_z.q1_parts(),
855 w.in_proj_a.f32_parts(),
856 w.in_proj_b.f32_parts(),
857 w.out_proj.q1_parts(),
858 );
859 let (Some(qkv), Some(z), Some(a), Some(b), Some(out)) = parts else {
860 if block_diag {
861 eprintln!(
862 "block-graph: L{scan} GDN parts refused (qkv={} z={} a_f32={} b_f32={} out={})",
863 w.in_proj_qkv.q1_parts().is_some(),
864 w.in_proj_z.q1_parts().is_some(),
865 w.in_proj_a.f32_parts().is_some(),
866 w.in_proj_b.f32_parts().is_some(),
867 w.out_proj.q1_parts().is_some(),
868 );
869 }
870 break;
871 };
872 if let QTensor::Mapped { model, .. } = &w.in_proj_qkv {
873 model_ref.get_or_insert_with(|| model.clone());
874 }
875 let gl = GdnGpuLayer {
876 attn_norm: &lw.input_norm,
877 post_norm: &lw.post_norm,
878 qkv,
879 z,
880 a,
881 b,
882 out,
883 ffn,
884 conv1d: &w.conv1d,
885 a_log: &w.a_log,
886 dt_bias: &w.dt_bias,
887 gnorm: &w.norm,
888 };
889 match plan.last_mut() {
890 Some(Item::Gdn { run, .. }) => run.push(gl),
891 _ => plan.push(Item::Gdn {
892 run: vec![gl],
893 first: scan,
894 }),
895 }
896 }
897 AttnKind::Full {
898 wq,
899 wk,
900 wv,
901 wo,
902 q_norm,
903 k_norm,
904 output_gate,
905 softplus_gate: None,
906 bias,
907 } if !self.kv_cache.layers[scan].o1_sealed() => {
908 let parts = (wq.q1_parts(), wk.q1_parts(), wv.q1_parts(), wo.q1_parts());
909 let (Some(pq), Some(pk), Some(pv), Some(po)) = parts else {
910 break;
911 };
912 if let QTensor::Mapped { model, .. } = wq {
913 model_ref.get_or_insert_with(|| model.clone());
914 }
915 let cache = &self.kv_cache.layers[scan];
916 let full_gpu = attend_contract
917 && cache.mode == crate::kv_cache::KvMode::F32
918 && cache.o1.is_none()
919 && bias.is_none()
920 && pq.1 == self.num_heads * self.head_dim * (1 + *output_gate as usize)
921 && pk.1 == self.num_kv_heads * self.head_dim
922 && pv.1 == self.num_kv_heads * self.head_dim
923 && po.2 == self.num_heads * self.head_dim;
924 plan.push(Item::Attn {
925 l: AttnGpuLayer {
926 attn_norm: &lw.input_norm,
927 post_norm: &lw.post_norm,
928 wq: pq,
929 wk: pk,
930 wv: pv,
931 wo: po,
932 ffn,
933 },
934 li: scan,
935 q_norm: q_norm.as_deref(),
936 k_norm: k_norm.as_deref(),
937 output_gate: *output_gate,
938 bias: bias
939 .as_ref()
940 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
941 full_gpu,
942 });
943 }
944 _ => break,
945 }
946 scan += 1;
947 }
948 let Some(model) = model_ref else {
949 if std::env::var("CMF_GRAPH_DBG").is_ok() {
950 eprintln!("q1-graph: no model ref (start {start}, scanned to {scan})");
951 }
952 return start;
953 };
954 if plan.is_empty() {
955 if std::env::var("CMF_GRAPH_DBG").is_ok() {
956 eprintln!("q1-graph: empty plan at layer {start}");
957 }
958 return start;
959 }
960 let has_moe = plan.iter().any(|it| match it {
961 Item::Gdn { run, .. } => run.iter().any(|l| matches!(l.ffn, MetalFfn::Moe(_))),
962 Item::Attn { l, .. } => matches!(l.ffn, MetalFfn::Moe(_)),
963 });
964 let dev_attend = attend_contract
965 && (self.head_dim <= 128
966 || has_moe
967 || attend_mode == "force"
968 || attend_mode == "256");
969 if !dev_attend {
970 for it in &mut plan {
971 if let Item::Attn { full_gpu, .. } = it {
972 *full_gpu = false;
973 }
974 }
975 }
976 if std::env::var("CMF_GRAPH_DBG").is_ok() {
977 use std::sync::atomic::{AtomicBool, Ordering};
978 static SAID: AtomicBool = AtomicBool::new(false);
979 if !SAID.swap(true, Ordering::Relaxed) {
980 let fg = plan
981 .iter()
982 .filter(|it| matches!(it, Item::Attn { full_gpu: true, .. }))
983 .count();
984 let att = plan
985 .iter()
986 .filter(|it| matches!(it, Item::Attn { .. }))
987 .count();
988 eprintln!(
989 "q1-graph: plan of {} items from layer {start} to {scan} | dev_attend={dev_attend} full_gpu {fg}/{att} | hd={} rd={} nkv={} nh={}",
990 plan.len(),
991 self.head_dim,
992 self.rotary_dim,
993 self.num_kv_heads,
994 self.num_heads,
995 );
996 }
997 }
998 let dims = GraphDims {
999 hidden: self.hidden_size,
1000 eps: self.rms_eps as f32,
1001 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1002 };
1003 let Some(mut graph) = TokenGraph::new(&model, dims, h) else {
1004 return start;
1005 };
1006 let gcfg = self.gdn_cfg.map(|cfg| GdnGpuCfg {
1007 nv: cfg.num_v_heads,
1008 nk: cfg.num_k_heads,
1009 dk: cfg.key_head_dim,
1010 dv: cfg.value_head_dim,
1011 kk: cfg.conv_kernel,
1012 hidden: self.hidden_size,
1013 inter: self.intermediate_size,
1014 c_dim: cfg.conv_dim(),
1015 eps: cfg.rms_eps as f32,
1016 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1017 });
1018 let mut valid = 0usize;
1022 let mut end = start;
1023 for item in &plan {
1024 let ok = match item {
1025 Item::Gdn { run, .. } => gcfg
1026 .as_ref()
1027 .map(|gc| run.iter().all(|l| graph.gdn_ok(l, gc)))
1028 .unwrap_or(false),
1029 Item::Attn { l, .. } => graph.attn_ok(l),
1030 };
1031 if !ok {
1032 if block_diag {
1033 eprintln!(
1034 "block-graph: plan item {} ({}) failed graph preflight",
1035 valid,
1036 match item {
1037 Item::Gdn { run, first } =>
1038 format!("GDN run L{first}+{}", run.len()),
1039 Item::Attn { li, .. } => format!("Attn L{li}"),
1040 }
1041 );
1042 }
1043 break;
1044 }
1045 valid += 1;
1046 end += match item {
1047 Item::Gdn { run, .. } => run.len(),
1048 Item::Attn { .. } => 1,
1049 };
1050 }
1051 plan.truncate(valid);
1052 if plan.is_empty() {
1053 return start;
1054 }
1055
1056 let inv_freq = self.inv_freq.clone();
1057 let pool = self.pool.clone();
1058 let (nh, nkv, hd, hs, rd, eps) = (
1059 self.num_heads,
1060 self.num_kv_heads,
1061 self.head_dim,
1062 self.hidden_size,
1063 self.rotary_dim,
1064 self.rms_eps,
1065 );
1066 let norm_style = self.norm_style;
1067 let gemma = norm_style == cortiq_core::NormStyle::Gemma;
1068 let want = self.gdn_cfg.map(|c| c.state_len()).unwrap_or(0);
1069 let kv_id = self.graph_kv_id;
1070 let mut pending: Vec<(usize, usize)> = Vec::new();
1073 let mut dev_attn: Vec<usize> = Vec::new();
1076 for item in &plan {
1077 if self.loop_final_norm {
1079 let item_start = match item {
1080 Item::Gdn { first, .. } => *first,
1081 Item::Attn { li, .. } => *li,
1082 };
1083 if item_start > start && self.is_loop_end(item_start - 1) {
1084 graph.encode_loop_norm(&self.weights.final_norm);
1085 }
1086 }
1087 match item {
1088 Item::Gdn { run, first } => {
1089 for l in &mut self.kv_cache.layers[*first..*first + run.len()] {
1090 if l.linear_state.len() != want {
1091 l.linear_state = vec![0f32; want];
1092 }
1093 }
1094 let ro: Vec<&[f32]> = self.kv_cache.layers[*first..*first + run.len()]
1095 .iter()
1096 .map(|l| l.linear_state.as_slice())
1097 .collect();
1098 if !graph.encode_gdn_run(run, &ro, gcfg.as_ref().unwrap()) {
1099 tracing::error!("q1 graph: GDN run refused after validation");
1101 return start;
1102 }
1103 graph.commit();
1106 pending.push((*first, run.len()));
1107 }
1108 Item::Attn {
1109 l,
1110 li,
1111 q_norm,
1112 k_norm,
1113 output_gate,
1114 bias,
1115 full_gpu,
1116 } => {
1117 if *full_gpu {
1119 let cache = &self.kv_cache.layers[*li];
1120 let cpu_k: Vec<&[f32]> = (0..nkv).map(|g| cache.head_keys(g)).collect();
1121 let cpu_v: Vec<&[f32]> = (0..nkv).map(|g| cache.head_values(g)).collect();
1122 let cpu_stored = cpu_k[0].len() / hd;
1123 let p = crate::gpu::AttnDeviceParams {
1124 kv_id,
1125 layer: *li,
1126 nh,
1127 nkv,
1128 hd,
1129 rd,
1130 position,
1131 eps: eps as f32,
1132 gemma,
1133 output_gate: *output_gate,
1134 q_norm: *q_norm,
1135 k_norm: *k_norm,
1136 inv_freq: &inv_freq,
1137 cpu_k,
1138 cpu_v,
1139 cpu_stored,
1140 };
1141 if graph.attn_device_ok(l, &p) && graph.encode_attn_device(l, &p) {
1142 graph.commit();
1143 dev_attn.push(*li);
1144 continue;
1145 }
1146 }
1148 graph.encode_attn_prefix(l);
1149 graph.sync();
1150 if !pending.is_empty() {
1151 let idxs: Vec<usize> =
1152 pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1153 let mut outs: Vec<&mut [f32]> = self
1154 .kv_cache
1155 .layers
1156 .iter_mut()
1157 .enumerate()
1158 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1159 .map(|(_, s)| s.linear_state.as_mut_slice())
1160 .collect();
1161 graph.read_states(&mut outs);
1162 }
1163 let mut q_raw = attention::take_buf(l.wq.1);
1164 let mut k = attention::take_buf(l.wk.1);
1165 let mut v = attention::take_buf(l.wv.1);
1166 graph.read_qkv(&mut q_raw, &mut k, &mut v);
1167 let cfg = QwenAttnCfg {
1168 num_heads: nh,
1169 num_kv_heads: nkv,
1170 head_dim: hd,
1171 hidden_size: hs,
1172 position,
1173 inv_freq: &inv_freq,
1174 rotary_dim: rd,
1175 scale: self.attn_scale,
1176 softcap: self.attn_softcap,
1177 window: None,
1178 v_norm: false,
1179 q_norm: *q_norm,
1180 k_norm: *k_norm,
1181 output_gate: *output_gate,
1182 softplus_gate: None,
1183 rope_scale: 1.0,
1184 bias: *bias,
1185 rms_eps: eps,
1186 norm_style,
1187 pool: pool.as_deref(),
1188 };
1189 let mut ao = attention::qwen_attention_core(
1190 q_raw,
1191 k,
1192 v,
1193 &mut self.kv_cache.layers[*li],
1194 &cfg,
1195 );
1196 graph.encode_attn_suffix(l, &ao);
1197 graph.commit();
1200 attention::recycle_buf(&mut ao);
1201 }
1202 }
1203 }
1204 let mut lm_rows = None;
1209 if self.graph_want_logits
1210 && upto.is_none()
1211 && end == self.num_layers
1212 && std::env::var("CMF_GPU_LMHEAD")
1213 .map(|v| v != "0")
1214 .unwrap_or(true)
1215 {
1216 if let Some(lm) = self.weights.lm_head.q1_parts() {
1217 if graph.lm_head_ok(lm) {
1218 graph.encode_lm_head(&self.weights.final_norm, lm);
1219 lm_rows = Some(lm.1);
1220 }
1221 }
1222 }
1223 graph.sync();
1224 if !pending.is_empty() {
1225 let idxs: Vec<usize> = pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1226 let mut outs: Vec<&mut [f32]> = self
1227 .kv_cache
1228 .layers
1229 .iter_mut()
1230 .enumerate()
1231 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1232 .map(|(_, s)| s.linear_state.as_mut_slice())
1233 .collect();
1234 graph.read_states(&mut outs);
1235 }
1236 if let Some(rows) = lm_rows {
1237 let mut lg = attention::take_buf(rows.min(self.vocab_size));
1238 graph.read_logits(&mut lg);
1239 lg.resize(self.vocab_size, 0.0);
1240 if let Some(c) = self.final_softcap {
1241 for l in lg.iter_mut() {
1242 *l = c * (*l / c).tanh();
1243 }
1244 }
1245 self.graph_logits = Some(lg);
1246 }
1247 graph.finish(h);
1248 for li in dev_attn {
1252 let mut krow = attention::take_buf(nkv * hd);
1253 let mut vrow = attention::take_buf(nkv * hd);
1254 if crate::gpu::kv_mirror_read_last(kv_id, li, nkv, hd, &mut krow, &mut vrow) {
1255 let cache = &mut self.kv_cache.layers[li];
1256 cache.append(&krow, &vrow, &[]);
1257 let n = cache.seq_len;
1258 let mut imp = attention::take_buf(n);
1259 crate::gpu::kv_mirror_take_imp(kv_id, li, &mut imp);
1260 cache.accumulate_imp(&imp);
1261 attention::recycle_buf(&mut imp);
1262 }
1263 attention::recycle_buf(&mut krow);
1264 attention::recycle_buf(&mut vrow);
1265 }
1266 end
1267 }
1268
1269 pub fn new(
1270 tokenizer: Tokenizer,
1271 weights: PipelineWeights,
1272 hidden_size: usize,
1273 intermediate_size: usize,
1274 num_heads: usize,
1275 num_kv_heads: usize,
1276 head_dim: usize,
1277 num_layers: usize,
1278 physical_layers: usize,
1279 loop_final_norm: bool,
1280 vocab_size: usize,
1281 rms_eps: f64,
1282 rope_base: f32,
1283 norm_style: NormStyle,
1284 max_seq_len: usize,
1285 sampler_config: SamplerConfig,
1286 ) -> Self {
1287 let rng = match sampler_config.seed {
1288 Some(s) => SplitMix64::new(s),
1289 None => SplitMix64::from_entropy(),
1290 };
1291 let inv_freq = std::sync::Arc::new(attention::rope_inv_freq(head_dim, rope_base));
1292 let pool = Pool::from_env();
1293 if let Some(p) = &pool {
1294 tracing::info!("worker pool: {} threads", p.n_workers());
1295 }
1296 Self {
1297 gpu_plan: None,
1298 tokenizer: std::sync::Arc::new(tokenizer),
1299 kv_cache: KvCache::new(num_layers, num_kv_heads, head_dim, max_seq_len),
1300 sampler_config,
1301 weights,
1302 hidden_size,
1303 intermediate_size,
1304 num_heads,
1305 num_kv_heads,
1306 head_dim,
1307 num_layers,
1308 physical_layers,
1309 loop_final_norm,
1310 vocab_size,
1311 rms_eps,
1312 rope_base,
1313 norm_style,
1314 rotary_dim: head_dim,
1315 attention_heads_per_layer: None,
1316 vmf_cfg: None,
1317 gdn_cfg: None,
1318 kda_cfg: None,
1319 g3n: None,
1320 dsv4: None,
1321 dsv4_mtp: Vec::new(),
1322 dspark: None,
1323 dspark_pending: Vec::new(),
1324 dspark_hist: Vec::new(),
1325 dspark_real: Vec::new(),
1326 dspark_trunk_picks: Vec::new(),
1327 dspark_exp: Vec::new(),
1328 dspark_draft_ns: 0,
1329 logit_multiplier: None,
1330 cancel: std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)),
1331 kv_history: Vec::new(),
1332 short_conv_cfg: None,
1333 mtp: None,
1334 speculative: std::env::var("CMF_MTP").map(|v| v != "0").unwrap_or(true),
1335 rng,
1336 sampler_scratch: SamplerScratch::default(),
1337 inv_freq,
1338 ws: ForwardScratch::new(hidden_size),
1339 pool,
1340 model: None,
1341 dyn_force_f32: false,
1342 dyn_skill_layers: Vec::new(),
1343 dyn_active: None,
1344 dyn_blend_loaded: false,
1345 dyn_phi_layer: None,
1346 dyn_phi_ema: Vec::new(),
1347 dyn_phi_seen: 0,
1348 dyn_router: None,
1349 o1_cfg: None,
1350 o1_epoch: 0,
1351 o1_flags: Vec::new(),
1352 trace: false,
1353 calib_temp: 1.0,
1354 confidence_on: true,
1355 embed_multiplier: 1.0,
1356 attn_scale: 1.0 / (head_dim as f32).sqrt(),
1357 swa: None,
1358 sliding_layers: None,
1359 inv_freq_local: None,
1360 rotary_dim_local: None,
1361 rope_scale: 1.0,
1362 rope_scale_local: 1.0,
1363 global_attn: None,
1364 inv_freq_global: None,
1365 attn_v_norm: false,
1366 final_softcap: None,
1367 attn_softcap: 0.0,
1368 graph_want_logits: false,
1369 graph_logits: None,
1370 graph_kv_id: {
1371 static NEXT: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(1);
1372 NEXT.fetch_add(1, std::sync::atomic::Ordering::Relaxed)
1373 },
1374 }
1375 }
1376
1377 pub fn set_o1(&mut self, cfg: Option<crate::nystrom::O1Cfg>) {
1384 self.o1_flags = match &cfg {
1385 Some(c) => {
1386 let mut flags = c.layer_flags(self.num_layers);
1387 for (li, f) in flags.iter_mut().enumerate() {
1388 if *f
1389 && !matches!(
1390 self.weights.layers[self.phys_layer(li)].attn,
1391 AttnKind::Full { .. }
1392 )
1393 {
1394 *f = false;
1395 }
1396 }
1397 flags
1398 }
1399 None => Vec::new(),
1400 };
1401 if let Some(c) = &cfg {
1402 let n = self.o1_flags.iter().filter(|&&f| f).count();
1403 tracing::info!(
1404 "o1 nystrom attention: {n}/{} layer(s), m={} w={} sink={} rect={:?}",
1405 self.num_layers,
1406 c.m,
1407 c.w,
1408 c.sink,
1409 c.rect
1410 );
1411 }
1412 self.o1_cfg = cfg;
1413 }
1414
1415 pub fn o1_active(&self) -> bool {
1417 self.o1_cfg.is_some() && self.o1_flags.iter().any(|&f| f)
1418 }
1419
1420 pub fn o1_begin(&mut self) {
1425 if let Some(c) = &self.o1_cfg {
1426 let (m, w, sink, rect) = (c.m, c.w, c.sink, c.rect);
1427 for (li, &f) in self.o1_flags.iter().enumerate() {
1428 if f {
1429 self.kv_cache.layers[li].o1_begin(m, w, sink, rect);
1430 }
1431 }
1432 }
1433 }
1434
1435 pub fn o1_seal(&mut self) {
1439 self.o1_epoch = self.o1_epoch.wrapping_add(1);
1440 if self.o1_cfg.is_none() {
1441 return;
1442 }
1443 for li in 0..self.num_layers {
1444 if self.o1_flags.get(li).copied().unwrap_or(false) {
1445 self.kv_cache.layers[li].o1_seal(self.num_heads);
1446 }
1447 }
1448 }
1449
1450 pub fn set_trace(&mut self, on: bool) {
1452 self.trace = on;
1453 }
1454
1455 pub fn set_sampler_config(&mut self, config: SamplerConfig) {
1458 self.rng = match config.seed {
1459 Some(seed) => SplitMix64::new(seed),
1460 None => SplitMix64::from_entropy(),
1461 };
1462 self.sampler_config = config;
1463 }
1464
1465 pub fn set_confidence(&mut self, on: bool) {
1470 self.confidence_on = on;
1471 }
1472
1473 pub fn set_calib_temp(&mut self, t: f32) {
1476 self.calib_temp = if t > 1e-3 { t } else { 1.0 };
1477 }
1478
1479 pub fn calib_temp(&self) -> f32 {
1481 self.calib_temp
1482 }
1483
1484 pub fn set_rotary(&mut self, rotary_dim: usize, base: f32) {
1487 self.rotary_dim = rotary_dim.min(self.head_dim);
1488 self.inv_freq = std::sync::Arc::new(attention::rope_inv_freq(self.rotary_dim, base));
1489 }
1490
1491 fn attn_cfg(&self, position: usize) -> QwenAttnCfg<'_> {
1492 QwenAttnCfg {
1493 num_heads: self.num_heads,
1494 num_kv_heads: self.num_kv_heads,
1495 head_dim: self.head_dim,
1496 hidden_size: self.hidden_size,
1497 position,
1498 inv_freq: &self.inv_freq,
1499 rotary_dim: self.rotary_dim,
1500 scale: self.attn_scale,
1501 softcap: self.attn_softcap,
1502 window: None,
1503 v_norm: false,
1504 q_norm: None,
1505 k_norm: None,
1506 output_gate: false,
1507 softplus_gate: None,
1508 rope_scale: self.rope_scale,
1509 bias: None,
1510 rms_eps: self.rms_eps,
1511 norm_style: self.norm_style,
1512 pool: self.pool.as_deref(),
1513 }
1514 }
1515
1516 pub fn generate(
1518 &mut self,
1519 prompt: &str,
1520 max_tokens: usize,
1521 task_mask: Option<&TaskMask>,
1522 on_token: Option<TokenCallback>,
1523 ) -> Result<GenerateResult, String> {
1524 let input_ids = self.tokenizer.with_bos(self.tokenizer.encode(prompt));
1525 self.generate_from_ids(&input_ids, max_tokens, task_mask, on_token)
1526 }
1527
1528 pub fn generate_from_ids(
1536 &mut self,
1537 input_ids: &[u32],
1538 max_tokens: usize,
1539 task_mask: Option<&TaskMask>,
1540 mut on_token: Option<TokenCallback>,
1541 ) -> Result<GenerateResult, String> {
1542 if std::env::var("CMF_TRACE_H").is_ok() {
1543 eprintln!("input_ids: {input_ids:?}");
1544 }
1545 if input_ids.is_empty() {
1546 return Err("empty prompt: nothing to generate from".to_string());
1547 }
1548
1549 let reuse_from = {
1557 let on = !std::env::var("CMF_KV_REUSE").is_ok_and(|v| v == "0");
1558 let h = &self.kv_history;
1559 if on
1560 && task_mask.is_none()
1561 && self.mtp.is_none()
1562 && self.o1_cfg.is_none()
1563 && !h.is_empty()
1564 && h.len() < input_ids.len()
1565 && input_ids[..h.len()] == h[..]
1566 {
1567 h.len()
1568 } else {
1569 0
1570 }
1571 };
1572 if reuse_from == 0 {
1573 self.kv_cache.clear();
1575 self.kv_history.clear();
1576 crate::gpu::graph_kv_reset(self.graph_kv_id);
1577 } else if std::env::var("CMF_PREFILL_PROF").is_ok() {
1578 eprintln!(
1579 "kv-reuse: {} of {} prompt positions already cached",
1580 reuse_from,
1581 input_ids.len()
1582 );
1583 }
1584 crate::gpu::graph_race_begin_generation();
1585 self.o1_begin();
1586
1587 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
1593 let graph_spec = self.speculative
1609 && graph_on
1610 && self.mtp.is_some()
1611 && task_mask.is_none()
1612 && !self.o1_active()
1613 && self.sampler_config.temperature < 1e-6
1614 && self.sampler_config.repetition_penalty == 1.0
1615 && std::env::var("CMF_GRAPH_SPEC").is_ok_and(|v| v != "0");
1616 let pair_pays = self.gdn_cfg.is_none()
1623 || std::env::var("CMF_MTP").as_deref() == Ok("1");
1624 let spec_active = self.speculative
1625 && self.mtp.is_some()
1626 && task_mask.is_none()
1627 && !self.o1_active()
1628 && ((!graph_on && pair_pays) || graph_spec)
1629 && self.sampler_config.temperature < 1e-6;
1630 let mut mtp = if spec_active { self.mtp.take() } else { None };
1633 if std::env::var("CMF_MTP_CHAIN_PROBE").is_ok() {
1634 eprintln!(
1635 "mtp-probe gate: spec_active={spec_active} mtp={} speculative={} graph_on={graph_on} temp_ok={}",
1636 mtp.is_some(),
1637 self.speculative,
1638 self.sampler_config.temperature < 1e-6,
1639 );
1640 }
1641 if let Some(m) = &mut mtp {
1642 m.kv.clear();
1643 }
1644 let mut router = if mtp.is_none() {
1648 self.dyn_router.take()
1649 } else {
1650 None
1651 };
1652 if let Some(r) = &mut router {
1653 r.reset(); self.dyn_phi_seen = 0; let _ = self.set_active_skill(None);
1656 }
1657
1658 let mut all_ids = input_ids.to_vec();
1659 let mut generated = 0usize;
1660 let mut finish_reason = "max_tokens".to_string();
1661 let mut drafted = 0usize;
1662 let mut accepted = 0usize;
1663 let mut confidence: Vec<f32> = Vec::new();
1664 let trace_on = self.trace;
1665 let calib_temp = self.calib_temp;
1666 let mut traces: Vec<TokenTrace> = Vec::new();
1667
1668 let mut hidden = vec![0.0f32; self.hidden_size];
1674 let mut pos = reuse_from;
1675 let fuse_lm = mtp.is_none()
1684 && router.is_none()
1685 && std::env::var("CMF_GPU_LMHEAD").as_deref() != Ok("0");
1686 self.graph_logits = None;
1687 self.graph_want_logits = false;
1688 let _tpf = std::time::Instant::now();
1689 let batch_k = std::env::var("CMF_BATCH_K")
1690 .ok()
1691 .and_then(|v| v.parse::<usize>().ok())
1692 .unwrap_or(0);
1693 while self.dsv4.is_some()
1704 && mtp.is_none()
1705 && pos < input_ids.len()
1706 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1707 {
1708 let end = (pos + prefill_chunk()).min(input_ids.len());
1709 let ids: Vec<u32> = input_ids[pos..end].to_vec();
1710 let mut lg = Vec::new();
1711 if let Some(b) = &mut self.dsv4 {
1712 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
1713 crate::dsv4::forward_chunk(
1714 g,
1715 layers,
1716 &cfg,
1717 st,
1718 &ids,
1719 pos,
1720 &self.inv_freq,
1721 self.pool.as_deref(),
1722 &mut lg,
1723 end == input_ids.len(),
1724 );
1725 }
1726 if end == input_ids.len() {
1727 self.graph_logits = Some(lg);
1728 }
1729 pos = end;
1730 hidden = vec![0.0; self.hidden_size];
1731 }
1732 let dyn_prefill = router.is_some();
1737 let graph_prefill = self.graph_prefill_preferred();
1743 if task_mask.is_none()
1744 && !dyn_prefill
1745 && !graph_prefill
1746 && self.can_prefill_batched()
1747 && self.g3n.is_none()
1748 && input_ids.len() > 2
1749 {
1750 let chunk = prefill_chunk();
1756 let hs = self.hidden_size;
1757 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
1758 let end = (pos + chunk).min(input_ids.len());
1759 let hb = self.prefill_batch(&input_ids[pos..end], pos);
1760 if let Some(m) = &mut mtp {
1761 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1762 .ok()
1763 .and_then(|v| v.parse().ok())
1764 .unwrap_or(0);
1765 for p in pos..end {
1766 if p + 1 < input_ids.len() {
1767 if probe >= 1 && p + 2 < input_ids.len() {
1768 let (d1, mut hx) = self.mtp_step_h(
1772 m,
1773 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1774 input_ids[p + 1],
1775 p,
1776 );
1777 let mut ok = d1 == input_ids[p + 2];
1778 Self::chain_probe_note(0, ok);
1779 let mut d_prev = d1;
1780 let mut extra = 0usize;
1781 for j in 1..probe {
1782 if p + 2 + j >= input_ids.len() {
1783 break;
1784 }
1785 let (dj, hj) =
1786 self.mtp_step_h(m, &hx, d_prev, p + 1 + j);
1787 extra += 1;
1788 ok = ok && dj == input_ids[p + 2 + j];
1789 Self::chain_probe_note(j, ok);
1790 d_prev = dj;
1791 hx = hj;
1792 }
1793 m.kv.truncate_last(extra);
1794 } else {
1795 let _ = self.mtp_step(
1796 m,
1797 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1798 input_ids[p + 1],
1799 p,
1800 );
1801 }
1802 }
1803 }
1804 }
1805 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
1806 pos = end;
1807 }
1808 }
1809 let pair_off = std::env::var("CMF_PAIR").is_ok_and(|v| v == "0");
1810 if task_mask.is_none()
1811 && !dyn_prefill
1812 && !graph_prefill
1813 && !pair_off
1814 && self.pair_supported()
1815 {
1816 while pos + 1 < input_ids.len()
1817 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1818 {
1819 let e1 = self.embed_single(input_ids[pos]);
1820 let e2 = self.embed_single(input_ids[pos + 1]);
1821 let (h1, h2) = self.forward_pair(&e1, &e2, pos);
1822 self.commit_linear_scratch();
1824 if let Some(m) = &mut mtp {
1825 let _ = self.mtp_step(m, &h1, input_ids[pos + 1], pos);
1826 if pos + 2 < input_ids.len() {
1827 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1828 .ok()
1829 .and_then(|v| v.parse().ok())
1830 .unwrap_or(0);
1831 if probe >= 1 && pos + 3 < input_ids.len() {
1832 let (d1, mut hx) =
1836 self.mtp_step_h(m, &h2, input_ids[pos + 2], pos + 1);
1837 let mut ok = d1 == input_ids[pos + 3];
1838 Self::chain_probe_note(0, ok);
1839 let mut d_prev = d1;
1840 let mut extra = 0usize;
1841 for j in 1..probe {
1842 if pos + 3 + j >= input_ids.len() {
1843 break;
1844 }
1845 let (dj, hj) =
1846 self.mtp_step_h(m, &hx, d_prev, pos + 2 + j);
1847 extra += 1;
1848 ok = ok && dj == input_ids[pos + 3 + j];
1849 Self::chain_probe_note(j, ok);
1850 d_prev = dj;
1851 hx = hj;
1852 }
1853 m.kv.truncate_last(extra);
1854 } else {
1855 let _ = self.mtp_step(m, &h2, input_ids[pos + 2], pos + 1);
1856 }
1857 }
1858 }
1859 hidden = h2;
1860 pos += 2;
1861 }
1862 }
1863 if batch_k > 0
1872 && graph_prefill
1873 && task_mask.is_none()
1874 && !self.o1_active()
1875 && mtp.is_none()
1876 && !dyn_prefill
1877 && pos + 1 < input_ids.len()
1878 {
1879 let hs = self.hidden_size;
1880 let chunk = batch_k;
1881 while pos < input_ids.len() {
1882 let end = (pos + chunk).min(input_ids.len());
1883 let bk = end - pos;
1884 let mut hiddens = vec![0f32; bk * hs];
1885 for (j, &id) in input_ids[pos..end].iter().enumerate() {
1886 hiddens[j * hs..(j + 1) * hs].copy_from_slice(&self.embed_single(id));
1887 }
1888 let positions: Vec<usize> = (pos..end).collect();
1889 let t_chunk = std::time::Instant::now();
1890 let ok_b = self.try_batch_graph_wgpu(&mut hiddens, &positions, bk, None);
1891 if std::env::var("CMF_GRAPH_PROF").is_ok() {
1892 let ms = t_chunk.elapsed().as_secs_f64() * 1000.0;
1893 eprintln!(
1894 "batch-chunk: k={bk} ok={ok_b} {ms:.1} ms ({:.1} tok/s)",
1895 bk as f64 / (ms / 1000.0)
1896 );
1897 }
1898 {
1899 use std::sync::atomic::{AtomicBool, Ordering};
1900 static SAID: AtomicBool = AtomicBool::new(false);
1901 if !SAID.swap(true, Ordering::Relaxed) {
1902 if ok_b {
1903 tracing::info!("batched prefill: ACTIVE (k={bk})");
1904 } else {
1905 tracing::warn!("batched prefill declined — per-position graph");
1906 }
1907 }
1908 }
1909 if ok_b {
1910 hidden.copy_from_slice(&hiddens[(bk - 1) * hs..]);
1911 pos = end;
1912 } else {
1913 break; }
1915 }
1916 }
1917 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
1918 self.graph_want_logits = fuse_lm && pos + 1 == input_ids.len();
1919 hidden = self.forward_layers(&self.embed_single(input_ids[pos]), pos, task_mask);
1920 if let Some(m) = &mut mtp {
1921 if pos + 1 < input_ids.len() {
1922 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1928 .ok()
1929 .and_then(|v| v.parse().ok())
1930 .unwrap_or(0);
1931 if probe >= 1 && pos + 2 < input_ids.len() {
1932 let (d1, mut hx) = self.mtp_step_h(m, &hidden, input_ids[pos + 1], pos);
1933 let mut ok = d1 == input_ids[pos + 2];
1934 Self::chain_probe_note(0, ok);
1935 let mut d_prev = d1;
1936 let mut extra = 0usize;
1937 for j in 1..probe {
1938 if pos + 2 + j >= input_ids.len() {
1939 break;
1940 }
1941 let (dj, hj) =
1942 self.mtp_step_h(m, &hx, d_prev, pos + 1 + j);
1943 extra += 1;
1944 ok = ok && dj == input_ids[pos + 2 + j];
1945 Self::chain_probe_note(j, ok);
1946 d_prev = dj;
1947 hx = hj;
1948 }
1949 m.kv.truncate_last(extra);
1952 } else {
1953 let _ = self.mtp_step(m, &hidden, input_ids[pos + 1], pos);
1954 }
1955 }
1956 }
1957 pos += 1;
1958 }
1959 if std::env::var("CMF_PREFILL_PROF").is_ok() {
1960 eprintln!(
1961 "prefill: {} tokens in {:.1} ms (batch_k={batch_k})",
1962 input_ids.len(),
1963 _tpf.elapsed().as_secs_f64() * 1000.0
1964 );
1965 }
1966 if self
1969 .cancel
1970 .swap(false, std::sync::atomic::Ordering::Relaxed)
1971 {
1972 self.kv_history.clear();
1973 if let Some(m) = mtp {
1974 self.mtp = Some(m);
1975 }
1976 return Ok(GenerateResult {
1977 text: String::new(),
1978 token_ids: Vec::new(),
1979 prompt_tokens: input_ids.len(),
1980 tokens_generated: 0,
1981 finish_reason: "cancelled".to_string(),
1982 mtp_drafted: 0,
1983 mtp_accepted: 0,
1984 token_confidence: Vec::new(),
1985 traces: Vec::new(),
1986 });
1987 }
1988
1989 self.o1_seal();
1992
1993 macro_rules! commit {
1995 ($id:expr) => {{
1996 all_ids.push($id);
1997 generated += 1;
1998 if self.tokenizer.is_eos($id) {
1999 finish_reason = "stop".to_string();
2000 false
2001 } else {
2002 let token_text = self.tokenizer.decode_token($id);
2003 let mut go = true;
2004 if let Some(ref mut cb) = on_token {
2005 if !cb(&token_text) {
2006 finish_reason = "cancelled".to_string();
2007 go = false;
2008 }
2009 }
2010 go
2011 }
2012 }};
2013 }
2014
2015 let mut next_pos = input_ids.len();
2017 'decode: while generated < max_tokens {
2018 if self
2019 .cancel
2020 .swap(false, std::sync::atomic::Ordering::Relaxed)
2021 {
2022 finish_reason = "cancelled".to_string();
2023 break 'decode;
2024 }
2025 let mut logits = match self.graph_logits.take() {
2026 Some(lg) => lg,
2027 None => {
2028 inference::rms_norm_into(
2029 &hidden,
2030 &self.weights.final_norm,
2031 self.rms_eps,
2032 self.norm_style,
2033 &mut self.ws.n1,
2034 );
2035 self.lm_head_forward(&self.ws.n1)
2036 }
2037 };
2038 let t_next = sampler::sample_with_scratch(
2039 &logits,
2040 &self.sampler_config,
2041 &all_ids,
2042 &mut self.rng,
2043 &mut self.sampler_scratch,
2044 );
2045 if self.confidence_on {
2046 confidence.push(top1_prob_t(&logits, t_next, calib_temp));
2047 }
2048 attention::recycle_buf(&mut logits);
2049 if trace_on {
2050 let skill = router.as_ref().and_then(|r| r.active_id());
2054 traces.push(TokenTrace {
2055 t: generated,
2056 token_id: t_next,
2057 confidence: confidence.last().copied().unwrap_or(0.0),
2058 active_skill: skill,
2059 recon: None,
2060 switched: false,
2061 });
2062 }
2063 if !commit!(t_next) {
2064 break 'decode;
2065 }
2066 if generated >= max_tokens {
2067 break 'decode;
2068 }
2069
2070 if self.kv_cache.needs_eviction() {
2071 static SAID: std::sync::Once = std::sync::Once::new();
2077 SAID.call_once(|| {
2078 tracing::warn!(
2079 "KV cache full at {} positions — evicting half; quality \
2080 will degrade. Raise CMF_MAX_SEQ.",
2081 self.kv_cache.max_seq_len,
2082 );
2083 });
2084 let keep = (self.kv_cache.max_seq_len / 2).max(1);
2085 self.kv_cache.evict(keep);
2086 }
2087
2088 match &mut mtp {
2089 #[cfg(feature = "gpu")]
2091 Some(m) if graph_spec && generated + 1 < max_tokens && next_pos > 0 => {
2092 if let Some((extra, n_pos, new_h)) = self.graph_spec_step(
2093 m,
2094 &hidden,
2095 t_next,
2096 next_pos,
2097 &mut drafted,
2098 &mut accepted,
2099 ) {
2100 next_pos = n_pos;
2101 hidden = new_h;
2102 let mut stopped = false;
2103 for &id in &extra {
2104 if self.confidence_on {
2105 confidence.push(0.0);
2106 }
2107 if !commit!(id) {
2108 stopped = true;
2109 break;
2110 }
2111 }
2112 if stopped {
2113 break 'decode;
2114 }
2115 continue 'decode;
2116 }
2117 hidden = self.forward_layers(&self.embed_single(t_next), next_pos, task_mask);
2119 next_pos += 1;
2120 continue 'decode;
2121 }
2122 Some(m) if !graph_spec && generated + 1 < max_tokens => {
2124 let draft = self.mtp_step(m, &hidden, t_next, next_pos - 1);
2125 drafted += 1;
2126 let emb1 = self.embed_single(t_next);
2127 let emb2 = self.embed_single(draft);
2128 let (h1, h2) = self.forward_pair(&emb1, &emb2, next_pos);
2129
2130 inference::rms_norm_into(
2131 &h1,
2132 &self.weights.final_norm,
2133 self.rms_eps,
2134 self.norm_style,
2135 &mut self.ws.n1,
2136 );
2137 let mut logits1 = self.lm_head_forward(&self.ws.n1);
2138 let t_after = sampler::sample_with_scratch(
2139 &logits1,
2140 &self.sampler_config,
2141 &all_ids,
2142 &mut self.rng,
2143 &mut self.sampler_scratch,
2144 );
2145 if self.confidence_on {
2146 confidence.push(top1_prob_t(&logits1, t_after, calib_temp));
2147 }
2148 attention::recycle_buf(&mut logits1);
2149 if trace_on {
2150 traces.push(TokenTrace {
2153 t: generated,
2154 token_id: t_after,
2155 confidence: confidence.last().copied().unwrap_or(0.0),
2156 active_skill: None,
2157 recon: None,
2158 switched: false,
2159 });
2160 }
2161 let stop = !commit!(t_after);
2162
2163 if t_after == draft {
2164 accepted += 1;
2165 self.commit_linear_scratch();
2166 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2167 hidden = h2;
2168 next_pos += 2;
2169 } else {
2170 for layer in &mut self.kv_cache.layers {
2172 layer.truncate_last(1);
2173 }
2174 if !stop {
2175 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2176 hidden = self.forward_layers(
2177 &self.embed_single(t_after),
2178 next_pos + 1,
2179 None,
2180 );
2181 }
2182 next_pos += 2;
2183 }
2184 if stop {
2185 break 'decode;
2186 }
2187 }
2188 _ => {
2190 #[cfg(feature = "gpu")]
2195 if Self::dsv4_spec_on() && self.dsv4.is_some() {
2196 static SAID: std::sync::Once = std::sync::Once::new();
2197 SAID.call_once(|| {
2198 eprintln!(
2199 "dsv4-spec гейт: mtp={} mask={} router={} trace={} temp={} rep={} ",
2200 !self.dsv4_mtp.is_empty(),
2201 task_mask.is_none(),
2202 router.is_none(),
2203 !trace_on,
2204 self.sampler_config.temperature < 1e-6,
2205 self.sampler_config.repetition_penalty == 1.0,
2206 );
2207 });
2208 }
2209 #[cfg(feature = "gpu")]
2210 if Self::dsv4_spec_on()
2211 && self.dsv4.is_some()
2212 && !self.dsv4_mtp.is_empty()
2213 && task_mask.is_none()
2214 && router.is_none()
2215 && !trace_on
2216 && self.sampler_config.temperature < 1e-6
2217 && self.sampler_config.repetition_penalty == 1.0
2218 && generated + 1 < max_tokens
2219 && all_ids.len() >= 2
2220 {
2221 let tip_token = all_ids[all_ids.len() - 2];
2222 if let Some((extra, n_pos)) = self.dsv4_spec_step(
2223 tip_token,
2224 t_next,
2225 next_pos,
2226 &mut drafted,
2227 &mut accepted,
2228 ) {
2229 next_pos = n_pos;
2230 let mut stopped = false;
2231 for &id in &extra {
2232 if self.confidence_on {
2233 confidence.push(0.0);
2234 }
2235 if !commit!(id) {
2236 stopped = true;
2237 break;
2238 }
2239 }
2240 if stopped {
2241 break 'decode;
2242 }
2243 continue 'decode;
2244 }
2245 }
2246 self.graph_want_logits = fuse_lm;
2247 let mut t_fwd = t_next;
2253 let pure_greedy = self.sampler_config.temperature < 1e-6
2254 && self.sampler_config.repetition_penalty == 1.0
2255 && self.sampler_config.suppress_tokens.is_empty();
2256 let burst_k = std::env::var("CMF_MULTISTEP")
2261 .ok()
2262 .and_then(|v| v.parse::<usize>().ok())
2263 .unwrap_or(0);
2264 if pure_greedy
2265 && burst_k >= 1
2266 && fuse_lm
2267 && task_mask.is_none()
2268 && router.is_none()
2269 && !trace_on
2270 && !self.confidence_on
2271 {
2272 let mut stopped = false;
2273 loop {
2274 let room = max_tokens.saturating_sub(generated);
2275 if room <= 2 {
2276 break;
2277 }
2278 let k = burst_k.min(room - 1);
2279 if k < 1 {
2280 break;
2281 }
2282 let Some(ids) = self.try_multi_burst(t_fwd, next_pos, k) else {
2283 break;
2284 };
2285 next_pos += k;
2286 for &id in &ids {
2287 if !commit!(id) {
2288 stopped = true;
2289 break;
2290 }
2291 }
2292 if stopped {
2293 break;
2294 }
2295 t_fwd = *ids.last().unwrap();
2296 }
2297 if stopped {
2298 break 'decode;
2299 }
2300 }
2301 hidden = self.forward_layers(&self.embed_single(t_fwd), next_pos, task_mask);
2302 next_pos += 1;
2303 if let Some(r) = &mut router {
2306 let phi = self.dyn_phi_ema.clone();
2307 let decision = r.step(&phi, generated);
2308 if let Some(new_active) = decision {
2309 let _ = self.set_active_skill(new_active);
2310 }
2311 if trace_on {
2314 if let Some(last) = traces.last_mut() {
2315 let e = r.last_best_e();
2316 last.recon = e.is_finite().then_some(e);
2317 last.switched = decision.is_some();
2318 }
2319 }
2320 }
2321 }
2322 }
2323 }
2324
2325 self.graph_want_logits = false;
2326 self.graph_logits = None;
2327 if router.is_some() {
2329 let _ = self.set_active_skill(None);
2330 }
2331 self.dyn_router = router.or(self.dyn_router.take());
2332 self.mtp = mtp.or(self.mtp.take());
2333
2334 let output_ids = &all_ids[input_ids.len()..];
2335 let forwarded = input_ids.len() + output_ids.len().saturating_sub(1);
2339 self.kv_history = all_ids[..forwarded.min(all_ids.len())].to_vec();
2340 confidence.truncate(output_ids.len()); traces.truncate(output_ids.len());
2342 Ok(GenerateResult {
2343 text: self.tokenizer.decode(output_ids),
2344 token_ids: output_ids.to_vec(),
2345 prompt_tokens: input_ids.len(),
2346 tokens_generated: generated,
2347 finish_reason,
2348 mtp_drafted: drafted,
2349 mtp_accepted: accepted,
2350 token_confidence: confidence,
2351 traces,
2352 })
2353 }
2354
2355 fn mtp_step(
2359 &mut self,
2360 m: &mut MtpModule,
2361 hidden: &[f32],
2362 next_token: u32,
2363 position: usize,
2364 ) -> u32 {
2365 self.mtp_step_h(m, hidden, next_token, position).0
2366 }
2367
2368 fn chain_probe_note(depth: usize, prefix_ok: bool) {
2372 use std::sync::Mutex;
2373 static T: Mutex<Vec<(u64, u64)>> = Mutex::new(Vec::new());
2374 let mut t = T.lock().unwrap();
2375 if t.len() <= depth {
2376 t.resize(depth + 1, (0, 0));
2377 }
2378 t[depth].0 += 1;
2379 t[depth].1 += prefix_ok as u64;
2380 if depth == 0 && t[0].0 % 128 == 0 {
2381 let line: Vec<String> = t
2382 .iter()
2383 .enumerate()
2384 .map(|(d, (n, k))| format!("d{}={:.0}%({n})", d + 1, 100.0 * *k as f64 / (*n).max(1) as f64))
2385 .collect();
2386 eprintln!("mtp-chain: {}", line.join(" "));
2387 }
2388 }
2389
2390 fn mtp_step_h(
2394 &mut self,
2395 m: &mut MtpModule,
2396 hidden: &[f32],
2397 next_token: u32,
2398 position: usize,
2399 ) -> (u32, Vec<f32>) {
2400 let e = self.embed_single(next_token);
2404 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2405 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2406 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2407 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2408 let mut x = vec![0.0f32; self.hidden_size];
2409 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2410
2411 let lw = &m.layer;
2413 inference::rms_norm_into(
2414 &x,
2415 &lw.input_norm,
2416 self.rms_eps,
2417 self.norm_style,
2418 &mut self.ws.n1,
2419 );
2420 let attn = match &lw.attn {
2421 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2423 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2424 AttnKind::Full {
2425 wq,
2426 wk,
2427 wv,
2428 wo,
2429 q_norm,
2430 k_norm,
2431 output_gate,
2432 softplus_gate,
2433 bias,
2434 } => {
2435 let mut cfg = self.attn_cfg(position);
2436 cfg.q_norm = q_norm.as_deref();
2437 cfg.k_norm = k_norm.as_deref();
2438 cfg.output_gate = *output_gate;
2439 cfg.softplus_gate = softplus_gate
2440 .as_ref()
2441 .map(|(gate, per_head)| (gate, *per_head));
2442 cfg.bias = bias
2443 .as_ref()
2444 .map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2445 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2446 }
2447 AttnKind::Linear(_) | AttnKind::LinearGdn(_) | AttnKind::ShortConv(_) => {
2448 unreachable!("MTP block is full attention")
2449 }
2450 };
2451 for (i, &a) in attn.iter().enumerate() {
2452 x[i] += a;
2453 }
2454 inference::rms_norm_into(
2455 &x,
2456 &lw.post_norm,
2457 self.rms_eps,
2458 self.norm_style,
2459 &mut self.ws.p1,
2460 );
2461 let ffn = ffn_forward(&lw.ffn, &self.ws.p1, self.pool.as_deref(), None);
2462 for (i, &f) in ffn.iter().enumerate() {
2463 x[i] += f;
2464 }
2465
2466 inference::rms_norm_into(
2467 &x,
2468 &m.final_norm,
2469 self.rms_eps,
2470 self.norm_style,
2471 &mut self.ws.n1,
2472 );
2473 let mut lg = self.lm_head_forward(&self.ws.n1);
2474 let draft = sampler::argmax(&lg);
2475 attention::recycle_buf(&mut lg);
2476 (draft, x)
2477 }
2478
2479 fn mtp_warm(&mut self, m: &mut MtpModule, hidden: &[f32], next_token: u32, position: usize) {
2483 let e = self.embed_single(next_token);
2484 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2485 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2486 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2487 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2488 let mut x = vec![0.0f32; self.hidden_size];
2489 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2490 inference::rms_norm_into(&x, &m.layer.input_norm, self.rms_eps, self.norm_style, &mut self.ws.n1);
2491 let attn = match &m.layer.attn {
2492 AttnKind::Full { wq, wk, wv, wo, q_norm, k_norm, output_gate, softplus_gate, bias } => {
2493 let mut cfg = self.attn_cfg(position);
2494 cfg.q_norm = q_norm.as_deref();
2495 cfg.k_norm = k_norm.as_deref();
2496 cfg.output_gate = *output_gate;
2497 cfg.softplus_gate = softplus_gate.as_ref().map(|(g, p)| (g, *p));
2498 cfg.bias = bias.as_ref().map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2499 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2500 }
2501 _ => return,
2502 };
2503 let _ = attn;
2504 }
2505
2506 #[cfg(feature = "gpu")]
2513 #[allow(clippy::too_many_arguments)]
2514 fn graph_spec_step(
2515 &mut self,
2516 m: &mut MtpModule,
2517 hidden: &[f32],
2518 t_next: u32,
2519 next_pos: usize,
2520 drafted: &mut usize,
2521 accepted: &mut usize,
2522 ) -> Option<(Vec<u32>, usize, Vec<f32>)> {
2523 let k_spec: usize = std::env::var("CMF_GRAPH_SPEC_K")
2529 .ok()
2530 .and_then(|v| v.parse().ok())
2531 .filter(|&v| (1..=8).contains(&v))
2532 .unwrap_or(3);
2533 if next_pos == 0 {
2534 return None;
2535 }
2536 let t_round = std::time::Instant::now();
2537 let subs = || crate::gpu_wgpu::SUBMITS.load(std::sync::atomic::Ordering::Relaxed);
2553 let sub0 = subs();
2554 let mut drafts = Vec::with_capacity(k_spec);
2559 let (d1, mut hx) = self.mtp_step_h(m, hidden, t_next, next_pos - 1);
2560 drafts.push(d1);
2561 for j in 1..k_spec {
2562 let (dj, hj) = self.mtp_step_h(m, &hx, drafts[j - 1], next_pos - 1 + j);
2563 drafts.push(dj);
2564 hx = hj;
2565 }
2566 *drafted += k_spec;
2567 let t_draft = t_round.elapsed();
2568 let sub_draft = subs();
2569 let b = k_spec + 1;
2572 let mut hiddens = vec![0.0f32; b * self.hidden_size];
2573 for (i, &t) in std::iter::once(&t_next).chain(drafts.iter()).enumerate() {
2574 let e = self.embed_single(t);
2575 hiddens[i * self.hidden_size..(i + 1) * self.hidden_size].copy_from_slice(&e);
2576 }
2577 let positions: Vec<usize> = (next_pos..next_pos + b).collect();
2578 let (lm_gw, lm_rows) = {
2579 let (_, i, kind, rs) = self.weights.lm_head.graph_weight()?;
2580 (
2581 crate::gpu::GraphW { idx: i, kind, row_scale: rs, data: &[] },
2582 self.weights.lm_head.rows(),
2583 )
2584 };
2585 let mut logits = Vec::new();
2586 let final_norm = self.weights.final_norm.clone();
2587 let ok = self.try_batch_graph_wgpu(
2588 &mut hiddens,
2589 &positions,
2590 b,
2591 Some(crate::gpu::SpecTail {
2592 lm: lm_gw,
2593 lm_rows,
2594 final_norm: &final_norm,
2595 logits_out: &mut logits,
2596 }),
2597 );
2598 if !ok {
2599 m.kv.truncate_last(k_spec);
2602 return None;
2603 }
2604 let t_verify = t_round.elapsed();
2605 let sub_verify = subs();
2606 let ids: Vec<u32> = (0..b)
2608 .map(|i| sampler::argmax(&logits[i * lm_rows..(i + 1) * lm_rows]))
2609 .collect();
2610 let mut a = 0usize;
2611 while a < k_spec && ids[a] == drafts[a] {
2612 a += 1;
2613 }
2614 if a + 1 < b {
2616 crate::gpu::gdn_spec_restore(self.graph_kv_id, a);
2617 }
2618 *accepted += a;
2619 m.kv.truncate_last(k_spec.saturating_sub(1));
2630 let warm_off = std::env::var("CMF_SPEC_WARM").is_ok_and(|v| v == "0");
2631 if !warm_off {
2632 for j in 0..a {
2633 let row = &hiddens[j * self.hidden_size..(j + 1) * self.hidden_size];
2634 let row = row.to_vec();
2635 self.mtp_warm(m, &row, ids[j], next_pos + j);
2636 }
2637 }
2638 let mut row = logits[a * lm_rows..(a + 1) * lm_rows].to_vec();
2640 row.resize(self.vocab_size, 0.0);
2641 if let Some(c) = self.final_softcap {
2642 for l in row.iter_mut() {
2643 *l = c * (*l / c).tanh();
2644 }
2645 }
2646 self.graph_logits = Some(row);
2647 let new_hidden = hiddens[a * self.hidden_size..(a + 1) * self.hidden_size].to_vec();
2648 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
2654 let end = subs();
2655 eprintln!(
2656 "spec-round: draft {:.1} ms/{} sub | verify {:.1} ms/{} sub | \
2657 commit {:.1} ms/{} sub (accepted {a} of {k_spec})",
2658 t_draft.as_secs_f64() * 1e3,
2659 sub_draft - sub0,
2660 (t_verify - t_draft).as_secs_f64() * 1e3,
2661 sub_verify - sub_draft,
2662 (t_round.elapsed() - t_verify).as_secs_f64() * 1e3,
2663 end - sub_verify,
2664 );
2665 }
2666 Some((drafts[..a].to_vec(), next_pos + a + 1, new_hidden))
2667 }
2668
2669 pub fn measure_pair_fusion(&mut self, iters: usize) -> (f64, f64) {
2678 if !self.pair_supported() {
2679 return (0.0, 0.0);
2680 }
2681 let emb1 = self.embed_single(1);
2682 let emb2 = self.embed_single(2);
2683 let pos = self.kv_cache.seq_len();
2684
2685 let t0 = std::time::Instant::now();
2686 for _ in 0..iters {
2687 let _ = self.forward_layers(&emb1, pos, None);
2688 let _ = self.forward_layers(&emb2, pos + 1, None);
2689 for l in &mut self.kv_cache.layers {
2690 l.truncate_last(2);
2691 }
2692 }
2693 let singles_ms = t0.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2694
2695 let t1 = std::time::Instant::now();
2696 for _ in 0..iters {
2697 let _ = self.forward_pair(&emb1, &emb2, pos);
2698 for l in &mut self.kv_cache.layers {
2699 l.truncate_last(2);
2700 }
2701 }
2702 let pair_ms = t1.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2703 (singles_ms, pair_ms)
2704 }
2705
2706 fn pair_supported(&self) -> bool {
2714 !self.weights.layers.is_empty()
2721 && self.g3n.is_none()
2722 && !self
2723 .weights
2724 .layers
2725 .iter()
2726 .any(|lw| matches!(&lw.attn, AttnKind::Mla(_) | AttnKind::Kda(_)))
2727 }
2728
2729 fn forward_pair(
2730 &mut self,
2731 emb1: &[f32],
2732 emb2: &[f32],
2733 position: usize,
2734 ) -> (Vec<f32>, Vec<f32>) {
2735 let mut h1 = emb1.to_vec();
2736 let mut h2 = emb2.to_vec();
2737 let (_nkv, _hd, hs, _rd, eps) = (
2738 self.num_kv_heads,
2739 self.head_dim,
2740 self.hidden_size,
2741 self.rotary_dim,
2742 self.rms_eps,
2743 );
2744 let pool = self.pool.clone();
2745
2746 for li in 0..self.num_layers {
2747 let lw = &self.weights.layers[self.phys_layer(li)];
2748 inference::rms_norm_into(
2751 &h1,
2752 &lw.input_norm,
2753 self.rms_eps,
2754 self.norm_style,
2755 &mut self.ws.n1,
2756 );
2757 inference::rms_norm_into(
2758 &h2,
2759 &lw.input_norm,
2760 self.rms_eps,
2761 self.norm_style,
2762 &mut self.ws.n2,
2763 );
2764
2765 let (a1, a2) = match &lw.attn {
2766 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2767 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2768 AttnKind::Linear(w) => {
2769 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
2770 let layer = &mut self.kv_cache.layers[li];
2771 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2772 vmf_phase_pair(
2773 &self.ws.n1,
2774 &self.ws.n2,
2775 w,
2776 &cfg,
2777 state,
2778 scratch,
2779 self.pool.as_deref(),
2780 )
2781 }
2782 AttnKind::LinearGdn(w) => {
2783 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
2784 let layer = &mut self.kv_cache.layers[li];
2785 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2786 gdn_pair(
2787 &self.ws.n1,
2788 &self.ws.n2,
2789 w,
2790 &cfg,
2791 state,
2792 scratch,
2793 self.pool.as_deref(),
2794 )
2795 }
2796 AttnKind::ShortConv(w) => {
2797 let cfg = self
2798 .short_conv_cfg
2799 .expect("short-conv layer without short_conv_cfg");
2800 let layer = &mut self.kv_cache.layers[li];
2801 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2802 short_conv_pair(
2803 &self.ws.n1,
2804 &self.ws.n2,
2805 w,
2806 &cfg,
2807 state,
2808 scratch,
2809 self.pool.as_deref(),
2810 )
2811 }
2812 AttnKind::Full {
2813 wq,
2814 wk,
2815 wv,
2816 wo,
2817 q_norm,
2818 k_norm,
2819 output_gate,
2820 softplus_gate,
2821 bias,
2822 } => {
2823 let inv_freq_l = self.layer_inv_freq(li);
2824 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
2825 let cfg = QwenAttnCfg {
2826 num_heads: self.layer_num_heads(li),
2827 num_kv_heads: nkv_l,
2828 head_dim: hd_l,
2829 hidden_size: hs,
2830 position,
2831 inv_freq: &inv_freq_l,
2832 rotary_dim: rd_l,
2833 scale: self.attn_scale,
2834 softcap: self.attn_softcap,
2835 window: self.layer_window(li),
2836 v_norm: self.attn_v_norm,
2837 q_norm: q_norm.as_deref(),
2838 k_norm: k_norm.as_deref(),
2839 output_gate: *output_gate,
2840 softplus_gate: softplus_gate
2841 .as_ref()
2842 .map(|(gate, per_head)| (gate, *per_head)),
2843 rope_scale: self.layer_rope_scale(li),
2844 bias: bias
2845 .as_ref()
2846 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
2847 rms_eps: eps,
2848 norm_style: self.norm_style,
2849 pool: pool.as_deref(),
2850 };
2851 attention::qwen_attention_pair(
2852 &self.ws.n1,
2853 &self.ws.n2,
2854 wq,
2855 wk,
2856 wv,
2857 wo,
2858 &mut self.kv_cache.layers[li],
2859 &cfg,
2860 )
2861 }
2862 };
2863 let (a1, a2) = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
2864 Some(w) => (
2865 inference::rms_norm(&a1, w, self.rms_eps, self.norm_style),
2866 inference::rms_norm(&a2, w, self.rms_eps, self.norm_style),
2867 ),
2868 None => (a1, a2),
2869 };
2870 for i in 0..self.hidden_size {
2871 h1[i] += a1[i];
2872 h2[i] += a2[i];
2873 }
2874 let (mut a1, mut a2) = (a1, a2);
2875 attention::recycle_buf(&mut a1);
2876 attention::recycle_buf(&mut a2);
2877
2878 let lw = &self.weights.layers[self.phys_layer(li)];
2879 inference::rms_norm_into(
2880 &h1,
2881 &lw.post_norm,
2882 self.rms_eps,
2883 self.norm_style,
2884 &mut self.ws.p1,
2885 );
2886 inference::rms_norm_into(
2887 &h2,
2888 &lw.post_norm,
2889 self.rms_eps,
2890 self.norm_style,
2891 &mut self.ws.p2,
2892 );
2893 let (f1, f2) = match &lw.ffn {
2894 FfnKind::DenseMoe(dm) => (
2897 dense_moe_ffn(
2898 dm,
2899 &self.ws.p1,
2900 &h1,
2901 self.rms_eps,
2902 self.norm_style,
2903 self.pool.as_deref(),
2904 ),
2905 dense_moe_ffn(
2906 dm,
2907 &self.ws.p2,
2908 &h2,
2909 self.rms_eps,
2910 self.norm_style,
2911 self.pool.as_deref(),
2912 ),
2913 ),
2914 _ => ffn_forward_pair(
2915 &lw.ffn,
2916 &self.ws.p1,
2917 &self.ws.p2,
2918 self.pool.as_deref(),
2919 None,
2920 ),
2921 };
2922 let (f1, f2) = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
2923 Some(w) => (
2924 inference::rms_norm(&f1, w, self.rms_eps, self.norm_style),
2925 inference::rms_norm(&f2, w, self.rms_eps, self.norm_style),
2926 ),
2927 None => (f1, f2),
2928 };
2929 for i in 0..self.hidden_size {
2930 h1[i] += f1[i];
2931 h2[i] += f2[i];
2932 }
2933 let (mut f1, mut f2) = (f1, f2);
2934 attention::recycle_buf(&mut f1);
2935 attention::recycle_buf(&mut f2);
2936 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
2937 for i in 0..self.hidden_size {
2938 h1[i] *= sc;
2939 h2[i] *= sc;
2940 }
2941 }
2942 if self.is_loop_end(li) && li + 1 < self.num_layers {
2944 h1 = inference::rms_norm(
2945 &h1,
2946 &self.weights.final_norm,
2947 self.rms_eps,
2948 self.norm_style,
2949 );
2950 h2 = inference::rms_norm(
2951 &h2,
2952 &self.weights.final_norm,
2953 self.rms_eps,
2954 self.norm_style,
2955 );
2956 }
2957 }
2958 (h1, h2)
2959 }
2960
2961 fn commit_linear_scratch(&mut self) {
2963 for layer in &mut self.kv_cache.layers {
2964 if !layer.linear_scratch.is_empty() {
2965 std::mem::swap(&mut layer.linear_state, &mut layer.linear_scratch);
2966 layer.linear_scratch.clear();
2967 }
2968 }
2969 }
2970
2971 pub fn forward_ids(
2974 &mut self,
2975 ids: &[u32],
2976 task_mask: Option<&TaskMask>,
2977 ) -> Result<Vec<f32>, String> {
2978 if ids.is_empty() {
2979 return Err("empty id sequence".to_string());
2980 }
2981 self.kv_cache.clear();
2982 self.kv_history.clear();
2983 self.o1_begin();
2984 let mut hidden = vec![0.0f32; self.hidden_size];
2985 let mut pos = 0usize;
2986 if self.can_prefill_batched() && !self.graph_prefill_preferred() && ids.len() > 2 {
2994 let chunk = prefill_chunk();
2998 let hs = self.hidden_size;
2999 while pos < ids.len() {
3000 let end = (pos + chunk).min(ids.len());
3001 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3002 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
3003 pos = end;
3004 }
3005 }
3006 if task_mask.is_none()
3015 && !self.graph_prefill_preferred()
3016 && !std::env::var("CMF_PAIR").is_ok_and(|v| v == "0")
3017 && self.pair_supported()
3018 {
3019 while pos + 1 < ids.len() {
3020 let e1 = self.embed_single(ids[pos]);
3021 let e2 = self.embed_single(ids[pos + 1]);
3022 let (_, h2) = self.forward_pair(&e1, &e2, pos);
3023 self.commit_linear_scratch();
3024 hidden = h2;
3025 pos += 2;
3026 }
3027 }
3028 while pos < ids.len() {
3029 hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3030 pos += 1;
3031 }
3032 self.o1_seal();
3036 let normed = inference::rms_norm(
3037 &hidden,
3038 &self.weights.final_norm,
3039 self.rms_eps,
3040 self.norm_style,
3041 );
3042 Ok(self.lm_head_forward(&normed))
3043 }
3044
3045 pub fn ppl_ids(&mut self, ids: &[u32]) -> f64 {
3052 let (nll, cnt) = self.nll_ids_from(ids, 0);
3053 (nll / cnt.max(1) as f64).exp()
3054 }
3055
3056 pub fn probe_ffn_mass(&mut self, ids: &[u32]) -> Vec<Vec<f64>> {
3061 self.kv_cache.clear();
3062 self.kv_history.clear();
3063 FFN_PROBE.with(|p| {
3064 *p.borrow_mut() = Some(vec![vec![0f64; self.intermediate_size]; self.num_layers]);
3065 });
3066 crate::gpu::cpu_scope(|| {
3067 for (pos, &id) in ids.iter().enumerate() {
3068 let emb = self.embed_single(id);
3069 let _ = self.forward_layers(&emb, pos, None);
3070 }
3071 });
3072 self.kv_cache.clear();
3073 self.kv_history.clear();
3074 FFN_PROBE
3075 .with(|p| p.borrow_mut().take())
3076 .unwrap_or_default()
3077 }
3078
3079 pub fn ppl_ids_masked(&mut self, ids: &[u32], mask: &TaskMask) -> f64 {
3083 self.kv_cache.clear();
3084 self.kv_history.clear();
3085 let mut nll = 0f64;
3086 let mut cnt = 0usize;
3087 let mut hidden = vec![0f32; self.hidden_size];
3088 for (pos, &id) in ids.iter().enumerate() {
3089 if pos > 0 {
3090 inference::rms_norm_into(
3091 &hidden,
3092 &self.weights.final_norm,
3093 self.rms_eps,
3094 self.norm_style,
3095 &mut self.ws.n1,
3096 );
3097 let mut logits = self.lm_head_forward(&self.ws.n1);
3098 let max = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
3099 let sum: f64 = logits.iter().map(|&v| ((v - max) as f64).exp()).sum();
3100 let p = ((logits[id as usize] - max) as f64).exp() / sum.max(1e-300);
3101 nll -= p.max(1e-300).ln();
3102 cnt += 1;
3103 attention::recycle_buf(&mut logits);
3104 }
3105 let emb = self.embed_single(id);
3106 hidden = self.forward_layers(&emb, pos, Some(mask));
3107 }
3108 self.kv_cache.clear();
3109 self.kv_history.clear();
3110 (nll / cnt.max(1) as f64).exp()
3111 }
3112
3113 pub fn nll_ids_masked(
3132 &mut self,
3133 ids: &[u32],
3134 start: usize,
3135 task_mask: Option<&TaskMask>,
3136 ) -> (f64, usize) {
3137 self.nll_ids_inner(ids, start, task_mask)
3138 }
3139
3140 pub fn nll_ids_from(&mut self, ids: &[u32], start: usize) -> (f64, usize) {
3141 self.nll_ids_inner(ids, start, None)
3142 }
3143
3144 fn nll_ids_inner(
3145 &mut self,
3146 ids: &[u32],
3147 start: usize,
3148 task_mask: Option<&TaskMask>,
3149 ) -> (f64, usize) {
3150 self.kv_cache.clear();
3151 self.kv_history.clear();
3152 let mut nll = 0f64;
3153 let mut cnt = 0usize;
3154 if self.can_prefill_batched() {
3155 const CHUNK: usize = 128;
3161 const LM_SUB: usize = 32;
3162 let n = ids.len().saturating_sub(1);
3163 let hs = self.hidden_size;
3164 let rows = self.weights.lm_head.rows();
3165 let mut pos = 0usize;
3166 while pos < n {
3167 let end = (pos + CHUNK).min(n);
3168 let bsz = end - pos;
3169 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3170 let mut k0 = 0usize;
3171 while k0 < bsz {
3172 let k1 = (k0 + LM_SUB).min(bsz);
3173 let sb = k1 - k0;
3174 if pos + k1 <= start {
3177 k0 = k1;
3178 continue;
3179 }
3180 let mut normed = vec![0.0f32; sb * hs];
3181 for k in 0..sb {
3182 let r = inference::rms_norm(
3183 &hb[(k0 + k) * hs..(k0 + k + 1) * hs],
3184 &self.weights.final_norm,
3185 self.rms_eps,
3186 self.norm_style,
3187 );
3188 normed[k * hs..(k + 1) * hs].copy_from_slice(&r);
3189 }
3190 let mut logits = vec![0.0f32; sb * rows];
3191 self.weights
3192 .lm_head
3193 .matmat(&normed, sb, &mut logits, self.pool.as_deref());
3194 for k in 0..sb {
3195 if pos + k0 + k < start {
3196 continue;
3197 }
3198 let lg = &mut logits[k * rows..k * rows + self.vocab_size.min(rows)];
3199 if let Some(mu) = self.logit_multiplier {
3200 for v in lg.iter_mut() {
3201 *v *= mu;
3202 }
3203 }
3204 if let Some(c) = self.final_softcap {
3208 for v in lg.iter_mut() {
3209 *v = c * (*v / c).tanh();
3210 }
3211 }
3212 let lg = &logits[k * rows..k * rows + self.vocab_size.min(rows)];
3213 let target = ids[pos + k0 + k + 1] as usize;
3214 let max = lg.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3215 let lse: f64 = lg
3216 .iter()
3217 .map(|&v| ((v - max) as f64).exp())
3218 .sum::<f64>()
3219 .ln()
3220 + max as f64;
3221 nll += lse - lg[target] as f64;
3222 cnt += 1;
3223 if std::env::var("CMF_PPL_TRACE").is_ok() {
3224 let top = lg
3225 .iter()
3226 .enumerate()
3227 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3228 .map(|(i, _)| i)
3229 .unwrap_or(0);
3230 eprintln!(
3231 "BTRACE pos {} target {} nll {:.4} top {} lg_t {:.3} lg_top {:.3}",
3232 pos + k0 + k,
3233 target,
3234 lse - lg[target] as f64,
3235 top,
3236 lg[target],
3237 lg[top]
3238 );
3239 }
3240 }
3241 k0 = k1;
3242 }
3243 pos = end;
3244 }
3245 self.kv_cache.clear();
3246 self.kv_history.clear();
3247 return (nll, cnt);
3248 }
3249 for pos in 0..ids.len().saturating_sub(1) {
3250 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3251 let out_of_band = self.graph_logits.take();
3259 if pos < start {
3260 continue;
3261 }
3262 let logits = match out_of_band {
3263 Some(lg) => lg,
3264 None => {
3265 let normed = inference::rms_norm(
3266 &hidden,
3267 &self.weights.final_norm,
3268 self.rms_eps,
3269 self.norm_style,
3270 );
3271 self.lm_head_forward(&normed)
3275 }
3276 };
3277 let target = ids[pos + 1] as usize;
3278 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3279 let lse: f64 = logits
3280 .iter()
3281 .map(|&v| ((v - max) as f64).exp())
3282 .sum::<f64>()
3283 .ln()
3284 + max as f64;
3285 let tok_nll = lse - logits[target] as f64;
3286 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3287 let top = logits
3288 .iter()
3289 .enumerate()
3290 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3291 .map(|(i, _)| i)
3292 .unwrap_or(0);
3293 eprintln!(
3294 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3295 logits[target], logits[top]
3296 );
3297 }
3298 nll += tok_nll;
3299 cnt += 1;
3300 }
3301 self.kv_cache.clear();
3302 self.kv_history.clear();
3303 (nll, cnt)
3304 }
3305
3306 pub fn nll_ids_o1(&mut self, ids: &[u32], prefill: usize) -> (f64, usize) {
3322 self.kv_cache.clear();
3323 self.kv_history.clear();
3324 self.o1_begin();
3325 let n = ids.len().saturating_sub(1);
3326 let p = prefill.min(n);
3327 let mut pos = 0usize;
3329 if self.can_prefill_batched() {
3330 const CHUNK: usize = 128;
3331 while pos < p {
3332 let end = (pos + CHUNK).min(p);
3333 let _ = self.prefill_batch(&ids[pos..end], pos);
3334 pos = end;
3335 }
3336 } else {
3337 while pos < p {
3338 let _ = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3339 pos += 1;
3340 }
3341 }
3342 self.o1_seal();
3343
3344 let mut nll = 0f64;
3345 let mut cnt = 0usize;
3346 for pos in p..n {
3347 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3348 let normed = inference::rms_norm(
3349 &hidden,
3350 &self.weights.final_norm,
3351 self.rms_eps,
3352 self.norm_style,
3353 );
3354 let logits = self.lm_head_forward(&normed);
3358 let target = ids[pos + 1] as usize;
3359 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3360 let lse: f64 = logits
3361 .iter()
3362 .map(|&v| ((v - max) as f64).exp())
3363 .sum::<f64>()
3364 .ln()
3365 + max as f64;
3366 let tok_nll = lse - logits[target] as f64;
3367 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3368 let top = logits
3369 .iter()
3370 .enumerate()
3371 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3372 .map(|(i, _)| i)
3373 .unwrap_or(0);
3374 eprintln!(
3375 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3376 logits[target], logits[top]
3377 );
3378 }
3379 nll += tok_nll;
3380 cnt += 1;
3381 }
3382 self.kv_cache.clear();
3383 self.kv_history.clear();
3384 (nll, cnt)
3385 }
3386
3387 pub fn calib_ids(&mut self, ids: &[u32], temps: &[f32]) -> (Vec<bool>, Vec<Vec<f32>>) {
3395 self.kv_cache.clear();
3396 self.kv_history.clear();
3397 let n = ids.len().saturating_sub(1);
3398 let mut correct = Vec::with_capacity(n);
3399 let mut pmax = Vec::with_capacity(n);
3400 for pos in 0..n {
3401 let emb = self.embed_single(ids[pos]);
3402 let hidden = self.forward_layers(&emb, pos, None);
3403 let normed = inference::rms_norm(
3404 &hidden,
3405 &self.weights.final_norm,
3406 self.rms_eps,
3407 self.norm_style,
3408 );
3409 let logits = self.lm_head_forward(&normed);
3413 let target = ids[pos + 1] as usize;
3414 let (mut amax, mut mval) = (0usize, f32::NEG_INFINITY);
3415 for (i, &v) in logits.iter().enumerate() {
3416 if v > mval {
3417 mval = v;
3418 amax = i;
3419 }
3420 }
3421 correct.push(amax == target);
3422 let row: Vec<f32> = temps
3423 .iter()
3424 .map(|&t| {
3425 let tt = t.max(1e-3);
3426 let s: f32 = logits.iter().map(|&v| ((v - mval) / tt).exp()).sum();
3427 1.0 / s.max(1e-12) })
3429 .collect();
3430 pmax.push(row);
3431 }
3432 self.kv_cache.clear();
3433 self.kv_history.clear();
3434 (correct, pmax)
3435 }
3436
3437 pub fn ppl_ids_dynamic(&mut self, ids: &[u32]) -> (f64, usize) {
3444 let mut router = match self.dyn_router.take() {
3445 Some(r) => r,
3446 None => return (self.ppl_ids(ids), 0),
3447 };
3448 router.reset();
3449 self.dyn_phi_seen = 0;
3450 let _ = self.set_active_skill(None);
3451
3452 self.kv_cache.clear();
3453
3454 self.kv_history.clear();
3455 let mut nll = 0f64;
3456 let mut cnt = 0usize;
3457 for pos in 0..ids.len().saturating_sub(1) {
3458 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3459 let normed = inference::rms_norm(
3460 &hidden,
3461 &self.weights.final_norm,
3462 self.rms_eps,
3463 self.norm_style,
3464 );
3465 let logits = self.lm_head_forward(&normed);
3469 let target = ids[pos + 1] as usize;
3470 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3471 let lse: f64 = logits
3472 .iter()
3473 .map(|&v| ((v - max) as f64).exp())
3474 .sum::<f64>()
3475 .ln()
3476 + max as f64;
3477 let tok_nll = lse - logits[target] as f64;
3478 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3479 let top = logits
3480 .iter()
3481 .enumerate()
3482 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3483 .map(|(i, _)| i)
3484 .unwrap_or(0);
3485 eprintln!(
3486 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3487 logits[target], logits[top]
3488 );
3489 }
3490 nll += tok_nll;
3491 cnt += 1;
3492 let phi = self.dyn_phi_ema.clone();
3494 if let Some(new_active) = router.step(&phi, pos) {
3495 let _ = self.set_active_skill(new_active);
3496 }
3497 }
3498 let switches = router.switches.len();
3499 let _ = self.set_active_skill(None);
3500 self.dyn_router = Some(router);
3501 self.kv_cache.clear();
3502 self.kv_history.clear();
3503 ((nll / cnt.max(1) as f64).exp(), switches)
3504 }
3505
3506 pub fn probe_phi(&mut self, ids: &[u32], layer: usize) -> Vec<f32> {
3508 self.kv_cache.clear();
3509 self.kv_history.clear();
3510 let mut acc = vec![0f32; self.hidden_size];
3511 for (pos, &id) in ids.iter().enumerate() {
3512 let h = self.forward_layers_upto(&self.embed_single(id), pos, None, Some(layer));
3513 for (a, v) in acc.iter_mut().zip(&h) {
3514 *a += v;
3515 }
3516 }
3517 let n = ids.len().max(1) as f32;
3518 for a in acc.iter_mut() {
3519 *a /= n;
3520 }
3521 self.kv_cache.clear();
3522 self.kv_history.clear();
3523 acc
3524 }
3525
3526 fn prefill_batch(&mut self, ids: &[u32], start_pos: usize) -> Vec<f32> {
3532 self.prefill_batch_masked(ids, start_pos, None)
3533 }
3534
3535 fn prefill_batch_masked(
3541 &mut self,
3542 ids: &[u32],
3543 start_pos: usize,
3544 task_mask: Option<&TaskMask>,
3545 ) -> Vec<f32> {
3546 self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, usize::MAX)
3547 }
3548
3549 fn prefill_batch_span(
3555 &mut self,
3556 input: PrefillIn<'_>,
3557 start_pos: usize,
3558 task_mask: Option<&TaskMask>,
3559 from: usize,
3560 upto_excl: usize,
3561 ) -> Vec<f32> {
3562 let hs = self.hidden_size;
3563 let b = match input {
3564 PrefillIn::Ids(ids) => ids.len(),
3565 PrefillIn::Hidden(hb) => hb.len() / hs,
3566 };
3567 let upto_excl = upto_excl.min(self.num_layers);
3568 let mut h: Vec<f32>;
3572 let mut h_ready;
3573 match input {
3574 PrefillIn::Ids(_) => {
3575 h = vec![0.0; b * hs];
3576 h_ready = false;
3577 }
3578 PrefillIn::Hidden(hb) => {
3579 h = hb.to_vec();
3580 h_ready = true;
3581 }
3582 }
3583 let fill_h = |h: &mut Vec<f32>, me: &Self| {
3584 if let PrefillIn::Ids(ids) = input {
3585 for (bi, &id) in ids.iter().enumerate() {
3586 let e = me.embed_single(id);
3587 h[bi * hs..(bi + 1) * hs].copy_from_slice(&e);
3588 }
3589 }
3590 };
3591 let (_nkv, _hd, _rd, eps) = (
3592 self.num_kv_heads,
3593 self.head_dim,
3594 self.rotary_dim,
3595 self.rms_eps,
3596 );
3597 let pool = self.pool.clone();
3598 let norm_style = self.norm_style;
3599
3600 #[cfg(target_os = "macos")]
3601 let mut chunk_skip_until = 0usize;
3602 for li in from..upto_excl {
3603 crate::gpu::set_layer(li as i64); #[cfg(target_os = "macos")]
3610 if task_mask.is_none() {
3611 if li < chunk_skip_until {
3612 continue;
3613 }
3614 if !h_ready && li == 0 && self.weights.embed_tokens.q8_row_parts().is_none() {
3620 fill_h(&mut h, self);
3621 h_ready = true;
3622 }
3623 let ids_for_embed = match input {
3624 PrefillIn::Ids(ids) => (!h_ready && li == 0).then_some(ids),
3625 PrefillIn::Hidden(_) => None,
3626 };
3627 let end = self.chunk_run_gpu(li, &mut h, b, start_pos, ids_for_embed, upto_excl);
3628 if end > li {
3629 h_ready = true;
3630 chunk_skip_until = end;
3631 if self.is_loop_end(end - 1) && end < self.num_layers {
3634 for bi in 0..b {
3635 let normed = inference::rms_norm(
3636 &h[bi * hs..(bi + 1) * hs],
3637 &self.weights.final_norm,
3638 eps,
3639 norm_style,
3640 );
3641 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
3642 }
3643 }
3644 continue;
3645 }
3646 }
3647 if !h_ready {
3648 fill_h(&mut h, self);
3649 h_ready = true;
3650 }
3651 let lw = &self.weights.layers[self.phys_layer(li)];
3652 match &lw.attn {
3654 AttnKind::Kda(w) => {
3655 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
3657 let mut normed = vec![0.0f32; b * hs];
3658 for bi in 0..b {
3659 inference::rms_norm_into(
3660 &h[bi * hs..(bi + 1) * hs],
3661 &lw.input_norm,
3662 eps,
3663 norm_style,
3664 &mut normed[bi * hs..(bi + 1) * hs],
3665 );
3666 }
3667 let attn = crate::linear_core::kda_forward_batch(
3668 &normed,
3669 b,
3670 w,
3671 &cfg,
3672 &mut self.kv_cache.layers[li].linear_state,
3673 pool.as_deref(),
3674 );
3675 for (dst, &a) in h.iter_mut().zip(&attn) {
3676 *dst += a;
3677 }
3678 }
3679 AttnKind::LinearGdn(w) => {
3680 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
3682 let mut normed = vec![0.0f32; b * hs];
3683 for bi in 0..b {
3684 let r = inference::rms_norm(
3685 &h[bi * hs..(bi + 1) * hs],
3686 &lw.input_norm,
3687 eps,
3688 norm_style,
3689 );
3690 normed[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3691 }
3692 let attn = crate::linear_core::gdn_forward_batch(
3693 &normed,
3694 b,
3695 w,
3696 &cfg,
3697 &mut self.kv_cache.layers[li].linear_state,
3698 pool.as_deref(),
3699 );
3700 for (dst, &a) in h.iter_mut().zip(&attn) {
3701 *dst += a;
3702 }
3703 }
3704 AttnKind::ShortConv(w) => {
3705 let cfg = self
3708 .short_conv_cfg
3709 .expect("short-conv layer without short_conv_cfg");
3710 let mut normed = vec![0.0f32; b * hs];
3711 for bi in 0..b {
3712 inference::rms_norm_into(
3713 &h[bi * hs..(bi + 1) * hs],
3714 &lw.input_norm,
3715 eps,
3716 norm_style,
3717 &mut normed[bi * hs..(bi + 1) * hs],
3718 );
3719 }
3720 let attn = short_conv_forward_batch(
3721 &normed,
3722 b,
3723 w,
3724 &cfg,
3725 &mut self.kv_cache.layers[li].linear_state,
3726 pool.as_deref(),
3727 );
3728 for (dst, &a) in h.iter_mut().zip(&attn) {
3729 *dst += a;
3730 }
3731 }
3732 AttnKind::Mla(w) => {
3733 let inv_freq_l = self.layer_inv_freq(li);
3736 let rs = self.layer_rope_scale(li);
3737 let mut normed = vec![0.0f32; hs];
3738 for bi in 0..b {
3739 inference::rms_norm_into(
3740 &h[bi * hs..(bi + 1) * hs],
3741 &lw.input_norm,
3742 eps,
3743 norm_style,
3744 &mut normed,
3745 );
3746 let ao = mla_attention(
3747 w,
3748 &normed,
3749 &mut self.kv_cache.layers[li],
3750 start_pos + bi,
3751 &inv_freq_l,
3752 rs,
3753 eps,
3754 pool.as_deref(),
3755 );
3756 for (dst, &a) in h[bi * hs..(bi + 1) * hs].iter_mut().zip(&ao) {
3757 *dst += a;
3758 }
3759 }
3760 }
3761 AttnKind::Full {
3762 wq,
3763 wk,
3764 wv,
3765 wo,
3766 q_norm,
3767 k_norm,
3768 output_gate,
3769 softplus_gate,
3770 bias,
3771 } => {
3772 let mut normed = vec![0.0f32; b * hs];
3776 for bi in 0..b {
3777 inference::rms_norm_into(
3778 &h[bi * hs..(bi + 1) * hs],
3779 &lw.input_norm,
3780 eps,
3781 norm_style,
3782 &mut normed[bi * hs..(bi + 1) * hs],
3783 );
3784 }
3785 let inv_freq_l = self.layer_inv_freq(li);
3786 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
3787 let cfg = QwenAttnCfg {
3788 num_heads: self.layer_num_heads(li),
3789 num_kv_heads: nkv_l,
3790 head_dim: hd_l,
3791 hidden_size: hs,
3792 position: start_pos,
3793 inv_freq: &inv_freq_l,
3794 rotary_dim: rd_l,
3795 scale: self.attn_scale,
3796 softcap: self.attn_softcap,
3797 window: self.layer_window(li),
3798 v_norm: self.attn_v_norm,
3799 q_norm: q_norm.as_deref(),
3800 k_norm: k_norm.as_deref(),
3801 output_gate: *output_gate,
3802 softplus_gate: softplus_gate
3803 .as_ref()
3804 .map(|(gate, per_head)| (gate, *per_head)),
3805 rope_scale: self.layer_rope_scale(li),
3806 bias: bias
3807 .as_ref()
3808 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
3809 rms_eps: eps,
3810 norm_style,
3811 pool: pool.as_deref(),
3812 };
3813 let mut attn = attention::qwen_attention_batch(
3814 &normed,
3815 b,
3816 wq,
3817 wk,
3818 wv,
3819 wo,
3820 &mut self.kv_cache.layers[li],
3821 &cfg,
3822 );
3823 if let Some(w) = &lw.attn_out_norm {
3824 for bi in 0..b {
3825 inference::rms_norm_into(
3826 &attn[bi * hs..(bi + 1) * hs],
3827 w,
3828 eps,
3829 norm_style,
3830 &mut normed[bi * hs..(bi + 1) * hs],
3831 );
3832 }
3833 attn.copy_from_slice(&normed);
3834 }
3835 for (dst, &a) in h.iter_mut().zip(&attn) {
3836 *dst += a;
3837 }
3838 }
3839 AttnKind::Linear(w) => {
3840 for bi in 0..b {
3841 let normed = inference::rms_norm(
3842 &h[bi * hs..(bi + 1) * hs],
3843 &lw.input_norm,
3844 eps,
3845 norm_style,
3846 );
3847 vmf_phase_forward(
3848 &normed,
3849 w,
3850 &self.vmf_cfg.expect("linear layer without vmf_cfg"),
3851 &mut self.kv_cache.layers[li].linear_state,
3852 pool.as_deref(),
3853 )
3854 .iter()
3855 .enumerate()
3856 .for_each(|(i, &a)| h[bi * hs + i] += a);
3857 }
3858 }
3859 }
3860
3861 let lw = &self.weights.layers[self.phys_layer(li)];
3863 let mut post = vec![0.0f32; b * hs];
3864 for bi in 0..b {
3865 let r =
3866 inference::rms_norm(&h[bi * hs..(bi + 1) * hs], &lw.post_norm, eps, norm_style);
3867 post[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3868 }
3869 let mask_row = task_mask
3872 .filter(|m| m.ffn_active_count(li) < self.intermediate_size)
3873 .and_then(|m| m.ffn_masks.get(li))
3874 .map(|v| v.as_slice());
3875 let mut ffn = match &lw.ffn {
3876 FfnKind::Dense(d) => dense_ffn_batch(d, &post, b, pool.as_deref(), mask_row),
3877 FfnKind::Moe(m) => moe_ffn_batch(m, &post, b, hs, pool.as_deref(), None),
3878 FfnKind::DenseMoe(dm) => {
3881 let mut out = vec![0.0f32; b * hs];
3882 for bi in 0..b {
3883 let r = dense_moe_ffn(
3884 dm,
3885 &post[bi * hs..(bi + 1) * hs],
3886 &h[bi * hs..(bi + 1) * hs],
3887 eps,
3888 norm_style,
3889 pool.as_deref(),
3890 );
3891 out[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3892 }
3893 out
3894 }
3895 };
3896 if let Some(w) = &lw.ffn_out_norm {
3897 for bi in 0..b {
3898 inference::rms_norm_into(
3899 &ffn[bi * hs..(bi + 1) * hs],
3900 w,
3901 eps,
3902 norm_style,
3903 &mut post[bi * hs..(bi + 1) * hs],
3904 );
3905 }
3906 ffn.copy_from_slice(&post);
3907 }
3908 for (dst, &f) in h.iter_mut().zip(&ffn) {
3909 *dst += f;
3910 }
3911 if let Some(sc) = lw.layer_scale {
3912 for v in h.iter_mut() {
3913 *v *= sc;
3914 }
3915 }
3916 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
3917 if let Some(t) = tp.parse::<usize>().ok() {
3918 if t >= start_pos && t < start_pos + b {
3919 let bi = t - start_pos;
3920 let row = &h[bi * hs..(bi + 1) * hs];
3921 let n: f32 = row.iter().map(|x| x * x).sum::<f32>().sqrt();
3922 eprintln!(
3923 "BATCH pos {t} after layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
3924 row[0], row[1]
3925 );
3926 }
3927 }
3928 }
3929 if std::env::var("CMF_DEBUG_LAYERS").is_ok() {
3933 let row = &h[(b - 1) * hs..b * hs];
3934 let rms =
3935 (row.iter().map(|&v| (v as f64) * (v as f64)).sum::<f64>() / hs as f64).sqrt();
3936 let mx = row.iter().fold(0f32, |m, &v| m.max(v.abs()));
3937 eprintln!(
3938 "layer {li:>3} {:>10} ffn={:<5} rms={rms:>12.4} max={mx:>12.4}",
3939 match &self.weights.layers[self.phys_layer(li)].attn {
3940 AttnKind::LinearGdn(_) => "gdn",
3941 AttnKind::Linear(_) => "vmf",
3942 AttnKind::ShortConv(_) => "conv",
3943 _ => "attn",
3944 },
3945 match &lw.ffn {
3946 FfnKind::Moe(_) => "moe",
3947 FfnKind::Dense(_) => "dense",
3948 FfnKind::DenseMoe(_) => "dense+moe",
3949 },
3950 );
3951 }
3952 if self.is_loop_end(li) && li + 1 < self.num_layers {
3954 for bi in 0..b {
3955 let normed = inference::rms_norm(
3956 &h[bi * hs..(bi + 1) * hs],
3957 &self.weights.final_norm,
3958 eps,
3959 norm_style,
3960 );
3961 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
3962 }
3963 }
3964 if std::env::var("CMF_TRACE_H").is_ok() {
3965 let n = h[..hs].iter().map(|v| v.abs()).sum::<f32>() / hs as f32;
3966 let mx = h[..hs].iter().fold(0.0f32, |a, &v| a.max(v.abs()));
3967 eprintln!(
3968 "layer {li}: mean|h|={n:.4} max|h|={mx:.2} scale={:?}",
3969 lw.layer_scale
3970 );
3971 }
3972 }
3973 crate::gpu::set_layer(-1); h
3975 }
3976
3977 fn embed_single(&self, id: u32) -> Vec<f32> {
3979 let mut out = vec![0.0f32; self.hidden_size];
3980 if (id as usize) < self.weights.embed_tokens.rows() {
3981 self.weights.embed_tokens.row_f32(id as usize, &mut out);
3982 }
3983 if self.embed_multiplier != 1.0 {
3984 for v in out.iter_mut() {
3985 *v *= self.embed_multiplier;
3986 }
3987 }
3988 if self.dsv4.is_some() {
3992 let mut v = vec![0.0f32; self.hidden_size.max(1)];
3993 v[0] = id as f32;
3994 return v;
3995 }
3996 if let Some(b) = &self.g3n {
3999 return b.0.extend_embedding(id, &out, self.pool.as_deref());
4000 }
4001 out
4002 }
4003
4004 #[cfg(target_os = "macos")]
4010 fn chunk_run_gpu(
4011 &mut self,
4012 li0: usize,
4013 h: &mut [f32],
4014 b: usize,
4015 pos0: usize,
4016 embed_ids: Option<&[u32]>,
4017 cap: usize,
4018 ) -> usize {
4019 if !crate::gpu::enabled_here()
4023 || std::env::var("CMF_GPU_CHUNK")
4024 .map(|v| v == "0")
4025 .unwrap_or(false)
4026 || b < 32
4027 || self.swa.is_some()
4028 || self.global_attn.is_some()
4029 || self.attn_v_norm
4030 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
4031 {
4032 return li0;
4033 }
4034 let Some(model) = self.model.clone() else {
4035 return li0;
4036 };
4037 let inv_freq = self.inv_freq.clone();
4038 let (nh, nkv, hd, hs) = (
4039 self.num_heads,
4040 self.num_kv_heads,
4041 self.head_dim,
4042 self.hidden_size,
4043 );
4044 let loop_end = if self.loop_final_norm {
4048 ((li0 / self.physical_layers) + 1) * self.physical_layers
4049 } else {
4050 self.num_layers
4051 };
4052 let mut layers: Vec<crate::gpu_metal::ChunkLayer> = Vec::new();
4053 let mut stored_at: Vec<usize> = Vec::new();
4054 for li in li0..self.num_layers.min(loop_end).min(cap) {
4055 let lw = &self.weights.layers[self.phys_layer(li)];
4056 if lw.attn_out_norm.is_some() || lw.ffn_out_norm.is_some() || lw.layer_scale.is_some() {
4057 break;
4058 }
4059 let AttnKind::Full {
4060 wq,
4061 wk,
4062 wv,
4063 wo,
4064 q_norm,
4065 k_norm,
4066 output_gate: false,
4067 softplus_gate: None,
4068 bias,
4069 } = &lw.attn
4070 else {
4071 break;
4072 };
4073 let FfnKind::Dense(d) = &lw.ffn else { break };
4074 if d.act != Act::Silu {
4075 break;
4076 }
4077 fn cw(t: &QTensor) -> Option<(usize, usize, usize, &[f32])> {
4082 t.q8_row_parts()
4083 .or_else(|| t.q4t_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4084 .or_else(|| t.q4tp_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4085 }
4086 let parts = (
4087 cw(wq),
4088 cw(wk),
4089 cw(wv),
4090 cw(wo),
4091 cw(&d.gate_proj),
4092 cw(&d.up_proj),
4093 cw(&d.down_proj),
4094 );
4095 let (Some(pq), Some(pk), Some(pv), Some(po), Some(pg), Some(pu), Some(pd)) = parts
4096 else {
4097 break;
4098 };
4099 let layer = &self.kv_cache.layers[li];
4100 if layer.mode != crate::kv_cache::KvMode::F32 || layer.o1.is_some() {
4101 break;
4102 }
4103 stored_at.push(layer.head_len(0));
4104 layers.push(crate::gpu_metal::ChunkLayer {
4105 model: &model,
4106 kv_id: self.graph_kv_id,
4107 layer: li,
4108 wq: pq,
4109 wk: pk,
4110 wv: pv,
4111 wo: po,
4112 gate: pg,
4113 up: pu,
4114 down: pd,
4115 input_norm: &lw.input_norm,
4116 post_norm: &lw.post_norm,
4117 bias: bias
4118 .as_ref()
4119 .map(|(a, bb, cc)| (a.as_slice(), bb.as_slice(), cc.as_slice())),
4120 q_norm: q_norm.as_deref(),
4121 k_norm: k_norm.as_deref(),
4122 inv_freq: &inv_freq,
4123 rd: self.rotary_dim,
4124 nh,
4125 nkv,
4126 hd,
4127 hs,
4128 inter: d.gate_proj.rows(),
4129 gemma: matches!(self.norm_style, cortiq_core::NormStyle::Gemma),
4130 eps: self.rms_eps as f32,
4131 });
4132 }
4133 if layers.is_empty() {
4134 return li0;
4135 }
4136 let row = nkv * hd;
4137 let mut store: Vec<(Vec<f32>, Vec<f32>, Vec<f32>)> = stored_at
4138 .iter()
4139 .map(|&st| (vec![0f32; b * row], vec![0f32; b * row], vec![0f32; st + b]))
4140 .collect();
4141 let mut io: Vec<crate::gpu_metal::ChunkIo> = Vec::with_capacity(layers.len());
4142 for (i, (ok, ov, oi)) in store.iter_mut().enumerate() {
4143 let li = layers[i].layer;
4144 let layer = &self.kv_cache.layers[li];
4145 io.push(crate::gpu_metal::ChunkIo {
4146 cpu_stored: stored_at[i],
4147 cpu_k: (0..nkv).map(|g| layer.head_keys(g)).collect(),
4148 cpu_v: (0..nkv).map(|g| layer.head_values(g)).collect(),
4149 out_k: ok,
4150 out_v: ov,
4151 imp: oi,
4152 });
4153 }
4154 let n_run = layers.len();
4155 let last = layers.last().map(|l| l.layer + 1).unwrap_or(li0);
4156 let ep = embed_ids.and_then(|ids| {
4159 self.weights
4160 .embed_tokens
4161 .q8_row_parts()
4162 .map(|(idx, rows, _c, rs)| crate::gpu_metal::ChunkEmbed {
4163 idx,
4164 rows,
4165 row_scale: rs,
4166 ids,
4167 mult: self.embed_multiplier,
4168 })
4169 });
4170 if embed_ids.is_some() && ep.is_none() {
4171 return li0;
4172 }
4173 if !crate::gpu_metal::chunk_run_gpu(&layers, &mut io, h, b, pos0, ep.as_ref()) {
4174 return li0;
4175 }
4176 drop(io);
4177 drop(layers);
4178 for (i, (ok, ov, oi)) in store.iter().enumerate().take(n_run) {
4181 let li = li0 + i;
4182 let layer = &mut self.kv_cache.layers[li];
4183 for bi in 0..b {
4184 layer.append(
4185 &ok[bi * row..(bi + 1) * row],
4186 &ov[bi * row..(bi + 1) * row],
4187 &[],
4188 );
4189 }
4190 layer.accumulate_imp(oi);
4191 }
4192 last
4193 }
4194
4195 fn layer_is_local(&self, li: usize) -> bool {
4198 if let Some(layers) = &self.sliding_layers {
4199 return layers.get(li).copied().unwrap_or(false);
4200 }
4201 match self.swa {
4202 Some((_, pattern)) => (li + 1) % pattern.max(1) != 0,
4203 None => false,
4204 }
4205 }
4206
4207 fn layer_inv_freq(&self, li: usize) -> std::sync::Arc<Vec<f32>> {
4210 if self.layer_is_local(li) {
4211 if let Some(f) = &self.inv_freq_local {
4212 return f.clone();
4213 }
4214 } else if let Some(f) = &self.inv_freq_global {
4215 return f.clone();
4216 }
4217 self.inv_freq.clone()
4218 }
4219
4220 fn layer_window(&self, li: usize) -> Option<usize> {
4222 self.swa
4223 .and_then(|(w, _)| self.layer_is_local(li).then_some(w))
4224 }
4225
4226 fn layer_num_heads(&self, li: usize) -> usize {
4227 self.attention_heads_per_layer
4228 .as_ref()
4229 .and_then(|v| v.get(li).copied())
4230 .unwrap_or(self.num_heads)
4231 }
4232
4233 fn layer_rope_scale(&self, li: usize) -> f32 {
4234 if self.layer_is_local(li) {
4235 self.rope_scale_local
4236 } else {
4237 self.rope_scale
4238 }
4239 }
4240
4241 fn layer_geom(&self, li: usize) -> (usize, usize, usize) {
4244 if !self.layer_is_local(li) {
4245 if let Some((ghd, gkv)) = self.global_attn {
4246 return (gkv, ghd, ghd);
4247 }
4248 }
4249 (
4250 self.num_kv_heads,
4251 self.head_dim,
4252 if self.layer_is_local(li) {
4253 self.rotary_dim_local.unwrap_or(self.rotary_dim)
4254 } else {
4255 self.rotary_dim
4256 },
4257 )
4258 }
4259
4260 fn forward_layers(
4262 &mut self,
4263 hidden: &[f32],
4264 position: usize,
4265 task_mask: Option<&TaskMask>,
4266 ) -> Vec<f32> {
4267 self.forward_layers_upto(hidden, position, task_mask, None)
4268 }
4269
4270 pub fn embed_id(&self, id: u32) -> Vec<f32> {
4278 self.embed_single(id)
4279 }
4280
4281 pub fn split_supported(&self) -> Result<(), String> {
4285 if self.dsv4.is_some() {
4286 return Err(
4287 "network split: DeepSeek-V4 runs its own fused stack (not splittable yet)".into(),
4288 );
4289 }
4290 if self.g3n.is_some() {
4291 return Err(
4292 "network split: Gemma-3n runs its own AltUp stack (not splittable yet)".into(),
4293 );
4294 }
4295 Ok(())
4296 }
4297
4298 pub fn forward_span(
4303 &mut self,
4304 hidden: &[f32],
4305 position: usize,
4306 from: usize,
4307 upto: usize,
4308 task_mask: Option<&TaskMask>,
4309 ) -> Result<Vec<f32>, String> {
4310 self.split_supported()?;
4311 if from > upto || upto >= self.num_layers {
4312 return Err(format!(
4313 "forward_span: layer range {from}..={upto} outside 0..{}",
4314 self.num_layers
4315 ));
4316 }
4317 if hidden.len() != self.hidden_size {
4318 return Err(format!(
4319 "forward_span: hidden len {} ≠ hidden_size {}",
4320 hidden.len(),
4321 self.hidden_size
4322 ));
4323 }
4324 Ok(self.forward_layers_span(hidden, position, task_mask, from, Some(upto)))
4325 }
4326
4327 pub fn logits_from_hidden(&mut self, hidden: &[f32]) -> Vec<f32> {
4330 let normed = inference::rms_norm(
4331 hidden,
4332 &self.weights.final_norm,
4333 self.rms_eps,
4334 self.norm_style,
4335 );
4336 self.lm_head_forward(&normed)
4337 }
4338
4339 pub fn sample_next(&mut self, logits: &[f32], past_tokens: &[u32]) -> u32 {
4341 sampler::sample_with_scratch(
4342 logits,
4343 &self.sampler_config,
4344 past_tokens,
4345 &mut self.rng,
4346 &mut self.sampler_scratch,
4347 )
4348 }
4349
4350 pub fn reset_session(&mut self) {
4352 self.kv_cache.clear();
4353 self.kv_history.clear();
4354 crate::gpu::graph_kv_reset(self.graph_kv_id);
4355 }
4356
4357 pub fn prefill_span_ids(
4363 &mut self,
4364 ids: &[u32],
4365 start_pos: usize,
4366 upto: usize,
4367 task_mask: Option<&TaskMask>,
4368 ) -> Result<Vec<f32>, String> {
4369 self.split_supported()?;
4370 if upto >= self.num_layers {
4371 return Err(format!(
4372 "prefill_span_ids: upto {upto} outside 0..{}",
4373 self.num_layers
4374 ));
4375 }
4376 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4380 Ok(self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, upto + 1))
4381 } else {
4382 let hs = self.hidden_size;
4383 let mut out = Vec::with_capacity(ids.len() * hs);
4384 for (i, &id) in ids.iter().enumerate() {
4385 let emb = self.embed_id(id);
4386 out.extend_from_slice(&self.forward_span(
4387 &emb,
4388 start_pos + i,
4389 0,
4390 upto,
4391 task_mask,
4392 )?);
4393 }
4394 Ok(out)
4395 }
4396 }
4397
4398 pub fn prefill_span_hidden(
4401 &mut self,
4402 hidden: &[f32],
4403 start_pos: usize,
4404 from: usize,
4405 upto: usize,
4406 task_mask: Option<&TaskMask>,
4407 ) -> Result<Vec<f32>, String> {
4408 self.split_supported()?;
4409 let hs = self.hidden_size;
4410 if hidden.is_empty() || hidden.len() % hs != 0 {
4411 return Err(format!(
4412 "prefill_span_hidden: {} floats is not a multiple of hidden {hs}",
4413 hidden.len()
4414 ));
4415 }
4416 if from > upto || upto >= self.num_layers {
4417 return Err(format!(
4418 "prefill_span_hidden: layer range {from}..={upto} outside 0..{}",
4419 self.num_layers
4420 ));
4421 }
4422 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4423 Ok(self.prefill_batch_span(
4424 PrefillIn::Hidden(hidden),
4425 start_pos,
4426 task_mask,
4427 from,
4428 upto + 1,
4429 ))
4430 } else {
4431 let b = hidden.len() / hs;
4432 let mut out = Vec::with_capacity(hidden.len());
4433 for i in 0..b {
4434 let h = self.forward_span(
4435 &hidden[i * hs..(i + 1) * hs],
4436 start_pos + i,
4437 from,
4438 upto,
4439 task_mask,
4440 )?;
4441 out.extend_from_slice(&h);
4442 }
4443 Ok(out)
4444 }
4445 }
4446
4447 fn try_token_graph_wgpu(
4451 &self,
4452 hidden: &[f32],
4453 position: usize,
4454 logits_out: &mut Vec<f32>,
4455 layers_run: &mut usize,
4456 ) -> Option<Vec<f32>> {
4457 self.try_token_graph_wgpu_steps(
4458 hidden,
4459 position,
4460 logits_out,
4461 1,
4462 None,
4463 Some(layers_run),
4464 0,
4465 self.num_layers,
4466 )
4467 }
4468
4469 fn try_token_graph_wgpu_span(
4473 &self,
4474 hidden: &[f32],
4475 position: usize,
4476 logits_out: &mut Vec<f32>,
4477 from: usize,
4478 upto_excl: usize,
4479 layers_run: &mut usize,
4480 ) -> Option<Vec<f32>> {
4481 self.try_token_graph_wgpu_steps(
4482 hidden,
4483 position,
4484 logits_out,
4485 1,
4486 None,
4487 Some(layers_run),
4488 from,
4489 upto_excl,
4490 )
4491 }
4492
4493 fn try_multi_burst(&self, t_next: u32, position: usize, k: usize) -> Option<Vec<u32>> {
4497 if self.o1_active() || self.attn_softcap > 0.0 {
4498 return None;
4499 }
4500 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
4501 if !graph_on || crate::gpu::graph_unsupported() {
4502 return None;
4509 }
4510 let emb = self.embed_single(t_next);
4511 let mut lg = Vec::new();
4512 let mut ids = Vec::new();
4513 self.try_token_graph_wgpu_steps(
4514 &emb,
4515 position,
4516 &mut lg,
4517 k,
4518 Some(&mut ids),
4519 None,
4520 0,
4521 self.num_layers,
4522 )?;
4523 (ids.len() == k).then_some(ids)
4524 }
4525
4526 fn try_token_graph_wgpu_steps(
4530 &self,
4531 hidden: &[f32],
4532 position: usize,
4533 logits_out: &mut Vec<f32>,
4534 steps: usize,
4535 ids_out: Option<&mut Vec<u32>>,
4536 layers_run: Option<&mut usize>,
4537 from: usize,
4538 upto_excl: usize,
4539 ) -> Option<Vec<f32>> {
4540 let o1_gpu = std::env::var("CMF_O1_GPU").as_deref() == Ok("1");
4543 if (self.o1_active() && !o1_gpu) || self.attn_softcap > 0.0 {
4544 return None;
4548 }
4549 let o1_views: Vec<Option<Vec<crate::nystrom::O1DeviceView<'_>>>> = (from..upto_excl)
4554 .map(|li| {
4555 if !o1_gpu {
4556 return None;
4557 }
4558 self.kv_cache.layers[self.phys_layer(li)].o1_views()
4559 })
4560 .collect();
4561 if self.o1_active() && o1_gpu {
4562 let want: usize = (from..upto_excl)
4565 .filter(|li| !matches!(self.kv_cache.layers[self.phys_layer(*li)].o1, None))
4566 .count();
4567 let have = o1_views.iter().filter(|v| v.is_some()).count();
4568 if want == 0 || have != want {
4569 return None;
4570 }
4571 }
4572 let nh = self.num_heads;
4573 let (nkv, hd, rd) = self.layer_geom(0);
4574 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
4575 let mut layers = Vec::with_capacity(upto_excl - from);
4576 let mut model = None;
4577 let dbg = std::env::var("CMF_GRAPH_DEBUG").is_ok();
4578 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
4579 if let Some((_, i, kind, rs)) = t.graph_weight() {
4580 return Some(crate::gpu::GraphW {
4581 idx: i,
4582 kind,
4583 row_scale: rs,
4584 data: &[],
4585 });
4586 }
4587 t.as_f32().map(|d| crate::gpu::GraphW {
4589 idx: 0,
4590 kind: 4,
4591 row_scale: &[],
4592 data: d,
4593 })
4594 }
4595 for li in from..upto_excl {
4596 let lw = &self.weights.layers[self.phys_layer(li)];
4597 if dbg {
4598 let ak = match &lw.attn {
4599 AttnKind::Mla(_) => "Mla".into(),
4600 AttnKind::Full {
4601 output_gate, bias, ..
4602 } => format!("Full gate={output_gate} bias={}", bias.is_some()),
4603 AttnKind::LinearGdn(_) => "LinearGdn".into(),
4604 AttnKind::Kda(_) => "Kda".into(),
4605 AttnKind::Linear(_) => "Linear".into(),
4606 AttnKind::ShortConv(_) => "ShortConv".into(),
4607 };
4608 let fk = match &lw.ffn {
4609 FfnKind::Dense(_) => "Dense",
4610 FfnKind::Moe(_) => "Moe",
4611 FfnKind::DenseMoe(_) => "DenseMoe",
4612 };
4613 eprintln!("graph L{li}: attn={ak} ffn={fk}");
4614 }
4615 let gffn = match &lw.ffn {
4616 FfnKind::DenseMoe(_) => return None, FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
4618 gate: gw(&d.gate_proj)?,
4619 up: gw(&d.up_proj)?,
4620 down: gw(&d.down_proj)?,
4621 },
4622 FfnKind::Moe(m) => {
4623 if m.router_sigmoid
4628 || m.expert_bias.is_some()
4629 || m.route_tau.is_some()
4630 || m.mask.is_some()
4631 {
4632 return None;
4633 }
4634 let (se, sg) = m.shared.as_ref()?;
4635 let sgate = gw(sg.as_ref()?)?;
4636 let router = gw(&m.router)?;
4637 let inter = m.experts.first()?.gate_proj.rows();
4638 let mut experts = Vec::with_capacity(m.experts.len() + 1);
4639 let mut q4tp: Option<bool> = None;
4642 let mut gu_q2: Option<bool> = None;
4645 for e in m.experts.iter().chain(std::iter::once(se)) {
4646 if !matches!(e.act, Act::Silu)
4647 || e.gate_proj.rows() != inter
4648 || e.up_proj.rows() != inter
4649 {
4650 return None;
4651 }
4652 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
4653 Some((mm, gi)) => (
4654 mm,
4655 gi,
4656 e.up_proj.mapped_q4t()?.1,
4657 e.down_proj.mapped_q4t()?.1,
4658 false,
4659 false,
4660 ),
4661 None => match e.gate_proj.mapped_q2tp() {
4662 Some((mm, gi)) => (
4663 mm,
4664 gi,
4665 e.up_proj.mapped_q2tp()?.1,
4666 e.down_proj.mapped_q4tp()?.1,
4667 true,
4668 true,
4669 ),
4670 None => {
4671 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
4672 (
4673 mm,
4674 gi,
4675 e.up_proj.mapped_q4tp()?.1,
4676 e.down_proj.mapped_q4tp()?.1,
4677 true,
4678 false,
4679 )
4680 }
4681 },
4682 };
4683 if *q4tp.get_or_insert(is_p) != is_p || *gu_q2.get_or_insert(is_q2) != is_q2
4684 {
4685 tracing::warn!(
4691 "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."
4692 );
4693 return None;
4694 }
4695 model.get_or_insert_with(|| mm.clone());
4696 experts.push((gi, ui, di));
4697 }
4698 crate::gpu::GraphFfn::Moe {
4699 router,
4700 shared_gate: sgate,
4701 experts,
4702 n_exp: m.experts.len(),
4703 top_k: std::env::var("CMF_TOPK_PROBE")
4709 .ok()
4710 .and_then(|v| v.parse::<usize>().ok())
4711 .filter(|k| *k > 0 && *k <= m.top_k)
4712 .unwrap_or(m.top_k),
4713 inter,
4714 norm_topk: m.norm_topk_prob,
4715 q4tp: q4tp?,
4716 gu_q2: gu_q2.unwrap_or(false),
4717 }
4718 }
4719 };
4720 let attn = match &lw.attn {
4721 AttnKind::Full {
4722 wq,
4723 wk,
4724 wv,
4725 wo,
4726 q_norm,
4727 k_norm,
4728 output_gate,
4729 softplus_gate,
4730 bias,
4731 } => {
4732 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
4733 return None;
4734 }
4735 let (m, _, _, _) = wq.graph_weight()?;
4736 model = Some(m.clone());
4737 crate::gpu::GraphAttn::Full {
4738 wq: gw(wq)?,
4739 wk: gw(wk)?,
4740 wv: gw(wv)?,
4741 wo: gw(wo)?,
4742 q_norm: q_norm.as_deref(),
4743 k_norm: k_norm.as_deref(),
4744 bias: bias
4745 .as_ref()
4746 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
4747 output_gate: *output_gate,
4748 cpu_k: self.kv_cache.layers[li].k_heads(),
4749 cpu_v: self.kv_cache.layers[li].v_heads(),
4750 }
4751 }
4752 AttnKind::LinearGdn(w) => {
4753 let cfg = self.gdn_cfg?;
4754 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
4755 model = Some(m.clone());
4756 crate::gpu::GraphAttn::Gdn {
4757 qkv: gw(&w.in_proj_qkv)?,
4758 z: gw(&w.in_proj_z)?,
4759 a: gw(&w.in_proj_a)?,
4760 b: gw(&w.in_proj_b)?,
4761 out: gw(&w.out_proj)?,
4762 conv1d: &w.conv1d,
4763 a_log: &w.a_log,
4764 dt_bias: &w.dt_bias,
4765 norm: &w.norm,
4766 nv: cfg.num_v_heads,
4767 nk: cfg.num_k_heads,
4768 dk: cfg.key_head_dim,
4769 dv: cfg.value_head_dim,
4770 kk: cfg.conv_kernel,
4771 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
4772 }
4773 }
4774 _ => return None,
4775 };
4776 layers.push(crate::gpu::GraphLayer {
4777 input_norm: &lw.input_norm,
4778 attn,
4779 post_norm: &lw.post_norm,
4780 ffn: gffn,
4781 });
4782 }
4783 let model = model?;
4784 let lm_gw = if upto_excl == self.num_layers
4790 && self.graph_want_logits
4791 && std::env::var("CMF_GPU_LMHEAD")
4792 .map(|v| v != "0")
4793 .unwrap_or(true)
4794 {
4795 self.weights.lm_head.graph_weight().map(|(_, i, kind, rs)| {
4796 (
4797 crate::gpu::GraphW {
4798 idx: i,
4799 kind,
4800 row_scale: rs,
4801 data: &[],
4802 },
4803 self.weights.lm_head.rows(),
4804 )
4805 })
4806 } else {
4807 None
4808 };
4809 let lm = lm_gw.as_ref().map(|(gw, rows)| (gw, *rows));
4810 let emb_gw = if steps > 1 {
4812 self.weights
4813 .embed_tokens
4814 .graph_weight()
4815 .map(|(_, i, kind, rs)| {
4816 (
4817 crate::gpu::GraphW {
4818 idx: i,
4819 kind,
4820 row_scale: rs,
4821 data: &[],
4822 },
4823 self.weights.embed_tokens.rows(),
4824 self.embed_multiplier as f32,
4825 )
4826 })
4827 } else {
4828 None
4829 };
4830
4831 let loop_norm_at: Vec<usize> = if self.loop_final_norm {
4837 (from..upto_excl.min(self.num_layers - 1))
4838 .filter(|&li| (li + 1) % self.physical_layers == 0)
4839 .map(|li| li - from)
4840 .collect()
4841 } else {
4842 Vec::new()
4843 };
4844 let mut h = hidden.to_vec();
4845 crate::gpu::forward_token_graph(
4846 &model,
4847 self.graph_kv_id,
4848 &layers,
4849 &o1_views,
4850 self.o1_epoch,
4851 &self.inv_freq,
4852 &mut h,
4853 nh,
4854 nkv,
4855 hd,
4856 rd,
4857 self.hidden_size,
4858 self.intermediate_size,
4859 position,
4860 self.kv_cache.max_seq_len,
4861 gemma,
4862 self.rms_eps as f32,
4863 lm,
4864 &self.weights.final_norm,
4865 logits_out,
4866 &loop_norm_at,
4867 steps,
4868 emb_gw.as_ref().map(|(gw, rows, m)| (gw, *rows, *m)),
4869 ids_out,
4870 layers_run,
4871 from,
4872 )
4873 .then_some(h)
4874 }
4875
4876 fn try_batch_graph_wgpu(
4881 &self,
4882 hiddens: &mut [f32],
4883 positions: &[usize],
4884 k: usize,
4885 spec: Option<crate::gpu::SpecTail<'_>>,
4886 ) -> bool {
4887 let _tb = std::time::Instant::now();
4888 if self.attn_softcap > 0.0 {
4889 return false; }
4891 if self.o1_active() {
4892 return false;
4893 }
4894 let nh = self.num_heads;
4895 let (nkv, hd, rd) = self.layer_geom(0);
4896 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
4897 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
4898 if let Some((_, i, kind, rs)) = t.graph_weight() {
4899 return Some(crate::gpu::GraphW {
4900 idx: i,
4901 kind,
4902 row_scale: rs,
4903 data: &[],
4904 });
4905 }
4906 t.as_f32().map(|d| crate::gpu::GraphW {
4907 idx: 0,
4908 kind: 4,
4909 row_scale: &[],
4910 data: d,
4911 })
4912 }
4913 let built: Option<(
4914 Vec<crate::gpu::GraphLayer<'_>>,
4915 std::sync::Arc<cortiq_core::CmfModel>,
4916 )> = (|| {
4917 let mut layers = Vec::with_capacity(self.num_layers);
4918 let mut model = None;
4919 for li in 0..self.num_layers {
4920 let lw = &self.weights.layers[self.phys_layer(li)];
4921 let gffn = match &lw.ffn {
4928 FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
4929 gate: gw(&d.gate_proj)?,
4930 up: gw(&d.up_proj)?,
4931 down: gw(&d.down_proj)?,
4932 },
4933 FfnKind::Moe(m) => {
4934 if m.router_sigmoid
4935 || m.expert_bias.is_some()
4936 || m.route_tau.is_some()
4937 || m.mask.is_some()
4938 {
4939 return None;
4940 }
4941 let (se, sg) = m.shared.as_ref()?;
4942 let sgate = gw(sg.as_ref()?)?;
4943 let router = gw(&m.router)?;
4944 let inter = m.experts.first()?.gate_proj.rows();
4945 let mut experts = Vec::with_capacity(m.experts.len() + 1);
4946 let mut q4tp: Option<bool> = None;
4947 let mut gu_q2: Option<bool> = None;
4948 for e in m.experts.iter().chain(std::iter::once(se)) {
4949 if !matches!(e.act, Act::Silu)
4950 || e.gate_proj.rows() != inter
4951 || e.up_proj.rows() != inter
4952 {
4953 return None;
4954 }
4955 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
4959 Some((mm, gi)) => (
4960 mm,
4961 gi,
4962 e.up_proj.mapped_q4t()?.1,
4963 e.down_proj.mapped_q4t()?.1,
4964 false,
4965 false,
4966 ),
4967 None => match e.gate_proj.mapped_q2tp() {
4968 Some((mm, gi)) => (
4969 mm,
4970 gi,
4971 e.up_proj.mapped_q2tp()?.1,
4972 e.down_proj.mapped_q4tp()?.1,
4973 true,
4974 true,
4975 ),
4976 None => {
4977 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
4978 (
4979 mm,
4980 gi,
4981 e.up_proj.mapped_q4tp()?.1,
4982 e.down_proj.mapped_q4tp()?.1,
4983 true,
4984 false,
4985 )
4986 }
4987 },
4988 };
4989 if *q4tp.get_or_insert(is_p) != is_p
4990 || *gu_q2.get_or_insert(is_q2) != is_q2
4991 {
4992 return None;
4993 }
4994 model.get_or_insert_with(|| mm.clone());
4995 experts.push((gi, ui, di));
4996 }
4997 crate::gpu::GraphFfn::Moe {
4998 router,
4999 shared_gate: sgate,
5000 experts,
5001 n_exp: m.experts.len(),
5002 top_k: m.top_k,
5003 inter,
5004 norm_topk: m.norm_topk_prob,
5005 q4tp: q4tp?,
5006 gu_q2: gu_q2.unwrap_or(false),
5007 }
5008 }
5009 _ => return None,
5010 };
5011 let attn = match &lw.attn {
5012 AttnKind::Full {
5013 wq,
5014 wk,
5015 wv,
5016 wo,
5017 q_norm,
5018 k_norm,
5019 output_gate,
5020 softplus_gate,
5021 bias,
5022 } => {
5023 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
5024 return None;
5025 }
5026 let (m, _, _, _) = wq.graph_weight()?;
5027 model = Some(m.clone());
5028 crate::gpu::GraphAttn::Full {
5029 wq: gw(wq)?,
5030 wk: gw(wk)?,
5031 wv: gw(wv)?,
5032 wo: gw(wo)?,
5033 q_norm: q_norm.as_deref(),
5034 k_norm: k_norm.as_deref(),
5035 bias: bias
5036 .as_ref()
5037 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
5038 output_gate: *output_gate,
5039 cpu_k: self.kv_cache.layers[li].k_heads(),
5040 cpu_v: self.kv_cache.layers[li].v_heads(),
5041 }
5042 }
5043 AttnKind::LinearGdn(w) => {
5044 let cfg = self.gdn_cfg?;
5045 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
5046 model = Some(m.clone());
5047 crate::gpu::GraphAttn::Gdn {
5048 qkv: gw(&w.in_proj_qkv)?,
5049 z: gw(&w.in_proj_z)?,
5050 a: gw(&w.in_proj_a)?,
5051 b: gw(&w.in_proj_b)?,
5052 out: gw(&w.out_proj)?,
5053 conv1d: &w.conv1d,
5054 a_log: &w.a_log,
5055 dt_bias: &w.dt_bias,
5056 norm: &w.norm,
5057 nv: cfg.num_v_heads,
5058 nk: cfg.num_k_heads,
5059 dk: cfg.key_head_dim,
5060 dv: cfg.value_head_dim,
5061 kk: cfg.conv_kernel,
5062 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
5063 }
5064 }
5065 _ => return None,
5066 };
5067 layers.push(crate::gpu::GraphLayer {
5068 input_norm: &lw.input_norm,
5069 attn,
5070 post_norm: &lw.post_norm,
5071 ffn: gffn,
5072 });
5073 }
5074 Some((layers, model?))
5075 })();
5076 let Some((layers, model)) = built else {
5077 {
5078 use std::sync::atomic::{AtomicBool, Ordering};
5079 static SAID: AtomicBool = AtomicBool::new(false);
5080 if !SAID.swap(true, Ordering::Relaxed) {
5081 tracing::warn!("batch graph: BUILDER refused (layer weights/kinds)");
5082 }
5083 }
5084 return false;
5085 };
5086 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
5087 eprintln!("batch-build: {:.1} ms", _tb.elapsed().as_secs_f64() * 1e3);
5088 }
5089 crate::gpu::forward_batch_graph(
5090 &model,
5091 self.graph_kv_id,
5092 &layers,
5093 &self.inv_freq,
5094 hiddens,
5095 nh,
5096 nkv,
5097 hd,
5098 rd,
5099 self.hidden_size,
5100 self.intermediate_size,
5101 positions,
5102 self.kv_cache.max_seq_len,
5103 gemma,
5104 self.rms_eps as f32,
5105 k,
5106 spec,
5107 )
5108 }
5109
5110 fn draft_probe() -> bool {
5114 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5115 *ON.get_or_init(|| std::env::var("CMF_DSV4_DRAFT_PROBE").is_ok_and(|v| v != "0"))
5116}
5117
5118 #[cfg(feature = "gpu")]
5130 fn dsv4_spec_on() -> bool {
5131 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5132 *ON.get_or_init(|| std::env::var("CMF_DSV4_SPEC").map(|v| v != "0").unwrap_or(true))
5133 }
5134
5135 #[cfg(feature = "gpu")]
5142 fn dsv4_spec_step(
5143 &mut self,
5144 tip_token: u32,
5145 t_next: u32,
5146 next_pos: usize,
5147 drafted: &mut usize,
5148 accepted_ctr: &mut usize,
5149 ) -> Option<(Vec<u32>, usize)> {
5150 let t_all = std::time::Instant::now();
5151 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5152 thread_local! {
5153 static LAST: std::cell::Cell<Option<std::time::Instant>> =
5154 const { std::cell::Cell::new(None) };
5155 }
5156 LAST.with(|l| {
5157 if let Some(prev) = l.get() {
5158 eprintln!("между раундами {:.1} мс", prev.elapsed().as_secs_f64() * 1e3);
5159 }
5160 l.set(Some(std::time::Instant::now()));
5161 });
5162 }
5163 if std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5164 eprintln!("spec_step: вход pos={next_pos}");
5165 }
5166 let n_layers = self.dsv4.as_ref().map(|b| b.1.len())?;
5167 let cfg = self.dsv4.as_ref().map(|b| b.2)?;
5168 if self.dspark.is_none() {
5170 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5171 if t.is_empty() {
5172 return None;
5173 }
5174 crate::dsv4::dspark_arm(&t, cfg.dim);
5175 self.dspark = Some(crate::dsv4::DsparkState::new(
5176 self.dsv4_mtp.len(),
5177 &cfg,
5178 t.len(),
5179 ));
5180 }
5181 let targets = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5182 let pack = crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg);
5183 if pack.is_none() && std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5184 eprintln!("spec_step: пак не построился (targets {targets:?})");
5185 }
5186 let pack = pack?;
5187 let block = crate::dsv4::dspark_block();
5188 let b_box = self.dsv4.as_mut()?;
5189 let (g, layers, st) = (&b_box.0, &b_box.1, &mut b_box.3);
5190 let ds = self.dspark.as_mut()?;
5191 let dbg = std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok();
5194 if !crate::dsv4::dspark_take(&mut ds.main_hidden) && !ds.have_hidden {
5195 if dbg {
5196 eprintln!("spec_step: нет захвата");
5197 }
5198 return None;
5199 }
5200 ds.have_hidden = true;
5201 let tip_pos = next_pos.checked_sub(1)?;
5202 let draft_started = std::time::Instant::now();
5203 let mut conf = Vec::new();
5204 let props = crate::dsv4::dspark_draft_gpu(
5205 g,
5206 &self.dsv4_mtp,
5207 &cfg,
5208 ds,
5209 pack,
5210 st.kv_id,
5211 tip_token,
5212 tip_pos,
5213 self.pool.as_deref(),
5214 &mut conf,
5215 );
5216 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5217 *drafted += block;
5218 if props.is_empty() || props[0] != t_next {
5219 if dbg {
5220 eprintln!(
5221 "spec_step: черновик {} (props0={:?} t_next={t_next})",
5222 if props.is_empty() { "пуст" } else { "мимо" },
5223 props.first()
5224 );
5225 }
5226 return None;
5227 }
5228 let mut k_verify = crate::dsv4::dspark_verify_k().min(props.len());
5229 let conf_min = {
5235 static M: std::sync::OnceLock<f32> = std::sync::OnceLock::new();
5236 *M.get_or_init(|| {
5237 std::env::var("CMF_DSPARK_CONF_MIN")
5238 .ok()
5239 .and_then(|v| v.parse().ok())
5240 .unwrap_or(0.0)
5241 })
5242 };
5243 if conf_min > 0.0 && conf.len() >= props.len() {
5244 let mut keep = 1usize;
5245 while keep < k_verify && conf.get(keep).copied().unwrap_or(0.0) >= conf_min {
5246 keep += 1;
5247 }
5248 k_verify = k_verify.min(keep.max(2));
5249 }
5250 if k_verify < 2 {
5251 return None;
5252 }
5253 let mut fed = Vec::with_capacity(k_verify);
5254 fed.push(t_next);
5255 fed.extend_from_slice(&props[1..k_verify]);
5256 let mut argmax = Vec::new();
5257 let mut logits_all = Vec::new();
5258 let mut walked = Vec::new();
5259 let txn = crate::dsv4::dsv4_verify_chunk(
5260 g,
5261 layers,
5262 &cfg,
5263 st,
5264 &fed,
5265 next_pos,
5266 &self.inv_freq,
5267 self.pool.as_deref(),
5268 &targets,
5269 &mut argmax,
5270 &mut logits_all,
5271 &mut walked,
5272 );
5273 if txn.is_none() && dbg {
5274 eprintln!("spec_step: verify отказал");
5275 }
5276 let txn = txn?;
5277 let b = fed.len();
5278 let mut accepted = 1usize;
5279 while accepted < b && fed[accepted] == argmax[accepted - 1] {
5280 accepted += 1;
5281 }
5282 if std::env::var("CMF_DSV4_SPEC_FORCE_REJECT").is_ok_and(|v| v != "0") {
5287 accepted = 1;
5288 }
5289 if std::env::var("CMF_DSV4_SPEC_TRACE").is_ok() {
5290 eprintln!(
5291 "spec@{next_pos}: fed={fed:?} argmax={argmax:?} accepted={accepted}"
5292 );
5293 }
5294 let t_fin = std::time::Instant::now();
5295 if !crate::dsv4::dsv4_spec_finish(
5296 g,
5297 layers,
5298 &cfg,
5299 st,
5300 txn,
5301 accepted,
5302 &fed,
5303 &self.inv_freq,
5304 self.pool.as_deref(),
5305 ) {
5306 tracing::warn!("dsv4: спекулятивный откат не удался — состояние подозрительно");
5307 return None;
5308 }
5309 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5310 eprintln!("finish(k={accepted}): {:.1} мс", t_fin.elapsed().as_secs_f64() * 1e3);
5311 }
5312 *accepted_ctr += accepted - 1;
5313 let (hc, dim) = (cfg.hc_mult, cfg.dim);
5318 let dev_caps: Vec<usize> = targets
5325 .iter()
5326 .copied()
5327 .filter(|&t| {
5328 st.dev_set.get(t).copied().unwrap_or(false)
5329 && !st.partial_set.get(t).copied().unwrap_or(false)
5330 })
5331 .collect();
5332 let mut caps_all = vec![0.0f32; dev_caps.len() * b * hc * dim];
5333 if !crate::gpu_wgpu::dsv4_spec_cap_read_all(b, dev_caps.len(), hc * dim, &mut caps_all) {
5334 return None;
5335 }
5336 for t in 0..accepted {
5337 let tip = t + 1 == accepted;
5338 for (slot, &tl) in targets.iter().enumerate() {
5339 if let Some(di) = dev_caps.iter().position(|&d| d == tl) {
5340 let lo = (di * b + t) * hc * dim;
5341 crate::dsv4::dspark_capture(
5342 &caps_all[lo..lo + hc * dim],
5343 &cfg,
5344 slot,
5345 &mut ds.main_hidden,
5346 );
5347 } else if tip
5348 && crate::dsv4::dspark_peek_slot(slot, dim, {
5349 let lo = slot * dim;
5350 &mut ds.main_hidden[lo..lo + dim]
5351 })
5352 {
5353 } else {
5358 crate::dsv4::dspark_capture(
5362 &walked[t * hc * dim..(t + 1) * hc * dim],
5363 &cfg,
5364 slot,
5365 &mut ds.main_hidden,
5366 );
5367 }
5368 }
5369 crate::dsv4::dspark_ring_append(g, &self.dsv4_mtp, &cfg, ds, next_pos + t, self.pool.as_deref());
5370 }
5371 let row = logits_all[(accepted - 1) * cfg.vocab..accepted * cfg.vocab].to_vec();
5372 self.graph_logits = Some(row);
5373 if std::env::var("CMF_DSV4_TRUNK_PICK_DUMP").is_ok() {
5378 crate::dsv4::trunk_freq_note(&crate::dsv4::pick_tally_take());
5379 crate::dsv4::pick_tally_arm();
5380 }
5381 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5382 eprintln!("spec_step total {:.1} мс (k={accepted})", t_all.elapsed().as_secs_f64() * 1e3);
5383 }
5384 Some((fed[1..accepted].to_vec(), next_pos + accepted))
5385 }
5386
5387 fn dspark_probe(&mut self, position: usize, token_id: u32) {
5388 if self.dsv4_mtp.is_empty() || !Self::draft_probe() {
5389 return;
5390 }
5391 let trunk_now = crate::dsv4::pick_tally_take();
5393 crate::dsv4::trunk_freq_note(&trunk_now);
5394 if !trunk_now.is_empty() {
5395 self.dspark_trunk_picks.push(trunk_now);
5396 let keep = crate::dsv4::dspark_block();
5397 if self.dspark_trunk_picks.len() > keep {
5398 self.dspark_trunk_picks.remove(0);
5399 }
5400 }
5401 for p in std::mem::take(&mut self.dspark_pending) {
5404 let Some(i) = position.checked_sub(p.0 + 1) else {
5405 continue;
5406 };
5407 let mut p = p;
5408 if i < p.1.len() {
5409 if p.2 && p.1[i] == token_id {
5410 p.3 = i + 1;
5411 } else {
5412 p.2 = false;
5413 }
5414 if i + 1 < p.1.len() {
5415 self.dspark_pending.push(p);
5416 continue;
5417 }
5418 }
5419 self.dspark_hist.push(p.3);
5420 self.dspark_real.push(token_id);
5421 }
5422 let Some(b) = &mut self.dsv4 else { return };
5423 let (g, layers, cfg) = (&b.0, &b.1, b.2);
5424 let n_layers = layers.len();
5425 if self.dspark.is_none() {
5426 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5427 if t.is_empty() {
5428 return;
5429 }
5430 eprintln!("DSpark: захват со слоёв {t:?}, блок {}", crate::dsv4::dspark_block());
5431 crate::dsv4::dspark_arm(&t, cfg.dim);
5432 self.dspark = Some(crate::dsv4::DsparkState::new(
5433 self.dsv4_mtp.len(),
5434 &cfg,
5435 t.len(),
5436 ));
5437 }
5438 let ds = self.dspark.as_mut().unwrap();
5439 if !crate::dsv4::dspark_take(&mut ds.main_hidden) {
5440 return; }
5442 let mut conf = Vec::new();
5443 crate::dsv4::pick_tally_arm();
5444 let draft_started = std::time::Instant::now();
5449 #[cfg(feature = "gpu")]
5450 let gpu_draft = crate::dsv4::dspark_gpu_on();
5451 #[cfg(not(feature = "gpu"))]
5452 let gpu_draft = false;
5453 let props = if gpu_draft {
5454 #[cfg(feature = "gpu")]
5455 {
5456 let kv_id = b.3.kv_id;
5457 match crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg) {
5458 Some(pk) => crate::dsv4::dspark_draft_gpu(
5459 g,
5460 &self.dsv4_mtp,
5461 &cfg,
5462 ds,
5463 pk,
5464 kv_id,
5465 token_id,
5466 position,
5467 self.pool.as_deref(),
5468 &mut conf,
5469 ),
5470 None => Vec::new(),
5471 }
5472 }
5473 #[cfg(not(feature = "gpu"))]
5474 Vec::new()
5475 } else {
5476 crate::gpu::cpu_scope(|| {
5477 crate::dsv4::dspark_draft(
5478 g,
5479 &self.dsv4_mtp,
5480 &cfg,
5481 ds,
5482 token_id,
5483 position,
5484 self.pool.as_deref(),
5485 &mut conf,
5486 )
5487 })
5488 };
5489 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5490 let draft_picks = crate::dsv4::pick_tally_take();
5491 crate::dsv4::dspark_freq_note(&draft_picks);
5492 crate::dsv4::pick_tally_arm();
5495 if !props.is_empty() {
5496 let (tu, tt) = {
5500 let flat: Vec<(usize, Vec<usize>)> = self
5501 .dspark_trunk_picks
5502 .iter()
5503 .flat_map(|v| v.iter().cloned())
5504 .collect();
5505 let mut per: std::collections::HashMap<usize, Vec<usize>> =
5507 std::collections::HashMap::new();
5508 for (li, picks) in flat {
5509 per.entry(li).or_default().extend(picks);
5510 }
5511 let n = per.len().max(1);
5512 let mut u = 0usize;
5513 let mut t = 0usize;
5514 for (_, v) in per {
5515 t += v.len();
5516 u += v.iter().collect::<std::collections::HashSet<_>>().len();
5517 }
5518 (u / n, t / n)
5519 };
5520 let (du, dt) = crate::dsv4::tally_unique(&draft_picks);
5521 self.dspark_exp.push((tu, tt, du, dt));
5522 self.dspark_pending.push((position, props, true, 0));
5523 }
5524 if self.dspark_hist.len() >= 8 && self.dspark_hist.len() % 8 == 0 {
5525 let n = self.dspark_hist.len() as f32;
5526 let mean: f32 = self.dspark_hist.iter().sum::<usize>() as f32 / n;
5527 let block = crate::dsv4::dspark_block();
5528 let mut at = vec![0usize; block + 1];
5529 for &k in &self.dspark_hist {
5530 at[k] += 1;
5531 }
5532 let mut surv = Vec::with_capacity(block);
5534 for i in 1..=block {
5535 let k = at[i..].iter().sum::<usize>() as f32 / n;
5536 surv.push(format!("{k:.2}"));
5537 }
5538 let distinct = self
5539 .dspark_real
5540 .iter()
5541 .collect::<std::collections::HashSet<_>>()
5542 .len();
5543 let (tu, tt, du, dt) = self.dspark_exp.iter().fold((0, 0, 0, 0), |a, b| {
5544 (a.0 + b.0, a.1 + b.1, a.2 + b.2, a.3 + b.3)
5545 });
5546 let m = self.dspark_exp.len().max(1);
5547 eprintln!(
5548 "DSpark: черновиков {}, принято в среднем {mean:.2} из {block} \
5549 (токенов за проход {:.2}), распределение {at:?}, выживание [{}]",
5550 self.dspark_hist.len(),
5551 mean + 1.0,
5552 surv.join(" ")
5553 );
5554 eprintln!(
5555 "DSpark: разных токенов {distinct} из {} (вырожденность), \
5556 эксперты ствол {}/{} на слой за {block} токенов, \
5557 черновик {}/{} за блок, draft {:.2} мс/блок",
5558 self.dspark_real.len(),
5559 tu / m,
5560 tt / m,
5561 du / m,
5562 dt / m,
5563 self.dspark_draft_ns as f64 / self.dspark_exp.len().max(1) as f64 / 1e6
5564 );
5565 }
5566 }
5567
5568 fn forward_layers_upto(
5569 &mut self,
5570 hidden: &[f32],
5571 position: usize,
5572 task_mask: Option<&TaskMask>,
5573 upto: Option<usize>,
5574 ) -> Vec<f32> {
5575 if let Some(plan) = self.gpu_plan.clone() {
5581 if upto.is_none() && plan.len() > 1 {
5582 let mut h = hidden.to_vec();
5583 for &(dev, from, upto_incl) in plan.iter() {
5584 h = crate::gpu::with_device(dev, || {
5585 self.forward_layers_span(&h, position, task_mask, from, Some(upto_incl))
5586 });
5587 }
5588 return h;
5589 }
5590 }
5591 self.forward_layers_span(hidden, position, task_mask, 0, upto)
5592 }
5593
5594 pub fn set_gpu_plan(&mut self, devices: Option<&[usize]>) -> Result<(), String> {
5599 self.set_gpu_plan_at(devices, None)
5600 }
5601
5602 pub fn set_gpu_plan_at(
5606 &mut self,
5607 devices: Option<&[usize]>,
5608 at: Option<usize>,
5609 ) -> Result<(), String> {
5610 let Some(devs) = devices.filter(|d| d.len() > 1) else {
5611 self.gpu_plan = None;
5612 return Ok(());
5613 };
5614 self.split_supported()?;
5615 let n = self.num_layers;
5616 if devs.len() > n {
5617 return Err(format!("{} devices for {n} layers", devs.len()));
5618 }
5619 if let Some(k) = at {
5620 if k == 0 || k >= n {
5621 return Err(format!("split at {k}: the model has {n} layers"));
5622 }
5623 if devs.len() == 2 {
5624 self.gpu_plan = Some(std::sync::Arc::new(vec![
5625 (devs[0], 0, k - 1),
5626 (devs[1], k, n - 1),
5627 ]));
5628 return Ok(());
5629 }
5630 return Err(format!(
5631 "an explicit split point takes exactly 2 devices, got {}",
5632 devs.len()
5633 ));
5634 }
5635 let per = n.div_ceil(devs.len());
5636 let mut plan = Vec::with_capacity(devs.len());
5637 let mut from = 0usize;
5638 for &d in devs {
5639 if from >= n {
5640 break;
5641 }
5642 let upto = (from + per - 1).min(n - 1);
5643 plan.push((d, from, upto));
5644 from = upto + 1;
5645 }
5646 self.gpu_plan = Some(std::sync::Arc::new(plan));
5647 Ok(())
5648 }
5649
5650 pub fn gpu_plan(&self) -> Option<Vec<(usize, usize, usize)>> {
5652 self.gpu_plan.as_ref().map(|p| p.as_ref().clone())
5653 }
5654
5655 fn forward_layers_span(
5661 &mut self,
5662 hidden: &[f32],
5663 position: usize,
5664 task_mask: Option<&TaskMask>,
5665 from: usize,
5666 upto: Option<usize>,
5667 ) -> Vec<f32> {
5668 debug_assert!(from == 0 || (self.dsv4.is_none() && self.g3n.is_none()));
5669 if let Some(b) = &mut self.dsv4 {
5675 let _ = (task_mask, upto);
5676 let token_id = hidden.first().copied().unwrap_or(0.0) as u32;
5677 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
5678 st.pos = position;
5679 let mut logits = Vec::new();
5680 crate::dsv4::forward_token(
5681 g,
5682 layers,
5683 &cfg,
5684 st,
5685 token_id,
5686 &self.inv_freq,
5687 self.pool.as_deref(),
5688 &mut logits,
5689 );
5690 self.graph_logits = Some(logits);
5691 self.dspark_probe(position, token_id);
5692 return vec![0.0; self.hidden_size];
5695 }
5696 if let Some(b) = &self.g3n {
5699 let _ = (task_mask, upto);
5700 return crate::g3n::g3n_forward(
5701 &b.0,
5702 &b.1,
5703 hidden,
5704 position,
5705 &mut self.kv_cache.layers,
5706 self.num_heads,
5707 self.num_kv_heads,
5708 self.head_dim,
5709 self.pool.as_deref(),
5710 );
5711 }
5712 let mut h = hidden.to_vec();
5713 let (nh, _nkv, _hd, hs, _rd, eps) = (
5716 self.num_heads,
5717 self.num_kv_heads,
5718 self.head_dim,
5719 self.hidden_size,
5720 self.rotary_dim,
5721 self.rms_eps,
5722 );
5723 let pool = self.pool.clone();
5724 let graph_env = std::env::var("CMF_GPU_WGPU_GRAPH").ok();
5736 let graph_on = match graph_env.as_deref() {
5737 Some("0") => false,
5738 Some("prefill") => false, Some(_) => true,
5740 None => crate::gpu::wgpu_graph_default(),
5746 };
5747 let graph_trusted =
5748 graph_env.is_some() || crate::gpu::wgpu_graph_default() || self.gdn_cfg.is_some();
5749 let race_eligible = graph_on
5750 && upto.is_none()
5751 && task_mask.is_none()
5752 && from == 0
5753 && !crate::gpu::graph_unsupported();
5754 let mut tail_start = 0usize;
5755 if race_eligible && crate::gpu::graph_race_use_graph(graph_trusted) {
5756 let t_graph = std::time::Instant::now();
5757 let mut lg = Vec::new();
5758 let mut gl = 0usize;
5759 let built = self.try_token_graph_wgpu(hidden, position, &mut lg, &mut gl);
5760 if built.is_none() && !self.o1_active() && self.attn_softcap == 0.0 {
5765 crate::gpu::graph_mark_unsupported();
5766 }
5767 graph_note(built.is_some());
5768 if let Some(hh) = built {
5769 let dur = t_graph.elapsed();
5770 if std::env::var("CMF_GRAPH_PROF").is_ok() {
5771 eprintln!("graph-call: {:.2} ms total", dur.as_secs_f64() * 1000.0);
5772 }
5773 if gl > 0 && gl < self.num_layers {
5774 h = hh;
5780 tail_start = gl;
5781 } else if graph_trusted || !crate::gpu::graph_race_first_token_hopeless(dur) {
5782 if !graph_trusted {
5783 crate::gpu::graph_race_record(true, dur);
5784 }
5785 if !lg.is_empty() {
5786 lg.resize(self.vocab_size, 0.0);
5789 if let Some(c) = self.final_softcap {
5790 for l in lg.iter_mut() {
5791 *l = c * (*l / c).tanh();
5792 }
5793 }
5794 self.graph_logits = Some(lg);
5795 }
5796 return hh;
5797 }
5798 }
5804 }
5805 let span = from > 0 || upto.is_some();
5829 if span && graph_on && task_mask.is_none() && graph_trusted {
5830 let upto_excl = upto.map_or(self.num_layers, |u| u + 1);
5831 let mut lg = Vec::new();
5832 let mut gl = 0usize;
5833 let span_res =
5834 self.try_token_graph_wgpu_span(hidden, position, &mut lg, from, upto_excl, &mut gl);
5835 graph_note(span_res.is_some() && gl == upto_excl - from);
5836 if std::env::var("CMF_GPU_DEBUG").is_ok() {
5837 static SEEN: std::sync::atomic::AtomicU32 =
5841 std::sync::atomic::AtomicU32::new(0);
5842 if SEEN.fetch_add(1, std::sync::atomic::Ordering::Relaxed) < 4 {
5843 eprintln!(
5844 "span graph: covered {gl} of {} layers [{from}..{upto_excl}) res={}",
5845 upto_excl - from,
5846 span_res.is_some()
5847 );
5848 }
5849 }
5850 if let Some(hh) = span_res {
5851 if gl == upto_excl - from {
5852 if !lg.is_empty() {
5853 lg.resize(self.vocab_size, 0.0);
5854 if let Some(c) = self.final_softcap {
5855 for l in lg.iter_mut() {
5856 *l = c * (*l / c).tanh();
5857 }
5858 }
5859 self.graph_logits = Some(lg);
5860 }
5861 crate::gpu::set_layer(-1);
5862 return hh;
5863 }
5864 h = hh;
5866 tail_start = from + gl;
5867 }
5868 }
5869 let t_race_cpu = (race_eligible && !graph_trusted).then(std::time::Instant::now);
5870
5871 #[cfg(target_os = "macos")]
5872 let mut gpu_skip_until = 0usize;
5873 for li in tail_start.max(from)..self.num_layers {
5874 crate::gpu::set_layer(li as i64); if let Some(u) = upto {
5876 if li > u {
5877 break;
5878 }
5879 }
5880 if let Some(mask) = task_mask {
5881 if !mask.layer_alive(li) {
5882 continue; }
5884 }
5885 #[cfg(target_os = "macos")]
5889 {
5890 if li < gpu_skip_until {
5891 continue;
5892 }
5893 if task_mask.is_none() {
5894 let end = self.q1_graph_gpu(li, upto, position, &mut h);
5895 if end > li {
5896 gpu_skip_until = end;
5897 if self.is_loop_end(end - 1) && end < self.num_layers {
5900 h = inference::rms_norm(
5901 &h,
5902 &self.weights.final_norm,
5903 self.rms_eps,
5904 self.norm_style,
5905 );
5906 }
5907 continue;
5908 }
5909 }
5910 }
5911
5912 let lw = &self.weights.layers[self.phys_layer(li)];
5913 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
5914 if tp.parse::<usize>().ok() == Some(position) {
5915 let n: f32 = h.iter().map(|x| x * x).sum::<f32>().sqrt();
5916 eprintln!(
5917 "TRACE pos {position} layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
5918 h[0], h[1]
5919 );
5920 }
5921 }
5922 inference::rms_norm_into(
5925 &h,
5926 &lw.input_norm,
5927 self.rms_eps,
5928 self.norm_style,
5929 &mut self.ws.n1,
5930 );
5931
5932 let attn_out = match &lw.attn {
5933 AttnKind::Mla(w) => {
5934 let inv_freq_l = self.layer_inv_freq(li);
5935 let rs = self.layer_rope_scale(li);
5936 let eps = self.rms_eps;
5937 let pool = self.pool.clone();
5938 mla_attention(
5939 w,
5940 &self.ws.n1,
5941 &mut self.kv_cache.layers[li],
5942 position,
5943 &inv_freq_l,
5944 rs,
5945 eps,
5946 pool.as_deref(),
5947 )
5948 }
5949 AttnKind::Linear(w) => {
5950 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
5951 vmf_phase_forward(
5952 &self.ws.n1,
5953 w,
5954 &cfg,
5955 &mut self.kv_cache.layers[li].linear_state,
5956 self.pool.as_deref(),
5957 )
5958 }
5959 AttnKind::Kda(w) => {
5960 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
5961 crate::linear_core::kda_forward(
5962 &self.ws.n1,
5963 w,
5964 &cfg,
5965 &mut self.kv_cache.layers[li].linear_state,
5966 self.pool.as_deref(),
5967 )
5968 }
5969 AttnKind::LinearGdn(w) => {
5970 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
5971 gdn_forward(
5972 &self.ws.n1,
5973 w,
5974 &cfg,
5975 &mut self.kv_cache.layers[li].linear_state,
5976 self.pool.as_deref(),
5977 )
5978 }
5979 AttnKind::ShortConv(w) => {
5980 let cfg = self
5981 .short_conv_cfg
5982 .expect("short-conv layer without short_conv_cfg");
5983 short_conv_forward(
5984 &self.ws.n1,
5985 w,
5986 &cfg,
5987 &mut self.kv_cache.layers[li].linear_state,
5988 self.pool.as_deref(),
5989 )
5990 }
5991 AttnKind::Full {
5992 wq,
5993 wk,
5994 wv,
5995 wo,
5996 q_norm,
5997 k_norm,
5998 output_gate,
5999 softplus_gate,
6000 bias,
6001 } if self.kv_cache.layers[li].o1_sealed() => {
6002 let inv_freq_l = self.layer_inv_freq(li);
6005 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6006 let cfg = QwenAttnCfg {
6007 num_heads: self.layer_num_heads(li),
6008 num_kv_heads: nkv_l,
6009 head_dim: hd_l,
6010 hidden_size: hs,
6011 position,
6012 inv_freq: &inv_freq_l,
6013 rotary_dim: rd_l,
6014 scale: self.attn_scale,
6015 softcap: self.attn_softcap,
6016 window: None,
6017 v_norm: self.attn_v_norm,
6018 q_norm: q_norm.as_deref(),
6019 k_norm: k_norm.as_deref(),
6020 output_gate: *output_gate,
6021 softplus_gate: softplus_gate
6022 .as_ref()
6023 .map(|(gate, per_head)| (gate, *per_head)),
6024 rope_scale: self.layer_rope_scale(li),
6025 bias: bias
6026 .as_ref()
6027 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6028 rms_eps: eps,
6029 norm_style: self.norm_style,
6030 pool: pool.as_deref(),
6031 };
6032 attention::qwen_attention_nystrom(
6033 &self.ws.n1,
6034 wq,
6035 wk,
6036 wv,
6037 wo,
6038 &mut self.kv_cache.layers[li],
6039 &cfg,
6040 )
6041 }
6042 AttnKind::Full {
6043 wq,
6044 wk,
6045 wv,
6046 wo,
6047 q_norm,
6048 k_norm,
6049 output_gate,
6050 softplus_gate,
6051 bias,
6052 } => 'attn: {
6053 if graph_on
6056 && !*output_gate
6057 && softplus_gate.is_none()
6058 && self.attention_heads_per_layer.is_none()
6059 && bias.is_none()
6060 && task_mask.is_none()
6061 {
6062 let inv_freq_l = self.layer_inv_freq(li);
6063 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6064 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
6065 if let (Some((gm, qi)), Some((_, ki)), Some((_, vi)), Some((_, oi))) = (
6066 wq.mapped_q1(),
6067 wk.mapped_q1(),
6068 wv.mapped_q1(),
6069 wo.mapped_q1(),
6070 ) {
6071 let gm = gm.clone();
6072 let mut out = vec![0f32; hs];
6073 let cache = &self.kv_cache.layers[li];
6074 if crate::gpu::attn_dropin(
6075 &gm,
6076 self.graph_kv_id,
6077 li,
6078 &self.ws.n1,
6079 qi,
6080 ki,
6081 vi,
6082 oi,
6083 q_norm.as_deref(),
6084 k_norm.as_deref(),
6085 &inv_freq_l,
6086 nh,
6087 nkv_l,
6088 hd_l,
6089 rd_l,
6090 hs,
6091 position,
6092 self.kv_cache.max_seq_len,
6093 gemma,
6094 eps as f32,
6095 cache.k_heads(),
6096 cache.v_heads(),
6097 &mut out,
6098 ) {
6099 break 'attn out;
6100 }
6101 }
6102 }
6103 let masked = task_mask
6104 .map(|m| m.head_flags(li, self.num_heads).iter().any(|&a| !a))
6105 .unwrap_or(false);
6106 let f32_view = (wq.as_f32(), wk.as_f32(), wv.as_f32(), wo.as_f32());
6107 match (masked, f32_view) {
6108 (true, (Some(q), Some(k), Some(v), Some(o))) => {
6111 let active_heads = task_mask.unwrap().head_flags(li, self.num_heads);
6112 attention::multi_head_attention(
6113 &self.ws.n1,
6114 q,
6115 k,
6116 v,
6117 o,
6118 &mut self.kv_cache.layers[li],
6119 self.num_heads,
6120 self.num_kv_heads,
6121 self.head_dim,
6122 self.hidden_size,
6123 position,
6124 &active_heads,
6125 &self.inv_freq,
6126 )
6127 }
6128 (masked, _) => {
6129 if masked {
6130 tracing::warn!(
6131 "layer {li}: head mask on quantized weights not \
6132 supported yet — executing dense"
6133 );
6134 }
6135 let inv_freq_l = self.layer_inv_freq(li);
6136 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6137 let cfg = QwenAttnCfg {
6138 num_heads: self.layer_num_heads(li),
6139 num_kv_heads: nkv_l,
6140 head_dim: hd_l,
6141 hidden_size: hs,
6142 position,
6143 inv_freq: &inv_freq_l,
6144 rotary_dim: rd_l,
6145 scale: self.attn_scale,
6146 softcap: self.attn_softcap,
6147 window: self.layer_window(li),
6148 v_norm: self.attn_v_norm,
6149 q_norm: q_norm.as_deref(),
6150 k_norm: k_norm.as_deref(),
6151 output_gate: *output_gate,
6152 softplus_gate: softplus_gate
6153 .as_ref()
6154 .map(|(gate, per_head)| (gate, *per_head)),
6155 rope_scale: self.layer_rope_scale(li),
6156 bias: bias
6157 .as_ref()
6158 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6159 rms_eps: eps,
6160 norm_style: self.norm_style,
6161 pool: pool.as_deref(),
6162 };
6163 attention::qwen_attention(
6164 &self.ws.n1,
6165 wq,
6166 wk,
6167 wv,
6168 wo,
6169 &mut self.kv_cache.layers[li],
6170 &cfg,
6171 )
6172 }
6173 }
6174 }
6175 };
6176 let attn_out = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
6179 Some(w) => inference::rms_norm(&attn_out, w, self.rms_eps, self.norm_style),
6180 None => attn_out,
6181 };
6182 let lw = &self.weights.layers[self.phys_layer(li)];
6183 inference::add_rmsnorm_fused_into(
6184 &mut h,
6185 &attn_out,
6186 &lw.post_norm,
6187 self.rms_eps,
6188 self.norm_style,
6189 &mut self.ws.p1,
6190 );
6191 let mut attn_out = attn_out;
6192 attention::recycle_buf(&mut attn_out);
6193 let post_normed = &self.ws.p1;
6194
6195 let ffn_masked = task_mask
6196 .map(|m| m.ffn_active_count(li) < self.intermediate_size)
6197 .unwrap_or(false);
6198 let ffn_out = match (ffn_masked, &lw.ffn) {
6210 (true, FfnKind::Dense(d)) => {
6211 let tm = task_mask.unwrap();
6212 let alive = tm.ffn_active_count(li);
6213 let deep = alive * 2 <= self.intermediate_size;
6214 if deep && d.down_proj.sparse_col_ok() {
6215 let active = tm.ffn_active_indices(li);
6216 sparse_ffn_quant(
6217 d,
6218 post_normed,
6219 &active,
6220 self.hidden_size,
6221 self.pool.as_deref(),
6222 )
6223 } else if deep
6224 && let (Some(g), Some(u), Some(dn)) = (
6225 d.gate_proj.as_f32(),
6226 d.up_proj.as_f32(),
6227 d.down_proj.as_f32(),
6228 )
6229 {
6230 let active = tm.ffn_active_indices(li);
6231 inference::sparse_ffn_forward(
6232 post_normed,
6233 g,
6234 u,
6235 dn,
6236 self.hidden_size,
6237 self.intermediate_size,
6238 &active,
6239 self.pool.as_deref(),
6240 )
6241 } else {
6242 let row = tm.ffn_masks.get(li).map(|v| v.as_slice());
6243 dense_ffn_batch(d, post_normed, 1, self.pool.as_deref(), row)
6244 }
6245 }
6246 (true, FfnKind::Moe(m)) => {
6247 let allowed = task_mask.and_then(|tm| tm.expert_flags(li, m.experts.len()));
6251 ffn_forward(
6252 &lw.ffn,
6253 post_normed,
6254 self.pool.as_deref(),
6255 allowed.as_deref(),
6256 )
6257 }
6258 (true, FfnKind::DenseMoe(dm)) => dense_moe_ffn(
6259 dm,
6260 post_normed,
6261 &h,
6262 self.rms_eps,
6263 self.norm_style,
6264 self.pool.as_deref(),
6265 ),
6266 (false, _) => match &lw.ffn {
6267 FfnKind::DenseMoe(dm) => dense_moe_ffn(
6268 dm,
6269 post_normed,
6270 &h,
6271 self.rms_eps,
6272 self.norm_style,
6273 self.pool.as_deref(),
6274 ),
6275 _ => {
6276 let allowed = match (&lw.ffn, task_mask) {
6277 (FfnKind::Moe(m), Some(tm)) => tm.expert_flags(li, m.experts.len()),
6278 _ => None,
6279 };
6280 ffn_forward(
6281 &lw.ffn,
6282 post_normed,
6283 self.pool.as_deref(),
6284 allowed.as_deref(),
6285 )
6286 }
6287 },
6288 };
6289 let ffn_out = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
6290 Some(w) => inference::rms_norm(&ffn_out, w, self.rms_eps, self.norm_style),
6291 None => ffn_out,
6292 };
6293 for (i, &f) in ffn_out.iter().enumerate() {
6294 h[i] += f;
6295 }
6296 let mut ffn_out = ffn_out;
6297 attention::recycle_buf(&mut ffn_out);
6298
6299 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
6301 for v in h.iter_mut() {
6302 *v *= sc;
6303 }
6304 }
6305
6306 if self.is_loop_end(li) && li + 1 < self.num_layers {
6309 h = inference::rms_norm(
6310 &h,
6311 &self.weights.final_norm,
6312 self.rms_eps,
6313 self.norm_style,
6314 );
6315 }
6316
6317 if self.dyn_phi_layer == Some(li) {
6321 self.update_dyn_phi(&h);
6322 }
6323 }
6324 crate::gpu::set_layer(-1); if let Some(t) = t_race_cpu {
6326 crate::gpu::graph_race_record(false, t.elapsed());
6327 }
6328
6329 h
6330 }
6331
6332 fn update_dyn_phi(&mut self, h: &[f32]) {
6335 const A: f32 = 0.2;
6336 if self.dyn_phi_ema.len() != h.len() {
6337 self.dyn_phi_ema = vec![0.0; h.len()];
6338 self.dyn_phi_seen = 0;
6339 }
6340 if self.dyn_phi_seen == 0 {
6341 self.dyn_phi_ema.copy_from_slice(h);
6342 } else {
6343 for (e, &v) in self.dyn_phi_ema.iter_mut().zip(h) {
6344 *e = (1.0 - A) * *e + A * v;
6345 }
6346 }
6347 self.dyn_phi_seen += 1;
6348 }
6349
6350 pub fn dyn_phi(&self) -> &[f32] {
6352 &self.dyn_phi_ema
6353 }
6354
6355 pub fn set_dyn_phi_layer(&mut self, layer: Option<usize>) {
6357 self.dyn_phi_layer = layer;
6358 self.dyn_phi_ema.clear();
6359 self.dyn_phi_seen = 0;
6360 }
6361
6362 pub fn dynamic_skills(&self) -> Vec<(usize, String, usize)> {
6364 let Some(model) = &self.model else {
6365 return Vec::new();
6366 };
6367 model
6368 .header
6369 .skills
6370 .iter()
6371 .enumerate()
6372 .filter_map(|(i, sk)| {
6373 let ok = matches!(self.dyn_skill_layers.get(i), Some(Some(_)));
6374 let sel = sk.selection.as_ref()?;
6375 (ok).then(|| (i, sk.id.clone(), sel.phi_layer))
6376 })
6377 .collect()
6378 }
6379
6380 pub fn active_skill(&self) -> Option<usize> {
6382 self.dyn_active
6383 }
6384
6385 pub fn enable_dynamic_routing(&mut self) -> usize {
6390 use crate::swarm::{DynRouter, RoutableSkill};
6391 let Some(model) = self.model.clone() else {
6392 return 0;
6393 };
6394 if self.dyn_blend_loaded {
6397 tracing::warn!("dynamic routing unavailable on a blend-loaded pipeline");
6398 return 0;
6399 }
6400 if let Some(a) = self.dyn_active {
6404 if !matches!(self.dyn_skill_layers.get(a), Some(Some(_))) {
6405 tracing::warn!("loaded skill is not FFN-eligible — dynamic routing unavailable");
6406 return 0;
6407 }
6408 }
6409 let hidden = self.hidden_size;
6410 let mut skills = Vec::new();
6411 for (idx, id, _phi) in self.dynamic_skills() {
6412 if let Some(sel) = model.header.skills[idx].selection.as_ref() {
6413 if let Some(rs) = RoutableSkill::from_descriptor(idx, id, sel, hidden) {
6414 skills.push(rs);
6415 }
6416 }
6417 }
6418 if skills.is_empty() {
6419 return 0;
6420 }
6421 let phi = skills[0].phi_layer;
6423 if skills.iter().any(|s| s.phi_layer != phi) {
6424 tracing::warn!("routable skills disagree on phi_layer; using {phi}");
6425 }
6426 let n = skills.len();
6427 self.set_dyn_phi_layer(Some(phi));
6428 self.dyn_router = Some(DynRouter::new(skills));
6429 n
6430 }
6431
6432 pub fn route_switches(&self) -> Vec<(usize, Option<String>, Option<String>)> {
6434 self.dyn_router
6435 .as_ref()
6436 .map(|r| r.switches.clone())
6437 .unwrap_or_default()
6438 }
6439
6440 fn lm_head_forward(&self, hidden: &[f32]) -> Vec<f32> {
6443 let rows = self.weights.lm_head.rows();
6444 let mut logits = attention::take_buf(rows.min(self.vocab_size));
6445 self.weights
6446 .lm_head
6447 .matvec(hidden, &mut logits, self.pool.as_deref());
6448 logits.resize(self.vocab_size, 0.0);
6449 if let Some(m) = self.logit_multiplier {
6450 for l in logits.iter_mut() {
6451 *l *= m;
6452 }
6453 }
6454 if let Some(c) = self.final_softcap {
6455 for l in logits.iter_mut() {
6456 *l = c * (*l / c).tanh();
6457 }
6458 }
6459 logits
6460 }
6461
6462 pub fn prefill_next_logits(&mut self, ids: &[u32], task_mask: Option<&TaskMask>) -> Vec<f32> {
6467 self.kv_cache.clear();
6468 self.kv_history.clear();
6469 let mut hidden = vec![0.0f32; self.hidden_size];
6470 for (pos, &id) in ids.iter().enumerate() {
6471 let emb = self.embed_single(id);
6472 hidden = self.forward_layers(&emb, pos, task_mask);
6473 }
6474 inference::rms_norm_into(
6475 &hidden,
6476 &self.weights.final_norm,
6477 self.rms_eps,
6478 self.norm_style,
6479 &mut self.ws.n1,
6480 );
6481 self.lm_head_forward(&self.ws.n1)
6482 }
6483}
6484
6485pub fn create_test_pipeline(
6487 hidden_size: usize,
6488 intermediate_size: usize,
6489 num_heads: usize,
6490 num_kv_heads: usize,
6491 head_dim: usize,
6492 num_layers: usize,
6493 vocab_size: usize,
6494) -> Pipeline {
6495 let synth = |n: usize, salt: usize| -> Vec<f32> {
6498 (0..n)
6499 .map(|i| (((i * 31 + salt * 17 + 7) % 97) as f32 / 97.0 - 0.5) * 0.2)
6500 .collect()
6501 };
6502 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
6503 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
6504 };
6505 let layer_weights: Vec<LayerWeights> = (0..num_layers)
6506 .map(|li| LayerWeights {
6507 input_norm: vec![1.0; hidden_size],
6508 post_norm: vec![1.0; hidden_size],
6509 attn_out_norm: None,
6510 ffn_out_norm: None,
6511 layer_scale: None,
6512 ffn: FfnKind::Dense(DenseFfn {
6513 gate_proj: qt(intermediate_size, hidden_size, li * 10 + 5),
6514 up_proj: qt(intermediate_size, hidden_size, li * 10 + 6),
6515 down_proj: qt(hidden_size, intermediate_size, li * 10 + 7),
6516 act: Act::Silu,
6517 }),
6518 attn: AttnKind::Full {
6519 bias: None,
6520 wq: qt(num_heads * head_dim, hidden_size, li * 10 + 1),
6521 wk: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 2),
6522 wv: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 3),
6523 wo: qt(hidden_size, num_heads * head_dim, li * 10 + 4),
6524 q_norm: None,
6525 k_norm: None,
6526 output_gate: false,
6527 softplus_gate: None,
6528 },
6529 })
6530 .collect();
6531
6532 Pipeline::new(
6533 Tokenizer::byte_level(),
6534 PipelineWeights {
6535 embed_tokens: qt(vocab_size, hidden_size, 100),
6536 layers: layer_weights,
6537 lm_head: qt(vocab_size, hidden_size, 200),
6538 final_norm: vec![1.0; hidden_size],
6539 },
6540 hidden_size,
6541 intermediate_size,
6542 num_heads,
6543 num_kv_heads,
6544 head_dim,
6545 num_layers,
6546 num_layers, false, vocab_size,
6549 1e-6,
6550 10_000.0,
6551 NormStyle::Qwen,
6552 4096,
6553 SamplerConfig {
6554 seed: Some(42),
6555 ..Default::default()
6556 },
6557 )
6558}
6559
6560#[inline]
6565fn mask_bit(row: &[u8], j: usize) -> bool {
6566 (row.get(j >> 3).copied().unwrap_or(0) >> (j & 7)) & 1 != 0
6567}
6568
6569fn zero_masked_cols(g: &mut [f32], rows: usize, inter: usize, row: &[u8]) {
6575 for r in 0..rows {
6576 let base = r * inter;
6577 for (bi, &byte) in row.iter().enumerate() {
6578 if byte == 0xFF {
6579 continue;
6580 }
6581 let j0 = bi * 8;
6582 for bit in 0..8 {
6583 let j = j0 + bit;
6584 if j < inter && byte & (1 << bit) == 0 {
6585 g[base + j] = 0.0;
6586 }
6587 }
6588 }
6589 }
6590}
6591
6592fn dense_ffn_batch(
6593 d: &DenseFfn,
6594 xs: &[f32],
6595 b: usize,
6596 pool: Option<&Pool>,
6597 mask_row: Option<&[u8]>,
6598) -> Vec<f32> {
6599 let inter = d.gate_proj.rows();
6600 let hidden = d.down_proj.rows();
6601 if mask_row.is_none()
6609 && d.act == Act::Silu
6610 && b >= 32
6611 && crate::gpu::enabled_here()
6612 && !crate::gpu::mm_killed()
6613 {
6614 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6615 d.gate_proj.mapped_q4t(),
6616 d.up_proj.mapped_q4t(),
6617 d.down_proj.mapped_q4t(),
6618 ) {
6619 let mut out = vec![0.0f32; b * hidden];
6620 if crate::gpu::q4t_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6621 return out;
6622 }
6623 }
6624 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6629 d.gate_proj.mapped_q4tp(),
6630 d.up_proj.mapped_q4tp(),
6631 d.down_proj.mapped_q4tp(),
6632 ) {
6633 let mut out = vec![0.0f32; b * hidden];
6634 if crate::gpu::q4tp_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6635 return out;
6636 }
6637 }
6638 }
6639 let mut g = vec![0.0f32; b * inter];
6640 d.gate_proj.matmat(xs, b, &mut g, pool);
6641 let mut u = vec![0.0f32; b * inter];
6642 d.up_proj.matmat(xs, b, &mut u, pool);
6643 for i in 0..b * inter {
6644 g[i] = d.act.combine(g[i], u[i]);
6645 }
6646 if let Some(row) = mask_row {
6647 zero_masked_cols(&mut g, b, inter, row);
6648 }
6649 let mut out = vec![0.0f32; b * hidden];
6650 d.down_proj.matmat(&g, b, &mut out, pool);
6651 out
6652}
6653
6654fn accumulate_act(m: &MoeFfn, xs: &[f32], b: usize) {
6659 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6660 static DUMP: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6661 let on = *ON.get_or_init(|| std::env::var("CMF_RMS_TRACE").is_ok());
6662 let dump = *DUMP.get_or_init(|| std::env::var("CMF_ACT_DUMP").is_ok());
6663 if (!on && !dump) || b == 0 {
6664 return;
6665 }
6666 let hidden = xs.len() / b;
6667 if on {
6668 let mut acc = m.act_sq.borrow_mut();
6669 if acc.len() < hidden {
6670 acc.resize(hidden, 0.0);
6671 }
6672 for t in 0..b {
6673 let row = &xs[t * hidden..(t + 1) * hidden];
6674 for (a, &v) in acc.iter_mut().zip(row) {
6675 *a += (v as f64) * (v as f64);
6676 }
6677 }
6678 }
6679 if dump {
6680 let cap: usize = std::env::var("CMF_ACT_DUMP_ROWS")
6683 .ok()
6684 .and_then(|v| v.parse().ok())
6685 .unwrap_or(4096);
6686 let mut rows = m.act_rows.borrow_mut();
6687 if rows.len() < cap * hidden {
6688 let take = b.min((cap * hidden - rows.len()) / hidden.max(1));
6689 rows.extend_from_slice(&xs[..take * hidden]);
6690 }
6691 }
6692}
6693
6694#[derive(Clone, Copy)]
6697struct SendVecs(*mut Vec<f32>);
6698unsafe impl Send for SendVecs {}
6699unsafe impl Sync for SendVecs {}
6700impl SendVecs {
6701 #[inline]
6702 fn at(self, i: usize) -> *mut Vec<f32> {
6703 unsafe { self.0.add(i) }
6704 }
6705}
6706
6707fn moe_ffn_batch(
6708 m: &MoeFfn,
6709 xs: &[f32],
6710 b: usize,
6711 hidden: usize,
6712 pool: Option<&Pool>,
6713 allowed: Option<&[bool]>,
6714) -> Vec<f32> {
6715 accumulate_act(m, xs, b);
6716 let ne = m.experts.len();
6717 let mut logits = vec![0.0f32; b * ne];
6718 m.router.matmat(xs, b, &mut logits, pool);
6719
6720 let mut assign: Vec<Vec<(usize, f32)>> = vec![Vec::new(); ne];
6723 {
6724 let mut st = m.stats.borrow_mut();
6725 if st.len() < ne {
6726 st.resize(ne, 0);
6727 }
6728 for bi in 0..b {
6729 let (idx, p, wsum) = moe_route(&logits[bi * ne..(bi + 1) * ne], m, allowed);
6730 for &e in &idx {
6731 st[e] += 1;
6732 assign[e].push((bi, p[e] / wsum));
6733 }
6734 }
6735 }
6736
6737 let mut out = vec![0.0f32; b * hidden];
6738 let cols = m.experts[0].gate_proj.cols();
6739 let run_expert = |d: &DenseFfn, list: &[(usize, f32)], out: &mut [f32]| {
6740 let sb = list.len();
6741 let mut sub = vec![0.0f32; sb * cols];
6742 for (k, &(bi, _)) in list.iter().enumerate() {
6743 sub[k * cols..(k + 1) * cols].copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6744 }
6745 let eo = dense_ffn_batch(d, &sub, sb, pool, None);
6746 for (k, &(bi, w)) in list.iter().enumerate() {
6747 for i in 0..hidden {
6748 out[bi * hidden + i] += w * eo[k * hidden + i];
6749 }
6750 }
6751 };
6752 let active: Vec<usize> = (0..ne).filter(|&e| !assign[e].is_empty()).collect();
6758 if pool.is_some() && active.len() >= 8 {
6759 let mut panels: Vec<Vec<f32>> = vec![Vec::new(); active.len()];
6760 {
6761 let panel_ptr = SendVecs(panels.as_mut_ptr());
6762 let experts = &m.experts;
6765 let (active_r, assign_r) = (&active, &assign);
6766 let run = |start: usize, end: usize| {
6767 for ai in start..end {
6768 let e = active_r[ai];
6769 let list = &assign_r[e];
6770 let sb = list.len();
6771 let mut sub = vec![0.0f32; sb * cols];
6772 for (k, &(bi, _)) in list.iter().enumerate() {
6773 sub[k * cols..(k + 1) * cols]
6774 .copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6775 }
6776 unsafe {
6778 *panel_ptr.at(ai) =
6779 dense_ffn_batch(&experts[e], &sub, sb, None, None);
6780 }
6781 }
6782 };
6783 match pool {
6784 Some(p) => p.run_rows(active.len(), &run),
6785 None => run(0, active.len()),
6786 }
6787 }
6788 for (ai, &e) in active.iter().enumerate() {
6789 for (k, &(bi, w)) in assign[e].iter().enumerate() {
6790 let eo = &panels[ai][k * hidden..(k + 1) * hidden];
6791 for i in 0..hidden {
6792 out[bi * hidden + i] += w * eo[i];
6793 }
6794 }
6795 }
6796 } else {
6797 for &e in &active {
6798 run_expert(&m.experts[e], &assign[e], &mut out);
6799 }
6800 }
6801 if let Some((se, gate)) = &m.shared {
6802 let all: Vec<(usize, f32)> = if let Some(gate) = gate {
6803 let mut gl = vec![0.0f32; b];
6804 gate.matmat(xs, b, &mut gl, pool);
6805 (0..b)
6806 .map(|bi| (bi, 1.0 / (1.0 + (-gl[bi]).exp())))
6807 .collect()
6808 } else {
6809 (0..b).map(|bi| (bi, 1.0)).collect()
6810 };
6811 run_expert(se, &all, &mut out);
6812 }
6813 out
6814}
6815
6816thread_local! {
6817 static FFN_SCRATCH: std::cell::RefCell<[Vec<f32>; 4]> =
6821 const { std::cell::RefCell::new([Vec::new(), Vec::new(), Vec::new(), Vec::new()]) };
6822}
6823
6824fn dense_ffn(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6826 if crate::gpu::enabled_here()
6837 && (d.gate_proj.rows() >= crate::gpu::min_rows() || d.gate_proj.is_q1())
6838 {
6839 let arm = if d.gate_proj.is_q1() && crate::gpu::q1_force() {
6840 crate::gpu::ProbeArm::Gpu
6841 } else {
6842 crate::gpu::probe_arm(crate::gpu::OpClass::Ffn)
6843 };
6844 match arm {
6845 crate::gpu::ProbeArm::Gpu => {
6846 let t0 = std::time::Instant::now();
6847 if let Some(out) = dense_ffn_gpu(d, x, pool) {
6848 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
6849 return out;
6850 }
6851 }
6852 crate::gpu::ProbeArm::CpuTimed => {
6853 let t0 = std::time::Instant::now();
6854 let out = crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6855 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
6856 return out;
6857 }
6858 crate::gpu::ProbeArm::Cpu => {
6859 return crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6860 }
6861 }
6862 }
6863 dense_ffn_cpu(d, x, pool)
6864}
6865
6866fn dense_ffn_cpu(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6868 let inter = d.gate_proj.rows();
6869 FFN_SCRATCH.with(|s| {
6870 let mut s = s.borrow_mut();
6871 let [g, u, ..] = &mut *s;
6872 g.resize(inter, 0.0);
6873 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
6876 } else {
6878 u.resize(inter, 0.0);
6879 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
6881 for i in 0..inter {
6882 g[i] = d.act.combine(g[i], u[i]);
6883 }
6884 }
6885 FFN_PROBE.with(|pr| {
6888 if let Some(acc) = pr.borrow_mut().as_mut() {
6889 let li = crate::gpu::cur_layer();
6890 if li >= 0 {
6891 if let Some(row) = acc.get_mut(li as usize) {
6892 for (a, &v) in row.iter_mut().zip(g.iter()) {
6893 *a += (v as f64).abs();
6894 }
6895 }
6896 }
6897 }
6898 });
6899 let mut out = attention::take_buf(d.down_proj.rows());
6900 d.down_proj.matvec(g, &mut out, pool);
6901 out
6902 })
6903}
6904
6905thread_local! {
6906 static FFN_PROBE: std::cell::RefCell<Option<Vec<Vec<f64>>>> =
6909 const { std::cell::RefCell::new(None) };
6910}
6911
6912fn dense_ffn_masked(
6917 d: &DenseFfn,
6918 x: &[f32],
6919 pool: Option<&Pool>,
6920 mask_row: &[u8],
6921) -> Vec<f32> {
6922 let inter = d.gate_proj.rows();
6923 FFN_SCRATCH.with(|s| {
6924 let mut s = s.borrow_mut();
6925 let [g, u, ..] = &mut *s;
6926 g.resize(inter, 0.0);
6927 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
6928 } else {
6930 u.resize(inter, 0.0);
6931 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
6932 for i in 0..inter {
6933 g[i] = d.act.combine(g[i], u[i]);
6934 }
6935 }
6936 zero_masked_cols(g, 1, inter, mask_row);
6937 let mut out = attention::take_buf(d.down_proj.rows());
6938 d.down_proj.matvec(g, &mut out, pool);
6939 out
6940 })
6941}
6942
6943fn dense_ffn_gpu(d: &DenseFfn, x: &[f32], _pool: Option<&Pool>) -> Option<Vec<f32>> {
6949 if d.act != Act::Silu {
6951 return None;
6952 }
6953 if d.gate_proj.rows() < crate::gpu::min_rows() && !d.gate_proj.is_q1() {
6956 return None;
6957 }
6958 let mut jobs: Vec<crate::gpu::MoeJob> = Vec::with_capacity(1);
6959 let mut model_ref = None;
6960 moe_push_job(d, x, 1.0, &mut jobs, &mut model_ref)?;
6961 let model = model_ref?;
6962 let hidden = jobs[0].down.1;
6963 let mut out = attention::take_buf(hidden);
6964 if crate::gpu::moe_block(&model, &jobs, &mut out) {
6965 Some(out)
6966 } else {
6967 let mut out = out;
6968 attention::recycle_buf(&mut out);
6969 None
6970 }
6971}
6972
6973#[allow(clippy::type_complexity)]
6978#[allow(clippy::type_complexity)]
6979pub(crate) fn moe_parts(
6980 t: &QTensor,
6981) -> Option<(
6982 &std::sync::Arc<cortiq_core::CmfModel>,
6983 usize,
6984 usize,
6985 usize,
6986 &[f32],
6987 &[f32],
6988 bool,
6989 bool,
6990 bool,
6991)> {
6992 match t {
6993 QTensor::Mapped {
6994 model,
6995 idx,
6996 dtype: dt @ (cortiq_core::TensorDtype::Q8_2f | cortiq_core::TensorDtype::Q8Row),
6997 rows,
6998 cols,
6999 row_scale,
7000 col_field,
7001 ..
7002 } if (*dt == cortiq_core::TensorDtype::Q8Row) || !col_field.is_empty() => Some((
7003 model, *idx, *rows, *cols, row_scale, col_field, false, false, false,
7004 )),
7005 QTensor::Mapped {
7007 model,
7008 idx,
7009 dtype: cortiq_core::TensorDtype::Q1,
7010 rows,
7011 cols,
7012 ..
7013 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], true, false, false)),
7014 QTensor::Mapped {
7016 model,
7017 idx,
7018 dtype: cortiq_core::TensorDtype::Q4Tiled,
7019 rows,
7020 cols,
7021 ..
7022 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7023 QTensor::Mapped {
7025 model,
7026 idx,
7027 dtype: cortiq_core::TensorDtype::Q4TiledP,
7028 rows,
7029 cols,
7030 ..
7031 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7032 QTensor::Mapped {
7036 model,
7037 idx,
7038 dtype: cortiq_core::TensorDtype::Q2TiledP,
7039 rows,
7040 cols,
7041 ..
7042 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, true)),
7043 _ => None,
7044 }
7045}
7046
7047#[cfg(target_os = "macos")]
7053fn metal_moe_graph_parts(m: &MoeFfn, hidden: usize) -> Option<crate::gpu::GpuMoe<'_>> {
7054 if m.router_sigmoid
7055 || m.router_input_norm
7056 || m.expert_bias.is_some()
7057 || m.route_tau.is_some()
7058 || m.mask.is_some()
7059 || m.per_expert_scale.is_some()
7060 || m.experts.is_empty()
7061 || m.top_k == 0
7062 {
7063 return None;
7064 }
7065 let (sh, sg) = match &m.shared {
7068 Some((sh, Some(sg))) => (sh, sg),
7069 _ => return None,
7070 };
7071 let (rf, rr, rc) = m.router.f32_parts()?;
7072 if rr != m.experts.len() || rc != hidden {
7073 return None;
7074 }
7075 let (sf, sr, sc) = sg.f32_parts()?;
7076 if sr * sc != hidden {
7077 return None;
7078 }
7079 let inter = m.experts[0].gate_proj.rows();
7080 let gu_q2 = m.experts[0].gate_proj.mapped_q2tp().is_some();
7083 let trio = |e: &DenseFfn| -> Option<(usize, usize, usize)> {
7084 if e.act != Act::Silu
7085 || e.gate_proj.rows() != inter
7086 || e.gate_proj.cols() != hidden
7087 || e.up_proj.rows() != inter
7088 || e.up_proj.cols() != hidden
7089 || e.down_proj.rows() != hidden
7090 || e.down_proj.cols() != inter
7091 {
7092 return None;
7093 }
7094 let pick = |t: &QTensor| -> Option<usize> {
7095 if gu_q2 {
7096 t.mapped_q2tp().map(|(_, i)| i)
7097 } else {
7098 t.mapped_q4tp().map(|(_, i)| i)
7099 }
7100 };
7101 Some((
7102 pick(&e.gate_proj)?,
7103 pick(&e.up_proj)?,
7104 e.down_proj.mapped_q4tp().map(|(_, i)| i)?,
7105 ))
7106 };
7107 let experts = m
7108 .experts
7109 .iter()
7110 .map(trio)
7111 .collect::<Option<Vec<_>>>()?;
7112 let shared = trio(sh)?;
7113 Some(crate::gpu::GpuMoe {
7114 router: rf,
7115 sgate: sf,
7116 experts,
7117 shared,
7118 n_exp: m.experts.len(),
7119 top_k: m.top_k,
7120 inter,
7121 norm_topk: m.norm_topk_prob,
7122 route_scale: m.routed_scaling,
7123 gu_q2,
7124 })
7125}
7126
7127pub(crate) fn moe_push_job_parts<'a>(
7131 gate: &'a QTensor,
7132 up: &'a QTensor,
7133 down: &'a QTensor,
7134 x: &[f32],
7135 w: f32,
7136 swiglu_limit: f32,
7137 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7138 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7139) -> Option<()> {
7140 use crate::qtensor::prescale;
7141 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(gate)?;
7142 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(up)?;
7143 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(down)?;
7144 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7145 return None; }
7147 if gq2 && (dq2 || !dq4 || down.mapped_q4tp().is_none()) {
7150 return None;
7151 }
7152 if !gq2 && dq2 {
7153 return None;
7154 }
7155 model_ref.get_or_insert_with(|| gm.clone());
7156 let dt = |cf: &[f32]| {
7157 if cf.is_empty() {
7158 cortiq_core::TensorDtype::Q8Row
7159 } else {
7160 cortiq_core::TensorDtype::Q8_2f
7161 }
7162 };
7163 jobs.push(crate::gpu::MoeJob {
7164 gate: (gi, gr, gc, grs),
7165 up: (ui, ur, uc, urs),
7166 down: (di, dr, dc, drs),
7167 xs_gate: prescale(x, gcf, dt(gcf)).into_owned(),
7168 xs_up: prescale(x, ucf, dt(ucf)).into_owned(),
7169 down_col: dcf,
7170 w,
7171 q1: gq1,
7172 q4t: gq4 && !gq2 && gate.mapped_q4tp().is_none(),
7173 q4tp: gq4 && (gq2 || gate.mapped_q4tp().is_some()),
7174 gu_q2: gq2,
7175 swiglu_limit,
7176 });
7177 Some(())
7178}
7179
7180fn moe_push_job<'a>(
7182 d: &'a DenseFfn,
7183 x: &[f32],
7184 w: f32,
7185 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7186 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7187) -> Option<()> {
7188 use crate::qtensor::prescale;
7189 if d.act != Act::Silu {
7190 return None; }
7192 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(&d.gate_proj)?;
7193 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(&d.up_proj)?;
7194 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(&d.down_proj)?;
7195 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7196 return None; }
7198 if gq2 && (dq2 || !dq4 || d.down_proj.mapped_q4tp().is_none()) {
7199 return None;
7200 }
7201 if !gq2 && dq2 {
7202 return None;
7203 }
7204 model_ref.get_or_insert_with(|| gm.clone());
7205 let gdt = if gcf.is_empty() {
7206 cortiq_core::TensorDtype::Q8Row
7207 } else {
7208 cortiq_core::TensorDtype::Q8_2f
7209 };
7210 let udt = if ucf.is_empty() {
7211 cortiq_core::TensorDtype::Q8Row
7212 } else {
7213 cortiq_core::TensorDtype::Q8_2f
7214 };
7215 jobs.push(crate::gpu::MoeJob {
7216 gate: (gi, gr, gc, grs),
7217 up: (ui, ur, uc, urs),
7218 down: (di, dr, dc, drs),
7219 xs_gate: prescale(x, gcf, gdt).into_owned(),
7220 xs_up: prescale(x, ucf, udt).into_owned(),
7221 down_col: dcf,
7222 w,
7223 q1: gq1,
7224 q4t: gq4 && !gq2 && d.gate_proj.mapped_q4tp().is_none(),
7225 q4tp: gq4 && (gq2 || d.gate_proj.mapped_q4tp().is_some()),
7226 gu_q2: gq2,
7227 swiglu_limit: 0.0,
7228 });
7229 Some(())
7230}
7231
7232fn sparse_ffn_quant(
7239 d: &DenseFfn,
7240 x: &[f32],
7241 active: &[u16],
7242 hidden: usize,
7243 pool: Option<&Pool>,
7244) -> Vec<f32> {
7245 let n = active.len();
7246 let inter = d.gate_proj.rows();
7247 let mut act = vec![0.0f32; n];
7248 let need_scratch = !(d.gate_proj.sparse_col_ok() && d.up_proj.sparse_col_ok());
7251 let compute = |ai: usize| -> f32 {
7252 let idx = active[ai] as usize;
7253 if idx >= inter {
7254 return 0.0; }
7256 let mut s = if need_scratch {
7257 vec![0.0f32; hidden]
7258 } else {
7259 Vec::new()
7260 };
7261 let gate = d.gate_proj.row_dot(idx, x, &mut s);
7262 let up = d.up_proj.row_dot(idx, x, &mut s);
7263 d.act.combine(gate, up)
7264 };
7265 match pool {
7266 Some(p) if n >= 256 => {
7267 let ptr = SendMut(act.as_mut_ptr());
7268 p.run(&|widx, nw| {
7269 let chunk = n.div_ceil(nw);
7270 let (s, e) = (widx * chunk, ((widx + 1) * chunk).min(n));
7271 for ai in s..e {
7272 unsafe { *ptr.at(ai) = compute(ai) };
7273 }
7274 });
7275 }
7276 _ => {
7277 for (ai, a) in act.iter_mut().enumerate() {
7278 *a = compute(ai);
7279 }
7280 }
7281 }
7282 let mut out = vec![0.0f32; hidden];
7284 for (ai, &idx) in active.iter().enumerate() {
7285 let w = act[ai];
7286 if w.abs() >= 1e-12 && (idx as usize) < inter {
7287 d.down_proj.add_col_scaled(idx as usize, w, &mut out);
7288 }
7289 }
7290 out
7291}
7292
7293#[doc(hidden)]
7295pub fn sparse_ffn_quant_for_test(
7296 d: &DenseFfn,
7297 x: &[f32],
7298 active: &[u16],
7299 hidden: usize,
7300) -> Vec<f32> {
7301 sparse_ffn_quant(d, x, active, hidden, None)
7302}
7303
7304fn dequant_dense_f32(d: &DenseFfn) -> (Vec<f32>, Vec<f32>, Vec<f32>) {
7308 let deq = |t: &QTensor| -> Vec<f32> {
7309 let (rows, cols) = (t.rows(), t.cols());
7310 let mut out = vec![0.0f32; rows * cols];
7311 for r in 0..rows {
7312 t.row_f32(r, &mut out[r * cols..(r + 1) * cols]);
7313 }
7314 out
7315 };
7316 (deq(&d.gate_proj), deq(&d.up_proj), deq(&d.down_proj))
7317}
7318
7319struct SendMut(*mut f32);
7321unsafe impl Send for SendMut {}
7322unsafe impl Sync for SendMut {}
7323impl SendMut {
7324 #[inline]
7325 #[allow(clippy::mut_from_ref)]
7328 unsafe fn at(&self, i: usize) -> &mut f32 {
7329 unsafe { &mut *self.0.add(i) }
7330 }
7331}
7332
7333fn moe_route(logits: &[f32], m: &MoeFfn, allowed: Option<&[bool]>) -> (Vec<usize>, Vec<f32>, f32) {
7343 let ne = logits.len();
7344 let p: Vec<f32> = if m.router_sigmoid {
7345 logits.iter().map(|&l| 1.0 / (1.0 + (-l).exp())).collect()
7346 } else {
7347 let mx = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
7348 let mut e: Vec<f32> = logits.iter().map(|&l| (l - mx).exp()).collect();
7349 let s: f32 = e.iter().sum();
7350 for v in &mut e {
7351 *v /= s;
7352 }
7353 e
7354 };
7355 let admit = |e: usize| {
7361 m.mask.as_ref().is_none_or(|mk| mk[e])
7362 && allowed.is_none_or(|a| a.get(e).copied().unwrap_or(false))
7363 };
7364 let mut idx: Vec<usize> = (0..ne).filter(|&e| admit(e)).collect();
7365 match &m.expert_bias {
7367 Some(b) => idx.sort_unstable_by(|&x, &y| {
7368 (p[y] + b[y])
7369 .partial_cmp(&(p[x] + b[x]))
7370 .unwrap()
7371 .then(x.cmp(&y))
7372 }),
7373 None => idx.sort_unstable_by(|&x, &y| p[y].partial_cmp(&p[x]).unwrap().then(x.cmp(&y))),
7374 }
7375 idx.truncate(m.top_k);
7376 if let Some(tau) = m.route_tau {
7380 let total: f32 = idx.iter().map(|&e| p[e]).sum();
7381 if total > 0.0 {
7382 let mut acc = 0.0f32;
7383 let mut keep = idx.len();
7384 for (i, &e) in idx.iter().enumerate() {
7385 acc += p[e];
7386 if acc >= tau * total {
7387 keep = i + 1;
7388 break;
7389 }
7390 }
7391 idx.truncate(keep);
7392 }
7393 }
7394 let wsum: f32 = if m.norm_topk_prob {
7395 let s: f32 = idx.iter().map(|&e| p[e]).sum();
7396 (if m.router_sigmoid { s + 1e-6 } else { s }) / m.routed_scaling
7399 } else {
7400 1.0 / m.routed_scaling
7401 };
7402 (idx, p, wsum)
7403}
7404
7405fn moe_ffn(m: &MoeFfn, x: &[f32], pool: Option<&Pool>, allowed: Option<&[bool]>) -> Vec<f32> {
7408 accumulate_act(m, x, 1);
7409 let ne = m.experts.len();
7410 let mut logits = vec![0.0f32; ne];
7411 m.router.matvec(x, &mut logits, pool);
7412 let (idx, p, wsum) = moe_route(&logits, m, allowed);
7413 {
7414 let mut st = m.stats.borrow_mut();
7415 if st.len() < ne {
7416 st.resize(ne, 0);
7417 }
7418 for &e in &idx {
7419 st[e] += 1;
7420 }
7421 }
7422 if crate::gpu::enabled_here() {
7427 match crate::gpu::probe_arm(crate::gpu::OpClass::Ffn) {
7428 crate::gpu::ProbeArm::Gpu => {
7429 let t0 = std::time::Instant::now();
7430 if let Some(out) = moe_ffn_gpu(m, x, &idx, &p, wsum, pool) {
7431 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
7432 return out;
7433 }
7434 }
7435 crate::gpu::ProbeArm::CpuTimed => {
7436 let t0 = std::time::Instant::now();
7437 let out = crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7438 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
7439 return out;
7440 }
7441 crate::gpu::ProbeArm::Cpu => {
7442 return crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7443 }
7444 }
7445 }
7446 moe_ffn_cpu(m, x, &idx, &p, wsum, pool)
7447}
7448
7449fn graph_note(built: bool) {
7453 use std::sync::atomic::{AtomicBool, Ordering};
7454 if built {
7455 GRAPH_TOK_OK.fetch_add(1, Ordering::Relaxed);
7456 } else {
7457 GRAPH_TOK_MISS.fetch_add(1, Ordering::Relaxed);
7458 }
7459 static SAID: AtomicBool = AtomicBool::new(false);
7460 if !SAID.swap(true, Ordering::Relaxed) {
7461 if built {
7462 tracing::info!("wgpu whole-token graph: ACTIVE");
7463 } else {
7464 tracing::warn!("wgpu whole-token graph refused — per-op path");
7465 }
7466 }
7467}
7468
7469pub static GRAPH_TOK_OK: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7473pub static GRAPH_TOK_MISS: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7474
7475fn moe_batch_enabled() -> bool {
7478 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
7479 *ON.get_or_init(|| std::env::var("CMF_MOE_BATCH").as_deref() != Ok("0"))
7480}
7481
7482fn moe_ffn_cpu_batched(
7488 m: &MoeFfn,
7489 x: &[f32],
7490 idx: &[usize],
7491 p: &[f32],
7492 wsum: f32,
7493 pool: Option<&Pool>,
7494) -> Option<Vec<f32>> {
7495 if idx.is_empty() || !moe_batch_enabled() {
7496 return None;
7497 }
7498 if FFN_PROBE.with(|pr| pr.borrow().is_some()) {
7502 return None;
7503 }
7504 let n = idx.len() + usize::from(m.shared.is_some());
7505 let mut pairs = Vec::with_capacity(n);
7506 let mut downs = Vec::with_capacity(n);
7507 let mut ws = Vec::with_capacity(n);
7508 for &e in idx {
7509 let d = &m.experts[e];
7510 if d.act != Act::Silu {
7511 return None;
7512 }
7513 pairs.push((&d.gate_proj, &d.up_proj));
7514 downs.push(&d.down_proj);
7515 ws.push(p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]));
7516 }
7517 if let Some((se, gate)) = &m.shared {
7520 if se.act != Act::Silu {
7521 return None;
7522 }
7523 let g = gate.as_ref().map_or(1.0, |gate| {
7524 let mut gl = [0.0f32; 1];
7525 gate.matvec(x, &mut gl, pool);
7526 1.0 / (1.0 + (-gl[0]).exp())
7527 });
7528 pairs.push((&se.gate_proj, &se.up_proj));
7529 downs.push(&se.down_proj);
7530 ws.push(g);
7531 }
7532 let inter = pairs[0].0.rows();
7533 let mut gs: Vec<Vec<f32>> = (0..pairs.len()).map(|_| vec![0f32; inter]).collect();
7534 if !QTensor::moe_gate_up_many(&pairs, x, &mut gs, pool) {
7535 return None;
7536 }
7537 let mut out = attention::take_buf(x.len());
7538 if !QTensor::moe_down_many(&downs, &gs, &ws, &mut out, pool) {
7539 attention::recycle_buf(&mut out);
7540 return None;
7541 }
7542 Some(out)
7543}
7544
7545fn moe_ffn_cpu(
7547 m: &MoeFfn,
7548 x: &[f32],
7549 idx: &[usize],
7550 p: &[f32],
7551 wsum: f32,
7552 pool: Option<&Pool>,
7553) -> Vec<f32> {
7554 if let Some(out) = moe_ffn_cpu_batched(m, x, idx, p, wsum, pool) {
7555 return out;
7556 }
7557 let mut out = attention::take_buf(x.len());
7558 for &e in idx {
7559 let mut eo = dense_ffn(&m.experts[e], x, pool);
7560 let w = p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]);
7561 for i in 0..out.len() {
7562 out[i] += w * eo[i];
7563 }
7564 attention::recycle_buf(&mut eo);
7565 }
7566 if let Some((se, gate)) = &m.shared {
7567 let mut so = dense_ffn(se, x, pool);
7568 let g = gate.as_ref().map_or(1.0, |gate| {
7569 let mut gl = [0.0f32; 1];
7570 gate.matvec(x, &mut gl, pool);
7571 1.0 / (1.0 + (-gl[0]).exp())
7572 });
7573 for i in 0..out.len() {
7574 out[i] += g * so[i];
7575 }
7576 attention::recycle_buf(&mut so);
7577 }
7578 out
7579}
7580
7581#[allow(clippy::too_many_arguments)]
7589fn mla_attention(
7590 w: &MlaWeights,
7591 normed: &[f32],
7592 cache: &mut crate::kv_cache::LayerKvCache,
7593 position: usize,
7594 inv_freq: &[f32],
7595 rope_scale: f32,
7596 eps: f64,
7597 pool: Option<&Pool>,
7598) -> Vec<f32> {
7599 let (nh, dr, dn, dv, lora) = (w.nh, w.qk_rope, w.qk_nope, w.v_dim, w.lora);
7600 let hd = dr + dn;
7601 let mut q = vec![0.0f32; nh * hd];
7602 match (&w.q_a, &w.q_a_norm) {
7603 (Some(qa), Some(qn)) => {
7604 let mut t = vec![0.0f32; qa.rows()];
7605 qa.matvec(normed, &mut t, pool);
7606 let tn = inference::rms_norm(&t, qn, eps, NormStyle::Qwen);
7607 w.q_proj.matvec(&tn, &mut q, pool);
7608 }
7609 _ => w.q_proj.matvec(normed, &mut q, pool),
7610 }
7611 let mut ca = vec![0.0f32; lora + dr];
7612 w.kv_a.matvec(normed, &mut ca, pool);
7613 let (c_lat, k_rope) = ca.split_at_mut(lora);
7614 let latn = inference::rms_norm(c_lat, &w.kv_a_norm, eps, NormStyle::Qwen);
7615 let mut kvb = vec![0.0f32; nh * (dn + dv)];
7616 w.kv_b.matvec(&latn, &mut kvb, pool);
7617 if !w.nope {
7618 attention::rope_rotate_scaled(k_rope, position, inv_freq, rope_scale);
7619 }
7620 for h in 0..nh {
7621 if !w.nope {
7622 attention::rope_rotate_scaled(
7623 &mut q[h * hd..h * hd + dr],
7624 position,
7625 inv_freq,
7626 rope_scale,
7627 );
7628 }
7629 }
7630 let mut k = vec![0.0f32; nh * hd];
7631 let mut v = vec![0.0f32; nh * hd];
7632 for h in 0..nh {
7633 k[h * hd..h * hd + dr].copy_from_slice(k_rope);
7634 k[h * hd + dr..(h + 1) * hd].copy_from_slice(&kvb[h * (dn + dv)..h * (dn + dv) + dn]);
7635 v[h * hd..h * hd + dv].copy_from_slice(&kvb[h * (dn + dv) + dn..(h + 1) * (dn + dv)]);
7636 }
7637 cache.append(&k, &v, &vec![true; nh]);
7638 let (ao, mut imp) = attention::attend_all_heads(&q, cache, nh, 1, hd, w.scale, None, 0.0);
7639 attention::recycle_buf(&mut imp);
7640 let mut ov = vec![0.0f32; nh * dv];
7641 for h in 0..nh {
7642 ov[h * dv..(h + 1) * dv].copy_from_slice(&ao[h * hd..h * hd + dv]);
7643 }
7644 let mut out = vec![0.0f32; w.o_proj.rows()];
7645 w.o_proj.matvec(&ov, &mut out, pool);
7646 out
7647}
7648
7649fn dense_moe_ffn(
7656 dm: &DenseMoeFfn,
7657 x_normed: &[f32],
7658 h_raw: &[f32],
7659 eps: f64,
7660 norm_style: NormStyle,
7661 pool: Option<&Pool>,
7662) -> Vec<f32> {
7663 let mut d = dense_ffn(&dm.dense, x_normed, pool);
7664 d = inference::rms_norm(&d, &dm.post_norm_1, eps, norm_style);
7665 let m = &dm.moe;
7666 let ne = m.experts.len();
7667 let mut logits = vec![0.0f32; ne];
7668 if m.router_input_norm {
7669 let ss: f32 = h_raw.iter().map(|v| v * v).sum::<f32>() / h_raw.len() as f32;
7670 let inv = 1.0 / (ss + eps as f32).sqrt();
7671 let xr: Vec<f32> = h_raw.iter().map(|v| v * inv).collect();
7672 m.router.matvec(&xr, &mut logits, pool);
7673 } else {
7674 m.router.matvec(h_raw, &mut logits, pool);
7675 }
7676 let (idx, p, wsum) = moe_route(&logits, m, None);
7677 {
7678 let mut st = m.stats.borrow_mut();
7679 if st.len() < ne {
7680 st.resize(ne, 0);
7681 }
7682 for &e in &idx {
7683 st[e] += 1;
7684 }
7685 }
7686 let x2 = inference::rms_norm(h_raw, &dm.pre_norm_2, eps, norm_style);
7687 let mo = moe_ffn_cpu(m, &x2, &idx, &p, wsum, pool);
7688 let mo = inference::rms_norm(&mo, &dm.post_norm_2, eps, norm_style);
7689 for (di, mi) in d.iter_mut().zip(&mo) {
7690 *di += mi;
7691 }
7692 d
7693}
7694
7695fn moe_gpu_refused(why: &'static str) {
7702 use std::sync::atomic::{AtomicBool, Ordering};
7703 static SAID: AtomicBool = AtomicBool::new(false);
7704 if !SAID.swap(true, Ordering::Relaxed) {
7705 tracing::warn!("MoE GPU block refused ({why}) — experts run on the CPU");
7706 }
7707}
7708
7709fn moe_ffn_gpu(
7710 m: &MoeFfn,
7711 x: &[f32],
7712 idx: &[usize],
7713 p: &[f32],
7714 wsum: f32,
7715 pool: Option<&Pool>,
7716) -> Option<Vec<f32>> {
7717 use crate::gpu::MoeJob;
7718
7719 let mut jobs: Vec<MoeJob> = Vec::with_capacity(idx.len() + 1);
7720 let mut model_ref = None;
7721 for &e in idx {
7722 if moe_push_job(&m.experts[e], x, p[e] / wsum, &mut jobs, &mut model_ref).is_none() {
7723 moe_gpu_refused("push_job(expert)");
7724 return None;
7725 }
7726 }
7727 if let Some((se, gate)) = &m.shared {
7728 let g = gate.as_ref().map_or(1.0, |gate| {
7729 let mut gl = [0.0f32; 1];
7730 gate.matvec(x, &mut gl, pool);
7731 1.0 / (1.0 + (-gl[0]).exp())
7732 });
7733 if moe_push_job(se, x, g, &mut jobs, &mut model_ref).is_none() {
7734 moe_gpu_refused("push_job(shared)");
7735 return None;
7736 }
7737 }
7738 let Some(model) = model_ref else {
7739 moe_gpu_refused("no model_ref");
7740 return None;
7741 };
7742 let hidden = jobs[0].down.1;
7743 let mut out = vec![0.0f32; hidden];
7744 if crate::gpu::moe_block(&model, &jobs, &mut out) {
7745 Some(out)
7746 } else {
7747 moe_gpu_refused("gpu::moe_block");
7748 None
7749 }
7750}
7751
7752fn ffn_forward(
7754 ffn: &FfnKind,
7755 x: &[f32],
7756 pool: Option<&Pool>,
7757 experts_allowed: Option<&[bool]>,
7758) -> Vec<f32> {
7759 match ffn {
7760 FfnKind::Dense(d) => dense_ffn(d, x, pool),
7761 FfnKind::Moe(m) => moe_ffn(m, x, pool, experts_allowed),
7762 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7766 }
7767}
7768
7769fn ffn_forward_pair(
7773 ffn: &FfnKind,
7774 x1: &[f32],
7775 x2: &[f32],
7776 pool: Option<&Pool>,
7777 experts_allowed: Option<&[bool]>,
7778) -> (Vec<f32>, Vec<f32>) {
7779 let d = match ffn {
7780 FfnKind::Dense(d) => d,
7781 FfnKind::Moe(m) => {
7782 return (
7783 moe_ffn(m, x1, pool, experts_allowed),
7784 moe_ffn(m, x2, pool, experts_allowed),
7785 );
7786 }
7787 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7788 };
7789 let inter = d.gate_proj.rows();
7790 FFN_SCRATCH.with(|s| {
7791 let mut s = s.borrow_mut();
7792 let [g1, g2, u1, u2] = &mut *s;
7793 g1.resize(inter, 0.0);
7794 g2.resize(inter, 0.0);
7795 u1.resize(inter, 0.0);
7796 u2.resize(inter, 0.0);
7797 QTensor::matvec2_many(
7800 [&d.gate_proj, &d.up_proj],
7801 x1,
7802 x2,
7803 [g1.as_mut_slice(), u1.as_mut_slice()],
7804 [g2.as_mut_slice(), u2.as_mut_slice()],
7805 pool,
7806 );
7807 for i in 0..inter {
7808 g1[i] = d.act.combine(g1[i], u1[i]);
7809 g2[i] = d.act.combine(g2[i], u2[i]);
7810 }
7811 let mut o1 = attention::take_buf(d.down_proj.rows());
7812 let mut o2 = attention::take_buf(d.down_proj.rows());
7813 d.down_proj.matvec2(g1, g2, &mut o1, &mut o2, pool);
7814 (o1, o2)
7815 })
7816}
7817
7818#[cfg(test)]
7819mod tests {
7820
7821 #[test]
7822 fn cancel_flag_stops_generation() {
7823 let mut p = create_test_pipeline(16, 32, 2, 2, 8, 2, 32);
7824 p.cancel.store(true, std::sync::atomic::Ordering::Relaxed);
7827 let r = p.generate_from_ids(&[1, 2, 3], 8, None, None).unwrap();
7828 assert_eq!(r.finish_reason, "cancelled");
7829 assert!(
7830 r.token_ids.is_empty(),
7831 "no tokens after cancel: {:?}",
7832 r.token_ids
7833 );
7834 let r2 = p.generate_from_ids(&[1, 2, 3], 4, None, None).unwrap();
7836 assert_ne!(r2.finish_reason, "cancelled");
7837 }
7838 use super::*;
7839
7840 #[test]
7846 fn sparse_ffn_quant_equals_dense_with_inactive_zeroed() {
7847 let (hidden, inter) = (16usize, 40usize);
7848 let synth = |n: usize, salt: usize| -> Vec<f32> {
7849 (0..n)
7850 .map(|i| (((i * 37 + salt * 11 + 3) % 101) as f32 / 101.0 - 0.5) * 0.4)
7851 .collect()
7852 };
7853 let d = DenseFfn {
7854 gate_proj: QTensor::from_f32(synth(inter * hidden, 1), inter, hidden),
7855 up_proj: QTensor::from_f32(synth(inter * hidden, 2), inter, hidden),
7856 down_proj: QTensor::from_f32(synth(hidden * inter, 3), hidden, inter),
7857 act: Act::Silu,
7858 };
7859 let x = synth(hidden, 9);
7860 let active: Vec<u16> = (0..inter as u16).filter(|i| i % 3 == 0).collect();
7862
7863 let sparse = sparse_ffn_quant(&d, &x, &active, hidden, None);
7864
7865 let mut g = vec![0.0f32; inter];
7867 d.gate_proj.matvec(&x, &mut g, None);
7868 let mut u = vec![0.0f32; inter];
7869 d.up_proj.matvec(&x, &mut u, None);
7870 let act_set: std::collections::HashSet<u16> = active.iter().copied().collect();
7871 for i in 0..inter {
7872 g[i] = if act_set.contains(&(i as u16)) {
7873 inference::silu(g[i]) * u[i]
7874 } else {
7875 0.0
7876 };
7877 }
7878 let mut reference = vec![0.0f32; hidden];
7879 d.down_proj.matvec(&g, &mut reference, None);
7880
7881 let max_d = sparse
7882 .iter()
7883 .zip(&reference)
7884 .map(|(a, b)| (a - b).abs())
7885 .fold(0.0f32, f32::max);
7886 assert!(max_d < 1e-5, "sparse != dense-zeroed: max|Δ| = {max_d}");
7887 }
7888
7889 fn attach_test_mtp(p: &mut Pipeline) {
7891 let (h, inter, heads, kv, hd) = (
7892 p.hidden_size,
7893 p.intermediate_size,
7894 p.num_heads,
7895 p.num_kv_heads,
7896 p.head_dim,
7897 );
7898 let synth = |n: usize, salt: usize| -> Vec<f32> {
7899 (0..n)
7900 .map(|i| (((i * 29 + salt * 23 + 5) % 101) as f32 / 101.0 - 0.5) * 0.2)
7901 .collect()
7902 };
7903 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
7904 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
7905 };
7906 p.mtp = Some(MtpModule {
7907 enorm: vec![1.0; h],
7908 hnorm: vec![1.0; h],
7909 eh_proj: qt(h, 2 * h, 301),
7910 layer: LayerWeights {
7911 input_norm: vec![1.0; h],
7912 post_norm: vec![1.0; h],
7913 attn_out_norm: None,
7914 ffn_out_norm: None,
7915 layer_scale: None,
7916 ffn: FfnKind::Dense(DenseFfn {
7917 gate_proj: qt(inter, h, 315),
7918 up_proj: qt(inter, h, 316),
7919 down_proj: qt(h, inter, 317),
7920 act: Act::Silu,
7921 }),
7922 attn: AttnKind::Full {
7923 bias: None,
7924 wq: qt(heads * hd, h, 311),
7925 wk: qt(kv * hd, h, 312),
7926 wv: qt(kv * hd, h, 313),
7927 wo: qt(h, heads * hd, 314),
7928 q_norm: None,
7929 k_norm: None,
7930 output_gate: false,
7931 softplus_gate: None,
7932 },
7933 },
7934 final_norm: vec![1.0; h],
7935 kv: crate::kv_cache::LayerKvCache::new(kv, hd),
7936 });
7937 }
7938
7939 #[test]
7940 fn speculative_equals_vanilla_greedy() {
7941 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
7945 let run = |spec: bool| {
7946 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
7947 p.sampler_config.temperature = 0.0;
7948 attach_test_mtp(&mut p);
7949 p.speculative = spec;
7950 let r = p.generate("abcdef", 12, None, None).unwrap();
7951 (r.token_ids, r.mtp_drafted, r.mtp_accepted)
7952 };
7953 let (vanilla, d0, _) = run(false);
7954 let (spec, d1, a1) = run(true);
7955 assert_eq!(d0, 0, "vanilla path must not draft");
7956 assert!(d1 > 0, "speculative path must draft");
7957 assert_eq!(
7958 vanilla, spec,
7959 "speculative must reproduce the exact greedy sequence (accepted {a1}/{d1})"
7960 );
7961 }
7962
7963 #[test]
7964 fn speculative_accepts_constant_oracle() {
7965 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
7967 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
7968 p.sampler_config.temperature = 0.0;
7969 p.sampler_config.repetition_penalty = 1.0;
7970 p.weights.lm_head = QTensor::from_f32(vec![0.01; 64 * 8], 64, 8);
7973 attach_test_mtp(&mut p);
7974 p.speculative = true;
7975 let r = p.generate("abcd", 10, None, None).unwrap();
7976 assert!(r.mtp_drafted > 0);
7977 assert_eq!(
7978 r.mtp_accepted, r.mtp_drafted,
7979 "constant logits → every draft accepted"
7980 );
7981 assert!(r.token_ids.windows(2).all(|w| w[0] == w[1]));
7984 }
7985
7986 #[test]
7987 fn empty_prompt_is_an_error_not_a_panic() {
7988 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
7989 let r = p.generate("", 4, None, None);
7990 assert!(r.is_err(), "empty prompt must be a clean error");
7991 }
7992
7993 #[test]
7994 fn every_token_enters_kv_exactly_once() {
7995 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
7996 p.sampler_config.temperature = 0.0;
7998 let r = p.generate("abc", 2, None, None).unwrap();
7999 assert_eq!(r.prompt_tokens, 3);
8000 assert_eq!(
8004 p.kv_cache.seq_len(),
8005 3 + r.tokens_generated - 1,
8006 "each token must be cached exactly once (v1 cached the last prompt token twice)"
8007 );
8008 }
8009
8010 #[test]
8011 fn generation_is_reproducible_with_seed() {
8012 let run = || {
8013 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8014 p.generate("hello", 8, None, None).unwrap().token_ids
8015 };
8016 assert_eq!(run(), run());
8017 }
8018
8019 #[test]
8020 fn resetting_sampler_restarts_the_seeded_stream() {
8021 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8022 let config = SamplerConfig {
8023 seed: Some(1234),
8024 ..SamplerConfig::default()
8025 };
8026 p.set_sampler_config(config.clone());
8027 let first = p.generate("hello", 8, None, None).unwrap().token_ids;
8028 p.set_sampler_config(config);
8029 let second = p.generate("hello", 8, None, None).unwrap().token_ids;
8030 assert_eq!(first, second);
8031 }
8032
8033 #[test]
8034 fn eviction_bounds_the_cache() {
8035 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
8036 p.kv_cache.max_seq_len = 6;
8037 p.sampler_config.temperature = 0.0;
8038 let _ = p.generate("abcd", 12, None, None).unwrap();
8039 assert!(
8040 p.kv_cache.seq_len() <= 6 + 1,
8041 "cache must stay bounded by max_seq_len (got {})",
8042 p.kv_cache.seq_len()
8043 );
8044 }
8045
8046 #[test]
8047 fn confidence_matches_tokens_and_is_a_probability() {
8048 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8049 p.sampler_config.temperature = 0.0;
8050 p.sampler_config.repetition_penalty = 1.0;
8051 let r = p.generate("abcd", 10, None, None).unwrap();
8052 assert_eq!(
8053 r.token_confidence.len(),
8054 r.token_ids.len(),
8055 "one confidence per emitted token"
8056 );
8057 for &c in &r.token_confidence {
8058 assert!((0.0..=1.0).contains(&c), "confidence out of [0,1]: {c}");
8059 }
8060 let logits = [1.0f32, 3.0, 0.5, 3.0];
8062 let p0 = top1_prob_t(&logits, 1, 1.0);
8063 let p1 = top1_prob_t(&logits, 3, 1.0);
8064 assert!((p0 - p1).abs() < 1e-6, "equal logits → equal prob");
8065 assert!(p0 > 0.0 && p0 < 1.0);
8066 let sharp = top1_prob_t(&logits, 1, 1.0);
8068 let soft = top1_prob_t(&logits, 1, 2.0);
8069 assert!(soft < sharp, "higher temperature lowers peak confidence");
8070 }
8071
8072 #[test]
8073 fn trace_is_opt_in_and_parallels_the_output() {
8074 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8076 p.sampler_config.temperature = 0.0;
8077 p.sampler_config.repetition_penalty = 1.0;
8078 let r = p.generate("abcd", 10, None, None).unwrap();
8079 assert!(r.traces.is_empty(), "trace must be empty unless enabled");
8080
8081 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8083 p.sampler_config.temperature = 0.0;
8084 p.sampler_config.repetition_penalty = 1.0;
8085 p.set_trace(true);
8086 let r = p.generate("abcd", 10, None, None).unwrap();
8087 assert_eq!(r.traces.len(), r.token_ids.len(), "one trace row per token");
8088 for (i, tr) in r.traces.iter().enumerate() {
8089 assert_eq!(tr.t, i, "trace index is sequential");
8090 assert_eq!(tr.token_id, r.token_ids[i], "trace token_id matches output");
8091 assert_eq!(
8092 tr.confidence, r.token_confidence[i],
8093 "trace confidence matches the confidence channel"
8094 );
8095 assert!(tr.active_skill.is_none() && tr.recon.is_none() && !tr.switched);
8097 }
8098 }
8099
8100 #[test]
8101 fn explain_prefill_logits_match_greedy_first_token() {
8102 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8106 p.sampler_config.temperature = 0.0;
8107 p.sampler_config.repetition_penalty = 1.0;
8108 let ids = p.tokenizer.encode("abcd");
8109 let logits = p.prefill_next_logits(&ids, None);
8110 let argmax = logits
8111 .iter()
8112 .enumerate()
8113 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
8114 .unwrap()
8115 .0 as u32;
8116 let r = p.generate("abcd", 1, None, None).unwrap();
8117 assert_eq!(
8118 argmax, r.token_ids[0],
8119 "explain preview must match greedy emit"
8120 );
8121 }
8122
8123 #[test]
8124 fn laguna_shared_expert_is_unconditionally_added() {
8125 let matrix = |values: Vec<f32>| QTensor::from_f32(values, 2, 2);
8126 let identity = || matrix(vec![1.0, 0.0, 0.0, 1.0]);
8127 let zero_dense = || DenseFfn {
8128 gate_proj: matrix(vec![0.0; 4]),
8129 up_proj: matrix(vec![0.0; 4]),
8130 down_proj: matrix(vec![0.0; 4]),
8131 act: Act::Silu,
8132 };
8133 let shared = DenseFfn {
8134 gate_proj: identity(),
8135 up_proj: identity(),
8136 down_proj: identity(),
8137 act: Act::Silu,
8138 };
8139 let x = [1.0, 2.0];
8140 let expected = dense_ffn(&shared, &x, None);
8141 let moe = MoeFfn {
8142 router: QTensor::from_f32(vec![0.0, 0.0], 1, 2),
8143 experts: vec![zero_dense()],
8144 top_k: 1,
8145 norm_topk_prob: true,
8146 router_sigmoid: true,
8147 expert_bias: None,
8148 routed_scaling: 1.0,
8149 route_tau: None,
8150 shared: Some((shared, None)),
8151 stats: std::cell::RefCell::new(Vec::new()),
8152 act_sq: std::cell::RefCell::new(Vec::new()),
8153 act_rows: std::cell::RefCell::new(Vec::new()),
8154 mask: None,
8155 per_expert_scale: None,
8156 router_input_norm: false,
8157 };
8158 let actual = moe_ffn_cpu(&moe, &x, &[0], &[0.0], 1.0, None);
8159 for (actual, expected) in actual.iter().zip(expected) {
8160 assert!((actual - expected).abs() < 1e-6);
8161 }
8162 }
8163}