+
    ÆÜ³jW¦  ã                   óð  € ^ RI Ht ^ RIt^ RIt^ RIHt ^ RIHt ^ RI	H
t
 ^ RIt^ RIHtHtHt ^RIHt ^RIHt R^R	 lt]! R_ U Uu. uF  p R` F  pR
V  RV 2NK  	  K  	  upp R_ U u. uF  p R
V  2NK
  	  up ,           ]P*                  .^,          ]P,                  ]P,                  ]P*                  .,           4       ]! Ra U u. uF  p RV  2NK
  	  up ]P*                  .^,          ]P,                  ]P,                  ]P,                  ]P,                  ]P,                  ]P*                  .,           4       ]! Ra U u. uF  p RV  2NK
  	  up ]P,                  .^,          ]P*                  .^,          ,           ]P,                  .^,          ,           ]P,                  ]P*                  .,           4       ]! R.]P*                  .]P,                  .^,          ,           ]P*                  .^,          ,           ]P,                  .^,          ,           ]P*                  .,           ]P,                  R7       ]! R.]P*                  .^,          ]P,                  ]P,                  ]P*                  .,           4       ]! R.]P*                  .^,          ]P.                  ]P,                  ]P,                  ]P*                  .,           4       ]! R_ U Uu. uF  p R` F  pRV  RV 2NK  	  K  	  upp R_ U u. uF  p RV  2NK
  	  up ,           ]P*                  .^,          ]P,                  ]P,                  .,           4       ]P0                  P2                  t]P6                  R 4       t]! RR4      R R l4       t]! RR4      R R l4       tR R lt]! RR4      RbR R ll4       t]! R R4      RcR! R" ll4       t]! R#R4      RcR$ R% ll4       tRcR& R' llt]! R(R4      R) R* l4       t]! R+R4      R, R- l4       t]! R.R4      R/ R0 l4       tR1 R2 lt ]! R3R4      R4 R5 l4       t]! R6R4      R7 R8 l4       t]! R9R4      R: R; l4       tR< R= lt!]! R>R4      R? R@ l4       t]! RAR4      RB RC l4       tRD RE lt"]P6                  RF 4       t#]P6                  RG 4       t$]P6                  RH 4       t%RdRI lt&RdRJ lt'^t(]PR                  PT                  f   ]#t+RKt,M]$t+RLt,]! RMR4      RdRN RO ll4       t RP]PZ                  ]P\                  ]P^                  3RQ]P`                  ]Pb                  3RR]Pd                  ]Pf                  3RS]Ph                  ]Pj                  ]Pl                  3RT]Pn                  ]Pp                  3RU]PZ                  ]P\                  ]P^                  3RV]Pr                  ]Pt                  ]Pv                  3RW]P`                  ]Pb                  3/t<RP]Pz                  ]P|                  ]P~                  3RQ]P€                  ]P‚                  ]P„                  3RR]P†                  ]Pˆ                  ]PŠ                  3RS]PŒ                  ]PŽ                  ]P�                  3RT]P’                  ]P”                  ]P–                  3RV]P˜                  ]Pš                  ]Pœ                  3/tOReRX RY lltPReRZ R[ lltQ]! R\R4      ! ]Q4       ]! R]R4      ! ]P4       R# u upp i u up i u up i u up i u upp i u up i )fé    )ÚSequenceN)Úprod)ÚOptional)Úwarn)ÚCUBLAS_ContextÚ_cuda_device_ofÚget_ptr)Úregister_kernel)Úlibc                 óP   € V  F  p\        \        V4      pWn        W$n        K!  	  R # ©N)Úgetattrr   ÚargtypesÚrestype)Únamesr   r   ÚnameÚfns   &&&  Ú_/home/luca/screen-checking/.venv/lib/python3.14/site-packages/bitsandbytes/backends/cuda/ops.pyÚ_setup_ctypesr      s"   € ÛˆÜ”S˜$ÓˆØŒØŽ
ó ó    Úcdequantize_blockwise_Ú_Úcgemm_4bit_Úcgemm_4bit_inference_naive_Úcigemmlt_32)r   Úcdequant_mm_int32_fp16Úcint8_vector_quantÚcquantize_blockwise_c                 ó†   € \         P                  P                  V 4      pVP                  VP                  VP
                  3# r   )ÚtorchÚcudaÚget_device_propertiesÚmulti_processor_countÚmajorÚminor)Údevice_indexÚpropss   & r   Ú_gpu_dispatch_propsr(   H   s1   € ä�J‰J×,Ñ,¨\Ó:€EØ×&Ñ&¨¯©°U·[±[Ð@Ð@r   z bitsandbytes::int8_linear_matmulr!   c                óX   € V ^8„  d   QhR\         P                  R\         P                  /# )é   ÚAÚB©r    ÚTensor)Úformats   "r   Ú__annotate__r0   O   s"   € ÷ /ñ /Œ�‰ð /œ%Ÿ,™,ñ /r   c                 óÂ   € \         P                  ! . V P                  R R OVP                  ^ ,          N5V P                  \         P                  R7      p\        WV4      # )N©ÚdeviceÚdtypeéÿÿÿÿ)r    ÚemptyÚshaper3   Úint32Ú_int8_linear_matmul_impl©r+   r,   Úouts   && r   r   r   N   sF   € ä
�+Š+Ð1˜Ÿ™  ˜Ð1 a§g¡g¨a¥jÑ1¸!¿(¹(Ì%Ï+É+Ô
V€CÜ# A¨#Ó.Ð.r   z$bitsandbytes::int8_linear_matmul.outc                óx   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  /# ©r*   r+   r,   r;   r-   )r/   s   "r   r0   r0   U   s-   € ÷ (ñ (Œ�‰ð (œ%Ÿ,™,ð (¬U¯\©\ñ (r   c                 ó   € \        WV4       R # r   )r9   r:   s   &&&r   r   r   T   s   € ä˜Q 3Ö'r   c                óx   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  /# r=   r-   )r/   s   "r   r0   r0   Y   s2   € ÷ Oñ O¤§¡ð O´·±ð OÄEÇLÁLñ Or   c                 ó  € YrV P                   pVP                   pV P                  \        P                  8w  d   \	        R 4      hVP                  \        P                  8w  d   \	        R4      hV P
                  ^8w  d   \	        R4      hVP
                  R9  d   \	        R4      h\        V4      ^ 8:  d   \	        RV 24      hVP                  \        P                  8w  d   \	        RVP                   24      h. VRR OV^ ,          N5pVP                   V8w  d   \	        RVP                    RV 24      hVw  rg\        VRR 4      pVR,          p	VR,          p
VR,          pWš8w  d   \	        R	V R
V 24      hV	^,          ^ 8w  dp   \        P                  ! VP                  4       V P                  4       P                  4       4      P                  \        P                  4      pVP                  V4      # \        V 4      ;_uu_ 4        \        P                  ! 4       P!                  V P"                  4      p\$        P&                  ! VVVVV P)                  4       VP)                  4       VP)                  4       RV	V
V\+        V P"                  P,                  4      4      pRRR4       X'       d¿   V^d8X  d•   ^ RIpVP1                  RV: RV: R2\2        ^R7       \        P                  ! VP                  4       V P                  4       P                  4       4      P                  \        P                  4      pVP                  V4      # \5        RV: RV: RV: RWšV3: RWxV3: 2
4      hV#   + '       g   i     LØ; i)zB must be int8zA must be int8z:Only two dimensional matrices are supported for argument BzCOnly two or three dimensional matrices are supported for argument Az(Input tensor dimensions need to be > 0: zout must be int32, got NzOutput shape z does not match expected shape zQint8_linear_matmul only supports B^T @ A. Inner dimensions do not match: B @ A = z @ zFint8_linear_matmul has no usable (hip|cu)blasLt algo for shape shapeA=z shapeB=z; falling back to fp32 matmul.)Ú
stacklevelz$cublasLt ran into an error!
	shapeA=z	, shapeB=z	, shapeC=z
	(lda, ldb, ldc)=z
	(m, n, k)=)r*   é   r5   )r7   r4   r    Úint8Ú
