+ TV-j^RIHtHt^RIHt^RIHt]PR4t Rt ] !4t RRlt R Rlt R RR lltRR R lltR#))OptionalTupleNcVP\P4p\P!W,4p\P !W^,V^,4#)r)astypemxfloat32nnsoftplusclip)dtdt_biastime_step_limits&&&b/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_lm/models/ssm.py compute_dtrs? 2:: B R\ "B 772q)?1+= >>c\PP4'gR#Rp\PP R.RORR.VR7#)Na auto n = thread_position_in_grid.z; auto h_idx = n % H; auto g_idx = n / G; constexpr int n_per_t = Ds / 32; auto x = X + n * Dh; out += n * Dh; auto i_state = state_in + n * Dh * Ds; auto o_state = state_out + n * Dh * Ds; // C and B have shape [batch, group, state_dim] // C and B need to be offset by group size auto C_ = C + g_idx * Ds; auto B_ = B + g_idx * Ds; auto ds_idx = thread_position_in_threadgroup.x; auto d_idx = thread_position_in_grid.y; auto dt_ = static_cast(dt[n]); auto A = -fast::exp(static_cast(A_log[h_idx])); auto dA = fast::exp(A * dt_); float acc = 0.0; auto x_ = static_cast(x[d_idx]); for (int i = 0; i < n_per_t; ++i) { auto s_idx = n_per_t * ds_idx + i; auto idx = d_idx * Ds + s_idx; auto dB_by_x = x_ * dt_ * static_cast(B_[s_idx]); auto state = dA * i_state[idx] + dB_by_x; o_state[idx] = static_cast(state); acc += state * C_[s_idx]; } acc = simd_sum(acc); if (thread_index_in_simdgroup == 0) { out[d_idx] = static_cast(acc + x_ * D[h_idx]); } ssm_kernelout state_out)name input_names output_namessource)XA_logBCDr state_in)rmetal is_availablefast metal_kernel)rs rmake_ssm_kernelr$sO 88 " "&FN 77   C[)  rcFV^8dQhR\PR\PR\PR\PR\PR\PR\PR\PR \\\3,/ #)  hidden_statesrrrrr r stater)rarrayrfloat)formats"r __annotate__r,Cs88 88 xx xx   xx   XX 885%<(rc VPwrrVPp VPpVPR Rwpp\WVV4p\WW#WEV.RV 3RV3RV 3RV3RV 3RW,3.^ WV ,3R V ^W3VP.W.R7#) r&NTUDhDsHG)inputstemplategrid threadgroup output_shapes output_dtypes) )shapedtyper _ssm_kernel)r'rrrrr r r(rn_hd input_type state_typehbdss&&&&&&&&& rssm_update_kernelrICs$$JA!$$JJ WWRS\FB B 1B aA59 *  *  1I 2J !H !'N  !U^1a|U[[1!. rcVPR,pVe \P!V^4pW,p\P!VR,VRR7p\P!VR4p\P !VRR7pVe8\P !VR,VR,,V\R4)4pV#)r=axisinf.Nr:).NNNN)r>r expand_dimsrepeattrilcumsumwherer*)xmasklx_segsums&& rsegsumrZds  A ~~dA& H !I,+A 2Ayy$H 88  i 0(U5\M  Orc*V^8dQhR\PR\PR\PR\PR\PR\PR\PR\\P,R \\\3,R \\P,R \\P,R \ R \\P\P3,/ #)r&rVrrrrr r r(rrWlengthsstepreturn)rr)rrr*int)r+s"rr,r,sscc xxc 88c xxc xx c  xx c  cXXc BHH c5%<(c 288 cbhh c c 288RXX crc &aa a aaaaaaSPwop ooVPwp oo\WVV4pSS,o\P!V4P VP 4)pW^P ^^R4,pVP SV S^4S,pVVVVVV VV V3 Rlp.p\^V S 4FpT!VRVVS ,13,VRVVS ,13,VRVVS ,13,VRVVS ,13,TV fRMV RVVS ,13,4wppS e S S , o VPV4K \P!V^R7SVP ^^S^4,,pVV3#)aSSD-SSM forward pass. Args: x: Input of shape (batch_size, seq_len, num_heads, head_dim). dt: Time deltas of shape (seq_len, num_heads,). A_log: State transition of shape (num_heads,). B: Input mixing of shape (batch_size, seq_len, num_groups, n). C: Output mixing of shape (batch_size, seq_len, num_groups, n). D: Residual connection. dt_bias: Bias for time deltas of shape (num_heads,). time_step_limit: Minimum and maximum value for time deltas. mask: Optional multiplicative mask. lengths: Optional lenghts of sequences, assumed to be the full length if unspecified. step: Step size for processing x. Code modified from https://github.com/cartesia-ai/edge/blob/main/cartesia-mlx/cartesia_mlx/layers/ssd/ops.py c< VP^,p\P!VR4p\P!V^^4V,p\P!VS^R7p\P !\ VP^^4VR74p\P!Wx,^4p WP^^4,p \P!V ^^4p Sed\P!\P!SS4^, ^4p \P!V R4p \P!W^R7pMVRRRR1R3,pVP^^^^4p\P!VSS,^R7P^^4pW,p V P^^4P^^4p W,p Ve\P !\P!VRR74pWR ,V,, p VPSVS^S^4pVPS^SSSS34V,PR4P^^4pWR ,V,, p Se4Ve0\P !\P!S^8R4WM4p V P#SP$4V 3#) r=rK)rWNrP)rr&r=)r=r&rbrNr:)rPrNrPNNrO)r>r transposeswapaxesrRexprZrSmaximumminimumrQtake_along_axisrTreshapesqueezeflattenrUrr?)dtxdtArrr(rWsCBdecaysurrogate_attention_matrixyposdtxdecay next_stateexp_dtA_cumsumy_prevbrDdhgrCr\repeatsr]rVs&&&&&& r_stepssm_attn.._stepsC IIaL LLL ) [[Aq !A % YYr7 +vcll1a0t<=%'WWRZ%;" &a); ; KK1a   **RZZ6:A>C..i0C&&u:E!QQ,'E1a+ IIaaa ( 1 1!Q 7;$$Q*33Aq9\  VVBIIc$;F FJ !Q1a+A1a"a89A=FFrJRRSTVWX   *V3 3A  5#4w{I6Jxx *,,rrPN.rKrN) r>rrrerr?rirangeappend concatenate)rVrrrrr r r(rrWr\r]rXrBArmrlr|ysirrrxrDryrzrCr{sf&&&&&&&&&ff @@@@@@rssm_attnrss|B''KAq!RJAq!Q B 1B1fG   bhh ''A yyAr" "C **Q1a 1 $C)-)-V B 1a  1q4x< 1q4x< aQXo  aQXo  LDd3AH +<&=  5  nG !  r"Q1aA)>%>>A e8OrcV^8dQhR\PR\PR\PR\PR\PR\PR\PR\\P,R \\\3,R \\P,R \\P,/ #) r&r'rrrrr r r(rrWr\)rr)rrr*)r+s"rr,r,s, , 88, 88,  xx,  xx ,  xx ,  , XX,  BHH , 5%<(,  288 , bhh , rc (VP^,p V ^8gQVeM\P!4\P8wg%\PP 4'g\ VVVVVVVVVV V R7 #\VVVVVVVVV4 #)r=)rWr\)r>rdefault_devicegpur r!rrI) r'rrrrr r r(rrWr\seq_lens &&&&&&&&&&& r ssm_updaters!!!$G! =   "&& (xx$$&&        !        r)N)NgMbP?gY@NN)NrNN)typingrrmlx.corecorermlx.nnr compilerr$r@rIrZrrrrrsR"? ? /d B cL, , r