Ë
    îÿæil  ã                   óè  — d dl mZ d dlmZmZ d dlmZ d dlm	Z	  eej                  d¬«      d„ «       Zdd„Z edd¬	«      Z edd¬	«      Z edd¬	«      Z edd¬	«      Zej                  	 dd
ej"                  dej"                  dej"                  dej"                  dej"                  dej"                  deej"                     deej"                  ej"                  f   fd„«       Z	 dd
ej"                  dej"                  dej"                  dej"                  dej"                  dej"                  deej"                     deej"                  ej"                  f   fd„Z	 	 dd
ej"                  dej"                  dej"                  dej"                  dej"                  deej"                     deej"                     deej"                  ej"                  f   fd„Z	 	 	 dd
ej"                  dej"                  dej"                  dej"                  dej"                  dej"                  dej"                  deej"                     deej"                     dedeej"                  ej"                  f   fd„Zy)é    )Úpartial)ÚOptionalÚTupleNT)Ú	shapelessc                 óò   — t        j                  t        j                  | j                  t         j                  «      «       t	        j
                  ||z   «      z  «      j                  |j                  «      S ©N)ÚmxÚexpÚastypeÚfloat32ÚnnÚsoftplusÚdtype)ÚA_logÚaÚdt_biass      ún/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/mlx_lm/models/gated_delta.pyÚ	compute_gr      sP   € ä�6‰6”2—6‘6˜%Ÿ,™,¤r§z¡zÓ2Ó3Ð3´b·k±kÀ!ÀgÁ+Ó6NÑNÓO×VÑVØ	�‰óð ó    Fc                 ó4  — t         j                  j                  «       sy | rdnd}|r	d}d}d}d}nd}d}d	}d
}d|› d|› d|› d|› d|› d�}g d¢}| r|j                  d«       d}	|r|	dz  }	| r|	dz  }	t         j                  j                  d|	› �|ddg|¬«      S )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);
            }
          }
          // 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<InT>(state[i]);
        }
    )ÚqÚkÚvÚgÚbetaÚstate_inÚTÚmaskÚ Ú_vecÚ_maskÚgated_delta_stepÚyÚ	state_out)ÚnameÚinput_namesÚoutput_namesÚsource)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_kernelr8      sþ   € Ü�8‰8× Ñ Ô"ØÙ+3Ñ'¸€Kñ Ø*ˆ	ØAˆØˆØ$‰	à&ˆ	Ø1ˆØˆØˆ	ð	ð8 
ˆð 	Ø	ˆ	ð ð ˆmð %ð &. Jð /ð. ˆ+ð ðyC€FòH ;€FÙØ�‰�fÔà€FÙØ�&ÑˆÙØ�'Ñˆä�7‰7×ÑØ ˜xÐ(ØØ˜;Ð'Øð	  ó ð r   )r/   r0   r   r   r   r   r   Ústater   Úreturnc                 ó¶  — |}|j                   dk(  r|d   }n2|j                   dk(  r|dddd…f   }nt        d|j                  › �«      ‚||z  }||dddd…f   z  j                  d¬«      }	||	z
  |d	   z  }
||dddd…f   |
d	   z  z   }|| dddd…f   z  j                  d¬«      }|�.t	        j
                  |d
¬«      }t	        j                  |||«      }||fS )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é   .NzUnsupported gating shape éÿÿÿÿ©Úaxis).N)é   r<   r=   )ÚndimÚ
ValueErrorÚshapeÚsumr	   Úexpand_dimsÚwhere)r   r   r   r   r   r9   r   Ú	old_stateÚdecayÚkv_memÚdeltar$   s               r   Ú_gated_delta_step_opsrL   ~   sý   € ð2 €IØ‡v�v�‚{Ø�/Ñ"‰Ø	
�‰�1ŠØ�#�tšQ�,‘‰äÐ4°Q·W±W°IÐ>Ó?Ð?Ø�E‰M€EØ�a˜˜T¢1˜‘oÑ%×*Ñ*°Ð*Ó3€FØ�‰Z˜4 	™?Ñ*€EØ�A�c˜4¢�l‘O e¨IÑ&6Ñ6Ñ6€Eà	��3˜ša�<‘Ñ	 ×%Ñ%¨2Ð%Ó.€AàÐÜ�~‰~˜d¨Ô3ˆÜ—‘˜˜u iÓ0ˆØˆeˆ8€Or   c           
      óŠ  — |j                   \  }}}	}
