+
    TV-jê  ã                   ó„   € ^ RI Ht ^ RIHt ^ RIHt ^ RIHtH	t	 R R lt
R t]! 4       t ! R R]P                  4      tR# )	é    N)ÚQuantizedLinear)Útree_flattenÚtree_unflattenc                ó$   € V ^8„  d   QhR\         /# )é   Úquantization_config)Údict)Úformats   "Úo/Volumes/fast/ai/experiments/ui-tars-smoke/.venv/lib/python3.14/site-packages/mlx_lm/models/bitlinear_layers.pyÚ__annotate__r   	   s   € ÷ ñ ´ñ ó    c                 óø  € . pVP                  R . 4      pVP                  RR4      R8g  p\        V P                  4       \        P                  P
                  R7       Fo  w  rVWS9  g   K  \        V\        P                  4      '       g   K/  VP                  pVP                  w  r‰RV9   p
\        V	VV
VR7      pVP                  W[34       Kq  	  \        V4      ^ 8”  d   V P                  \        V4      4       V # )Úmodules_to_not_convertÚlinear_classÚ Úautobitlinear)Úis_leafÚbias)r   Úinvert_weight_scales)Úgetr   Úleaf_modulesÚnnÚModuleÚ	is_moduleÚ
isinstanceÚLinearÚweightÚshapeÚ	BitLinearÚappendÚlenÚupdate_modulesr   )Úmodelr   Úquantize_layersr   r   ÚnameÚmoduleÚ
old_weightÚout_featuresÚin_featuresr   Ú	new_layers   &&          r   Úbitnet_quantizer+   	   sç   € Ø€OØ0×4Ñ4Ð5MÈrÓRÐà×Ñ °Ó3°ÑFð ô % U×%7Ñ%7Ó%9Ä2Ç9Á9×CVÑCV×W‰ˆð Ö-´*¸VÄRÇYÁY×2OÔ2OØŸ™ˆJØ(2×(8Ñ(8Ñ%ˆLØ˜VÑ#ˆDÜ!ØØØØ%9ô	ˆIð ×"Ñ" DÐ#4Ö5ñ Xô ˆ?Ó˜aÔØ×Ñœ^¨OÓ<Ô=Ø€Lr   c                 óR   € Rp \         P                  P                  R. ROR.V R7      # )z©
Custom Metal kernel that performs matrix multiplication directly on
packed weights and scales the output. This eliminates the need to
store unpacked weights in memory.
aŸ  
    constexpr int M = 4;
    constexpr int BLOCK = 32;

    uint tid = thread_position_in_grid.y;
    uint in_offset = thread_position_in_grid.x;

    uint batch_idx = tid / (out_features / 4);
    uint row_idx = tid % (out_features / 4);

    float sum[4] = {0.0};

    for (uint i = in_offset * M; i < in_features; i += BLOCK * M) {
        float v[M];
        for (int j=0; j<M; j++) {
            v[j] = x[batch_idx * in_features + i + j];
        }

        for (int j=0; j<M; j++) {
            uint8_t w = packed_weights[row_idx * in_features + i + j];
            sum[0] += v[j] * ((w & 3) - 1);
            sum[1] += v[j] * (((w >> 2) & 3) - 1);
            sum[2] += v[j] * (((w >> 4) & 3) - 1);
            sum[3] += v[j] * (((w >> 6) & 3) - 1);
        }
    }

    for (int j=0; j<4; j++) {
        sum[j] = simd_sum(sum[j]);
    }

    // Apply weight scaling by diving them or multiplying them
    if (in_offset == 0) {
        float scale = invert_weight_scales ? 1 / weight_scale[0] : weight_scale[0];
        for (int i=0; i<4; i++) {
            out[batch_idx * out_features + row_idx + i * (out_features/4)] = static_cast<T>(sum[i] * scale);
        }
    }
    Úbitlinear_matmulÚout)r%   Úinput_namesÚoutput_namesÚsource)ÚxÚpacked_weightsÚweight_scale)ÚmxÚfastÚmetal_kernel)r1   s    r   Úmake_bitlinear_kernelr8   #   s6   € ð&€FôP �7‰7×ÑØÚ;Ø�WØð	  ó ð r   c                   óF   a a€ ] tR t^\t oRtRV 3R lltR tR tRtVt	V ;t
