iddlmZmZddlmZddlmZejdZ 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))OptionalTupleNcltj||z}tj||d|dS)Nr)nnsoftplusmxclip)dtdt_biastime_step_limits f/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/mlx_lm/models/ssm.py compute_dtrs1 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)r metal is_availablefast metal_kernel)rs rmake_ssm_kernelr# sL 88 "&FN 77   C[)  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 WWRS\FB B 1B uaAr59 #dAYr S!HsAQSGnU!QU^1a|U[[1!:.  rcP|jd}|tj|d}||z}tj|d|d}tj|d}tj |d}|/tj |ddddf|dz|td }|S)Nr.Naxisr'.inf)r5r expand_dimsrepeattrilcumsumwherefloat)xmasklx_segsums rsegsumrP[s  A ~~dA& H !I,+A 2Ayy$H 88 dA i 0(U5\M  OrrLrMlengthsstepreturnc  |j\} |j\} } t|||}ztj|j |j  }||j dddz}|j | d|z}  fd}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 rrAc|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|| } | | fS) Nr)rrrVrC)rMr)rrVrWrArWr'rB)r5r transposeswapaxesrGexprPrHmaximumminimumrFtake_along_axisrIreshapesqueezeflattenrJ)dtxdtArrr%rMsCBdecaysurrogate_attention_matrixyposdtxdecay next_stateexp_dtA_cumsumy_prevbr;dhgr:rQrepeatsrRs r_stepzssm_attn.._stepsO 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{I6zJ*}rrN.rC) r5rr rZastyper6r^rangeappend concatenate)rLrrrrr r r%r rMrQrRrNr9Arbrarqysirgrmr;rnror:rps `` @@@@@@rssm_attnryjsB''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)rMrQ)r5r default_devicegpurr ryr?) r$rrrrr r r%r rMrQseq_lens r ssm_updater~s!!!$G! =   "&& (xx$$&        !        r)N)NgMbP?gY@NN)NrNN)typingrrmlx.corecorer mlx.nnrcompilerr#r7arrayrKr?rPintryr~rrrsq"? ? /d 88 88 xx xx   xx   XX 885%<(2 .!%+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