
    (HJj 7                        d dl mZ d dlmZ d dlmZ d dlmZ d dl	m
Z
 ddlmZ ddlmZ e G d d	e             Z eej                   d
      d        Z eej                   d
      d        Zej                   d        Zd Z e       Zdej,                  dej,                  dej,                  dej,                  dej,                  dej,                  dej,                  fdZ G d de
j0                        Z G d de
j0                        Z G d de
j0                        Z G d de
j0                        Z G d  d!e
j0                        Z G d" d#e
j0                        Z G d$ d%e
j0                        Z G d& d'e
j0                        Z y)(    )	dataclass)partial)OptionalN   )BaseModelArgs)ArraysCachec                       e Zd ZU eed<   eed<   eed<   eed<   eed<   eed<   eed<   eed<   eed	<   eed
<   eed<   dZeed<   y)	ModelArgs
model_type
vocab_sizehidden_sizeintermediate_sizenorm_epshead_dimnum_hidden_layersa_low_rank_dimv_low_rank_dimgate_low_rank_dimdecay_low_rank_dimFtie_word_embeddingsN)	__name__
__module____qualname__str__annotations__intfloatr   bool     ]/Users/ahmed/devFolder/claude-voice/.venv/lib/python3.12/site-packages/mlx_lm/models/rwkv7.pyr
   r
      sK    OOOM %%r    r
   T)	shapelessc                     | ||z  z   S Nr   )xyzs      r!   addcmulr(      s    q1u9r    c                 t    | t        j                  t         j                  j                  | dd      d      z  S )NTaxiskeepdimsgHz>)mxmaximumlinalgnorm)r%   s    r!   l2_normr2   #   s+    rzz"))..d.CTJJJr    c                     ||d   z  |dd d d f   z  }||d d d d d d d f   z  |d   |dd d d f   z  z   |z   }|| d   z  }||fS )N).N.r   )	rwkvabstatesabr&   s	            r!   _wkv7_step_opsr<   (   sj    1Y<1S$\?
2CAaD!m$$q|aT1o'EEKE)Ae8Or    c                      t         j                  j                         sy d} g d}t         j                  j	                  d|ddg|       S )Na  
        auto n = thread_position_in_grid.z;
        auto b_idx = n / H;
        auto h_idx = n % H;
        constexpr int n_per_t = D / 32;

        // [B, T, H, D]
        auto r_ = r + b_idx * T * H * D + h_idx * D;
        auto w_ = w + b_idx * T * H * D + h_idx * D;
        auto k_ = k + b_idx * T * H * D + h_idx * D;
        auto v_ = v + b_idx * T * H * D + h_idx * D;
        auto a_ = a + b_idx * T * H * D + h_idx * D;
        auto b_ = b + b_idx * T * H * D + h_idx * D;
        y += b_idx * T * H * D + h_idx * D;

        auto dk_idx = thread_position_in_threadgroup.x;
        auto dv_idx = thread_position_in_grid.y;

        // state_in, state_out: [B, H, D, D]
        auto i_state = state_in  + (n * D + dv_idx) * D;
        auto o_state = state_out + (n * D + dv_idx) * D;

        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]);
        }

        for (int t = 0; t < T; ++t) {
          float sa = 0.0f;
          for (int i = 0; i < n_per_t; ++i) {
            auto s_idx = n_per_t * dk_idx + i;
            sa += state[i] * a_[s_idx];
            state[i] = state[i] * w_[s_idx];
          }
          sa = simd_sum(sa);

          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] * v_[dv_idx] + sa * b_[s_idx];
            out += state[i] * r_[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
          r_ += H * D;
          w_ += H * D;
          k_ += H * D;
          v_ += H * D;
          a_ += H * D;
          b_ += H * D;
          y  += H * D;
        }
        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]);
        }
    )r4   r5   r6   r7   r8   r9   state_inTwkv7_kernelr&   	state_out)nameinput_namesoutput_namessource)r.   metalis_availablefastmetal_kernel)rE   inputss     r!   _make_wkv7_kernelrK   0   sQ    88  "=F| =F77;'	    r    r4   r5   r6   r7   r8   r9   r:   c           
          | j                   \  }}}	}