# )r   z9
BitLinear module with memory-efficient weight handling.
c                óR  <€ \         SV `  4        Wn        W n        V^,           ^,          p\        P
                  ! WQ3\        P                  R7      V n        W@n        \        P                  ! R.4      V n
        V'       d   \        P
                  ! V34      V n        R# RV n        R# )é   )Údtypeg      ð?N)ÚsuperÚ__init__r)   r(   r5   ÚzerosÚuint8r   r   Úarrayr4   r   )Úselfr)   r(   r   r   Úpacked_out_featuresÚ	__class__s   &&&&& €r   r>   ÚBitLinear.__init__a   s|   ø€ ô 	‰ÑÔØ&ÔØ(Ôð  ,¨aÕ/°AÕ5ÐÜ—h’hÐ 3ÐAÌÏÉÔRˆŒà$8Ô!ÜŸHšH c U›OˆÔçÜŸš , Ó1ˆDŽIàˆDŽIr   c           	     óì  € VP                   p\        V4      ^8”  d   VP                  RVR,          4      pVP                   w  rEV P                  pV P                  P
                  pVP
                  V8X  g   Q R4       h\        VVV P                  .RV3RV P                  3RV3RV3.^ WF,          ^,          ^3R	WF3.V.R7      ^ ,          p\        V4      ^8”  d   VP                  ! . VRR OVN5!  pV# )
r   zWrong type for input.ÚTr   r)   r(   )ÚinputsÚtemplateÚgridÚthreadgroupÚoutput_shapesÚoutput_dtypesNéÿÿÿÿ)é    é   rP   )r   r!   Úreshaper(   r4   r<   Ú_bitlinear_kernelr   )	rB   r2   r3   Úoriginal_shapeÚtotal_batch_elementsr)   r(   r<   r.   s	   &&&      r   Úexecute_matmul_kernelÚBitLinear.execute_matmul_kernely   s  € ØŸ™ˆÜˆ~Ó Ô"Ø—	‘	˜"˜n¨RÕ0Ó1ˆAØ,-¯G©GÑ)Ðà×(Ñ(ˆà×!Ñ!×'Ñ'ˆØ�w‰w˜%ÔÐ8Ð!8Ó8ÐÜàØØ×!Ñ!ðð �e�Ø'¨×)BÑ)BÐCØ Ð,Ø Ð.ð	ð Ð*Õ9¸QÕ>ÀÐBØ"Ø0Ð?Ð@Ø ˜'ô
ð  õ!ˆô$ ˆ~Ó Ô"Ø—+’+ÐA˜~¨c¨rÐ2ÐA°LÓAˆCØˆ
r   c                ó˜   € V P                  WP                  4      pV P                  e!   \        P                  ! W P                  4      pV# )N)rU   r   r   r5   Úadd)rB   r2   Úys   && r   Ú__call__ÚBitLinear.__call__™   s7   € Ø×&Ñ& q¯+©+Ó6ˆà�9‰9Ò Ü—’�qŸ)™)Ó$ˆAØˆr   )r   r)   r   r(   r   r4   )TF)Ú__name__Ú
__module__Ú__qualname__Ú__firstlineno__Ú__doc__r>   rU   rZ   Ú__static_attributes__Ú__classdictcell__Ú__classcell__)rD   Ú__classdict__s   @@r   r   r   \   s   ù‡ € ñ÷ò0÷@ò r   r   )Úmlx.coreÚcorer5   Úmlx.nnr   Úmlx.nn.layers.quantizedr   Ú	mlx.utilsr   r   r+   r8   rR   r   r   © r   r   Ú<module>rk      s;   ðõ Ý Ý 3ß 2õò43ñl *Ó+Ð ôB�—	‘	ö Br   