
    ;i                        S SK r S SKrS SKJr  \" 5       (       d  S\R                  S\R                  4S jrgS SKrS SKJr	  \R                  " \R                  " 0 SSS9\R                  " 0 S	SS9\R                  " 0 S
SS9\R                  " 0 SSS9\R                  " 0 SS9\R                  " 0 S	S9\R                  " 0 S
S9\R                  " 0 SS9\R                  " 0 SS9\R                  " 0 S	S9\R                  " 0 S
S9\R                  " 0 SS9/S/S9\R                  S\	R                  S\	R                  4S j5       5       rS\R                  S\R                  4S jrg)    N)is_triton_availablexstate_xc                     g )N )r   r   s     v/home/dmtnaga/Documents/work/airagagent/rag_env/lib/python3.13/site-packages/bitsandbytes/triton/dequantize_rowwise.pydequantize_rowwiser	   	   s              )
num_stages	num_warps      )r   )r   
n_elements)configskey
BLOCK_SIZEP2c                    [         R                  " SS9nXu-  n[         R                  " SU5      n	X-   n
X:  n[         R                  " X
-   US9n[         R                  " X-   5      nX-  U-  n[         R                  " X*-   XS9  g )Nr   )axis)mask)tl
program_idarangeloadstore)x_ptrr   
output_ptrinv_127r   r   r   pidblock_startr   offsetsrow_maskr   max_valoutputs                  r   _dequantize_rowwiser'      sy    6 mm#&1b!&&GGEO(3'''-(w&
%v=r
   c           
        ^  [         R                  " T R                  T R                  [         R                  S.6n[        S[        R                  " [        R                  " T R                  S   5      5      -  5      nT R                  (       a  UR                  (       d   eUR                  5       nU 4S jn[        U   " T XSUT R                  S   US9  U$ )N)devicedtyper   r   c                 $   > TR                   S   4$ )Nr   )shape)metar   s    r   <lambda>$dequantize_rowwise.<locals>.<lambda>>   s    QWWQZMr
   g@ ?)r   r   )torchemptyr,   r)   float16intmathceillog2is_cudanumelr'   )r   r   r&   r   r   grids   `     r   r	   r	   7   s    aggahhemmLtyy1771:!6789yyV^^++\\^
)D!!WiXYX_X_`aXbgijr
   )r4   r0    bitsandbytes.triton.triton_utilsr   Tensorr	   tritontriton.languagelanguager   autotuneConfigjit	constexprr'   r   r
   r   <module>rC      sX     @ell U\\   
 __MM"a8MM"a8MM"a8MM"a8MM"+MM"+MM"+MM"+MM"*MM"*MM"*MM"*
 N" ZZ> LL> LL> #$>&	ell 	U\\ 	r
   