)jddlmZmZddlmZddlmZejdZ dZ e Z dej dej dej dej d ej d ej d ej d ej d ee e ffdZddZ ddej dej dej dej d ej d ej d ej d eej d ee e fdeej deej dedeej ej ffdZ ddej dej dej dej d ej d ej d ej d eej d ee e fdeej deej fdZdS))OptionalTupleNcttj||z}tj||d|dS)Nr)nnsoftplusmxclip)dtdt_biastime_step_limits [/lsinfo/ai/hellotax_ai/base_platform/venv/lib/python3.11/site-packages/mlx_lm/models/ssm.py compute_dtrs3 R'\ " "B 72q)?1+= > >>ctjsdSd}tjdgdddg|S)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_kernel)XA_logBCDr state_inout state_out)name input_names output_namessource)r metal is_availablefast metal_kernel)rs rmake_ssm_kernelr# s] 8 " "t&FN 7   CCC[)   r hidden_statesrrrrr r stater c |j\} } } } |j} |jdd\}}t|||}t|||||||gd| fd| fd|fd| fd| |zfgd| | | zfd| d | | f|jg| | g S) NTDhDsHG )r-rr)inputstemplategrid threadgroup output_shapes output_dtypes)shapedtyper _ssm_kernel)r$rrrrr r r%r n_hd input_typehbdss rssm_update_kernelr?Bs$JAq!Q$J WRSS\FB B 1 1B uaAr59 #dAYr S!HsAQSGnU!QU^1a|U[1!:.    rc\|jd}|tj|d}||z}tj|d|d}tj|d}tj|d}|8tj|ddddf|dz|td }|S)Nr.Naxisr'.inf)r5r expand_dimsrepeattrilcumsumwherefloat)xmasklx_segsums rsegsumrP[s  A ~dA&& H !I,+++A 2Ay$$$H 8 dAAA i 0(U5\\M   OrgMbP?gY@rLrMlengthsstepreturnc  |j\} |j\} } t|||}ztj||j }||dddz}|| d|z}  fd}g}td| D]~}||dd|| zf|dd|| zf|dd|| zf|dd|| zf|| dn| d|| zf\}}  z ||tj |d||dddzz}||fS)a5SSD-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 rrAc:|jd}tj|d}tj|dd|z}tj|d}tjt |dd|}tj||zd} | |ddz} tj| dd} Xtjtj dz d} tj | d} tj || d}n|ddddddddf}|dd dd}tj|zddd }||z} | dddd } | |z} |tjtj |d }| |dddddddf|zz } | |dd}| df|zddd }| |d |zz } /|-tjtj dkd|| } | | fS) Nr)rrrXrC)rMr)rrXrYrArYr'rB)r5r transposeswapaxesrGexprPrHmaximumminimumrFtake_along_axisrIreshapesqueezeflattenrJ)dtxdtArrr%rMsCBdecaysurrogate_attention_matrixyposdtxdecay next_stateexp_dtA_cumsumy_prevbr;dhgr:rSrepeatsrTs r_stepzssm_attn.._steps IaL LL ) ) [Aq ! !A % Yr7 + + +vcll1a00t<<<==%'WR%Z%;%;" &a);); ; K1a   *RZ66:A>>C.i00C&uc:::EE!!!QQQQQQ,'E1a++ Iaaa ( ( ( 1 1!Q 7 7;$$Q**33Aq99\  VBIc$;$;$;<F FJ !Q1a++A1a"a899A=FFrJJRRSTVWXX   *V3 3A  5#4w{I66zJ*}rrN.rC) r5rr r\astyper6r`rangeappend concatenate)rLrrrrr r r%r rMrSrTrNr9Ardrcrsysiriror;rprqr:rrs `` @@@@@@rssm_attnr{jsB'KAq!RJAq!Q B 1 1B1fG   bh ' ''A qyyAr"" "C **Q1a 1 $C))))))))))))V B 1a    5 1q4x< 1q4x< aaaQXo  aaaQXo  LDDd3AH +<&=   5  nG !  r"""Q1aA)>)>%>>A e8Orc  |jd} | dksA|?tjtjkstjst |||||||||| |  St||||||||| S)Nr)rMrS)r5r default_devicegpurr r{r?) r$rrrrr r r%r rMrSseq_lens r ssm_updaters!!$G! =   "& ( (x$$&& )          !         r)N)NrQNNrR)NrQNN)typingrrmlx.corecorer mlx.nnrcompilerr#r7arrayrKr?rPintr{rrrrs""""""""?? ? ///do 8 8 x x   x   X 85%<(2    .!%+9#"&cc xc 8c xc x c  x c  cXc BH c5%<(c 28 cbh c c 28RX cccc\!%+9#"&, , 8, 8,  x,  x ,  x ,  , X,  BH , 5%<(,  28 , bh , , , , , , r