Ë
    4^(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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jT                  deeejV                  f   fd„Z,de-dej\                  fd„Z/de-fd„Z0d „ Z1d&d!„Z2d&d"„Z3ejh                  jk                   e2«         e3«       ¬#«      Z6y)'é    N)Ú_path_to_binary)Ú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   c                  óH   — t         j                  j                  dd«      dk(  S )NÚTRITON_INTERPRETÚ0Ú1)ÚosÚenvironÚget© ó    úV/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/triton/_internal_testing.pyÚis_interpreterr#      s   € Ü�:‰:�>‰>Ð,¨cÓ2°cÑ9Ð9r!   c                  ó|   — t        «       ry t        j                  j                  j                  j                  «       S ©N)r#   ÚtritonÚruntimeÚdriverÚactiveÚget_current_targetr    r!   r"   r*   r*      s*   € ÜÔØÜ�>‰>× Ñ ×'Ñ'×:Ñ:Ó<Ð<r!   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚcuda©r*   Úbackend©Útargets    r"   Úis_cudar1   #   s"   € ÜÓ!€FØ�Nˆ5Ð@¨¯©¸&Ñ(@Ð@r!   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk\  S )Nr   é	   )r1   Útorchr,   Úget_device_capabilityr    r!   r"   Ú	is_hopperr6   (   s&   € Ü‹9ÒCœŸ™×9Ñ9Ó;¸AÑ>À!ÑCÐCr!   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚhipr-   r/   s    r"   Úis_hipr9   ,   ó"   € ÜÓ!€FØ�Nˆ5Ð?¨¯©¸%Ñ(?Ð?r!   c                  óX   — t        «       } | �| j                  dk7  ry| j                  dk(  S )Nr8   FÚgfx90a©r*   r.   Úarchr/   s    r"   Úis_hip_mi200r?   1   s,   € ÜÓ!€FØ€~˜Ÿ™¨5Ò0ØØ�;‰;˜(Ñ"Ð"r!   c                  óV   — t        «       } | �| j                  dk7  ry| j                  dv S )Nr8   F)Úgfx940Úgfx941Úgfx942r=   r/   s    r"   Úis_hip_mi300rD   8   s-   € ÜÓ!€FØ€~˜Ÿ™¨5Ò0ØØ�;‰;Ð8Ð8Ð8r!   c                  óV   — t        «       } | �| j                  dk7  ry| j                  dv S )Nr8   FÚgfx950r=   r/   s    r"   Úis_hip_mi350rG   ?   s,   € ÜÓ!€FØ€~˜Ÿ™¨5Ò0ØØ�;‰;˜8Ð$Ð$r!   c                  óF   — t        «       xs t        «       xs
 t        «       S r%   )r?   rD   rG   r    r!   r"   Úis_hip_cdnarI   F   s   € Ü‹>Ò=œ\›^Ò=¬|«~Ð=r!   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚxpur-   r/   s    r"   Úis_xpurL   J   r:   r!   c                  óH   — t        «       } | €dS t        | j                  «      S )NÚ )r*   Ústrr>   r/   s    r"   Úget_archrP   O   s"   € ÜÓ!€FØ�ˆ2Ð5¤S¨¯©Ó%5Ð5r!   Úrsc                 óF  — 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
        |«      }|j                  ||| |¬«      }d||dk(  <   |S |r)d|v r%|j                  dd| t
        j                  ¬«      }|S |t        v r"|j                  dd| «      j                  |«      S |d	k(  rV|j                  dd| «      j                  d
