
    o9i                     "   d dl Z d dlZd dlZd dlmZ  e       sdej
                  dej
                  fdZyd dlZd dlm	Z
 d dlmZmZ  ej                   ej                  i dd	       ej                  i d
d	       ej                  i dd	       ej                  i dd	       ej                  i d       ej                  i d
       ej                  i d       ej                  i d       ej                  i d       ej                  i d
       ej                  i d       ej                  i d      gdg      ej                   de
j"                  de
j"                  fd              Zdej
                  dej
                  fdZy)    N)is_triton_availablexstate_xc                      y )N )r   r   s     p/var/www/html/backtest/airagagent/rag_env/lib/python3.12/site-packages/bitsandbytes/triton/dequantize_rowwise.pydequantize_rowwiser	      s    4    )early_config_pruneestimate_matmul_time      )
num_stages	num_warps      )r   )r   
n_elements)configskey
BLOCK_SIZEP2c                 $   t        j                  d      }||z  }t        j                  d|      }	||	z   }
|	|k  }t        j                  | |
z   |      }t        j                  ||z         }||z  |z  }t        j                  ||
z   ||       y )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)      s    6 mm#J&1b!&J&GGEGO(3'''C-(1w&
g%vH=r
   c           	          t        j                   j                   j                  t         j                  d}t        dt        j                  t        j                   j                  d               z        } j                  r|j                  sJ |j                         } fd}t        |    ||d| j                  d   |       |S )N)devicedtyper   r   c                 $    j                   d   fS )Nr   )shape)metar   s    r   <lambda>z$dequantize_rowwise.<locals>.<lambda>>   s    QWWQZM r
   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!!WffjUVU\U\]^U_dfgr
   )r5   r1   time bitsandbytes.triton.triton_utilsr   Tensorr	   tritontriton.languagelanguager   triton.ops.matmul_perf_modelr   r   autotuneConfigjit	constexprr)   r   r
   r   <module>rF      so      @OellOU\\O  U
 V__bQ!<bQ!<bQ!<bQ!<bQ/bQ/bQ/bQ/bA.bA.bA.bA. " ZZ> LL> LL> #$>(	ell 	U\\ 	r
   