§
    6˜ºh¶  ã                   ó®  — d dl Z d dlZd dlZd dlZd dlZd dlmZ d dlm	Z	 d dl
Z
d dlmZ d dlmZmZ d dlmZmZmZ g d¢Zg d¢Zeez   Zg d¢Zed	gz   Zeez   Zed	gz   Zd
dgZdgez   dgz   ez   d	gz   Z e ee¦  «        h d£z
  ¦  «        Zd„ Z d„ Z!d„ Z"d„ Z#d„ Z$d„ Z%d„ Z&d„ Z'd„ Z(d„ Z)d„ Z*d„ Z+d,dee         fd„Z,d-dej-        deeej.        f         fd„Z/de0dej1        fd „Z2de0fd!„Z3d"„ Z4d.d$„Z5d.d%„Z6e
j7         8                     e5¦   «           e6¦   «         ¬&¦  «        Z9d'e:d(e:fd)„Z;d*eej.        ej<        j=        j        f         dej.        fd+„Z>dS )/é    N)Úknobs)ÚRandomState)ÚOptionalÚUnion)ÚTensorWrapperÚreinterpretÚtype_canonicalisation_dict)Úint8Úint16Úint32Úint64)Úuint8Úuint16Úuint32Úuint64)Úfloat16Úfloat32Úfloat64Úbfloat16Úfloat8_e4m3fnÚfloat8_e5m2Úboolr   >   r   r   r   c                  óJ   — t           j                             dd¦  «        dk    S )NÚTRITON_INTERPRETÚ0Ú1)ÚosÚenvironÚget© ó    ú`/var/www/html/services/rag-service/venv/lib/python3.11/site-packages/triton/_internal_testing.pyÚis_interpreterr#      s   € ÝŒ:�>Š>Ð,¨cÑ2Ô2°cÒ9Ð9r!   c                  ór   — t          ¦   «         rd S t          j        j        j                             ¦   «         S ©N)r#   ÚtritonÚruntimeÚdriverÚactiveÚget_current_targetr    r!   r"   r*   r*      s/   € ÝÑÔð ØˆtÝŒ>Ô Ô'×:Ò:Ñ<Ô<Ð<r!   c                  ó<   — t          ¦   «         } | €dn
| j        dk    S )NFÚcuda©r*   Úbackend©Útargets    r"   Úis_cudar1   $   s"   € ÝÑ!Ô!€FØ�Nˆ5ˆ5¨¬¸&Ò(@Ð@r!   c                  ón   — t          ¦   «         o't          j                             ¦   «         d         dk    S )Nr   é	   )r1   Útorchr,   Úget_device_capabilityr    r!   r"   Ú	is_hopperr6   )   s)   € Ý‰9Œ9ÐC�œ×9Ò9Ñ;Ô;¸AÔ>À!ÒCÐCr!   c                  ó<   — t          ¦   «         } | €dn
| j        dk    S )NFÚhipr-   r/   s    r"   Úis_hipr9   -   ó"   € ÝÑ!Ô!€FØ�Nˆ5ˆ5¨¬¸%Ò(?Ð?r!   c                  óR   — t          ¦   «         } | d uo| j        dk    o
| j        dk    S )Nr8   Úgfx90a©r*   r.   Úarchr/   s    r"   Úis_hip_cdna2r?   2   ó0   € ÝÑ!Ô!€FØ˜ÐÐU &¤.°EÒ"9ÐU¸f¼kÈXÒ>UÐUr!   c                  óR   — t          ¦   «         } | d uo| j        dk    o
| j        dk    S )Nr8   Úgfx942r=   r/   s    r"   Úis_hip_cdna3rC   7   r@   r!   c                  óR   — t          ¦   «         } | d uo| j        dk    o
| j        dk    S )Nr8   Úgfx950r=   r/   s    r"   Úis_hip_cdna4rF   <   r@   r!   c                  óv   — t          ¦   «         } t          | j        ¦  «         | d uo| j        dk    od| j        v S )Nr8   Úgfx12)r*   Úprintr>   r.   r/   s    r"   Úis_hip_gfx12rJ   A   s@   € ÝÑ!Ô!€FÝ	ˆ&Œ+ÑÔÐØ˜ÐÐT &¤.°EÒ"9ÐT¸gÈÌÐ>TÐTr!   c                  óV   — t          ¦   «         pt          ¦   «         pt          ¦   «         S r%   )r?   rC   rF   r    r!   r"   Úis_hip_cdnarL   G   s   € Ý‰>Œ>Ð=�\™^œ^Ð=­|©~¬~Ð=r!   c                  ó<   — t          ¦   «         } | €dn
| j        dk    S )NFÚxpur-   r/   s    r"   Úis_xpurO   K   r:   r!   c                  óN   — t          ¦   «         } | €dnt          | j        ¦  «        S )NÚ )r*   Ústrr>   r/   s    r"   Úget_archrS   P   s%   € ÝÑ!Ô!€FØ�ˆ2ˆ2¥S¨¬Ñ%5Ô%5Ð5r!   Úrsc                 ó¤  — t          | t          ¦  «        r| f} |€t          d¬¦  «        }|t          t          z   v rœt          j        t          t
          |¦  «        ¦  «        }|€|j        nt          ||j        ¦  «        }|€|j	        nt          ||j	        ¦  «        }t          t
          |¦  «        }| 
                    ||| |¬¦  «        }d||dk    <   |S |r)d|v r%| 
                    dd	| t
          j        ¬¦  «        }|S |t          v r*|                     dd| ¦  «                             |¦  «        S |d