| j                  }t        | |||||||gd|fd|	fd|
fgd|
||	z  fd|||	|
f|j                   g||g      S )NInTHD    )rP      r   )rJ   templategridthreadgroupoutput_shapesoutput_dtypes)shapedtype_wkv7_kernel)r4   r5   r6   r7   r8   r9   r:   Br?   rN   rO   input_dtypes               r!   r@   r@   }   s     JAq!Q''K1aAq%+K !H!H

 !QU^1a|U[[1"K0 r    c                   $     e Zd Z fdZd Z xZS )LayerNormPerHeadc                     t         |           t        j                  ||f      | _        t        j                  ||f      | _        || _        y r$   )super__init__r.   zerosweightbiaseps)selfr   	num_headsrd   	__class__s       r!   r`   zLayerNormPerHead.__init__   s?    hh	845HHi23	r    c                     | j                   t        j                  j                  |d d | j                        z  | j
                  z   S r$   )rb   r.   rH   
layer_normrd   rc   re   r%   s     r!   __call__zLayerNormPerHead.__call__   s3    {{RWW//4txxHH499TTr    )r   r   r   r`   rk   __classcell__rg   s   @r!   r]   r]      s    Ur    r]   c                   f     e Zd Z	 	 d	dedededee   dee   f
 fdZdej                  fdZ
 xZS )
LoRA	input_dim
output_dimlow_rank_dimrc   
activationc                 N   t         |           || _        || _        || _        || _        |t        j                         | _        nl|dk(  rt        j                         | _        nM|dk(  rt        j                         | _        n.|dk(  rt        j                         | _        nt        d| d      t        j                  | j                  | j                  d      | j                  t        j                  | j                  | j                  | j
                        g| _        y )NsigmoidtanhreluzUnsupported activation type: .Frc   )r_   r`   rp   rq   rr   rc   nnIdentityrs   SigmoidTanhReLU
ValueErrorLinearlora)re   rp   rq   rr   rc   rs   rg   s         r!   r`   zLoRA.__init__   s     	"$(	 kkmDO9$ jjlDO6! ggiDO6! ggiDO<ZLJKK IIdnnd&7&7eDOOIId''tyyI
	r    returnc                 x     | j                   d    | j                   d    | j                   d   |                  S )N   r   r   )r   rj   s     r!   rk   zLoRA.__call__   s4    tyy|LDIIaL1a9::r    )Trv   )r   r   r   r   r   r   r   r`   r.   arrayrk   rl   rm   s   @r!   ro   ro      sX      $$*

 
 	

 tn
 SM
>;RXX ;r    ro   c                       e Zd Zd Zy)
TokenShiftc                     |j                   \  }}}|#t        j                  |d|f|j                        }|dk(  r|S t        j                  ||d d d dd d f   gd      S )Nr   r*   r,   )rW   r.   ra   rX   concatenate)re   r%   r:   rZ   LrO   s         r!   rk   zTokenShift.__call__   s`    ''1a=HHaAY0E6L>>5!AssAI,"7a@@r    N)r   r   r   rk   r   r    r!   r   r      s    Ar    r   c                   D     e Zd Zdef fdZdej                  fdZ xZS )Rwkv7ChannelMixingargsc                    t         |           |j                  }|j                  }t	        j
                  ||d      | _        t	        j
                  ||d      | _        t        j                  |      | _
        t               | _        y NFry   )r_   r`   r   r   rz   r   keyvaluer.   ra   x_kr   token_shift)re   r   
hidden_dimr   rg   s       r!   r`   zRwkv7ChannelMixing.__init__   si    %%
 2299Z):GYY0*5I
88Z)%<r    r   c                     ||d   nd }| j                  ||      }t        |||z
  | j                        }||d d dd d d f   |d<   | j                  t	        j
                  | j                  |                  S )Nr   r*   )r   r(   r   r   rz   relu2r   )re   r%   cacher:   x_prevxxs         r!   rk   zRwkv7ChannelMixing.__call__   su    !-a4!!!U+Q
