Ë
    ÄÓ^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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.m/Z/ d dl0m1Z1m2Z2m3Z3 d d	lm4Z4 d dl5Z5d dl6Z6d dl7m8c m9Z: d d
l;m<Z< d dl=m>Z> d dl?m@Z@ d dlAmBZB d dlCmDZD d dl7mEZEmFZF d dlGmHZH ddgZId dlJmKZK d dlLmMZMmNZNmOZOmPZPmQZQ e,rhd dlmRZRmSZSmTZT d dlUmVZV d dl6mWZWmXZXmYZY 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 lhmiZi dd!ljmkZkmlZlmmZmmnZnmoZompZp dd"lqmrZr dd#lsmtZtmuZu g d$¢Zv e/d%«      Zwejð                  �d_d&„«       Zyd 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… d d-l†m‡Z‡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d0l“m”Z” dd1l•m–Z— ej                  d2k(  Z˜ e�j2                  eš«      Z›d3Zœ e/d4«      Z�eže5�j>                  e5�j>                  f   Z e6�jB                  e¢z  e6j²                  z  dz  Z£e˜rd5n e�jH                  d6d7«      Z¥d8d9d:e¥› �d;œZ¦d<Z§d<Z¨d<Z©d=Zª eDe6�jV                  e6�jX                  e6�jZ                  e6�j\                  e6�j^                  e6�j`                  e6�jb                  e6�jd                  e6�jf                  e6�jh                  e6�jj                  e6�jl                  e6�jn                  e6�jp                  e6�jr                  g«      Zºd>e»d?<   d@Z¼e¼e¼dz
  z  d k(  re¼dAk\  sJ dB«       ‚�d`dC„Z½�dadD„Z¾ G dE„ dFe5�j~                  «      ZÀ e�j‚                  dG¬H«       G dI„ dJ«      «       ZÂ�db�dcdK„ZÃ	 	 	 �dd	 	 	 	 	 	 	 	 	 �dedM„ZÄ	 	 	 	 	 	 	 	 	 	 �dfdN„ZÅ	 	 	 �dd	 	 	 	 	 	 	 	 	 �dedO„ZÆejð                  �dgdP„«       ZÇ�dhdQ„ZÈ�didR„ZÉ�djdS„ZÊ	 	 	 	 	 	 �dkdT„ZË	 	 	 	 	 	 �dldU„ZÌ�dmdV„ZÍ�dndW„Z–�dodX„ZÎ	 	 	 	 �dpdY„ZÏ�dqdZ„ZÐ�drd[„ZÑ	 	 	 	 �dsd\„ZÒ�dtd]„ZÓd^„ f	 	 	 	 	 �dud_„ZÔ	 	 	 	 	 	 	 	 �dvda„ZÕ�dw�dxdb„ZÖ	 	 �dy	 	 	 	 	 	 	 	 	 �dzdc„Z×	 	 	 	 	 �d{	 	 	 	 	 	 	 	 	 	 	 	 	 �d|dd„ZØ�d}de„ZÙ�d~df„ZÚ�ddg„ZÛ�d€dh„ZÜ�d�di„ZÝ e2dj«      ZÞ e/dkdG¬l«      Zßee'e%eÞf   eßf   Zà G dm„ dne+e(eÞeßf   «      Zá�d‚do„Zâ	 	 	 	 �d‚dp„Zã	 	 	 	 �dƒdq„Zä	 	 	 	 �d„dr„Zå	 	 	 	 	 	 �d…ds„Zæ	 	 	 	 	 	 �d†dt„Zç	 �d‡	 	 	 	 	 �dˆdu„Zè	 	 	 	 	 	 �d‰dv„Zé�dŠdw„Zê�d‹dx„Zë�dŒdy„Zì�dŒdz„Zí�d�d{„Zî�dŽd|„Zï�d�d}„Zð�d�d~„Zñ�d‘d„Zò�d’d€„Zó eôg d�¢e6�jê                  �jì                  �d‚gng z   «      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dLdŒœ�d›d�„�ZdLdŒœ	 	 	 	 	 	 	 �dœdŽ„�Z ejD                  dA«      �d�d�„«       �Z G d�„ d‘e*«      �Z	e�j‚                   G d’„ d“«      «       �Z
 G d”„ d•«      �Z G d–„ d—�e«      �Z G d˜„ d™�e«      �Z�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jD                  �d¢d£„«       �ZejD                  �dgd¤„«       �Z�dgd¥„�Z�d¢d¦„�Z	 �d‡	 	 	 	 	 	 	 �d£d§„�Z�d¤�d¥d¨„�Z	 	 	 	 	 	 �d¦d©„�Z�d§dª„�Z�d§d«„�Z�d§d¬„�Z�d§d­„�ZdLdLdGd®œ	 	 	 	 	 	 	 	 	 �d¨d¯„�ZddLd°œ	 	 	 	 	 	 	 �d©d±„�Z 	 �dª	 	 	 	 	 �d«d²„�Z!dLd³œ	 	 	 	 	 	 	 �d¬d´„�Z"dLd³œ	 	 	 	 	 	 	 �d¬dµ„�Z#	 	 	 	 	 	 	 	 �d­d¶„�Z$ ejD                  d¬·«      �dgd¸„«       �Z% ejD                  d¬·«      �dgd¹„«       �Z&dº„ �Z' ejD                  d¬·«      �dgd»„«       �Z(	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d®d¼„�Z)�d¯d½„�Z*	 	 �d°	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d±d¿„�Z+�d²dÀ„�Z,e¢e5�j>                  z  �Z-dÁe»d¾<   ejð                  	 �d³	 	 	 	 	 	 	 	 	 �d´dÂ„«       �Z.ejð                  �dµdÃ„«       �Z/ejð                  �d¶dÄ„«       �Z0ejð                  �d·dÅ„«       �Z1ejð                  �d¸dÆ„«       �Z2�d¹dÇ„�Z3�d¯dÈ„�Z4�d¯dÉ„�Z5�d¹dÊ„�Z6�d¹dË„�Z7	 	 	 	 	 	 	 	 �dºdÌ„�Z8	 	 	 	 �d»	 	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d¼dÍ„�Z9�dgdÎ„�Z: G dÏ„ dÐ«      �Z;	 	 	 	 	 	 	 	 �d½dÑ„�Z<	 	 	 	 	 	 	 	 �d½dÒ„�Z=�d¾dÓ„�Z>�d¿dÔ„�Z?�dÀdÕ„�Z@	 	 	 	 	 	 	 	 �dÀdÖ„�ZA	 	 	 	 	 	 	 	 �dÁd×„�ZBe�j                   	 	 	 	 	 	 �dÂdØ„«       �ZC	 �d‡	 	 	 	 	 �dÃdÙ„�ZD�dÄdÚ„�ZE�dÅdÛ„�ZF�dÆdÜ„�ZG�dÆdÝ„�ZH�dÇdÞ„�ZI�dÈdß„�ZJe�j                   �dÉdà„«       �ZK�d¢dá„�ZLejð                  �d¢dâ„«       �ZMejð                  �dÊdã„«       �ZNejð                  �d¢dä„«       �ZO�d¢då„�ZP�d¢dæ„�ZQ�dËdç„�ZR�dÌdè„�ZS�dgdé„�ZT�dgdê„�ZU�dÍdë„�ZV�d’dì„�ZW G dí„ dîe�j°                  «      �ZY	 	 	 	 	 	 	 	 	 	 �dÎdï„�ZZ�dÏdð„�Z[	 	 	 	 �dÏdñ„�Z\	 �d‡	 	 	 	 	 �dÐdò„�Z]�dÑdó„�Z^	 �d‡	 	 	 	 	 �dÒdô„�Z_�dÓdõ„�Z`	 	 	 	 	 	 �dÔdö„�Za	 	 	 	 	 	 	 	 �dÕd÷„�Zbdø„ f	 	 	 	 	 	 	 	 	 	 	 �dÖdù„�Zcdú„ f	 	 	 	 	 	 	 	 	 	 	 �dÖdû„�Zd�d×dü„�Ze�dØdý„�Zf�dÙdþ„�Zge�j‚                   G dÿ„ �d «      «       �Zhe�j                   �dÚ�d„«       �Zi�dÛ�d„�Zj�dÜ�d„�Zk�dg�d„�Zl ejD                  d«      	 	 	 	 	 	 	 	 �dÝ�d„«       �Zm	 �d‡ddd�dœ	 	 	 	 	 	 	 	 	 	 	 �dÞ�d„�Zn�dß�d„�Zo�dà�d	„�Zp	 	 	 	 	 	 	 	 	 	 	 	 	 	 �dá�d
„�Zq�dâ�d„�Zr�dã�d„�Zs�dä�d„�Zt�då�d„�Zu	 	 	 	 	 	 	 	 �dæ�d„�Zv�dç�d„�Zw	 	 	 	 	 	 	 	 �dè�d„�Zx�dé�d„�Zy	 �d‡	 	 	 	 	 	 	 �dê�d„�Zz	 	 	 	 	 	 �dë�d„�Z{�dì�d„�Z|	 	 	 	 	 	 �dí�d„�Z}�dg�d„�Z~�dÛ�d„�Z�d�d�d�d�d�d�d�dœ�Z€�e€�j                  «       D � �ci c]  \  } }|| “Œ
 c}} �Z‚ e�j                  �d «      �Z„�dî�d!„�Z…�dï�d"„�Z†�dð�d#„�Z‡�dð�d$„�Zˆejð                  �dñ�d%„«       �Z‰e�j‚                   G �d&„ �d'«      «       �ZŠi �Z‹�d(e»�d)<   	 	 	 	 	 	 	 	 �dò�d*„�ZŒ eD«       �Z��d+e»�d,<   �dó�d-„�ZŽ�d‡�dô�d.„�Z��dõ�d/„�Z��dö�d0„�Z‘ e/�d1«      �Z’ e/�d2«      �Z“ G �d3„ �d4e�e’�e“f   «      �Z” e1dG�¬5«      �d‡dGdHœ�d÷�d6„«       �Z•�dø�d7„�Z– G �d8„ �d9e�j°                  «      �Z—ejð                  �dù�d:„«       �Z˜�dg�d;„�Z™�dú�d<„�Zš�dû�d=„�Z›�dû�d>„�Zœ�dü�d?„�Z��d_�d@„�Zž�dý�dA„�ZŸ�dg�dB„�Z �dþ�dC„�Z¡�dD�Z¢�dÿ�dE„�Z£�dÿ�dF„�Z¤�d �dG„�Z¥	 	 �d°	 	 	 	 	 	 	 	 	 �d�dH„�Z¦�d�dI„�Z§�d�dJ„�Z¨�dg�dK„�Z©�d�dL„�Zª�d�dM„�Z« e�j‚                  dG¬H«       G �dN„ �dO«      «       �Z¬e�dPe%f   �Z­e�e­�e¬g�e­f   �Z® G �dQ„ �dR«      �Z¯ �e¯«       �Z°�d�dS„�Z±�d�dT„�Z²�d�dU„�Z³�d	�dV„�Z´�d
