(HJjddlmZddlmZmZddlmZddlm Z eejddZ ddZ e dd Z e dd Ze dd Ze dd Zej dd ej"d ej"d ej"d ej"dej"dej"deej"deej"ej"ffdZ dd ej"d ej"d ej"d ej"dej"dej"deej"deej"ej"ffdZ dd ej"d ej"d ej"d ej"dej"deej"deej"deej"ej"ffdZ dd ej"d ej"d ej"dej"dej"dej"dej"deej"deej"dedeej"ej"ffdZy))partial)OptionalTupleNT) shapelessctjtj|jtj t j ||zzSN)mxexpastypefloat32nnsoftplus)A_logadt_biass c/Users/ahmed/devFolder/claude-voice/.venv/lib/python3.12/site-packages/mlx_lm/models/gated_delta.py compute_grs< 66266%,,rzz233bkk!g+6NN OOFc 4tjjsy|rdnd}|r d}d}d}d}nd}d}d }d }d |d |d |d|d|d }gd}|r|jdd} |r| dz } |r| dz } tjj d| |ddg|S)Nzmask[b_idx * T + t]truez// g: [B, T, Hv, Dk]z-auto g_ = g + (b_idx * T * Hv + hv_idx) * Dk;z g_[s_idx]zg_ += Hv * Dk;z// g: [B, T, Hv]zauto g_ = g + b_idx * T * Hv;z g_[hv_idx]z g_ += Hv;a auto n = thread_position_in_grid.z; auto b_idx = n / Hv; auto hv_idx = n % Hv; auto hk_idx = hv_idx / (Hv / Hk); constexpr int n_per_t = Dk / 32; // q, k: [B, T, Hk, Dk] auto q_ = q + b_idx * T * Hk * Dk + hk_idx * Dk; auto k_ = k + b_idx * T * Hk * Dk + hk_idx * Dk; // v, y: [B, T, Hv, Dv] auto v_ = v + b_idx * T * Hv * Dv + hv_idx * Dv; y += b_idx * T * Hv * Dv + hv_idx * Dv; auto dk_idx = thread_position_in_threadgroup.x; auto dv_idx = thread_position_in_grid.y; // state_in, state_out: [B, Hv, Dv, Dk] auto i_state = state_in + (n * Dv + dv_idx) * Dk; auto o_state = state_out + (n * Dv + dv_idx) * Dk; float state[n_per_t]; for (int i = 0; i < n_per_t; ++i) { auto s_idx = n_per_t * dk_idx + i; state[i] = static_cast(i_state[s_idx]); } z zb auto beta_ = beta + b_idx * T * Hv; for (int t = 0; t < T; ++t) { if (z) { float kv_mem = 0.0f; for (int i = 0; i < n_per_t; ++i) { auto s_idx = n_per_t * dk_idx + i; state[i] = state[i] * a; kv_mem += state[i] * k_[s_idx]; } kv_mem = simd_sum(kv_mem); auto delta = (v_[dv_idx] - kv_mem) * beta_[hv_idx]; float out = 0.0f; for (int i = 0; i < n_per_t; ++i) { auto s_idx = n_per_t * dk_idx + i; state[i] = state[i] + k_[s_idx] * delta; out += state[i] * q_[s_idx]; } out = simd_sum(out); if (thread_index_in_simdgroup == 0) { y[dv_idx] = static_cast(out); } } else { y[dv_idx] = static_cast(0); } // Increment data pointers to next time step q_ += Hk * Dk; k_ += Hk * Dk; v_ += Hv * Dv; y += Hv * Dv; z beta_ += Hv; } for (int i = 0; i < n_per_t; ++i) { auto s_idx = n_per_t * dk_idx + i; o_state[s_idx] = static_cast(state[i]); } )qkvgbetastate_inTmask_vec_maskgated_delta_stepy state_out)name input_names output_namessource)r metal is_availableappendfast metal_kernel) has_mask vectorized mask_source g_commentg_setupg_access g_advancer(inputssuffixs r_make_gated_delta_kernelr7 s 88 "+3'K* A$ & 1  8    m%&.J/ 2 +}EFL;F f F&' 77  x (;'  r)r.r/rrrrrstaterreturnc|}|jdk(r|d}n2|jdk(r |ddddf}ntd|j||z}||ddddfzjd} || z |d z} ||ddddf| d zz}||ddddfzjd} |.t j |d }t j |||}| j|j|fS) a Ops-based reference implementation for a single recurrent step. Shapes: - q, k: [B, H, Dk] - v: [B, H, Dv] - g: [B, H] or [B, H, Dk] - beta: [B, H] - state: [B, H, Dv, Dk] Returns: - y: [B, H, Dv] - new_state: [B, H, Dv, Dk] ).NN.NzUnsupported gating shape axis).N)r;r<) ndim ValueErrorshapesumr expand_dimswherer dtype) rrrrrr8r old_statedecaykv_memdeltar#s r_gated_delta_step_opsrL~s 2Ivv{/" 1#tQ,4QWWI>?? EMEaT1 o% * * * 3F Z4 ? *E Ac4lOeI&66 6E 3a< %%2%.A ~~d3ui0 88AGG e ##rc |j\}}} } |jdd\} } |j} |j}|jdk(r)t}|||||||g}|@t}|j |n(t }|||||||g}|t}|j |||d| fd|fd| fd| fd| fd| fgd | || zfd ||| | f|jg| |g S) Nr;InTStTDkDvHkHv )rUrNr@)r5templategrid threadgroup output_shapes output_dtypes)rCrGrA_gated_delta_kernel_vec_gated_delta_kernel_vec_maskedr+_gated_delta_kernel_gated_delta_kernel_masked)rrrrrr8rBrrSrQrTrR input_type state_typekernelr5s rgated_delta_kernelrcs77LAq"b WWQR[FBJJvv{(Q1dE1-  3F MM$ $Q1dE1-  /F MM$   J  J  2J 2J 2J 2J  "a"f 1b"~u{{3!:. rc |j\}}} } |jdd\} } |)tj|| | | ftj}| | zx} dkDr.tj|| d}tj|| d}g}t |D]U}t |dd|f|dd|f|dd|f|dd|f|dd|f||dn|dd|f\}}|j|Wtj|d}||fS)a Ops-based reference implementation for prompt prefill (sequential loop). Supports both scalar and vectorized gating. Shapes: - q, k: [B, T, Hk, Dk] - v: [B, T, Hv, Dv] - g: [B, T, Hv] (scalar) or [B, T, Hv, Dk] (vectorized) - beta: [B, T, Hv] - state: [B, Hv, Dv, Dk] Returns: - y: [B, T, Hv, Dv] - state: [B, Hv, Dv, Dk] NrGr@r>) rCr zerosr repeatrangerLr+stack)rrrrrr8rr_rrSrQrTrR repeat_factorystr#s rgated_delta_opsrns.77LAq"b WWRS\FB }!RR ;r! Q& IIa + IIa + B 1X( adG adG adG adG AJ LDd1a4j 5 !  !A e8Orrbrr use_kernelc tj|} t|||} |L|j\} } }}|jdd\}}tj| |||ftj }| rCtj tjk7stjjst|||| | ||St|||| | ||S)Nrerf) r sigmoidrrCrgr default_devicegpur)r*rnrc)rrrrrorrr8rrprrr__rSrQrTrRs rgated_delta_updatervs ::a=D%G$A }ww 1b"B!RR ; **,6bhh>S>S>Uq!Q4== aAq$t <rs" t$P%PfR/%P5tPUV2EdS!9 d"  $)$ xx)$ xx)$ xx)$ xx )$ (( )$ 88 )$ 288 )$ 288RXX )$ )$f $( xx( xx( xx( xx ( (( ( 88 ( 288 ( 288RXX (b!%#- xx- xx- xx- xx - (( - BHH  - 288 - 288RXX -p!%#= xx= xx= xx= xx =  xx = 88 =XX= BHH = 288 == 288RXX =r