Ë
    L²j*[ ã                  óÐ  — U d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dlZd dlZd dlZ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mZmZmZmZm Z m!Z! d dl"m"Z" d dlm#Z# d dl	m$Z$ d dl%m&Z&m'Z'm(Z(m)Z)m*Z*m+Z+m,Z,m-Z-m.Z.m/Z/m0Z0 d dl1m2Z2m3Z3m4Z4 d d	lm5Z5 d dl6Z6d dl7Z7d dl8m9c m:Z; d d
l<m=Z= d dl>m?Z? d dl@mAZA d dlBmCZC d dlDmEZE d dl8mFZFmGZG ddgZHd dlImJZJ d dlKmLZLmMZMmNZNmOZO e-rhd dlmPZPmQZQmRZR d dlSmTZT d dl7mUZUmVZVmWZW d dlXmYZY d dlZm[Z[ d dl\m]Z] d dl^m_Z_ ddl`maZa ddlbmcZc ddldmeZe ddlfmgZg dd lhmiZimjZjmkZkmlZlmmZmmnZn dd!lompZp dd"lqmrZrmsZs g d#¢Zt e0d$«      Zuejì                  �dLd%„«       Zwd d&lxmyZy d d'lzm{Z{ d d(l|m}Z} d d)l~mZ d d*l€m�Z� d d+l‚mƒZƒ d d,l„m…Z…m†Z†m‡Z‡mˆZˆm‰Z‰ d d-lŠm‹Z‹mŒZŒ d d.l�mŽZŽm�Z� dd/l�m‘Z‘ dd0l’m“Z” ej                  d1k(  Z• e�j,                  e—«      Z˜ e0d2«      Z™eše6�j6                  e6�j6                  f   Zœe7�j:                  ežz  e7j®                  z  dz  ZŸe•rd3n e�j@                  d4d5«      Z¡d6d7d8e¡› �d9œZ¢d:Z£d:Z¤d:Z¥d;Z¦ eEe7�jN                  e7�jP                  e7�jR                  e7�jT                  e7�jV                  e7�jX                  e7�jZ                  e7�j\                  e7�j^                  e7�j`                  e7�jb                  e7�jd                  e7�jf                  e7�jh                  e7�jj                  g«      Z¶d<e·d=<   d>Z¸e¸e¸dz
  z  d k(  re¸d?k\  sJ d@«       ‚�dMdA„Z¹�dNdB„Zº G dC„ dDe6�jv                  «      Z¼ e�jz                  dE¬F«       G dG„ dH«      «       Z¾�dO�dPdI„Z¿	 	 	 �dQ	 	 	 	 	 	 	 	 	 �dRdK„ZÀ	 	 	 �dQ	 	 	 	 	 	 	 	 	 �dRdL„ZÁejì                  �dSdM„«       ZÂ�dTdN„ZÃ�dUdO„ZÄ�dVdP„ZÅ�dWdQ„ZÆ�dXdR„Z“�dYdS„ZÇ	 	 	 	 �dZdT„ZÈ�d[dU„ZÉ�d\dV„ZÊ	 	 	 	 �d]dW„ZË�d^dX„ZÌdY„ f	 	 	 	 	 �d_dZ„ZÍ	 	 	 	 	 	 	 	 �d`d\„ZÎ�da�dbd]„ZÏ	 	 �dc	 	 	 	 	 	 	 	 	 �ddd^„ZÐ	 	 	 	 	 �de	 	 	 	 	 	 	 	 	 	 	 	 	 �dfd_„ZÑ�dgd`„ZÒ�dhda„ZÓ�didb„ZÔ�djdc„ZÕ�dkdd„ZÖ e3de«      Z× e0dfdE¬g«      ZØee(e&e×f   eØf   ZÙ G dh„ die,e)e×eØf   «      ZÚ�dldj„ZÛ	 	 	 	 �dldk„ZÜ	 	 	 	 �dmdl„ZÝ	 	 	 	 �dndm„ZÞ	 	 	 	 	 	 �dodn„Zß	 	 	 	 	 	 �dpdo„Zà	 �dq	 	 	 	 	 �drdp„Zá	 	 	 	 	 	 �dsdq„Zâ�dtdr„Zã�duds„Zä�dvdt„Zå�dwdu„Zæ�dxdv„Zç�dydw„Zè�dzdx„Zé�d{dy„Zê�d|dz„Zë eìg d{¢«      Zí	 	 	 	 �d}d|„Zî�d~d}„Zï�dd~„Zðd dlñZñ�d€d„Zòg Zód[e·d€<   �d�d�„Zô�d€d‚„Zõe�jì                  �d‚dƒ„«       Z÷e�jì                  	 	 	 �dƒ	 	 	 	 	 	 	 �d„d„„«       ZøeôZùeõZúeøZûdJd…œ�d…d†„ZüdJd…œ	 	 	 	 	 	 	 �d†d‡„Zý ejF                  d?«      �d‡dˆ„«       Zþ G d‰„ dŠe+«      Zÿe�jz                   G d‹„ dŒ«      «       �Z  G d�„ dŽ«      �Z G d�„ d��e«      �Ze�jì                  �dˆd‘„«       �Z G d’„ d“«      �Z G d”„ d•�e«      �Zejì                  �d‰�dŠd–„«       �ZejF                  �d‹d—„«       �ZejF                  �dSd˜„«       �Z�d‹d™„�Z		 �dq	 	 	 	 	 	 	 �dŒdš„�Z
�d��dŽd›„�Z	 	 	 	 	 	 �d�dœ„�Z�d�d�„�Z�d�dž„�ZdJdJdEdŸœ	 	 	 	 	 	 	 	 	 �d‘d „�ZddJd¡œ	 	 	 	 	 	 	 �d’d¢„�Z	 �d“	 	 	 	 	 �d”d£„�ZdJd¤œ	 	 	 	 	 	 	 �d•d¥„�ZdJd¤œ	 	 	 	 	 	 	 �d•d¦„�Z	 	 	 	 	 	 	 	 �d–d§„�Z ejF                  d¬¨«      �dSd©„«       �Z ejF                  d¬¨«      �dSdª„«       �Zd«„ �Z ejF                  d¬¨«      �dSd¬„«       �Z	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d—d­„�Z�d˜d®„�Z	 	 �d™	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �dšd°„�Z�d›d±„�Zeže6�j6                  z  �Zd²e·d¯<   ejì                  	 �dœ	 	 	 	 	 	 	 	 	 �d�d³„«       �Zejì                  �džd´„«       �Zejì                  �dŸdµ„«       �Z ejì                  �d d¶„«       �Z!ejì                  �d¡d·„«       �Z"�d¢d¸„�Z#�d˜d¹„�Z$�d˜dº„�Z%�d¢d»„�Z&�d¢d¼„�Z'	 	 	 	 	 	 	 	 �d£d½„�Z(	 	 	 	 �d¤	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d¥d¾„�Z)�dSd¿„�Z* G dÀ„ dÁ«      �Z+	 	 	 	 	 	 	 	 �d¦dÂ„�Z,	 	 	 	 	 	 	 	 �d¦dÃ„�Z-�d§dÄ„�Z.�d¨dÅ„�Z/�d©dÆ„�Z0	 	 	 	 	 	 	 	 �d©dÇ„�Z1	 	 	 	 	 	 	 	 �dªdÈ„�Z2e�jì                  	 	 	 	 	 	 �d«dÉ„«       �Z3	 �dq	 	 	 	 	 �d¬dÊ„�Z4�d­dË„�Z5�d®dÌ„�Z6�d¯dÍ„�Z7�d¯dÎ„�Z8�d°dÏ„�Z9�d±dÐ„�Z:e�jì                  �d²dÑ„«       �Z;�d‹dÒ„�Z<ejì                  �d‹dÓ„«       �Z=ejì                  �d³dÔ„«       �Z>ejì                  �d‹dÕ„«       �Z?�d‹dÖ„�Z@�d‹d×„�ZA�d´dØ„�ZB�dµdÙ„�ZC�dSdÚ„�ZD�dSdÛ„�ZE�d¶dÜ„�ZF�d|dÝ„�ZG G dÞ„ dße�j�                  «      �ZI	 	 	 	 	 	 	 	 	 	 �d·dà„�ZJ�d¸dá„�ZK	 	 	 	 �d¸dâ„�ZL	 �dq	 	 	 	 	 �d¹dã„�ZM�dºdä„�ZN	 �dq	 	 	 	 	 �d»då„�ZO�d¼dæ„�ZP	 	 	 	 	 	 �d½dç„�ZQ	 	 	 	 	 	 	 	 �d¾dè„�ZRdé„ f	 	 	 	 	 	 	 	 	 	 	 �d¿dê„�ZSdë„ f	 	 	 	 	 	 	 	 	 	 	 �d¿dì„�ZT�dÀdí„�ZU�dÁdî„�ZV�dÂdï„�ZWe�jz                   G dð„ dñ«      «       �ZXe�jì                  �dÃdò„«       �ZY�dÄdó„�ZZ�dÅdô„�Z[�dSdõ„�Z\�dÆdö„�Z]�dÇd÷„�Z^	 	 	 	 	 	 	 	 	 	 	 	 	 	 �dÈdø„�Z_�dÉdù„�Z`�dÊdú„�Za�dËdû„�Zb�dÌdü„�Zc	 	 	 	 	 	 	 	 �dÍdý„�Zd�dÎdþ„�Ze	 	 	 	 	 	 	 	 �dÏdÿ„�Zf�dÐ�d „�Zg	 �dq	 	 	 	 	 	 	 �dÑ�d„�Zh	 	 	 	 	 	 �dÒ�d„�Zi�dÓ�d„�Zj	 	 	 	 	 	 �dÔ�d„�Zk�dS�d„�Zl�dÄ�d„�Zm�d�d�d	�d
�d�d�d�dœ�Zn�en�jß                  «       D � �ci c]  \  } }|| “Œ
 c}} �Zp e�jâ                  �d«      �Zr�dÕ�d„�Zs�dÖ�d„�Zt�d×�d„�Zu�d×�d„�Zvejì                  �dØ�d„«       �Zwe�jz                   G �d„ �d«      «       �Zxi �Zy�de·�d<   	 	 	 	 	 	 	 	 �dÙ�d„�Zz eE«       �Z{�de·�d<   �dÚ�d„�Z|�dq�dÛ�d„�Z}�dÜ�d„�Z~ e0�d«      �Z e0�d«      �Z€ G �d „ �d!e �e�e€f   «      �Z� e2dE�¬"«      �dqdEdFœ�dÝ�d#„«       �Z‚�dÞ�d$„�Zƒ G �d%„ �d&e�j�                  «      �Z„ejì                  �dß�d'„«       �Z…�dS�d(„�Z†�dà�d)„�Z‡�dá�d*„�Zˆ�dá�d+„�Z‰�dâ�d,„�ZŠ�dL�d-„�Z‹�dã�d.„�ZŒ�dS�d/„�Z��dä�d0„�ZŽ�d1�Z��då�d2„�Z��då�d3„�Z‘�dæ�d4„�Z’	 	 �d™	 	 	 	 	 	 	 	 	 �dç�d5„�Z“�dè�d6„�Z”�dé�d7„�Z•�dS�d8„�Z–�dê�d9„�Z—�dë�d:„�Z˜ e�jz                  dE¬F«       G �d;„ �d<«      «       �Z™e�d=e&f   �Zše�eš�e™g�ešf   �Z› G �d>„ �d?«      �Zœ �eœ«       �Z��dì�d@„�Zž�dí�dA„�ZŸ�dî�dB„�Z �dï�dC„�Z¡�dð�dD„�Z¢ eEg �dE¢«      �Z£�d›�dF„�Z¤e#�dñ�dG„«       �Z¥�dò�dH„�Z¦	 	 	 	 	 	 	 	 	 	 	 	 	 	 �dó�dI„�Z§	 	 	 	 	 	 �dô�dJ„�Z¨	 �d“	 	 	 	 	 	 	 �dõ�dK„�Z©yc c}} w (ö  é    )ÚannotationsN)ÚCallableÚ
CollectionÚ	GeneratorÚIteratorÚMappingÚMutableMappingÚ
MutableSet)Údatetime)Ú	lru_cache)ÚStringIO)ÚAnyÚcastÚConcatenateÚGenericÚLiteralÚ
NamedTupleÚProtocolÚTYPE_CHECKINGÚ	TypeAliasÚ	TypeGuardÚTypeVar)Údataclass_transformÚ	ParamSpecÚSelf)Úmock)Údatasheet_tops)ÚDeviceProperties)Ú_needs_inductor_compile)Údtype_abbrs)Ú
OrderedSet)Útree_flattenÚtree_map_onlyÚ!activation_quantization_aten_passÚinductor_autotune_lookup_table©Ú_sympy_subs)Úfree_symbolsÚfree_unbacked_symbolsÚIterateExprsÚShapeEnv)ÚIterableÚSequenceÚ
ValuesView)ÚPath)ÚSymBoolÚSymFloatÚSymInt)ÚELEMENTWISE_TYPE_PROMOTION_KIND)ÚGraphModule)ÚNode)ÚScalingTypeé   )ÚWorkspaceArg©ÚPythonWrapperCodegen)ÚDep©ÚGraphLowering)ÚBufferÚExternKernelÚIRNodeÚLayoutÚ	OperationÚReinterpretView©ÚCompiledFxGraph)ÚBaseSchedulerNodeÚSchedulerBuffer)ÚcudaÚmpsÚxpuÚmtiaÚTc                 óØ   — t         D � cg c]#  } t        t        | «      j                  «       sŒ"| ‘Œ% }} t	        |«      dk  sJ ‚t	        |«      dk(  rd}|S |j                  «       }|S c c} w )Nr7   r   rH   )Ú	GPU_TYPESÚgetattrÚtorchÚis_availableÚlenÚpop)ÚxÚ
avail_gpusÚgpu_types      úT/var/www/html/ia-enhanced/venv/lib/python3.12/site-packages/torch/_inductor/utils.pyÚget_gpu_typerX   j   sg   € ä&ÖK˜¬'´%¸Ó*;×*HÑ*HÕ*J’!ÐK€JÐKÜˆz‹?˜aÒÐÐÜ˜Z›¨AÒ-ˆv€HØ€Oð 4>·>±>Ó3C€HØ€Oùò Ls
   ‰#A'­A')Úget_interface_for_device)Údetect_fake_mode)Ú
DeviceType)Ú	EventList)ÚGraphTransformObserver)Ú	ShapeProp)ÚCeilDivÚCleanDivÚFloorDivÚIdentityÚModularIndexing)Úmake_symbolÚSymT)Úbound_sympyÚValueRanges©Úconfig)ÚceildivÚwin32Ú_TÚspvÚTORCHINDUCTOR_XPU_KERNEL_FORMATÚzebinz.cubinz.hsacoú.)rH   ÚhiprJ   é   é€   zOrderedSet[torch.dtype]Ú_TMA_SUPPORTED_DTYPESé@   é   zmust be power of 2c                ó*   — | t         z   dz
  t          z  S )z/Round up to the nearest multiple of ALIGN_BYTESr7   )ÚALIGN_BYTES)Únbytess    rW   Ú_alignrz   º   s   € à”[Ñ  1Ñ$¬¨Ñ4Ð4ó    c                ó  — t        | t        j                  t        j                  f«      r#t	        t        t        | j                  «      «      S t        | t        «      xs! t        j                  | t        «      t        k(  S )z:v can be statically proven to be a multiple of ALIGN_BYTES)Ú
isinstanceÚsympyÚAddÚMaxÚallÚmapÚ_is_alignedÚargsÚalignÚgcdrx   )Úvs    rW   rƒ   rƒ   ¿   sQ   € ä�!”e—i‘i¤§¡Ð+Ô,Ü”3”{ A§F¡FÓ+Ó,Ð,Ü�aœÓÒK¤5§9¡9¨Q´Ó#<ÄÑ#KÐKr{   c                  ó*   — e Zd ZdZdZdZedd„«       Zy)r…   z<Symbolically round up to the nearest multiple of ALIGN_BYTES©r7   Tc                ó†   — t        |t        t        j                  f«      rt	        t        |«      «      S t        |«      r|S y ©N)r}   Úintr~   ÚIntegerrz   rƒ   )ÚclsÚvalues     rW   Úevalz
align.evalÌ   s6   € ä�eœc¤5§=¡=Ð1Ô2Üœ#˜e›*Ó%Ð%Ü�uÔØˆLð r{   N)r�   ú
sympy.ExprÚreturnzsympy.Expr | None)Ú__name__Ú
__module__Ú__qualname__Ú__doc__ÚnargsÚ
is_integerÚclassmethodr�   © r{   rW   r…   r…   Æ   s!   „ ÙFà€EØ€Jàòó ñr{   r…   T)Úfrozenc                  ó:   — e Zd ZU dZded<   ded<   ded<   ded<   y	)
ÚGraphPartitionMapzP
    Mapping from the partition info (e.g., input/output) to the graph info
    rŒ   Úidzlist[int | None]Úinput_index_mappingÚoutput_index_mappingú	list[str]Úconstant_namesN©r“   r”   r•   r–   Ú__annotations__rš   r{   rW   r�   r�   Ô   s$   … ñð
 	ƒGð *Ó)Ø*Ó*ð Ôr{   r�   c           
     óˆ  —  | «        t         j                  j                  «        t        j                  t	        d«      t         j
                  d¬«      }t         j                  j                  d¬«      }t         j                  j                  d¬«      }|j                  «        t        d«      D ]  }|j                  «         | «        Œ |j                  «        t         j                  j                  «        |j                  |«      dz  }t        dt	        ||z  «      «      }t        dt	        ||z  «      «      }	t        |«      D ]	  } | «        Œ t        |	«      D �cg c]"  }t         j                  j                  d¬«      ‘Œ$ }}t        |	«      D �cg c]"  }t         j                  j                  d¬«      ‘Œ$ }}t         j                  j                  t         j                  j                  j                  g¬«      5 }
t         j                  j                  «        t        |	«      D ]q  }|j                  «        ||   j                  «        t         j                  j                   j                  d	«      5   | «        d
d
d
«       ||   j                  «        Œs t         j                  j                  «        t        j"                  t%        ||«      D ��cg c]  \  }}|j                  |«      ‘Œ c}}«      }d
d
d
«       t        j&                  «      j)                  «       }t*        j-                  d«       t*        j-                  
j/                  «       j1                  dd¬«      «       t3        |
j5                  «       D �cg c]A  }|j6                  t8        j                  k(  r"t;        j<                  d|j>                  «      �|‘ŒC c}«      }|r"|tA        j&                  d„ |D «       «      dz  z  }t*        j-                  d|«       |S c c}w c c}w # 1 sw Y   �ŒœxY wc c}}w # 1 sw Y   �Œ3xY wc c}w )áR  
    Returns benchmark results by examining torch profiler events.
    This could be more accurate as it doesn't count CPU side overhead.
    However, this also requires manually excluding irrelevant event, e.g.
    vectorized_elementwise_kernel which is used to fill L2 cache,
    various CUDA events, etc, so could also be fragile.
    ç    €„ŽArH   ©ÚdtypeÚdeviceT©Úenable_timingé   r7   ©Ú
activitiesÚRunCudaModuleNú
raw eventsÚself_device_time_totaléÿÿÿÿ©Úsort_byÚ	row_limitzfused_abs_max_\dc              3  ó4   K  — | ]  }|j                   –— Œ y ­wr‹   ©Údevice_time_total©Ú.0Úevents     rW   ú	<genexpr>zfp8_bench.<locals>.<genexpr>(  s   è ø€ ÒQ¸˜E×3Õ3ÑQùó   ‚ç     @�@úprofiling results: %s ms)!rP   rH   ÚsynchronizeÚemptyrŒ   Úfloat16ÚEventÚrecordÚrangeÚzero_Úelapsed_timeÚmaxÚprofilerÚprofileÚProfilerActivityÚCUDAÚnvtxÚtensorÚzipÚmeanÚitemÚlogÚdebugÚkey_averagesÚtabler\   ÚeventsÚdevice_typer[   ÚreÚmatchÚnameÚ
statistics)ÚfnÚwarmupÚrepÚcacheÚstart_eventÚ	end_eventÚ_Úestimate_msÚn_warmupÚn_repeatÚpÚiÚsÚeÚtimesÚresr¼   Úfiltered_eventss                     rW   Ú	fp8_benchrî   æ   sT  € ñ „DÜ	‡J�J×ÑÔÜ�K‰Kœ˜J›¬u¯}©}ÀVÔL€Eô —*‘*×"Ñ"°Ð"Ó6€KÜ—
‘
× Ñ ¨tÐ Ó4€IØ×ÑÔÜ�1‹Xò ˆØ�‰ŒÙ
�ðð ×ÑÔÜ	‡J�J×ÑÔØ×*Ñ*¨9Ó5¸Ñ9€Kô �1”c˜& ;Ñ.Ó/Ó0€HÜ�1”c˜# Ñ+Ó,Ó-€Hô �8‹_ò ˆÙ
�ðô BGÀxÃÖQ¸A”5—:‘:×#Ñ#°$Ð#Õ7ÐQ€KÐQÜ?DÀX»ÖO¸!”—‘×!Ñ!°Ð!Õ5ÐO€IÐOÜ	�‰×	Ñ	ä�N‰N×+Ñ+×0Ñ0ð
ð 
 ó 
ð 
ð 
Ü�
‰
×ÑÔ Ü�x“ò 	"ˆAØ�K‰KŒMØ˜‰N×!Ñ!Ô#Ü—‘—‘×&Ñ& Ó7ñ Ù”÷à�a‰L×ÑÕ!ð	"ô 	�
‰
×ÑÔ Ü—‘Ü+.¨{¸IÓ+F×G¡4 1 aˆQ�^‰^˜AÕÓGó
ˆ÷
ô" �*‰*�UÓ
×
 Ñ
 Ó
"€CÜ‡I�IˆlÔÜ‡I�Iˆa�n‰nÓ×$Ñ$Ð-EÐQSÐ$ÓTÔUÜð Ÿ™›ö	
àà×!Ñ!¤Z§_¡_Ò4Ü—H‘HÐ0°%·*±*Ó=ÐIò	 ò	
ó	€Oñ ØÜ�O‰OÑQÀÔQÓQØññ	
ˆô
 ‡I�IÐ(¨#Ô.Ø€JùòO RùÚO÷ñ üó
 H÷
ñ 
üò*	
sE   Å"'PÆ'PÈA9P2É=PÊAP2ËP,Ë9P2ÎAP?ÐP)Ð$P2Ð2P<Fc                ó8   — ddl m}   |t        «      | |||«      S )Nr   )Úmay_distort_benchmarking_result)Ú$torch._inductor.runtime.benchmarkingrð   Ú_do_bench_using_profiling)rÝ   rÞ   rß   Úis_vetted_benchmarkingrð   s        rW   Údo_bench_using_profilingrô   0  s(   € õ" UàEÑ*Ô+DÓEØ
ˆF�CÐ/óð r{   c           
     óè  — |sddl m}  |«        t        «       }|j                  «       }t	        |«      } | «        |j                  «        t        j                  t        d«      t        j                  |¬«      }|j                  d¬«      }	|j                  d¬«      }
|	j                  «        t        d«      D ]  }|j                  «         | «        Œ |
j                  «        |j                  «        |	j                  |
«      dz  }t        dt        ||z  «      «      }t        dt        ||z  «      «      }t        |«      D ]	  } | «        Œ |j                  «        t        j                  j!                  t#        t        j                  j$                  |«      g¬	«      5 }t        |«      D ]  }|j                  «         | «        Œ |j                  «        d
d
d
«       t&        j)                  d«       t&        j)                  j+                  «       j-                  dd¬«      «       t/        |j1                  «       D �cg c]0  }|j2                  t#        t4        |«      k(  r|j6                  dk7  r|‘Œ2 c}«      }t/        |D �cg c]  }d|j6                  vsŒ|‘Œ c}«      }t9        |«      dk(  rt;        d|› dt9        |«      › d|› �«      ‚|j=                  «        |j+                  «       }t&        j)                  d«       t&        j)                  |j-                  d¬«      «       t?        d„ |D «       «      dz  |z  }t&        j)                  d|«       |S # 1 sw Y   �ŒƒxY wc c}w c c}w )r¦   r   )Úmay_ban_benchmarkingr§   r¨   Tr«   r­   r7   r®   Nr±   r²   r³   r´   zContext SyncÚFillFunctorzDFailed to capture any events after filtering cache clearing events. z	 events: z, repeats: zprofiling time breakdown)r¶   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr‹   r¸   rº   s     rW   r½   z,_do_bench_using_profiling.<locals>.<genexpr>¡  s   è ø€ ÒA¨%ˆe×%Õ%ÑAùr¾   r¿   rÀ   ) rñ   rö   rX   ÚupperrY   rÁ   rP   rÂ   rŒ   rÄ   rÅ   rÆ   rÇ   rÈ   rÉ   rÊ   rË   rO   rÌ   rÓ   rÔ   rÕ   rÖ   r\   r×   rØ   r[   rÛ   rR   ÚRuntimeErrorÚ_build_treeÚsum)rÝ   rÞ   rß   ró   rö   rØ   Údevice_type_upperÚdevice_interfacerà   rá   râ   rã   rä   rå   ræ   rç   r¼   rí   Úactual_eventsrì   s                       rW   rò   rò   H  s
  € ñ "ÝMáÔä“.€KØ#×)Ñ)Ó+ÐÜ/°Ó<ÐÙ„DØ× Ñ Ô"Ü�K‰Kœ˜J›¬u¯y©yÀÔM€Eð #×(Ñ(°tÐ(Ó<€KØ ×&Ñ&°TÐ&Ó:€IØ×ÑÔÜ�1‹Xò ˆØ�‰ŒÙ
�ðð ×ÑÔØ× Ñ Ô"Ø×*Ñ*¨9Ó5¸Ñ9€Kô �1”c˜& ;Ñ.Ó/Ó0€HÜ�1”c˜# Ñ+Ó,Ó-€Hô �8‹_ò ˆÙ
�ðð × Ñ Ô"Ü	�‰×	Ñ	ä”E—N‘N×3Ñ3Ð5FÓGð
ð 
 ó 
ð 'ð 
ä�x“ò 	ˆAà�K‰KŒMá�Dð		ð 	×$Ñ$Ô&÷'ô ‡I�IˆlÔÜ‡I�Iˆa�n‰nÓ×$Ñ$Ð-EÐQSÐ$ÓTÔUäð Ÿ™›ö	
àØ× Ñ ¤G¬JÐ8IÓ$JÒJØ—
‘
˜nÒ,ò ò	
ó€Oô Ø+ÖO�5¨}ÀEÇJÁJÒ/NŠÒOó€Mô ˆ=Ó˜QÒÜØRØˆm˜9¤S¨Ó%9Ð$:¸+ÀhÀZðQó
ð 	
ð ×ÑÔØ!×.Ñ.Ó0€Mä‡I�IÐ(Ô)Ü‡I�Iˆm×!Ñ!¨BÐ!Ó/Ô0ä
ÑA°=ÔAÓ
AÀFÑ
JÈXÑ
U€CÜ‡I�IÐ(¨#Ô.Ø€J÷['ñ 'üò$	
ùò 	Ps   Æ$8MÉ5M*Ê	M/ÊM/ÍM'c                 ó  — 	 ddl m}  t        j                  j	                  dd«       | d uxr% t        t        t        j                  dd «      d«      S # t        $ r Y yt        $ r}dt        |«      v sJ ‚Y d }~yd }~ww xY w)	Nr   )Ú	roi_alignztorchvision::nmsÚMetaÚtorchvisionr  Fztorchvision::nms does not exist)Útorchvision.opsr  rP   Ú_CÚ%_dispatch_has_kernel_for_dispatch_keyÚhasattrrO   ÚopsÚImportErrorrú   Ústr)r  rê   s     rW   Úhas_torchvision_roi_alignr  ¦  s|   € ðÝ-ä�‰×6Ñ6Ð7IÈ6ÔRØ Ð$ò 
¬Ü”E—I‘I˜}¨dÓ3°[ó*
ð 	
øô ò ÙÜò Ø0´C¸³FÑ:Ð:Ð:Üûðús   ‚AA Á	A?ÁA?Á&A:Á:A?c                ób  — | €t        j                  d«      j                  S t        | t        «      rt        j                  | «      } | j
                  dvrZ| j                  €Nt        | j
                  «      }t        j                  | j
                  |j                  j                  «       ¬«      S | S )Ng        )ÚcpuÚmeta)Úindex)
rP   rÏ   rª   r}   r
  Útyper  rY   ÚWorkerÚcurrent_device©rª   rþ   s     rW   Údecode_devicer  ¶  s‚   € Ø€~Ü�|‰|˜CÓ ×'Ñ'Ð'Ü�&œ#ÔÜ—‘˜fÓ%ˆØ‡{�{˜/Ñ)¨f¯l©lÐ.BÜ3°F·K±KÓ@ÐÜ�|‰|˜FŸK™KÐ/?×/FÑ/F×/UÑ/UÓ/WÔXÐXØ€Mr{   c                ó|   — t        j                  t        j                  | t        j
                  j                  «      S r‹   )Ú	functoolsÚreduceÚoperatorÚmulr~   ÚSÚOne)Úits    rW   Úsympy_productr  Á  s#   € Ü×ÑœHŸL™L¨"¬e¯g©g¯k©kÓ:Ð:r{   c           	     ó’   — t        | «      t        |«      k(  sJ ‚t        j                  t        d„ t	        | |«      D «       «      «      S )Nc              3  ó,   K  — | ]  \  }}||z  –— Œ y ­wr‹   rš   )r»   ÚaÚbs      rW   r½   zsympy_dot.<locals>.<genexpr>Ç  s   è ø€ Ò>¡d a¨˜A �EÑ>ùs   ‚)rR   r~   Úexpandrü   rÐ   )Úseq1Úseq2s     rW   Ú	sympy_dotr%  Å  s8   € Üˆt‹9œ˜D›	Ò!Ð!Ð!Ü�<‰<œÑ>¬c°$¸«oÔ>Ó>Ó?Ð?r{   c                ó\   — | D �ci c]  }t        |«      |“Œ c}j                  «       S c c}w r‹   )rž   Úvalues)r  rT   s     rW   Úuniquer(  Ê  s'   € Ø Ö!˜ŒBˆq‹E�1‰HÒ!×(Ñ(Ó*Ð*ùÒ!s   …)c           
     ón  — t        | t        j                  «      st        |t        j                  «      r2t        t        j                  | «      t        j                  |«      «      S t        | t
        «      rt        |t
        «      s$J | › dt        | «      › d|› dt        |«      › �«       ‚t        | |«      S )Nz: ú, )r}   r~   ÚExprr_   ÚsympifyrŒ   r  Úruntime_ceildiv)ÚnumberÚdenoms     rW   rj   rj   Î  s�   € Ü�&œ%Ÿ*™*Ô%¬°E¼5¿:¹:Ô)FÜ”u—}‘} VÓ,¬e¯m©m¸EÓ.BÓCÐCô �fœcÔ"¤z°%¼Ô'=ð Øˆ(�"”T˜&“\�N " U G¨2¬d°5«k¨]Ð;óÐ=ô ˜6 5Ó)Ð)r{   c                ój  — | €yt        | «      j                  d«      d   }i dd“dd“dd	“d
d“dd“dd“dd	“dd“dd“dd“dd“dd“dd“dd“dd“dd “d!d"“d#d$dd%d&d'd(œ¥}|j                  t        |j	                  «       «      D �ci c]  }||“Œ c}«       t        | t         «      r| S d)||   › �S c c}w )*Nz*i8rp   r³   ÚboolÚi1Ú
float8e4nvÚfp8e4nvÚfloat8e5Úfp8e5Úfloat8e4b15Úfp8e4b15Úfloat8e4b15x4Ú
fp8e4b15x4Úfloat8_e4m3fnÚfloat8_e5m2Úfloat8_e4m3fnuzÚfp8e4b8Úfloat8_e5m2fnuzÚfp8e5b16Úfloat8_e8m0fnuÚu8Úfloat4_e2m1fn_x2rÃ   Úfp16Úbfloat16Úbf16Úfloat32Úfp32Úfloat64Úfp64Úint8Úi8Úint16Úi16Úi32Úi64Úu16Úu32Úu64)Úint32Úint64Úuint8Úuint16Úuint32Úuint64Ú*)r
  ÚsplitÚupdateÚlistr'  r}   )ÚkeyÚ	dtype_strÚtysr‡   s       rW   Ú_type_ofra  Ú  sX  € ð
 €{ØÜ�C“—‘˜sÓ# BÑ'€IðØ�ðà�iðð 	�Gðð 	�zð	ð
 	˜ðð 	˜ðð 	�wðð 	˜9ðð 	˜:ðð 	˜$ðð 	˜Dðð 	�6ðð 	�Fðð  	�6ð!ð" 	�6ð#ð$ 	�ð%ð& 	�ð'ð( ØØØØØò3€Cð8 ‡J�Jœd 3§:¡:£<Ó0Ö1˜��1‘Ò1Ô2Ü˜S¤#Ô&ˆ3Ð@¨a°°I±Ð/?Ð,@Ð@ùò 2s   Â
B0c                óR   — | D �cg c]  }t        j                  |«      ‘Œ c}S c c}w )z­
    Gets the shape and stride of a tensor. For non-symbolic tensors, this is
    trivial. But for symbolic tensors, we need to map from SymIntNode into
    sympy.Expr.
    )r~   r,  ©Úlstrè   s     rW   Úconvert_shape_to_inductorre    s!   € ð '*Ö* ŒE�M‰M˜!ÕÒ*Ð*ùÒ*s   …$c                óf   — t        | t        j                  «      r| j                  j                  S | S )zÂ
    Convert SymInt to sympy.Expr, leave int as is.

    Unlike sympy.sympify() which converts int to sympy.Integer,
    this function preserves int as int and only converts SymInt to Expr.
    )r}   rP   r2   ÚnodeÚexpr©Úvals    rW   Úconvert_symint_to_exprrk    s%   € ô �#”u—|‘|Ô$Ø�x‰x�}‰}ÐØ€Jr{   c                óÞ   — ddl m} t        | t        «      r| S t        | t        j
                  «      rt        | «      S |j                  j                  j                  j                  | d¬«      S )zL
    Like convert_shape_to_symint, but operates on a single expression.
    r7   ©ÚVN)Úhint)