�dW„�Zµ eDg �dX¢«      �Z¶�d²�dY„�Z·e"�d�dZ„«       �Z¸�d�d[„�Z¹	 	 	 	 	 	 	 	 	 	 	 	 	 	 �d�d\„�Zº	 	 	 	 	 	 �d�d]„�Z»	 �dª	 	 	 	 	 	 	 �d�d^„�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)Úhas_triton_packageÚ!activation_quantization_aten_passÚinductor_autotune_lookup_table©Ú_sympy_subs)Úfree_symbolsÚfree_unbacked_symbolsÚGuardOnDataDependentSymNodeÚ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 )Nr9   r   rJ   )Ú	GPU_TYPESÚgetattrÚtorchÚis_availableÚlenÚpop)ÚxÚ
avail_gpusÚgpu_types      ú`/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/utils.pyÚget_gpu_typerZ   k   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ÚMaxÚModularIndexing)Úmake_symbolÚSymT)Úbound_sympyÚValueRanges©Úconfig)ÚceildivÚwin32Ú!inductor_do_bench_using_profilingÚ_TÚspvÚTORCHINDUCTOR_XPU_KERNEL_FORMATÚzebinz.cubinz.hsacoú.)rJ   ÚhiprL   é   é€   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_BYTESr9   )ÚALIGN_BYTES)Únbytess    rY   Ú_alignr~   ¿   s   € à”[Ñ  1Ñ$¬¨Ñ4Ð4ó    c                ó  — t        | t        j                  t        j                  t        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ÚAddre   ÚallÚmapÚ_is_alignedÚargsÚalignÚgcdr|   )Úvs    rY   r†   r†   Ä   sS   € ä�!”e—i‘i¤§¡¬CÐ0Ô1Ü”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©r9   Tc                ó†   — t        |t        t        j                  f«      rt	        t        |«      «      S t        |«      r|S y ©N)r�   Úintr‚   ÚIntegerr~   r†   )ÚclsÚvalues     rY   Ú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   rY   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   rY   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 )aR  
    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.
    ç    €„ŽArJ   ©ÚdtypeÚdeviceT©Úenable_timingé   r9   ©Ú
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     rY   ú	<genexpr>zfp8_bench.<locals>.<genexpr>-  s   è ø€ ÒQ¸˜E×3Õ3ÑQùó   ‚ç     @�@úprofiling results: %s ms)!rR   rJ   Ú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                     rY   Ú	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                óL   — ddl m}m}  |t        «      }  ||«      | |||«      S )Nr   )Úgpu_benchmark_lockÚmay_distort_benchmarking_result)Ú$torch._inductor.runtime.benchmarkingrð   rñ   Ú_do_bench_using_profiling)rÝ   rÞ   rß   Úis_vetted_benchmarkingrð   rñ   Úlocked_benchs          rY   Údo_bench_using_profilingrö   5  s3   € ÷"ñ
 &Ô&?Ó@€LØ8Ñ*¨<Ó8Ø
ˆF�CÐ/óð r   c           
     ó¨  ‡
‡— t        «       Š
dˆ
ˆfd„Š|D �cg c]4  }|j                  t        k(  r|j                  t        j
                  k(  r|‘Œ6 }}t        |«      |k7  r"t        d|› dt        › dt        |«      › d�«      ‚|D ]
  } ‰|«       Œ d}| D ]–  }|j                  «       }|j                  «       }|j                  «       }	|j                  «       |k(  sŒG|	dk7  sŒM|‰
v s|dk(  sŒW|‰
v sŒ\|j                  «       d	k7  sŒp||j                  «       |j                  «       z
  d
z  z  }Œ˜ |dk  rt        dt        › d�«      ‚|d
z  |z  S c c}w )Nc                óª   •— | j                   t        j                  k7  ry ‰j                  | j                  «       | j
                  D ]
  } ‰|«       Œ y rŽ   )rØ   r]   ÚCPUÚaddr¡   Úcpu_children)r¼   ÚchildÚbenchmark_event_idsÚcollect_cpu_event_idss     €€rY   rþ   z;_get_do_bench_profile_result.<locals>.collect_cpu_event_idsY  sG   ø€ Ø×Ñ¤
§¡Ò.Øà×Ñ §¡Ô)Ø×'Ñ'ò 	)ˆEÙ! %Õ(ñ	)r   z	Expected ú z profiling events. Found z events.ç        Úgpu_user_annotationr   zContext Syncr¿   z$Failed to capture device events for rt   )r¼   r   r•   ÚNone)r!   rÛ   Ú_DO_BENCH_PROFILE_EVENT_NAMErØ   r]   rù   rT   ÚRuntimeErrorÚlinked_correlation_idÚcorrelation_idÚactivity_typeÚend_nsÚstart_ns)Úkineto_eventsÚprofiler_eventsræ   Úexpected_device_typer¼   Úbenchmark_eventsÚdevice_time_usr  r  r  rý   rþ   s             @@rY   Ú_get_do_bench_profile_resultr  Q  s•  ù€ ô ,6«<Ðö)ð %öàØ�:‰:Ô5Ò5Ø×Ñ¤§¡Ò/ò 	ðÐð ô ÐÓ Ò(ÜØ˜�z Ô#?Ð"@ð AÜÐ)Ó*Ð+¨8ð5ó
ð 	
ð
 "ò %ˆÙ˜eÕ$ð%ð €NØò KˆØ %× ;Ñ ;Ó =ÐØ×-Ñ-Ó/ˆØ×+Ñ+Ó-ˆà×ÑÓÐ#7Ó7ØÐ!6Ó6à%Ð)<Ñ<à)¨QÓ.°>ÐEXÒ3Xð —
‘
“ Ó.à˜uŸ|™|›~°·±Ó0@Ñ@ÀFÑJÑJ‰NðKð" ˜ÒÜØ2Ô3OÐ2PÐPQÐRó
ð 	
ð ˜FÑ" XÑ-Ð-ùòMs   ˜9Ec           	     óf  — |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        t        j                   j"                  |«      }t        j                   j%                  t        j                   j"                  j&                  |g¬	«      5 }t        |«      D ]E  }|j                  «        t        j                   j)                  t*        «      5   | «        d
d
d
«       ŒG |j                  «        d
d
d
«       t,        j/                  d«       t,        j1                  t2        j4                  «      r4t,        j/                  j7                  «       j9                  dd¬«      «       t;        j                   j<                  j?                  «       |j?                  «       |t        t@        |«      «      }t,        j/                  d«       t,        j/                  d|«       |S # 1 sw Y   �ŒOxY w# 1 sw Y   ŒûxY w)aB  
    Returns benchmark results by examining torch profiler events.
    This could be more accurate as it doesn't count CPU side overhead.
    The benchmarked function is wrapped in a profiler record_function so
    cache-clearing kernels can be excluded without relying on raw CUDA event
    grouping or kernel names.
    r   )Úmay_ban_benchmarkingr©   rª   Tr­   r¯   r9   r°   Nr³   r´   rµ   r¶   zprofiling time breakdownrÀ   )!rò   r  rZ   Úupperr[   rÁ   rR   rÂ   r�   rÄ   rÅ   rÆ   rÇ   rÈ   rÉ   rQ   rÊ   rÌ   rË   rù   Úrecord_functionr  rÓ   rÔ   ÚisEnabledForÚloggingÚDEBUGrÕ   rÖ   r  Úkineto_resultsr×   r]   )rÝ   rÞ   rß   rô   r  rØ   Údevice_type_upperÚdevice_interfacerà   rá   râ   rã   rä   rå   ræ   Úprofile_activityrç   Úresults                     rY   ró   ró   Š  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‹_ò ˆÙ
�ðð × Ñ Ô"ÜœuŸ~™~×>Ñ>Ð@QÓRÐÜ	�‰×	Ñ	ä�N‰N×+Ñ+×/Ñ/Øð
ð 
 ó 
ð 'ð
 
ä�x“ò 	ˆAà�K‰KŒMä—‘×/Ñ/Ô0LÓMñ Ù”÷ð ð		ð 	×$Ñ$Ô&÷'ô  ‡I�IˆlÔÜ
×ÑœŸ™Ô&Ü�	‰	Ø�N‰NÓ×"Ñ"Ð+CÈrÐ"ÓRô	
ô *Ø	�
‰
×!Ñ!×(Ñ(Ó*Ø	�‰‹
ØÜ”
Ð-Ó.ó	€Fô ‡I�IÐ(Ô)Ü‡I�IÐ(¨&Ô1Ø€M÷)ñ ú÷'ð 'ús%   Ç	AL'ÈLÈL'ÌL$ÌL'Ì'L0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  rR   Ú_CÚ%_dispatch_has_kernel_for_dispatch_keyÚhasattrrQ   ÚopsÚImportErrorr  Ústr)r  rê   s     rY   Ú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 )Nr   )ÚcpuÚmeta)Úindex)
rR   rÏ   r¬   r�   r&  Útyper+  r[   ÚWorkerÚcurrent_device©r¬   r  s     rY   Údecode_devicer0  ë  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    rY   Úsympy_productr9  ö  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      rY   r½   zsympy_dot.<locals>.<genexpr>ü  s   è ø€ Ò>¡d a¨˜A �EÑ>ùs   ‚)rT   r‚   ÚexpandÚsumrÐ   )Úseq1Úseq2s     rY   Ú	sympy_dotrB  ú  s8   € Üˆt‹9œ˜D›	Ò!Ð!Ð!Ü�<‰<œÑ>¬c°$¸«oÔ>Ó>Ó?Ð?r   c                ó¤   — t        | «      t        |«      k(  sJ ‚t        j                  j                  }t	        | |«      D ]  \  }}||z  |z   }Œ |S )z7Row-major flatten: per-dimension indices -> flat index.)rT   r‚   r6  ÚZerorÐ   )ÚindicesÚsizesÚflatr+  Úsizes        rY   Úflatten_indexrI  ÿ  sT   € ô
 ˆw‹<œ3˜u›:Ò%Ð%Ð%Ü�7‰7�<‰<€DÜ˜7 EÓ*ò #‰ˆˆtØ�d‰{˜UÑ"‰ð#à€Kr   c                ó~   — t        |«      D ��cg c]!  \  }}t        | t        ||dz   d «      |«      ‘Œ# c}}S c c}}w )z=Row-major decomposition: flat index -> per-dimension indices.r9   N)Ú	enumeraterf   r9  )r+  rF  rè   rH  s       rY   Údecompose_indexrL    sE   € ô ! Ó'÷áˆAˆtô 	˜œ}¨U°1°q±5°7¨^Ó<¸dÕCóð ùó s   �&9c                ó\   — | D �ci c]  }t        |«      |“Œ c}j                  «       S c c}w rŽ   )r¡   Úvalues)r8  rV   s     rY   ÚuniquerO    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‚   ÚExprra   Úsympifyr�   r,  Úruntime_ceildiv)ÚnumberÚdenoms     rY   rm   rm     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*i8rt   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ÚlistrN  r�   )ÚkeyÚ	dtype_strÚtysrŠ   s       rY   Ú_type_ofrˆ  &  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‚   rS  ©Úlstrè   s     rY   Úconvert_shape_to_inductorrŒ  N  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�   rR   r4   ÚnodeÚexpr©Úvals    rY   Úconvert_symint_to_exprr’  Y  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.
    r9   ©ÚVN)Úhint)
Úvirtualizedr•  r�   r�   r‚   r�   ÚgraphÚsizevarsÚ	shape_envÚcreate_symintnode)rè   r•  s     rY   Úconvert_to_symintrœ  e  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).
    )rœ  rŠ  s     rY   Úconvert_shape_to_symintrž  v  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     rY   r½   zis_view.<locals>.<genexpr>„  s   è ø€ ÒF¨Aˆq�|‰| 4Ô'ÑFùs   ‚)ÚanyÚ_schemaÚ	arguments©Úops    rY   Úis_viewr¨  €  s   € ô ÑF°·±×1EÑ1EÔFÓFÐFr   c                 ó   — y©NFr�   )rã   s    rY   ú<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     €rY   r½   z#is_pointwise_use.<locals>.<genexpr>š  s   øè ø€ ÒK¸AÔ# A ×7ÑKùs   ƒ)r§  r�   ÚtargetrR   Ú_opsÚ
OpOverloadr4  Úgetitemr   r¨  r„   ÚusersÚTagÚ	pointwiseÚtags)Úuser²  r³  s    ` rY   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ÚplaceholderrT   )r¿  ÚgÚ
graph_argss    €€rY   Úadd_tensor_argz)gen_gm_and_inputs.<locals>.add_tensor_arg¥  s,   ø€ Ø×Ñ˜#ÔØ�}‰}˜s¤3 z£?Ð"3Ð4Ó5Ð5r   r9   r   ÚTensor)r¿  útorch.Tensorr•   r7   )rR   ÚfxÚGraphr®  r#   rÅ  rT   r¤  Úreturnsr&  r,  Úoutputr6   )r³  r‡   ÚkwargsrÄ  rŽ  ÚgmrÂ  rÃ  s         @@rY   Úgen_gm_and_inputsrÍ  Ÿ  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)  )r[   rS   rÁ   r/  s     rY   rÁ   rÁ   ·  s4   € Ø�‚ØÜ/°Ó7ÐØ×$Ñ$Ô&Ø×$Ñ$Õ&ð 'r   c                óä   — t        |«       t        j                  d«       t        j                  «       }t        |«      D ]  } | |Ž }t        |«       Œ t        j                  «       }€J ‚||z
  S )Ni9  )rÁ   rR   Úmanual_seedÚtimeÚperf_counterrÆ   )ÚmodelÚexample_inputsrë   r¬   Út0rã   r  Út1s           rY   ÚtimedrØ  ¿  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)rR   rÏ   rÆ   rØ  ÚmedianÚprintrÒ   )	rÔ  rÕ  rë   ÚrepeatÚbaseliner¬   rã   ÚtimingsÚtooks	            rY   Ú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   €rY   r«  z#precompute_method.<locals>.<lambda>ä  s   ø€  € r   N)rQ   Úsetattr)ÚobjÚmethodr  s     @rY   Ú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      rY   Úprecompute_methodsré  ç  s   € àò 'ˆÜ˜#˜vÕ&ñ'r   c                ó<   — t        | |kD  «      t        | |k  «      z
  S rŽ   )r�   ©r<  r=  s     rY   Ú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 )Nr9   r   )r�   r�   rT   r,  )rV   rH  s     rY   Ú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 )Nr9   )rH   )r�   r&  Ú	schedulerrH   Úget_name)ÚelemrH   s     rY   Ú	sort_funcztuple_sorted.<locals>.sort_funcþ  s1   € Ü�dœCÔ ØˆKå0ä˜$Ð 1Ô2Ð2Ð2Ø�}‰}‹Ðr   ©r…  )ró  rp   r•   r&  )rT   Úsorted)rV   rô  s     rY   Ú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    rY   Úclear_cachezCachedMethod.clear_cache  s   € Ø),r   c                 ó   — y rŽ   r�   ©Úselfr‡   rË  s      rY   Ú__call__zCachedMethod.__call__  r¬  r   N)rà   r   r•   r  )r‡   úP.argsrË  úP.kwargsr•   rù  )r–   r—   r˜   Ústaticmethodrþ  r  r�   r   rY   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    €rY   rþ  z"cache_on_self.<locals>.clear_cache,  s   ø€ Ü�4˜ÔÜ�D˜#Õð r   ©r  r   r•   r  ©r–   ÚexecÚlstripr2  Úwrapsrþ  )rÝ   rÛ   ÚctxÚwrapperrþ  r…  s        @rY   Ú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    rY   Úcache_property_on_selfr  4  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    €rY   rþ  z<cache_on_self_and_args.<locals>.wrapper.<locals>.clear_cacheb  s   ø€ Ü�t˜SÔ!Ü˜˜cÕ"ð "r   r  r  )rÝ   r  r  rþ  r…  Ú
class_names       @€rY   r  z'cache_on_self_and_args.<locals>.wrapperC  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   ` rY   Ú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 )Nr9   ©ÚirrŽ  )Ú r!  r�   r„  r2  r3  r4  Úor_r#  rŽ  Úoriginsr!   rA   )Únode_scheduler!  rŽ  s      rY   Úaggregate_originsr&  l  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�   rR   r´  rµ  Ú_overloadpacketr–   ÚHigherOrderOperatorr&  rÛ   )Úoriginr(  r…  s      rY   Ú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®  rR   r"  Úsource_fn_stackrµ   Úfwd_source_fn_stackÚbackwardr9   Ú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           rY   Úget_fused_kernel_namer:  ‚  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Ž   )r˜  )r»   Úns     rY   r½   z&get_kernel_metadata.<locals>.<genexpr>Ö  s   è ø€ Ò"C¨q 1§7¥7Ñ"Cùr¾   r9   r   Ú)_inductor_kernel_metadata_node_to_idx_mapc                ó"   •— ‰j                   |    S rŽ   )r>  )r=  Úsingle_graphs    €rY   r«  z%get_kernel_metadata.<locals>.<lambda>Þ  s   ø€ ˜l×TÑTÐUVÑW€ r   rõ  r(  Ú	from_nodeÚpartitioner_tagÚis_backwardzTopologically SortedÚUnsortedrÿ   z Source Nodes: [rQ  z], Original ATen: [ú]z" Source node to ATen node mapping:ú   z => r   z Graph fragment:r”  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_layoutr3  )ÚbufferÚrw_namerK  rÛ   Úlayoutr!  s        €rY   Ú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ú[rQ  rE  )r4  r&  )ÚshaperV   s     rY   Ú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"  ú")rH  Ústrider¬   r    r«   )rO  Úshape_annotationÚstride_annotationÚdevice_annotationrT  s       €rY   Ú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)rF  N)Úcomment)r»   Úliner  s     €rY   r½   z&get_kernel_metadata.<locals>.<genexpr>O  s)   øè ø€ ò )àð —‘Ð' s¨4¨&Ô1ñ)ùs   ƒz
   return ú,ú
)rM  z-ir.TensorBox | ir.Buffer | ir.TorchBindObjectrN  r&  r•   ztuple[str, ir.Layout | None])rS  zIterable[int]r•   r&  )rO  zir.Layout | Noner•   r&  ),r&  r§  ÚcollectionsÚdefaultdictr„  r!   rT   r˜  r#  rK  Únodesr>  Úsortr*  r�   rR   r´  rµ  r&  r*  r+  rÛ   rÀ  Úgetrb  r4  Úkeysrö  Úitemsr"  r!  rA   r—  r•  r\  r]  rú   Útry_get_bufferr^  Úformat_nodeÚversionru   ÚextendÚ
splitlines)#r%  r  r6  r,  Úinductor_nodesÚfrom_node_dictÚoriginal_aten_dictÚunique_graphsÚidxr=  Únode_to_idx_maprŽ  r(  r…  Úsort_strÚmetadataÚdetailed_metadataÚoriginal_noderh  Ú	all_readsÚ
all_writesr•  rP  r[  ÚrrM  Ú
input_namerO  ÚwÚoutput_namerã   Úformatted_noder!  r@  rT  s#    `                              @@@rY   Úget_kernel_metadatarƒ  ¹  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!   rU   r·  rú   rÀ  )Úinitial_queueÚskip_filterÚdominated_setrŽ  Úusers        rY   Údominated_nodesr‰  [  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 )Nr9   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�   rH  rI  rJ  rB   ÚComputedBufferÚInputsKernelÚInputBufferÚTemplateBuffer)r=  r!  Úis_unrealized_nodes    €€rY   r�  z*gather_origins.<locals>.is_unrealized_nodet  s‰   ø€ Ü�a˜Ÿ™Ô&Ù% a§f¡fÓ-Ð-Ü�a˜Ÿ™Ô'Ù% a§f¡fÓ-Ð-Ü˜!˜RŸY™YÓ'ò 
´
Øà×!Ñ!Ø—‘Ø—‘Ø×!Ñ!ð	ó1
ð -
ð 	
r   )r=  rB   r•   rX  )r"  r!  r"   r$  r!   Ú	itertoolsÚchain)
r‡   rË  Úkwargs_flattenrã   r‘  Úkwargs_originsÚargs_flattenÚargs_originsr!  r�  s
           @@rY   Úgather_originsr—  o  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‚   ÚMulrT   r‡   )r�  s    rY   Úis_neg_leadzsympy_str.<locals>.is_neg_lead“  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 )Nrš  r9   r   z - z + )r�   r‚   rƒ   rT   r‡   r4  r…   )r�  rœ  Úsympy_str_muls    €€rY   Úsympy_str_addz sympy_str.<locals>.sympy_str_add˜  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ú-r9   z * )r�   r‚   r›  r‡   r4  r…   )r�  rœ  Úsympy_str_atoms    €€rY   rž  z sympy_str.<locals>.sympy_str_mul£  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ú(ú)rQ  )r�   r‚   ÚSymbolrÛ   rƒ   r›  rf   rb   rc   rd   Úfuncr–   r4  r…   Ú	sympy_strr‡   r&  )r�  rŸ  s    €rY   r¢  z!sympy_str.<locals>.sympy_str_atom®  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   )r�  r”   r•   rX  ©r�  r”   r•   r&  r�   )r�  rœ  rŸ  r¢  rž  s    @@@@rY   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 )Nr9   r”  Úcurrent_nodeÚ
index_expr)
r—  r•  rl   Úcompute_all_boundsrQ   Úinterpreterr³  ri   rj   Úunknown)r+  r•  Úfx_nodes      rY   Úget_bounds_index_exprr±  »  sN   € Ýô 	×!Ò!Ü §¡¨~¸tÓDÐDˆWÐDØ�N‰N˜lÒ*ä˜5Ó!Ð!ä×"Ñ"Ó$Ð$r   c                ó   — | d   dk(  S )Nr   r~  r�   ©Úprefixs    rY   Úprefix_is_reductionrµ  É  s   € Ø�!‰9˜ÑÐr   c                ó   — | d   dv S )Nr   )rV   ÚyÚzr�   r³  s    rY   Úprefix_is_pointwiser¹  Í  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)rh   ÚSIZErg   )r´  rv  s     rY   Ú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Ž   )rl   Údebug_index_assertsÚassert_indirect_indexing)Úchecks    rY   Ú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    rY   Ú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'   )r�  Úreplacementss     rY   Ú
sympy_subsrË  í  s   € ô
 �t˜\Ó*Ð*r   c                óŠ   — t        | t        j                  «      xs( t        | t        j                  «      xr | j                  S rŽ   )r�   rR   r4   rÅ  Ú_has_symbolic_sizes_strides)r<  s    rY   Ú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     rY   r½   z"any_is_symbolic.<locals>.<genexpr>ü  s   è ø€ Ò, !Œ{˜1�~Ñ,ùó   ‚©r£  )r‡   s    rY   Ú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_scalarzaten.topk.defaultc                ó¸   — ddl m} | j                  j                  D ];  }t	        |«      r|c S |j
                  j                  d«      x}€Œ0 ||«      sŒ9|c S  y )Nr   )r*   r‘  )Ú%torch.fx.experimental.symbolic_shapesr*   r˜  rh  Úis_cudagraph_unsafe_fx_noder*  rj  )rÌ  r*   rŽ  r‘  s       rY   Ú%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Úreversedr˜  rh  r§  )rÌ  Ú	last_nodes     rY   Ú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©r‘  N)r�   r*  rj  rR   rÅ  r¬   )r»   rŽ  s     rY   r½   z"get_all_devices.<locals>.<genexpr>4  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�   rR   rÇ  r7   r*  rj  rÅ  r¬   )r»   r¿  s     rY   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<)r˜  Ú