ValueErrorÚndimr   r8   ÚmatmulÚfloatÚtÚtoÚcopy_r   r   Úget_instanceÚget_contextr3   r   r   Údata_ptrÚ_get_raw_streamÚindexÚwarningsr   ÚRuntimeWarningÚRuntimeError)r+   r,   r;   ÚshapeAÚshapeBÚshapeCÚkÚmÚnÚldaÚldbÚldcÚresultÚctxÚ	has_errorrP   s   &&&             r   r9   r9   Y   s  € Ø€qà�W‰W€FØ�W‰W€Fà‡w�w”%—*‘*ÔÜÐ)Ó*Ð*Ø‡w�w”%—*‘*ÔÜÐ)Ó*Ð*Ø‡v�v�„{ÜÐUÓVÐVØ‡v�v�VÔÜÐ^Ó_Ð_ÜˆFƒ|�qÔÜÐCÀFÀ8ÐLÓMÐMØ
‡y�y”E—K‘KÔÜÐ2°3·9±9°+Ð>Ó?Ð?à&ˆv�c�rˆ{Ð&˜F 1�IÑ&€FØ
‡y�y�FÔÜ˜=¨¯©¨Ð3RÐSYÐRZÐ[Ó\Ð\à�D€AÜˆV�C�Rˆ[Ó€AØ
��*€CØ
��*€CØ
��*€Cà
„zÜØ_Ð`fÐ_gÐgjÐkqÐjrÐsó
ð 	
ð ˆQ…w�!„|Ü—’˜aŸg™g›i¨¯©«¯©«Ó7×:Ñ:¼5¿;¹;ÓGˆØ�y‰y˜Ó Ð ä	˜×	Õ	Ü×)Ò)Ó+×7Ñ7¸¿¹ÓAˆÜ—O’OØØØØØ�J‰J‹LØ�J‰J‹LØ�L‰L‹NØØØØÜ˜AŸH™HŸN™NÓ+ó
ˆ	÷ 
÷" Ø˜Ôó
 à�M‰MðØ‘)˜9˜V™IÐ%CðEäØð	 ô ô —\’\ !§'¡'£)¨Q¯W©W«Y¯[©[«]Ó;×>Ñ>¼u¿{¹{ÓKˆFØ—9‘9˜VÓ$Ð$äØ9°&±¸*¸V¹IÀZÈÁyÐPeÐVYÐ`cÐUdÑTfÐfuÐlmÐrsÐktÑjvÐwóð ð €J÷M 
×	ús   ÈBM4Í4N	zbitsandbytes::int8_mm_dequantc                ó  € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        \         P                  ,          R\        \         P                  ,          R\         P                  /# )r*   r+   Ú	row_statsÚ	col_statsr4   ÚbiasÚreturn)r    r.   r   r4   )r/   s   "r   r0   r0   ¬   sf   € ÷ &*ñ &*Ü‡|�|ð&*ä�|‰|ð&*ô �|‰|ð&*ô ”E—K‘KÕ ð	&*ô
 ”5—<‘<Õ
 ð&*ô ‡\�\ñ&*r   c                 ól  € V P                   \        P                  8w  d   \        R V P                    24      hVP                   \        P                  8w  d   \        RVP                    24      hVP                   \        P                  8w  d   \        RVP                    24      h\        P
                  ! V \        P                  R7      pVe0   VP                   \        P                  8X  d   VP                  4       MRp\        V 4      ;_uu_ 4        \        P                  ! V P                  4       VP                  4       VP                  4       VP                  4       VV P                  4       V P                  R,          ,          V P                  R,          \        V P                  P                  4      4       RRR4       Ve1   VP                   \        P                  8w  d   VP!                  V4       TP#                  T;'       g    \        P                  4      #   + '       g   i     Lm; i)zA must be int32, got zrow_stats must be float32, got zcol_stats must be float32, got ©r4   Nr5   )r4   r    r8   rD   Úfloat32Ú
empty_likeÚfloat16rM   r   r   r   Únumelr7   rN   r3   rO   Úadd_rI   )r+   r`   ra   r4   rb   r;   Úbias_ptrs   &&&&&  r   r   r   «   sy  € ð 	‡w�w”%—+‘+ÔÜÐ0°·±°	Ð:Ó;Ð;Ø‡�œ%Ÿ-™-Ô'ÜÐ:¸9¿?¹?Ð:KÐLÓMÐMØ‡�œ%Ÿ-™-Ô'ÜÐ:¸9¿?¹?Ð:KÐLÓMÐMô ×
Ò
˜1¤E§M¡MÔ
2€Cð #'Ò"2°t·z±zÄUÇ]Á]Ô7Rˆt�}‰}ŒÐX\€Hä	˜×	Õ	Ü×"Ò"Ø�J‰J‹LØ×ÑÓ Ø×ÑÓ Ø�L‰L‹NØØ�G‰G‹I˜Ÿ™ �Õ$Ø�G‰G�B�KÜ˜AŸH™HŸN™NÓ+ô		
÷ 
ð Ò˜DŸJ™J¬%¯-©-Ô7Ø�‰�Œà�6‰6�%×(Ð(œ5Ÿ=™=Ó)Ð)÷! 
×	ús   ÄB*H#È#H3	z#bitsandbytes::int8_vectorwise_quantc                ó8   € V ^8„  d   QhR\         P                  /# )r*   r+   r-   )r/   s   "r   r0   r0   Ö   s   € ÷ (,ñ (,Œ�‰ñ (,r   c                 ób  € V P                   \        P                  8w  d   \        R V P                    24      hVR8  d   \        R4      hV P	                  4       V P
                  R,          ,          pV P
                  R,          p\        P                  ! W P                  \        P                  R7      p\        P                  ! V P
                  V P                  \        P                  R7      pRpVR8”  d‘   V P                  4       V8¬  pVP                  4       '       d7   \        P                  ! VP                  ^ R7      4      P                  R4      pM1\        P                  ! ^ V P                  \        P                  R7      p\        V 4      ;_uu_ 4        \         P"                  ! V P%                  4       VP%                  4       VP%                  4       VVV\'        V P                  P(                  4      4       RRR4       V^8”  d   Ve   ^ VRV3&   WTV3#   + '       g   i     L&; i)zA must be float16, got ç        zthreshold must be non-negativer2   N)Údim:NNNr5   )r4   r    rh   rD   ri   r7   r6   r3   rf   rC   ÚabsÚanyÚargwhereÚviewÚint64r   r   r   rM   rN   rO   )r+   Ú	thresholdÚrowsÚcolsr`   Úout_rowÚoutlier_colsÚoutlierss   &&      r   r   r   Õ   s†  € à‡w�w”%—-‘-ÔÜÐ2°1·7±7°)Ð<Ó=Ð=Ø�3„ÜÐ9Ó:Ð:à�7‰7‹9˜Ÿ™ �Õ#€DØ�7‰7�2�;€Dä—’˜D¯©¼¿¹ÔG€IÜ�kŠk˜!Ÿ'™'¨!¯(©(¼%¿*¹*ÔE€Gà€Là�3„à—5‘5“7˜iÑ'ˆà�<‰<�>Š>Ü Ÿ>š>¨(¯,©,¸1¨,Ó*=Ó>×CÑCÀBÓG‰Lô !Ÿ;š; q°·±ÄÇÁÔMˆLä	˜×	Õ	Ü×ÒØ�J‰J‹LØ×ÑÓØ×ÑÓ ØØØÜ˜AŸH™HŸN™NÓ+ô	
÷ 
ð ˆa„x�LÒ,Ø#$ˆ��<�Ñ à˜|Ð+Ð+÷! 
×	ús   ÆA$HÈH.	zbitsandbytes::int8_double_quantc          
      óþ   € V ^8„  d   QhR\         P                  R\        \         P                  \         P                  \         P                  \         P                  \        \         P                  ,          3,          /# ©r*   r+   rc   ©r    r.   Útupler   )r/   s   "r   r0   r0     sO   € ÷ Vñ VÜ‡|�|ðVô Œ5�<‰<œŸ™¤u§|¡|´U·\±\Ä8ÌEÏLÉLÕCYÐYÕZñVr   c                 óº  € \         P                  P                  P                  P	                  V VR 7      w  r#p\        WR 7      w  rVVR8”  d   Ve   V P                  VR4      p \         P                  ! V P                  R4      VP                  ^ 4      ,          4      P                  \         P                  4      pW'W5P                  4       P                  4       V3# ))ru   rn   g     À_@)r    ÚopsÚbitsandbytesÚint8_vectorwise_quantÚdefaultÚ_get_col_absmaxÚmasked_fillÚroundÚmulÚ	unsqueezerI   rC   ÚflattenrG   )r+   ru   Ú	quant_rowr`   ry   ra   Úoutlier_maskÚ	quant_cols   &&      r   r   r     s´   € ô */¯©×)?Ñ)?×)UÑ)U×)]Ñ)]Ø	Øð *^ó *Ñ&€I˜,ô .¨aÔEÑ€IØ�3„˜<Ò3Ø�M‰M˜,¨Ó,ˆÜ—’˜AŸE™E %›L¨9×+>Ñ+>¸qÓ+AÕAÓB×EÑEÄeÇjÁjÓQ€Ià ×,=Ñ,=Ó,?×,EÑ,EÓ,GÈÐUÐUr   c                ó¤   € V ^8„  d   QhR\         P                  R\        \         P                  \        \         P                  ,          3,          /# r|   r}   )r/   s   "r   r0   r0     s9   € ÷ #ñ #Ü‡|�|ð#ô Œ5�<‰<œ¤%§,¡,Õ/Ð/Õ0ñ#r   c                 óF  € V P                  4       '       g   \        R V P                   24      hRpV P                  4       P	                  RV P
                  R,          4      pVR8”  d   W18¬  pVP                  VR4       VP                  ^ RR7      P                  4       pWB3# )z'A must be a floating point tensor, got Nrn   F)ro   Úkeepdimr5   )	Úis_floating_pointrD   r4   rp   rs   r7   Úmasked_fill_ÚamaxrG   )r+   ru   r‹   ÚabsAra   s   &&   r   r„   r„     s’   € ð ×Ñ× Ò ÜÐBÀ1Ç7Á7À)ÐLÓMÐMà€Là�5‰5‹7�<‰<˜˜AŸG™G B�KÓ(€Dà�3„àÑ(ˆØ×Ñ˜,¨Ô,ð —	‘	˜a¨�	Ó/×5Ñ5Ó7€IàÐ"Ð"r   z bitsandbytes::quantize_blockwisec          
      óº   € V ^8„  d   QhR\         P                  R\         P                  R\        R\        \         P                  \         P                  3,          /# )r*   r+   ÚcodeÚ	blocksizerc   )r    r.   Úintr~   )r/   s   "r   r0   r0   ,  sE   € ÷  ñ  Œ�‰ð  œUŸ\™\ð  ´cð  ¼eÄEÇLÁLÔRW×R^ÑR^ÐD^Õ>_ñ  r   c           
      óà  € V P                  4       p VP                  \        P                  8w  d   \	        R VP                   24      hVR9  d   \	        RV 24      hV P                  4       pW2) ,          ) p\        P                  ! V3V P                  \        P                  R7      p\        P                  ! V \        P                  R7      pV P                  \        P                  8X  d   \        P                  pMxV P                  \        P                  8X  d   \        P                  pMHV P                  \        P                  8X  d   \        P                  pM\	        RV P                   24      h\!        V 4      ;_uu_ 4        V! VP#                  4       V P#                  4       VP#                  4       VP#                  4       VV4       RRR4       We3#   + '       g   i     Ye3# ; i)zcode must be float32, got úinvalid blocksize r2   re   z?Blockwise quantization only supports 16/32-bit floats, but got N)é@   é€   é   é   é   é   é   )Ú
