Ë
    .²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
mZmZmZ d dlZd dl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„ Z-d„ Z.d„ Z/d„ Z0d„ Z1d „ Z2d!„ Z3d"„ Z4d#„ Z5d7d$ee   fd%„Z6d8d&ejn                  d'eeejp                  f   fd(„Z9d&e:d'ejv                  fd)„Z<d'e:fd*„Z=d+„ Z>d9d,„Z?d-„ Z@d9d.„ZAej„                  j‡                   e?«         eA«       ¬/«      ZDd0eEd1eEfd2„ZFd3eejp                  ejŽ                  j�                  j$                  f   d'ejp                  fd4„ZId8d5eee:      fd6„ZJy):é    N©Úknobs)ÚOptionalÚSetÚUnion)ÚRandomState)Ú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© ó    úW/var/www/html/ia-enhanced/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_cudar3   $   s"   € ÜÓ!€FØ�Nˆ5Ð@¨¯©¸&Ñ(@Ð@r#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk\  S )Nr   é   ©r3   Útorchr.   Úget_device_capabilityr"   r#   r$   Úis_ampere_or_newerr9   )   ó&   € Ü‹9ÒCœŸ™×9Ñ9Ó;¸AÑ>À!ÑCÐCr#   c                  ó`   — t        «       xr# t        j                  j                  «       d   dv S )Nr   )é
   é   r6   r"   r#   r$   Úis_blackwellr>   -   s&   € Ü‹9ÒJœŸ™×9Ñ9Ó;¸AÑ>À(ÐJÐJr#   c                  ób   — t        «       xr$ t        j                  j                  «       dd dk(  S )Nr   é   )r<   é   r6   r"   r#   r$   Úis_blackwell_ultrarB   1   s(   € Ü‹9ÒKœŸ™×9Ñ9Ó;¸A¸aÐ@ÀGÑKÐKr#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk\  S ©Nr   é	   r6   r"   r#   r$   Úis_hopper_or_newerrF   5   r:   r#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk(  S rD   r6   r"   r#   r$   Ú	is_hopperrH   9   r:   r#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk(  S )Nr   é   r6   r"   r#   r$   Úis_sm12xrK   =   s&   € Ü‹9ÒDœŸ™×9Ñ9Ó;¸AÑ>À"ÑDÐDr#   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚhipr/   r1   s    r$   Úis_hiprN   A   ó"   € ÜÓ!€FØ�Nˆ5Ð?¨¯©¸%Ñ(?Ð?r#   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )NrM   Úgfx90a©r,   r0   Úarchr1   s    r$   Úis_hip_cdna2rT   F   ó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 )NrM   Úgfx942rR   r1   s    r$   Úis_hip_cdna3rX   K   rU   r#   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )NrM   Úgfx950rR   r1   s    r$   Úis_hip_cdna4r[   P   rU   r#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrM   Úgfx11rR   r1   s    r$   Úis_hip_rdna3r^   U   ó1   € ÜÓ!€FØ˜ÐÒT &§.¡.°EÑ"9ÒT¸gÈÏÉÐ>TÐTr#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrM   Úgfx12rR   r1   s    r$   Úis_hip_rdna4rb   Z   r_   r#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrM   Úgfx1250rR   r1   s    r$   Úis_hip_gfx1250re   _   s1   € ÜÓ!€FØ˜ÐÒV &§.¡.°EÑ"9ÒV¸iÈ6Ï;É;Ð>VÐVr#   c                  óF   — t        «       xs t        «       xs
 t        «       S r'   )rT   rX   r[   r"   r#   r$   Úis_hip_cdnarg   d   s   € Ü‹>Ò=œ\›^Ò=¬|«~Ð=r#   c                  ó.   — t        «       xs
 t        «       S r'   )r^   rb   r"   r#   r$   Úis_hip_rdnari   h   s   € Ü‹>Ò+œ\›^Ð+r#   c                  ó   — t        «       rdS dS )Ni € i   )r[   r"   r#   r$   Úget_hip_lds_sizerk   l   s   € Ü!”^ˆ6Ð.¨Ð.r#   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚxpur/   r1   s    r$   Úis_xpurn   p   rO   r#   c                  óH   — t        «       } | €dS t        | j                  «      S )NÚ )r,   ÚstrrS   r1   s    r$   Úget_archrr   u   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_strrs   ÚlowÚhighrƒ   rw   Úxs           r$   Únumpy_randomr’   z   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#   r‘   Ú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)Údevicery   r   r   )rw   Únamer�   ÚlstriprŠ   r„   r‚   r
   r7   ÚtensorÚtlr   )r‘   r–   Údst_typeÚtÚsigned_type_nameÚx_signeds         r$   Ú	to_tritonrŸ   ˜   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        |    d «      S r'   )rš   Ú	str_to_tyr   ©r‘   s    r$   Ústr_to_triton_dtyper£   «   s   € Ü�<‰<Ô2°1Ñ5°tÓ<Ð<r#   c                 ó$  — t        | t        j                  j                  «      r| j                  S t        | t
        j                  «      r0t        j                  dt        | «      «      }|j                  d«      S t        dt        | «      › �«      ‚)Nz^torch\.(\w+)$rx   znot a triton or torch dtype: )r~   r(   Úlanguagerw   r—   r7   ÚreÚmatchrq   ÚgroupÚ	TypeErrorÚtype)rw   Ú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ÚnumpyrŠ   r„   r‚   r¬   rw   r7   Ú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)rJ   r   )rJ   rA   ú.r@   r   rE   )r%   r3   r   ÚnvidiaÚptxasÚversionÚtupleÚmapr   ÚsplitÚlenr7   r.   r8   )Ú
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                  óv   — t        «       ryt        «       syt        j                  j	                  «       d   dk\  S )NTFr   rE   )r%   r3   r7   r.   r8   r"   r#   r$   Úsupports_wsrÄ   Ñ   s0   € ÜÔØÜŒ9ØÜ�:‰:×+Ñ+Ó-¨aÑ0°AÑ5Ð5r#   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.   )rw   r–   )r7   Ú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#   Úskipped_attrc                 óX  ‡‡‡‡‡	— ddl mŠ | €
t        «       } t        j                  «       Š‰j
                  j                  «       D ��ci c]1  \  }}t        |‰j                  «      r|‰j                  k7  r|| vr||“Œ3 c}}Šg Š‰j                  Š	ˆˆˆˆfd„}ˆˆˆˆˆ	fd„}||fS c c}}w )Nr   r   c                  ó€  •— ‰j                  «       D ]¢  \  } }t        ‰| |j                  «       j                  «       «       |j                  j                  «       D ]W  }|j                  t        j                  v r‰j                  |j                  d¬«       Œ=‰j                  |j                  «       ŒY Œ¤ d‰_        ‰S )NF)ÚraisingT)ÚitemsÚsetattrÚcopyÚresetÚknob_descriptorsÚvaluesÚkeyr   r    ÚdelenvÚappendÚpropagate_env)r—   ÚknobsetÚknobÚenv_to_unsetr   Ú	knobs_mapÚmonkeypatchs      €€€€r$   Úfresh_functionz)_fresh_knobs_impl.<locals>.fresh_function  sž   ø€ à&Ÿ_™_Ó.ò 	2‰MˆD�'Ü�E˜4 §¡£×!5Ñ!5Ó!7Ô8Ø×0Ñ0×7Ñ7Ó9ò 2�Ø—8‘8œrŸz™zÑ)Ø×&Ñ& t§x¡x¸Ð&Õ?à ×'Ñ'¨¯©Õ1ñ	2ð	2ð #ˆÔØˆr#   c                  óÔ   •— ‰j                  «       D ]  \  } }t        ‰| |«       Œ ‰j                  «        ‰D ]&  }|t        j                  v sŒt        j                  |= Œ( ‰‰_        y r'   )rÕ   rÖ   Úundor   r    rÞ   )r—   rß   Úkrá   r   râ   rã   Úprev_propagate_envs      €€€€€r$   Úreset_functionz)_fresh_knobs_impl.<locals>.reset_function  sf   ø€ Ø&Ÿ_™_Ó.ò 	*‰MˆD�'Ü�E˜4 Õ)ð	*ð 	×ÑÔØò 	"ˆAØ”B—J‘JŠÜ—J‘J˜q‘Mð	"ð 1ˆÕr#   )