|j                   dd  \  }}| j                  }|j                  dk(  r)t        }| ||||||g}|�@t        }|j                  |«       n(t        }| ||||||g}|�t        }|j                  |«        ||d|fd|
fd|fd|	fd|fgd|||z  fd	||||f|j                   g||g¬
«      S )Nr<   é   ÚInTÚDkÚDvÚHkÚHvé    )rT   rN   rA   )r6   ÚtemplateÚgridÚthreadgroupÚoutput_shapesÚoutput_dtypes)rD   r   rB   Ú_gated_delta_kernel_vecÚ_gated_delta_kernel_vec_maskedr,   Ú_gated_delta_kernelÚ_gated_delta_kernel_masked)r   r   r   r   r   r9   r   ÚBr   rR   rP   rS   rQ   Ú
input_typeÚkernelr6   s                   r   Úgated_delta_kernelra   «   s  € ð —7‘7�L€A€qˆ"ˆbØ�W‰W�Q�Rˆ[�F€BˆØ—‘€JØ‡v�v�‚{Ü(ˆØ�Q˜˜1˜d E¨1Ð-ˆØÐÜ3ˆFØ�M‰M˜$Õä$ˆØ�Q˜˜1˜d E¨1Ð-ˆØÐÜ/ˆFØ�M‰M˜$ÔáØà�JÐØ�2ˆJØ�2ˆJØ�2ˆJØ�2ˆJð
ð �"�a˜"‘fÐØØ˜1˜b "�~ u§{¡{Ð3Ø! :Ð.ôð r   c                 ó  — | j                   \  }}}	}
|j                   dd \  }}|€%t        j                  ||||
f| j                  ¬«      }||	z  x}dkD  r.t        j                  | |d«      } t        j                  ||d«      }g }t        |«      D ]U  }t        | dd…|f   |dd…|f   |dd…|f   |dd…|f   |dd…|f   ||€dn|dd…|f   «      \  }}|j                  |«       ŒW t        j                  |d¬«      }||fS )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]
    éþÿÿÿN©r   rA   r?   )	rD   r	   Úzerosr   ÚrepeatÚrangerL   r,   Ústack)r   r   r   r   r   r9   r   r^   r   rR   rP   rS   rQ   Úrepeat_factorÚysÚtr$   s                    r   Úgated_delta_opsrl   Ô   s  € ð. —7‘7�L€A€qˆ"ˆbØ�W‰W�R�Sˆ\�F€BˆØ€}Ü—‘˜!˜R  R˜°·±Ô8ˆà˜r™Ð!ˆ 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   r   Úbr   r   Ú
use_kernelc
           	      ó¨  — t        j                  |«      }
t        |||«      }|€H| j                  \  }}}}|j                  dd  \  }}t        j                  ||||f| j
                  ¬«      }|	rCt        j                  «       t         j                  k7  st         j                  j                  «       st        | ||||
||«      S t        | ||||
||«      S )Nrc   rd   )r	   Úsigmoidr   rD   re   r   Údefault_deviceÚgpur*   r+   rl   ra   )r   r   r   r   rm   r   r   r9   r   rn   r   r   r^   Ú_rR   rP   rS   rQ   s                     r   Úgated_delta_updatert     s¸   € ô �:‰:�a‹=€DÜ�%˜˜GÓ$€AØ€}Ø—w‘w‰ˆˆ1ˆb�"Ø—‘˜˜�‰ˆˆBÜ—‘˜!˜R  R˜°·±Ô8ˆáœ×*Ñ*Ó,´·±Ò6¼b¿h¹h×>SÑ>SÔ>UÜ˜q ! Q¨¨4°¸Ó=Ð=Ü˜a  A q¨$°°tÓ<Ð<r   )FFr   )NN)NNT)Ú	functoolsr   Útypingr   r   Úmlx.coreÚcorer	   Úmlx.nnr   Úcompiler   r8   r\   r]   rZ   r[   ÚarrayrL   ra   rl   Úboolrt   © r   r   Ú<module>r~      sä  ðÝ ß "å Ý ñ 	ˆ�‰˜tÔ$ñó %ðódñN /¸È%ÔPÐ Ù5¸tÐPUÔVÐ Ù2¸EÈdÔSÐ Ù!9Ø˜dô"Ð ð
 ‡�ð  $ñ)Ø	‡x�xð)à	‡x�xð)ð 
‡x�xð)ð 
‡x�xð	)ð
 �(‰(ð)ð �8‰8ð)ð �2—8‘8Ñ
ð)ð ˆ2�8‰8�R—X‘XÐÑò)ó ð)ðf  $ñ&Ø	‡x�xð&à	‡x�xð&ð 
‡x�xð&ð 
‡x�xð	&ð
 �(‰(ð&ð �8‰8ð&ð �2—8‘8Ñ
ð&ð ˆ2�8‰8�R—X‘XÐÑó&ð^ !%Ø#ñ-Ø	‡x�xð-à	‡x�xð-ð 
‡x�xð-ð 
‡x�xð	-ð
 �(‰(ð-ð �B—H‘HÑð-ð �2—8‘8Ñ
ð-ð ˆ2�8‰8�R—X‘XÐÑó-ðp !%Ø#Øñ=Ø	‡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   