1use crate::attention::{self, QwenAttnCfg};
10use crate::inference;
11use crate::kv_cache::KvCache;
12use crate::linear_core::{
13 GdnCfg, GdnWeights, ShortConvCfg, ShortConvWeights, VmfPhaseCfg, VmfPhaseWeights, gdn_forward,
14 gdn_pair, short_conv_forward, short_conv_forward_batch, short_conv_pair, vmf_phase_forward,
15 vmf_phase_pair,
16};
17use crate::pool::Pool;
18use crate::qtensor::QTensor;
19use crate::sampler::{self, SamplerConfig, SamplerScratch, SplitMix64};
20use crate::tokenizer::Tokenizer;
21use cortiq_core::mask::TaskMask;
22use cortiq_core::types::NormStyle;
23
24pub static GLOBAL_USE_GPU: std::sync::atomic::AtomicBool =
25 std::sync::atomic::AtomicBool::new(false);
26
27struct ForwardScratch {
31 n1: Vec<f32>,
32 n2: Vec<f32>,
33 p1: Vec<f32>,
34 p2: Vec<f32>,
35}
36
37impl ForwardScratch {
38 fn new(hidden: usize) -> Self {
39 Self {
40 n1: vec![0.0; hidden],
41 n2: vec![0.0; hidden],
42 p1: vec![0.0; hidden],
43 p2: vec![0.0; hidden],
44 }
45 }
46}
47
48pub struct Pipeline {
50 gpu_plan: Option<std::sync::Arc<Vec<(usize, usize, usize)>>>,
55 pub tokenizer: std::sync::Arc<Tokenizer>,
58 pub kv_cache: KvCache,
59 pub sampler_config: SamplerConfig,
60 pub weights: PipelineWeights,
61 pub hidden_size: usize,
62 pub intermediate_size: usize,
63 pub num_heads: usize,
64 pub num_kv_heads: usize,
65 pub head_dim: usize,
66 pub num_layers: usize,
68 pub physical_layers: usize,
70 pub loop_final_norm: bool,
72 pub vocab_size: usize,
73 pub rms_eps: f64,
74 pub rope_base: f32,
75 pub norm_style: NormStyle,
76 pub rotary_dim: usize,
78 pub attention_heads_per_layer: Option<Vec<usize>>,
80 pub vmf_cfg: Option<VmfPhaseCfg>,
82 pub gdn_cfg: Option<GdnCfg>,
84 pub logit_multiplier: Option<f32>,
86 pub cancel: std::sync::Arc<std::sync::atomic::AtomicBool>,
91 pub kv_history: Vec<u32>,
96 pub kda_cfg: Option<crate::linear_core::KdaCfg>,
98 pub g3n: Option<Box<(crate::g3n::G3nGlobals, Vec<crate::g3n::G3nLayer>)>>,
101 pub dsv4: Option<
105 Box<(
106 crate::dsv4::Dsv4Globals,
107 Vec<crate::dsv4::Dsv4Layer>,
108 crate::dsv4::Dsv4Cfg,
109 crate::dsv4::Dsv4State,
110 )>,
111 >,
112 pub dsv4_mtp: Vec<crate::dsv4::Dsv4Mtp>,
116 pub dspark: Option<crate::dsv4::DsparkState>,
118 pub dspark_pending: Vec<(usize, Vec<u32>, bool, usize)>,
121 pub dspark_hist: Vec<usize>,
123 pub dspark_real: Vec<u32>,
127 pub dspark_trunk_picks: Vec<Vec<(usize, Vec<usize>)>>,
131 pub dspark_exp: Vec<(usize, usize, usize, usize)>,
133 pub dspark_draft_ns: u128,
137 pub short_conv_cfg: Option<ShortConvCfg>,
140 pub mtp: Option<MtpModule>,
142 pub speculative: bool,
144 rng: SplitMix64,
145 sampler_scratch: SamplerScratch,
146 pub(crate) inv_freq: std::sync::Arc<Vec<f32>>,
150 ws: ForwardScratch,
154 pool: Option<std::sync::Arc<Pool>>,
156 pub(crate) model: Option<std::sync::Arc<cortiq_core::CmfModel>>,
160 pub(crate) dyn_force_f32: bool,
162 pub(crate) dyn_skill_layers: Vec<Option<Vec<usize>>>,
167 pub(crate) dyn_active: Option<usize>,
173 pub(crate) dyn_blend_loaded: bool,
177 pub(crate) dyn_phi_layer: Option<usize>,
180 dyn_phi_ema: Vec<f32>,
182 dyn_phi_seen: usize,
183 pub dyn_router: Option<crate::swarm::DynRouter>,
186 o1_cfg: Option<crate::nystrom::O1Cfg>,
189 o1_epoch: u64,
192 o1_flags: Vec<bool>,
194 trace: bool,
197 calib_temp: f32,
200 #[cfg_attr(not(target_os = "macos"), allow(dead_code))]
202 graph_kv_id: u64,
203 #[cfg_attr(not(target_os = "macos"), allow(dead_code))]
206 graph_want_logits: bool,
207 graph_logits: Option<Vec<f32>>,
210 pub embed_multiplier: f32,
212 pub attn_scale: f32,
215 pub swa: Option<(usize, usize)>,
218 pub sliding_layers: Option<Vec<bool>>,
221 pub inv_freq_local: Option<std::sync::Arc<Vec<f32>>>,
224 pub rotary_dim_local: Option<usize>,
225 pub rope_scale: f32,
226 pub rope_scale_local: f32,
227 pub global_attn: Option<(usize, usize)>,
230 pub inv_freq_global: Option<std::sync::Arc<Vec<f32>>>,
233 pub attn_v_norm: bool,
235 pub final_softcap: Option<f32>,
237 pub attn_softcap: f32,
239 confidence_on: bool,
243}
244
245#[cfg(target_os = "macos")]
246impl Drop for Pipeline {
247 fn drop(&mut self) {
248 crate::gpu::kv_mirror_drop(self.graph_kv_id);
249 }
250}
251
252pub struct PipelineWeights {
257 pub embed_tokens: QTensor,
259 pub layers: Vec<LayerWeights>,
261 pub lm_head: QTensor,
263 pub final_norm: Vec<f32>,
265}
266
267pub struct LayerWeights {
269 pub input_norm: Vec<f32>,
270 pub post_norm: Vec<f32>,
273 pub attn_out_norm: Option<Vec<f32>>,
276 pub layer_scale: Option<f32>,
278 pub ffn_out_norm: Option<Vec<f32>>,
281 pub ffn: FfnKind,
282 pub attn: AttnKind,
283}
284
285#[derive(Clone, Copy, PartialEq, Debug, Default)]
288pub enum Act {
289 #[default]
290 Silu,
291 GeluTanh,
292 Situ {
295 beta: f32,
296 linear_beta: f32,
297 },
298}
299
300impl Act {
301 pub fn from_arch(name: &str) -> Self {
302 if name == "gelu_tanh" {
303 Self::GeluTanh
304 } else {
305 Self::Silu
306 }
307 }
308
309 pub fn from_arch_full(arch: &cortiq_core::ModelArch) -> Self {
311 match arch.hidden_act.as_str() {
312 "situ" => Self::Situ {
313 beta: arch.activation_situ_beta.unwrap_or(1.0) as f32,
314 linear_beta: arch.activation_situ_linear_beta.unwrap_or(0.0) as f32,
315 },
316 other => Self::from_arch(other),
317 }
318 }
319
320 #[inline]
321 pub fn apply(self, x: f32) -> f32 {
322 match self {
323 Self::Silu => inference::silu(x),
324 Self::GeluTanh => inference::gelu_tanh(x),
325 Self::Situ { beta, .. } => beta * (x / beta).tanh() * (1.0 / (1.0 + (-x).exp())),
326 }
327 }
328
329 #[inline]
332 pub fn combine(self, g: f32, u: f32) -> f32 {
333 match self {
334 Self::Situ { linear_beta, .. } if linear_beta > 0.0 => {
335 self.apply(g) * (linear_beta * (u / linear_beta).tanh())
336 }
337 _ => self.apply(g) * u,
338 }
339 }
340}
341
342pub struct DenseFfn {
344 pub gate_proj: QTensor,
345 pub up_proj: QTensor,
346 pub down_proj: QTensor,
347 pub act: Act,
349}
350
351pub enum FfnKind {
354 Dense(DenseFfn),
355 Moe(MoeFfn),
359 DenseMoe(Box<DenseMoeFfn>),
366}
367
368pub struct DenseMoeFfn {
370 pub dense: DenseFfn,
371 pub moe: MoeFfn,
372 pub post_norm_1: Vec<f32>,
374 pub pre_norm_2: Vec<f32>,
377 pub post_norm_2: Vec<f32>,
379}
380
381pub struct MoeFfn {
382 pub router: QTensor,
384 pub experts: Vec<DenseFfn>,
385 pub top_k: usize,
386 pub norm_topk_prob: bool,
387 pub router_sigmoid: bool,
390 pub expert_bias: Option<Vec<f32>>,
394 pub routed_scaling: f32,
397 pub route_tau: Option<f32>,
403 pub shared: Option<(DenseFfn, Option<QTensor>)>,
406 pub stats: std::cell::RefCell<Vec<u64>>,
410 pub act_sq: std::cell::RefCell<Vec<f64>>,
417 pub act_rows: std::cell::RefCell<Vec<f32>>,
423 pub mask: Option<Vec<bool>>,
428 pub per_expert_scale: Option<Vec<f32>>,
431 pub router_input_norm: bool,
435}
436
437pub enum AttnKind {
440 Full {
442 wq: QTensor,
443 wk: QTensor,
444 wv: QTensor,
445 wo: QTensor,
446 q_norm: Option<Vec<f32>>,
447 k_norm: Option<Vec<f32>>,
448 output_gate: bool,
449 softplus_gate: Option<(QTensor, bool)>,
453 bias: Option<(Vec<f32>, Vec<f32>, Vec<f32>)>,
455 },
456 Linear(VmfPhaseWeights),
458 LinearGdn(GdnWeights),
460 ShortConv(ShortConvWeights),
463 Mla(Box<MlaWeights>),
471 Kda(Box<crate::linear_core::KdaWeights>),
475}
476
477pub struct MlaWeights {
479 pub q_proj: QTensor,
483 pub q_a: Option<QTensor>,
486 pub q_a_norm: Option<Vec<f32>>,
487 pub kv_a: QTensor,
489 pub kv_a_norm: Vec<f32>,
491 pub kv_b: QTensor,
493 pub o_proj: QTensor,
495 pub nh: usize,
496 pub qk_rope: usize,
497 pub qk_nope: usize,
498 pub v_dim: usize,
499 pub lora: usize,
500 pub scale: f32,
502 pub nope: bool,
504}
505
506pub struct MtpModule {
511 pub enorm: Vec<f32>,
512 pub hnorm: Vec<f32>,
513 pub eh_proj: QTensor,
515 pub layer: LayerWeights,
516 pub final_norm: Vec<f32>,
517 pub kv: crate::kv_cache::LayerKvCache,
518}
519
520pub struct GenerateResult {
522 pub text: String,
523 pub token_ids: Vec<u32>,
524 pub prompt_tokens: usize,
525 pub tokens_generated: usize,
526 pub finish_reason: String,
527 pub mtp_drafted: usize,
529 pub mtp_accepted: usize,
530 pub token_confidence: Vec<f32>,
535 pub traces: Vec<TokenTrace>,
538}
539
540#[derive(Clone, Debug)]
545pub struct TokenTrace {
546 pub t: usize,
548 pub token_id: u32,
550 pub confidence: f32,
552 pub active_skill: Option<String>,
554 pub recon: Option<f32>,
558 pub switched: bool,
561}
562
563fn top1_prob_t(logits: &[f32], id: u32, temp: f32) -> f32 {
568 let t = if temp > 1e-3 { temp } else { 1.0 };
569 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
570 let sum: f32 = logits.iter().map(|&v| ((v - max) / t).exp()).sum();
571 if sum > 0.0 {
572 (((logits[id as usize] - max) / t).exp()) / sum
573 } else {
574 0.0
575 }
576}
577
578fn prefill_batched() -> bool {
581 std::env::var("CMF_PREFILL")
582 .map(|v| v != "seq")
583 .unwrap_or(true)
584}
585
586#[derive(Clone, Copy)]
590enum PrefillIn<'a> {
591 Ids(&'a [u32]),
592 Hidden(&'a [f32]),
593}
594
595impl Pipeline {
602 fn can_prefill_batched(&self) -> bool {
603 prefill_batched() && !self.weights.layers.is_empty()
604 }
605}
606
607pub fn prefill_chunk() -> usize {
614 if let Some(n) = std::env::var("CMF_PREFILL_CHUNK")
615 .ok()
616 .and_then(|v| v.parse::<usize>().ok())
617 {
618 return n.max(1);
619 }
620 if cfg!(target_os = "macos") {
621 512
622 } else if cfg!(target_arch = "aarch64") {
623 256
626 } else {
627 48
628 }
629}
630
631pub type TokenCallback = Box<dyn FnMut(&str) -> bool + Send>;
633
634impl Pipeline {
635 #[inline]
639 pub fn phys_layer(&self, virtual_idx: usize) -> usize {
640 virtual_idx % self.physical_layers
641 }
642
643 #[inline]
646 pub fn is_loop_end(&self, virtual_idx: usize) -> bool {
647 self.loop_final_norm && (virtual_idx + 1) % self.physical_layers == 0
648 }
649
650 #[allow(clippy::too_many_arguments)]
652
653 #[cfg(target_os = "macos")]
672 fn graph_prefill_preferred(&self) -> bool {
673 if !crate::gpu::enabled_here()
674 || !crate::gpu::q1_force()
675 || std::env::var("CMF_GPU_BLOCK")
676 .map(|v| v == "0")
677 .unwrap_or(false)
678 {
679 return false;
680 }
681 self.weights
682 .layers
683 .iter()
684 .any(|lw| matches!(&lw.attn, AttnKind::LinearGdn(w) if w.in_proj_qkv.is_q1()))
685 }
686
687 #[cfg(not(target_os = "macos"))]
688 fn graph_prefill_preferred(&self) -> bool {
689 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Prefill);
697 if !graph_on || !crate::gpu::enabled_here() {
698 return false;
699 }
700 if self.o1_active() {
709 return false;
710 }
711 self.weights
712 .layers
713 .iter()
714 .any(|lw| matches!(&lw.attn, AttnKind::LinearGdn(_)))
715 }
716
717 #[cfg(target_os = "macos")]
718 fn q1_graph_gpu(
719 &mut self,
720 start: usize,
721 upto: Option<usize>,
722 position: usize,
723 h: &mut [f32],
724 ) -> usize {
725 let _mt0 = std::time::Instant::now(); use crate::gpu::{AttnGpuLayer, GdnGpuCfg, GdnGpuLayer, GraphDims, MetalFfn, TokenGraph};
727 if self.attn_softcap > 0.0 || !crate::gpu::enabled_here()
729 || !crate::gpu::q1_force()
730 || std::env::var("CMF_GPU_BLOCK")
731 .map(|v| v == "0")
732 .unwrap_or(false)
733 {
734 if std::env::var("CMF_GRAPH_DBG").is_ok() {
735 eprintln!(
736 "block-graph: front gate (softcap={} enabled_here={} q1_force={})",
737 self.attn_softcap > 0.0,
738 crate::gpu::enabled_here(),
739 crate::gpu::q1_force(),
740 );
741 }
742 return start;
743 }
744 if self.swa.is_some()
749 || self.global_attn.is_some()
750 || self.attention_heads_per_layer.is_some()
751 || self.attn_v_norm
752 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
753 || self.weights.layers.iter().any(|lw| {
754 lw.attn_out_norm.is_some()
755 || lw.ffn_out_norm.is_some()
756 || lw.layer_scale.is_some()
757 || matches!(&lw.ffn, FfnKind::Dense(d) if d.act != Act::Silu)
758 })
759 {
760 if std::env::var("CMF_GRAPH_DBG").is_ok() {
761 eprintln!(
762 "block-graph: arch ineligible (swa={} gattn={} hpl={} vnorm={} scale_delta={:.2e})",
763 self.swa.is_some(),
764 self.global_attn.is_some(),
765 self.attention_heads_per_layer.is_some(),
766 self.attn_v_norm,
767 (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs(),
768 );
769 }
770 return start;
771 }
772 let limit = upto
775 .map(|u| u + 1)
776 .unwrap_or(self.num_layers)
777 .min(self.num_layers);
778
779 enum Item<'a> {
780 Gdn {
781 run: Vec<GdnGpuLayer<'a>>,
782 first: usize,
783 },
784 Attn {
785 l: AttnGpuLayer<'a>,
786 li: usize,
787 q_norm: Option<&'a [f32]>,
788 k_norm: Option<&'a [f32]>,
789 output_gate: bool,
790 bias: Option<(&'a [f32], &'a [f32], &'a [f32])>,
791 full_gpu: bool,
794 },
795 }
796
797 let attend_mode = std::env::var("CMF_GPU_ATTEND").unwrap_or_else(|_| "auto".into());
804 let attend_contract = attend_mode != "0"
805 && attend_mode != "off"
806 && self.head_dim % 4 == 0
807 && self.head_dim <= 256
808 && self.rotary_dim >= 2
809 && self.rotary_dim <= self.head_dim
810 && (self.rotary_dim / 2) % 32 == 0
811 && self.num_kv_heads > 0
812 && self.num_heads % self.num_kv_heads == 0;
813
814 let mut plan: Vec<Item> = Vec::new();
815 let mut model_ref: Option<std::sync::Arc<cortiq_core::CmfModel>> = None;
816 let block_diag = std::env::var("CMF_GRAPH_DBG").is_ok();
818 let mut scan = start;
819 while scan < limit {
820 let lw = &self.weights.layers[self.phys_layer(scan)];
821 let ffn = match &lw.ffn {
822 FfnKind::Dense(d) => {
823 let (Some(g), Some(u), Some(dn)) = (
824 d.gate_proj.q1_parts(),
825 d.up_proj.q1_parts(),
826 d.down_proj.q1_parts(),
827 ) else {
828 if block_diag {
829 eprintln!(
830 "block-graph: L{scan} FFN trio not graph-mappable — run ends"
831 );
832 }
833 break;
834 };
835 MetalFfn::Dense {
836 gate: g,
837 up: u,
838 down: dn,
839 }
840 }
841 FfnKind::Moe(m) => {
842 let Some(moe) = metal_moe_graph_parts(m, self.hidden_size) else {
843 if block_diag {
844 eprintln!(
845 "block-graph: L{scan} MoE outside the graph contract — run ends"
846 );
847 }
848 break;
849 };
850 if let QTensor::Mapped { model, .. } = &m.experts[0].gate_proj {
851 model_ref.get_or_insert_with(|| model.clone());
852 }
853 MetalFfn::Moe(moe)
854 }
855 _ => {
856 if block_diag {
857 eprintln!("block-graph: L{scan} non-graph FFN — run ends");
858 }
859 break;
860 }
861 };
862 match &lw.attn {
863 AttnKind::LinearGdn(w) if self.gdn_cfg.is_some() => {
864 let parts = (
865 w.in_proj_qkv.q1_parts(),
866 w.in_proj_z.q1_parts(),
867 w.in_proj_a.f32_parts(),
868 w.in_proj_b.f32_parts(),
869 w.out_proj.q1_parts(),
870 );
871 let (Some(qkv), Some(z), Some(a), Some(b), Some(out)) = parts else {
872 if block_diag {
873 eprintln!(
874 "block-graph: L{scan} GDN parts refused (qkv={} z={} a_f32={} b_f32={} out={})",
875 w.in_proj_qkv.q1_parts().is_some(),
876 w.in_proj_z.q1_parts().is_some(),
877 w.in_proj_a.f32_parts().is_some(),
878 w.in_proj_b.f32_parts().is_some(),
879 w.out_proj.q1_parts().is_some(),
880 );
881 }
882 break;
883 };
884 if let QTensor::Mapped { model, .. } = &w.in_proj_qkv {
885 model_ref.get_or_insert_with(|| model.clone());
886 }
887 let gl = GdnGpuLayer {
888 attn_norm: &lw.input_norm,
889 post_norm: &lw.post_norm,
890 qkv,
891 z,
892 a,
893 b,
894 out,
895 ffn,
896 conv1d: &w.conv1d,
897 a_log: &w.a_log,
898 dt_bias: &w.dt_bias,
899 gnorm: &w.norm,
900 };
901 match plan.last_mut() {
902 Some(Item::Gdn { run, .. }) => run.push(gl),
903 _ => plan.push(Item::Gdn {
904 run: vec![gl],
905 first: scan,
906 }),
907 }
908 }
909 AttnKind::Full {
910 wq,
911 wk,
912 wv,
913 wo,
914 q_norm,
915 k_norm,
916 output_gate,
917 softplus_gate: None,
918 bias,
919 } if !self.kv_cache.layers[scan].o1_sealed()
920 || std::env::var("CMF_O1_METAL").as_deref() == Ok("1") => {
925 let parts = (wq.q1_parts(), wk.q1_parts(), wv.q1_parts(), wo.q1_parts());
926 let (Some(pq), Some(pk), Some(pv), Some(po)) = parts else {
927 break;
928 };
929 if let QTensor::Mapped { model, .. } = wq {
930 model_ref.get_or_insert_with(|| model.clone());
931 }
932 let cache = &self.kv_cache.layers[scan];
933 let o1_metal = cache.o1.is_some()
937 && std::env::var("CMF_O1_METAL").as_deref() == Ok("1")
938 && cache.o1_views().is_some();
939 let full_gpu = attend_contract
940 && cache.mode == crate::kv_cache::KvMode::F32
941 && (cache.o1.is_none() || o1_metal)
942 && bias.is_none()
943 && pq.1 == self.num_heads * self.head_dim * (1 + *output_gate as usize)
944 && pk.1 == self.num_kv_heads * self.head_dim
945 && pv.1 == self.num_kv_heads * self.head_dim
946 && po.2 == self.num_heads * self.head_dim;
947 plan.push(Item::Attn {
948 l: AttnGpuLayer {
949 attn_norm: &lw.input_norm,
950 post_norm: &lw.post_norm,
951 wq: pq,
952 wk: pk,
953 wv: pv,
954 wo: po,
955 ffn,
956 },
957 li: scan,
958 q_norm: q_norm.as_deref(),
959 k_norm: k_norm.as_deref(),
960 output_gate: *output_gate,
961 bias: bias
962 .as_ref()
963 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
964 full_gpu,
965 });
966 }
967 _ => break,
968 }
969 scan += 1;
970 }
971 let Some(model) = model_ref else {
972 if std::env::var("CMF_GRAPH_DBG").is_ok() {
973 eprintln!("q1-graph: no model ref (start {start}, scanned to {scan})");
974 }
975 return start;
976 };
977 if plan.is_empty() {
978 if std::env::var("CMF_GRAPH_DBG").is_ok() {
979 eprintln!("q1-graph: empty plan at layer {start}");
980 }
981 return start;
982 }
983 let has_moe = plan.iter().any(|it| match it {
984 Item::Gdn { run, .. } => run.iter().any(|l| matches!(l.ffn, MetalFfn::Moe(_))),
985 Item::Attn { l, .. } => matches!(l.ffn, MetalFfn::Moe(_)),
986 });
987 let has_gdn = plan.iter().any(|it| matches!(it, Item::Gdn { .. }));
988 let dev_attend = attend_contract
989 && (self.head_dim <= 128
990 || has_moe
991 || (self.head_dim <= 256 && has_gdn)
997 || attend_mode == "force"
998 || attend_mode == "256");
999 if !dev_attend {
1000 for it in &mut plan {
1001 if let Item::Attn { li, full_gpu, .. } = it {
1002 let keep_o1 = self.kv_cache.layers[*li].o1.is_some()
1005 && std::env::var("CMF_O1_METAL").as_deref() == Ok("1");
1006 if !keep_o1 {
1007 *full_gpu = false;
1008 }
1009 }
1010 }
1011 }
1012 if std::env::var("CMF_GRAPH_DBG").is_ok() {
1013 use std::sync::atomic::{AtomicBool, Ordering};
1014 static SAID: AtomicBool = AtomicBool::new(false);
1015 if !SAID.swap(true, Ordering::Relaxed) {
1016 let fg = plan
1017 .iter()
1018 .filter(|it| matches!(it, Item::Attn { full_gpu: true, .. }))
1019 .count();
1020 let att = plan
1021 .iter()
1022 .filter(|it| matches!(it, Item::Attn { .. }))
1023 .count();
1024 eprintln!(
1025 "q1-graph: plan of {} items from layer {start} to {scan} | dev_attend={dev_attend} full_gpu {fg}/{att} | hd={} rd={} nkv={} nh={}",
1026 plan.len(),
1027 self.head_dim,
1028 self.rotary_dim,
1029 self.num_kv_heads,
1030 self.num_heads,
1031 );
1032 }
1033 }
1034 let dims = GraphDims {
1035 hidden: self.hidden_size,
1036 eps: self.rms_eps as f32,
1037 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1038 };
1039 let Some(mut graph) = TokenGraph::new(&model, dims, h) else {
1040 return start;
1041 };
1042 let gcfg = self.gdn_cfg.map(|cfg| GdnGpuCfg {
1043 nv: cfg.num_v_heads,
1044 nk: cfg.num_k_heads,
1045 dk: cfg.key_head_dim,
1046 dv: cfg.value_head_dim,
1047 kk: cfg.conv_kernel,
1048 hidden: self.hidden_size,
1049 inter: self.intermediate_size,
1050 c_dim: cfg.conv_dim(),
1051 eps: cfg.rms_eps as f32,
1052 gemma: self.norm_style == cortiq_core::NormStyle::Gemma,
1053 });
1054 let mut valid = 0usize;
1058 let mut end = start;
1059 crate::gpu::stageprof(1, _mt0.elapsed()); if std::env::var("CMF_PLAN_DUMP").is_ok() {
1061 static ONCE: std::sync::Once = std::sync::Once::new();
1062 ONCE.call_once(|| {
1063 for it in &plan {
1064 match it {
1065 Item::Gdn { first, run } => {
1066 eprintln!("plan: Gdn first={first} len={}", run.len())
1067 }
1068 Item::Attn { li, full_gpu, .. } => {
1069 eprintln!("plan: Attn li={li} full_gpu={full_gpu}")
1070 }
1071 }
1072 }
1073 });
1074 }
1075 for item in &plan {
1076
1077 let ok = match item {
1078 Item::Gdn { run, .. } => gcfg
1079 .as_ref()
1080 .map(|gc| run.iter().all(|l| graph.gdn_ok(l, gc)))
1081 .unwrap_or(false),
1082 Item::Attn { l, .. } => graph.attn_ok(l),
1083 };
1084 if !ok {
1085 if block_diag {
1086 eprintln!(
1087 "block-graph: plan item {} ({}) failed graph preflight",
1088 valid,
1089 match item {
1090 Item::Gdn { run, first } =>
1091 format!("GDN run L{first}+{}", run.len()),
1092 Item::Attn { li, .. } => format!("Attn L{li}"),
1093 }
1094 );
1095 }
1096 break;
1097 }
1098 valid += 1;
1099 end += match item {
1100 Item::Gdn { run, .. } => run.len(),
1101 Item::Attn { .. } => 1,
1102 };
1103 }
1104 plan.truncate(valid);
1105 if plan.is_empty() {
1106 return start;
1107 }
1108
1109 let inv_freq = self.inv_freq.clone();
1110 let pool = self.pool.clone();
1111 let (nh, nkv, hd, hs, rd, eps) = (
1112 self.num_heads,
1113 self.num_kv_heads,
1114 self.head_dim,
1115 self.hidden_size,
1116 self.rotary_dim,
1117 self.rms_eps,
1118 );
1119 let norm_style = self.norm_style;
1120 let gemma = norm_style == cortiq_core::NormStyle::Gemma;
1121 let want = self.gdn_cfg.map(|c| c.state_len()).unwrap_or(0);
1122 let kv_id = self.graph_kv_id;
1123 let mut pending: Vec<(usize, usize)> = Vec::new();
1126 let mut dev_attn: Vec<usize> = Vec::new();
1129 for item in &plan {
1130 let _xt0 = std::time::Instant::now();
1131 let _xkind: u32 = match item { Item::Gdn { .. } => 2, Item::Attn { .. } => 3 };
1132 if self.loop_final_norm {
1134 let item_start = match item {
1135 Item::Gdn { first, .. } => *first,
1136 Item::Attn { li, .. } => *li,
1137 };
1138 if item_start > start && self.is_loop_end(item_start - 1) {
1139 graph.encode_loop_norm(&self.weights.final_norm);
1140 }
1141 }
1142 match item {
1143 Item::Gdn { run, first } => {
1144 for l in &mut self.kv_cache.layers[*first..*first + run.len()] {
1145 if l.linear_state.len() != want {
1146 l.linear_state = vec![0f32; want];
1147 }
1148 }
1149 let ro: Vec<&[f32]> = self.kv_cache.layers[*first..*first + run.len()]
1150 .iter()
1151 .map(|l| l.linear_state.as_slice())
1152 .collect();
1153 let _ig = std::time::Instant::now();
1154 if !graph.encode_gdn_run(run, &ro, gcfg.as_ref().unwrap()) {
1155 tracing::error!("q1 graph: GDN run refused after validation");
1157 return start;
1158 }
1159 graph.commit_kind = 2;
1162 graph.commit();
1163 crate::gpu::stageprof(0, _ig.elapsed());
1164 pending.push((*first, run.len()));
1165 }
1166 Item::Attn {
1167 l,
1168 li,
1169 q_norm,
1170 k_norm,
1171 output_gate,
1172 bias,
1173 full_gpu,
1174 } => {
1175 let _ia = std::time::Instant::now();
1176 if *full_gpu {
1178 let cache = &self.kv_cache.layers[*li];
1179 let o1p = if cache.o1.is_some() {
1180 match cache.o1_views() {
1181 Some(views) => Some(crate::gpu::O1AttnParams {
1182 views,
1183 epoch: self.o1_epoch,
1184 }),
1185 None => None,
1187 }
1188 } else {
1189 None
1190 };
1191 let o1_layer = cache.o1.is_some();
1192 if o1_layer && o1p.is_none() {
1193 }
1195 let cpu_k: Vec<&[f32]> = (0..nkv).map(|g| cache.head_keys(g)).collect();
1196 let cpu_v: Vec<&[f32]> = (0..nkv).map(|g| cache.head_values(g)).collect();
1197 let cpu_stored = if o1_layer { 0 } else { cpu_k[0].len() / hd };
1198 let p = crate::gpu::AttnDeviceParams {
1199 kv_id,
1200 layer: *li,
1201 nh,
1202 nkv,
1203 hd,
1204 rd,
1205 position,
1206 eps: eps as f32,
1207 gemma,
1208 output_gate: *output_gate,
1209 q_norm: *q_norm,
1210 k_norm: *k_norm,
1211 inv_freq: &inv_freq,
1212 cpu_k,
1213 cpu_v,
1214 cpu_stored,
1215 o1: o1p,
1216 };
1217 let o1_bad = o1_layer && p.o1.is_none();
1218 if !o1_bad
1219 && graph.attn_device_ok(l, &p)
1220 && graph.encode_attn_device(l, &p)
1221 {
1222 if p.o1.is_none() {
1224 dev_attn.push(*li);
1225 }
1226 graph.commit_kind = 3;
1227 graph.commit();
1228 crate::gpu::stageprof(_xkind, _xt0.elapsed());
1232 continue;
1233 }
1234 }
1236 graph.encode_attn_prefix(l);
1237 graph.sync();
1238 if !pending.is_empty() {
1239 let idxs: Vec<usize> =
1240 pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1241 let mut outs: Vec<&mut [f32]> = self
1242 .kv_cache
1243 .layers
1244 .iter_mut()
1245 .enumerate()
1246 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1247 .map(|(_, s)| s.linear_state.as_mut_slice())
1248 .collect();
1249 graph.read_states(&mut outs);
1250 }
1251 let mut q_raw = attention::take_buf(l.wq.1);
1252 let mut k = attention::take_buf(l.wk.1);
1253 let mut v = attention::take_buf(l.wv.1);
1254 graph.read_qkv(&mut q_raw, &mut k, &mut v);
1255 let cfg = QwenAttnCfg {
1256 num_heads: nh,
1257 num_kv_heads: nkv,
1258 head_dim: hd,
1259 hidden_size: hs,
1260 position,
1261 inv_freq: &inv_freq,
1262 rotary_dim: rd,
1263 scale: self.attn_scale,
1264 softcap: self.attn_softcap,
1265 window: None,
1266 v_norm: false,
1267 q_norm: *q_norm,
1268 k_norm: *k_norm,
1269 output_gate: *output_gate,
1270 softplus_gate: None,
1271 rope_scale: 1.0,
1272 bias: *bias,
1273 rms_eps: eps,
1274 norm_style,
1275 pool: pool.as_deref(),
1276 };
1277 let mut ao = attention::qwen_attention_core(
1278 q_raw,
1279 k,
1280 v,
1281 &mut self.kv_cache.layers[*li],
1282 &cfg,
1283 );
1284 graph.encode_attn_suffix(l, &ao);
1285 graph.commit();
1288 attention::recycle_buf(&mut ao);
1289 }
1290 }
1291
1292 crate::gpu::stageprof(_xkind, _xt0.elapsed());
1293 }
1294 let mut lm_rows = None;
1299 if self.graph_want_logits
1300 && upto.is_none()
1301 && end == self.num_layers
1302 && std::env::var("CMF_GPU_LMHEAD")
1303 .map(|v| v != "0")
1304 .unwrap_or(true)
1305 {
1306 if let Some(lm) = self.weights.lm_head.q1_parts() {
1307 if graph.lm_head_ok(lm) {
1308 graph.encode_lm_head(&self.weights.final_norm, lm);
1309 lm_rows = Some(lm.1);
1310 }
1311 }
1312 }
1313 let _sy0 = std::time::Instant::now();
1314 graph.sync();
1315 let _rs0 = std::time::Instant::now();
1316 if !pending.is_empty() {
1317 let idxs: Vec<usize> = pending.drain(..).flat_map(|(f, n)| f..f + n).collect();
1318 let mut outs: Vec<&mut [f32]> = self
1319 .kv_cache
1320 .layers
1321 .iter_mut()
1322 .enumerate()
1323 .filter(|(i, _)| idxs.binary_search(i).is_ok())
1324 .map(|(_, s)| s.linear_state.as_mut_slice())
1325 .collect();
1326 graph.read_states(&mut outs);
1327 }
1328 if std::env::var("CMF_GRAPH_HOSTPROF").as_deref() == Ok("1") {
1329 use std::sync::atomic::{AtomicU64, Ordering};
1330 static SY: AtomicU64 = AtomicU64::new(0);
1331 static RS: AtomicU64 = AtomicU64::new(0);
1332 static N: AtomicU64 = AtomicU64::new(0);
1333 SY.fetch_add((_rs0 - _sy0).as_nanos() as u64, Ordering::Relaxed);
1334 RS.fetch_add(_rs0.elapsed().as_nanos() as u64, Ordering::Relaxed);
1335 let n = N.fetch_add(1, Ordering::Relaxed) + 1;
1336 if n % 100 == 0 {
1337 eprintln!(
1338 "postprof: sync-wait {:.1} ms/ток | read_states {:.1} ms/ток ({n})",
1339 SY.load(Ordering::Relaxed) as f64 / n as f64 / 1e6,
1340 RS.load(Ordering::Relaxed) as f64 / n as f64 / 1e6
1341 );
1342 }
1343 }
1344 if let Some(rows) = lm_rows {
1345 crate::gpu::hostprof_encode_done(_mt0);
1346 let mut lg = attention::take_buf(rows.min(self.vocab_size));
1347 graph.read_logits(&mut lg);
1348 crate::gpu::hostprof_total(_mt0);
1349 lg.resize(self.vocab_size, 0.0);
1350 if let Some(c) = self.final_softcap {
1351 for l in lg.iter_mut() {
1352 *l = c * (*l / c).tanh();
1353 }
1354 }
1355 self.graph_logits = Some(lg);
1356 }
1357 graph.finish(h);
1358 for li in dev_attn {
1362 let mut krow = attention::take_buf(nkv * hd);
1363 let mut vrow = attention::take_buf(nkv * hd);
1364 if crate::gpu::kv_mirror_read_last(kv_id, li, nkv, hd, &mut krow, &mut vrow) {
1365 let cache = &mut self.kv_cache.layers[li];
1366 cache.append(&krow, &vrow, &[]);
1367 let n = cache.seq_len;
1368 let mut imp = attention::take_buf(n);
1369 crate::gpu::kv_mirror_take_imp(kv_id, li, &mut imp);
1370 cache.accumulate_imp(&imp);
1371 attention::recycle_buf(&mut imp);
1372 }
1373 attention::recycle_buf(&mut krow);
1374 attention::recycle_buf(&mut vrow);
1375 }
1376 end
1377 }
1378
1379 pub fn new(
1380 tokenizer: Tokenizer,
1381 weights: PipelineWeights,
1382 hidden_size: usize,
1383 intermediate_size: usize,
1384 num_heads: usize,
1385 num_kv_heads: usize,
1386 head_dim: usize,
1387 num_layers: usize,
1388 physical_layers: usize,
1389 loop_final_norm: bool,
1390 vocab_size: usize,
1391 rms_eps: f64,
1392 rope_base: f32,
1393 norm_style: NormStyle,
1394 max_seq_len: usize,
1395 sampler_config: SamplerConfig,
1396 ) -> Self {
1397 let rng = match sampler_config.seed {
1398 Some(s) => SplitMix64::new(s),
1399 None => SplitMix64::from_entropy(),
1400 };
1401 let inv_freq = std::sync::Arc::new(attention::rope_inv_freq(head_dim, rope_base));
1402 let pool = Pool::from_env();
1403 if let Some(p) = &pool {
1404 tracing::info!("worker pool: {} threads", p.n_workers());
1405 }
1406 Self {
1407 gpu_plan: None,
1408 tokenizer: std::sync::Arc::new(tokenizer),
1409 kv_cache: KvCache::new(num_layers, num_kv_heads, head_dim, max_seq_len),
1410 sampler_config,
1411 weights,
1412 hidden_size,
1413 intermediate_size,
1414 num_heads,
1415 num_kv_heads,
1416 head_dim,
1417 num_layers,
1418 physical_layers,
1419 loop_final_norm,
1420 vocab_size,
1421 rms_eps,
1422 rope_base,
1423 norm_style,
1424 rotary_dim: head_dim,
1425 attention_heads_per_layer: None,
1426 vmf_cfg: None,
1427 gdn_cfg: None,
1428 kda_cfg: None,
1429 g3n: None,
1430 dsv4: None,
1431 dsv4_mtp: Vec::new(),
1432 dspark: None,
1433 dspark_pending: Vec::new(),
1434 dspark_hist: Vec::new(),
1435 dspark_real: Vec::new(),
1436 dspark_trunk_picks: Vec::new(),
1437 dspark_exp: Vec::new(),
1438 dspark_draft_ns: 0,
1439 logit_multiplier: None,
1440 cancel: std::sync::Arc::new(std::sync::atomic::AtomicBool::new(false)),
1441 kv_history: Vec::new(),
1442 short_conv_cfg: None,
1443 mtp: None,
1444 speculative: std::env::var("CMF_MTP").map(|v| v != "0").unwrap_or(true),
1445 rng,
1446 sampler_scratch: SamplerScratch::default(),
1447 inv_freq,
1448 ws: ForwardScratch::new(hidden_size),
1449 pool,
1450 model: None,
1451 dyn_force_f32: false,
1452 dyn_skill_layers: Vec::new(),
1453 dyn_active: None,
1454 dyn_blend_loaded: false,
1455 dyn_phi_layer: None,
1456 dyn_phi_ema: Vec::new(),
1457 dyn_phi_seen: 0,
1458 dyn_router: None,
1459 o1_cfg: None,
1460 o1_epoch: 0,
1461 o1_flags: Vec::new(),
1462 trace: false,
1463 calib_temp: 1.0,
1464 confidence_on: true,
1465 embed_multiplier: 1.0,
1466 attn_scale: 1.0 / (head_dim as f32).sqrt(),
1467 swa: None,
1468 sliding_layers: None,
1469 inv_freq_local: None,
1470 rotary_dim_local: None,
1471 rope_scale: 1.0,
1472 rope_scale_local: 1.0,
1473 global_attn: None,
1474 inv_freq_global: None,
1475 attn_v_norm: false,
1476 final_softcap: None,
1477 attn_softcap: 0.0,
1478 graph_want_logits: false,
1479 graph_logits: None,
1480 graph_kv_id: {
1481 static NEXT: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(1);
1482 NEXT.fetch_add(1, std::sync::atomic::Ordering::Relaxed)
1483 },
1484 }
1485 }
1486
1487 pub fn set_o1(&mut self, cfg: Option<crate::nystrom::O1Cfg>) {
1494 self.o1_flags = match &cfg {
1495 Some(c) => {
1496 let mut flags = c.layer_flags(self.num_layers);
1497 for (li, f) in flags.iter_mut().enumerate() {
1498 if *f
1499 && !matches!(
1500 self.weights.layers[self.phys_layer(li)].attn,
1501 AttnKind::Full { .. }
1502 )
1503 {
1504 *f = false;
1505 }
1506 }
1507 flags
1508 }
1509 None => Vec::new(),
1510 };
1511 if let Some(c) = &cfg {
1512 let n = self.o1_flags.iter().filter(|&&f| f).count();
1513 tracing::info!(
1514 "o1 nystrom attention: {n}/{} layer(s), m={} w={} sink={} rect={:?}",
1515 self.num_layers,
1516 c.m,
1517 c.w,
1518 c.sink,
1519 c.rect
1520 );
1521 }
1522 self.o1_cfg = cfg;
1523 }
1524
1525 pub fn o1_active(&self) -> bool {
1527 self.o1_cfg.is_some() && self.o1_flags.iter().any(|&f| f)
1528 }
1529
1530 pub fn o1_begin(&mut self) {
1535 if let Some(c) = &self.o1_cfg {
1536 let (m, w, sink, rect) = (c.m, c.w, c.sink, c.rect);
1537 for (li, &f) in self.o1_flags.iter().enumerate() {
1538 if f {
1539 self.kv_cache.layers[li].o1_begin(m, w, sink, rect);
1540 }
1541 }
1542 }
1543 }
1544
1545 pub fn o1_seal(&mut self) {
1549 self.o1_epoch = self.o1_epoch.wrapping_add(1);
1550 if self.o1_cfg.is_none() {
1551 return;
1552 }
1553 for li in 0..self.num_layers {
1554 if self.o1_flags.get(li).copied().unwrap_or(false) {
1555 self.kv_cache.layers[li].o1_seal(self.num_heads);
1556 }
1557 }
1558 }
1559
1560 pub fn set_trace(&mut self, on: bool) {
1562 self.trace = on;
1563 }
1564
1565 pub fn set_sampler_config(&mut self, config: SamplerConfig) {
1568 self.rng = match config.seed {
1569 Some(seed) => SplitMix64::new(seed),
1570 None => SplitMix64::from_entropy(),
1571 };
1572 self.sampler_config = config;
1573 }
1574
1575 pub fn set_confidence(&mut self, on: bool) {
1580 self.confidence_on = on;
1581 }
1582
1583 pub fn set_calib_temp(&mut self, t: f32) {
1586 self.calib_temp = if t > 1e-3 { t } else { 1.0 };
1587 }
1588
1589 pub fn calib_temp(&self) -> f32 {
1591 self.calib_temp
1592 }
1593
1594 pub fn set_rotary(&mut self, rotary_dim: usize, base: f32) {
1597 self.rotary_dim = rotary_dim.min(self.head_dim);
1598 self.inv_freq = std::sync::Arc::new(attention::rope_inv_freq(self.rotary_dim, base));
1599 }
1600
1601 fn attn_cfg(&self, position: usize) -> QwenAttnCfg<'_> {
1602 QwenAttnCfg {
1603 num_heads: self.num_heads,
1604 num_kv_heads: self.num_kv_heads,
1605 head_dim: self.head_dim,
1606 hidden_size: self.hidden_size,
1607 position,
1608 inv_freq: &self.inv_freq,
1609 rotary_dim: self.rotary_dim,
1610 scale: self.attn_scale,
1611 softcap: self.attn_softcap,
1612 window: None,
1613 v_norm: false,
1614 q_norm: None,
1615 k_norm: None,
1616 output_gate: false,
1617 softplus_gate: None,
1618 rope_scale: self.rope_scale,
1619 bias: None,
1620 rms_eps: self.rms_eps,
1621 norm_style: self.norm_style,
1622 pool: self.pool.as_deref(),
1623 }
1624 }
1625
1626 pub fn generate(
1628 &mut self,
1629 prompt: &str,
1630 max_tokens: usize,
1631 task_mask: Option<&TaskMask>,
1632 on_token: Option<TokenCallback>,
1633 ) -> Result<GenerateResult, String> {
1634 let input_ids = self.tokenizer.with_bos(self.tokenizer.encode(prompt));
1635 self.generate_from_ids(&input_ids, max_tokens, task_mask, on_token)
1636 }
1637
1638 pub fn generate_from_ids(
1646 &mut self,
1647 input_ids: &[u32],
1648 max_tokens: usize,
1649 task_mask: Option<&TaskMask>,
1650 mut on_token: Option<TokenCallback>,
1651 ) -> Result<GenerateResult, String> {
1652 if std::env::var("CMF_TRACE_H").is_ok() {
1653 eprintln!("input_ids: {input_ids:?}");
1654 }
1655 if input_ids.is_empty() {
1656 return Err("empty prompt: nothing to generate from".to_string());
1657 }
1658
1659 let reuse_from = {
1667 let on = !std::env::var("CMF_KV_REUSE").is_ok_and(|v| v == "0");
1668 let h = &self.kv_history;
1669 if on
1670 && task_mask.is_none()
1671 && self.mtp.is_none()
1672 && self.o1_cfg.is_none()
1673 && !h.is_empty()
1674 && h.len() < input_ids.len()
1675 && input_ids[..h.len()] == h[..]
1676 {
1677 h.len()
1678 } else {
1679 0
1680 }
1681 };
1682 if reuse_from == 0 {
1683 self.kv_cache.clear();
1685 self.kv_history.clear();
1686 crate::gpu::graph_kv_reset(self.graph_kv_id);
1687 } else if std::env::var("CMF_PREFILL_PROF").is_ok() {
1688 eprintln!(
1689 "kv-reuse: {} of {} prompt positions already cached",
1690 reuse_from,
1691 input_ids.len()
1692 );
1693 }
1694 crate::gpu::graph_race_begin_generation();
1695 self.o1_begin();
1696
1697 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
1703 let graph_spec = self.speculative
1719 && graph_on
1720 && self.mtp.is_some()
1721 && task_mask.is_none()
1722 && !self.o1_active()
1723 && self.sampler_config.temperature < 1e-6
1724 && self.sampler_config.repetition_penalty == 1.0
1725 && std::env::var("CMF_GRAPH_SPEC").is_ok_and(|v| v != "0");
1726 let pair_pays = self.gdn_cfg.is_none()
1733 || std::env::var("CMF_MTP").as_deref() == Ok("1");
1734 let spec_active = self.speculative
1735 && self.mtp.is_some()
1736 && task_mask.is_none()
1737 && !self.o1_active()
1738 && ((!graph_on && pair_pays) || graph_spec)
1739 && self.sampler_config.temperature < 1e-6;
1740 let mut mtp = if spec_active { self.mtp.take() } else { None };
1743 if std::env::var("CMF_MTP_CHAIN_PROBE").is_ok() {
1744 eprintln!(
1745 "mtp-probe gate: spec_active={spec_active} mtp={} speculative={} graph_on={graph_on} temp_ok={}",
1746 mtp.is_some(),
1747 self.speculative,
1748 self.sampler_config.temperature < 1e-6,
1749 );
1750 }
1751 if let Some(m) = &mut mtp {
1752 m.kv.clear();
1753 }
1754 let mut router = if mtp.is_none() {
1758 self.dyn_router.take()
1759 } else {
1760 None
1761 };
1762 if let Some(r) = &mut router {
1763 r.reset(); self.dyn_phi_seen = 0; let _ = self.set_active_skill(None);
1766 }
1767
1768 let mut all_ids = input_ids.to_vec();
1769 let mut generated = 0usize;
1770 let mut finish_reason = "max_tokens".to_string();
1771 let mut drafted = 0usize;
1772 let mut accepted = 0usize;
1773 let mut confidence: Vec<f32> = Vec::new();
1774 let trace_on = self.trace;
1775 let calib_temp = self.calib_temp;
1776 let mut traces: Vec<TokenTrace> = Vec::new();
1777
1778 let mut hidden = vec![0.0f32; self.hidden_size];
1784 let mut pos = reuse_from;
1785 let fuse_lm = mtp.is_none()
1794 && router.is_none()
1795 && std::env::var("CMF_GPU_LMHEAD").as_deref() != Ok("0");
1796 self.graph_logits = None;
1797 self.graph_want_logits = false;
1798 let _tpf = std::time::Instant::now();
1799 let batch_k = std::env::var("CMF_BATCH_K")
1800 .ok()
1801 .and_then(|v| v.parse::<usize>().ok())
1802 .unwrap_or(0);
1803 while self.dsv4.is_some()
1814 && mtp.is_none()
1815 && pos < input_ids.len()
1816 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1817 {
1818 let end = (pos + prefill_chunk()).min(input_ids.len());
1819 let ids: Vec<u32> = input_ids[pos..end].to_vec();
1820 let mut lg = Vec::new();
1821 if let Some(b) = &mut self.dsv4 {
1822 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
1823 crate::dsv4::forward_chunk(
1824 g,
1825 layers,
1826 &cfg,
1827 st,
1828 &ids,
1829 pos,
1830 &self.inv_freq,
1831 self.pool.as_deref(),
1832 &mut lg,
1833 end == input_ids.len(),
1834 );
1835 }
1836 if end == input_ids.len() {
1837 self.graph_logits = Some(lg);
1838 }
1839 pos = end;
1840 hidden = vec![0.0; self.hidden_size];
1841 }
1842 let dyn_prefill = router.is_some();
1847 let graph_prefill = self.graph_prefill_preferred();
1853 if task_mask.is_none()
1854 && !dyn_prefill
1855 && !graph_prefill
1856 && self.can_prefill_batched()
1857 && self.g3n.is_none()
1858 && input_ids.len() > 2
1859 {
1860 let chunk = prefill_chunk();
1866 let hs = self.hidden_size;
1867 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
1868 let end = (pos + chunk).min(input_ids.len());
1869 let hb = self.prefill_batch(&input_ids[pos..end], pos);
1870 if let Some(m) = &mut mtp {
1871 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1872 .ok()
1873 .and_then(|v| v.parse().ok())
1874 .unwrap_or(0);
1875 for p in pos..end {
1876 if p + 1 < input_ids.len() {
1877 if probe >= 1 && p + 2 < input_ids.len() {
1878 let (d1, mut hx) = self.mtp_step_h(
1882 m,
1883 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1884 input_ids[p + 1],
1885 p,
1886 );
1887 let mut ok = d1 == input_ids[p + 2];
1888 Self::chain_probe_note(0, ok);
1889 let mut d_prev = d1;
1890 let mut extra = 0usize;
1891 for j in 1..probe {
1892 if p + 2 + j >= input_ids.len() {
1893 break;
1894 }
1895 let (dj, hj) =
1896 self.mtp_step_h(m, &hx, d_prev, p + 1 + j);
1897 extra += 1;
1898 ok = ok && dj == input_ids[p + 2 + j];
1899 Self::chain_probe_note(j, ok);
1900 d_prev = dj;
1901 hx = hj;
1902 }
1903 m.kv.truncate_last(extra);
1904 } else {
1905 let _ = self.mtp_step(
1906 m,
1907 &hb[(p - pos) * hs..(p - pos + 1) * hs],
1908 input_ids[p + 1],
1909 p,
1910 );
1911 }
1912 }
1913 }
1914 }
1915 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
1916 pos = end;
1917 }
1918 }
1919 let pair_off = std::env::var("CMF_PAIR").is_ok_and(|v| v == "0");
1920 if task_mask.is_none()
1921 && !dyn_prefill
1922 && !graph_prefill
1923 && !pair_off
1924 && self.pair_supported()
1925 {
1926 while pos + 1 < input_ids.len()
1927 && !self.cancel.load(std::sync::atomic::Ordering::Relaxed)
1928 {
1929 let e1 = self.embed_single(input_ids[pos]);
1930 let e2 = self.embed_single(input_ids[pos + 1]);
1931 let (h1, h2) = self.forward_pair(&e1, &e2, pos);
1932 self.commit_linear_scratch();
1934 if let Some(m) = &mut mtp {
1935 let _ = self.mtp_step(m, &h1, input_ids[pos + 1], pos);
1936 if pos + 2 < input_ids.len() {
1937 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
1938 .ok()
1939 .and_then(|v| v.parse().ok())
1940 .unwrap_or(0);
1941 if probe >= 1 && pos + 3 < input_ids.len() {
1942 let (d1, mut hx) =
1946 self.mtp_step_h(m, &h2, input_ids[pos + 2], pos + 1);
1947 let mut ok = d1 == input_ids[pos + 3];
1948 Self::chain_probe_note(0, ok);
1949 let mut d_prev = d1;
1950 let mut extra = 0usize;
1951 for j in 1..probe {
1952 if pos + 3 + j >= input_ids.len() {
1953 break;
1954 }
1955 let (dj, hj) =
1956 self.mtp_step_h(m, &hx, d_prev, pos + 2 + j);
1957 extra += 1;
1958 ok = ok && dj == input_ids[pos + 3 + j];
1959 Self::chain_probe_note(j, ok);
1960 d_prev = dj;
1961 hx = hj;
1962 }
1963 m.kv.truncate_last(extra);
1964 } else {
1965 let _ = self.mtp_step(m, &h2, input_ids[pos + 2], pos + 1);
1966 }
1967 }
1968 }
1969 hidden = h2;
1970 pos += 2;
1971 }
1972 }
1973 if batch_k > 0
1982 && graph_prefill
1983 && task_mask.is_none()
1984 && !self.o1_active()
1985 && mtp.is_none()
1986 && !dyn_prefill
1987 && pos + 1 < input_ids.len()
1988 {
1989 let hs = self.hidden_size;
1990 let chunk = batch_k;
1991 while pos < input_ids.len() {
1992 let end = (pos + chunk).min(input_ids.len());
1993 let bk = end - pos;
1994 let mut hiddens = vec![0f32; bk * hs];
1995 for (j, &id) in input_ids[pos..end].iter().enumerate() {
1996 hiddens[j * hs..(j + 1) * hs].copy_from_slice(&self.embed_single(id));
1997 }
1998 let positions: Vec<usize> = (pos..end).collect();
1999 let t_chunk = std::time::Instant::now();
2000 let ok_b = self.try_batch_graph_wgpu(&mut hiddens, &positions, bk, None);
2001 if std::env::var("CMF_GRAPH_PROF").is_ok() {
2002 let ms = t_chunk.elapsed().as_secs_f64() * 1000.0;
2003 eprintln!(
2004 "batch-chunk: k={bk} ok={ok_b} {ms:.1} ms ({:.1} tok/s)",
2005 bk as f64 / (ms / 1000.0)
2006 );
2007 }
2008 {
2009 use std::sync::atomic::{AtomicBool, Ordering};
2010 static SAID: AtomicBool = AtomicBool::new(false);
2011 if !SAID.swap(true, Ordering::Relaxed) {
2012 if ok_b {
2013 tracing::info!("batched prefill: ACTIVE (k={bk})");
2014 } else {
2015 tracing::warn!("batched prefill declined — per-position graph");
2016 }
2017 }
2018 }
2019 if ok_b {
2020 hidden.copy_from_slice(&hiddens[(bk - 1) * hs..]);
2021 pos = end;
2022 } else {
2023 break; }
2025 }
2026 }
2027 while pos < input_ids.len() && !self.cancel.load(std::sync::atomic::Ordering::Relaxed) {
2028 self.graph_want_logits = fuse_lm && pos + 1 == input_ids.len();
2029 hidden = self.forward_layers(&self.embed_single(input_ids[pos]), pos, task_mask);
2030 if let Some(m) = &mut mtp {
2031 if pos + 1 < input_ids.len() {
2032 let probe: usize = std::env::var("CMF_MTP_CHAIN_PROBE")
2038 .ok()
2039 .and_then(|v| v.parse().ok())
2040 .unwrap_or(0);
2041 if probe >= 1 && pos + 2 < input_ids.len() {
2042 let (d1, mut hx) = self.mtp_step_h(m, &hidden, input_ids[pos + 1], pos);
2043 let mut ok = d1 == input_ids[pos + 2];
2044 Self::chain_probe_note(0, ok);
2045 let mut d_prev = d1;
2046 let mut extra = 0usize;
2047 for j in 1..probe {
2048 if pos + 2 + j >= input_ids.len() {
2049 break;
2050 }
2051 let (dj, hj) =
2052 self.mtp_step_h(m, &hx, d_prev, pos + 1 + j);
2053 extra += 1;
2054 ok = ok && dj == input_ids[pos + 2 + j];
2055 Self::chain_probe_note(j, ok);
2056 d_prev = dj;
2057 hx = hj;
2058 }
2059 m.kv.truncate_last(extra);
2062 } else {
2063 let _ = self.mtp_step(m, &hidden, input_ids[pos + 1], pos);
2064 }
2065 }
2066 }
2067 pos += 1;
2068 }
2069 if std::env::var("CMF_PREFILL_PROF").is_ok() {
2070 eprintln!(
2071 "prefill: {} tokens in {:.1} ms (batch_k={batch_k})",
2072 input_ids.len(),
2073 _tpf.elapsed().as_secs_f64() * 1000.0
2074 );
2075 }
2076 if self
2079 .cancel
2080 .swap(false, std::sync::atomic::Ordering::Relaxed)
2081 {
2082 self.kv_history.clear();
2083 if let Some(m) = mtp {
2084 self.mtp = Some(m);
2085 }
2086 return Ok(GenerateResult {
2087 text: String::new(),
2088 token_ids: Vec::new(),
2089 prompt_tokens: input_ids.len(),
2090 tokens_generated: 0,
2091 finish_reason: "cancelled".to_string(),
2092 mtp_drafted: 0,
2093 mtp_accepted: 0,
2094 token_confidence: Vec::new(),
2095 traces: Vec::new(),
2096 });
2097 }
2098
2099 self.o1_seal();
2102
2103 macro_rules! commit {
2105 ($id:expr) => {{
2106 all_ids.push($id);
2107 generated += 1;
2108 if self.tokenizer.is_eos($id) {
2109 finish_reason = "stop".to_string();
2110 false
2111 } else {
2112 let token_text = self.tokenizer.decode_token($id);
2113 let mut go = true;
2114 if let Some(ref mut cb) = on_token {
2115 if !cb(&token_text) {
2116 finish_reason = "cancelled".to_string();
2117 go = false;
2118 }
2119 }
2120 go
2121 }
2122 }};
2123 }
2124
2125 let mut next_pos = input_ids.len();
2127 'decode: while generated < max_tokens {
2128 if self
2129 .cancel
2130 .swap(false, std::sync::atomic::Ordering::Relaxed)
2131 {
2132 finish_reason = "cancelled".to_string();
2133 break 'decode;
2134 }
2135 let mut logits = match self.graph_logits.take() {
2136 Some(lg) => lg,
2137 None => {
2138 inference::rms_norm_into(
2139 &hidden,
2140 &self.weights.final_norm,
2141 self.rms_eps,
2142 self.norm_style,
2143 &mut self.ws.n1,
2144 );
2145 self.lm_head_forward(&self.ws.n1)
2146 }
2147 };
2148 let t_next = sampler::sample_with_scratch(
2149 &logits,
2150 &self.sampler_config,
2151 &all_ids,
2152 &mut self.rng,
2153 &mut self.sampler_scratch,
2154 );
2155 if self.confidence_on {
2156 confidence.push(top1_prob_t(&logits, t_next, calib_temp));
2157 }
2158 attention::recycle_buf(&mut logits);
2159 if trace_on {
2160 let skill = router.as_ref().and_then(|r| r.active_id());
2164 traces.push(TokenTrace {
2165 t: generated,
2166 token_id: t_next,
2167 confidence: confidence.last().copied().unwrap_or(0.0),
2168 active_skill: skill,
2169 recon: None,
2170 switched: false,
2171 });
2172 }
2173 if !commit!(t_next) {
2174 break 'decode;
2175 }
2176 if generated >= max_tokens {
2177 break 'decode;
2178 }
2179
2180 if self.kv_cache.needs_eviction() {
2181 static SAID: std::sync::Once = std::sync::Once::new();
2187 SAID.call_once(|| {
2188 tracing::warn!(
2189 "KV cache full at {} positions — evicting half; quality \
2190 will degrade. Raise CMF_MAX_SEQ.",
2191 self.kv_cache.max_seq_len,
2192 );
2193 });
2194 let keep = (self.kv_cache.max_seq_len / 2).max(1);
2195 self.kv_cache.evict(keep);
2196 }
2197
2198 match &mut mtp {
2199 #[cfg(feature = "gpu")]
2201 Some(m) if graph_spec && generated + 1 < max_tokens && next_pos > 0 => {
2202 if let Some((extra, n_pos, new_h)) = self.graph_spec_step(
2203 m,
2204 &hidden,
2205 t_next,
2206 next_pos,
2207 &mut drafted,
2208 &mut accepted,
2209 ) {
2210 next_pos = n_pos;
2211 hidden = new_h;
2212 let mut stopped = false;
2213 for &id in &extra {
2214 if self.confidence_on {
2215 confidence.push(0.0);
2216 }
2217 if !commit!(id) {
2218 stopped = true;
2219 break;
2220 }
2221 }
2222 if stopped {
2223 break 'decode;
2224 }
2225 continue 'decode;
2226 }
2227 hidden = self.forward_layers(&self.embed_single(t_next), next_pos, task_mask);
2229 next_pos += 1;
2230 continue 'decode;
2231 }
2232 Some(m) if !graph_spec && generated + 1 < max_tokens => {
2234 let draft = self.mtp_step(m, &hidden, t_next, next_pos - 1);
2235 drafted += 1;
2236 let emb1 = self.embed_single(t_next);
2237 let emb2 = self.embed_single(draft);
2238 let (h1, h2) = self.forward_pair(&emb1, &emb2, next_pos);
2239
2240 inference::rms_norm_into(
2241 &h1,
2242 &self.weights.final_norm,
2243 self.rms_eps,
2244 self.norm_style,
2245 &mut self.ws.n1,
2246 );
2247 let mut logits1 = self.lm_head_forward(&self.ws.n1);
2248 let t_after = sampler::sample_with_scratch(
2249 &logits1,
2250 &self.sampler_config,
2251 &all_ids,
2252 &mut self.rng,
2253 &mut self.sampler_scratch,
2254 );
2255 if self.confidence_on {
2256 confidence.push(top1_prob_t(&logits1, t_after, calib_temp));
2257 }
2258 attention::recycle_buf(&mut logits1);
2259 if trace_on {
2260 traces.push(TokenTrace {
2263 t: generated,
2264 token_id: t_after,
2265 confidence: confidence.last().copied().unwrap_or(0.0),
2266 active_skill: None,
2267 recon: None,
2268 switched: false,
2269 });
2270 }
2271 let stop = !commit!(t_after);
2272
2273 if t_after == draft {
2274 accepted += 1;
2275 self.commit_linear_scratch();
2276 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2277 hidden = h2;
2278 next_pos += 2;
2279 } else {
2280 for layer in &mut self.kv_cache.layers {
2282 layer.truncate_last(1);
2283 }
2284 if !stop {
2285 let _ = self.mtp_step(m, &h1, t_after, next_pos);
2286 hidden = self.forward_layers(
2287 &self.embed_single(t_after),
2288 next_pos + 1,
2289 None,
2290 );
2291 }
2292 next_pos += 2;
2293 }
2294 if stop {
2295 break 'decode;
2296 }
2297 }
2298 _ => {
2300 #[cfg(feature = "gpu")]
2305 if Self::dsv4_spec_on() && self.dsv4.is_some() {
2306 static SAID: std::sync::Once = std::sync::Once::new();
2307 SAID.call_once(|| {
2308 eprintln!(
2309 "dsv4-spec гейт: mtp={} mask={} router={} trace={} temp={} rep={} ",
2310 !self.dsv4_mtp.is_empty(),
2311 task_mask.is_none(),
2312 router.is_none(),
2313 !trace_on,
2314 self.sampler_config.temperature < 1e-6,
2315 self.sampler_config.repetition_penalty == 1.0,
2316 );
2317 });
2318 }
2319 #[cfg(feature = "gpu")]
2320 if Self::dsv4_spec_on()
2321 && self.dsv4.is_some()
2322 && !self.dsv4_mtp.is_empty()
2323 && task_mask.is_none()
2324 && router.is_none()
2325 && !trace_on
2326 && self.sampler_config.temperature < 1e-6
2327 && self.sampler_config.repetition_penalty == 1.0
2328 && generated + 1 < max_tokens
2329 && all_ids.len() >= 2
2330 {
2331 let tip_token = all_ids[all_ids.len() - 2];
2332 if let Some((extra, n_pos)) = self.dsv4_spec_step(
2333 tip_token,
2334 t_next,
2335 next_pos,
2336 &mut drafted,
2337 &mut accepted,
2338 ) {
2339 next_pos = n_pos;
2340 let mut stopped = false;
2341 for &id in &extra {
2342 if self.confidence_on {
2343 confidence.push(0.0);
2344 }
2345 if !commit!(id) {
2346 stopped = true;
2347 break;
2348 }
2349 }
2350 if stopped {
2351 break 'decode;
2352 }
2353 continue 'decode;
2354 }
2355 }
2356 self.graph_want_logits = fuse_lm;
2357 let mut t_fwd = t_next;
2363 let pure_greedy = self.sampler_config.temperature < 1e-6
2364 && self.sampler_config.repetition_penalty == 1.0
2365 && self.sampler_config.suppress_tokens.is_empty();
2366 let burst_k = std::env::var("CMF_MULTISTEP")
2371 .ok()
2372 .and_then(|v| v.parse::<usize>().ok())
2373 .unwrap_or(0);
2374 if pure_greedy
2375 && burst_k >= 1
2376 && fuse_lm
2377 && task_mask.is_none()
2378 && router.is_none()
2379 && !trace_on
2380 && !self.confidence_on
2381 {
2382 let mut stopped = false;
2383 loop {
2384 let room = max_tokens.saturating_sub(generated);
2385 if room <= 2 {
2386 break;
2387 }
2388 let k = burst_k.min(room - 1);
2389 if k < 1 {
2390 break;
2391 }
2392 let Some(ids) = self.try_multi_burst(t_fwd, next_pos, k) else {
2393 break;
2394 };
2395 next_pos += k;
2396 for &id in &ids {
2397 if !commit!(id) {
2398 stopped = true;
2399 break;
2400 }
2401 }
2402 if stopped {
2403 break;
2404 }
2405 t_fwd = *ids.last().unwrap();
2406 }
2407 if stopped {
2408 break 'decode;
2409 }
2410 }
2411 hidden = self.forward_layers(&self.embed_single(t_fwd), next_pos, task_mask);
2412 next_pos += 1;
2413 if let Some(r) = &mut router {
2416 let phi = self.dyn_phi_ema.clone();
2417 let decision = r.step(&phi, generated);
2418 if let Some(new_active) = decision {
2419 let _ = self.set_active_skill(new_active);
2420 }
2421 if trace_on {
2424 if let Some(last) = traces.last_mut() {
2425 let e = r.last_best_e();
2426 last.recon = e.is_finite().then_some(e);
2427 last.switched = decision.is_some();
2428 }
2429 }
2430 }
2431 }
2432 }
2433 }
2434
2435 self.graph_want_logits = false;
2436 self.graph_logits = None;
2437 if router.is_some() {
2439 let _ = self.set_active_skill(None);
2440 }
2441 self.dyn_router = router.or(self.dyn_router.take());
2442 self.mtp = mtp.or(self.mtp.take());
2443
2444 let output_ids = &all_ids[input_ids.len()..];
2445 let forwarded = input_ids.len() + output_ids.len().saturating_sub(1);
2449 self.kv_history = all_ids[..forwarded.min(all_ids.len())].to_vec();
2450 confidence.truncate(output_ids.len()); traces.truncate(output_ids.len());
2452 Ok(GenerateResult {
2453 text: self.tokenizer.decode(output_ids),
2454 token_ids: output_ids.to_vec(),
2455 prompt_tokens: input_ids.len(),
2456 tokens_generated: generated,
2457 finish_reason,
2458 mtp_drafted: drafted,
2459 mtp_accepted: accepted,
2460 token_confidence: confidence,
2461 traces,
2462 })
2463 }
2464
2465 fn mtp_step(
2469 &mut self,
2470 m: &mut MtpModule,
2471 hidden: &[f32],
2472 next_token: u32,
2473 position: usize,
2474 ) -> u32 {
2475 self.mtp_step_h(m, hidden, next_token, position).0
2476 }
2477
2478 fn chain_probe_note(depth: usize, prefix_ok: bool) {
2482 use std::sync::Mutex;
2483 static T: Mutex<Vec<(u64, u64)>> = Mutex::new(Vec::new());
2484 let mut t = T.lock().unwrap();
2485 if t.len() <= depth {
2486 t.resize(depth + 1, (0, 0));
2487 }
2488 t[depth].0 += 1;
2489 t[depth].1 += prefix_ok as u64;
2490 if depth == 0 && t[0].0 % 128 == 0 {
2491 let line: Vec<String> = t
2492 .iter()
2493 .enumerate()
2494 .map(|(d, (n, k))| format!("d{}={:.0}%({n})", d + 1, 100.0 * *k as f64 / (*n).max(1) as f64))
2495 .collect();
2496 eprintln!("mtp-chain: {}", line.join(" "));
2497 }
2498 }
2499
2500 fn mtp_step_h(
2504 &mut self,
2505 m: &mut MtpModule,
2506 hidden: &[f32],
2507 next_token: u32,
2508 position: usize,
2509 ) -> (u32, Vec<f32>) {
2510 let e = self.embed_single(next_token);
2514 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2515 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2516 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2517 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2518 let mut x = vec![0.0f32; self.hidden_size];
2519 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2520
2521 let lw = &m.layer;
2523 inference::rms_norm_into(
2524 &x,
2525 &lw.input_norm,
2526 self.rms_eps,
2527 self.norm_style,
2528 &mut self.ws.n1,
2529 );
2530 let attn = match &lw.attn {
2531 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2533 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2534 AttnKind::Full {
2535 wq,
2536 wk,
2537 wv,
2538 wo,
2539 q_norm,
2540 k_norm,
2541 output_gate,
2542 softplus_gate,
2543 bias,
2544 } => {
2545 let mut cfg = self.attn_cfg(position);
2546 cfg.q_norm = q_norm.as_deref();
2547 cfg.k_norm = k_norm.as_deref();
2548 cfg.output_gate = *output_gate;
2549 cfg.softplus_gate = softplus_gate
2550 .as_ref()
2551 .map(|(gate, per_head)| (gate, *per_head));
2552 cfg.bias = bias
2553 .as_ref()
2554 .map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2555 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2556 }
2557 AttnKind::Linear(_) | AttnKind::LinearGdn(_) | AttnKind::ShortConv(_) => {
2558 unreachable!("MTP block is full attention")
2559 }
2560 };
2561 for (i, &a) in attn.iter().enumerate() {
2562 x[i] += a;
2563 }
2564 inference::rms_norm_into(
2565 &x,
2566 &lw.post_norm,
2567 self.rms_eps,
2568 self.norm_style,
2569 &mut self.ws.p1,
2570 );
2571 let ffn = ffn_forward(&lw.ffn, &self.ws.p1, self.pool.as_deref(), None);
2572 for (i, &f) in ffn.iter().enumerate() {
2573 x[i] += f;
2574 }
2575
2576 inference::rms_norm_into(
2577 &x,
2578 &m.final_norm,
2579 self.rms_eps,
2580 self.norm_style,
2581 &mut self.ws.n1,
2582 );
2583 let mut lg = self.lm_head_forward(&self.ws.n1);
2584 let draft = sampler::argmax(&lg);
2585 attention::recycle_buf(&mut lg);
2586 (draft, x)
2587 }
2588
2589 fn mtp_warm(&mut self, m: &mut MtpModule, hidden: &[f32], next_token: u32, position: usize) {
2593 let e = self.embed_single(next_token);
2594 let mut cat = vec![0.0f32; 2 * self.hidden_size];
2595 let (cat_e, cat_h) = cat.split_at_mut(self.hidden_size);
2596 inference::rms_norm_into(&e, &m.enorm, self.rms_eps, self.norm_style, cat_e);
2597 inference::rms_norm_into(hidden, &m.hnorm, self.rms_eps, self.norm_style, cat_h);
2598 let mut x = vec![0.0f32; self.hidden_size];
2599 m.eh_proj.matvec(&cat, &mut x, self.pool.as_deref());
2600 inference::rms_norm_into(&x, &m.layer.input_norm, self.rms_eps, self.norm_style, &mut self.ws.n1);
2601 let attn = match &m.layer.attn {
2602 AttnKind::Full { wq, wk, wv, wo, q_norm, k_norm, output_gate, softplus_gate, bias } => {
2603 let mut cfg = self.attn_cfg(position);
2604 cfg.q_norm = q_norm.as_deref();
2605 cfg.k_norm = k_norm.as_deref();
2606 cfg.output_gate = *output_gate;
2607 cfg.softplus_gate = softplus_gate.as_ref().map(|(g, p)| (g, *p));
2608 cfg.bias = bias.as_ref().map(|(q, k, v)| (q.as_slice(), k.as_slice(), v.as_slice()));
2609 attention::qwen_attention(&self.ws.n1, wq, wk, wv, wo, &mut m.kv, &cfg)
2610 }
2611 _ => return,
2612 };
2613 let _ = attn;
2614 }
2615
2616 #[cfg(feature = "gpu")]
2623 #[allow(clippy::too_many_arguments)]
2624 fn graph_spec_step(
2625 &mut self,
2626 m: &mut MtpModule,
2627 hidden: &[f32],
2628 t_next: u32,
2629 next_pos: usize,
2630 drafted: &mut usize,
2631 accepted: &mut usize,
2632 ) -> Option<(Vec<u32>, usize, Vec<f32>)> {
2633 let k_spec: usize = std::env::var("CMF_GRAPH_SPEC_K")
2639 .ok()
2640 .and_then(|v| v.parse().ok())
2641 .filter(|&v| (1..=8).contains(&v))
2642 .unwrap_or(3);
2643 if next_pos == 0 {
2644 return None;
2645 }
2646 let t_round = std::time::Instant::now();
2647 let subs = || crate::gpu_wgpu::SUBMITS.load(std::sync::atomic::Ordering::Relaxed);
2663 let sub0 = subs();
2664 let mut drafts = Vec::with_capacity(k_spec);
2669 let (d1, mut hx) = self.mtp_step_h(m, hidden, t_next, next_pos - 1);
2670 drafts.push(d1);
2671 for j in 1..k_spec {
2672 let (dj, hj) = self.mtp_step_h(m, &hx, drafts[j - 1], next_pos - 1 + j);
2673 drafts.push(dj);
2674 hx = hj;
2675 }
2676 *drafted += k_spec;
2677 let t_draft = t_round.elapsed();
2678 let sub_draft = subs();
2679 let b = k_spec + 1;
2682 let mut hiddens = vec![0.0f32; b * self.hidden_size];
2683 for (i, &t) in std::iter::once(&t_next).chain(drafts.iter()).enumerate() {
2684 let e = self.embed_single(t);
2685 hiddens[i * self.hidden_size..(i + 1) * self.hidden_size].copy_from_slice(&e);
2686 }
2687 let positions: Vec<usize> = (next_pos..next_pos + b).collect();
2688 let (lm_gw, lm_rows) = {
2689 let (_, i, kind, rs) = self.weights.lm_head.graph_weight()?;
2690 (
2691 crate::gpu::GraphW { idx: i, kind, row_scale: rs, data: &[] },
2692 self.weights.lm_head.rows(),
2693 )
2694 };
2695 let mut logits = Vec::new();
2696 let final_norm = self.weights.final_norm.clone();
2697 let ok = self.try_batch_graph_wgpu(
2698 &mut hiddens,
2699 &positions,
2700 b,
2701 Some(crate::gpu::SpecTail {
2702 lm: lm_gw,
2703 lm_rows,
2704 final_norm: &final_norm,
2705 logits_out: &mut logits,
2706 }),
2707 );
2708 if !ok {
2709 m.kv.truncate_last(k_spec);
2712 return None;
2713 }
2714 let t_verify = t_round.elapsed();
2715 let sub_verify = subs();
2716 let ids: Vec<u32> = (0..b)
2718 .map(|i| sampler::argmax(&logits[i * lm_rows..(i + 1) * lm_rows]))
2719 .collect();
2720 let mut a = 0usize;
2721 while a < k_spec && ids[a] == drafts[a] {
2722 a += 1;
2723 }
2724 if a + 1 < b {
2726 crate::gpu::gdn_spec_restore(self.graph_kv_id, a);
2727 }
2728 *accepted += a;
2729 m.kv.truncate_last(k_spec.saturating_sub(1));
2740 let warm_off = std::env::var("CMF_SPEC_WARM").is_ok_and(|v| v == "0");
2741 if !warm_off {
2742 for j in 0..a {
2743 let row = &hiddens[j * self.hidden_size..(j + 1) * self.hidden_size];
2744 let row = row.to_vec();
2745 self.mtp_warm(m, &row, ids[j], next_pos + j);
2746 }
2747 }
2748 let mut row = logits[a * lm_rows..(a + 1) * lm_rows].to_vec();
2750 row.resize(self.vocab_size, 0.0);
2751 if let Some(c) = self.final_softcap {
2752 for l in row.iter_mut() {
2753 *l = c * (*l / c).tanh();
2754 }
2755 }
2756 self.graph_logits = Some(row);
2757 let new_hidden = hiddens[a * self.hidden_size..(a + 1) * self.hidden_size].to_vec();
2758 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
2764 let end = subs();
2765 eprintln!(
2766 "spec-round: draft {:.1} ms/{} sub | verify {:.1} ms/{} sub | \
2767 commit {:.1} ms/{} sub (accepted {a} of {k_spec})",
2768 t_draft.as_secs_f64() * 1e3,
2769 sub_draft - sub0,
2770 (t_verify - t_draft).as_secs_f64() * 1e3,
2771 sub_verify - sub_draft,
2772 (t_round.elapsed() - t_verify).as_secs_f64() * 1e3,
2773 end - sub_verify,
2774 );
2775 }
2776 Some((drafts[..a].to_vec(), next_pos + a + 1, new_hidden))
2777 }
2778
2779 pub fn measure_pair_fusion(&mut self, iters: usize) -> (f64, f64) {
2788 if !self.pair_supported() {
2789 return (0.0, 0.0);
2790 }
2791 let emb1 = self.embed_single(1);
2792 let emb2 = self.embed_single(2);
2793 let pos = self.kv_cache.seq_len();
2794
2795 let t0 = std::time::Instant::now();
2796 for _ in 0..iters {
2797 let _ = self.forward_layers(&emb1, pos, None);
2798 let _ = self.forward_layers(&emb2, pos + 1, None);
2799 for l in &mut self.kv_cache.layers {
2800 l.truncate_last(2);
2801 }
2802 }
2803 let singles_ms = t0.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2804
2805 let t1 = std::time::Instant::now();
2806 for _ in 0..iters {
2807 let _ = self.forward_pair(&emb1, &emb2, pos);
2808 for l in &mut self.kv_cache.layers {
2809 l.truncate_last(2);
2810 }
2811 }
2812 let pair_ms = t1.elapsed().as_secs_f64() * 1000.0 / iters as f64;
2813 (singles_ms, pair_ms)
2814 }
2815
2816 fn pair_supported(&self) -> bool {
2824 !self.weights.layers.is_empty()
2831 && self.g3n.is_none()
2832 && !self
2833 .weights
2834 .layers
2835 .iter()
2836 .any(|lw| matches!(&lw.attn, AttnKind::Mla(_) | AttnKind::Kda(_)))
2837 }
2838
2839 fn forward_pair(
2840 &mut self,
2841 emb1: &[f32],
2842 emb2: &[f32],
2843 position: usize,
2844 ) -> (Vec<f32>, Vec<f32>) {
2845 let mut h1 = emb1.to_vec();
2846 let mut h2 = emb2.to_vec();
2847 let (_nkv, _hd, hs, _rd, eps) = (
2848 self.num_kv_heads,
2849 self.head_dim,
2850 self.hidden_size,
2851 self.rotary_dim,
2852 self.rms_eps,
2853 );
2854 let pool = self.pool.clone();
2855
2856 for li in 0..self.num_layers {
2857 let lw = &self.weights.layers[self.phys_layer(li)];
2858 inference::rms_norm_into(
2861 &h1,
2862 &lw.input_norm,
2863 self.rms_eps,
2864 self.norm_style,
2865 &mut self.ws.n1,
2866 );
2867 inference::rms_norm_into(
2868 &h2,
2869 &lw.input_norm,
2870 self.rms_eps,
2871 self.norm_style,
2872 &mut self.ws.n2,
2873 );
2874
2875 let (a1, a2) = match &lw.attn {
2876 AttnKind::Mla(_) => unreachable!("MLA has no MTP/pair path"),
2877 AttnKind::Kda(_) => unreachable!("KDA has no MTP/pair path"),
2878 AttnKind::Linear(w) => {
2879 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
2880 let layer = &mut self.kv_cache.layers[li];
2881 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2882 vmf_phase_pair(
2883 &self.ws.n1,
2884 &self.ws.n2,
2885 w,
2886 &cfg,
2887 state,
2888 scratch,
2889 self.pool.as_deref(),
2890 )
2891 }
2892 AttnKind::LinearGdn(w) => {
2893 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
2894 let layer = &mut self.kv_cache.layers[li];
2895 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2896 gdn_pair(
2897 &self.ws.n1,
2898 &self.ws.n2,
2899 w,
2900 &cfg,
2901 state,
2902 scratch,
2903 self.pool.as_deref(),
2904 )
2905 }
2906 AttnKind::ShortConv(w) => {
2907 let cfg = self
2908 .short_conv_cfg
2909 .expect("short-conv layer without short_conv_cfg");
2910 let layer = &mut self.kv_cache.layers[li];
2911 let (state, scratch) = (&mut layer.linear_state, &mut layer.linear_scratch);
2912 short_conv_pair(
2913 &self.ws.n1,
2914 &self.ws.n2,
2915 w,
2916 &cfg,
2917 state,
2918 scratch,
2919 self.pool.as_deref(),
2920 )
2921 }
2922 AttnKind::Full {
2923 wq,
2924 wk,
2925 wv,
2926 wo,
2927 q_norm,
2928 k_norm,
2929 output_gate,
2930 softplus_gate,
2931 bias,
2932 } => {
2933 let inv_freq_l = self.layer_inv_freq(li);
2934 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
2935 let cfg = QwenAttnCfg {
2936 num_heads: self.layer_num_heads(li),
2937 num_kv_heads: nkv_l,
2938 head_dim: hd_l,
2939 hidden_size: hs,
2940 position,
2941 inv_freq: &inv_freq_l,
2942 rotary_dim: rd_l,
2943 scale: self.attn_scale,
2944 softcap: self.attn_softcap,
2945 window: self.layer_window(li),
2946 v_norm: self.attn_v_norm,
2947 q_norm: q_norm.as_deref(),
2948 k_norm: k_norm.as_deref(),
2949 output_gate: *output_gate,
2950 softplus_gate: softplus_gate
2951 .as_ref()
2952 .map(|(gate, per_head)| (gate, *per_head)),
2953 rope_scale: self.layer_rope_scale(li),
2954 bias: bias
2955 .as_ref()
2956 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
2957 rms_eps: eps,
2958 norm_style: self.norm_style,
2959 pool: pool.as_deref(),
2960 };
2961 attention::qwen_attention_pair(
2962 &self.ws.n1,
2963 &self.ws.n2,
2964 wq,
2965 wk,
2966 wv,
2967 wo,
2968 &mut self.kv_cache.layers[li],
2969 &cfg,
2970 )
2971 }
2972 };
2973 let (a1, a2) = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
2974 Some(w) => (
2975 inference::rms_norm(&a1, w, self.rms_eps, self.norm_style),
2976 inference::rms_norm(&a2, w, self.rms_eps, self.norm_style),
2977 ),
2978 None => (a1, a2),
2979 };
2980 for i in 0..self.hidden_size {
2981 h1[i] += a1[i];
2982 h2[i] += a2[i];
2983 }
2984 let (mut a1, mut a2) = (a1, a2);
2985 attention::recycle_buf(&mut a1);
2986 attention::recycle_buf(&mut a2);
2987
2988 let lw = &self.weights.layers[self.phys_layer(li)];
2989 inference::rms_norm_into(
2990 &h1,
2991 &lw.post_norm,
2992 self.rms_eps,
2993 self.norm_style,
2994 &mut self.ws.p1,
2995 );
2996 inference::rms_norm_into(
2997 &h2,
2998 &lw.post_norm,
2999 self.rms_eps,
3000 self.norm_style,
3001 &mut self.ws.p2,
3002 );
3003 let (f1, f2) = match &lw.ffn {
3004 FfnKind::DenseMoe(dm) => (
3007 dense_moe_ffn(
3008 dm,
3009 &self.ws.p1,
3010 &h1,
3011 self.rms_eps,
3012 self.norm_style,
3013 self.pool.as_deref(),
3014 ),
3015 dense_moe_ffn(
3016 dm,
3017 &self.ws.p2,
3018 &h2,
3019 self.rms_eps,
3020 self.norm_style,
3021 self.pool.as_deref(),
3022 ),
3023 ),
3024 _ => ffn_forward_pair(
3025 &lw.ffn,
3026 &self.ws.p1,
3027 &self.ws.p2,
3028 self.pool.as_deref(),
3029 None,
3030 ),
3031 };
3032 let (f1, f2) = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
3033 Some(w) => (
3034 inference::rms_norm(&f1, w, self.rms_eps, self.norm_style),
3035 inference::rms_norm(&f2, w, self.rms_eps, self.norm_style),
3036 ),
3037 None => (f1, f2),
3038 };
3039 for i in 0..self.hidden_size {
3040 h1[i] += f1[i];
3041 h2[i] += f2[i];
3042 }
3043 let (mut f1, mut f2) = (f1, f2);
3044 attention::recycle_buf(&mut f1);
3045 attention::recycle_buf(&mut f2);
3046 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
3047 for i in 0..self.hidden_size {
3048 h1[i] *= sc;
3049 h2[i] *= sc;
3050 }
3051 }
3052 if self.is_loop_end(li) && li + 1 < self.num_layers {
3054 h1 = inference::rms_norm(
3055 &h1,
3056 &self.weights.final_norm,
3057 self.rms_eps,
3058 self.norm_style,
3059 );
3060 h2 = inference::rms_norm(
3061 &h2,
3062 &self.weights.final_norm,
3063 self.rms_eps,
3064 self.norm_style,
3065 );
3066 }
3067 }
3068 (h1, h2)
3069 }
3070
3071 fn commit_linear_scratch(&mut self) {
3073 for layer in &mut self.kv_cache.layers {
3074 if !layer.linear_scratch.is_empty() {
3075 std::mem::swap(&mut layer.linear_state, &mut layer.linear_scratch);
3076 layer.linear_scratch.clear();
3077 }
3078 }
3079 }
3080
3081 pub fn forward_ids(
3084 &mut self,
3085 ids: &[u32],
3086 task_mask: Option<&TaskMask>,
3087 ) -> Result<Vec<f32>, String> {
3088 if ids.is_empty() {
3089 return Err("empty id sequence".to_string());
3090 }
3091 self.kv_cache.clear();
3092 self.kv_history.clear();
3093 self.o1_begin();
3094 let mut hidden = vec![0.0f32; self.hidden_size];
3095 let mut pos = 0usize;
3096 if self.can_prefill_batched() && !self.graph_prefill_preferred() && ids.len() > 2 {
3104 let chunk = prefill_chunk();
3108 let hs = self.hidden_size;
3109 while pos < ids.len() {
3110 let end = (pos + chunk).min(ids.len());
3111 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3112 hidden.copy_from_slice(&hb[(end - pos - 1) * hs..]);
3113 pos = end;
3114 }
3115 }
3116 if task_mask.is_none()
3125 && !self.graph_prefill_preferred()
3126 && !std::env::var("CMF_PAIR").is_ok_and(|v| v == "0")
3127 && self.pair_supported()
3128 {
3129 while pos + 1 < ids.len() {
3130 let e1 = self.embed_single(ids[pos]);
3131 let e2 = self.embed_single(ids[pos + 1]);
3132 let (_, h2) = self.forward_pair(&e1, &e2, pos);
3133 self.commit_linear_scratch();
3134 hidden = h2;
3135 pos += 2;
3136 }
3137 }
3138 while pos < ids.len() {
3139 hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3140 pos += 1;
3141 }
3142 self.o1_seal();
3146 let normed = inference::rms_norm(
3147 &hidden,
3148 &self.weights.final_norm,
3149 self.rms_eps,
3150 self.norm_style,
3151 );
3152 Ok(self.lm_head_forward(&normed))
3153 }
3154
3155 pub fn ppl_ids(&mut self, ids: &[u32]) -> f64 {
3162 let (nll, cnt) = self.nll_ids_from(ids, 0);
3163 (nll / cnt.max(1) as f64).exp()
3164 }
3165
3166 pub fn probe_ffn_mass(&mut self, ids: &[u32]) -> Vec<Vec<f64>> {
3171 self.kv_cache.clear();
3172 self.kv_history.clear();
3173 FFN_PROBE.with(|p| {
3174 *p.borrow_mut() = Some(vec![vec![0f64; self.intermediate_size]; self.num_layers]);
3175 });
3176 crate::gpu::cpu_scope(|| {
3177 for (pos, &id) in ids.iter().enumerate() {
3178 let emb = self.embed_single(id);
3179 let _ = self.forward_layers(&emb, pos, None);
3180 }
3181 });
3182 self.kv_cache.clear();
3183 self.kv_history.clear();
3184 FFN_PROBE
3185 .with(|p| p.borrow_mut().take())
3186 .unwrap_or_default()
3187 }
3188
3189 pub fn ppl_ids_masked(&mut self, ids: &[u32], mask: &TaskMask) -> f64 {
3193 self.kv_cache.clear();
3194 self.kv_history.clear();
3195 let mut nll = 0f64;
3196 let mut cnt = 0usize;
3197 let mut hidden = vec![0f32; self.hidden_size];
3198 for (pos, &id) in ids.iter().enumerate() {
3199 if pos > 0 {
3200 inference::rms_norm_into(
3201 &hidden,
3202 &self.weights.final_norm,
3203 self.rms_eps,
3204 self.norm_style,
3205 &mut self.ws.n1,
3206 );
3207 let mut logits = self.lm_head_forward(&self.ws.n1);
3208 let max = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
3209 let sum: f64 = logits.iter().map(|&v| ((v - max) as f64).exp()).sum();
3210 let p = ((logits[id as usize] - max) as f64).exp() / sum.max(1e-300);
3211 nll -= p.max(1e-300).ln();
3212 cnt += 1;
3213 attention::recycle_buf(&mut logits);
3214 }
3215 let emb = self.embed_single(id);
3216 hidden = self.forward_layers(&emb, pos, Some(mask));
3217 }
3218 self.kv_cache.clear();
3219 self.kv_history.clear();
3220 (nll / cnt.max(1) as f64).exp()
3221 }
3222
3223 pub fn nll_ids_masked(
3242 &mut self,
3243 ids: &[u32],
3244 start: usize,
3245 task_mask: Option<&TaskMask>,
3246 ) -> (f64, usize) {
3247 self.nll_ids_inner(ids, start, task_mask)
3248 }
3249
3250 pub fn nll_ids_from(&mut self, ids: &[u32], start: usize) -> (f64, usize) {
3251 self.nll_ids_inner(ids, start, None)
3252 }
3253
3254 fn nll_ids_inner(
3255 &mut self,
3256 ids: &[u32],
3257 start: usize,
3258 task_mask: Option<&TaskMask>,
3259 ) -> (f64, usize) {
3260 self.kv_cache.clear();
3261 self.kv_history.clear();
3262 let mut nll = 0f64;
3263 let mut cnt = 0usize;
3264 if self.can_prefill_batched() {
3265 const CHUNK: usize = 128;
3271 const LM_SUB: usize = 32;
3272 let n = ids.len().saturating_sub(1);
3273 let hs = self.hidden_size;
3274 let rows = self.weights.lm_head.rows();
3275 let mut pos = 0usize;
3276 while pos < n {
3277 let end = (pos + CHUNK).min(n);
3278 let bsz = end - pos;
3279 let hb = self.prefill_batch_masked(&ids[pos..end], pos, task_mask);
3280 let mut k0 = 0usize;
3281 while k0 < bsz {
3282 let k1 = (k0 + LM_SUB).min(bsz);
3283 let sb = k1 - k0;
3284 if pos + k1 <= start {
3287 k0 = k1;
3288 continue;
3289 }
3290 let mut normed = vec![0.0f32; sb * hs];
3291 for k in 0..sb {
3292 let r = inference::rms_norm(
3293 &hb[(k0 + k) * hs..(k0 + k + 1) * hs],
3294 &self.weights.final_norm,
3295 self.rms_eps,
3296 self.norm_style,
3297 );
3298 normed[k * hs..(k + 1) * hs].copy_from_slice(&r);
3299 }
3300 let mut logits = vec![0.0f32; sb * rows];
3301 self.weights
3302 .lm_head
3303 .matmat(&normed, sb, &mut logits, self.pool.as_deref());
3304 for k in 0..sb {
3305 if pos + k0 + k < start {
3306 continue;
3307 }
3308 let lg = &mut logits[k * rows..k * rows + self.vocab_size.min(rows)];
3309 if let Some(mu) = self.logit_multiplier {
3310 for v in lg.iter_mut() {
3311 *v *= mu;
3312 }
3313 }
3314 if let Some(c) = self.final_softcap {
3318 for v in lg.iter_mut() {
3319 *v = c * (*v / c).tanh();
3320 }
3321 }
3322 let lg = &logits[k * rows..k * rows + self.vocab_size.min(rows)];
3323 let target = ids[pos + k0 + k + 1] as usize;
3324 let max = lg.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3325 let lse: f64 = lg
3326 .iter()
3327 .map(|&v| ((v - max) as f64).exp())
3328 .sum::<f64>()
3329 .ln()
3330 + max as f64;
3331 nll += lse - lg[target] as f64;
3332 cnt += 1;
3333 if std::env::var("CMF_PPL_TRACE").is_ok() {
3334 let top = lg
3335 .iter()
3336 .enumerate()
3337 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3338 .map(|(i, _)| i)
3339 .unwrap_or(0);
3340 eprintln!(
3341 "BTRACE pos {} target {} nll {:.4} top {} lg_t {:.3} lg_top {:.3}",
3342 pos + k0 + k,
3343 target,
3344 lse - lg[target] as f64,
3345 top,
3346 lg[target],
3347 lg[top]
3348 );
3349 }
3350 }
3351 k0 = k1;
3352 }
3353 pos = end;
3354 }
3355 self.kv_cache.clear();
3356 self.kv_history.clear();
3357 return (nll, cnt);
3358 }
3359 for pos in 0..ids.len().saturating_sub(1) {
3360 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, task_mask);
3361 let out_of_band = self.graph_logits.take();
3369 if pos < start {
3370 continue;
3371 }
3372 let logits = match out_of_band {
3373 Some(lg) => lg,
3374 None => {
3375 let normed = inference::rms_norm(
3376 &hidden,
3377 &self.weights.final_norm,
3378 self.rms_eps,
3379 self.norm_style,
3380 );
3381 self.lm_head_forward(&normed)
3385 }
3386 };
3387 let target = ids[pos + 1] as usize;
3388 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3389 let lse: f64 = logits
3390 .iter()
3391 .map(|&v| ((v - max) as f64).exp())
3392 .sum::<f64>()
3393 .ln()
3394 + max as f64;
3395 let tok_nll = lse - logits[target] as f64;
3396 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3397 let top = logits
3398 .iter()
3399 .enumerate()
3400 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3401 .map(|(i, _)| i)
3402 .unwrap_or(0);
3403 eprintln!(
3404 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3405 logits[target], logits[top]
3406 );
3407 }
3408 nll += tok_nll;
3409 cnt += 1;
3410 }
3411 self.kv_cache.clear();
3412 self.kv_history.clear();
3413 (nll, cnt)
3414 }
3415
3416 pub fn nll_ids_o1(&mut self, ids: &[u32], prefill: usize) -> (f64, usize) {
3432 self.kv_cache.clear();
3433 self.kv_history.clear();
3434 self.o1_begin();
3435 let n = ids.len().saturating_sub(1);
3436 let p = prefill.min(n);
3437 let mut pos = 0usize;
3439 if self.can_prefill_batched() {
3440 const CHUNK: usize = 128;
3441 while pos < p {
3442 let end = (pos + CHUNK).min(p);
3443 let _ = self.prefill_batch(&ids[pos..end], pos);
3444 pos = end;
3445 }
3446 } else {
3447 while pos < p {
3448 let _ = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3449 pos += 1;
3450 }
3451 }
3452 self.o1_seal();
3453
3454 let mut nll = 0f64;
3455 let mut cnt = 0usize;
3456 for pos in p..n {
3457 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3458 let normed = inference::rms_norm(
3459 &hidden,
3460 &self.weights.final_norm,
3461 self.rms_eps,
3462 self.norm_style,
3463 );
3464 let logits = self.lm_head_forward(&normed);
3468 let target = ids[pos + 1] as usize;
3469 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3470 let lse: f64 = logits
3471 .iter()
3472 .map(|&v| ((v - max) as f64).exp())
3473 .sum::<f64>()
3474 .ln()
3475 + max as f64;
3476 let tok_nll = lse - logits[target] as f64;
3477 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3478 let top = logits
3479 .iter()
3480 .enumerate()
3481 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3482 .map(|(i, _)| i)
3483 .unwrap_or(0);
3484 eprintln!(
3485 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3486 logits[target], logits[top]
3487 );
3488 }
3489 nll += tok_nll;
3490 cnt += 1;
3491 }
3492 self.kv_cache.clear();
3493 self.kv_history.clear();
3494 (nll, cnt)
3495 }
3496
3497 pub fn calib_ids(&mut self, ids: &[u32], temps: &[f32]) -> (Vec<bool>, Vec<Vec<f32>>) {
3505 self.kv_cache.clear();
3506 self.kv_history.clear();
3507 let n = ids.len().saturating_sub(1);
3508 let mut correct = Vec::with_capacity(n);
3509 let mut pmax = Vec::with_capacity(n);
3510 for pos in 0..n {
3511 let emb = self.embed_single(ids[pos]);
3512 let hidden = self.forward_layers(&emb, pos, None);
3513 let normed = inference::rms_norm(
3514 &hidden,
3515 &self.weights.final_norm,
3516 self.rms_eps,
3517 self.norm_style,
3518 );
3519 let logits = self.lm_head_forward(&normed);
3523 let target = ids[pos + 1] as usize;
3524 let (mut amax, mut mval) = (0usize, f32::NEG_INFINITY);
3525 for (i, &v) in logits.iter().enumerate() {
3526 if v > mval {
3527 mval = v;
3528 amax = i;
3529 }
3530 }
3531 correct.push(amax == target);
3532 let row: Vec<f32> = temps
3533 .iter()
3534 .map(|&t| {
3535 let tt = t.max(1e-3);
3536 let s: f32 = logits.iter().map(|&v| ((v - mval) / tt).exp()).sum();
3537 1.0 / s.max(1e-12) })
3539 .collect();
3540 pmax.push(row);
3541 }
3542 self.kv_cache.clear();
3543 self.kv_history.clear();
3544 (correct, pmax)
3545 }
3546
3547 pub fn ppl_ids_dynamic(&mut self, ids: &[u32]) -> (f64, usize) {
3554 let mut router = match self.dyn_router.take() {
3555 Some(r) => r,
3556 None => return (self.ppl_ids(ids), 0),
3557 };
3558 router.reset();
3559 self.dyn_phi_seen = 0;
3560 let _ = self.set_active_skill(None);
3561
3562 self.kv_cache.clear();
3563
3564 self.kv_history.clear();
3565 let mut nll = 0f64;
3566 let mut cnt = 0usize;
3567 for pos in 0..ids.len().saturating_sub(1) {
3568 let hidden = self.forward_layers(&self.embed_single(ids[pos]), pos, None);
3569 let normed = inference::rms_norm(
3570 &hidden,
3571 &self.weights.final_norm,
3572 self.rms_eps,
3573 self.norm_style,
3574 );
3575 let logits = self.lm_head_forward(&normed);
3579 let target = ids[pos + 1] as usize;
3580 let max = logits.iter().fold(f32::NEG_INFINITY, |m, &v| m.max(v));
3581 let lse: f64 = logits
3582 .iter()
3583 .map(|&v| ((v - max) as f64).exp())
3584 .sum::<f64>()
3585 .ln()
3586 + max as f64;
3587 let tok_nll = lse - logits[target] as f64;
3588 if std::env::var("CMF_PPL_TRACE").is_ok() && pos < 48 {
3589 let top = logits
3590 .iter()
3591 .enumerate()
3592 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
3593 .map(|(i, _)| i)
3594 .unwrap_or(0);
3595 eprintln!(
3596 "pos {pos:3} tgt {target:6} nll {tok_nll:7.3} | top1 {top:6} lg[t]={:.2} lg[top]={:.2}",
3597 logits[target], logits[top]
3598 );
3599 }
3600 nll += tok_nll;
3601 cnt += 1;
3602 let phi = self.dyn_phi_ema.clone();
3604 if let Some(new_active) = router.step(&phi, pos) {
3605 let _ = self.set_active_skill(new_active);
3606 }
3607 }
3608 let switches = router.switches.len();
3609 let _ = self.set_active_skill(None);
3610 self.dyn_router = Some(router);
3611 self.kv_cache.clear();
3612 self.kv_history.clear();
3613 ((nll / cnt.max(1) as f64).exp(), switches)
3614 }
3615
3616 pub fn probe_phi(&mut self, ids: &[u32], layer: usize) -> Vec<f32> {
3618 self.kv_cache.clear();
3619 self.kv_history.clear();
3620 let mut acc = vec![0f32; self.hidden_size];
3621 for (pos, &id) in ids.iter().enumerate() {
3622 let h = self.forward_layers_upto(&self.embed_single(id), pos, None, Some(layer));
3623 for (a, v) in acc.iter_mut().zip(&h) {
3624 *a += v;
3625 }
3626 }
3627 let n = ids.len().max(1) as f32;
3628 for a in acc.iter_mut() {
3629 *a /= n;
3630 }
3631 self.kv_cache.clear();
3632 self.kv_history.clear();
3633 acc
3634 }
3635
3636 fn prefill_batch(&mut self, ids: &[u32], start_pos: usize) -> Vec<f32> {
3642 self.prefill_batch_masked(ids, start_pos, None)
3643 }
3644
3645 fn prefill_batch_masked(
3651 &mut self,
3652 ids: &[u32],
3653 start_pos: usize,
3654 task_mask: Option<&TaskMask>,
3655 ) -> Vec<f32> {
3656 self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, usize::MAX)
3657 }
3658
3659 fn prefill_batch_span(
3665 &mut self,
3666 input: PrefillIn<'_>,
3667 start_pos: usize,
3668 task_mask: Option<&TaskMask>,
3669 from: usize,
3670 upto_excl: usize,
3671 ) -> Vec<f32> {
3672 let hs = self.hidden_size;
3673 let b = match input {
3674 PrefillIn::Ids(ids) => ids.len(),
3675 PrefillIn::Hidden(hb) => hb.len() / hs,
3676 };
3677 let upto_excl = upto_excl.min(self.num_layers);
3678 let mut h: Vec<f32>;
3682 let mut h_ready;
3683 match input {
3684 PrefillIn::Ids(_) => {
3685 h = vec![0.0; b * hs];
3686 h_ready = false;
3687 }
3688 PrefillIn::Hidden(hb) => {
3689 h = hb.to_vec();
3690 h_ready = true;
3691 }
3692 }
3693 let fill_h = |h: &mut Vec<f32>, me: &Self| {
3694 if let PrefillIn::Ids(ids) = input {
3695 for (bi, &id) in ids.iter().enumerate() {
3696 let e = me.embed_single(id);
3697 h[bi * hs..(bi + 1) * hs].copy_from_slice(&e);
3698 }
3699 }
3700 };
3701 let (_nkv, _hd, _rd, eps) = (
3702 self.num_kv_heads,
3703 self.head_dim,
3704 self.rotary_dim,
3705 self.rms_eps,
3706 );
3707 let pool = self.pool.clone();
3708 let norm_style = self.norm_style;
3709
3710 #[cfg(target_os = "macos")]
3711 let mut chunk_skip_until = 0usize;
3712 for li in from..upto_excl {
3713 crate::gpu::set_layer(li as i64); #[cfg(target_os = "macos")]
3720 if task_mask.is_none() {
3721 if li < chunk_skip_until {
3722 continue;
3723 }
3724 if !h_ready && li == 0 && self.weights.embed_tokens.q8_row_parts().is_none() {
3730 fill_h(&mut h, self);
3731 h_ready = true;
3732 }
3733 let ids_for_embed = match input {
3734 PrefillIn::Ids(ids) => (!h_ready && li == 0).then_some(ids),
3735 PrefillIn::Hidden(_) => None,
3736 };
3737 let end = self.chunk_run_gpu(li, &mut h, b, start_pos, ids_for_embed, upto_excl);
3738 if end > li {
3739 h_ready = true;
3740 chunk_skip_until = end;
3741 if self.is_loop_end(end - 1) && end < self.num_layers {
3744 for bi in 0..b {
3745 let normed = inference::rms_norm(
3746 &h[bi * hs..(bi + 1) * hs],
3747 &self.weights.final_norm,
3748 eps,
3749 norm_style,
3750 );
3751 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
3752 }
3753 }
3754 continue;
3755 }
3756 }
3757 if !h_ready {
3758 fill_h(&mut h, self);
3759 h_ready = true;
3760 }
3761 let lw = &self.weights.layers[self.phys_layer(li)];
3762 match &lw.attn {
3764 AttnKind::Kda(w) => {
3765 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
3767 let mut normed = vec![0.0f32; b * hs];
3768 for bi in 0..b {
3769 inference::rms_norm_into(
3770 &h[bi * hs..(bi + 1) * hs],
3771 &lw.input_norm,
3772 eps,
3773 norm_style,
3774 &mut normed[bi * hs..(bi + 1) * hs],
3775 );
3776 }
3777 let attn = crate::linear_core::kda_forward_batch(
3778 &normed,
3779 b,
3780 w,
3781 &cfg,
3782 &mut self.kv_cache.layers[li].linear_state,
3783 pool.as_deref(),
3784 );
3785 for (dst, &a) in h.iter_mut().zip(&attn) {
3786 *dst += a;
3787 }
3788 }
3789 AttnKind::LinearGdn(w) => {
3790 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
3792 let mut normed = vec![0.0f32; b * hs];
3793 for bi in 0..b {
3794 let r = inference::rms_norm(
3795 &h[bi * hs..(bi + 1) * hs],
3796 &lw.input_norm,
3797 eps,
3798 norm_style,
3799 );
3800 normed[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3801 }
3802 let attn = crate::linear_core::gdn_forward_batch(
3803 &normed,
3804 b,
3805 w,
3806 &cfg,
3807 &mut self.kv_cache.layers[li].linear_state,
3808 pool.as_deref(),
3809 );
3810 for (dst, &a) in h.iter_mut().zip(&attn) {
3811 *dst += a;
3812 }
3813 }
3814 AttnKind::ShortConv(w) => {
3815 let cfg = self
3818 .short_conv_cfg
3819 .expect("short-conv layer without short_conv_cfg");
3820 let mut normed = vec![0.0f32; b * hs];
3821 for bi in 0..b {
3822 inference::rms_norm_into(
3823 &h[bi * hs..(bi + 1) * hs],
3824 &lw.input_norm,
3825 eps,
3826 norm_style,
3827 &mut normed[bi * hs..(bi + 1) * hs],
3828 );
3829 }
3830 let attn = short_conv_forward_batch(
3831 &normed,
3832 b,
3833 w,
3834 &cfg,
3835 &mut self.kv_cache.layers[li].linear_state,
3836 pool.as_deref(),
3837 );
3838 for (dst, &a) in h.iter_mut().zip(&attn) {
3839 *dst += a;
3840 }
3841 }
3842 AttnKind::Mla(w) => {
3843 let inv_freq_l = self.layer_inv_freq(li);
3846 let rs = self.layer_rope_scale(li);
3847 let mut normed = vec![0.0f32; hs];
3848 for bi in 0..b {
3849 inference::rms_norm_into(
3850 &h[bi * hs..(bi + 1) * hs],
3851 &lw.input_norm,
3852 eps,
3853 norm_style,
3854 &mut normed,
3855 );
3856 let ao = mla_attention(
3857 w,
3858 &normed,
3859 &mut self.kv_cache.layers[li],
3860 start_pos + bi,
3861 &inv_freq_l,
3862 rs,
3863 eps,
3864 pool.as_deref(),
3865 );
3866 for (dst, &a) in h[bi * hs..(bi + 1) * hs].iter_mut().zip(&ao) {
3867 *dst += a;
3868 }
3869 }
3870 }
3871 AttnKind::Full {
3872 wq,
3873 wk,
3874 wv,
3875 wo,
3876 q_norm,
3877 k_norm,
3878 output_gate,
3879 softplus_gate,
3880 bias,
3881 } => {
3882 let mut normed = vec![0.0f32; b * hs];
3886 for bi in 0..b {
3887 inference::rms_norm_into(
3888 &h[bi * hs..(bi + 1) * hs],
3889 &lw.input_norm,
3890 eps,
3891 norm_style,
3892 &mut normed[bi * hs..(bi + 1) * hs],
3893 );
3894 }
3895 let inv_freq_l = self.layer_inv_freq(li);
3896 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
3897 let cfg = QwenAttnCfg {
3898 num_heads: self.layer_num_heads(li),
3899 num_kv_heads: nkv_l,
3900 head_dim: hd_l,
3901 hidden_size: hs,
3902 position: start_pos,
3903 inv_freq: &inv_freq_l,
3904 rotary_dim: rd_l,
3905 scale: self.attn_scale,
3906 softcap: self.attn_softcap,
3907 window: self.layer_window(li),
3908 v_norm: self.attn_v_norm,
3909 q_norm: q_norm.as_deref(),
3910 k_norm: k_norm.as_deref(),
3911 output_gate: *output_gate,
3912 softplus_gate: softplus_gate
3913 .as_ref()
3914 .map(|(gate, per_head)| (gate, *per_head)),
3915 rope_scale: self.layer_rope_scale(li),
3916 bias: bias
3917 .as_ref()
3918 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
3919 rms_eps: eps,
3920 norm_style,
3921 pool: pool.as_deref(),
3922 };
3923 let mut attn = attention::qwen_attention_batch(
3924 &normed,
3925 b,
3926 wq,
3927 wk,
3928 wv,
3929 wo,
3930 &mut self.kv_cache.layers[li],
3931 &cfg,
3932 );
3933 if let Some(w) = &lw.attn_out_norm {
3934 for bi in 0..b {
3935 inference::rms_norm_into(
3936 &attn[bi * hs..(bi + 1) * hs],
3937 w,
3938 eps,
3939 norm_style,
3940 &mut normed[bi * hs..(bi + 1) * hs],
3941 );
3942 }
3943 attn.copy_from_slice(&normed);
3944 }
3945 for (dst, &a) in h.iter_mut().zip(&attn) {
3946 *dst += a;
3947 }
3948 }
3949 AttnKind::Linear(w) => {
3950 for bi in 0..b {
3951 let normed = inference::rms_norm(
3952 &h[bi * hs..(bi + 1) * hs],
3953 &lw.input_norm,
3954 eps,
3955 norm_style,
3956 );
3957 vmf_phase_forward(
3958 &normed,
3959 w,
3960 &self.vmf_cfg.expect("linear layer without vmf_cfg"),
3961 &mut self.kv_cache.layers[li].linear_state,
3962 pool.as_deref(),
3963 )
3964 .iter()
3965 .enumerate()
3966 .for_each(|(i, &a)| h[bi * hs + i] += a);
3967 }
3968 }
3969 }
3970
3971 let lw = &self.weights.layers[self.phys_layer(li)];
3973 let mut post = vec![0.0f32; b * hs];
3974 for bi in 0..b {
3975 let r =
3976 inference::rms_norm(&h[bi * hs..(bi + 1) * hs], &lw.post_norm, eps, norm_style);
3977 post[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
3978 }
3979 let mask_row = task_mask
3982 .filter(|m| m.ffn_active_count(li) < self.intermediate_size)
3983 .and_then(|m| m.ffn_masks.get(li))
3984 .map(|v| v.as_slice());
3985 let mut ffn = match &lw.ffn {
3986 FfnKind::Dense(d) => dense_ffn_batch(d, &post, b, pool.as_deref(), mask_row),
3987 FfnKind::Moe(m) => moe_ffn_batch(m, &post, b, hs, pool.as_deref(), None),
3988 FfnKind::DenseMoe(dm) => {
3991 let mut out = vec![0.0f32; b * hs];
3992 for bi in 0..b {
3993 let r = dense_moe_ffn(
3994 dm,
3995 &post[bi * hs..(bi + 1) * hs],
3996 &h[bi * hs..(bi + 1) * hs],
3997 eps,
3998 norm_style,
3999 pool.as_deref(),
4000 );
4001 out[bi * hs..(bi + 1) * hs].copy_from_slice(&r);
4002 }
4003 out
4004 }
4005 };
4006 if let Some(w) = &lw.ffn_out_norm {
4007 for bi in 0..b {
4008 inference::rms_norm_into(
4009 &ffn[bi * hs..(bi + 1) * hs],
4010 w,
4011 eps,
4012 norm_style,
4013 &mut post[bi * hs..(bi + 1) * hs],
4014 );
4015 }
4016 ffn.copy_from_slice(&post);
4017 }
4018 for (dst, &f) in h.iter_mut().zip(&ffn) {
4019 *dst += f;
4020 }
4021 if let Some(sc) = lw.layer_scale {
4022 for v in h.iter_mut() {
4023 *v *= sc;
4024 }
4025 }
4026 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
4027 if let Some(t) = tp.parse::<usize>().ok() {
4028 if t >= start_pos && t < start_pos + b {
4029 let bi = t - start_pos;
4030 let row = &h[bi * hs..(bi + 1) * hs];
4031 let n: f32 = row.iter().map(|x| x * x).sum::<f32>().sqrt();
4032 eprintln!(
4033 "BATCH pos {t} after layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
4034 row[0], row[1]
4035 );
4036 }
4037 }
4038 }
4039 if std::env::var("CMF_DEBUG_LAYERS").is_ok() {
4043 let row = &h[(b - 1) * hs..b * hs];
4044 let rms =
4045 (row.iter().map(|&v| (v as f64) * (v as f64)).sum::<f64>() / hs as f64).sqrt();
4046 let mx = row.iter().fold(0f32, |m, &v| m.max(v.abs()));
4047 eprintln!(
4048 "layer {li:>3} {:>10} ffn={:<5} rms={rms:>12.4} max={mx:>12.4}",
4049 match &self.weights.layers[self.phys_layer(li)].attn {
4050 AttnKind::LinearGdn(_) => "gdn",
4051 AttnKind::Linear(_) => "vmf",
4052 AttnKind::ShortConv(_) => "conv",
4053 _ => "attn",
4054 },
4055 match &lw.ffn {
4056 FfnKind::Moe(_) => "moe",
4057 FfnKind::Dense(_) => "dense",
4058 FfnKind::DenseMoe(_) => "dense+moe",
4059 },
4060 );
4061 }
4062 if self.is_loop_end(li) && li + 1 < self.num_layers {
4064 for bi in 0..b {
4065 let normed = inference::rms_norm(
4066 &h[bi * hs..(bi + 1) * hs],
4067 &self.weights.final_norm,
4068 eps,
4069 norm_style,
4070 );
4071 h[bi * hs..(bi + 1) * hs].copy_from_slice(&normed);
4072 }
4073 }
4074 if std::env::var("CMF_TRACE_H").is_ok() {
4075 let n = h[..hs].iter().map(|v| v.abs()).sum::<f32>() / hs as f32;
4076 let mx = h[..hs].iter().fold(0.0f32, |a, &v| a.max(v.abs()));
4077 eprintln!(
4078 "layer {li}: mean|h|={n:.4} max|h|={mx:.2} scale={:?}",
4079 lw.layer_scale
4080 );
4081 }
4082 }
4083 crate::gpu::set_layer(-1); h
4085 }
4086
4087 fn embed_single(&self, id: u32) -> Vec<f32> {
4089 let mut out = vec![0.0f32; self.hidden_size];
4090 if (id as usize) < self.weights.embed_tokens.rows() {
4091 self.weights.embed_tokens.row_f32(id as usize, &mut out);
4092 }
4093 if self.embed_multiplier != 1.0 {
4094 for v in out.iter_mut() {
4095 *v *= self.embed_multiplier;
4096 }
4097 }
4098 if self.dsv4.is_some() {
4102 let mut v = vec![0.0f32; self.hidden_size.max(1)];
4103 v[0] = id as f32;
4104 return v;
4105 }
4106 if let Some(b) = &self.g3n {
4109 return b.0.extend_embedding(id, &out, self.pool.as_deref());
4110 }
4111 out
4112 }
4113
4114 #[cfg(target_os = "macos")]
4120 fn chunk_run_gpu(
4121 &mut self,
4122 li0: usize,
4123 h: &mut [f32],
4124 b: usize,
4125 pos0: usize,
4126 embed_ids: Option<&[u32]>,
4127 cap: usize,
4128 ) -> usize {
4129 if !crate::gpu::enabled_here()
4133 || std::env::var("CMF_GPU_CHUNK")
4134 .map(|v| v == "0")
4135 .unwrap_or(false)
4136 || b < 32
4137 || self.swa.is_some()
4138 || self.global_attn.is_some()
4139 || self.attn_v_norm
4140 || (self.attn_scale - 1.0 / (self.head_dim as f32).sqrt()).abs() > 1e-9
4141 {
4142 return li0;
4143 }
4144 let Some(model) = self.model.clone() else {
4145 return li0;
4146 };
4147 let inv_freq = self.inv_freq.clone();
4148 let (nh, nkv, hd, hs) = (
4149 self.num_heads,
4150 self.num_kv_heads,
4151 self.head_dim,
4152 self.hidden_size,
4153 );
4154 let loop_end = if self.loop_final_norm {
4158 ((li0 / self.physical_layers) + 1) * self.physical_layers
4159 } else {
4160 self.num_layers
4161 };
4162 let mut layers: Vec<crate::gpu_metal::ChunkLayer> = Vec::new();
4163 let mut stored_at: Vec<usize> = Vec::new();
4164 for li in li0..self.num_layers.min(loop_end).min(cap) {
4165 let lw = &self.weights.layers[self.phys_layer(li)];
4166 if lw.attn_out_norm.is_some() || lw.ffn_out_norm.is_some() || lw.layer_scale.is_some() {
4167 break;
4168 }
4169 let AttnKind::Full {
4170 wq,
4171 wk,
4172 wv,
4173 wo,
4174 q_norm,
4175 k_norm,
4176 output_gate: false,
4177 softplus_gate: None,
4178 bias,
4179 } = &lw.attn
4180 else {
4181 break;
4182 };
4183 let FfnKind::Dense(d) = &lw.ffn else { break };
4184 if d.act != Act::Silu {
4185 break;
4186 }
4187 fn cw(t: &QTensor) -> Option<(usize, usize, usize, &[f32])> {
4192 t.q8_row_parts()
4193 .or_else(|| t.q4t_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4194 .or_else(|| t.q4tp_parts().map(|(i, r, c)| (i, r, c, &[][..])))
4195 }
4196 let parts = (
4197 cw(wq),
4198 cw(wk),
4199 cw(wv),
4200 cw(wo),
4201 cw(&d.gate_proj),
4202 cw(&d.up_proj),
4203 cw(&d.down_proj),
4204 );
4205 let (Some(pq), Some(pk), Some(pv), Some(po), Some(pg), Some(pu), Some(pd)) = parts
4206 else {
4207 break;
4208 };
4209 let layer = &self.kv_cache.layers[li];
4210 if layer.mode != crate::kv_cache::KvMode::F32 || layer.o1.is_some() {
4211 break;
4212 }
4213 stored_at.push(layer.head_len(0));
4214 layers.push(crate::gpu_metal::ChunkLayer {
4215 model: &model,
4216 kv_id: self.graph_kv_id,
4217 layer: li,
4218 wq: pq,
4219 wk: pk,
4220 wv: pv,
4221 wo: po,
4222 gate: pg,
4223 up: pu,
4224 down: pd,
4225 input_norm: &lw.input_norm,
4226 post_norm: &lw.post_norm,
4227 bias: bias
4228 .as_ref()
4229 .map(|(a, bb, cc)| (a.as_slice(), bb.as_slice(), cc.as_slice())),
4230 q_norm: q_norm.as_deref(),
4231 k_norm: k_norm.as_deref(),
4232 inv_freq: &inv_freq,
4233 rd: self.rotary_dim,
4234 nh,
4235 nkv,
4236 hd,
4237 hs,
4238 inter: d.gate_proj.rows(),
4239 gemma: matches!(self.norm_style, cortiq_core::NormStyle::Gemma),
4240 eps: self.rms_eps as f32,
4241 });
4242 }
4243 if layers.is_empty() {
4244 return li0;
4245 }
4246 let row = nkv * hd;
4247 let mut store: Vec<(Vec<f32>, Vec<f32>, Vec<f32>)> = stored_at
4248 .iter()
4249 .map(|&st| (vec![0f32; b * row], vec![0f32; b * row], vec![0f32; st + b]))
4250 .collect();
4251 let mut io: Vec<crate::gpu_metal::ChunkIo> = Vec::with_capacity(layers.len());
4252 for (i, (ok, ov, oi)) in store.iter_mut().enumerate() {
4253 let li = layers[i].layer;
4254 let layer = &self.kv_cache.layers[li];
4255 io.push(crate::gpu_metal::ChunkIo {
4256 cpu_stored: stored_at[i],
4257 cpu_k: (0..nkv).map(|g| layer.head_keys(g)).collect(),
4258 cpu_v: (0..nkv).map(|g| layer.head_values(g)).collect(),
4259 out_k: ok,
4260 out_v: ov,
4261 imp: oi,
4262 });
4263 }
4264 let n_run = layers.len();
4265 let last = layers.last().map(|l| l.layer + 1).unwrap_or(li0);
4266 let ep = embed_ids.and_then(|ids| {
4269 self.weights
4270 .embed_tokens
4271 .q8_row_parts()
4272 .map(|(idx, rows, _c, rs)| crate::gpu_metal::ChunkEmbed {
4273 idx,
4274 rows,
4275 row_scale: rs,
4276 ids,
4277 mult: self.embed_multiplier,
4278 })
4279 });
4280 if embed_ids.is_some() && ep.is_none() {
4281 return li0;
4282 }
4283 if !crate::gpu_metal::chunk_run_gpu(&layers, &mut io, h, b, pos0, ep.as_ref()) {
4284 return li0;
4285 }
4286 drop(io);
4287 drop(layers);
4288 for (i, (ok, ov, oi)) in store.iter().enumerate().take(n_run) {
4291 let li = li0 + i;
4292 let layer = &mut self.kv_cache.layers[li];
4293 for bi in 0..b {
4294 layer.append(
4295 &ok[bi * row..(bi + 1) * row],
4296 &ov[bi * row..(bi + 1) * row],
4297 &[],
4298 );
4299 }
4300 layer.accumulate_imp(oi);
4301 }
4302 last
4303 }
4304
4305 fn layer_is_local(&self, li: usize) -> bool {
4308 if let Some(layers) = &self.sliding_layers {
4309 return layers.get(li).copied().unwrap_or(false);
4310 }
4311 match self.swa {
4312 Some((_, pattern)) => (li + 1) % pattern.max(1) != 0,
4313 None => false,
4314 }
4315 }
4316
4317 fn layer_inv_freq(&self, li: usize) -> std::sync::Arc<Vec<f32>> {
4320 if self.layer_is_local(li) {
4321 if let Some(f) = &self.inv_freq_local {
4322 return f.clone();
4323 }
4324 } else if let Some(f) = &self.inv_freq_global {
4325 return f.clone();
4326 }
4327 self.inv_freq.clone()
4328 }
4329
4330 fn layer_window(&self, li: usize) -> Option<usize> {
4332 self.swa
4333 .and_then(|(w, _)| self.layer_is_local(li).then_some(w))
4334 }
4335
4336 fn layer_num_heads(&self, li: usize) -> usize {
4337 self.attention_heads_per_layer
4338 .as_ref()
4339 .and_then(|v| v.get(li).copied())
4340 .unwrap_or(self.num_heads)
4341 }
4342
4343 fn layer_rope_scale(&self, li: usize) -> f32 {
4344 if self.layer_is_local(li) {
4345 self.rope_scale_local
4346 } else {
4347 self.rope_scale
4348 }
4349 }
4350
4351 fn layer_geom(&self, li: usize) -> (usize, usize, usize) {
4354 if !self.layer_is_local(li) {
4355 if let Some((ghd, gkv)) = self.global_attn {
4356 return (gkv, ghd, ghd);
4357 }
4358 }
4359 (
4360 self.num_kv_heads,
4361 self.head_dim,
4362 if self.layer_is_local(li) {
4363 self.rotary_dim_local.unwrap_or(self.rotary_dim)
4364 } else {
4365 self.rotary_dim
4366 },
4367 )
4368 }
4369
4370 fn forward_layers(
4372 &mut self,
4373 hidden: &[f32],
4374 position: usize,
4375 task_mask: Option<&TaskMask>,
4376 ) -> Vec<f32> {
4377 self.forward_layers_upto(hidden, position, task_mask, None)
4378 }
4379
4380 pub fn embed_id(&self, id: u32) -> Vec<f32> {
4388 self.embed_single(id)
4389 }
4390
4391 pub fn split_supported(&self) -> Result<(), String> {
4395 if self.dsv4.is_some() {
4396 return Err(
4397 "network split: DeepSeek-V4 runs its own fused stack (not splittable yet)".into(),
4398 );
4399 }
4400 if self.g3n.is_some() {
4401 return Err(
4402 "network split: Gemma-3n runs its own AltUp stack (not splittable yet)".into(),
4403 );
4404 }
4405 Ok(())
4406 }
4407
4408 pub fn forward_span(
4413 &mut self,
4414 hidden: &[f32],
4415 position: usize,
4416 from: usize,
4417 upto: usize,
4418 task_mask: Option<&TaskMask>,
4419 ) -> Result<Vec<f32>, String> {
4420 self.split_supported()?;
4421 if from > upto || upto >= self.num_layers {
4422 return Err(format!(
4423 "forward_span: layer range {from}..={upto} outside 0..{}",
4424 self.num_layers
4425 ));
4426 }
4427 if hidden.len() != self.hidden_size {
4428 return Err(format!(
4429 "forward_span: hidden len {} ≠ hidden_size {}",
4430 hidden.len(),
4431 self.hidden_size
4432 ));
4433 }
4434 Ok(self.forward_layers_span(hidden, position, task_mask, from, Some(upto)))
4435 }
4436
4437 pub fn logits_from_hidden(&mut self, hidden: &[f32]) -> Vec<f32> {
4440 let normed = inference::rms_norm(
4441 hidden,
4442 &self.weights.final_norm,
4443 self.rms_eps,
4444 self.norm_style,
4445 );
4446 self.lm_head_forward(&normed)
4447 }
4448
4449 pub fn sample_next(&mut self, logits: &[f32], past_tokens: &[u32]) -> u32 {
4451 sampler::sample_with_scratch(
4452 logits,
4453 &self.sampler_config,
4454 past_tokens,
4455 &mut self.rng,
4456 &mut self.sampler_scratch,
4457 )
4458 }
4459
4460 pub fn reset_session(&mut self) {
4462 self.kv_cache.clear();
4463 self.kv_history.clear();
4464 crate::gpu::graph_kv_reset(self.graph_kv_id);
4465 }
4466
4467 pub fn prefill_span_ids(
4473 &mut self,
4474 ids: &[u32],
4475 start_pos: usize,
4476 upto: usize,
4477 task_mask: Option<&TaskMask>,
4478 ) -> Result<Vec<f32>, String> {
4479 self.split_supported()?;
4480 if upto >= self.num_layers {
4481 return Err(format!(
4482 "prefill_span_ids: upto {upto} outside 0..{}",
4483 self.num_layers
4484 ));
4485 }
4486 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4490 Ok(self.prefill_batch_span(PrefillIn::Ids(ids), start_pos, task_mask, 0, upto + 1))
4491 } else {
4492 let hs = self.hidden_size;
4493 let mut out = Vec::with_capacity(ids.len() * hs);
4494 for (i, &id) in ids.iter().enumerate() {
4495 let emb = self.embed_id(id);
4496 out.extend_from_slice(&self.forward_span(
4497 &emb,
4498 start_pos + i,
4499 0,
4500 upto,
4501 task_mask,
4502 )?);
4503 }
4504 Ok(out)
4505 }
4506 }
4507
4508 pub fn prefill_span_hidden(
4511 &mut self,
4512 hidden: &[f32],
4513 start_pos: usize,
4514 from: usize,
4515 upto: usize,
4516 task_mask: Option<&TaskMask>,
4517 ) -> Result<Vec<f32>, String> {
4518 self.split_supported()?;
4519 let hs = self.hidden_size;
4520 if hidden.is_empty() || hidden.len() % hs != 0 {
4521 return Err(format!(
4522 "prefill_span_hidden: {} floats is not a multiple of hidden {hs}",
4523 hidden.len()
4524 ));
4525 }
4526 if from > upto || upto >= self.num_layers {
4527 return Err(format!(
4528 "prefill_span_hidden: layer range {from}..={upto} outside 0..{}",
4529 self.num_layers
4530 ));
4531 }
4532 if self.can_prefill_batched() && !self.graph_prefill_preferred() {
4533 Ok(self.prefill_batch_span(
4534 PrefillIn::Hidden(hidden),
4535 start_pos,
4536 task_mask,
4537 from,
4538 upto + 1,
4539 ))
4540 } else {
4541 let b = hidden.len() / hs;
4542 let mut out = Vec::with_capacity(hidden.len());
4543 for i in 0..b {
4544 let h = self.forward_span(
4545 &hidden[i * hs..(i + 1) * hs],
4546 start_pos + i,
4547 from,
4548 upto,
4549 task_mask,
4550 )?;
4551 out.extend_from_slice(&h);
4552 }
4553 Ok(out)
4554 }
4555 }
4556
4557 fn try_token_graph_wgpu(
4561 &self,
4562 hidden: &[f32],
4563 position: usize,
4564 logits_out: &mut Vec<f32>,
4565 layers_run: &mut usize,
4566 ) -> Option<Vec<f32>> {
4567 self.try_token_graph_wgpu_steps(
4568 hidden,
4569 position,
4570 logits_out,
4571 1,
4572 None,
4573 Some(layers_run),
4574 0,
4575 self.num_layers,
4576 )
4577 }
4578
4579 fn try_token_graph_wgpu_span(
4583 &self,
4584 hidden: &[f32],
4585 position: usize,
4586 logits_out: &mut Vec<f32>,
4587 from: usize,
4588 upto_excl: usize,
4589 layers_run: &mut usize,
4590 ) -> Option<Vec<f32>> {
4591 self.try_token_graph_wgpu_steps(
4592 hidden,
4593 position,
4594 logits_out,
4595 1,
4596 None,
4597 Some(layers_run),
4598 from,
4599 upto_excl,
4600 )
4601 }
4602
4603 fn try_multi_burst(&self, t_next: u32, position: usize, k: usize) -> Option<Vec<u32>> {
4607 if self.o1_active() || self.attn_softcap > 0.0 {
4608 return None;
4609 }
4610 let graph_on = crate::gpu::wgpu_graph_on(crate::gpu::GraphPhase::Decode);
4611 if !graph_on || crate::gpu::graph_unsupported() {
4612 return None;
4619 }
4620 let emb = self.embed_single(t_next);
4621 let mut lg = Vec::new();
4622 let mut ids = Vec::new();
4623 self.try_token_graph_wgpu_steps(
4624 &emb,
4625 position,
4626 &mut lg,
4627 k,
4628 Some(&mut ids),
4629 None,
4630 0,
4631 self.num_layers,
4632 )?;
4633 (ids.len() == k).then_some(ids)
4634 }
4635
4636 fn try_token_graph_wgpu_steps(
4640 &self,
4641 hidden: &[f32],
4642 position: usize,
4643 logits_out: &mut Vec<f32>,
4644 steps: usize,
4645 ids_out: Option<&mut Vec<u32>>,
4646 layers_run: Option<&mut usize>,
4647 from: usize,
4648 upto_excl: usize,
4649 ) -> Option<Vec<f32>> {
4650 let o1_gpu = std::env::var("CMF_O1_GPU").as_deref() == Ok("1");
4653 if (self.o1_active() && !o1_gpu) || self.attn_softcap > 0.0 {
4654 return None;
4658 }
4659 let o1_views: Vec<Option<Vec<crate::nystrom::O1DeviceView<'_>>>> = (from..upto_excl)
4664 .map(|li| {
4665 if !o1_gpu {
4666 return None;
4667 }
4668 self.kv_cache.layers[self.phys_layer(li)].o1_views()
4669 })
4670 .collect();
4671 if self.o1_active() && o1_gpu {
4672 let want: usize = (from..upto_excl)
4675 .filter(|li| !matches!(self.kv_cache.layers[self.phys_layer(*li)].o1, None))
4676 .count();
4677 let have = o1_views.iter().filter(|v| v.is_some()).count();
4678 if want == 0 || have != want {
4679 use std::sync::atomic::{AtomicUsize, Ordering};
4689 static LAST: AtomicUsize = AtomicUsize::new(usize::MAX);
4690 let code = have * 1000 + want;
4691 if LAST.swap(code, Ordering::Relaxed) != code {
4692 tracing::warn!(
4693 "o1 graph: {have} of {want} layers sealed — per-op until all seal"
4694 );
4695 }
4696 return None;
4697 }
4698 }
4699 let nh = self.num_heads;
4700 let (nkv, hd, rd) = self.layer_geom(0);
4701 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
4702 let mut layers = Vec::with_capacity(upto_excl - from);
4703 let mut model = None;
4704 let dbg = std::env::var("CMF_GRAPH_DEBUG").is_ok();
4705 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
4706 if let Some((_, i, kind, rs)) = t.graph_weight() {
4707 return Some(crate::gpu::GraphW {
4708 idx: i,
4709 kind,
4710 row_scale: rs,
4711 data: &[],
4712 });
4713 }
4714 t.as_f32().map(|d| crate::gpu::GraphW {
4716 idx: 0,
4717 kind: 4,
4718 row_scale: &[],
4719 data: d,
4720 })
4721 }
4722 for li in from..upto_excl {
4723 let lw = &self.weights.layers[self.phys_layer(li)];
4724 if dbg {
4725 let ak = match &lw.attn {
4726 AttnKind::Mla(_) => "Mla".into(),
4727 AttnKind::Full {
4728 output_gate, bias, ..
4729 } => format!("Full gate={output_gate} bias={}", bias.is_some()),
4730 AttnKind::LinearGdn(_) => "LinearGdn".into(),
4731 AttnKind::Kda(_) => "Kda".into(),
4732 AttnKind::Linear(_) => "Linear".into(),
4733 AttnKind::ShortConv(_) => "ShortConv".into(),
4734 };
4735 let fk = match &lw.ffn {
4736 FfnKind::Dense(_) => "Dense",
4737 FfnKind::Moe(_) => "Moe",
4738 FfnKind::DenseMoe(_) => "DenseMoe",
4739 };
4740 eprintln!("graph L{li}: attn={ak} ffn={fk}");
4741 }
4742 let gffn = match &lw.ffn {
4743 FfnKind::DenseMoe(_) => return None, FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
4745 gate: gw(&d.gate_proj)?,
4746 up: gw(&d.up_proj)?,
4747 down: gw(&d.down_proj)?,
4748 },
4749 FfnKind::Moe(m) => {
4750 if m.router_sigmoid
4755 || m.expert_bias.is_some()
4756 || m.route_tau.is_some()
4757 || m.mask.is_some()
4758 {
4759 return None;
4760 }
4761 let (se, sg) = m.shared.as_ref()?;
4762 let sgate = gw(sg.as_ref()?)?;
4763 let router = gw(&m.router)?;
4764 let inter = m.experts.first()?.gate_proj.rows();
4765 let mut experts = Vec::with_capacity(m.experts.len() + 1);
4766 let mut q4tp: Option<bool> = None;
4769 let mut gu_q2: Option<bool> = None;
4772 for e in m.experts.iter().chain(std::iter::once(se)) {
4773 if !matches!(e.act, Act::Silu)
4774 || e.gate_proj.rows() != inter
4775 || e.up_proj.rows() != inter
4776 {
4777 return None;
4778 }
4779 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
4780 Some((mm, gi)) => (
4781 mm,
4782 gi,
4783 e.up_proj.mapped_q4t()?.1,
4784 e.down_proj.mapped_q4t()?.1,
4785 false,
4786 false,
4787 ),
4788 None => match e.gate_proj.mapped_q2tp() {
4789 Some((mm, gi)) => (
4790 mm,
4791 gi,
4792 e.up_proj.mapped_q2tp()?.1,
4793 e.down_proj.mapped_q4tp()?.1,
4794 true,
4795 true,
4796 ),
4797 None => {
4798 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
4799 (
4800 mm,
4801 gi,
4802 e.up_proj.mapped_q4tp()?.1,
4803 e.down_proj.mapped_q4tp()?.1,
4804 true,
4805 false,
4806 )
4807 }
4808 },
4809 };
4810 if *q4tp.get_or_insert(is_p) != is_p || *gu_q2.get_or_insert(is_q2) != is_q2
4811 {
4812 tracing::warn!(
4818 "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."
4819 );
4820 return None;
4821 }
4822 model.get_or_insert_with(|| mm.clone());
4823 experts.push((gi, ui, di));
4824 }
4825 crate::gpu::GraphFfn::Moe {
4826 router,
4827 shared_gate: sgate,
4828 experts,
4829 n_exp: m.experts.len(),
4830 top_k: std::env::var("CMF_TOPK_PROBE")
4836 .ok()
4837 .and_then(|v| v.parse::<usize>().ok())
4838 .filter(|k| *k > 0 && *k <= m.top_k)
4839 .unwrap_or(m.top_k),
4840 inter,
4841 norm_topk: m.norm_topk_prob,
4842 q4tp: q4tp?,
4843 gu_q2: gu_q2.unwrap_or(false),
4844 }
4845 }
4846 };
4847 let attn = match &lw.attn {
4848 AttnKind::Full {
4849 wq,
4850 wk,
4851 wv,
4852 wo,
4853 q_norm,
4854 k_norm,
4855 output_gate,
4856 softplus_gate,
4857 bias,
4858 } => {
4859 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
4860 return None;
4861 }
4862 let (m, _, _, _) = wq.graph_weight()?;
4863 model = Some(m.clone());
4864 crate::gpu::GraphAttn::Full {
4865 wq: gw(wq)?,
4866 wk: gw(wk)?,
4867 wv: gw(wv)?,
4868 wo: gw(wo)?,
4869 q_norm: q_norm.as_deref(),
4870 k_norm: k_norm.as_deref(),
4871 bias: bias
4872 .as_ref()
4873 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
4874 output_gate: *output_gate,
4875 cpu_k: self.kv_cache.layers[li].k_heads(),
4876 cpu_v: self.kv_cache.layers[li].v_heads(),
4877 }
4878 }
4879 AttnKind::LinearGdn(w) => {
4880 let cfg = self.gdn_cfg?;
4881 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
4882 model = Some(m.clone());
4883 crate::gpu::GraphAttn::Gdn {
4884 qkv: gw(&w.in_proj_qkv)?,
4885 z: gw(&w.in_proj_z)?,
4886 a: gw(&w.in_proj_a)?,
4887 b: gw(&w.in_proj_b)?,
4888 out: gw(&w.out_proj)?,
4889 conv1d: &w.conv1d,
4890 a_log: &w.a_log,
4891 dt_bias: &w.dt_bias,
4892 norm: &w.norm,
4893 nv: cfg.num_v_heads,
4894 nk: cfg.num_k_heads,
4895 dk: cfg.key_head_dim,
4896 dv: cfg.value_head_dim,
4897 kk: cfg.conv_kernel,
4898 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
4899 }
4900 }
4901 _ => return None,
4902 };
4903 layers.push(crate::gpu::GraphLayer {
4904 input_norm: &lw.input_norm,
4905 attn,
4906 post_norm: &lw.post_norm,
4907 ffn: gffn,
4908 });
4909 }
4910 let model = model?;
4911 let lm_gw = if upto_excl == self.num_layers
4917 && self.graph_want_logits
4918 && std::env::var("CMF_GPU_LMHEAD")
4919 .map(|v| v != "0")
4920 .unwrap_or(true)
4921 {
4922 self.weights.lm_head.graph_weight().map(|(_, i, kind, rs)| {
4923 (
4924 crate::gpu::GraphW {
4925 idx: i,
4926 kind,
4927 row_scale: rs,
4928 data: &[],
4929 },
4930 self.weights.lm_head.rows(),
4931 )
4932 })
4933 } else {
4934 None
4935 };
4936 let lm = lm_gw.as_ref().map(|(gw, rows)| (gw, *rows));
4937 let emb_gw = if steps > 1 {
4939 self.weights
4940 .embed_tokens
4941 .graph_weight()
4942 .map(|(_, i, kind, rs)| {
4943 (
4944 crate::gpu::GraphW {
4945 idx: i,
4946 kind,
4947 row_scale: rs,
4948 data: &[],
4949 },
4950 self.weights.embed_tokens.rows(),
4951 self.embed_multiplier as f32,
4952 )
4953 })
4954 } else {
4955 None
4956 };
4957
4958 let loop_norm_at: Vec<usize> = if self.loop_final_norm {
4964 (from..upto_excl.min(self.num_layers - 1))
4965 .filter(|&li| (li + 1) % self.physical_layers == 0)
4966 .map(|li| li - from)
4967 .collect()
4968 } else {
4969 Vec::new()
4970 };
4971 let mut h = hidden.to_vec();
4972 crate::gpu::forward_token_graph(
4973 &model,
4974 self.graph_kv_id,
4975 &layers,
4976 &o1_views,
4977 self.o1_epoch,
4978 &self.inv_freq,
4979 &mut h,
4980 nh,
4981 nkv,
4982 hd,
4983 rd,
4984 self.hidden_size,
4985 self.intermediate_size,
4986 position,
4987 self.kv_cache.max_seq_len,
4988 gemma,
4989 self.rms_eps as f32,
4990 lm,
4991 &self.weights.final_norm,
4992 logits_out,
4993 &loop_norm_at,
4994 steps,
4995 emb_gw.as_ref().map(|(gw, rows, m)| (gw, *rows, *m)),
4996 ids_out,
4997 layers_run,
4998 from,
4999 )
5000 .then_some(h)
5001 }
5002
5003 fn try_batch_graph_wgpu(
5008 &self,
5009 hiddens: &mut [f32],
5010 positions: &[usize],
5011 k: usize,
5012 spec: Option<crate::gpu::SpecTail<'_>>,
5013 ) -> bool {
5014 let _tb = std::time::Instant::now();
5015 if self.attn_softcap > 0.0 {
5016 return false; }
5018 if self.o1_active() {
5019 return false;
5020 }
5021 let nh = self.num_heads;
5022 let (nkv, hd, rd) = self.layer_geom(0);
5023 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
5024 fn gw(t: &QTensor) -> Option<crate::gpu::GraphW<'_>> {
5025 if let Some((_, i, kind, rs)) = t.graph_weight() {
5026 return Some(crate::gpu::GraphW {
5027 idx: i,
5028 kind,
5029 row_scale: rs,
5030 data: &[],
5031 });
5032 }
5033 t.as_f32().map(|d| crate::gpu::GraphW {
5034 idx: 0,
5035 kind: 4,
5036 row_scale: &[],
5037 data: d,
5038 })
5039 }
5040 let built: Option<(
5041 Vec<crate::gpu::GraphLayer<'_>>,
5042 std::sync::Arc<cortiq_core::CmfModel>,
5043 )> = (|| {
5044 let mut layers = Vec::with_capacity(self.num_layers);
5045 let mut model = None;
5046 for li in 0..self.num_layers {
5047 let lw = &self.weights.layers[self.phys_layer(li)];
5048 let gffn = match &lw.ffn {
5055 FfnKind::Dense(d) => crate::gpu::GraphFfn::Dense {
5056 gate: gw(&d.gate_proj)?,
5057 up: gw(&d.up_proj)?,
5058 down: gw(&d.down_proj)?,
5059 },
5060 FfnKind::Moe(m) => {
5061 if m.router_sigmoid
5062 || m.expert_bias.is_some()
5063 || m.route_tau.is_some()
5064 || m.mask.is_some()
5065 {
5066 return None;
5067 }
5068 let (se, sg) = m.shared.as_ref()?;
5069 let sgate = gw(sg.as_ref()?)?;
5070 let router = gw(&m.router)?;
5071 let inter = m.experts.first()?.gate_proj.rows();
5072 let mut experts = Vec::with_capacity(m.experts.len() + 1);
5073 let mut q4tp: Option<bool> = None;
5074 let mut gu_q2: Option<bool> = None;
5075 for e in m.experts.iter().chain(std::iter::once(se)) {
5076 if !matches!(e.act, Act::Silu)
5077 || e.gate_proj.rows() != inter
5078 || e.up_proj.rows() != inter
5079 {
5080 return None;
5081 }
5082 let (mm, gi, ui, di, is_p, is_q2) = match e.gate_proj.mapped_q4t() {
5086 Some((mm, gi)) => (
5087 mm,
5088 gi,
5089 e.up_proj.mapped_q4t()?.1,
5090 e.down_proj.mapped_q4t()?.1,
5091 false,
5092 false,
5093 ),
5094 None => match e.gate_proj.mapped_q2tp() {
5095 Some((mm, gi)) => (
5096 mm,
5097 gi,
5098 e.up_proj.mapped_q2tp()?.1,
5099 e.down_proj.mapped_q4tp()?.1,
5100 true,
5101 true,
5102 ),
5103 None => {
5104 let (mm, gi) = e.gate_proj.mapped_q4tp()?;
5105 (
5106 mm,
5107 gi,
5108 e.up_proj.mapped_q4tp()?.1,
5109 e.down_proj.mapped_q4tp()?.1,
5110 true,
5111 false,
5112 )
5113 }
5114 },
5115 };
5116 if *q4tp.get_or_insert(is_p) != is_p
5117 || *gu_q2.get_or_insert(is_q2) != is_q2
5118 {
5119 return None;
5120 }
5121 model.get_or_insert_with(|| mm.clone());
5122 experts.push((gi, ui, di));
5123 }
5124 crate::gpu::GraphFfn::Moe {
5125 router,
5126 shared_gate: sgate,
5127 experts,
5128 n_exp: m.experts.len(),
5129 top_k: m.top_k,
5130 inter,
5131 norm_topk: m.norm_topk_prob,
5132 q4tp: q4tp?,
5133 gu_q2: gu_q2.unwrap_or(false),
5134 }
5135 }
5136 _ => return None,
5137 };
5138 let attn = match &lw.attn {
5139 AttnKind::Full {
5140 wq,
5141 wk,
5142 wv,
5143 wo,
5144 q_norm,
5145 k_norm,
5146 output_gate,
5147 softplus_gate,
5148 bias,
5149 } => {
5150 if softplus_gate.is_some() || self.attention_heads_per_layer.is_some() {
5151 return None;
5152 }
5153 let (m, _, _, _) = wq.graph_weight()?;
5154 model = Some(m.clone());
5155 crate::gpu::GraphAttn::Full {
5156 wq: gw(wq)?,
5157 wk: gw(wk)?,
5158 wv: gw(wv)?,
5159 wo: gw(wo)?,
5160 q_norm: q_norm.as_deref(),
5161 k_norm: k_norm.as_deref(),
5162 bias: bias
5163 .as_ref()
5164 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
5165 output_gate: *output_gate,
5166 cpu_k: self.kv_cache.layers[li].k_heads(),
5167 cpu_v: self.kv_cache.layers[li].v_heads(),
5168 }
5169 }
5170 AttnKind::LinearGdn(w) => {
5171 let cfg = self.gdn_cfg?;
5172 let (m, _, _, _) = w.in_proj_qkv.graph_weight()?;
5173 model = Some(m.clone());
5174 crate::gpu::GraphAttn::Gdn {
5175 qkv: gw(&w.in_proj_qkv)?,
5176 z: gw(&w.in_proj_z)?,
5177 a: gw(&w.in_proj_a)?,
5178 b: gw(&w.in_proj_b)?,
5179 out: gw(&w.out_proj)?,
5180 conv1d: &w.conv1d,
5181 a_log: &w.a_log,
5182 dt_bias: &w.dt_bias,
5183 norm: &w.norm,
5184 nv: cfg.num_v_heads,
5185 nk: cfg.num_k_heads,
5186 dk: cfg.key_head_dim,
5187 dv: cfg.value_head_dim,
5188 kk: cfg.conv_kernel,
5189 cpu_state: &self.kv_cache.layers[self.phys_layer(li)].linear_state,
5190 }
5191 }
5192 _ => return None,
5193 };
5194 layers.push(crate::gpu::GraphLayer {
5195 input_norm: &lw.input_norm,
5196 attn,
5197 post_norm: &lw.post_norm,
5198 ffn: gffn,
5199 });
5200 }
5201 Some((layers, model?))
5202 })();
5203 let Some((layers, model)) = built else {
5204 {
5205 use std::sync::atomic::{AtomicBool, Ordering};
5206 static SAID: AtomicBool = AtomicBool::new(false);
5207 if !SAID.swap(true, Ordering::Relaxed) {
5208 tracing::warn!("batch graph: BUILDER refused (layer weights/kinds)");
5209 }
5210 }
5211 return false;
5212 };
5213 if std::env::var("CMF_GRAPH_SPEC_TIME").is_ok() {
5214 eprintln!("batch-build: {:.1} ms", _tb.elapsed().as_secs_f64() * 1e3);
5215 }
5216 crate::gpu::forward_batch_graph(
5217 &model,
5218 self.graph_kv_id,
5219 &layers,
5220 &self.inv_freq,
5221 hiddens,
5222 nh,
5223 nkv,
5224 hd,
5225 rd,
5226 self.hidden_size,
5227 self.intermediate_size,
5228 positions,
5229 self.kv_cache.max_seq_len,
5230 gemma,
5231 self.rms_eps as f32,
5232 k,
5233 spec,
5234 )
5235 }
5236
5237 fn draft_probe() -> bool {
5241 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5242 *ON.get_or_init(|| std::env::var("CMF_DSV4_DRAFT_PROBE").is_ok_and(|v| v != "0"))
5243}
5244
5245 #[cfg(feature = "gpu")]
5257 fn dsv4_spec_on() -> bool {
5258 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
5259 *ON.get_or_init(|| std::env::var("CMF_DSV4_SPEC").map(|v| v != "0").unwrap_or(true))
5260 }
5261
5262 #[cfg(feature = "gpu")]
5269 fn dsv4_spec_step(
5270 &mut self,
5271 tip_token: u32,
5272 t_next: u32,
5273 next_pos: usize,
5274 drafted: &mut usize,
5275 accepted_ctr: &mut usize,
5276 ) -> Option<(Vec<u32>, usize)> {
5277 let t_all = std::time::Instant::now();
5278 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5279 thread_local! {
5280 static LAST: std::cell::Cell<Option<std::time::Instant>> =
5281 const { std::cell::Cell::new(None) };
5282 }
5283 LAST.with(|l| {
5284 if let Some(prev) = l.get() {
5285 eprintln!("между раундами {:.1} мс", prev.elapsed().as_secs_f64() * 1e3);
5286 }
5287 l.set(Some(std::time::Instant::now()));
5288 });
5289 }
5290 if std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5291 eprintln!("spec_step: вход pos={next_pos}");
5292 }
5293 let n_layers = self.dsv4.as_ref().map(|b| b.1.len())?;
5294 let cfg = self.dsv4.as_ref().map(|b| b.2)?;
5295 if self.dspark.is_none() {
5297 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5298 if t.is_empty() {
5299 return None;
5300 }
5301 crate::dsv4::dspark_arm(&t, cfg.dim);
5302 self.dspark = Some(crate::dsv4::DsparkState::new(
5303 self.dsv4_mtp.len(),
5304 &cfg,
5305 t.len(),
5306 ));
5307 }
5308 let targets = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5309 let pack = crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg);
5310 if pack.is_none() && std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok() {
5311 eprintln!("spec_step: пак не построился (targets {targets:?})");
5312 }
5313 let pack = pack?;
5314 let block = crate::dsv4::dspark_block();
5315 let b_box = self.dsv4.as_mut()?;
5316 let (g, layers, st) = (&b_box.0, &b_box.1, &mut b_box.3);
5317 let ds = self.dspark.as_mut()?;
5318 let dbg = std::env::var("CMF_DSV4_SPEC_DEBUG").is_ok();
5321 if !crate::dsv4::dspark_take(&mut ds.main_hidden) && !ds.have_hidden {
5322 if dbg {
5323 eprintln!("spec_step: нет захвата");
5324 }
5325 return None;
5326 }
5327 ds.have_hidden = true;
5328 let tip_pos = next_pos.checked_sub(1)?;
5329 let draft_started = std::time::Instant::now();
5330 let mut conf = Vec::new();
5331 let props = crate::dsv4::dspark_draft_gpu(
5332 g,
5333 &self.dsv4_mtp,
5334 &cfg,
5335 ds,
5336 pack,
5337 st.kv_id,
5338 tip_token,
5339 tip_pos,
5340 self.pool.as_deref(),
5341 &mut conf,
5342 );
5343 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5344 *drafted += block;
5345 if props.is_empty() || props[0] != t_next {
5346 if dbg {
5347 eprintln!(
5348 "spec_step: черновик {} (props0={:?} t_next={t_next})",
5349 if props.is_empty() { "пуст" } else { "мимо" },
5350 props.first()
5351 );
5352 }
5353 return None;
5354 }
5355 let mut k_verify = crate::dsv4::dspark_verify_k().min(props.len());
5356 let conf_min = {
5362 static M: std::sync::OnceLock<f32> = std::sync::OnceLock::new();
5363 *M.get_or_init(|| {
5364 std::env::var("CMF_DSPARK_CONF_MIN")
5365 .ok()
5366 .and_then(|v| v.parse().ok())
5367 .unwrap_or(0.0)
5368 })
5369 };
5370 if conf_min > 0.0 && conf.len() >= props.len() {
5371 let mut keep = 1usize;
5372 while keep < k_verify && conf.get(keep).copied().unwrap_or(0.0) >= conf_min {
5373 keep += 1;
5374 }
5375 k_verify = k_verify.min(keep.max(2));
5376 }
5377 if k_verify < 2 {
5378 return None;
5379 }
5380 let mut fed = Vec::with_capacity(k_verify);
5381 fed.push(t_next);
5382 fed.extend_from_slice(&props[1..k_verify]);
5383 let mut argmax = Vec::new();
5384 let mut logits_all = Vec::new();
5385 let mut walked = Vec::new();
5386 let txn = crate::dsv4::dsv4_verify_chunk(
5387 g,
5388 layers,
5389 &cfg,
5390 st,
5391 &fed,
5392 next_pos,
5393 &self.inv_freq,
5394 self.pool.as_deref(),
5395 &targets,
5396 &mut argmax,
5397 &mut logits_all,
5398 &mut walked,
5399 );
5400 if txn.is_none() && dbg {
5401 eprintln!("spec_step: verify отказал");
5402 }
5403 let txn = txn?;
5404 let b = fed.len();
5405 let mut accepted = 1usize;
5406 while accepted < b && fed[accepted] == argmax[accepted - 1] {
5407 accepted += 1;
5408 }
5409 if std::env::var("CMF_DSV4_SPEC_FORCE_REJECT").is_ok_and(|v| v != "0") {
5414 accepted = 1;
5415 }
5416 if std::env::var("CMF_DSV4_SPEC_TRACE").is_ok() {
5417 eprintln!(
5418 "spec@{next_pos}: fed={fed:?} argmax={argmax:?} accepted={accepted}"
5419 );
5420 }
5421 let t_fin = std::time::Instant::now();
5422 if !crate::dsv4::dsv4_spec_finish(
5423 g,
5424 layers,
5425 &cfg,
5426 st,
5427 txn,
5428 accepted,
5429 &fed,
5430 &self.inv_freq,
5431 self.pool.as_deref(),
5432 ) {
5433 tracing::warn!("dsv4: спекулятивный откат не удался — состояние подозрительно");
5434 return None;
5435 }
5436 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5437 eprintln!("finish(k={accepted}): {:.1} мс", t_fin.elapsed().as_secs_f64() * 1e3);
5438 }
5439 *accepted_ctr += accepted - 1;
5440 let (hc, dim) = (cfg.hc_mult, cfg.dim);
5445 let dev_caps: Vec<usize> = targets
5452 .iter()
5453 .copied()
5454 .filter(|&t| {
5455 st.dev_set.get(t).copied().unwrap_or(false)
5456 && !st.partial_set.get(t).copied().unwrap_or(false)
5457 })
5458 .collect();
5459 let mut caps_all = vec![0.0f32; dev_caps.len() * b * hc * dim];
5460 if !crate::gpu_wgpu::dsv4_spec_cap_read_all(b, dev_caps.len(), hc * dim, &mut caps_all) {
5461 return None;
5462 }
5463 for t in 0..accepted {
5464 let tip = t + 1 == accepted;
5465 for (slot, &tl) in targets.iter().enumerate() {
5466 if let Some(di) = dev_caps.iter().position(|&d| d == tl) {
5467 let lo = (di * b + t) * hc * dim;
5468 crate::dsv4::dspark_capture(
5469 &caps_all[lo..lo + hc * dim],
5470 &cfg,
5471 slot,
5472 &mut ds.main_hidden,
5473 );
5474 } else if tip
5475 && crate::dsv4::dspark_peek_slot(slot, dim, {
5476 let lo = slot * dim;
5477 &mut ds.main_hidden[lo..lo + dim]
5478 })
5479 {
5480 } else {
5485 crate::dsv4::dspark_capture(
5489 &walked[t * hc * dim..(t + 1) * hc * dim],
5490 &cfg,
5491 slot,
5492 &mut ds.main_hidden,
5493 );
5494 }
5495 }
5496 crate::dsv4::dspark_ring_append(g, &self.dsv4_mtp, &cfg, ds, next_pos + t, self.pool.as_deref());
5497 }
5498 let row = logits_all[(accepted - 1) * cfg.vocab..accepted * cfg.vocab].to_vec();
5499 self.graph_logits = Some(row);
5500 if std::env::var("CMF_DSV4_TRUNK_PICK_DUMP").is_ok() {
5505 crate::dsv4::trunk_freq_note(&crate::dsv4::pick_tally_take());
5506 crate::dsv4::pick_tally_arm();
5507 }
5508 if std::env::var("CMF_DSV4_SPEC_TIME").is_ok() {
5509 eprintln!("spec_step total {:.1} мс (k={accepted})", t_all.elapsed().as_secs_f64() * 1e3);
5510 }
5511 Some((fed[1..accepted].to_vec(), next_pos + accepted))
5512 }
5513
5514 fn dspark_probe(&mut self, position: usize, token_id: u32) {
5515 if self.dsv4_mtp.is_empty() || !Self::draft_probe() {
5516 return;
5517 }
5518 let trunk_now = crate::dsv4::pick_tally_take();
5520 crate::dsv4::trunk_freq_note(&trunk_now);
5521 if !trunk_now.is_empty() {
5522 self.dspark_trunk_picks.push(trunk_now);
5523 let keep = crate::dsv4::dspark_block();
5524 if self.dspark_trunk_picks.len() > keep {
5525 self.dspark_trunk_picks.remove(0);
5526 }
5527 }
5528 for p in std::mem::take(&mut self.dspark_pending) {
5531 let Some(i) = position.checked_sub(p.0 + 1) else {
5532 continue;
5533 };
5534 let mut p = p;
5535 if i < p.1.len() {
5536 if p.2 && p.1[i] == token_id {
5537 p.3 = i + 1;
5538 } else {
5539 p.2 = false;
5540 }
5541 if i + 1 < p.1.len() {
5542 self.dspark_pending.push(p);
5543 continue;
5544 }
5545 }
5546 self.dspark_hist.push(p.3);
5547 self.dspark_real.push(token_id);
5548 }
5549 let Some(b) = &mut self.dsv4 else { return };
5550 let (g, layers, cfg) = (&b.0, &b.1, b.2);
5551 let n_layers = layers.len();
5552 if self.dspark.is_none() {
5553 let t = crate::dsv4::dspark_targets(&self.dsv4_mtp, &cfg, n_layers);
5554 if t.is_empty() {
5555 return;
5556 }
5557 eprintln!("DSpark: захват со слоёв {t:?}, блок {}", crate::dsv4::dspark_block());
5558 crate::dsv4::dspark_arm(&t, cfg.dim);
5559 self.dspark = Some(crate::dsv4::DsparkState::new(
5560 self.dsv4_mtp.len(),
5561 &cfg,
5562 t.len(),
5563 ));
5564 }
5565 let ds = self.dspark.as_mut().unwrap();
5566 if !crate::dsv4::dspark_take(&mut ds.main_hidden) {
5567 return; }
5569 let mut conf = Vec::new();
5570 crate::dsv4::pick_tally_arm();
5571 let draft_started = std::time::Instant::now();
5576 #[cfg(feature = "gpu")]
5577 let gpu_draft = crate::dsv4::dspark_gpu_on();
5578 #[cfg(not(feature = "gpu"))]
5579 let gpu_draft = false;
5580 let props = if gpu_draft {
5581 #[cfg(feature = "gpu")]
5582 {
5583 let kv_id = b.3.kv_id;
5584 match crate::dsv4::dspark_pack_get(&self.dsv4_mtp, &cfg) {
5585 Some(pk) => crate::dsv4::dspark_draft_gpu(
5586 g,
5587 &self.dsv4_mtp,
5588 &cfg,
5589 ds,
5590 pk,
5591 kv_id,
5592 token_id,
5593 position,
5594 self.pool.as_deref(),
5595 &mut conf,
5596 ),
5597 None => Vec::new(),
5598 }
5599 }
5600 #[cfg(not(feature = "gpu"))]
5601 Vec::new()
5602 } else {
5603 crate::gpu::cpu_scope(|| {
5604 crate::dsv4::dspark_draft(
5605 g,
5606 &self.dsv4_mtp,
5607 &cfg,
5608 ds,
5609 token_id,
5610 position,
5611 self.pool.as_deref(),
5612 &mut conf,
5613 )
5614 })
5615 };
5616 self.dspark_draft_ns += draft_started.elapsed().as_nanos();
5617 let draft_picks = crate::dsv4::pick_tally_take();
5618 crate::dsv4::dspark_freq_note(&draft_picks);
5619 crate::dsv4::pick_tally_arm();
5622 if !props.is_empty() {
5623 let (tu, tt) = {
5627 let flat: Vec<(usize, Vec<usize>)> = self
5628 .dspark_trunk_picks
5629 .iter()
5630 .flat_map(|v| v.iter().cloned())
5631 .collect();
5632 let mut per: std::collections::HashMap<usize, Vec<usize>> =
5634 std::collections::HashMap::new();
5635 for (li, picks) in flat {
5636 per.entry(li).or_default().extend(picks);
5637 }
5638 let n = per.len().max(1);
5639 let mut u = 0usize;
5640 let mut t = 0usize;
5641 for (_, v) in per {
5642 t += v.len();
5643 u += v.iter().collect::<std::collections::HashSet<_>>().len();
5644 }
5645 (u / n, t / n)
5646 };
5647 let (du, dt) = crate::dsv4::tally_unique(&draft_picks);
5648 self.dspark_exp.push((tu, tt, du, dt));
5649 self.dspark_pending.push((position, props, true, 0));
5650 }
5651 if self.dspark_hist.len() >= 8 && self.dspark_hist.len() % 8 == 0 {
5652 let n = self.dspark_hist.len() as f32;
5653 let mean: f32 = self.dspark_hist.iter().sum::<usize>() as f32 / n;
5654 let block = crate::dsv4::dspark_block();
5655 let mut at = vec![0usize; block + 1];
5656 for &k in &self.dspark_hist {
5657 at[k] += 1;
5658 }
5659 let mut surv = Vec::with_capacity(block);
5661 for i in 1..=block {
5662 let k = at[i..].iter().sum::<usize>() as f32 / n;
5663 surv.push(format!("{k:.2}"));
5664 }
5665 let distinct = self
5666 .dspark_real
5667 .iter()
5668 .collect::<std::collections::HashSet<_>>()
5669 .len();
5670 let (tu, tt, du, dt) = self.dspark_exp.iter().fold((0, 0, 0, 0), |a, b| {
5671 (a.0 + b.0, a.1 + b.1, a.2 + b.2, a.3 + b.3)
5672 });
5673 let m = self.dspark_exp.len().max(1);
5674 eprintln!(
5675 "DSpark: черновиков {}, принято в среднем {mean:.2} из {block} \
5676 (токенов за проход {:.2}), распределение {at:?}, выживание [{}]",
5677 self.dspark_hist.len(),
5678 mean + 1.0,
5679 surv.join(" ")
5680 );
5681 eprintln!(
5682 "DSpark: разных токенов {distinct} из {} (вырожденность), \
5683 эксперты ствол {}/{} на слой за {block} токенов, \
5684 черновик {}/{} за блок, draft {:.2} мс/блок",
5685 self.dspark_real.len(),
5686 tu / m,
5687 tt / m,
5688 du / m,
5689 dt / m,
5690 self.dspark_draft_ns as f64 / self.dspark_exp.len().max(1) as f64 / 1e6
5691 );
5692 }
5693 }
5694
5695 fn forward_layers_upto(
5696 &mut self,
5697 hidden: &[f32],
5698 position: usize,
5699 task_mask: Option<&TaskMask>,
5700 upto: Option<usize>,
5701 ) -> Vec<f32> {
5702 if let Some(plan) = self.gpu_plan.clone() {
5708 if upto.is_none() && plan.len() > 1 {
5709 let mut h = hidden.to_vec();
5710 for &(dev, from, upto_incl) in plan.iter() {
5711 h = crate::gpu::with_device(dev, || {
5712 self.forward_layers_span(&h, position, task_mask, from, Some(upto_incl))
5713 });
5714 }
5715 return h;
5716 }
5717 }
5718 self.forward_layers_span(hidden, position, task_mask, 0, upto)
5719 }
5720
5721 pub fn set_gpu_plan(&mut self, devices: Option<&[usize]>) -> Result<(), String> {
5726 self.set_gpu_plan_at(devices, None)
5727 }
5728
5729 pub fn set_gpu_plan_at(
5733 &mut self,
5734 devices: Option<&[usize]>,
5735 at: Option<usize>,
5736 ) -> Result<(), String> {
5737 let Some(devs) = devices.filter(|d| d.len() > 1) else {
5738 self.gpu_plan = None;
5739 return Ok(());
5740 };
5741 self.split_supported()?;
5742 let n = self.num_layers;
5743 if devs.len() > n {
5744 return Err(format!("{} devices for {n} layers", devs.len()));
5745 }
5746 if let Some(k) = at {
5747 if k == 0 || k >= n {
5748 return Err(format!("split at {k}: the model has {n} layers"));
5749 }
5750 if devs.len() == 2 {
5751 self.gpu_plan = Some(std::sync::Arc::new(vec![
5752 (devs[0], 0, k - 1),
5753 (devs[1], k, n - 1),
5754 ]));
5755 return Ok(());
5756 }
5757 return Err(format!(
5758 "an explicit split point takes exactly 2 devices, got {}",
5759 devs.len()
5760 ));
5761 }
5762 let per = n.div_ceil(devs.len());
5763 let mut plan = Vec::with_capacity(devs.len());
5764 let mut from = 0usize;
5765 for &d in devs {
5766 if from >= n {
5767 break;
5768 }
5769 let upto = (from + per - 1).min(n - 1);
5770 plan.push((d, from, upto));
5771 from = upto + 1;
5772 }
5773 self.gpu_plan = Some(std::sync::Arc::new(plan));
5774 Ok(())
5775 }
5776
5777 pub fn gpu_plan(&self) -> Option<Vec<(usize, usize, usize)>> {
5779 self.gpu_plan.as_ref().map(|p| p.as_ref().clone())
5780 }
5781
5782 fn forward_layers_span(
5788 &mut self,
5789 hidden: &[f32],
5790 position: usize,
5791 task_mask: Option<&TaskMask>,
5792 from: usize,
5793 upto: Option<usize>,
5794 ) -> Vec<f32> {
5795 debug_assert!(from == 0 || (self.dsv4.is_none() && self.g3n.is_none()));
5796 if let Some(b) = &mut self.dsv4 {
5802 let _ = (task_mask, upto);
5803 let token_id = hidden.first().copied().unwrap_or(0.0) as u32;
5804 let (g, layers, cfg, st) = (&b.0, &b.1, b.2, &mut b.3);
5805 st.pos = position;
5806 let mut logits = Vec::new();
5807 crate::dsv4::forward_token(
5808 g,
5809 layers,
5810 &cfg,
5811 st,
5812 token_id,
5813 &self.inv_freq,
5814 self.pool.as_deref(),
5815 &mut logits,
5816 );
5817 self.graph_logits = Some(logits);
5818 self.dspark_probe(position, token_id);
5819 return vec![0.0; self.hidden_size];
5822 }
5823 if let Some(b) = &self.g3n {
5826 let _ = (task_mask, upto);
5827 return crate::g3n::g3n_forward(
5828 &b.0,
5829 &b.1,
5830 hidden,
5831 position,
5832 &mut self.kv_cache.layers,
5833 self.num_heads,
5834 self.num_kv_heads,
5835 self.head_dim,
5836 self.pool.as_deref(),
5837 );
5838 }
5839 let mut h = hidden.to_vec();
5840 let (nh, _nkv, _hd, hs, _rd, eps) = (
5843 self.num_heads,
5844 self.num_kv_heads,
5845 self.head_dim,
5846 self.hidden_size,
5847 self.rotary_dim,
5848 self.rms_eps,
5849 );
5850 let pool = self.pool.clone();
5851 let graph_env = std::env::var("CMF_GPU_WGPU_GRAPH").ok();
5863 let graph_on = match graph_env.as_deref() {
5864 Some("0") => false,
5865 Some("prefill") => false, Some(_) => true,
5867 None => crate::gpu::wgpu_graph_default(),
5873 };
5874 let graph_trusted =
5875 graph_env.is_some() || crate::gpu::wgpu_graph_default() || self.gdn_cfg.is_some();
5876 let race_eligible = graph_on
5877 && upto.is_none()
5878 && task_mask.is_none()
5879 && from == 0
5880 && !crate::gpu::graph_unsupported();
5881 let mut tail_start = 0usize;
5882 if race_eligible && crate::gpu::graph_race_use_graph(graph_trusted) {
5883 let t_graph = std::time::Instant::now();
5884 let mut lg = Vec::new();
5885 let mut gl = 0usize;
5886 let built = self.try_token_graph_wgpu(hidden, position, &mut lg, &mut gl);
5887 if built.is_none() && !self.o1_active() && self.attn_softcap == 0.0 {
5892 crate::gpu::graph_mark_unsupported();
5893 }
5894 graph_note(built.is_some());
5895 if let Some(hh) = built {
5896 let dur = t_graph.elapsed();
5897 if std::env::var("CMF_GRAPH_PROF").is_ok() {
5898 eprintln!("graph-call: {:.2} ms total", dur.as_secs_f64() * 1000.0);
5899 }
5900 if gl > 0 && gl < self.num_layers {
5901 h = hh;
5907 tail_start = gl;
5908 } else if graph_trusted || !crate::gpu::graph_race_first_token_hopeless(dur) {
5909 if !graph_trusted {
5910 crate::gpu::graph_race_record(true, dur);
5911 }
5912 if !lg.is_empty() {
5913 lg.resize(self.vocab_size, 0.0);
5916 if let Some(c) = self.final_softcap {
5917 for l in lg.iter_mut() {
5918 *l = c * (*l / c).tanh();
5919 }
5920 }
5921 self.graph_logits = Some(lg);
5922 }
5923 return hh;
5924 }
5925 }
5931 }
5932 let span = from > 0 || upto.is_some();
5956 if span && graph_on && task_mask.is_none() && graph_trusted {
5957 let upto_excl = upto.map_or(self.num_layers, |u| u + 1);
5958 let mut lg = Vec::new();
5959 let mut gl = 0usize;
5960 let span_res =
5961 self.try_token_graph_wgpu_span(hidden, position, &mut lg, from, upto_excl, &mut gl);
5962 graph_note(span_res.is_some() && gl == upto_excl - from);
5963 if std::env::var("CMF_GPU_DEBUG").is_ok() {
5964 static SEEN: std::sync::atomic::AtomicU32 =
5968 std::sync::atomic::AtomicU32::new(0);
5969 if SEEN.fetch_add(1, std::sync::atomic::Ordering::Relaxed) < 4 {
5970 eprintln!(
5971 "span graph: covered {gl} of {} layers [{from}..{upto_excl}) res={}",
5972 upto_excl - from,
5973 span_res.is_some()
5974 );
5975 }
5976 }
5977 if let Some(hh) = span_res {
5978 if gl == upto_excl - from {
5979 if !lg.is_empty() {
5980 lg.resize(self.vocab_size, 0.0);
5981 if let Some(c) = self.final_softcap {
5982 for l in lg.iter_mut() {
5983 *l = c * (*l / c).tanh();
5984 }
5985 }
5986 self.graph_logits = Some(lg);
5987 }
5988 crate::gpu::set_layer(-1);
5989 return hh;
5990 }
5991 h = hh;
5993 tail_start = from + gl;
5994 }
5995 }
5996 let t_race_cpu = (race_eligible && !graph_trusted).then(std::time::Instant::now);
5997
5998 #[cfg(target_os = "macos")]
5999 let mut gpu_skip_until = 0usize;
6000 for li in tail_start.max(from)..self.num_layers {
6001 crate::gpu::set_layer(li as i64); if let Some(u) = upto {
6003 if li > u {
6004 break;
6005 }
6006 }
6007 if let Some(mask) = task_mask {
6008 if !mask.layer_alive(li) {
6009 continue; }
6011 }
6012 #[cfg(target_os = "macos")]
6016 {
6017 if li < gpu_skip_until {
6018 continue;
6019 }
6020 if task_mask.is_none() {
6021 let end = self.q1_graph_gpu(li, upto, position, &mut h);
6022 if end > li {
6023 gpu_skip_until = end;
6024 if self.is_loop_end(end - 1) && end < self.num_layers {
6027 h = inference::rms_norm(
6028 &h,
6029 &self.weights.final_norm,
6030 self.rms_eps,
6031 self.norm_style,
6032 );
6033 }
6034 continue;
6035 }
6036 }
6037 }
6038
6039 let lw = &self.weights.layers[self.phys_layer(li)];
6040 if let Ok(tp) = std::env::var("CMF_TRACE_POS") {
6041 if tp.parse::<usize>().ok() == Some(position) {
6042 let n: f32 = h.iter().map(|x| x * x).sum::<f32>().sqrt();
6043 eprintln!(
6044 "TRACE pos {position} layer {li}: |h| = {n:.6} h0 {:.6} h1 {:.6}",
6045 h[0], h[1]
6046 );
6047 }
6048 }
6049 inference::rms_norm_into(
6052 &h,
6053 &lw.input_norm,
6054 self.rms_eps,
6055 self.norm_style,
6056 &mut self.ws.n1,
6057 );
6058
6059 let attn_out = match &lw.attn {
6060 AttnKind::Mla(w) => {
6061 let inv_freq_l = self.layer_inv_freq(li);
6062 let rs = self.layer_rope_scale(li);
6063 let eps = self.rms_eps;
6064 let pool = self.pool.clone();
6065 mla_attention(
6066 w,
6067 &self.ws.n1,
6068 &mut self.kv_cache.layers[li],
6069 position,
6070 &inv_freq_l,
6071 rs,
6072 eps,
6073 pool.as_deref(),
6074 )
6075 }
6076 AttnKind::Linear(w) => {
6077 let cfg = self.vmf_cfg.expect("linear layer without vmf_cfg");
6078 vmf_phase_forward(
6079 &self.ws.n1,
6080 w,
6081 &cfg,
6082 &mut self.kv_cache.layers[li].linear_state,
6083 self.pool.as_deref(),
6084 )
6085 }
6086 AttnKind::Kda(w) => {
6087 let cfg = self.kda_cfg.expect("kda layer without kda_cfg");
6088 crate::linear_core::kda_forward(
6089 &self.ws.n1,
6090 w,
6091 &cfg,
6092 &mut self.kv_cache.layers[li].linear_state,
6093 self.pool.as_deref(),
6094 )
6095 }
6096 AttnKind::LinearGdn(w) => {
6097 let cfg = self.gdn_cfg.expect("gdn layer without gdn_cfg");
6098 gdn_forward(
6099 &self.ws.n1,
6100 w,
6101 &cfg,
6102 &mut self.kv_cache.layers[li].linear_state,
6103 self.pool.as_deref(),
6104 )
6105 }
6106 AttnKind::ShortConv(w) => {
6107 let cfg = self
6108 .short_conv_cfg
6109 .expect("short-conv layer without short_conv_cfg");
6110 short_conv_forward(
6111 &self.ws.n1,
6112 w,
6113 &cfg,
6114 &mut self.kv_cache.layers[li].linear_state,
6115 self.pool.as_deref(),
6116 )
6117 }
6118 AttnKind::Full {
6119 wq,
6120 wk,
6121 wv,
6122 wo,
6123 q_norm,
6124 k_norm,
6125 output_gate,
6126 softplus_gate,
6127 bias,
6128 } if self.kv_cache.layers[li].o1_sealed() => {
6129 let inv_freq_l = self.layer_inv_freq(li);
6132 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6133 let cfg = QwenAttnCfg {
6134 num_heads: self.layer_num_heads(li),
6135 num_kv_heads: nkv_l,
6136 head_dim: hd_l,
6137 hidden_size: hs,
6138 position,
6139 inv_freq: &inv_freq_l,
6140 rotary_dim: rd_l,
6141 scale: self.attn_scale,
6142 softcap: self.attn_softcap,
6143 window: None,
6144 v_norm: self.attn_v_norm,
6145 q_norm: q_norm.as_deref(),
6146 k_norm: k_norm.as_deref(),
6147 output_gate: *output_gate,
6148 softplus_gate: softplus_gate
6149 .as_ref()
6150 .map(|(gate, per_head)| (gate, *per_head)),
6151 rope_scale: self.layer_rope_scale(li),
6152 bias: bias
6153 .as_ref()
6154 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6155 rms_eps: eps,
6156 norm_style: self.norm_style,
6157 pool: pool.as_deref(),
6158 };
6159 attention::qwen_attention_nystrom(
6160 &self.ws.n1,
6161 wq,
6162 wk,
6163 wv,
6164 wo,
6165 &mut self.kv_cache.layers[li],
6166 &cfg,
6167 )
6168 }
6169 AttnKind::Full {
6170 wq,
6171 wk,
6172 wv,
6173 wo,
6174 q_norm,
6175 k_norm,
6176 output_gate,
6177 softplus_gate,
6178 bias,
6179 } => 'attn: {
6180 if graph_on
6183 && !*output_gate
6184 && softplus_gate.is_none()
6185 && self.attention_heads_per_layer.is_none()
6186 && bias.is_none()
6187 && task_mask.is_none()
6188 {
6189 let inv_freq_l = self.layer_inv_freq(li);
6190 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6191 let gemma = self.norm_style == cortiq_core::NormStyle::Gemma;
6192 if let (Some((gm, qi)), Some((_, ki)), Some((_, vi)), Some((_, oi))) = (
6193 wq.mapped_q1(),
6194 wk.mapped_q1(),
6195 wv.mapped_q1(),
6196 wo.mapped_q1(),
6197 ) {
6198 let gm = gm.clone();
6199 let mut out = vec![0f32; hs];
6200 let cache = &self.kv_cache.layers[li];
6201 if crate::gpu::attn_dropin(
6202 &gm,
6203 self.graph_kv_id,
6204 li,
6205 &self.ws.n1,
6206 qi,
6207 ki,
6208 vi,
6209 oi,
6210 q_norm.as_deref(),
6211 k_norm.as_deref(),
6212 &inv_freq_l,
6213 nh,
6214 nkv_l,
6215 hd_l,
6216 rd_l,
6217 hs,
6218 position,
6219 self.kv_cache.max_seq_len,
6220 gemma,
6221 eps as f32,
6222 cache.k_heads(),
6223 cache.v_heads(),
6224 &mut out,
6225 ) {
6226 break 'attn out;
6227 }
6228 }
6229 }
6230 let masked = task_mask
6231 .map(|m| m.head_flags(li, self.num_heads).iter().any(|&a| !a))
6232 .unwrap_or(false);
6233 let f32_view = (wq.as_f32(), wk.as_f32(), wv.as_f32(), wo.as_f32());
6234 match (masked, f32_view) {
6235 (true, (Some(q), Some(k), Some(v), Some(o))) => {
6238 let active_heads = task_mask.unwrap().head_flags(li, self.num_heads);
6239 attention::multi_head_attention(
6240 &self.ws.n1,
6241 q,
6242 k,
6243 v,
6244 o,
6245 &mut self.kv_cache.layers[li],
6246 self.num_heads,
6247 self.num_kv_heads,
6248 self.head_dim,
6249 self.hidden_size,
6250 position,
6251 &active_heads,
6252 &self.inv_freq,
6253 )
6254 }
6255 (masked, _) => {
6256 if masked {
6257 tracing::warn!(
6258 "layer {li}: head mask on quantized weights not \
6259 supported yet — executing dense"
6260 );
6261 }
6262 let inv_freq_l = self.layer_inv_freq(li);
6263 let (nkv_l, hd_l, rd_l) = self.layer_geom(li);
6264 let cfg = QwenAttnCfg {
6265 num_heads: self.layer_num_heads(li),
6266 num_kv_heads: nkv_l,
6267 head_dim: hd_l,
6268 hidden_size: hs,
6269 position,
6270 inv_freq: &inv_freq_l,
6271 rotary_dim: rd_l,
6272 scale: self.attn_scale,
6273 softcap: self.attn_softcap,
6274 window: self.layer_window(li),
6275 v_norm: self.attn_v_norm,
6276 q_norm: q_norm.as_deref(),
6277 k_norm: k_norm.as_deref(),
6278 output_gate: *output_gate,
6279 softplus_gate: softplus_gate
6280 .as_ref()
6281 .map(|(gate, per_head)| (gate, *per_head)),
6282 rope_scale: self.layer_rope_scale(li),
6283 bias: bias
6284 .as_ref()
6285 .map(|(a, b, c)| (a.as_slice(), b.as_slice(), c.as_slice())),
6286 rms_eps: eps,
6287 norm_style: self.norm_style,
6288 pool: pool.as_deref(),
6289 };
6290 attention::qwen_attention(
6291 &self.ws.n1,
6292 wq,
6293 wk,
6294 wv,
6295 wo,
6296 &mut self.kv_cache.layers[li],
6297 &cfg,
6298 )
6299 }
6300 }
6301 }
6302 };
6303 let attn_out = match &self.weights.layers[self.phys_layer(li)].attn_out_norm {
6306 Some(w) => inference::rms_norm(&attn_out, w, self.rms_eps, self.norm_style),
6307 None => attn_out,
6308 };
6309 let lw = &self.weights.layers[self.phys_layer(li)];
6310 inference::add_rmsnorm_fused_into(
6311 &mut h,
6312 &attn_out,
6313 &lw.post_norm,
6314 self.rms_eps,
6315 self.norm_style,
6316 &mut self.ws.p1,
6317 );
6318 let mut attn_out = attn_out;
6319 attention::recycle_buf(&mut attn_out);
6320 let post_normed = &self.ws.p1;
6321
6322 let ffn_masked = task_mask
6323 .map(|m| m.ffn_active_count(li) < self.intermediate_size)
6324 .unwrap_or(false);
6325 let ffn_out = match (ffn_masked, &lw.ffn) {
6337 (true, FfnKind::Dense(d)) => {
6338 let tm = task_mask.unwrap();
6339 let alive = tm.ffn_active_count(li);
6340 let deep = alive * 2 <= self.intermediate_size;
6341 if deep && d.down_proj.sparse_col_ok() {
6342 let active = tm.ffn_active_indices(li);
6343 sparse_ffn_quant(
6344 d,
6345 post_normed,
6346 &active,
6347 self.hidden_size,
6348 self.pool.as_deref(),
6349 )
6350 } else if deep
6351 && let (Some(g), Some(u), Some(dn)) = (
6352 d.gate_proj.as_f32(),
6353 d.up_proj.as_f32(),
6354 d.down_proj.as_f32(),
6355 )
6356 {
6357 let active = tm.ffn_active_indices(li);
6358 inference::sparse_ffn_forward(
6359 post_normed,
6360 g,
6361 u,
6362 dn,
6363 self.hidden_size,
6364 self.intermediate_size,
6365 &active,
6366 self.pool.as_deref(),
6367 )
6368 } else {
6369 let row = tm.ffn_masks.get(li).map(|v| v.as_slice());
6370 dense_ffn_batch(d, post_normed, 1, self.pool.as_deref(), row)
6371 }
6372 }
6373 (true, FfnKind::Moe(m)) => {
6374 let allowed = task_mask.and_then(|tm| tm.expert_flags(li, m.experts.len()));
6378 ffn_forward(
6379 &lw.ffn,
6380 post_normed,
6381 self.pool.as_deref(),
6382 allowed.as_deref(),
6383 )
6384 }
6385 (true, FfnKind::DenseMoe(dm)) => dense_moe_ffn(
6386 dm,
6387 post_normed,
6388 &h,
6389 self.rms_eps,
6390 self.norm_style,
6391 self.pool.as_deref(),
6392 ),
6393 (false, _) => match &lw.ffn {
6394 FfnKind::DenseMoe(dm) => dense_moe_ffn(
6395 dm,
6396 post_normed,
6397 &h,
6398 self.rms_eps,
6399 self.norm_style,
6400 self.pool.as_deref(),
6401 ),
6402 _ => {
6403 let allowed = match (&lw.ffn, task_mask) {
6404 (FfnKind::Moe(m), Some(tm)) => tm.expert_flags(li, m.experts.len()),
6405 _ => None,
6406 };
6407 ffn_forward(
6408 &lw.ffn,
6409 post_normed,
6410 self.pool.as_deref(),
6411 allowed.as_deref(),
6412 )
6413 }
6414 },
6415 };
6416 let ffn_out = match &self.weights.layers[self.phys_layer(li)].ffn_out_norm {
6417 Some(w) => inference::rms_norm(&ffn_out, w, self.rms_eps, self.norm_style),
6418 None => ffn_out,
6419 };
6420 for (i, &f) in ffn_out.iter().enumerate() {
6421 h[i] += f;
6422 }
6423 let mut ffn_out = ffn_out;
6424 attention::recycle_buf(&mut ffn_out);
6425
6426 if let Some(sc) = self.weights.layers[self.phys_layer(li)].layer_scale {
6428 for v in h.iter_mut() {
6429 *v *= sc;
6430 }
6431 }
6432
6433 if self.is_loop_end(li) && li + 1 < self.num_layers {
6436 h = inference::rms_norm(
6437 &h,
6438 &self.weights.final_norm,
6439 self.rms_eps,
6440 self.norm_style,
6441 );
6442 }
6443
6444 if self.dyn_phi_layer == Some(li) {
6448 self.update_dyn_phi(&h);
6449 }
6450 }
6451 crate::gpu::set_layer(-1); if let Some(t) = t_race_cpu {
6453 crate::gpu::graph_race_record(false, t.elapsed());
6454 }
6455
6456 h
6457 }
6458
6459 fn update_dyn_phi(&mut self, h: &[f32]) {
6462 const A: f32 = 0.2;
6463 if self.dyn_phi_ema.len() != h.len() {
6464 self.dyn_phi_ema = vec![0.0; h.len()];
6465 self.dyn_phi_seen = 0;
6466 }
6467 if self.dyn_phi_seen == 0 {
6468 self.dyn_phi_ema.copy_from_slice(h);
6469 } else {
6470 for (e, &v) in self.dyn_phi_ema.iter_mut().zip(h) {
6471 *e = (1.0 - A) * *e + A * v;
6472 }
6473 }
6474 self.dyn_phi_seen += 1;
6475 }
6476
6477 pub fn dyn_phi(&self) -> &[f32] {
6479 &self.dyn_phi_ema
6480 }
6481
6482 pub fn set_dyn_phi_layer(&mut self, layer: Option<usize>) {
6484 self.dyn_phi_layer = layer;
6485 self.dyn_phi_ema.clear();
6486 self.dyn_phi_seen = 0;
6487 }
6488
6489 pub fn dynamic_skills(&self) -> Vec<(usize, String, usize)> {
6491 let Some(model) = &self.model else {
6492 return Vec::new();
6493 };
6494 model
6495 .header
6496 .skills
6497 .iter()
6498 .enumerate()
6499 .filter_map(|(i, sk)| {
6500 let ok = matches!(self.dyn_skill_layers.get(i), Some(Some(_)));
6501 let sel = sk.selection.as_ref()?;
6502 (ok).then(|| (i, sk.id.clone(), sel.phi_layer))
6503 })
6504 .collect()
6505 }
6506
6507 pub fn active_skill(&self) -> Option<usize> {
6509 self.dyn_active
6510 }
6511
6512 pub fn enable_dynamic_routing(&mut self) -> usize {
6517 use crate::swarm::{DynRouter, RoutableSkill};
6518 let Some(model) = self.model.clone() else {
6519 return 0;
6520 };
6521 if self.dyn_blend_loaded {
6524 tracing::warn!("dynamic routing unavailable on a blend-loaded pipeline");
6525 return 0;
6526 }
6527 if let Some(a) = self.dyn_active {
6531 if !matches!(self.dyn_skill_layers.get(a), Some(Some(_))) {
6532 tracing::warn!("loaded skill is not FFN-eligible — dynamic routing unavailable");
6533 return 0;
6534 }
6535 }
6536 let hidden = self.hidden_size;
6537 let mut skills = Vec::new();
6538 for (idx, id, _phi) in self.dynamic_skills() {
6539 if let Some(sel) = model.header.skills[idx].selection.as_ref() {
6540 if let Some(rs) = RoutableSkill::from_descriptor(idx, id, sel, hidden) {
6541 skills.push(rs);
6542 }
6543 }
6544 }
6545 if skills.is_empty() {
6546 return 0;
6547 }
6548 let phi = skills[0].phi_layer;
6550 if skills.iter().any(|s| s.phi_layer != phi) {
6551 tracing::warn!("routable skills disagree on phi_layer; using {phi}");
6552 }
6553 let n = skills.len();
6554 self.set_dyn_phi_layer(Some(phi));
6555 self.dyn_router = Some(DynRouter::new(skills));
6556 n
6557 }
6558
6559 pub fn route_switches(&self) -> Vec<(usize, Option<String>, Option<String>)> {
6561 self.dyn_router
6562 .as_ref()
6563 .map(|r| r.switches.clone())
6564 .unwrap_or_default()
6565 }
6566
6567 fn lm_head_forward(&self, hidden: &[f32]) -> Vec<f32> {
6570 let rows = self.weights.lm_head.rows();
6571 let mut logits = attention::take_buf(rows.min(self.vocab_size));
6572 self.weights
6573 .lm_head
6574 .matvec(hidden, &mut logits, self.pool.as_deref());
6575 logits.resize(self.vocab_size, 0.0);
6576 if let Some(m) = self.logit_multiplier {
6577 for l in logits.iter_mut() {
6578 *l *= m;
6579 }
6580 }
6581 if let Some(c) = self.final_softcap {
6582 for l in logits.iter_mut() {
6583 *l = c * (*l / c).tanh();
6584 }
6585 }
6586 logits
6587 }
6588
6589 pub fn prefill_next_logits(&mut self, ids: &[u32], task_mask: Option<&TaskMask>) -> Vec<f32> {
6594 self.kv_cache.clear();
6595 self.kv_history.clear();
6596 let mut hidden = vec![0.0f32; self.hidden_size];
6597 for (pos, &id) in ids.iter().enumerate() {
6598 let emb = self.embed_single(id);
6599 hidden = self.forward_layers(&emb, pos, task_mask);
6600 }
6601 inference::rms_norm_into(
6602 &hidden,
6603 &self.weights.final_norm,
6604 self.rms_eps,
6605 self.norm_style,
6606 &mut self.ws.n1,
6607 );
6608 self.lm_head_forward(&self.ws.n1)
6609 }
6610}
6611
6612pub fn create_test_pipeline(
6614 hidden_size: usize,
6615 intermediate_size: usize,
6616 num_heads: usize,
6617 num_kv_heads: usize,
6618 head_dim: usize,
6619 num_layers: usize,
6620 vocab_size: usize,
6621) -> Pipeline {
6622 let synth = |n: usize, salt: usize| -> Vec<f32> {
6625 (0..n)
6626 .map(|i| (((i * 31 + salt * 17 + 7) % 97) as f32 / 97.0 - 0.5) * 0.2)
6627 .collect()
6628 };
6629 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
6630 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
6631 };
6632 let layer_weights: Vec<LayerWeights> = (0..num_layers)
6633 .map(|li| LayerWeights {
6634 input_norm: vec![1.0; hidden_size],
6635 post_norm: vec![1.0; hidden_size],
6636 attn_out_norm: None,
6637 ffn_out_norm: None,
6638 layer_scale: None,
6639 ffn: FfnKind::Dense(DenseFfn {
6640 gate_proj: qt(intermediate_size, hidden_size, li * 10 + 5),
6641 up_proj: qt(intermediate_size, hidden_size, li * 10 + 6),
6642 down_proj: qt(hidden_size, intermediate_size, li * 10 + 7),
6643 act: Act::Silu,
6644 }),
6645 attn: AttnKind::Full {
6646 bias: None,
6647 wq: qt(num_heads * head_dim, hidden_size, li * 10 + 1),
6648 wk: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 2),
6649 wv: qt(num_kv_heads * head_dim, hidden_size, li * 10 + 3),
6650 wo: qt(hidden_size, num_heads * head_dim, li * 10 + 4),
6651 q_norm: None,
6652 k_norm: None,
6653 output_gate: false,
6654 softplus_gate: None,
6655 },
6656 })
6657 .collect();
6658
6659 Pipeline::new(
6660 Tokenizer::byte_level(),
6661 PipelineWeights {
6662 embed_tokens: qt(vocab_size, hidden_size, 100),
6663 layers: layer_weights,
6664 lm_head: qt(vocab_size, hidden_size, 200),
6665 final_norm: vec![1.0; hidden_size],
6666 },
6667 hidden_size,
6668 intermediate_size,
6669 num_heads,
6670 num_kv_heads,
6671 head_dim,
6672 num_layers,
6673 num_layers, false, vocab_size,
6676 1e-6,
6677 10_000.0,
6678 NormStyle::Qwen,
6679 4096,
6680 SamplerConfig {
6681 seed: Some(42),
6682 ..Default::default()
6683 },
6684 )
6685}
6686
6687#[inline]
6692fn mask_bit(row: &[u8], j: usize) -> bool {
6693 (row.get(j >> 3).copied().unwrap_or(0) >> (j & 7)) & 1 != 0
6694}
6695
6696fn zero_masked_cols(g: &mut [f32], rows: usize, inter: usize, row: &[u8]) {
6702 for r in 0..rows {
6703 let base = r * inter;
6704 for (bi, &byte) in row.iter().enumerate() {
6705 if byte == 0xFF {
6706 continue;
6707 }
6708 let j0 = bi * 8;
6709 for bit in 0..8 {
6710 let j = j0 + bit;
6711 if j < inter && byte & (1 << bit) == 0 {
6712 g[base + j] = 0.0;
6713 }
6714 }
6715 }
6716 }
6717}
6718
6719fn dense_ffn_batch(
6720 d: &DenseFfn,
6721 xs: &[f32],
6722 b: usize,
6723 pool: Option<&Pool>,
6724 mask_row: Option<&[u8]>,
6725) -> Vec<f32> {
6726 let inter = d.gate_proj.rows();
6727 let hidden = d.down_proj.rows();
6728 if mask_row.is_none()
6736 && d.act == Act::Silu
6737 && b >= 32
6738 && crate::gpu::enabled_here()
6739 && !crate::gpu::mm_killed()
6740 {
6741 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6742 d.gate_proj.mapped_q4t(),
6743 d.up_proj.mapped_q4t(),
6744 d.down_proj.mapped_q4t(),
6745 ) {
6746 let mut out = vec![0.0f32; b * hidden];
6747 if crate::gpu::q4t_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6748 return out;
6749 }
6750 }
6751 if let (Some((model, w1)), Some((_, w3)), Some((_, w2))) = (
6756 d.gate_proj.mapped_q4tp(),
6757 d.up_proj.mapped_q4tp(),
6758 d.down_proj.mapped_q4tp(),
6759 ) {
6760 let mut out = vec![0.0f32; b * hidden];
6761 if crate::gpu::q4tp_ffn(model, w1, w3, w2, xs, b, hidden, inter, &mut out) {
6762 return out;
6763 }
6764 }
6765 }
6766 let mut g = vec![0.0f32; b * inter];
6767 d.gate_proj.matmat(xs, b, &mut g, pool);
6768 let mut u = vec![0.0f32; b * inter];
6769 d.up_proj.matmat(xs, b, &mut u, pool);
6770 for i in 0..b * inter {
6771 g[i] = d.act.combine(g[i], u[i]);
6772 }
6773 if let Some(row) = mask_row {
6774 zero_masked_cols(&mut g, b, inter, row);
6775 }
6776 let mut out = vec![0.0f32; b * hidden];
6777 d.down_proj.matmat(&g, b, &mut out, pool);
6778 out
6779}
6780
6781fn accumulate_act(m: &MoeFfn, xs: &[f32], b: usize) {
6786 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6787 static DUMP: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
6788 let on = *ON.get_or_init(|| std::env::var("CMF_RMS_TRACE").is_ok());
6789 let dump = *DUMP.get_or_init(|| std::env::var("CMF_ACT_DUMP").is_ok());
6790 if (!on && !dump) || b == 0 {
6791 return;
6792 }
6793 let hidden = xs.len() / b;
6794 if on {
6795 let mut acc = m.act_sq.borrow_mut();
6796 if acc.len() < hidden {
6797 acc.resize(hidden, 0.0);
6798 }
6799 for t in 0..b {
6800 let row = &xs[t * hidden..(t + 1) * hidden];
6801 for (a, &v) in acc.iter_mut().zip(row) {
6802 *a += (v as f64) * (v as f64);
6803 }
6804 }
6805 }
6806 if dump {
6807 let cap: usize = std::env::var("CMF_ACT_DUMP_ROWS")
6810 .ok()
6811 .and_then(|v| v.parse().ok())
6812 .unwrap_or(4096);
6813 let mut rows = m.act_rows.borrow_mut();
6814 if rows.len() < cap * hidden {
6815 let take = b.min((cap * hidden - rows.len()) / hidden.max(1));
6816 rows.extend_from_slice(&xs[..take * hidden]);
6817 }
6818 }
6819}
6820
6821#[derive(Clone, Copy)]
6824struct SendVecs(*mut Vec<f32>);
6825unsafe impl Send for SendVecs {}
6826unsafe impl Sync for SendVecs {}
6827impl SendVecs {
6828 #[inline]
6829 fn at(self, i: usize) -> *mut Vec<f32> {
6830 unsafe { self.0.add(i) }
6831 }
6832}
6833
6834fn moe_ffn_batch(
6835 m: &MoeFfn,
6836 xs: &[f32],
6837 b: usize,
6838 hidden: usize,
6839 pool: Option<&Pool>,
6840 allowed: Option<&[bool]>,
6841) -> Vec<f32> {
6842 accumulate_act(m, xs, b);
6843 let ne = m.experts.len();
6844 let mut logits = vec![0.0f32; b * ne];
6845 m.router.matmat(xs, b, &mut logits, pool);
6846
6847 let mut assign: Vec<Vec<(usize, f32)>> = vec![Vec::new(); ne];
6850 {
6851 let mut st = m.stats.borrow_mut();
6852 if st.len() < ne {
6853 st.resize(ne, 0);
6854 }
6855 for bi in 0..b {
6856 let (idx, p, wsum) = moe_route(&logits[bi * ne..(bi + 1) * ne], m, allowed);
6857 for &e in &idx {
6858 st[e] += 1;
6859 assign[e].push((bi, p[e] / wsum));
6860 }
6861 }
6862 }
6863
6864 let mut out = vec![0.0f32; b * hidden];
6865 let cols = m.experts[0].gate_proj.cols();
6866 let run_expert = |d: &DenseFfn, list: &[(usize, f32)], out: &mut [f32]| {
6867 let sb = list.len();
6868 let mut sub = vec![0.0f32; sb * cols];
6869 for (k, &(bi, _)) in list.iter().enumerate() {
6870 sub[k * cols..(k + 1) * cols].copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6871 }
6872 let eo = dense_ffn_batch(d, &sub, sb, pool, None);
6873 for (k, &(bi, w)) in list.iter().enumerate() {
6874 for i in 0..hidden {
6875 out[bi * hidden + i] += w * eo[k * hidden + i];
6876 }
6877 }
6878 };
6879 let active: Vec<usize> = (0..ne).filter(|&e| !assign[e].is_empty()).collect();
6885 if pool.is_some() && active.len() >= 8 {
6886 let mut panels: Vec<Vec<f32>> = vec![Vec::new(); active.len()];
6887 {
6888 let panel_ptr = SendVecs(panels.as_mut_ptr());
6889 let experts = &m.experts;
6892 let (active_r, assign_r) = (&active, &assign);
6893 let run = |start: usize, end: usize| {
6894 for ai in start..end {
6895 let e = active_r[ai];
6896 let list = &assign_r[e];
6897 let sb = list.len();
6898 let mut sub = vec![0.0f32; sb * cols];
6899 for (k, &(bi, _)) in list.iter().enumerate() {
6900 sub[k * cols..(k + 1) * cols]
6901 .copy_from_slice(&xs[bi * cols..(bi + 1) * cols]);
6902 }
6903 unsafe {
6905 *panel_ptr.at(ai) =
6906 dense_ffn_batch(&experts[e], &sub, sb, None, None);
6907 }
6908 }
6909 };
6910 match pool {
6911 Some(p) => p.run_rows(active.len(), &run),
6912 None => run(0, active.len()),
6913 }
6914 }
6915 for (ai, &e) in active.iter().enumerate() {
6916 for (k, &(bi, w)) in assign[e].iter().enumerate() {
6917 let eo = &panels[ai][k * hidden..(k + 1) * hidden];
6918 for i in 0..hidden {
6919 out[bi * hidden + i] += w * eo[i];
6920 }
6921 }
6922 }
6923 } else {
6924 for &e in &active {
6925 run_expert(&m.experts[e], &assign[e], &mut out);
6926 }
6927 }
6928 if let Some((se, gate)) = &m.shared {
6929 let all: Vec<(usize, f32)> = if let Some(gate) = gate {
6930 let mut gl = vec![0.0f32; b];
6931 gate.matmat(xs, b, &mut gl, pool);
6932 (0..b)
6933 .map(|bi| (bi, 1.0 / (1.0 + (-gl[bi]).exp())))
6934 .collect()
6935 } else {
6936 (0..b).map(|bi| (bi, 1.0)).collect()
6937 };
6938 run_expert(se, &all, &mut out);
6939 }
6940 out
6941}
6942
6943thread_local! {
6944 static FFN_SCRATCH: std::cell::RefCell<[Vec<f32>; 4]> =
6948 const { std::cell::RefCell::new([Vec::new(), Vec::new(), Vec::new(), Vec::new()]) };
6949}
6950
6951fn dense_ffn(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6953 if crate::gpu::enabled_here()
6964 && (d.gate_proj.rows() >= crate::gpu::min_rows() || d.gate_proj.is_q1())
6965 {
6966 let arm = if d.gate_proj.is_q1() && crate::gpu::q1_force() {
6967 crate::gpu::ProbeArm::Gpu
6968 } else {
6969 crate::gpu::probe_arm(crate::gpu::OpClass::Ffn)
6970 };
6971 match arm {
6972 crate::gpu::ProbeArm::Gpu => {
6973 let t0 = std::time::Instant::now();
6974 if let Some(out) = dense_ffn_gpu(d, x, pool) {
6975 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
6976 return out;
6977 }
6978 }
6979 crate::gpu::ProbeArm::CpuTimed => {
6980 let t0 = std::time::Instant::now();
6981 let out = crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6982 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
6983 return out;
6984 }
6985 crate::gpu::ProbeArm::Cpu => {
6986 return crate::gpu::cpu_scope(|| dense_ffn_cpu(d, x, pool));
6987 }
6988 }
6989 }
6990 dense_ffn_cpu(d, x, pool)
6991}
6992
6993fn dense_ffn_cpu(d: &DenseFfn, x: &[f32], pool: Option<&Pool>) -> Vec<f32> {
6995 let inter = d.gate_proj.rows();
6996 FFN_SCRATCH.with(|s| {
6997 let mut s = s.borrow_mut();
6998 let [g, u, ..] = &mut *s;
6999 g.resize(inter, 0.0);
7000 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
7003 } else {
7005 u.resize(inter, 0.0);
7006 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
7008 for i in 0..inter {
7009 g[i] = d.act.combine(g[i], u[i]);
7010 }
7011 }
7012 FFN_PROBE.with(|pr| {
7015 if let Some(acc) = pr.borrow_mut().as_mut() {
7016 let li = crate::gpu::cur_layer();
7017 if li >= 0 {
7018 if let Some(row) = acc.get_mut(li as usize) {
7019 for (a, &v) in row.iter_mut().zip(g.iter()) {
7020 *a += (v as f64).abs();
7021 }
7022 }
7023 }
7024 }
7025 });
7026 let mut out = attention::take_buf(d.down_proj.rows());
7027 d.down_proj.matvec(g, &mut out, pool);
7028 out
7029 })
7030}
7031
7032thread_local! {
7033 static FFN_PROBE: std::cell::RefCell<Option<Vec<Vec<f64>>>> =
7036 const { std::cell::RefCell::new(None) };
7037}
7038
7039fn dense_ffn_masked(
7044 d: &DenseFfn,
7045 x: &[f32],
7046 pool: Option<&Pool>,
7047 mask_row: &[u8],
7048) -> Vec<f32> {
7049 let inter = d.gate_proj.rows();
7050 FFN_SCRATCH.with(|s| {
7051 let mut s = s.borrow_mut();
7052 let [g, u, ..] = &mut *s;
7053 g.resize(inter, 0.0);
7054 if d.act == Act::Silu && QTensor::matvec_silu_mul(&d.gate_proj, &d.up_proj, x, g, pool) {
7055 } else {
7057 u.resize(inter, 0.0);
7058 QTensor::matvec_many([&d.gate_proj, &d.up_proj], x, [g, u], pool);
7059 for i in 0..inter {
7060 g[i] = d.act.combine(g[i], u[i]);
7061 }
7062 }
7063 zero_masked_cols(g, 1, inter, mask_row);
7064 let mut out = attention::take_buf(d.down_proj.rows());
7065 d.down_proj.matvec(g, &mut out, pool);
7066 out
7067 })
7068}
7069
7070fn dense_ffn_gpu(d: &DenseFfn, x: &[f32], _pool: Option<&Pool>) -> Option<Vec<f32>> {
7076 if d.act != Act::Silu {
7078 return None;
7079 }
7080 if d.gate_proj.rows() < crate::gpu::min_rows() && !d.gate_proj.is_q1() {
7083 return None;
7084 }
7085 let mut jobs: Vec<crate::gpu::MoeJob> = Vec::with_capacity(1);
7086 let mut model_ref = None;
7087 moe_push_job(d, x, 1.0, &mut jobs, &mut model_ref)?;
7088 let model = model_ref?;
7089 let hidden = jobs[0].down.1;
7090 let mut out = attention::take_buf(hidden);
7091 if crate::gpu::moe_block(&model, &jobs, &mut out) {
7092 Some(out)
7093 } else {
7094 let mut out = out;
7095 attention::recycle_buf(&mut out);
7096 None
7097 }
7098}
7099
7100#[allow(clippy::type_complexity)]
7105#[allow(clippy::type_complexity)]
7106pub(crate) fn moe_parts(
7107 t: &QTensor,
7108) -> Option<(
7109 &std::sync::Arc<cortiq_core::CmfModel>,
7110 usize,
7111 usize,
7112 usize,
7113 &[f32],
7114 &[f32],
7115 bool,
7116 bool,
7117 bool,
7118)> {
7119 match t {
7120 QTensor::Mapped {
7121 model,
7122 idx,
7123 dtype: dt @ (cortiq_core::TensorDtype::Q8_2f | cortiq_core::TensorDtype::Q8Row),
7124 rows,
7125 cols,
7126 row_scale,
7127 col_field,
7128 ..
7129 } if (*dt == cortiq_core::TensorDtype::Q8Row) || !col_field.is_empty() => Some((
7130 model, *idx, *rows, *cols, row_scale, col_field, false, false, false,
7131 )),
7132 QTensor::Mapped {
7134 model,
7135 idx,
7136 dtype: cortiq_core::TensorDtype::Q1,
7137 rows,
7138 cols,
7139 ..
7140 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], true, false, false)),
7141 QTensor::Mapped {
7143 model,
7144 idx,
7145 dtype: cortiq_core::TensorDtype::Q4Tiled,
7146 rows,
7147 cols,
7148 ..
7149 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7150 QTensor::Mapped {
7152 model,
7153 idx,
7154 dtype: cortiq_core::TensorDtype::Q4TiledP,
7155 rows,
7156 cols,
7157 ..
7158 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, false)),
7159 QTensor::Mapped {
7163 model,
7164 idx,
7165 dtype: cortiq_core::TensorDtype::Q2TiledP,
7166 rows,
7167 cols,
7168 ..
7169 } => Some((model, *idx, *rows, *cols, &[][..], &[][..], false, true, true)),
7170 _ => None,
7171 }
7172}
7173
7174#[cfg(target_os = "macos")]
7180fn metal_moe_graph_parts(m: &MoeFfn, hidden: usize) -> Option<crate::gpu::GpuMoe<'_>> {
7181 if m.router_sigmoid
7182 || m.router_input_norm
7183 || m.expert_bias.is_some()
7184 || m.route_tau.is_some()
7185 || m.mask.is_some()
7186 || m.per_expert_scale.is_some()
7187 || m.experts.is_empty()
7188 || m.top_k == 0
7189 {
7190 return None;
7191 }
7192 let (sh, sg) = match &m.shared {
7195 Some((sh, Some(sg))) => (sh, sg),
7196 _ => return None,
7197 };
7198 let (rf, rr, rc) = m.router.f32_parts()?;
7199 if rr != m.experts.len() || rc != hidden {
7200 return None;
7201 }
7202 let (sf, sr, sc) = sg.f32_parts()?;
7203 if sr * sc != hidden {
7204 return None;
7205 }
7206 let inter = m.experts[0].gate_proj.rows();
7207 let gu_q2 = m.experts[0].gate_proj.mapped_q2tp().is_some();
7210 let trio = |e: &DenseFfn| -> Option<(usize, usize, usize)> {
7211 if e.act != Act::Silu
7212 || e.gate_proj.rows() != inter
7213 || e.gate_proj.cols() != hidden
7214 || e.up_proj.rows() != inter
7215 || e.up_proj.cols() != hidden
7216 || e.down_proj.rows() != hidden
7217 || e.down_proj.cols() != inter
7218 {
7219 return None;
7220 }
7221 let pick = |t: &QTensor| -> Option<usize> {
7222 if gu_q2 {
7223 t.mapped_q2tp().map(|(_, i)| i)
7224 } else {
7225 t.mapped_q4tp().map(|(_, i)| i)
7226 }
7227 };
7228 Some((
7229 pick(&e.gate_proj)?,
7230 pick(&e.up_proj)?,
7231 e.down_proj.mapped_q4tp().map(|(_, i)| i)?,
7232 ))
7233 };
7234 let experts = m
7235 .experts
7236 .iter()
7237 .map(trio)
7238 .collect::<Option<Vec<_>>>()?;
7239 let shared = trio(sh)?;
7240 Some(crate::gpu::GpuMoe {
7241 router: rf,
7242 sgate: sf,
7243 experts,
7244 shared,
7245 n_exp: m.experts.len(),
7246 top_k: m.top_k,
7247 inter,
7248 norm_topk: m.norm_topk_prob,
7249 route_scale: m.routed_scaling,
7250 gu_q2,
7251 })
7252}
7253
7254pub(crate) fn moe_push_job_parts<'a>(
7258 gate: &'a QTensor,
7259 up: &'a QTensor,
7260 down: &'a QTensor,
7261 x: &[f32],
7262 w: f32,
7263 swiglu_limit: f32,
7264 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7265 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7266) -> Option<()> {
7267 use crate::qtensor::prescale;
7268 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(gate)?;
7269 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(up)?;
7270 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(down)?;
7271 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7272 return None; }
7274 if gq2 && (dq2 || !dq4 || down.mapped_q4tp().is_none()) {
7277 return None;
7278 }
7279 if !gq2 && dq2 {
7280 return None;
7281 }
7282 model_ref.get_or_insert_with(|| gm.clone());
7283 let dt = |cf: &[f32]| {
7284 if cf.is_empty() {
7285 cortiq_core::TensorDtype::Q8Row
7286 } else {
7287 cortiq_core::TensorDtype::Q8_2f
7288 }
7289 };
7290 jobs.push(crate::gpu::MoeJob {
7291 gate: (gi, gr, gc, grs),
7292 up: (ui, ur, uc, urs),
7293 down: (di, dr, dc, drs),
7294 xs_gate: prescale(x, gcf, dt(gcf)).into_owned(),
7295 xs_up: prescale(x, ucf, dt(ucf)).into_owned(),
7296 down_col: dcf,
7297 w,
7298 q1: gq1,
7299 q4t: gq4 && !gq2 && gate.mapped_q4tp().is_none(),
7300 q4tp: gq4 && (gq2 || gate.mapped_q4tp().is_some()),
7301 gu_q2: gq2,
7302 swiglu_limit,
7303 });
7304 Some(())
7305}
7306
7307fn moe_push_job<'a>(
7309 d: &'a DenseFfn,
7310 x: &[f32],
7311 w: f32,
7312 jobs: &mut Vec<crate::gpu::MoeJob<'a>>,
7313 model_ref: &mut Option<std::sync::Arc<cortiq_core::CmfModel>>,
7314) -> Option<()> {
7315 use crate::qtensor::prescale;
7316 if d.act != Act::Silu {
7317 return None; }
7319 let (gm, gi, gr, gc, grs, gcf, gq1, gq4, gq2) = moe_parts(&d.gate_proj)?;
7320 let (_, ui, ur, uc, urs, ucf, uq1, uq4, uq2) = moe_parts(&d.up_proj)?;
7321 let (_, di, dr, dc, drs, dcf, dq1, dq4, dq2) = moe_parts(&d.down_proj)?;
7322 if gq1 != uq1 || uq1 != dq1 || gq4 != uq4 || uq4 != dq4 || gq2 != uq2 {
7323 return None; }
7325 if gq2 && (dq2 || !dq4 || d.down_proj.mapped_q4tp().is_none()) {
7326 return None;
7327 }
7328 if !gq2 && dq2 {
7329 return None;
7330 }
7331 model_ref.get_or_insert_with(|| gm.clone());
7332 let gdt = if gcf.is_empty() {
7333 cortiq_core::TensorDtype::Q8Row
7334 } else {
7335 cortiq_core::TensorDtype::Q8_2f
7336 };
7337 let udt = if ucf.is_empty() {
7338 cortiq_core::TensorDtype::Q8Row
7339 } else {
7340 cortiq_core::TensorDtype::Q8_2f
7341 };
7342 jobs.push(crate::gpu::MoeJob {
7343 gate: (gi, gr, gc, grs),
7344 up: (ui, ur, uc, urs),
7345 down: (di, dr, dc, drs),
7346 xs_gate: prescale(x, gcf, gdt).into_owned(),
7347 xs_up: prescale(x, ucf, udt).into_owned(),
7348 down_col: dcf,
7349 w,
7350 q1: gq1,
7351 q4t: gq4 && !gq2 && d.gate_proj.mapped_q4tp().is_none(),
7352 q4tp: gq4 && (gq2 || d.gate_proj.mapped_q4tp().is_some()),
7353 gu_q2: gq2,
7354 swiglu_limit: 0.0,
7355 });
7356 Some(())
7357}
7358
7359fn sparse_ffn_quant(
7366 d: &DenseFfn,
7367 x: &[f32],
7368 active: &[u16],
7369 hidden: usize,
7370 pool: Option<&Pool>,
7371) -> Vec<f32> {
7372 let n = active.len();
7373 let inter = d.gate_proj.rows();
7374 let mut act = vec![0.0f32; n];
7375 let need_scratch = !(d.gate_proj.sparse_col_ok() && d.up_proj.sparse_col_ok());
7378 let compute = |ai: usize| -> f32 {
7379 let idx = active[ai] as usize;
7380 if idx >= inter {
7381 return 0.0; }
7383 let mut s = if need_scratch {
7384 vec![0.0f32; hidden]
7385 } else {
7386 Vec::new()
7387 };
7388 let gate = d.gate_proj.row_dot(idx, x, &mut s);
7389 let up = d.up_proj.row_dot(idx, x, &mut s);
7390 d.act.combine(gate, up)
7391 };
7392 match pool {
7393 Some(p) if n >= 256 => {
7394 let ptr = SendMut(act.as_mut_ptr());
7395 p.run(&|widx, nw| {
7396 let chunk = n.div_ceil(nw);
7397 let (s, e) = (widx * chunk, ((widx + 1) * chunk).min(n));
7398 for ai in s..e {
7399 unsafe { *ptr.at(ai) = compute(ai) };
7400 }
7401 });
7402 }
7403 _ => {
7404 for (ai, a) in act.iter_mut().enumerate() {
7405 *a = compute(ai);
7406 }
7407 }
7408 }
7409 let mut out = vec![0.0f32; hidden];
7411 for (ai, &idx) in active.iter().enumerate() {
7412 let w = act[ai];
7413 if w.abs() >= 1e-12 && (idx as usize) < inter {
7414 d.down_proj.add_col_scaled(idx as usize, w, &mut out);
7415 }
7416 }
7417 out
7418}
7419
7420#[doc(hidden)]
7422pub fn sparse_ffn_quant_for_test(
7423 d: &DenseFfn,
7424 x: &[f32],
7425 active: &[u16],
7426 hidden: usize,
7427) -> Vec<f32> {
7428 sparse_ffn_quant(d, x, active, hidden, None)
7429}
7430
7431fn dequant_dense_f32(d: &DenseFfn) -> (Vec<f32>, Vec<f32>, Vec<f32>) {
7435 let deq = |t: &QTensor| -> Vec<f32> {
7436 let (rows, cols) = (t.rows(), t.cols());
7437 let mut out = vec![0.0f32; rows * cols];
7438 for r in 0..rows {
7439 t.row_f32(r, &mut out[r * cols..(r + 1) * cols]);
7440 }
7441 out
7442 };
7443 (deq(&d.gate_proj), deq(&d.up_proj), deq(&d.down_proj))
7444}
7445
7446struct SendMut(*mut f32);
7448unsafe impl Send for SendMut {}
7449unsafe impl Sync for SendMut {}
7450impl SendMut {
7451 #[inline]
7452 #[allow(clippy::mut_from_ref)]
7455 unsafe fn at(&self, i: usize) -> &mut f32 {
7456 unsafe { &mut *self.0.add(i) }
7457 }
7458}
7459
7460fn moe_route(logits: &[f32], m: &MoeFfn, allowed: Option<&[bool]>) -> (Vec<usize>, Vec<f32>, f32) {
7470 let ne = logits.len();
7471 let p: Vec<f32> = if m.router_sigmoid {
7472 logits.iter().map(|&l| 1.0 / (1.0 + (-l).exp())).collect()
7473 } else {
7474 let mx = logits.iter().cloned().fold(f32::NEG_INFINITY, f32::max);
7475 let mut e: Vec<f32> = logits.iter().map(|&l| (l - mx).exp()).collect();
7476 let s: f32 = e.iter().sum();
7477 for v in &mut e {
7478 *v /= s;
7479 }
7480 e
7481 };
7482 let admit = |e: usize| {
7488 m.mask.as_ref().is_none_or(|mk| mk[e])
7489 && allowed.is_none_or(|a| a.get(e).copied().unwrap_or(false))
7490 };
7491 let mut idx: Vec<usize> = (0..ne).filter(|&e| admit(e)).collect();
7492 match &m.expert_bias {
7494 Some(b) => idx.sort_unstable_by(|&x, &y| {
7495 (p[y] + b[y])
7496 .partial_cmp(&(p[x] + b[x]))
7497 .unwrap()
7498 .then(x.cmp(&y))
7499 }),
7500 None => idx.sort_unstable_by(|&x, &y| p[y].partial_cmp(&p[x]).unwrap().then(x.cmp(&y))),
7501 }
7502 idx.truncate(m.top_k);
7503 if let Some(tau) = m.route_tau {
7507 let total: f32 = idx.iter().map(|&e| p[e]).sum();
7508 if total > 0.0 {
7509 let mut acc = 0.0f32;
7510 let mut keep = idx.len();
7511 for (i, &e) in idx.iter().enumerate() {
7512 acc += p[e];
7513 if acc >= tau * total {
7514 keep = i + 1;
7515 break;
7516 }
7517 }
7518 idx.truncate(keep);
7519 }
7520 }
7521 let wsum: f32 = if m.norm_topk_prob {
7522 let s: f32 = idx.iter().map(|&e| p[e]).sum();
7523 (if m.router_sigmoid { s + 1e-6 } else { s }) / m.routed_scaling
7526 } else {
7527 1.0 / m.routed_scaling
7528 };
7529 (idx, p, wsum)
7530}
7531
7532fn moe_ffn(m: &MoeFfn, x: &[f32], pool: Option<&Pool>, allowed: Option<&[bool]>) -> Vec<f32> {
7535 accumulate_act(m, x, 1);
7536 let ne = m.experts.len();
7537 let mut logits = vec![0.0f32; ne];
7538 m.router.matvec(x, &mut logits, pool);
7539 let (idx, p, wsum) = moe_route(&logits, m, allowed);
7540 {
7541 let mut st = m.stats.borrow_mut();
7542 if st.len() < ne {
7543 st.resize(ne, 0);
7544 }
7545 for &e in &idx {
7546 st[e] += 1;
7547 }
7548 }
7549 if crate::gpu::enabled_here() {
7554 match crate::gpu::probe_arm(crate::gpu::OpClass::Ffn) {
7555 crate::gpu::ProbeArm::Gpu => {
7556 let t0 = std::time::Instant::now();
7557 if let Some(out) = moe_ffn_gpu(m, x, &idx, &p, wsum, pool) {
7558 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, true, t0.elapsed());
7559 return out;
7560 }
7561 }
7562 crate::gpu::ProbeArm::CpuTimed => {
7563 let t0 = std::time::Instant::now();
7564 let out = crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7565 crate::gpu::probe_record(crate::gpu::OpClass::Ffn, false, t0.elapsed());
7566 return out;
7567 }
7568 crate::gpu::ProbeArm::Cpu => {
7569 return crate::gpu::cpu_scope(|| moe_ffn_cpu(m, x, &idx, &p, wsum, pool));
7570 }
7571 }
7572 }
7573 moe_ffn_cpu(m, x, &idx, &p, wsum, pool)
7574}
7575
7576fn graph_note(built: bool) {
7580 use std::sync::atomic::{AtomicBool, Ordering};
7581 if built {
7582 GRAPH_TOK_OK.fetch_add(1, Ordering::Relaxed);
7583 } else {
7584 GRAPH_TOK_MISS.fetch_add(1, Ordering::Relaxed);
7585 }
7586 static SAID: AtomicBool = AtomicBool::new(false);
7587 if !SAID.swap(true, Ordering::Relaxed) {
7588 if built {
7589 tracing::info!("wgpu whole-token graph: ACTIVE");
7590 } else {
7591 tracing::warn!("wgpu whole-token graph refused — per-op path");
7592 }
7593 }
7594}
7595
7596pub static GRAPH_TOK_OK: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7600pub static GRAPH_TOK_MISS: std::sync::atomic::AtomicU64 = std::sync::atomic::AtomicU64::new(0);
7601
7602fn moe_batch_enabled() -> bool {
7605 static ON: std::sync::OnceLock<bool> = std::sync::OnceLock::new();
7606 *ON.get_or_init(|| std::env::var("CMF_MOE_BATCH").as_deref() != Ok("0"))
7607}
7608
7609fn moe_ffn_cpu_batched(
7615 m: &MoeFfn,
7616 x: &[f32],
7617 idx: &[usize],
7618 p: &[f32],
7619 wsum: f32,
7620 pool: Option<&Pool>,
7621) -> Option<Vec<f32>> {
7622 if idx.is_empty() || !moe_batch_enabled() {
7623 return None;
7624 }
7625 if FFN_PROBE.with(|pr| pr.borrow().is_some()) {
7629 return None;
7630 }
7631 let n = idx.len() + usize::from(m.shared.is_some());
7632 let mut pairs = Vec::with_capacity(n);
7633 let mut downs = Vec::with_capacity(n);
7634 let mut ws = Vec::with_capacity(n);
7635 for &e in idx {
7636 let d = &m.experts[e];
7637 if d.act != Act::Silu {
7638 return None;
7639 }
7640 pairs.push((&d.gate_proj, &d.up_proj));
7641 downs.push(&d.down_proj);
7642 ws.push(p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]));
7643 }
7644 if let Some((se, gate)) = &m.shared {
7647 if se.act != Act::Silu {
7648 return None;
7649 }
7650 let g = gate.as_ref().map_or(1.0, |gate| {
7651 let mut gl = [0.0f32; 1];
7652 gate.matvec(x, &mut gl, pool);
7653 1.0 / (1.0 + (-gl[0]).exp())
7654 });
7655 pairs.push((&se.gate_proj, &se.up_proj));
7656 downs.push(&se.down_proj);
7657 ws.push(g);
7658 }
7659 let inter = pairs[0].0.rows();
7660 let mut gs: Vec<Vec<f32>> = (0..pairs.len()).map(|_| vec![0f32; inter]).collect();
7661 if !QTensor::moe_gate_up_many(&pairs, x, &mut gs, pool) {
7662 return None;
7663 }
7664 let mut out = attention::take_buf(x.len());
7665 if !QTensor::moe_down_many(&downs, &gs, &ws, &mut out, pool) {
7666 attention::recycle_buf(&mut out);
7667 return None;
7668 }
7669 Some(out)
7670}
7671
7672fn moe_ffn_cpu(
7674 m: &MoeFfn,
7675 x: &[f32],
7676 idx: &[usize],
7677 p: &[f32],
7678 wsum: f32,
7679 pool: Option<&Pool>,
7680) -> Vec<f32> {
7681 if let Some(out) = moe_ffn_cpu_batched(m, x, idx, p, wsum, pool) {
7682 return out;
7683 }
7684 let mut out = attention::take_buf(x.len());
7685 for &e in idx {
7686 let mut eo = dense_ffn(&m.experts[e], x, pool);
7687 let w = p[e] / wsum * m.per_expert_scale.as_ref().map_or(1.0, |v| v[e]);
7688 for i in 0..out.len() {
7689 out[i] += w * eo[i];
7690 }
7691 attention::recycle_buf(&mut eo);
7692 }
7693 if let Some((se, gate)) = &m.shared {
7694 let mut so = dense_ffn(se, x, pool);
7695 let g = gate.as_ref().map_or(1.0, |gate| {
7696 let mut gl = [0.0f32; 1];
7697 gate.matvec(x, &mut gl, pool);
7698 1.0 / (1.0 + (-gl[0]).exp())
7699 });
7700 for i in 0..out.len() {
7701 out[i] += g * so[i];
7702 }
7703 attention::recycle_buf(&mut so);
7704 }
7705 out
7706}
7707
7708#[allow(clippy::too_many_arguments)]
7716fn mla_attention(
7717 w: &MlaWeights,
7718 normed: &[f32],
7719 cache: &mut crate::kv_cache::LayerKvCache,
7720 position: usize,
7721 inv_freq: &[f32],
7722 rope_scale: f32,
7723 eps: f64,
7724 pool: Option<&Pool>,
7725) -> Vec<f32> {
7726 let (nh, dr, dn, dv, lora) = (w.nh, w.qk_rope, w.qk_nope, w.v_dim, w.lora);
7727 let hd = dr + dn;
7728 let mut q = vec![0.0f32; nh * hd];
7729 match (&w.q_a, &w.q_a_norm) {
7730 (Some(qa), Some(qn)) => {
7731 let mut t = vec![0.0f32; qa.rows()];
7732 qa.matvec(normed, &mut t, pool);
7733 let tn = inference::rms_norm(&t, qn, eps, NormStyle::Qwen);
7734 w.q_proj.matvec(&tn, &mut q, pool);
7735 }
7736 _ => w.q_proj.matvec(normed, &mut q, pool),
7737 }
7738 let mut ca = vec![0.0f32; lora + dr];
7739 w.kv_a.matvec(normed, &mut ca, pool);
7740 let (c_lat, k_rope) = ca.split_at_mut(lora);
7741 let latn = inference::rms_norm(c_lat, &w.kv_a_norm, eps, NormStyle::Qwen);
7742 let mut kvb = vec![0.0f32; nh * (dn + dv)];
7743 w.kv_b.matvec(&latn, &mut kvb, pool);
7744 if !w.nope {
7745 attention::rope_rotate_scaled(k_rope, position, inv_freq, rope_scale);
7746 }
7747 for h in 0..nh {
7748 if !w.nope {
7749 attention::rope_rotate_scaled(
7750 &mut q[h * hd..h * hd + dr],
7751 position,
7752 inv_freq,
7753 rope_scale,
7754 );
7755 }
7756 }
7757 let mut k = vec![0.0f32; nh * hd];
7758 let mut v = vec![0.0f32; nh * hd];
7759 for h in 0..nh {
7760 k[h * hd..h * hd + dr].copy_from_slice(k_rope);
7761 k[h * hd + dr..(h + 1) * hd].copy_from_slice(&kvb[h * (dn + dv)..h * (dn + dv) + dn]);
7762 v[h * hd..h * hd + dv].copy_from_slice(&kvb[h * (dn + dv) + dn..(h + 1) * (dn + dv)]);
7763 }
7764 cache.append(&k, &v, &vec![true; nh]);
7765 let (ao, mut imp) = attention::attend_all_heads(&q, cache, nh, 1, hd, w.scale, None, 0.0);
7766 attention::recycle_buf(&mut imp);
7767 let mut ov = vec![0.0f32; nh * dv];
7768 for h in 0..nh {
7769 ov[h * dv..(h + 1) * dv].copy_from_slice(&ao[h * hd..h * hd + dv]);
7770 }
7771 let mut out = vec![0.0f32; w.o_proj.rows()];
7772 w.o_proj.matvec(&ov, &mut out, pool);
7773 out
7774}
7775
7776fn dense_moe_ffn(
7783 dm: &DenseMoeFfn,
7784 x_normed: &[f32],
7785 h_raw: &[f32],
7786 eps: f64,
7787 norm_style: NormStyle,
7788 pool: Option<&Pool>,
7789) -> Vec<f32> {
7790 let mut d = dense_ffn(&dm.dense, x_normed, pool);
7791 d = inference::rms_norm(&d, &dm.post_norm_1, eps, norm_style);
7792 let m = &dm.moe;
7793 let ne = m.experts.len();
7794 let mut logits = vec![0.0f32; ne];
7795 if m.router_input_norm {
7796 let ss: f32 = h_raw.iter().map(|v| v * v).sum::<f32>() / h_raw.len() as f32;
7797 let inv = 1.0 / (ss + eps as f32).sqrt();
7798 let xr: Vec<f32> = h_raw.iter().map(|v| v * inv).collect();
7799 m.router.matvec(&xr, &mut logits, pool);
7800 } else {
7801 m.router.matvec(h_raw, &mut logits, pool);
7802 }
7803 let (idx, p, wsum) = moe_route(&logits, m, None);
7804 {
7805 let mut st = m.stats.borrow_mut();
7806 if st.len() < ne {
7807 st.resize(ne, 0);
7808 }
7809 for &e in &idx {
7810 st[e] += 1;
7811 }
7812 }
7813 let x2 = inference::rms_norm(h_raw, &dm.pre_norm_2, eps, norm_style);
7814 let mo = moe_ffn_cpu(m, &x2, &idx, &p, wsum, pool);
7815 let mo = inference::rms_norm(&mo, &dm.post_norm_2, eps, norm_style);
7816 for (di, mi) in d.iter_mut().zip(&mo) {
7817 *di += mi;
7818 }
7819 d
7820}
7821
7822fn moe_gpu_refused(why: &'static str) {
7829 use std::sync::atomic::{AtomicBool, Ordering};
7830 static SAID: AtomicBool = AtomicBool::new(false);
7831 if !SAID.swap(true, Ordering::Relaxed) {
7832 tracing::warn!("MoE GPU block refused ({why}) — experts run on the CPU");
7833 }
7834}
7835
7836fn moe_ffn_gpu(
7837 m: &MoeFfn,
7838 x: &[f32],
7839 idx: &[usize],
7840 p: &[f32],
7841 wsum: f32,
7842 pool: Option<&Pool>,
7843) -> Option<Vec<f32>> {
7844 use crate::gpu::MoeJob;
7845
7846 let mut jobs: Vec<MoeJob> = Vec::with_capacity(idx.len() + 1);
7847 let mut model_ref = None;
7848 for &e in idx {
7849 if moe_push_job(&m.experts[e], x, p[e] / wsum, &mut jobs, &mut model_ref).is_none() {
7850 moe_gpu_refused("push_job(expert)");
7851 return None;
7852 }
7853 }
7854 if let Some((se, gate)) = &m.shared {
7855 let g = gate.as_ref().map_or(1.0, |gate| {
7856 let mut gl = [0.0f32; 1];
7857 gate.matvec(x, &mut gl, pool);
7858 1.0 / (1.0 + (-gl[0]).exp())
7859 });
7860 if moe_push_job(se, x, g, &mut jobs, &mut model_ref).is_none() {
7861 moe_gpu_refused("push_job(shared)");
7862 return None;
7863 }
7864 }
7865 let Some(model) = model_ref else {
7866 moe_gpu_refused("no model_ref");
7867 return None;
7868 };
7869 let hidden = jobs[0].down.1;
7870 let mut out = vec![0.0f32; hidden];
7871 if crate::gpu::moe_block(&model, &jobs, &mut out) {
7872 Some(out)
7873 } else {
7874 moe_gpu_refused("gpu::moe_block");
7875 None
7876 }
7877}
7878
7879fn ffn_forward(
7881 ffn: &FfnKind,
7882 x: &[f32],
7883 pool: Option<&Pool>,
7884 experts_allowed: Option<&[bool]>,
7885) -> Vec<f32> {
7886 match ffn {
7887 FfnKind::Dense(d) => dense_ffn(d, x, pool),
7888 FfnKind::Moe(m) => moe_ffn(m, x, pool, experts_allowed),
7889 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7893 }
7894}
7895
7896fn ffn_forward_pair(
7900 ffn: &FfnKind,
7901 x1: &[f32],
7902 x2: &[f32],
7903 pool: Option<&Pool>,
7904 experts_allowed: Option<&[bool]>,
7905) -> (Vec<f32>, Vec<f32>) {
7906 let d = match ffn {
7907 FfnKind::Dense(d) => d,
7908 FfnKind::Moe(m) => {
7909 return (
7910 moe_ffn(m, x1, pool, experts_allowed),
7911 moe_ffn(m, x2, pool, experts_allowed),
7912 );
7913 }
7914 FfnKind::DenseMoe(_) => unreachable!("DenseMoe dispatches via dense_moe_ffn"),
7915 };
7916 let inter = d.gate_proj.rows();
7917 FFN_SCRATCH.with(|s| {
7918 let mut s = s.borrow_mut();
7919 let [g1, g2, u1, u2] = &mut *s;
7920 g1.resize(inter, 0.0);
7921 g2.resize(inter, 0.0);
7922 u1.resize(inter, 0.0);
7923 u2.resize(inter, 0.0);
7924 QTensor::matvec2_many(
7927 [&d.gate_proj, &d.up_proj],
7928 x1,
7929 x2,
7930 [g1.as_mut_slice(), u1.as_mut_slice()],
7931 [g2.as_mut_slice(), u2.as_mut_slice()],
7932 pool,
7933 );
7934 for i in 0..inter {
7935 g1[i] = d.act.combine(g1[i], u1[i]);
7936 g2[i] = d.act.combine(g2[i], u2[i]);
7937 }
7938 let mut o1 = attention::take_buf(d.down_proj.rows());
7939 let mut o2 = attention::take_buf(d.down_proj.rows());
7940 d.down_proj.matvec2(g1, g2, &mut o1, &mut o2, pool);
7941 (o1, o2)
7942 })
7943}
7944
7945#[cfg(test)]
7946mod tests {
7947
7948 #[test]
7949 fn cancel_flag_stops_generation() {
7950 let mut p = create_test_pipeline(16, 32, 2, 2, 8, 2, 32);
7951 p.cancel.store(true, std::sync::atomic::Ordering::Relaxed);
7954 let r = p.generate_from_ids(&[1, 2, 3], 8, None, None).unwrap();
7955 assert_eq!(r.finish_reason, "cancelled");
7956 assert!(
7957 r.token_ids.is_empty(),
7958 "no tokens after cancel: {:?}",
7959 r.token_ids
7960 );
7961 let r2 = p.generate_from_ids(&[1, 2, 3], 4, None, None).unwrap();
7963 assert_ne!(r2.finish_reason, "cancelled");
7964 }
7965 use super::*;
7966
7967 #[test]
7973 fn sparse_ffn_quant_equals_dense_with_inactive_zeroed() {
7974 let (hidden, inter) = (16usize, 40usize);
7975 let synth = |n: usize, salt: usize| -> Vec<f32> {
7976 (0..n)
7977 .map(|i| (((i * 37 + salt * 11 + 3) % 101) as f32 / 101.0 - 0.5) * 0.4)
7978 .collect()
7979 };
7980 let d = DenseFfn {
7981 gate_proj: QTensor::from_f32(synth(inter * hidden, 1), inter, hidden),
7982 up_proj: QTensor::from_f32(synth(inter * hidden, 2), inter, hidden),
7983 down_proj: QTensor::from_f32(synth(hidden * inter, 3), hidden, inter),
7984 act: Act::Silu,
7985 };
7986 let x = synth(hidden, 9);
7987 let active: Vec<u16> = (0..inter as u16).filter(|i| i % 3 == 0).collect();
7989
7990 let sparse = sparse_ffn_quant(&d, &x, &active, hidden, None);
7991
7992 let mut g = vec![0.0f32; inter];
7994 d.gate_proj.matvec(&x, &mut g, None);
7995 let mut u = vec![0.0f32; inter];
7996 d.up_proj.matvec(&x, &mut u, None);
7997 let act_set: std::collections::HashSet<u16> = active.iter().copied().collect();
7998 for i in 0..inter {
7999 g[i] = if act_set.contains(&(i as u16)) {
8000 inference::silu(g[i]) * u[i]
8001 } else {
8002 0.0
8003 };
8004 }
8005 let mut reference = vec![0.0f32; hidden];
8006 d.down_proj.matvec(&g, &mut reference, None);
8007
8008 let max_d = sparse
8009 .iter()
8010 .zip(&reference)
8011 .map(|(a, b)| (a - b).abs())
8012 .fold(0.0f32, f32::max);
8013 assert!(max_d < 1e-5, "sparse != dense-zeroed: max|Δ| = {max_d}");
8014 }
8015
8016 fn attach_test_mtp(p: &mut Pipeline) {
8018 let (h, inter, heads, kv, hd) = (
8019 p.hidden_size,
8020 p.intermediate_size,
8021 p.num_heads,
8022 p.num_kv_heads,
8023 p.head_dim,
8024 );
8025 let synth = |n: usize, salt: usize| -> Vec<f32> {
8026 (0..n)
8027 .map(|i| (((i * 29 + salt * 23 + 5) % 101) as f32 / 101.0 - 0.5) * 0.2)
8028 .collect()
8029 };
8030 let qt = |rows: usize, cols: usize, salt: usize| -> QTensor {
8031 QTensor::from_f32(synth(rows * cols, salt), rows, cols)
8032 };
8033 p.mtp = Some(MtpModule {
8034 enorm: vec![1.0; h],
8035 hnorm: vec![1.0; h],
8036 eh_proj: qt(h, 2 * h, 301),
8037 layer: LayerWeights {
8038 input_norm: vec![1.0; h],
8039 post_norm: vec![1.0; h],
8040 attn_out_norm: None,
8041 ffn_out_norm: None,
8042 layer_scale: None,
8043 ffn: FfnKind::Dense(DenseFfn {
8044 gate_proj: qt(inter, h, 315),
8045 up_proj: qt(inter, h, 316),
8046 down_proj: qt(h, inter, 317),
8047 act: Act::Silu,
8048 }),
8049 attn: AttnKind::Full {
8050 bias: None,
8051 wq: qt(heads * hd, h, 311),
8052 wk: qt(kv * hd, h, 312),
8053 wv: qt(kv * hd, h, 313),
8054 wo: qt(h, heads * hd, 314),
8055 q_norm: None,
8056 k_norm: None,
8057 output_gate: false,
8058 softplus_gate: None,
8059 },
8060 },
8061 final_norm: vec![1.0; h],
8062 kv: crate::kv_cache::LayerKvCache::new(kv, hd),
8063 });
8064 }
8065
8066 #[test]
8067 fn speculative_equals_vanilla_greedy() {
8068 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
8072 let run = |spec: bool| {
8073 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8074 p.sampler_config.temperature = 0.0;
8075 attach_test_mtp(&mut p);
8076 p.speculative = spec;
8077 let r = p.generate("abcdef", 12, None, None).unwrap();
8078 (r.token_ids, r.mtp_drafted, r.mtp_accepted)
8079 };
8080 let (vanilla, d0, _) = run(false);
8081 let (spec, d1, a1) = run(true);
8082 assert_eq!(d0, 0, "vanilla path must not draft");
8083 assert!(d1 > 0, "speculative path must draft");
8084 assert_eq!(
8085 vanilla, spec,
8086 "speculative must reproduce the exact greedy sequence (accepted {a1}/{d1})"
8087 );
8088 }
8089
8090 #[test]
8091 fn speculative_accepts_constant_oracle() {
8092 unsafe { std::env::set_var("CMF_GPU_WGPU_GRAPH", "0") };
8094 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8095 p.sampler_config.temperature = 0.0;
8096 p.sampler_config.repetition_penalty = 1.0;
8097 p.weights.lm_head = QTensor::from_f32(vec![0.01; 64 * 8], 64, 8);
8100 attach_test_mtp(&mut p);
8101 p.speculative = true;
8102 let r = p.generate("abcd", 10, None, None).unwrap();
8103 assert!(r.mtp_drafted > 0);
8104 assert_eq!(
8105 r.mtp_accepted, r.mtp_drafted,
8106 "constant logits → every draft accepted"
8107 );
8108 assert!(r.token_ids.windows(2).all(|w| w[0] == w[1]));
8111 }
8112
8113 #[test]
8114 fn empty_prompt_is_an_error_not_a_panic() {
8115 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
8116 let r = p.generate("", 4, None, None);
8117 assert!(r.is_err(), "empty prompt must be a clean error");
8118 }
8119
8120 #[test]
8121 fn every_token_enters_kv_exactly_once() {
8122 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8123 p.sampler_config.temperature = 0.0;
8125 let r = p.generate("abc", 2, None, None).unwrap();
8126 assert_eq!(r.prompt_tokens, 3);
8127 assert_eq!(
8131 p.kv_cache.seq_len(),
8132 3 + r.tokens_generated - 1,
8133 "each token must be cached exactly once (v1 cached the last prompt token twice)"
8134 );
8135 }
8136
8137 #[test]
8138 fn generation_is_reproducible_with_seed() {
8139 let run = || {
8140 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8141 p.generate("hello", 8, None, None).unwrap().token_ids
8142 };
8143 assert_eq!(run(), run());
8144 }
8145
8146 #[test]
8147 fn resetting_sampler_restarts_the_seeded_stream() {
8148 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 2, 260);
8149 let config = SamplerConfig {
8150 seed: Some(1234),
8151 ..SamplerConfig::default()
8152 };
8153 p.set_sampler_config(config.clone());
8154 let first = p.generate("hello", 8, None, None).unwrap().token_ids;
8155 p.set_sampler_config(config);
8156 let second = p.generate("hello", 8, None, None).unwrap().token_ids;
8157 assert_eq!(first, second);
8158 }
8159
8160 #[test]
8161 fn eviction_bounds_the_cache() {
8162 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 260);
8163 p.kv_cache.max_seq_len = 6;
8164 p.sampler_config.temperature = 0.0;
8165 let _ = p.generate("abcd", 12, None, None).unwrap();
8166 assert!(
8167 p.kv_cache.seq_len() <= 6 + 1,
8168 "cache must stay bounded by max_seq_len (got {})",
8169 p.kv_cache.seq_len()
8170 );
8171 }
8172
8173 #[test]
8174 fn confidence_matches_tokens_and_is_a_probability() {
8175 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8176 p.sampler_config.temperature = 0.0;
8177 p.sampler_config.repetition_penalty = 1.0;
8178 let r = p.generate("abcd", 10, None, None).unwrap();
8179 assert_eq!(
8180 r.token_confidence.len(),
8181 r.token_ids.len(),
8182 "one confidence per emitted token"
8183 );
8184 for &c in &r.token_confidence {
8185 assert!((0.0..=1.0).contains(&c), "confidence out of [0,1]: {c}");
8186 }
8187 let logits = [1.0f32, 3.0, 0.5, 3.0];
8189 let p0 = top1_prob_t(&logits, 1, 1.0);
8190 let p1 = top1_prob_t(&logits, 3, 1.0);
8191 assert!((p0 - p1).abs() < 1e-6, "equal logits → equal prob");
8192 assert!(p0 > 0.0 && p0 < 1.0);
8193 let sharp = top1_prob_t(&logits, 1, 1.0);
8195 let soft = top1_prob_t(&logits, 1, 2.0);
8196 assert!(soft < sharp, "higher temperature lowers peak confidence");
8197 }
8198
8199 #[test]
8200 fn trace_is_opt_in_and_parallels_the_output() {
8201 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8203 p.sampler_config.temperature = 0.0;
8204 p.sampler_config.repetition_penalty = 1.0;
8205 let r = p.generate("abcd", 10, None, None).unwrap();
8206 assert!(r.traces.is_empty(), "trace must be empty unless enabled");
8207
8208 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8210 p.sampler_config.temperature = 0.0;
8211 p.sampler_config.repetition_penalty = 1.0;
8212 p.set_trace(true);
8213 let r = p.generate("abcd", 10, None, None).unwrap();
8214 assert_eq!(r.traces.len(), r.token_ids.len(), "one trace row per token");
8215 for (i, tr) in r.traces.iter().enumerate() {
8216 assert_eq!(tr.t, i, "trace index is sequential");
8217 assert_eq!(tr.token_id, r.token_ids[i], "trace token_id matches output");
8218 assert_eq!(
8219 tr.confidence, r.token_confidence[i],
8220 "trace confidence matches the confidence channel"
8221 );
8222 assert!(tr.active_skill.is_none() && tr.recon.is_none() && !tr.switched);
8224 }
8225 }
8226
8227 #[test]
8228 fn explain_prefill_logits_match_greedy_first_token() {
8229 let mut p = create_test_pipeline(8, 16, 2, 1, 4, 1, 64);
8233 p.sampler_config.temperature = 0.0;
8234 p.sampler_config.repetition_penalty = 1.0;
8235 let ids = p.tokenizer.encode("abcd");
8236 let logits = p.prefill_next_logits(&ids, None);
8237 let argmax = logits
8238 .iter()
8239 .enumerate()
8240 .max_by(|a, b| a.1.partial_cmp(b.1).unwrap())
8241 .unwrap()
8242 .0 as u32;
8243 let r = p.generate("abcd", 1, None, None).unwrap();
8244 assert_eq!(
8245 argmax, r.token_ids[0],
8246 "explain preview must match greedy emit"
8247 );
8248 }
8249
8250 #[test]
8251 fn laguna_shared_expert_is_unconditionally_added() {
8252 let matrix = |values: Vec<f32>| QTensor::from_f32(values, 2, 2);
8253 let identity = || matrix(vec![1.0, 0.0, 0.0, 1.0]);
8254 let zero_dense = || DenseFfn {
8255 gate_proj: matrix(vec![0.0; 4]),
8256 up_proj: matrix(vec![0.0; 4]),
8257 down_proj: matrix(vec![0.0; 4]),
8258 act: Act::Silu,
8259 };
8260 let shared = DenseFfn {
8261 gate_proj: identity(),
8262 up_proj: identity(),
8263 down_proj: identity(),
8264 act: Act::Silu,
8265 };
8266 let x = [1.0, 2.0];
8267 let expected = dense_ffn(&shared, &x, None);
8268 let moe = MoeFfn {
8269 router: QTensor::from_f32(vec![0.0, 0.0], 1, 2),
8270 experts: vec![zero_dense()],
8271 top_k: 1,
8272 norm_topk_prob: true,
8273 router_sigmoid: true,
8274 expert_bias: None,
8275 routed_scaling: 1.0,
8276 route_tau: None,
8277 shared: Some((shared, None)),
8278 stats: std::cell::RefCell::new(Vec::new()),
8279 act_sq: std::cell::RefCell::new(Vec::new()),
8280 act_rows: std::cell::RefCell::new(Vec::new()),
8281 mask: None,
8282 per_expert_scale: None,
8283 router_input_norm: false,
8284 };
8285 let actual = moe_ffn_cpu(&moe, &x, &[0], &[0.0], 1.0, None);
8286 for (actual, expected) in actual.iter().zip(expected) {
8287 assert!((actual - expected).abs() < 1e-6);
8288 }
8289 }
8290}