§
    Ý˜ºh;  ã                   óä  — d dl mZ ddlmZmZmZmZ  e¦   «         rd dlZd dlm	Z	 d dl
Z
d dlmZ d dlmZ  e¦   «         rd dlmZ  ej        e¦  «        Ze
j        dej        fd„¦   «         Zd#d
ej        dedeej        ej        f         fd„Ze
j        dej        dej        dej        dej        fd„¦   «         Zej        fdej        dej        dej        dej        dee         dej        dej        fd„Z ej!        dej        fdej        dej        dej        dej        deeeef                  dej        dej        fd„¦   «         Z" G d„ de	j#        ¦  «        Z$	 	 	 	 	 d$d!„Z%	 	 d%d"„Z&dS )&é    )ÚOptionalé   )Úis_accelerate_availableÚis_torch_accelerator_availableÚis_torch_availableÚloggingN)Ú
functional)Úinit_empty_weightsÚ
BLOCK_SIZEc                 óÈ  — t          j        d¬¦  «        }||z  t          j        d|¦  «        z   }t          j        | |z   ¦  «                             t           j        ¦  «        }t          j        t          j        |¦  «        ¦  «        dz  }||z  }|                     |j        j	        ¦  «        }t          j
        ||z   |¦  «         t          j
        ||z   |¦  «         d S )Nr   ©Úaxisg      |@)ÚtlÚ
program_idÚarangeÚloadÚtoÚfloat32ÚmaxÚabsÚdtypeÚ
element_tyÚstore)	Úx_ptrÚy_ptrÚs_ptrr   ÚpidÚoffsÚxÚsÚys	            úq/var/www/html/services/rag-service/venv/lib/python3.11/site-packages/transformers/integrations/finegrained_fp8.pyÚact_quant_kernelr#   $   s¿   € å
Œ-˜QÐ
Ñ
Ô
€CØ�Ñ�bœi¨¨:Ñ6Ô6Ñ6€DÝ
Œ�˜‘ÑÔ× Ò ¥¤Ñ,Ô,€AÝ
Œ�rŒv�a‰yŒyÑÔ˜EÑ!€AØ	ˆA‰€AØ	�ŠˆUŒ[Ô#Ñ$Ô$€AÝ„HˆU�T‰\˜1ÑÔÐÝ„HˆU�S‰[˜!ÑÔÐÐÐó    é€   r   Ú
block_sizeÚreturnc                 óx  ‡ — ‰                       ¦   «         sJ ‚‰ j        d         |z  dk    sJ ‚t          j        ‰ t          j        ¬¦  «        } ‰ j        g ‰                      ¦   «         d d…         ¢‰                      d¦  «        |z  ‘R dt          j        iŽ}ˆ fd„}t          |         ‰ |||¬¦  «         ||fS )Néÿÿÿÿr   ©r   r   c                 ó`   •— t          j        ‰                     ¦   «         | d         ¦  «        fS )Nr   )ÚtritonÚcdivÚnumel)Úmetar   s    €r"   Úgridzact_quant.<locals>.grid6   s%   ø€ Ý”˜AŸGšG™IœI t¨LÔ'9Ñ:Ô:Ð<Ð<r$   )r   )	Úis_contiguousÚshapeÚtorchÚ
empty_likeÚfloat8_e4m3fnÚ	new_emptyÚsizer   r#   )r   r&   r!   r    r0   s   `    r"   Ú	act_quantr8   0   sÒ   ø€ Ø�?Š?ÑÔÐÐÐØŒ7�2Œ;˜Ñ# qÒ(Ð(Ð(Ð(ÝÔ˜¥%Ô"5Ð6Ñ6Ô6€AØˆŒÐR�Q—V’V‘X”X˜c˜r˜c”]ÐR A§F¢F¨2¡J¤J°*Ñ$<ÐRÐRÐRÅEÄMÐRÐR€Að=ð =ð =ð =ð =õ �TÔ˜1˜a ¨zÐ:Ñ:Ô:Ð:Øˆaˆ4€Kr$   ÚBLOCK_SIZE_MÚBLOCK_SIZE_NÚBLOCK_SIZE_KÚGROUP_SIZE_Mc                 ó¾  — t          j        d¬¦  «        }t          j        ||¦  «        }t          j        ||¦  «        }||z  }||z  }||z  }t          ||z
  |¦  «        }|||z  z   }||z  |z  } ||z  t          j        d|¦  «        z   |z  }!| |z  t          j        d|¦  «        z   |z  }"t          j        d|¦  «        }#| |!dd…df         |
