+
    TV-jã  ã                   ó&  € ^ RI Ht ^ RIHtHt ^ RIHt ^ RIH	t	 ]! ]P                  RR7      R 4       tRR lt]! RRR	7      t]! RRR	7      t]! RRR	7      t]! RRR	7      t]P                  RR
 R ll4       tRR R lltRR R lltRR R lltR# )é    )Úpartial)ÚOptionalÚTupleNT)Ú	shapelessc                 óÔ   € \         P                  ! \         P                  ! V P                  \         P                  4      4      ) \        P
                  ! W,           4      ,          4      # ©N)ÚmxÚexpÚastypeÚfloat32ÚnnÚsoftplus)ÚA_logÚaÚdt_biass   &&&Új/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_lm/models/gated_delta.pyÚ	compute_gr      s:   € ä�6Š6”2—6’6˜%Ÿ,™,¤r§z¡zÓ2Ó3Ð3´b·k²kÀ!Å+Ó6NÕNÓOÐOó    Fc                 óŽ  € \         P                  P                  4       '       g   R # V '       d   RMRpV'       d
   RpRpRpRpMRpRpR	pR
pRV RV RV RV RV R2p. ROpV '       d   VP                  R4       Rp	V'       d
   V	R,          p	V '       d
   V	R,          p	\         P                  P                  RV	 2VRR.VR7      # )Nzmask[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]zg_ += Hv * Dk;z// g: [B, T, Hv]zauto 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<float>(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<InT>(out);
            }
          } else {
            y[dv_idx] = static_cast<InT>(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<StT>(state[i]);
        }
    ÚmaskÚ Ú_vecÚ_maskÚgated_delta_stepÚyÚ	state_out)ÚnameÚinput_namesÚoutput_namesÚsource)ÚqÚkÚvÚgÚbetaÚstate_inÚT)r	   ÚmetalÚis_availableÚappendÚfastÚmetal_kernel)
Úhas_maskÚ
vectorizedÚmask_sourceÚ	g_commentÚg_setupÚg_accessÚ	g_advancer!   ÚinputsÚsuffixs
   &&        r   Ú_make_gated_delta_kernelr7      s  € Ü�8‰8× Ñ ×"Ò"Ùß+3Ñ'¸€K÷ Ø*ˆ	ØAˆØˆØ$‰	à&ˆ	Ø1ˆØˆØˆ	ð	ð8 
ˆð 	Ø	ˆ	ð ð ˆmð %ð &. Jð /ð2 ˆ+ð ð}E€FòL ;€FßØ�‰�fÔà€FßØ�&ÕˆßØ�'Õˆä�7‰7×ÑØ ˜xÐ(ØØ˜;Ð'Øð	  ó ð r   )r.   r/   c                ód  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\        \         P                  ,          R\        \         P                  \         P                  3,          /# ©	é   r"   r#   r$   r%   r&   Ústater   Úreturn©r	   Úarrayr   r   )Úformats   "r   Ú__annotate__r@      s‹   € ÷ )$ñ )$Ü	‡x�xð)$ä	‡x�xð)$ô 
‡x�xð)$ô 
‡x�xð	)$ô
 �(‰(ð)$ô �8‰8ð)$ô ”2—8‘8Õ
ð)$ô Œ2�8‰8”R—X‘XÐÕñ)$r   c                ó4  € TpVP                   ^8X  d   VR,          pM3VP                   ^8X  d   VR,          pM\        RVP                   24      hWX,          pWQR,          ,          P                  RR7      p	W),
          VR,          ,          p