DHH-BC|E!Hzz"((488B<011r    	r   r   r   r
   r`   r.   r   rk   rl   rm   s   @r!   r   r      s    (Y (2BHH 2r    r   c                   4     e Zd Zdedef fdZd Zd Z xZS )Rwkv7TimeMixingr   	layer_idxc                    t         |           || _        || _        |j                  | _        |j
                  | _        | j                  | j
                  z  | _        |j                  | _        |j                  | _        |j                  | _	        |j                  | _
        t               | _        t        j                  dd| j                  f      | _        t        j                  dd| j                  f      | _        t        j                  dd| j                  f      | _        t        j                  dd| j                  f      | _        t        j                  dd| j                  f      | _        t        j                  dd| j                  f      | _        t        j                  | j                  | j
                  f      | _        t        j                  | j                  | j
                  f      | _        t        j                  | j                  | j
                  f      | _        t1        j2                  | j                  | j                  d      | _        t1        j2                  | j                  | j                  d      | _        t1        j2                  | j                  | j                  d      | _        t1        j2                  | j                  | j                  d      | _        t=        | j
                  | j                  d      | _        tA        | j                  | j                  | j                  d      | _!        | j                  dkD  r2tA        | j                  | j                  | j                  d       | _"        tA        | j                  | j                  | j                  d       | _#        tA        | j                  | j                  | j                  d	d
      | _$        y )Nr   Fry   gh㈵D?rd   rv   )rr   rs   r   ru   )rr   rs   rc   )%r_   r`   r   r   r   r   rf   r   r   r   r   r   r   r.   ra   x_rx_wr   x_vx_ax_gk_kk_ar_krz   r   r_projk_projv_projo_projr]   g_normro   w_lorav_loraa_lorag_lorare   r   r   rg   s      r!   r`   zRwkv7TimeMixing.__init__   s   "	++))T]]:"11"11!%!7!7"&"9"9%<88Q4#3#34588Q4#3#34588Q4#3#34588Q4#3#34588Q4#3#34588Q4#3#34588T^^T]];<88T^^T]];<88T^^T]];<ii 0 0$2B2BOii 0 0$2B2BOii 0 0$2B2BOii 0 0$2B2BO&t}}dnn%P00	
 >>A    !00	DK ,,	
 // 
r    c                    |j                   \  }}	}
}
|Ct        j                  || j                  | j                  | j                  f|j
                        }t        j                         t        j                  k(  r/t        j                  j                         rt        |||||||      S g }t        |	      D ]Q  }t        |d d |f   |d d |f   |d d |f   |d d |f   |d d |f   |d d |f   |      \  }}|j                  |       S t        j                  |d      j                  |j
                        }||fS )N)rX   r   r   )rW   r.   ra   rf   r   rX   default_devicegpurF   rG   r@   ranger<   appendstackastype)re   r4   r5   r6   r7   r8   r9   r:   rZ   r   _ystr&   s                 r!   _wkv7zRwkv7TimeMixing._wkv7'  s   WW
1a=HHDNNDMM4==AE "&&(RXX-B-B-Dq!Q1a77B1X)adGQq!tWa1gqAw!Q$1a4%5 		!	  !$++AGG4Ae8Or    c           	      x   |d\  }}n
|d   |d   }}|j                   \  }}}| j                  ||      }	|	|z
  }
t        ||
| j                        }t        ||
| j                        }t        ||
| j
                        }t        ||
| j                        }t        ||
| j                        }t        ||
| j                        }| j                  |      j                  ||| j                  | j                        }| j                  |      j                  ||| j                  | j                        }| j                  |      j                  ||| j                  | j                        }t        j                   | j#                  |            j                  ||| j                  | j                        }| j%                  |      }| j&                  dk(  r|}nZt        j                   | j)                  |            j                  ||| j                  | j                        }t        |||z
  |      }t        j                   | j+                  |      j                  ||| j                  | j                              j-                  t        j.                        }t        j0                  d|z        j-                  |j2                        }t5        || j6                  z        }|d|dz
  | j8                  z  z   z  }| }||z  }| j;                  |||||||      \  }}| j=                  |j                  ||| j                  | j                              }|||z  | j>                  z  jA                  dd      |z  z   j                  |||g      }||d d dd d d f   |d<   ||d<   | jC                  ||z        }||fS )N)NNr   r   g]S hr*   Tr+   )"rW   r   r(   r   r   r   r   r   r   r   reshaperf   r   r   r   r.   ru   r   r   r   r   r   r   float32exprX   r2   r   r   r   r   r   sumr   )re   r%   v_firstr   token_shift_cachestate_cacherZ   r   rO   r   r   xrxwxkxvxaxgr   r   
receptanceiclrgatevvdecaykkr8   r9   outnew_state_caches                                r!   rk   zRwkv7TimeMixing.__call__;  s2   =-7*{-21XuQx{''1a!!!%67aZQDHH%QDHH%QDHH%QDHH%QDHH%QDHH%kk"o%%aDNNDMMJB''1dnndmmL[[_,,Q4>>4==Q
zz$++b/*221aW{{2>>QGDKKO,441dnndmmB E7U?B7E

KKO##Aq$..$--H

&
 	 y5()001A1ABcDHHn&Q$(dhh../CI#zzsE1a 
_ kk#++aDNNDMMJK:#dhh.33d3KeSS
'1a)
 	 BC|E!H&E!Hkk#*%G|r    )	r   r   r   r
   r   r`   r   rk   rl   rm   s   @r!   r   r      s!    <