z  |#ddd…f         |z  z   z   }$||#dd…df         |z  |"ddd…f         |z  z   z   }%||!|z  z   }&|"|z  }'||'|z  z   }(t          j        ||ft           j        ¬¦  «        })t          dt          j        ||¦  «        ¦  «        D ]Ø}*t          j        |$|#ddd…f         ||*|z  z
  k     d¬¦  «        }+t          j        |%|#dd…df         ||*|z  z
  k     d¬¦  «        },|*|z  }-|-|	z  }.t          j        |&|.|z  z   ¦  «        }/t          j        |(|.|z  z   ¦  «        }0|)t          j	        |+|,¦  «        |/dd…df         z  |0ddd…f         z  z  })|$||z  z  }$|%||z  z  }%ŒÙ|j
        j        t           j        k    r |)                     t           j        ¦  «        }1nY|j
        j        t           j        k    r |)                     t           j        ¦  «        }1n|)                     t           j        ¦  «        }1||z  t          j        d|¦  «        z   }2| |z  t          j        d|¦  «        z   }3|||2dd…df         z  z   ||3ddd…f         z  z   }4|2dd…df         |k     |3ddd…f         |k     z  }5t          j        |4|1|5¬¦  «         dS )z¾Triton-accelerated function used to perform linear operations (dot
    product) on input tensors `A` and `B` with block-wise quantization, and
    store the result in output tensor `C`.
    r   r   Nr*   g        )ÚmaskÚother)r>   )r   r   r-   Úminr   Úzerosr   Úranger   Údotr   r   Úbfloat16r   Úfloat16r   )6ÚAÚBÚCÚAsÚBsÚMÚNÚKÚgroup_nÚgroup_kÚ	stride_amÚ	stride_akÚ	stride_bkÚ	stride_bnÚ	stride_cmÚ	stride_cnÚstride_As_mÚstride_As_kÚstride_Bs_kÚstride_Bs_nr9   r:   r;   r<   r   Ú	num_pid_mÚ	num_pid_nÚnum_pid_in_groupÚgroup_idÚfirst_pid_mÚgroup_size_mÚpid_mÚpid_nÚoffs_amÚoffs_bnÚoffs_kÚa_ptrsÚb_ptrsÚAs_ptrsÚoffs_bsnÚBs_ptrsÚaccumulatorÚkÚaÚbÚk_startÚoffs_ksÚa_sÚb_sÚcÚoffs_cmÚoffs_cnÚc_ptrsÚc_masks6                                                         r"   Ú_w8a8_block_fp8_matmulrw   >   s¼  € õJ Œ-˜QÐ