«      j                  d«      t        j                   d«      z  j                  d
«      S |dv r|j                  dd| «      dkD  S t#        d|› �«      ‚)zp
    Override `rs` if you're calling this function twice and don't want the same
    result for both calls.
    é   )Ú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_strrQ   ÚlowÚhighra   rU   Úxs           r"   Únumpy_randomrp   T   s|  € ô
 �%œÔØ�	ˆØ	€zÜ˜bÔ!ˆØ”J¤Ñ,Ñ,Ü—‘œ¤ YÓ/Ó0ˆØ˜;ˆe�iŠi¬C°°U·Y±YÓ,?ˆØ ˜Lˆu�yŠy¬c°$¸¿	¹	Ó.BˆÜœ˜IÓ&ˆØ�J‰J�s˜D %¨uˆJÓ5ˆØˆˆ!ˆq‰&‰	ØˆÙ	�x 9Ñ,Ø�J‰J�r˜2˜u¬B¯G©GˆJÓ4ˆØˆØ	”lÑ	"Ø�y‰y˜˜A˜uÓ%×,Ñ,¨YÓ7Ð7Ø	�jÒ	 Ø—	‘	˜!˜Q Ó&×-Ñ-¨iÓ8×=Ñ=¸hÓGÌ"Ï)É)ÐT^ÓJ_Ñ_×eÑeÐfoÓpÐpØ	Ð/Ñ	/Ø�y‰y˜˜A˜uÓ%¨Ñ+Ð+ä˜^¨I¨;Ð7Ó8Ð8r!   ro   Úreturnc                 óò  — | j                   j                  }|t        v r_|j                  d«      }| j	                  t        t        |«      «      }t        t        j                  ||¬«      t        t        |«      «      S |r3d|v r/t        t        j                  | |¬«      t        t        |«      «      S |dk(  r*|dk(  r%t        j                  | |¬«      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)ÚdevicerW   r   r   )rU   Únamer_   Úlstriprh   rb   r`   r   r4   ÚtensorÚtlr   )ro   rt   Údst_typeÚtÚsigned_type_nameÚx_signeds         r"   Ú	to_tritonr}   r   s½   € ð 	
�‰�‰€AØŒKÑØŸ8™8 C›=ÐØ—8‘8œG¤BÐ(8Ó9Ó:ˆÜœ5Ÿ<™<¨¸Ô@Ä'Ì"ÈaÃ.ÓQÐQá˜ HÑ,ÜœuŸ|™|¨A°fÔ=¼wÄrÈ8Ó?TÓUÐUØ�	Š>˜h¨*Ò4Ü—<‘< ¨&Ô1×:Ñ:Ó<Ð<Ü�|‰|˜A fÔ-Ð-r!   c                 ó:   — t        j                  t        |    «      S r%   )rx   Ú	str_to_tyr	   ©ro   s    r"   Ústr_to_triton_dtyper�   …   s   € Ü�<‰<Ô2°1Ñ5Ó6Ð6r!   c                 ó$  — t        | t        j                  j                  «      r| j                  S t        | t
        j                  «      r0t        j                  dt        | «      «      }|j                  d«      S t        dt        | «      › �«      ‚)Nz^torch\.(\w+)$rV   znot a triton or torch dtype: )r\   r&   ÚlanguagerU   ru   r4   ÚreÚmatchrO   ÚgroupÚ	TypeErrorÚtype)rU   Úms     r"   Útorch_dtype_namerŠ   ‰   sh   € Ü�%œŸ™×.Ñ.Ô/Ø�z‰zÐÜ	�Eœ5Ÿ;™;Ô	'ä�H‰HÐ&¬¨E«
Ó3ˆØ�w‰w�q‹zÐäÐ7¼¸U»°}ÐEÓFÐFr!   c                 óî  — t        | t        «      rX| j                  j                  «       j	                  «       j                  t        t        t        | j                  «      «      «      S t        | t        j                  «      rf| j                  t        j                  u r,| j                  «       j                  «       j	                  «       S | j                  «       j	                  «       S t        d| › �«      ‚)Nz Not a triton-compatible tensor: )r\   r   ÚbaseÚcpuÚnumpyrh   rb   r`   rŠ   rU   r4   ÚTensorr   ÚfloatÚ
ValueErrorr€   s    r"   Úto_numpyr’   ”   sŸ   € Ü�!”]Ô#Ø�v‰v�z‰z‹|×!Ñ!Ó#×*Ñ*¬7´2Ô7GÈÏÉÓ7PÓ+QÓRÐRÜ	�A”u—|‘|Ô	$Ø�7‰7”e—n‘nÑ$Ø—5‘5“7—=‘=“?×(Ñ(Ó*Ð*Ø�u‰u‹w�}‰}‹ÐäÐ;¸A¸3Ð?Ó@Ð@r!   c                 ó&  — t        «       ryt        «       syt        d«      \  }}| rdnd}t        t	        t
        |j                  d«      «      «      }t        |«      dk(  sJ |«       ‚t        j                  j                  «       d   d	k\  xr ||k\  S )
NTFÚptxas)é   r   )r•   é   ú.é   r   r3   )r#   r1   r   ÚtupleÚmapr]   ÚsplitÚlenr4   r,   r5   )Ú
byval_onlyÚ_Úcuda_versionÚmin_cuda_versionÚcuda_version_tuples        r"   Úsupports_tmar¢   Ÿ   s‹   € ÜÔØÜŒ9ØÜ% gÓ.�O€A€|Ù",‘w°'ÐÜœs¤3¨×(:Ñ(:¸3Ó(?Ó@ÓAÐÜÐ!Ó" aÒ'Ð;Ð);Ó;Ð'Ü�:‰:×+Ñ+Ó-¨aÑ0°AÑ5Ò`Ð:LÐP`Ñ:`Ð`r!   c                 ó
   — | ryy)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à]r!   )Úreason)NNNr%   )F)7r   r„   rŽ   r`   r4   r&   Útriton.languagerƒ   rx   Útriton.backends.nvidia.compilerr   ÚpytestÚnumpy.randomr   Útypingr   r   Útriton.runtime.jitr   r   r	   r^   r_   Úintegral_dtypesrf   Úfloat_dtypes_with_bfloat16ÚdtypesÚdtypes_with_bfloat16Útorch_float8_dtypesÚtorch_dtypesr#   r*   r1   r6   r9   r?   rD   rG   rI   rL   rP   rp   Úndarrayr�   r}   rO   rU   r�   rŠ   r’   r¢   r¤   ÚmarkÚskipifÚrequires_tmar    r!   r"   ú<module>r¶      sZ  ðÛ 	Û 	Û Û Û Ý Ý ;Û å $ß "ß UÑ Uâ0€
Ú5€Ø˜{Ñ*€Ú0€Ø)¨Z¨LÑ8Ð Ø	˜<Ñ	'€Ø  Ñ,Ð Ø&¨Ð6Ð Øˆx˜*Ñ$¨ yÑ0°<Ñ?À:À,ÑN€ò:ò=òAò
Dò@ò
#ò9ò%ò>ò@ò
6ñ
9 x°Ñ'<ó 9ñ<.�—‘ð .°u¸]ÈEÏLÉLÐ=XÑ7Yó .ð&7˜3ð 7 2§8¡8ó 7ðG˜só GòAó	aó^ð �{‰{×!Ñ!¡l£nÐ"4¹\»^Ð!ÓL�r!   