(HJjddlmZmZddlmZddlmZejdZ dZ e Z dejdejdejdejd ejd ejd ejd ejd ee e ffdZddZ ddejdejdejdejd ejd ejd ejd eejd ee e fdeejdeejdedeejejffdZ ddejdejdejdejd ejd ejd ejd eejd ee e fdeejdeejfdZy))OptionalTupleNc|jtj}tj||z}tj ||d|dS)Nr)astypemxfloat32nnsoftplusclip)dtdt_biastime_step_limits [/Users/ahmed/devFolder/claude-voice/.venv/lib/python3.12/site-packages/mlx_lm/models/ssm.py compute_dtrsC 2:: B R'\ "B 772q)?1+= >>ctjjsyd}tjj dgdddg|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)rmetal is_availablefast metal_kernel)r s rmake_ssm_kernelr%sL 88 "&FN 77   C[)  r hidden_statesrrrrr rstaterc |j\} } } } |j} |j}|jdd\}}t|||}t|||||||gd| fd|fd| fd|fd| fd| |zfgd| | | zfd | d | | f|jg| |g S) NTUDhDsHG )r0rr)inputstemplategrid threadgroup output_shapes output_dtypes)shapedtyper _ssm_kernel)r&rrrrr rr'rn_hd input_type state_typehbdss rssm_update_kernelrCCs$$JAq!Q$$JJ WWRS\FB B 1B uaAr59 *  *  1I 2J !H !r'N  !QU^1a|U[[1!:. rcP|jd}|tj|d}||z}tj|d|d}tj|d}tj |d}|/tj |ddddf|dz|td }|S)Nr.Naxisr).inf)r8r expand_dimsrepeattrilcumsumwherefloat)xmasklx_segsums rsegsumrTds  A ~~dA& H !I,+A 2Ayy$H 88 dA i 0(U5\M  OrrPrQlengthsstepreturnc  j\} |j\} } t|||}ztj|j |j  }||j dddz}|j | dz}  f d}g}td| D]h}||dd|| zf|dd|| zf|dd|| zf|dd|| zf|| dn | d|| zf\}}  z |j|jtj|d|j 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 rrEc" |jd}tj|d}tj|dd|z}tj|d}tj t |jdd|}tj||zd} | |jddz} tj| dd} \tjtjdz d} tj| d} tj|| d}n|ddddddddf}|jdd dd}tj|zdjdd }||z} | jddjdd } | |z} |tj tj|d }| |dddddddf|zz } |j|dd}|jdf|zjdjdd }| |d |zz } 0|.tj tjdkd|| } | j#j$| fS) Nr)rrrZrG)rQr)rrZr[rEr[r)rF)r8r transposeswapaxesrKexprTrLmaximumminimumrJtake_along_axisrMreshapesqueezeflattenrNrr9)dtxdtArrr'rQsCBdecaysurrogate_attention_matrixyposdtxdecay next_stateexp_dtA_cumsumy_prevbr>dhgr=rUrepeatsrVrPs r_stepzssm_attn.._steps] IIaL LLL ) [[Aq !A % YYr7 +vcll1a0t<=%'WWR%Z%;" &a); ; KK1a   **RZZ6:A>C..i0C&&uc:E!QQ,'E1a+ IIaaa ( 1 1!Q 7;$$Q*33Aq9\  VVBIIc$;F FJ !Q1a+A1a"a89A=FFrJRRSTVWX   *V3 3A  5#4w{I6zJxx *,,rrN.rG) r8rrr^rr9rbrangeappend concatenate)rPrrrrr rr'rrQrUrVrRr<Arfreruysirkrqr>rrrsr=rts` `` @@@@@@rssm_attnr|ssB''KAq!RJAq!Q B 1B1fG   bhh ''A qyyAr" "C **Q1a 1 $C)-)-V B 1a  1q4x< 1q4x< aQXo  aQXo  LDd3AH +<&=  5  nG !  r"Q1aA)>%>>A e8Orc |jd} | dkDsE|Ctjtjk7stjj st |||||||||| |  St||||||||| S)Nr)rQrU)r8rdefault_devicegpur!r"r|rC) r&rrrrr rr'rrQrUseq_lens r ssm_updaters!!!$G! =   "&& (xx$$&        !        r)N)NgMbP?gY@NN)NrNN)typingrrmlx.corecorermlx.nnr compilerr%r:arrayrOrCrTintr|rrrrsr"? ? /d 88 88 xx xx   xx   XX 885%<(B .!%+9#"&c xxc 88c xxc xx c  xx c  cXXc BHH c5%<(c 288 cbhh c c 288RXX c\!%+9#"&, 88, 88,  xx,  xx ,  xx ,  , XX,  BHH , 5%<(,  288 , bhh , r