Ñ
Ô
€CÝ”˜˜<Ñ(Ô(€IÝ”˜˜<Ñ(Ô(€IØ# iÑ/ÐØÐ&Ñ&€HØ˜\Ñ)€KÝ�y ;Ñ.°Ñ=Ô=€LØ˜3 Ñ-Ñ.€EØÐ#Ñ#¨Ñ4€Eà�|Ñ#¥b¤i°°<Ñ&@Ô&@Ñ@ÀAÑE€GØ�|Ñ#¥b¤i°°<Ñ&@Ô&@Ñ@ÀAÑE€GÝŒY�q˜,Ñ'Ô'€FØ�'˜!˜!˜!˜T˜'Ô" YÑ.°¸¸a¸a¸a¸´À9Ñ1LÑLÑM€FØ�&˜˜˜˜D˜”/ IÑ-°¸¸a¸a¸a¸Ô0@À9Ñ0LÑLÑM€Fà�7˜[Ñ(Ñ(€GØ˜'Ñ!€HØ�8˜kÑ)Ñ)€Gå”(˜L¨,Ð7½r¼zÐJÑJÔJ€KÝ�1•b”g˜a Ñ.Ô.Ñ/Ô/ð +ð +ˆÝŒG�F ¨¨a¨a¨a¨¤°1°q¸<Ñ7GÑ3GÒ!GÈsÐSÑSÔSˆÝŒG�F ¨¨¨¨4¨¤°1°q¸<Ñ7GÑ3GÒ!GÈsÐSÑSÔSˆà�lÑ"ˆØ˜WÑ$ˆÝŒg�g ¨+Ñ 5Ñ5Ñ6Ô6ˆÝŒg�g ¨+Ñ 5Ñ5Ñ6Ô6ˆà•r”v˜a ‘|”| c¨!¨!¨!¨T¨'¤lÑ2°S¸¸q¸q¸q¸´\ÑAÑAˆØ�, Ñ*Ñ*ˆØ�, Ñ*Ñ*ˆˆà„wÔ�Rœ[Ò(Ð(Ø�NŠN�2œ;Ñ'Ô'ˆˆØ	
ŒÔ	�rœzÒ	)Ð	)Ø�NŠN�2œ:Ñ&Ô&ˆˆà�NŠN�2œ:Ñ&Ô&ˆà�lÑ"¥R¤Y¨q°,Ñ%?Ô%?Ñ?€GØ�lÑ"¥R¤Y¨q°,Ñ%?Ô%?Ñ?€GØ�˜W Q Q Q¨ WÔ-Ñ-Ñ-°	¸GÀDÈ!È!È!ÀGÔ<LÑ0LÑL€FØ�a�a�a˜�gÔ Ò" w¨t°Q°Q°Q¨wÔ'7¸!Ò';Ñ<€FÝ„HˆV�Q˜VÐ$Ñ$Ô$Ð$Ð$Ð$r$   rF   rG   rI   rJ   Úoutput_dtypec                 óR  ‡‡— t          |¦  «        dk    sJ ‚|d         |d         }}| j        d         |j        d         k    sJ ‚| j        dd…         |j        dd…         k    r|                      ¦   «         sJ ‚t          j        | j        d         |¦  «        |j        d         k    sJ ‚|                      ¦   «         | j        d         z  Š|j        dk    r|                     ¦   «         r|j        dk    sJ ‚|j        \  Š}t          j        ‰|¦  «        |j        d         k    sJ ‚t          j        ||¦  «        |j        d         k    sJ ‚| j        dd…         ‰fz   }	|                      |	|¬¦  «        }
d}‰|k     r$t          j        ‰¦  «        }t          |d¦  «        }|}||z  dk    sJ ‚|}ˆˆfd	„}t          |         | ||
||‰‰||||                      d
¦  «        |                      d¦  «        |                     d¦  «        |                     d¦  «        |
                     d
¦  «        |
                     d¦  «        |                     d