contiguousr4   r    rf   rD   ri   r6   r3   rg   Úuint8r   Úcquantize_blockwise_fp32rh   Úcquantize_blockwise_fp16Úbfloat16Úcquantize_blockwise_bf16r   rM   )r+   r•   r–   rX   ÚblocksÚabsmaxr;   r   s   &&&     r   r   r   +  se  € à	�‰‹€Aà‡z�z”U—]‘]Ô"ÜÐ5°d·j±j°\ÐBÓCÐCØÐ=Ô=ÜÐ-¨i¨[Ð9Ó:Ð:à	�‰‹	€AØ�J�Ð€FÜ�[Š[˜&˜¨1¯8©8¼5¿=¹=ÔI€FÜ
×
Ò
˜1¤E§K¡KÔ
0€Cà‡w�w”%—-‘-ÔÜ×)Ñ)‰Ø	
�‰”E—M‘MÔ	!Ü×)Ñ)‰Ø	
�‰”E—N‘NÔ	"Ü×)Ñ)‰äÐZÐ[\×[bÑ[bÐZcÐdÓeÐeä	˜×	Õ	Ù
Ø�M‰M‹OØ�J‰J‹LØ�O‰OÓØ�L‰L‹NØØô	
÷ 
ð ˆ;Ð÷ 
Ö	ð ˆ;Ðús   Æ
AGÇG-	z"bitsandbytes::dequantize_blockwisec                óÄ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        R\         P                  R\         P                  /# )r*   r+   r¨   r•   r–   r4   rc   ©r    r.   r—   r4   )r/   s   "r   r0   r0   P  sP   € ÷ ñ Œ�‰ð œuŸ|™|ð ´5·<±<ð ÌCð ÔX]×XcÑXcð Ôhm×htÑhtñ r   c           	      óP   € \         P                  ! WR 7      p\        WW#WER7       V# )re   ©r;   )r    rg   Ú_dequantize_blockwise_impl©r+   r¨   r•   r–   r4   r;   s   &&&&& r   r   r   O  s#   € ä
×
Ò
˜1Ô
*€CÜ˜q¨$¸5ÕJØ€Jr   z&bitsandbytes::dequantize_blockwise.outc                óÈ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        R\         P                  R\         P                  RR/# ©	r*   r+   r¨   r•   r–   r4   r;   rc   Nrª   )r/   s   "r   r0   r0   W  sk   € ÷ Kñ KÜ‡|�|ðKä�L‰LðKô �,‰,ðKô ð	Kô
 �;‰;ðKô 
�‰ðKð 
ñKr   c           	      óú   € VP                   V8w  d   \        R V RVP                    24      hVP                  V P                  8w  d&   \        RV P                   RVP                   24      h\        WW#WER7       R# )úExpected out.dtype == ú, got úExpected out.shape == r¬   N)r4   rD   r7   r­   r®   s   &&&&&&r   r   r   V  sj   € ð ‡y�y�EÔÜÐ1°%°¸¸s¿y¹y¸kÐJÓKÐKØ
‡y�y�A—G‘GÔÜÐ1°!·'±'°¸&ÀÇÁÀÐLÓMÐMÜ˜q¨$¸5×Jr   c                óÈ   € V ^8„  d   QhR\         P                  R\         P                  R\         P                  R\        R\         P                  R\         P                  RR/# r°   rª   )r/   s   "r   r0   r0   f  s\   € ÷ 
ñ 
Ü‡|�|ð
Ü"Ÿ\™\ð
Ü16·±ð
ÜJMð
ÜV[×VaÑVað
Ühm×htÑhtð
à	ñ
r   c                 óh  € V P                  4       p V\        P                  8X  d   \        P                  pMZV\        P
                  8X  d   \        P                  pM4V\        P                  8X  d   \        P                  pM\        R V 24      h\        V 4      ;_uu_ 4        V! VP                  4       V P                  4       VP                  4       VP                  4       VV P                  4       \        V P                  P                  4      4       RRR4       R#   + '       g   i     R# ; i)zABlockwise dequantization only supports 16/32-bit floats, but got N)r¡   r    rf   r   Úcdequantize_blockwise_fp32rh   Úcdequantize_blockwise_fp16r¥   Úcdequantize_blockwise_bf16rD   r   rM   ri   rN   r3   rO   )r+   r¨   r•   r–   r4   r;   r   s   &&&&&& r   r­   r­   f  sÌ   € ð 	
�‰‹€Aà”—‘ÔÜ×+Ñ+‰Ø	”%—-‘-Ô	Ü×+Ñ+‰Ø	”%—.‘.Ô	 Ü×+Ñ+‰äÐ\Ð]bÐ\cÐdÓeÐeä	˜×	Õ	Ù
Ø�M‰M‹OØ�J‰J‹LØ�O‰OÓØ�L‰L‹NØØ�G‰G‹IÜ˜AŸH™HŸN™NÓ+ô	
÷ 
×	×	Ò	ús   Â$A2D Ä D1	zbitsandbytes::quantize_4bitc                óÆ   € V ^8„  d   QhR\         P                  R\        R\        R\         P                  R\
        \         P                  \         P                  3,          /# )r*   r+   r–   Ú
quant_typeÚquant_storagerc   )r    r.   r—   Ústrr4   r~   )r/   s   "r   r0   r0   �  sN   € ÷ #ñ #Ü‡|�|ð#Ü #ð#Ü14ð#ÜEJÇ[Á[ð#ä
Œ5�<‰<œŸ™Ð%Õ&ñ#r   c           
      óÆ  € V P                  4       p V P                  4       pWA) ,          ) p\        P                  ! V3V P                  \        P
                  R 7      p\        P                  ! V^,           VP                  ^,          ,          ^3V P                  VR 7      pV P                  \        P                  8X  d*   VR8X  d   \        P                  pM \        P                  pM�V P                  \        P                  8X  d*   VR8X  d   \        P                  pMX\        P                  pMGV P                  \        P
                  8X  d)   VR8X  d   \        P                  pM\        P                   p\#        V 4      ;_uu_ 4        X! RV P%                  4       VP%                  4       VP%                  4       VV4       RRR4       Wv3#   + '       g   i     Yv3# ; i)r2   Úfp4N)r¡   ri   r    r6   r3   rf   Úitemsizer4   r¥   r   Úcquantize_blockwise_bf16_fp4Úcquantize_blockwise_bf16_nf4rh   Úcquantize_blockwise_fp16_fp4Úcquantize_blockwise_fp16_nf4Úcquantize_blockwise_fp32_fp4Úcquantize_blockwise_fp32_nf4r   rM   )	r+   r–   r»   r¼   rX   r§   r¨   r;   r   s	   &&&&     r   r   r   €  sV  € ð 	