find_nodesr!   rß  r‡   r�   Útuple)rÌ  Úplaceholder_nodesÚinput_devicesÚout_argÚout_argsÚout_devicess         rY   Úget_all_devicesrë  2  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Úmodulesrk  Ú
startswithÚ__dict__rQ   r�   rR   Ú	_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         rY   Úunload_xpu_triton_pydsr  H  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    rY   Úclear_on_fresh_cacher  i  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    rY   Úclear_cachesr  u  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Úenvironrj  rU   )r…  r’   Úolds      rY   Ú_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      €rY   r«  zfresh_cache.<locals>.<lambda>À  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  r4  r�   ÚdictrT   ÚexistsÚlistdirrƒ  ÚgetsizeÚ
is_windowsrR   rL   rS   r  ÚshutilÚrmtreeÚ	Exceptionr  )Úcache_entriesr  Údeleter  Útriton_cache_dirÚfilesÚfr  s          @rY   Úfresh_cacher1  ’  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…  r2  )Ú__getitem__rÆ   rT   r„  rö  rÝ  )Úseqr2  ÚgetterÚa_rÚsort_idxs        rY   Úargsortr:  Ó  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 )zç
    Return a symbolic sort order for optimization-only layout heuristics.

    Data-dependent comparisons can fall back to non-guarding optimization hints,
    so callers must not use this order to make correctness decisions.
    c                ó¤   •‡‡‡— | \  }Š|\  }Š	 	 	 	 	 	 dˆfd„}d	ˆfd„Š |‰‰k  ˆˆˆfd„«      ry |‰‰kD  ˆˆˆfd„«      ry||k  ry||kD  ryy)
Nc                óx   •— t        | t        «      r| S 	 ‰j                  | «      S # t        $ r
  |«       cY S w xY wrŽ   )r�   rX  Úevaluate_exprr+   )r�  Úfallbackrš  s     €rY   Úevaluatez*argsort_sym.<locals>.cmp.<locals>.evaluateô  sA   ø€ ô ˜$¤Ô%Ø�ð"Ø ×.Ñ.¨tÓ4Ð4øÜ.ò "Ù“zÒ!ð"ús   •& ¦9¸9c                óJ   •— ‰j                  | «      ‰j                  |«      k  S rŽ   )Úoptimization_hint)ÚlhsÚrhsrš  s     €rY   Úhint_ltz)argsort_sym.<locals>.cmp.<locals>.hint_ltÿ  s'   ø€ ð ×.Ñ.¨sÓ3°i×6QÑ6QÐRUÓ6VÑVÐVr   c                 ó   •—  ‰‰ ‰«      S rŽ   r�   ©Úa_valÚb_valrE  s   €€€rY   r«  z*argsort_sym.<locals>.cmp.<locals>.<lambda>  ó   ø€ ©7°5¸%Ó+@€ r   rµ   c                 ó   •—  ‰‰‰ «      S rŽ   r�   rG  s   €€€rY   r«  z*argsort_sym.<locals>.cmp.<locals>.<lambda>  rJ  r   r9   r   )r�  z bool | torch.SymInt | sympy.Exprr?  zCallable[[], bool]r•   rX  )rC  r”   rD  r”   r•   rX  r�   )	r<  r=  Úa_idxÚb_idxr@  rH  rI  rE  rš  s	        @@@€rY   rì  zargsort_sym.<locals>.cmpð  sy   û€ Ø‰ˆˆuØ‰ˆˆuð		"Ø2ð		"à(ð		"ð õ		"õ	Wñ �E˜E‘MÕ#@ÔAØÙ�E˜E‘MÕ#@ÔAØð
 �5Š=ØØ�5Š=ØØr   r4  )r<  útuple[int, sympy.Expr]r=  rN  r•   r�   )	rK  r�   rR   r4   rŽ  r�  rö  r2  Ú
cmp_to_key)	rš  r6  r2  rì  rv  ré   Úexprsrã   r  s	   `        rY   Úargsort_symrQ  ã  s�   ø€ õ!ôL   “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 )Nrz   r�   ©r«   )rR   r€  rÂ   Úelement_sizerS  s    rY   Úget_dtype_sizerU    s-   € ð ”—‘ÒØÜ�;‰;�r Ô'×4Ñ4Ó6Ð6r   c                  ó   — e Zd ZU ded<   y)ÚLineContextr   ÚcontextN©r–   r—   r˜   r§   r�   r   rY   rW  rW  &  s   … Ø„Lr   rW  c                  ó"   — e Zd ZU ded<   ded<   y)ÚValueWithLineMapr&  r’   zlist[tuple[int, LineContext]]Úline_mapNrY  r�   r   rY   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„Zd#d$d„Zd#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     rY   Ú__init__zIndentedBuffer.__init__3  s   € ØBDˆŒØ%ˆ�r   c              #  ób   K  — | j                   }	 || _         d –— || _         y # || _         w xY w­wrŽ   )Útabwidth)r  rf  Úprevs      rY   Úset_tabwidthzIndentedBuffer.set_tabwidth7  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 )Nr9   re  )r   ra  r�   ÚDeferredLineBaserW  rÀ  rX  r&  ÚwriteÚcountr[  Úgetvalue)r  Úbufrç   ÚlinemapÚlirc  s         rY   Ú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Ž   )rq  r’   ©r  s    rY   rm  zIndentedBuffer.getvalueT  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µ   re  )	r   ra  r�   rj  rW  r&  Úendswithrk  rm  )r  rn  rp  rc  s       rY   ÚgetrawvaluezIndentedBuffer.getrawvalueW  sš   € Ü‹jˆØ—+‘+ò 	 ˆBÜ˜"Ô.Ô/Ù“t�Ø�<ØÜ˜B¤Ô,Øà�Ü˜d¤CÔ(Ð(Ð(à�}‰}˜TÔ"Ø—	‘	˜$˜s ˜)Õ$à—	‘	˜$”Ø—	‘	˜$•ð	 ð  �|‰|‹~Ðr   c                ó   — | j                   S rŽ   ©ra  rs  s    rY   Úget_lines_refzIndentedBuffer.get_lines_refk  s   € Ø�{‰{Ðr   c                ó8   — | j                   j                  «        y rŽ   )ra  Úclearrs  s    rY   r|  zIndentedBuffer.clearn  s   € Ø�‰×ÑÕr   c                ó,   — t        | j                  «      S rŽ   )rX  ra  rs  s    rY   Ú__bool__zIndentedBuffer.__bool__q  s   € Ü�D—K‘KÓ Ð r   c                ó:   — d| j                   | j                  z  z  S )Nrÿ   )rb  rf  rs  s    rY   r´  zIndentedBuffer.prefixt  s   € Ø�d—l‘l T§]¡]Ñ2Ñ3Ð3r   c                ó&   — | j                  d«       y )Nre  ©Ú	writeliners  s    rY   ÚnewlinezIndentedBuffer.newlinew  ó   € Ø�‰�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�   rW  ra  rÀ  rj  Úwith_prefixr´  Ústrip©r  rc  s     rY   r‚  zIndentedBuffer.writelinez  sˆ   € Ü�dœKÔ(Ø�K‰K×Ñ˜tÕ$Ü˜Ô.Ô/Ø�K‰K×Ñ˜t×/Ñ/°·±³Ó>Õ?Ø�Z‰ZŒ\Ø�K‰K×Ñ $§+¡+£- °°Ð7Õ8à�K‰K×Ñ˜rÕ"r   c                ó&   — | j                  |«       y)zGWrite to JIT buffer only. On a plain IndentedBuffer, same as writeline.Nr�  r‰  s     rY   Úwriteline_jitzIndentedBuffer.writeline_jit„  s   € à�‰�tÕr   c                 ó   — y)z;Write to AOTI buffer only. No-op on a plain IndentedBuffer.Nr�   r‰  s     rY   Úwriteline_aotzIndentedBuffer.writeline_aotˆ  r¬  r   c                ó*   — | j                  ||¬«       y)zESplice to JIT buffer only. On a plain IndentedBuffer, same as splice.©rˆ  N©Úsplice©r  Ú
other_coderˆ  s      rY   Ú
splice_jitzIndentedBuffer.splice_jit‹  s   € à�‰�J eˆÕ,r   c                 ó   — y)z<Splice to AOTI buffer only. No-op on a plain IndentedBuffer.Nr�   r’  s      rY   Ú
splice_aotzIndentedBuffer.splice_aot�  r¬  r   c                ó4   — |D ]  }| j                  |«       Œ y rŽ   r�  ©r  Úlinesrc  s      rY   Ú
writelineszIndentedBuffer.writelines’  ó   € Øò 	!ˆ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Ž   ©rb  ©Úoffsetr  s   €€rY   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   `` rY   ÚindentzIndentedBuffer.indent–  s$   ù€ Ü	×	"Ñ	"õ	'ó 
#ð	'ñ ‹uˆr   c                ó.   — | xj                   |z  c_         y rŽ   rž  ©r  r   s     rY   Ú	do_indentzIndentedBuffer.do_indent¡  ó   € Ø�Š˜ÑŽr   c                ó.   — | xj                   |z  c_         y rŽ   rž  r©  s     rY   Údo_unindentzIndentedBuffer.do_unindent¤  r«  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 ]  }t        j                  | |«       Œ y )NÚinfr   re  )r�   r^  Úfloatra  rW  ÚminrT   r  ÚmathÚisinfrÀ  r‚  r�   ÚtextwrapÚdedentÚrstripr‚  )r  r“  rˆ  rµ  rc  ré   s         rY   r‘  z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Ó+ò 2�Ü×(Ñ(¨¨qÕ1ñ2r   c                ó†   — t        | j                  ¬«      }| j                  D �cg c]
  } ||«      ‘Œ c}|_        |S c c}w ©N)rc  )r^  rb  ra  )r  r§  rì   rc  s       rY   r…   zIndentedBuffer.map¿  s4   € Ü¨D¯L©LÔ9ˆØ-1¯[©[Ö9 T‘d˜4•jÒ9ˆŒ
Øˆ
ùò :s   ¥>c                ó@   — t        | «      › d| j                  «       › d�S )Nr¤  r¥  )r,  rm  rs  s    rY   Ú__repr__zIndentedBuffer.__repr__Ä  s    € Ü�t“*�˜Q˜tŸ}™}›Ð/¨qÐ1Ð1r   c                óÔ   — | j                   |j                   k(  sJ ‚t        | j                   ¬«      }|j                  | j                  «       |j                  |j                  «       |S r¸  )rb  r^  rš  ra  )r  Úotherrì   s      rY   Ú__add__zIndentedBuffer.__add__Ç  sK   € Ø�|‰|˜uŸ}™}Ò,Ð,Ð,Ü¨D¯L©LÔ9ˆà�‰�t—{‘{Ô#Ø�‰�u—|‘|Ô$Øˆ
r   c                ó   — || j                   v S rŽ   ry  )r  Únew_lines     rY   ÚcontainszIndentedBuffer.containsÏ  s   € Ø˜4Ÿ;™;Ð&Ð&r   N©r   ©rc  r�   r•   r  )rf  r�   r•   r¢  )r•   r[  ©r•   r&  ©r•   r  ©r•   rX  )rc  z$LineContext | DeferredLineBase | strr•   r  ©F)r“  zIndentedBuffer | strrˆ  rX  r•   r  )r™  z.Sequence[LineContext | DeferredLineBase | str]r•   r  rŒ   )r   r�   r•   ú'contextlib.AbstractContextManager[None]©r   r�   r•   r  )r§  zCallable[[Any], Any]r•   r^  )r¼  r   r•   r^  )r¿  z$DeferredLineBase | LineContext | strr•   rX  )r–   r—   r˜   rf  rd  r¤  r¥  rh  rq  rm  rw  rz  r|  r~  r´  rƒ  r‚  r‹  r�  r”  r–  rš  r§  rª  r­  r‘  r…   rº  r½  rÀ  r�   r   rY   r^  r^  0  s–   „ Ø€Hô&ð ×Ñò!ó ð!ó9ó(0óò(óó!ó4óó#óóJô-ôKó!ô	ôôô2ó0ó
2óô'r   r^  c                  óœ   ‡ — e Zd ZdZddˆ fd„Zedd„«       Zd„ Zdd„Zdd„Z	d„ Z
ddd„Zddd	„Zddd
„Zddd„Zddd„Zddd„Zdˆ fd„Zˆ xZS )ÚDualIndentedBufferzÿIndentedBuffer that simultaneously accumulates JIT and AOTI output.

    The base class (self._lines) holds JIT content. self.aot holds AOTI content.
    By default, writeline/splice write to both. Use _jit/_aot variants for
    mode-specific writes.
    c                óD   •— t         ‰| �  |«       t        |«      | _        y rŽ   )Úsuperrd  r^  Úaot)r  rc  Ú	__class__s     €rY   rd  zDualIndentedBuffer.__init__Û  s   ø€ Ü‰Ñ˜Ô(Ü! .Ó1ˆ�r   c                ó   — | S )uë   Read-only accessor for JIT buffer (for splicing into result).

        WARNING: Do NOT use jit for writing â€” writes would go to both buffers
        because self.writeline is overridden. Use writeline_jit/splice_jit instead.
        r�   rs  s    rY   ÚjitzDualIndentedBuffer.jitß  s	   € ð ˆr   c                óf   — t         j                  | |«       | j                  j                  |«       y rŽ   )r^  r‚  rÍ  r‰  s     rY   r‚  zDualIndentedBuffer.writelineè  s$   € Ü× Ñ   tÔ,Ø�‰×Ñ˜4Õ r   c                ó0   — t         j                  | |«       y rŽ   )r^  r‚  r‰  s     rY   r‹  z DualIndentedBuffer.writeline_jitì  s   € Ü× Ñ   tÕ,r   c                ó:   — | j                   j                  |«       y rŽ   )rÍ  r‚  r‰  s     rY   r�  z DualIndentedBuffer.writeline_aotï  s   € Ø�‰×Ñ˜4Õ r   c                ó4   — |D ]  }| j                  |«       Œ y rŽ   r�  r˜  s      rY   rš  zDualIndentedBuffer.writelinesò  r›  r   c                óª   — t         j                  | ||¬«       t        |t        «      r|j                  n|}| j                  j                  ||¬«       y ©Nr�  )r^  r‘  r�   rÊ  rÍ  ©r  r“  rˆ  Ú	aot_others       rY   r‘  zDualIndentedBuffer.spliceö  sG   € ô 	×Ñ˜d J°eÐÔ<ä(¨Ô5GÔHˆJ�NŠNÈjð 	ð 	�‰�‰˜	¨ˆÕ/r   c                ó4   — t         j                  | ||¬«       y rÖ  )r^  r‘  r’  s      rY   r”  zDualIndentedBuffer.splice_jitÿ  s   € Ü×Ñ˜d J°eÐÕ<r   c                óz   — t        |t        «      r|j                  n|}| j                  j                  ||¬«       y rÖ  )r�   rÊ  rÍ  r‘  r×  s       rY   r–  zDualIndentedBuffer.splice_aot  s0   € ä(¨Ô5GÔHˆJ�NŠNÈjð 	ð 	�‰�‰˜	¨ˆÕ/r   c                óF   ‡ ‡— t         j                  ˆˆ fd„«       } |«       S )Nc               3  óZ  •K  — ‰xj                   ‰ z  c_         ‰j                  xj                   ‰ z  c_         	 d –— ‰xj                   ‰ z  c_         ‰j                  xj                   ‰ z  c_         y # ‰xj                   ‰ z  c_         ‰j                  xj                   ‰ z  c_         w xY w­wrŽ   ©rb  rÍ  rŸ  s   €€rY   r  z&DualIndentedBuffer.indent.<locals>.ctx	  s{   øè ø€ à�LŠL˜FÑ"�LØ�H‰H×Ò Ñ&Õð+Ûà—’ Ñ&•Ø—‘× Ò  FÑ*Ö øð —’ Ñ&•Ø—‘× Ò  FÑ*Ö üs   ƒ5B+¹A2 ½5B+Á26B(Â(B+r£  r¦  s   `` rY   r§  zDualIndentedBuffer.indent  s$   ù€ Ü	×	"Ñ	"ô	+ó 
#ð	+ñ ‹uˆr   c                ól   — | xj                   |z  c_         | j                  xj                   |z  c_         y rŽ   rÝ  r©  s     rY   rª  zDualIndentedBuffer.do_indent  ó%   € Ø�Š˜Ñ�Ø�‰×Ò˜FÑ"Ör   c                ól   — | xj                   |z  c_         | j                  xj                   |z  c_         y rŽ   rÝ  r©  s     rY   r­  zDualIndentedBuffer.do_unindent  rß  r   c                óV   •— t         ‰| �  «        | j                  j                  «        y rŽ   )rÌ  r|  rÍ  ©r  rÎ  s    €rY   r|  zDualIndentedBuffer.clear  s   ø€ Ü‰‰ŒØ�‰�‰Õr   rÁ  rÂ  ©r•   r^  rÄ  rÆ  ©rˆ  rX  r•   r  rŒ   )r   r�   rÈ  )r–   r—   r˜   r™   rd  ÚpropertyrÐ  r‚  r‹  r�  rš  r‘  r”  r–  r§  rª  r­  r|  Ú__classcell__©rÎ  s   @rY   rÊ  rÊ  Ó  s]   ø„ ñö2ð òó ðò!ó-ó!ò!ô0ô=ô0ôô#ô#÷ñ r   rÊ  c                  ó4   — e Zd ZdZdd„Zdd„Zdd	d„Zdd	d„Zy)
ÚAotOnlyBuffera*  IndentedBuffer for pure-AOTI codegen.

    Mirror of the base class's pure-JIT defaults: writeline_aot/splice_aot
    write to the buffer; writeline_jit/splice_jit are no-ops. Lets call
    sites use writeline_jit/writeline_aot uniformly across pure-JIT,
    pure-AOTI, and dual-wrapper modes.
    c                 ó   — y rŽ   r�   r‰  s     rY   r‹  zAotOnlyBuffer.writeline_jit+  ó   € Ør   c                ó&   — | j                  |«       y rŽ   r�  r‰  s     rY   r�  zAotOnlyBuffer.writeline_aot.  r„  r   c                 ó   — y rŽ   r�   r’  s      rY   r”  zAotOnlyBuffer.splice_jit1  rë  r   c                ó*   — | j                  ||¬«       y rÖ  r�  r’  s      rY   r–  zAotOnlyBuffer.splice_aot4  s   € Ø�‰�J eˆÕ,r   NrÄ  rÆ  rä  )r–   r—   r˜   r™   r‹  r�  r”  r–  r�   r   rY   ré  ré  "  s   „ ñóóôõ-r   ré  c                 ób   — ddl m}  | j                  j                  r
t	        «       S t        «       S )zèConstruct the IndentedBuffer subclass matching the current codegen mode.

    Pure AOTI -> AotOnlyBuffer (writeline_aot writes; writeline_jit drops).
    Pure JIT  -> IndentedBuffer  (writeline_jit writes; writeline_aot drops).
    r9   r”  )r—  r•  r˜  Úaot_moderé  r^  r”  s    rY   Úmake_codegen_bufferrñ  8  s%   € õ à‡w�w×ÒÜ‹ÐÜÓÐr   c                  ó(   ‡ — e Zd Zdˆ fd„Zdd„Zˆ xZS )ÚFakeIndentedBufferc                ó"   •— t         ‰| �  «        y rŽ   )rÌ  rd  râ  s    €rY   rd  zFakeIndentedBuffer.__init__F  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     rY   r÷  z#FakeIndentedBuffer.__getattribute__I  s;   € Ø�;ÒÜ×*Ñ*¨4°Ó6Ð6ÜØ! $ ð (=ð =ó
ð 	
r   rÄ  )rÛ   r&  r•   r   )r–   r—   r˜   rd  r÷  ræ  rç  s   @rY   ró  ró  E  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     rY   Úrestore_stdout_stderrrý  T  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
)rj  z.A line that can be 'unwritten' at a later timec                ó6   — |j                  «       sd}|| _        y r†  )rˆ  rc  r‰  s     rY   rd  zDeferredLineBase.__init__`  s   € Ø�z‰zŒ|ØˆDØˆ�	r   c                ó   — t         ‚)zJReturns either self.line or None to indicate the line has been 'unwritten'©r3  rs  s    rY   r  zDeferredLineBase.__call__e  ó   € ä!Ð!r   c                ó   — t         ‚)z3Returns a new deferred line with the same conditionr  r‰  s     rY   Ú	_new_linezDeferredLineBase._new_linei  r  r   c                ó@   — | j                  |› | j                  › �«      S rŽ   ©r  rc  )r  r´  s     rY   r‡  zDeferredLineBase.with_prefixm  s   € Ø�~‰~  ¨¯©¨Ð4Ó5Ð5r   c                óT   — | j                  | j                  j                  «       «      S rŽ   )r  rc  r  rs  s    rY   r  zDeferredLineBase.lstripp  s   € Ø�~‰~˜dŸi™i×.Ñ.Ó0Ó1Ð1r   c                ó>   — | j                  | j                  |   «      S rŽ   r  )r  r+  s     rY   r5  zDeferredLineBase.__getitem__s  s   € Ø�~‰~˜dŸi™i¨Ñ.Ó/Ð/r   c                ó,   — t        | j                  «      S rŽ   )rX  rc  rs  s    rY   r~  zDeferredLineBase.__bool__v  s   € Ü�D—I‘I‹Ðr   c                ó,   — t        | j                  «      S rŽ   )rT   rc  rs  s    rY   Ú__len__zDeferredLineBase.__len__y  s   € Ü�4—9‘9‹~Ðr   N)rc  r&  ©r•   ú
str | None)rc  r&  r•   r   )r´  r&  r•   r   )r•   r   )r+  zint | slicer•   r   rÅ  ©r•   r�   )r–   r—   r˜   r™   rd  r  r  r‡  r  r5  r~  r  r�   r   rY   rj  rj  ]  s-   „ Ù8óó
"ó"ó6ó2ó0óôr   rj  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Ì  rd  r…  Úvalue_fn)r  r…  r  rc  rÎ  s       €rY   rd  zDelayReplaceLine.__init__€  s   ø€ Ü‰Ñ˜ÔØˆŒØ ˆ�r   c                ój   — | j                   j                  | j                  | j                  «       «      S rŽ   )rc  Úreplacer…  r  rs  s    rY   r  zDelayReplaceLine.__call__…  s#   € Ø�y‰y× Ñ  §¡¨4¯=©=«?Ó;Ð;r   c                óD   — t        | j                  | j                  |«      S rŽ   )r  r…  r  r‰  s     rY   r  zDelayReplaceLine._new_lineˆ  s   € Ü §¡¨$¯-©-¸Ó>Ð>r   )r…  r&  r  zCallable[[], str]rc  r&  rÃ  )rc  r&  r•   r  )r–   r—   r˜   r™   rd  r  r  ræ  rç  s   @rY   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 usageFTrL   rv   éD   z,Not enough SMs to use max_autotune_gemm mode)Úmin_smsÚ	avail_sms)Úextra)r�   rR   r¬   rZ   r   Úcreatero  ru   ÚmajorrÓ   r  r,  Úmulti_processor_count)Úindex_or_devicer¬   Úpropr  r  s        rY   Ú
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 )NrJ   )rR   rL   rS   Úget_device_propertiesÚgpu_subslice_countrJ   r  r�   r   rY   Ú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  )rR   rJ   rS   r$  r.  r  )Údevice_propertiess    rY   Ú
using_b200r)  °  sJ   € ô �:‰:×"Ñ"Ô$ØäŸ
™
×8Ñ8¼¿¹×9RÑ9RÓ9TÓUÐØ×"Ñ" bÑ(Ð(r   c                 óÀ   — t         j                  j                  «       xr? t         j                  j                  du xr! t         j                  j                  «       dk\  S )zCReturn true for NVIDIA CUDA devices with compute capability SM100+.N)r  r   )rR   rJ   rS   ro  ru   Úget_device_capabilityr�   r   rY   Úis_nvidia_sm100_or_laterr,  º  sJ   € ô 	�
‰
×ÑÓ!ò 	:Ü�M‰M×Ñ Ð%ò	:ä�J‰J×,Ñ,Ó.°'Ñ9ð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   )rR   rL   rS   r&  r!  Ú_get_sm_carveout_experimental)Úcarveouts    rY   Úget_num_smsr0  Ã  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.r9   )r:   ÚWorkspaceZeroModeF)rl  Ú	zero_moder¬   Ú
outer_name)Úcodegen.commonr:   r2  r0  Ú	from_boolÚTMA_DESCRIPTOR_SIZEÚunique_name)Únum_tma_descriptorsr¬   Únum_programsr:   r2  r3  rH  s          rY   Ú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   Úgfx942rv   r9   rš  )rR   ro  ru   rJ   r$  ÚgcnArchName)Úblock_ks    rY   Ú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"  )rO  Úallowed_layout_dtypess     rY   Ú_use_template_for_gpurD  è  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 ©Nrd  )r  rl   Úmax_autotune_gemm_backendsr‚  rˆ  ©ÚbackendrV   s     rY   Ú_use_autotune_backendrJ  ø  ó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 rF  )r  rl   Úmax_autotune_conv_backendsr‚  rˆ  rH  s     rY   Ú_use_conv_autotune_backendrO  þ  rK  rL  c                ó¾   — | j                  «       t        j                  j                  «       j                  d«      D �cg c]  }|j	                  «       ‘Œ c}v S c c}w rF  )r  rl   Ú%max_autotune_conv_bwd_weight_backendsr‚  rˆ  rH  s     rY   Ú%_use_conv_bwd_weight_autotune_backendrR    sP   € Ø�=‰=‹?ä×=Ñ=×CÑCÓE×KÑKÈCÓPöàð 	
�‰�	òð ð ùò rL  c                ó¾   — | j                  «       t        j                  j                  «       j                  d«      D �cg c]  }|j	                  «       ‘Œ c}v S c c}w rF  )r  rl   Ú$max_autotune_conv_bwd_input_backendsr‚  rˆ  rH  s     rY   Ú$_use_conv_bwd_input_autotune_backendrU    sP   € Ø�=‰=‹?ä×<Ñ<×BÑBÓD×JÑJÈ3ÓOöàð 	
�‰�	òð ð ùò rL  )Ú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 )Nr9   )ÚBackendFeatureÚhas_backend_featurer)  ÚTRITON)r5  rZ  r[  rR   rÃ   rl  rn  r{  rp  rb  rc  rB  r¬   r,  rD  r«   rl   Úmax_autotuneÚmax_autotune_gemmrJ  ÚTRITON_TEMPLATES)rO  rV  rW  rX  rZ  r[  Úlayout_dtypess          rY   Úuse_triton_templatera    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_devicer9   r”  c                óX   •— ‰j                   j                  j                  | t        «      S rŽ   )r˜  r™  Ústatically_known_multiple_ofÚTMA_ALIGNMENT)Ú
expr_bytesr•  s    €rY   Ú_alignedzcan_use_tma.<locals>._alignedD  s    ø€ Ø�w‰w×Ñ×<Ñ<¸ZÌÓWÐWr   c                óŒ   •— | €y| j                   }| j                  }| j                  } ‰| j                  «      sy ‰|||«      S )NTF)rH  rW  r«   r   )rO  rF  Ústridesr«   rk  Ú_is_tma_compatibles       €€rY   Ú_is_tma_compatible_layoutz.can_use_tma.<locals>._is_tma_compatible_layoutG  sF   ø€ Øˆ>ØØ—‘ˆØ—-‘-ˆØ—‘ˆñ ˜Ÿ™Ô&Øá! %¨°%Ó8Ð8r   c                ó  •— | j                  «       }| j                  «       }| j                  «       }| j                  «       ‰j                  j
                  v ry| j                  «       x}�|j                  dk(  r
 ‰|||«      S  ‰|||«      S )NFrL   )Úget_sizeÚ