¦  «        |                     d¦  «        |                     d¦  «        |                     d¦  «        |||d¬¦  «         |
S )a‰  This function performs matrix multiplication with block-wise
    quantization.
    It takes two input tensors `A` and `B` with scales `As` and `Bs`.
    The output is returned in the specified `output_dtype`.
    Args:
        A: The input tensor, e.g., activation.
        B: The input tensor, e.g., weight.
        As: The per-token-group quantization scale for `A`.
        Bs: The per-block quantization scale for `B`.
        block_size: The block size for per-block quantization. It should
        be 2-dim, e.g., [128, 128].
        output_dytpe: The dtype of the returned tensor.
    Returns:
        torch.Tensor: The result of matmul.
    r   r   é   r)   Nr*   r%   é   c                 ót   •— t          j        ‰| d         ¦  «        t          j        ‰| d         ¦  «        z  fS )Nr9   r:   )r,   r-   )ÚMETArK   rL   s    €€r"   r0   z*w8a8_block_fp8_matmul_triton.<locals>.gridÂ   s3   ø€ Ý”˜A˜t NÔ3Ñ4Ô4µv´{À1ÀdÈ>ÔFZÑ7[Ô7[Ñ[Ð]Ð]r$   éþÿÿÿé   )r9   r:   r;   r<   )Úlenr2   r1   r,   r-   r.   Úndimr6   Únext_power_of_2r   rw   Ústride)rF   rG   rI   rJ   r&   rx   Úblock_nÚblock_krM   ÚC_shaperH   r9   r;   r:   r0   rK   rL   s                  @@r"   Úw8a8_block_fp8_matmul_tritonr‡   “   s§  øø€ õ. ˆz‰?Œ?˜aÒÐÐÐØ! !”} j°¤mˆW€GàŒ7�2Œ;˜!œ' "œ+Ò%Ð%Ð%Ð%ØŒ7�3�B�3Œ<˜2œ8 C R Cœ=Ò(Ð(¨Q¯_ª_Ñ->Ô->Ð(Ð(Ð(ÝŒ;�q”w˜r”{ GÑ,Ô,°´¸´Ò<Ð<Ð<Ð<Ø	�Š‰	Œ	�Q”W˜R”[Ñ €AàŒ6�QŠ;ˆ;˜1Ÿ?š?Ñ,Ô,ˆ;°´¸A²°°°ØŒ7�D€A€qÝŒ;�q˜'Ñ"Ô" b¤h¨q¤kÒ1Ð1Ð1Ð1ÝŒ;�q˜'Ñ"Ô" b¤h¨q¤kÒ1Ð1Ð1Ð1àŒg�c�r�cŒl˜a˜TÑ!€GØ	�Š�G <ˆÑ0Ô0€Aà€LØˆ<ÒÐÝÔ-¨aÑ0Ô0ˆÝ˜<¨Ñ,Ô,ˆØ€LØ�\Ñ! QÒ&Ð&Ð&Ð&Ø€Lð^ð ^ð ^ð ^ð ^ð ^õ ˜4Ô Ø	Ø	Ø	Ø
Ø
Ø	Ø	Ø	ØØØ	�Š�‰ŒØ	�Š�‰ŒØ	�Š�‰ŒØ	�Š�‰ŒØ	�Š�‰ŒØ	�Š�‰ŒØ
�	Š	�"‰ŒØ
�	Š	�"‰ŒØ
�	Š	�!‰ŒØ
�	Š	�!‰ŒØ!Ø!Ø!Øð1ñ ô ð ð6 €Hr$   Úinput_qÚweight_qÚinput_scaleÚweight_scalec                 ó¤  — | j         dk    r| j        nd| j        d         | j        d         f\  }}}|j        d         }	|                      d|¦  «        }
|                     |j        d         d¦  «        }|	|d         z  }||d         z  }t          j        ||z  |	ft          j        | j        ¬¦  «        }t          |¦  «        D ]ß}||d         z  }||d         z   }t          |¦  «        D ]·}||d         z  }||d         z   }|
dd…||…f         }|||…||…f         }|dd…||dz   …f         }|||f         }t          j        || 	                    ¦   «         t          j
        dt          j        | j        ¬¦  «        ||¬¦  «        |z  }|dd…||…fxx         |z  cc<   Œ¸Œà|                     |||	¦  «        }|                     |¦  «        S )aÁ  
    Performs blocked matrix multiplication with FP8 quantized matrices.

    Args:
        input_q: Quantized input tensor with 1x128 block quantization
        weight_q: Quantized weight tensor with 128x128 block quantization
        input_scale: Scaling factors for input blocks
        weight_scale: Scaling factors for weight blocks
        block_size: Tuple of (M, N) for weight block dimensions
        output_dtype: Desired output dtype
    é   rz   r   r)   ©r   ÚdeviceN)Úscale_aÚscale_bÚ	out_dtype)r�   r2   Úviewr3   rA   r   r�   rB   Ú
