
    oi8              
       P   % S SK r S SKrS SKrSr\\S'   \R                  rS SKrS SKJ	r	  SSK
JrJrJrJrJrJr  SSKJrJr  S SKrS SKr\R*                  rS S	KJr  \S
:X  a$  \" \R2                  5      \" S5      :  a  \" S5      e\" \R2                  5      \" S5      :  aA  \R6                  R8                  R:                  r\R6                  R8                  R>                  r O2\R8                  R;                  SS9r\R8                  R?                  SS9r \S
:X  a2  \R8                  R;                  S
S9r\R8                  R?                  S
S9r S SKrS SK!J"r#  \" \R2                  5      \" S5      :  aP  \S
:X  a+  \#RH                  RJ                  RL                  RN                  r(OS SK)J&r&  \&RN                  r(\#RT                  r+O(\#RX                  RN                  r(\RZ                  S 5       r+\R\                  " S5      S 5       r/S\S\\44S jr0Sr1S SK2r3\" \3R2                  5      \" S5      :  r1\3Rh                  Rj                  r5\S
:X  a  Sr6\S:  a:  \S;   a  \R6                  Rn                  r8O&\S
:X  a  \Rr                  Rn                  r8O	S SK:J;r;  S r8\S
:X  a  \Rx                  Rz                  r>O\Rx                  R~                  r>\R                  r@S\S\@4S jrA\S
:X  a  \B" \5       V Vs0 s HW  n \Rr                  Ro                  U 5      R                  =qD\R                  \Rx                  R{                  [        5      5      _MY     snn rES/\F" \ER                  5       5      S-   -  qHS/\F" \ER                  5       5      S-   -  qIS/\F" \ER                  5       5      S-   -  qJ\ER                  5        H  u  rLrM\M[        \L'   M     \N" [        5      qHCEO\B" \5       V Vs0 s HW  n \R6                  Ro                  U 5      R                  =qD\R                  \Rx                  R                  [        5      5      _MY     snn rOS/\F" \OR                  5       5      S-   -  qPS/\F" \OR                  5       5      S-   -  qIS/\F" \OR                  5       5      S-   -  qJ\OR                  5        H  u  rLrM\M[        \L'   M     \N" [        5      qPCO\R                  rR\R                  rT\3Rh                  R                  R                  rV\3Rh                  R                  R                  rW\3Rh                  R                  R                  rX\S
:X  aA  \3Rh                  R                  R                  rZ\3Rh                  R                  R                  r\O@\3Rh                  R                  R                  rZ\3Rh                  R                  R                  r\\S
:X  a  \Rr                  R                  O\R6                  R                  r^\R                  r`\R                  rb\R                  rd\R                  rf\R                  rh\R                  rj\R                  rl\R                  rn\ R                  R                  S5      b   S S KqJrrr  O\u" S5      rrS# rvS$ rwS% rxS&\R*                  S'\R                  R                  S\R*                  4S( jr{\S
:X  a  \6(       a  \R                  S1S) j5       r}O7\S;   a  \1(       a  \R                  S1S* j5       r}O\R                  S1S+ j5       r} \S
:X  a  \6(       a  S2S, jr~O\S;   a  \1(       a  S2S- jr~OS2S. jr~ S3S/ jrS2S0 jrgs  snn f s  snn f !   S SKsrs\" \sR2                  5      \" S!5      :  a  \t" S"\sR2                   35        \u" S5      rr GN= f)4    Ni   MAX_FUSED_SIZE)Optional   )is_hipget_device_typeDEVICE_TYPEDEVICE_TYPE_TORCHDEVICE_COUNTALLOW_PREQUANTIZED_MODELS   )weight_dequant
fp8_linear)Versionxpuz2.6.0z@Intel xpu currently supports unsloth with torch.version >= 2.6.0z2.4.0cuda)device_typez3.0.0)	libdevicec                 $    U R                  U5      $ N)to)xdtypes     O/home/james-whalen/.local/lib/python3.13/site-packages/unsloth/kernels/utils.pytriton_castr   J   s    ttE{    c                      [        5       =(       a?    [        R                  R                  R                  R                  5       R                  S;   $ )N)gfx940gfx941gfx942)r   tritonruntimedriveractiveget_current_targetarch r   r   is_cdnar'   O   s=    8 --44GGINN S  r   nreturnc                     [        U 5      nU[        :  a  [        SU  S[         S35      eSnUS:  a  SnX4$ US:  a  SnX4$ US	:  a  S
nX4$ )Nz&Cannot launch Triton kernel since n = z& exceeds the maximum CUDA blocksize = .   i       i       i      )next_power_of_2r   RuntimeError)r(   
BLOCK_SIZE	num_warpss      r   calculate_settingsr4   X   s     &a(JN"4QC 8,,:+;1>
 	
 IU	
   	 
t		    
t		  r   Fz0.43.3T)r   hipnullcontextc                     [        5       $ r   r6   )devices    r   torch_gpu_devicer:      s
    }r   tensorc                 R    [        [        U R                  R                  5      5      $ r   )c_void_p_gpu_getCurrentRawStreamr9   index)r;   s    r   _get_tensor_streamr@      s    ,V]]-@-@ABBr   torchao)Float8Tensorz0.15.0zKUnsloth: `from torchao.quantization import Float8Tensor` failed on version=c                     [        U SS 5      $ )Nquant_stategetattr)Ws    r   QUANT_STATErH      s    1mT**r   c                    [        U SU 5      nUR                  n[        US5      (       a  [        USS5      nUb  U" U5      n[        USS5      nUc  [        USS5      nUc  [        USS5      n[        USS5      S:X  a%  [        US	S
S
/5      Ul        UR                  Ul        [        U SS5      (       d  U R                  (       a  X$SSS4$ [        U SS5      nUc  [        U SS5      nUS   nU R
                  U   nU R                  U   nUR                  nUR                  n	[        US5      (       a  [        USS5      n
U
b  U
" U5      n[        US5      (       a  [        USS5      nUb  U" U	5      n	UUUU	U R                  U   4$ )z
Return a 5-tuple of (weight, weight quant_state, lora A, lora B, and lora scale).
If QAT is enabled, additionally fake quantize the base layer and lora weights.

base_layerweight_fake_quantizerNrD   weight_scale_invweight_scalequant_methodfp8
block_size   disable_adaptersTactive_adaptersactive_adapterdefaultr   )rF   weighthasattrrP   mergedlora_Alora_Bscaling)projrJ   rG   rK   W_quantadapterlora_A_linearlora_B_linearABlora_A_fake_quantizerlora_B_fake_quantizers               r   get_lora_parametersre      s    lDJ 	A z233 '
4KT R ,%a(A a-G*&8$??j.$?Gz>40E9z<#sD\\ t'..$++4t++d-t4G$ 09>ajG KK(MKK(MAA}566 '7NPT U ,%a(A}566 '7NPT U ,%a(A 	
		W r   c                 @   [        U SU 5      nUR                  n[        USS 5      nUc  [        USS 5      nUc  [        USS 5      n[        U SS5      (       d  U R                  (       a  X#S S S UR                  4$ [        USS 5      S:X  a%  [        US	S
S
/5      Ul        UR                  Ul        [        U SS 5      nUc  [        U SS5      nUS   nUUU R
                  U   R                  U R                  U   R                  U R                  U   UR                  4$ )NrJ   rD   rL   rM   rR   TrN   rO   rP   rQ   rS   rT   rU   r   )rF   rV   rX   biasrP   rY   rZ   r[   )r\   rJ   rG   r]   r^   s        r   get_lora_parameters_biasrh   '  s*   lDJ 	A a-G*&8$??j.$?G t'..$++4tZ__<<z>40E9z<#sD\\d-t4G$ 09>ajG 	
G##G##W r   Xr\   c                 P    [        USU5      n[        USS5      nUb  U" U 5      n U $ )z
If QAT is enabled, fake quantize the input activations.
Otherwise, just return the input activations as is.
Weights are fake quantized separately in `get_lora_parameters`.
rJ   activation_fake_quantizerNrE   )ri   r\   rJ   rk   s       r    _maybe_fake_quantize_activationsrl   M  s8     |T2J '
4OQU V ,%a(Hr   c                 @   [        U [        5      (       a  U R                  5       $ Uc  U $ U R                  [        R
                  :X  a  [        X5      $ [        U5      [        Lam  UR                  nUR                  nUR                  nUR                  nUR                  nUR                  n	U	R                  n
U	R                  nU	R                  nOUu  pEpgn  nUu  pU	u  pn      nU R                  nUR                   n["        U   nUR%                  5       nU(       a  US   US   -  n[&        U   n[(        U   nUc7  [+        UXoSS9=[&        U'   n[+        U[        R,                  USS9=[(        U'   nUUR%                  5       :  a  UR/                  U5        UUR%                  5       :  a  UR/                  U5        US U R1                  U5      nUS U nOCUc  [+        XVUSS9nO$UR                  U:X  d   eUR                  U:X  d   e[+        U[2        USS9n[5        U5      n[7        U5         [9        [5        U5      [5        U5      [5        U
5      U[;        U5      [;        U5      U5        UU-  nU[<        :X  a  [>        O[@        nU" [5        S 5      [5        U 5      U[5        U5      [;        U5      [;        UR%                  5       5      U5        S S S 5        U R                  S   S:X  a  SOSnU(       a  URC                  5       $ U$ ! , (       d  f       N>= fNr   r   Fr   r9   requires_gradT)"
isinstancerB   
dequantizer   torchfloat8_e4m3fnr   typelistabsmaxshape	blocksizeoffsetstate2coder9   r?   XPU_STREAMSnumelWEIGHT_BUFFERSABSMAX_BUFFERStorch_emptyfloat32resize_viewtorch_float32get_ptrr:   cdequantize_blockwise_fp32ctypes_c_inttorch_float16cdequantize_blockwise_fp16_nf4cdequantize_blockwise_bf16_nf4t)rG   rD   outuse_global_bufferrw   rx   r   ry   rz   r{   absmax2code2
blocksize2compressed_stats_r9   device_index
XPU_STREAMn_elements_absmaxsizeWEIGHT_BUFFERABSMAX_BUFFER
out_absmaxptr_out_absmaxfxis_transposeds                             r   fast_dequantizer   _  s    a&&<<>!H77e)))!!11D( !''F%%E%%E#--I ''F ''FmmGKKE))J GRCF5-=q!-NF5;2GJ1a|| .
"LLN8eAh&D +<8M*<8M$?J%%@ |,} @K%!MM#$)	@ |,} m))++%%d+ =#6#6#88%%&78&++E2C&'9(9:J{!65 yyE)))yyE)))$!% %	J !,f%& Z(./ & J
 M) /3 
 
Y'SYY[)% &8 !"
aU'suuw0S0; &%s   4B#L
Lc                 2   [        U [        5      (       a  U R                  5       $ Uc  U $ U R                  [        R
                  :X  a  [        X5      $ [        U5      [        Lam  UR                  nUR                  nUR                  nUR                  nUR                  nUR                  n	U	R                  n
U	R                  nU	R                  nOUu  pEpgn  nUu  pU	u  pn      n U R                  nUR                   n["        U   nUR%                  5       nU(       a  US   US   -  n[&        U   n[(        U   nUc-  [+        UXoSS9=[&        U'   n[+        U[,        USS9=[(        U'   nUUR%                  5       :  a  UR/                  U5        UUR%                  5       :  a  UR/                  U5        US U R1                  U5      nUS U nOCUc  [+        XVUSS9nO$UR                  U:X  d   eUR                  U:X  d   e[+        U[,        USS9n [3        U5      n[5        U5         [7        [3        U5      [3        U5      [3        U
5      U[9        U5      [9        U5      U5        UU-  nU[:        :X  a  [<        O[>        nU" [3        S 5      [3        U 5      U[3        U5      [9        U5      [9        UR%                  5       5      U5        S S S 5         U R                  S   S:X  a  SOSnU(       a  URA                  5       $ U$ ! , (       d  f       N?= frn   )!rq   rB   rr   r   rs   rt   r   ru   rv   rw   rx   ry   rz   r{   r|   r9   r?   CUDA_STREAMSr~   r   r   r   r   r   r   r   r:   r   r   r   r   r   r   )rG   rD   r   r   rw   rx   r   ry   rz   r{   r   r   r   r   r   r9   r   CUDA_STREAMr   r   r   r   r   r   r   r   s                             r   r   r     s!   a&&<<>!H77e)))!!11D( !''F%%E%%E#--I ''F ''FmmGKKE))J GRCF5-=q!-NF5;2GJ1a||"<0"LLN 8eAh&D +<8M*<8M$?J%%@ |,} @K%)#$)	@ |,} m))++%%d+ =#6#6#88%%&78&++E2C&'9(9:J{!65 yyE)))yyE)))$!% %	J 	 !,f%& Z(./ & J
 M) /3 
 
Y'SYY[)% &6 	 !
aU'suuw0S0= &%s   ,B#L
Lc                 F   [        U [        5      (       a  U R                  5       $ Uc  U $ U R                  [        R
                  :X  a  [        X5      $ [        U5      [        Lam  UR                  nUR                  nUR                  nUR                  nUR                  nUR                  n	U	R                  n
U	R                  nU	R                  nOUu  pEpgn  nUu  pU	u  pn      n UR                  5       nU R                   nUc  [#        XVUSS9nO$UR                  U:X  d   eUR                  U:X  d   e[#        U[$        USS9n['        U5      n[)        ['        U5      ['        U5      ['        U
5      U[+        U5      [+        U5      5        UU-  nU[,        :X  a  [.        O[0        nU" ['        S 5      ['        U 5      U['        U5      [+        U5      [+        UR                  5       5      5        U R                  S   S:X  a  SOSnU(       a  UR3                  5       $ U$ )NFro   r   r   T)rq   rB   rr   r   rs   rt   r   ru   rv   rw   rx   ry   rz   r{   r|   r~   r9   r   r   r   r   r   r   r   r   r   )rG   rD   r   r   rw   rx   r   ry   rz   r{   r   r   r   r   r   r   r9   r   r   r   r   s                        r   r   r   <  s   a&&<<>!H77e)))!!11D( !''F%%E%%E#--I ''F ''FmmGKKE))J GRCF5-=q!-NF5;2GJ1a"LLN ;eVUZ[C99%%%99%%% }vW\


 !,"ENFOG$*+	
 	f
 % +/ 	
 	DMAJCL#%	
 !"
aU'suuw0S0r   c                    Uc
  [        XUS9$ U R                  u  pEn[        U5      [        Lay  UR                  nUR                  nUR
                  n	UR                  n
UR                  nUR                  nUR                  nUR                  nUR                  nUR                  nOUu  pxpnnnUu  pUu  pn      nUR                  nUR                  n[        U   nUS   nUc  [        SSU4U	US9n[        S:X  a  SnUS   nOSnUS   nUS   nUS   nUS   nUS-   S-  n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        UR                  [         US9n[#        U5         [%        ['        U5      ['        U5      ['        U5      ['        U5      [)        U5      [)        UR+                  5       5      U5        UU-  nUnU	[,        :X  a  [.        O[0        n[        U
5      n
U" UUU['        U 5      ['        U5      ['        U5      ['        U5      ['        U5      UUUU
U5        S S S 5        U$ ! , (       d  f       U$ = f)Nr   r   r   r   r9   r   r   )torch_matmulrx   ru   rv   rw   r   ry   r|   rz   r{   r9   r?   r}   r   r   ctypes_c_int32r   r:   r   r   r   r~   r   cgemm_4bit_inference_naive_fp16cgemm_4bit_inference_naive_bf16)ri   rG   rD   r   r   q_lenhdrw   rx   r   ry   statsrz   r{   r   r   r   r   
quant_typer9   r   r   boutmr(   kldaldcldbdfr   s                                  r   	fast_gemvr     s   C00 ww" D( ''F%%E%%E#--I$$E ''F ''FmmGKKE))J  QF5-=z5 .NF5;2GJ1a|| .
 Qx;
 C %AaAAaA!HAhAhAv!m111S!S!S!}vNf%& Z(RXXZ( &LBF M) 04  'y1I

) &H 
I &%H 
s   
B?I
I"c                    Uc
  [        XUS9$ U R                  u  pEn[        U5      [        Lay  UR                  nUR                  nUR
                  n	UR                  n
UR                  nUR                  nUR                  nUR                  nUR                  nUR                  nOUu  pxpnnnUu  pUu  pn      n UR                  nUR                  n[        U   nUS   nUc  [        SSU4U	US9nSnUS   nUS   nUS   nUS   nUS-   S-  n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        UR                  [        US9n[!        U5         [#        [%        U5      [%        U5      [%        U5      [%        U5      ['        U5      ['        UR)                  5       5      U5        UU-  nUnU	[*        :X  a  [,        O[.        n[        U
5      n
U" UUU[%        U 5      [%        U5      [%        U5      [%        U5      [%        U5      UUUU
U5        S S S 5         U$ ! , (       d  f        U$ = fNr   r   r   r   r   )r   rx   ru   rv   rw   r   ry   r|   rz   r{   r9   r?   r   r   r   r   r:   r   r   r   r~   r   r   r   )ri   rG   rD   r   r   r   r   rw   rx   r   ry   r   rz   r{   r   r   r   r   r   r9   r   r   r   r(   r   r   r   r   r   r   r   s                                  r   r   r     s   C00 ww" D( ''F%%E%%E#--I$$E ''F ''FmmGKKE))J  QF5-=z5 .NF5;2GJ1a||"<0 Qx;
 C !H!HAhAhAv!m111S!S!S!}vNf%& Z(RXXZ( &LBF M) 04  'y1I

) &F 	
K &%F 	
s   9B?I
Ic                 .   Uc
  [        XUS9$ U R                  u  pEn[        U5      [        Lay  UR                  nUR                  nUR
                  n	UR                  n
UR                  nUR                  nUR                  nUR                  nUR                  nUR                  nOUu  pxpnnnUu  pUu  pn      n US   nUR                  nUc  [        SSU4U	US9nSnUS   nUS   nUS   nUS   nUS-   S-  n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        U5      n[        UR                  [        US9n[        [        U5      [        U5      [        U5      [        U5      [!        U5      [!        UR#                  5       5      5        UU-  nUnU	[$        :X  a  [&        O[(        n[        U
5      n
U" UUU[        U 5      [        U5      [        U5      [        U5      [        U5      UUUU
5        U$ r   )r   rx   ru   rv   rw   r   ry   r|   rz   r{   r9   r   r   r   r   r   r   r~   r   r   r   )ri   rG   rD   r   r   r   r   rw   rx   r   ry   r   rz   r{   r   r   r   r   r   r   r9   r(   r   r   r   r   r   r   r   s                                r   r   r   [  sC   C00 ww" D( ''F%%E%%E#--I$$E ''F ''FmmGKKE))J  QF5-=z5 .NF5;2GJ1aQx;
 C !H!HAhAhAv!m111S!S!S!}vN"ENFOGBK$$	
 	f % ,0 	 #9-	
AJAJFOENCL	
 
r   c                    [        U 5      u  pEpgpUR                  u  pnUS:w  a  [        XXVXx5      $ Uc  [        XR	                  5       US9nOeUR
                  [        R                  :X  a  [        XXY5      nO:U
S:X  a  US:X  a  [        XXSS9nO#[        UR	                  5       USS9n[        XUS9nUGb+  UR                  S   nUR
                  n[        US5      (       d,  UR                  U5      Ul        UR                  U5      Ul        U
S:X  aO  UR                  U5      n[        UR                  UR!                  5       US9nUR#                  UR                  X(S9  OkUR                  X5      n[%        UR                  X5      UR                  R	                  5       US9nUR'                  X'R                  R	                  5       US9  UR                  U
SU5      nU	b  X9-  nU$ )Nr   r   Tr   r   
_fast_loraalpha)rh   rx   matmul_lorar   r   r   rs   rt   r   r   r   rW   r   r   r   torch_mvraveladdmv_torch_mmaddmm_)r\   ri   	temp_lorar   rG   r]   rY   rZ   lora_Srg   bszr   in_dimout_dimr   s                  r   fast_linear_forwardr     s   /G/M,ACz1&AA1cce3/	
E''	'w-	eqjg1ACCE7E1s+ ))A,v|,, &		% 0F &		% 0F!8((7#C !2!2AGGIYOIJJv(()JD((3(C s#V%6%6%8%8%:)I JJy"3"3"5"5"7JHhhsAw'Jr   c                 `   U R                   nU R                  5       S:X  a1  U R                  u  pn
U R                  SU R                  S   5      n SnOSn[	        U[
        5      (       a  UR                  S:X  d   eUR                  S   UR                  S   :X  a$  UR                  S   S:X  a  UR                  5       nOUR                  5       n[        XR                  5       US9nOMUR                   [        R                  :X  a  [        XU5      nO"[        XSS	9n[        XR                  5       US9nUb  AUbX  UR                  5       UR                  5       pC[        XR!                  U5      5      nUR#                  XR!                  U5      US
9  U(       a  UR                  WW	S5      $ U$ )N   TFr   r   r   r   r   r   )r   dimrx   r   rq   rB   ndimrP   rr   
contiguousr   r   rs   rt   r   r   r   r   )ri   rG   r]   ra   rb   sr   r   batchseq_lendreshapeXAs                r   r   r     sY   GGEuuw!|GGFF2qwwr{#!\""vv{{<<?aggaj(Q\\!_-A AA1cce3/	
E''	'w'ADA1cce3/}ssuacce1!TT%[)

2ttE{A
. ,3388E7B';;r   )NNFr   )NN)	importlibr    ctypesr   int__annotations__r0   	functoolstypingr   r   r   r   r   r	   r
   r   rO   r   r   rs   Tensortorch_Tensorunsloth_zoo.utilsr   __version__r1   r   amp
custom_fwdtorch_amp_custom_fwd
custom_bwdtorch_amp_custom_bwdtriton.languagelanguagetlextraintelr   tanhtriton_tanhtriton.language.extracastr   mathjit	lru_cacher'   r4   HAS_CUDA_STREAMbitsandbytesbnb
functionalr   HAS_XPU_STREAMr9   r:   r   
contextlibr7   _C_xpu_getCurrentRawStreamr>   _cuda_getCurrentRawStreamr=   r@   rangeidxr?   _XPU_STREAMSmaxkeysr}   r   r   itemsr   vtuple_CUDA_STREAMSr   c_intr   c_int32r   libr   r   r   cgemv_4bit_inference_fp16r   cgemv_4bit_inference_bf16r   current_streamtorch_device_streammmr   mvr   matmulr   addmmtorch_addmmemptyr   r   r   float16r   bfloat16torch_bfloat16util	find_spectorchao.quantizationrB   rA   printru   rH   re   rh   nnModulerl   inference_moder   r   r   r   )ir?   s   00r   <module>r     s1      ((    ,  || %%GE$5$569II
J  5 00 ::>>44 ::>>44 99//f/E 99//f/E% 99//e/D 99//e/D  
6''"22ehhnn..333nn''K'',,K ZZ  Q !
! !,   #//*WX->>
..
 
 %N!o% ::,,		 99++&
 %$xx@@  %xxAA??C| C C %
 |$	 %A ))""1%))	)FOOHH--e4-
 

 %	L &C 1 1 34q89KVs<#4#4#67!;<NVs<#4#4#67!;<N""$1A %$K |$	 %A **##A&**	*V__HH..u5.
 

 %	M 6S!3!3!56:;LVs=#5#5#781<=NVs=#5#5#781<=N##%1Q &&L || ^^//JJ !$!3!3!R!R !$!3!3!R!R % '*nn&8&8&R&R#&)nn&8&8&R&R#&)nn&8&8&X&X#&)nn&8&8&X&X# !,u 4EII%**:S:S  8888||kkkk >>I&2	"5 :L+:z#L|| 88??
\\  %N
g1 g1T O#
j1 j1X 	 C1 C1J 	 %NfP O#dL 	\| 	(V#<[b"7&&'78+<<]^e^q^q]rs Dzs   0A]%A]<]# #?^%