get_strideÚ	get_dtyperò  r˜  Úunaligned_buffersÚ
get_devicer,  )r  rF  rm  r«   Úm_devicer•  rn  Ú_is_tma_compatible_xpus        €€€rY   Ú_is_tma_compatible_matrixz.can_use_tma.<locals>._is_tma_compatible_matrixT  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 )Nr9   r¯   Fr   é    T)
rT   Úitemsizerx   r˜  r™  Úguard_int_seqÚ!replace_backed_symbols_with_hintsrK  Ústatically_known_equalsÚstatically_known_geq)rF  rm  r«   Úrankr{  Úsizes_iÚ	strides_iré   Ústrè   r  Ú	inner_idxÚ	inner_dimr•  rk  rd  s                €€€rY   rn  z'can_use_tma.<locals>._is_tma_compatibleb  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µ   r9   Fl   ÿÿ T)r˜  r™  r}  r~  Ústatically_known_gt)	rF  rm  r«   Úlast_strideÚlast_stride_hintÚ
MAX_UINT32rH  Ú	size_hintr•  s	           €rY   rw  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  rx  s     €rY   r½   zcan_use_tma.<locals>.<genexpr>¯  s   øè ø€ Ò?°Ñ)¨!×,Ñ?ùó   ƒ)rj  úint | sympy.Exprr•   rX  )rO  úLayout | Noner•   rX  )r  rB   r•   rX  )rF  úSequence[sympy.Expr]rm  zSequence[_IntLike]r«   útorch.dtyper•   rX  )Útorch.utils._tritonrf  r—  r•  r„   )
rc  rd  Úmatricesrf  ro  r•  rk  rn  rx  rw  s
    `   @@@@@rY   Ú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 )NFTr9   r”  r   )rR   Úiinfor{  rÉ   r�   r�   r‚   r�   rÀ  ÚLer—  r•  rT   ÚAndr˜  r™  Úguard_or_falseÚstatically_known_true)rF  rd  Ú	int32_maxÚ
conditionsrH  r•  Ú	conditions          rY   Ú_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   ©rd  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)rš  N)rT   rq  )r»   r  s     rY   r½   z*use_triton_tma_template.<locals>.<genexpr>Ò  s    è ø€ Ò8¨!Œs�1—:‘:“<Ó  AÕ%Ñ8ùs   ‚&(c              3  óT   •K  — | ]  }t        |j                  «       ‰¬ «      –— Œ! y­w)rŸ  N)rž  rq  )r»   r  rd  s     €rY   r½   z*use_triton_tma_template.<locals>.<genexpr>Ô  s)   øè ø€ ò àô 	(¨¯
©
«À×LÐLñùs   ƒ%(rŸ  Trb  )rl   r  Úenable_persistent_tma_matmulr„   Úenable_template_tma_storerž  rH  rR   ro  ru   r”  )rc  rd  r“  rO  s    `  rY   Úuse_triton_tma_templater¥  Í  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 )Nrb  Fr   )Ú%has_triton_tensor_descriptor_host_tmar9   ©Úis_datacenter_blackwell_arch)r¥  r’  r§  Úcodegen.cuda.cuda_envr©  )rc  rd  r“  r§  r©  s        rY   Ú!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      rY   Úuse_triton_scaling_templater°  õ  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   rY   Úensure_cute_availabler¸  ý  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)rµ  r¶  r·  r%  Ú_ensure_fp4_dtype_registered)Ú	availables    rY   Ú"ensure_nv_universal_gemm_availabler¾  
	  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Ž   )rR   rj  ÚFloat4E2M1FN)r«   Ú_origr³  s    €€rY   Ú_patchedz._ensure_fp4_dtype_registered.<locals>._patched*	  s'   ø€ Øœ×.Ñ.Ò.Ø×+Ñ+Ð+Ù˜“<Ðr   )Úcutlass_api.utilsrþ  Úcutlass_type_from_torch_typerR   rj  ÚKeyErrorr
  r³  )r»  rÃ  rÂ  r³  s     @@rY   r¼  r¼  	  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.
    ÚnvMatmulHeuristicsNFr´  r�   r   rY   Ú$ensure_nvmatmul_heuristics_availablerÉ  2	  s4   € ðÜ�~‰~×'Ñ'Ð(<Ó=ÀTÐIÐIøÜò Ùðúr¹  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ÚCUTEDSLr9   r¨  )rc  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrŽ   )Ú