Úvirtualizedrn  r}   rŒ   r~   r�   ÚgraphÚsizevarsÚ	shape_envÚcreate_symintnode)rè   rn  s     rW   Úconvert_to_symintru    se   € õ ô �aœÔð 	
ðô
 ˜!œUŸ]™]Ô+ô �‹Fð	ð —‘×!Ñ!×+Ñ+×=Ñ=¸aÀdÐ=ÓKðr{   c                ó>   — | D �cg c]  }t        |«      ‘Œ c}S c c}w )zz
    Takes a list of shapes from Inductor and converts them into symints (or just
    ints if all shapes are static).
    )ru  rc  s     rW   Úconvert_shape_to_symintrw  *  s   € ð +.Ö. QÔ˜aÕ Ò.Ð.ùÒ.s   …c                óN   — t        d„ | j                  j                  D «       «      S )z-
    Does this op overload have aliasing
    c              3  ó8   K  — | ]  }|j                   d u–— Œ y ­wr‹   )Ú
alias_info©r»   r   s     rW   r½   zis_view.<locals>.<genexpr>8  s   è ø€ ÒF¨Aˆq�|‰| 4Ô'ÑFùs   ‚)ÚanyÚ_schemaÚ	arguments©Úops    rW   Úis_viewr�  4  s   € ô ÑF°·±×1EÑ1EÔFÓFÐFr{   c                 ó   — y©NFrš   )rã   s    rW   ú<lambda>r„  =  ó   � r{   c                óì  ‡— | j                   dk7  ryt        | j                  t        j                  j
                  «      s| j                  t        j                  u syt        t        j                  j
                  | j                  «      }|t        j                  u st        |«      rt        ˆfd„| j                  D «       «      S t        j                  j                  |j                  v xs  ‰|«      S )z™
    Do all uses of this op have torch.Tag.pointwise or return True for optional `is_pointwise_fn`

    Uses in views ops will follow the views uses
    Úcall_functionFc              3  ó6   •K  — | ]  }t        |‰«      –— Œ y ­wr‹   )Úis_pointwise_use)r»   ÚuÚis_pointwise_fns     €rW   r½   z#is_pointwise_use.<locals>.<genexpr>N  s   øè ø€ ÒK¸AÔ# A ×7ÑKùs   ƒ)r€  r}   ÚtargetrP   Ú_opsÚ
OpOverloadr  Úgetitemr   r�  r�   ÚusersÚTagÚ	pointwiseÚtags)Úuser‹  rŒ  s    ` rW   r‰  r‰  ;  s¦   ø€ ð ‡v�v�Ò Øä�3—:‘:œuŸz™z×4Ñ4Ô5¸¿¹Äx×GWÑGWÑ9Wàä”%—*‘*×'Ñ'¨¯©Ó4€FØ”×!Ñ!Ñ!¤W¨V¤_ÜÓKÀÇÁÔKÓKÐKä�9‰9×Ñ &§+¡+Ð-ÒH±ÀÓ1HÐHr{   ú	list[Any]c           	     óÂ  ‡‡— t         j                  j                  «       Šg Šdˆˆfd„} ‰j                  | gt	        t         j
                  |||f«      ¢­Ž }t        | j                  j                  «      dk(  r2t        | j                  j                  d   j                  «      dk(  r|f}‰j                  |«       t         j                  j                  i ‰«      }|‰fS )Nc                ó`   •— ‰j                  | «       ‰j                  dt        ‰«      › �«      S )NÚarg)ÚappendÚplaceholderrR   )r˜  ÚgÚ
graph_argss    €€rW   Úadd_tensor_argz)gen_gm_and_inputs.<locals>.add_tensor_argY  s,   ø€ Ø×Ñ˜#ÔØ�}‰}˜s¤3 z£?Ð"3Ð4Ó5Ð5r{   r7   r   ÚTensor)r˜  útorch.Tensorr’   r5   )rP   ÚfxÚGraphr‡  r#   rž  rR   r}  Úreturnsr
  r  Úoutputr4   )rŒ  r„   Úkwargsr�  rg  Úgmr›  rœ  s         @@rW   Úgen_gm_and_inputsr¦  S  sº   ù€ ô 	�‰�‰Ó€AØ%'€Jö6ð ˆ1�?‰?ØðÜœuŸ|™|¨^¸dÀF¸^ÓLò€Dô 	ˆF�N‰N×"Ñ"Ó# qÒ(Ü�—‘×&Ñ& qÑ)×.Ñ.Ó/°8Ò;àˆwˆØ‡H�HˆT„Nä	�‰×	Ñ	˜b !Ó	$€BØˆzˆ>Ðr{   c                óh   — | dk(  ry t        | «      }|j                  «       r|j                  «        y y ©Nr  )rY   rQ   rÁ   r  s     rW   rÁ   rÁ   k  s4   € Ø�‚ØÜ/°Ó7ÐØ×$Ñ$Ô&Ø×$Ñ$Õ&ð 'r{   c                óä   — t        |«       t        j                  d«       t        j                  «       }t        |«      D ]  } | |Ž }t        |«       Œ t        j                  «       }€J ‚||z
  S )Ni9  )rÁ   rP   Úmanual_seedÚtimeÚperf_counterrÆ   )ÚmodelÚexample_inputsrë   rª   Út0rã   ÚresultÚt1s           rW   Útimedr²  s  sr   € ô �ÔÜ	×Ñ�dÔÜ	×	Ñ	Ó	€BÜ�5‹\ò ˆÙ˜Ð'ˆÜ�FÕðô 
×	Ñ	Ó	€BàÐÐÐØ�‰7€Nr{   c                óì   — t        j                  t        |«      D �cg c]  }t        | |||«      ‘Œ c}«      }t        j                  |«      |z  }t        ||z  d›«       |j                  «       S c c}w )Nz.6f)rP   rÏ   rÆ   r²  ÚmedianÚprintrÒ   )	r­  r®  rë   ÚrepeatÚbaselinerª   rã   ÚtimingsÚtooks	            rW   Úprint_performancerº  …  sg   € ô �l‰lÜ>CÀF»mÖL¸Œˆu�n e¨VÕ	4ÒLó€Gô �<‰<˜Ó  5Ñ(€DÜ	ˆT�H‰_˜SÐ!Ô#Ø�9‰9‹;Ðùò	 	Ms   �A1c                óH   ‡—  t        | |«      «       Št        | |ˆfd„«       y)zKReplace obj.method() with a new method that returns a precomputed constant.c                 ó   •— ‰ S r‹   rš   )r°  s   €rW   r„  z#precompute_method.<locals>.<lambda>˜  s   ø€  € r{   N)rO   Úsetattr)ÚobjÚmethodr°  s     @rW   Úprecompute_methodrÀ  •  s    ø€ à!ŒW�S˜&Ó!Ó#€FÜˆC�›Õ(r{   c                ó*   — |D ]  }t        | |«       Œ y)zFReplace methods with new methods that returns a precomputed constants.N)rÀ  )r¾  Úmethodsr¿  s      rW   Úprecompute_methodsrÃ  ›  s   € àò 'ˆÜ˜#˜vÕ&ñ'r{   c                ó<   — t        | |kD  «      t        | |k  «      z
  S r‹   )rŒ   ©r   r!  s     rW   ÚcmprÆ  ¡  s   € Üˆq�1‰u‹:œ˜A ™E›
Ñ"Ð"r{   c                ó~   — t        | t        «      r| g|z  S t        | «      dk(  r t        | «      | d   g«      |z  S | S )Nr7   r   )r}   rŒ   rR   r  )rT   Úsizes     rW   Úpad_listlikerÉ  ¥  sC   € Ü�!”SÔØˆs�T‰zÐÜ
ˆ1ƒv�‚{ØŒt�A‹w˜˜!™�v‹ Ñ%Ð%Ø€Hr{   c                óD   — t        | «      dk(  rg S dd„}t        | |¬«      S )Nr   c                ón   — t        | t        «      r| S ddlm} t        | |«      sJ ‚| j	                  «       S )Nr7   )rF   )r}   r
  Ú	schedulerrF   Úget_name)ÚelemrF   s     rW   Ú	sort_funcztuple_sorted.<locals>.sort_func²  s1   € Ü�dœCÔ ØˆKå0ä˜$Ð 1Ô2Ð2Ð2Ø�}‰}‹Ðr{   ©r^  )rÎ  rl   r’   r
  )rR   Úsorted)rT   rÏ  s     rW   Útuple_sortedrÒ  ®  s&   € Ü
ˆ1ƒv�‚{Øˆ	óô �!˜Ô#Ð#r{   ÚPÚRV)Ú	covariantc                  ó&   — e Zd Zedd„«       Zdd„Zy)ÚCachedMethodc                 ó   — y r‹   rš   )rà   s    rW   Úclear_cachezCachedMethod.clear_cacheÄ  s   € Ø),r{   c                 ó   — y r‹   rš   ©Úselfr„   r¤  s      rW   Ú__call__zCachedMethod.__call__Ç  r…  r{   N)rà   r   r’   ÚNone)r„   úP.argsr¤  úP.kwargsr’   rÔ  )r“   r”   r•   ÚstaticmethodrÙ  rÝ  rš   r{   rW   r×  r×  Ã  s   „ ØÚ,ó Ø,äDr{   r×  c           	     óØ   ‡— | j                   }d|› d�Šd| i}t        d|› d‰› d‰› d�j                  «       |«        t        j                  | «      ||› d�   «      }d
ˆfd	„}||_        |S )NÚ__Ú_cacherÝ   z        def zC_cache_on_self(self):
            try:
                return self.zy
            except AttributeError:
                pass
            rv = fn(self)
            object.__setattr__(self, "z%", rv)
            return rv
        Ú_cache_on_selfc                ó8   •— t        | ‰«      rt        | ‰«       y y r‹   ©r  Údelattr©rÜ  r^  s    €rW   rÙ  z"cache_on_self.<locals>.clear_cacheà  s   ø€ Ü�4˜ÔÜ�D˜#Õð r{   ©rÜ  r   r’   rÞ  ©r“   ÚexecÚlstripr  ÚwrapsrÙ  )rÝ   rÛ   ÚctxÚwrapperrÙ  r^  s        @rW   Úcache_on_selfrñ  Ë  sš   ø€ Ø�;‰;€DØˆtˆf�FÐ
€Cð �ˆ*€CÜðØˆFð à ˜Eð "'ð (+ eð ,	ð		÷ ‰F‹HØôð "Œi�o‰o˜bÓ! #¨¨¨nÐ&=Ñ">Ó?€Gõð &€GÔØ€Nr{   c                ó   — t        | «      S )z]
    Variant of cache_on_self for properties. The only difference is the type signature.
    )rñ  )rÝ   s    rW   Úcache_property_on_selfró  è  s   € ô ˜ÓÐr{   c                ó   ‡ — 	 	 	 	 dˆ fd„}|S )Nc           	     óÖ   •‡— d‰› d| j                   › d�Šd| i}t        d‰› d‰› d‰› d�j                  «       |«        t        j                  | «      |d	   «      }dˆfd
„}||_        |S )Nrã  rã   rä  rÝ   z¶            def inner(self: Any, *args: P.args, **kwargs: P.kwargs) -> RV:
                args_kwargs = (args, tuple(sorted(kwargs.items())))

                if not hasattr(self, "z2"):
                    object.__setattr__(self, "z%", {})

                cache = self.zþ

                try:
                    return cache[args_kwargs]
                except KeyError:
                    pass

                rv = fn(self, *args, **kwargs)

                cache[args_kwargs] = rv
                return rv
            Úinnerc                ó8   •— t        | ‰«      rt        | ‰«       y y r‹   rç  ré  s    €rW   rÙ  z<cache_on_self_and_args.<locals>.wrapper.<locals>.clear_cache  s   ø€ Ü�t˜SÔ!Ü˜˜cÕ"ð "r{   rê  rë  )rÝ   rï  rö  rÙ  r^  Ú
class_names       @€rW   rð  z'cache_on_self_and_args.<locals>.wrapper÷  s•   ù€ ð �:�,˜a §¡˜}¨FÐ3ˆð �RˆjˆÜð'ð (+ eð ,/Ø/2¨eð 4à!˜Uð #ð÷$ ‘“Øô)	
ð, $”	—‘ Ó# C¨¡LÓ1ˆõ	#ð (ˆÔØˆr{   )rÝ   úFN_TYPE[P, RV]r’   rù  rš   )rø  rð  s   ` rW   Úcache_on_self_and_argsrú  ò  s    ø€ ð
$Øð$à	õ$ðL €Nr{   c           
     ó^  — ddl m} t        | t        «      rgt	        j
                  t        j                  | D �cg c]0  }t        |d«      r"|j                  r|j                  j                  ‘Œ2 c}t        «       «      S t        | |j                  «      r| j                  S t        «       S c c}w )Nr7   ©Úirrg  )Ú rý  r}   r]  r  r  r  Úor_r  rg  Úoriginsr!   r?   )Únode_schedulerý  rg  s      rW   Úaggregate_originsr     s�   € õ ä�-¤Ô&Ü×ÑÜ�L‰Lð *öð Ü˜4 Ô(¨T¯YªYð —	‘	×!Ó!òô ‹Ló	
ð 		
ô 
�M 2§?¡?Ô	3Ø×$Ñ$Ð$ä‹|Ðùòs   ¹5B*
c                ó  — t        | «      }|dk(  rYd„ }|D �cg c]6  }|j                  dk(  r%d|j                  v r|j                  d   � ||«      ‘Œ8 }}t        t	        |«      «      }nþ|dk(  rÄg }|D ]¨  }|j                  dk(  sŒd }d}d|j                  v r|j                  d   d   }n"d|j                  v r|j                  d   d   }d	}|sŒ]t        |d
   t        «      r|j                  |d
   |z   «       Œˆ|j                  |d
   j                  |z   «       Œª t        t	        |«      «      }n5|dk(  r*|D �cg c]  }|j                  dk(  sŒ|j                  ‘Œ  }}nt        ‚dj                  dg|z   «      S c c}w c c}w )NÚoriginal_atenc                ó  — | j                   d   }d}t        |t        j                  j                  «      r|j
                  j                  }|S t        |t        j                  j                  «      rt        |j                  «       «      }|S )Nr  rþ  )
r  r}   rP   r�  rŽ  Ú_overloadpacketr“   ÚHigherOrderOperatorr
  rÛ   )Úoriginr  r^  s      rW   Úget_origin_meta_strz2get_fused_kernel_name.<locals>.get_origin_meta_str=  so   € Ø"ŸK™K¨Ñ8ˆMØˆCÜ˜-¬¯©×)>Ñ)>Ô?Ø#×3Ñ3×<Ñ<�ð ˆJô ˜M¬5¯:©:×+IÑ+IÔJÜ˜-×,Ñ,Ó.Ó/�ØˆJr{   r‡  rP   rþ  Úsource_fn_stackr³   Úfwd_source_fn_stackÚbackwardr7   Úinductor_noderã   Úfused)r  r€  r  rÑ  r!   r}   r
  r™  r“   rÛ   ÚNotImplementedErrorÚjoin)r  Údescriptive_namesÚall_originsr	  r  ÚsourcesÚ	source_fnÚsuffixs           rW   Úget_fused_kernel_namer  6  sŸ  € ô $ MÓ2€KØ˜OÒ+ò	ð &ö
àØ�y‰y˜OÒ+Ø 6§;¡;Ñ.Ø—‘˜OÑ,Ð8ñ	   Õ'ð
ˆð 
ô œ GÓ,Ó-‰Ø	˜gÒ	%àˆØ!ò 	CˆFØ�y‰y˜OÓ+Ø �	Ø�Ø$¨¯©Ñ3Ø &§¡Ð,=Ñ >¸rÑ B‘IØ*¨f¯k©kÑ9à &§¡Ð,AÑ BÀ2Ñ F�IØ'�FÙ ØÜ˜i¨™l¬CÔ0Ø—N‘N 9¨Q¡<°&Ñ#8Õ9à—N‘N 9¨Q¡<×#8Ñ#8¸6Ñ#AÕBð	Cô" œ GÓ,Ó-‰Ø	˜oÒ	-à&1ö
Ø"°V·Y±YÀ/Ó5QˆF�K‹Kð
ˆñ 
ô "Ð!Ø�8‰8�W�I Ñ'Ó(Ð(ùòG
ùò<
s   ˜;E=Ä<FÅFc                ó¨  ‡‡ ‡!‡"— t        | «      }|D �cg c]  }|j                  dk(  sŒ|‘Œ }}t        j                  t        «      }t        j                  t        «      }dŠ!|rt        d„ |D «       «      }t        |«      dk(  r_|d   j                  Š!t        ‰!d«      s/t        ‰!j                  «      D ��	ci c]  \  }}	|	|“Œ
 }
}}	|
‰!_        |j                  ˆ!fd„¬«       |D �]Z  }d	|j                  v r·|j                  d	   �¨|j                  d	   }d}t        |t        j                   j"                  «      rt%        |j&                  «      }n=t        |t        j                   j(                  «      rt%        |j+                  «       «      }|r||   j-                  |j*                  «       d
|j                  v r<|j                  d
   d   j*                  }||   j-                  |j*                  «       �Œ|j                  j/                  d«      dk(  s�Œ3||j*                     j-                  |j*                  «       �Œ] ‰!�dnd}‰j0                  › d|› ddj3                  |j5                  «       «      › ddj3                  |j5                  «       «      › d�}‰j0                  › d�g}t7        |j9                  «       «      D ]@  \  }}|j-                  ‰j0                  › d|› ddj3                  t7        |«      «      › �«       ŒB ‰!��«ddlmŠ  |j-                  ‰j0                  › d�«       t        «       }g }t        | ‰ j>                  «      �s·ddl m!} 	 	 	 	 	 	 d*ˆ fd„}d+d„Š"d,ˆ"fd„}| D �]•  }	t        |	d«      r|	jD                  €Œt        |	jD                  d«      rÆ|	jD                  jF                  �°|	jD                  jF                  D ]—  }|j*                  |v rŒ|jI                  |j*                  «       |j                  jK                  |j*                  «      }|€ŒU |||j*                  «      \  }}|j-                  ‰j0                  › d|› d  ||«      › d!|› d�«       Œ™ t        |	jD                  d"«      s�Œ|	jD                  jL                  €�Œ)|	jD                  jL                  D ]T  }|j                  jK                  |j*                  «      }|€Œ+ |||j*                  «      \  }}|j-                  d#|z   «       ŒV �Œ˜ |D ]y  }|jO                  d$¬%«      }|�Ct        jP                  jR                  r)|jU                  ˆfd&„|jW                  «       D «       «       ŒZ|j-                  ‰j0                  › d|› �«       Œ{ |j-                  ‰j0                  › d'd(j3                  |«      › �«       |d)j3                  |«      fS c c}w c c}	}w )-aH  
    Retrieves metadata information for a kernel.
    Args:
        node_schedule (Union[Sequence[BaseSchedulerNode], ExternKernel]):
            Either a sequence of BaseSchedulerNode objects or an ExternKernel instance.
        wrapper (PythonWrapperCodegen):
            An instance of PythonWrapperCodegen, used to define the code comment format.
    Returns:
        tuple[str, str]:
            A tuple containing two strings:
                - The first string represents the kernel's metadata.
                - The second string represent the kernel's detailed metadata.
    r‡  Nc              3  ó4   K  — | ]  }|j                   –— Œ y ­wr‹   )rq  )r»   Úns     rW   r½   z&get_kernel_metadata.<locals>.<genexpr>Š  s   è ø€ Ò"C¨q 1§7¥7Ñ"Cùr¾   r7   r   Ú)_inductor_kernel_metadata_node_to_idx_mapc                ó"   •— ‰j                   |    S r‹   )r  )r  Úsingle_graphs    €rW   r„  z%get_kernel_metadata.<locals>.<lambda>’  s   ø€ ˜l×TÑTÐUVÑW€ r{   rÐ  r  Ú	from_nodeÚpartitioner_tagÚis_backwardzTopologically SortedÚUnsortedú z Source Nodes: [r*  z], Original ATen: [ú]z" Source node to ATen node mapping:ú   z => rü  z Graph fragment:rm  c                ó>  •— t        | ‰j                  «      rAt        | j                  ‰j                  «      r!| j                  j                  j                  }n| j                  }|€|}n|j
                  }	 | j                  «       }||fS # t        $ r d }Y ||fS w xY wr‹   )r}   Ú	TensorBoxÚdataÚ
StorageBoxÚorigin_noderÛ   Ú
get_layoutr  )ÚbufferÚrw_namer(  rÛ   Úlayoutrý  s        €rW   Úget_buffer_infoz,get_kernel_metadata.<locals>.get_buffer_info¼  sž   ø€ ô ˜f b§l¡lÔ3¼
Ø—K‘K §¡ô9ð #)§+¡+×"2Ñ"2×">Ñ">‘Kà"(×"4Ñ"4�KØÐ&à"‘Dà&×+Ñ+�Dð"Ø#×.Ñ.Ó0�Fð ˜V�|Ð#øô +ò "Ø!‘FØ˜V�|Ð#ð"ús   Á7B ÂBÂBc           	     ód   — ddj                  | D �cg c]  }t        |«      ‘Œ c}«      › d�S c c}w )Nú[r*  r"  )r  r
  )ÚshaperT   s     rW   Ústringify_shapez,get_kernel_metadata.<locals>.stringify_shapeÐ  s-   € Ø˜4Ÿ9™9°eÖ%<°¤c¨!¥fÒ%<Ó=Ð>¸aÐ@Ð@ùÒ%<s   ‘-
c                óª   •— | €y ‰| j                   «      › } ‰| j                  «      › }| j                  › }dt        | j                     › |› |› |› d�S )Nrþ  ú")rÈ  Ústriderª   r    r©   )r,  Úshape_annotationÚstride_annotationÚdevice_annotationr1  s       €rW   Ústringfy_layoutz,get_kernel_metadata.<locals>.stringfy_layoutÓ  sl   ø€ Ø�>ØÙ&5°f·k±kÓ&BÐ%CÐ Ù'6°v·}±}Ó'EÐ&FÐ!Ø'-§}¡} oÐ!ð œ F§L¡LÑ1Ð2Ð3CÐ2DØ(Ð)Ð*;Ð)<¸Að?ðr{   Úread_writesÚreadsz   %z
 : Tensor z = PlaceHolder[target=Úwritesú%T)Úinclude_tensor_metadatac              3  ó@   •K  — | ]  }‰j                   › d |› �–— Œ y­w)r#  N)Úcomment)r»   Úlinerð  s     €rW   r½   z&get_kernel_metadata.<locals>.<genexpr>  s)   øè ø€ ò )àð —‘Ð' s¨4¨&Ô1ñ)ùs   ƒz
   return ú,ú
)r*  z-ir.TensorBox | ir.Buffer | ir.TorchBindObjectr+  r
  r’   ztuple[str, ir.Layout | None])r0  zIterable[int]r’   r
  )r,  zir.Layout | Noner’   r
  ),r  r€  ÚcollectionsÚdefaultdictr]  r!   rR   rq  r  Ú	enumerateÚnodesr  Úsortr  r}   rP   r�  rŽ  r
  r  r  rÛ   r™  Úgetr?  r  ÚkeysrÑ  Úitemsrþ  rý  r?   rp  rn  r9  r:  ÚaddÚtry_get_bufferr;  Úformat_nodeÚversionrq   ÚextendÚ
splitlines)#r  rð  r  r  Úinductor_nodesÚfrom_node_dictÚoriginal_aten_dictÚunique_graphsÚidxr  Únode_to_idx_maprg  r  r^  Úsort_strÚmetadataÚdetailed_metadataÚoriginal_noderF  Ú	all_readsÚ
all_writesrn  r-  r8  Úrr*  Ú
input_namer,  ÚwÚoutput_namerã   Úformatted_noderý  r  r1  s#    `                              @@@rW   Úget_kernel_metadatarb  m  sZ  û€ ô$ $ MÓ2€KØ+6ÖW ¸&¿)¹)ÀÓ:V’fÐW€NÐWä ×,Ñ,¬TÓ2€NÜ$×0Ñ0´Ó6Ðð
 €LÙÜ"Ñ"C°NÔ"CÓCˆÜˆ}Ó Ò"Ø)¨!Ñ,×2Ñ2ˆLä˜<Ð)TÔUÜ8AÀ,×BTÑBTÓ8U×"V©f¨c°1 1 c¡6Ð"V�Ñ"VØIX�ÔFØ×ÑÛWð  ô ð ó 8ˆØ˜dŸi™iÑ'¨D¯I©I°oÑ,FÐ,RØ ŸI™I oÑ6ˆMØˆCÜ˜-¬¯©×)>Ñ)>Ô?Ü˜-×7Ñ7Ó8‘Ü˜M¬5¯:©:×+IÑ+IÔJÜ˜-×,Ñ,Ó.Ó/�ÙØ" 3Ñ'×.Ñ.¨t¯y©yÔ9Ø˜$Ÿ)™)Ñ#Ø—)‘)˜KÑ(¨Ñ+×0Ñ0ˆCØ˜3Ñ×&Ñ& t§y¡yÖ1Ø�Y‰Y�]‰]Ð,Ó-°Ô>à˜4Ÿ9™9Ñ%×,Ñ,¨T¯Y©YÖ7ð8ð  *6Ð)AÑ%Àz€Hà�?‰?Ð
˜1˜X˜JÐ&6°t·y±yÀ×ATÑATÓAVÓ7WÐ6Xð YØŸ9™9Ð%7×%<Ñ%<Ó%>Ó?Ð@Àð	Cð ð $ŸO™OÐ,Ð,NÐOÐPÐÜ & ~×';Ñ';Ó'=Ó >ò 
Ñˆ�uØ× Ñ Ø�‰Ð˜s = /°°d·i±iÄÀuÃÓ6NÐ5OÐPõ	
ð
ð ÑÝà× Ñ  G§O¡OÐ#4Ð4DÐ!EÔFÜ%/£\ˆ	Ø "ˆ
Ü˜-¨¯©Õ9Ý&ð$ØEð$ØPSð$à-õ$ó(Aõ
ð #ó =�Ü˜q -Ô0°A·M±MÐ4IØÜ˜1Ÿ=™=¨'Ô2°q·}±}×7JÑ7JÐ7VØŸ]™]×0Ñ0ò ˜àŸ6™6 YÑ.Ø$Ø!Ÿ™ a§f¡fÔ-Ø!"§¡×!7Ñ!7¸¿¹Ó!?˜Ø!˜>Ø$Ù-<¸VÀQÇVÁVÓ-LÑ*˜
 FØ)×0Ñ0Ø&Ÿ™Ð/¨t°J°<¸zÙ.¨vÓ6Ð7Ð7MÈjÈ\ÐYZð\õðô ˜AŸM™M¨8Ö4ØŸ™×,Ñ,Ò8àŸ]™]×1Ñ1ò =˜Ø!"§¡×!7Ñ!7¸¿¹Ó!?˜Ø!˜>Ø$Ù)8¸ÀÇÁÓ)H™˜ Qà"×)Ñ)¨#°Ñ*;Õ<ò=ð-=ð< #ò 	RˆDØ!×-Ñ-ÀdÐ-ÓKˆNØÐ)¬e¯m©m×.?Ò.?ð "×(Ñ(ó )à .× 9Ñ 9Ó ;ô)õ ð
 "×(Ñ(¨G¯O©OÐ+<¸CÀÐ?OÐ)PÕQð	Rð 	× Ñ  G§O¡OÐ#4°J¸s¿x¹xÈ
Ó?SÐ>TÐ!UÔVà�T—Y‘YÐ0Ó1Ð1Ð1ùòY Xùó #Ws   ”W	©W	Â8Wc                óâ   — t        | «      } t        | «      }| rV| j                  «       }|j                  D ]4  }|r	 ||«      rŒ||vsŒ|j	                  |«       | j                  |«       Œ6 | rŒV|S )zJReturns the set of nodes whose values depend on those within initial_queue)r]  r!   rS   r�  rK  r™  )Úinitial_queueÚskip_filterÚdominated_setrg  Úusers        rW   Údominated_nodesrh    sz   € ô
 ˜Ó'€MÜ˜}Ó-€Má
Ø× Ñ Ó"ˆØ—J‘Jò 	+ˆDÙ™{¨4Ô0ØØ˜=Ò(Ø×!Ñ! $Ô'Ø×$Ñ$ TÕ*ð	+ò ð Ðr{   c                ó4  ‡‡	— ddl mŠ dˆˆ	fd„Š	t        |«      \  }}|D �cg c]  } ‰	|«      sŒ|j                  ‘Œ }}t        | «      \  }}|D �cg c]  } ‰	|«      sŒ|j                  ‘Œ }}t	        t        j                  g |¢|¢­Ž «      S c c}w c c}w )Nr7   rü  c                óF  •— t        | ‰j                  «      r ‰| j                  «      S t        | ‰j                  «      r ‰| j                  «      S t        | ‰j                  «      xr9 t        | ‰j
                  ‰j                  ‰j                  ‰j                  f«       S r‹   )	r}   r%  r&  r'  r@   ÚComputedBufferÚInputsKernelÚInputBufferÚTemplateBuffer)r  rý  Úis_unrealized_nodes    €€rW   ro  z*gather_origins.<locals>.is_unrealized_node(  s‰   ø€ Ü�a˜Ÿ™Ô&Ù% a§f¡fÓ-Ð-Ü�a˜Ÿ™Ô'Ù% a§f¡fÓ-Ð-Ü˜!˜RŸY™YÓ'ò 
´
Øà×!Ñ!Ø—‘Ø—‘Ø×!Ñ!ð	ó1
ð -
ð 	
r{   )r  r@   r’   r1  )rþ  rý  r"   r   r!   Ú	itertoolsÚchain)
r„   r¤  Úkwargs_flattenrã   rj  Úkwargs_originsÚargs_flattenÚargs_originsrý  ro  s
           @@rW   Úgather_originsrv  #  s�   ù€ õ ö
ô" % VÓ,Ñ€N�AØ-;ÖW cÑ?QÐRUÕ?V�c—k“kÐW€NÐWÜ" 4Ó(�O€L�!Ø+7ÖS CÑ;MÈcÕ;R�C—K“KÐS€LÐSÜ”i—o‘oÐE |ÐE°nÒEÓFÐFùò XùâSs   ¢B°BÁBÁ Bc                óJ   ‡‡‡‡— dd„Šdˆˆfd„Šdˆˆfd„Šdˆfd„Š ‰| «      S )zµ
    Normal sympy str is very slow, this is a lot faster.  The result are
    somewhat worse, as it doesn't do as much simplification.  So don't
    use this for final codegen.
    c                ó’   — t        | t        j                  «      xr, t        | j                  «      dk(  xr | j                  d   dk(  S )Né   r   r³   )r}   r~   ÚMulrR   r„   )rh  s    rW   Úis_neg_leadzsympy_str.<locals>.is_neg_leadG  s:   € ä�tœUŸY™YÓ'ÒV¬C°·	±	«N¸aÑ,?ÒVÀDÇIÁIÈaÁLÐTVÑDVð	