_scaled_mmÚtÚtensorr   )rˆ   r‰   rŠ   r‹   r&   rx   Ú
batch_sizeÚseq_lenÚ
hidden_dimÚout_featuresÚinput_reshapedÚinput_scale_reshapedÚnum_weight_blocks_mÚnum_weight_blocks_nÚoutputÚiÚm_startÚm_endÚjÚn_startÚn_endÚinput_blockÚweight_blockÚcurr_input_scaleÚcurr_weight_scaleÚblock_results                             r"   Úw8a8_block_fp8_matmul_compiler«   ä   s4  € ð( 8?´|ÀqÒ7HÐ7H g¤m mÈqÐRYÔR_Ð`aÔRbÐdkÔdqÐrsÔdtÐNuÑ#€J�˜Ø”> !Ô$€Lð —\’\ " jÑ1Ô1€NØ&×+Ò+¨KÔ,=¸aÔ,@À"ÑEÔEÐà&¨*°Q¬-Ñ7ÐØ$¨
°1¬Ñ5ÐåŒ[˜* wÑ.°Ð=ÅUÄ]Ð[bÔ[iÐjÑjÔj€FåÐ&Ñ'Ô'ð 5ð 5ˆØ�j ”mÑ#ˆØ˜* Qœ-Ñ'ˆåÐ*Ñ+Ô+ð 	5ð 	5ˆAØ˜* Qœ-Ñ'ˆGØ˜j¨œmÑ+ˆEð )¨¨¨¨G°E¨MÐ)9Ô:ˆKØ# G¨E M°7¸5°=Ð$@ÔAˆLð  4°A°A°A°q¸1¸q¹5°y°LÔAÐØ ,¨Q°¨TÔ 2Ðõ Ô ØØ —N’NÑ$Ô$Ý!œL¨µ%´-ÈÌÐWÑWÔWØ-Ø*ðñ ô ð #ñ#ð ð �1�1�1�g˜e�mÐ#Ð$Ð$Ô$¨Ñ4Ð$Ð$Ñ$Ð$ð/	5ð2 �[Š[˜ W¨lÑ;Ô;€Fà�9Š9�\Ñ"Ô"Ð"r$   c                   óŠ   ‡ — e Zd Zej        Z	 	 	 	 	 ddedededee	eef                  fˆ fd„Z
d	ej        d
ej        fd„Zˆ xZS )Ú	FP8LinearFNÚdynamicÚin_featuresrš   Úbiasr&   c                 óÖ  •— t          ¦   «                              ||¦  «         || _        || _        t          j                             t	          j        ||t          j	        |¬¦  «        ¦  «        | _
        | j
                             ¦   «         dk    rh||d         z   dz
  |d         z  }||d         z   dz
  |d         z  }	t          j        t	          j        ||	t          j        |¬¦  «        ¦  «        | _        n|                      dd ¦  «         || _        || _        |r2t          j        t	          j        | j        ¦  «        ¦  «        | _        d S |                      dd ¦  «         d S )NrŽ   rz   r   Úweight_scale_invr°   )ÚsuperÚ__init__r¯   rš   r3   ÚnnÚ	ParameterÚemptyr­   r   ÚweightÚelement_sizer   r²   Úregister_parameterr&   Úactivation_schemer°   )Úselfr¯   rš   r°   r   r&   r�   r»   Úscale_out_featuresÚscale_in_featuresÚ	__class__s             €r"   r´   zFP8Linear.__init__)  sS  ø€ õ 	‰Œ×Ò˜ lÑ3Ô3Ð3Ø&ˆÔØ(ˆÔå”h×(Ò(­¬°\À;ÕV_ÔVeÐntÐ)uÑ)uÔ)uÑvÔvˆŒàŒ;×#Ò#Ñ%Ô%¨Ò*Ð*Ø".°¸A´Ñ">ÀÑ"BÀzÐRSÄ}Ñ!TÐØ!,¨z¸!¬}Ñ!<¸qÑ!@ÀZÐPQÄ]Ñ RÐÝ$&¤LÝ”Ð.Ð0AÍÌÐ_eÐfÑfÔfñ%ô %ˆDÔ!Ð!ð ×#Ò#Ð$6¸Ñ=Ô=Ð=à$ˆŒà!2ˆÔàð 	2Ýœ¥U¤[°Ô1BÑ%CÔ%CÑDÔDˆDŒIˆIˆIà×#Ò# F¨DÑ1Ô1Ð1Ð1Ð1r$   Úinputr'   c           	      ó”  — | j                              ¦   «         dk    r t          j        || j         | j        ¦  «        S t          ¦   «         r#t          j                             ¦   «         j	        nd}t          t          |t          j        ¦  «        }|                     |j        ¦  «        5  t          || j        d         ¦  «        \  }}t          || j         || j        | j        |j        ¬¦  «        }d d d ¦  «         n# 1 swxY w Y   |                     ¦   «          | j        �
