
    ;i                        S SK r S SKrS SKJr  \" 5       (       d  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                  4S jrg)    N)is_triton_availablexc                     g )N )r   s    t/home/dmtnaga/Documents/work/airagagent/rag_env/lib/python3.13/site-packages/bitsandbytes/triton/quantize_rowwise.pyquantize_rowwiser   	   s              )
num_stages	num_warps      )r   )r   
n_elements)configskey
BLOCK_SIZEP2c                    [         R                  " SS9nXd-  n[         R                  " SU5      nXx-   n	X:  n
[         R                  " X	-   U
S9n[         R                  " U5      n[         R
                  " [         R                  " XS5      SS9n[         R                  R                  SX-  -  5      n[         R                  " X-   XS9  [         R                  " X&-   U5        g )Nr   )axis)maskg     _@)
tl
program_idarangeloadabsmaxwhere	libdevicellrintstore)x_ptr
output_ptroutput_maxsr   r   r   pidblock_startr   offsetsrow_maskr   abs_xmax_valoutputs                  r   _quantize_rowwiser,      s    4 mm#&1b!&&GGEO(3q	&&(15A>$$Uak%:;
%v=
"G,r	   c           	      (  ^  [         R                  " T R                  T R                  [         R                  S.6n[         R                  " T R                  S   T R                  [         R
                  S9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UT R                  S   US9  X4$ )N)devicedtyper   r   r
   c                 $   > TR                   S   4$ )Nr   )shape)metar   s    r   <lambda>"quantize_rowwise.<locals>.<lambda>A   s    QWWQZMr	   )r   r   )torchemptyr1   r.   int8float16intmathceillog2is_cudanumelr,   )r   r+   r$   r   r   grids   `     r   r   r   9   s    aggahhejjIkk!''!*QXXU]]Styy1771:!6789yyV^^++\\^
)$6
qwwWXz^`a""r	   )r:   r5    bitsandbytes.triton.triton_utilsr   Tensorr   tritontriton.languagelanguager   autotuneConfigjit	constexprr,   r   r	   r   <module>rI      sB     @ELL   
 __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 
#r	   