r{   c                ó`  •— t        | t        j                  «      rŒt        | j                  «      dk(  rO ‰| j                  d   «      r: ‰| j                  d   «      › d ‰| j                  d   j                  d   «      › �S dj                  t        ‰| j                  «      «      S  ‰| «      S )Nry  r7   r   z - z + )r}   r~   r   rR   r„   r  r‚   )rh  r{  Úsympy_str_muls    €€rW   Úsympy_str_addz sympy_str.<locals>.sympy_str_addL  s‘   ø€ Ü�dœEŸI™IÔ&ô �4—9‘9‹~ Ò"¡{°4·9±9¸Q±<Ô'@Ù'¨¯	©	°!©Ó5Ð6°c¹-ÈÏ	É	ÐRSÉ×HYÑHYÐZ[ÑH\Ó:]Ð9^Ð_Ð_à—z‘z¤# m°T·Y±YÓ"?Ó@Ð@á  Ó&Ð&r{   c                óÒ   •— t        | t        j                  «      rE ‰| «      rd ‰| j                  d   «      › �S dj	                  t        ‰| j                  «      «      S  ‰| «      S )Nú-r7   z * )r}   r~   rz  r„   r  r‚   )rh  r{  Úsympy_str_atoms    €€rW   r}  z sympy_str.<locals>.sympy_str_mulW  s[   ø€ Ü�dœEŸI™IÔ&Ù˜4Ô ð ™>¨$¯)©)°A©,Ó7Ð8Ð9Ð9à—z‘z¤# n°d·i±iÓ"@ÓAÐAá! $Ó'Ð'r{   c                ó˜  •— t        | t        j                  «      r| j                  S t        | t        j                  t        j
                  f«      rd ‰| «      › d�S t        | t        t        t        t        f«      rC| j                  j                  › ddj                  t        t        | j                  «      «      › d�S t!        | «      S )Nú(ú)r*  )r}   r~   ÚSymbolrÛ   r   rz  rc   r`   ra   rb   Úfuncr“   r  r‚   Ú	sympy_strr„   r
  )rh  r~  s    €rW   r�  z!sympy_str.<locals>.sympy_str_atomb  s˜   ø€ Ü�dœEŸL™LÔ)Ø—9‘9ÐÜ˜œuŸy™y¬%¯)©)Ð4Ô5Ø‘} TÓ*Ð+¨1Ð-Ð-Ü˜œ´¼(ÄHÐMÔNØ—i‘i×(Ñ(Ð)¨¨4¯9©9´S¼ÀDÇIÁIÓ5NÓ+OÐ*PÐPQÐRÐRä�t“9Ðr{   )rh  r‘   r’   r1  ©rh  r‘   r’   r
  rš   )rh  r{  r~  r�  r}  s    @@@@rW   r‡  r‡  @  s$   û€ ó
ö
	'ö	(õñ ˜ÓÐr{   c                ó¼   — ddl m} t        j                  r3t	        |j
                  dd «      x}r|j                  dk7  rt        | «      S t        j                  «       S )Nr7   rm  Úcurrent_nodeÚ
index_expr)
rp  rn  ri   Úcompute_all_boundsrO   ÚinterpreterrŒ  rf   rg   Úunknown)r  rn  Úfx_nodes      rW   Úget_bounds_index_exprr�  o  sN   € Ýô 	×!Ò!Ü §¡¨~¸tÓDÐDˆWÐDØ�N‰N˜lÒ*ä˜5Ó!Ð!ä×"Ñ"Ó$Ð$r{   c                ó   — | d   dk(  S )Nr   r]  rš   )Úprefixs    rW   Úprefix_is_reductionr“  }  s   € Ø�!‰9˜ÑÐr{   c                óJ   — | t         j                  k7  sJ ‚t        | |dd¬«      S )ú9
    Used to generate an integer-nonnegative symbol.
    T©ÚintegerÚnonnegative)re   ÚSIZErd   )r’  rU  s     rW   Úsympy_index_symbol_with_prefixrš  �  s)   € ð ”T—Y‘YÒÐÐô �v˜s¨D¸dÔCÐCr{   c                óN   — | xs t         j                  xr t         j                  S r‹   )ri   Údebug_index_assertsÚassert_indirect_indexing)Úchecks    rW   Úgenerate_assertrŸ  �  s   € ØÒ/”V×/Ñ/ÒT´V×5TÑ5TÐTr{   c                óF   — | d   dk7  sJ ‚t        j                  | dd¬«      S )r•  r   ré   Tr–  )r~   r…  ©rÛ   s    rW   Úsympy_index_symbolr¢  ‘  s)   € ð �‰7�cŠ>Ðˆ>ô �<‰<˜ d¸Ô=Ð=r{   c                ó   — t        | |«      S )z²
    When the passed replacement symbol v is a string, it is converted to a symbol with name v that
    have the same replaced expression integer and nonnegative properties.
    r&   )rh  Úreplacementss     rW   Ú
sympy_subsr¥  �  s   € ô
 �t˜\Ó*Ð*r{   c                óŠ   — t        | t        j                  «      xs( t        | t        j                  «      xr | j                  S r‹   )r}   rP   r2   rž  Ú_has_symbolic_sizes_strides)r   s    rW   Úis_symbolicr¨  ¥  s3   € Ü�aœŸ™Ó&ò Ü�1”e—l‘lÓ#ÒE¨×(EÑ(Eðr{   c                 ó&   — t        d„ | D «       «      S )Nc              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr‹   )r¨  r{  s     rW   r½   z"any_is_symbolic.<locals>.<genexpr>¬  s   è ø€ Ò, !Œ{˜1�~Ñ,ùó   ‚©r|  )r„   s    rW   Úany_is_symbolicr­  «  s   € ÜÑ, tÔ,Ó,Ð,r{   )z,aten._fused_moving_avg_obs_fq_helper.defaultz7aten._fused_moving_avg_obs_fq_helper_functional.defaultzfbgemm.dense_to_jagged.defaultz%fbgemm.jagged_to_padded_dense.defaultÚrun_and_save_rng_stateÚrun_with_rng_statezaten._local_scalar_densezaten._assert_scalarc                ó¸   — ddl m} | j                  j                  D ];  }t	        |«      r|c S |j
                  j                  d«      x}€Œ0 ||«      sŒ9|c S  y )Nr   )r)   rj  )Ú%torch.fx.experimental.symbolic_shapesr)   rq  rF  Úis_cudagraph_unsafe_fx_noder  rH  )r¥  r)   rg  rj  s       rW   Ú%get_first_incompatible_cudagraph_noder³  Ã  sV   € õ Là—‘—‘ò ˆÜ& tÔ,ØŠKà—9‘9—=‘= Ó'Ð'ˆCÑ4Ñ9NÈsÕ9SØŠKðð r{   c                óŠ   — t        t        t        | j                  j                  «      «      «      }|j
                  dk(  sJ ‚|S )z$Get the output node from an FX graphr£  )ÚnextÚiterÚreversedrq  rF  r€  )r¥  Ú	last_nodes     rW   Úoutput_noder¹  Ò  s6   € ä”Tœ( 2§8¡8§>¡>Ó2Ó3Ó4€IØ�<‰<˜8Ò#Ð#Ð#ØÐr{   c                óæ   — | j                   j                  d¬«      }t        d„ |D «       «      }t        | «      j                  d   }t        |t        «      r|n|f}t        d„ |D «       «      }||z  S )Nrš  r  c              3  ó´   K  — | ]P  }t        |j                  j                  d «      t        j                  «      r|j                  d    j
                  –— ŒR y­w©rj  N)r}   r  rH  rP   rž  rª   )r»   rg  s     rW   r½   z"get_all_devices.<locals>.<genexpr>Û  sB   è ø€ ò 9àÜ�d—i‘i—m‘m EÓ*¬E¯L©LÔ9ð 	�	‰	�%Ñ×Õñ9ùs   ‚AAr   c              3  óü   K  — | ]t  }t        |t        j                  j                  «      rNt        |j                  j                  d «      t        j                  «      r|j                  d    j                  –— Œv y­wr¼  )r}   rP   r   r5   r  rH  rž  rª   )r»   r˜  s     rW   r½   z"get_all_devices.<locals>.<genexpr>ã  sS   è ø€ ò 7àÜ�cœ5Ÿ8™8Ÿ=™=Ô)Ü�s—x‘x—|‘| EÓ*¬E¯L©LÔ9ð 	�‰�‰×Õñ7ùs   ‚A:A<)rq  Ú
find_nodesr!   r¹  r„   r}   Útuple)r¥  Úplaceholder_nodesÚinput_devicesÚout_argÚout_argsÚout_devicess         rW   Úget_all_devicesrÅ  Ù  s}   € ØŸ™×+Ñ+¨}Ð+Ó=ÐÜ.8ñ 9à%ô9ó /€Mô ˜"‹o×"Ñ" 1Ñ%€GÜ$ W¬eÔ4‰w¸7¸*€HÜ,6ñ 7àô7ó -€Kð ˜;Ñ&Ð&r{   c                 ó|  — t        t        j                  j                  «       «      D �]  } | j	                  d«      sŒt        j                  |    }|j
                  D ]Ñ  }|j	                  d«      sŒt        ||«      }t        |t        j                  j                  j                  j                  «      sŒZ|j                  D ]i  }t        |t        j                  j                  j                  j                  «      sŒ<|j                  j                   j"                  j%                  «        Œk ŒÓ t        j                  | = �Œ dt        j                  v rRt        j                  d   }t'        |j(                  j*                  j,                  «      `|j(                  j*                  `t1        j2                  «        y )Nz&torch._inductor.runtime.compile_tasks.Útriton_ztriton.runtime.driver)r]  ÚsysÚmodulesrI  Ú
startswithÚ__dict__rO   r}   rP   Ú	_inductorÚruntimeÚtriton_heuristicsÚCachingAutotunerÚcompile_resultsÚTritonCompileResultÚkernelÚrunÚmodÚ__del__r  ÚdriverÚactiveÚutilsÚinstanceÚgcÚcollect)Úmodule_nameÚmÚ	attr_namerÒ  r°  rÔ  s         rW   Úunload_xpu_triton_pydsrß  ï  sK  € äœCŸK™K×,Ñ,Ó.Ó/ó %ˆØ×%Ñ%Ð&NÔOØÜ�K‰K˜Ñ$ˆØŸ™ò 	<ˆIØ×#Ñ# IÕ.Ü   IÓ.�ÜØœEŸO™O×3Ñ3×EÑE×VÑVõð #)×"8Ñ"8ò <˜Ü%Ø"Ü!ŸO™O×3Ñ3×EÑE×YÑYõð
 #ŸM™M×-Ñ-×1Ñ1×9Ñ9Õ;ñ<ð	<ô �K‰K˜Ò$ð#%ð( ¤#§+¡+Ñ-Ü�k‰kÐ1Ñ2ˆÜ�—‘×"Ñ"×(Ñ(Ó)Ð2Ø�J‰J×ÑÐ#ä‡J�J…Lr{   Ú_registered_cachesc                óŽ   — t        | d«      rt        | j                  «      st        | › d�«      ‚t        j                  | «       | S )zh
    Use this decorator to register any caches that should be cache_clear'd
    with fresh_cache().
    Úcache_clearz# does not have a cache_clear method)r  Úcallablerâ  ÚAttributeErrorrà  r™  ©r¾  s    rW   Úclear_on_fresh_cacheræ    s?   € ô
 �3˜Ô&¬h°s·±Ô.GÜ ˜uÐ$GÐHÓIÐIä×Ñ˜cÔ"Ø€Jr{   c                 ó:   — t         D ]  } | j                  «        Œ y)z&
    Clear all registered caches.
    N)rà  râ  rå  s    rW   Úclear_cachesrè    s   € ô "ò ˆØ�‰Õñr{   c              #  ó\  K  — t         j                  j                  | «      }	 |t         j                  | <   d–— |€!t         j                  j                  | d«       y|t         j                  | <   y# |€!t         j                  j                  | d«       w |t         j                  | <   w xY w­w)a×  Thread-safe env var set/restore using atomic C-level lookups.

    We avoid mock.patch.dict(os.environ, ...) because it internally calls
    os.environ.copy(), which iterates all env var keys then fetches values in
    separate steps. That approach is not atomic and can race with background threads
    (e.g. Triton async compilation) modifying the environment, causing KeyError,
    so we use os.environ.get() for individual keys which is an atomic C-level lookup.
    N)ÚosÚenvironrH  rS   )r^  r�   Úolds      rW   Ú_set_envrí  $  sz   è ø€ ô �*‰*�.‰.˜Ó
€Cð"ØŒ�
‰
�3‰Ûàˆ;Ü�J‰J�N‰N˜3 Õ%à!ŒB�J‰J�sŠOøð ˆ;Ü�J‰J�N‰N˜3 Õ%à!ŒB�J‰J�sŠOüs   ‚ B,£A1 º7B,Á18B)Â)B,c              #  ó  ‡K  — t        «        ddlm}  |t        j                  |¬«      «      Š	 t        d‰«      5  t        j                  d‰«        |t        j                  j                  ‰d«      «      }t        d|«      5  d–— t        | t        «      r¨t        | «      dk(  sJ d	«       ‚t        j                  j                  |«      rtt        j                  |«      }| j!                  |D �ci c]D  }d
|vr>|t        j                  j#                  t        j                  j                  ||«      «      “ŒF c}«       ddd«       ddd«       |rUt%        «       r(t&        j(                  j+                  «       r
t-        «        t/        j0                  ‰t%        «       ˆfd„¬«       t        «        yc c}w # 1 sw Y   ŒxxY w# 1 sw Y   Œ|xY w# t2        $ r t        j5                  d‰«       ‚ w xY w# t        «        w xY w­w)zÈ
    Contextmanager that provides a clean tmp cachedir for pt2 caches.

    Optionally, pass a dict as 'cache_entries' to get a list of filenames and sizes
    generated with this cache instance.
    r   )Únormalize_path_separator)ÚdirÚTORCHINDUCTOR_CACHE_DIRzUsing inductor cache dir %sÚtritonÚTRITON_CACHE_DIRNz!expected empty cache_entries dictz.lockc                ó4   •— t         j                  d‰|¬«      S )Nz*Failed to remove temporary cache dir at %s)Úexc_info)rÓ   Úwarning)r†  Úpathrõ  Úinductor_cache_dirs      €rW   r„  zfresh_cache.<locals>.<lambda>g  s   ø€ ´S·[±[Ø@Ø&Ø%ð 6Aó 6€ r{   )Úignore_errorsÚonerrorz(on error, temporary cache dir kept at %s)rè  Útorch._inductor.cpp_builderrï  ÚtempfileÚmkdtemprí  rÓ   rÔ   rê  r÷  r  r}   ÚdictrR   ÚexistsÚlistdirr\  ÚgetsizeÚ
is_windowsrP   rJ   rQ   rß  ÚshutilÚrmtreeÚ	Exceptionrö  )Úcache_entriesrð  Údeleterï  Útriton_cache_dirÚfilesÚfrø  s          @rW   Úfresh_cacher  9  s¦  øè ø€ ô „NåDá1´(×2BÑ2BÀsÔ2KÓLÐð'ÜÐ/Ð1CÓDñ 	Ü�I‰IÐ3Ð5GÔHÙ7Ü—‘—‘Ð/°Ó:ó Ðô Ð,Ð.>Ó?ñ ÛÜ˜m¬TÔ2Ü˜}Ó-°Ò2ÐWÐ4WÓWÐ2Ü—w‘w—~‘~Ð&6Ô7Ü "§
¡
Ð+;Ó <˜Ø%×,Ñ,ð */öà$%Ø#*°!Ñ#3ð !"¤2§7¡7§?¡?´2·7±7·<±<Ð@PÐRSÓ3TÓ#UÑ Uòô÷÷	ñ$ ÜŒ|¤§	¡	× 6Ñ 6Ô 8Ü&Ô(ä�M‰MØ"ô )›lóõô  	�ùò5÷ð ú÷	ð 	ûôD ò Ü�‰Ð>Ð@RÔSØðûô 	�üsm   ƒ-H±G ½A	GÂA-F:Ã3A	F5Ä<F:ÅGÅAG Æ*HÆ5F:Æ:G	Æ?GÇGÇG Ç!G3Ç3G6 Ç6HÈH)Úreversec               ó    — | j                   }t        t        | «      «      }t        t	        ||d¬«      «      }|st        t        |«      «      S |S )NT©r^  r  )Ú__getitem__rÆ   rR   r]  rÑ  r·  )Úseqr  ÚgetterÚa_rÚsort_idxs        rW   Úargsortr  z  sE   € Ø�_‰_€FÜ
”�C“‹/€Cô ”F˜3 F°DÔ9Ó:€HÙÜ”H˜XÓ&Ó'Ð'Ø€Or{   c          	     ó4  ‡ — dˆ fd„}t        |«      D ��cg c]9  \  }}|t        |t        j                  «      r|j                  j
                  n|f‘Œ; }}}t        |t        j                  |«      |¬«      }|D ��cg c]  \  }}|‘Œ	 }}}|S c c}}w c c}}w )Nc                ón   •— | \  }}|\  }}dˆfd„} |||k  «      ry |||kD  «      ry||k  ry||kD  ryy)Nc                óN   •— t        | t        «      r| S ‰j                  | d¬«      S )NT)Úsize_oblivious)r}   r1  Úevaluate_expr)rh  rs  s    €rW   Úevaluatez*argsort_sym.<locals>.cmp.<locals>.evaluate”  s(   ø€ Ü˜$¤Ô%Ø�Ø×*Ñ*¨4ÀÐ*ÓEÐEr{   r³   r7   r   )rh  z bool | torch.SymInt | sympy.Exprr’   r1  rš   )r   r!  Úa_idxÚa_valÚb_idxÚb_valr  rs  s          €rW   rÆ  zargsort_sym.<locals>.cmp�  sT   ø€ Ø‰ˆˆuØ‰ˆˆuõ	Fñ
 �E˜E‘MÔ"ØÙ�E˜E‘MÔ"Øð
 �5Š=ØØ�5Š=ØØr{   r  )r   útuple[int, sympy.Expr]r!  r  r’   rŒ   )	rE  r}   rP   r2   rg  rh  rÑ  r  Ú
cmp_to_key)	rs  r  r  rÆ  rU  ré   Úexprsrã   r°  s	   `        rW   Úargsort_symr"  Š  sŒ   ø€ õô4   “n÷áˆC�ð 
œZ¨¬5¯<©<Ô8ˆa�f‰f�kŠk¸aÒ@ð€Eñ ô �5œi×2Ñ2°3Ó7ÀÔI€EØ %×&‘f�c˜1ŠcÐ&€FÑ&Ø€Mùóùó
 's   –>BÁ=Bc                ót   — | t         j                  k(  ryt        j                  d| ¬«      j                  «       S )Nrv   rš   ©r©   )rP   rY  rÂ   Úelement_sizer$  s    rW   Úget_dtype_sizer&  ±  s-   € ð ”—‘ÒØÜ�;‰;�r Ô'×4Ñ4Ó6Ð6r{   c                  ó   — e Zd ZU ded<   y)ÚLineContextr   ÚcontextN©r“   r”   r•   r¤   rš   r{   rW   r(  r(  º  s   … Ø„Lr{   r(  c                  ó"   — e Zd ZU ded<   ded<   y)ÚValueWithLineMapr
  r�   zlist[tuple[int, LineContext]]Úline_mapNr*  rš   r{   rW   r,  r,  ¾  s   … àƒJØ+Ô+r{   r,  c                  óÖ   — e Zd ZdZddd„Zej                  dd„«       Zdd„Zdd„Z	dd„Z
d„ Zdd„Zdd	„Zdd
„Zdd„Zdd„Zdd„Zd d!d„Zd d"d„Zd d"d„Zd#d$d„Zd%d„Zdd„Zd&d„Zd'd„Zy)(ÚIndentedBufferé   c                ó    — g | _         || _        y r‹   )Ú_linesÚ_indent)rÜ  Úinitial_indents     rW   Ú__init__zIndentedBuffer.__init__Ç  s   € ØBDˆŒØ%ˆ�r{   c              #  ób   K  — | j                   }	 || _         d –— || _         y # || _         w xY w­wr‹   )Útabwidth)rÜ  r7  Úprevs      rW   Úset_tabwidthzIndentedBuffer.set_tabwidthË  s,   è ø€ à�}‰}ˆð	!Ø$ˆDŒMÛà ˆD�Mø˜DˆD�Müs   ‚/�# ›/£	,¬/c                óž  — t        «       }d}g }| j                  D ]—  }t        |t        «      r
 |«       }|€1Œt        |t        «      r|j                  ||j                  f«       ŒK|}t        |t        «      sJ ‚|j                  |«       |j                  d«       |d|j                  d«      z   z  }Œ™ t        |j                  «       |«      S )Nr7   rB  )r   r2  r}   ÚDeferredLineBaser(  r™  r)  r
  ÚwriteÚcountr,  Úgetvalue)rÜ  Úbufrç   ÚlinemapÚlir@  s         rW   Úgetvaluewithlinemapz"IndentedBuffer.getvaluewithlinemapÔ  s¹   € Ü‹jˆØˆØ13ˆØ—+‘+ò 	&ˆBÜ˜"Ô.Ô/Ù“t�Ø�<ØÜ˜B¤Ô,Ø—‘  2§:¡:˜Ô/Øà�Ü˜d¤CÔ(Ð(Ð(Ø�I‰I�dŒOØ�I‰I�dŒOØ��T—Z‘Z Ó%Ñ%Ñ%‰Að	&ô   §¡£°Ó8Ð8r{   c                ó6   — | j                  «       j                  S r‹   )rB  r�   ©rÜ  s    rW   r>  zIndentedBuffer.getvalueè  s   € Ø×'Ñ'Ó)×/Ñ/Ð/r{   c                óf  — t        «       }| j                  D ]‰  }t        |t        «      r
 |«       }|€Œt        |t        «      rŒ.|}t        |t
        «      sJ ‚|j                  d«      r|j                  |d d «       Œh|j                  |«       |j                  d«       Œ‹ |j                  «       S )Nú\r³   rB  )	r   r2  r}   r;  r(  r
  Úendswithr<  r>  )rÜ  r?  rA  r@  s       rW   ÚgetrawvaluezIndentedBuffer.getrawvalueë  sš   € Ü‹jˆØ—+‘+ò 	 ˆBÜ˜"Ô.Ô/Ù“t�Ø�<ØÜ˜B¤Ô,Øà�Ü˜d¤CÔ(Ð(Ð(à�}‰}˜TÔ"Ø—	‘	˜$˜s ˜)Õ$à—	‘	˜$”Ø—	‘	˜$•ð	 ð  �|‰|‹~Ðr{   c                ó   — | j                   S r‹   ©r2  rD  s    rW   Úget_lines_refzIndentedBuffer.get_lines_refÿ  s   € Ø�{‰{Ðr{   c                ó8   — | j                   j                  «        y r‹   )r2  ÚclearrD  s    rW   rM  zIndentedBuffer.clear  s   € Ø�‰×ÑÕr{   c                ó,   — t        | j                  «      S r‹   )r1  r2  rD  s    rW   Ú__bool__zIndentedBuffer.__bool__  s   € Ü�D—K‘KÓ Ð r{   c                ó:   — d| j                   | j                  z  z  S )Nr!  )r3  r7  rD  s    rW   r’  zIndentedBuffer.prefix  s   € Ø�d—l‘l T§]¡]Ñ2Ñ3Ð3r{   c                ó&   — | j                  d«       y )NrB  ©Ú	writelinerD  s    rW   ÚnewlinezIndentedBuffer.newline  s   € Ø�‰�tÕr{   c                ó   — t        |t        «      r| j                  j                  |«       y t        |t        «      r9| j                  j                  |j                  | j                  «       «      «       y |j                  «       r.| j                  j                  | j                  «       › |› �«       y | j                  j                  d«       y ©Nrþ  )r}   r(  r2  r™  r;  Úwith_prefixr’  Ústrip©rÜ  r@  s     rW   rS  zIndentedBuffer.writeline  sˆ   € Ü�dœKÔ(Ø�K‰K×Ñ˜tÕ$Ü˜Ô.Ô/Ø�K‰K×Ñ˜t×/Ñ/°·±³Ó>Õ?Ø�Z‰ZŒ\Ø�K‰K×Ñ $§+¡+£- °°Ð7Õ8à�K‰K×Ñ˜rÕ"r{   c                ó4   — |D ]  }| j                  |«       Œ y r‹   rR  )rÜ  Úlinesr@  s      rW   Ú
writelineszIndentedBuffer.writelines  s   € Øò 	!ˆDØ�N‰N˜4Õ ñ	!r{   c                óH   ‡ ‡— t         j                  dˆˆ fd„«       } |«       S )Nc               3  ó    •K  — ‰xj                   ‰ z  c_         	 d –— ‰xj                   ‰ z  c_         y # ‰xj                   ‰ z  c_         w xY w­wr‹   ©r3  )ÚoffsetrÜ  s   €€rW   rï  z"IndentedBuffer.indent.<locals>.ctx  s9   øè ø€ à�LŠL˜FÑ"�Lð'Ûà—’ Ñ&–ø�—’ Ñ&–üs   ƒAš4 žA´AÁA©r’   úIterator[None])Ú
contextlibÚcontextmanager)rÜ  r`  rï  s   `` rW   ÚindentzIndentedBuffer.indent  s$   ù€ Ü	×	"Ñ	"õ	'ó 
#ð	'ñ ‹uˆr{   c                ó.   — | xj                   |z  c_         y r‹   r_  ©rÜ  r`  s     rW   Ú	do_indentzIndentedBuffer.do_indent'  ó   € Ø�Š˜ÑŽr{   c                ó.   — | xj                   |z  c_         y r‹   r_  rg  s     rW   Údo_unindentzIndentedBuffer.do_unindent*  ri  r{   c           	     ó’  — t        |t        «      rÖt        d«      }|j                  D ]E  }t        |t        «      rŒ|sŒt        |t        |«      t        |j                  «       «      z
  «      }ŒG t        j                  |«      rd}|j                  D ]P  }t        |t        «      r| j                  j                  |«       Œ/t        j                  | |t        |«      d  «       ŒR y t        j                  |«      }|r|j                  «       }|sy |j                  «       }|j!                  d«      D ]  }| j                  |«       Œ y )NÚinfr   rB  )r}   r/  Úfloatr2  r(  ÚminrR   rí  ÚmathÚisinfr™  rS  rŒ   ÚtextwrapÚdedentÚrstripr[  )rÜ  Ú
other_coderX  rs  r@  ré   s         rW   ÚsplicezIndentedBuffer.splice-  s  € Ü�j¤.Ô1Ü˜5“\ˆFà"×)Ñ)ò I�Ü! $¬Õ4ºÜ  ¬¨T«´S¸¿¹»Ó5GÑ)GÓH‘FðIô �z‰z˜&Ô!Ø�Ø"×)Ñ)ò H�Ü˜d¤KÔ0Ø—K‘K×&Ñ& tÕ,ä"×,Ñ,¨T°4¼¸F»¸Ð3FÕGñ	Hô "Ÿ™¨Ó4ˆJÙØ'×.Ñ.Ó0�
ÙØØ#×*Ñ*Ó,ˆJØ×%Ñ% dÓ+ò "�Ø—‘˜qÕ!ñ"r{   c                ó†   — t        | j                  ¬«      }| j                  D �cg c]
  } ||«      ‘Œ c}|_        |S c c}w ©N)r4  )r/  r3  r2  )rÜ  r†  rì   r@  s       rW   r‚   zIndentedBuffer.mapE  s4   € Ü¨D¯L©LÔ9ˆØ-1¯[©[Ö9 T‘d˜4•jÒ9ˆŒ
Øˆ
ùò :s   ¥>c                ó@   — t        | «      › d| j                  «       › d�S )Nrƒ  r„  )r  r>  rD  s    rW   Ú__repr__zIndentedBuffer.__repr__J  s    € Ü�t“*�˜Q˜tŸ}™}›Ð/¨qÐ1Ð1r{   c                óÔ   — | j                   |j                   k(  sJ ‚t        | j                   ¬«      }|j                  | j                  «       |j                  |j                  «       |S rx  )r3  r/  r\  r2  )rÜ  Úotherrì   s      rW   Ú__add__zIndentedBuffer.__add__M  sK   € Ø�|‰|˜uŸ}™}Ò,Ð,Ð,Ü¨D¯L©LÔ9ˆà�‰�t—{‘{Ô#Ø�‰�u—|‘|Ô$Øˆ
r{   c                ó   — || j                   v S r‹   rJ  )rÜ  Únew_lines     rW   ÚcontainszIndentedBuffer.containsU  s   € Ø˜4Ÿ;™;Ð&Ð&r{   N©r   )r4  rŒ   r’   rÞ  )r7  rŒ   r’   rb  )r’   r,  ©r’   r
  ©r’   rÞ  ©r’   r1  )r@  z$LineContext | DeferredLineBase | strr’   rÞ  )r[  z.Sequence[LineContext | DeferredLineBase | str]r’   rÞ  r‰   )r`  rŒ   r’   ú'contextlib.AbstractContextManager[None])r`  rŒ   r’   rÞ  ©F)ru  zIndentedBuffer | strrX  r1  r’   rÞ  )r†  zCallable[[Any], Any]r’   r/  )r|  r   r’   r/  )r  z$DeferredLineBase | LineContext | strr’   r1  )r“   r”   r•   r7  r5  rc  rd  r9  rB  r>  rH  rK  rM  rO  r’  rT  rS  r\  re  rh  rk  rv  r‚   rz  r}  r€  rš   r{   rW   r/  r/  Ä  s€   „ Ø€Hô&ð ×Ñò!ó ð!ó9ó(0óò(óó!ó4óó#ó!ô	ôôô"ó0ó
2óô'r{   r/  c                  ó(   ‡ — e Zd Zdˆ fd„Zdd„Zˆ xZS )ÚFakeIndentedBufferc                ó"   •— t         ‰| �  «        y r‹   )Úsuperr5  )rÜ  Ú	__class__s    €rW   r5  zFakeIndentedBuffer.__init__Z  s   ø€ Ü‰ÑÕr{   c                óV   — |dk(  rt         j                  | |«      S t        d|› d�«      ‚)Nr‹  zTried to call self.zÊ on FakeIndentedBuffer. This bufferis currently used on TritonTemplateKernel to prevent actualwrites to the body without explicitly specifying the body with`TritonTemplateKernel.set_subgraph_body(name)`)ÚobjectÚ__getattribute__rú   )rÜ  rÛ   s     rW   rŽ  z#FakeIndentedBuffer.__getattribute__]  s;   € Ø�;ÒÜ×*Ñ*¨4°Ó6Ð6ÜØ! $ ð (=ð =ó
ð 	
r{   rƒ  )rÛ   r
  r’   r   )r“   r”   r•   r5  rŽ  Ú__classcell__©r‹  s   @rW   rˆ  rˆ  Y  s   ø„ õ÷
r{   rˆ  c               #  ó¼   K  — t         j                  t         j                  }} 	 d –— | |ct         _        t         _        y # | |ct         _        t         _        w xY w­wr‹   )rÈ  ÚstdoutÚstderr)Úinitial_stdoutÚinitial_stderrs     rW   Úrestore_stdout_stderrr–  h  s@   è ø€ ä%(§Z¡Z´·±�N€Nð@Ûà!/°ÐŒŒ
”C•Jø °ÐŒŒ
”C•Jüs   ‚!A¤A  ¨AÁ AÁAc                  óP   — e Zd ZdZdd„Zdd„Zdd„Zdd„Zdd„Zdd„Z	dd„Z
dd	„Zy
)r;  z.A line that can be 'unwritten' at a later timec                ó6   — |j                  «       sd}|| _        y rV  )rX  r@  rY  s     rW   r5  zDeferredLineBase.__init__t  s   € Ø�z‰zŒ|ØˆDØˆ�	r{   c                ó   — t         ‚)zJReturns either self.line or None to indicate the line has been 'unwritten'©r  rD  s    rW   rÝ  zDeferredLineBase.__call__y  ó   € ä!Ð!r{   c                ó   — t         ‚)z3Returns a new deferred line with the same conditionrš  rY  s     rW   Ú	_new_linezDeferredLineBase._new_line}  r›  r{   c                ó@   — | j                  |› | j                  › �«      S r‹   ©r�  r@  )rÜ  r’  s     rW   rW  zDeferredLineBase.with_prefix�  s   € Ø�~‰~  ¨¯©¨Ð4Ó5Ð5r{   c                óT   — | j                  | j                  j                  «       «      S r‹   )r�  r@  rí  rD  s    rW   rí  zDeferredLineBase.lstrip„  s   € Ø�~‰~˜dŸi™i×.Ñ.Ó0Ó1Ð1r{   c                ó>   — | j                  | j                  |   «      S r‹   rŸ  )rÜ  r  s     rW   r  zDeferredLineBase.__getitem__‡  s   € Ø�~‰~˜dŸi™i¨Ñ.Ó/Ð/r{   c                ó,   — t        | j                  «      S r‹   )r1  r@  rD  s    rW   rO  zDeferredLineBase.__bool__Š  s   € Ü�D—I‘I‹Ðr{   c                ó,   — t        | j                  «      S r‹   )rR   r@  rD  s    rW   Ú__len__zDeferredLineBase.__len__�  s   € Ü�4—9‘9‹~Ðr{   N)r@  r
  ©r’   ú