�‰‹€AØ	�‰‹	€AØ�J�Ð€FÜ�[Š[˜&˜¨1¯8©8¼5¿=¹=ÔI€FÜ
�+Š+˜˜A� =×#9Ñ#9¸AÕ#=Õ>ÀÐBÈ1Ï8É8Ð[hÔ
i€Cà‡w�w”%—.‘.Ô Ø˜ÔÜ×1Ñ1‰Bä×1Ñ1‰BØ	
�‰”E—M‘MÔ	!Ø˜ÔÜ×1Ñ1‰Bä×1Ñ1‰BØ	
�‰”E—M‘MÔ	!Ø˜ÔÜ×1Ñ1‰Bä×1Ñ1ˆBä	˜×	Õ	Ù
ØØ�J‰J‹LØ�O‰OÓØ�L‰L‹NØØô	
÷ 
ð ˆ;Ð÷ 
Ö	ð ˆ;Ðús   Æ8GÇG 	zbitsandbytes::dequantize_4bitc                óÒ   € V ^8„  d   QhR\         P                  R\         P                  R\        R\        R\        \        ,          R\         P
                  R\         P                  /# )r*   r+   r¨   r–   r»   r7   r4   rc   ©r    r.   r—   r½   r   r4   )r/   s   "r   r0   r0   ¨  sb   € ÷ 
ñ 
Ü‡|�|ð
ä�L‰Lð
ô ð
ô ð	
ô
 ”C�=ð
ô �;‰;ð
ô ‡\�\ñ
r   c           	      óf   € \         P                  ! WEV P                  R 7      p\        WW#WVR7       V# )©r4   r3   r¬   )r    r6   r3   Ú_dequantize_4bit_impl©r+   r¨   r–   r»   r7   r4   r;   s   &&&&&& r   r   r   §  s)   € ô �+Š+�e°·±Ô
:€CÜ˜! Y¸EÕKØ€Jr   z!bitsandbytes::dequantize_4bit.outc                óÖ   € V ^8„  d   QhR\         P                  R\         P                  R\        R\        R\        \        ,          R\         P
                  R\         P                  RR	/# )
r*   r+   r¨   r–   r»   r7   r4   r;   rc   NrÈ   )r/   s   "r   r0   r0   ¶  sv   € ÷ Lñ LÜ‡|�|ðLä�L‰LðLô ðLô ð	Lô
 ”C�=ðLô �;‰;ðLô 
�‰ðLð 
ñLr   c           	      óä   € VP                   \        V4      8w  d   \        R V RVP                    24      hVP                  V8w  d   \        RV RVP                   24      h\	        WW#WVR7       R# )r´   r³   r²   r¬   N)r7   r~   rD   r4   rË   rÌ   s   &&&&&&&r   r   r   µ  sf   € ð ‡y�y”E˜%“LÔ ÜÐ1°%°¸¸s¿y¹y¸kÐJÓKÐKØ
‡y�y�EÔÜÐ1°%°¸¸s¿y¹y¸kÐJÓKÐKÜ˜! Y¸E×Kr   c                ó´   € V ^8„  d   QhR\         P                  R\         P                  R\        R\        R\         P                  R\         P                  RR/# )	r*   r+   r¨   r–   r»   r4   r;   rc   N)r    r.   r—   r½   r4   )r/   s   "r   r0   r0   Æ  s^   € ÷ %
ñ %
Ü‡|�|ð%
ä�L‰Lð%
ô ð%
ô ð	%
ô
 �;‰;ð%
ô 
�‰ð%
ð 
ñ%
r   c                 óÜ  € V P                  4       p V\        P                  8X  d*   VR 8X  d   \        P                  pM›\        P
                  pMŠV\        P                  8X  d*   VR 8X  d   \        P                  pM]\        P                  pMLV\        P                  8X  d*   VR 8X  d   \        P                  pM\        P                  pM\        RV 24      h\        V 4      ;_uu_ 4        V! RV P                  4       VP                  4       VP                  4       VVP                  4       \!        V P"                  P$                  4      4       RRR4       R#   + '       g   i     R# ; i)r¿   zFBlockwise 4bit dequantization only supports 16/32-bit floats, but got N)r¡   r    r¥   r   Úcdequantize_blockwise_bf16_fp4Úcdequantize_blockwise_bf16_nf4rh   Úcdequantize_blockwise_fp16_fp4Úcdequantize_blockwise_fp16_nf4rf   Úcdequantize_blockwise_fp32_fp4Úcdequantize_blockwise_fp32_nf4rD   r   rM   ri   rN   r3   rO   )r+   r¨   r–   r»   r4   r;   r   s   &&&&&& r   rË   rË   Æ  sÿ   € ð 	