is_dynamic©r»   rV   s     rY   r½   z3use_blackwell_cutedsl_grouped_mm.<locals>.<genexpr>t	  s   è ø€ Ò
1˜QŒ:�a�=Ñ
1ùrÑ  T)r¸  rJ  rª  r©  rB  r¬   r,  rR   rl  rD  rl   r]  r^  r”  r£  )
Úmat_aÚmat_brO  Úa_is_2dÚb_is_2dÚoffsÚbiasÚscale_resultr©  r`  s
             rY   Ú use_blackwell_cutedsl_grouped_mmrÖ  B	  s­   € ô2 !Ô"Øä  Ô+ØåCä�&—-‘-×$Ñ$Ô%Øá'Ô)Øä—^‘^Ð$€MÜ  ¨Ô7Øä×Ò¤6×#;Ò#;Øô �u˜e°6Õ:Øä
Ñ
1 5¨% .Ô
1Ô1Øá‘gØà€|ØàÐ˜<Ð3Øàr   c                ó<  — ddl m} |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/t(        j+                  d	t
        j                  j,                  «       y|S )
Nr9   r”  rµ   ©r?  r   F)Ú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.)r—  r•  r˜  r™  rB  rl   r³  Úcutlass_backend_min_gemm_sizeÚcodegen.cutlass.utilsrÙ  rR   ro  ru   rÃ   rl  r{  rD  r]  r^  rJ  rÓ   r  Úcutlass_dir)	rO  r  r=  Úkr•  Ú	gemm_sizerÙ  r`  rì   s	            rY   Úuse_cutlass_templaterà  ƒ	  sà   € Ýà—‘× Ñ ×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ÚNVGEMMr9   r”  rJ   c              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wrŽ   r�   )r»   Údimrã  s     €rY   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ò  r˜  rt  )r»   Útr•  s     €rY   r½   z1use_nv_universal_gemm_template.<locals>.<genexpr>æ	  s&   øè ø€ Ò
O¸ˆ1�:‰:‹<˜1Ÿ7™7×4Ñ4Ô4Ñ
Oùs   ƒ03T)r×  rã  r¸  r¾  rJ  r—  r•  Úaot_compilationr¬   r,  rR   ro  ru   rl   r]  r^  rÀ  r£  )rO  r  r=  rÞ  rÏ  rÐ  rÓ  rÂ  Údims_to_checkÚtensors_to_checkr•  rã  s             @@rY   Ú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.ÚALLTrd  )rl   r³  Úcutlass_enabled_opsr  r‚  rˆ  )Úop_nameÚenabled_opsrV   s      rY   Ú_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   r”  )Útorch._inductor.virtualizedr•  rl   r  Údecompose_k_thresholdr˜  r™  rš  r‚   r˜  ÚGerð  Úcpp_wrapperÚnum_decompose_k_splits)r  r=  rÞ  Úthreshold_multipler•  rõ  s         rY   Ú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   r”  )rl   ÚrocmÚcontiguous_thresholdrô  r•  rX  rR   ro  ru   r˜  r™  rš  r‚   r˜  rö  rð  r÷  )r  r=  rÞ  rý  r•  s        rY   Ú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)rv   rz  ry   rw   é   r   r   rš  rw   r9   rz  Ú
EXHAUSTIVE)rl   r  rø  r�   r‚   rR  Ú	is_numberr±  ÚdivisorsrÀ  Úmax_autotune_gemm_search_space)r  r=  rÞ  Ú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                  rY   Úget_k_splitsr  $
  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Ž   )rR   rJ   r$  r>  ©r¬   s    rY   Ú_rocm_native_device_arch_namer  ]
  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   rY   r  z*try_import_ck_lib.<locals>.gen_ops_librarys
  ó   € ØˆIr   c                 ó   — g S rŽ   r�   r�   r   rY   r  z.try_import_ck_lib.<locals>.gen_ops_preselectedv
  r  r   c                  ó   — e Zd Zy)ú*try_import_ck_lib.<locals>.CKGemmOperationN)r–   r—   r˜   r�   r   rY   r  r  y
  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        rY   Útry_import_ck_libr#  b
  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 )NFrJ   ú:r   z,Please pip install Composable Kernel packageT)rl   r]  r^  rR   ro  ru   r¬   r,  r  rü  Úarchr‚  rk  Úck_supported_archr«   rÃ   rl  rn  r#  rÓ   r  Úck_dir)rO  Únative_archrÞ  Úrequested_archsÚrequested_supported_archsÚck_package_dirnamerã   s          rY   Ú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 )Nr9   r”  ÚCKrµ   rØ  r   ©r—  r•  rJ  r-  r˜  r™  rB  ©rO  r  r=  rÞ  r•  s        rY   Úuse_ck_gemm_templater2  ¥
  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 )Nr9   r”  ÚCKTILErµ   rØ  r   r0  r1  s        rY   Úuse_ck_tile_gemm_templater5  ¯
  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/  )rO  r-  ©rO  s    rY   Úuse_ck_conv_templater8  ¹
  s   € Ü% dÓ+ÒG´ÀÓ0GÐGr   c                ó|   — t         j                  xs t         j                  xr | j                  j                  dk(  S rÏ  )rl   r]  r^  r¬   r,  r7  s    rY   Ú_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 )Nr9   )rC   é   rš  F)Úrequire_constant_mat2)r!  rC   r�   rO  rH  rW  r:  rs  rR   rn  rT   Úuse_cpp_gemm_templateÚis_contiguous)rO  Úmat1Úmat2rC   Ú	mat1_sizeÚmat1_strideÚmat1_each_batch_is_contiguouss          rY   Úuse_cpp_bmm_templaterE  Ã
  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 )Nr9   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µ   r9   )Úfreeze_layoutrr  ©rV   s    rY   Úis_last_dim_stride1z2use_cpp_gemm_template.<locals>.is_last_dim_stride1  s"   € Ø	�‰ÔØ�|‰|‹~˜bÑ! QÑ&Ð&r   )rV   rB   r•   rX  )r"  r!  Úcodegen.cpp_micro_gemmrG  Úcodegen.cpp_utilsrH  Úkernel.mm_commonrI  r:  rJ  rl   ÚcppÚweight_prepackrs  rR   r}  rr  rn  rl  Úhalfr«   Úhas_free_symbolsr�   ÚBaseViewÚunwrap_viewÚparallel_num_threadsrJ  Úis_module_buffer)rO  r@  rA  rL  r=  Úis_woq_int4rT  r!  rG  rH  rI  Ú	int8_gemmr`  r  r=  rÞ  rQ  rã   rN  rX  s                       rY   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)rl   r]  r^  rJ  r�   r   rY   Úuse_aten_gemm_kernelsrh    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Û  rj  Úcounterr¡   rs  s    rY   rd  zDebugDirManager.__init__&  s   € Ü”×.Ñ.Ó/ˆ�r   c                óè   — t         j                  j                  j                  | _        | j                  › d| j
                  › �| _        | j                  t         j                  j                  _        y )NÚ_tmp_)rR   Ú_dynamorl   Údebug_dir_rootrk  r¡   Únew_namers  s    rY   Ú	__enter__zDebugDirManager.__enter__)  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*  rr  rk  rR   rp  rl   rq  )r  r‡   s     rY   Ú__exit__zDebugDirManager.__exit__.  s*   € Ü�‰�d—m‘mÔ$Ø.2×.BÑ.BŒ�‰×ÑÕ+r   NrÄ  )r‡   rö  r•   r  )
r–   r—   r˜   r‘  rl  rm  r§   rd  rs  ru  r�   r   rY   rj  rj  "  s(   … Øˆi�o‰o˜aÓ €GØÓó0ó<ô
Cr   rj  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)Nr9   r>   c                ó(   •— ‰j                  | «       y rŽ   )rú   ©ÚcodeÚsource_codess    €rY   Úsave_output_codez*run_and_get_code.<locals>.save_output_code<  s   ø€ Ø×Ñ˜Õr   r{  ©ry  r&  r•   r  )
r˜  r?   r!   r   Úpatchrö  rR   rp  Úresetr„  )rÝ   r‡   rË  r?   r{  r  rz  s         @rY   Úrun_and_get_coder  3  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<  éýÿÿÿ)
rU   r  rl   r÷  r  Úautotune_at_compile_timerp  rÙ   ÚfindallÚDOTALL)	rÝ   r‡   rË  r�  r  rz  Úkernelsry  rø  s	            rY   Úrun_and_get_kernelsr‡  E  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?  r0  )r  rÝ   s    €rY   Úrun_with_backwardz1run_fw_bw_and_get_code.<locals>.run_with_backwardY  s!   ø€ Ù“ˆØ�
‰
‹×ÑÔØˆr   )r•   r   )r  )rÝ   rŠ  s   ` rY   Úrun_fw_bw_and_get_coder‹  X  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.r9   r>   c                ó(   •— ‰j                  | «       y rŽ   ©rÀ  rx  s    €rY   r{  z"get_code.<locals>.save_output_codeg  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�   rs  s    rY   rd  zIget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.__init__n  s   € Ør   c                 ó   — y rŽ   r�   r   s      rY   ÚcallzEget_code.<locals>.patched_compile_to_module.<locals>.DummyModule.callq  s   € àr   NrÄ  ©r‡   r   rË  r   r•   r  )r–   r—   r˜   r™   rd  r”  r�   r   rY   ÚDummyModuler‘  k  s   „ ÙFóôr   r–  )r÷  Úcodegen_with_cpp_wrapperÚcodegenr’   )r  r–  Úwrapper_codeÚkernel_coder{  s       €rY   Úpatched_compile_to_modulez+get_code.<locals>.patched_compile_to_modulej  s]   ø€ ÷	ñ 	ð 04×/?Ò/?ˆD×)Ñ)Ô+ÀTÇ\Á\Ã^ñ 	"ˆ�kñ 	˜×+Ñ+Ô,ÙÙ˜[×.Ñ.Ô/á‹}Ðr   Úcompile_to_moduler{  Nr|  )r  r?   r•   r   )r˜  r?   r   r}  rö  rR   rp  r~  )rÝ   r‡   rË  r?   r›  rã   r{  rz  s         @@rY   Úget_coder�  a  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 ©Nr9   rš  z%expected one or two code outputs got r   )r�  rT   )rÝ   r‡   rË  rz  s       rY   Úget_triton_coder   Œ  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 rŸ  )r  rT   )rÝ   r‡   rË  rã   rz  s        rY   Úrun_and_get_triton_coder¢  –  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>   rF   c                 ó^   •—  ‰| i |¤Ž | d   }t        |‰«      sJ ‚‰j                  |«       y )Nrš  )r�   rÀ  )r‡   rË  r˜  r?   Úgraph_loweringsÚ	real_inits      €€€rY   Ú	fake_initz-run_and_get_graph_lowering.<locals>.fake_init«  s7   ø€ Ù�4Ð"˜6Ò"Ø�Q‘ˆÜ˜% Ô/Ð/Ð/Ø×Ñ˜uÕ%r   rd  r•  )Útorch._inductor.graphr?   Útorch._inductor.output_coderG   rd  r   r}  rö  )	rÝ   r‡   rË  rG   r§  r  r?   r¥  r¦  s	         @@@rY   Ú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._inductorr¬  Ú	loweringsr2  Úpartial)Úaten_opÚoverride_fnr¬  Úorig_fns       rY   Úoverride_loweringr³  ·  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ñ  rh  Úoutr²  Úpost_fnÚpre_fns      €€€rY   r  z(add_scheduler_init_hook.<locals>.wrapperÔ  s+   ø€ Ùˆy˜%Ô Ù�i Ó'ˆÙÙ�I˜uÔ%Øˆ
r   rd  )rñ  r   rh  r   r•   r   )Útorch._inductor.schedulerrµ  rd  Úunittestr   r}  rö  )r¹  r¸  rµ  r  r²  s   ``  @rY   Úadd_scheduler_init_hookr¼  É  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)rl   Údeveloper_warningsrÓ   r  Úinfo)Úmsgs    rY   Údeveloper_warningrÁ  Þ  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 to handle:
    1. --only model_name
    2. --only=model_name
    z--onlyr9   r   r¡  z--only=N)rî  Úargvr+  rT   Ú
ValueErrorrð  )rv  r¿  s     rY   Ú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©r9   Nr�   rÎ  s     rY   r½   zis_ones.<locals>.<genexpr>  ó   è ø€ Ò%˜!ˆq�A�vÑ%ùó   ‚©r„   ©rl  s    rY   Úis_onesrÍ    ó   € ÜÑ%˜uÔ%Ó%Ð%r   c                ó&   — t        d„ | D «       «      S )Nc              3  ó&   K  — | ]	  }|d k(  –— Œ y­w)r   Nr�   rÎ  s     rY   r½   zis_zeros.<locals>.<genexpr>  rÉ  rÊ  rË  rÌ  s    rY   Úis_zerosrÑ    rÎ  r   c                ó&   — t        d„ | D «       «      S )Nc              3  ó”   K  — | ]@  }t        |t        j                  «      r$|j                  t        j                  d «      k(  –— ŒB y­w)r)  N)r�   rR   rÅ  r¬   )r»   rÒ   s     rY   r½   z is_cpu_device.<locals>.<genexpr>  s8   è ø€ ò àÜ�dœEŸL™LÔ)ð 	�‰”u—|‘| EÓ*Õ*ñùs   ‚AArË  )Úinputss    rY   Úis_cpu_devicerÕ    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‚   rR  r›   rR   r|  rp  r�  s    rY   Úget_sympy_Expr_dtyper×    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Ž   )rR   rÊ   rË   )Úshould_profiler‡   rË  rç   s       rY   Úmaybe_profilerÚ  %  sE   è ø€ áÜ�^‰^×#Ñ# TÐ4¨VÑ4ð 	¸ØŠG÷	ð 	ô 	÷	ð 	üs   ‚"A¤7©A·A ¼Ac                 ól   — t         j                  j                  } | dk  rt        j                  «       } | S ©Nr9   )rl   r\  ÚthreadsrR   Úget_num_threads)rÝ  s    rY   rb  rb  .  s+   € Ü�j‰j× Ñ €GØ�‚{Ü×'Ñ'Ó)ˆØ€Nr   c                 ó€   — ddl m}   | «       }|j                  dt        j                  j
                  rd«      S d«      S )Nr9   )Úget_backend_optionsÚ
num_stagesrš  r<  )Úruntime.triton_helpersrà  rj  rR   ro  ru   )rà  Úoptionss     rY   Úget_backend_num_stagesrä  5  s2   € å;á!Ó#€GØ�;‰;�|¬%¯-©-×*;Ò*; QÓCÐCÀÓCÐCr   c                ó‚  — t        | t        j                  j                  j                  j
                  dk(  ¬«      }|�|S t        j                  j                  «       st        j                  d«       y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_tf32z‡get_device_tflops: no Triton fallback available for non-CUDA devices. Returning 0.0; roofline estimates will use memory bandwidth only.r   r   )Úget_max_simd_tflopsÚget_max_tensorcore_tflops)rz   r   Ú
clock_rate)Úmax_clock_rate)r   rR   ÚbackendsrJ   ÚmatmulÚfp32_precisionrS   rÓ   r  Útriton.testingrè  ré  r+  rÃ   rl  rn  ÚinspectÚ	signatureÚ
parametersrj  Útorch._utils_internalrë  )r«   Úds_topsrè  ré  ÚSM80OrLaterrë  Úsm_clocks          rY   Úget_device_tflopsr÷  =  s¦  € ô Ø”u—~‘~×*Ñ*×1Ñ1×@Ñ@ÀFÑJô€Gð ÐØˆä�:‰:×"Ñ"Ô$Ü�‰ðPô	
ð ç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    rY   Úget_gpu_dram_gbpsrû  o  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$  rj  rý  s    rY   Úget_gpu_shared_memoryr   v  s.   € Ý%à�=‰=×Ñ×4Ñ4°QÓ7×;Ñ;Ð<LÈaÓPÐPr   c                 ó   — t         j                  j                  «       rht        j                  dt         j                  j	                  «       «      } t        j                  | «      }|j                  }|j                  }|€J ‚||z  S d}d}||z  S )NrJ   rz  i   )	rR   rJ   rS   r¬   r.  r   r  Úwarp_size_or_defaultÚmax_threads_per_block)r¬   ÚpropsÚ	warp_sizer  s       rY   Úget_max_numwarpsr  |  s‡   € Ü‡z�z×ÑÔ Ü—‘˜f¤e§j¡j×&?Ñ&?Ó&AÓBˆÜ ×'Ñ'¨Ó/ˆØ×.Ñ.ˆ	Ø %× ;Ñ ;ÐØ$Ð0Ð0Ð0ð
 ! IÑ-Ð-ð ˆ	Ø $ÐØ  IÑ-Ð-r   c                ó$   — | j                  d«      S )NÚwelford)rð  ©Úreduction_types    rY   Úis_welford_reductionr  Š  s   € Ø×$Ñ$ YÓ/Ð/r   c                ó&   — t        | «      ry| dv ryy)Nr<  )Úargmax_with_valueÚargmin_with_valueÚonline_softmax_reducerš  r9   )r  r	  s    rY   Úreduction_num_outputsr  Ž  s"   € Ü˜NÔ+ØØ	ð ñ 
ð
 àr   c                 ó0   — t        j                  «       dk(  S )NÚLinux)ÚplatformÚsystemr�   r   rY   Úis_linuxr  ›  s   € Ü�?‰?Ó Ñ'Ð'r   c                 ó(   — t         j                  dk(  S )Nrn   )rî  r  r�   r   rY   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‚   rR  r  rÎ  s     rY   r½   z#has_free_symbols.<locals>.<genexpr>¤  s)   è ø€ ÒJÀŒz˜!œUŸZ™ZÓ(Ò<°·±¨_Ó<ÑJùs   ‚35rÒ  )Úitrs    rY   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)Nr9   r   r�   Tzunexpected type for is_dynamic F)r"  r!  r�   rH  rJ  r`  rŒ  r@   r_  Úmaybe_get_sizeÚmaybe_get_striderB   Ú	TypeErrorr,  )r‡   r!  rè  s      rY   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   rY   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)Nr9   )Ú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`   r\   Ú	propagaterÛ  r˜  r   Únowr_   ÚlintÚ	recompilerm  rÓ   r¿  rÛ   )r§  rÌ  ÚinprÀ  r#  r0  Ú	before_ioÚafter_ioÚ