str | None)r@  r
  r’   r   )r’  r
  r’   r   )r’   r   )r  zint | slicer’   r   r„  ©r’   rŒ   )r“   r”   r•   r–   r5  rÝ  r�  rW  rí  r  rO  r¤  rš   r{   rW   r;  r;  q  s-   „ Ù8óó
"ó"ó6ó2ó0óôr{   r;  c                  ó4   ‡ — e Zd ZdZdˆ fd„Zdd„Zdd„Zˆ xZS )ÚDelayReplaceLinez6At end of codegen call `line.replace(key, value_fn())`c                ó@   •— t         ‰| �  |«       || _        || _        y r‹   )rŠ  r5  r^  Úvalue_fn)rÜ  r^  r«  r@  r‹  s       €rW   r5  zDelayReplaceLine.__init__”  s   ø€ Ü‰Ñ˜ÔØˆŒØ ˆ�r{   c                ój   — | j                   j                  | j                  | j                  «       «      S r‹   )r@  Úreplacer^  r«  rD  s    rW   rÝ  zDelayReplaceLine.__call__™  s#   € Ø�y‰y× Ñ  §¡¨4¯=©=«?Ó;Ð;r{   c                óD   — t        | j                  | j                  |«      S r‹   )r©  r^  r«  rY  s     rW   r�  zDelayReplaceLine._new_lineœ  s   € Ü §¡¨$¯-©-¸Ó>Ð>r{   )r^  r
  r«  zCallable[[], str]r@  r
  r‚  )r@  r
  r’   r©  )r“   r”   r•   r–   r5  rÝ  r�  r�  r�  s   @rW   r©  r©  ‘  s   ø„ Ù@õ!ó
<÷?r{   r©  c                óÜ  — t        | t        j                  «      r| }nt        j                  t        «       | «      }t	        j
                  |«      }t        j                  j                  rC|j                  €J ‚|j                  dk  s|j                  dk(  rt        j                  d«       yy|j                  dk(  rdnd}|j                  }||k  rt        j                  d	||d
œ¬«       yy)Né	   é
   z6GPU arch does not support max_autotune_gemm mode usageFTrJ   rr   éD   z,Not enough SMs to use max_autotune_gemm mode)Úmin_smsÚ	avail_sms)Úextra)r}   rP   rª   rX   r   ÚcreaterN  rq   ÚmajorrÓ   rö  r  Úmulti_processor_count)Úindex_or_devicerª   Úpropr³  r´  s        rW   Ú
is_big_gpur»     sÈ   € ä�/¤5§<¡<Ô0Ø ‰ä—‘œl›n¨oÓ>ˆä×"Ñ" 6Ó*€Dô ‡}�}×ÒØ�z‰zÐ%Ð%Ð%Ø�:‰:˜Š>˜TŸZ™Z¨2Ò-Ü�K‰KÐPÔQØØà—K‘K 5Ò(‰b¨b€GØ×*Ñ*€IØ�7ÒÜ�‰Ø:Ø%°IÑ>ð 	ô 	
ð Ør{   c                 óà   — t         j                  j                  «       r(t         j                  j                  «       j                  S t         j
                  j                  d«      j                  S )NrH   )rP   rJ   rQ   Úget_device_propertiesÚgpu_subslice_countrH   r¸  rš   r{   rW   Úget_max_num_smsr¿  ½  sF   € ä‡y�y×ÑÔÜ�y‰y×.Ñ.Ó0×CÑCÐCÜ�:‰:×+Ñ+¨FÓ3×IÑIÐIr{   c                 óÔ   — t         j                  j                  «       syt         j                  j                  t         j                  j	                  «       «      } | j
                  dk(  S )zEReturns true if the device is a NVIDIA B200, otherwise returns false.Fr±  )rP   rH   rQ   r½  r  r·  )Údevice_propertiess    rW   Ú
using_b200rÂ  Ä  sJ   € ô �:‰:×"Ñ"Ô$ØäŸ
™
×8Ñ8¼¿¹×9RÑ9RÓ9TÓUÐØ×"Ñ" bÑ(Ð(r{   c                 ó´   — t         j                  j                  «       r
t        «       S t         j                  j                  «       } t        «       | �| z
  S dz
  S )zFHandle experimental carveout if set otherwise return hardware SM countr   )rP   rJ   rQ   r¿  r  Ú_get_sm_carveout_experimental)Úcarveouts    rW   Úget_num_smsrÆ  Î  sJ   € ô ‡y�y×ÑÔÜÓ Ð Ü�x‰x×5Ñ5Ó7€HÜÓ¨HÐ,@ ÑHÐHÀaÑHÐHr{   c                óš   — ddl m}m} |€
t        «       }|j	                  d«      }|| z  t
        z  } |||| |j                  «       ¬«      S )zKBuilds and returns a WorkspaceArg for the device side TMA workspace buffer.r7   )r8   ÚWorkspaceZeroModeF)r=  Ú	zero_moderª   Ú
outer_name)Úcodegen.commonr8   rÈ  rÆ  Ú	from_boolÚTMA_DESCRIPTOR_SIZEÚunique_name)Únum_tma_descriptorsrª   Únum_programsr8   rÈ  rÉ  rÈ  s          rW   Úget_tma_workspace_argrÑ  ×  sZ   € ÷ @àÐÜ"“}ˆØ!×+Ñ+¨EÓ2€IØÐ-Ñ-Ô0CÑC€DÙØØØØ+�<×+Ñ+Ó-ô	ð r{   c                óœ   — t         j                  j                  sydt         j                  j	                  d«      j
                  v r| dk  ryy)Nr   Úgfx942rr   r7   ry  )rP   rN  rq   rH   r½  ÚgcnArchName)Úblock_ks    rW   Úget_default_kpackrÖ  ë  s<   € Ü�=‰=×ÒØØ”5—:‘:×3Ñ3°AÓ6×BÑBÑBÀwÐRTÂ}ØØr{   c                óì   — | j                   |vr!t        j                  d| j                   |«       t        | j                  j
                  «      xr% | j                   |v xr t        | j                  «      S )NzDNot using template since dtype %s is not in allowed layout dtypes %s)r©   rÓ   rÔ   Úis_gpurª   r  r»  )r,  Úallowed_layout_dtypess     rW   Ú_use_template_for_gpurÚ  ó  sf   € ð ‡|�|Ð0Ñ0Ü�	‰	ØRØ�L‰LØ!ô	
ô 	ˆv�}‰}×!Ñ!Ó"ò 	&Ø�L‰LÐ1Ð1ò	&ä�v—}‘}Ó%ðr{   c                ó¾   — | j                  «       t        j                  j                  «       j                  d«      D �cg c]  }|j	                  «       ‘Œ c}v S c c}w ©NrA  )rù   ri   Úmax_autotune_gemm_backendsr[  rX  ©ÚbackendrT   s     rW   Ú_use_autotune_backendrà    óM   € Ø�=‰=‹?Ü!×<Ñ<×BÑBÓD×JÑJÈ3ÓOöØˆ�‰�	òð ð ùò ó   ¿Ac                ó¾   — | j                  «       t        j                  j                  «       j                  d«      D �cg c]  }|j	                  «       ‘Œ c}v S c c}w rÜ  )rù   ri   Úmax_autotune_conv_backendsr[  rX  rÞ  s     rW   Ú_use_conv_autotune_backendrå  	  rá  râ  )Úenable_int32Úenable_float8Úcheck_max_autotunec               ó¬  — ddl m}m} t        j                  t        j
                  t        j                  g}|r>t        j                  t        j
                  t        j                  t        j                  g}|r/|j                  t        j                  t        j                  g«       t        | j                  j                  «      xr t        | |«      xs) | j                  j                  dk(  xr | j                  |v xrS t         j"                  xs t         j$                  xs | xr* t'        d«      xr  || j                  |j(                  «      S )Nr7   )ÚBackendFeatureÚhas_backend_featurer  ÚTRITON)rË  rê  rë  rP   rÃ   rE  rG  rT  rO  r;  r<  rØ  rª   r  rÚ  r©   ri   Úmax_autotuneÚmax_autotune_gemmrà  ÚTRITON_TEMPLATES)r,  ræ  rç  rè  rê  rë  Úlayout_dtypess          rW   Úuse_triton_templaterñ    s  € ÷ Dä—]‘]¤E§N¡N´E·M±MÐB€MÙÜŸ™¬¯©¼¿¹ÄuÇ{Á{ÐSˆÙØ×Ñœe×1Ñ1´5×3DÑ3DÐEÔFô �v—}‘}×)Ñ)Ó*ò AÜ)¨&°-Ó@òOð —‘×"Ñ" eÑ+ÒM°·±ÀÐ0Mò
	Pô × Ñ ÒV¤F×$<Ñ$<ÒVÐDVÐ@Vò
	Pô " (Ó+ò
	Pñ   §¡¨~×/NÑ/NÓOðr{   ©Úoutput_layoutÚ
add_guardsc                óÚ   ‡‡‡‡‡‡	— ddl m} ddlmŠ dˆfd„Šdˆˆfd„}dˆˆˆ	fd„Š	 	 	 	 	 	 	 	 dˆˆˆfd„Š	 	 	 	 	 	 	 	 dˆfd	„Š	 |«       xr t	        ˆfd
„|D «       «      xr  || «      S )u^  
    Return True iff *all* supplied tensors satisfy the CUDA TMA constraints
    that Triton relies on today.
    * https://docs.nvidia.com/cuda/cuda-driver-api/group__CUDA__TENSOR__MEMORY.html

    A tensor is accepted when:
      * 1 â‰¤ rank â‰¤ 5 (cuTensorMapEncodeTiled)
      * dtype in _TMA_SUPPORTED_DTYPES (CUtensorMapDataType enum)
      * Base pointer 16-byte aligned
      * Exactly one contiguous ("inner") dim with stride 1
      * All "outer" dims have 16-byte aligned strides
      * Inner dim size Ã— itemsize is a multiple of 16
      * For 1-byte dtypes (e.g. FP8), inner dim â‰¥ 32
    r   )Úhas_triton_tma_devicer7   rm  c                óX   •— ‰j                   j                  j                  | t        «      S r‹   )rq  rr  Ústatically_known_multiple_ofÚTMA_ALIGNMENT)Ú
expr_bytesrn  s    €rW   Ú_alignedzcan_use_tma.<locals>._alignedA  s    ø€ Ø�w‰w×Ñ×<Ñ<¸ZÌÓWÐWr{   c                óŒ   •— | €y| j                   }| j                  }| j                  } ‰| j                  «      sy ‰|||«      S )NTF)rÈ  r4  r©   r`  )r,  ÚsizesÚstridesr©   rû  Ú_is_tma_compatibles       €€rW   Ú_is_tma_compatible_layoutz.can_use_tma.<locals>._is_tma_compatible_layoutD  sF   ø€ Øˆ>ØØ—‘ˆØ—-‘-ˆØ—‘ˆñ ˜Ÿ™Ô&Øá! %¨°%Ó8Ð8r{   c                ó  •— | j                  «       }| j                  «       }| j                  «       }| j                  «       ‰j                  j
                  v ry| j                  «       x}�|j                  dk(  r
 ‰|||«      S  ‰|||«      S )NFrJ   )Úget_sizeÚ
get_strideÚ	get_dtyperÍ  rq  Úunaligned_buffersÚ
get_devicer  )rÝ  rý  rþ  r©   Úm_devicern  rÿ  Ú_is_tma_compatible_xpus        €€€rW   Ú_is_tma_compatible_matrixz.can_use_tma.<locals>._is_tma_compatible_matrixQ  s{   ø€ Ø—
‘
“ˆØ—,‘,“.ˆØ—‘“ˆð �:‰:‹<˜1Ÿ7™7×4Ñ4Ñ4ØàŸ™›Ð&ˆHÐ3¸¿¹ÈÒ8NÙ)¨%°¸%Ó@Ð@á! %¨°%Ó8Ð8r{   c                óR  •— t        | «      }|j                  }|dk  s|dkD  ry|t        vry‰rK‰j                  j                  j                  | «      }‰j                  j                  j                  |«      }nd| D �cg c]'  }‰j                  j                  j                  |«      ‘Œ) }}|D �cg c]'  }‰j                  j                  j                  |«      ‘Œ) }}t        |«      D �	�cg c]-  \  }	}‰j                  j                  j                  |d«      r|	‘Œ/ }
}	}t        |
«      dk7  ry|
d   }t        |«      D ]  \  }	}|	|k(  rŒ ‰||z  «      rŒ y ||   } ‰||z  «      sy|dk(  r'‰j                  j                  j                  |d«      syyc c}w c c}w c c}}	w )Nr7   r­   Fr   é    T)
rR   Úitemsizert   rq  rr  Úguard_int_seqÚ!replace_backed_symbols_with_hintsrE  Ústatically_known_equalsÚstatically_known_geq)rý  rþ  r©   Úrankr  Úsizes_iÚ	strides_iré   Ústrè   rö  Ú	inner_idxÚ	inner_dimrn  rû  rô  s                €€€rW   rÿ  z'can_use_tma.<locals>._is_tma_compatible_  s³  ø€ ô
 �5‹zˆØ—>‘>ˆà�!Š8�t˜a’xØàÔ-Ñ-ØáØ—g‘g×&Ñ&×4Ñ4°UÓ;ˆGØŸ™×(Ñ(×6Ñ6°wÓ?‰Ið PUöØJK�—‘× Ñ ×BÑBÀ1ÕEðˆGð ð RYöØKM�—‘× Ñ ×BÑBÀ2ÕFðˆIð ô # 9Ó-÷
á��2Ø�w‰w×Ñ×7Ñ7¸¸AÔ>ò ð
ˆñ 
ô
 ˆu‹:˜Š?ØØ˜!‘Hˆ	ô ˜yÓ)ò 	‰EˆAˆrØ�IŠ~ØÙ˜B ™MÕ*Ùð		ð ˜IÑ&ˆ	Ù˜	 HÑ,Ô-Øð �qŠ= §¡×!1Ñ!1×!FÑ!FÀyÐRTÔ!UØàùòCùòùó

s   Á>,FÂ0,FÃ,2F#c                óT  •— |d   }‰j                   j                  j                  |«      }‰j                   j                  j                  |d«      syd}| D ]O  }‰j                   j                  j                  |«      }‰j                   j                  j	                  ||«      sŒO y y)Nr³   r7   Fl   ÿÿ T)rq  rr  r  r  Ústatically_known_gt)	rý  rþ  r©   Úlast_strideÚlast_stride_hintÚ
MAX_UINT32rÈ  Ú	size_hintrn  s	           €rW   r  z+can_use_tma.<locals>._is_tma_compatible_xpu”  sž   ø€ ð ˜b‘kˆØŸ7™7×+Ñ+×MÑMØó
Ðð �w‰w×Ñ×7Ñ7Ð8HÈ!ÔLØð ˆ
Øò 	ˆDØŸ™×(Ñ(×JÑJÈ4ÓPˆIØ�w‰w×Ñ×3Ñ3°I¸zÕJÙð	ð
 r{   c              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wr‹   rš   )r»   rÝ  r	  s     €rW   r½   zcan_use_tma.<locals>.<genexpr>¬  s   øè ø€ Ò?°Ñ)¨!×,Ñ?ùó   ƒ)rú  úint | sympy.Exprr’   r1  )r,  úLayout | Noner’   r1  )rÝ  r@   r’   r1  )rý  úSequence[sympy.Expr]rþ  zSequence[_IntLike]r©   útorch.dtyper’   r1  )Útorch.utils._tritonrö  rp  rn  r�   )
ró  rô  Úmatricesrö  r   rn  rû  rÿ  r	  r  s
    `   @@@@@rW   Úcan_use_tmar%  ,  s›   ý€ õ" :åõXö9÷9ð3Ø#ð3à#ð3ð ð3ð 
÷	3ðjØ#ðà#ðð ðð 
õ	ñ. 	Óò 	5ÜÓ?°hÔ?Ó?ò	5á% mÓ4ðr{   c                óü  — t        j                  t         j                  «      j                  }g }| D ]O  }t	        |t
        t        j                  f«      r||kD  sŒ) y|j                  t        j                  ||«      «       ŒQ |syddl
m} t        |«      dk(  r|d   nt        j                  |Ž }|r%|j                  j                  j!                  |«      S |j                  j                  j#                  |«      S )NFTr7   rm  r   )rP   ÚiinforT  rÉ   r}   rŒ   r~   r�   r™  ÚLerp  rn  rR   ÚAndrq  rr  Úguard_or_falseÚstatically_known_true)rý  rô  Ú	int32_maxÚ
conditionsrÈ  rn  Ú	conditions          rW   Ú_descriptor_shape_fits_in_int32r/  ±  sÔ   € ô —‘œEŸK™KÓ(×,Ñ,€IØ€JØò 9ˆÜ�dœS¤%§-¡-Ð0Ô1Ø�iÓÙà×ÑœeŸh™h t¨YÓ7Õ8ð9ñ Øåä!$ Z£°AÒ!5�
˜1’¼5¿9¹9ÀjÐ;Q€Iñ ð 	
�‰×Ñ×'Ñ'¨	Ó2ðð �W‰W×Ñ×3Ñ3°IÓ>ðr{   ©rô  c                óx  ‡— t         j                  j                  syt        d„ |D «       «      syt        ˆfd„|D «       «      syt         j                  j                  rt        | j                  ‰¬«      syt        j                  j                  �yt         j                  j                  r| nd }t        ||‰dœŽS )NFc              3  óT   K  — | ]   }t        |j                  «       «      d k(  –— Œ" y­w)ry  N)rR   r  )r»   rÝ  s     rW   r½   z*use_triton_tma_template.<locals>.<genexpr>Ï  s    è ø€ Ò8¨!Œs�1—:‘:“<Ó  AÕ%Ñ8ùs   ‚&(c              3  óT   •K  — | ]  }t        |j                  «       ‰¬ «      –— Œ! y­w)r0  N)r/  r  )r»   rÝ  rô  s     €rW   r½   z*use_triton_tma_template.<locals>.<genexpr>Ñ  s)   øè ø€ ò àô 	(¨¯
©
«À×LÐLñùs   ƒ%(r0  Trò  )ri   rò  Úenable_persistent_tma_matmulr�   Úenable_template_tma_storer/  rÈ  rP   rN  rq   r%  )ró  rô  r$  r,  s    `  rW   Úuse_triton_tma_templater6  Ê  sœ   ø€ ô �=‰=×5Ò5ØÜÑ8¨xÔ8Ô8ØÜó àôô ð Ü‡}�}×.Ò.Ô7VØ×Ñ zõ8ð ô ‡}�}×ÑÐ$ØÜ$Ÿm™m×EÒE‰]È4€FÜ˜°À:ÒNÐNr{   c                óT   — t        || |dœŽsyddlm} ddlm}  |«       xr  |«       S )Nrò  Fr   )Ú%has_triton_tensor_descriptor_host_tmar7   ©Úis_datacenter_blackwell_arch)r6  r#  r8  Úcodegen.cuda.cuda_envr:  )ró  rô  r$  r8  r:  s        rW   Ú!use_triton_blackwell_tma_templater<  â  s2   € ô #Ø	 ¸:óð åIåCñ 1Ó2ÒUÑ7SÓ7UÐUr{   c                ó   — | |v xr ||v S r‹   rš   )Úscale_option_aÚscale_option_bÚscaling_typess      rW   Úuse_triton_scaling_templaterA  ò  s   € ð
 ˜]Ð*ÒN¨~ÀÐ/NÐNr{   )Úmaxsizec                 ód   — 	 t         j                  j                  d«      duS # t        $ r Y yw xY w)z¸Check if CuTeDSL is importable; cache the result for reuse.

    Call ensure_cute_available.cache_clear() after installing CuTeDSL
    in the same interpreter to retry the import.
    ÚcutlassNF©Ú	importlibÚutilÚ	find_specr	  rš   r{   rW   Úensure_cute_availablerI  ú  s3   € ðÜ�~‰~×'Ñ'¨	Ó2¸$Ð>Ð>øÜò Ùðúó   ‚ # £	/®/c                 ó€   — 	 t         j                  j                  d«      du} | r
t	        «        | S # t        $ r Y yw xY w)zåCheck if NVIDIA Universal GEMM (cutlass_api) is importable; cache the result for reuse.

    Call ensure_nv_universal_gemm_available.cache_clear() after installing cutlass_api
    in the same interpreter to retry the import.
    Úcutlass_apiNF)rF  rG  rH  r	  Ú_ensure_fp4_dtype_registered)Ú	availables    rW   Ú"ensure_nv_universal_gemm_availablerO    sE   € ðÜ—N‘N×,Ñ,¨]Ó;À4ÐGˆ	ñ Ü$Ô&ØÐøô	 ò Ùðús   ‚!1 ±	=¼=c                 óð   ‡‡— ddl } 	 | j                  j                  t        j                  «       y# t
        t        f$ r4 ddlŠ| j                  j                  Šˆˆfd„}|| j                  _        Y yw xY w)a$  Patch cutlass_api to handle torch.float4_e2m1fn_x2 -> cutlass.Float4E2M1FN.

    NOTE: cutlass_api doesn't natively map this dtype. We patch the lookup function
    in-place so all callers (including TensorWrapper) pick up the change.
    Remove once cutlass_api adds native FP4 support.
    r   Nc                óR   •— | t         j                  k(  r‰j                  S  ‰| «      S r‹   )rP   rC  ÚFloat4E2M1FN)r©   Ú_origrD  s    €€rW   Ú_patchedz._ensure_fp4_dtype_registered.<locals>._patched'  s'   ø€ Øœ×.Ñ.Ò.Ø×+Ñ+Ð+Ù˜“<Ðr{   )Úcutlass_api.utilsrØ  Úcutlass_type_from_torch_typerP   rC  ÚKeyErrorrä  rD  )rL  rT  rS  rD  s     @@rW   rM  rM    sh   ù€ ó ðBØ×Ñ×6Ñ6´u×7MÑ7MÕNøÜ”nÐ%ò 
BÛà×!Ñ!×>Ñ>ˆõ	 ð
 :Bˆ×ÑÖ6ð
Bús   ˆ)2 ²A A5Á4A5c                 ód   — 	 t         j                  j                  d«      duS # t        $ r Y yw xY w)aG  Check if nvMatmulHeuristics is importable; cache the result for reuse.

    nvMatmulHeuristics provides performance model-based kernel selection
    for NVIDIA GEMM operations.

    Call ensure_nvmatmul_heuristics_available.cache_clear() after installing
    nvMatmulHeuristics in the same interpreter to retry the import.
    ÚnvMatmulHeuristicsNFrE  rš   r{   rW   Ú$ensure_nvmatmul_heuristics_availablerZ  /  s4   € ðÜ�~‰~×'Ñ'Ð(<Ó=ÀTÐIÐIøÜò ÙðúrJ  c                ón  — t        «       syt        d«      syddlm} t	        |j
                  j                  «      sy |«       syt        j                  g}	t        ||	«      syt        j                  st        j                  syt        | ||¬«      syt        d„ | |fD «       «      ry|r|ry|€y|€|�yy)añ  
    Returns True if we can use the blackwell kernel for grouped mm.
    Required conditions:
        1. CuTeDSL backend is enabled
        2. CuTeDSL is available
        3. We are on a blackwell arch
        4. The dtype is bf16
        5. Max autotune or max autotune gemm is enabled
        6. A, B, and the output are 16B aligned
        7. We are not using dynamic shapes
        8. A is 2d
        9. B is 3d
        10. Offsets are provided
        11. Bias and Scale are not provided
    FÚCUTEDSLr7   r9  )ró  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr‹   )Ú
is_dynamic©r»   rT   s     rW   r½   z3use_blackwell_cutedsl_grouped_mm.<locals>.<genexpr>q  s   è ø€ Ò
1˜QŒ:�a�=Ñ
1ùr«  T)rI  rà  r;  r:  rØ  rª   r  rP   rE  rÚ  ri   rí  rî  r%  r|  )
Úmat_aÚmat_br,  Úa_is_2dÚb_is_2dÚoffsÚbiasÚscale_resultr:  rð  s
             rW   Ú use_blackwell_cutedsl_grouped_mmrg  ?  s­   € ô2 !Ô"Øä  Ô+ØåCä�&—-‘-×$Ñ$Ô%Øá'Ô)Øä—^‘^Ð$€MÜ  ¨Ô7Øä×Ò¤6×#;Ò#;Øô �u˜e°6Õ:Øä
Ñ
1 5¨% .Ô
1Ô1Øá‘gØà€|ØàÐ˜<Ð3Øàr{   c                óÀ  — ddl m} |j                  j                  r,|j                  j                  st        j                  d«       y|j                  j                  j                  ||z  |z  d¬«      }|dk  s|t        j                  j                  k  ryddlm} t        j                  j                   ryt        j"                  t        j$                  t        j&                  g}t)        | |«      xr/ t        j*                  xs t        j,                  xr t/        d	«      }|r6 |«       s/t0        j3                  d
t        j                  j4                  «       y|S )Nr7   rm  zYCUTLASS backend is not supported with non-AOT cpp_wrapper mode. Skipping CUTLASS backend.Fr³   ©Úfallbackr   )Útry_import_cutlassÚCUTLASSzŽFailed to import CUTLASS lib. Please check whether _inductor.config.cutlass.cutlass_dir %s is set correctly. Skipping CUTLASS backend for now.)rp  rn  rq  Úcpp_wrapperÚaot_modeÚwarningsÚwarnrr  Úoptimization_hintri   rD  Úcutlass_backend_min_gemm_sizeÚcodegen.cutlass.utilsrk  rP   rN  rq   rÃ   rE  rT  rÚ  rí  rî  rà  rÓ   rö  Úcutlass_dir)	r,  rÝ  r  Úkrn  Ú	gemm_sizerk  rð  rì   s	            rW   Úuse_cutlass_templaterw  €  s  € Ýð 	‡w�w×Ò 1§7¡7×#3Ò#3Ü�‰ð(ô	
ð à—‘× Ñ ×2Ñ2°1°q±5¸1±9ÀrÐ2ÓJ€IØ�A‚~˜¤V§^¡^×%QÑ%QÒQØÝ9ô ‡}�}×ÒØô —]‘]¤E§N¡N´E·K±KÐ@€Mä˜f mÓ4ò 	-Ü× Ñ Ò<¤F×$<Ñ$<ò	-ä! )Ó,ð ñ Ù!Ô#Ü�K‰Kð4ô —‘×*Ñ*ô	ð Ø€Jr{   Ú_IntLikec                óÚ  ‡
‡— ddl mŠ t        «       syt        «       syt	        d«      syddlmŠ
 ‰
j                  ry| j                  j                  dk7  st        j                  j                  ryt        j                  st        j                  sy|||g}|�|j!                  |«       t#        ˆfd„|D «       «      ry||g}	|�|	j!                  |«       t#        ˆ
fd	„|	D «       «      ryy
)a3  
    Return True if we can use the NVIDIA Universal GEMM Template.

    Required conditions:
        1. NVGEMM backend is enabled
        2. cutlass_api is available
        3. We are on a NVIDIA GPU
        4. Max autotune or max autotune gemm is enabled
        5. Not in AOT Inductor mode (requires runtime JIT compilation)
        6. Base pointers are 16-byte aligned
        7. Shape dimensions are not unbacked symbols

    Note:
        - Shape and stride constraints are handled internally by
          cutlass_api.get_kernels() which filters incompatible kernels.
        - GroupedGemm currently only supports TN layout (column-major B).
          Any other layout will act as a noop and fall back to ATen.
        - Dynamic shapes are supported as long as they have hints
          (from example inputs).
    r   )Úhas_free_unbacked_symbolsFÚNVGEMMr7   rm  rH   c              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wr‹   rš   )r»   Údimrz  s     €rW   r½   z1use_nv_universal_gemm_template.<locals>.<genexpr>ê  s   øè ø€ Ò
C¨cÑ$ S×)Ñ
Cùr  c              3  ój   •K  — | ]*  }|j                  «       ‰j                  j                  v –— Œ, y ­wr‹   )rÍ  rq  r  )r»   Útrn  s     €rW   r½   z1use_nv_universal_gemm_template.<locals>.<genexpr>ò  s&   øè ø€ Ò
O¸ˆ1�:‰:‹<˜1Ÿ7™7×4Ñ4Ô4Ñ
Oùs   ƒ03T)r±  rz  rI  rO  rà  rp  rn  Úaot_compilationrª   r  rP   rN  rq   ri   rí  rî  r™  r|  )r,  rÝ  r  ru  r`  ra  rd  r›  Údims_to_checkÚtensors_to_checkrn  rz  s             @@rW   Úuse_nv_universal_gemm_templaterƒ  °  sÐ   ù€ õ< Pä Ô"Øä-Ô/Øä  Ô*Øåà×ÒØà‡}�}×Ñ˜VÒ#¤u§}¡}×'8Ò'8Øä×Ò¤6×#;Ò#;Øð
 ˜˜1�I€MØ€}Ø×Ñ˜QÔÜ
Ó
C°]Ô
CÔCØð ˜u�~ÐØÐØ×Ñ Ô%Ü
Ó
OÐ>NÔ
OÔOØàr{   c                óâ   — t         j                  j                  j                  «       }|dk(  ry| j                  «       |j	                  d«      D �cg c]  }|j                  «       ‘Œ c}v S c c}w )z8Check if CUTLASS should be used for the given operation.ÚALLTrA  )ri   rD  Úcutlass_enabled_opsrù   r[  rX  )Úop_nameÚenabled_opsrT   s      rW   Ú_use_cutlass_for_opr‰  ø  sU   € ä—.‘.×4Ñ4×:Ñ:Ó<€KØ�eÒØØ�=‰=‹?°+×2CÑ2CÀCÓ2HÖI¨Q˜qŸw™w�yÒIÐIÐIùÒIs   ÁA,r   c           
     ó¸  — ddl m} t        j                  j                  |z  }|j
                  j                  j                  t        j                  t        j                  ||| z  «      t        j                  |||z  «      «      «      xrO |j
                  j                   xr6 |j
                  j                   xr t        j                  j                  dkD  S )Nr   rm  )Útorch._inductor.virtualizedrn  ri   rò  Údecompose_k_thresholdrq  rr  r+  r~   r)  ÚGern  rm  Únum_decompose_k_splits)rÝ  r  ru  Úthreshold_multiplern  rŒ  s         rW   Úuse_decompose_k_choicer�  	  s·   € õ .ä"ŸM™M×?Ñ?ÐBTÑTÐð 	
�‰×Ñ×.Ñ.Ü�I‰IÜ—‘˜Ð1°AÑ5Ó6Ü—‘˜Ð1°AÑ5Ó6óó	
ò 	5ð —‘× Ñ Ð ò	5ð —‘×#Ñ#Ð#ò	5ô �M‰M×0Ñ0°1Ñ4ð
r{   c           
     ó¾  — t         j                  j                  }ddlm} t        t        j                  j                  «      xr™ |j                  j                  j                  t        j                  t        j                  ||| z  «      t        j                  |||z  «      «      «      xr0 |j                  j                   xr |j                  j                    S )zŒ
    Check if we should use the contiguous subgraph transform.
    This transform makes the second matrix contiguous before the matmul.
    r   rm  )ri   ÚrocmÚcontiguous_thresholdr‹  rn  r1  rP   rN  rq   rq  rr  r+  r~   r)  r�  rn  rm  )rÝ  r  ru  r“  rn  s        rW   Úuse_contiguousr”  	  s±   € ô "Ÿ;™;×;Ñ;Ðõ .ô 	ŒU�]‰]×ÑÓò 	$Ø�G‰G×Ñ×2Ñ2Ü�I‰IÜ—‘˜Ð0°1Ñ4Ó5Ü—‘˜Ð0°1Ñ4Ó5óó
