Ë
    ÍÍ:j¶  ã                   óÔ  — 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jZ                  deeej\                  f   fd„Z/de0dejb                  fd „Z2de0fd!„Z3d"„ Z4d-d#„Z5d-d$„Z6e
jn                  jq                   e5«         e6«       ¬%«      Z9d&e:d'e:fd(„Z;d)eej\                  ejx                  jz                  j"                  f   dej\                  fd*„Z>y).é    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                  óH   — t         j                  j                  dd«      dk(  S )NÚTRITON_INTERPRETÚ0Ú1)ÚosÚenvironÚget© ó    úm/home/mcse/projects/srt_converter/srt-converter-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                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )Nr8   Úgfx90a©r*   r.   Úarchr/   s    r"   Úis_hip_cdna2r?   2   ó1   € ÜÓ!€FØ˜ÐÒU &§.¡.°EÑ"9ÒU¸f¿k¹kÈXÑ>UÐUr!   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )Nr8   Úgfx942r=   r/   s    r"   Úis_hip_cdna3rC   7   r@   r!   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )Nr8   Úgfx950r=   r/   s    r"   Úis_hip_cdna4rF   <   r@   r!   c                  óŠ   — t        «       } t        | j                  «       | d uxr | j                  dk(  xr 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                  óF   — t        «       xs t        «       xs
 t        «       S r%   )r?   rC   rF   r    r!   r"   Úis_hip_cdnarL   G   s   € Ü‹>Ò=œ\›^Ò=¬|«~Ð=r!   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚxpur-   r/   s    r"   Úis_xpurO   K   r:   r!   c                  óH   — t        «       } | €dS t        | j                  «      S )NÚ )r*   Ústrr>   r/   s    r"   Úget_archrS   P   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_strrT   ÚlowÚhighrd   rX   Úxs           r"   Únumpy_randomrs   U   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!   rr   Ú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)Ú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Ó:ˆÜœ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%   )r{   Ú	str_to_tyr	   ©rr   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+)$rY   znot a triton or torch dtype: )r_   r&   ÚlanguagerX   rx   r4   ÚreÚmatchrR   ÚgroupÚ	TypeErrorÚtype)rX   Ú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Únumpyrk   re   rc   r�   rX   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                 óR  — t        «       ryt        «       syt        j                  j                  j
                  }| 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)é   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’   € ÜÔØÜŒ9ØÜ—<‘<×%Ñ%×-Ñ-€LÙ",‘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ÚsizeÚalignc                 óN   — 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                 óz   — t        | t        j                  j                  j                  «      r| j
                  S | S r%   )r_   r&   r'   Újitr   r�   )r}   s    r"   Úunwrap_tensorr²   º   s*   € Ü�!”V—^‘^×'Ñ'×5Ñ5Ô6Ø�v‰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â0€
Ú5€Ø˜{Ñ*€Ú0€Ø)¨Z¨LÑ8Ð Ø	˜<Ñ	'€Ø  Ñ,Ð Ø&¨Ð6Ð Øˆx˜*Ñ$¨ yÑ0°<Ñ?À:À,ÑN€Ù‘CÐ,Ó-Ò0NÑNÓO€
ò:ò=òAò
Dò@ò
Vò
Vò
Vò
Uò>ò@ò
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€ð>˜3ð > só >ð�U˜5Ÿ<™<¨¯©×);Ñ);×)IÑ)IÐIÑJð ÈuÏ|É|ô r!   