start_timeÚtime_elapsedrè  s              rY   Úpass_execution_and_saver4  Å  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
    r9   r   )r"  r!  r�   r�  Úis_multi_outputs_template©Ú	input_bufr!  s     rY   r6  r6  è  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
    r9   r   r   )r"  r!  r�   ÚMultiOutputrT   rÔ  r6  r7  s     rY   Ú#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 )NFr9   r   Úall_to_all_singleÚall_gather_into_tensorÚreduce_scatter_tensor)r"  r!  r�   Ú_CollectiveKernelÚ_WaitKernelÚop_overloadr,  ÚFallbackKernelr#  rR   r$  Útorchrecr=  Údefaultr>  r?  ©rŽ  r§  r!  s      rY   Úis_collectiverG    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 ©Nr9   r   )r"  r!  r,  rA  )rŽ  r!  s     rY   Úis_waitrJ  )  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     rY   r½   z&contains_collective.<locals>.<genexpr>6  s   è ø€ Ò@¨aÔ& q×)Ñ@ùrÑ  )rº  rM  r�   r£  ÚsnodesrG  rŽ  )ÚsnodeÚ	filter_fnrM  s      rY   rO  rO  /  sJ   € õ ?ä�%Ð-Ô.ÜÑ@°5·<±<Ô@Ó@Ð@ä˜Ÿ™Ó$ÒP¨)°tÐ*;Ò*O¹yÈÓ?OÐPr   c                óˆ   — ddl m} t        | |«      rt        d„ | j                  D «       «      S t        | j                  «      S )Nr   rL  c              3  ó2   K  — | ]  }t        |«      –— Œ y ­wrŽ   )Úcontains_waitrÎ  s     rY   r½   z contains_wait.<locals>.<genexpr>?  s   è ø€ Ò:¨”= ×#Ñ:ùrÑ  )rº  rM  r�   r£  rP  rJ  rŽ  )rQ  rM  s     rY   rU  rU  ;  s4   € Ý>ä�%Ð-Ô.ÜÑ:¨U¯\©\Ô:Ó:Ð:ä�u—z‘zÓ"Ð"r   c                ó¨   — ddl m} t        |t        j                  j
                  «      r|g}t        | |j                  «      xr | j                  |v S rI  )r"  r!  r�   rR   r´  rµ  rC  rB  rF  s      rY   Úis_fallback_oprW  D  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      rY   Úbuf_name_to_fused_snoder]  O  s#   € ð ˜k¨(Ñ3×?Ñ?×HÑHÓJÑKÐKr   c                 ó   — yrª  r�   ©rQ  s    rY   r«  r«  Z  r¬  r   c                ó²   —  || «      ry |j                  | «       | j                  D ].  }t        |j                  ||«      }||v rŒt	        |||||¬«       Œ0 y )N©Úcriteria_cb)rú   Úunmet_dependenciesr]  rÛ   Úfind_recursive_deps_of_node)rQ  Úcollected_node_setr[  r\  rb  ÚdepÚdefining_op_for_deps          rY   rd  rd  U  sn   € ñ �5ÔØØ×Ñ˜5Ô!Ø×'Ñ'ò 
ˆÜ5Ø�H‰H�kÐ#5ó
Ðð Ð"4Ñ4ØÜ#ØØØØØ#ö	
ñ
r   c                 ó   — yrª  r�   r_  s    rY   r«  r«  s  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ÚOUTPUTra  )rú   Úget_outputsr·  rŽ  rò  Úfind_recursive_users_of_node)rQ  re  r[  r\  rb  Úorˆ  Úuser_ops           rY   rl  rl  n  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)rš  r   )rR   Ú
_functorchrl   Úfunctionalize_rng_ops)Údynamo_gm_num_inputsÚaot_fw_gm_num_inputsÚnum_rng_seed_offset_inputss      rY   Únum_fw_fixed_argumentsru  ‹  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Ç  rW  s    rY   Úis_saved_tensorz'count_tangents.<locals>.is_saved_tensor›  sH   € à˜aŸf™fÐ$ò .Ø !§&¡&Ð(ò.à!¨¯©Ð/ò.ð   q§v¡vÐ-ð		
r   r   rÁ  r9   )rV   r7   r•   rX  )r˜  rh  r§  rÀ  r„  rÆ   rT   )Úfx_gr|  Ú	arg_countÚstatic_arg_idxsr=  s        rY   Ú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_)rK  r˜  rh  r§  rÛ   rð  rÀ  )r}  Ústatic_idxsrv  r=  s       rY   Ú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)Ú	BoxedBoolrX  r’   c                ó   — | j                   S rŽ   )r’   rs  s    rY   r~  zBoxedBool.__bool__Ã  s   € Ø�z‰zÐr   c                ó6   — t        | t        «      r	d| _        | S yrª  )r�   r†  r’   r  s    rY   ÚdisablezBoxedBool.disableÆ  s   € ä�cœ9Ô%ØˆCŒIØˆJØr   NrÅ  )rä  r   r•   zBoxedBool | bool)r–   r—   r˜   r§   r~  r  r‰  r�   r   rY   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)Nr9   r;   c                ó@   •— ‰j                  |«        ‰| |||||«      S rŽ   rŽ  )r  Úkernel_namerš  ry  ÚgpuÚcpp_definitionÚkernel_listÚorig_define_kernels         €€rY   Údefine_kernelz.collect_defined_kernels.<locals>.define_kernelÔ  s-   ø€ ð 	×Ñ˜;Ô'Ù!Ø�+˜{¨H°c¸>ó
ð 	
r   r‘  )NTN)r  r<   rŒ  r&  rš  r&  ry  r  r�  rX  rŽ  r  r•   r   )Úcodegen.wrapperr<   r‘  r   r}  rö  )r�  r<   r‘  r�  s   `  @rY   Úcollect_defined_kernelsr“  Î  s˜   ùè ø€ å5à-×;Ñ;Ðð  $ØØ%)ð
Ø"ð
àð
ð ð
ð ð	
ð
 ð
ð #ð
ð 
ö
ô 
�‰×	Ñ	Ð/°À-Ó	Pñ Û÷÷ ñ üs   „AA*ÁAÁ	A*ÁA'Á#A*c                ó   — | dz   S )NÚ__original__r�   rÇ  s    rY   Ú get_cloned_parameter_buffer_namer–  å  s   € Ø�.Ñ Ð r   c                ó   — | t         v S rŽ   )rP   r  s    rY   rB  rB  é  s   € Ø”YÐÐr   c                 ó:   — t         j                  j                  duS )z,Check if we're running on ROCm/HIP platform.N)rR   ro  ru   r�   r   rY   Úis_rocmr™  í  s   € ä�=‰=×Ñ DÐ(Ð(r   c                óè   — t        «       sy 	 ddlm} ddlm}  || ||«      } ||«      j                  i «      }t        |dd «      }|€y t        |«      S # t        t        t        t        f$ r Y y w xY w)Nr   )Ú	GPUTarget)Úmake_backendÚsupported_fp8_dtypes)r$   Útriton.backends.compilerr›  Útriton.compiler.compilerrœ  Úparse_optionsr
  r%  rÆ  r  rQ   rå  )Úbackend_namer&  r  r›  rœ  r³  rã  r�  s           rY   Ú_triton_supported_fp8_dtypesr¢  ò  s�   € ô ÔØðÝ6Ý9ñ ˜<¨¨yÓ9ˆÙ˜vÓ&×4Ñ4°RÓ8ˆô # 7Ð,BÀDÓIÐØÐ#ØÜÐ%Ó&Ð&øô œK¬´9Ð=ò Ùðús   �-A ÁA1Á0A1)Útriton_backendÚtriton_archr  c               óä  — | t         j                  t         j                  t         j                  t         j                  fvryt        | «      j                  d«      }|€è|€yt        j                  |«      }|j                  dk7  ryt         j                  j                  «       sy|j                  }|€t         j                  j                  «       }t         j                  j                  |«      }t        «       r!d}|j                  j!                  d«      d   }nd}|j"                  dz  |j$                  z   }t'        |dd «      }|€t)        d	«      ‚|€	|dk(  rd
nd}t+        |||«      }|€y||v S )NTr�  rJ   ru   r%  r   r  r  z0triton_arch must be provided with triton_backendry   rz  )rR   rb  rc  rd  rf  rˆ  Úremoveprefixr¬   r,  rJ   rS   r+  r.  r$  r™  r>  r‚  r  ÚminorrQ   rÄ  r¢  )	r«   r¬   r£  r¤  r  Útriton_dtypeÚdevice_indexÚ
propertiesr�  s	            rY   Úis_triton_fp8_dtype_supportedr«  	  sc  € ð Ü×ÑÜ×ÑÜ×ÑÜ×Ñð	ñ ð ä˜E“?×/Ñ/°Ó4€LàÐØˆ>ØÜ—‘˜fÓ%ˆØ�;‰;˜&Ò ØÜ�z‰z×&Ñ&Ô(Øà—|‘|ˆØÐÜ Ÿ:™:×4Ñ4Ó6ˆLÜ—Z‘Z×5Ñ5°lÓCˆ
ÜŒ9Ø"ˆNØ$×0Ñ0×6Ñ6°sÓ;¸AÑ>‰Kà#ˆNØ$×*Ñ*¨RÑ/°*×2BÑ2BÑBˆKÜ˜J¨°TÓ:ˆ	àÐÜÐKÓLÐLØÐØ(¨EÒ1‘B°rˆ	ä7Ø˜ YóÐð Ð#ØØÐ/Ð/Ð/r   c                ó&   — | dk7  xr t        | «      S )NrK   )rB  r  s    rY   Ú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)rR   rl  rJ   rS   r+  rL   r|  rX  rS  s    rY   Ú,needs_fallback_due_to_atomic_add_limitationsr¯  A  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 )NFrú   r?  r)  r9   )ÚoverloadpacketrR   r$  ÚatenÚscatter_reduce_Úscatter_reduceÚscatter_rB  r¯  rl   r\  Úfallback_scatter_reduce_sumÚdynamic_threadsrb  rX  r|  Ú$are_deterministic_algorithms_enabled)rB  r
  Ú
self_dtypeÚ	src_dtypeÚsrc_device_typeÚsrc_is_tensorÚ	reduce_tys          rY   Úuse_scatter_fallbackr¾  J  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.simdrÀ  rÁ  rº  rÂ  rÛ  rT   rK  r�   Úis_reductionrŽ  rI  Úreduction_hintr\  r]  r^  r  r,  )r%  rÀ  rÁ  rÂ  rv  rŽ  Úis_redrf  s           rY   Údump_node_schedulerÊ  q  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_offsetrU  r«   ÚGPU_ALIGN_BYTES)rÏ   rš  s     rY   Útensor_is_alignedrÎ  ’  s:   € õ Lá Ø	×	Ñ	Ó	 ¤>°&·,±,Ó#?Ñ	?Ä?ÑRÐVWÑWóð r   c                ó|   — t        | j                  j                  «      syt        j                  xs t        | «      S rª  )rB  r¬   r,  rl   Úassume_aligned_inputsrÎ  )Úexample_inputs    rY   Úshould_assume_input_alignedrÒ     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Ž   )	rR   Ú_guardsÚTracingContextÚtry_getr¤  Únullcontextr&  rš  Úsuppress_guards)Útracing_contextrš  s     rY   Ú#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)r»  r   r}  rö  rl   rR   rp  r~  r*  r  r   ÚStreamHandlerÚtorch._inductor.codecacherÜ  Ú
addHandlerÚlevelÚsetLevelr  rm  ÚremoveHandler)rÝ   r‡   rË  r*  r  Úlog_capture_stringÚchrÜ  Ú
prev_levelr  ré   s              rY   Úrun_and_get_cpp_coderæ  ¹  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Ž   )	r\   rš  r�   rR   r4   rŽ  rÅ  rH  rW  )rÔ  r&  ÚinputrH  rW  s        rY   Úshape_env_from_inputsré  Ò  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_inputsrT   rR   Ú_foreach_copy_)Ú
new_inputsÚold_tensorsÚnew_tensorsr·  Úinputs_to_checkrÔ  Úmutated_input_idxss       €€€rY   rù  z)align_inputs_from_check_idxs.<locals>.run÷  sE   ø€ Ü#9Ø˜Ð);ó$
Ñ ˆ�[ñ �JÓˆô ˆ{ÔÜ× Ñ  ¨kÔ:àˆ
r   )rî  úlist[InputType]r•   r   )rT   )rÔ  rñ  rò  rù  s   ``` rY   Úalign_inputs_from_check_idxsrô  ï  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»   rS  rW  s      rY   r½   z)clone_preserve_strides.<locals>.<genexpr>  s   è ø€ ÒT©¨°�˜‘˜fÕ$ÑTùrÑ  r9   rŒ   )rH  r?  rÐ   rW  rR   Ú
as_stridedÚclone)rV   Úneeded_sizerM  s      rY   Ú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�   rR   rÅ  r,  Údata_ptrÚ	ALIGNMENTrú  rÀ  )rî  Úcheck_inputs_idxsÚreturn_pair_idxsrï  rð  Úret_pair_definedrè   Ú_inps           rY   rì  rì    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�   rR   rÅ  rü  rý  rÀ  rT   )rÔ  Ústatic_input_idxsÚaligned_static_input_idxsrv  rè  s        rY   Úremove_unaligned_input_idxsr  1  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.
    r9   r”  Tg@Œµx¯DF)r—  r•  rR   r–  r{  rÉ   r˜  r™  Úguarding_hint_or_throwrš  Úhas_guarding_hintrl   Úassume_32bit_indexingÚ	check_leqrš  ré  )rê   r•  Úint_maxr  r  s        rY   Úexpr_fits_within_32bitr  C  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_callrš  s    €€rY   Úmap_exprz4set_tracing_context_output_strides.<locals>.map_expr�  s7   ø€ Ø Ð(Ü" 1›v˜Ù(Ø(×<Ñ<¸QÓ?Ð?Ø$×5Ñ5°aÓ8Ð8r   c              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wrŽ   r�   )r»   rê   r  s     €rY   r½   z5set_tracing_context_output_strides.<locals>.<genexpr>‰  s   øè ø€ Ò5¨!™( 1Ÿ+Ñ5ùr�  )rê   r   r•   z)float | int | SymInt | SymFloat | SymBool)
rR   rÔ  rÕ  rÖ  Úoutput_stridesrT   ré  rÀ  r  rå  )rÕ  Úcompiled_graphrX  rP  r  r  r  rš  s        @@@rY   Ú"set_tracing_context_output_stridesr  p  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)
rl   Úfx_graph_remote_cacheÚ	is_fbcoderR   Ú_utils_internalÚis_fb_unit_testÚtorch._inductor.fb.remote_cacher  ÚModuleNotFoundErrorÚjustknobs_getval_intr  s    rY   Ú should_use_remote_fx_graph_cacher!  �  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    rY   Únormalize_namer$     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_mappingrj  )r«   Útriton_type_names     rY   Útriton_typer*  ¶  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_mappingrj  r  rQ   rR   r�   r«   )r«   Úadjusted_typeÚ	type_namerK  s       rY   Útriton_type_to_torchr/  ¼  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_mkldnnrH  rW  r«   r¬   Úuntyped_storagerü  rÌ  ©rI  r’   s     rY   Úis_same_tensorr4  Ä  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Ž   )r1  rH  r«   r¬   rR   r$  Úmkldnnrü  r3  s     rY   Úis_same_mkldnn_tensorr7  Ð  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)r³  ÚisnanÚlogical_notÚlogical_andÚsignbitÚand_ÚleÚltÚgeÚgtÚeqÚner#  Úxorr�   r�   r   rY   Úboolean_opsrE  Ú  s   € ðr   c                  ó"   — e Zd ZU ded<   ded<   y)ÚOpDtypeRuler5   Útype_promotion_kindútorch.dtype | NoneÚoverride_return_dtypeNrY  r�   r   rY   rG  rG  î  s   … à8Ó8Ø-Ô-r   rG  zdict[str, OpDtypeRule]Úop_dtype_propagation_rulesc                ó*   — t        ||«      t        | <   y rŽ   )rG  rK  )rÛ   rH  rJ  s      rY   Ú#register_op_dtype_propagation_rulesrM  ÷  s   € ô
 (3ØÐ2ó(Ô˜tÒ$r   zOrderedSet[str]Úop_requires_libdevice_fp64c                ó.   — t         j                  | «       y rŽ   )rN  rú   rÇ  s    rY   Ú#register_op_requires_libdevice_fp64rP    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   r”  r)  rK   rL   Útpu)
rô  r•  r˜  Úget_current_device_or_throwr,  rl   Úcpu_backendÚxpu_backendÚtpu_backendÚcuda_backend)rØ   r•  s     rY   Úget_current_backendrX    ss   € Ý-áØ—g‘g×9Ñ9Ó;×@Ñ@ˆØ�eÒÜ×!Ñ!Ð!Ø	˜Ò	ØØ	˜Ò	Ü×!Ñ!Ð!Ø	˜Ò	Ü×!Ñ!Ð!ä×"Ñ"Ð"r   c                óx   — | �8| j                   dk(  r)t        j                  j                  | «      j                  S y)z+Check if the given device supports float64.rL   T)r,  rR   rL   r$  Úhas_fp64r  s    rY   Údevice_supports_fp64r[    s1   € àÐ˜fŸk™k¨UÒ2Ü�y‰y×.Ñ.¨vÓ6×?Ñ?Ð?Ø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  )rR   rÃ   rl  rl   r  Úcodegen_upcast_to_fp32rX  rn  rS  s    rY   Ú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  re  s     rY   rd  zScopedDict.__init__7  s   € Ø*ˆÔØ13ˆ�r   c                óZ   — || j                   v r| j                   |   S | j                  |   S rŽ   ©rf  re  r  s     rY   r5  zScopedDict.__getitem__;  s.   € Ø�$—.‘.Ñ Ø—>‘> #Ñ&Ð&Ø×!Ñ! #Ñ&Ð&r   c                ó"   — || j                   |<   y rŽ   )rf  )r  r…  r’   s      rY   Ú__setitem__zScopedDict.__setitem__@  s   € Ø#ˆ�‰�sÒr   c                ó>   — || j                   v xs || j                  v S rŽ   rh  r  s     rY   Ú__contains__zScopedDict.__contains__C  s!   € Ø�d—n‘nÐ$ÒA¨¨t×/AÑ/AÐ(AÐAr   Nc                ót   — || j                   v r| j                   |   S | j                  j                  ||«      S rŽ   )rf  re  rj  )r  r…  rE  s      rY   rj  zScopedDict.getF  s6   € Ø�$—.‘.Ñ Ø—>‘> #Ñ&Ð&Ø×!Ñ!×%Ñ% c¨7Ó3Ð3r   c                óz   — t        | j                  «      }| j                  D ]  }|| j                  vsŒ|dz  }Œ |S rÜ  )rT   re  rf  )r  r=  rÞ  s      rY   r  zScopedDict.__len__K  sC   € Ü�×"Ñ"Ó#ˆØ—‘ò 	ˆAØ˜×*Ñ*Ò*Ø�Q‘‘ð	ð ˆr   c              #  ó€   K  — | j                   E d {  –—†  | j                  D ]  }|| j                   vsŒ|–— Œ y 7 Œ)­wrŽ   rd  )r  rÞ  s     rY   Ú__iter__zScopedDict.__iter__R  s@   è ø€ Ø×%Ñ%×%Ð%Ø—‘ò 	ˆAØ˜×*Ñ*Ò*Ø“ñ	ð 	&ús   ‚>’<“!>µ>c                óH   — t        | j                  xs | j                  «      S rŽ   )rX  re  rf  rs  s    rY   r~  zScopedDict.__bool__X  s   € Ü�D×&Ñ&Ò8¨$¯.©.Ó9Ð9r   c                ó   — t         ‚rŽ   r  r  s     rY   Ú__delitem__zScopedDict.__delitem__[  s   € Ü!Ð!r   )re  úMapping[KeyType, ValType])r…  r_  r•   r`  )r…  r_  r’   r`  r•   r  )r…  rö  r•   rX  rŽ   )r…  r_  rE  úValType | Noner•   ru  r  )r•   zIterator[KeyType]rÅ  )r…  r_  r•   r  )r–   r—   r˜   r™   rd  r5  rj  rl  rj  r  rp  r~  rs  r�   r   rY   rb  rb  /  s5   „ ñó4ó'ó
$óBô4ó
óó:ô"r   rb  )Úfrozen_defaultc              ó(   ‡— dˆfd„}| €|S  || «      S )Nc                ó4   •— t        j                  | d‰¬«      S )NT)Úkw_onlyrž   )ÚdataclassesÚ	dataclass)r‘   rž   s    €rY   Úwrapzir_dataclass.<locals>.wrapa  s   ø€ Ü×$Ñ$ S°$¸vÔFÐFr   )r‘   rp   r•   rp   r�   )r‘   rž   r|  s    ` rY   Úir_dataclassr}  _  s   ø€ õGð €{ØˆÙ�‹9Ðr   c                 óœ   — t         j                  j                  j                  «       } | �"| j                  r| j                  j
                  S y rŽ   )rR   rÔ  rÕ  rÖ  Úfw_metadataÚbw_donated_idxs)rÙ  s    rY   Úget_donated_idxsr�  i  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   r9   rš  r<  r_  N)r–   r—   r˜   ÚV0_NO_TRITONÚV1_COMPILERÚV2_BACKENDSÚV3_BACKENDS_TUPLEÚV4_DICTr�   r   rY   rƒ  rƒ  p  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)rµ  r¶  r·  rƒ  r„  rž  rŸ  r#  rì  Úcompilerr†  r…  rˆ  )r  s    rY   Ú#get_triton_attrs_descriptor_versionrŒ  z  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   rY   Útriton_version_uses_attrs_dictrŽ  ”  s   € Ü.Ó0Ô4P×4XÑ4XÑXÐXr   c                ó˜   — | j                  «       }t        | t        j                  j                  «      r|› d| j
                  › �n|}||fS )Nrt   )rÛ   r�   rR   r´  rµ  Ú_overloadname)r§  Úop_overload_packet_nameÚop_overload_names      rY   Úget_op_namesr“  ˜  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_kwargsrE  r‘  )Útorch.fx.operator_schemasr•  r³  r�   rR   r´  rµ  r$  r²  Ú	index_putrE  Ú