ò	$ð —‘× Ñ Ð ò	$ð —‘×#Ñ#Ð#ð
r{   c                óâ  — t         j                  j                  }g d¢}t        |t        j
                  «      r|j                  s|S |dk(  rg S t        | t        j
                  «      r| j                  r&t        |t        j
                  «      r|j                  sd}nt        || z  ||z  «      }d}t	        j                  |«      }|D �cg c]  }||k  r||k\  r|‘Œ }}g g g }}
}	|D ]Z  }||z  }|dk  rŒ||dz
  z  dk(  r|dk\  r|	j                  |«       Œ0|dz  dk(  r|
j                  |«       ŒJ|j                  |«       Œ\ t         j                  dk(  r|	|
z   |z   S |	|
z   |z   }|d | S c c}w )	N)rr   r  ru   rs   é   r   r–  ry  rs   r7   r  Ú
EXHAUSTIVE)ri   rò  rŽ  r}   r~   r+  Ú	is_numberro  Údivisorsr™  Úmax_autotune_gemm_search_space)rÝ  r  ru  Úk_splits_limitÚdefault_k_splitsÚmax_k_splitÚmin_k_splitr™  ÚdivisorÚpow_of_2_divisorsÚmul_of_32_divisorsÚrest_of_splitsÚdÚkPartÚbest_splitss                  rW   Úget_k_splitsr¦  0	  sŒ  € ô —]‘]×9Ñ9€Nò .Ðä�!”U—Z‘ZÔ ¨¯ªØÐØ	˜1Ò	Øˆ	ä�1”e—j‘jÔ!¨!¯+ª+Ü�1”e—j‘jÔ!¨!¯+ª+à‰ä˜!˜q™& ! q¡&Ó)ˆà€Kä�~‰~˜aÓ €Hð  öàØ�kÒ! g°Ò&<ò 	ð€Hð ð =?ÀÀB¨>Ð)Ðàò %ˆØ�Q‘ˆð �3Š;Øð �E˜A‘IÑ !Ò#¨°ªØ×$Ñ$ QÕ'à�R‰Z˜1Š_Ø×%Ñ% aÕ(ð ×!Ñ! !Õ$ð%ô" ×,Ñ,°Ò<Ø Ð#5Ñ5¸ÑFÐFà#Ð&8Ñ8¸>ÑI€Kà�˜Ð'Ð'ùò=s   Ã
E,c                óT   — t         j                  j                  | «      j                  S r‹   )rP   rH   r½  rÔ  ©rª   s    rW   Ú_rocm_native_device_arch_namer©  i	  s   € ä�:‰:×+Ñ+¨FÓ3×?Ñ?Ð?r{   c                 óÌ   — 	 dd l } ddlm}m} ddlm} t        j                  j                  | j                  «      }||||fS # t        $ r dd„}dd„} G d„ d«      }d }Y Œ%w xY w)	Nr   )Úgen_ops_libraryÚgen_ops_preselected)ÚCKGemmOperationc                 ó   — g S r‹   rš   rš   r{   rW   r«  z*try_import_ck_lib.<locals>.gen_ops_library	  ó   € ØˆIr{   c                 ó   — g S r‹   rš   rš   r{   rW   r¬  z.try_import_ck_lib.<locals>.gen_ops_preselected‚	  r¯  r{   c                  ó   — e Zd Zy)ú*try_import_ck_lib.<locals>.CKGemmOperationN)r“   r”   r•   rš   r{   rW   r­  r²  …	  s   „ Ør{   r­  )r’   r•  )Úck4inductorÚ(ck4inductor.universal_gemm.gen_instancesr«  r¬  Úck4inductor.universal_gemm.opr­  rê  r÷  ÚdirnameÚ__file__r	  )r³  r«  r¬  r­  Úpackage_dirnames        rW   Útry_import_ck_libr¹  n	  sl   € ðÛ÷	
õ	
ô Ÿ'™'Ÿ/™/¨+×*>Ñ*>Ó?ˆð ˜OÐ-@À/ÐQÐQøô ò ó	ó	÷	ñ 	ð Šðús   ‚;A ÁA#Á"A#c                ó   — t         j                  st         j                  syt        j                  j
                  sy| j                  j                  dk7  ryt        | j                  «      }t         j                  j                  D �ci c]  }|j                  d«      d   |“Œ c}xs |j                  d«      d   |i}|j                  «       t         j                  j                  z  D �cg c]  }||   ‘Œ	 }}|sy| j                  t        j                  t        j                   t        j"                  fvryt%        «       \  }}}}|st&        j)                  d«       y|t         j                  _        yc c}w c c}w )NFrH   ú:r   z,Please pip install Composable Kernel packageT)ri   rí  rî  rP   rN  rq   rª   r  r©  r’  Úarchr[  rI  Úck_supported_archr©   rÃ   rE  rG  r¹  rÓ   rö  Úck_dir)r,  Únative_archru  Úrequested_archsÚrequested_supported_archsÚck_package_dirnamerã   s          rW   Úuse_ck_templaterÃ  Œ	  s<  € ä×Ò¤6×#;Ò#;Øä�=‰=×ÒØà‡}�}×Ñ˜VÒ#Øô 0°·±Ó>€KÜ39·;±;×3CÑ3CÖD¨a�q—w‘w˜s“| A‘¨Ñ)ÒDò Ø×Ñ˜#Ó˜qÑ! ;ðI€Oð
 !×%Ñ%Ó'¬&¯+©+×*GÑ*GÑGö!àð 	˜Óð!Ðð !ñ %Øà‡|�|œEŸM™M¬5¯>©>¼5¿=¹=ÐIÑIØä"3Ó"5ÑÐ˜˜1˜aáÜ�‰ÐBÔCØà+„F‡K�KÔàùò+ Eùò!s   ÂE6Ã,E;c                ó¢   — ddl m} t        d«      xr= t        | «      xr0 |j                  j
                  j                  ||z  |z  d¬«      dkD  S )Nr7   rm  ÚCKr³   ri  r   ©rp  rn  rà  rÃ  rq  rr  rq  ©r,  rÝ  r  ru  rn  s        rW   Úuse_ck_gemm_templaterÈ  ±	  sR   € Ýô 	˜dÓ#ò 	KÜ˜FÓ#ò	Kà�G‰G×Ñ×.Ñ.¨q°1©u°q©yÀ2Ð.ÓFÈÑJðr{   c                ó¢   — ddl m} t        d«      xr= t        | «      xr0 |j                  j
                  j                  ||z  |z  d¬«      dkD  S )Nr7   rm  ÚCKTILEr³   ri  r   rÆ  rÇ  s        rW   Úuse_ck_tile_gemm_templaterË  »	  sR   € Ýô 	˜hÓ'ò 	KÜ˜FÓ#ò	Kà�G‰G×Ñ×.Ñ.¨q°1©u°q©yÀ2Ð.ÓFÈÑJðr{   c                ó2   — t        d«      xr t        | «      S )NrÅ  )rå  rÃ  ©r,  s    rW   Úuse_ck_conv_templaterÎ  Å	  s   € Ü% dÓ+ÒG´ÀÓ0GÐGr{   c                ó|   — t         j                  xs t         j                  xr | j                  j                  dk(  S r¨  )ri   rí  rî  rª   r  rÍ  s    rW   Ú_use_template_for_cpurÐ  É	  s2   € ä×ÑÒ7œv×7Ñ7ò&à
�-‰-×
Ñ
 Ñ
%ð&r{   c                ó¾  — ddl m} t        |j                  |«      sJ ‚|j                  j                  }|j                  j
                  }t        | «      xrX |j                  «       t        j                  k(  xr5 t        |«      dk(  xr% t        |«      dk(  xr |d   |d   k(  xr |d   dk(  }t        | ||d¬«      xr |j                  j                  «       xs |S )Nr7   )rA   é   ry  F)Úrequire_constant_mat2)rý  rA   r}   r,  rÈ  r4  rÐ  r  rP   rG  rR   Úuse_cpp_gemm_templateÚis_contiguous)r,  Úmat1Úmat2rA   Ú	mat1_sizeÚmat1_strideÚmat1_each_batch_is_contiguouss          rW   Úuse_cpp_bmm_templaterÛ  Ï	  sß   € õ ä�d—k‘k 6Ô*Ð*Ð*ð
 —‘× Ñ €IØ—+‘+×$Ñ$€Kä˜fÓ%ò 	"Ø�N‰NÓ¤§¡Ñ-ò	"ä�‹^˜qÑ ò	"ô �Ó Ñ"ò	"ð ˜‰^˜y¨™|Ñ+ò		"ð
 ˜‰^˜qÑ ð "ô ! ¨¨tÈ5ÔQò Ø�‰×!Ñ!Ó#ÒDÐ'Dðr{   c                ó~  — ddl m} ddlm} ddlm}	 ddlm}
 t        | «      rt        d«      syt        j                  j                  sy|j                  «       t        j                  t        j                   fv }t        j"                  t        j$                  t        j&                  t        j                  t        j                   g} |
|||r| j(                  nd ||¬«      \  }}}} }}t+        ||f«      ryt-        ||j.                  «      r|j1                  «       } |	|j                  «       «      \  }} |d	||||j                  «       |j                  «       |t3        «       | |¬
«
      }dd„}| j(                  |v xr= |d uxr7  ||«      xr- t-        ||j4                  «      xr |j7                  «       xs | S )Nr7   rü  )Úcreate_micro_gemm)Ú*get_gemm_template_output_and_compute_dtype)Úmm_argsÚCPPF)Ú	out_dtypeÚmat2_transposedÚuse_4x2_dimÚ
micro_gemm)Úinput_dtypeÚinput2_dtypeÚoutput_dtypeÚnum_threadsÚuse_refÚq_group_sizec                óN   — | j                  «        | j                  «       d   dk(  S )Nr³   r7   )Úfreeze_layoutr  ©rT   s    rW   Úis_last_dim_stride1z2use_cpp_gemm_template.<locals>.is_last_dim_stride1
  s"   € Ø	�‰ÔØ�|‰|‹~˜bÑ! QÑ&Ð&r{   )rT   r@   r’   r1  )rþ  rý  Úcodegen.cpp_micro_gemmrÝ  Úcodegen.cpp_utilsrÞ  Úkernel.mm_commonrß  rÐ  rà  ri   ÚcppÚweight_prepackr  rP   rV  rK  rG  rE  Úhalfr©   Úhas_free_symbolsr}   ÚBaseViewÚunwrap_viewÚparallel_num_threadsr'  Úis_module_buffer)r,  rÖ  r×  râ  rÓ  Úis_woq_int4rê  rý  rÝ  rÞ  rß  Ú	int8_gemmrð  rÝ  r  ru  rç  rã   rä  rî  s                       rW   rÔ  rÔ  è	  s‡  € õ Ý9ÝMÝ)ä  Ô(Ô0EÀeÔ0LØä�:‰:×$Ò$Øà—‘Ó ¤U§[¡[´%·*±*Ð$=Ð=€IÜ—]‘]¤E§N¡N´E·J±JÄÇÁÌUÏZÉZÐX€MÙ")ØØÙ"+�&—,’,°Ø'Øô#Ñ€A€qˆ!ˆV�T˜4ô ˜˜A˜ÔØä�$˜Ÿ™Ô$Ø×ÑÓ!ˆá@ÀÇÁÓAQÓR�O€L�!Ù"ØØ	Ø	Ø	Ø—N‘NÓ$Ø—^‘^Ó%Ø!Ü(Ó*Ø�Ø!ô€Jó'ð
 	�‰˜Ð%ò 	CØ˜dÐ"ò	Cá Ó%ò	Cô �t˜RŸ]™]Ó+ò	Cð ×"Ñ"Ó$ÒAÐ,AÐ(Aðr{   c                 ób   — t         j                  xs t         j                   xs t        d«      S )NÚATEN)ri   rí  rî  rà  rš   r{   rW   Úuse_aten_gemm_kernelsrþ  (
  s-   € ä×ÑÒ7œv×7Ñ7ðò 'ä	˜vÓ	&ð'r{   c                  óT   — e Zd ZU  ej                  d«      Zded<   dd„Zdd„Zd	d„Z	y)
ÚDebugDirManagerr   r
  Úprev_debug_namec                ó@   — t        t        j                  «      | _        y r‹   )rµ  r   Úcounterrž   rD  s    rW   r5  zDebugDirManager.__init__2
  s   € Ü”×.Ñ.Ó/ˆ�r{   c                óè   — t         j                  j                  j                  | _        | j                  › d| j
                  › �| _        | j                  t         j                  j                  _        y )NÚ_tmp_)rP   Ú_dynamori   Údebug_dir_rootr  rž   Únew_namerD  s    rW   Ú	__enter__zDebugDirManager.__enter__5
  sM   € Ü$Ÿ}™}×3Ñ3×BÑBˆÔØ×/Ñ/Ð0°°d·g±g°YÐ?ˆŒØ.2¯m©mŒ�‰×ÑÕ+r{   c                ó”   — t        j                  | j                  «       | j                  t        j
                  j                  _        y r‹   )r  r  r  r  rP   r  ri   r  )rÜ  r„   s     rW   Ú__exit__zDebugDirManager.__exit__:
  s*   € Ü�‰�d—m‘mÔ$Ø.2×.BÑ.BŒ�‰×ÑÕ+r{   Nrƒ  )r„   r   r’   rÞ  )
r“   r”   r•   rp  r=  r  r¤   r5  r	  r  rš   r{   rW   r   r   .
  s(   … Øˆi�o‰o˜aÓ €GØÓó0ó<ô
Cr{   r   c                ó  ‡— ddl m} t        «       Šdˆfd„}t        j                  j                  |d|«      5  t        j                  j                  «         | |i |¤Ž}d d d «       t        ‰«      fS # 1 sw Y   ŒxY w)Nr7   r<   c                ó(   •— ‰j                  | «       y r‹   )rK  ©ÚcodeÚsource_codess    €rW   Úsave_output_codez*run_and_get_code.<locals>.save_output_codeH
  s   ø€ Ø×Ñ˜Õr{   r  ©r  r
  r’   rÞ  )
rq  r=   r!   r   Úpatchr�  rP   r  Úresetr]  )rÝ   r„   r¤  r=   r  r°  r  s         @rW   Úrun_and_get_coder  ?
  st   ø€ õ
 %ä$.£L€Lõô 
�‰×	Ñ	˜=Ð*<Ð>NÓ	Oñ %Ü�‰×ÑÔÙ�TÐ$˜VÑ$ˆ÷%ð ”4˜Ó%Ð%Ð%÷%ð %ús   ¹'A5Á5A>c                óÂ  — |j                  dd«      }t        | g|¢­i |¤Ž\  }}g }|D ]¬  }t        j                  rQt        j                  j
                  dur5|j                  t        j                  d|t        j                  «      «       n4|j                  t        j                  d|t        j                  «      «       |sŒ›|D �cg c]  }|dd ‘Œ	 }}Œ® ||fS c c}w )NÚremove_quoteFTzR"TRITON\((.*?)\)TRITON"z	'''.*?'''rÒ  éýÿÿÿ)
rS   r  ri   rm  rò  Úautotune_at_compile_timerO  rÙ   ÚfindallÚDOTALL)	rÝ   r„   r¤  r  r°  r  Úkernelsr  rÒ  s	            rW   Úrun_and_get_kernelsr  Q
  sÆ   € ð —:‘:˜n¨eÓ4€Lä+¨BÐ@°Ò@¸Ñ@Ñ€FˆLØ€GØò ;ˆÜ×Ò¤&§-¡-×"HÑ"HÐPTÑ"Tð �N‰Nœ2Ÿ:™:Ð&AÀ4ÌÏÉÓSÕTà�N‰Nœ2Ÿ:™: l°D¼"¿)¹)ÓDÔEÚØ29Ö:¨�v˜a ’|Ð:ˆGÑ:ð;ð �7ˆ?Ðùò ;s   ÃCc                ó&   ‡ — dˆ fd„}t        |«      S )Nc                 óR   •—  ‰«       } | j                  «       j                  «        | S r‹   )rü   r  )r°  rÝ   s    €rW   Úrun_with_backwardz1run_fw_bw_and_get_code.<locals>.run_with_backwarde
  s!   ø€ Ù“ˆØ�
‰
‹×ÑÔØˆr{   )r’   r   )r  )rÝ   r   s   ` rW   Úrun_fw_bw_and_get_coder!  d
  s   ø€ õô
 Ð-Ó.Ð.r{   c                óX  ‡‡— ddl m} g Šdˆfd„Šd	ˆfd„}t        j                  j	                  |d|«      5  t        j                  j	                  |d‰«      5  t
        j                  j                  «         | |i |¤Ž}ddd«       ddd«       ‰S # 1 sw Y   ŒxY w# 1 sw Y   ‰S xY w)
zLGet the inductor-generated code, but skip any actual compilation or running.r7   r<   c                ó(   •— ‰j                  | «       y r‹   ©r™  r  s    €rW   r  z"get_code.<locals>.save_output_codes
  s   ø€ Ø×Ñ˜DÕ!r{   c                óÐ   •—  G d„ d«      }| j                   r| j                  «       n| j                  «       \  }} ‰|j                  «       |r ‰|j                  «        |«       S )Nc                  ó    — e Zd ZdZdd„Zdd„Zy)ú@get_code.<locals>.patched_compile_to_module.<locals>.DummyModulez4This is empty to replace the generated triton modulec                 ó   — y r‹   rš   rD  s    rW   r5  zIget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.__init__z
  s   € Ør{   c                 ó   — y r‹   rš   rÛ  s      rW   ÚcallzEget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.call}
  s   € àr{   Nrƒ  ©r„   r   r¤  r   r’   rÞ  )r“   r”   r•   r–   r5  r*  rš   r{   rW   ÚDummyModuler'  w
  s   „ ÙFóôr{   r,  )rm  Úcodegen_with_cpp_wrapperÚcodegenr�   )rÜ  r,  Úwrapper_codeÚkernel_coder  s       €rW   Úpatched_compile_to_modulez+get_code.<locals>.patched_compile_to_modulev
  s]   ø€ ÷	ñ 	ð 04×/?Ò/?ˆD×)Ñ)Ô+ÀTÇ\Á\Ã^ñ 	"ˆ�kñ 	˜×+Ñ+Ô,ÙÙ˜[×.Ñ.Ô/á‹}Ðr{   Úcompile_to_moduler  Nr  )rÜ  r=   r’   r   )rq  r=   r   r  r�  rP   r  r  )rÝ   r„   r¤  r=   r1  rã   r  r  s         @@rW   Úget_coder3  m
  s¢   ù€ å$à €Lõ"õô, 	�
‰
×ÑØÐ.Ð0Ió	
ñ ô 	�
‰
×Ñ˜-Ð);Ð=MÓNñ	 ô 	�‰×ÑÔá�Ð˜Ñˆ÷ ÷  ð Ð÷ ð  ú÷  ð Ðús#   ¸"BÁ'BÂBÂB	ÂBÂB)c                ó|   — t        | g|¢­i |¤Ž}dt        |«      cxk  rdk  sn J dt        |«      › �«       ‚|d   S ©Nr7   ry  z%expected one or two code outputs got r   )r3  rR   )rÝ   r„   r¤  r  s       rW   Úget_triton_coder6  ˜
  sQ   € ä˜BÐ0 Ò0¨Ñ0€Là”�LÓ!Ô& QÔ&ð Ø
/´°LÓ0AÐ/BÐCóÐ&ð ˜‰?Ðr{   c                ó‚   — t        | g|¢­i |¤Ž\  }}dt        |«      cxk  rdk  sn J dt        |«      › �«       ‚|d   S r5  )r  rR   )rÝ   r„   r¤  rã   r  s        rW   Úrun_and_get_triton_coder8  ¢
  sW   € ô ' rÐ;¨DÒ;°FÑ;�O€A€|à”�LÓ!Ô& QÔ&ð Ø
/´°LÓ0AÐ/BÐCóÐ&ð ˜‰?Ðr{   c                óÖ   ‡‡‡— ddl mŠ ddlm} |j                  Šg Šdˆˆˆfd„}t
        j                  j                  |d|«      5   | |i |¤Ž}d d d «       |‰fS # 1 sw Y   ‰fS xY w)Nr   r<   rD   c                 ó^   •—  ‰| i |¤Ž | d   }t        |‰«      sJ ‚‰j                  |«       y )Nry  )r}   r™  )r„   r¤  rq  r=   Úgraph_loweringsÚ	real_inits      €€€rW   Ú	fake_initz-run_and_get_graph_lowering.<locals>.fake_init·
  s7   ø€ Ù�4Ð"˜6Ò"Ø�Q‘ˆÜ˜% Ô/Ð/Ð/Ø×Ñ˜uÕ%r{   r5  r+  )Útorch._inductor.graphr=   Útorch._inductor.output_coderE   r5  r   r  r�  )	rÝ   r„   r¤  rE   r=  r°  r=   r;  r<  s	         @@@rW   Úrun_and_get_graph_loweringr@  ®
  sq   ú€ õ 4Ý;à×(Ñ(€IØ€O÷&ô 
�‰×	Ñ	˜?¨J¸	Ó	Bñ %Ù�TÐ$˜VÑ$ˆ÷%ð �?Ð"Ð"÷%ð �?Ð"Ð"ús   Á	AÁA(c              #  óÌ   K  — ddl m} |j                  |    }	 t        j                  ||«      |j                  | <   d–— ||j                  | <   y# ||j                  | <   w xY w­w)z
    Override the lowering of aten_op with override_fn.
    The first argument of override_fn is the original lowering fn.
    r   )ÚloweringN)Útorch._inductorrB  Ú	loweringsr  Úpartial)Úaten_opÚoverride_fnrB  Úorig_fns       rW   Úoverride_loweringrI  Ã
  s`   è ø€ õ )à× Ñ  Ñ)€Gð.Ü&/×&7Ñ&7¸ÀWÓ&Mˆ×Ñ˜7Ñ#Ûà&-ˆ×Ñ˜7Ò#ø gˆ×Ñ˜7Ò#üs   ‚A$™'A Á A$ÁA!Á!A$c                ó’   ‡ ‡‡— ddl m} |j                  Šdˆˆˆ fd„}t        j                  j
                  j                  |d|«      S )zr
    Add hook functions to be called at the beginning and end of Scheduler.__init__.
    Used for unit tests.
    r   )Ú	Schedulerc                óB   •—  ‰| |«        ‰| |«      }‰r	 ‰| |«       |S r‹   rš   )rÌ  rF  ÚoutrH  Úpost_fnÚpre_fns      €€€rW   rð  z(add_scheduler_init_hook.<locals>.wrapperà
  s+   ø€ Ùˆy˜%Ô Ù�i Ó'ˆÙÙ�I˜uÔ%Øˆ
r{   r5  )rÌ  r   rF  r   r’   r   )Útorch._inductor.schedulerrK  r5  Úunittestr   r  r�  )rO  rN  rK  rð  rH  s   ``  @rW   Úadd_scheduler_init_hookrR  Õ
  s9   ú€ õ 4à× Ñ €G÷ô �=‰=×Ñ×%Ñ% i°¸WÓEÐEr{   c                óz   — t         j                  rt        j                  | «       yt        j	                  | «       y)z·
    Warnings that will be actionable for PyTorch developers, but not
    end users.  Allows us to easily disable them in stable releases but
    keep them on for nightly builds.
    N)ri   Údeveloper_warningsrÓ   rö  Úinfo)Úmsgs    rW   Údeveloper_warningrW  ê
  s$   € ô × Ò Ü�‰�CÕä�‰��r{   c                 ó´  — 	 t         j                  j                  d«      } | dz   t        t         j                  «      k  rTt        t         j                  | dz      «      dkD  r2t         j                  | dz      d   dk7  rt         j                  | dz      S t         j                  D ]#  }|j                  d«      sŒ|t        d«      d c S  y# t        $ r Y ŒBw xY w)a  
    An experimental API used only when config.benchmark_kernel is true.

    The benchmark name is only available at codegen time. So we can not
    directly call it in benchmark_all_kernels which is run after codegen.

    The function assumes the argument after --only is the benchmark name.
    It works for torchbench.py/hugginface.py/timm_models.py. But for ad-hoc
    scripts, this function may return None.

    There are 2 flavors of --only argument we need handle:
    1. --only model_name
    2. --only=model_name
    z--onlyr7   r   r€  z--only=N)rÈ  Úargvr  rR   Ú
ValueErrorrÊ  )rU  r˜  s     rW   Úget_benchmark_namer[  ö
  sÀ   € ð	Ü�h‰h�n‰n˜XÓ&ˆà�!‰G”cœ#Ÿ(™(“mÒ#Ü”C—H‘H˜S 1™WÑ%Ó&¨Ò*Ü—‘˜˜q™Ñ! !Ñ$¨Ò+ä—8‘8˜C !™GÑ$Ð$ô �x‰xò )ˆØ�>‰>˜)Õ$Ø”s˜9“~Ð'Ð(Ò(ð)ð øô ò Ùðús   ‚BC Ã	CÃCc                ó&   — t        d„ | D «       «      S )Nc              3  ó&   K  — | ]	  }|d k(  –— Œ y­w©r7   Nrš   r_  s     rW   r½   zis_ones.<locals>.<genexpr>  ó   è ø€ Ò%˜!ˆq�A�vÑ%ùó   ‚©r�   ©rJ  s    rW   Úis_onesrc    ó   € ÜÑ%˜uÔ%Ó%Ð%r{   c                ó&   — t        d„ | D «       «      S )Nc              3  ó&   K  — | ]	  }|d k(  –— Œ y­w)r   Nrš   r_  s     rW   r½   zis_zeros.<locals>.<genexpr>  r_  r`  ra  rb  s    rW   Úis_zerosrg    rd  r{   c                ó&   — t        d„ | D «       «      S )Nc              3  ó”   K  — | ]@  }t        |t        j                  «      r$|j                  t        j                  d «      k(  –— ŒB y­w)r  N)r}   rP   rž  rª   )r»   rÒ   s     rW   r½   z is_cpu_device.<locals>.<genexpr>   s8   è ø€ ò àÜ�dœEŸL™LÔ)ð 	�‰”u—|‘| EÓ*Õ*ñùs   ‚AAra  )Úinputss    rW   Úis_cpu_devicerk    s   € Üñ àôó ð r{   c                óœ   — t        | t        j                  «      sJ d«       ‚| j                  rt        j
                  S t        j                  S )Nz8only support sympy.Expr as input to get_sympy_Expr_dtype)r}   r~   r+  r˜   rP   rU  rI  ri  s    rW   Úget_sympy_Expr_dtyperm  '  s=   € Ü�cœ5Ÿ:™:Ô&ð ØBóÐ&ð ‡~‚~Ü�{‰{Ðä�}‰}Ðr{   c              /  óŠ   K  — | r-t        j                  j                  |i |¤Ž5 }|–— d d d «       y d –— y # 1 sw Y   y xY w­wr‹   )rP   rÊ   rË   )Úshould_profiler„   r¤  rç   s       rW   Úmaybe_profilerp  1  sE   è ø€ áÜ�^‰^×#Ñ# TÐ4¨VÑ4ð 	¸ØŠG÷	ð 	ô 	÷	ð 	üs   ‚"A¤7©A·A ¼Ac                 ól   — t         j                  j                  } | dk  rt        j                  «       } | S ©Nr7   )ri   rò  ÚthreadsrP   Úget_num_threads)rs  s    rW   rø  rø  :  s+   € Ü�j‰j× Ñ €GØ�‚{Ü×'Ñ'Ó)ˆØ€Nr{   c                 ó€   — ddl m}   | «       }|j                  dt        j                  j
                  rd«      S d«      S )Nr7   )Úget_backend_optionsÚ
num_stagesry  rÒ  )Úruntime.triton_helpersrv  rH  rP   rN  rq   )rv  Úoptionss     rW   Úget_backend_num_stagesrz  A  s2   € å;á!Ó#€GØ�;‰;�|¬%¯-©-×*;Ò*; QÓCÐCÀÓCÐCr{   c                ó  — t        | t        j                  j                  j                  j
                  dk(  ¬«      }|�|S ddlm}m} t        j                  j                  «       xr! t        j                  j                  «       dk\  }| t        j                  t        j                  t        j                  fv sJ ‚t        j                  |«      j                   j#                  d«      r™ddlm}  |«       }| t        j                  t        j                  fv r|r	 || |«      S t        j                  j                  j                  j
                  dk(  r |t        j                  |«      S  |t        j                  |«      S | t        j                  t        j                  fv r
|r || «      S t        j                  j                  j                  j
                  dk(  r |t        j                  «      S  |t        j                  «      S )z¨
    We don't want to throw errors in this function. First check to see if the device is in device_info.py,
    then fall back to the inaccurate triton estimation.
    Útf32)Úis_tf32r   )Úget_max_simd_tflopsÚget_max_tensorcore_tflops)rv   r   Ú
clock_rate)Úmax_clock_rate)r   rP   ÚbackendsrH   ÚmatmulÚfp32_precisionÚtriton.testingr~  r  rQ   Úget_device_capabilityrÃ   rE  rG  ÚinspectÚ	signatureÚ
parametersrH  Útorch._utils_internalr�  )r©   Úds_topsr~  r  ÚSM80OrLaterr�  Úsm_clocks          rW   Úget_device_tflopsrŽ  I  s  € ô Ø”u—~‘~×*Ñ*×1Ñ1×@Ñ@ÀFÑJô€Gð ÐØˆçMä—*‘*×)Ñ)Ó+ò ´·
±
×0PÑ0PÓ0Rð Wñ 1€Kð
 ”U—]‘]¤E§N¡N´E·M±MÐBÑBÐBÐBä×ÑÐ,Ó-×8Ñ8×<Ñ<¸\ÔJå8á!Ó#ˆØ”U—]‘]¤E§N¡NÐ3Ñ3¹Ù,¨U°HÓ=Ð=ä�>‰>×Ñ×%Ñ%×4Ñ4¸Ò>Ù,¬U¯]©]¸HÓEÐEá&¤u§}¡}°hÓ?Ð?à”U—]‘]¤E§N¡NÐ3Ñ3¹Ù,¨UÓ3Ð3ä�>‰>×Ñ×%Ñ%×4Ñ4¸Ò>Ù,¬U¯]©]Ó;Ð;á&¤u§}¡}Ó5Ð5r{   c                 ó   — ddl m}   | «       S )Nr   ©Úget_dram_gbps)r…  r‘  r�  s    rW   Úget_gpu_dram_gbpsr’  t  s   € å,á‹?Ðr{   c                 óx   — ddl m}  | j                  j                  j	                  d«      j                  dd«      S )Nr   ©rÖ  Úmax_shared_mem)Útriton.runtimerÖ  r×  rØ  r½  rH  r”  s    rW   Úget_gpu_shared_memoryr—  {  s.   € Ý%à�=‰=×Ñ×4Ñ4°QÓ7×;Ñ;Ð<LÈaÓPÐPr{   c                 óú   — t         j                  j                  «       rUt         j                  j                  «       j                  } t         j                  j                  «       j
                  }|| z  S d} d}|| z  S )Nr  i   )rP   rH   rQ   r½  Ú	warp_sizeÚmax_threads_per_block)r™  rš  s     rW   Úget_max_numwarpsr›  �  sh   € Ü‡z�z×ÑÔ Ü—J‘J×4Ñ4Ó6×@Ñ@ˆ	ä %§
¡
× @Ñ @Ó B× XÑ XÐð
 ! IÑ-Ð-ð ˆ	Ø $ÐØ  IÑ-Ð-r{   c                ó$   — | j                  d«      S )NÚwelford)rÊ  ©Úreduction_types    rW   Úis_welford_reductionr   �  s   € Ø×$Ñ$ YÓ/Ð/r{   c                ó(   — t        | «      ry| dk(  ryy)NrÒ  Úonline_softmax_reducery  r7   )r   rž  s    rW   Úreduction_num_outputsr£  ‘  s   € Ü˜NÔ+ØØ	Ð2Ò	2Øàr{   c                 ó0   — t        j                  «       dk(  S )NÚLinux)ÚplatformÚsystemrš   r{   rW   Úis_linuxr¨  š  s   € Ü�?‰?Ó Ñ'Ð'r{   c                 ó(   — t         j                  dk(  S )Nrk   )rÈ  r¦  rš   r{   rW   r  r  ž  s   € Ü�<‰<˜7Ñ"Ð"r{   c                ó&   — t        d„ | D «       «      S )Nc              3  ón   K  — | ]-  }t        |t        j                  «      xr |j                   –— Œ/ y ­wr‹   )r}   r~   r+  r˜  r_  s     rW   r½   z#has_free_symbols.<locals>.<genexpr>£  s)   è ø€ ÒJÀŒz˜!œUŸZ™ZÓ(Ò<°·±¨_Ó<ÑJùs   ‚35r¬  )Úitrs    rW   rõ  rõ  ¢  s   € ÜÑJÀcÔJÓJÐJr{   c            	     óx  — ddl m} | D ]¯  }t        ||j                  |j                  |j
                  |j                  |j                  f«      r=t        |j                  «       xs d«      st        |j                  «       xs d«      sŒ� yt        ||j                  «      sŒšt        dt        |«      › �«      ‚ y)Nr7   rü  rš   Tzunexpected type for is_dynamic F)rþ  rý  r}   r%  r'  rö  rk  r>   rõ  Úmaybe_get_sizeÚmaybe_get_strider@   Ú	TypeErrorr  )r„   rý  r  s      rW   r^  r^  ¦  s§   € Ýàò IˆÜØ�—‘˜bŸm™m¨R¯[©[¸"×:KÑ:KÈRÏYÉYÐWô