�‰‹€Aà”—‘ÔØ˜ÔÜ×3Ñ3‰Bä×3Ñ3‰BØ	”%—-‘-Ô	Ø˜ÔÜ×3Ñ3‰Bä×3Ñ3‰BØ	”%—-‘-Ô	Ø˜ÔÜ×3Ñ3‰Bä×3Ñ3‰BäÐaÐbgÐahÐiÓjÐjä	˜×	Õ	Ù
ØØ�J‰J‹LØ�O‰OÓØ�L‰L‹NØØ�I‰I‹KÜ˜AŸH™HŸN™NÓ+ô	
÷ 
×	×	Ò	ús   Ã,A$EÅE+	zbitsandbytes::gemv_4bitc                óæ   € V ^8„  d   QhR\         P                  R\         P                  R\        \        ,          R\         P                  R\         P                  R\        R\         P                  /# )r*   r+   r,   rT   r¨   r•   r–   rc   ©r    r.   r   r—   )r/   s   "r   r0   r0   ï  s^   € ÷ ñ Ü‡|�|ðÜŸ™ðÜ.6´s­mðÜEJÇ\Á\ðÜY^×YeÑYeðÜruðä
‡\�\ñr   c           
      ó²   € . V P                   R R OV^ ,          N5p\        P                  ! W`P                  V P                  R7      p\        WW#WEVR7       V# )Nr2   r¬   r5   )r7   r    r6   r3   r4   Ú_gemv_4bit_impl)r+   r,   rT   r¨   r•   r–   r7   r;   s   &&&&&&  r   r   r   î  sL   € ð 'ˆa�g‰g�c�rˆlÐ&˜F 1�IÑ&€EÜ
�+Š+�e§H¡H°A·G±GÔ
<€CÜ�A˜&¨$¸sÕCØ€Jr   zbitsandbytes::gemv_4bit.outc                óê   € V ^8„  d   QhR\         P                  R\         P                  R\        \        ,          R\         P                  R\         P                  R\        R\         P                  RR	/# ©
r*   r+   r,   rT   r¨   r•   r–   r;   rc   NrØ   )r/   s   "r   r0   r0   ù  sz   € ÷ Dñ DÜ‡|�|ðDä‡|�|ðDô ”S�MðDô �L‰Lð	Dô
 �,‰,ðDô ðDô 
�‰ðDð 
ñDr   c           
      ó2  € . V P                   R R OV^ ,          N5pVP                   V8w  d   \        RV RVP                    24      hVP                  V P                  8w  d&   \        RV P                   RVP                   24      h\        WW#WEVR7       R # )Nr´   r³   r²   r¬   r5   )r7   rD   r4   rÚ   )r+   r,   rT   r¨   r•   r–   r;   Úexpected_shapes   &&&&&&& r   r   r   ø  sŠ   € ð 0�q—w‘w˜s �|Ð/ V¨A¥YÑ/€NØ
‡y�y�NÔ"ÜÐ1°.Ð1AÀÈÏ	É	À{ÐSÓTÐTØ
‡y�y�A—G‘GÔÜÐ1°!·'±'°¸&ÀÇÁÀÐLÓMÐMÜ�A˜&¨$¸s×Cr   c                óê   € V ^8„  d   QhR\         P                  R\         P                  R\        \        ,          R\         P                  R\         P                  R\        R\         P                  RR	/# rÜ   rØ   )r/   s   "r   r0   r0   
  sp   € ÷ :
ñ :
Ü‡|�|ð:
ä‡|�|ð:
ô ”S�Mð:
ô �L‰Lð	:
ô
 �,‰,ð:
ô ð:
ô 
�‰ð:
ð 
ñ:
r   c                 ó  € VR9  d   \        RV 24      hV^ ,          p^pV^,          p	Tp
V P                  R,          ^,           ^,          pTpV P                  \        P                  8X  d   \
        P                  pM_V P                  \        P                  8X  d   \
        P                  pM/V P                  \        P                  8X  d   \
        P                  p\        V 4      ;_uu_ 4        X! VVV	V P                  4       VP                  4       VP                  4       VP                  4       VP                  4       V
VVV\        V P                  P                  4      4       RRR4       R#   + '       g   i     R# ; i)é    r™   N)rá   rš   r›   rœ   r�   rž   rŸ   r    r5   )rD   r7   r4   r    rh   r   Úcgemm_4bit_inference_naive_fp16r¥   Úcgemm_4bit_inference_naive_bf16rf   Úcgemm_4bit_inference_naive_fp32r   rM   rN   r3   rO   )r+   r,   rT   r¨   r•   r–   r;   rW   rX   rV   rY   rZ   r[   r   s   &&&&&&&       r   rÚ   rÚ   
  s  € ð ÐAÔAÜÐ-¨i¨[Ð9Ó:Ð:ð$ 	ˆq�	€AØ	€AØˆq�	€Aà
€CØ�7‰7�2�;˜�?˜qÕ
 €CØ
€Cà‡w�w”%—-‘-ÔÜ×0Ñ0‰Ø	
�‰”E—N‘NÔ	"Ü×0Ñ0‰Ø	
�‰”E—M‘MÔ	!Ü×0Ñ0ˆä	˜×	Õ	Ù
ØØØØ�J‰J‹LØ�J‰J‹LØ�O‰OÓØ�M‰M‹OØ�L‰L‹NØØØØÜ˜AŸH™HŸN™NÓ+ô	
÷ 
×	×	Ò	ús   Ã1A8E3Å3F	c                ó$  € V\         8:  d   R# \        V 4      w  rVpV^?,           ^@,          pV\        P                  8X  g   V^8  d   V^8  # V^8X  d"   V^8  g   V\        P                  8w  d   V^8  # V^8X  d   V^8X  g   V^8X  d   R# V^
8X  d(   W…^,          8¼  d   V^ 8*  # W…8¼  d   WC8¼  d   R# V^8*  # R# V^8H  ;'       d    V^8H  p	V^8H  ;'       d    V^ 8H  p
V^8H  ;'       d    V^8H  pV^	8H  pV^8H  ;'       d    V^ 8H  pV^8H  ;'       d    V^8H  pT
;'       g    TpWC^,          8„  pW4,          R
8  Ed8   V^,          V8  dŸ   V'       d7   V^,          V8  pV
'       d   TV'       d   R8*  # R8*  # TV'       d   R8*  # R8*  # V	'       d0   W…^,          8¼  d   VR8*  # VR8¼  d   V^@8*  # VR8¼  d   V^`8*  # VR8*  # W…8¼  d   VR8*  # W„,          VR,          8”  d   V^À8*  # VR8*  # V^,          V8*  pV
'       d*   V'       d"   WC8¼  d   Y(^,          V8:  d   ^ 8*  # ^€8*  # VR8*  # V	'       d   W…8  d   V^8”  d   V^@8*  # V'       d   Y$V8  g	   V'       d   ^ 8*  # ^8*  # V^,          V8:  d   V^8*  # R# VR8”  d   R# V^8  d�   V'       d   ^$M^0pV'       dw   V^8¼  dp   VV8¼  di   V^,          V^,          8  pV^,          V^,          8  pWC8”  d   V'       d   R# WC8X  d-   V
'       d   V'       d   R# V'       d   V'       d   V'       g   R# R# V	'       d4   W…^,          8¼  d   Y$V8  d   ^€8*  # ^@8*  # W…^,          8¼  d   V^@8*  # V^ 8*  # V
'       d¯   W…^,          8¼  d   V^€8*  # W…8¼  d   Y$V8  d   ^@8*  # ^ 8*  # WC^,          8¼  d   V^,          V8”  d   R# WC8X  d   V^08¼  d   V^,          V^,          8  d   R# V^,          V8  d   WC8:  d   V^€8*  # V^,          V8¼  d   R# W…^,          8¼  d   WC8  d   V^€8*  # V^8*  # V'       d#   W…8¼  d   V^€8*  # W…^,          8¼  d   V^@8*  # V^8*  # V'       da   W…^,          8¼  d   V^@8*  # W…^,          8¼  d   V^08*  # W…8¼  d   V^ 8*  # W…^,          8¼  d#   WC8¼  d   V^,          V^,          8  d   R# V^8*  # R# V'       d   W…8¼  d   VR	8*  # V'       d—   W…^,          8¼  d   VR	8*  # W…^,          8¼  d   V^€8*  # V^,          V^,          8¼  d   Y$V8¼  d   ^`8*  # ^@8*  # W…8¼  d   V^@8*  # W…^,          8¼  d   Y$^,          V8¼  d   V^€8¼  d   ^@8*  # ^8*  # V'       d   V^@8”  d   V^8*  # V^8*  # W…^,          8¼  d   VR	8*  # W…^,          8¼  d   V^€8*  # V^,          V^,          8¼  d/   V'       g   V^È8¼  d   V^€8*  # WC8¼  d   Y%^<8¼  d   ^€8*  # ^ 8*  # V^@8*  # W…^,          8¼  d   V'       d   V^@8*  # V^@8¼  d   V^8*  # V^ 8*  # WC8”  d   V^08:  d   Y$V^,          8¼  d   ^8*  # ^
8*  # TV'       g   V^08  d   ^8*  # ^8*  # )aÙ  Custom kernel vs dequant+F.linear heuristic for M in [5, 1536].

Per-arch notes (bf16/fp16, M >= 8, large weight):
  sm75 (T4, ~300 GB/s GDDR6):    fp16 MMA only; GDDR makes dequant expensive.
  sm80 (A100, ~2 TB/s HBM2e):    mma.sync; HBM thresholds; K-heavy shapes handled explicitly.
  sm86 (A10, ~600 GB/s GDDR6):   dedicated block; wider M caps than sm89 at medium N.
  sm89 (4090, L40S, GDDR6X):     default fallback; tall-K and large-N get higher M caps.
  sm90 (H100/H200, HBM3/HBM3e):  dequant+linear is much faster; thresholds are tight.
  sm100 (B200/B300, HBM3e):       exits early at top of function.
  sm120 (RTX 5000, GDDR7):        dedicated block; medium-N tiers differ from sm89.
  sm121 (GB10 DGX Spark, LPDDR5X):  dedicated block at >=1 wave only; the low
                                 memory bandwidth keeps the custom kernel ahead
                                 through M=256 there. Sub-wave shapes use the
                                 sm89 tiers below.
