
    (HJj                         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               h/Users/ahmed/devFolder/claude-voice/.venv/lib/python3.12/site-packages/mlx_lm/models/bitlinear_layers.pybitnet_quantizer%   	   s    O0445MrR3F  %U%7%7%9299CVCVWf --*VRYY2O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    &FP 77;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hh 3[AR$8!HHcUO,1DIDIr&   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!   ~"		"nR01A,-GG)k((!!''ww%8!88!! e')B)BC,.	 *\9Q>B"0,?@ '
  !$ ~"#++A~cr2ALAC
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99 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&   