|| j        z   }|                     |j        ¬¦  «        S )Nrz   Úcuda)rx   r*   )r¸   r¹   ÚFÚlinearr°   r   r3   ÚacceleratorÚcurrent_acceleratorÚtypeÚgetattrrÂ   r�   r8   r&   r‡   r²   r   Úsynchronizer   )r¼   rÀ   Údevice_typeÚtorch_accelerator_moduleÚqinputÚscalerŸ   s          r"   ÚforwardzFP8Linear.forwardK  se  € ØŒ;×#Ò#Ñ%Ô%¨Ò)Ð)Ý”8˜E 4¤;°´	Ñ:Ô:Ð:õ KiÑJjÔJjÐv�%Ô+×?Ò?ÑAÔAÔFÐFÐpvˆKÝ'.­u°kÅ5Ä:Ñ'NÔ'NÐ$Ø)×0Ò0°´Ñ>Ô>ð 	ð 	Ý )¨%°´ÀÔ1CÑ DÔ D‘�˜Ý5ØØ”KØØÔ)Ø”OØ!&¤ðñ ô �ð	ð 	ð 	ñ 	ô 	ð 	ð 	ð 	ð 	ð 	ð 	øøøð 	ð 	ð 	ð 	ð %×0Ò0Ñ2Ô2Ð2ØŒyÐ$Ø $¤)Ñ+�Ø—9’9 5¤;�9Ñ/Ô/Ð/s   Â+AC?Ã?DÄD)FNNNr®   )Ú__name__Ú
__module__Ú__qualname__r3   r5   r   ÚintÚboolr   Útupler´   ÚTensorrÎ   Ú__classcell__)r¿   s   @r"   r­   r­   &  sµ   ø€ € € € € ØÔ€Eð ØØ04ØØ#ð 2ð  2àð 2ð ð 2ð ð	 2ð ˜U 3¨ 8œ_Ô-ð 2ð  2ð  2ð  2ð  2ð  2ðD0˜Uœ\ð 0¨e¬lð 0ð 0ð 0ð 0ð 0ð 0ð 0ð 0r$   r­   Fc                 óÂ  ‡	— |€g }|                       ¦   «         D �]A\  }}|                     |¦  «         t          |t          j        ¦  «        r­||pg vr§d                     |¦  «        Š	t          ˆ	fd„|pg D ¦   «         ¦  «        sut          ¦   «         5  t          |j	        |j
        |j        du|j        j        |j        j        |j        |j        ¬¦  «        | j        |<   d}ddd¦  «         n# 1 swxY w Y   t%          t'          |                     ¦   «         ¦  «        ¦  «        dk    rt+          ||||||¬¦  «        \  }}|                     d¦  «         �ŒC| |fS )	z%Replace Linear layers with FP8Linear.Nú.c              3   ó    •K  — | ]}|‰v V — Œ	d S )N© )Ú.0ÚkeyÚcurrent_key_name_strs     €r"   ú	<genexpr>z+_replace_with_fp8_linear.<locals>.<genexpr>u  s)   øè è € Ð]Ð]°s�sÐ2Ð2Ð]Ð]Ð]Ð]Ð]Ð]r$   )r¯   rš   r°   r�   r   r»   r&   Tr   )Úhas_been_replacedr)   )Únamed_childrenÚappendÚ
isinstancerµ   ÚLinearÚjoinÚanyr
   r­   r¯   rš   r°   r¸   r�   r   r»   Úweight_block_sizeÚ_modulesr€   ÚlistÚchildrenÚ_replace_with_fp8_linearÚpop)
ÚmodelÚtp_planÚmodules_to_not_convertÚcurrent_key_nameÚquantization_configrß   ÚnameÚmoduleÚ_rÝ   s
            @r"   rê   rê   d  sÖ  ø€ ð ÐØÐà×,Ò,Ñ.Ô.ð !ñ !‰ˆˆfØ×Ò Ñ%Ô%Ð%å�f�bœiÑ(Ô(ð 	-¨TÐ:PÐ:VÐTVÐ-WÐ-WØ#&§8¢8Ð,<Ñ#=Ô#=Ð ÝÐ]Ð]Ð]Ð]Ð?UÐ?[ÐY[Ð]Ñ]Ô]Ñ]Ô]ð -Ý'Ñ)Ô)ð 
