+
    ÆÜ³jŽ  ã                   óä   € ^ RI t ^ RIt^ RIHt ]P
                  R R l4       tRR R llt]P
                  R R l4       tRR lt	]P
                  R	 R
 l4       t
]P
                  R R l4       tR# )é    Nc                óX   € V ^8„  d   QhR\         P                  R\         P                  /# )é   ÚQUANT_BLOCKÚ
SPLIT_SIZE©ÚtlÚ	constexpr)Úformats   "Úp/home/luca/screen-checking/.venv/lib/python3.14/site-packages/bitsandbytes/backends/triton/kernels_8bit_quant.pyÚ__annotate__r      s(   € ÷ .ñ .ô —‘ð.ô —‘ñ.ó    c                 óÜ   € \         P                  ! ^ R7      pWv,          pV\         P                  ! ^ V4      ,           p	W”8  p
\        W	W#W¥4      p\         P                  ! W,           Wº4       R# )r   ©ÚaxisN)r   Ú
program_idÚarangeÚ"dequant_8bit_blockwise_kernel_utilÚstore)Úa_ptrÚout_ptrÚcode_ptrÚ
absmax_ptrÚnr   r   ÚpidÚblock_startÚoffsetsÚmaskÚout_dqs   &&&&&&&     r   Údequant_8bit_kernelr      sV   € ô �-Š-˜QÔ
€CØÕ"€KØœBŸIšI a¨Ó4Õ4€GØ‰;€DÜ/°ÀÐVZÓh€FÜ‡H‚HˆWÕ Ö-r   c                óÄ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        R\         P                  R\         P                  /# )r   ÚaÚabsmaxÚquant_state_codeÚquant_blocksizeÚdtypeÚout)ÚtorchÚTensorÚintr%   )r
   s   "r   r   r   -   sX   € ÷ ñ Ü‡|�|ðä�L‰Lðô —l‘lðô ð	ô
 �;‰;ðô 
�‰ñr   c           	      óô   € V P                  4       pVf2   Vf   \        R4      h\        P                  ! WV P                  R7      pRp\
        P                  ! Wg4      3p\        V,          ! V VVVVVV4       V# )Nz'If out is None, dtype must be specified)r%   Údeviceé   )ÚnumelÚ