k    re|                     dd| ¦  «                             d¦  «                             d¦  «        t          j        d¦  «        z                       d¦  «        S |dv r|                     dd| ¦  «        dk    S t#          d|› �¦  «        ‚)zp
    Override `rs` if you're calling this function twice and don't want the same
    result for both calls.
    Né   )Úseed)Údtypeé   r   Úfloat8é   é(   r   r   r   l     þ )r   Úint1Úbool_g        zUnknown dtype )Ú
isinstanceÚintr   Ú
int_dtypesÚuint_dtypesÚnpÚiinfoÚgetattrÚminÚmaxÚrandintr
   Úfloat_dtypesÚnormalÚastypeÚviewr   ÚRuntimeError)ÚshapeÚ	dtype_strrT   ÚlowÚhighrd   rX   Úxs           r"   Únumpy_randomrs   U   sÆ  € õ
 �%�ÑÔð Ø�	ˆØ	€zÝ˜bÐ!Ñ!Ô!ˆØ•J¥Ñ,Ð,Ð,Ý”�¥ YÑ/Ô/Ñ0Ô0ˆØ˜;ˆeŒiˆi­C°°U´YÑ,?Ô,?ˆØ ˜LˆuŒyˆy­c°$¸¼	Ñ.BÔ.BˆÝ�˜IÑ&Ô&ˆØ�JŠJ�s˜D %¨uˆJÑ5Ô5ˆØˆˆ!ˆqŠ&‰	ØˆØ	ð 
9�x 9Ð,Ð,Ø�JŠJ�r˜2˜u­B¬GˆJÑ4Ô4ˆØˆØ	•lÐ	"Ð	"Ø�yŠy˜˜A˜uÑ%Ô%×,Ò,¨YÑ7Ô7Ð7Ø	�jÒ	 Ð	 Ø—	’	˜!˜Q Ñ&Ô&×-Ò-¨iÑ8Ô8×=Ò=¸hÑGÔGÍ"Ì)ÐT^ÑJ_ÔJ_Ñ_×eÒeÐfoÑpÔpÐpØ	Ð/Ð	/Ð	/Ø�yŠy˜˜A˜uÑ%Ô%¨Ò+Ð+åÐ7¨IÐ7Ð7Ñ8Ô8Ð8r!   rr   Úreturnc                 ó"  — | j         j        }|t          v rt|                     d¦  «        }|                      t          t          |¦  «        ¦  «        }t          t          j	        ||¬¦  «        t          t          |¦  «        ¦  «        S |r;d|v r7t          t          j	        | |¬¦  «        t          t          |¦  «        ¦  «        S |dk    r.|dk    r(t          j	        | |¬¦  «                             ¦   «         S t          j	        | |¬¦  «        S )zÖ
    Note: We need dst_type because the type of x can be different from dst_type.
          For example: x is of type `float32`, dst_type is `bfloat16`.
          If dst_type is None, we infer dst_type from x.
    Úu)ÚdevicerZ   r   r   )rX   Únamerb   Úlstriprk   re   rc   r   r4   ÚtensorÚtlr   )rr   rw   Údst_typeÚtÚsigned_type_nameÚx_signeds         r"   Ú	to_tritonr€   s   sð   € ð 	