ô   × 0Ñ 0Ó 2Ò 8°bÔ9Ô=MØ×"Ñ"Ó$Ò*¨õ>ñ Ü˜A˜rŸy™yÔ)ØäÐ=¼dÀ1»g¸YÐGÓHÐHðIð r{   c                  ó   — e Zd ZdZdZy)ÚPlaceholderÚKERNEL_NAMEÚDESCRIPTIVE_NAMEN)r“   r”   r•   r³  r´  rš   r{   rW   r²  r²  º  s   „ ð  €Kð *Ñr{   r²  c                ó~  — ddl m} t        j                  dd¬«      5 }t	        j
                  «       }t	        j
                  «       } t        |t        |«      ¬«      j                  |Ž  t        d|j                  › �|¬«       t        |j                  |¬«       t        j                  «       }t        ||«      5   | |j                  «       d d d «       t        j                  «       |z
  }	 ||j                  «       |j                  j                  «        |j                  «        t        d	|j                  › �|¬«       t        |j                  |¬«       |j!                  «       |j!                  «       k(  }
t"        j%                  d
||j&                  |
|	«       d d d «       y # 1 sw Y   ŒÚxY w# 1 sw Y   y xY w)Nr7   )Ústable_topological_sortr_  zutf-8)ÚmodeÚencoding)r¥  Ú	fake_modezBefore:
)ÚfilezAfter:
zZ%s, save before/after graph to %s, graph before/after are the same = %s, time elapsed = %s)Úpattern_matcherr¶  rü  ÚNamedTemporaryFileÚior   r^   rZ   Ú	propagaterµ  rq  r   Únowr]   ÚlintÚ	recompiler>  rÓ   rU  rÛ   )r†  r¥  ÚinprV  r¶  r
  Ú	before_ioÚafter_ioÚ
start_timeÚtime_elapsedr  s              rW   Úpass_execution_and_saverÇ  Ä  sU  € õ 9ä	×	$Ñ	$ØØô
ð 
ð 
Ü—K‘K“Mˆ	Ü—;‘;“=ˆØCŒ	�RÔ#3°CÓ#8Ô9×CÑCÀSÑIÜ�	˜"Ÿ(™(˜Ð$¨1Õ-Üˆb�h‰h˜YÕ'Ü—\‘\“^ˆ
Ü# B¨Ó,ñ 	Ù�—‘ŒN÷	ä—|‘|“~¨
Ñ2ˆá §¡Ô)Ø
�‰�‰ŒØ
�‰Œä�˜Ÿ™˜
Ð#¨!Õ,Üˆb�h‰h˜XÕ&Ø×ÑÓ  H×$5Ñ$5Ó$7Ñ7ˆÜ�‰ØhØØ�F‰FØØô	
÷+
ð 
÷	ð 	ú÷
ð 
ús%   žBF3Â;F'ÃCF3Æ'F0	Æ,F3Æ3F<c                ó^   — ddl m} t        | |j                  «      xr | j	                  «       S )zB
    Check if input buffer is a multi-outputs template buffer
    r7   rü  )rþ  rý  r}   rn  Úis_multi_outputs_template©Ú	input_bufrý  s     rW   rÉ  rÉ  ç  s/   € õ ô 	�9˜b×/Ñ/Ó0ò 	2Ø×/Ñ/Ó1ðr{   c                ó¢   — ddl m} t        | |j                  «      xr2 t	        | j
                  «      dk(  xr t        | j
                  d   «      S )zL
    Check if input buffer is a output of multi-outputs template buffer
    r7   rü  r   )rþ  rý  r}   ÚMultiOutputrR   rj  rÉ  rÊ  s     rW   Ú#is_output_of_multi_outputs_templaterÎ  ó  sL   € õ ô 	�9˜bŸn™nÓ-ò 	;Ü�	× Ñ Ó! QÑ&ò	;ä% i×&6Ñ&6°qÑ&9Ó:ðr{   c                ó&  — | €yddl m} t        | |j                  «      xr- t        | |j                  «       xr |d u xs | j
                  |u x�sA t        | «      |j                  u x�r' t        t        j                  j                  d«      xr; | j
                  t        j                  j                  j                  j                  k(  xsÄ t        t        j                  j                  d«      xr; | j
                  t        j                  j                  j                  j                  k(  xsa t        t        j                  j                  d«      xr; | j
                  t        j                  j                  j                  j                  k(  S )NFr7   rü  Úall_to_all_singleÚall_gather_into_tensorÚreduce_scatter_tensor)rþ  rý  r}   Ú_CollectiveKernelÚ_WaitKernelÚop_overloadr  ÚFallbackKernelr  rP   r  ÚtorchrecrÐ  ÚdefaultrÑ  rÒ  ©rg  r€  rý  s      rW   Úis_collectiverÚ    sG  € ð €|Øåô 	�4˜×-Ñ-Ó.ò 	3Ü˜4 §¡Ó0Ð0ò	3à�4ˆZÒ1˜4×+Ñ+¨rÐ1óô 	ˆT‹
�b×'Ñ'Ð'ó 	
ô
 œŸ	™	×*Ñ*Ð,?Ó@ò UØ×$Ñ$¬¯	©	×(:Ñ(:×(LÑ(L×(TÑ(TÑTòô
 œŸ	™	×*Ñ*Ð,DÓEò EØ×$Ñ$Ü—9‘9×%Ñ%×<Ñ<×DÑDñEòô œŸ	™	×*Ñ*Ð,CÓDò YØ×$Ñ$¬¯	©	×(:Ñ(:×(PÑ(P×(XÑ(XÑXð/r{   c                ó<   — ddl m} t        | «      |j                  u S ©Nr7   rü  )rþ  rý  r  rÔ  )rg  rý  s     rW   Úis_waitrÝ  (  s   € Ýä�‹:˜Ÿ™Ð'Ð'r{   c                ó¨   — ddl m} t        | |«      rt        d„ | j                  D «       «      S t        | j                  «      xr |d u xs  || «      S )Nr   ©ÚGroupedSchedulerNodec              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr‹   )Úcontains_collectiver_  s     rW   r½   z&contains_collective.<locals>.<genexpr>5  s   è ø€ Ò@¨aÔ& q×)Ñ@ùr«  )rP  rà  r}   r|  ÚsnodesrÚ  rg  )ÚsnodeÚ	filter_fnrà  s      rW   râ  râ  .  sJ   € õ ?ä�%Ð-Ô.ÜÑ@°5·<±<Ô@Ó@Ð@ä˜Ÿ™Ó$ÒP¨)°tÐ*;Ò*O¹yÈÓ?OÐPr{   c                óˆ   — ddl m} t        | |«      rt        d„ | j                  D «       «      S t        | j                  «      S )Nr   rß  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wr‹   )Úcontains_waitr_  s     rW   r½   z contains_wait.<locals>.<genexpr>>  s   è ø€ Ò:¨”= ×#Ñ:ùr«  )rP  rà  r}   r|  rã  rÝ  rg  )rä  rà  s     rW   rè  rè  :  s4   € Ý>ä�%Ð-Ô.ÜÑ:¨U¯\©\Ô:Ó:Ð:ä�u—z‘zÓ"Ð"r{   c                ó¨   — ddl m} t        |t        j                  j
                  «      r|g}t        | |j                  «      xr | j                  |v S rÜ  )rþ  rý  r}   rP   r�  rŽ  rÖ  rÕ  rÙ  s      rW   Úis_fallback_oprê  C  sE   € õ ä�"”e—j‘j×+Ñ+Ô,ØˆTˆÜ�d˜B×-Ñ-Ó.ÒI°4×3CÑ3CÀrÐ3IÐIr{   c                óB   — |||    j                   j                  «          S r‹   )Údefining_oprÍ  )Úbuf_nameÚname_to_bufÚname_to_fused_nodes      rW   Úbuf_name_to_fused_snoderð  N  s#   € ð ˜k¨(Ñ3×?Ñ?×HÑHÓJÑKÐKr{   c                 ó   — yrƒ  rš   ©rä  s    rW   r„  r„  Y  r…  r{   c                ó²   —  || «      ry |j                  | «       | j                  D ].  }t        |j                  ||«      }||v rŒt	        |||||¬«       Œ0 y )N©Úcriteria_cb)rK  Úunmet_dependenciesrð  rÛ   Úfind_recursive_deps_of_node)rä  Úcollected_node_setrî  rï  rõ  ÚdepÚdefining_op_for_deps          rW   r÷  r÷  T  sn   € ñ �5ÔØØ×Ñ˜5Ô!Ø×'Ñ'ò 
ˆÜ5Ø�H‰H�kÐ#5ó
Ðð Ð"4Ñ4ØÜ#ØØØØØ#ö	
ñ
r{   c                 ó   — yrƒ  rš   rò  s    rW   r„  r„  r  r…  r{   c           	     óz  —  || «      ry |j                  | «       | j                  «       D ]Ž  }|j                  D ]}  }|j                  €J ‚|j                  j	                  «       dk(  rŒ/|j                  j	                  «       |vrŒL||j                  j	                  «          }||v rŒnt        |||||¬«       Œ Œ� y )NÚOUTPUTrô  )rK  Úget_outputsr�  rg  rÍ  Úfind_recursive_users_of_node)rä  rø  rî  rï  rõ  Úorg  Úuser_ops           rW   rÿ  rÿ  m  sÃ   € ñ �5ÔØØ×Ñ˜5Ô!Ø×ÑÓ ò ˆØ—G‘Gò 	ˆDØ—9‘9Ð(Ð(Ð(Ø�y‰y×!Ñ!Ó# xÒ/ØØ�y‰y×!Ñ!Ó#Ð+=Ñ=ØØ(¨¯©×);Ñ);Ó)=Ñ>ˆGØÐ,Ñ,ØÜ(ØØ"ØØ"Ø'öñ	ñr{   c                ób   — t         j                  j                  j                  rdnd}|| z
  |z
  S )zaComputes the number of inputs to the aot fw graph which have fixed addresses (params and buffers)ry  r   )rP   Ú
_functorchri   Úfunctionalize_rng_ops)Údynamo_gm_num_inputsÚaot_fw_gm_num_inputsÚnum_rng_seed_offset_inputss      rW   Únum_fw_fixed_argumentsr  Š  s6   € ô ×Ñ×$Ñ$×:Ò:‰Àð ð  Ð"6Ñ6Ð9SÑSÐSr{   c                óþ   — dd„}d}g }| j                   j                  D ]0  }|j                  dk(  sŒ ||«      r|j                  |«       |dz  }Œ2 |t	        t        t        |«      «      «      k(  sJ ‚t        |«      S )z>
    Infers which inputs are static for a backwards graph
    c                ó~   — d| j                   vxr. d| j                   vxr d| j                   vxr d| j                   vS )NÚtangentsÚbwd_seedÚbwd_base_offsetÚbwd_rng_stater¡  rí  s    rW   Úis_saved_tensorz'count_tangents.<locals>.is_saved_tensorš  sH   € à˜aŸf™fÐ$ò .Ø !§&¡&Ð(ò.à!¨¯©Ð/ò.ð   q§v¡vÐ-ð		
r{   r   rš  r7   )rT   r5   r’   r1  )rq  rF  r€  r™  r]  rÆ   rR   )Úfx_gr  Ú	arg_countÚstatic_arg_idxsr  s        rW   Úcount_tangentsr  •  s�   € ó

ð €IØ€OØ�Z‰Z×Ñò ˆØ�4‰4�=Ó Ù˜qÔ!Ø×&Ñ& yÔ1Ø˜‰N‰Ið	ð œd¤5¬¨_Ó)=Ó#>Ó?Ò?Ð?Ð?ÜˆÓÐr{   c                óÖ   — g }t        | j                  j                  «      D ]D  \  }}|j                  dk7  r |S |j                  j                  d«      sŒ4|j                  |«       ŒF |S )a  
    Returns indices of backward graph inputs that are always at fixed
    addresses: primals (parameters/buffers/user inputs saved for backward).
    Excludes saved activations which may not be at fixed addresses when
    the forward is partitioned for CUDA graphs.
    rš  Úprimals_)rE  rq  rF  r€  rÛ   rÊ  r™  )r  Ústatic_idxsrU  r  s       rW   Úget_static_bw_input_idxsr  ®  sl   € ð €KÜ˜DŸJ™J×,Ñ,Ó-ò $‰ˆˆQØ�4‰4�=Ò Øð Ðð �6‰6×Ñ˜ZÕ(Ø×Ñ˜sÕ#ð	$ð
 Ðr{   c                  ó2   — e Zd ZU ded<   dd„Zedd„«       Zy)Ú	BoxedBoolr1  r�   c                ó   — | j                   S r‹   )r�   rD  s    rW   rO  zBoxedBool.__bool__Â  s   € Ø�z‰zÐr{   c                ó6   — t        | t        «      r	d| _        | S yrƒ  )r}   r  r�   rå  s    rW   ÚdisablezBoxedBool.disableÅ  s   € ä�cœ9Ô%ØˆCŒIØˆJØr{   Nr„  )r¾  r   r’   zBoxedBool | bool)r“   r”   r•   r¤   rO  rá  r  rš   r{   rW   r  r  ¾  s    … àƒKóð òó ñr{   r  c              #  óØ   ‡ ‡K  — ddl m} |j                  Š	 	 	 d	 	 	 	 	 	 	 	 	 	 	 	 	 dˆ ˆfd„}t        j                  j                  |d|«      5  d –— d d d «       y # 1 sw Y   y xY w­w)Nr7   r9   c                ó@   •— ‰j                  |«        ‰| |||||«      S r‹   r$  )rÜ  Úkernel_namer0  rX  ÚgpuÚcpp_definitionÚkernel_listÚorig_define_kernels         €€rW   Údefine_kernelz.collect_defined_kernels.<locals>.define_kernelÓ  s-   ø€ ð 	×Ñ˜;Ô'Ù!Ø�+˜{¨H°c¸>ó
ð 	
r{   r$  )NTN)rÜ  r:   r  r
  r0  r
  rX  r¦  r   r1  r!  r¦  r’   r   )Úcodegen.wrapperr:   r$  r   r  r�  )r"  r:   r$  r#  s   `  @rW   Úcollect_defined_kernelsr&  Í  s˜   ùè ø€ å5à-×;Ñ;Ðð  $ØØ%)ð
Ø"ð
àð
ð ð
ð ð	
ð
 ð
ð #ð
ð 
ö
ô 
�‰×	Ñ	Ð/°À-Ó	Pñ Û÷÷ ñ üs   „AA*ÁAÁ	A*ÁA'Á#A*c                ó   — | dz   S )NÚ__original__rš   r¡  s    rW   Ú get_cloned_parameter_buffer_namer)  ä  s   € Ø�.Ñ Ð r{   c                ó   — | t         v S r‹   )rN   r¨  s    rW   rØ  rØ  è  s   € Ø”YÐÐr{   c                 ó:   — t         j                  j                  duS )z,Check if we're running on ROCm/HIP platform.N)rP   rN  rq   rš   r{   rW   Úis_rocmr,  ì  s   € ä�=‰=×Ñ DÐ(Ð(r{   c                ó&   — | dk7  xr t        | «      S )NrI   )rØ  r¨  s    rW   Údevice_need_guardr.  ñ  s   € Ø�U‰?Ò-œv f›~Ð-r{   c                óN  — | t         j                  k(  r?t         j                  j                  «       r!t         j                  j	                  «       dk  S | t         j                  k(  rt         j
                  j                  «       ry| t         j                  t         j                  fv S )N)r°  r   T)rP   rE  rH   rQ   r†  rJ   rU  r1  r$  s    rW   Ú,needs_fallback_due_to_atomic_add_limitationsr0  õ  sk   € Ø”—‘Ò¤5§:¡:×#:Ñ#:Ô#<Ü�z‰z×/Ñ/Ó1°FÑ:Ð:Ø	”%—.‘.Ò	 ¤U§Y¡Y×%;Ñ%;Ô%=ØàœŸ™¤e§j¡jÐ1Ð1Ð1r{   c                óü  — | j                   t        j                  j                  j                  t        j                  j                  j
                  fv r|€y| j                   t        j                  j                  j                  k(  rdnd}|d |fvxsé |xr t        |«      xr t        |«      xsË | j                   t        j                  j                  j                  k(  xrW |dk(  xrP |xrL |dk(  xrE t        j                  j                  xr) t        j                  j                  xs t        «       dk7  xs? ||k(  xr" |t        j                  t        j                  fv xs t        j                   «       S )NFrK  rü   r  r7   )ÚoverloadpacketrP   r  ÚatenÚscatter_reduce_Úscatter_reduceÚscatter_rØ  r0  ri   rò  Úfallback_scatter_reduce_sumÚdynamic_threadsrø  r1  rU  Ú$are_deterministic_algorithms_enabled)rÕ  rŸ  Ú
self_dtypeÚ	src_dtypeÚsrc_device_typeÚsrc_is_tensorÚ	reduce_tys          rW   Úuse_scatter_fallbackr?  þ  sZ  € ð 	×"Ñ"Ü�I‰I�N‰N×*Ñ*¬E¯I©I¯N©N×,IÑ,IÐJñ	KàÐ"àð ×+Ñ+¬u¯y©y¯~©~×/FÑ/FÒF‰ÈEð ð
 	˜t YÐ/Ð/ò 	8àò HÜ�Ó'òHä<¸YÓGò		8ð ×&Ñ&¬%¯)©)¯.©.×*HÑ*HÑHò LØ %Ñ'òLàòLð   5Ñ(òLô —
‘
×6Ñ6ò	Lô
 —‘×+Ñ+ÒJÔ/CÓ/EÈÑ/Jò	8ð ˜iÑ'ÒS¨J¼5¿:¹:ÄuÇ{Á{Ð:SÐ,Sò	8ô ×5Ñ5Ó7ð!r{   c                ó   — ddl m}m} ddlm} t        dt        | «      › d�«       t        | «      D �]  \  }}t        d|d›d�«       ||u rt        d	«       Œ'||u rt        d
«       Œ7t        ||«      rÂ|j                  «       }t        |rdnd› d�«       |r:|j                  €J ‚t        d|j                  j                  j                  › �«       t        d«       |j                  j                  D ]  }t        |«       Œ t        d«       |j                  j                  D ]  }t        |«       Œ �Œt!        dt#        |«      › �«      ‚ y)z�
    An API that can be used in pdb to dump a node_schedule.
    Right mainly dump the read/write dependencies but can add more as needed.
    r   )ÚDisableReductionÚEnableReduction)ÚSchedulerNodezNode schedule with z nodesr!  Ú3r»  zenable reductionzdisable reductionÚredÚpwz scheduler nodeNzoriginal reduction hint zReadDep:z	WriteDep:zUnrecognized node type: )Útorch._inductor.codegen.simdrA  rB  rP  rC  rµ  rR   rE  r}   Úis_reductionrg  r&  Úreduction_hintr9  r:  r;  rú   r  )r  rA  rB  rC  rU  rg  Úis_redrù  s           rW   Údump_node_schedulerK  %  s=  € ÷
 OÝ7ä	Ð¤ MÓ 2Ð3°6Ð
:Ô;Ü˜}Ó-ó H‰	ˆˆTÜ��#�a�˜ˆlÔà�?Ñ"ÜÐ$Õ%àÐ%Ñ%ÜÐ%Õ&Ü˜˜mÔ,Ø×&Ñ&Ó(ˆFÜ™f‘U¨$Ð/¨Ð?Ô@ÙØ—y‘yÐ,Ð,Ð,ÜÐ0°·±·±×1NÑ1NÐ0OÐPÔQÜ�*ÔØ×'Ñ'×-Ñ-ò �Ü�c•
ðä�+ÔØ×'Ñ'×.Ñ.ò �Ü�c•
òô Ð!9¼$¸t»*¸ÐFÓGÐGñ+Hr{   c                óz   — ddl m}  || j                  «       t        | j                  «      z  t
        z  dk(  «      S )Nr   )r+  )r±  r+  Ústorage_offsetr&  r©   ÚGPU_ALIGN_BYTES)rÏ   r+  s     rW   Útensor_is_alignedrO  F  s:   € õ Lá Ø	×	Ñ	Ó	 ¤>°&·,±,Ó#?Ñ	?Ä?ÑRÐVWÑWóð r{   c                ó|   — t        | j                  j                  «      syt        j                  xs t        | «      S rƒ  )rØ  rª   r  ri   Úassume_aligned_inputsrO  )Úexample_inputs    rW   Úshould_assume_input_alignedrS  T  s2   € ô �-×&Ñ&×+Ñ+Ô,ØÜ×'Ñ'ÒKÔ+<¸]Ó+KÐKr{   c                 ó6  — t         j                  j                  j                  «       } | st	        j
                  «       S | j                  r| j                  j                  st	        j
                  «       S | j                  j                  }|j                  «       S r‹   )	rP   Ú_guardsÚTracingContextÚtry_getrc  Únullcontextr¹  rs  Úsuppress_guards)Útracing_contextrs  s     rW   Ú#maybe_get_suppress_shape_guards_ctxr[  ]  sv   € ô
 —m‘m×2Ñ2×:Ñ:Ó<€OÙÜ×%Ñ%Ó'Ð'ð ×$Ò$¨O×,EÑ,E×,OÒ,OÜ×%Ñ%Ó'Ð'Ø×)Ñ)×3Ñ3€IØ×$Ñ$Ó&Ð&r{   c                ó  — t         j                  j                  j                  t        dd«      5  t
        j                  j                  «        dd l}dd l	} |j                  «       } |j                  |«      }ddlm} |j                  |«       |j                  }|j!                  |j"                  «        | |i |¤Ž}	|j%                  «       }
|j!                  |«       |j'                  |«       d d d «       |	|
fS # 1 sw Y   	
fS xY w)NrÔ   Tr   )Úoutput_code_log)rQ  r   r  r�  ri   rP   r  r  r½  Úloggingr   ÚStreamHandlerÚtorch._inductor.codecacher]  Ú
addHandlerÚlevelÚsetLevelÚDEBUGr>  ÚremoveHandler)rÝ   r„   r¤  r½  r^  Úlog_capture_stringÚchr]  Ú
prev_levelr°  ré   s              rW   Úrun_and_get_cpp_coderi  m  sæ   € ô 
�‰×	Ñ	×	#Ñ	#¤F¨G°TÓ	:ñ *Ü�‰×ÑÔÛÛà(˜RŸ[™[›]ÐØ"ˆW×"Ñ"Ð#5Ó6ˆÝ=à×"Ñ" 2Ô&Ø$×*Ñ*ˆ
Ø× Ñ  §¡Ô/Ù�TÐ$˜VÑ$ˆØ×'Ñ'Ó)ˆØ× Ñ  Ô,Ø×%Ñ% bÔ)÷*ð  �1ˆ9Ð÷!*ð  �1ˆ9Ðús   °CC>Ã>D
c                ó  — t        | «      }|�|j                  S | D ]á  }t        |t        j                  «      r|j
                  j                  c S t        |t        j                  «      sŒP|j                  «       D ]6  }t        |t        j                  «      sŒ|j
                  j                  c c S  |j                  «       D ]6  }t        |t        j                  «      sŒ|j
                  j                  c c S  Œã y r‹   )	rZ   rs  r}   rP   r2   rg  rž  rÈ  r4  )rj  r¹  ÚinputrÈ  r4  s        rW   Úshape_env_from_inputsrl  †  sÌ   € Ü  Ó(€Ið ÐØ×"Ñ"Ð"ð ò 1ˆÜ�eœUŸ\™\Ô*Ø—:‘:×'Ñ'Ò'ô �eœUŸ\™\Õ*ØŸ
™
›ò /�Ü˜d¤E§L¡LÕ1ØŸ9™9×.Ñ.Ô.ð/ð  Ÿ,™,›.ò 1�Ü˜f¤e§l¡lÕ3Ø!Ÿ;™;×0Ñ0Ô0ñ1ð1ð r{   c                ó<   ‡ ‡‡— t        ‰«      dk(  r‰ S dˆˆ ˆfd„}|S )Nr   c                óz   •— t        | ‰‰«      \  }} ‰| «      }t        |«      rt        j                  ||«       |S r‹   )Úcopy_misaligned_inputsrR   rP   Ú_foreach_copy_)Ú
new_inputsÚold_tensorsÚnew_tensorsrM  Úinputs_to_checkr­  Úmutated_input_idxss       €€€rW   rÓ  z)align_inputs_from_check_idxs.<locals>.run«  sE   ø€ Ü#9Ø˜Ð);ó$
Ñ ˆ�[ñ �JÓˆô ˆ{ÔÜ× Ñ  ¨kÔ:àˆ
r{   )rq  úlist[InputType]r’   r   )rR   )r­  rt  ru  rÓ  s   ``` rW   Úalign_inputs_from_check_idxsrw  £  s#   ú€ ô
 ˆ?Ó˜qÒ Øˆ÷ð €Jr{   c                óT  — d| j                  «       v rd}n;t        d„ t        | j                  «       | j                  «       «      D «       «      dz   }t	        j
                  | |fd«      j                  «       }t	        j
                  || j                  «       | j                  «       «      S )Nr   c              3  ó2   K  — | ]  \  }}|d z
  |z  –— Œ y­wr^  rš   )r»   r0  r4  s      rW   r½   z)clone_preserve_strides.<locals>.<genexpr>Á  s   è ø€ ÒT©¨°�˜‘˜fÕ$ÑTùr«  r7   r‰   )rÈ  rü   rÐ   r4  rP   Ú
as_stridedÚclone)rT   Úneeded_sizer*  s      rW   Úclone_preserve_stridesr}  »  s„   € ØˆA�F‰F‹H�}à‰ô ÑT¼#¸a¿f¹f»hÈÏÉË
Ó:SÔTÓTÐWXÑXð 	ô ×Ñ˜a + °Ó6×<Ñ<Ó>€FÜ×Ñ˜F A§F¡F£H¨a¯h©h«jÓ9Ð9r{   c                ó2  — g }g }|du}|D ]‡  }| |   }t        |t        j                  «      sJ dt        |«      › �«       ‚|j	                  «       t
        z  sŒMt        |«      | |<   |sŒ^||v sŒc|j                  |«       |j                  | |   «       Œ‰ ||fS )z¯
    Clones misaligned tensors which we inferred were aligned. Returns a tuple of [old_tensors], [new_tensors] for every
    cloned tensor which is in `return_pair_idxs`.
    Nz Expected tensors only, but got: )r}   rP   rž  r  Údata_ptrÚ	ALIGNMENTr}  r™  )rq  Úcheck_inputs_idxsÚreturn_pair_idxsrr  rs  Úret_pair_definedrè   Ú_inps           rW   ro  ro  Ç  s²   € ð ')€KØ&(€Kð (¨tÐ3ÐØò 
2ˆØ˜!‰}ˆÜ˜$¤§¡Ô-ð 	
Ø.¬t°D«z¨lÐ;ó	
Ð-ð �=‰=‹?œYÓ&Ü2°4Ó8ˆJ�q‰Mâ AÐ)9Ò$9Ø×"Ñ" 4Ô(Ø×"Ñ" :¨a¡=Õ1ð
2ð ˜Ð#Ð#r{   c                óâ   — g }|D ]N  }| |   }t        |t        j                  «      sŒ#|j                  «       t        z  dk(  sŒ>|j                  |«       ŒP t        |«      t        |«      k7  r|S |S )z[
    We require all inputs to be aligned, so introduce a copy for any
    that aren't.
    r   )r}   rP   rž  r  r€  r™  rR   )rj  Ústatic_input_idxsÚaligned_static_input_idxsrU  rk  s        rW   Úremove_unaligned_input_idxsrˆ  å  st   € ð !#ÐØ ò 2ˆØ�s‘ˆÜ�eœUŸ\™\Õ*°·±Ó0@Ä9Ñ0LÐQRÓ/RØ%×,Ñ,¨SÕ1ð2ô Ð$Ó%¬Ð->Ó)?Ò?Ø(Ð(ØÐr{   c                óP  — ddl m} t        j                  t        j                  «      j
                  }|j                  j                  j                  }|j                  j                  j                  j                  }t        j                  r'|j                  j                  j                  | |«       y|j                  j                  j                  | |k  «      ry|j                  r)|j                  j                  j                  | dk  «      ry || «      xr  || «      |k  S )zýCheck if an expression fits within 32-bit integer range.

    NOTE: This function intentionally does not install guards. Callers are
    responsible for guarding (e.g. via check_leq) when they decide to use
    32-bit indexing based on this result.
    r7   rm  Tg@Œµx¯DF)rp  rn  rP   r'  rT  rÉ   rq  rr  Úguarding_hint_or_throwrs  Úhas_guarding_hintri   Úassume_32bit_indexingÚ	check_leqr+  r€  )rê   rn  Úint_maxrŠ  r‹  s        rW   Úexpr_fits_within_32bitr�  ÷  sã   € õ ä�k‰kœ%Ÿ+™+Ó&×*Ñ*€GØŸW™W×-Ñ-×DÑDÐØŸ™×(Ñ(×2Ñ2×DÑDÐä×#Ò#Ø	�‰×Ñ×"Ñ" 1 gÔ.Øð 	‡w�w×Ñ×-Ñ-¨a°7©lÔ;Øð 	×Òð �7‰7×Ñ×1Ñ1°!°d±(Ô;ð ñ ˜QÓÒHÑ$:¸1Ó$=ÀÑ$HÐHr{   c                ó  ‡‡‡— t         j                  j                  j                  «       }|�Ü|j                  �Ït        |j                  «      dk(  sJ ‚t        | «      Š|j                  €J ‚|j                  D ]Œ  }|€|j                  j                  d «       Œ!dŠt         j                  j                  j                  «       x}r|j                  Šdˆˆfd„Š|j                  j                  t        ˆfd„|D «       «      «       ŒŽ y y y )Nr   Fc                óf   •— ‰€t        | «      S ‰r‰j                  | «      S ‰j                  | «      S r‹   )rŒ   Údeserialize_symexprÚevaluate_symexpr)rê   Úfakify_first_callrs  s    €€rW   Úmap_exprz4set_tracing_context_output_strides.<locals>.map_expr5  s7   ø€ Ø Ð(Ü" 1›v˜Ù(Ø(×<Ñ<¸QÓ?Ð?Ø$×5Ñ5°aÓ8Ð8r{   c              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wr‹   rš   )r»   rê   r•  s     €rW   r½   z5set_tracing_context_output_strides.<locals>.<genexpr>=  s   øè ø€ Ò5¨!™( 1Ÿ+Ñ5ùr  )rê   r   r’   z)float | int | SymInt | SymFloat | SymBool)
rP   rU  rV  rW  Úoutput_stridesrR   rl  r™  r”  r¿  )r®  Úcompiled_graphr)  r!  rï  r”  r•  rs  s        @@@rW   Ú"set_tracing_context_output_stridesr™  $  só   ú€ ô �m‰m×*Ñ*×2Ñ2Ó4€GØÐ˜w×5Ñ5ÐAÜ�7×)Ñ)Ó*¨aÒ/Ð/Ð/Ü)¨.Ó9ˆ	Ø×,Ñ,Ð8Ð8Ð8Ø#×2Ñ2ò 	ˆEØˆ}Ø×&Ñ&×-Ñ-¨dÕ3à$)Ð!ÜŸ-™-×6Ñ6×>Ñ>Ó@Ð@�3Ð@Ø(+×(=Ñ(=Ð%ö9ð ×&Ñ&×-Ñ-ÜÓ5¨uÔ5Ó5õñ	ð	  BÐr{   c                 ó  — t         j                  �t         j                  S t        j                  «       syt        j                  j                  «       ry	 ddlm}  | t        j                  j                  d«      k\  S # t        $ r Y yw xY w)NFr   ©ÚREMOTE_CACHE_VERSIONz.pytorch/remote_cache:fx_graph_memcache_version)
ri   Úfx_graph_remote_cacheÚ	is_fbcoderP   Ú_utils_internalÚis_fb_unit_testÚtorch._inductor.fb.remote_cacherœ  ÚModuleNotFoundErrorÚjustknobs_getval_intr›  s    rW   Ú should_use_remote_fx_graph_cacher¤  A  s   € Ü×#Ñ#Ð/Ü×+Ñ+Ð+Ü×ÑÔØä×Ñ×,Ñ,Ô.ØðÝHð  ¤5×#8Ñ#8×#MÑ#MØ8ó$ñ ð øô ò Ùðús   ÁA> Á>	B
Â	B
c                ó0   — t        j                  dd| «      S )Nz[^a-zA-Z0-9_]rã   )rÙ   Úsubr¡  s    rW   Únormalize_namer§  T  s   € Ü�6‰6Ð" C¨Ó.Ð.r{   ztl.int1ztl.float8e4nvztl.float8e5ztl.float8e4b8ztl.float8e5b16ztl.uint8)ztl.boolztl.float8_e4m3fnztl.float8_e5m2ztl.float8_e4m3fnuzztl.float8_e5m2fnuzztl.float8_e8m0fnuztl.float4_e2m1fn_x2z^.*[.]c                ól   — t         j                  dt        | «      «      }t        j	                  ||«      S )z"Convert torch.dtype to triton typeútl.)Ú_triton_type_rer¦  r
  Ú_triton_type_mappingrH  )r©   Útriton_type_names     rW   Útriton_typer­  j  s.   € ä&×*Ñ*¨5´#°e³*Ó=ÐÜ×#Ñ#Ð$4Ð6FÓGÐGr{   c                ó®   — t         j                  | | «      }|j                  dd«      }t        t        |«      }t        |t        j                  «      sJ ‚|S )Nr©  rþ  )Ú_torch_triton_mappingrH  r­  rO   rP   r}   r©   )r©   Úadjusted_typeÚ	type_namerá  s       rW   Útriton_type_to_torchr²  p  sL   € Ü)×-Ñ-¨e°UÓ;€MØ×%Ñ% e¨RÓ0€IÜœ˜yÓ)€IÜ�i¤§¡Ô-Ð-Ð-ØÐr{   c                óØ  — | j                    xrÜ | j                  «       |j                  «       k(  xr¹ | j                  «       |j                  «       k(  xr– | j                  |j                  k(  xr{ | j                  |j                  k(  xr` | j                  «       j                  «       |j                  «       j                  «       k(  xr! | j                  «       |j                  «       k(  S r‹   )Ú	is_mkldnnrÈ  r4  r©   rª   Úuntyped_storager  rM  ©r&  r�   s     rW   Úis_same_tensorr·  x  sÅ   € à�N‰NÐò 	<Ø�I‰I‹K˜5Ÿ:™:›<Ñ'ò	<à�K‰K‹M˜UŸ\™\›^Ñ+ò	<ð �J‰J˜%Ÿ+™+Ñ%ò	<ð �K‰K˜5Ÿ<™<Ñ'ò		<ð
 × Ñ Ó"×+Ñ+Ó-°×1FÑ1FÓ1H×1QÑ1QÓ1SÑSò	<ð ×ÑÓ! U×%9Ñ%9Ó%;Ñ;ðr{   c                óv  — | j                   xr¬ | j                  «       |j                  «       k(  xr‰ | j                  |j                  k(  xrn | j                  |j                  k(  xrS t        j
                  j                  j                  | «      t        j
                  j                  j                  |«      k(  S r‹   )r´  rÈ  r©   rª   rP   r  Úmkldnnr  r¶  s     rW   Úis_same_mkldnn_tensorrº  „  s�   € à�‰ò 	PØ�I‰I‹K˜5Ÿ:™:›<Ñ'ò	Pà�J‰J˜%Ÿ+™+Ñ%ò	Pð �K‰K˜5Ÿ<™<Ñ'ò	Pô �I‰I×Ñ×%Ñ% dÓ+¬u¯y©y×/?Ñ/?×/HÑ/HÈÓ/OÑOðr{   c                  ó   — y)N)rq  ÚisnanÚlogical_notÚlogical_andÚsignbitÚand_ÚleÚltÚgeÚgtÚeqÚnerÿ  Úxorrš   rš   r{   rW   Úboolean_opsrÈ  Ž  s   € ðr{   c                  ó"   — e Zd ZU ded<   ded<   y)ÚOpDtypeRuler3   Útype_promotion_kindútorch.dtype | NoneÚoverride_return_dtypeNr*  rš   r{   rW   rÊ  rÊ  ¢  s   … à8Ó8Ø-Ô-r{   rÊ  zdict[str, OpDtypeRule]Úop_dtype_propagation_rulesc                ó*   — t        ||«      t        | <   y r‹   )rÊ  rÎ  )rÛ   rË  rÍ  s      rW   Ú#register_op_dtype_propagation_rulesrÐ  «  s   € ô
 (3ØÐ2ó(Ô˜tÒ$r{   zOrderedSet[str]Úop_requires_libdevice_fp64c                ó.   — t         j                  | «       y r‹   )rÑ  rK  r¡  s    rW   Ú#register_op_requires_libdevice_fp64rÓ  ¸  s   € Ü×"Ñ" 4Õ(r{   c                ó  — ddl m} | s$|j                  j                  «       j                  } | dk(  rt
        j                  S | dk(  ry| dk(  rt
        j                  S | dk(  rt
        j                  S t
        j                  S )Nr   rm  r  rI   rJ   Útpu)
r‹  rn  rq  Úget_current_device_or_throwr  ri   Úcpu_backendÚxpu_backendÚtpu_backendÚcuda_backend)rØ   rn  s     rW   Úget_current_backendrÛ  ¼  ss   € Ý-áØ—g‘g×9Ñ9Ó;×@Ñ@ˆØ�eÒÜ×!Ñ!Ð!Ø	˜Ò	ØØ	˜Ò	Ü×!Ñ!Ð!Ø	˜Ò	Ü×!Ñ!Ð!ä×"Ñ"Ð"r{   c                ó¸   — | t         j                  t         j                  fv r7t        j                  j
                  rt        «       dk(  rt         j                  S | S )z"Maybe upcast [b]float16 to float32rò  )rP   rÃ   rE  ri   rò  Úcodegen_upcast_to_fp32rÛ  rG  r$  s    rW   Úupcast_compute_typerÞ  Í  s@   € ð 	”%—-‘-¤§¡Ð0Ñ0Ü�M‰M×0Ò0ÜÓ! XÒ-ä�}‰}ÐØ€Lr{   ÚKeyTypeÚValTypec                  óZ   — e Zd ZdZdd„Zdd„Zdd„Zdd„Zddd„Zdd„Z	dd	„Z
dd
„Zdd„Zy)Ú
ScopedDictzè
    A dictionary-like object that allows for scoped updates. It maintains
    an original dictionary and a set of new items that can override
    the original items within the scope.  The original dictionary is
    unmodified.
    c                ó    — || _         i | _        y r‹   ©Úoriginal_dictÚ	new_items)rÜ  rå  s     rW   r5  zScopedDict.__init__ä  s   € Ø*ˆÔØ13ˆ�r{   c                óZ   — || j                   v r| j                   |   S | j                  |   S r‹   ©ræ  rå  ré  s     rW   r  zScopedDict.__getitem__è  s.   € Ø�$—.‘.Ñ Ø—>‘> #Ñ&Ð&Ø×!Ñ! #Ñ&Ð&r{   c                ó"   — || j                   |<   y r‹   )ræ  )rÜ  r^  r�   s      rW   Ú__setitem__zScopedDict.__setitem__í  s   € Ø#ˆ�‰�sÒr{   c                ó>   — || j                   v xs || j                  v S r‹   rè  ré  s     rW   Ú__contains__zScopedDict.__contains__ð  s!   € Ø�d—n‘nÐ$ÒA¨¨t×/AÑ/AÐ(AÐAr{   Nc                ót   — || j                   v r| j                   |   S | j                  j                  ||«      S r‹   )ræ  rå  rH  )rÜ  r^  rØ  s      rW   rH  zScopedDict.getó  s6   € Ø�$—.‘.Ñ Ø—>‘> #Ñ&Ð&Ø×!Ñ!×%Ñ% c¨7Ó3Ð3r{   c                óz   — t        | j                  «      }| j                  D ]  }|| j                  vsŒ|dz  }Œ |S rr  )rR   rå  ræ  )rÜ  r  ru  s      rW   r¤  zScopedDict.__len__ø  sC   € Ü�×"Ñ"Ó#ˆØ—‘ò 	ˆAØ˜×*Ñ*Ò*Ø�Q‘‘ð	ð ˆr{   c              #  ó€   K  — | j                   E d {  –—†  | j                  D ]  }|| j                   vsŒ|–— Œ y 7 Œ)­wr‹   rä  )rÜ  ru  s     rW   Ú__iter__zScopedDict.__iter__ÿ  s@   è ø€ Ø×%Ñ%×%Ð%Ø—‘ò 	ˆAØ˜×*Ñ*Ò*Ø“ñ	ð 	&ús   ‚>’<“!>µ>c                óH   — t        | j                  xs | j                  «      S r‹   )r1  rå  ræ  rD  s    rW   rO  zScopedDict.__bool__  s   € Ü�D×&Ñ&Ò8¨$¯.©.Ó9Ð9r{   c                ó   — t         ‚r‹   rš  ré  s     rW   Ú__delitem__zScopedDict.__delitem__  s   € Ü!Ð!r{   )rå  úMapping[KeyType, ValType])r^  rß  r’   rà  )r^  rß  r�   rà  r’   rÞ  )r^  r�  r’   r1  r‹   )r^  rß  rØ  úValType | Noner’   rõ  r§  )r’   zIterator[KeyType]r„  )r^  rß  r’   rÞ  )r“   r”   r•   r–   r5  r  rê  rì  rH  r¤  rð  rO  ró  rš   r{   rW   râ  râ  Ü  s5   „ ñó4ó'ó
$óBô4ó
óó:ô"r{   râ  )Úfrozen_defaultc              ó(   ‡— dˆfd„}| €|S  || «      S )Nc                ó4   •— t        j                  | d‰¬«      S )NT)Úkw_onlyr›   )ÚdataclassesÚ	dataclass)rŽ   r›   s    €rW   Úwrapzir_dataclass.<locals>.wrap  s   ø€ Ü×$Ñ$ S°$¸vÔFÐFr{   )rŽ   rl   r’   rl   rš   )rŽ   r›   rü  s    ` rW   Úir_dataclassrý    s   ø€ õGð €{ØˆÙ�‹9Ðr{   c                 óœ   — t         j                  j                  j                  «       } | �"| j                  r| j                  j
                  S y r‹   )rP   rU  rV  rW  Úfw_metadataÚbw_donated_idxs)rZ  s    rW   Úget_donated_idxsr    s=   € Ü—m‘m×2Ñ2×:Ñ:Ó<€OØÐ" ×'BÒ'BØ×*Ñ*×:Ñ:Ð:Ør{   c                  ó    — e Zd ZdZdZdZdZdZy)ÚTritonAttrsDescriptorVersionr   r7   ry  rÒ  r0  N)r“   r”   r•   ÚV0_NO_TRITONÚV1_COMPILERÚV2_BACKENDSÚV3_BACKENDS_TUPLEÚV4_DICTrš   r{   rW   r  r    s    „ Ø€LØ€KØ€Kà	ð ð �Gr{   r  c                 óP  — t         j                  j                  d«      €t        j                  S dd l} dd l} t        | j                  j                  d«      rt        j                  S t        | j                  j                  d«      rt        j                  S t        j                  S )Nrò  r   ÚAttrsDescriptor)rF  rG  rH  r  r  Útriton.backends.compilerÚtriton.compiler.compilerr  r‚  Úcompilerr  r  r  )rò  s    rW   Ú#get_triton_attrs_descriptor_versionr  '  s{   € ä‡~�~×Ñ Ó)Ð1Ü+×8Ñ8Ð8ã#Û#äˆv�‰×'Ñ'Ð):Ô;ô ,×7Ñ7Ð7Ü	�—‘×)Ñ)Ð+<Ô	=ä+×7Ñ7Ð7ô ,×3Ñ3Ð3r{   c                 ó8   — t        «       t        j                  k(  S r‹   )r  r  r  rš   r{   rW   Útriton_version_uses_attrs_dictr  A  s   € Ü.Ó0Ô4P×4XÑ4XÑXÐXr{   c                ó˜   — | j                  «       }t        | t        j                  j                  «      r|› d| j
                  › �n|}||fS )Nrp   )rÛ   r}   rP   r�  rŽ  Ú_overloadname)r€  Úop_overload_packet_nameÚop_overload_names      rW   Úget_op_namesr  E  sR   € Ø#%§7¡7£9Ðô �bœ%Ÿ*™*×/Ñ/Ô0ð #Ð
