
    bi;"                     `   d dl Z d dlmZ  e       sd Zyd dlZd dlmZ ddlm	Z	m
Z
 d Zd Z ej                   ej                  dd	d
dddd       ej                  d	dd
dddd       ej                  d	dd
dddd       ej                  dd	d
dddd       ej                  ddd
dddd       ej                  ddd
dddd       ej                  ddd
dddd       ej                  dd
d
dddd       ej                  dd
d
dddd       ej                  dd	ddddd       ej                  d	dddddd       ej                  d	dddddd       ej                  dd	ddddd       ej                  ddddddd       ej                  ddddddd       ej                  ddddddd       ej                  dd
ddddd       ej                  dd
ddddd      g e       g de	e
dd       ej                  dd i      ej                   dej"                  dej"                  dej"                  dej"                  dej"                  dej"                  dej"                  dej"                  d ej"                  fd!                     Zd" Zy)#    N)is_triton_availablec                      y N )abstate_xstate_wbiass        k/home/cdr/jupyterlab/.venv/lib/python3.12/site-packages/bitsandbytes/triton/int8_matmul_mixed_dequantize.pyint8_matmul_mixed_dequantizer      s           )early_config_pruneestimate_matmul_timec                       fdS )Nc                 *    |    j                         S r   )zero_)nargsnames    r   <lambda>zinit_to_zero.<locals>.<lambda>   s    U4[..0 r   r   )r   s   `r   init_to_zeror      s	    00r   c                  &   g } dD ]  }dD ]  }dD ]{  }dD ]t  }|dk  rdnd}| j                  t        j                  |||dd	||
             dD ]8  }| j                  t        j                  ||||d	||t        d                   : v }   | S )N)               )       )r    @   )r    r!         r!   r   r   r   BLOCK_MBLOCK_NBLOCK_KSPLIT_K
num_stages	num_warps)r   r      r   C)r*   r+   pre_hook)appendtritonConfigr   )configsr*   block_mblock_kblock_nr+   split_ks          r   get_configs_io_boundr7      s    ) 	J# ' G#5 )0BAA	"MM,3T[hi j+5*3 (5 G#NN &07GX_ls$t/9.7-9#->	!"	, r   r"   r#   r    r$   r   r,   r)   r!   r   r   r   )MNK
   )r   
perf_modeltop_k)r2   keyprune_configs_byEVEN_Kc                 *    | d   | d   | d   z  z  dk(  S )Nr:   r'   r(   r   r   )argss    r   r   r   N   s"    49Y$y/0Q#RVW#W r   	divfactorhas_biasr%   r&   r'   GROUP_Mr(   ACC_TYPEc                 n   t        j                  d      }t        j                  d      }t        j                  ||      }t        j                  ||      }||z  }||z  }t        |||z  z
  |      }||z  ||z  z   }||z  |z  } ||z  t        j                  d|      z   }!| |z  t        j                  d|      z   }"t        j
                  t        j                  |!|z  |      |      }#t        j
                  t        j                  |"|z  |      |      }$||z  t        j                  d|      z   }%| |#d d d f   |z  |%d d d f   |z  z   z   } ||%d d d f   |z  |$d d d f   |z  z   z   }||z  t        j                  d|      z   }!| |z  t        j                  d|      z   }"t        j                  |      }&t        j                  ||#z         d d d f   }'t        j                  ||ft         j                        }(t        dt        j                  |||z              D ]  })|r+t        j                  |       }*t        j                  |      }+nO||)||z  z  z
  },t        j                  | |%d d d f   |,k  d      }*t        j                  ||%d d d f   |,k  d      }+|(t        j                  |*|+      z  }(| ||z  |z  z  } |||z  |z  z  } |&|'|(|	z  z  z  }(|(j                  |j                  j                        }(|
rGt        j                  ||"z         j                  |j                  j                        }|(|d d d f   z   }(||!d d d f   |z  |"d d d f   |z  z   z   }|!|k  d d d f   |"|k  d d d f   z  }-|dk(  rt        j                  ||(|-       y t        j                   ||(|-       y )Nr   r   )dtypeg        )maskother)rI   )tl
program_idcdivminarangemax_contiguousmultiple_ofloadzerosint32rangedottorH   
element_tystore
atomic_add).ABr-   r   state_x_ptrstate_w_ptrr8   r9   r:   rC   rD   	stride_am	stride_ak	stride_bk	stride_bn	stride_cm	stride_cnr%   r&   r'   rE   r(   r@   rF   pidpid_zgrid_mgrid_nwidthgroup_id
group_sizepid_mpid_nrmrnramrbnrkw_factorx_factoracckr   r   k_remainingrI   s.                                                 r   _int8_matmul_mixed_dequantizerx   1   s   x mmAa G$G$& %<(W"44g>
7"cJ&67u*-W_ryyG44W_ryyG44rAvw ?IrAvw ?IW_ryyG44QW	)BtQwK),CCDAtGy(3tQw<)+CCD W_ryyG44W_ryyG4477;'77;,-ag6 hh):q"''!Ww%678 
	/AGGAJGGAJ!w'8"99GGABtQwK+$=SIGGABq$wK+$=SI266!Q<C7"Y..A7"Y..A
	/ (cIo67ffQWW''( 774"9%((););<DT1W%CAtGy(2dAg;+BBCQ4 BFD!G#44a<HHQ$'MM!St,r   c                    | j                   }d}|dnd}| j                  d      dkD  r$| j                  d      dkD  r| j                         } |j                  d      dkD  r$|j                  d      dkD  r|j                         }| j                  d   |j                  d   k(  sJ d       | j                  \  }|j                  \  }	t	        j
                  f|t        j                        }
t        j                  }fd}t        |   | ||
||||||| j                  d      | j                  d      |j                  d      |j                  d      |
j                  d      |
j                  d      d|       |
S )	NgA@?r   r   zincompatible dimensions)devicerH   c                 t    t        j                  | d         t        j                  | d         z  | d   fS )Nr%   r&   r(   )r0   rM   )METAr8   r9   s    r   r   z.int8_matmul_mixed_dequantize.<locals>.<lambda>   s7    V[[DO<v{{1dS\o?^^`den`op r   r,   )rE   rF   )
rz   stride
contiguousshapetorchemptyfloat16rK   float32rx   )r   r   r	   r
   r   rz   rC   rD   r:   _crF   gridr8   r9   s                @@r   r   r      sO   )	1!88A;?qxx{QA88A;?qxx{QAwwqzQWWQZ'B)BB'ww1ww1KKAvU]]C::p%d+HHQKHHQKHHQKHHQKHHQKHHQK'	
* r   )r    bitsandbytes.triton.triton_utilsr   r   r0   triton.languagelanguagerK   matmul_perf_modelr   r   r   r7   autotuner1   
heuristicsjit	constexprrx   r   r   r   <module>r      sK    @  K14 V__ FMMccbUVWdeqrsFMMccbUVWdeqrsFMMcbRTUVcdpqrFMMbSRTUVcdpqrFMMccbUVWdeqrsFMMcbRTUVcdpqrFMMbSRTUVcdpqrFMMcbRTUVcdpqrFMMbRBSTUbcopqFMMcccVWXefrstFMMcccVWXefrstFMMcbSUVWdeqrsFMMbSSUVWdeqrsFMMcccVWXefrstFMMcbRTUVcdpqrFMMbSRTUVcdpqrFMMcbRTUVcdpqrFMMbRBSTUbcopq)
* "#+
. 0BRfqst36 VW	

 ZZR- <<R- ,,R-$ %R-& 'R-( )R-* +R-, -R-. /R-0 ,,1R- 7BR-h(r   