ŒŒ€AØ�KÐÐØŸ8š8 C™=œ=ÐØ—8’8�G¥BÐ(8Ñ9Ô9Ñ:Ô:ˆÝ�5œ<¨¸Ð@Ñ@Ô@Å'Í"ÈaÁ.Ä.ÑQÔQÐQàð 	V˜ HÐ,Ð,Ý�uœ|¨A°fÐ=Ñ=Ô=½wÅrÈ8Ñ?TÔ?TÑUÔUÐUØ�	Š>ˆ>˜h¨*Ò4Ð4Ý”< ¨&Ð1Ñ1Ô1×:Ò:Ñ<Ô<Ð<ÝŒ|˜A fÐ-Ñ-Ô-Ð-r!   c                 ó@   — t          j        t          |          ¦  «        S r%   )r{   Ú	str_to_tyr	   ©rr   s    r"   Ústr_to_triton_dtyper„   †   s   € ÝŒ<Õ2°1Ô5Ñ6Ô6Ð6r!   c                 ó.  — t          | t          j        j        ¦  «        r| j        S t          | t
          j        ¦  «        r7t          j        dt          | ¦  «        ¦  «        }| 	                    d¦  «        S t          dt          | ¦  «        › �¦  «        ‚)Nz^torch\.(\w+)$rY   znot a triton or torch dtype: )r_   r&   ÚlanguagerX   rx   r4   ÚreÚmatchrR   ÚgroupÚ	TypeErrorÚtype)rX   Úms     r"   Útorch_dtype_namer�   Š   s|   € Ý�%�œÔ.Ñ/Ô/ð GØŒzÐÝ	�E�5œ;Ñ	'Ô	'ð GåŒHÐ&­¨E©