r(   r   ÚsetÚpytestÚMonkeyPatchÚ__dict__rÕ   r~   Ú
base_knobsrÞ   )
rÑ   r—   rß   rä   ré   rá   r   râ   rã   rè   s
        @@@@@r$   Ú_fresh_knobs_implrï   í   s«   ü€ ÝàÐÜ“uˆä×$Ñ$Ó&€Kð #Ÿ^™^×1Ñ1Ó3÷áˆD�'Ü�g˜u×/Ñ/Ô0°WÀ×@PÑ@PÒ5PÐUYÐamÑUmð 	ˆg‰ó€Ið €LØ×,Ñ,Ð÷
÷	1ð 	1ð ˜>Ð)Ð)ùóGs   Á	6B&)NNNr'   )F)Kr   r¦   r°   r‚   r7   r(   Útriton.languager¥   rš   r   Útypingr   r   r   rë   Únumpy.randomr   Útriton.runtime.jitr	   r
   r   r€   r�   Úintegral_dtypesrˆ   Úfloat_dtypes_with_bfloat16ÚdtypesÚdtypes_with_bfloat16Útorch_float8_dtypesÚtorch_dtypesÚsortedrê   Ú
tma_dtypesr%   r,   r3   r9   r>   rB   rF   rH   rK   rN   rT   rX   r[   r^   rb   re   rg   ri   rk   rn   rr   r’   Úndarrayr±   rŸ   rq   rw   r£   r¬   r´   rÂ   rÄ   rÆ   ÚmarkÚskipifÚrequires_tmar   rÍ   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òKòLòDòDòEò@ò
Vò
Vò
Vò
Uò
Uò
Wò
>ò,ò/ò@ò
6ñ
9 x°Ñ'<ó 9ñ<.�—‘ð .°u¸]ÈEÏLÉLÐ=XÑ7Yó .ð&=˜3ð = 2§8¡8ó =ðG˜só GòAó	aò6ó^ð �{‰{×!Ñ!¡l£nÐ"4¹\»^Ð!ÓL€ð>˜3ð > só >ð�U˜5Ÿ<™<¨¯©×);Ñ);×)IÑ)IÐIÑJð ÈuÏ|É|ó ñ+* H¨S°©XÑ$6ô +*r#   