Ë
    Öéñi`"  ã                   ó"  — 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!„ Z3d5d"ee   fd#„Z4d6d$ejj                  d%eeejl                  f   fd&„Z7d$e8d%ejr                  fd'„Z:d%e8fd(„Z;d)„ Z<d7d*„Z=d+„ Z>d7d,„Z?ej€                  jƒ                   e=«         e?«       ¬-«      ZBd.eCd/eCfd0„ZDd1eejl                  ejŠ                  jŒ                  j$                  f   d%ejl                  fd2„ZGd6d3eee8      fd4„ZHy)8é    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© ó    úZ/var/www/pod-logistic/pod-ai/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                  ób   — t        «       xr$ t        j                  j                  «       d   dk(  S )Nr   é
   r6   r"   r#   r$   Úis_blackwellr=   -   ó&   € Ü‹9ÒDœŸ™×9Ñ9Ó;¸AÑ>À"ÑDÐDr#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk\  S ©Nr   é	   r6   r"   r#   r$   Úis_hopper_or_newerrB   1   r:   r#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk(  S r@   r6   r"   r#   r$   Ú	is_hopperrD   5   r:   r#   c                  ób   — t        «       xr$ t        j                  j                  «       d   dk(  S )Nr   é   r6   r"   r#   r$   Úis_sm12xrG   9   r>   r#   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚhipr/   r1   s    r$   Úis_hiprJ   =   ó"   € ÜÓ!€FØ�Nˆ5Ð?¨¯©¸%Ñ(?Ð?r#   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )NrI   Úgfx90a©r,   r0   Úarchr1   s    r$   Úis_hip_cdna2rP   B   ó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 )NrI   Úgfx942rN   r1   s    r$   Úis_hip_cdna3rT   G   rQ   r#   c                  ób   — t        «       } | d uxr  | j                  dk(  xr | j                  dk(  S )NrI   Úgfx950rN   r1   s    r$   Úis_hip_cdna4rW   L   rQ   r#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrI   Úgfx11rN   r1   s    r$   Úis_hip_gfx11rZ   Q   ó1   € ÜÓ!€FØ˜ÐÒT &§.¡.°EÑ"9ÒT¸gÈÏÉÐ>TÐTr#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrI   Úgfx12rN   r1   s    r$   Úis_hip_gfx12r^   V   r[   r#   c                  ó`   — t        «       } | d uxr | j                  dk(  xr d| j                  v S )NrI   Úgfx1250rN   r1   s    r$   Úis_hip_gfx1250ra   [   s1   € ÜÓ!€FØ˜ÐÒV &§.¡.°EÑ"9ÒV¸iÈ6Ï;É;Ð>VÐVr#   c                  óF   — t        «       xs t        «       xs
 t        «       S r'   )rP   rT   rW   r"   r#   r$   Úis_hip_cdnarc   `   s   € Ü‹>Ò=œ\›^Ò=¬|«~Ð=r#   c                  ó   — t        «       rdS dS )Ni € i   )rW   r"   r#   r$   Úget_hip_lds_sizere   d   s   € Ü!”^ˆ6Ð.¨Ð.r#   c                  ó<   — t        «       } | €dS | j                  dk(  S )NFÚxpur/   r1   s    r$   Úis_xpurh   h   rK   r#   c                  óH   — t        «       } | €dS t        | j                  «      S )NÚ )r,   ÚstrrO   r1   s    r$   Úget_archrl   m   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_strrm   ÚlowÚhighr}   rq   Úxs           r$   Únumpy_randomrŒ   r   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)Údevicers   r   r   )rq   Ú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+)$rr   znot a triton or torch dtype: )rx   r(   Úlanguagerq   r‘   r7   ÚreÚmatchrk   ÚgroupÚ	TypeErrorÚtype)rq   Ú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: )rx   r	   ÚbaseÚcpuÚnumpyr„   r~   r|   r¦   rq   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)rF   r   )rF   é   ú.é   r   rA   )r%   r3   r   ÚnvidiaÚptxasÚversionÚtupleÚmapry   Ú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   rA   )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.   )rq   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'   )rx   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Ñ   rx   Ú
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)Ir   r    rª   r|   r7   r(   Útriton.languagerŸ   r”   r   Útypingr   r   r   rç   Únumpy.randomr   Útriton.runtime.jitr	   r
   r   rz   r{   Úintegral_dtypesr‚   Úfloat_dtypes_with_bfloat16ÚdtypesÚdtypes_with_bfloat16Útorch_float8_dtypesÚtorch_dtypesÚsortedræ   Ú
tma_dtypesr%   r,   r3   r9   r=   rB   rD   rG   rJ   rP   rT   rW   rZ   r^   ra   rc   re   rh   rl   rŒ   Úndarrayr«   r™   rk   rq   r�   r¦   r®   r¾   rÀ   rÂ   ÚmarkÚskipifÚrequires_tmary   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òEò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#   