TFrž   é   r�   i@  iÀ  i€  rœ   i  @ )Ú_GEMM_4BIT_CUSTOM_FLOOR_Mr(   r    rf   rh   )r&   r4   ÚMÚNÚKÚnum_smsr$   r%   Ún_blocksÚis_sm75Úis_sm80Úis_sm86Úis_sm90Úis_sm120Úis_sm121Úis_hbmÚ
tall_k_2xnÚlow_waveÚquarter_waveÚhbm_m67_threshÚlt_75pct_waveÚlt_60pct_waves   &&&&&                 r   Ú_gemm_4bit_use_custom_cudarú   G  s4  € ð" 	Ô%Ô%Ùä/°Ó=Ñ€G�EØ�B•˜2�~€Hð ”—‘Ô ¨¤Ø�1‰uˆØ�„z�u˜q”y E¬U¯]©]Ô$:Ø�1‰uˆð 	�Œ
�u ”z e¨r¤kÙð �„{Ø •{Ô"Ø˜‘7ˆNØÔØœF�5Ð.¨¨Q©Ð.Ùà�q‰j×'Ð'˜U a™Z€GØ�q‰j×'Ð'˜U a™Z€GØ�q‰j×'Ð'˜U a™Z€GØ�q‰j€GØ˜‰{×)Ð)˜u¨™z€HØ˜‰{×)Ð)˜u¨™z€HØ×Ð˜€FØ˜•U‘€Jð 	…uˆÕØˆq�5�1Œ9çð $ a�<¨'Ñ1�ßØ¯ Ñ;Ð;°sÑ;Ð;Ø§H˜SÑ6Ð6°#Ñ6Ð6ßà¨�{Ô*Ø ™8�OØ˜”9Ø ™7�NØ˜”8Ø ™7�NØ˜C‘x�àÔ"Ø˜C‘x�Ø'�|¨g¸­mÔ;�1˜‘8ÐIÀÀcÁÐIà !•| wÑ.ˆß—|àŒvØ¨a¥<°7Ô#:˜RÑDÐDÀÑDÐDØ˜‘8ˆOç�xÔ)¨a°!¬eØ˜‘7ˆN÷ Ø A¤¯˜Ñ9Ð9°bÑ9Ð9Ø�a�<˜7Ô"Ø˜‘7ˆNÙð 	ˆ3„wÙð 	ˆ1„uß&™¨Bˆß�a˜1”f ¨^Ô!;Ø$ q�L¨7°Q­;Ñ6ˆMØ$ q�L¨7°Q­;Ñ6ˆMàŒuŸÙð ŒvßŸ}Ù ßŸ}·]Ù Ù÷ à •{Ô"Ø A¤˜Ñ.Ð.¨2Ñ.Ð.Ø !•|Ô#Ø˜‘7ˆNØ�B‰wˆçà •{Ô"Ø˜‘8ˆOØÔØ 1œu˜Ñ-Ð-¨"Ñ-Ð-ð �A•Œ:˜( Q�,¨Ô0Ùð Œ6�h "”n¨°A­¸À!½Ô)CÙà�a�<˜'Ô!ØŒvØ˜C‘x�Ø˜!�|˜wÔ&Ùà !•|Ô#¨¬Ø˜‘8ˆOØ�B‰wˆçð ÔØ˜‘8ˆOØ !•|Ô#Ø˜‘7ˆNØ�B‰wˆçð  •{Ô"Ø˜‘7ˆNØ •{Ô"Ø˜‘7ˆNØÔØ˜‘7ˆNØ !•|Ô#àŒv˜( Q�,¨°1­Ô4ÙØ˜‘7ˆNÙçð ÔØ˜‘8ˆOçà •{Ô"Ø˜‘8ˆOØ •{Ô"Ø˜‘8ˆOà�a�<˜7 Q�;Ô&Ø Aœv˜Ñ.Ð.¨2Ñ.Ð.ØÔØ˜‘7ˆNØ !•|Ô#à A¥¨¤
¨x¸3¬˜ÑGÐGÀQÑGÐGß˜( Rœ-Ø˜‘7ˆNØ�A‰vˆð ˜Q•;ÔØ�C‰xˆØ˜Q•;ÔØ�C‰xˆð �!…|�w •{Ô"ß˜ SœØ˜‘8ˆOàŒ6Ø¨2¤˜Ñ6Ð6°2Ñ6Ð6Ø�B‰wˆØ˜a•<ÔßØ˜‘7ˆNØ�rŒ>Ø˜‘6ˆMØ�B‰wˆð 	„u�˜R”Ø  A¥œ:�RÑ.Ð.¨2Ñ.Ð.ØŸ
 h°¤m�Ñ<Ð<¸!Ñ<Ð<r   c                ó4  € V\         8:  d   V\        P                  8w  d   R# \        V 4      pVP	                  R4      '       g   VP	                  R4      '       d   V^8*  # VP	                  R4      '       d   Y!\        P                  8X  d   ^8*  # ^8*  # V^8*  # )a%  
Fused SIMT kernel vs dequant+F.linear heuristic for ROCm.

RDNA3/RDNA4 calibration keeps the SIMT kernel through ~M=8.
gfx125x (CDNA5) is Wave32 but lacks the v_dot2 SIMT
instructions, so its kernel takes the fp32-FMA CDNA math path. The crossover
below is left at the gfx12 default (M<=8) pending dedicated calibration.
CDNA/gfx9 is calibrated on MI308X (gfx942): bf16/fp16 win through M<=4
after the SIMT math-path tuning, while fp32 only has a broad win through M<=2.

TODO: revisit once WMMA/MFMA kernels land; calibrate the gfx125x crossover.
TÚgfx11Úgfx12Úgfx9)rç   r    rf   Ú_rocm_gfx_archÚ
startswith)r&   r4   rè   ré   rê   Úarchs   &&&&& r   Ú_gemm_4bit_use_custom_rocmr  .  s   € ð 	Ô%Ô%¨%´5·=±=Ô*@Ùä˜,Ó'€DØ‡��w×Ò 4§?¡?°7×#;Ò#;Ø�A‰vˆØ‡��v×ÒØ¤5§=¡=Ô0�QÑ8Ð8°aÑ8Ð8Ø�‰6€Mr   c                óš   € \        \        P                  P                  V 4      RR4      ;'       g    RpVP	                  R4      ^ ,          # )zKgfx arch string (e.g. 'gfx1100') for a ROCm device, feature flags stripped.ÚgcnArchNameÚ Ú:)r   r    r!   r"   Úsplit)r&   r   s   & r   rÿ   rÿ   G  s>   € ô ”5—:‘:×3Ñ3°LÓAÀ=ÐRTÓU×[Ð[ÐY[€DØ�:‰:�c‹?˜1ÕÐr   c
                ó,  € V P                   R,          p
V P                  4       V
,          pV^ ,          pW¢^,          8w  d   \        RV
 RV^,           R24      hVP                  \        P
                  8w  d   \        RVP                   24      hVel   VP                  ^8w  d   \        RVP                   R24      hVP                  V P                  8w  d'   \        RVP                   R	V P                   R24      hV P                  4       p VR
8X  d   ^M^p\        P                  ! . V P                   RR OVN5V P                  V P                  R7      p\        V P                  P                  4      pV P                  \        P                  8X  d   \        P                  pMxV P                  \        P                  8X  d   \        P                   pMHV P                  \        P
                  8X  d   \        P"                  pM\        RV P                   24      hV	e!   V	P%                  \        P
                  R7      MRp\'        V 4      ;_uu_ 4        T! V P)                  4       VP)                  4       VP)                  4       Ve   VP)                  4       MRVe   VP)                  4       MRVe   VP)                  4       MRVP)                  4       Ve   VP)                  4       MRVVV
VVV4       RRR4       V#   + '       g   i     T# ; i)z×Invoke the fused cgemm_4bit_* kernel (shared by the CUDA and ROCm dispatch; the
C dispatch in gemm_4bit.cu picks SIMT vs MMA per arch/shape). A is made contiguous
because the kernel reads it as row-major (stride K).zA inner dim (z) does not match weight (Ú)zabsmax must be float32, got Nzbias must be 1D, got ÚDzbias dtype (z) must match A dtype (r¿   rÊ   zunsupported dtype re   r5   )r7   ri   rR   r4   r    rf   rE   r¡   r6   r3   rN   rO   r¥   r   Úcgemm_4bit_bf16rh   Úcgemm_4bit_fp16Úcgemm_4bit_fp32rI   r   rM   )r+   r,   rT   r¨   r–   r»   rb   Úabsmax_8bitÚabsmax_codeÚabsmax_offsetrê   rè   ré   Úquant_type_intr;   Ústreamr   Úabsmax_offset_f32s   &&&&&&&&&&        r   Ú_gemm_4bit_kernel_implr  N  sp  € ð 	
�‰��€AØ	�‰‹	�Q�€AØˆq�	€Aà�1�I„~Ü˜]¨1¨#Ð-FÀvÈaÅyÀkÐQRÐSÓTÐTØ‡|�|”u—}‘}Ô$ÜÐ9¸&¿,¹,¸ÐHÓIÐIØÒØ�9‰9˜Œ>ÜÐ!6°t·y±y°kÀÐCÓDÐDØ�:‰:˜Ÿ™Ô Ü ¨d¯j©j¨\Ð9OÐPQ×PWÑPWÈyÐXYÐZÓ[Ð[à	�‰‹€AØ$¨Ô-‘Q°1€NÜ
�+Š+Ð(˜Ÿ™  ˜Ð( aÑ(°·±ÀÇÁÔ
I€CÜ˜QŸX™XŸ^™^Ó,€Fà‡w�w”%—.‘.Ô Ü× Ñ ‰Ø	
�‰”E—M‘MÔ	!Ü× Ñ ‰Ø	
�‰”E—M‘MÔ	!Ü× Ñ ‰äÐ/°·±¨yÐ9Ó:Ð:ð BOÒAZ˜×(Ñ(¬u¯}©}Ð(Ô=Ð`dÐä	˜×	Õ	Ù
Ø�J‰J‹LØ�J‰J‹LØ�O‰OÓØ&1Ò&=ˆK× Ñ Ô"À4Ø&1Ò&=ˆK× Ñ Ô"À4Ø,=Ò,IÐ×&Ñ&Ô(ÈtØ�L‰L‹NØ#Ò/ˆD�M‰MŒO°TØØØØØØô	
÷ 
ð$ €J÷% 
Ö	ð$ €Jús   ÉBLÌL	c
           	     ó€  € VeM   \         P                  ! V\         P                  R7      p
\        WsVR\         P                  V
R7       W©,           p\         P                  ! W P
                  V P                  R7      p\        WWEV P
                  VR7       \         P                  P                  P                  WV4      # )z¾Unfused fallback shared by CUDA and ROCm: reconstruct the (optionally nested)
absmax, dequantize the 4-bit weight via the backend dequant impls (reusing
preallocated buffers), then F.linear.re   rœ   r¬   rÊ   )r    rg   rf   r­   r6   r4   r3   rË   ÚnnÚ
functionalÚlinear)r+   r,   rT   r¨   r–   r»   rb   r  r  r  Ú	absmax_dqÚB_dqs   &&&&&&&&&&  r   Ú_dequant_linear_fallbackr  ˆ  s„   € ð ÒÜ×$Ò$ [¼¿¹ÔFˆ	Ü" ;¸ÀSÌ%Ï-É-Ð]fÕgØÕ*ˆÜ�;Š;�v§W¡W°Q·X±XÔ>€DÜ˜! Y¸A¿G¹GÈÕNÜ�8‰8×Ñ×%Ñ% a¨tÓ4Ð4r   ræ   rœ   zbitsandbytes::gemm_4bitc                óª  € V ^8„  d   QhR\         P                  R\         P                  R\        \        ,          R\         P                  R\        R\        R\
        \         P                  ,          R\
        \         P                  ,          R	\
        \         P                  ,          R
\
        \         P                  ,          R\         P                  /# )r*   r+   r,   rT   r¨   r–   r»   rb   r  r  r  rc   )r    r.   r   r—   r½   r   )r/   s   "r   r0   r0   §  s®   € ÷ /ñ /Ü‡|�|ð/ä‡|�|ð/ô ”S�Mð/ô �L‰Lð	/ô
 ð/ô ð/ô ”5—<‘<Õ
 ð/ô œ%Ÿ,™,Õ'ð/ô œ%Ÿ,™,Õ'ð/ô œEŸL™LÕ)ð/ô ‡\�\ñ/r   c
                 ó‚  € V P                   R,          p
V P                  4       V
,          pV^ ,          pV\        8”  d   RpMSW¤,          ^ 8w  d   \        RV
 RV R2\        4       RpM,\        V P                  P                  V P                  W¼V
4      pV'       g   \        V VVVVVVVVV	R7
      # \        WW#WEWgW‰4
      # )é   Fzinner dimension (z0) is not aligned for fast kernel with blocksize=z(, falling back to slower implementation.)r  r  r  r5   )r7   ri   Ú_gemm_4bit_custom_max_mr   ÚUserWarningÚ_gemm_4bit_use_custom_fnr3   rO   r4   r  r  )r+   r,   rT   r¨   r–   r»   rb   r  r  r  rê   rè   ré   Ú
use_customs   &&&&&&&&&&    r   r   r   ¦  sÓ   € ð 	
�‰��€AØ	�‰‹	�Q�€AØˆq�	€Að
 	Ô"Ô"Ø‰
Ø	
�˜!Ô	ÜØ ˜sð #Ø'˜[Ð(PðRäô	
ð
 ‰
ä-¨a¯h©h¯n©n¸a¿g¹gÀqÈQÓOˆ
çÜ'ØØØØØØØØ#Ø#Ø'ô
ð 	
ô "Ø	ˆf˜i°TÈóð r   ÚadamÚmomentumÚrmspropÚlionÚadagradÚlambÚademamixÚlarsc          $      óx  € V ^8„  d   QhR\         R\        P                  R\        P                  R\        P                  R\        \        P                  ,          R\        \        P                  ,          R\        R\        R	\        R
\        R\        R\        R\        R\        R\
        R\        R\        RR/# )r*   Úoptimizer_nameÚgÚpÚstate1Ústate2Ú	unorm_vecÚ	max_unormÚ
param_normÚbeta1Úbeta2Úbeta3ÚalphaÚepsÚweight_decayÚstepÚlrÚgnorm_scalerc   N©r½   r    r.   r   rG   r—   )r/   s   "r   r0   r0   #  sÚ   € ÷ 8
ñ 8
Üð8
ä‡|�|ð8
ô ‡|�|ð8
ô �L‰Lð	8
ô
 ”U—\‘\Õ"ð8
ô œŸ™Õ%ð8
ô ð8
ô ð8
ô ð8
ô ð8
ô ð8
ô ð8
ô 
ð8
ô ð8
ô ð8
ô  	ð!8
ô" ð#8
ð& 
ñ'8
r   c                 ó  € \         P                  V R 4      pVf-   \        RV  R\        \         P	                  4       4       24      hVP
                  \        P                  8X  d   V^ ,          pM‡VP
                  \        P                  8X  d   V^,          pM^VP
                  \        P                  8X  d   \        V4      ^8X  d   V^,          pM%\        RVP
                   RVP
                   24      h\        V4      ;_uu_ 4        V! \        V4      \        V4      \        V4      \        V4      \        V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V	4      \        P                  ! V
4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                   ! V4      \        P                  ! VP#                  4       4      4       R R R 4       R #   + '       g   i     R # ; i)NúUnsupported optimizer name: ú. Supported optimizers: zAGradient+optimizer bit data type combination not supported: grad z, optimizer )Ústr2optimizer32bitÚgetrD   ÚlistÚkeysr4   r    rf   rh   r¥   Úlenr   r	   ÚctÚc_floatÚc_int32Úc_boolri   )r,  r-  r.  r/  r0  r1  r2  r3  r4  r5  r6  r7  r8  r9  r:  r;  r<  Ú
skip_zerosÚ	optim_fnsÚ
optim_funcs   &&&&&&&&&&&&&&&&&&  r   Ú_optimizer_update_32bit_implrM  #  s¹  € ô( #×&Ñ& ~°tÓ<€IØÒÜØ*¨>Ð*:Ð:RÔSWÔXj×XoÑXoÓXqÓSrÐRsÐtó
ð 	
ð 	‡w�w”%—-‘-ÔØ˜q•\‰
Ø	
�‰”E—M‘MÔ	!Ø˜q•\‰
Ø	
�‰”E—N‘NÔ	"¤s¨9£~¸Ô':Ø˜q•\‰
äØOÐPQ×PWÑPWÈyÐXdÐek×eqÑeqÐdrÐsó
ð 	
ô 
˜×	Õ	ÙÜ�A‹JÜ�A‹JÜ�F‹OÜ�F‹OÜ�IÓÜ�JŠJ�yÓ!Ü�JŠJ�zÓ"Ü�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�s‹OÜ�JŠJ�|Ó$Ü�JŠJ�tÓÜ�JŠJ�r‹NÜ�JŠJ�{Ó#Ü�IŠI�jÓ!Ü�JŠJ�q—w‘w“yÓ!ô%	
÷ 
×	×	Ò	ús   Ä
EI-É-I>	c          &      óÖ  € V ^8„  d   QhR\         R\        P                  R\        P                  R\        P                  R\        \        P                  ,          R\        R\        R\        R	\        R
\        R\
        R\        R\        P                  R\        \        P                  ,          R\        P                  R\        \        P                  ,          R\        R\        RR/# )r*   r,  r-  r.  r/  r0  r4  r5  r6  r7  r8  r:  r;  Úqmap1Úqmap2Úabsmax1Úabsmax2r9  r<  rc   Nr=  )r/   s   "r   r0   r0   ^  s
  € ÷ Z
ñ Z
ÜðZ
ä‡|�|ðZ
ô ‡|�|ðZ
ô �L‰Lð	Z
ô
 ”U—\‘\Õ"ðZ
ô ðZ
ô ðZ
ô ðZ
ô ðZ
ô 
ðZ
ô ðZ
ô 	ðZ
ô �<‰<ðZ
ô ”E—L‘LÕ!ðZ
ô �\‰\ðZ
ô  ”e—l‘lÕ#ð!Z
ô" ð#Z
ô$ ð%Z
ð( 
ñ)Z
r   c                 ó°  € \         P                  V 4      pVf-   \        RV  R\        \         P	                  4       4       24      hVP
                  \        P                  8X  d   V^ ,          pMkVP
                  \        P                  8X  d   V^,          pMBVP
                  \        P                  8X  d   V^,          pM\        RVP
                   R24      h\        V4      ;_uu_ 4        V! \        V4      \        V4      \        V4      \        V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V	4      \        P                  ! V
4      \        P                  ! V4      \        V4      \        V4      \        V4      \        V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! V4      \        P                  ! VP!                  4       4      4       R R R 4       R #   + '       g   i     R # ; i)Nr?  r@  zUnsupported gradient dtype: z@. Supported dtypes: torch.float32, torch.float16, torch.bfloat16)Ústr2optimizer8bit_blockwiserB  rD   rC  rD  r4   r    rf   rh   r¥   r   r	   rF  rG  rH  rI  ri   )r,  r-  r.  r/  r0  r4  r5  r6  r7  r8  r:  r;  rO  rP  rQ  rR  r9  r<  rJ  Úoptimizer_fnsÚoptimizer_fns   &&&&&&&&&&&&&&&&&&&  r   Ú%_optimizer_update_8bit_blockwise_implrW  ^  sŸ  € ôh 0×3Ñ3°NÓC€MØÒÜØ*¨>Ð*:Ð:RÔSWÔXs×XxÑXxÓXzÓS{ÐR|Ð}ó
ð 	
ð 	‡w�w”%—-‘-ÔØ$ QÕ'‰Ø	
�‰”E—M‘MÔ	!Ø$ QÕ'‰Ø	
�‰”E—N‘NÔ	"Ø$ QÕ'‰äØ*¨1¯7©7¨)Ð3sÐtó
ð 	
ô 
˜×	Õ	ÙÜ�A‹JÜ�A‹JÜ�F‹OÜ�F‹OÜ�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�uÓÜ�JŠJ�s‹OÜ�JŠJ�tÓÜ�JŠJ�r‹NÜ�E‹NÜ�E‹NÜ�GÓÜ�GÓÜ�JŠJ�|Ó$Ü�JŠJ�{Ó#Ü�IŠI�jÓ!Ü�JŠJ�q—w‘w“yÓ!ô'	
÷ 
×	×	Ò	ús   Ã-EIÉI	z-bitsandbytes::optimizer_update_8bit_blockwisez$bitsandbytes::optimizer_update_32bitr   )Úfp32Úbf16Úfp16)Únf4r¿   )rY  rZ  rX  )NN)rn   )NNNN)F)RÚcollections.abcr   ÚctypesrF  Ú	functoolsÚmathr   Útypingr   rP   r   r    Úbitsandbytes.functionalr   r   r	   Ú_opsr
   Ú
cextensionr   r   Úc_void_prH  rG  Ú_CÚ_cuda_getCurrentRawStreamrN   Úcacher(   r   r9   r„   r­   rË   rÚ   rú   r  rÿ   r  r  rç   ÚversionÚhipr!  r  Úcadam32bit_grad_fp32Úcadam32bit_grad_fp16Úcadam32bit_grad_bf16Úcmomentum32bit_grad_32Úcmomentum32bit_grad_16Úcrmsprop32bit_grad_32Úcrmsprop32bit_grad_16Úclion32bit_grad_fp32Úclion32bit_grad_fp16Úclion32bit_grad_bf16Úcadagrad32bit_grad_32Úcadagrad32bit_grad_16Úcademamix32bit_grad_fp32Úcademamix32bit_grad_fp16Úcademamix32bit_grad_bf16rA  Úcadam_8bit_blockwise_grad_fp32Úcadam_8bit_blockwise_grad_fp16Úcadam_8bit_blockwise_grad_bf16Ú"cmomentum_8bit_blockwise_grad_fp32Ú"cmomentum_8bit_blockwise_grad_fp16Ú"cmomentum_8bit_blockwise_grad_bf16Ú!crmsprop_8bit_blockwise_grad_fp32Ú!crmsprop_8bit_blockwise_grad_fp16Ú!crmsprop_8bit_blockwise_grad_bf16Úclion_8bit_blockwise_grad_fp32Úclion_8bit_blockwise_grad_fp16Úclion_8bit_blockwise_grad_bf16Ú!cadagrad_8bit_blockwise_grad_fp32Ú!cadagrad_8bit_blockwise_grad_fp16Ú!cadagrad_8bit_blockwise_grad_bf16Ú"cademamix_8bit_blockwise_grad_fp32Ú"cademamix_8bit_blockwise_grad_fp16Ú"cademamix_8bit_blockwise_grad_bf16rT  rM  rW  )ÚdÚqs   00r   Ú<module>r�     sd  ðÝ $Û Û Ý Ý Ý ã ç LÑ Lå #Ý ôñ Ù/GÔ`Ñ/G¨!ÓQ_ÈAÐ˜a˜S  ! Ó%ÑQ_Ñ%Ñ/GÒ`Ù-EÓFÑ-E¨Ð	 ˜sÓ#Ñ-EÑFõGà‡[�[€M�AÕ˜Ÿ™ R§Z¡Z°·±Ð=Õ=ôñ Ù 8Ó9Ñ 8˜1€{�1�#ÓÑ 8Ñ9Ø‡[�[€M�AÕ˜Ÿ™ R§Z¡Z°·±¸R¿Z¹ZÈÏÉÐUW×U`ÑU`ÐaÕaôñ Ù0HÓIÑ0H¨1Ð" 1 #Ó&Ñ0HÑIØ‡Z�Z€L�1Õ˜Ÿ™�} qÕ(Õ(¨B¯J©J¨<¸!Õ+;Õ;¸r¿z¹zÈ2Ï;É;Ð>WÕWôñ Ø€OØ‡[�[€M�R—Z‘Z�L 1Õ$Õ$¨¯© }°qÕ'8Õ8¸B¿J¹J¸<È!Õ;KÕKÈrÏ{É{ÈmÕ[Ø�J‰Jõñ ØÐØ‡[�[€M�AÕ˜Ÿ™ R§Z¡Z°·±Ð=Õ=ôñ ØÐØ‡[�[€M�AÕ˜Ÿ™ R§Z¡Z°·±¸R¿[¹[ÐIÕIôñ Ù-EÔ^Ñ-E¨Ë~È!Ð˜A˜3˜a ˜sÓ#É~Ñ#Ñ-EÒ^Ù+CÓDÑ+C aÐ	˜a˜SÓ!Ñ+CÑDõEà‡[�[€M�AÕ˜Ÿ™ R§Z¡ZÐ0Õ0ôð —(‘(×4Ñ4€ð ‡�ñAó ðAñ
 Ð3°VÓ<ô/ó =ð/ñ
 Ð7¸Ó@ô(ó Að(õOñd Ð0°&Ó9ö&*ó :ð&*ñR Ð6¸Ó?ö(,ó @ð(,ñV Ð2°FÓ;öVó <ðV÷&#ñ, Ð3°VÓ<ô ó =ð ñF Ð5°vÓ>ôó ?ðñ Ð9¸6ÓBôKó CðKõ
ñ4 Ð.°Ó7ô#ó 8ð#ñL Ð0°&Ó9ô
ó :ð
ñ Ð4°fÓ=ôLó >ðLõ %
ñP Ð*¨FÓ3ôó 4ðñ Ð.°Ó7ôDó 8ðDõ":
ðz ‡�ñc=ó ðc=ðL ‡�ñó ðð0 ‡�ñó ðô7ôt5ð" Ð Ø‡=�=×ÑÒØ9Ðð #Ñà9Ðð "Ðñ Ð*¨FÓ3ö/ó 4ð/ðd !à
Ø× Ñ Ø× Ñ Ø× Ñ ðð
 Ø×"Ñ"Ø×"Ñ"ðð Ø×!Ñ!Ø×!Ñ!ðð Ø× Ñ Ø× Ñ Ø× Ñ ðð
 Ø×!Ñ!Ø×!Ñ!ðð Ø× Ñ Ø× Ñ Ø× Ñ ðð
 Ø×$Ñ$Ø×$Ñ$Ø×$Ñ$ðð
 Ø×"Ñ"Ø×"Ñ"ððC%Ð ðP Ø×*Ñ*Ø×*Ñ*Ø×*Ñ*ðð
 Ø×.Ñ.Ø×.Ñ.Ø×.Ñ.ðð
 Ø×-Ñ-Ø×-Ñ-Ø×-Ñ-ðð
 Ø×*Ñ*Ø×*Ñ*Ø×*Ñ*ðð
 Ø×-Ñ-Ø×-Ñ-Ø×-Ñ-ðð
 Ø×.Ñ.Ø×.Ñ.Ø×.Ñ.ðð5Ð ÷D8
÷vZ
ñz Ð?ÀÔ HÐInÔ oÙ Ð6¸Ô ?Ð@\Ö ]ùóG% aùÚFùò :ùò Jùó2 _ùÚDs$   Á[
Á([
Ã[#Ä;[(Ë
[-
Ë+[3