index_put_Ú_unsafe_index_putr‡   rË  r*  r«   rX  r}  )r°  r•  r³  Ú
normalizedrã   rË  rE  rv  s           rY   Ú,_fx_node_is_input_dependent_cudagraph_unsaferœ  ¢  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
    Tr‘  F)r³  r&  ÚFORBIDDEN_CUDAGRAPH_OPSr�   rR   r´  rµ  r!  r¸  Úcudagraph_unsaferº  rœ  r*  rj  r„  rå  rÅ  Ú	is_sparse)r°  r³  r‘  ÚvalsrŠ   s        rY   rØ  rØ  Æ  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
    r9   r   TFr°  N)	r"  r!  r�   ÚConditionalÚ	WhileLooprC  rA   rQ   rØ  )rŽ  r!  r°  s      rY   Ú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  rj  rl   r  Úlibfb.py.parutilr¨  r  r4  Úpathsep)r  r¨  Úruntime_pathÚlib_paths       rY   Úget_ld_library_pathr®    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.wrapperr°  r�   Úpartition_signatures)r  r°  s     rY   Ú#is_codegen_graph_partition_subgraphr³    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Ž   )rR   rò  rl   r  Ú
cudagraphsÚ&_unstable_customized_partition_wrapperr  Úgraph_partitionr�   r   rY   Ú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 )Nr9   r”  l        i   €)	r—  r•  r˜  r™  Ústatically_known_ltr  rR   r{  r|  )rH  r•  s     rY   Údtype_from_sizer»  "  sP   € Ýà‡w�w×Ñ×+Ñ+Øˆeôà
�'‰'×
Ñ
×
/Ñ
/°°hÔ
?Ü�{‰{Ðä�{‰{Ðr   )r)  rL   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)  rL   TF)rR   r$  r6  Ú_is_mkldnn_bf16_supported©rØ   s    rY   Úis_mkldnn_bf16_supportedr¿  0  ó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)  rL   TF)rR   r$  r6  Ú_is_mkldnn_fp16_supportedr¾  s    rY   Úis_mkldnn_fp16_supportedrÃ  <  rÀ  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      rY   r½   ztabulate_2d.<locals>.<genexpr>O  s$   è ø€ ÒH©4¨1¨a˜A˜a   0˜t˜W Aœ,ÑHùó   ‚rš  r9   r¡  c              3  ó6   K  — | ]  \  }}d ||› d�›d �–— Œ y­wrÇ  r�   )r»   rê   r€  s      rY   r½   ztabulate_2d.<locals>.<genexpr>T  s$   è ø€ ÒH©t¨q°!  ! Q C p 4 ¨œlÑHùrÉ  re  )rT   r&  rK  rÉ   rÀ  r4  rÐ   r?  )ÚelementsÚheadersrê   ÚwidthsÚrowrè   r™  Útotal_widths           rY   Útabulate_2drÐ  H  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!   rk  rj  )Údict1Údict2Ú
d1_defaultÚ
d2_defaultÚall_keysr…  Úvalue1Úvalue2s           rY   Ú	zip_dictsrÙ  X  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.)rj  rQ   rl   r  ©Úconfig_patchesÚconfig_nameÚconfig_valuer’   s       rY   Úpatch_configz2maybe_aoti_standalone_config.<locals>.patch_configŒ  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.)rj  rQ   rl   rÓ   r  rÝ  s       rY   Úforce_patch_configz8maybe_aoti_standalone_config.<locals>.force_patch_config—  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.)rÞ  údict[str, Any]rß  r&  rà  r   r•   r  )rj  rl   Úaot_inductor_modeÚcompile_standaloneÚcopyrR   ro  ru   Útest_configsÚuse_libtorchÚaot_inductorÚcross_target_platformÚpackage_constants_in_sor  )rÞ  rá  rã  rè  rí  rî  s         rY   Úmaybe_aoti_standalone_configrï  {  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.rå  zKwhen cross_target_platform is windows, use_mmap_weights should not be true.TFi ”5w)FF)rl   rì  Úforce_mmap_weightsÚ package_constants_on_disk_formatr  rí  r  )Úconsts_sizeÚuse_mmap_weightsÚuse_external_weightss      rY   Údetermine_aoti_mmap_flagsr÷  Ð  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   rk   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)	r­  rl   rì  Úmodel_name_for_generated_filesr�   r&  rÄ  rÙ   rÚ   )rl   Ú
model_names     rY   Ú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)   )rV   Úunbacked_onlys     rY   Ú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_PYTHONPATHrI  Ú
PYTHONHOME)r  r  rj  r«  r4  rî  r  rl   r  Ú	sysconfigÚget_path)Úenvs    rY   Ú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   rY   r  r  6  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   rY   r  r  M  s   … Ø+/€GÐ(Ô/r   r  c                ó   — | t         _        y rŽ   )r¶  r  )r  s    rY   Ú!set_customized_partition_wrappersr  _  s   € Ø5<Ô*Õ2r   c                óÜ  ‡— t        d| j                  «      }t        |t        j                  j
                  j                  «      r*|j                  |j                  |j                  «      \  }}n&g |j                  ¢|j                  ¢}|j                  }|j                  ||«      }t        |«      }t        |dd «      }t        |t        j                  j                  «      r_|j                   j"                  D �cg c]  }|j$                  rŒ|j&                  ‘Œ }}|d t)        |«       D ]  }|j+                  |d «       Œ t-        j.                  ||f«      \  }}	dd„}
|D �cg c]7  } |
|«      r+t        j                  j
                  j1                  |d¬«      n|‘Œ9 }}d	d„Šd
ˆfd„}|D �cg c]
  } ||«      ‘Œ }}t-        j2                  ||	«      \  }}||fS c c}w c c}w c c}w )NrA   rB  c                ó  — t        | t        j                  j                  j                  «      xrS t        | t        j                  j                  j
                  t        j                  j                  j                  f«       S rŽ   )r�   rR   rò  r!  rB   ÚGeneratorStateÚOpaqueObjectStaterW  s    rY   Ú_is_tensor_irz(snode_args_kwargs.<locals>._is_tensor_irv  s^   € Ü˜!œUŸ_™_×/Ñ/×6Ñ6Ó7ò 
Ä
ØÜ�_‰_×Ñ×.Ñ.´·±×0BÑ0B×0TÑ0TÐUóA
ð =
ð 	
r   T)Úreplace_symbols_with_hintsc                ó2   — t        j                  | ||¬«      S )Nrª   )rR   rÂ   )rH  r«   r¬   s      rY   Ú_tensorz"snode_args_kwargs.<locals>._tensorƒ  s   € Ü�{‰{˜4 u°VÔ<Ð<r   c                ó˜   •— t        | t        j                  «      s| S  ‰| j                  «       | j                  | j
                  «      }|S rŽ   )r�   rR   rÅ  rH  r«   r¬   )rê   r·  r  s     €rY   Ú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   )r   rŽ  r�   rR   rò  r!  rC  Úunflatten_argsrÔ  Úconstant_argsrË  Úfill_non_provided_argsr$  rQ   r´  rµ  r¤  r¥  Ú
kwarg_onlyrÛ   rT   rU   Úpytreer"   Úir_node_to_tensorÚtree_unflatten)rQ  rŽ  r‡   rË  rB  r¿  Úpositional_arg_namesÚarg_nameÚ	flat_argsÚflat_args_pytree_specr  r<  r  r  s                @rY   Úsnode_args_kwargsr$  c  sÇ  ø€ Ü� §
¡
Ó+€DÜ�$œŸ™×*Ñ*×9Ñ9Ô:Ø×*Ñ*¨4¯;©;¸×8JÑ8JÓK‰ˆ‰fà2�—‘Ð2˜t×1Ñ1Ð2ˆØ—‘ˆà×&Ñ& t¨VÓ4€DÜ�&‹\€FÜ˜$ ¨tÓ4€KÜ�+œuŸz™z×4Ñ4Ô5à +× 3Ñ 3× =Ñ =ö 
ØÀSÇ^Ã^ˆC�H‹Hð 
Ðð  
ð -¨[¬s°4«yÐ9ò 	'ˆHØ�J‰J�x Õ&ð	'ä'-×':Ñ':¸DÀ&¸>Ó'JÑ$€IÐ$ó
ð ö	ð ñ ˜Ôô 	�‰×Ñ×,Ñ,¨QÈ4Ð,ÔPàñ	ð€Ið ó=õð -6Ö6 q‘ Õ"Ð6€IÐ6Ü×(Ñ(¨Ð4IÓJ�L€Dˆ&Ø�ˆ<Ðùò? 
ùòùò  7s   Ã<GÄGÅ%<G$Æ1G)c                óÄ   — ddl m} | j                  }|j                  j                  r(|j	                  |j                  j                  dz   «      }|j                  d«      S )Nr9   r”  rã   )r‚  r¿  Úfwd_rng_stater{  rx  )r—  r•  rÛ   r˜  r¦  rð  )rf  r•  Údep_names      rY   Úis_nonfreeable_buffersr(  ‘  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_dirr0  s      rY   Úload_templater-  Ÿ  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�   rR   r´  rµ  r+  r,  rl   Úfallback_by_defaultr!   r$  r²  Ú_assert_scalarrE  Úlift_fresh_copyÚhigher_orderÚ triton_kernel_wrapper_functionalr   )rŽ  r³  Ú"skip_fallback_due_to_dynamic_shapeÚfallback_hopss       rY   Úshould_fallback_by_defaultr6  ¥  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    rY   Úis_collective_opr9  Ý  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)rl   r  Ú)torch._inductor.fb.tlx_templates.registryr<  r%  r;  s    rY   r<  r<  â  s9   € ä×ÑÔð	ÝWà(Ð(ð ˆ	øô	 ò 	ØŠIð	ús   –   .­.c                ó   — | |z   dz
  |z  |z  S )z(Round x up to the nearest multiple of y.r9   r�   )rV   r·  s     rY   Ú	_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   )r8   ÚSwizzleTyper9   rš  rw   rv   r_  rz  ©NN)Útorch.nn.functionalr8   rA  Ú