¬
Ñ3Ô3ˆØ�wŠw�q‰zŒzÐåÐE½¸U¹¼ÐEÐEÑFÔFÐFr!   c                 ó,  — t          | t          ¦  «        rc| j                             ¦   «                              ¦   «                              t          t          t          | j	        ¦  «        ¦  «        ¦  «        S t          | t          j        ¦  «        rq| j	        t          j        u r8|                      ¦   «                              ¦   «                              ¦   «         S |                      ¦   «                              ¦   «         S t          d| › �¦  «        ‚)Nz Not a triton-compatible tensor: )r_   r   ÚbaseÚcpuÚnumpyrk   re   rc   r�   rX   r4   ÚTensorr   ÚfloatÚ
ValueErrorrƒ   s    r"   Úto_numpyr•   •   sÇ   € Ý�!•]Ñ#Ô#ð AØŒv�zŠz‰|Œ|×!Ò!Ñ#Ô#×*Ò*­7µ2Õ7GÈÌÑ7PÔ7PÑ+QÔ+QÑRÔRÐRÝ	�A•u”|Ñ	$Ô	$ð AØŒ7•e”nÐ$Ð$Ø—5’5‘7”7—=’=‘?”?×(Ò(Ñ*Ô*Ð*Ø�uŠu‰wŒw�}Š}‰ŒÐåÐ?¸AÐ?Ð?Ñ@Ô@Ð@r!   Fc                 óz  — t          ¦   «         rdS t          ¦   «         sdS t          j        j        j        }| rdnd}t          t          t          | 	                    d¦  «        ¦  «        ¦  «        }t          |¦  «        dk    s
J |¦   «         ‚t          j                             ¦   «         d         dk    o||k    S )	NTF)é   r   )r—   é   ú.é   r   r3   )r#   r1   r   ÚnvidiaÚptxasÚversionÚtupleÚmapr`   ÚsplitÚlenr4   r,   r5   )Ú
byval_onlyÚcuda_versionÚmin_cuda_versionÚcuda_version_tuples       r"   Úsupports_tmar¦       sµ   € ÝÑÔð ØˆtÝ‰9Œ9ð ØˆuÝ”<Ô%Ô-€LØ",Ð9�w�w°'ÐÝ�s¥3¨×(:Ò(:¸3Ñ(?Ô(?Ñ@Ô@ÑAÔAÐÝÐ!Ñ"Ô" aÒ'Ð'Ð'Ð);Ñ'Ô'Ð'ÝŒ:×+Ò+Ñ-Ô-¨aÔ0°AÒ5Ð`Ð:LÐP`Ò:`Ð`r!   c                 ó   — | rdS dS )NzURequires __grid_constant__ TMA support (NVIDIA Hopper or higher, CUDA 12.0 or higher)zLRequires advanced TMA support (NVIDIA Hopper or higher, CUDA 12.3 or higher)r    )r¢   s    r"   Útma_skip_msgr¨   ¬   s   € Øð ^ØfÐfà]Ð]r!   )ÚreasonÚsizeÚalignc                 óD   — t          j        | t           j        d¬¦  «        S )Nr,   )rX   rw   )r4   Úemptyr
   )rª   r«   Ú_s      r"   Údefault_alloc_fnr¯   ¶   s   € ÝŒ;�t¥5¤:°fÐ=Ñ=Ô=Ð=r!   r}   c                 ó\   — t          | t          j        j        j        ¦  «        r| j        S | S r%   )r_   r&   r'   Újitr   r�   )r}   s    r"   Úunwrap_tensorr²   º   s(   € Ý�!•V”^Ô'Ô5Ñ6Ô6ð ØŒvˆØ€Hr!   )NNNr%   )F)?r   r‡   r‘   rc   r4   r&   Útriton.languager†   r{   r   ÚpytestÚnumpy.randomr   Útypingr   r   Útriton.runtime.jitr   r   r	   ra   rb   Úintegral_dtypesri   Úfloat_dtypes_with_bfloat16ÚdtypesÚdtypes_with_bfloat16Útorch_float8_dtypesÚtorch_dtypesÚsortedÚsetÚ
tma_dtypesr#   r*   r1   r6   r9   r?   rC   rF   rJ   rL   rO   rS   rs   Úndarrayr’   r€   rR   rX   r„   r�   r•   r¦   r¨   ÚmarkÚskipifÚrequires_tmar`   r¯   r'   r±   r²   r    r!   r"   ú<module>rÅ      s„  ðØ 	€	€	€	Ø 	€	€	€	Ø Ð Ð Ð Ø €€€Ø €€€Ø Ð Ð Ð Ð Ð Ø Ð Ð Ð Ð Ð Ø €€€à $Ð $Ð $Ð $Ð $Ð $Ø "Ð "Ð "Ð "Ð "Ð "Ð "Ð "Ø UÐ UÐ UÐ UÐ UÐ UÐ UÐ UÐ UÐ Uà0Ð0Ð0€
Ø5Ð5Ð5€Ø˜{Ñ*€Ø0Ð0Ð0€Ø)¨Z¨LÑ8Ð Ø	˜<Ñ	'€Ø  Ñ,Ð Ø&¨Ð6Ð Øˆx˜*Ñ$¨ yÑ0°<Ñ?À:À,ÑN€ØˆV�C�CÐ,Ñ-Ô-Ð0NÐ0NÐ0NÑNÑOÔO€
ð:ð :ð :ð=ð =ð =ðAð Að Að
Dð Dð Dð@ð @ð @ð
Vð Vð Vð
Vð Vð Vð
Vð Vð Vð
Uð Uð Uð>ð >ð >ð@ð @ð @ð
6ð 6ð 6ð
9ð 9 x°Ô'<ð 9ð 9ð 9ð 9ð<.ð .�”ð .°u¸]ÈEÌLÐ=XÔ7Yð .ð .ð .ð .ð&7˜3ð 7 2¤8ð 7ð 7ð 7ð 7ðG˜sð Gð Gð Gð GðAð Að Að	að 	að 	að 	að^ð ^ð ^ð ^ð Œ{×!Ò! l l¡n¤nÐ"4¸\¸\¹^¼^Ð!ÑLÔL€ð>˜3ð > sð >ð >ð >ð >ð�U˜5œ<¨¬Ô);Ô)IÐIÔJð ÈuÌ|ð ð ð ð ð ð r!   