WQR,          V
R,          ,          ,           pWPR,          ,          P                  RR7      pVe0   \        P
                  ! VRR7      p\        P                  ! WeV4      pVP                  V P                  4      V3# )zñ
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]
zUnsupported gating shape ©Úaxis).NN).NºNNNéÿÿÿÿ).N)é   r:   é   )	ÚndimÚ
ValueErrorÚshapeÚsumr	   Úexpand_dimsÚwherer   Údtype)r"   r#   r$   r%   r&   r;   r   Ú	old_stateÚdecayÚkv_memÚdeltar   s   &&&&&&&     r   Ú_gated_delta_step_opsrS   ~   sè   € ð2 €IØ‡v�v�„{Ø�/Õ"‰Ø	
�‰�1ŒØ�,•‰äÐ4°Q·W±W°IÐ>Ó?Ð?Ø�M€EØ˜•oÕ%×*Ñ*°Ð*Ó3€FØ�Z˜4 	�?Õ*€EØ�l•O e¨IÕ&6Õ6Õ6€Eà	�<•Õ	 ×%Ñ%¨2Ð%Ó.€AàÒÜ�~Š~˜d¨Ô3ˆÜ—’˜ iÓ0ˆØ�8‰8�A—G‘GÓ˜eÐ#Ð#r   c                ód  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\        \         P                  ,          R\        \         P                  \         P                  3,          /# r9   r=   )r?   s   "r   r@   r@   «   s‹   € ÷ (ñ (Ü	‡x�xð(ä	‡x�xð(ô 
‡x�xð(ô 
‡x�xð	(ô
 �(‰(ð(ô �8‰8ð(ô ”2—8‘8Õ
ð(ô Œ2�8‰8”R—X‘XÐÕñ(r   c           
      óª  € VP                   w  rxršVP                   R ,          w  r¼V P                  pVP                  pVP                  ^8X  d)   \        pWW#WEV.pVe   \        pVP                  V4       M'\        pWW#WEV.pVe   \        pVP                  V4       V! VRV3RV3RV
3RV3RV	3RV3.^ WÇV,          3RWxW¼3VP                   .WÞ.R7      # )	:r:   NNÚInTÚStTÚDkÚDvÚHkÚHv)r5   ÚtemplateÚgridÚthreadgroupÚoutput_shapesÚoutput_dtypes)é    é   rF   )rJ   rN   rH   Ú_gated_delta_kernel_vecÚ_gated_delta_kernel_vec_maskedr+   Ú_gated_delta_kernelÚ_gated_delta_kernel_masked)r"   r#   r$   r%   r&   r;   r   ÚBr(   rZ   rX   r[   rY   Ú
input_typeÚ
state_typeÚkernelr5   s   &&&&&&&          r   Úgated_delta_kernelrk   «   sú   € ð —7‘7�L€Aˆ"Ø�W‰W�R�[�F€BØ—‘€JØ—‘€JØ‡v�v�„{Ü(ˆØ˜˜d¨1Ð-ˆØÒÜ3ˆFØ�M‰M˜$Ôøä$ˆØ˜˜d¨1Ð-ˆØÒÜ/ˆFØ�M‰M˜$ÔáØà�JÐØ�JÐØ�2ˆJØ�2ˆJØ�2ˆJØ�2ˆJð
ð �"˜"•fÐØØ˜b�~ u§{¡{Ð3Ø!Ð.ôð r   c                óz  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\        \         P                  ,          R\        \         P                  ,          R\        \         P                  \         P                  3,          /# r9   r=   )r?   s   "r   r@   r@   Ö   s�   € ÷ -ñ -Ü	‡x�xð-ä	‡x�xð-ô 
‡x�xð-ô 
‡x�xð	-ô
 �(‰(ð-ô ”B—H‘HÕð-ô ”2—8‘8Õ
ð-ô Œ2�8‰8”R—X‘XÐÕñ-r   c                ó6  € V P                   w  rxršVP                   RR w  r¼Vf)   \        P                  ! W{WÊ3\        P                  R7      pW¹,          ;p^8”  d/   \        P                  ! WR4      p \        P                  ! WR4      p. p\        V4       Fd  p\        V RV3,          VRV3,          VRV3,          VRV3,          VRV3,          TVf   RM
VRV3,          4      w  ppVP                  V4       Kf  	  \        P                  ! V^R7      pVV3# )aR  
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]
N©rN   rD   rB   éþÿÿÿ)	rJ   r	   Úzerosr   ÚrepeatÚrangerS   r+   Ústack)r"   r#   r$   r%   r&   r;   r   rg   r(   rZ   rX   r[   rY   Úrepeat_factorÚysÚtr   s   &&&&&&&          r   Úgated_delta_opsrw   Ö   s  € ð. —7‘7�L€Aˆ"Ø�W‰W�R�Sˆ\�F€BØ‚}Ü—’˜! ˜´·
±
Ô;ˆà�Ð!ˆ QÔ&Ü�IŠI�a¨Ó+ˆÜ�IŠI�a¨Ó+ˆà	€BÜ�1ŽXˆÜ(Øˆa�ˆd�GØˆa�ˆd�GØˆa�ˆd�GØˆa�ˆd�GØ��A��JØØ’L‰D d¨1¨a¨4¥jó
‰ˆˆ5ð 	�	‰	�!Žñ ô 	�Š�˜!Ô€AØˆeˆ8€Or   c                óÆ  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\         P                  R\        \         P                  ,          R	\        \         P                  ,          R
\        R\        \         P                  \         P                  3,          /# )r:   r"   r#   r$   r   Úbr   r   r;   r   Ú
use_kernelr<   )r	   r>   r   Úboolr   )r?   s   "r   r@   r@     s¶   € ÷ =ñ =Ü	‡x�xð=ä	‡x�xð=ô 
‡x�xð=ô 
‡x�xð	=ô
 
‡x�xð=ô �8‰8ð=ô �X‰Xð=ô ”B—H‘HÕð=ô ”2—8‘8Õ
ð=ô ð=ô Œ2�8‰8”R—X‘XÐÕñ=r   c
           	      óÄ  € \         P                  ! V4      p
\        WSV4      pVfL   V P                  w  rÍrïVP                  RR  w  pp\         P                  ! VVVV3\         P
                  R7      pV	'       dM   \         P                  ! 4       \         P                  8w  g%   \         P                  P                  4       '       g   \        WW+W§V4      # \        WW+W§V4      # )Nrn   ro   )r	   Úsigmoidr   rJ   rp   r   Údefault_deviceÚgpur)   r*   rw   rk   )r"   r#   r$   r   ry   r   r   r;   r   rz   r&   r%   rg   Ú_rZ   rX   r[   rY   s   &&&&&&&&&&        r   Úgated_delta_updater�     s©   € ô �:Š:�a‹=€DÜ�%˜GÓ$€AØ‚}Ø—w‘w‰ˆˆbØ—‘˜˜�‰ˆˆBÜ—’˜!˜R  R˜´·
±
Ô;ˆçœ×*Ò*Ó,´·±Ô6¼b¿h¹h×>SÑ>S×>UÒ>UÜ˜q Q¨4¸Ó=Ð=Ü˜a A¨$°tÓ<Ð<r   )FFr   )NN)NNT)Ú	functoolsr   Útypingr   r   Úmlx.coreÚcorer	   Úmlx.nnr   Úcompiler   r7   re   rf   rc   rd   rS   rk   rw   r�   © r   r   Ú<module>r‰      s¢   ðÝ ß "å Ý ñ 	ˆ�‰˜tÔ$ñPó %ðPôfñR /¸È%ÔPÐ Ù5¸tÐPUÔVÐ Ù2¸EÈdÔSÐ Ù!9Ø˜dô"Ð ð
 ‡�ö)$ó ð)$÷X(÷V-÷`=ñ =r   