ValueErrorr'   Ú
empty_liker+   ÚtritonÚcdivr   )	r!   r"   r#   r$   r%   r&   r   r   Úgrids	   &&&&&&   r   Údequant_8bit_blockwiser3   -   sz   € ð 	
�‰‹	€AØ
‚{ØŠ=ÜÐFÓGÐGÜ×Ò˜q°a·h±hÔ?ˆà€Jä�KŠK˜Ó&Ð(€DÜ˜ÖØ	ØØØØ	ØØôð €Jr   c                óx   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  /# )r   Ú
BLOCK_SIZEÚ	CODE_SIZEÚSPLIT_NUM_BLOCKSr   )r
   s   "r   r   r   T   s6   € ÷ 6ñ 6ô —‘ð6ô �|‰|ð6ô —l‘lñ6r   c                 óª  € \         P                  ! ^ 4      V,          p\         P                  ! ^ Wu,          4      p	W…,          V	,           p
W¤8  p\         P                  ! W
,           VRR7      p\	        WÁWeV4      w  rÞ\         P
                  ! W(,           \         P                  ! ^ V4      ,           V4       \         P
                  ! W:,           WÛR7       R# )r   ç        )r   Úother)r   N)r   r   r   ÚloadÚ#quantize_8bit_blockwise_kernel_utilr   )ÚA_ptrr   r   r   Ú
n_elementsr5   r6   r7   Úblock_start_idxÚ
thread_idxr   r   ÚAÚ	quantizedr"   s   &&&&&&&&       r   Úquantize_8bit_blockwise_kernelrC   S   sš   € ô —m’m AÓ&Ð)9Õ9€OÜ—’˜1Ð.Õ;Ó<€JàÕ*¨ZÕ7€GØÑ€Dä
�Š�• d°#Ô6€Aä;¸AÈÐ`pÓqÑ€IÜ‡H‚HˆZÕ)¬B¯IªI°aÐ9IÓ,JÕJÈFÔSÜ‡H‚HˆWÕ 	×5r   c                 óÊ  € V P                  4       pWR) ,          ) pVf/   \        P                  ! V3V P                  V P                  R7      pVf5   \        P
                  ! V P                  4       \        P                  R7      p^p\        P                  ! Wg4      3p\        V,          ! V VVVVVVP                  4       VR7       VP                  V P                  4      pWC3# )N)r+   r%   ©r%   )r=   r   r   r   r>   r5   r6   r7   )r-   r'   Úemptyr+   r%   r/   ÚflattenÚuint8r0   r1   rC   ÚreshapeÚshape)	rA   ÚcodeÚ	blocksizer"   r&   r   ÚblocksÚsplit_num_blocksr2   s	   &&&&&    r   Úquantize_blockwise_tritonrO   k   s¹   € Ø	�‰‹	€AØ�J�Ð€Fà‚~Ü—’˜f˜Y¨q¯x©x¸q¿w¹wÔGˆØ
‚{Ü×Ò˜qŸy™y›{´%·+±+Ô>ˆàÐÜ�KŠK˜Ó1Ð3€Dä" 4Ö(ØØØØØØØ—*‘*“,Ø)õð �+‰+�a—g‘gÓ
€Càˆ;Ðr   c                óx   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  /# )r   r6   r5   ÚN_PER_THr   )r
   s   "r   r   r   ‰   s6   € ÷ '"ñ '"ô �|‰|ð'"ô —‘ð	'"ô
 �l‰lñ'"r   c                 óô  € \         P                  ! WV34      p\         P                  ! \         P                  ! V4      ^R7      pWVR,          ,          p\         P                  ! VRR4      p\         P
                  ! WC3\         P                  R7      p\         P                  ! WC3V^,
          \         P                  R7      p	\        ^4       F`  p
W‰,           ^,          p\         P                  ! W,           4      pW|8„  p\         P                  ! WÛV4      p\         P                  ! WÙV4      p	Kb  	  \         P                  ! W,           4      p\         P                  ! W,           4      p\         P                  ! W~,
          4      p\         P                  ! W,
          4      p\         P                  ! VV8*  W‰4      P                  \         P                  4      p\         P                  ! VW4,          34      pVV3# )é   r   g      ð?rE   ):NNNNg      ð¿)r   rI   ÚmaxÚabsÚclampÚzerosÚint32ÚfullÚranger;   ÚwhereÚtorH   )r!   r   r6   r5   rQ   Ú
a_reshapedr"   Úa_normalizedÚlower_pivotÚupper_pivotÚ_ÚpivotÚvalÚ	is_higherÚ	lower_valÚ	upper_valÚ
lower_distÚ
upper_distrB   Úquantized_flats   &&&&&               r   r<   r<   ˆ   sj  € ô —’˜A¨*Ð5Ó6€Jô �VŠV”B—F’F˜:Ó&¨QÔ/€Fà w¥Õ/€LÜ—8’8˜L¨$°Ó4€Lä—(’(˜HÐ1¼¿¹ÔB€KÜ—'’'˜8Ð0°)¸aµ-ÄrÇxÁxÔP€Kô �1ŽXˆØÕ*¨qÕ0ˆÜ�gŠg�hÕ&Ó'ˆØ Ñ&ˆ	Ü—h’h˜y°Ó=ˆÜ—h’h˜y°uÓ=Šñ ô —’˜Õ.Ó/€IÜ—’˜Õ.Ó/€IÜ—’˜Õ0Ó1€JÜ—’˜Õ0Ó1€JÜ—’˜ zÑ1°;ÓL×OÑOÔPR×PXÑPXÓY€Iô —Z’Z 	¨JÕ,AÐ+CÓD€NØ˜6Ð!Ð!r   c                ó8   € V ^8„  d   QhR\         P                  /# )r   r5   r   )r
   s   "r   r   r   ´   s   € ÷ ñ ô —‘ñr   c                 ó  € \         P                  ! W,           V^ R7      P                  \         P                  4      p\         P                  ! W&,           V4      pW,          p\         P                  ! W8,           VRRR7      p	Wy,          p
V
# )r   )r:   r9   Ú
evict_last)r   r:   Úeviction_policy)r   r;   r\   rH   )r   r   r   r   r   r5   r!   Úscaled_int8Úabsmax_offsetsr"   r   s   &&&&&&     r   r   r   ³   sf   € ô 	�Š�• ¨QÔ/×2Ñ2´2·8±8Ó<€AÜ—'’'˜(�,¨Ó-€KàÕ*€NÜ�WŠW�ZÕ0°tÀ3ÐXdÔe€FàÕ!€FØ€Mr   )é@   NN)NN)r'   r0   Útriton.languageÚlanguager   Újitr   r3   rC   rO   r<   r   © r   r   Ú<module>ru      sz   ðÛ ã Ý ð. ‡�ô.ó ð.÷"ðL ‡�ô6ó ð6ô.ð: ‡�ô'"ó ð'"ðT ‡�ôó òr   