# 1 R×%5Ñ%5Ð$6Ñ7à$ð ð
 #Ð$4Ð4Ð4r{   c                ój  — ddl m} | j                  }t        |t        j
                  j                  «      sy|t        j                  j                  j                  j                  t        j                  j                  j                  j                  t        j                  j                  j                  j                  fv rq ||| j                  | j                  d¬«      }|�O|\  }}|d   }|D ]@  }|€Œ|j                  d   j                   t        j"                  t        j$                  fv sŒ@ y y)a  
    Check if an FX node is cudagraph-unsafe based on its input arguments.

    Some ops are only cudagraph-unsafe depending on their inputs (e.g., index_put
    with boolean indices triggers .nonzero() during capture, but integer indices
    are safe).
    r   )Únormalize_functionFT)Únormalize_to_only_use_kwargsÚindicesrj  )Útorch.fx.operator_schemasr  rŒ  r}   rP   r�  rŽ  r  r3  Ú	index_putrØ  Ú
index_put_Ú_unsafe_index_putr„   r¤  r  r©   r1  rV  )r�  r  rŒ  Ú
normalizedrã   r¤  r  rU  s           rW   Ú,_fx_node_is_input_dependent_cudagraph_unsafer  O  sú   € õ =à�^‰^€FÜ�fœeŸj™j×3Ñ3Ô4Øð Ü�	‰	�‰× Ñ ×(Ñ(Ü�	‰	�‰×!Ñ!×)Ñ)Ü�	‰	�‰×(Ñ(×0Ñ0ðñ ñ
 (Ø�G—L‘L '§.¡.Ètô
ˆ
ð Ð!Ø"‰IˆAˆvØ˜YÑ'ˆGØò  �Ø‘? s§x¡x°¡×'<Ñ'<Ü—J‘JÜ—K‘KðAò (ñ  ð ð r{   c                óÒ  — | j                   }t        |«      t        v ryt        |t        j
                  j                  «      r1t        j                  j                  j                  |j                  v ryt        | «      ry| j                  j                  d«      x}�Kt        |t        t        f«      s|gn|}|D ]+  }t        |t        j                   «      sŒ|j"                  sŒ+ y y)a   
    Check if an FX node is cudagraph-unsafe.

    This includes:
    - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.)
    - Ops with the cudagraph_unsafe tag
    - Input-dependent unsafe ops (e.g., index_put with boolean indices)
    - Ops with sparse tensor outputs
    Trj  F)rŒ  r
  ÚFORBIDDEN_CUDAGRAPH_OPSr}   rP   r�  rŽ  r  r‘  Úcudagraph_unsafer“  r  r  rH  r]  r¿  rž  Ú	is_sparse)r�  rŒ  rj  Úvalsr‡   s        rW   r²  r²  s  s¼   € ð �^‰^€Fô ˆ6ƒ{Ô-Ñ-Øô 	�6œ5Ÿ:™:×0Ñ0Ô1Ü�H‰H�L‰L×)Ñ)¨V¯[©[Ñ8àô 4°GÔ<Øð �|‰|×Ñ Ó&Ð&ˆÐ3Ü& s¬T´5¨MÔ:�‰uÀˆØò 	ˆAÜ˜!œUŸ\™\Õ*¨q¯{«{Ùð	ð r{   c                óÒ   — ddl m} t        | |j                  |j                  f«      ryt        | |j
                  |j                  f«      syt        | dd«      }|�t        |«      ryy)ah  
    Returns True if the node is an op that is not cudagraphable.
    This includes:
    - Ops in FORBIDDEN_CUDAGRAPH_OPS (CPU sync, dynamic alloc, etc.)
    - Ops with the cudagraph_unsafe tag
    - index_put_ with boolean indices (triggers .nonzero() during capture)
    - Control flow nodes (Conditional, WhileLoop)
    - Ops with sparse tensor outputs
    r7   rü  TFr�  N)	rþ  rý  r}   ÚConditionalÚ	WhileLooprÖ  r?   rO   r²  )rg  rý  r�  s      rW   Úis_cudagraph_unsafe_opr(  ˜  sa   € õ ô �$˜Ÿ™¨¯©Ð6Ô7Øä�d˜R×.Ñ.°·±Ð@ÔAØä�d˜I tÓ,€GØÐÔ:¸7ÔCØàr{   c                 ó  — t         j                  j                  dd«      } t        j                  «       rUddlm}  |«       }|rFt         j                  j                  |dd«      }| r!t         j                  j                  || g«      n|} | S )NÚLD_LIBRARY_PATHrþ  r   )Úget_runtime_pathrÍ  Úlib)
rê  rë  rH  ri   rž  Úlibfb.py.parutilr+  r÷  r  Úpathsep)r÷  r+  Úruntime_pathÚlib_paths       rW   Úget_ld_library_pathr1  ²  sg   € Ü�:‰:�>‰>Ð+¨RÓ0€DÜ×ÑÔÝ5á'Ó)ˆÙÜ—w‘w—|‘| L°)¸UÓCˆHÙ8<”2—:‘:—?‘? H¨dÐ#3Ô4À(ˆDà€Kr{   c                óF   — ddl m} t        | |«      xr | j                  d uS )Nr   )ÚSubgraphPythonWrapperCodegen)Útorch._inductor.codegen.wrapperr3  r}   Úpartition_signatures)rð  r3  s     rW   Ú#is_codegen_graph_partition_subgraphr6  ¿  s*   € ÝLô 	�7Ð8Ó9ò 	5Ø×(Ñ(°Ð4ðr{   c                 óÒ   — t         j                  j                  j                  j                  xs t
        j                  d uxr$ t         j                  j                  j                  S r‹   )rP   rÌ  ri   rò  Ú
cudagraphsÚ&_unstable_customized_partition_wrapperrð  Úgraph_partitionrš   r{   rW   Úis_using_cudagraph_partitionr;  È  sN   € ä�‰×Ñ×%Ñ%×0Ñ0ò 	FÜ1×9Ñ9ÀÐEò1ô �/‰/×
 Ñ
 ×
0Ñ
0ð1r{   c                óæ   — ddl m} |j                  j                  j	                  | d«      r6|j                  j                  j                  | d«      rt        j                  S t        j                  S )Nr7   rm  l        i   €)	rp  rn  rq  rr  Ústatically_known_ltr  rP   rT  rU  )rÈ  rn  s     rW   Údtype_from_sizer>  Ï  sP   € Ýà‡w�w×Ñ×+Ñ+Øˆeôà
�'‰'×
Ñ
×
/Ñ
/°°hÔ
?Ü�{‰{Ðä�{‰{Ðr{   )r  rJ   c                óh   — | dk(  r(t         j                  j                  j                  «       S d| v ryy)z;
    Returns True if the device supports MKL-DNN BF16.
    r  rJ   TF)rP   r  r¹  Ú_is_mkldnn_bf16_supported©rØ   s    rW   Úis_mkldnn_bf16_supportedrB  Ý  ó3   € ð �eÒÜ�y‰y×Ñ×9Ñ9Ó;Ð;Ø	�+Ñ	àØr{   c                óh   — | dk(  r(t         j                  j                  j                  «       S d| v ryy)z;
    Returns True if the device supports MKL-DNN FP16.
    r  rJ   TF)rP   r  r¹  Ú_is_mkldnn_fp16_supportedrA  s    rW   Úis_mkldnn_fp16_supportedrF  é  rC  r{   c           
     ón  — |D �cg c]  }t        t        |«      «      ‘Œ }}| D ]R  }t        |«      t        |«      k(  sJ ‚t        |«      D ])  \  }}t        ||   t        t        |«      «      «      ||<   Œ+ ŒT g }|j	                  dj                  d„ t        ||«      D «       «      «       t        |«      t        |«      dz  z   t        |«      dz
  z   }|j	                  d|z  «       | D ]3  }|j	                  dj                  d„ t        ||«      D «       «      «       Œ5 dj                  |«      S c c}w )Nú|c              3  ó6   K  — | ]  \  }}d ||› d�›d �–— Œ y­w©r!  rþ  Nrš   )r»   Úhr_  s      rW   r½   ztabulate_2d.<locals>.<genexpr>ü  s$   è ø€ ÒH©4¨1¨a˜A˜a   0˜t˜W Aœ,ÑHùó   ‚ry  r7   r€  c              3  ó6   K  — | ]  \  }}d ||› d�›d �–— Œ y­wrJ  rš   )r»   rê   r_  s      rW   r½   ztabulate_2d.<locals>.<genexpr>  s$   è ø€ ÒH©t¨q°!  ! Q C p 4 ¨œlÑHùrL  rB  )rR   r
  rE  rÉ   r™  r  rÐ   rü   )ÚelementsÚheadersrê   ÚwidthsÚrowrè   r[  Útotal_widths           rW   Útabulate_2drS  õ  s  € Ø#*Ö+˜aŒc”#�a“&�kÐ+€FÐ+Øò 4ˆÜ�3‹xœ3˜w›<Ò'Ð'Ð'Ü˜c“Nò 	4‰DˆAˆqÜ˜F 1™I¤s¬3¨q«6£{Ó3ˆF�1ŠIñ	4ð4ð €EØ	‡L�L�—‘ÑH´3°wÀÓ3GÔHÓHÔIä�f“+¤ V£¨q¡Ñ1´S¸³[À1±_ÑE€KØ	‡L�L��{Ñ"Ô#Øò JˆØ�‰�S—X‘XÑH´s¸3ÀÓ7GÔHÓHÕIðJà�9‰9�UÓÐùò ,s   …D2c              #  óâ   K  — t        | j                  «       «      t        |j                  «       «      z  }|D ]3  }| j                  |«      }|j                  |«      }||�|n||�|n|f–— Œ5 y­w)að  
    Zip two dictionaries together, replacing missing keys with default values.

    Args:
        dict1 (dict): The first dictionary.
        dict2 (dict): The second dictionary.
        d1_default (Any): the default value for the first dictionary
        d2_default (Any): the default value for the second dictionary

    Yields:
        tuple: A tuple containing the key, the value from dict1 (or d1_default if missing),
               and the value from dict2 (or d2_default if missing).
    N)r!   rI  rH  )Údict1Údict2Ú
d1_defaultÚ
d2_defaultÚall_keysr^  Úvalue1Úvalue2s           rW   Ú	zip_dictsr\    sv   è ø€ ô( ˜%Ÿ*™*›,Ó'¬*°U·Z±Z³\Ó*BÑB€Hð ò 	
ˆà—‘˜3“ˆØ—‘˜3“ˆð ØÐ(‰F¨jØÐ(‰F¨jð
ó 	
ñ	
ùs   ‚A-A/c                óT  — 	 	 	 	 	 	 	 	 dd„}	 	 	 	 	 	 	 	 dd„}| j                  dt        j                  j                  «      }| j	                  «       } |rm || dd«        || dd«        || dt
        j                  j                   «        || dd	«        || d
t        j                  j                  «        || dd«       | j                  dt        j                  j                  «      }| j                  dt        j                  j                  «      }|dk(  r|rt        d«      ‚| S )a6  
    Ensures the configuration is internally consistent for standalone AOTInductor.

    If `aot_inductor_mode.compile_standalone` is set to True in the provided
    `config_patches` (or falls back to the global config), this function ensures
    that the following configs are also enabled:
        - `aot_inductor.package_cpp_only`

    Args:
        config_patches (dict[str, Any]): A dictionary of user-provided config
            overrides for AOTInductor compilation.

    Returns:
        dict[str, Any]: The possibly-updated `config_patches` dictionary.
    c                óˆ   — | j                  |t        t        |«      «      }|€|| |<   y |s||k7  rt        d|› d|› d�«      ‚y y )NzInvalid config: ú=z3 when aot_inductor_mode.compile_standalone is True.)rH  rO   ri   rú   ©Úconfig_patchesÚconfig_nameÚconfig_valuer�   s       rW   Úpatch_configz2maybe_aoti_standalone_config.<locals>.patch_config9  s]   € ð ×"Ñ" ;´¼ÀÓ0LÓMˆØˆ=Ø*6ˆN˜;Ò'Ù˜5 LÒ0ÜØ" ; -¨q°°Ð>qÐróð ð 1�r{   c                ó†   — | j                  |t        t        |«      «      }||k7  rt        j	                  d||«       || |<   y )NzDOverriding: %s=%s when aot_inductor_mode.compile_standalone is True.)rH  rO   ri   rÓ   rö  r`  s       rW   Úforce_patch_configz8maybe_aoti_standalone_config.<locals>.force_patch_configD  sF   € ð ×"Ñ" ;´¼ÀÓ0LÓMˆØ�LÒ Ü�K‰KØVØØôð
 '3ˆ�{Ò#r{   z$aot_inductor_mode.compile_standalonezaot_inductor.package_cpp_onlyTz aot_inductor.embed_kernel_binaryz#aot_inductor.emit_multi_arch_kernelz+aot_inductor.model_name_for_generated_filesÚ
aoti_modelzaot_inductor.link_libtorchzaot_inductor.dynamic_linkageFz"aot_inductor.cross_target_platformz$aot_inductor.package_constants_in_soÚwindowszªconfig.aot_inductor.package_constants_in_so is not supported for windows cross-compilation. Please use config.aot_inductor.package_constants_on_disk_format = binary_blob.)ra  údict[str, Any]rb  r
  rc  r   r’   rÞ  )rH  ri   Úaot_inductor_modeÚcompile_standaloneÚcopyrP   rN  rq   Útest_configsÚuse_libtorchÚaot_inductorÚcross_target_platformÚpackage_constants_in_sorú   )ra  rd  rf  rk  rp  rq  s         rW   Úmaybe_aoti_standalone_configrr  (  sk  € ð"	Ø&ð	Ø58ð	ØHKð	à	ó	ð
3Ø&ð
3Ø58ð
3ØHKð
3à	ó
3ð (×+Ñ+Ø.Ü× Ñ ×3Ñ3óÐð
 $×(Ñ(Ó*€NÙá�^Ð%DÀdÔKá�^Ð%GÈÔNáØÐAÄuÇ}Á}×GXÑGXÐCXô	
ñ 	ØÐIÈ<ô	
ñ 	ØØ(Ü×Ñ×,Ñ,ô	
ñ
 	˜>Ð+IÈ5ÔQà*×.Ñ.Ø,Ü×Ñ×1Ñ1óÐð
 -×0Ñ0Ø.Ü×Ñ×3Ñ3óÐð
  	Ò)Ñ.EÜð]ó
ð 	
ð
 Ðr{   c                ó¦  — t         j                  j                  r(t         j                  j                  dk(  rt	        d«      ‚t         j                  j                  r0t         j                  j
                  dk(  rt	        d«      ‚d}d}||fS t         j                  j                  dk(  rd}d}||fS | dk  ryd}t        j                  «        }||fS )	a  
    Decide whether we should mmap weights, and whether to store the weights with .so.

    If force_mmap_weights or package_constants_on_disk_format == "binary_blob" configs are set, respect the config.

    Returns tuple (use_external_weights, use_mmap_weights).
    Úbinary_blobz‚config.aot_inductor.package_constants_on_disk_format = binary_blob and config.aot_inductor.force_mmap_weights cannot both be True.rh  zKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi ”5w)FF)ri   ro  Úforce_mmap_weightsÚ package_constants_on_disk_formatrú   rp  rž  )Úconsts_sizeÚuse_mmap_weightsÚuse_external_weightss      rW   Údetermine_aoti_mmap_flagsrz  }  sä   € ô 	×Ñ×.Ò.Ü×Ñ×@Ñ@ÀMÒQäðJó
ð 	
ô
 ×Ñ×-Ò-Ü×Ñ×4Ñ4¸	ÒAÜØ]óð ð  ÐØ$ÐØ#Ð%5Ð5Ð5ä×Ñ×;Ñ;¸}ÒLØ#ÐØ ÐØ#Ð%5Ð5Ð5à�mÒ#Øà ÐÜ!×+Ñ+Ó-Ð-ÐàÐ!1Ð1Ð1r{   c                 óÆ   — ddl m}  | j                  j                  }|€yt	        |t
        «      st        d«      ‚|dk(  ryt        j                  d|«      st        d«      ‚y)zL
    Validates if a model name is suitable for use in code generation.

    r   rh   Tz4Invalid AOTI model name: Model name must be a stringrþ  z^[a-zA-Z_][a-zA-Z0-9_]*$zVInvalid AOTI model name: Model name can only contain letters, numbers, and underscores)	rC  ri   ro  Úmodel_name_for_generated_filesr}   r
  rZ  rÙ   rÚ   )ri   Ú
model_names     rW   Úis_valid_aoti_model_namer~  ¦  sh   € õ
 'à×$Ñ$×CÑC€JàÐØä�j¤#Ô&ÜÐOÓPÐPà�RÒØô �8‰8Ð/°Ô<ÜØdó
ð 	
ð r{   c                ó2   — |rt        | «      S t        | «      S r‹   )r)   r(   )rT   Úunbacked_onlys     rW   Úget_free_symbolsr�  Á  s   € ÙÜ$ QÓ'Ð'ä˜A‹Ðr{   c                 ó  — i t         j                  ¥dt         j                  j                  dt         j                  j	                  t
        j                  «      «      i¥} t        j                  «       rt        j                  d«      | d<   | S )zA
    Get a base environment for running Python subprocesses.
    Ú
PYTHONPATHÚTORCH_CUSTOM_PYTHONPATHr&  Ú
PYTHONHOME)rê  rë  rH  r.  r  rÈ  r÷  ri   rž  Ú	sysconfigÚget_path)Úenvs    rW   Úpython_subprocess_envr‰  È  sl   € ð
ä
�*‰*ðð 	”b—j‘j—n‘nØ%¤r§z¡z§¡´s·x±xÓ'@ó
ñ	€Cô  ×ÑÔÜ%×.Ñ.¨vÓ6ˆˆLÑà€Jr{   c                  ó&   — e Zd ZU dZded<   ded<   y)ÚCUDAGraphWrapperMetadataz˜
    Metadata for Customized CUDAGraphWrapper.

    Currently assumes there is 1 dynamo graph and will extend to
    multiple graphs in the future.
    rŒ   Únum_partitionsÚpartition_indexNr£   rš   r{   rW   r‹  r‹  ã  s   … ñð Óð Ôr{   r‹  .c                  ó   — e Zd ZU dZded<   y)ÚCUDAGraphWrapperNzCUDAGraphWrapperType | Nonerð  )r“   r”   r•   rð  r¤   rš   r{   rW   r�  r�  ú  s   … Ø+/€GÐ(Ô/r{   r�  c                ó   — | t         _        y r‹   )r9  rð  )rð  s    rW   Ú!set_customized_partition_wrappersr‘    s   € Ø5<Ô*Õ2r{   c                ó8  ‡— | j                   j                  }| j                   j                  g |¢| j                   j                  ¢| j                   j                  «      }| j                   j                  }t        j                  ||f«      \  }}dd„}|D �cg c]7  } ||«      r+t        j                  j                  j                  |d¬«      n|‘Œ9 }}dd„Šdˆfd„}|D �cg c]
  } ||«      ‘Œ }}t        j                  ||«      \  }}||fS c c}w c c}w )	Nc                ó  — t        | t        j                  j                  j                  «      xrS t        | t        j                  j                  j
                  t        j                  j                  j                  f«       S r‹   )r}   rP   rÌ  rý  r@   ÚGeneratorStateÚOpaqueObjectStaterí  s    rW   Ú_is_tensor_irz(snode_args_kwargs.<locals>._is_tensor_ir  s^   € Ü˜!œUŸ_™_×/Ñ/×6Ñ6Ó7ò 
Ä
ØÜ�_‰_×Ñ×.Ñ.´·±×0BÑ0B×0TÑ0TÐUóA
ð =
ð 	
r{   T)Úreplace_symbols_with_hintsc                ó2   — t        j                  | ||¬«      S )Nr¨   )rP   rÂ   )rÈ  r©   rª   s      rW   Ú_tensorz"snode_args_kwargs.<locals>._tensor&  s   € Ü�{‰{˜4 u°VÔ<Ð<r{   c                ó˜   •— t        | t        j                  «      s| S  ‰| j                  «       | j                  | j
                  «      }|S r‹   )r}   rP   rž  rÈ  r©   rª   )rê   rM  r™  s     €rW   Úto_real_tensorz)snode_args_kwargs.<locals>.to_real_tensor)  s7   ø€ Ü˜!œUŸ\™\Ô*ØˆHÙ�a—f‘f“h §¡¨¯©Ó2ˆØˆ
