Ë
    îÿæiê  ã                   ó‚   — d dl mZ d dlmZ d dlmZ d dlmZm	Z	 de
fd„Zd„ Z e«       Z G d„ dej                  «      Zy)	é    N)ÚQuantizedLinear)Útree_flattenÚtree_unflattenÚquantization_configc                 óà  — g }|j                  dg «      }|j                  dd«      dk7  }t        | j                  «       t        j                  j
                  ¬«      D ]f  \  }}||vsŒt        |t        j                  «      sŒ&|j                  }|j                  \  }}	d|v }
t        |	||
|¬«      }|j                  ||f«       Œh t        |«      dkD  r| j                  t        |«      «       | S )	NÚmodules_to_not_convertÚlinear_classÚ Úautobitlinear)Úis_leafÚbias)r   Úinvert_weight_scalesr   )Ú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               ús/Volumes/fast/ai/experiments/voice-extract-mac/.venv/lib/python3.12/site-packages/mlx_lm/models/bitlinear_layers.pyÚbitnet_quantizer%   	   sí   € Ø€OØ0×4Ñ4Ð5MÈrÓRÐà×Ñ °Ó3°ÑFð ô % U×%7Ñ%7Ó%9Ä2Ç9Á9×CVÑCV×W‰ˆˆfð Ð-Ò-´*¸VÄRÇYÁYÕ2OØŸ™ˆJØ(2×(8Ñ(8Ñ%ˆL˜+Ø˜VÐ#ˆDÜ!ØØØØ%9ô	ˆIð ×"Ñ" D¨)Ð#4Õ5ð Xô ˆ?Ó˜aÒØ×Ñœ^¨OÓ<Ô=Ø€Ló    c                  óR   — d} t         j                  j                  dg d¢dg| ¬«      S )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)ÚxÚpacked_weightsÚweight_scaleÚout)r   Úinput_namesÚoutput_namesÚsource)ÚmxÚfastÚmetal_kernel)r/   s    r$   Úmake_bitlinear_kernelr3   #   s6   € ð&€FôP �7‰7×ÑØÚ;Ø�WØð	  ó ð r&   c                   ó4   ‡ — e Zd ZdZ	 	 dˆ fd„	Zd„ Zd„ Zˆ xZS )r   zA
    BitLinear module with memory-efficient weight handling.
    c                 ó4  •— t         ‰| �  «        || _        || _        |dz   dz  }t	        j
                  ||ft        j                  ¬«      | _        || _        t	        j                  dg«      | _
        |rt	        j
                  |f«      | _        y d | _        y )Né   é   )Údtypeg      ð?)ÚsuperÚ__init__r"   r!   r0   ÚzerosÚuint8r   r   Úarrayr+   r   )Úselfr"   r!   r   r   Úpacked_out_featuresÚ	__class__s         €r$   r:   zBitLinear.__init__a   s„   ø€ ô 	‰ÑÔØ&ˆÔØ(ˆÔð  ,¨aÑ/°AÑ5ÐÜ—h‘hÐ 3°[ÐAÌÏÉÔRˆŒà$8ˆÔ!ÜŸH™H c U›OˆÔáÜŸ™ , Ó1ˆD�IàˆD�Ir&   c                 óÄ  — |j                   }t        |«      dkD  r|j                  d|d   «      }|j                   \  }}| j                  }| j                  j
                  }|j
                  |k(  sJ d«       ‚t        ||| j                  gd|fd| j                  fd|fd|fgd||z  d	z  d
fd||fg|g¬«      d   }t        |«      dkD  r |j                  g |d d ¢|‘­Ž }|S )Né   éÿÿÿÿzWrong type for input.ÚTr   r"   r!   é    r7   é   )rE   rF   rF   )ÚinputsÚtemplateÚgridÚthreadgroupÚoutput_shapesÚoutput_dtypesr   )r   r   Úreshaper!   r+   r8   Ú_bitlinear_kernelr   )	r>   r)   r*   Úoriginal_shapeÚtotal_batch_elementsr"   r!   r8   r,   s	            r$   Úexecute_matmul_kernelzBitLinear.execute_matmul_kernely   s!  € ØŸ™ˆÜˆ~Ó Ò"Ø—	‘	˜"˜n¨RÑ0Ó1ˆAØ,-¯G©GÑ)Ð˜kà×(Ñ(ˆà×!Ñ!×'Ñ'ˆØ�w‰w˜%ÒÐ8Ð!8Ó8ÐÜàØØ×!Ñ!ðð �e�Ø'¨×)BÑ)BÐCØ Ð,Ø Ð.ð	ð Ð*¨\Ñ9¸QÑ>ÀÐBØ"Ø0°,Ð?Ð@Ø ˜'ô
ð  ñ!ˆô$ ˆ~Ó Ò"Ø�#—+‘+ÐA˜~¨c¨rÐ2ÐA°LÒAˆCØˆ
r&   c                 ó–   — | j                  || j                  «      }| j                  � t        j                  || j                  «      }|S )N)rQ   r   r   r0   Úadd)r>   r)   Úys      r$   Ú__call__zBitLinear.__call__™   s;   € Ø×&Ñ& q¨$¯+©+Ó6ˆà�9‰9Ð Ü—‘�q˜$Ÿ)™)Ó$ˆAØˆr&   )TF)Ú__name__Ú
__module__Ú__qualname__Ú__doc__r:   rQ   rU   Ú__classcell__)r@   s   @r$   r   r   \   s    ø„ ñð Ø"õò0ö@r&   r   )Úmlx.coreÚcorer0   Úmlx.nnr   Úmlx.nn.layers.quantizedr   Ú	mlx.utilsr   r   Údictr%   r3   rN   r   r   © r&   r$   Ú<module>rb      sB   ðõ Ý Ý 3ß 2ð°ó ò43ñl *Ó+Ð ôB�—	‘	õ Br&   