TensorWiseÚ
NO_SWIZZLErT   ÚRowWiserm   ÚBlockWise1x128ÚBlockWise128x128rR   rj  rb  r?  ÚBlockWise1x16ÚSWIZZLE_32_4_4rh  ro  ru   ÚBlockWise1x32)Úmat_sizeÚ
scale_sizeÚscale_numelÚ	mat_dtypeÚscale_dtypeÚeq_fnr8   rA  ÚK_multiplierÚexpected_numel_aÚexpected_numel_bs              rY   Ú_infer_scale_swizzle_implrU  ö  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   r9   c                ó   — | |k(  S rŽ   r�   rë  s     rY   r«  z%infer_scale_swizzle.<locals>.<lambda>`  s
   € ˜1 ™6€ r   ©rL  rM  rN  rO  rP  rQ  )rU  rS  rå  Únumelr«   )ÚmatÚscales     rY   Úinfer_scale_swizzler\  F  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   r”  r9   c                óP   •— ‰j                   j                  j                  | |«      S )z5Compare values using symbolic equality when possible.)r˜  r™  r~  )r<  r=  r•  s     €rY   Úsymbolic_eqz+infer_scale_swizzle_ir.<locals>.symbolic_eq{  s    ø€ à�w‰w×Ñ×7Ñ7¸¸1Ó=Ð=r   rš  rX  )r<  r   r=  r   r•   rX  )rô  r•  rq  r2  r3  r4  r5  rU  rT   rå  r«   )rZ  r[  Ú	transposerL  rM  rN  r_  r•  s          @rY   Úinfer_scale_swizzle_irra  d  s¯   ø€ õ .à�|‰|‹~€HØ—‘Ó!€Jñ Ø˜Q‘K ¨!¡Ð-ˆñ DN”)×"Ñ"¤8§<¡<°¸QÔ?ÐST€Kõ>ô %Ü/2°8«}ÀÒ/A�(˜1‘+˜x¨™{Ñ+ÈÐQRÉÐUVÐGWÜ˜Ó$ØØ—)‘)Ø—K‘KØôð r   rÃ  )r}   r�   r•   r�   )rŠ   r”   r•   rX  )é   éd   )rÝ   úCallable[[], Any]rÞ   r�   rß   r�   r•   r°  )rb  rc  F)
rÝ   rd  rÞ   r�   rß   r�   rô   rX  r•   r°  )
r
  úIterable[Any]r  re  ræ   r�   r  r]   r•   r°  rÅ  )r¬   ztorch.device | None | strr•   útorch.device)r8  zIterable[sympy.Expr]r•   r”   )r@  r�  rA  r�  r•   r”   )rE  r�  rF  r�  r•   r”   )r+  r”   rF  r�  r•   úlist[sympy.Expr])r8  zIterable[_T]r•   zValuesView[_T])rU  rŽ  rV  rŽ  r•   rŽ  )r…  rI  r•   r&  )r‹  zIterable[int | torch.SymInt]r•   rg  )r‘  úint | torch.SymIntr•   rŽ  )rè   rŽ  r•   rh  )r‹  zIterable[int | sympy.Expr]r•   zlist[int | torch.SymInt])r§  útorch._ops.OpOverloadr•   rX  )r»  r7   r²  z'Callable[[torch._ops.OpOverload], bool]r•   rX  )r³  r   r‡   r¼  rË  ræ  r•   z&tuple[GraphModule, list[torch.Tensor]])rJ   )r¬   r&  r•   r  )r9   rJ   )
rÔ  úCallable[..., Any]rÕ  úSequence[Any]rë   r�   r¬   r&  r•   r°  )r�   r  r  g      ð?rJ   )rÔ  rj  rÕ  rk  rë   r�   rÜ  r�   rÝ  r°  r¬   r&  r•   r°  )rä  r   rå  r&  r•   r  )rä  r   rè  r¤   r•   r  )r<  r�   r=  r�   r•   r�   )rV   zint | Sequence[int]rH  r�   r•   úSequence[int])rV   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]r5  z8Literal[True, 'torch', 'original_aten', 'inductor_node']r•   r&  )r%  rm  r  r<   r•   útuple[str, str]rŽ   )r…  zIterable[torch.fx.Node]r†  zCallable[[Any], bool] | Noner•   úOrderedSet[torch.fx.Node])r‡   zSequence[IRNode]rË  zdict[str, IRNode]r•   rp  r©  )r+  r”   r•   zValueRanges[Any])r´  r&  r•   rX  )r´  rh   rv  r�   r•   úsympy.Symbol)rÄ  rX  r•   rX  )rÛ   r&  r•   rq  )r�  r”   rÊ  zdict[sympy.Expr, Any]r•   r”   )r<  r   r•   z&TypeGuard[torch.SymInt | torch.Tensor])r‡   r   r•   rX  )rÌ  útorch.fx.GraphModuler•   ztorch.fx.Node | None)rÌ  rr  r•   r7   )rÌ  rr  r•   zOrderedSet[torch.device]rÄ  )rä  r   r•   r   )r…  r&  r’   r&  r•   r¢  )NNT)r,  zdict[str, Any] | Noner  r  r-  rX  r•   r¢  )r6  rk  r2  rX  r•   ú	list[int])rš  r-   r6  z)Sequence[int | torch.SymInt | sympy.Expr]r2  rX  r•   rs  )r«   r‘  r•   r�   rã  r¡  rÁ  )r   zint | torch.devicer•   rX  r  )r9  r�   r¬   rf  r:  ú
int | Noner•   r:   )rv   )r?  r�   r•   r�   )rO  rC   rC  zlist[torch.dtype]r•   rX  )rI  r&  r•   rX  )
rO  rC   rV  rX  rW  rX  rX  rX  r•   rX  )r“  rB   rc  r�  rd  rX  r•   rX  rÆ  )rF  r�  rd  rX  r•   rX  )r“  rB   rc  rC   rd  rX  r•   rX  )r­  r8   r®  r8   r¯  zlist[ScalingType]r•   rX  )rÏ  r   rÐ  r   rO  rC   rÑ  rX  rÒ  rX  rÓ  ú
Any | NonerÔ  ru  rÕ  ru  r•   rX  )
rO  rC   r  r�   r=  r�   rÞ  r�   r•   rX  rB  )rO  rC   r  rá  r=  rá  rÞ  rá  rÏ  rB   rÐ  rB   rÓ  zIRNode | NonerÂ  z_IntLike | Noner•   rX  )rð  r&  r•   rX  rŒ   )
r  rá  r=  rá  rÞ  rá  rù  r�   r•   rX  )r  rá  r=  rá  rÞ  rá  r•   rX  )r  rá  r=  rá  rÞ  rá  r•   rs  )r¬   r&  r•   r&  )r•   zNtuple[str | None, Callable[[], list[Any]], Callable[[], list[Any]], type[Any]])rO  rC   r•   rX  )rO  rC   r@  zReinterpretView | BufferrA  rB   r•   rX  )FTFN)rO  rC   r@  rB   rA  rB   rL  rX  r=  rX  rd  rX  rT  rt  r•   rX  )rÝ   úCallable[P, _T]r‡   r  rË  r  r•   ztuple[_T, list[str]])rÝ   rj  r•   ztuple[Any, list[str]])rÝ   rv  r‡   r  rË  r  r•   r¤   )rÝ   rv  r‡   r  rË  r  r•   r&  )rÝ   rv  r‡   r  rË  r  r•   ztuple[Any, list[GraphLowering]])r°  rj  r±  rj  r•   r¢  )r¹  rj  r¸  zCallable[..., Any] | Noner•   r   )rÀ  r&  r•   r  r  )rl  rk  r•   rX  )rÔ  zSequence[torch.Tensor]r•   rX  )r‘  r”   r•   r‘  )rÙ  rX  r‡   r   rË  r   r•   zIterator[Any])r«   r‘  r•   r°  )r
  r&  r•   rX  )r
  r&  r•   r�   )r  re  r•   rX  )
r§  rj  rÌ  r6   r/  rk  rÀ  r&  r•   r  )r8  zBuffer | Operation | Noner•   rX  )rŽ  zNode | Operation | Noner§  ztorch._ops.OperatorBase | Noner•   rX  )rŽ  zIRNode | Operation | Noner•   rX  )rQ  rH   rR  z*Callable[[BaseSchedulerNode], bool] | Noner•   rX  )rQ  rH   r•   rX  )rŽ  zOperation | Noner§  z9torch._ops.OpOverload | Collection[torch._ops.OpOverload]r•   rX  )rZ  r&  r[  ræ  r\  ræ  r•   r   )rQ  rH   re  zMutableSet[BaseSchedulerNode]r[  zdict[str, SchedulerBuffer]r\  zdict[str, BaseSchedulerNode]rb  zCallable[[Any], bool]r•   r  )rr  r�   rs  r�   r•   r�   )r}  rr  r•   r�   )r}  rr  r•   rs  )r�  r¤   r•   r¢  )rÛ   r&  r•   r&  )r¬   r  r•   rX  )r¡  r&  r&  z	int | strr  r�   r•   ztuple[str, ...] | None)r«   r‘  r¬   ztorch.device | str | Noner£  r  r¤  zint | str | Noner  rt  r•   rX  )r¬   r&  r•   rX  )r«   r‘  r•   rX  )rB  ri  r
  r  r¹  r‘  rº  r‘  r»  r&  r¼  rX  r•   rX  )r%  rn  r•   r  )rÏ   rÆ  r•   rX  )rÑ  rÆ  r•   rX  )r•   rÇ  )rÝ   rv  r‡   r  rË  r  r•   ztuple[_T, str])rÔ  úSequence[InputType]r•   zShapeEnv | None)rÔ  úCallable[[list[InputType]], _T]rñ  rl  rò  zOrderedSet[int]r•   rx  )rV   rÆ  r•   rÆ  )rî  ró  rþ  rl  rÿ  zOrderedSet[int] | Noner•   z-tuple[list[torch.Tensor], list[torch.Tensor]])rÔ  rw  r  rl  r•   rl  )rê   r”   r•   rX  )rÕ  rk  r  rG   r•   r  )r«   r‘  r•   r&  )r«   r&  r•   r‘  )rI  rÆ  r’   rÆ  r•   rX  )r•   ztuple[str, ...])rÛ   r&  rH  r5   rJ  rI  r•   r  )rÛ   r&  r•   r  )rØ   r  r•   r&  )r¬   ztorch.device | Noner•   rX  )r«   r‘  r•   r‘  )r‘   ztype[Any] | Nonerž   rX  r•   r   )r•   zlist[int] | None)r•   rƒ  )r§  ztorch._ops.OperatorBaser•   ro  )r°  útorch.fx.Noder•   rX  )rŽ  rD   r•   rX  )r  r<   r•   rX  )rH  r�   r•   r‘  )rØ   r&  r•   rX  )rË  zSequence[Sequence[T]]rÌ  zSequence[T]r•   r&  )
rÒ  rt  rÓ  rt  rÔ  ru  rÕ  ru  r•   zEGenerator[tuple[KeyType, ValType | None, ValType | None], None, None])rÞ  ræ  r•   ræ  )rô  r�   r•   ztuple[bool, bool])rV   r,   rý  rX  r•   zOrderedSet[sympy.Symbol])r•   zdict[str, str])r  ÚCUDAGraphWrapperTyper•   r  )rQ  rH   r•   z tuple[list[Any], dict[str, Any]])rf  r=   r•   rX  )rÛ   r&  r,  r1   r•   r&  )rŽ  ry  r•   rX  )r•   r¤   )rV   r�   r·  r�   r•   r�   )rL  ztuple[Any, Any]rM  ztuple[Any, ...]rN  r   rO  r‘  rP  r‘  rQ  zCallable[[Any, Any], bool]r•   útuple[Any | None, Any | None])rZ  rÆ  r[  rÆ  r•   r{  )rZ  r@   r[  r@   r`  rX  r•   r{  (½  Ú
__future__r   rf  r¤  rz  Úenumr2  rµ  rð  r*  r‘  r  r²  r4  r  r  rÙ   r)  rÜ   rî  r  r"  r´  rÒ  r»  Ú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‚   rR   Ú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#   r’  r$   ÚOPTIMUS_EXCLUDE_POST_GRADÚ#torch.fx.experimental._size_hintingr(   r×  r)   r*   r+   r,   r-   r.   r/   r0   Úpathlibr1   r2   r3   r4   Útorch._prims_commonr5   Útorch.fxr6   Útorch.fx.noder7   rC  r8   r5  r:   r’  r<   Údependenciesr=   r˜  r?   r!  r@   rA   rB   rC   rD   rE   Úoutput_coderG   rñ  rH   rI   rP   rN   rà   rZ   Útorch._dynamo.device_interfacer[   Útorch._dynamo.utilsr\   Útorch.autogradr]   Útorch.autograd.profiler_utilr^   Ú(torch.fx.passes.graph_transform_observerr_   Útorch.fx.passes.shape_propr`   Útorch.utils._sympy.functionsra   rb   rc   rd   re   rf   Útorch.utils._sympy.symbolrg   rh   Útorch.utils._sympy.value_rangesri   rj   r"  rl   Úruntime.runtime_utilsrm   rT  Ú_IS_WINDOWSÚ	getLoggerr–   rÓ   r  rp   r$  rR  Ú	VarRangesrÅ  r�   Ú	InputTypeÚgetenvÚXPU_KERNEL_FORMATÚGPU_KERNEL_BIN_EXTSrÍ  rý  ri  r7  r}  rr  r~  rt  r  r{  r|  rÃ   rl  rn  rp  rb  rc  rd  rf  rx   r§   r|   r~   r†   ÚFunctionrˆ   r{  r    rî   rö   r  ró   r'  r0  r9  rB  rI  rL  rO  rˆ  rŒ  r’  rœ  rž  r¨  r°  rÍ  rÁ   rØ  rà  ræ  ré  rì  rî  r÷  rø  rù  ÚFN_TYPErü  r  r  r  r&  r:  rƒ  r‰  r—  r¨  r±  rµ  r¹  rÀ  rÅ  rÈ  rË  rÎ  rÓ  Ú	frozensetro  ru   rž  rÙ  rß  rë  r   r  r  r  r  r¥  r  r1  Úclear_on_fresh_inductor_cacheÚclear_inductor_cachesÚfresh_inductor_cacher:  rQ  rU  rW  r[  r^  rÊ  ré  rñ  ró  rý  rj  r  r"  r&  r)  r,  r0  r;  r@  rD  rJ  rO  rR  rU  ra  r”  rž  r¥  r«  r°  r¸  r¾  r¼  rÉ  rÖ  rà  rì  rò  rá  rú  rþ  r  r  r#  r-  r2  r5  r8  r:  rE  r>  rh  rj  r  r‡  r‹  r�  r   r¢  rª  r³  r¼  rÁ  rÅ  rÍ  rÑ  rÕ  r×  rÚ  rb  rä  r÷  rû  r   r  r  r  r  r(  r_  rÍ  ÚEnumr  r4  r6  r;  rG  rJ  rO  rU  rW  r]  rd  rl  ru  r€  r„  r†  r“  r–  rB  r™  r¢  r«  r­  r¯  r¾  rÊ  rÎ  rÒ  rÚ  ræ  ré  rô  rú  rì  r  r  r  r!  r$  r(  rl  r,  Úcompiler'  r*  r/  r4  r7  rE  rG  rK  rM  rN  rP  rX  r[  r^  r_  r`  rb  r}  r�  rƒ  rŒ  rŽ  r“  rœ  rØ  r¥  r®  r³  r¸  r»  ÚSUPPORTED_MKLDNN_DEVICESr¿  rÃ  rÐ  rÙ  rï  r÷  rû  rþ  r  r  ÚPartitionFnTyperz  r  r¶  r  r$  r(  r-  r6  r8  r9  r<  r?  rU  r\  ra  )rÞ  rŠ   s   00rY   ú<module>r«     sõ  ðÞ "ã Û Û Û Û Û Û Û 	Û Û Û Û Û 	Û Û 	Û Û Û 
Û Û Û Û Û ÷÷ ñ õ Ý Ý ÷÷ ÷ ñ ÷ CÑ BÝ ã ã ß $Ð $Ý ?Ý :Ý EÝ 0Ý /ß ;Ý 2ð (Ø$ðÐ õ
 <÷õ ñ ß>Ñ>Ýç/Ñ/ÝCÝ$Ý"Ý/å,Ý5Ý!Ý$ßT×TÝ,ß=ò +€	ÙˆCƒL€ð
 ‡�óó ðõ DÝ 0Ý %Ý 2Ý KÝ 0÷÷ ÷ 8ß Då Ý =ð �l‰l˜gÑ%€à€g×Ò˜Ó!€ð  CÐ ñ ˆ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 ØØ#(ñ	Øðàðð 
ðð !ð	ð
 ôð86.Ø ð6.à"ð6.ð ð6.ð %ð	6.ð
 ô6.ðv ØØ#(ñ	NØðNàðNð 
ðNð !ð	Nð
 ôNðb ‡�óó ðôô;ô@ð
	Ø!ð	àð	ð ô	ðØðàðð ôô+ô	*ô%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ô	>ô+ôô-ñ $òð( �=Š=×ÒÐ(ð  ñ	
ð ñ+óÐ ð6Øðàôôô'ó& 
ôð< !#Ð �IÓ "ô	ôð ×Òó"ó ñ"ð( ×Òà+/ØØñ7Ø(ð7à	ð7ð ð7ð ó	7ó ñ7ðv !5Ñ Ø$Ñ Ù"Ñ ð 49÷ ð( ñ	7Øð7à	2ð7ð ð	7ð
 õ7ðt €×Ñ�QÓó7ó ñ7ô�*õ ð ×Ò÷,ð ,ó ñ,÷
`'ò `'ôFL™õ Lô^-‘Nõ -õ,
ô
™õ 
ð ×Òó@ó ñ@÷ò ô@?Ñ'õ ?ð ‡�õó ñð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ðP Øñ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ðb ‡�óó ñõ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ð  ×Ò÷ñ ó ñð ×Òôó ñö,!öö)ð
 €×Ñ�TÓð'Øð'Ø&ð'Ø36ð'àô'ó ñ'ð0 )-ñ10ð "&Ø$(Ø ò10Øð10à%ð10ð ð	10ð
 "ð10ð ð10ð 
ö10öh.ö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ò *:Ó);Ñ &ö=ö+ö\
öö"-ñL óó�ö$%ð
 ôó ñö"ð
MØðMàðMð ðMð ð	Mð
 ðMð &ðMð #öMð`Ø	ðØ*ðà"öðB ñ"Ø	ð"àð"ð ð"ð #÷	"ùói Hs   ð3{