r{   r„  )r’   rŸ  )rê   r   r’   r   )rg  rj  Úfill_non_provided_argsÚconstant_argsr¤  Úpytreer"   rP   rÌ  rý  Úir_node_to_tensorÚtree_unflatten)	rä  r„   r¤  Ú	flat_argsÚflat_args_pytree_specr–  r   r›  r™  s	           @rW   Úsnode_args_kwargsr£    s  ø€ Ø�:‰:×Ñ€DØ�:‰:×,Ñ,Ø*ˆ$Ð*�—‘×)Ñ)Ð*Ø�
‰
×Ñó€Dð �Z‰Z×Ñ€FÜ'-×':Ñ':¸DÀ&¸>Ó'JÑ$€IÐ$ó
ð ö	ð ñ ˜Ôô 	�‰×Ñ×,Ñ,¨QÈ4Ð,ÔPàñ	ð€Ið ó=õð -6Ö6 q‘ Õ"Ð6€IÐ6Ü×(Ñ(¨Ð4IÓJ�L€Dˆ&Ø�ˆ<Ðùò%ùò  7s   Â<DÃ$Dc                óÄ   — ddl m} | j                  }|j                  j                  r(|j	                  |j                  j                  dz   «      }|j                  d«      S )Nr7   rm  rã   )r  r˜  Úfwd_rng_stater  r  )rp  rn  rÛ   rq  ÚremoveprefixrÊ  )rù  rn  Údep_names      rW   Úis_nonfreeable_buffersr¨  4  sN   € Ýà�x‰x€Hð 	‡w�w‡|‚|Ø×(Ñ(¨¯©¯©¸Ñ);Ó<ˆØ×ÑØIóð r{   c                óp   — t        || › d�z  «      5 }|j                  «       cddd«       S # 1 sw Y   yxY w)z,Load a template file and return its content.z	.py.jinjaN)ÚopenÚread)rÛ   Útemplate_dirr
  s      rW   Úload_templater­  B  s6   € ä	ˆl ˜v YÐ/Ñ/Ó	0ð °AØ�v‰v‹x÷÷ ò ús   ’,¬5c                óv  — | j                   }t        |t        j                  j                  t        j                  j
                  f«      sJ dt        |«      › �«       ‚t        j                  syt        t        j                  j                  j                  j                  t        j                  j                  j                  j                  g«      }||v ryt        t        j                  j                  j                   g«      }t        |t        j                  j
                  «      r||v S t#        | «       S )zLDecide whether fallback for a node. This is only used in inductor lite mode.z6Expected OpOverload or HigherOrderOperator, but found F)rŒ  r}   rP   r�  rŽ  r  r  ri   Úfallback_by_defaultr!   r  r3  Ú_assert_scalarrØ  Úlift_fresh_copyÚhigher_orderÚ triton_kernel_wrapper_functionalr   )rg  rŒ  Ú"skip_fallback_due_to_dynamic_shapeÚfallback_hopss       rW   Úshould_fallback_by_defaultr¶  H  sý   € à�[‰[€FäØ”—‘×&Ñ&¬¯
©
×(FÑ(FÐGôð Oà	?ÄÀVÃ¸~ÐNóOð ô ×%Ò%Øô *4ä�I‰I�N‰N×)Ñ)×1Ñ1Ü�I‰I�N‰N×*Ñ*×2Ñ2ð	
ó*Ð&ð Ð3Ñ3Øô Ü	�‰×	Ñ	×	@Ñ	@ÐAó€Mô �&œ%Ÿ*™*×8Ñ8Ô9Ø˜Ð&Ð&ä& tÓ,Ð,Ð,r{   )z-torch.ops._c10d_functional.all_reduce.defaultz.torch.ops._c10d_functional.all_reduce_.defaultz9torch.ops._c10d_functional.all_gather_into_tensor.defaultz8torch.ops._c10d_functional.reduce_scatter_tensor.defaultz4torch.ops._c10d_functional.all_to_all_single.defaultz6torch.ops._c10d_functional_autograd.all_reduce.defaultzBtorch.ops._c10d_functional_autograd.all_gather_into_tensor.defaultzAtorch.ops._c10d_functional_autograd.reduce_scatter_tensor.defaultz=torch.ops._c10d_functional_autograd.all_to_all_single.defaultz(torch.ops._c10d_functional.isend.defaultz(torch.ops._c10d_functional.irecv.defaultz0torch.ops._c10d_functional.batch_p2p_ops.defaultc                ó   — | t         v S )z0Check if an operation is a collective operation.)ÚCOLLECTIVE_OPS)r‡  s    rW   Úis_collective_opr¹  €  s   € à”nÐ$Ð$r{   c                 ób   — t        j                  «       r		 ddlm}  | S g S # t        $ r g cY S w xY w)Nr   ©Útlx_only_cuda_options)ri   rž  Ú)torch._inductor.fb.tlx_templates.registryr¼  r	  r»  s    rW   r¼  r¼  …  s9   € ä×ÑÔð	ÝWà(Ð(ð ˆ	øô	 ò 	ØŠIð	ús   –   .­.c                ó   — | |z   dz
  |z  |z  S )z(Round x up to the nearest multiple of y.r7   rš   )rT   Úys     rW   Ú	_round_uprÀ  ”  s   € à�‰U�Q‰Y˜1Ñ Ñ!Ð!r{   c                ó  — ddl m}m}  ||d«      r|j                  |j                  fS t        |«      dk\  �r  ||d   | d   «      r ||d   d«      s ||d   d«      r' ||d   | d   «      r|j                  |j                  fS  ||d   | d   «      r ||d   t        | d   d«      «      s( ||d   | d   «      r1 ||d   t        | d   d«      «      r|j                  |j                  fS  ||d   t        | d   d«      «      r1 ||d   t        | d   d«      «      r|j                  |j                  fS |t        j                  k(  rdnd}|t        j                  k(  r•|t        j                  k(  r‚t        | d   d«      t        t        || d   z  d«      d«      z  }	t        | d   d«      t        t        || d   z  d«      d«      z  }
 |||	«      s	 |||
«      r|j                  |j                  fS |t        j                   k(  r÷t        j"                  j$                  sƒt        | d   d«      t        t        || d   z  d«      d«      z  }	t        | d   d«      t        t        || d   z  d«      d«      z  }
 |||	«      s	 |||
«      r|j&                  |j                  fS y	t        | d   d«      |z  | d   z  }	t        || d   z  d«      | d   z  }
 |||	«      s	 |||
«      r|j&                  |j                  fS y	)
z:
    Core implementation for scale/swizzle inference.
    r   )r6   ÚSwizzleTyper7   ry  rs   rr   r0  r  ©NN)Útorch.nn.functionalr6   rÂ  Ú
TensorWiseÚ
NO_SWIZZLErR   ÚRowWiserj   ÚBlockWise1x128ÚBlockWise128x128rP   rC  r;  rÀ  ÚBlockWise1x16ÚSWIZZLE_32_4_4rA  rN  rq   ÚBlockWise1x32)Úmat_sizeÚ
scale_sizeÚscale_numelÚ	mat_dtypeÚscale_dtypeÚeq_fnr6   rÂ  ÚK_multiplierÚexpected_numel_aÚexpected_numel_bs              rW   Ú_infer_scale_swizzle_implrÖ  ™  sD  € ÷ =ñ ˆ[˜!ÔØ×%Ñ% {×'=Ñ'=Ð=Ð=ô ˆ:ƒ˜!ÓÙ�*˜Q‘- ¨!¡Ô-±%¸
À1¹ÀqÔ2IÙ�*˜Q‘- Ô#©¨j¸©m¸XÀa¹[Ô(Ià×&Ñ&¨×(>Ñ(>Ð>Ð>ñ �*˜Q‘- ¨!¡Ô-Ù�j ‘m¤W¨X°a©[¸#Ó%>Ô?á�*˜Q‘- ¨!¡Ô-Ù�j ‘m¤W¨X°a©[¸#Ó%>Ô?à×-Ñ-¨{×/EÑ/EÐEÐEñ �˜A‘¤¨°©°SÓ 9Ô:¹uØ�q‰Mœ7 8¨A¡;°Ó4ô@
ð ×/Ñ/°×1GÑ1GÐGÐGð "¤U×%;Ñ%;Ò;‘1À€Lð ”E×*Ñ*Ò*¨{¼e×>QÑ>QÒ/QÜ$ X¨a¡[°#Ó6¼Ü�L 8¨A¡;Ñ.°Ó3°Qó:
ñ 
Ðô % X¨a¡[°#Ó6¼Ü�L 8¨A¡;Ñ.°Ó3°Qó:
ñ 
Ðñ �Ð.Ô/±5¸ÐFVÔ3WØ×,Ñ,¨k×.HÑ.HÐHÐHð ”e×*Ñ*Ò*Ü�}‰}× Ò ä(¨°!©°cÓ:¼YÜ˜ x°¡{Ñ2°BÓ7¸ó>ñ  Ðô  )¨°!©°cÓ:¼YÜ˜ x°¡{Ñ2°BÓ7¸ó>ñ  Ðñ �[Ð"2Ô3±uØÐ-ô8ð #×0Ñ0°+×2LÑ2LÐLÐLð ô  ' x°¡{°BÓ7¸,ÑFÈÐRSÉÑTÐÜ& |°h¸q±kÑ'AÀ2ÓFÈÐRSÉÑTÐÙ�[Ð"2Ô3±uØÐ-ô8ð #×0Ñ0°+×2HÑ2HÐHÐHàr{   c                óÈ   — t        | j                  d   | j                  d   ft        |j                  «      |j                  «       | j                  |j                  d„ ¬«      S )a†  
    Infer the scaling type and swizzle mode from matrix and scale tensor shapes/dtypes.

    This function determines how scale factors are laid out relative to the matrix:
    - TensorWise: Single scale for entire tensor
    - RowWise: One scale per row
    - BlockWise1x128/128x128: Block-scaled with float32 scales
    - BlockWise1x32: MXFP8 with float8_e8m0fnu scales (swizzled on NVIDIA)
    - BlockWise1x16: NVFP4 with float8_e4m3fn scales (swizzled)

    Args:
        mat: The matrix tensor (FP8 or FP4)
        scale: The scale factor tensor

    Returns:
        Tuple of (ScalingType, SwizzleType) or (None, None) if unrecognized
    r   r7   c                ó   — | |k(  S r‹   rš   rÅ  s     rW   r„  z%infer_scale_swizzle.<locals>.<lambda>  s
   € ˜1 ™6€ r{   ©rÍ  rÎ  rÏ  rÐ  rÑ  rÒ  )rÖ  r0  r¿  Únumelr©   )ÚmatÚscales     rW   Úinfer_scale_swizzlerÝ  é  sO   € ô( %Ø—)‘)˜A‘, §	¡	¨!¡Ð-Ü˜Ÿ™Ó%Ø—K‘K“MØ—)‘)Ø—K‘KÙ!ôð r{   c                ó\  ‡— ddl mŠ | j                  «       }|j                  «       }|r
|d   |d   f}|r%t        j                  t
        j                  |d«      nd}dˆfd„}t        t        |«      dk\  r
|d   |d   fn|d   dft        |«      || j                  |j                  |¬«      S )zí
    Infer the scaling type and swizzle mode for IR nodes (used during graph lowering).

    This is the IR-compatible version of infer_scale_swizzle, using symbolic
    size comparisons via V.graph.sizevars.statically_known_equals.
    r   rm  r7   c                óP   •— ‰j                   j                  j                  | |«      S )z5Compare values using symbolic equality when possible.)rq  rr  r  )r   r!  rn  s     €rW   Úsymbolic_eqz+infer_scale_swizzle_ir.<locals>.symbolic_eq  s    ø€ à�w‰w×Ñ×7Ñ7¸¸1Ó=Ð=r{   ry  rÙ  )r   r   r!  r   r’   r1  )r‹  rn  r  r  r  r  r  rÖ  rR   r¿  r©   )rÛ  rÜ  Ú	transposerÍ  rÎ  rÏ  rà  rn  s          @rW   Úinfer_scale_swizzle_irrâ    s¯   ø€ õ .à�|‰|‹~€HØ—‘Ó!€Jñ Ø˜Q‘K ¨!¡Ð-ˆñ DN”)×"Ñ"¤8§<¡<°¸QÔ?ÐST€Kõ>ô %Ü/2°8«}ÀÒ/A�(˜1‘+˜x¨™{Ñ+ÈÐQRÉÐUVÐGWÜ˜Ó$ØØ—)‘)Ø—K‘KØôð r{   r‚  )ry   rŒ   r’   rŒ   )r‡   r‘   r’   r1  )é   éd   )rÝ   úCallable[[], Any]rÞ   rŒ   rß   rŒ   r’   rn  )rã  rä  F)
rÝ   rå  rÞ   rŒ   rß   rŒ   ró   r1  r’   rn  r„  )rª   ztorch.device | None | strr’   útorch.device)r  zIterable[sympy.Expr]r’   r‘   )r#  r!  r$  r!  r’   r‘   )r  zIterable[_T]r’   zValuesView[_T])r.  r  r/  r  r’   r  )r^  rÌ  r’   r
  )rd  zIterable[int | torch.SymInt]r’   zlist[sympy.Expr])rj  úint | torch.SymIntr’   r  )rè   r  r’   rç  )rd  zIterable[int | sympy.Expr]r’   zlist[int | torch.SymInt])r€  útorch._ops.OpOverloadr’   r1  )r”  r5   r‹  z'Callable[[torch._ops.OpOverload], bool]r’   r1  )rŒ  r   r„   r•  r¤  ri  r’   z&tuple[GraphModule, list[torch.Tensor]])rH   )rª   r
  r’   rÞ  )r7   rH   )
r­  úCallable[..., Any]r®  úSequence[Any]rë   rŒ   rª   r
  r’   rn  )rš   r±  r±  g      ð?rH   )r­  ré  r®  rê  rë   rŒ   r¶  rŒ   r·  rn  rª   r
  r’   rn  )r¾  r   r¿  r
  r’   rÞ  )r¾  r   rÂ  r¡   r’   rÞ  )r   rŒ   r!  rŒ   r’   rŒ   )rT   zint | Sequence[int]rÈ  rŒ   r’   úSequence[int])rT   ztuple[_T, ...]r’   zlist[_T])rÝ   z!Callable[Concatenate[Any, P], RV]r’   zCachedMethod[P, RV])rø  r
  r’   z*Callable[[FN_TYPE[P, RV]], FN_TYPE[P, RV]])r  ú*Sequence[BaseSchedulerNode] | ExternKernelr’   zOrderedSet[Node])r  úSequence[BaseSchedulerNode]r  z8Literal[True, 'torch', 'original_aten', 'inductor_node']r’   r
  )r  rì  rð  r:   r’   útuple[str, str]r‹   )rd  zIterable[torch.fx.Node]re  zCallable[[Any], bool] | Noner’   úOrderedSet[torch.fx.Node])r„   zSequence[IRNode]r¤  zdict[str, IRNode]r’   rï  rˆ  )r  r‘   r’   zValueRanges[Any])r’  r
  r’   r1  )r’  re   rU  rŒ   r’   úsympy.Symbol)rž  r1  r’   r1  )rÛ   r
  r’   rð  )rh  r‘   r¤  zdict[sympy.Expr, Any]r’   r‘   )r   r   r’   z&TypeGuard[torch.SymInt | torch.Tensor])r„   r   r’   r1  )r¥  útorch.fx.GraphModuler’   ztorch.fx.Node | None)r¥  rñ  r’   r5   )r¥  rñ  r’   zOrderedSet[torch.device]rƒ  )r¾  r   r’   r   )r^  r
  r�   r
  r’   rb  )NNT)r  zdict[str, Any] | Nonerð  r¦  r  r1  r’   rb  )r  rê  r  r1  r’   ú	list[int])rs  r+   r  z)Sequence[int | torch.SymInt | sympy.Expr]r  r1  r’   rò  )r©   r"  r’   rŒ   ra  r�  )r¹  zint | torch.devicer’   r1  r§  )rÏ  rŒ   rª   ræ  rÐ  ú
int | Noner’   r8   )rr   )rÕ  rŒ   r’   rŒ   )r,  rA   rÙ  zlist[torch.dtype]r’   r1  )rß  r
  r’   r1  )
r,  rA   ræ  r1  rç  r1  rè  r1  r’   r1  )r$  r@   ró  r   rô  r1  r’   r1  r†  )rý  r!  rô  r1  r’   r1  )r$  r@   ró  rA   rô  r1  r’   r1  )r>  r6   r?  r6   r@  zlist[ScalingType]r’   r1  )r`  r   ra  r   r,  rA   rb  r1  rc  r1  rd  ú
Any | Nonere  rô  rf  rô  r’   r1  )
r,  rA   rÝ  rŒ   r  rŒ   ru  rŒ   r’   r1  rÃ  )r,  rA   rÝ  rx  r  rx  ru  rx  r`  r@   ra  r@   rd  zIRNode | Noner›  z_IntLike | Noner’   r1  )r‡  r
  r’   r1  r‰   )
rÝ  rx  r  rx  ru  rx  r�  rŒ   r’   r1  )rÝ  rx  r  rx  ru  rx  r’   r1  )rÝ  rx  r  rx  ru  rx  r’   rò  )rª   r
  r’   r
  )r’   zNtuple[str | None, Callable[[], list[Any]], Callable[[], list[Any]], type[Any]])r,  rA   r’   r1  )r,  rA   rÖ  zReinterpretView | Bufferr×  r@   r’   r1  )FTFN)r,  rA   rÖ  r@   r×  r@   râ  r1  rÓ  r1  rú  r1  rê  ró  r’   r1  )rÝ   úCallable[P, _T]r„   rß  r¤  rà  r’   ztuple[_T, list[str]])rÝ   ré  r’   ztuple[Any, list[str]])rÝ   rõ  r„   rß  r¤  rà  r’   r¡   )rÝ   rõ  r„   rß  r¤  rà  r’   r
  )rÝ   rõ  r„   rß  r¤  rà  r’   ztuple[Any, list[GraphLowering]])rF  ré  rG  ré  r’   rb  )rO  ré  rN  zCallable[..., Any] | Noner’   r   )rV  r
  r’   rÞ  r¥  )rJ  rê  r’   r1  )rj  zSequence[torch.Tensor]r’   r1  )rj  r‘   r’   r"  )ro  r1  r„   r   r¤  r   r’   zIterator[Any])r©   r"  r’   rn  )rŸ  r
  r’   r1  )rŸ  r
  r’   rŒ   )r¬  zIterable[Any]r’   r1  )
r†  ré  r¥  r4   rÂ  rê  rV  r
  r’   rÞ  )rË  zBuffer | Operation | Noner’   r1  )rg  zNode | Operation | Noner€  ztorch._ops.OperatorBase | Noner’   r1  )rg  zIRNode | Operation | Noner’   r1  )rä  rF   rå  z*Callable[[BaseSchedulerNode], bool] | Noner’   r1  )rä  rF   r’   r1  )rg  zOperation | Noner€  z9torch._ops.OpOverload | Collection[torch._ops.OpOverload]r’   r1  )rí  r
  rî  ri  rï  ri  r’   r   )rä  rF   rø  zMutableSet[BaseSchedulerNode]rî  zdict[str, SchedulerBuffer]rï  zdict[str, BaseSchedulerNode]rõ  zCallable[[Any], bool]r’   rÞ  )r  rŒ   r  rŒ   r’   rŒ   )r  rñ  r’   rŒ   )r  rñ  r’   rò  )r"  r¡   r’   rb  )rÛ   r
  r’   r
  )rª   r¦  r’   r1  )rª   r
  r’   r1  )r©   r"  r’   r1  )rÕ  rè  rŸ  r¦  r:  r"  r;  r"  r<  r
  r=  r1  r’   r1  )r  rí  r’   rÞ  )rÏ   rŸ  r’   r1  )rR  rŸ  r’   r1  )r’   r…  )rÝ   rõ  r„   rß  r¤  rà  r’   ztuple[_T, str])rj  úSequence[InputType]r’   zShapeEnv | None)r­  úCallable[[list[InputType]], _T]rt  rë  ru  zOrderedSet[int]r’   r÷  )rT   rŸ  r’   rŸ  )rq  rv  r�  rë  r‚  zOrderedSet[int] | Noner’   z-tuple[list[torch.Tensor], list[torch.Tensor]])rj  rö  r†  rë  r’   rë  )rê   r‘   r’   r1  )r®  rê  r˜  rE   r’   rÞ  )r©   r"  r’   r
  )r©   r
  r’   r"  )r&  rŸ  r�   rŸ  r’   r1  )r’   ztuple[str, ...])rÛ   r
  rË  r3   rÍ  rÌ  r’   rÞ  )rÛ   r
  r’   rÞ  )rØ   r¦  r’   r
  )r©   r"  r’   r"  )rŽ   ztype[Any] | Noner›   r1  r’   r   )r’   zlist[int] | None)r’   r  )r€  ztorch._ops.OperatorBaser’   rî  )r�  útorch.fx.Noder’   r1  )rg  rB   r’   r1  )rð  r:   r’   r1  )rÈ  rŒ   r’   r"  )rØ   r
  r’   r1  )rN  zSequence[Sequence[T]]rO  zSequence[T]r’   r
  )
rU  rô  rV  rô  rW  rõ  rX  rõ  r’   zEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None])ra  ri  r’   ri  )rw  rŒ   r’   ztuple[bool, bool])rT   r*   r€  r1  r’   zOrderedSet[sympy.Symbol])r’   zdict[str, str])rð  ÚCUDAGraphWrapperTyper’   rÞ  )rä  rF   r’   z tuple[list[Any], dict[str, Any]])rù  r;   r’   r1  )rÛ   r
  r¬  r/   r’   r
  )rg  rø  r’   r1  )r’   r¡   )rT   rŒ   r¿  rŒ   r’   rŒ   )rÍ  ztuple[Any, Any]rÎ  ztuple[Any, ...]rÏ  r   rÐ  r"  rÑ  r"  rÒ  zCallable[[Any, Any], bool]r’   útuple[Any | None, Any | None])rÛ  rŸ  rÜ  rŸ  r’   rú  )rÛ  r>   rÜ  r>   rá  r1  r’   rú  (ª  Ú
__future__r   rC  rc  rú  Úenumr  rF  r‡  r½  rp  r^  rp  r  rê  r¦  rÙ   r  rÜ   rÈ  r†  rü  rr  r«  rQ  ro  Úcollections.abcr   r   r   r   r   r	   r
   r   r   r   Útypingr   r   r   r   r   r   r   r   r   r   r   Útyping_extensionsr   r   r   r   r~   rP   Útorch.utils._pytreerØ  Ú_pytreerž  Ú$torch._inductor.analysis.device_infor   Útorch._inductor.runtime.hintsr   Ú!torch.fx.passes.regional_inductorr   Útorch.utils._dtype_abbrsr    Útorch.utils._ordered_setr!   r"   r#   ÚOPTIMUS_EXCLUDE_POST_GRADÚ#torch.fx.experimental._size_hintingr'   r±  r(   r)   r*   r+   r,   r-   r.   Úpathlibr/   r0   r1   r2   Útorch._prims_commonr3   Útorch.fxr4   Útorch.fx.noder5   rÄ  r6   rË  r8   r%  r:   Údependenciesr;   rq  r=   rý  r>   r?   r@   rA   rB   rC   Úoutput_coderE   rÌ  rF   rG   rN   rL   rà   rX   Útorch._dynamo.device_interfacerY   Útorch._dynamo.utilsrZ   Útorch.autogradr[   Útorch.autograd.profiler_utilr\   Ú(torch.fx.passes.graph_transform_observerr]   Útorch.fx.passes.shape_propr^   Útorch.utils._sympy.functionsr_   r`   ra   rb   rc   Útorch.utils._sympy.symbolrd   re   Útorch.utils._sympy.value_rangesrf   rg   rþ  ri   Úruntime.runtime_utilsrj   r-  Ú_IS_WINDOWSÚ	getLoggerr“   rÓ   rl   rþ  r+  Ú	VarRangesrž  rŒ   Ú	InputTypeÚgetenvÚXPU_KERNEL_FORMATÚGPU_KERNEL_BIN_EXTSrN  r€  rù  rÍ  rV  rK  rW  rM  rX  rT  rU  rÃ   rE  rG  rI  r;  r<  r=  r?  rt   r¤   rx   rz   rƒ   ÚFunctionr…   rû  r�   rî   rô   rò   r  r  r  r%  r(  ra  re  rk  ru  rw  r�  r‰  r¦  rÁ   r²  rº  rÀ  rÃ  rÆ  rÉ  rÒ  rÓ  rÔ  ÚFN_TYPEr×  rñ  ró  rú  r  r  rb  rh  rv  r‡  r�  r“  rš  rŸ  r¢  r¥  r¨  r­  Ú	frozensetr!  r³  r¹  rÅ  rÚ  rß  rà  ræ  rè  rd  rí  r  Úclear_on_fresh_inductor_cacheÚclear_inductor_cachesÚfresh_inductor_cacher  r"  r&  r(  r,  r/  rˆ  r–  r;  r©  r»  r¿  rÂ  rÆ  rÑ  rÖ  rÚ  rà  rå  rñ  r%  r/  r6  r<  rA  rI  rO  rM  rZ  rg  rw  rƒ  r‰  rx  r�  r”  r¦  r©  r¹  rÃ  rÈ  rË  rÎ  rÐ  rÛ  rÔ  rþ  r   r  r  r!  r3  r6  r8  r@  rI  rR  rW  r[  rc  rg  rk  rm  rp  rø  rz  rŽ  r’  r—  r›  r   r£  r¨  r  rõ  r^  ÚEnumr²  rÇ  rÉ  rÎ  rÚ  rÝ  râ  rè  rê  rð  r÷  rÿ  r  r  r  r  r&  r)  rØ  r,  r.  r0  r?  rK  rO  rS  r[  ri  rl  rw  r}  ro  rˆ  r�  r™  r¤  r§  r«  rJ  r¯  Úcompilerª  r­  r²  r·  rº  rÈ  rÊ  rÎ  rÐ  rÑ  rÓ  rÛ  rÞ  rß  rà  râ  rý  r  r  r  r  r  r  r²  r(  r1  r6  r;  r>  ÚSUPPORTED_MKLDNN_DEVICESrB  rF  rS  r\  rr  rz  r~  r�  r‰  r‹  ÚPartitionFnTyperù  r�  r9  r‘  r£  r¨  r­  r¶  r¸  r¹  r¼  rÀ  rÖ  rÝ  râ  )ru  r‡   s   00rW   ú<module>r*     s‡  ðÞ "ã Û Û Û Û Û Û Û 	Û Û Û Û Û 	Û Û 	Û Û Û 
Û Û Û Û Û Û ÷÷ ñ õ Ý Ý ÷÷ ÷ ñ ÷ CÑ BÝ ã ã ß $Ð $Ý ?Ý :Ý EÝ 0Ý /ß ;ð (Ø$ðÐ õ
 <÷ó ñ ß>Ñ>Ýç/Ñ/ÝCÝ$Ý"Ý/å,Ý5Ý!Ý$ßT×TÝ,ß=ò +€	ÙˆCƒL€ð
 ‡�óó ðõ DÝ 0Ý %Ý 2Ý KÝ 0÷õ ÷ 8ß Då Ý =ð �l‰l˜gÑ%€à€g×Ò˜Ó!€ñ ˆTƒ]€Ø�—’˜UŸZšZÐ'Ñ(€	Ø�LŠL˜3Ñ §¡Ñ-°Ñ4€	ñ �E˜i˜bŸišiÐ(IÈ7ÓSð ð
 ØØÐ Ð!Ð"ñÐ ð €Ø€	à€ØÐ ñ
 2<à�ŠØ�
Š
Ø�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ�ŠØ×ÒØ×ÒØ×ÒØ×Òðó2Ð Ð.ó ð( €Ø�{ Q‘Ñ'¨AÒ-°+ÀÒ2BÐ XÐDXÓ XÐBô5ô
LôˆE�NŠNô ð €×Ò˜dÔ#÷ð ó $ðö"GðX ØØ#(ñ	Øðàðð 
ðð !ð	ð
 ôð4 ØØ#(ñ	[Øð[àð[ð 
ð[ð !ð	[ð
 ô[ð| ‡�óó ðôô;ô@ô
+ô	*ô%AðP+Ø	%ð+àô+ô	ôð"/Ø	#ð/àô/ôGñ @OðIØ	ðIà<ðIð 
ôIð0ØðØ ðØ*8ðà+ôö0'ð Øñ	Øðà!ðð ðð ð	ð
 ôð( %'ØØØØñØðà!ðð ðð ð	ð
 ðð ðð ôô )ô'ô#ôô$ñ  ˆcƒN€ÙˆT˜TÔ"€Ø
�;˜s A˜vÑ&¨Ð*Ñ
+€ôE�8˜W Q¨ U™^ô Eôð:Ø)ðàôð+Øð+à/ô+ð\Ø=ðàôð,4)Ø.ð4)àOð4)ð 	ô4)ðn_2Ø=ð_2à!ð_2ð ô_2ðH 15ñØ*ðà-ðð ôð(GØ
ðGØ$5ðGàôGô:,ô^%ôô	DôUô	>ô+ôô-ñ $òóÐ ð$Øðàôôô'ó& 
ôð< !#Ð �IÓ "ô	ôð ×Òó"ó ð"ð( ×Òà+/ØØñ7Ø(ð7à	ð7ð ð7ð ó	7ó ð7ðv !5Ð Ø$Ð Ø"Ð ð 49ö ð( ñ	$Øð$à	2ð$ð ð	$ð
 ô$ðN €×Ñ�QÓó7ó ð7ô�*ô ð ×Ò÷,ð ,ó ñ,÷
R'ò R'ôj
™õ 
ð ×Òó@ó ñ@÷ò ô@?Ñ'õ ?ð ‡�õó ñð8 ×ÑóJó ñJð ×Ñó)ó ñ)õIð  $ñØðàðð ðð õ	÷(ðØðØ+<ðà	õõ õð ØØ#ñØðð ðð ð	ð
 ðð 
õð< 7;ÈuñBØðBØ&3ðBØHLðBà	õBðL 5:ñØðØ-1ðà	õð4 BGñOØðOØ&,ðOØ:>ðOà	õOð2 BGñVØðVØ&,ðVØ:>ðVà	õVð OØðOàðOð %ðOð 
õ	Oð €×Ñ˜QÔó	ó  ñ	ð €×Ñ˜QÔóó  ñóBð0 €×Ñ˜QÔóó  ñð>Øð>àð>ð ð>ð ð	>ð
 ð>ð ð>ð ð>ð ð>ð 
õ>õB-ðn ØñEØðEàðEð ðEð ð	Eð
 ðEð ðEð ðEð ðEð 
õEõPJð ˜EŸJšJÑ&�ˆ)Ó &ð ‡�àEFñØðØðØ!)ðØ?Bðà	óó ñð( ‡�óó ñð. ‡�ó5(ó ñ5(ðp ‡�ó@ó ñ@ð ‡�óRó ñRõ:"õJõõHõ&ðØðØ2ðØ:@ðà	õð: "Ø"&ØØ#ñ=Øð=à
ð=ð ð=ð ð	=ð
  ð=ð ð=ð ð=ð 
õ=õ@'÷Cò Cð"&Øð&àð&ð ð&ð õ	&ð$ØðØ &ðØ2:ðàõõ&/õ(õVð	Øð	Ø &ð	Ø2:ð	àõ	ð#Øð#Ø &ð#Ø2:ð#à$õ#ð* ×Òð.Øð.Ø.@ð.àó.ó ñ.ð$ FJñFØðFØ)BðFàõFõ*	õõB&õ&õõð ×Òóó ñõð ‡�óDó ñDð ‡�ó'6ó ñ'6ðT ‡�óó ñõQõ	.õ0õõ(õ#õKõô(*�$—)’)õ *ð 
Ø
ð 
Ø"-ð 
Ø4Að 
ØHKð 
à	õ 
õF	ðØ(ðà	õð" *.ñ#Ø
!ð#à&ð#ð 
õ#õL(ð =Añ	QØð	Qà9ð	Qð 
õ	Qõ#ðJØ
ðJàAðJð 
õJðLØðLØ .ðLØDRðLàõLñ *=ð
Øð
à5ð
ð ,ð
ð 5ð	
ð
 'ð
ð 
õ
ñ< *=ðØðà5ðð ,ðð 5ð	ð
 'ðð 
õõ:Tõ õ2ð  ×Ò÷ð ó ñð ×Òóó ñõ,!õõ)õ
.õ2ð$Ø&ð$àð$ð ð$ð ð	$ð
 ð$ð ð$ð 
õ$õNHõBõLõ'ð ØðØ &ðØ2:ðàõõ2ð:Ø*ðà"ðð (ðð %õ	ö0	:ð 04ñ$Øð$à$ð$ð -ð$ð 3ö	$ð<Øðà$ðð öö$*IðZØ!ðØ3Bðà	öö:ö&/ñ Ù'Ù#Ù)Ù*ñ $Ù%ò
Ñ ñ +?×*DÒ*DÓ*F×G¡$ ! Q˜˜A™ÓGÑ ð �"—*’*™YÓ'�öHöö	öð ‡�ôó ñð& ×Ò÷.ò .ó ñ.ð
 68Ñ Ñ2Ô 7ðØ
ðà8ðð .ðð 
ö	ñ /9«lÑ ™OÔ :ö)÷#ð #ö"ñ ‘)Ó
�Ù
‘)Ó
�ö-"�¡©Ð 0Ñ1õ -"ñ`  DÕ)ñÀDö ó *ñöö 4§9¢9õ ð ‡�ô4ó ñ4ö2Yö5ö!öH"öJö4
öö1öñ *Ñ ö	ö	öð& "&Ø!%ñ	 
Ø$ð 
à$ð 
ð ð 
ð ð	 
ð
 Kö 
öFRöj&2öRö6öð6 €×Ò˜dÔ#÷ò ó $ñð ™3 ˜8Ñ$�ØÙÑ.Ð/±Ð@ñÑ ÷0ô 0ò *:Ó);Ñ &ö=ö!öH
öö"-ñL óó�ö$%ð
 ôó ñö"ð
MØðMàðMð ðMð ð	Mð
 ðMð &ðMð #öMð`Ø	ðØ*ðà"öðB ñ"Ø	ð"àð"ð ð"ð #÷	"ùóG Hs   íw"