Y <
3 <
|(5r    r   c                   .     e Zd Zdedef fdZd Z xZS )
Rwkv7Layerr   r   c                    t         |           || _        | j                  dk(  r0t        j                  |j
                  |j                        | _        t        || j                        | _	        t        |      | _        t        j                  |j
                  |j                        | _        t        j                  |j
                  |j                        | _        y )Nr   r   r   )r_   r`   r   rz   	LayerNormr   r   pre_normr   attnr   ffn	attn_normffn_normr   s      r!   r`   zRwkv7Layer.__init__t  s    ">>QLL)9)9t}}MDM#DDNNC	%d+d&6&6DMMJT%5%54==Ir    c                     | j                   dk(  r| j                  |      }| j                  | j                  |      ||      \  }}||z   }|| j	                  | j                  |      |      z   }||fS )Nr   )r   r   r   r   r   r   )re   r%   r   r   hr   s         r!   rk   zRwkv7Layer.__call__~  sl    >>Qa AYYt~~a0'5A
7E$((4==+U33G|r    )r   r   r   r
   r   r`   rk   rl   rm   s   @r!   r   r   s  s    JY J3 Jr    r   c                   D     e Zd Zdef fdZdej                  fdZ xZS )
Rwkv7Modelr   c                 P   t         |           t        j                  |j                  |j
                        | _        t        |j                        D cg c]  }t        ||       c}| _
        t        j                  |j
                  |j                        | _        y c c}w )Nr   r   )r_   r`   rz   	Embeddingr   r   
embeddingsr   r   r   layersr   r   r1   )re   r   irg   s      r!   r`   zRwkv7Model.__init__  s{    ,,t8H8HI389O9O3P
3PaJtq)3P
 LL!1!1t}}E	
s   B#r%   c                     | j                  |      }|d gt        | j                        z  }d }t        | j                  |      D ]  \  }} ||||      \  }} | j	                  |      S r$   )r   lenr   zipr1   )re   r%   r   r   layercs         r!   rk   zRwkv7Model.__call__  sh    OOA=FS--EDKK/HE1q'1-JAw 0yy|r    r   rm   s   @r!   r   r     s!    FY F"(( r    r   c                   r     e Zd Zdef fdZd	dej                  fdZd Ze	d        Z
d Ze	d        Z xZS )
Modelr   c                     t         |           || _        |j                  | _        t	        |      | _        |j                  s2t        j                  |j                  |j                  d      | _        y y r   )r_   r`   r   r   r   modelr   rz   r   r   r   lm_head)re   r   rg   s     r!   r`   zModel.__init__  sX    	//%
''99T%5%5tUSDL (r    rJ   c                     | j                  ||      }| j                  j                  r'| j                   j                  j	                  |      }|S | j                  |      }|S r$   )r   r   r   r   	as_linearr   )re   rJ   r   r%   logitss        r!   rk   zModel.__call__  sT    JJvu%99((ZZ**44Q7F  \\!_Fr    c                 x    t        t        | j                              D cg c]  }t        d       c}S c c}w )N   )size)r   r   r   r   )re   r   s     r!   
make_cachezModel.make_cache  s/    -23t{{3C-DE-D#-DEEEs   7c                 .    | j                   j                  S r$   )r   r   )re   s    r!   r   zModel.layers  s    zz   r    c                     |j                         D ]i  \  }}d|v s	d|v sd|v s||   j                  | j                  j                  | j                  j                  z  | j                  j                        ||<   k |S )Nr   r   r   )itemsr   r   r   r   )re   weightsr6   r7   s       r!   sanitizezModel.sanitize  sp    MMODAqzUaZ8q=$QZ//II))TYY-?-??ASAS
 $
 r    c                     d }|S )Nc                     d| v sd| v rddiS y)Nzlora.2r   bits   Tr   )pathr   s     r!   	predicatez(Model.quant_predicate.<locals>.predicate  s    4<4#7{"r    r   )re   r  s     r!   quant_predicatezModel.quant_predicate  s    	
 r    r$   )r   r   r   r
   r`   r.   r   rk   r   propertyr   r   r  rl   rm   s   @r!   r   r     sT    TY Trxx F ! !  r    r   )!dataclassesr   	functoolsr   typingr   mlx.corecorer.   mlx.nnrz   baser   r   r   r
   compiler(   r2   r<   rK   rY   r   r@   Moduler]   ro   r   r   r   r   r   r   r   r    r!   <module>r     s   "       & & & 	t$ % 	t$K %K  GT !"	xx	xx 
xx 
xx	
 
xx 
xx 884Uryy U!;299 !;HA A2 2.Hbii HV * ()BII )r    