-ð 
-Ý+4Ø$*Ô$6Ø%+Ô%8Ø#œ[°Ð4Ø%œ}Ô3Ø$œmÔ1Ø*=Ô*OØ#6Ô#Hð,ñ ,ô ,�E”N 4Ñ(ð )-Ð%ð
-ð 
-ð 
-ñ 
-ô 
-ð 
-ð 
-ð 
-ð 
-ð 
-ð 
-øøøð 
-ð 
-ð 
-ð 
-õ �t�F—O’OÑ%Ô%Ñ&Ô&Ñ'Ô'¨!Ò+Ð+Ý#;ØØØ&Ø Ø#Ø"3ð$ñ $ô $Ñ ˆAÐ ð 	×Ò˜RÑ Ô Ð Ñ àÐ#Ð#Ð#s   ÂAC1Ã1C5	Ã8C5	c                 óü   — |€dgn|}|j         �|                     |j         ¦  «         t          t          |¦  «        ¦  «        }t	          | | j        ||¬¦  «        \  } }|st                               d¦  «         | S )z:Helper function to replace model layers with FP8 versions.NÚlm_head)rí   rî   rð   z�You are loading your model using fp8 but no linear modules were found in your model. Please double check your model architecture.)rî   Úextendrè   Úsetrê   Ú_tp_planÚloggerÚwarning)rì   rî   rð   rß   s       r"   Úreplace_with_fp8_linearrû   ’  s£   € ð -CÐ,J˜i˜[˜[ÐPfÐàÔ1Ð=Ø×%Ò%Ð&9Ô&PÑQÔQÐQÝ!¥#Ð&<Ñ"=Ô"=Ñ>Ô>ÐÝ7ØØ”Ø5Ø/ð	 ñ  ô  Ñ€EÐð ð 
Ý�Šð<ñ	
ô 	
ð 	
ð
 €Lr$   )r%   )NNNNF)NN)'Útypingr   Úutilsr   r   r   r   r3   Útorch.nnrµ   r,   Útriton.languageÚlanguager   r	   rÃ   Ú
accelerater
   Ú
get_loggerrÏ   rù   ÚjitÚ	constexprr#   rÕ   rÒ   rÔ   r8   rw   r   rè   r   r‡   Úcompiler«   rã   r­   rê   rû   rÚ   r$   r"   ú<module>r     s  ðð  Ð Ð Ð Ð Ð à hÐ hÐ hÐ hÐ hÐ hÐ hÐ hÐ hÐ hÐ hÐ hð ÐÑÔð )Ø€L€L€LØÐÐÐÐÐØ€M€M€MØ Ð Ð Ð Ð Ð Ø(Ð(Ð(Ð(Ð(Ð(àÐÑÔð .Ø-Ð-Ð-Ð-Ð-Ð-ð 
ˆÔ	˜HÑ	%Ô	%€ð „ð°b´lð ð ð ñ „ðð
ð 
�”ð 
¨3ð 
¸¸u¼|ÈUÌ\Ð?YÔ9Zð 
ð 
ð 
ð 
ð „ðQ%ð4 ”,ð5Q%ð6 ”,ð7Q%ð8 ”,ð9Q%ð: ”,ð;Q%ð Q%ð Q%ñ „ðQ%ðt !&¤ðMð MØ„|ðMà„|ðMð 	ŒðMð 	Œð	Mð
 �S”	ðMð ”+ðMð „\ðMð Mð Mð Mðb „ð -1Ø %¤ð>#ð >#ØŒ\ð>#àŒlð>#ð ”ð>#ð ”,ð	>#ð
 ˜˜s C˜xœÔ)ð>#ð ”+ð>#ð „\ð>#ð >#ð >#ñ „ð>#ðB;0ð ;0ð ;0ð ;0ð ;0�”	ñ ;0ô ;0ð ;0ð@ ØØØØð+$ð +$ð +$ð +$ð`  Øðð ð ð ð ð r$   