Ë
    ÄÓ^j¢Î ã                  óš  — 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mZ d dlmZmZ d dlmZmZmZ d dlZd dlmZ d dlZd dlZd dlmc mZ d dlmZ  d dl!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- d dl.m/Z/ d dl0m1Z1m2Z2m3Z3m4Z4m5Z5 d dl6m7Z7 d dl8m9Z9 d dl:m;Z; d dl<m=Z=m>Z> d dl?m@Z@ d dlAmBZBmCZC ddlDmEZEmFZFmGZG ddlHmIZI ddlGmJZJmKZK ddlLmMZM ddlNmOZO ddlPmQZQmRZRmSZS ddlTmUZUmVZV ddlmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZd ddlemfZf dd lgmhZhmiZimjZjmkZkmlZlmmZm dd!lnmoZo dd"lpmqZq dd#lrmsZsmtZtmuZu er&d d$lmvZvmwZwmxZx d dlyZydd%lzm{Z{ dd&lGm|Z| dd'l}m~Z~ dd(lm€Z€  e�j                  e‚«      Zƒ ek«       �j                  Z…e†e�j                  ej>                  eˆe‰eŠf   Z‹e†e�j                  ej>                  eˆd)eˆf   ZŒeG�j                  elz  ZŽed*gdf   Z�e�j                    G d+„ d,«      «       Z‘d}d-„Z’d~d.„Z“dd/„Z”d0„ Z•e–eˆeŠf   Z—e†eŠej0                  z  d1f   ee—ge†eŠd1f   f   z  Z˜	 	 d€	 	 	 	 	 	 	 	 	 	 	 d�d2„Z™	 d‚	 	 	 dƒd3„Zše�j                    G d4„ d5«      «       Z› G d6„ d7«      Zœ G d8„ d9e«      Z� G d:„ d*«      Zže�j                    G d;„ d<ež«      «       ZŸe�j                    G d=„ d>ež«      «       Z e�j                    G d?„ d@ež«      «       Z¡e�j                    G dA„ dBež«      «       Z¢e�j                    G dC„ dDež«      «       Z£e�j                    G dE„ dFež«      «       Z¤ G dG„ dHež«      Z¥e�j                    G dI„ dJež«      «       Z¦e�j                    G dK„ dLež«      «       Z§e�j                    G dM„ dNež«      «       Z¨e�j                    G dO„ dPež«      «       Z©e�j                    G dQ„ dRež«      «       Zªe�j                    G dS„ dTež«      «       Z«e�j                    G dU„ dVež«      «       Z¬e�j                    G dW„ dXe¤«      «       Z­e�j                    G dY„ dZe¥«      «       Z®e�j                    G d[„ d\ež«      «       Z¯e�j                    G d]„ d^ež«      «       Z° G d_„ d`«      Z±e�j                    G da„ dbe¬«      «       Z²e�j                    G dc„ dde¬«      «       Z³e�j                    G de„ dfe¬«      «       Z´e�j                    G dg„ dhe¬«      «       Zµ G di„ dje¬«      Z¶e�j                    G dk„ dlež«      «       Z·e�j                    G dm„ dnež«      «       Z¸e�j                    G do„ dpež«      «       Z¹e�j                    G dq„ drež«      «       Zºe�j                    G ds„ dtež«      «       Z»e�j                    G du„ dvež«      «       Z¼e�j                    G dw„ dxež«      «       Z½eˆZ¾e¬e_z  Z¿ G dy„ dzei«      ZÀ G d{„ d|eÀ«      ZÁy)„é    )ÚannotationsN)ÚCallable)ÚchainÚcount)ÚAnyÚProtocolÚTYPE_CHECKING)ÚExpr©Údtype)ÚcountersÚdynamo_timedÚget_debug_dir)ÚDebugPrinterManager)ÚMultiKernelState)Ú	cache_dir)Úget_opaque_obj_reprÚis_opaque_value_type)Útrace_structured)ÚCallMethodKeyÚConvertIntKeyÚDivideByKeyÚresolve_unbacked_bindingsÚSymTypes)Ú_get_qualified_name)ÚStorageWeakRef)Ú
OrderedSet)ÚMaxÚMin)ÚSingletonInt©Úsymbol_is_typeÚSymTé   )Úasync_compileÚconfigÚir)Úoutput_code_log)ÚIRNodeÚReinterpretView)Útriton_heuristics)ÚDeviceProperties)ÚDEFAULT_STREAMÚDEFAULT_STREAM_IDXÚSTREAM_NAME_TEMPLATE)Úget_raw_stream_nameÚget_stream_name)Úcache_on_selfÚDeferredLineBaseÚDelayReplaceLineÚget_benchmark_nameÚget_dtype_sizeÚIndentedBufferÚ#is_codegen_graph_partition_subgraphÚis_using_cudagraph_partitionÚLineContextÚmake_codegen_bufferÚsympy_productÚ	sympy_strÚ
sympy_subsÚtriton_version_uses_attrs_dict)ÚVé   )ÚArgNameÚCodeGenÚDeferredLineÚPythonPrinterÚWorkspaceArgÚWorkspaceZeroMode)Úcexpr)ÚCUSTOM_EXTERN_KERNEL_CODEGEN)Ú	config_ofÚshould_unwrap_unspec_argÚsignature_to_meta)ÚIterableÚIteratorÚSequence)ÚGraphLowering)ÚExternKernel)ÚBaseSchedulerNode)ÚFxConverterzir.CommBufferTypeÚWrapperLinec                  ó@   — e Zd ZU ded<   ded<   ded<   ded<   ded	<   y
)ÚBenchmarkStorageGroupÚstrÚbuffer_nameÚintÚnbytesr   Údevicer   z&dict[str, tuple[list[int], list[int]]]ÚinputsN)Ú__name__Ú
__module__Ú__qualname__Ú__annotations__© ó    új/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/codegen/wrapper.pyrV   rV   j   s   … àÓØƒKØƒKØƒJØ2Ô2rb   rV   c                ó¨  — t         j                  j                  | «      }| j                  «       t         j                  j                  v}t         j                  j
                  j                  | j                  «       «      }| j                  «       | j                  «       t        t         j                  j                  j                  |«      «      ||fS ©N)r@   ÚgraphÚget_allocation_storage_sizeÚget_nameÚunaligned_buffersÚ	schedulerÚget_buf_streamÚget_device_or_errorÚ	get_dtyper=   ÚsizevarsÚsimplify)ÚnodeÚstorage_sizeÚ	alignmentÚstreams       rc   Úbuffer_reuse_keyrt   s   s•   € Ü—7‘7×6Ñ6°tÓ<€LØ—‘“¤q§w¡w×'@Ñ'@Ð@€IÜ�W‰W×Ñ×-Ñ-¨d¯m©m«oÓ>€Fà× Ñ Ó"Ø�‰Óô 	”!—'‘'×"Ñ"×+Ñ+¨LÓ9Ó:ØØð	ð 	rb   c                óf  — t         j                  j                  | «      }| j                  «       }t	        |t
        j                  «      sJ ‚| j                  «       | j                  «       t        t         j                  j                  j                  |«      «      |j                  |j                  fS re   )r@   rf   rg   Úget_output_specÚ
isinstancer'   ÚCommBufferLayoutrl   rm   r=   rn   ro   Úcomm_buffer_typeÚ
group_name)rp   rq   Úlayouts      rc   Úcomm_buffer_reuse_keyr|   ƒ   s‰   € ä—7‘7×6Ñ6°tÓ<€LØ×!Ñ!Ó#€FÜ�fœb×1Ñ1Ô2Ð2Ð2à× Ñ Ó"Ø�‰ÓÜ”!—'‘'×"Ñ"×+Ñ+¨LÓ9Ó:Ø×ÑØ×Ñðð rb   c                ó‚  — | j                  «       |j                  «       k7  ry| j                  «       |j                  «       k7  ryt        j                  j                  j                  t        j                  j                  | «      «      }t        j                  j                  j                  t        j                  j                  |«      «      }t        |«      t        |«      k(  sWt        j                  j                  j                  |d|z  «      r+t        j                  j                  j                  ||«      ryy)NFgffffffî?T)
rl   rm   r@   rf   rn   ro   rg   r=   Ústatically_known_geqÚstatically_known_leq)Ú	input_bufÚ
output_bufÚ
input_sizeÚoutput_sizes       rc   Úcan_match_buffer_sizer„   ‘   sí   € ð ×$Ñ$Ó&¨*×*HÑ*HÓ*JÒJØà×ÑÓ 
× 4Ñ 4Ó 6Ò6Øä—‘×!Ñ!×*Ñ*Ü	�‰×+Ñ+¨IÓ6ó€Jô —'‘'×"Ñ"×+Ñ+Ü	�‰×+Ñ+¨JÓ7ó€Kô 	�*Ó¤¨;Ó!7Ò7ô 	
�‰×Ñ×-Ñ-¨k¸4À*Ñ;LÔMÜ�G‰G×Ñ×1Ñ1°+¸zÔJààrb   c                ó  — t        | t        j                  «      r?| j                  «       }|j                  |j
                  |j                  |j                  dfS g }| }t        |t        j                  t        j                  t        j                  f«      rj|j                  «       }|€y|j                  |«       |j                  }t        |t        j                  t        j                  t        j                  f«      rŒjt        |t        j                  «      sy|D ]+  }|j                  |j                  «       j                  k7  sŒ+ y |j                  «       }|j                  |j
                  |j                  |j                  dfS )a  
    Collapse a chain of ReinterpretView <- StorageBox
    <- ReinterpretView <- StorageBox.... <- buffer wrappers if every layer
    has the same offset as the innermost (base) buffer.

    Returns:
        (size, stride, offset, dtype, collapsible: bool)
    T)NNNNF)rw   r'   ÚBufferÚ
get_layoutÚsizeÚstrideÚoffsetr   Ú	TensorBoxÚ
StorageBoxr*   ÚappendÚdata)rŽ   ÚlayÚlayoutsÚcurÚbase_lays        rc   Úcodegen_reinterpret_view_helperr“   °   s!  € ô �$œŸ	™	Ô"Ø�o‰oÓˆØ�x‰x˜Ÿ™ S§Z¡Z°·±¸DÐ@Ð@à€GØ
€CÜ
�Sœ2Ÿ<™<¬¯©¼×8JÑ8JÐKÔ
LØ�n‰nÓˆØˆ;Ø0Ø�‰�sÔØ�h‰hˆô �Sœ2Ÿ<™<¬¯©¼×8JÑ8JÐKÕ
Lô �cœ2Ÿ9™9Ô%Ø,ð ò 1ˆØ�:‰:˜Ÿ™Ó)×0Ñ0Ó0Ù0ð1ð �~‰~Ó€HØ�=‰=˜(Ÿ/™/¨8¯?©?¸H¿N¹NÈDÐPÐPrb   .c                óÚ  ‡ ‡‡‡— t        «       Šdd„Š	 d	 	 	 dˆˆfd„}ddˆ ˆˆfd„}d‰ › �} |d|› d�«       ‰r4t        j                  j                  r‰j                  j                  «       nt        j                  «       }‰j                  «       5  |5  t        j                  j                  rV|rTt        j                  j                  r:|t        j                  j                  v rt        j                  j                  |   }	nd gt        |«      z  }	t        |«      dk(  r" ||d   |	d   «      \  }
} |d	|
› �d	|› �«       nít        |«      dkD  sJ ‚t        |«      t        |«      k(  sJ ‚t        «       }t        t        |||	«      d
„ d¬«      D ]�  \  }
}}g }|j                  r:|j                  D ]+  }|dvsŒ|j!                  d|› d|j                  |   › �«       Œ- |rdj#                  |«      }nd} ||
|«      \  }
}d|› d|
› �}||v rŒ~|j%                  |«        ||d|› d|› �«       ŒŸ d d d «       d d d «       |‰j'                  «       fS # 1 sw Y   Œ#xY w# 1 sw Y   Œ'xY w)Nc                ód   — t        | t        j                  «      r| S t        j                  | «      S re   )rw   Úsympyr
   ÚInteger)Úitems    rc   Ú_convert_to_sympy_exprz@user_defined_kernel_grid_fn_code.<locals>._convert_to_sympy_exprâ   s#   € Ü! $¬¯
©
Ô3ˆtÐL¼¿¹ÀtÓ9LÐLrb   c                óø   •— ‰�t        | «      r| | fS t        ˆfd„| D «       «      }|s|}‰j                  |«      t        j                  j
                  r$‰j                  t        ˆfd„|D «       «      «      fS dfS )a'  
        This function return a tuple of two values: the first one is for the real grid
        which is used in the generated code; the second one is an example grid with
        concreate values which is used in the autotune block to run the generated
        kernels at compile time.
        Nc              3  ó.   •K  — | ]  } ‰|«      –— Œ y ­wre   ra   )Ú.0Úgr™   s     €rc   ú	<genexpr>zKuser_defined_kernel_grid_fn_code.<locals>.determine_grid.<locals>.<genexpr>ó   s   øè ø€ ÒC¸Ñ1°!×4ÑCùs   ƒc              3  óT   •K  — | ]  }‰j                  |t        |«      «      –— Œ! y ­wre   ©Úgenerate_example_arg_valueÚtype)rœ   r�   Úwrappers     €rc   rž   zKuser_defined_kernel_grid_fn_code.<locals>.determine_grid.<locals>.<genexpr>ú   s*   øè ø€ ò àð  ×:Ñ:¸1¼dÀ1»g×Fñùó   ƒ%()ÚcallableÚtupleÚcodegen_python_shape_tupler&   ÚtritonÚautotune_at_compile_time)ÚgridÚexample_gridÚ
sympy_gridr™   r£   s      €€rc   Údetermine_gridz8user_defined_kernel_grid_fn_code.<locals>.determine_gridå   s�   ø€ ð ˆ?œh tœnà˜�:ÐäÓC¸dÔCÓCˆ
ÙØ%ˆLà×.Ñ.¨zÓ:ô —=‘=×9Ò9ð ×2Ñ2Üó à!-ôó óð
ð 	
ð ð
ð 	
rb   c                óÀ   •— ‰j                  | «       ‰rJt        j                  j                  r/‰‰j                  vr ‰j
                  j                  |xs | «       y y y y re   )Ú	writeliner&   r¨   r©   Úkernel_autotune_namesÚkernel_autotune_calls)Úliner«   ÚnameÚoutputr£   s     €€€rc   r¯   z3user_defined_kernel_grid_fn_code.<locals>.writeline  sW   ø€ Ø×Ñ˜ÔáÜ—‘×6Ò6Ø˜G×9Ñ9Ñ9à×)Ñ)×3Ñ3°LÒ4HÀDÕIð :ð 7ð rb   Úgrid_wrapper_for_údef z(meta):rA   r   zreturn c                ó2   — t        | d   j                  «      S ©NrA   ©ÚlenÚkwargs©Úxs    rc   ú<lambda>z2user_defined_kernel_grid_fn_code.<locals>.<lambda>,  s   € œc ! A¡$§+¡+Ó.€ rb   T©ÚkeyÚreverse)Úmatrix_instr_nonkdimÚwaves_per_euÚkpackzmeta['z'] == z and ÚTrueúif z	: return )r˜   úint | sympy.ExprÚreturnú
sympy.Exprre   )rª   Ú
TritonGridr«   zTritonGrid | None)r²   rW   r«   ú
str | None)r7   r&   r¨   r©   r±   ÚindentÚ
contextlibÚnullcontextr@   rf   Úautotuning_gridsrº   r   ÚsortedÚzipr»   r�   ÚjoinÚaddÚgetvalue)r³   ÚconfigsÚgridsr£   Úoriginal_fxnode_namer­   r¯   Úfn_nameÚkernel_autotune_calls_indentÚexample_gridsrª   r«   ÚseenÚcÚ
guardslistÚkwargÚguardsÚ	statementr™   r´   s   `  `              @@rc   Ú user_defined_kernel_grid_fn_coderá   Ù   s’  û€ ô Ó€FóMð
 +/ð
Øð
à'ö
÷>Jð Jð " $ Ð(€GÙ��W�I˜WÐ%Ô&ñ ”v—}‘}×=Ò=ð 	×%Ñ%×,Ñ,Ô.ä×#Ñ#Ó%ð !ð
 
�‰‹ñ .LÐ6ñ .Lä�M‰M×2Ò2Ù$Ü—‘×(Ò(Ø$¬¯©×(@Ñ(@Ñ@äŸG™G×4Ñ4Ð5IÑJ‰Mà!˜F¤S¨£ZÑ/ˆMÜˆu‹:˜Š?Ù!/°°a±¸-ÈÑ:JÓ!KÑˆD�,Ù˜ ˜vÐ&¨'°,°Ð(@ÕAä�u“: ’>Ð!�>Ü�u“:¤ W£Ò-Ð-Ð-Ü$.£LˆDô *0Ü�E˜7 MÓ2Ù.Øô*ò LÑ%��a˜ð
  �
Ø—8’8à!"§¡ò W˜Ø ð )ò ð
 '×-Ñ-°°u°g¸VÀAÇHÁHÈUÁOÐCTÐ.UÕVðWñ Ø$Ÿ\™\¨*Ó5‘Fà#�FÙ%3°D¸,Ó%GÑ"��lØ! & ¨°4°&Ð9�	Ø Ñ$ØØ—‘˜Ô#Ù˜) s¨6¨(°)¸L¸>Ð%JÕKð1L÷-.L÷ .Lð` �F—O‘OÓ%Ð%Ð%÷a.Lð .Lú÷ .Lð .Lús,   ÂI!ÂD%IÆ<A7IÈ3I!ÉI	ÉI!É!I*c                ó  ‡‡‡‡‡‡‡	— t        «       Š| j                  }|r|d   }‰j                  |d¬«       ddlŠddlmŠ ddlmŠ ddlmŠ	 t        | j                  g«      Šˆˆˆˆˆˆˆ	fd	„Š ‰| «       ‰j                  «       S )
z¶
    Given a triton kernel function pointer collect the transitive closure of
    its dependencies

    epilogue_fusion: Optional[(fused epilogue node, modified kerel src code)]
    rA   T©Ústripr   N)ÚJITFunction)Ú	constexprr   c           	     ó  •— t        d„ t        j                  | j                  «      D «       «      }| j                  j                  j                  di «      }| j                  j                  j                  D �]…  }|‰v rŒ	|| j                  j                  v sŒ"| j                  j                  |   }t        |‰	«      rX‰
j                  «        ‰
j                  d«       ‰
j                  |j                  d¬«       ‰j                  |«        ‰|«       ŒŸt        ‰d«      �rWt        |‰j                  j                   j"                  «      �r,d|j                  v r‚d‰vr~|j                  j                  j                  d«      }|rWt        |d«      rK|j$                  j'                  d	«      r0‰
j                  d
|j$                  › d�«       ‰j                  d«       ‰
j                  «        ‰
j                  d«       ‰
j                  |j                  d¬«       ||j                  j(                  k7  r*‰
j                  |› d|j                  j(                  › �«       ‰j                  |«        ‰|«       �Œt        |t*        t,        t.        ‰f«      r·‰
j                  «        t        |‰«      rd|j0                  ›d�}n|›}|j                  |«      x}rKt        |t2        «      rd|j$                  › d|j(                  › �}nd|›�}‰
j                  |› |› d|› �«       n‰
j                  |› d|› �«       ‰j                  |«       �ŒÖ||v s�ŒÜ|dk7  s�Œãt        |d«      s�Œñ|j$                  j'                  d	«      s�Œt        |‰«      r!‰
j                  |› d|j4                  › �«       n:t        |d«      r.‰
j                  d
|j$                  › d|j(                  › d|› �«       ‰j                  |«       �Œˆ y )Nc              3  óR   K  — | ]  }|j                   d k(  r|j                  –— Œ! y­w)ÚLOAD_GLOBALN)ÚopnameÚargval)rœ   Úinsts     rc   rž   z^user_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse.<locals>.<genexpr>d  s(   è ø€ ò '
àØ�{‰{˜mÒ+ð �K�Kñ'
ùs   ‚%'r`   z@triton.jitTrã   Úconstexpr_functionr   r^   r¨   zfrom z import dtype as dtypez@triton.constexpr_functionú = ztl.constexpr(ú)ú: ú.Útlz = tl.r]   z import z as )r   ÚdisÚBytecodeÚfnÚ__globals__ÚgetÚ__code__Úco_namesrw   Únewliner¯   ÚspliceÚsrcrÓ   ÚhasattrÚruntimeÚjitÚConstexprFunctionr^   Ú
startswithr]   rY   rW   ÚboolÚvaluer¢   r³   )Ú
cur_kernelÚunqualified_loadsÚglobal_annotationsÚsymbol_nameÚsymbolÚdtype_symbolÚ
symbol_strÚ
annotationÚannotation_coderå   Úcompile_wrapperræ   Úsymbols_includedÚtraverser¨   Útriton_dtypes            €€€€€€€rc   r  zKuser_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse_  s¢  ø€ ô
 'ñ '
äŸ™ Z§]¡]Ó3ô'
ó 
Ðð
 (Ÿ]™]×6Ñ6×:Ñ:Ð;LÈbÓQÐØ%Ÿ=™=×1Ñ1×:Ñ:ó K	6ˆKØÐ.Ñ.ØØ˜jŸm™m×7Ñ7Ò7Ø#Ÿ™×2Ñ2°;Ñ?�Ü˜f kÔ2Ø#×+Ñ+Ô-Ø#×-Ñ-¨mÔ<Ø#×*Ñ*¨6¯:©:¸TÐ*ÔBØ$×(Ñ(¨Ô5Ù˜VÕ$Ü˜VÐ%9Õ:¼zØØ—N‘N×&Ñ&×8Ñ8õ@ð
  &§*¡*Ñ,°Ð@PÑ1PØ'-§y¡y×'<Ñ'<×'@Ñ'@ÀÓ'I˜á(Ü '¨°lÔ CØ ,× 7Ñ 7× BÑ BÀ8Ô Là+×5Ñ5Ø"'¨×(?Ñ(?Ð'@Ð@VÐ Wôð -×0Ñ0°Ô9Ø#×+Ñ+Ô-Ø#×-Ñ-Ð.JÔKØ#×*Ñ*¨6¯:©:¸TÐ*ÔBØ" f§i¡i×&8Ñ&8Ò8Ø'×1Ñ1Ø*˜m¨3¨v¯y©y×/AÑ/AÐ.BÐCôð %×(Ñ(¨Ô5Ù˜VÖ$Ü ¬¬c´4¸Ð(CÔDØ#×+Ñ+Ô-Ü! &¨)Ô4Ø'4°V·\±\Ð4DÀAÐ%F™
à(. z˜
Ø%7×%;Ñ%;¸KÓ%HÐH�zÐHÜ% j´$Ô7à"$ Z×%:Ñ%:Ð$;¸1¸Z×=PÑ=PÐ<QÐ Rñ ,ð 13°:°.Ð.A˜OØ'×1Ñ1Ø*˜m¨OÐ+<¸CÀ
¸|ÐLõð (×1Ñ1°[°MÀÀZÀLÐ2QÔRØ$×(Ñ(¨Ö5àÐ#4Ó4Ø# tÔ+Ü ¨Ö5ð ×)Ñ)×4Ñ4°XÖ>ô " &¨,Ô7Ø'×1Ñ1°[°MÀÈÏÉÀ}Ð2UÕVÜ  ¨Ô4Ø'×1Ñ1Ø# F×$5Ñ$5Ð#6°h¸v¿¹Ð>OÈtÐT_ÐS`Ðaôð %×(Ñ(¨Ö5ñWK	6rb   )r7   rü   rû   r¨   rå   Útriton.languageræ   Útriton.language.corer   r   r]   rÔ   )
ÚkernelÚepilogue_fusionÚ
kernel_srcrå   r  ræ   r  r  r¨   r  s
      @@@@@@@rc   Ú9user_defined_triton_kernel_transitive_closure_source_coder  G  s~   þ€ ô %Ó&€OØ—‘€JÙØ$ QÑ'ˆ
Ø×Ñ˜:¨TÐÔ2ó Ý"Ý)Ý:ô " 6§?¡?Ð"3Ó4Ð÷V6ò V6ñp ˆVÔØ×#Ñ#Ó%Ð%rb   c                  ó(   — e Zd ZU ded<   ded<   d„ Zy)ÚSymbolicCallArgúsympy.SymbolÚinnerrÉ   Ú
inner_exprc                ó,   — t        | j                  «      S re   )rW   r  ©Úselfs    rc   Ú__str__zSymbolicCallArg.__str__Á  s   € Ü�4—:‘:‹Ðrb   N)r]   r^   r_   r`   r  ra   rb   rc   r  r  »  s   … àÓàÓórb   r  c                  óZ   ‡ — e Zd Zˆ fd„Zdd„Zd	d„Zd
d„Zdd„Zdd„Z	 	 	 	 	 	 dd„Z	ˆ xZ
S )ÚMemoryPlanningStatec                ó¨   •— t         ‰| �  «        t        j                  t        «      | _        t        j                  t        «      | _        d| _        y ©Nr   )ÚsuperÚ__init__ÚcollectionsÚdefaultdictÚlistÚ
reuse_poolÚcomm_buffer_reuse_poolÚtotal_allocated_buffer_size)r  Ú	__class__s    €rc   r%  zMemoryPlanningState.__init__Æ  sG   ø€ Ü‰ÑÔô ×#Ñ#¤DÓ)ð 	Œô ×#Ñ#¤DÓ)ð 	Ô#ð 12ˆÕ(rb   c                óL   — t        | j                  j                  |d «      «      S re   )r  r)  r÷   ©r  rÀ   s     rc   Ú__contains__z MemoryPlanningState.__contains__Ò  s   € Ü�D—O‘O×'Ñ'¨¨TÓ2Ó3Ð3rb   c                ó\   — | j                   |   j                  «       }|j                  rJ ‚|S re   )r)  ÚpopÚ	is_reused©r  rÀ   r˜   s      rc   r1  zMemoryPlanningState.popÕ  s+   € Ø�‰˜sÑ#×'Ñ'Ó)ˆØ—>’>Ð!Ð!Øˆrb   c                ó\   — |j                   rJ ‚| j                  |   j                  |«       y re   )r2  r)  r�   r3  s      rc   ÚpushzMemoryPlanningState.pushÚ  s&   € Ø—>’>Ð!Ð!Ø�‰˜Ñ×#Ñ# DÕ)rb   c                óL   — t        | j                  j                  |d «      «      S re   )r  r*  r÷   r.  s     rc   Úcomm_buffer_containsz(MemoryPlanningState.comm_buffer_containsÞ  s    € Ü�D×/Ñ/×3Ñ3°C¸Ó>Ó?Ð?rb   c                ó\   — | j                   |   j                  «       }|j                  rJ ‚|S re   )r*  r1  r2  r3  s      rc   Úcomm_buffer_popz#MemoryPlanningState.comm_buffer_popá  s-   € Ø×*Ñ*¨3Ñ/×3Ñ3Ó5ˆØ—>’>Ð!Ð!Øˆrb   c                ó\   — |j                   rJ ‚| j                  |   j                  |«       y re   )r2  r*  r�   r3  s      rc   Úcomm_buffer_pushz$MemoryPlanningState.comm_buffer_pushæ  s*   € ð —>’>Ð!Ð!Ø×#Ñ# CÑ(×/Ñ/°Õ5rb   )rÀ   ÚReuseKeyrÈ   r  )rÀ   r<  rÈ   ÚFreeIfNotReusedLine)rÀ   r<  r˜   r=  rÈ   ÚNone)rÀ   ÚCommBufferReuseKeyrÈ   r  )rÀ   r?  rÈ   r=  )rÀ   r?  r˜   r=  rÈ   r>  )r]   r^   r_   r%  r/  r1  r5  r7  r9  r;  Ú__classcell__©r,  s   @rc   r!  r!  Å  s?   ø„ ô
2ó4óó
*ó@óð
6Ø%ð6Ø-@ð6à	÷6rb   r!  c                  ó   — e Zd Zdd„Zy)ÚHasWriteLinec                 ó   — y re   ra   ©r  r²   s     rc   r¯   zHasWriteLine.writelineî  ó   � rb   N)r²   z$LineContext | DeferredLineBase | strrÈ   r>  )r]   r^   r_   r¯   ra   rb   rc   rC  rC  í  s   „ ÜPrb   rC  c                  ó   — e Zd Zdd„Zy)rT   c                ó0   — t        dt        | «      › �«      ‚)Nz&FX codegen not yet supported for type )ÚNotImplementedErrorr¢   ©r  Ú	converters     rc   Ú
codegen_fxzWrapperLine.codegen_fxò  s   € Ü!Ð$JÌ4ÐPTË:È,Ð"WÓXÐXrb   N©rK  rS   rÈ   ÚFxConversionFunc©r]   r^   r_   rL  ra   rb   rc   rT   rT   ñ  s   „ ôYrb   c                  ó:   — e Zd ZU ded<   ded<   d	d„Zd
d„Zdd„Zy)ÚEnterSubgraphLineÚPythonWrapperCodegenr£   rP   rf   c                ób   — | j                   j                  | j                   j                  «       y re   )r£   Úpush_computed_sizesÚcomputed_sizesr  s    rc   Ú__post_init__zEnterSubgraphLine.__post_init__û  s   € Ø�‰×(Ñ(¨¯©×)DÑ)DÕErb   c                ón   — | j                   j                  | j                  «       |j                  «        y re   )r£   Úpush_codegened_graphrf   Ú	do_indent©r  Úcodes     rc   ÚcodegenzEnterSubgraphLine.codegenþ  s"   € Ø�‰×)Ñ)¨$¯*©*Ô5Ø�‰Õrb   c                ó   — |j                   S re   )Ú_generate_enter_subgraphrJ  s     rc   rL  zEnterSubgraphLine.codegen_fx  s   € Ø×1Ñ1Ð1rb   N©rÈ   r>  ©r[  r7   rÈ   r>  rM  ©r]   r^   r_   r`   rV  r\  rL  ra   rb   rc   rQ  rQ  ö  s   … à!Ó!ØÓóFóô2rb   rQ  c                  ó<   — e Zd ZU ded<   ded<   dd„Zed	d„«       Zy)
ÚConditionalLinerR  r£   zir.Conditionalrp   c                ó   — t        d«      ‚)NzOnly supports FX codegen©rI  rZ  s     rc   r\  zConditionalLine.codegen  s   € Ü!Ð"<Ó=Ð=rb   c                ó   — | j                   S re   )Ú_generate_conditional©rK  s    rc   rL  zConditionalLine.codegen_fx  s   € à×.Ñ.Ð.rb   Nr`  rM  ©r]   r^   r_   r`   r\  ÚstaticmethodrL  ra   rb   rc   rc  rc    s'   … à!Ó!Ø
Óó>ð ò/ó ñ/rb   rc  c                  ó2   — e Zd ZU ded<   dd„Zedd„«       Zy)ÚCommentLiner:   r²   c                ó:   — |j                  | j                  «       y re   )r¯   r²   rZ  s     rc   r\  zCommentLine.codegen  s   € Ø�‰�t—y‘yÕ!rb   c                ó   — | j                   S re   )Ú_generate_commentrh  s    rc   rL  zCommentLine.codegen_fx  s   € à×*Ñ*Ð*rb   Nr`  rM  ri  ra   rb   rc   rl  rl    s!   … à
Óó"ð ò+ó ñ+rb   rl  c                  ó<   — e Zd ZU ded<   ded<   dd„Zed	d„«       Zy)
ÚDynamicScalarLinerR  r£   zir.DynamicScalarrp   c                óN   — | j                   j                  | j                  «       y re   )r£   Ú_codegen_dynamic_scalarrp   rZ  s     rc   r\  zDynamicScalarLine.codegen$  s   € Ø�‰×,Ñ,¨T¯Y©YÕ7rb   c                ó   — | j                   S re   )Ú_generate_dynamic_scalarrh  s    rc   rL  zDynamicScalarLine.codegen_fx'  s   € à×1Ñ1Ð1rb   Nr`  rM  ri  ra   rb   rc   rq  rq    s'   … à!Ó!Ø
Óó8ð ò2ó ñ2rb   rq  c                  ó0   — e Zd ZU ded<   dd„Zdd„Zd	d„Zy)
ÚExitSubgraphLinerR  r£   c                óV   — | j                   j                  «       | j                   _        y re   )r£   Úpop_computed_sizesrU  r  s    rc   rV  zExitSubgraphLine.__post_init__0  s   € Ø&*§l¡l×&EÑ&EÓ&Gˆ�‰Õ#rb   c                óX   — | j                   j                  «        |j                  «        y re   )r£   Úpop_codegened_graphÚdo_unindentrZ  s     rc   r\  zExitSubgraphLine.codegen3  s   € Ø�‰×(Ñ(Ô*Ø×ÑÕrb   c                ó   — |j                   S re   )Ú_generate_exit_subgraphrJ  s     rc   rL  zExitSubgraphLine.codegen_fx7  s   € Ø×0Ñ0Ð0rb   Nr_  r`  rM  ra  ra   rb   rc   rw  rw  ,  s   … à!Ó!óHóô1rb   rw  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚEnterDeviceContextManagerLinerY   Ú
device_idxú
int | NoneÚlast_seen_device_guard_indexc                óx  — t         j                  j                  �r|j                  d«       t         j                  j                  rg| j
                  €;|j                  t         j                  j                  j                  «       › d�«       y | j
                  | j                  k(  s{J d«       ‚| j
                  €H|j                  t         j                  j                  j                  «       › d| j                  › d�«       y |j                  d| j                  › d�«       y y |j                  dt         j                  j                  j                  | j                  «      › d�«       |j                  «        |j                  t         j                  j                  j                  | j                  «      «       y )	Nú
z) stream_guard(stream, this->device_idx_);z4AOTInductor only supports running on one CUDA devicez device_guard(z);zdevice_guard.set_index(úwith ú:)r@   rf   Úcpp_wrapperr¯   Úaot_moderƒ  Ú
device_opsÚcpp_aoti_stream_guardr�  Úcpp_aoti_device_guardÚdevice_guardrY  Ú
set_devicerZ  s     rc   r\  z%EnterDeviceContextManagerLine.codegen@  sP  € Ü�7‰7×ÓØ�N‰N˜4Ô Ü�w‰w×Òð ×4Ñ4Ð<Ø—N‘NÜŸ7™7×-Ñ-×CÑCÓEÐFÐFoÐpõð  ×<Ñ<ÀÇÁÒOð ØNóÐOð ×4Ñ4Ð<Ø—N‘NÜŸ7™7×-Ñ-×CÑCÓEÐFÀnÐUY×UdÑUdÐTeÐegÐhõð —N‘NÐ%<¸T¿_¹_Ð<MÈRÐ#PÕQð Pð �N‰N˜U¤1§7¡7×#5Ñ#5×#BÑ#BÀ4Ç?Á?Ó#SÐ"TÐTUÐVÔWØ�N‰NÔØ�N‰Nœ1Ÿ7™7×-Ñ-×8Ñ8¸¿¹ÓIÕJrb   c                ó   — |j                   S re   )Ú&_generate_enter_device_context_managerrJ  s     rc   rL  z(EnterDeviceContextManagerLine.codegen_fx]  s   € Ø×?Ñ?Ð?rb   Nr`  rM  ©r]   r^   r_   r`   r\  rL  ra   rb   rc   r€  r€  ;  s   … àƒOØ",Ó,óKô:@rb   r€  c                  ó   — e Zd Zdd„Zdd„Zy)ÚExitDeviceContextManagerLinec                óZ   — t         j                  j                  s|j                  «        y y re   ©r@   rf   rˆ  r|  rZ  s     rc   r\  z$ExitDeviceContextManagerLine.codegenb  s    € Ü�w‰w×"Ò"Ø×ÑÕð #rb   c                ó   — |j                   S re   )Ú%_generate_exit_device_context_managerrJ  s     rc   rL  z'ExitDeviceContextManagerLine.codegen_fxf  s   € Ø×>Ñ>Ð>rb   Nr`  rM  )r]   r^   r_   r\  rL  ra   rb   rc   r“  r“  a  s   „ óô?rb   r“  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚExternKernelAllocLinerR  r£   úir.ExternKernelAllocrp   c                ó¬   — | j                   }g |j                  «       ¢|j                  «       ¢}| j                  j	                  | j                   |«       y re   )rp   Úcodegen_argsÚcodegen_kwargsr£   Ú$_generate_extern_kernel_alloc_helper)r  r[  rp   Úargss       rc   r\  zExternKernelAllocLine.codegeno  sD   € Ø�y‰yˆØ=�×"Ñ"Ó$Ð= t×':Ñ':Ó'<Ð=ˆØ�‰×9Ñ9¸$¿)¹)ÀTÕJrb   c                ó   — |j                   S re   )Ú_generate_extern_kernel_allocrJ  s     rc   rL  z ExternKernelAllocLine.codegen_fxt  s   € Ø×6Ñ6Ð6rb   Nr`  rM  r‘  ra   rb   rc   r™  r™  j  s   … à!Ó!Ø
ÓóKô
7rb   r™  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚExternKernelOutLinerR  r£   úir.ExternKernelOutrp   c           	     óD  — | j                   }g |j                  «       ¢|j                  d¬«      ¢}|j                  «       }t        j
                  j                  r|j                  dk(  rd}n|j                  «       }|j                  «       x}r|j                  nt        j
                  j                  }| j                  j                  ||j                  «       |j                  r|j                  j                  «       nd ||| j                   j                  «       «       y )NT)Úskip_outztorch::inductor::_mm_plus_mmÚaoti_torch__mm_plus_mm_out)rp   rœ  r�  Úget_kernel_namer@   rf   rˆ  Úcpp_kernel_nameÚ
get_devicer¢   Údevice_typer£   Ú"_generate_extern_kernel_out_helperÚcodegen_referenceÚoutput_viewÚget_stack_traces)r  r[  rp   rŸ  Úkernel_nameÚdr[   s          rc   r\  zExternKernelOutLine.codegen}  sé   € Ø�y‰yˆØJ�×"Ñ"Ó$ÐJ t×':Ñ':ÀDÐ':Ó'IÐJˆØ×*Ñ*Ó,ˆä�G‰G×ÒØ×$Ñ$Ð(FÒFð 7‰Kà×.Ñ.Ó0ˆKØ!%§¡Ó!2Ð2˜AÐ2�—’¼¿¹×9LÑ9LˆØ�‰×7Ñ7ØØ×"Ñ"Ó$Ø48×4DÒ4DˆD×Ñ×.Ñ.Ô0È$ØØØ�I‰I×&Ñ&Ó(õ	
rb   c                ó   — |j                   S re   )Ú_generate_extern_kernel_outrJ  s     rc   rL  zExternKernelOutLine.codegen_fx“  ó   € Ø×4Ñ4Ð4rb   Nr`  rM  r‘  ra   rb   rc   r£  r£  x  s   … à!Ó!Ø
Óó
ô,5rb   r£  c                  ó.   — e Zd ZU dZded<   ded<   dd„Zy)	ÚExternKernelMultiOutLinez¾Codegen line for multi-output .out() variant calls.

    Generates a kernel call with pre-allocated output buffers passed as
    keyword arguments. E.g. kernel(x, out0=buf0, out1=buf1).
    rR  r£   úir.ExternKernelMultiOutrp   c           	     ó†  — | j                   }|j                  «       }g |j                  «       ¢|j                  «       ¢}t	        |j
                  |j                  «      D ])  \  }}|j                  |› d|j                  «       › �«       Œ+ |j                  |j                  «       › d|› ddj                  |«      › d�«       y )Nú=rî   ú(ú, rï   )rp   r¨  rœ  r�  rÑ   Úout_arg_namesÚout_variant_output_nodesr�   rh   r¯   rÒ   )r  r[  rp   r°  rŸ  Úout_nameÚout_nodes          rc   r\  z ExternKernelMultiOutLine.codegen¢  s¸   € Ø�y‰yˆØ×*Ñ*Ó,ˆà=�×"Ñ"Ó$Ð= t×':Ñ':Ó'<Ð=ˆÜ"%Ø×Ñ × =Ñ =ó#
ò 	=ÑˆH�hð �K‰K˜8˜* A h×&7Ñ&7Ó&9Ð%:Ð;Õ<ð	=ð
 	�‰˜$Ÿ-™-›/Ð*¨#¨k¨]¸!¸D¿I¹IÀd»OÐ;LÈAÐNÕOrb   Nr`  ©r]   r^   r_   Ú__doc__r`   r\  ra   rb   rc   r¶  r¶  —  s   … ñð "Ó!Ø
!Ó!ô
Prb   r¶  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚFreeLinerR  r£   úBufferLike | ir.TorchBindObjectrp   c                óØ   — | j                   j                  «       t        j                  j                  vsJ ‚|j                  | j                  j                  | j                   «      «       y re   )rp   rh   r@   rf   Úremoved_buffersr¯   r£   Úmake_buffer_freerZ  s     rc   r\  zFreeLine.codegen´  sF   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�‰�t—|‘|×4Ñ4°T·Y±YÓ?Õ@rb   c                ó   — |j                   S re   )Ú_generate_freerJ  s     rc   rL  zFreeLine.codegen_fx¸  ó   € Ø×'Ñ'Ð'rb   Nr`  rM  r‘  ra   rb   rc   rÃ  rÃ  ¯  s   … à!Ó!Ø
)Ó)óAô(rb   rÃ  c                  ó¤   — e Zd ZU ded<   ded<   ded<   ded<   ded<   d	ed
<   ded<   ded<   ded<   ded<   ded<   ded<   dZded<   dd„Zdd„Zy)ÚKernelCallLinerR  r£   rW   r°  ztuple[Any, ...]Ú	call_argsÚraw_keysÚraw_argsú	list[str]Ú	arg_typesr  r¨   zdict[str, Any]Útriton_metazdict[str, Any] | NoneÚinductor_metaztorch.devicer[   Ú
graph_namer×   Nr‚  Úcurrent_stream_idxc                óB  — | j                   j                  | j                  | j                  | j                  | j
                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  ¬«       y )N)
r¨   rÑ  rÎ  rÏ  rÒ  rÓ  r[   rÔ  r×   rÕ  )r£   Ú_generate_kernel_call_helperr°  rÍ  r¨   rÑ  rÎ  rÏ  rÒ  rÓ  r[   rÔ  r×   rÕ  rZ  s     rc   r\  zKernelCallLine.codegenÌ  sw   € Ø�‰×1Ñ1Ø×ÑØ�N‰NØ—;‘;Ø—n‘nØ—]‘]Ø—]‘]Ø×(Ñ(Ø×,Ñ,Ø—;‘;Ø—‘Ø!%×!:Ñ!:Ø#×6Ñ6ð 	2õ 	
rb   c                ó   — |j                   S re   )Ú_generate_kernel_callrJ  s     rc   rL  zKernelCallLine.codegen_fxÜ  ó   € Ø×.Ñ.Ð.rb   r`  rM  )r]   r^   r_   r`   rÕ  r\  rL  ra   rb   rc   rÌ  rÌ  ¼  s]   … à!Ó!ØÓØÓØÓØÓØÓØƒLØÓØ(Ó(ØÓØƒOØÓØ%)Ð˜
Ó)ó
ô /rb   rÌ  c                  óf   — e Zd ZU ded<   ded<   ded<   dZded<   d	Zd
ed<   dZded<   dd„Zdd„Zy)ÚKernelDefinitionLinerR  r£   rW   r°  Úkernel_bodyNrË   ÚmetadataTr  ÚgpuÚcpp_definitionc                ó¨   — | j                   j                  | j                  | j                  | j                  | j
                  | j                  ¬«       y ©N)rÞ  rß  rà  )r£   Ú_define_kernel_helperr°  rÝ  rÞ  rß  rà  rZ  s     rc   r\  zKernelDefinitionLine.codegené  sB   € Ø�‰×*Ñ*Ø×ÑØ×ÑØ—]‘]Ø—‘Ø×.Ñ.ð 	+õ 	
rb   c                ó   — |j                   S re   )Ú_generate_kernel_definitionrJ  s     rc   rL  zKernelDefinitionLine.codegen_fxò  r´  rb   r`  rM  )	r]   r^   r_   r`   rÞ  rß  rà  r\  rL  ra   rb   rc   rÜ  rÜ  à  s<   … à!Ó!ØÓØÓØ€HˆjÓØ€CˆÓØ!%€N�JÓ%ó
ô5rb   rÜ  c                  ó0   — e Zd ZU ded<   dd„Zdd„Zd	d„Zy)
ÚMemoryPlanningLinerR  r£   c                ó   — | S )zFirst pass to find reusera   ©r  Ústates     rc   ÚplanzMemoryPlanningLine.planú  s   € àˆrb   c                 ó   — y)zSecond pass to output codeNra   rZ  s     rc   r\  zMemoryPlanningLine.codegenþ  rF  rb   c                ór  — g }t        j                  | «      D ]t  }|j                  dk(  rŒt        | |j                  «      }|j	                  |j                  › d|j
                  t        j                  u r|j                  «       n|› �«       Œv t        | «      j                  › ddj                  |«      › d�S )zF
        Emits a string representation that fits on one line.
        r£   r¹  rº  r»  rï   )ÚdataclassesÚfieldsr³   Úgetattrr�   r¢   r'   r†   rh   r]   rÒ   )r  rŸ  ÚfieldÚvals       rc   r  zMemoryPlanningLine.__str__  s¡   € ð ˆÜ ×'Ñ'¨Ó-ò 	ˆEØ�z‰z˜YÒ&ØÜ˜$ §
¡
Ó+ˆCØ�K‰KØ—:‘:�,˜a°%·*±*ÄÇ	Á	Ñ2I §¡¤ÈsÐSÐTõð		ô �t“*×%Ñ%Ð& a¨¯	©	°$«Ð'8¸Ð:Ð:rb   N©rê  r!  rÈ   rç  r`  ©rÈ   rW   )r]   r^   r_   r`   rë  r\  r  ra   rb   rc   rç  rç  ö  s   … à!Ó!óó)ô;rb   rç  c                  ód   ‡ — e Zd ZU dZdZded<    ej                  e¬«      Z	ded<   d	ˆ fd„Z
ˆ xZS )
Ú+EnterDeviceContextManagerWithStreamInfoLineu;  Enter a device context and retrieve user stream objects.

    Attributes:
        num_streams: Number of streams (determined by user annotations on nodes).
        stream_idx_to_user_obj_idx: Maps stream_idx â†’ user_object_index for
            retrieving user stream objects via get_external_object_by_index.
    rA   rY   Únum_streams)Údefault_factoryzdict[int, int]Ústream_idx_to_user_obj_idxc                óÂ  •— t         j                  j                  rt        ‰| �  |«       yt        ‰| �  |«       |j                  t        › dt         j                  j                  j                  «       › �«       | j                  dkD  rVt        d| j                  «      D ]<  }| j                  |   }|j                  t        j                  |¬«      › d|› d�«       Œ> yy)z5Generate context switching and stream retrieval code.rî   rA   ©Ú
stream_idxz  = get_external_object_by_index(rï   N)r@   rf   rˆ  r$  r\  r¯   r-   rŠ  Úcurrent_streamr÷  Úrangerù  r/   Úformat)r  r[  ÚiÚuser_obj_idxr,  s       €rc   r\  z3EnterDeviceContextManagerWithStreamInfoLine.codegen  sÃ   ø€ ä�7‰7×ÒÜ‰G‰O˜DÕ!ä‰G‰O˜DÔ!Ø�N‰NœnÐ-¨S´·±×1CÑ1C×1RÑ1RÓ1TÐ0UÐVÔWà×Ñ !Ò#Ü˜q $×"2Ñ"2Ó3ò �AØ#'×#BÑ#BÀ1Ñ#E�LØ—N‘NÜ/×6Ñ6À!ÔDÐEð F:Ø:F¸ÀqðJõñð $rb   r`  )r]   r^   r_   rÁ  r÷  r`   rî  rñ  Údictrù  r\  r@  rA  s   @rc   rö  rö    s7   ø… ñð €K�ÓØ1B°×1BÑ1BÐSWÔ1XÐ ÓX÷ñ rb   rö  c                  ó(   — e Zd ZU dZdZded<   dd„Zy)Ú*ExitDeviceContextManagerWithStreamInfoLinez{Exit a device context.

    Attributes:
        num_streams: Number of streams that were allocated (must match Enter).
    rA   rY   r÷  c                óZ   — t         j                  j                  s|j                  «        yy)zGenerate context exit code.Nr•  rZ  s     rc   r\  z2ExitDeviceContextManagerWithStreamInfoLine.codegen8  s    € ä�w‰w×"Ò"Ø×ÑÕð #rb   Nr`  )r]   r^   r_   rÁ  r÷  r`   r\  ra   rb   rc   r  r  .  s   … ñð €K�Óôrb   r  c                  ó$   — e Zd ZU dZded<   dd„Zy)ÚEnterCudaStreamContextLinezœEnter a context executed by the respective device stream.

    Attributes:
        stream_idx: The index number for the entering device stream context.
    rY   rü  c                ót   — |j                  dt        | j                  «      › d�«       |j                  «        y )Nr†  r‡  )r¯   r1   rü  rY  rZ  s     rc   r\  z"EnterCudaStreamContextLine.codegenH  s+   € Ø�‰˜œ¨t¯©Ó?Ð@ÀÐBÔCØ�‰Õrb   Nr`  rÀ  ra   rb   rc   r  r  >  s   … ñð ƒOôrb   r  c                  ó   — e Zd ZdZdd„Zy)ÚExitCudaStreamContextLinez1Generate code to exit the current stream context.c                ó$   — |j                  «        y re   )r|  rZ  s     rc   r\  z!ExitCudaStreamContextLine.codegenQ  s   € Ø×ÑÕrb   Nr`  )r]   r^   r_   rÁ  r\  ra   rb   rc   r
  r
  M  s
   „ á;ôrb   r
  c                  ó*   — e Zd Zd„ Zdd„Zdd„Zdd„Zy)ÚEfficientPeakEstimatec                ó’  — ddl m}m} t        j                  j
                  j                  }t        t        j                  j                  j                  «       «      }t        t        j                  j                  «       «      } |||«      } ||||«      \  | _        }ddlm}  ||t        j                  t         d«      | _        y )Nr$   )Úestimate_peak_memoryÚget_freeable_input_bufrA   )ÚSegmentedTreer   )Úmemoryr  r  r@   rf   rj   Únodesr   Úgraph_inputsÚkeysÚget_output_namesÚoverall_peak_memoryÚsegmented_treer  ÚoperatorrÓ   Úmax)	r  r  r  Úscheduler_nodesr  Úgraph_outputsÚnames_to_freeable_bufsÚpeak_by_scheduler_noder  s	            rc   r%  zEfficientPeakEstimate.__init__V  s–   € ßIäŸ'™'×+Ñ+×1Ñ1ˆÜ!¤!§'¡'×"6Ñ"6×";Ñ";Ó"=Ó>ˆÜ"¤1§7¡7×#;Ñ#;Ó#=Ó>ˆÙ!7¸ÈÓ!VÐÙ;OØØ"Øó<
Ñ8ˆÔ Ð"8õ 	2á+Ø"¤H§L¡L´#°qó
ˆÕrb   c                óÆ   — t         j                  j                  j                  t         j                  j	                  |«      d¬«      t        |j                  «       «      z  S )Nr   ©Úfallback)r@   rf   rn   Úoptimization_hintrg   r6   rm   ©r  rp   s     rc   Ú	_get_sizezEfficientPeakEstimate._get_sizei  sL   € Ü�w‰w×Ñ×1Ñ1Ü�G‰G×/Ñ/°Ó5Àð 2ó 
ä˜4Ÿ>™>Ó+Ó,ñ-ð 	-rb   c                ón   — | j                   j                  |j                  dz   |j                  dz
  «      S r¸   )r  Úsummarize_rangeÚscheduler_node_index©r  Úline_aÚline_bs      rc   Úpeak_betweenz"EfficientPeakEstimate.peak_betweenn  s6   € Ø×"Ñ"×2Ñ2Ø×'Ñ'¨!Ñ+¨V×-HÑ-HÈ1Ñ-Ló
ð 	
rb   c                óÞ   — |j                   dz   |j                   k(  ry | j                  j                  |j                   dz   |j                   dz
  | j                  |j                  «      «       y r¸   )r'  r  Úupdate_ranger$  rp   r(  s      rc   Úupdate_peak_betweenz)EfficientPeakEstimate.update_peak_betweens  s^   € Ø×&Ñ&¨Ñ*¨f×.IÑ.IÒIØØ×Ñ×(Ñ(Ø×'Ñ'¨!Ñ+Ø×'Ñ'¨!Ñ+Ø�N‰N˜6Ÿ;™;Ó'õ	
rb   N)rp   Ú
BufferLikerÈ   rY   )r)  r=  r*  ÚAllocateLine)r]   r^   r_   r%  r$  r+  r.  ra   rb   rc   r  r  U  s   „ ò
ó&-ó

ô

rb   r  c                  óX   — e Zd ZU dZded<   dZded<   d„ Zdd„Zdd	„Zdd
„Z	dd„Z
dd„Zy)r0  z6Represents a buffer allocation during memory planning.r/  rp   Fr  Úcomm_bufferc                ó  — t         j                  j                  j                  €J ‚t         j                  j                  j                  j                  t         j                  j                  j                  «      | _        y re   ©r@   rf   rj   Úcurrent_noder  Úindexr'  r  s    rc   rV  zAllocateLine.__post_init__„  óT   € Ü�w‰w× Ñ ×-Ñ-Ð9Ð9Ð9Ü$%§G¡G×$5Ñ$5×$;Ñ$;×$AÑ$AÜ�G‰G×Ñ×*Ñ*ó%
ˆÕ!rb   c                óö   — | j                   ry|j                  dz   | j                  k(  ry| j                  j                  j                  }| j                  j                  j                  || «      }||z   }||k  S )NTrA   )r2  r'  r£   Úestimate_peakr  r+  )r  Ú	free_linerˆ   r  Úpeak_memory_in_rangeÚnew_peak_memorys         rc   Úshould_reuse_bufferz AllocateLine.should_reuse_bufferŠ  su   € Ø×ÒØØ×)Ñ)¨AÑ-°×1JÑ1JÒJØØ"Ÿl™l×8Ñ8×LÑLÐØ#Ÿ|™|×9Ñ9×FÑFÀyÐRVÓWÐØÐ!5Ñ5ˆØÐ"5Ñ5Ð5rb   c           	     ó(  — | j                   j                  «       t        j                  j                  v rt        | j                  «      S | j                  r}t        | j                   «      }t        j                  rV|j                  |«      rE|j                  |«      }d|_        t        | j                  |j                   | j                   d¬«      S | S t        | j                   «      }t        j                  �r
||v �r|j!                  |«      }t        j                  j"                  j%                  t        j                  j'                  | j                   «      d¬«      t)        | j                   j+                  «       «      z  }| j-                  ||«      rXd|_        | j                  j.                  j1                  || «       t        | j                  |j                   | j                   «      S |j3                  ||«       | S | j                   j5                  «       j6                  dk(  rh| j                  j9                  | j                   «      }|�A|xj:                  t=        t?        j@                  tB        jD                  |d«      «      z  c_        | S )NT©r2  r   r   ÚcpurA   )#rp   rh   r@   rf   rÆ  ÚNullLiner£   r2  r|   r&   Úallow_buffer_reuser7  r9  r2  Ú	ReuseLinert   r1  rn   r"  rg   r6   rm   r=  r9  r.  r5  rl   r¢   Ústatic_shape_for_buffer_or_noner+  rY   Ú	functoolsÚreducer  Úmul)r  rê  rÀ   r:  rˆ   Ústatic_shapes         rc   rë  zAllocateLine.plan”  së  € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ñ:Ü˜DŸL™LÓ)Ð)à×Òä'¨¯	©	Ó2ˆCÜ×(Ò(¨U×-GÑ-GÈÔ-LØ!×1Ñ1°#Ó6�	Ø&*�	Ô#Ü Ø—L‘L )§.¡.°$·)±)Èôð ð ˆKô ˜tŸy™yÓ)ˆÜ×$Ó$¨°ªØŸ	™	 #›ˆIÜ—7‘7×#Ñ#×5Ñ5Ü—‘×3Ñ3°D·I±IÓ>Èð 6ó ä˜tŸy™y×2Ñ2Ó4Ó5ñ6ˆDð ×'Ñ'¨	°4Ô8Ø&*�	Ô#Ø—‘×*Ñ*×>Ñ>¸yÈ$ÔOÜ  §¡¨y¯~©~¸t¿y¹yÓIÐIà—
‘
˜3 	Ô*Ø�à�9‰9×(Ñ(Ó*×/Ñ/°5Ò8ØŸ<™<×GÑGÈÏ	É	ÓRˆLØÐ'Ø×1Ò1´SÜ×$Ñ$¤X§\¡\°<ÀÓCó6ñ Õ1ð ˆrb   c                ó  — | j                   j                  «       t        j                  j                  vsJ ‚| j
                  r| j                  |«       y | j                  j                  | j                   «      }|j                  |«       y re   )
rp   rh   r@   rf   rÆ  r2  Ú_codegen_comm_bufferr£   Úmake_buffer_allocationr¯   )r  r[  r²   s      rc   r\  zAllocateLine.codegen¼  sb   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ×ÒØ×%Ñ% dÕ+à—<‘<×6Ñ6°t·y±yÓAˆDØ�N‰N˜4Õ rb   c                ó@  — | j                   j                  «       }| j                   j                  «       }|�|j                  €
J d|› �«       ‚| j                   j	                  «       }t        | j                   j                  «       «      }t        | j                   j                  «       «      }| j                   j                  «       }t        |t        j                  «      sJ ‚|j                  }|j                  }	|t        j                  j                  k(  rh|› d| j                   j#                  |«      › d| j                   j#                  |«      › d|› d|j                  › d|	› dt%        j&                  dd	«      › d
�}
nt)        d|› �«      ‚|j+                  |
«       y)z*Generate allocation code for comm buffers.Nz9Comm buffer requires a valid CUDA device with index, got z = empty_strided_p2p(r»  z, torch.device("cuda:z"), group_name="z", alloc_id=r   l   ÿÿÿÿ rï   zUnsupported comm buffer type: )rp   rh   rª  r6  rm   r¦   Úget_sizeÚ
get_striderv   rw   r'   rx   ry   rz   ÚCommBufferTypeÚSYMM_MEMr£   Úcodegen_shape_tupleÚrandomÚrandintrI  r¯   )r  r[  r³   r[   r   Úshaper‰   r{   ry   rz   r²   s              rc   rJ  z!AllocateLine._codegen_comm_bufferÄ  s  € à�y‰y×!Ñ!Ó#ˆØ—‘×%Ñ%Ó'ˆØÐ! f§l¡lÐ&>ð 	
ØGÈÀxÐPó	
Ð>ð —	‘	×#Ñ#Ó%ˆÜ�d—i‘i×(Ñ(Ó*Ó+ˆÜ�t—y‘y×+Ñ+Ó-Ó.ˆØ—‘×*Ñ*Ó,ˆÜ˜&¤"×"5Ñ"5Ô6Ð6Ð6Ø!×2Ñ2ÐØ×&Ñ&ˆ
àœr×0Ñ0×9Ñ9Ò9à�&Ð-Ø—<‘<×3Ñ3°EÓ:Ð;¸2Ø—<‘<×3Ñ3°FÓ;Ð<¸BØ�'ð &Ø&,§l¡l ^ð 4Ø)˜lð +Ü"ŸN™N¨1¨iÓ8Ð9¸ð<ñ ô &Ø0Ð1AÐ0BÐCóð ð 	�‰�tÕrb   c                óJ   — | j                   r|j                  S |j                  S re   )r2  Ú_generate_comm_buffer_allocateÚ_generate_allocaterJ  s     rc   rL  zAllocateLine.codegen_fxã  s#   € Ø×ÒØ×;Ñ;Ð;Ø×+Ñ+Ð+rb   N)r:  r=  rˆ   rY   rÈ   r  ró  r`  rM  )r]   r^   r_   rÁ  r`   r2  rV  r=  rë  r\  rJ  rL  ra   rb   rc   r0  r0  }  s4   … á@à
ÓØ€K�Óò
ó6ó&óP!óô>,rb   r0  c                  óR   — e Zd ZU ded<   dZded<   dZded<   d„ Zdd„Zdd	„Zdd
„Z	y)r=  r/  rp   Fr  r2  r2  c                ó  — t         j                  j                  j                  €J ‚t         j                  j                  j                  j                  t         j                  j                  j                  «      | _        y re   r4  r  s    rc   rV  z!FreeIfNotReusedLine.__post_init__ï  r7  rb   c                ó<  — t        | j                  j                  «       «      dkD  r| S t        | j                  j                  t
        j                  «      r| S | j                  rJ ‚| j                  j                  «       t        j                  j                  v rt        | j                  «      S t        j                  r\| j                   r)t#        | j                  «      }|j%                  || «       | S t'        | j                  «      }|j)                  || «       | S r#  )rº   rp   Úget_inputs_that_alias_outputrw   r{   r'   ÚMultiOutputLayoutr2  rh   r@   rf   rÆ  rA  r£   r&   rB  r2  r|   r;  rt   r5  )r  rê  rÀ   s      rc   rë  zFreeIfNotReusedLine.planõ  sÕ   € Üˆt�y‰y×5Ñ5Ó7Ó8¸1Ò<ØˆKÜ�d—i‘i×&Ñ&¬×(<Ñ(<Ô=ØˆKØ—>’>Ð!Ð!Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ñ:Ü˜DŸL™LÓ)Ð)Ü×$Ò$Ø×Òä+¨D¯I©IÓ6�Ø×&Ñ& s¨DÔ1ð ˆô ' t§y¡yÓ1�Ø—
‘
˜3 Ô%Øˆrb   c                óÒ  — | j                   j                  «       t        j                  j                  vsJ ‚| j
                  s¥| j                  j                  | j                   «      }| j                  rb| j                   j                  «       }t        |t        j                  «      sJ ‚|j                  |› d|j                  j                  › d�«       y |j                  |«       y y )Nz # z buffer free)rp   rh   r@   rf   rÆ  r2  r£   rÇ  r2  rv   rw   r'   rx   r¯   ry   r  )r  r[  r²   r{   s       rc   r\  zFreeIfNotReusedLine.codegen  s®   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�~Š~Ø—<‘<×0Ñ0°·±Ó;ˆDØ×ÒØŸ™×2Ñ2Ó4�Ü! &¬"×*=Ñ*=Ô>Ð>Ð>Ø—‘ $  s¨6×+BÑ+B×+HÑ+HÐ*IÈÐVÕWà—‘˜tÕ$ð rb   c                óJ   — | j                   r|j                  S |j                  S re   )r2  Ú_generate_comm_buffer_freeÚ_generate_free_if_not_reusedrJ  s     rc   rL  zFreeIfNotReusedLine.codegen_fx  s#   € Ø×ÒØ×7Ñ7Ð7Ø×5Ñ5Ð5rb   Nró  r`  rM  )
r]   r^   r_   r`   r2  r2  rV  rë  r\  rL  ra   rb   rc   r=  r=  é  s0   … à
ÓØ€IˆtÓØ€K�Óò
óó$	%ô6rb   r=  c                  óD   — e Zd ZU ded<   ded<   ded<   d
d„Zdd„Zdd„Zy	)ÚReinterpretLiner/  rp   Ú	reused_asz	ir.Layoutr{   c                ó   — | S re   ra   ré  s     rc   rë  zReinterpretLine.plan  s   € Øˆrb   c                ó@  — t        | j                  t        j                  «      sJ ‚t        | j                  j                  t        j
                  «      sJ ‚| j                  j                  | j                  j                  «       | j                  j                  «       y re   )
rw   r{   r'   ÚNonOwningLayoutÚviewr*   r£   Úcodegen_deferred_allocationrc  rh   rZ  s     rc   r\  zReinterpretLine.codegen!  sj   € Ü˜$Ÿ+™+¤r×'9Ñ'9Ô:Ð:Ð:Ü˜$Ÿ+™+×*Ñ*¬B×,>Ñ,>Ô?Ð?Ð?Ø�‰×0Ñ0Ø�N‰N×#Ñ#Ó% t§{¡{×'7Ñ'7õ	
rb   c                ó   — |j                   S re   )Ú_generate_reinterpretrJ  s     rc   rL  zReinterpretLine.codegen_fx(  rÚ  rb   Nró  r`  rM  )r]   r^   r_   r`   rë  r\  rL  ra   rb   rc   rb  rb    s#   … à
ÓØÓØÓóó
ô/rb   rb  c                  óV   — e Zd ZU ded<   ded<   dZded<   dZded<   dd	„Zdd
„Zdd„Zy)rC  r/  rp   rc  Tr  Ú
delete_oldFr2  c                óp  — | j                   j                  «       t        j                  j                  v rK| j
                  j                  «       t        j                  j                  v sJ ‚t        | j                  «      S | j
                  j                  «       t        j                  j                  vsJ ‚| S re   )rp   rh   r@   rf   rÆ  rc  rA  r£   ré  s     rc   rë  zReuseLine.plan3  s   € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ñ:Ø—>‘>×*Ñ*Ó,´·±×0GÑ0GÑGÐGÐGÜ˜DŸL™LÓ)Ð)Ø�~‰~×&Ñ&Ó(´·±×0GÑ0GÑGÐGÐGØˆrb   c                óp  — | j                   j                  «       t        j                  j                  vsJ ‚| j
                  j                  «       t        j                  j                  vsJ ‚|j                  | j                  j                  | j                   | j
                  | j                  «      «       y re   )
rp   rh   r@   rf   rÆ  rc  r¯   r£   Úmake_buffer_reuserl  rZ  s     rc   r\  zReuseLine.codegen:  sz   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�~‰~×&Ñ&Ó(´·±×0GÑ0GÑGÐGÐGØ�‰Ø�L‰L×*Ñ*¨4¯9©9°d·n±nÀdÇoÁoÓVõ	
rb   c                ó   — |j                   S re   )Ú_generate_reuserJ  s     rc   rL  zReuseLine.codegen_fxA  s   € Ø×(Ñ(Ð(rb   Nró  r`  rM  )	r]   r^   r_   r`   rl  r2  rë  r\  rL  ra   rb   rc   rC  rC  ,  s1   … à
ÓØÓØ€J�ÓØ€K�Óóó
ô)rb   rC  c                  ó   — e Zd Zdd„Zy)rA  c                ó   — |j                   S re   )Ú_generate_nullrJ  s     rc   rL  zNullLine.codegen_fxF  rÊ  rb   NrM  rO  ra   rb   rc   rA  rA  E  s   „ ô(rb   rA  c                  óJ   — e Zd ZU dZded<   ded<   ded<   ded<   dd	„Zdd
„Zy)ÚMultiOutputLinezU
    Given a MultiOutputLayout buffer, indexes actual buffer(s) from the result.
    rR  r£   rW   Úresult_nameÚarg_nameúSequence[Any]Úindicesc                óæ   ‡ ‡— ˆˆ fd„Š ‰‰ j                   ‰ j                  «      }|j                  ‰ j                  j                  › ‰ j
                  › d|› ‰ j                  j                  › �«       y )Nc                ól  •— t        |«      dkD  r¤|d   \  }}t        |t        «      r ‰| › d|› d�|dd  «      S t        |t        «      r<‰j                  j                  | ‰j                  t        |«      «      } ‰||dd  «      S t        |t        «      r ‰| › d|› d�|dd  «      S t        d|«      ‚| S )Nr   ú[ú]rA   z['z']znon supported index type: )
rº   Ú
issubclassr(  r¦   r£   Úcodegen_tuple_accessrw  rW   r  ÚAssertionError)Úbasenamerz  Úityper   Útuple_accessÚcodegen_list_tuple_accessr  s        €€rc   r…  z:MultiOutputLine.codegen.<locals>.codegen_list_tuple_accessV  sË   ø€ Ü�7‹|˜aÒØ" 1™:‘��qÜ˜e¤TÔ*Ù4¸°zÀÀ1À#ÀQÐ5GÈÐQRÐQSÈÓUÐUÜ ¤uÔ-à#'§<¡<×#DÑ#DØ  $×"2Ñ"2´C¸³Fó$�Lñ 5°\À7È1È2À;ÓOÐOÜ ¤tÔ,Ù4¸°zÀÀAÀ3ÀbÐ5IÈ7ÐSTÐSUÈ;ÓWÐWä(Ð)EÀuÓMÐMà�rb   rî   )rx  rz  r¯   r£   Údeclarerw  Úending)r  r[  r  r…  s   `  @rc   r\  zMultiOutputLine.codegenU  s]   ù€ õ	 ñ$ *¨$¯-©-¸¿¹ÓFˆØ�‰Ø�|‰|×#Ñ#Ð$ T×%5Ñ%5Ð$6°c¸%¸ÀÇÁ×ATÑATÐ@UÐVõ	
rb   c                ó   — |j                   S re   )Ú_generate_multi_outputrJ  s     rc   rL  zMultiOutputLine.codegen_fxm  s   € Ø×/Ñ/Ð/rb   Nr`  rM  )r]   r^   r_   rÁ  r`   r\  rL  ra   rb   rc   rv  rv  J  s*   … ñð "Ó!ØÓØƒMØÓó
ô00rb   rv  c                  ó<   — e Zd ZU ded<   ded<   ded<   d
d„Zdd„Zy	)ÚIndexPutFallbackLinerR  r£   úir.IndexPutFallbackrp   zlist[ir.IRNode | None]rz  c                ó�  — | j                   }t        j                  |j                  «      sJ ‚d„ |j                  d d D «       \  }}| j                  D �cg c]*  }|r|j                  «       n| j                  j                  ‘Œ, }} | j                  j                  |j                  «       |||g|j                  «       ¢­Ž  y c c}w )Nc              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wre   ©r­  ©rœ   Úts     rc   rž   z/IndexPutFallbackLine.codegen.<locals>.<genexpr>z  s   è ø€ ÒF°�q×*Ñ*×,ÑFùó   ‚r$   )rp   r'   Úis_node_sequencer\   rz  r­  r£   Únone_strÚ_generate_index_put_fallbackr¨  Úcodegen_const_args)r  r[  rp   r½   ÚvaluesÚidxrz  s          rc   r\  zIndexPutFallbackLine.codegenw  s¶   € Ø�y‰yˆÜ×"Ñ" 4§;¡;Ô/Ð/Ð/ÙF°d·k±kÀ"À1°oÔF‰ˆˆFð —|‘|ö
àñ (+ˆC×!Ñ!Ô#°·±×0EÑ0EÑEð
ˆð 
ð
 	2ˆ�‰×1Ñ1Ø× Ñ Ó" A w°ð	
Ø9=×9PÑ9PÓ9Rô	
ùò
s   Á/Cc                ó   — |j                   S re   )r•  rJ  s     rc   rL  zIndexPutFallbackLine.codegen_fx„  s   € Ø×5Ñ5Ð5rb   Nr`  rM  r‘  ra   rb   rc   r‹  r‹  q  s   … à!Ó!Ø
ÓØ#Ó#ó
ô6rb   r‹  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚScatterFallbackLinerR  r£   úir.ScatterFallbackrp   c           
     ó8  — | j                   }t        j                  |j                  «      sJ ‚|j                  rd„ |j                  D «       \  }}}n%d„ |j                  D «       \  }}|j
                  d   }|j                  «       x}r|j                  nt        j                  j                  }| j                  j                  |||j
                  d   ||g|j                  |j                  |j                  |j                  d   |j!                  «       |«       y )Nc              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wre   r�  r�  s     rc   rž   z.ScatterFallbackLine.codegen.<locals>.<genexpr>‘  s   è ø€ ÒJ¸˜q×2Ñ2×4ÑJùr’  c              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wre   r�  r�  s     rc   rž   z.ScatterFallbackLine.codegen.<locals>.<genexpr>“  s   è ø€ ÒE°A˜!×-Ñ-×/ÑEùr’  rA   r   rF  )rp   r'   r“  r\   Úsrc_is_tensorÚconstant_argsrª  r¢   r@   rf   r«  r£   Ú_generate_scatter_fallbackr©  Úpython_kernel_namer»   r�  )r  r[  rp   r½   r6  rü   r±  r[   s           rc   r\  zScatterFallbackLine.codegen�  sê   € Ø�y‰yˆÜ×"Ñ" 4§;¡;Ô/Ð/Ð/Ø×ÒÙJ¸d¿k¹kÔJ‰OˆQ�‘sáE¸¿¹ÔE‰JˆQ�Ø×$Ñ$ QÑ'ˆCØ!%§¡Ó!2Ð2˜AÐ2�—’¼¿¹×9LÑ9LˆØ�‰×/Ñ/ØØ�×"Ñ" 1Ñ% u¨cÐ2Ø× Ñ Ø×#Ñ#Ø×ÑØ�K‰K˜Ñ!Ø×ÑÓ!Øõ		
rb   c                ó   — |j                   S re   )r¢  rJ  s     rc   rL  zScatterFallbackLine.codegen_fx¡  s   € Ø×3Ñ3Ð3rb   Nr`  rM  r‘  ra   rb   rc   r›  r›  ˆ  s   … à!Ó!Ø
Óó
ô(4rb   r›  c                  ó<   — e Zd ZU ded<   ded<   ded<   d
d„Zdd„Zy	)ÚSymbolicCallArgLinerR  r£   r  ÚargrP   rf   c                ód   — | j                   j                  | j                  | j                  «       y re   )r£   Ú"_generate_symbolic_call_arg_helperr§  rf   rZ  s     rc   r\  zSymbolicCallArgLine.codegen«  s   € Ø�‰×7Ñ7¸¿¹À$Ç*Á*ÕMrb   c                ó   — |j                   S re   )Ú_generate_symbolic_call_argrJ  s     rc   rL  zSymbolicCallArgLine.codegen_fx®  r´  rb   Nr`  rM  r‘  ra   rb   rc   r¦  r¦  ¥  s   … à!Ó!Ø	ÓØÓóNô5rb   r¦  c                  óF   — e Zd ZU ded<   ded<   ded<   ded<   dd	„Zdd
„Zy)ÚUnbackedSymbolDefsLinerR  r£   rW   Úoutput_namer   Úoutputsú)dict[sympy.Symbol, pytree.KeyPath] | NoneÚunbacked_bindingsc                óz   — | j                   j                  | j                  | j                  | j                  «       y re   )r£   Ú)_codegen_unbacked_symbol_defs_for_outputsr®  r¯  r±  rZ  s     rc   r\  zUnbackedSymbolDefsLine.codegen¹  s+   € Ø�‰×>Ñ>Ø×Ñ˜dŸl™l¨D×,BÑ,Bõ	
rb   c                ó   — |j                   S re   )Ú_generate_unbacked_symbol_defsrJ  s     rc   rL  z!UnbackedSymbolDefsLine.codegen_fx¾  s   € Ø×7Ñ7Ð7rb   Nr`  rM  r‘  ra   rb   rc   r­  r­  ²  s#   … à!Ó!ØÓØƒLØ@Ó@ó
ô
8rb   r­  c                  ó^   — e Zd ZU ded<   ded<   ded<   ded<   dZded<   dd	„Zedd
„«       Zy)ÚAssertSizeStrideLinerR  r£   rW   r³   rˆ   r‰   ÚinputÚop_namec                ó’   — | j                   j                  || j                  | j                  | j                  | j
                  «       y re   )r£   Ú_codegen_assert_size_strider³   rˆ   r‰   r¹  rZ  s     rc   r\  zAssertSizeStrideLine.codegenÊ  s/   € Ø�‰×0Ñ0Ø�$—)‘)˜TŸY™Y¨¯©°T·\±\õ	
rb   c                ó   — | j                   S re   )Ú_generate_assert_size_striderh  s    rc   rL  zAssertSizeStrideLine.codegen_fxÏ  s   € à×5Ñ5Ð5rb   Nr`  rM  )r]   r^   r_   r`   r¹  r\  rj  rL  ra   rb   rc   r·  r·  Â  s:   … à!Ó!Ø
ƒIØ
ƒIØƒKØ€GˆSÓó
ð
 ò6ó ñ6rb   r·  c                  ó8   — e Zd ZU dZded<   ded<   ded<   d	d„Zy)
ÚAssertDivByZeroLinea  Deferred AOTI runtime check that a sizevar divisor is non-zero.

    Queued from CppWrapperCpu.codegen_cpp_sizevar during the build phase
    and replayed at the position of use so the check lands after the
    declarations of any unbacked symbols it references.
    rR  r£   rW   Údivisorr¹  c                óf   — | j                   j                  || j                  | j                  «       y re   )r£   Ú_codegen_assert_div_by_zerorÀ  r¹  rZ  s     rc   r\  zAssertDivByZeroLine.codegená  s    € Ø�‰×0Ñ0°°t·|±|ÀTÇ\Á\ÕRrb   Nr`  rÀ  ra   rb   rc   r¿  r¿  Ô  s   … ñð "Ó!ØƒLØƒLôSrb   r¿  c                  óÊ  ‡ — e Zd ZU dZdZded<   ˆ fd„Ze	 d¥	 	 	 	 	 	 	 d¦d„«       Zd§d„Z	d¨d	„Z
d§d
„Zd©d„Zd§d„Zed§d„«       Zd§d„Zed§d„«       Zdªd„Zed«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d«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$d°d#„Z%dµd$„Z&d%„ Z'd&„ Z(d'„ Z)d(„ Z*d)„ Z+d¶d*„Z,	 	 d·	 	 	 	 	 	 	 d¸d+„Z-d§d,„Z.	 	 	 	 d¹d-„Z/d§d.„Z0dºd/„Z1d»d0„Z2d»d1„Z3d»d2„Z4d¼d3„Z5d½d4„Z6d¾d5„Z7d6„ Z8	 	 	 	 d¿d7„Z9	 d¥	 	 	 	 	 	 	 	 	 	 	 	 	 dÀd8„Z:dÁd9„Z;dÁd:„Z<dÁd;„Z=d<„ Z>dÂd=„Z?d>„ Z@dÃd?„ZAd@„ ZB	 	 	 	 	 	 	 	 	 	 	 	 	 	 dÄdA„ZCdB„ ZDd¯dC„ZEeFjŽ                  	 	 	 	 	 	 dÅdD„«       ZHd§dE„ZIdF„ ZJdG„ ZKdH„ ZLdI„ ZMdÆdJ„ZN	 	 	 	 	 	 dÇdK„ZOdL„ ZPdÈdM„ZQdN„ ZRddOœdÉdP„ZSddOœdÉdQ„ZTdÊdR„ZUdËdS„ZVdÌdT„ZWdÌdU„ZX	 	 dÍdV„ZY	 d¥	 	 	 dÎdW„ZZdÏdX„Z[dÐdY„Z\dÑdZ„Z]d[„ Z^d\„ Z_d]„ Z`d^„ Zad_„ Zbd`„ Zc	 	 	 dÒ	 	 	 	 	 	 	 	 	 dÓda„Zde	 d¥	 	 	 	 	 dÔdb„«       Ze	 	 	 dÒ	 	 	 	 	 	 	 	 	 dÓdc„ZfdÕdd„Zgehde„ «       Zi	 	 	 	 dÖdf„Zjd¥d×dg„Zk	 	 	 	 	 	 dØdh„ZldÙdi„ZmdÙdj„Zndk„ Zodl„ Zpdm„ Zqdn„ Zrdo„ Zsdp„ Ztdq„ Zudr„ Zvds„ Zwd¥dt„Zxdu„ Zydddddddddvœ	 dÚdw„Zzddddddddxdddyœ
	 dÚdz„Z{d{„ Z|d|„ Z}d}„ Z~d¥d~„ZdÛd„Z€ed€„ «       Z�	 dÜ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‹„ ZŒd¥dŒ„Z�d�„ ZŽdådŽ„Z�d�„ Z�	 	 	 	 	 	 	 	 dæd�„Z‘	 	 	 	 	 	 	 	 dæd‘„Z’d’„ Z“	 	 	 	 dçd“„Z”dèd”„Z•d•„ Z–d–„ Z—d—„ Z˜d˜„ Z™d™„ Zšdš„ Z›d§d›„Zœdœ„ Z�ed�„ «       Zžedž„ «       ZŸedŸ„ «       Z ed „ «       Z¡ed¡„ «       Z¢	 	 	 	 déd¢„Z£d£„ Z¤d¤„ Z¥ˆ xZ¦S )êrR  zB
    Generate outer wrapper in Python that calls the kernels.
    Tr  Úsupports_cachingc                óÎ  •‡ — t         ‰‰ �  «        i ‰ _        t        «       ‰ _        t        «       ‰ _        i ‰ _        t        «       ‰ _	        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        t        «       ‰ _        i ‰ _        d‰ _        i ‰ _        t        «       ‰ _        g ‰ _        d‰ _        d‰ _        d‰ _        d‰ _        d‰ _        t<        j>                  j@                  rdnd‰ _!        t<        j>                  j@                  rdnd‰ _"        d ‰ _#        d‰ _$        i ‰ _%        t        «       ‰ _&        t        «       ‰ _'        d ‰ _(        ‰ jS                  «        g ‰ _*        g ‰ _+        ‰ jY                  «        t[        ‰ «      s‰ j]                  «        ‰ j_                  «        t<        j>                  j`                  sBt<        j>                  jb                  je                  «       D ]  \  }}‰ jg                  ||«       Œ t        th           «       ‰ _5        t        th           «       ‰ _6        i ‰ _7         tq        jr                  d «      ‰ jt                  «      ‰ _:        tp        jv                  d
ˆ fd„«       }|‰ _<        i ‰ _=        t        «       ‰ _>        t        «       ‰ _@        t        «       ‰ _A        i ‰ _B        t‡        tˆ        jŠ                  jŒ                  tˆ        jŠ                  jŽ                  ¬	«      ‰ _H        g ‰ _I        t        «       ‰ _J        y )Nr   Ú ú#r>  z
std::move(rï   Tc                ó¨   •— ‰j                   j                  | «       t        j                  j                  r‰j
                  j                  | «       y y re   )Úimportsr¯   r&   r¨   r©   r±   )r²   r  s    €rc   Úadd_import_oncez6PythonWrapperCodegen.__init__.<locals>.add_import_once=  s;   ø€ à�L‰L×"Ñ" 4Ô(Ü�}‰}×5Ò5Ø×*Ñ*×4Ñ4°TÕ:ð 6rb   )Údebug_printer_levelÚuse_array_ref)r²   rW   rÈ   r>  )Kr$  r%  Ú_pending_input_assertsr   Ú_pending_alignment_copiesr   Ú_names_iterÚargs_to_buffersr7   rÉ  r;   ÚheaderÚprefixÚsuffixÚkernel_declarationsÚwrapper_callÚkernel_autotune_defsr±   Úsubgraph_definitionsr°   Úkernel_autotune_example_argsÚkernel_autotune_tmp_arg_idxÚsrc_to_kernelÚkernel_numel_exprÚlinesr†  Údeclare_maybe_referencer‡  Úcommentr”  r@   rf   rˆ  Ú
move_beginÚmove_endrƒ  Úsupports_intermediate_hooksÚuser_defined_kernel_cacheÚunbacked_symbol_declsrU  Úlauncher_fn_nameÚset_launcher_fn_nameÚcodegened_graph_stackÚcomputed_sizes_stackÚwrite_headerr8   Úwrite_prefixÚ!write_kernel_autotune_defs_headerr‰  Úconstant_reprsÚitemsÚwrite_constantÚ
BufferNameÚ	allocatedÚfreedÚreusesrE  Ú	lru_cacheÚwrite_get_raw_streamÚcacherÊ  Ú_metasÚ
_meta_varsr   Úmulti_kernel_stateÚalready_codegened_subgraphsÚallocated_workspacesr   r&   Úaot_inductorÚ debug_intermediate_value_printerÚallow_stack_allocationÚdebug_printerÚadditional_filesÚexternal_kernel_libs)r  r³   ÚhashedrÊ  r,  s   `   €rc   r%  zPythonWrapperCodegen.__init__ð  sÌ  ù€ Ü‰ÑÔØBDˆÔ#Ü:D»,ˆÔ&Ü*/«'ˆÔð ð 	Ôô &Ó'ˆŒÜ)Ó+ˆŒÜ)Ó+ˆŒÜ$Ó&ˆŒÜ#1Ó#3ˆÔ Ü/Ó1ˆÔÜ$2Ó$4ˆÔ!Ü%3Ó%5ˆÔ"Ü$2Ó$4ˆÔ!Ü6@³lˆÔ"ð IKˆÔ)Ø01ˆÔ(ð .0ˆÔÜHRËˆÔØ!#ˆŒ
ØˆŒØ')ˆÔ$ØˆŒØˆŒØˆŒÜ*+¯'©'×*=Ò*=™,À2ˆŒÜ Ÿw™w×2Ò2™¸ˆŒØ8<ˆÔ)Ø+/ˆÔ(ð ð 	Ô&ô ‹Lð 	Ô"ô 9C»ˆÔØ $ˆÔà×!Ñ!Ô#ð &(ˆÔ"Ø$&ˆÔ!à×ÑÔä2°4Ô8à×ÑÔà×.Ñ.Ô0ä�w‰w×ÒÜ !§¡× 6Ñ 6× <Ñ <Ó >ò 2‘��fà×#Ñ# D¨&Õ1ð2ô $¤JÑ/Ó1ˆŒÜ¤
Ñ+Ó-ˆŒ
ð 57ˆŒà$=¤I×$7Ñ$7¸Ó$=Ø×%Ñ%ó%
ˆÔ!ô 
�‰ô	;ó 
ð	;ð
  /ˆÔØ&(ˆŒÜ+5«<ˆŒÜ"2Ó"4ˆÔÜ<F»LˆÔ(Ø46ˆÔ!ô 1Ü &× 3Ñ 3× TÑ TÜ ×-Ñ-×DÑDô
ˆÔð !#ˆÔô
 6@³\ˆÕ!rb   Nc                óD   — | r|€J ‚|€J ‚t        |||«      S t        «       S re   )ÚSubgraphPythonWrapperCodegenrR  )Úis_subgraphÚsubgraph_nameÚparent_wrapperÚpartition_signaturess       rc   ÚcreatezPythonWrapperCodegen.createX  s?   € ñ Ø Ð,Ð,Ð,Ø!Ð-Ð-Ð-Ü/Ø˜~Ð/Cóð ô $Ó%Ð%rb   c                ó   — d| _         y )NÚcall)rä  r  s    rc   rå  z)PythonWrapperCodegen.set_launcher_fn_nameg  s
   € à &ˆÕrb   c                óD   — | j                   j                  |› d|› �«       y )Nz = None  # )rÑ  r¯   )r  r³   r   s      rc   rí  z#PythonWrapperCodegen.write_constantk  s   € Ø�‰×Ñ   k°&°Ð:Õ;rb   c                óÜ  — t         j                  j                  j                  «       }d}|�|j                  �d|j                  › �}d}t        t        j                  j                  «      dkD  rd}n0t         j                  j                  j                  j                  rd}| j                  j                  d|› dt        j                  › d	|› d
�d¬«       | j                   j                  dd¬«       	 ddlm} | j                   j                  dd¬«       t        j*                  r| j                   j-                  d«       t        j.                  j0                  �rø| j                   j-                  d«       | j                   j-                  d«       | j                   j-                  d«       | j                   j-                  d«       | j                   j-                  d«       | j                   j-                  d«       | j                   j-                  d«       t        j.                  j2                  xs( t4        j6                  j9                  t;        «       d«      }| j                   j-                  d|› d�«       d|› d�}d|› d�}t        j.                  j<                  }t        j.                  j>                  }	t        j.                  j@                  }
| j                   j-                  d«       | j                   j                  d|› d |› d!|	› d"|
› d#�	«       | j                   j-                  d$|› d%�«       | j                   j-                  d&«       yy# t&        t(        f$ r Y �ŒQw xY w)'z>Write the header section of the generated Python wrapper code.rÆ  Nz
# AOT ID: r   zRfrom torch._inductor.codegen.debug_utils import _print_debugging_tensor_value_infozFfrom torch._inductor.runtime.debug_utils import tracked_empty_strided
z
                aH  
                from ctypes import c_void_p, c_long, c_int
                import torch
                import math
                import random
                import os
                import tempfile
                from math import inf, nan
                from cmath import nanj
                from torch._inductor.hooks import run_intermediate_hooks
                from torch._inductor.utils import maybe_profile
                from torch._inductor.codegen.memory_planning import _align as align
                from torch import device, empty_strided
                from zq import AsyncCompile
                from torch._inductor.select_algorithm import extern_kernels
                z
            Trã   a–  
                aten = torch.ops.aten
                inductor_ops = torch.ops.inductor
                _quantized = torch.ops._quantized
                assert_size_stride = torch._C._dynamo.guards.assert_size_stride
                assert_alignment = torch._C._dynamo.guards.assert_alignment
                empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
                empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
                empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
                reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
                alloc_from_pool = torch.ops.inductor._alloc_from_pool
                async_compile = AsyncCompile()
            )Ú_SymmetricMemoryzs
                empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
                zfrom torch.cuda import nvtxz import triton.profiler as protonz%import triton.profiler.language as plzCfrom triton.profiler.hooks import HookManager as _ProtonHookManagerzimport tritonzimport atexitz	import oszetriton.set_allocator(lambda size, align, stream: torch.empty(size, dtype=torch.uint8, device='cuda'))Úprotonzos.makedirs("z", exist_ok=True)zos.path.join("z", "inductor")z", "inductor.chrome_trace")z^from torch._inductor.runtime.proton_utils import process_proton_trace as _proton_process_tracez�
                def _proton_finalize_and_postprocess():
                    proton.finalize()
                    _trace_path = zµ
                    if os.path.exists(_trace_path):
                        _proton_process_trace(
                            _trace_path,
                            group_by_sm=z0,
                            split_invocations=z0,
                            per_cta_occupancy=z,,
                        )
                z5if not _ProtonHookManager.active_hooks: proton.start(z], backend="instrumentation", data="trace"); atexit.register(_proton_finalize_and_postprocess)úpl.enable_semantic("triton"))!ÚtorchÚ_guardsÚTracingContextÚtry_getÚaot_graph_namerY   r&   rú  rû  Ú	_inductorÚtest_configsÚtrack_memory_lifecyclerÉ  rû   r%   r]   rÑ  Útorch._C._distributed_c10dr  ÚAttributeErrorÚImportErrorÚannotate_trainingr¯   r¨   Úproton_profilingÚproton_output_dirÚosÚpathrÒ   r   Úproton_group_by_smÚproton_split_invocationsÚproton_per_cta_occupancy)r  ÚcontextÚaot_config_commentÚinductor_debug_utilsr  Ú
output_dirÚproton_nameÚ
trace_pathÚgroup_by_smÚsplit_invocationsÚper_cta_occupancys              rc   rè  z!PythonWrapperCodegen.write_headern  s4  € ä—-‘-×.Ñ.×6Ñ6Ó8ˆØÐØÐ 7×#9Ñ#9Ð#EØ#-¨g×.DÑ.DÐ-EÐ!FÐØ!ÐÜŒv×"Ñ"×CÑCÓDÀqÒHØ#wÑ Ü�_‰_×#Ñ#×0Ñ0×GÒGØ#lÐ à�‰×ÑðØ#Ð$ð %ô $×,Ñ,Ð-ð .à%Ð&ð 'ð!ð$ ð' 	ô 	
ð* 	�‰×Ñðð ð! 	ô 	
ð$	õ Dà�K‰K×Ñðð ð	 ô ô ×#Ò#Ø�K‰K×!Ñ!Ð"?Ô@Ü�=‰=×)Ó)Ø�K‰K×!Ñ!Ð"DÔEØ�K‰K×!Ñ!Ð"IÔJØ�K‰K×!Ñ!ØUôð �K‰K×!Ñ! /Ô2Ø�K‰K×!Ñ! /Ô2Ø�K‰K×!Ñ! +Ô.Ø�K‰K×!Ñ!ðGôô  Ÿ™×8Ñ8ò ¼B¿G¹G¿L¹LÜ“ ó=ˆJð �K‰K×!Ñ! M°*°Ð=NÐ"OÔPØ*¨:¨,°nÐEˆKØ)¨*¨Ð5PÐQˆJÜ Ÿ-™-×:Ñ:ˆKÜ &§¡× FÑ FÐÜ &§¡× FÑ FÐØ�K‰K×!Ñ!Øpôð �K‰K×Ñð#ð $. ,ð /)ð *5¨ð 6/Ø/@Ð.Að B/Ø/@Ð.Að Bðôð �K‰K×!Ñ!ð Ø +˜}ð -DðDôð
 �K‰K×!Ñ!Ð"@ÕAð[ *øô	 ¤Ð,ò 	Úð	ús   Ã6#M ÍM+Í*M+c                 ó   — y re   ra   )r  rÑ  s     rc   Úinclude_extra_headerz)PythonWrapperCodegen.include_extra_headerß  ó   € Ørb   c                óÐ   — | j                   j                  dt        j                  › d�«       	 ddlm} | j                   j                  dd¬«       y # t        t        f$ r Y y w xY w)Na3  
                import torch
                from math import inf, nan
                from torch._dynamo.testing import rand_strided
                from torch._dynamo.utils import preserve_rng_state
                from torch._inductor.select_algorithm import AlgorithmSelectorCache
                from aH   import AsyncCompile

                async_compile = AsyncCompile()
                generate_example_value = AlgorithmSelectorCache.generate_example_value
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
            r   )Ú_cuda_getCurrentRawStreamzU
                get_raw_stream = torch._C._cuda_getCurrentRawStream
                Trã   )rÖ  rû   r%   r]   Útorch._Cr/  r  r  )r  r/  s     rc   rê  z6PythonWrapperCodegen.write_kernel_autotune_defs_headerâ  sw   € Ø×!Ñ!×(Ñ(ðô $×,Ñ,Ð-ð .ðô	
ð 
	Ý:à×%Ñ%×,Ñ,ðð ð	 -õ øô œ^Ð,ò 	Ùð	ús   ¯#A ÁA%Á$A%c                ó  — dt         j                  › d�}t        j                  j                  r]| j
                  j                  |«       | j
                  j                  t        j                  j                  j                  d«      «       t        j                  j                  s`| j                  j                  |d¬«       | j                  j                  t        j                  j                  j                  d«      «       y y )NzU
            import triton
            import triton.language as tl
            from z+ import start_graph, end_graph
            Úget_raw_streamTrã   )r+   r]   r&   r¨   r©   r±   rû   r¯   r@   rf   rŠ  Úimport_get_raw_stream_asrˆ  rÉ  ©r  Ú
import_strs     rc   Úwrite_triton_header_oncez-PythonWrapperCodegen.write_triton_header_onceÿ  sÇ   € ðô $×,Ñ,Ð-ð .ðˆ
ô
 �=‰=×1Ò1Ø×&Ñ&×-Ñ-¨jÔ9Ø×&Ñ&×0Ñ0Ü—‘×"Ñ"×;Ñ;Ð<LÓMôô �w‰w×"Ò"Ø�L‰L×Ñ 
°$ÐÔ7Ø�L‰L×"Ñ"Ü—‘×"Ñ"×;Ñ;Ð<LÓMõð #rb   c                óš  — t         j                  j                  j                  d«      }t        j
                  j                  r6| j                  j                  |«      s| j                  j                  |«       t         j                  j                  s8| j                  j                  |«      s| j                  j                  |«       y y y )Nr2  )r@   rf   rŠ  r3  r&   r¨   r©   r±   Úcontainsr¯   rˆ  rÉ  )r  Úimport_get_raw_stream_strs     rc   Úwrite_get_raw_stream_headerz0PythonWrapperCodegen.write_get_raw_stream_header  s—   € Ü$%§G¡G×$6Ñ$6×$OÑ$OØó%
Ð!ô �=‰=×1Ò1Ø×-Ñ-×6Ñ6Ð7PÔQØ×*Ñ*×4Ñ4Ð5NÔOÜ�w‰w×"Ò"Ø—<‘<×(Ñ(Ð)BÔCØ—‘×&Ñ&Ð'@ÕAð Dð #rb   c                ó$   — | j                  «        y re   )r:  r  s    rc   Ú write_get_raw_stream_header_oncez5PythonWrapperCodegen.write_get_raw_stream_header_once  s   € à×(Ñ(Õ*rb   c                óŠ  — t        |«      }|| j                  vrœdt        | j                  «      › �}|| j                  |<   | j                  j	                  |› d|› �«       t
        j                  j                  r;| j                  j	                  |› d|› �«       | j                  j                  |«       | j                  |   S )NÚmetarî   )Úreprrõ  rº   rÑ  r¯   r&   r¨   r©   r±   rö  rÓ   )r  r>  Úvars      rc   Úadd_meta_oncez"PythonWrapperCodegen.add_meta_once   s¥   € ä�D‹zˆØ�t—{‘{Ñ"Øœ˜TŸ[™[Ó)Ð*Ð+ˆCà #ˆD�K‰K˜ÑØ�K‰K×!Ñ! S E¨¨T¨FÐ"3Ô4Ü�}‰}×5Ò5Ø×*Ñ*×4Ñ4¸°u¸CÀ¸vÐ5FÔGØ—‘×#Ñ# CÔ(à�{‰{˜4Ñ Ð rb   c                óz   — | j                  «       D �cg c]  }|j                  | j                  «      ‘Œ c}S c c}w re   )Úget_graph_outputsr­  rÕ  ©r  r½   s     rc   Úget_output_refsz$PythonWrapperCodegen.get_output_refs.  s<   € ð =A×<RÑ<RÓ<Tö
Ø78ˆA×Ñ × 1Ñ 1Õ2ò
ð 	
ùò 
s   “"8c                 ó   — y re   ra   r  s    rc   Úmark_output_typez%PythonWrapperCodegen.mark_output_type4  ó   € Ørb   c                ó6   — t         j                  j                  S re   )r@   rf   r  r  s    rc   Úget_graph_inputsz%PythonWrapperCodegen.get_graph_inputs7  s   € ô �w‰w×#Ñ#Ð#rb   c                ó6   — t         j                  j                  S re   )r@   rf   r  r  s    rc   rC  z&PythonWrapperCodegen.get_graph_outputs<  s   € Ü�w‰w×$Ñ$Ð$rb   c                óô  — | j                  «       j                  «       D ]×  \  }}t        |t        j                  t
        j                  t
        j                  t
        j                  f«      rŒO|t        j                  j                  vrŒlt        |j                  «       «      dk(  rŒ‰| j                  |j                  «       «      }| j                  |j                  «       «      }||f| j                   |<   ŒÙ y r#  )rJ  rì  rw   r–   ÚBasicr'   ÚTorchBindObjectÚGeneratorStateÚOpaqueObjectStater@   rf   Úgraph_input_namesr<   rM  rQ  rN  rÍ  )r  r³   Úbufrˆ   r‰   s        rc   Úcodegen_input_size_assertsz/PythonWrapperCodegen.codegen_input_size_asserts?  sÎ   € Ø×.Ñ.Ó0×6Ñ6Ó8ò 	?‰IˆD�#ÜØä—K‘KÜ×&Ñ&Ü×%Ñ%Ü×(Ñ(ð	ôð ð œ1Ÿ7™7×4Ñ4Ñ4Øô ˜SŸ\™\›^Ó,°Ò1ØØ×+Ñ+¨C¯L©L«NÓ;ˆDØ×-Ñ-¨c¯n©nÓ.>Ó?ˆFØ15°v°ˆD×'Ñ'¨Ò-ñ+	?rb   c                ó`  — | j                   j                  d«       | j                  «       j                  «       D ]r  \  }}t	        |t
        j                  t        j                  f«      rŒ1d|› d�}| j                   j                  |«       d|› d�}| j                   j                  |«       Œt y )Nz(# make sure graph inputs are not nan/infzassert not z.isnan().any().item()z.isinf().any().item())	rÒ  r¯   rJ  rì  rw   r–   rM  r'   rN  )r  r³   rR  r²   s       rc   Úcodegen_input_nan_assertsz.PythonWrapperCodegen.codegen_input_nan_assertsW  s˜   € Ø�‰×ÑÐHÔIØ×.Ñ.Ó0×6Ñ6Ó8ò 	(‰IˆD�#Ü˜#¤§¡¬R×-?Ñ-?Ð@ÔAØØ   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Ô'Ø   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Õ'ñ	(rb   c                ó:   — | j                   j                  d«       y )NzV

            async_compile.wait(globals())
            del async_compile
            )rÒ  rû   r  s    rc   Úwrite_async_compile_waitz-PythonWrapperCodegen.write_async_compile_waita  s   € Ø�‰×Ñðõ	
rb   c                ó¾   — dj                  |«      }t        |«      dk(  r|dz  }| j                  j                  |› d�«       | j                  j                  d«       y )Nr»  rA   ú,z = argszargs.clear())rÒ   rº   rÒ  r¯   )r  Úinput_namesÚlhss      rc   Ú
write_argszPythonWrapperCodegen.write_argsj  sP   € Ø�i‰i˜Ó$ˆÜˆ{Ó˜qÒ Ø�3‰JˆCØ�‰×Ñ   W˜oÔ.Ø�‰×Ñ˜nÕ-rb   c                óº   — t         j                  r| j                  j                  d«       d}|S | j                  j                  d| j                  › d�«       d}|S )NaÙ  
                class Runner:
                    def __init__(self, partitions):
                        self.partitions = partitions

                    def recursively_apply_fns(self, fns):
                        new_callables = []
                        for fn, c in zip(fns, self.partitions):
                            new_callables.append(fn(c))
                        self.partitions = new_callables

                    def call(self, args):
                r$   z
                def z(args):
                rA   )r&   Úgraph_partitionrÒ  rû   rä  ©r  Úprefix_indents     rc   Ú!write_launcher_fn_call_get_indentz6PythonWrapperCodegen.write_launcher_fn_call_get_indentq  sm   € Ü×!Ò!Ø�K‰K×Ñðôð ˆMð Ðð �K‰K×ÑðØ×*Ñ*Ð+ð ,ðôð
 ˆMàÐrb   c                ó6   — t         j                  j                  S re   )r@   rf   rQ  r  s    rc   Úget_graph_input_namesz*PythonWrapperCodegen.get_graph_input_names�  s   € Ü�w‰w×(Ñ(Ð(rb   c                óÀ  — | j                   €J ‚| j                  «        | j                  «       }| j                  j	                  |«      5  t
        j                  j                  rA| j                  j                  t        j                  j                  j                  «       «       t        j                  j                  «       }t
        j                  r| j                  j                  d|› d�«       | j                  «       x}r| j!                  |«       | j#                  «        t%        «       rt'        | «      r| j)                  «        d d d «       y # 1 sw Y   y xY w)Nz0training_annotation = nvtx._device_range_start('ú'))rä  rW  ra  rÒ  rÌ   r&   r¨   Údebug_sync_graphr¯   r@   rf   rŠ  ÚsynchronizeÚget_training_phaser  rc  r\  Úcodegen_inputsr9   r8   Ú"codegen_input_size_and_nan_asserts)r  r`  ÚphaserQ  s       rc   ré  z!PythonWrapperCodegen.write_prefix�  s  € Ø×$Ñ$Ð0Ð0Ð0Ø×%Ñ%Ô'Ø×>Ñ>Ó@ˆà�[‰[×Ñ Ó.ñ 	:Ü�}‰}×-Ò-Ø—‘×%Ñ%¤a§g¡g×&8Ñ&8×&DÑ&DÓ&FÔGÜ—G‘G×.Ñ.Ó0ˆEÜ×'Ò'Ø—‘×%Ñ%ØFÀuÀgÈRÐPôð %)×$>Ñ$>Ó$@Ð@Ð Ð@Ø—‘Ð 1Ô2à×ÑÔ!ô
 -Ô.Ü<¸TÔBà×7Ñ7Ô9÷)	:÷ 	:ñ 	:ús   Á
DEÅEc                ó†   — t         j                  r| j                  «        t         j                  r| j	                  «        y y re   )r&   Úsize_assertsrS  Únan_assertsrU  r  s    rc   rj  z7PythonWrapperCodegen.codegen_input_size_and_nan_asserts«  s1   € Ü×ÒØ×+Ñ+Ô-Ü×ÒØ×*Ñ*Õ,ð rb   c                ó”   — |D ]C  }|| j                   v sŒ| j                   j                  |«      \  }}| j                  |||d«       ŒE y )Nr¸  )rÍ  r1  Úwrite_assert_size_stride)r  rZ  r³   rˆ   r‰   s        rc   Úcodegen_deferred_input_assertsz3PythonWrapperCodegen.codegen_deferred_input_assertsµ  sO   € Øò 	KˆDØ�t×2Ñ2Ò2Ø#×:Ñ:×>Ñ>¸tÓD‘��fØ×-Ñ-¨d°D¸&À'ÕJñ	Krb   c           	     ó@   — | j                  t        | ||||«      «       y)z7Queue an assert_size_stride for emission during replay.N)r¯   r·  )r  r³   rˆ   r‰   r¹  s        rc   rp  z-PythonWrapperCodegen.write_assert_size_stride»  s   € ð 	�‰Ô+¨D°$¸¸fÀgÓNÕOrb   c                ó@   — |j                  d|› d|› d|› d|›d�	«       y)z½Emit one assert_size_stride line to `code` (replay-phase target).

        Subclasses override to change the emitted form (e.g., C++ assert with
        an AOTI runtime env guard).
        zassert_size_stride(r»  rï   N©r¯   )r  r[  r³   rˆ   r‰   r¹  s         rc   r»  z0PythonWrapperCodegen._codegen_assert_size_strideÁ  s-   € ð 	�‰Ð,¨T¨F°"°T°F¸"¸V¸HÀBÀwÀkÐQRÐSÕTrb   c                ó<   — | j                  t        | ||«      «       y)zÖQueue a div-by-zero AOTI check for emission during replay.

        codegen_cpp_sizevar is called in both build and replay phases;
        set_writeline routes the WrapperLine appropriately in either case.
        N)r¯   r¿  )r  Údivisor_strr¹  s      rc   Úwrite_assert_div_by_zeroz-PythonWrapperCodegen.write_assert_div_by_zeroÐ  s   € ð 	�‰Ô*¨4°¸gÓFÕGrb   c                ó   — t        d«      ‚)zÙEmit one div-by-zero AOTI check to `code` (replay-phase target).

        Only emitted by C++ wrappers; the Python wrapper relies on Python's
        native ZeroDivisionError instead of pre-checking divisors.
        z6AOTI div-by-zero check is only emitted by C++ wrappersre  )r  r[  rv  r¹  s       rc   rÂ  z0PythonWrapperCodegen._codegen_assert_div_by_zeroØ  s   € ô "ØDó
ð 	
rb   c                ó¼  — t         j                  j                  ryt         j                  j                  }|syt	        t         j                  j
                  «      }|D ]?  }||vsŒt         j                  j                  |   }| j                  j                  |«       ŒA | j                  r1dt         j                  _	        | j                  j                  d«       yy)z�Populate pending alignment copies for non-mutated inputs.
        Called from the scheduler after mutated_input_idxs is computed.NTz6from torch._C._dynamo.guards import copy_if_misaligned)r@   rf   rˆ  Úinputs_to_checkr   Úmutated_input_idxsrQ  rÎ  rÓ   Ú_defers_input_alignmentrÉ  r¯   )r  rz  Úmutated_idxsr˜  r³   s        rc   Úregister_alignment_check_inputsz4PythonWrapperCodegen.register_alignment_check_inputsä  s¯   € ô �7‰7×ÒØÜŸ'™'×1Ñ1ˆÙØô "¤!§'¡'×"<Ñ"<Ó=ˆØ"ò 	9ˆCØ˜,Ò&Ü—w‘w×0Ñ0°Ñ5�Ø×.Ñ.×2Ñ2°4Õ8ð	9ð ×)Ò)Ø.2ŒA�G‰GÔ+Ø�L‰L×"Ñ"ØHõð *rb   c                óÊ   — t         j                  j                  ry|D ]C  }|| j                  v sŒ| j                  j	                  |«       | j                  |› d|› d�«       ŒE y)z~Emit alignment check + clone just before the first kernel
        that reads each input, hiding the cost behind GPU execution.Nz = copy_if_misaligned(rï   )r@   rf   rˆ  rÎ  Údiscardr¯   )r  rZ  r³   s      rc   Ú!codegen_deferred_alignment_copiesz6PythonWrapperCodegen.codegen_deferred_alignment_copiesù  sa   € ô �7‰7×ÒØØò 	GˆDØ�t×5Ñ5Ò5Ø×.Ñ.×6Ñ6°tÔ<Ø—‘ $ Ð'=¸d¸VÀ1ÐEÕFñ	Grb   c                ó  — | j                  «        t        |«      }t        j                  j                  r=| j
                  j                  |› d|› d�«       t        j                  j                  r|S | j                  |› d|› d�«       |S )Nú = get_raw_stream(rï   )
r:  r0   r&   r¨   r©   r±   r¯   r@   rf   rˆ  )r  r�  rÔ  r³   s       rc   ró  z)PythonWrapperCodegen.write_get_raw_stream  s}   € Ø×(Ñ(Ô*Ü" :Ó.ˆÜ�=‰=×1Ò1Ø×&Ñ&×0Ñ0Ø�&Ð*¨:¨,°aÐ8ôô �w‰w×"Ò"à�Ø�‰˜$˜Ð1°*°¸QÐ?Ô@Øˆrb   c                ó    — | j                   d   S )Néÿÿÿÿ)ræ  r  s    rc   Úget_codegened_graphz(PythonWrapperCodegen.get_codegened_graph  s   € Ø×)Ñ)¨"Ñ-Ð-rb   c                ó:   — | j                   j                  |«       y re   )ræ  r�   )r  rf   s     rc   rX  z)PythonWrapperCodegen.push_codegened_graph  s   € Ø×"Ñ"×)Ñ)¨%Õ0rb   c                ó6   — | j                   j                  «       S re   )ræ  r1  r  s    rc   r{  z(PythonWrapperCodegen.pop_codegened_graph  s   € Ø×)Ñ)×-Ñ-Ó/Ð/rb   c                óP   — ddl m} | j                  j                   ||«      «      S )Nr   )Údeepcopy)ÚcopyrŠ  rç  r�   )r  rU  rŠ  s      rc   rT  z(PythonWrapperCodegen.push_computed_sizes  s!   € Ý!à×(Ñ(×/Ñ/±¸Ó0HÓIÐIrb   c                ó6   — | j                   j                  «       S re   )rç  r1  r  s    rc   ry  z'PythonWrapperCodegen.pop_computed_sizes!  s   € Ø×(Ñ(×,Ñ,Ó.Ð.rb   c                ó.   — t        | j                  «      › S re   )ÚnextrÏ  r  s    rc   Únext_kernel_suffixz'PythonWrapperCodegen.next_kernel_suffix$  s   € Ü�t×'Ñ'Ó(Ð)Ð*rb   c                óÚ  — |dkD  rd|€J ‚d}| j                   j                  |«      s| j                   j                  |«       | j                  t        || j                  ||«      «       n%| j                  t        || j                  «      «       t        j                  j                  rµ| j                  «        | j                  j                  dt        j                  j                  j                  |«      › d�«       | j                  j                  «        t!        | «      r| j#                  «        | j                  j                  t%        |«      › d|› d�«       || _        || _        y )NrA   zLfrom torch._dynamo.graph_bytecode_inputs import get_external_object_by_indexr†  r‡  rƒ  rï   )rÉ  r8  r¯   rö  rƒ  r€  r&   r¨   r©   r6  r±   r@   rf   rŠ  r�  rY  r8   r:  r0   Ú_num_streams)r  r�  r÷  rù  Úimport_lines        rc   Úcodegen_device_guard_enterz/PythonWrapperCodegen.codegen_device_guard_enter'  sE  € ð ˜Š?Ø-Ð9Ð9Ð9ð/ð ð —<‘<×(Ñ(¨Ô5Ø—‘×&Ñ& {Ô3Ø�N‰NÜ;ØØ×5Ñ5ØØ.ó	õð �N‰NÜ-Ø × AÑ Aóôô
 �=‰=×1Ò1à×)Ñ)Ô+Ø×&Ñ&×0Ñ0ØœŸ™×*Ñ*×7Ñ7¸
ÓCÐDÀAÐFôð ×&Ñ&×0Ñ0Ô2Ü2°4Ô8à×0Ñ0Ô2Ø×&Ñ&×0Ñ0Ü& zÓ2Ð3Ð3EÀjÀ\ÐQRÐSôð -7ˆÔ)Ø!,ˆÕrb   c                ó   — t        | d«      r4| j                  dkD  r%| j                  t        | j                  «      «       n| j                  t	        «       «       t
        j                  j                  r| j                  j                  «        y y )Nr‘  rA   )
rý   r‘  r¯   r  r“  r&   r¨   r©   r±   r|  r  s    rc   Úcodegen_device_guard_exitz.PythonWrapperCodegen.codegen_device_guard_exitS  si   € Ü�4˜Ô(¨T×->Ñ->ÀÒ-BØ�N‰NÜ:¸4×;LÑ;LÓMõð �N‰NÔ7Ó9Ô:Ü�=‰=×1Ò1Ø×&Ñ&×2Ñ2Õ4ð 2rb   c                óÀ   — t         j                  j                  j                  x}�t	        d|› dt        |«      › �«      ‚t        |¬«      }| j                  |«       |S )a  Generate data structure for entering a CUDA Stream context.

        Args:
            stream_idx: The index number of the entering CUDA Stream context.

        Raises:
            ValueError: If this function is called while the previous stream context isn't exited.
        z!Nested stream context switching: z -> rû  )r@   rf   rj   Úcurrent_stream_nameÚ
ValueErrorr1   r  r¯   )r  rü  r—  Úctx_entrances       rc   Úcodegen_cuda_stream_enterz.PythonWrapperCodegen.codegen_cuda_stream_enter]  sf   € ô $%§7¡7×#4Ñ#4×#HÑ#HÐHÐÐUÜØ3Ð4GÐ3HÈÜ" :Ó.Ð/ð1óð ô 2¸ZÔHˆØ�‰�|Ô$ØÐrb   c                ó6   — | j                  t        «       «       y)z:Generate data structure for exiting a CUDA Stream context.N)r¯   r
  r  s    rc   Úcodegen_cuda_stream_exitz-PythonWrapperCodegen.codegen_cuda_stream_exits  s   € à�‰Ô0Ó2Õ3rb   c                ó˜  — |�r,t         j                  rë| j                  j                  ddj	                  |«      z   dz   «       | j                  j                  d«       | j                  j                  «        | j                  j                  d«       | j                  j                  «        | j                  j                  d«       | j                  j                  d«       | j                  j                  d«       | j                  j                  d	dj	                  |«      z   dz   «       y | j                  j                  d
«       y )Nzreturn_vars = (r»  ú, )zfor var in return_vars:z!if isinstance(var, torch.Tensor):z#assert not var.isnan().any().item()z#assert not var.isinf().any().item()r$   zreturn (z	return ())r&   rn  rÕ  r¯   rÒ   rY  r|  )r  Úoutput_refss     rc   Úgenerate_returnz$PythonWrapperCodegen.generate_returnw  s  € ÚÜ×!Ò!Ø×!Ñ!×+Ñ+Ø%¨¯	©	°+Ó(>Ñ>ÀÑFôð ×!Ñ!×+Ñ+Ð,EÔFØ×!Ñ!×+Ñ+Ô-Ø×!Ñ!×+Ñ+Ð,OÔPØ×!Ñ!×+Ñ+Ô-Ø×!Ñ!×+Ñ+Ð,QÔRØ×!Ñ!×+Ñ+Ð,QÔRØ×!Ñ!×-Ñ-¨aÔ0à×Ñ×'Ñ'¨
°T·Y±Y¸{Ó5KÑ(KÈeÑ(SÕTà×Ñ×'Ñ'¨Õ4rb   c                 ó   — y re   ra   ©r  Úresults     rc   Úgenerate_before_suffixz+PythonWrapperCodegen.generate_before_suffix‰  rH  rb   c                óÀ   — t         j                  rNdj                  | j                  «      t	        | j                  «      dk(  rdndz   }|j                  d|› d�«       y y )Nr»  rA   rY  rÆ  z-
                runner = Runner(partitions=[z{])
                call = runner.call
                recursively_apply_fns = runner.recursively_apply_fns
                )r&   r^  rÒ   Úall_partition_namesrº   rû   )r  r£  Úall_partition_name_lists      rc   Úgenerate_after_suffixz*PythonWrapperCodegen.generate_after_suffixŒ  se   € Ü×!Ò!Ø&*§i¡i°×0HÑ0HÓ&IÜ˜4×3Ñ3Ó4¸Ò9‘¸rñ'Ð#ð �M‰Mð-Ø-DÐ,Eð Fðõð "rb   c                 ó   — y re   ra   r¢  s     rc   Úgenerate_endz!PythonWrapperCodegen.generate_endš  rH  rb   c                ó¸   — |j                   }|�1|t        v r)t        |   j                  }|� ||| j                  «       y | j                  t	        | |«      «       y re   )r£  rI   Úpythonr¯   r™  )r  rp   r¹  Úcustom_codegens       rc   Úgenerate_fallback_kernelz-PythonWrapperCodegen.generate_fallback_kernel�  sV   € à×)Ñ)ˆØÐ 7Ô.JÑ#JÜ9¸'ÑB×IÑIˆNØÐ)Ù˜t T§^¡^Ô4ØØ�‰Ô,¨T°4Ó8Õ9rb   c                ó  — |j                   D ]  }| j                  |«       Œ | j                  t        | |«      «       |j                   D ]8  }t	        |j
                  t        j                  «      sŒ(|j                  | «       Œ: y)z7Generate .out() call with pre-allocated output buffers.N)	r½  Úcodegen_allocationr¯   r¶  rw   r{   r'   ÚLayoutÚcodegen_size_asserts)r  rp   r¿  s      rc   Ú generate_extern_kernel_multi_outz5PythonWrapperCodegen.generate_extern_kernel_multi_out§  sn   € à×5Ñ5ò 	.ˆHØ×#Ñ# HÕ-ð	.à�‰Ô/°°dÓ;Ô<Ø×5Ñ5ò 	4ˆHÜ˜(Ÿ/™/¬2¯9©9Õ5Ø×-Ñ-¨dÕ3ñ	4rb   c                óÈ   — |j                  | «       | j                  t        | |«      «       t        |j                  t
        j                  «      r|j                  | «       y y re   )Úcodegen_commentr¯   r™  rw   r{   r'   r±  r²  r#  s     rc   Úgenerate_extern_kernel_allocz1PythonWrapperCodegen.generate_extern_kernel_alloc°  sI   € Ø×Ñ˜TÔ"Ø�‰Ô,¨T°4Ó8Ô9Ü�d—k‘k¤2§9¡9Ô-Ø×%Ñ% dÕ+ð .rb   c           
     ó€  — t        |j                  t        j                  «      }|j	                  «       }|j                  «       }|j                  «       }| j                  }t        j                  r	d|v rd|› �}|r5| j                  | j                  › |› ddj                  |«      › d|› �«       y | j                  | j                  › |› d|› ddj                  |«      › d|› �«       | j                  rKt        j                  r:|�7t        d   dxx   d	z  cc<   | j                  d
|j                   ›d|› d�«       y y y y )NÚview_as_complexú.clone()rº  r»  rï   rî   ÚinductorÚintermediate_hooksrA   zrun_intermediate_hooks()rw   r{   r'   Ú
NoneLayoutrh   Úget_origin_noder¨  r‡  r&   Úmemory_planningr¯   r†  rÒ   rá  Úgenerate_intermediate_hooksr   r³   )r  Úextern_kernelrŸ  Ú	no_returnr®  Úorigin_noder°  r‡  s           rc   rž  z9PythonWrapperCodegen._generate_extern_kernel_alloc_helper¶  s;  € ô ˜}×3Ñ3´R·]±]ÓCˆ	Ø#×,Ñ,Ó.ˆØ#×3Ñ3Ó5ˆØ#×3Ñ3Ó5ˆØ—‘ˆÜ×!Ò!Ð&7¸;Ñ&Fð   ˜xÐ(ˆFáØ�N‰N˜dŸl™l˜^¨K¨=¸¸$¿)¹)ÀD»/Ð9JÈ!ÈFÈ8ÐTÕUà�N‰NØ—<‘<�.  ¨S°°¸Q¸t¿y¹yÈ»Ð>OÈqÐQWÐPXÐYôð ×0Ò0Ü×6Ò6ØÐ+ä˜Ñ$Ð%9Ó:¸aÑ?Ó:Ø—‘Ø-¨k×.>Ñ.>Ð-AÀÀKÀ=ÐPQÐRõð ,ð 7ð 1rb   c                ó\   — |j                  | «       | j                  t        | |«      «       y re   )rµ  r¯   r£  r#  s     rc   Úgenerate_extern_kernel_outz/PythonWrapperCodegen.generate_extern_kernel_outÓ  s&   € ð 	×Ñ˜TÔ"Ø�‰Ô*¨4°Ó6Õ7rb   c                ó   — t         j                  j                  j                  }|j	                  ||d d d«       |j                  d|r|n|› �«       |5  | j                  |› ddj                  |«      › d�«       d d d «       y # 1 sw Y   y xY w)NÚexternzout=rº  r»  rï   )r@   rf   Úwrapper_coderý  Úset_printer_argsr�   r¯   rÒ   )r  r  ÚoutÚout_viewrŸ  r[   Ústack_tracesÚdebug_printer_managers           rc   r¬  z7PythonWrapperCodegen._generate_extern_kernel_out_helperÚ  s‚   € ô !"§¡× 4Ñ 4× BÑ BÐØ×.Ñ.¨t°V¸TÀ4ÈÔRØ�‰�d¡x™8°SÐ9Ð:Ô;Ø"ñ 	;Ø�N‰N˜f˜X Q t§y¡y°£Ð&7°qÐ9Ô:÷	;÷ 	;ñ 	;ús   Á'BÂBc                ó  ‡ — |j                   }|j                  }|rRt        j                  j                  j                  |«      }t        j                  j                  j                  |«      }|j                  j                  «       › d�}dj                  ˆ fd„|D «       «      }dj                  ˆ fd„|D «       «      }t        j                  ‰ |j                  «      }d}|› d|j                  › d�}|› d|› d|› d|› �}	|› d|	› d	�}
|
S )
Nz.data_ptr()r»  c              3  óJ   •K  — | ]  }t         j                  ‰|«      –— Œ y ­wre   ©rR  Úval_to_arg_str©rœ   Údimr  s     €rc   rž   zRPythonWrapperCodegen._generate_tma_descriptor_call_experimental.<locals>.<genexpr>ó  s   øè ø€ ÒXÈCÔ-×<Ñ<¸TÀ3×GÑXùó   ƒ #c              3  óJ   •K  — | ]  }t         j                  ‰|«      –— Œ y ­wre   rÏ  rÑ  s     €rc   rž   zRPythonWrapperCodegen._generate_tma_descriptor_call_experimental.<locals>.<genexpr>ô  s$   øè ø€ ò 
Ø?BÔ ×/Ñ/°°c×:ñ
ùrÓ  z$triton.tools.experimental_descriptorz.create_Úd_tma_descriptorrº  rï   )ÚdimsÚ
block_dimsr@   rf   rn   r"  Úoptimization_hintsÚtensorr­  rÒ   rR  rÐ  Úelement_sizeÚrank)r  ÚdescÚapply_size_hintsrÖ  r×  ÚptrrÚ  rÒ  rõ   rŸ  r	  s   `          rc   Ú*_generate_tma_descriptor_call_experimentalz?PythonWrapperCodegen._generate_tma_descriptor_call_experimentalê  sþ   ø€ Ø�y‰yˆØ—_‘_ˆ
ÙÜ—7‘7×#Ñ#×5Ñ5°dÓ;ˆDÜŸ™×)Ñ)×<Ñ<¸ZÓHˆJà—‘×.Ñ.Ó0Ð1°Ð=ˆà�y‰yÓXÐSWÔXÓXˆØ—Y‘Yó 
ØFPô
ó 
ˆ
ô ,×:Ñ:¸4À×ARÑARÓSˆØ7ˆØˆx�x §	¡	˜{Ð*:Ð;ˆØ��b˜˜˜b  ¨B¨|¨nÐ=ˆØ��Q�t�f˜AˆˆØˆrb   c                óÐ   — |j                   }|r)t        j                  j                  j	                  |«      }d}|› d�}|j
                  j                  «       › d|› �}|› d|› d�}|S )Nz/triton.tools.tensor_descriptor.TensorDescriptorz.from_tensorr»  rº  rï   )Úblock_shaper@   rf   rn   rØ  rÙ  r­  )r  rÜ  rÝ  rá  rÒ  rõ   rŸ  r	  s           rc   Ú$_generate_tma_descriptor_call_stablez9PythonWrapperCodegen._generate_tma_descriptor_call_stableþ  so   € Ø×&Ñ&ˆÙÜŸ'™'×*Ñ*×=Ñ=¸kÓJˆKàBˆØˆx�|Ð$ˆØ—+‘+×/Ñ/Ó1Ð2°"°[°MÐBˆØ��Q�t�f˜AˆˆØˆrb   c                ó¶   — t        |t        j                  «      r| j                  ||«      S t        |t        j                  «      sJ ‚| j                  ||«      S re   )rw   r'   ÚTMADescriptorExperimentalrß  ÚTMADescriptorStablerâ  )r  rÜ  rÝ  s      rc   Ú_generate_tma_descriptor_callz2PythonWrapperCodegen._generate_tma_descriptor_call	  sW   € Ü�dœB×8Ñ8Ô9Ø×BÑBØÐ&óð ô ˜d¤B×$:Ñ$:Ô;Ð;Ð;Ø×<Ñ<¸TÐCSÓTÐTrb   c                ó‚   — | j                  |«      }|j                  › d|› | j                  › �}| j                  |«       y ©Nrî   )ræ  r³   r‡  r¯   )r  rÜ  r	  r²   s       rc   Úgenerate_tma_descriptorz,PythonWrapperCodegen.generate_tma_descriptor  s:   € Ø×1Ñ1°$Ó7ˆØ—)‘)�˜C ˜v d§k¡k ]Ð3ˆØ�‰�tÕrb   c                ó:   — | j                  t        | |«      «       y re   )r¯   r›  r#  s     rc   Úgenerate_scatter_fallbackz.PythonWrapperCodegen.generate_scatter_fallback  s   € Ø�‰Ô*¨4°Ó6Õ7rb   c	                óò   — |› ddj                  t        t        |«      «      › �}	|j                  d«      r|	dj                  dg|z   «      z  }	n|r|	dt	        |«      › �z  }	|	dz  }	| j                  |	«       y )Nrº  rY  zaten.scatter_reducer»  rÆ  z	, reduce=rï   )rÒ   ÚmaprW   r  r?  r¯   )
r  r´   r\   r©  r£  r   rF  r»   r[   r²   s
             rc   r¢  z/PythonWrapperCodegen._generate_scatter_fallback  s{   € ð %Ð% Q s§x¡x´´C¸Ó0@Ó'AÐ&BÐCˆØ×(Ñ(Ð)>Ô?Ø�D—I‘I˜r˜d V™mÓ,Ñ,‰DáØ˜)¤D¨£L >Ð2Ñ2�Ø�‰ˆØ�‰�tÕrb   c                ó`  — g }|j                   dd  }t        |«      }t        |j                  «      D ]^  \  }}|j                  |   �9t	        |«      }t        |t        j                  «      sJ ‚|j                  |«       ŒN|j                  d «       Œ` | j                  t        | ||«      «       y )Nr$   )r\   ÚiterÚ	enumeraterz  rŽ  rw   r'   r)   r�   r¯   r‹  )r  rp   rz  Úvalid_indicesÚiter_valid_indicesr   Ú_r6  s           rc   Úgenerate_index_put_fallbackz0PythonWrapperCodegen.generate_index_put_fallback.  sš   € à*,ˆØŸ™ A B˜ˆÜ! -Ó0ÐÜ˜dŸl™lÓ+ò 	%‰DˆAˆqØ�|‰|˜A‰Ð*ÜÐ/Ó0�Ü! %¬¯©Ô3Ð3Ð3Ø—‘˜uÕ%à—‘˜tÕ$ð	%ð 	�‰Ô+¨D°$¸Ó@ÕArb   c                ó|   — ddj                  |«      › d�}||||g}| j                  | j                  ||«      «       y )Nr}  r»  r~  )rÒ   r¯   Úwrap_kernel_call)r  r  r½   rz  r—  Ú
accumulateÚindices_strrŸ  s           rc   r•  z1PythonWrapperCodegen._generate_index_put_fallback=  sA   € Ø˜$Ÿ)™) GÓ,Ð-¨QÐ/ˆØ�; ¨
Ð3ˆØ�‰�t×,Ñ,¨V°TÓ:Õ;rb   c           
     ó`   — | j                  |› d|› ddj                   |«       «      › d�«       y )Nrî   rº  r»  rï   )r¯   rÒ   )r  Úbuf_namer£  Úget_argsÚop_overloadrÏ  r¯  s          rc   Ú,generate_fallback_kernel_with_runtime_lookupzAPythonWrapperCodegen.generate_fallback_kernel_with_runtime_lookupB  s2   € ð 	�‰˜(˜ 3Ð'9Ð&:¸!¸D¿I¹IÁhÃjÓ<QÐ;RÐRSÐTÕUrb   c                óf   — t        d«      5  | j                  |«      cd d d «       S # 1 sw Y   y xY w)NzPythonWrapperCodegen.generate)r   Ú	_generate©r  Úis_inferences     rc   ÚgeneratezPythonWrapperCodegen.generateM  s,   € ÜÐ9Ó:ñ 	0Ø—>‘> ,Ó/÷	0÷ 	0ò 	0ús   Œ'§0c                ó&   — t         j                  ryy)Nr$   rA   )r&   r^  r  s    rc   Úget_wrapper_call_indentz,PythonWrapperCodegen.get_wrapper_call_indentQ  s   € Ü×!Ò!Øàrb   c              #  ót   ‡‡K  — | j                   }dˆˆfd„}	 || _         |–— || _         y# || _         w xY w­w)z®Temporarily redirect self.writeline to write into `target`.

        Plain strings go through `writeline`.
        WrapperLines are codegen'd into `target` directly.
        c                óZ   •— t        | t        «      r| j                  ‰«       y  ‰| «       y re   )rw   rT   r\  )r²   Útargetr¯   s    €€rc   Údispatchz4PythonWrapperCodegen.set_writeline.<locals>.dispatchd  s    ø€ Ü˜$¤Ô,à—‘˜VÕ$á˜$•rb   N)r²   z2LineContext | DeferredLineBase | WrapperLine | strrÈ   r>  rt  )r  r  r¯   Úoldr  s    ``  rc   Úset_writelinez"PythonWrapperCodegen.set_writelineW  s6   ùè ø€ ð �n‰nˆö	 ð	!Ø%ˆDŒNØŠNà ˆD�Nø˜SˆD�Nüs   „8™, ¤8¬	5µ8c                óÒ   — | j                   j                  }t        j                  j                  r| j
                  j                  |«       y | j                  j                  |«       y re   )r÷  Úkernel_defsr&   r¨   r©   rÖ  rû   rÑ  )r  r  s     rc   Ú_write_multi_kernel_defsz-PythonWrapperCodegen._write_multi_kernel_defsq  sF   € Ø×-Ñ-×9Ñ9ˆÜ�=‰=×1Ò1Ø×%Ñ%×,Ñ,¨[Õ9à�K‰K×Ñ˜{Õ+rb   c                ó
  — t         j                  r| j                  «        t        j                  «       5 }|j                  | j                  j                  «       «       t         j                  r| j                  |«       t         j                  r| j                  «        | j                  |«       t         j                  j                  r*t         j                  j                  s| j                  «        | j!                  | j                  | j                  j"                  «      5  | j$                  D ]I  }t'        |t(        «      r|j+                  | j                  «       Œ/| j                  j#                  |«       ŒK 	 d d d «       | j-                  «        | j/                  «       }| j1                  «        t         j                  j2                  rA| j                  j#                  t4        j6                  j8                  j;                  «       «       t         j                  r| j=                  «        t         j                  j>                  r| jA                  «        t         j                  j                  r*t         j                  j                  s| jC                  «        t         j                  j                  r| jE                  «        t         jF                  r+t         jH                  s| j                  j#                  d«       | jK                  |«       d d d «       tM        «       }|jO                  | jP                  «       |j#                  d«       |jO                  | jR                  «       t4        j6                  jT                  r>t4        j6                  jH                  r$t4        j6                  jV                  r
tM        «       }|jO                  | jX                  «       | j[                  «        |jO                  | j\                  «       | j_                  «       }|j                  |«      5  |jO                  | j                  «       d d d «       | ja                  |«       |jO                  | jb                  «       | je                  |«       | jg                  |«       | ji                  |«       |jk                  «       | jl                  jk                  «       fS # 1 sw Y   �Œ^xY w# 1 sw Y   �ŒÔxY w# 1 sw Y   Œ¬xY w)Nz+nvtx._device_range_end(training_annotation)rÆ  )7r&   Úprofile_bandwidthr6  rÍ   Ú	ExitStackÚenter_contextrÕ  rÌ   Úprofiler_mark_wrapper_callÚ#generate_profiler_mark_wrapper_callÚgenerate_start_graphÚrun_wrapper_ir_passesr¨   Ústore_cubinr©   Ú!generate_reset_kernel_saved_flagsr
  r¯   rÜ  rw   rT   r\  r  rE  rG  rf  r@   rf   rŠ  rg  Úgenerate_end_graphr  Úgenerate_proton_finalizeÚ generate_save_uncompiled_kernelsÚgenerate_and_run_autotune_blockr  rˆ  r   r7   rû   rÉ  rÑ  r‰  Úis_const_graphr×  Úfinalize_prefixrÒ  r  r¤  rÓ  r¨  rª  Úadd_benchmark_harnessÚgetvaluewithlinemaprÔ  )r  r  Ústackr²   rŸ  r£  Úwrapper_call_indents          rc   rÿ  zPythonWrapperCodegen._generatex  s†  € Ü×#Ò#Ø×)Ñ)Ô+ä×!Ñ!Ó#ð .	. uØ×Ñ × 1Ñ 1× 8Ñ 8Ó :Ô;Ü×0Ò0Ø×8Ñ8¸Ô?Ü×'Ò'Ø×)Ñ)Ô+à×&Ñ& |Ô4ä�}‰}×(Ò(´·±×1WÒ1WØ×6Ñ6Ô8ð ×#Ñ# D×$5Ñ$5°t×7HÑ7H×7RÑ7RÓSñ :Ø ŸJ™Jò :�DÜ! $¬Ô4àŸ™ T×%6Ñ%6Õ7à×)Ñ)×3Ñ3°DÕ9ñ:÷:ð ×)Ñ)Ô+à×.Ñ.Ó0ˆKØ×!Ñ!Ô#Ü�}‰}×-Ò-Ø×!Ñ!×+Ñ+¬A¯G©G×,>Ñ,>×,JÑ,JÓ,LÔMä×'Ò'Ø×'Ñ'Ô)ä�}‰}×-Ò-Ø×-Ñ-Ô/ä�}‰}×(Ò(´·±×1WÒ1WØ×5Ñ5Ô7ä�}‰}×5Ò5Ø×4Ñ4Ô6ô ×'Ò'´×0BÒ0BØ×!Ñ!×+Ñ+ØAôð × Ñ  Ô-÷].	.ôb  Ó!ˆØ�‰�d—l‘lÔ#Ø×Ñ˜ÔØ�‰�d—k‘kÔ"ô �7‰7×Ò¤§¡× 3Ò 3¼¿¹×8NÒ8NÜ#Ó%ˆFð 	�‰�d×/Ñ/Ô0Ø×ÑÔØ�‰�d—k‘kÔ"à"×:Ñ:Ó<Ðà�]‰]Ð.Ó/ñ 	-Ø�M‰M˜$×+Ñ+Ô,÷	-ð 	×#Ñ# FÔ+Ø�‰�d—k‘kÔ"Ø×"Ñ" 6Ô*à×Ñ˜&Ô!à×"Ñ" 6Ô*ð ×&Ñ&Ó(Ø×$Ñ$×8Ñ8Ó:ð
ð 	
÷}:ñ :ú÷.	.ñ .	.ú÷B	-ð 	-ús2   µC0S(Ä%ASÅ?FS(Ð.S5ÓS%	Ó S(Ó(S2Ó5S>c                ó\  ‡— | j                   j                  d«       i }t        j                  j                  r_t
        j                  j                  rEt        t
        j                  j                  «      D ��ci c]  \  }}| j                  |«      |“Œ }}}| j                   j                  «       dz   | j                  j                  «       z   Št        j                  t        j                  k(  rkt!        j"                  t%        «       dd¬«      5 }|j'                  ‰j)                  d«      «       |j*                  }ddd«       t        j,                  d«       t/        d	d
„ ˆfd„¬«       	 t1        ‰|«       yc c}}w # 1 sw Y   ŒExY w# t2        $ r}t5        d|› �«      |‚d}~ww xY w)zº
        Compose self.kernel_autotune_defs and self.kernel_autotune_calls into a single block of
        code and execute it to trigger Triton kernel compilation and auto-tuning
        zQ
            async_compile.wait(globals())
            del async_compile
        r…  z.pyF)ÚdirrÓ  Údeletezutf-8NzAuto-tuning code written to %sÚartifactc                 ó   — dddœS )NÚ&inductor_autotune_at_compile_time_codeÚstring)r³   Úencodingra   ra   rb   rc   r¾   zFPythonWrapperCodegen.generate_and_run_autotune_block.<locals>.<lambda>ñ  s   € Ø@Ø$ñ!€ rb   c                 ó   •— ‰ S re   ra   )Útuning_codes   €rc   r¾   zFPythonWrapperCodegen.generate_and_run_autotune_block.<locals>.<lambda>õ  s   ø€ ˜{€ rb   )Úmetadata_fnÚ
payload_fnz%Failed to run autotuning code block: )rÖ  rû   r&   r¨   r©   r@   rf   Úautotuning_inputsrð  Úget_autotuning_input_namerÔ   r±   r(   ÚlevelÚloggingÚDEBUGÚtempfileÚNamedTemporaryFiler   ÚwriteÚencoder³   Údebugr   ÚexecÚ	ExceptionÚRuntimeError)r  Úscoper˜  ÚvÚfÚ	file_pathÚer+  s          @rc   r  z4PythonWrapperCodegen.generate_and_run_autotune_blockÍ  sˆ  ø€ ð
 	×!Ñ!×(Ñ(ðô	
ð ˆÜ�=‰=×1Ò1´a·g±g×6OÒ6Oô (¬¯©×(AÑ(AÓB÷á�C˜ð ×.Ñ.¨sÓ3°QÑ6ðˆEñ ð
 ×%Ñ%×.Ñ.Ó0Øñà×(Ñ(×1Ñ1Ó3ñ4ð 	ô
 × Ñ ¤G§M¡MÒ1ô ×,Ñ,Ü“K¨°eôð #àØ—‘˜×*Ñ*¨7Ó3Ô4ØŸF™F�	÷	#ô
 ×!Ñ!Ø0Øôô 	Øñó +õ	
ð	SÜ�˜eÕ$ùó?÷#ð #ûô( ò 	SÜÐ!FÀqÀcÐJÓKÐQRÐRûð	Sús*   Á9E<Ä-FÅ/F ÆFÆ	F+ÆF&Æ&F+c                ó\   — ddl m}  || «      j                  | j                  «      | _        y )NrA   )ÚMemoryPlanner)r¾  rA  rë  rÜ  )r  rA  s     rc   Úmemory_planz PythonWrapperCodegen.memory_planý  s    € Ý2á" 4Ó(×-Ñ-¨d¯j©jÓ9ˆ�
rb   c                ó  — | j                  «       }t        j                  j                  |«      }| j                  r«t        | j                  d   t        «      rŽ| j                  d   j                  j                  |vri| j                  j                  «        | j                  rCt        | j                  d   t        «      r&| j                  d   j                  j                  |vrŒit        «       g}g }g }t        t        | j                  «      «      D �](  }| j                  |   }t        |t        «      r"|j                  |d   «      | j                  |<   ŒEt        |t        «      r}|j                  t        «       «       t         j"                  sŒ|j                  | j$                  «       t        j&                  |j                  «      5  t)        «       | _        d d d «       ŒÒt        |t*        «      sŒã|j                  |j                  «       «       t         j"                  s�Œ|j                  «       | _        �Œ+ |j                  |j                  «       «       t        |«      dk(  sJ ‚t-        d„ |D «       «      }y # 1 sw Y   �ŒxxY w)Nr…  r   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wre   )r+  )rœ   Úss     rc   rž   z9PythonWrapperCodegen.memory_plan_reuse.<locals>.<genexpr>'	  s   è ø€ ò +
Ø./ˆA×)Õ)ñ+
ùs   ‚)rC  r@   rf   Ú_get_output_namesrÜ  rw   rç  rp   r³   r1  r!  rþ  rº   rë  rQ  r�   r&   rB  r9  Úset_graph_handlerr  rw  Úsum)	r  r¯  Ú	out_namesÚplanning_statesÚpast_planning_statesÚpeak_estimate_stackr   r²   Ú_total_allocated_buffer_sizes	            rc   Úmemory_plan_reusez&PythonWrapperCodegen.memory_plan_reuse	  s  € Ø×(Ñ(Ó*ˆÜ—G‘G×-Ñ-¨gÓ6ˆ	ð �JŠJÜ˜4Ÿ:™: b™>Ô+=Ô>à—
‘
˜2‘×#Ñ#×(Ñ(°	Ñ9ð �J‰J�N‰NÔð �JŠJÜ˜4Ÿ:™: b™>Ô+=Ô>à—
‘
˜2‘×#Ñ#×(Ñ(°	Ò9ô /Ó0Ð1ˆØ!ÐØ;=ÐÜ”s˜4Ÿ:™:“Ó'ó 	CˆAØ—:‘:˜a‘=ˆDÜ˜$Ô 2Ô3Ø $§	¡	¨/¸"Ñ*=Ó >�—
‘
˜1’Ü˜DÔ"3Ô4Ø×&Ñ&Ô':Ó'<Ô=Ü×,Ó,Ø'×.Ñ.¨t×/AÑ/AÔBÜ×,Ñ,¨T¯Z©ZÓ8ñ EÜ-BÓ-D˜Ô*÷Eð Eä˜DÔ"2Õ3Ø$×+Ñ+¨O×,?Ñ,?Ó,AÔBÜ×,Ô,Ø)<×)@Ñ)@Ó)B�DÖ&ð	Cð 	×#Ñ# O×$7Ñ$7Ó$9Ô:Ü�?Ó# qÒ(Ð(Ð(ô
 (+ñ +
Ø3Gô+
ó (
Ñ$÷Eñ Eús   ÇJÊJ	c                ó¨   — |r!t         j                  r| j                  «        y t         j                  rt	        «       | _        | j                  «        y re   )r&   r¾  rB  rB  r  r9  rN  r   s     rc   r  z*PythonWrapperCodegen.run_wrapper_ir_passes+	  s9   € áœF×2Ò2Ø×ÑÕä×(Ò(Ü%:Ó%<�Ô"Ø×"Ñ"Õ$rb   c           	     óX  ‡‡
— | j                   Š
t        j                  ˆ
fd„«       }t        j                  ˆ
fd„«       }dˆˆ
fd„}t        |t        j
                  «      rOt        |t        j                  «      r|‰v ry ‰
j                  |› d|› �«       ‰j                  |«        ||«       y t        |t        j                  «      rít        |j                  «       «      D ]^  \  }}t        |t        j                  «      sŒ!|‰vsŒ&‰
j                  |› d ||«      › d|› d�«       ‰j                  |«        ||«       Œ` t        |j                  «       «      D ]V  \  }}	t        |	t        j                  «      sŒ!|	‰vsŒ&‰
j                  |	› d ||«      › d|› d�«       ‰j                  |	«       ŒX y t        |t        j                  t        j                  t        j                   f«      ry t"        j$                  j&                  j(                  ry t+        dt-        |«      › �«      ‚)	Nc                ó<   •— ‰j                  | › d| › d�«       | › d�S )Nz_size = z.size()Ú_sizert  ©r³   r[  s    €rc   ÚsizeofzDPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.sizeof<	  s(   ø€ à�N‰N˜d˜V 8¨D¨6°Ð9Ô:Ø�V˜5�>Ð!rb   c                ó<   •— ‰j                  | › d| › d�«       | › d�S )Nz
_stride = z	.stride()Ú_stridert  rS  s    €rc   ÚstrideofzFPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.strideofA	  s)   ø€ à�N‰N˜d˜V :¨d¨V°9Ð=Ô>Ø�V˜7Ð#Ð#rb   c                óx  •— ddl m}m} t        j                  j
                  j                  j                  j                  «       D ]r  \  }}|| k(  sŒt        |t        j                  «      sŒ'|‰vsŒ, |||j                  |j                  f«      rŒL‰j                  |› d| › �«       ‰j                  |«       Œt y )Nr   r!   rî   )Útorch.utils._sympy.symbolr"   r#   r@   rf   rn   Ú	shape_envÚreplacementsrì  rw   r–   ÚSymbolÚUNBACKED_INTÚUNBACKED_FLOATr¯   rÓ   )Úsymr"   r#   rü   ÚtgtÚ
bound_varsr[  s        €€rc   Úmaybe_emit_replacement_aliasesz\PythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.maybe_emit_replacement_aliasesF	  s›   ø€ ÷ GäŸG™G×,Ñ,×6Ñ6×CÑC×IÑIÓKò 
(‘��Sà˜3“JÜ" 3¬¯©Õ5Ø :Ò-Ù*Ø˜d×/Ñ/°×1DÑ1DÐEõð —N‘N c U¨#¨c¨UÐ#3Ô4Ø—N‘N 3Õ'ñ
(rb   rî   r}  r~  zUnknown value type: )r_  r  rÈ   r>  )rÒ  rE  rô  rw   r–   r
   r\  r¯   rÓ   r'   r‹   rð  rM  rN  rN  rO  rP  r  r  r&   r^  r�  r¢   )r  r³   r  ra  rT  rW  rb  rÒ  rˆ   r‰   r[  s      `      @rc   Úcodegen_input_symbol_assignmentz4PythonWrapperCodegen.codegen_input_symbol_assignment4	  sÏ  ù€ ð �{‰{ˆä	�‰ó	"ó 
ð	"ô 
�‰ó	$ó 
ð	$ö	(ô( �eœUŸZ™ZÔ(Ü˜e¤U§\¡\Ô2°e¸zÑ6IØØ�N‰N˜e˜W C¨ vÐ.Ô/Ø�N‰N˜5Ô!Ù*¨5Õ1Ü˜œrŸ|™|Ô,Ü& u§~¡~Ó'7Ó8ò 9‘	��TÜ˜d¤E§L¡LÕ1°dÀ*Ò6LØ—N‘N d V¨3©v°d«|¨n¸A¸c¸UÀ!Ð#DÔEØ—N‘N 4Ô(Ù2°4Õ8ð	9ô
  )¨×)9Ñ)9Ó);Ó<ò +‘��VÜ˜f¤e§l¡lÕ3¸ÀjÒ8PØ—N‘N f X¨S±¸$³Ð0@ÀÀ#ÀÀaÐ#HÔIØ—N‘N 6Õ*ñ+ô Ø”B×&Ñ&¬×(9Ñ(9¼2×;OÑ;OÐPô
ð ä�‰×%Ñ%×5Ò5Øä$Ð';¼DÀ»K¸=Ð%IÓJÐJrb   c           	     ó  — t        t        j                     «       }| j                  «       }|j	                  «       D ��cg c]$  \  }}t        |t        j                  «      sŒ!||f‘Œ& c}}|j	                  «       D ��cg c]$  \  }}t        |t        j                  «      rŒ!||f‘Œ& c}}z   }|D ]  \  }}| j                  |||«       Œ 	 	 	 	 dd„}|D ])  \  }	}t        |t        j                  «      sŒ! |||«       Œ+ yc c}}w c c}}w )z$Assign all symbolic shapes to localsc                óP  — t        j                  | j                  «       | j                  «       g«      D ]k  }t	        |t
        «      rt	        |t        j                  «      rŒ.|j                  D �cg c]	  }||vsŒ|‘Œ }}t        |«      dkD  sŒ[t        d|› d|› d�«      ‚ y c c}w )Nr   zFor z, expected z to have been codegen-ed.)r   Úfrom_iterablerM  rN  rw   r
   r–   r\  Úfree_symbolsrº   r�  )r  ra  Úexprr_  Úundefined_symbolss        rc   Ú_verify_input_symbol_assignmentzLPythonWrapperCodegen.codegen_inputs.<locals>._verify_input_symbol_assignment…	  s®   € ô ×+Ñ+¨U¯^©^Ó-=¸u×?OÑ?OÓ?QÐ,RÓSò 
�Ü! $¬Ô-´¸DÄ%Ç,Á,Ô1OØð $(×#4Ñ#4ö%Ø¸À:Ò8M’Cð%Ð!ð %ô Ð(Ó)¨AÓ-Ü(Ø˜t˜f KÐ0AÐ/BÐB[Ð\óð ñ
ùò%s   Á0	B#Á:B#N)r  úir.TensorBoxra  úOrderedSet[sympy.Symbol])	r   r–   r\  rJ  rì  rw   rc  r'   r‹   )
r  ra  r  Úkr<  r\   r³   r  rj  ró  s
             rc   ri  z#PythonWrapperCodegen.codegen_inputst	  sþ   € ä¤§¡Ñ-Ó/ˆ
ð ×,Ñ,Ó.ˆà+×1Ñ1Ó3÷
Ù�q˜!´zÀ!ÄUÇ\Á\Õ7RˆQ�ŠFó
à ,× 2Ñ 2Ó 4×X™˜˜1¼JÀqÌ%Ï,É,Õ<Wˆa�ŠVÓXñYˆð "ò 	J‰KˆD�%Ø×0Ñ0°°u¸jÕIð	Jð	Øð	à0ó	ð& ò 	?‰HˆAˆuÜ˜e¤R§\¡\Ô2ØÙ+¨E°:Õ>ñ	?ùó3
ùãXs   ¿"DÁ"DÁ>"DÂ!Dc                ó~  — t        |t        j                  «      r£t        |t        j
                  «      rˆ|| j                  v ry | j                  j                  |«       t        j                  j                  j                  |   }t        ||«      }| j                  t        | |t        j                  «      «       y y y re   )rw   r–   r\  r"   r#   ÚPRECOMPUTED_SIZErU  rÓ   r@   rf   rn   Úinv_precomputed_replacementsr  r¯   r¦  )r  r_  rh  r§  s       rc   Úensure_size_computedz)PythonWrapperCodegen.ensure_size_computed�	  s�   € Ü�cœ5Ÿ<™<Ô(¬^¸CÄ×AVÑAVÔ-WØ�d×)Ñ)Ñ)ØØ×Ñ×#Ñ# CÔ(Ü—7‘7×#Ñ#×@Ñ@ÀÑEˆDÜ! # tÓ,ˆCØ�N‰NÔ.¨t°S¼!¿'¹'ÓBÕCð .XÐ(rb   c                 ó   — y re   ra   r  s    rc   r  z$PythonWrapperCodegen.finalize_prefix¦	  r-  rb   ©ro   c               ó   — t        d«      ‚)Nz8codegen_cpp_sizevar is only implemented for cpp_wrapper!)r:  ©r  r½   ro   s      rc   Úcodegen_cpp_sizevarz(PythonWrapperCodegen.codegen_cpp_sizevar©	  s   € ÜÐUÓVÐVrb   c               ó   — t        ||¬«      S )Nrs  )Úpexprru  s      rc   Úcodegen_python_sizevarz+PythonWrapperCodegen.codegen_python_sizevar¬	  s   € Ü�Q Ô*Ð*rb   c                ó$   — | j                  |«      S re   )ry  rD  s     rc   Úcodegen_sizevarz$PythonWrapperCodegen.codegen_sizevar¯	  s   € Ø×*Ñ*¨1Ó-Ð-rb   c                ó   — |› d|› d�S )Nr}  r~  ra   )r  r‚  r³   r6  s       rc   r€  z)PythonWrapperCodegen.codegen_tuple_access²	  s   € Ø�˜1˜U˜G 1Ð%Ð%rb   c                ó¨   — g t        | j                  |«      ¢}t        |«      dk(  ryt        |«      dk(  r	d|d   › d�S ddj                  |«      › d�S )Nr   ú()rA   rº  rž  r»  rï   )rí  ry  rº   rÒ   )r  rT  Úpartss      rc   r§   z/PythonWrapperCodegen.codegen_python_shape_tupleµ	  s^   € Ø:”#�d×1Ñ1°5Ó9Ð:ˆÜˆu‹:˜Š?ØÜˆu‹:˜Š?Ø�u˜Q‘x�j Ð$Ð$Ø�4—9‘9˜UÓ#Ð$ AÐ&Ð&rb   c                ó$   — | j                  |«      S re   )r§   )r  rT  s     rc   rQ  z(PythonWrapperCodegen.codegen_shape_tuple½	  s   € Ø×.Ñ.¨uÓ5Ð5rb   c                ó°   — dj                  dj                  |t        |«      t        |«      | j	                  |«      | j	                  |«      g«      «      g fS )Nzalloc_from_pool({})r»  )rÿ  rÒ   rx  rW   r§   )r  r³   rŠ   r   rT  r‰   s         rc   Úcodegen_alloc_from_poolz,PythonWrapperCodegen.codegen_alloc_from_poolÀ	  s_   € ð %×+Ñ+Ø�I‰IàÜ˜&“MÜ˜“JØ×3Ñ3°EÓ:Ø×3Ñ3°FÓ;ðóó

ð ð
ð 
	rb   c                ó~  ‡ — t        |«      \  }}}	}
}ˆ fd„}|j                  «       }|xr ||	k(  }|r||k(  xr ||k(  }|
}n[||j                  j                  k(  xr4 ||j                  j                  k(  xr ||j                  j
                  k(  }|j                  }|r|�||k7  r	d|› d|› d�S |› S  |||||||«      S )Nc           	     ó¬   •— ‰
j                  |«      }‰
j                  |«      }‰
j                  |«      }d| › d|› d|› d|› d�	}	|�||k7  r	d|	› d|› d�S |	S )Nzreinterpret_tensor(r»  rï   úaten.view.dtype()r§   r{  )r³   Útgt_sizeÚ
tgt_strideÚ
tgt_offsetÚ
cast_dtypeÚ
base_dtyperE  ÚstÚoffrh  r  s             €rc   Úapply_reinterpretzHPythonWrapperCodegen.codegen_reinterpret_view.<locals>.apply_reinterpretî	  s|   ø€ ð ×/Ñ/°Ó9ˆAØ×0Ñ0°Ó<ˆBØ×&Ñ& zÓ2ˆCØ(¨¨¨b°°°2°b°T¸¸C¸5ÀÐBˆDØÐ%¨*¸
Ò*BØ)¨$¨¨r°*°¸QÐ?Ð?ØˆKrb   r…  r»  rï   )r“   rh   r{   rˆ   r‰   rŠ   r   )r  rŽ   rˆ   r‰   rŠ   r¯   r   Úd_sizeÚd_strideÚd_offsetÚd_dtypeÚcollapsibler�  r³   Ú	collapsedÚsame_layoutrŠ  s   `                rc   Úcodegen_reinterpret_viewz-PythonWrapperCodegen.codegen_reinterpret_viewÏ	  sì   ø€ ô8 ,¨DÓ1ñ 	9ˆ�˜( G¨[ô		ð �}‰}‹ˆØÒ6 F¨hÑ$6ˆ	ÙØ &™.Ò?¨V°xÑ-?ˆKØ ‰Jð ˜Ÿ™×(Ñ(Ñ(ò 1Ø˜dŸk™k×0Ñ0Ñ0ò1à˜dŸk™k×0Ñ0Ñ0ð ð
 Ÿ™ˆJáØÐ  U¨jÒ%8Ø)¨$¨¨r°%°¸Ð:Ð:Ø�VÐá   t¨V°V¸UÀJÓOÐOrb   c                ó8   — | j                  |› d|› d|› d�«       y )Nz.copy_(r»  rï   rt  )r  rü   ÚdstÚnon_blockings       rc   Úcodegen_device_copyz(PythonWrapperCodegen.codegen_device_copy
  s!   € Ø�‰˜#˜˜g c U¨"¨\¨N¸!Ð<Õ=rb   c                ó†   — d|› �}| j                  |› d�«       | j                  |› d�«       | j                  |› d�«       y )NÚ_d2h_event_z = torch.Event()z	.record()z.synchronize()rt  )r  rX   Ú	event_vars      rc   Úsync_d2h_copyz"PythonWrapperCodegen.sync_d2h_copy
  sH   € Ø! + Ð/ˆ	Ø�‰˜)˜Ð$4Ð5Ô6Ø�‰˜)˜ IÐ.Ô/Ø�‰˜)˜ NÐ3Õ4rb   c                ó”   — |j                  «       }|j                  d«      }| j                  t        | |||j                  «      «       y r#  )rh   Ú
input_namer¯   rv  rz  )r  rp   rw  rx  s       rc   Úcodegen_multi_outputz)PythonWrapperCodegen.codegen_multi_output
  s6   € Ø—m‘m“oˆØ—?‘? 1Ó%ˆØ�‰” t¨[¸(ÀDÇLÁLÓQÕRrb   c           
     ój  — |j                   › d|j                  › d|j                   › d|j                   › �}|rd|j                  › d|› d�}| j                  |j                  › d|j                  › d|j
                  › d|› d	�«       | j                  j                  t        |j                  «      «       y )
Nú + ú if z
 < 0 else zmax(0, min(r»  z))rî   z * (rï   )	r6  rˆ   r¯   Úunbacked_offset_symbolÚbase_offsetÚbase_dim_striderã  rÓ   rW   )r  rp   ÚclampÚ	index_strs       rc   Úcodegen_dynamic_select_indexz1PythonWrapperCodegen.codegen_dynamic_select_index
  s«   € Ø—z‘z�l # d§i¡i [°°T·Z±Z°LÀ
È4Ï:É:È,ÐWˆ	ÙØ% d§i¡i [°°9°+¸RÐ@ˆIØ�‰Ø×*Ñ*Ð+¨3¨t×/?Ñ/?Ð.@ÀÀD×DXÑDXÐCYÐY]Ð^gÐ]hÐhiÐjô	
ð 	×"Ñ"×&Ñ&¤s¨4×+FÑ+FÓ'GÕHrb   c                ó”  ‡ ‡— ˆˆ fd„}ˆ fd„}‰j                   } |‰j                  «      } |‰j                  «      }‰ j                  |› d|› �«       ‰ j                  |› d|› �«        ||› d�|› d�‰j                  «      }‰ j                  |› d|› d�«       ‰ j
                  j                  t        ‰j                   «      «       y )	Nc           	     ó  •— ‰j                  t        dt        | ‰j                  «      «      «      }‰j                  t        dt        | ‰j                  z   ‰j                  «      «      «      }‰j                  | «      }|› d|› d|› �S )Nr   r£  z >= 0 else )r{  r   r   rˆ   )r½   ÚposÚnegÚx_condrp   r  s       €€rc   Úclamp_indexzDPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.clamp_index&
  sv   ø€ Ø×&Ñ&¤s¨1¬c°!°T·Y±YÓ.?Ó'@ÓAˆCØ×&Ñ&¤s¨1¬c°!°d·i±i±-ÀÇÁÓ.KÓ'LÓMˆCØ×)Ñ)¨!Ó,ˆFØ�U˜$˜v˜h k°#°Ð7Ð7rb   c                óZ   •— |dk(  r|› d| › �S ‰j                  |«      }d|› d| › d|› d|› �S )NrA   z - rº  r¢  z	 - 1) // )r{  )Ú	start_varÚend_varÚstepÚstep_r  s       €rc   Úcodegen_with_stepzJPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.codegen_with_step,
  sJ   ø€ Ø�qŠyØ!˜ # i [Ð1Ð1Ø×(Ñ(¨Ó.ˆEØ�w�i˜s 9 +¨S°°°yÀÀÐHÐHrb   z	_start = z_end = Ú_startÚ_endz
 = max(0, rï   )Úunbacked_size_symbolÚstartÚendr¯   r³  rã  rÓ   rW   )r  rp   r¯  rµ  r_  r¹  rº  Ú	with_steps   ``      rc   Úcodegen_dynamic_slice_sizez/PythonWrapperCodegen.codegen_dynamic_slice_size%
  s»   ù€ õ	8ô	Ið ×'Ñ'ˆÙ˜DŸJ™JÓ'ˆÙ˜$Ÿ(™(Ó#ˆØ�‰˜#˜˜i¨ wÐ/Ô0Ø�‰˜#˜˜g c UÐ+Ô,Ù%¨¨¨V n¸¸¸T°lÀDÇIÁIÓNˆ	Ø�‰˜#˜˜j¨¨°1Ð5Ô6Ø×"Ñ"×&Ñ&¤s¨4×+DÑ+DÓ'EÕFrb   c                ó:   — | j                  t        | |«      «       y re   )r¯   rq  r#  s     rc   Úcodegen_dynamic_scalarz+PythonWrapperCodegen.codegen_dynamic_scalar<
  s   € Ø�‰Ô(¨¨tÓ4Õ5rb   c                ó¾  — d„ |j                   D «       \  }t        |j                  «      dk(  r#| j                  |j                  › d|› d�«       �nkt        |j                  «      dk(  r@t        |j                  d   t        «      r#| j                  |j                  › d|› d�«       �nt        |j                  «      dk(  rãt        |j                  d   t        «      rÆ| j                  |j                  › d|› d�«       | j                  d	|j                  › d
|j                  d   j                  › d|j                  › d|j                  d   j                  › d�	«       | j                  |j                  › d|j                  › d|j                  d   j                  › �«       nt        d|j                  › �«      ‚| j                  |j                  «       › d�«       y )Nc              3  ó<   K  — | ]  }|j                  «       –— Œ y ­wre   r�  r�  s     rc   rž   z?PythonWrapperCodegen._codegen_dynamic_scalar.<locals>.<genexpr>@
  s   è ø€ Ò>¨Q�1×&Ñ&×(Ñ>ùr’  r   rî   ú.item()rA   z = 1 if z.item() else 0z_undivided = zassert z_undivided % z
 == 0, f'{z_undivided} not divisible by ú'z_undivided // úunrecognized keypath ú = None)r\   rº   Úkeypathr¯   r_  rw   r   r   rÀ  r�  rh   )r  rp   rŽ   s      rc   rs  z,PythonWrapperCodegen._codegen_dynamic_scalar?
  s†  € Ù>°$·+±+Ô>‰ˆÜˆt�|‰|Ó Ò!Ø�N‰N˜dŸh™h˜Z s¨4¨&°Ð8Ö9Ü�—‘Ó !Ò#¬
°4·<±<À±?ÄMÔ(RØ�N‰N˜dŸh™h˜Z x°¨v°^ÐDÖEÜ�—‘Ó !Ò#¬
°4·<±<À±?ÄKÔ(PØ�N‰N˜dŸh™h˜Z }°T°F¸'ÐBÔCØ�N‰NØ˜$Ÿ(™(˜ =°·±¸a±×1HÑ1HÐ0Ið JØ—x‘x�jÐ >¸t¿|¹|ÈA¹×?VÑ?VÐ>WÐWXðZôð �N‰NØ—8‘8�*˜C §¡˜z¨¸¿¹ÀQ¹×8OÑ8OÐ7PÐQõô !Ð#8¸¿¹¸Ð!GÓHÐHð 	�‰˜$Ÿ-™-›/Ð*¨'Ð2Õ3rb   c                ó  ‡ ‡— ˆˆ fd„}ˆfd„}ˆfd„}i }i }t         j                  j                  }|��t        t         j                  j                  j                  «       |«      D ]Î  \  }}	t        |	t        j                  «      sŒ!|	j                  «       }
|
j                  «       }t        |
«      }|j                  |«      }|€5t        dt        |«      › �||	j                  |	j                   i ¬«      }|||<   t#        |	j$                  «      t#        |	j'                  «       «      f|j(                  |<   ŒÐ ‰j+                  g d¢«       ‰j-                  «       5  ‰j/                  dd	¬
«       t         j                  j0                  j3                  «       D ]U  \  }}‰j5                  d|› �«        |||j7                  «       |j'                  «       |j                  |j                   «       ŒW t        t         j                  j8                  «      dkD  r^‰j5                  d«       t         j                  j8                  j3                  «       D ]"  \  }}‰j5                  d|› �«        |||«       Œ$ |j;                  «       D ]Ã  }t        |j(                  «      dk  rŒ|j                  t        j<                  j?                  |j                   «      z  }‰j5                  |j@                  › d|› d|j                  › d|j                   › d�«       |j(                  j3                  «       D ]  \  }\  }}|j@                  ||f||<   Œ ŒÅ t         j                  j                  j3                  «       D �]U  \  }}t        |tB        jD                  «      rCt        t         j                  jF                  jH                  j                  |d«      tJ        «      rŒdt        |tL        jN                  «      r‰j5                  |› d�«       Œ“t        |tB        jP                  «      r3 ||t         j                  jF                  jS                  |d¬«      «       Œàt        |tB        jT                  «      r ||d«       �Œt        |tL        jV                  «      r# ||d|j                  jX                  › d�«       �ŒBt        |tL        jZ                  «      r‰j5                  |› d�«       �Œr||v rF||   \  }}}‰j5                  |› d|› d‰ j]                  |«      › d‰ j]                  |«      › d�«       �Œ¼t         j                  jF                  j_                  |ja                  «       d¬«      }t         j                  jF                  j_                  |jc                  «       d¬«      } |||||je                  «       |jg                  «       «       �ŒX ‰j5                  ddji                  t         j                  j                  j                  «       «      › d�«       ddd«       ‰j+                  g d¢«       ‰j-                  «       5  ‰j/                  dt         j                  jj                  › d�d	¬
«       ddd«       y# 1 sw Y   ŒexY w# 1 sw Y   yxY w) z2Write out codegen for benchmarking the output codec                ó‚   •— ‰j                  | › d‰j                  |«      › d‰j                  |«      › d|› d|› d�
«       y )Nz = rand_strided(r»  ú
, device='ú	', dtype=rï   )r¯   r§   )r³   rT  r‰   r[   r   r´   r  s        €€rc   Úadd_fake_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_fake_inputW
  sT   ø€ Ø×ÑØ�&Ð(Ø×2Ñ2°5Ó9Ð:¸"Ø×2Ñ2°6Ó:Ð;ð <Ø!˜( )¨E¨7°!ð5õrb   c                ó2   •— ‰j                  | › d|› �«       y rè  rt  )r³   rò  r´   s     €rc   Úadd_expr_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_expr_input_
  s   ø€ Ø×Ñ ˜v S¨¨Ð.Õ/rb   c                ó2  •— |€‰j                  | › d�«       y dd l}	 ‰j                  | › d|j                  |«      ›d�«       y # t        t        |j
                  f$ r7}‰j                  dt        |«      › d| › dt        |«      › d�«       Y d }~y d }~ww xY w)	NrÄ  r   z = pickle.loads(rï   z.raise TypeError("Failed to pickle opaque type z for variable rð   z"))r¯   ÚpickleÚdumpsÚ	TypeErrorr  ÚPicklingErrorr¢   rW   )r³   r  rÎ  r?  r´   s       €rc   Úadd_torchbind_inputzKPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_torchbind_inputb
  s¨   ø€ Øˆ}Ø× Ñ  D 6¨Ð!1Ô2ØãðØ× Ñ  D 6Ð)9¸&¿,¹,ÀuÓ:MÐ9PÐPQÐ!RÕSøÜœ~¨v×/CÑ/CÐDò Ø× Ñ ØDÄTÈ%Ã[ÀMÐQ_Ð`dÐ_eÐegÔhkÐlmÓhnÐgoÐoqÐr÷ñ ûðús   ž&A ÁBÁ-BÂBNÚ_shared_storage_)rX   rZ   r[   r   r\   )rÆ  rÆ  zdef get_args():zP
                from torch._dynamo.testing import rand_strided
                Trã   zglobal r   zimport pickler$   z = rand_strided((z,), (1,), device='rÉ  rï   rÄ  é*   r   Fztorch.cuda.default_generators[z].graphsafe_get_state()z = torch.as_strided(r»  zreturn [r~  )rÆ  rÆ  z9def benchmark_compiled_module(args, times=10, repeat=10):zÄ
                from torch._inductor.utils import print_performance
                fn = lambda: call(list(args))
                return print_performance(fn, times=times, repeat=repeat, device='z')
                )6r@   rf   Úexample_inputsrÑ   r  r  rw   r  ÚTensorÚuntyped_storagerZ   r   r÷   rV   rº   r[   r   r(  rT  r‰   r\   Ú
writelinesrÌ   rû   Ú	constantsrì  r¯   rˆ   Útorchbind_constantsr—  Ú_utilsÚ_element_sizerX   r–   r\  rn   Úbacked_var_to_valr    r'   rN  r
   r"  rM  rO  r6  rP  r§   rØ  rM  rN  rª  rm   rÒ   r«  )r  r´   rÊ  rÌ  rÒ  Ústorage_groupsÚaliased_input_specsrÕ  r³   ÚexÚstoragerZ   Ústorage_keyÚgroupr  Útorchbind_objÚnumelrT  r‰   rú  s   ``                  rc   Úbenchmark_compiled_modulez.PythonWrapperCodegen.benchmark_compiled_moduleT
  sÞ  ù€ õ	ô	0ô	ð  GIˆØ?AÐÜŸ™×/Ñ/ˆàÑ%Ü¤§¡× 4Ñ 4× 9Ñ 9Ó ;¸^ÓLò I‘��bÜ! "¤e§l¡lÔ3ØØ×,Ñ,Ó.�Ø Ÿ™Ó)�Ü,¨WÓ5�Ø&×*Ñ*¨;Ó7�Ø�=Ü1Ø&6´s¸>Ó7JÐ6KÐ$LØ%Ø!Ÿy™yØ Ÿh™hØ!ô�Eð 38�N ;Ñ/ä&*¨2¯8©8£n´d¸2¿9¹9»;Ó6GÐ%H�—‘˜TÒ"ð#Ið( 	×ÑÒ5Ô6Ø�]‰]‹_ñ X	SØ�M‰Mðð ð	 ô ô  !Ÿw™w×0Ñ0×6Ñ6Ó8ò ‘��eð × Ñ  7¨4¨&Ð!1Ô2ÙØ˜%Ÿ*™*›,¨¯©«¸¿¹ÀeÇkÁkõð	ô ”1—7‘7×.Ñ.Ó/°!Ò3Ø× Ñ  Ô1Ü+,¯7©7×+FÑ+F×+LÑ+LÓ+Nò =Ñ'�D˜-ð ×$Ñ$ w¨t¨fÐ%5Ô6Ù'¨¨mÕ<ð	=ð (×.Ñ.Ó0ò S�Ü�u—|‘|Ó$ qÒ(ØØŸ™¬¯©×(BÑ(BÀ5Ç;Á;Ó(OÑO�Ø× Ñ Ø×(Ñ(Ð)Ð):¸5¸'ÐASÐTY×T`ÑT`ÐSaÐajÐkp×kvÑkvÐjwÐwxÐyôð .3¯\©\×-?Ñ-?Ó-Aò SÑ)�D™/˜5 &Ø16×1BÑ1BÀEÈ6Ð0RÐ'¨Ò-ñSðSô  !Ÿw™w×3Ñ3×9Ñ9Ó;ó 3‘��eÜ˜e¤U§\¡\Ô2´zÜ—G‘G×$Ñ$×6Ñ6×:Ñ:¸5À$ÓGÌô8ð Ü˜e¤R×%7Ñ%7Ô8Ø×$Ñ$¨ v¨WÐ%5Õ6Ü ¤u§z¡zÔ2ñ
 #ØœaŸg™g×.Ñ.×@Ñ@ÀÐQSÐ@ÓTõô   ¤u§{¡{Ô3ñ # 4¨Ö/Ü ¤r×'8Ñ'8Ô9Ù"ØØ8¸¿¹×9KÑ9KÐ8LÐLcÐdöô   ¤r×';Ñ';Ô<Ø×$Ñ$¨ v¨WÐ%5Ö6ØÐ0Ñ0Ø.AÀ$Ñ.GÑ+�H˜e VØ×$Ñ$Ø˜&Ð 4°X°J¸bØ×:Ñ:¸5ÓAÐBÀ"Ø×:Ñ:¸6ÓBÐCÀ1ðFöô ŸG™G×,Ñ,×?Ñ?ØŸ™Ó(°2ð @ó �Eô ŸW™W×-Ñ-×@Ñ@Ø×(Ñ(Ó*°Rð Aó �Fñ #ØØØØ×(Ñ(Ó*ØŸ™Ó)öð[3ðj ×Ñ˜x¨¯	©	´!·'±'×2FÑ2F×2KÑ2KÓ2MÓ(NÐ'OÈqÐQÔR÷qX	Sðv 	×ÑÚQô	
ð �]‰]‹_ñ 	Ø�M‰MðRô ST×RYÑRY×ReÑReÐQfð gðð
 ð ô ÷	ð 	÷}X	Sð X	Sú÷|	ð 	ús   ÅR9Y7Ø>0ZÙ7Z ÚZc                óø   — t         j                  sy| j                  |«       |j                  g d¢«       |j	                  «       5  |j                  dddt        «       › d�g«       ddd«       y# 1 sw Y   yxY w)zL
        Append a benchmark harness to generated code for debugging
        N)rÆ  rÆ  zif __name__ == "__main__":zBfrom torch._inductor.wrapper_benchmark import compiled_module_mainzargs = get_args()zcompiled_module_main('zU', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat)))r&   Úbenchmark_harnessræ  rØ  rÌ   r5   ©r  r´   s     rc   r  z*PythonWrapperCodegen.add_benchmark_harnessô
  s}   € ô ×'Ò'Øà×&Ñ& vÔ.à×ÑÒ@ÔAØ�]‰]‹_ñ 
	Ø×ÑàXØ'à0Ô1CÓ1EÐ0Fð Gmð mð	ô	÷
	÷ 
	ñ 
	ús   Á!A0Á0A9c           
     óD   — | j                  t        | |||||¬«      «       y râ  )r¯   rÜ  )r  r°  rÝ  rÞ  rß  rà  s         rc   Údefine_kernelz"PythonWrapperCodegen.define_kernel
  s*   € ð 	�‰Ü ØØØØ!ØØ-ôõ		
rb   c                ó²   — t         j                  j                  r)|r't        j                  dd|t        j
                  ¬«      }|r|› d�nd}d|› | › d|› �}|S )Nz^// z# )Úflagsr…  rÆ  z

rî   )r&   r¨   r©   ÚreÚsubÚ	MULTILINE)r°  rÝ  rÞ  Úmetadata_commentÚbodys        rc   Ú_format_kernel_definitionz.PythonWrapperCodegen._format_kernel_definition  sZ   € ô �=‰=×1Ò1±hô —v‘v˜g t¨X¼R¿\¹\ÔJˆHÙ.6˜h˜Z r™?¸BÐØÐ&Ð'¨ }°C¸°}ÐEˆØˆrb   c                ó.  — t         j                  j                  rL|rJ| j                  |||¬«      }| j                  j                  |«       t        j                  j                  ry | j                  |||¬«      }| j                  j                  |«       y )N)rÞ  )
r&   r¨   r©   ró  rÖ  rû   r@   rf   rˆ  rÑ  )r  r°  rÝ  rÞ  rß  rà  rò  s          rc   rã  z*PythonWrapperCodegen._define_kernel_helper)  s…   € ô �=‰=×1Ò1±cØ×1Ñ1Ø˜[°8ð 2ó ˆDð ×%Ñ%×,Ñ,¨TÔ2Ü�w‰w×"Ò"àà×-Ñ-Ø˜¨xð .ó 
ˆð 	�‰×Ñ˜4Õ rb   c                óN   — | j                   j                  |j                  «       y re   )r×  rû   r  )r  r³   Úsubgraph_codes      rc   Údefine_subgraph_launcher_fnz0PythonWrapperCodegen.define_subgraph_launcher_fn?  s   € Ø×!Ñ!×(Ñ(¨×)<Ñ)<Õ=rb   c                ó   — ddl m} |S )NrA   )ÚTritonKernel)r¨   rù  )Úclsrù  s     rc   Ú_get_triton_info_kernel_clsz0PythonWrapperCodegen._get_triton_info_kernel_clsB  s   € õ 	)àÐrb   c                óþ  ‡‡1‡2‡3‡4‡5‡6— ddl m}m}	m}
 ddlmŠ1m}m}m}m	} |j                  }g Š6i Š4g Š3g }ˆ3ˆ6fd„Š2d5ˆ1ˆ2ˆ4ˆfd„	}|j                  D �cg c]  }|j                  ‘Œ }}|j                  D �cg c]  }|j                  sŒ|j                  ‘Œ }}t        |«      D �]Ä  \  }}||v r || ‰1|¬«      d¬	«       Œ|‰vrŒ#‰|   }‰|   € || ‰1|¬«      d¬«       Œ@t!        |t"        j$                  «      r[t!        |t"        j&                  «      r'd|j(                  |j*                  j-                  «       fnd\  }}} || |||||¬«      «       Œµt!        |t"        j.                  «      r/ || |||j1                  «       |j-                  «       ¬«      «       Œþt!        |t"        j2                  «      rO || |||j4                  j1                  «       |j-                  «       |j6                  j8                  ¬«      «       �Œgt!        |t:        t<        j>                  f«      xr* t@        jB                  jD                  jG                  |d«      } || |||«      |¬«       �ŒÇ tI        ‰6d
‰3|jJ                  D �cg c]  }tM        |«      ‘Œ c}¬«      }|tO        jP                  t@        jB                  jS                  «       «      i ‰4¥tT        jW                  |d«      ¥tY        ‰6‰3d¬«      gdœ}|rt[        |«      |d<   |rt[        |«      |d<   t]        |«      dk(  r0|	j_                  «       }g ta        t<        jb                  |d   «      ¢} níd6ˆ5fd„}!i Š5|D �"cg c]  }"g ta        |!|"«      ¢‘Œ }}"|rt]        |«      t]        |«      k(  sJ ‚g }#te        tg        ||«      d„ d¬«      D ]Q  \  }"}$|#ji                   ||$«      g ta        tj        |"«      ¢g ta        tl        |"«      ¢g ta        tj        |"«      ¢dœ«       ŒS |
j                  |#g ta        tn        ‰5jq                  «       «      ¢dœ}g ‰5js                  «       ¢} |r|D �%cg c]  }%||%   ‘Œ	 c}%|d<   tu        |jv                  «      g}&t]        |«      dkD  rQ‰jq                  «       D ]>  }t!        |t"        j.                  t"        j2                  f«      rŒ.|&ji                  |«       Œ@ |&ji                  to        |«      «       |&jy                  to        |«      «       |�'|&ji                  |d   j1                  «       |d   f«       t[        |&«      }&|&| jz                  v r| jz                  |&   \  }'}}(|'||(| fS |› dt]        | jz                  «      › �}'|'j}                  d «      r|'j                  d «      s|'dd
 }'t�        «       })t‚        j„                  j†                  r|)j‰                  d!|'›d"�«       n|)j‰                  d!|›d"�«       |'|d#<   | j‹                  «       }*|j�                  |*j�                  «       «       |)j‘                  |*j“                  «       «       t‚        j„                  j”                  r|)j‰                  d$«       |)j‘                  d%g ta        ||«      ¢›d&|›d'|›d(�«       t—        ||«      }+t‚        j„                  j†                  r|+j™                  d)|› d*�d)|'› d*�«      }+|+j™                  d+d,«      }+t‚        jš                  r|+j™                  d-d.«      }+|+j™                  d/d0«      }+|)j‘                  |+«       t@        jB                  jS                  «       },|)j‰                  d1|,jœ                  › d2�«       tŸ        j                   |jv                  «      \  }-}.tŸ        j¢                  |jv                  «      }/d3|/› d4|.› �}0| j¥                  |'|)j§                  «       |0«       |'||f| jz                  |&<   |'||| fS c c}w c c}w c c}w c c}"w c c}%w )7aø  Codegen a user-defined Triton kernel and return its cache entry.

        Emits the ``async_compile.triton(...)`` wrapper, assigns a graph-unique
        name (with a leading dunder stripped to avoid Python class-based name
        mangling at the call site), and records the kernel in
        ``user_defined_kernel_cache``. Returns ``(name, triton_meta,
        inductor_meta, extra_launcher_call_args)``; subsequent calls with the
        same ``cache_key`` reuse the previously assigned name.
        r$   )Úconfig_to_dictÚ	FixedGridÚPrecomputedGridrA   )ÚConstexprArgÚKernelArgTypeÚSizeArgÚ	TensorArgÚTMADescriptorArgc                óJ   •— ‰j                  |«       ‰j                  | «       y re   )r�   )r˜  r§  Úarg_indicesÚ	signatures     €€rc   Úadd_to_signaturezPPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_to_signaturer  s   ø€ Ø×Ñ˜SÔ!Ø×Ñ˜sÕ#rb   c                ó¦  •— |r?t        «       r	 ‰| |«       |j                  ‰v r‰|j                     ‰|j                  <   y y |j                  ‰v sJ ‚|r>t        «       r ‰|  ‰|j                  ¬«      «       n	 ‰| |«       d‰|j                  <   y |r4t        «       r ‰|  ‰|j                  ¬«      «       d ‰|j                  <   y  ‰| |«       y )N©r³   rA   )r?   r³   )	r˜  r§  Úis_constexprÚequals_1Úequals_noner   r  rÙ  r»   s	        €€€€rc   Úadd_argzGPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_argv  sÂ   ø€ ÙÜ1Ô3ñ % S¨#Ô.à—8‘8˜vÑ%ð +1°·±Ñ*:�I˜cŸh™hÒ'ð &ð —x‘x 6Ñ)Ð)Ð)áÜ5Ô7ñ
 )¨©lÀÇÁÔ.IÕJá(¨¨cÔ2Ø*+�I˜cŸh™hÒ'Ù Ü5Ô7ñ )¨©lÀÇÁÔ.IÔJØ*.�I˜cŸh™hÒ'á$ S¨#Õ.rb   r
  T)r  N)r  Ústable)ÚexperimentalNN)r³   Úapi_typerá  r   )r³   Úbufferr   )r³   r  r   rŠ   )r  )Ú
size_dtyperz  Úargdefsra   )rz  Úpointer_range_override)r  r[   rÙ  rÕ   Úrestore_valueÚreset_to_zeror   c                óN  •— t        | t        j                  «      rdg | j                  ¢}|s| S |j	                  t
        ¬«       |D ]+  }|‰v rŒt        j                  dt        ‰«      › �«      ‰|<   Œ- t        | ‰«      S t        | t        «      sJ ‚t        j                  | «      S )N)rÀ   Ú_launcher_s)rw   r–   r
   rg  ÚsortrW   r\  rº   r>   rY   r—   )rh  Úsymbolsr_  Úextra_launcher_argss      €rc   Úrename_sizes_for_launcherzYPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.rename_sizes_for_launcher   s¦   ø€ Ü˜d¤E§J¡JÔ/Ø2 × 1Ñ 1Ð2�GÙ"Ø#˜Ø—L‘L¤S�LÔ)Ø&ò ˜ØÐ"5Ñ5Ø$Ü38·<±<Ø)¬#Ð.AÓ*BÐ)CÐDó4Ð+¨CÒ0ðô & dÐ,?Ó@Ð@Ü! $¬Ô,Ð,Ð,Ü—}‘} TÓ*Ð*rb   c                ó2   — t        | d   j                  «      S r¸   r¹   r¼   s    rc   r¾   zHPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.<lambda>  s   € ´3°q¸±t·{±{Ó3C€ rb   r¿   )r&   r¬  ÚcppÚpython_slow)Ú	grid_typeÚprecomputed_gridsr  Údeclared_constexpr_namesró  Ú__zasync_compile.triton(z, '''r°  r  zG
            @triton_heuristics.user_autotune(
                configs=z ,
                inductor_meta=z,
                triton_meta=z{,
                filename=__file__,
                custom_kernel=True,
            )
            @triton.jit
            r¶   rº  ú\z\\z"""z\"\"\"z'''z\'\'\'z''', device_str='re  z# Original path: r‡  )FFF)rh  rÇ   rÈ   rÉ   )TÚruntime.triton_heuristicsrý  rþ  rÿ  Úcommonr   r  r  r  r  r]   Úparamsr³   r  Únumrð  rw   r'   ÚTMADescriptorrå  rá  rÙ  rm   r†   rh   r*   rŽ   r{   rŠ   rY   r–   r—   r@   rf   rn   Ústatically_known_equalsrL   Ú	arg_namesrB   r,   r  Úget_current_device_or_throwr  ÚfromkeysrJ   r¦   rº   Úsetup_grid_as_argsrí  ÚsympifyrÐ   rÑ   r�   rx  rH   rW   r—  r  Úidrõ   Úextendrâ  r  Úendswithr7   r&   r¨   Úunique_user_kernel_namesr¯   rû  ÚupdateÚinductor_meta_commonrû   Úgen_common_triton_importsr  r  Úreplacerˆ  r¢   ÚinspectÚgetsourcelinesÚgetsourcefilerë  rÔ   )7r  r  rÕ   r»   Úrestore_value_argsÚreset_to_zero_argsrÖ   r  rý  rþ  rÿ  r  r  r  r  Úoriginal_nameÚequal_to_1_argsr  Úpr,  Ú
constexprsr˜  rÀ   r§  r  rá  r   r  r½   Útriton_signaturerÒ  rÓ  Úextra_launcher_call_argsr  rª   r"  Úcfgr   Ú	cache_keyr³   Úcached_inductor_metar  Útriton_info_kernel_clsr  Úcurrent_deviceró  ÚlinenoÚsrcfilerÞ  r   r  r  rÙ  r  r  s7      `                                             @@@@@@rc   Ú!define_user_defined_triton_kernelz6PythonWrapperCodegen.define_user_defined_triton_kernelL  sÕ  þ€ ÷&	
ñ 	
÷
	
õ 	
ð Ÿ™ˆØ)+ˆ	Ø$&ˆ	Ø!#ˆØ%'ˆõ	$÷"	/ð "	/ðH &,§]¡]Ö3 �Q—V“VÐ3ˆ	Ð3Ø%+§]¡]ÖE °a·n³n�a—e“eÐEˆ
ÐEÜ! )Ó,ó 9	G‰HˆC�Ø�jÑ Ù˜™\¨sÔ3À$ÕGØà˜&Ñ Øà˜‘+ˆCà�c‰{Ð"Ù˜™\¨sÔ3ÀÖFä˜c¤2×#3Ñ#3Ô4ô & c¬2×+AÑ+AÔBð " 3§?¡?°C·J±J×4HÑ4HÓ4JÑKà9ñ 1�H˜k¨5ñ
 ØÙ(Ø!$Ø%-Ø(3Ø"'ô	õô   ¤R§Y¡YÔ/ÙØÙ!Ø!$Ø#&§<¡<£>Ø"%§-¡-£/ôõô   ¤R×%7Ñ%7Ô8ñ ØÙ!Ø!$Ø#&§8¡8×#4Ñ#4Ó#6Ø"%§-¡-£/Ø#&§:¡:×#4Ñ#4ô	öô  *Øœc¤5§=¡=Ð1ó ò  äŸ'™'×*Ñ*×BÑBØØóð ñ ˜C¡¨¨cÓ!2¸X×Fðs9	Gôv -ØØØØ)/×)9Ñ)9Ö: A”W˜Q•ZÒ:ô	
Ðð *Ü&×-Ñ-¬a¯g©g×.QÑ.QÓ.SÓTðØðä—-‘- °Ó3ðô
 ØØ'Ø+-ôðñ'
ˆñ. Ü+0Ð1CÓ+DˆK˜Ñ(áÜ+0Ð1CÓ+DˆK˜Ñ(äˆu‹:˜Š?à,5×,HÑ,HÓ,JˆMØ'F¬¬U¯]©]¸EÀ!¹HÓ)EÐ'FÑ$õ+ð  EGÐØINÖOÀÐ<”sÐ4°dÓ;Ò<ÐOˆEÐOáœS ›Z¬3¨w«<Ò7Ð7Ð7Ø "ÐÜ#Ü�E˜7Ó#Ñ)CÈTôò 
‘	��cð "×(Ñ(á"0°Ó"5Ø"5¤C¬¨tÓ$4Ð"5Ø2¤¤U¨DÓ!1Ð2Ø':¬¬U°DÓ)9Ð':ñ	õð
ð -×5Ñ5Ø%6Ø'P¬¬SÐ2E×2LÑ2LÓ2NÓ)OÐ'PñˆMð
 (EÐ)<×)AÑ)AÓ)CÐ'DÐ$áà&0ö9Ø!"�	˜!“ò9ˆMÐ4Ñ5ô
 ˜VŸY™Y›-˜ˆ	Üˆw‹<˜!ÒØ—}‘}“ò *�ä! #¬¯	©	´2×3EÑ3EÐ'FÕGØ×$Ñ$ SÕ)ð*ð 	×Ñœ˜[Ó)Ô*Ø×Ñœ˜]Ó+Ô,àÐ&Ø×Ñ˜o¨aÑ0×9Ñ9Ó;¸_ÈQÑ=OÐPÔQä˜)Ó$ˆ	Ø˜×6Ñ6Ñ6Ø6:×6TÑ6TØñ7Ñ3ˆD�+Ð3ð ØØ$Ø(ð	ð ð  � ¤# d×&DÑ&DÓ"EÐ!FÐGˆð �?‰?˜4Ô ¨¯©°tÔ)<Ø˜˜�8ˆDä(Ó*ˆÜ�=‰=×1Ò1Ø×%Ñ%Ð(=¸d¸XÀUÐ&KÕLà×%Ñ%Ð(=¸mÐ=NÈeÐ&TÔUà'+ˆ�mÑ$Ø!%×!AÑ!AÓ!CÐØ×ÑÐ3×HÑHÓJÔKà×ÑÐ5×OÑOÓQÔRÜ�=‰=×)Ò)Ø×%Ñ%Ð&DÔEà×Ñðà8œ3˜~¨wÓ7Ð8Ð;ð <Ø,Ð/ð 0Ø(˜Oð ,ð		ô	
ô OØ�Oó
ˆ
ô �=‰=×1Ò1à#×+Ñ+¨d°=°/ÀÐ,CÀtÈDÈ6ÐQRÀ^ÓTˆJØ×'Ñ'¨¨fÓ5ˆ
Ü×Òð $×+Ñ+¨E°;Ó?ˆJØ×'Ñ'¨¨{Ó;ˆ
Ø×Ñ˜zÔ*äŸ™×<Ñ<Ó>ˆØ×!Ñ!Ð$5°n×6IÑ6IÐ5JÈ"Ð"MÔNÜ×*Ñ*¨6¯9©9Ó5‰	ˆˆ6Ü×'Ñ'¨¯	©	Ó2ˆØ& w i¨q°°Ð9ˆØ×ÑØØ×$Ñ$Ó&Øô	
ð 6:¸;ÈÐ4Vˆ×&Ñ& yÑ1Ø�[ -Ð1IÐIÐIùòO 4ùÚEùò@ ;ùòl Pùò.9s$   Á_&Á4_+Â_+É?_0Í_5Ñ_:c                óü   — |› d|j                   › d�}|�|d|› �z  }t        j                  |dd¬«      }t        ||j                  «      }|dk(  }|s*| j                  t        | |t        j                  «      «       |S )Nró  rå  T)Ú
is_integerÚis_positiverÆ  )	rÒ  r–   r\  r  rå  r¯   r¦  r@   rf   )r  r°  ÚtreerÓ  Úsym_namer_  r§  Úis_benchmark_kernels           rc   Úgenerate_numel_exprz(PythonWrapperCodegen.generate_numel_exprƒ  s   € Ø!�] ! D§K¡K =°Ð6ˆØÐØ˜!˜F˜8˜Ñ$ˆHÜ�l‰l˜8°À$ÔGˆô ˜c 4§:¡:Ó.ˆà)¨RÑ/ÐÙ"Ø�N‰NÔ.¨t°S¼!¿'¹'ÓBÔCàˆ
rb   c                ój   — | j                  |j                  › dt        |j                  «      › �«       y rè  )r¯   r  rx  r  )r  r§  rf   s      rc   r©  z7PythonWrapperCodegen._generate_symbolic_call_arg_helper˜  s)   € ð 	�‰˜#Ÿ)™)˜ C¬¨c¯n©nÓ(=Ð'>Ð?Õ@rb   c                óÂ  — |j                  «       }t        | |«      }|j                  t        j                  k(  r| j                  |«       �n1|j                  t        j                  k(  r2| j                  |«       | j                  | j                  |«      «       nâ|j                  t        j                  k(  r°| j                  j                  |«      }|rRt        |t        «      rt        |j                  t        «      sJ ‚t        j                  |j                  |«      |_        nV| j                  |«       | j                  | j                  |«      «       || j                  |<   nt        |j                  «      ‚t         j"                  j$                  rÈ| j&                  j                  t(        j+                  | ||j,                  |j.                  t0        j2                  j4                  j7                  |j8                  «      fd¬«      «       |j                  t        j                  k7  r0| j&                  j                  t(        j                  | |«      «       y y y )N)rA   )rT  r‰   )rh   r0  Ú	zero_moderG   ÚUNINITIALIZEDr¯   ÚZERO_ON_CALLÚmake_zero_bufferÚZERO_PER_GRAPHrù  r÷   rw   rp   rF   Úmaximumr�  r&   r¨   r©   r±   rR  Úmake_allocationr[   r   r@   rf   rn   r"  r   )r  Úwsr³   r²   Úpriors        rc   Úgenerate_workspace_allocationz2PythonWrapperCodegen.generate_workspace_allocation�  sÇ  € Ø�{‰{‹}ˆÜ˜D "Ó%ˆØ�<‰<Ô,×:Ñ:Ò:Ø�N‰N˜4Ö Ø�\‰\Ô.×;Ñ;Ò;Ø�N‰N˜4Ô Ø�N‰N˜4×0Ñ0°Ó6Õ7Ø�\‰\Ô.×=Ñ=Ò=Ø×-Ñ-×1Ñ1°$Ó7ˆEÙÜ! %¬Ô6¼:Ø—J‘J¤ô<ð ð ô *×1Ñ1°%·*±*¸bÓA�•
à—‘˜tÔ$Ø—‘˜t×4Ñ4°TÓ:Ô;Ø26�×)Ñ)¨$Ò/ä  §¡Ó.Ð.ä�=‰=×1Ò1Ø×&Ñ&×0Ñ0Ü$×4Ñ4ØØØ—I‘IØ—H‘HÜŸ7™7×+Ñ+×=Ñ=¸b¿h¹hÓGÐIØð 5ó ô	ð �|‰|Ô0×>Ñ>Ò>Ø×*Ñ*×4Ñ4Ü(×9Ñ9¸$ÀÓEõð ?ð 2rb   c                óv   — |j                   t        j                  k7  r| j                  t	        | |«      «       y y re   )rU  rG   rY  r¯   r=  )r  r\  s     rc   Úgenerate_workspace_deallocationz4PythonWrapperCodegen.generate_workspace_deallocationÄ  s.   € Ø�<‰<Ô,×;Ñ;Ò;Ø�N‰NÔ.¨t°RÓ8Õ9ð <rb   c                ó$   — |› d| j                   › �S )Nz.zero_())r‡  )r  r³   s     rc   rX  z%PythonWrapperCodegen.make_zero_bufferÈ  s   € Ø��x §¡˜}Ð-Ð-rb   c                óH   — |› ddj                  |«      › d| j                  › �S )Nrº  r»  rï   )rÒ   r‡  )r  r³   rÍ  s      rc   rö  z%PythonWrapperCodegen.wrap_kernel_callË  s'   € Ø��q˜Ÿ™ 9Ó-Ð.¨a°·±¨}Ð=Ð=rb   c                óú   — | j                   j                  d«       | j                   j                  dt        j                  j                  › d�«       |j                  | j                   j                  «       «       y )Nz*from torch.profiler import record_functionzwith record_function('graph_z_inductor_wrapper_call'):)rÕ  r¯   r@   rf   Úgraph_idr  rÌ   )r  r   s     rc   r  z8PythonWrapperCodegen.generate_profiler_mark_wrapper_callÎ  sb   € Ø×Ñ×#Ñ#Ð$PÔQØ×Ñ×#Ñ#Ø*¬1¯7©7×+;Ñ+;Ð*<Ð<UÐVô	
ð 	×Ñ˜D×-Ñ-×4Ñ4Ó6Õ7rb   c                ó:   — | j                   j                  d«       y )Nzstart_graph())rÕ  r¯   r  s    rc   r  z)PythonWrapperCodegen.generate_start_graphÕ  s   € Ø×Ñ×#Ñ# OÕ4rb   c                ó^   — | j                   j                  dt        j                  ›d�«       y )Nz
end_graph(rï   )rÕ  r¯   r&   Úprofile_bandwidth_outputr  s    rc   r  z'PythonWrapperCodegen.generate_end_graphØ  s'   € Ø×Ñ×#Ñ# j´×1PÑ1PÐ0SÐSTÐ$UÕVrb   c                ó†   — | j                   j                  t        j                  j                  j                  «       «       y)z<Synchronize GPU to ensure proton captures all kernel events.N)rÕ  r¯   r@   rf   rŠ  rg  r  s    rc   r  z-PythonWrapperCodegen.generate_proton_finalizeÛ  s)   € à×Ñ×#Ñ#¤A§G¡G×$6Ñ$6×$BÑ$BÓ$DÕErb   c                ó^   — | j                   j                  dt        j                  › d�«       y )NúU
            for kernel in globals().values():
                if isinstance(kernel, z‰.CachingAutotuner):
                    kernel.cuda_kernel_saved = False
                    kernel.cpu_kernel_saved = False
            ©rÕ  rû   r+   r]   r  s    rc   r  z6PythonWrapperCodegen.generate_reset_kernel_saved_flagsß  s2   € Ø×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð Cðõ	
rb   c                ó^   — | j                   j                  dt        j                  › d�«       y)a¨  
        Precompile and save the per-config kernel artifacts (CUBINs on GPU,
        ``.so`` + launcher ``.so`` on CPU) for Triton kernels that haven't
        been precompiled and saved as a side effect of running the generated
        JIT model (Python wrapper). This can happen when the model contains
        control flow: only one pass through the control flow operators covers
        the kernels that are saved, the remaining kernels are not launched,
        hence not saved. The main purpose of this codegen is to compile and
        save the Triton kernels outside the active control flow path for
        subsequent AOTInductor code generation and compilation.
        rj  aé  .CachingAutotuner):
                    if kernel.device_props.type == "cpu":
                        if not kernel.cpu_kernel_saved:
                            if len(kernel.launchers) == 0:
                                kernel.precompile()
                            kernel.save_cpu_kernel(launcher=kernel.launchers[0])
                    else:
                        if not kernel.cuda_kernel_saved:
                            if len(kernel.launchers) == 0:
                                kernel.precompile()
                            kernel.save_gpu_kernel(
                                stream="stream",  # use dummy stream
                                launcher=kernel.launchers[0],
                            )
            Nrk  r  s    rc   r  z5PythonWrapperCodegen.generate_save_uncompiled_kernelsé  s4   € ð 	×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð Cðõ	
rb   c                ó>   — d„ }|D �cg c]
  } ||«      ‘Œ c}S c c}w )Nc                ó   — t        | t        «      rt        | «      r| dz   S | S t        | t        t        t
        t        f«      rt        | «      S t        t        j                  j                  j                  | «      «      S )NrÁ  )rw   rW   rK   rY   Úfloatr  r  rx  r@   rf   rn   ro   )r§  s    rc   Úwrap_argzAPythonWrapperCodegen.prepare_triton_kernel_call.<locals>.wrap_arg
  s^   € Ü˜#œsÔ#ä*BÀ3Ô*G�s˜Y‘ÐPÈSÐPÜ˜C¤#¤u¬d´OÐ!DÔEÜ˜3“x�äœQŸW™W×-Ñ-×6Ñ6°sÓ;Ó<Ð<rb   ra   )r  rÍ  rp  r§  s       rc   Úprepare_triton_kernel_callz/PythonWrapperCodegen.prepare_triton_kernel_call	  s!   € ò	=ð *3Ö3 #‘˜•Ò3Ð3ùÒ3s   ˆc                óp  ‡ — t        |t        «      �rKt        |t        j                  «      r.|j	                  «       j                  «       }‰ j                  |   }n\‰ j                  j                  |«      r|}‰ j                  |   }n/|€J d«       ‚d‰ j                  › �}|}‰ xj                  dz  c_        |€
J d|› �«       ‚t        j                  j                  j                  |j                  «       «      }t        j                  j                  j                  t        j                  j                  |«      «      }t        j                  j                  j                  |j                  «       «      }|j!                  «       }	|j#                  «       }
t        j                  j                  j%                  |j'                  «       j(                  «      }d|› d|› d|	› d|
› d|› d|› d	�}‰ j*                  j-                  |› d
|› �«       t        |t        j                  «      r5‰ j/                  |d¬«      }|}‰ j*                  j-                  |› d
|› �«       |S t1        |t2        j4                  «      st        |t6        «      rÀt        |t8        «      r|‰ j:                  v r|S |€y|}t        |t6        «      r|j<                  }|t        j                  j                  j>                  v r't        j                  j                  j>                  |   }t9        t        j                  j                  j%                  |«      «      S t        |t8        t@        tB        tD        f«      rt9        |«      S t        |tF        «      rddjI                  ˆ fd„|D «       «      › d�S tK        dtM        |«      › �«      ‚)NzBV.graph.get_buffer(arg) and raw_arg can't be None at the same timeÚtmp_arg_rA   z Failed to find a buffer for arg zgenerate_example_value(r»  z, 'z', rï   rî   T)rÜ  rÝ  r>  r}  c              3  óT   •K  — | ]  }‰j                  |t        |«      «      –— Œ! y ­wre   r    )rœ   Úar  s     €rc   rž   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>P  s#   øè ø€ Ò ZÐQR ×!@Ñ!@ÀÄDÈÃG×!LÑ Zùr¤   r~  zUnsupported type )'rw   Útorch_dtyper'   r*  Ú
get_tensorrh   rÐ  r÷   rÙ  r@   rf   rn   rØ  rM  Úget_allocation_sizerN  rª  rm   r"  r‡   rŠ   r±   r¯   ræ  r  r–   rM  r  rW   rö  r  rp  rY   ro  r  r(  rÒ   rI  r¢   )r  r§  Úarg_typeÚraw_argrú  rR  rˆ   Úallocation_sizer‰   r[   r   rŠ   r  s   `            rc   r¡   z/PythonWrapperCodegen.generate_example_arg_value  s  ø€ Ü�h¤Õ,Ü˜'¤2×#3Ñ#3Ô4à"×-Ñ-Ó/×8Ñ8Ó:�Ø×*Ñ*¨3Ñ/‘Ø×%Ñ%×)Ñ)¨#Ô.Ø�Ø×*Ñ*¨3Ñ/‘àÐ*ð ØXóÐ*ð & d×&FÑ&FÐ%GÐH�Ø�Ø×0Ò0°AÑ5Õ0à�?ÐLÐ&FÀsÀeÐ$LÓL�?Ü—7‘7×#Ñ#×6Ñ6°s·|±|³~ÓFˆDÜŸg™g×.Ñ.×AÑAÜ—‘×+Ñ+¨CÓ0óˆOô —W‘W×%Ñ%×8Ñ8¸¿¹Ó9IÓJˆFà—^‘^Ó%ˆFØ—M‘M“OˆEÜ—W‘W×%Ñ%×7Ñ7¸¿¹Ó8H×8OÑ8OÓPˆFØ-¨d¨V°2°f°X¸SÀÀÈÈEÈ7ÐRTÐU[ÐT\Ð\^Ð_nÐ^oÐopÐqˆEØ×&Ñ&×0Ñ0°H°:¸SÀÀÐ1HÔIä˜'¤2×#3Ñ#3Ô4ð ×:Ñ:Ø Ø%)ð ;ó �ð �Ø×*Ñ*×4Ñ4¸°zÀÀUÀGÐ5LÔMàˆOÜ˜¤%§+¡+Ô.´*¸SÄ/Ô2Rä˜#œsÔ#Ø˜$Ÿ/™/Ñ)Ø�JØ�?Ø!Ø�Ü˜#œÔ/Ø—n‘n�Ø”a—g‘g×&Ñ&×CÑCÑCÜ—g‘g×&Ñ&×CÑCÀCÑH�ä”q—w‘w×'Ñ'×9Ñ9¸#Ó>Ó?Ð?ä˜œc¤3¬¬tÐ4Ô5Ü�s“8ˆOÜ˜œTÔ"Ø�t—y‘yÓ ZÐVYÔ ZÓZÐ[Ð[\Ð]Ð]ä%Ð(9¼$¸s»)¸Ð&EÓFÐFrb   c                óz   ‡ — t        |t        «      r ddj                  ˆ fd„|D «       «      z   dz   S t        |«      S )Nr}  r»  c              3  ó@   •K  — | ]  }‰j                  |«      –— Œ y ­wre   )Ú_grid_dim_str)rœ   r˜   r  s     €rc   rž   z5PythonWrapperCodegen._grid_dim_str.<locals>.<genexpr>W  s   øè ø€ ÒR¸T × 2Ñ 2°4× 8ÑRùs   ƒr~  )rw   r(  rÒ   rx  )r  Úgrid_per_dims   ` rc   r~  z"PythonWrapperCodegen._grid_dim_strT  s<   ø€ Ü�l¤DÔ)à�d—i‘iÓRÀ\ÔRÓRÑRÐUXÑXðô ˜Ó&Ð&rb   )r[   r¨   rÑ  rÎ  rÏ  rÒ  rÓ  r×   c               óÆ  — | j                   j                  |D �ci c]2  }t        |t        «      r |t        j
                  j                  |«      “Œ4 c}«       |xs t        j
                  j                  «       }t        j
                  j                  j                  }| j                  t        | ||||||||	|t        j
                  j                  |
|¬«      «       yc c}w )zÎ
        Generates kernel call code.

        triton: Defines whether the backend uses Triton for codegen. Otherwise it uses the CUDA language when gpu=True,
                and C++ when gpu=False.
        )r°  rÍ  rÎ  rÏ  rÑ  r¨   rÒ  rÓ  r[   rÔ  r×   rÕ  N)rÐ  r5  rw   rW   r@   rf   Útry_get_bufferr-  rj   rÕ  r¯   rÌ  r³   )r  r°  rÍ  r[   r¨   rÑ  rÎ  rÏ  rÒ  rÓ  r×   r§  rÕ  s                rc   Úgenerate_kernel_callz)PythonWrapperCodegen.generate_kernel_call\  sÅ   € ð. 	×Ñ×#Ñ#ð %öàÜ˜c¤3Ô'ð ”Q—W‘W×+Ñ+¨CÓ0Ñ0òô	
ð Ò@œ1Ÿ7™7×>Ñ>Ó@ˆÜŸW™W×.Ñ.×AÑAÐØ�‰ÜØØ'Ø#à!à!à#Øà'Ø+ØÜŸ7™7Ÿ<™<à%9Ø#5ô%õ	
ùòs   š7CrÆ  )
r[   r¨   rÑ  rÎ  rÏ  rÒ  rÓ  rÔ  r×   rÕ  c       
   
     óÊ  ‡ ‡‡— |xs t         j                  j                  «       }|s‰|j                  dvr{|j                  dk(  r"‰ j	                  ‰ j                  ‰|«      «       y |j                  dk(  r"‰ j	                  ‰ j                  ‰|«      «       y t        d|j                  › d�«      ‚‰ j                  |«      }dj                  |«      }|�=|t        k7  r4‰ j                  «        d}‰ j	                  |› d|j                  › d	�«       n!t        j                  ‰ |j                  |
«      }|s$d
|› d	�}‰ j	                  ‰› d‰› d|› d|› d	�«       y ‰ j                  «        t        j                   j"                  �rç‰‰ j$                  v�rØ|�t'        |«      t'        |«      k(  sJ d«       ‚d Š|rDt         j                  j(                  r*t         j                  j(                  j+                  |d «      Št        j,                  j.                  }dˆˆ fd„}ˆˆ fd„}g }g }|€(|�J d«       ‚d gt'        |«      z  }d gt'        |«      z  }nt'        |«      t'        |«      k(  sJ d«       ‚i }t1        t3        ||||«      «      D �]™  \  }\  }}}}d }t5        |t6        «      r!dt7        |«      v r|j9                  d«      \  }}d }‰r|‰v r‰ j;                  ‰|   «      }|rB|}t5        |t<        «      sñt?        |t@        jB                  «      st5        |tD        «      rÇ|||<   nÁ|dk(  r |||||«      r||   }n«t5        |t<        «      rˆtG        jH                  d|«      r|}n[|r%‰ jK                  |||«      }|jM                  |«       n4|‰ jN                  vr‰ jK                  |||«      }n‰ jN                  |   d   }|s%|‰f‰ jN                  |<   n‰ jK                  |||«      }t5        |t6        «      rtQ        |«      r|dz  }|jM                  |€|n|› d|› �«       �Œœ ‰ jR                  j	                  dt         j                  jT                  jW                  |j                  «      › d�«       ‰ jR                  jY                  «        ‰ jR                  j	                  ‰› ddj                  |«      › d|› d	�«       ‰ jR                  j[                  «        |r0|r.‰ jR                  j	                  ddj                  |«      › �«       n(|s&‰ jR                  j	                  t]        d|d«      «       ‰ j$                  j_                  ‰«       t         j                  j`                  ry t         j                  jb                  jd                  }|jg                  |‰|d «       |5  ‰ j	                  ‰› d|› d|› d	�«       d d d «       ‰ j                  «        y # 1 sw Y   ŒxY w)N)ÚcudaÚxpur@  Úmpszdevice z nyir»  Ú
raw_streamrƒ  rï   z	c_void_p(rñ   rº  z$call_args and arg_types do not matchc                 ó¢   •— ‰j                   j                  «       D � �cg c]  \  } }|‰k(  r| ‘Œ }} }|rddj                  |«      › d�S yc c}} w )z;Returns del for tensors whose last consumer is this kernel.údel r»  r…  rÆ  )rØ  r—  rÒ   )rÙ  ÚknÚtensors_to_deleter°  r  s      €€rc   Úget_autotune_deletion_callzUPythonWrapperCodegen._generate_kernel_call_helper.<locals>.get_autotune_deletion_callÖ  se   ø€ ð '+×&GÑ&G×&NÑ&NÓ&P÷%á"˜ Ø˜[Ò(ò ð%Ð!ñ %ñ
 %Ø! $§)¡)Ð,=Ó">Ð!?¸rÐBÐBØùó%s   ŸAc                ój  •— ||   }||v ryt        t        | |«      «      D ]�  \  }\  }}||k(  st        |t        «      sŒd}‰r|‰v r‰j	                  ‰|   «      }|dk(  rŒA	 |j                  «       }	t        |	j                  «      D ]  \  }
}||k(  sŒ|› d|
› d�||<     y Œƒ y# t        $ r Y Œ�w xY w)zÍWe try to infer raw_arg (i.e. raw_args[idx]) from remaining raw_args.
                This is particularly useful for jagged cases, where the dimension is often
                being passed in as an input.TrÆ  z.shape[r~  F)rð  rÑ   rw   r)   r/  r‡   rˆ   rI  )rÎ  rÏ  r˜  Úreused_argsÚ
target_argr   Úraw_keyrz  Útriton_inputr{   rÒ  rE  Úautotune_argsr  s               €€rc   Úinfer_arg_by_inputszNPythonWrapperCodegen._generate_kernel_call_helper.<locals>.infer_arg_by_inputsá  só   ø€ ð
 & c™]�
Ø Ñ,Øä-6´s¸8ÀXÓ7NÓ-Oò !Ñ)�AÑ)˜ Ø˜C’x¤z°'¼6Ô'BØ à#%�LÙ$¨°MÑ)AØ'+×'EÑ'EØ)¨'Ñ2ó(˜ð $ rÒ)Ø ð	!Ø!(×!3Ñ!3Ó!5˜Ü&/°·±Ó&<ò ,™F˜C Ø  J›Ø=I¸NÈ'ÐRUÐQVÐVWÐ:X ¨JÑ 7Ú'+ñ,ð!ð, øô	 /ò !ñ !ð!ús   Á#0B&ÂB&Â"B&Â&	B2Â1B2zkeys are not None but args arez#call_args and raw_args do not matchr¹  rÆ  z^(workspace|semaphore)r   rÁ  r†  r‡  z.run(z	, stream=r‰  z
<del_call>rô  )4r@   rf   r-  r¢   r¯   rö  r:  rq  rÒ   r.   r:  r6  rR  ró  r6  r&   r¨   r©   r°   rº   Úautotuning_mappingr÷   rú  Úautotune_per_kernel_allocrð  rÑ   rw   rW   Úsplitr/  rv  r  r–   rM  r  rî  Úmatchr¡   r�   rØ  rK   r±   rŠ  r�  rY  r|  r4   rÓ   rˆ  rÇ  rý  rÈ  ) r  r°  rÍ  r[   r¨   rÑ  rÎ  rÏ  rÒ  rÓ  rÔ  r×   rÕ  Úcall_args_strÚstream_nameÚ
stream_ptrÚ_per_kernelrŒ  r“  Úall_argsÚtensor_arg_strsrŽ  r   r§  ry  r�  rz  rÀ   r‘  Úarg_strrÌ  r’  s    ``                             @rc   r×  z1PythonWrapperCodegen._generate_kernel_call_helper”  s¹  ú€ ð  Ò@œ1Ÿ7™7×>Ñ>Ó@ˆÙ˜&Ÿ+™+¨_Ñ<Ø�{‰{˜eÒ#Ø—‘˜t×4Ñ4°[À)ÓLÔMð ð —‘ Ò%à—‘˜t×4Ñ4°[À)ÓLÔMð ô # W¨V¯[©[¨M¸Ð#>Ó?Ð?ð ×7Ñ7¸	ÓBˆØŸ	™	 -Ó0ˆØÐ)Ð.@ÔDVÒ.Vð ×,Ñ,Ô.Ø&ˆKØ�N‰N˜k˜]Ð*<¸V¿\¹\¸NÈ!ÐLÕMä.×CÑCØ�f—l‘l JóˆKñ Ø$ [ M°Ð3ˆJØ�N‰NØ�-˜q  ¨Q¨}¨o¸RÀ
¸|È1ÐMôð à×%Ñ%Ô'ô �M‰M×2Ó2Ø 4×#=Ñ#=Ò=ð Ð(¬S°«^¼sÀ9»~Ò-Mð Ø6óÐMð !ˆMÙ#¬¯©×(BÒ(BÜ !§¡× :Ñ :× >Ñ >Ø(¨$ó!�ô !×-Ñ-×GÑGˆKö	õðB ˆHØ ˆOØÐàÐ'ÐIÐ)IÓIÐ'Ø ˜6¤C¨	£NÑ2�Ø ˜6¤C¨	£NÑ2‘ä˜8“}¬¨I«Ò6ð Ø9óÐ6ð ˆKÜ8Aä�I˜y¨(°HÓ=ó9ó 3PÑ4�Ñ4�C˜ 7¨Gð �Ü˜c¤3Ô'¨C´3°s³8©Oà"Ÿy™y¨›~‘H�C˜à+/�Ù  W°Ñ%=Ø#'×#AÑ#AØ% gÑ.ó$�Lñ  Ø*�GÜ% h´Ô<Ü" 8¬U¯[©[Ô9Ü% c¬?Ô;à/6˜ GÒ,Ø ’]Ñ':Ø˜h¨¨;ô(ð
 *¨'Ñ2‘GÜ ¬+Ô6ô —x‘xÐ 9¸3Ô?Ø"%™Ù$Ø"&×"AÑ"AØ ¨7ó#˜ð (×.Ñ.¨wÕ7Ø D×$EÑ$EÑEØ"&×"AÑ"AØ ¨7ó#™ð #'×"CÑ"CÀCÑ"HÈÑ"K˜Ù&ØBIÈ;ÐAW˜×9Ñ9¸#Ò>à"×=Ñ=¸cÀ8ÈWÓU�Gä˜c¤3Ô'Ô,DÀSÔ,IØ˜yÑ(�GØ—‘¨3¨;¡¸s¸eÀ1ÀWÀIÐ<NÖOðg3Pðl ×&Ñ&×0Ñ0ØœŸ™×*Ñ*×7Ñ7¸¿¹ÓEÐFÀaÐHôð ×&Ñ&×0Ñ0Ô2Ø×&Ñ&×0Ñ0Ø�-˜u T§Y¡Y¨xÓ%8Ð$9¸À;À-ÈqÐQôð ×&Ñ&×2Ñ2Ô4á™Ø×*Ñ*×4Ñ4Ø˜4Ÿ9™9 _Ó5Ð6Ð7õñ !Ø×*Ñ*×4Ñ4Ü$Ø$Ø2Ø$óôð ×&Ñ&×*Ñ*¨;Ô7Ü�w‰w×"Ò"àô !"§¡× 4Ñ 4× BÑ BÐØ×.Ñ.¨y¸+ÀyÐRVÔWØ"ñ 	XØ�N‰N˜k˜]¨%°¨¸iÈÀ}ÐTUÐVÔW÷	Xà×%Ñ%Õ'÷	Xð 	Xús   Ö%W×W"c                ó:   — | j                   j                  |«       y re   )rÜ  r�   rE  s     rc   r¯   zPythonWrapperCodegen.writelinef  s   € Ø�
‰
×Ñ˜$Õrb   c                ó4   — |D ]  }| j                  |«       Œ y re   rt  )r  rÜ  r²   s      rc   rØ  zPythonWrapperCodegen.writelinesi  s   € Øò 	!ˆDØ�N‰N˜4Õ ñ	!rb   c                óL   — | j                   j                  t        |«      «       y re   )rÜ  r�   r:   )r  Úctxs     rc   r  z"PythonWrapperCodegen.enter_contextm  s   € Ø�
‰
×Ñœ+ cÓ*Õ+rb   c                óú  ‡ ‡	— ddl m}  |«       rdd l}t        |t        «      rt        |j                  j                  «      S t        |t        j                  «      rt        |«      S t        |t        t        f«      rAt        j                   G d„ d«      «       Š	t         t        |«      ˆ	ˆ fd„|D «       «      «      S t        |t         j"                  j$                  «      rt'        |«      S t        |t(        j*                  t(        j,                  t.        f«      r|j1                  «       S  |«       r+t        |j2                  j4                  «      rt        |«      S t        |t(        j6                  t(        j8                  f«      r|j1                  «       S t;        t        |«      «      rEt=        |«      \  }}|j?                  «       D ]"  \  }}|t@        jB                  jD                  |<   Œ$ |S t        |«      S )Nr   )Úhas_triton_packagec                  ó   — e Zd ZU ded<   d„ Zy)ú1PythonWrapperCodegen.val_to_arg_str.<locals>.Shimr   Úrefc                ó   — | j                   S re   )r§  r  s    rc   Ú__repr__z:PythonWrapperCodegen.val_to_arg_str.<locals>.Shim.__repr__€  s   € ØŸ8™8�Orb   N)r]   r^   r_   r`   r©  ra   rb   rc   ÚShimr¦  |  s   … à“ó$rb   rª  c              3  óV   •K  — | ]   } ‰t         j                  ‰|«      «      –— Œ" y ­wre   rÏ  )rœ   ru  rª  r  s     €€rc   rž   z6PythonWrapperCodegen.val_to_arg_str.<locals>.<genexpr>…  s$   øè ø€ ÒVÈq™Ô1×@Ñ@ÀÀqÓI×JÑVùs   ƒ&))#Útorch.utils._tritonr¤  r¨   rw   r   rx  rp   rh  r–   r
   r¦   r(  rî  Ú	dataclassr?  r¢   r  Ú_opsÚ
OpOverloadr   r'   r†   Ú
MutableBoxr*   r­  Úlanguager   rO  rP  r   r   rì  r@   rf   Úopaque_value_type_classes)
r  rE  Útype_r¤  r¨   Úobj_reprÚopaque_typesÚnr‘  rª  s
   `        @rc   rÐ  z#PythonWrapperCodegen.val_to_arg_strp  s€  ù€ Ý:áÔÛä�aœÔ"Ü˜Ÿ™Ÿ™Ó%Ð%Ü˜œ5Ÿ:™:Ô&Ü˜“8ˆOÜ˜œE¤4˜=Ô)ä×"Ñ"÷$ð $ó #ð$ô Ø”�Q“ÔVÐTUÔVÓVóð ô ˜œ5Ÿ:™:×0Ñ0Ô1Ü& qÓ)Ð)Ü˜œBŸI™I¤r§}¡}´oÐFÔGØ×&Ñ&Ó(Ð(ÙÔ!¤j°°F·O±O×4IÑ4IÔ&JÜ˜“7ˆNÜ˜œB×-Ñ-¬r×/CÑ/CÐDÔEØ×&Ñ&Ó(Ð(Ü!¤$ q£'Ô*Ü%8¸Ó%;Ñ"ˆH�lØ$×*Ñ*Ó,ò 9‘��1Ø78”—‘×1Ñ1°!Ò4ð9àˆOä˜“7ˆNrb   c           	     ó`  — |j                  «       }|j                  «       }t        |j                  «       «      }t        t        j
                  j                  |«      «      }t        |j                  «       «      }|j                  «       }| j                  |j                  «       ||||||«      S re   )rª  rm   r¦   rM  r@   rf   rx  rN  Úget_is_pinnedr[  rh   )r  r  r[   r   rT  Úallocation_shaper‰   Ú	is_pinneds           rc   rK  z+PythonWrapperCodegen.make_buffer_allocation˜  s”   € Ø×"Ñ"Ó$ˆØ× Ñ Ó"ˆÜ�f—o‘oÓ'Ó(ˆÜ ¤§¡×!<Ñ!<¸VÓ!DÓEÐÜ�v×(Ñ(Ó*Ó+ˆØ×(Ñ(Ó*ˆ	Ø×#Ñ#Ø�O‰OÓ˜v u¨e°VÐ=MÈyó
ð 	
rb   c                óx   — d}t         j                  j                  s| j                  j	                  |d¬«       y y )Nzi
            from torch._inductor.runtime.debug_utils import check_memory_step, track_tensor
            Trã   )r@   rf   rˆ  rÉ  rû   r4  s     rc   Ú"write_memory_track_allocation_oncez7PythonWrapperCodegen.write_memory_track_allocation_once£  s4   € ðˆ
ô �w‰w×"Ò"Ø�L‰L×Ñ 
°$ÐÕ7ð #rb   c                óú  — |€|}| j                  |«      }| j                  |«      }	| j                  |«      }
t        j                  j                  j                  j
                  r|› d|	› d|
› d|› d|j                  › d|› d�}ne|j                  dk(  r|r|› d|	› d|
› d|› d	�}nE|j                  d
v r|› d|j                  › d|	› d|
› d|› d	�
}n|› d|	› d|
› d|j                  › d|› d	�
}||	k7  r|d|› d|
› d	�z   }|S )Nz = tracked_empty_strided(r»  z, dtype=rÈ  z	', name='re  r@  z = empty_strided_cpu_pinned(rï   )r@  r„  r…  Úmtiaz = empty_strided_rº  z = empty_strided(rÉ  z.as_strided()r§   r  r  r&   r  r  r¢   )r  r³   r[   r   rT  r‰   r¹  rº  rQ  Úcodegen_allocation_shape_tupleÚcodegen_stride_tuplerÉ  s               rc   r[  z$PythonWrapperCodegen.make_allocation«  sˆ  € ð Ð#Ø$Ðà"×=Ñ=¸eÓDÐØ)-×)HÑ)HØó*
Ð&ð  $×>Ñ>¸vÓFÐÜ�?‰?×!Ñ!×.Ñ.×EÒEà�&Ð1Ø1Ð2°"Ø'Ð(ð )Ø˜ð  Ø!Ÿ;™;˜-ð (Ø˜˜bð"ñ ð �[‰[˜EÒ!¡ià�&Ð4Ø1Ð2°"Ø'Ð(¨Ø�'˜ðñ ð �[‰[Ð:Ñ:ð �&Ð)¨&¯+©+¨°aØ1Ð2°"Ø'Ð(¨Ø�'˜ðñ ð �&Ð)Ø1Ð2°"Ø'Ð(ð )Ø!Ÿ;™;˜- y°°°qð:ð ð Ð"@Ò@à˜,Ð':Ð&;¸2Ð>RÐ=SÐSTÐUÑUˆCØˆ
rb   c                ó8   — | j                  t        |«      «       y re   )r¯   rl  rE  s     rc   Úmake_commentz!PythonWrapperCodegen.make_commentÛ  s   € Ø�‰”{ 4Ó(Õ)rb   c           	     ó`   — | j                   › |› d|› | j                  › d| j                  › d|› �	S )Nrî   ú  ú )r†  r‡  rÞ  )r  Únew_nameÚold_namerÞ  s       rc   Úmake_tensor_aliasz&PythonWrapperCodegen.make_tensor_aliasÞ  s6   € Ø—,‘,� ˜z¨¨X¨J°t·{±{°mÀ2ÀdÇlÁlÀ^ÐSTÐU\ÐT]Ð^Ð^rb   c                ó(   — d|j                  «       › �S )Nr‰  )rh   )r  r  s     rc   rÇ  z%PythonWrapperCodegen.make_buffer_freeá  s   € Ø�f—o‘oÓ'Ð(Ð)Ð)rb   c                ó8   — ddj                  d„ |D «       «      › �S )Nr‰  r»  c              3  ó    K  — | ]  }|–— Œ y ­wre   ra   )rœ   r³   s     rc   rž   z:PythonWrapperCodegen.make_free_by_names.<locals>.<genexpr>å  s   è ø€ Ò>¨¤Ñ>ùs   ‚)rÒ   )r  Únames_to_dels     rc   Úmake_free_by_namesz'PythonWrapperCodegen.make_free_by_namesä  s   € Ø�d—i‘iÑ>°Ô>Ó>Ð?Ð@Ð@rb   c           	     ó`   — | j                   › |› d|› |› | j                  › d| j                  › d�	S )Nrî   rÄ  ú reuse)rÝ  r‡  rÞ  )r  rÇ  rÆ  Údel_lines       rc   Úcodegen_exact_buffer_reusez/PythonWrapperCodegen.codegen_exact_buffer_reuseç  s@   € Ø×.Ñ.Ð/°¨z¸¸X¸JÀxÀjÐQU×Q\ÑQ\ÐP]Ð]_Ð`d×`lÑ`lÐ_mÐmsÐtÐtrb   c                óP   — |�$| j                  | j                  › d|› d|› �«       y y )Nz [Provenance debug handles] r‡  )r¯   rÞ  )r  r°  Údebug_handles      rc   Úwrite_provenance_debug_handlez2PythonWrapperCodegen.write_provenance_debug_handleê  s4   € ð
 Ð#Ø�N‰NØ—<‘<�.Ð <¸[¸MÈÈ<È.ÐYõð $rb   c                ór  — |j                  «       |j                  «       k(  sJ ‚|j                  «       }|j                  «       }d}|t        j                  j	                  «       vr|rd| j                  |«      › �}|j                  «       |j                  «       k(  r4|j                  «       |j                  «       k(  r| j                  |||«      S | j                  ||j                  «       |j                  «       d| j                  j                  «      }| j                  › |› d|› |› d| j                  › d�S )Nú;z; r   rî   rÄ  rÏ  )rm   rh   r@   rf   r  rÇ  rM  rN  rÑ  r•  rÕ  r¯   r†  rÞ  )r  r	  Únewrl  rÇ  rÆ  rÐ  Úreinterpret_views           rc   ro  z&PythonWrapperCodegen.make_buffer_reuseô  s  € Ø�}‰}‹ #§-¡-£/Ò1Ð1Ð1Ø—<‘<“>ˆØ—<‘<“>ˆØˆØœ1Ÿ7™7×3Ñ3Ó5Ñ5¹*Ø˜D×1Ñ1°#Ó6Ð7Ð8ˆHà�<‰<‹>˜SŸ\™\›^Ò+°·±Ó0@ÀCÇNÁNÓDTÒ0TØ×2Ñ2°8¸XÀxÓPÐPà×8Ñ8Ø�—‘“ §¡Ó!1°1°d×6GÑ6G×6QÑ6Qó
Ðð —,‘,� ˜z¨Ð-=Ð,>¸x¸jÈÈ4Ï<É<È.ÐX^Ð_Ð_rb   c                ó¬   — | j                  t        || j                  › |› d|j                  «       › | j                  › d| j
                  › d�«      «       y )Nrî   rÄ  z alias)r¯   rD   r†  r­  r‡  rÞ  )r  r³   rg  s      rc   rh  z0PythonWrapperCodegen.codegen_deferred_allocation  sS   € Ø�‰ÜØØ—<‘<�.   c¨$×*@Ñ*@Ó*BÐ)CÀDÇKÁKÀ=ÐPRÐSW×S_ÑS_ÐR`Ð`fÐgóõ	
rb   c                ót  ‡— |j                  «       }|t        j                  j                  v sG|| j                  v s9t        |t        j                  t        j                  t        j                  f«      ry | j                  j                  |«       t        |j                  «       t        j                  t        j                  f«      r|j                  «       sy |j                  «       }t        |t        j                   «      ry t        |t        j"                  «      ry t        |t        j$                  «      �r+t        |j&                  t        j(                  «      s*J dt+        |j&                  «      › d|j&                  › �«       ‚|j&                  j,                  }t        |t        j.                  «      sJ t+        |«      «       ‚|j,                  }t        |t        j0                  t        j(                  f«      sJ t+        |«      «       ‚t        |t        j(                  «      rdˆfd„Š ‰|«      }| j3                  |«       | j5                  t7        | |||«      «       y t        |t        j8                  «      r| j5                  t;        | |d¬«      «       y | j5                  t;        | |«      «       y )Nzunexpected rð   c                ó  •— t        | t        j                  «      r ‰| j                  «       «      S t        | t        j                  «      r ‰| j
                  «      S t        | t        j                  «      sJ t        | «      «       ‚| S re   )rw   r'   ÚBaseViewÚunwrap_viewr°  rŽ   r†   r¢   )r  Úunwrap_viewss    €rc   rÞ  z=PythonWrapperCodegen.codegen_allocation.<locals>.unwrap_views0  sd   ø€ Ü! &¬"¯+©+Ô6Ù+¨F×,>Ñ,>Ó,@ÓAÐAÜ! &¬"¯-©-Ô8Ù+¨F¯K©KÓ8Ð8Ü% f¬b¯i©iÔ8ÐF¼$¸v»,ÓFÐ8Ø!�Mrb   Tr?  )rÈ   ú	ir.Buffer)rh   r@   rf   rÆ  rï  rw   r'   ÚDonatedBufferÚSubgraphBufferÚInputBufferrÓ   Úget_defining_opÚExternKernelAllocÚMultiOutputÚshould_allocaterv   ÚMutationLayoutSHOULDREMOVEr¼  rf  rg  r*   r¢   rŽ   rŒ   r†   r°  r¯   rb  rx   r0  )r  r  r³   r{   ÚboxÚinput_bufferrÞ  s         @rc   r°  z'PythonWrapperCodegen.codegen_allocation  s  ø€ Ø�‰Ó ˆð ”A—G‘G×+Ñ+Ñ+Ø�t—~‘~Ñ%Ü˜&¤2×#3Ñ#3´R×5FÑ5FÌÏÉÐ"WÔXàØ�‰×Ñ˜4Ô äØ×&Ñ&Ó(Ü×%Ñ%¤r§~¡~Ð6ôð ×*Ñ*Ô,àà×'Ñ'Ó)ˆÜ�fœb×;Ñ;Ô<ØÜ�fœbŸm™mÔ,ØÜ�fœb×0Ñ0Õ1Ü˜fŸk™k¬2×+=Ñ+=Ô>ð Øœd 6§;¡;Ó/Ð0°°6·;±;°-Ð@óÐ>ð —+‘+×"Ñ"ˆCÜ˜c¤2§=¡=Ô1Ð<´4¸³9Ó<Ð1ØŸ8™8ˆLÜ˜l¬R¯Y©Y¼×8JÑ8JÐ,KÔLð ÌdØóOó ÐLô ˜,¬×(:Ñ(:Ô;õ"ñ  ,¨LÓ9�Ø×#Ñ# LÔ1Ø�N‰Nœ?¨4°¸vÀvÓNÔOØä�fœb×1Ñ1Ô2Ø�N‰Nœ<¨¨fÀ$ÔGÔHØà�‰”| D¨&Ó1Õ2rb   c                óÌ  — |j                  «       }t        |t        j                  t        j                  f«      r| j                  t        | |«      «       y t        |j                  «       t        j                  «      r| j                  t        | |d¬«      «       y | j                  |«      sy | j                  j                  |«       | j                  t        | |«      «       y )NTr?  )rh   rw   r'   râ  rN  r¯   rÃ  rv   rx   r=  Ú	can_reuserð  rÓ   )r  r  r³   s      rc   Úcodegen_freez!PythonWrapperCodegen.codegen_freeC  s¦   € Ø�‰Ó ˆô �fœrŸ~™~¬r×/AÑ/AÐBÔCØ�N‰Nœ8 D¨&Ó1Ô2Øä�f×,Ñ,Ó.´×0CÑ0CÔDð �N‰NÔ.¨t°VÈÔNÔOØà�~‰~˜fÔ%ØØ�
‰
�‰�tÔà�‰Ô*¨4°Ó8Õ9rb   c                óÜ  — |j                  «       }|t        j                  j                  v xs¾ |t        j                  j                  v xr6 t        t        j                  j                  |   t        j                  «       xsh |t        j                  j                  v xsJ |t        j                  j                  v xs, |t        j                  j                  v xs || j                  v  S re   )rh   r@   rf   rÆ  r  rw   Úgraph_inputs_originalr'   rà  rÙ  rÚ  Únever_reuse_buffersrð  )r  ré  Úoutput_bufferr³   s       rc   rë  zPythonWrapperCodegen.can_reuseW  sÍ   € Ø×$Ñ$Ó&ˆà”A—G‘G×+Ñ+Ð+ò 
"àœŸ™×,Ñ,Ð,ò Ü"Ü—G‘G×1Ñ1°$Ñ7¼×9IÑ9Ióð ò
"ð ”q—w‘w×(Ñ(Ð(ò
"ð ”q—w‘w×2Ñ2Ð2ò
"ð ”q—w‘w×2Ñ2Ð2ò
"ð �t—z‘zÐ!ð
ð 	
rb   c                óš   — |j                  «       | j                  v xr. | j                  |j                  «          |j                  «       k(  S re   )rh   rñ  )r  r  Úreused_buffers      rc   Ú	did_reusezPythonWrapperCodegen.did_reuseg  sC   € ð �O‰OÓ §¡Ð,ò KØ—‘˜FŸO™OÓ-Ñ.°-×2HÑ2HÓ2JÑJð	
rb   c                ót  — t        ||«      sJ ‚| j                  |«       | j                  j                  |j	                  «       «       | j
                  j                  |j	                  «       «       |j	                  «       | j                  |j	                  «       <   | j                  t        | ||«      «       y re   )	r„   r°  rð  rÓ   rh   rï  rñ  r¯   rC  )r  ré  rð  s      rc   Úcodegen_inplace_reusez*PythonWrapperCodegen.codegen_inplace_reuseo  s‰   € Ü$ \°=ÔAÐAÐAØ×Ñ Ô-Ø�
‰
�‰�|×,Ñ,Ó.Ô/Ø�‰×Ñ˜=×1Ñ1Ó3Ô4Ø0<×0EÑ0EÓ0Gˆ�‰�M×*Ñ*Ó,Ñ-Ø�‰”y  |°]ÓCÕDrb   c                óŒ   — t        |«      }|| j                  v r|S | j                  j                  |«       | j                  |z   S re   )rW   rã  rÓ   r†  )r  r  r³   s      rc   Úcodegen_unbacked_symbol_declz1PythonWrapperCodegen.codegen_unbacked_symbol_declw  sC   € Ü�6‹{ˆØ�4×-Ñ-Ñ-ØˆKð ×&Ñ&×*Ñ*¨4Ô0Ø—<‘< $Ñ&Ð&rb   c                óš   — t        t        j                  j                  j                  |«      }| j                  t        | |||«      «       y re   )r   r@   rf   rn   rZ  r¯   r­  )r  r®  r¯  r±  s       rc   Ú(codegen_unbacked_symbol_defs_for_outputsz=PythonWrapperCodegen.codegen_unbacked_symbol_defs_for_outputs€  sB   € ô 6Ü�G‰G×Ñ×&Ñ&Ð(9ó
Ðð 	�‰Ü" 4¨°gÐ?PÓQõ	
rb   c                óÊ   ‡‡‡‡— |sy |j                  «       D ]I  \  }Šdˆfd„Šˆˆˆˆfd„}| j                  | j                  |«      › d |«       › | j                  › �«       ŒK y )Nc                ó0  •— |dk(  r| S t        |«      dk\  r_t        |d   t        «      rLt        |d   t        j                  «      r/ ‰| › d|d   j
                  › d|d   j                  › d�|dd  «      S t        |d   t        «      r ‰| › d|d   j
                  › d�|dd  «      S t        |d   t        j                  «      rYt        j                  j                  r  ‰d	|d   j                  › d
| › d�|dd  «      S  ‰| › d|d   j                  › d�|dd  «      S t        |d   t        «      rYt        j                  j                  r  ‰d| › d|d   j                  › d�|dd  «      S  ‰| › d|d   j                  › d�|dd  «      S t        d|› �«      ‚)Nra   r$   r   rA   rñ   rº  rï   r~  z	std::get<z>(r}  r~  z / z.__floordiv__(rÃ  )rº   rw   r   ÚpytreeÚSequenceKeyr³   r˜  r@   rf   rˆ  r   rÀ  r�  )rh  rÅ  Úgos     €rc   rþ  zJPythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs.<locals>.go   s¾  ø€ Ø˜b’=Ø�Kô ˜“L AÒ%Ü" 7¨1¡:¬}Ô=Ü" 7¨1¡:¬v×/AÑ/AÔBáØ˜&  '¨!¡*§/¡/Ð!2°!°G¸A±J·N±NÐ3CÀ1ÐEÀwÈqÈrÀ{óð ô   ¨¡
¬MÔ:Ù   a¨°©
¯©Ð'8¸Ð;¸WÀQÀR¸[ÓIÐIÜ ¨¡
¬F×,>Ñ,>Ô?ô Ÿ7™7×.Ò.ñ ˜Y w¨q¡z§~¡~Ð&6°b¸¸¸aÐ@À'È!È"À+ÓNðñ   4 &¨¨'°!©*¯.©.Ð)9¸Ð ;¸WÀQÀR¸[ÓIðô
   ¨¡
¬KÔ8ä—w‘w×*Ò*Ù! A d V¨3¨w°q©z×/AÑ/AÐ.BÀ!Ð"DÀgÈaÈbÀkÓRÐRÙ   n°W¸Q±Z×5GÑ5GÐ4HÈÐJÈGÐTUÐTVÈKÓXÐXä(Ð+@ÀÀ	Ð)JÓKÐKrb   c                 ó¬  •— t         j                  j                  r±t        ‰«      dk(  rZ‰d   }  ‰‰d   j	                  «       t        | t        j                  «      r!t        | j                  «      dk7  r	‰dd  «      S ‰«      S t        ‰d   t        j                  «      sJ ‚ ‰‰‰d   j                     j	                  «       ‰dd  «      S  ‰‰‰«      S )NrA   r   )r@   rf   rˆ  rº   rh   rw   r'   rå  rz  rü  rý  r˜  )rÉ  rþ  rÅ  r®  r¯  s    €€€€rc   Úgo_outerzPPythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs.<locals>.go_outer¿  sÑ   ø€ Ü—7‘7×&Ò&ô
 ˜7“| qÒ(Ø% a™j˜ñ  "Ø# A™J×/Ñ/Ó1ä)¨#¬r¯~©~Ô>Ä3ÀsÇ{Á{ÓCSÐWXÒCXð $ A B˜Kó ð ð ")ó	 ð ô  *¨'°!©*´f×6HÑ6HÔIÐIÐIÙ! '¨'°!©*¯.©.Ñ"9×"BÑ"BÓ"DÀgÈaÈbÀkÓRÐRá˜k¨7Ó3Ð3rb   rî   )rh  rW   rÅ  zpytree.KeyPath)rì  r¯   r÷  r‡  )r  r®  r¯  r±  rE  r   rþ  rÅ  s    ``   @@rc   r³  z>PythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs�  sg   û€ ñ !Øð ,×1Ñ1Ó3ò =	‰JˆAˆwõ
L÷>4ð. �N‰NØ×4Ñ4°QÓ7Ð8¸¹H»J¸<ÈÏÉÀ}ÐUõñw=	rb   c                óÚ  ‡ ‡‡‡— ˆˆ ˆfd„}ˆˆ ˆfd„}	 ‰ j                  ‰j                  «       ‰ j                  ‰ j                  › d‰j                  › �«        |«        t
        j                  }t        j                  ‰j                  «      5  ‰j                  j                  |¬«       d d d «        |«        ‰ j                  «        y # 1 sw Y   Œ!xY w# ‰ j                  «        w xY w)Nc                 ó  •— t        ‰j                  j                  «      t        ‰«      k(  sJ ‚t        ‰j                  j                  ‰«      D ]3  \  } }‰j	                  ‰j
                  › | › d|› ‰j                  › �«       Œ5 y rè  )rº   rf   r  rÑ   r¯   r†  r‡  )Úinner_inputÚouter_inputÚouter_inputsr  Úsubgraphs     €€€rc   Ú_codegen_subgraph_prefixzSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_prefixè  sy   ø€ Ü�x—~‘~×2Ñ2Ó3´s¸<Ó7HÒHÐHÐHÜ,/Ø—‘×+Ñ+¨\ó-ò Ñ(�˜[ð —‘Ø—|‘|�n [ M°°[°MÀ$Ç+Á+ÀÐOõñrb   c                 ó  •— t        ‰j                  j                  «      t        ‰«      k(  sJ ‚t        ‰j                  j                  ‰«      D ]5  \  } }‰j	                  |› d| j                  «       › ‰j                  › �«       Œ7 y rè  )rº   rf   r  rÑ   r¯   r­  r‡  )Úinner_outputÚouter_outputÚouter_outputsr  r  s     €€€rc   Ú_codegen_subgraph_suffixzSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_suffixñ  s{   ø€ Ü�x—~‘~×3Ñ3Ó4¼¸MÓ8JÒJÐJÐJÜ.1Ø—‘×,Ñ,¨mó/ò Ñ*�˜lð —‘Ø#�n C¨×(FÑ(FÓ(HÐ'IÈ$Ï+É+ÈÐWõñrb   ú subgraph: )Úparent_graph)	rX  rf   r¯   rÞ  r³   r@   rG  Úcodegen_subgraphr{  )r  r  r  r  r  r  r  s   ````   rc   Úcodegen_subgraph_by_inliningz1PythonWrapperCodegen.codegen_subgraph_by_inliningÚ  s¶   û€ ö	ö	ð	'Ø×%Ñ% h§n¡nÔ5Ø�N‰N˜dŸl™l˜^¨;°x·}±}°oÐFÔGÙ$Ô&ÜŸ7™7ˆLÜ×$Ñ$ X§^¡^Ó4ñ Ø—‘×/Ñ/Ø!-ð 0ô ÷ñ %Ô&à×$Ñ$Õ&÷ð ûð ×$Ñ$Õ&ús$   ”A;C ÂCÂ,C ÃCÃC ÃC*c           	     ó¾  — |j                   }|j                  }t        |j                  «       «      |j                  D �cg c]  }|j
                  ‘Œ c}z   }dj                  |«      t        |«      dk(  rdndz   }|D �cg c]  }|j                  «       ‘Œ }	}dj                  |	«      t        |«      dk(  rdndz   }
| j                  d|› d|› d�«       |j                  «       D ��cg c]
  \  }}|sŒ	|‘Œ }}}|r#| j                  ddj                  |«      › �«       | j                  d	|
› d
|› d|› d�«       | j                  d|› d�«       yc c}w c c}w c c}}w )z'Generate code to call a graph partitionr»  rA   rY  rÆ  Ú	partitionú	_args = [r~  r‰  rº  z) = self.partitions[z](partitionú_args)zdel partitionÚ_argsN)Úinput_deallocationÚoutput_nodesr(  r  Úsymbol_inputsr³   rÒ   rº   rh   r¯   rì  )r  Úpartition_idr  r  r  Úsymbol_inputrZ  r\   rp   Úoutput_namesr¯  r³   Ú
deallocaterÌ  s                 rc   Úcodegen_partition_callz+PythonWrapperCodegen.codegen_partition_call  sj  € ð 2×DÑDÐØ+×8Ñ8ˆäÐ-×2Ñ2Ó4Ó5Ø2F×2TÑ2Tö9
Ø".ˆL×Óò9
ñ 
ˆð —‘˜;Ó'´#°kÓ2BÀaÒ2G©3ÈRÑPˆà4@ÖA¨D˜Ÿ™�ÐAˆÐAØ—)‘)˜LÓ)´C¸Ó4EÈÒ4J©SÐPRÑSˆð 	�‰˜ < .°	¸&¸ÀÐCÔDð *<×)AÑ)AÓ)C÷
Ù%�T˜:ÂzŠDð
ˆñ 
ñ Ø�N‰N˜T $§)¡)¨LÓ"9Ð!:Ð;Ô<ð 	�‰Ø�ˆyÐ,¨\¨N¸+ÀlÀ^ÐSYÐZô	
ð 	�‰˜ | n°EÐ:Õ;ùò-9
ùò Bùó
s   ¿EÁ?EÃ'
EÃ2Ec                óP   — t        |«      D �cg c]  }d|› �‘Œ	 c}| _        y c c}w )NÚ
partition_)rþ  r¦  )r  Únum_partitionsr˜  s      rc   Úset_all_partition_namesz,PythonWrapperCodegen.set_all_partition_names(  s$   € ÜBGÈÓBWÖ#X¸3 j°°Ò$6Ò#XˆÕ ùÒ#Xs   Ž#c           	     óp  — dj                  |«      t        |«      dk(  rdndz   }dj                  |«      t        |«      dk(  rdndz   }| j                  |j                  j                  › d|› d�«       | j                  d|› d|j                  j                  › d|j                  j                  › d	�«       y )
Nr»  rA   rY  rÆ  r  r~  rº  z) = r  )rÒ   rº   r¯   rf   r³   )r  r  r  Úouter_flattened_outputsÚouter_output_namesÚouter_input_namess         rc   Ú,codegen_subgraph_call_with_flattened_outputszAPythonWrapperCodegen.codegen_subgraph_call_with_flattened_outputs+  s¹   € ð "ŸY™YÐ'>Ó?ÜÐ.Ó/°1Ò4‰C¸"ñ
Ðð !ŸI™I lÓ3Ü�|Ó$¨Ò)‰C¨rñ
Ðð 	�‰˜(Ÿ.™.×-Ñ-Ð.¨iÐ8IÐ7JÈ!ÐLÔMð 	�‰ØÐ"Ð# 4¨¯©×(;Ñ(;Ð'<¸A¸h¿n¹n×>QÑ>QÐ=RÐRXÐYõ	
rb   c                óv  — dj                  |«      t        |«      dk(  rdndz   }| j                  |j                  j                  › d|› d�«       t
        j                  j                  j                  «        | j                  |› d|j                  j                  › d|j                  j                  › d	�«       y )
Nr»  rA   rY  rÆ  r  r~  rî   rº  r  )rÒ   rº   r¯   rf   r³   r@   rj   Úfree_buffers)r  r  r  Úouter_buffer_namer%  s        rc   Úcodegen_subgraph_callz*PythonWrapperCodegen.codegen_subgraph_call=  s¢   € à ŸI™I lÓ3Ü�|Ó$¨Ò)‰C¨rñ
Ðð 	�‰˜(Ÿ.™.×-Ñ-Ð.¨iÐ8IÐ7JÈ!ÐLÔMô 	
�‰×Ñ×&Ñ&Ô(ð 	�‰Ø Ð!  X§^¡^×%8Ñ%8Ð$9¸¸8¿>¹>×;NÑ;NÐ:OÈvÐVõ	
rb   c                ó   — | j                  |j                  «       | j                  d«       | j                  | j                  › d|j                  › �«       t
        j                  }|j                  |j                  _        |j                  |j                  _        |j                  j                  | j                  vr¨t        j                  |j                  «      5  t        j                  dd«      5  |j                  j                  «       \  }}d d d «       d d d «       |j                  j                  }| j                  j                  |«       | j                  |«       y y # 1 sw Y   ŒVxY w# 1 sw Y   ŒZxY w)NrÆ  r  r^  F)rX  rf   rÂ  rÞ  r³   r@   rˆ  Ú
fx_wrapperrø  rG  r&   Úpatchr\  rÓ   r÷  )r  r  r  rö  ró  r  s         rc   Úcodegen_subgraph_commonz,PythonWrapperCodegen.codegen_subgraph_commonN  s/  € Ø×!Ñ! (§.¡.Ô1Ø×Ñ˜"ÔØ×Ñ˜TŸ\™\˜N¨+°h·m±m°_ÐEÔFä—w‘wˆØ%1×%=Ñ%=ˆ�‰Ô"Ø$0×$;Ñ$;ˆ�‰Ô!à�>‰>×Ñ d×&FÑ&FÑFô ×$Ñ$ X§^¡^Ó4ñ @ä—\‘\Ð"3°UÓ;ñ @à'/§~¡~×'=Ñ'=Ó'?Ñ$�M 1÷@÷@ð
 %ŸN™N×/Ñ/ˆMØ×,Ñ,×0Ñ0°Ô?Ø×,Ñ,¨]¸MÕJð G÷
@ð @ú÷@ð @ús$   ÃE4Ã5E(ÄE4Å(E1	Å-E4Å4E=c                óL   — | j                  |«       | j                  |||«       y re   )r.  r&  )r  r  r  r#  s       rc   Ú'codegen_subgraph_with_flattened_outputsz<PythonWrapperCodegen.codegen_subgraph_with_flattened_outputsc  s(   € ð 	×$Ñ$ XÔ.Ø×9Ñ9Ø�lÐ$;õ	
rb   c                óL   — | j                  |«       | j                  |||«       y re   )r.  r*  )r  r  r  r)  s       rc   r  z%PythonWrapperCodegen.codegen_subgraphk  s%   € ð 	×$Ñ$ XÔ.Ø×"Ñ" 8¨\Ð;LÕMrb   c                óä  — |j                  «       }| j                  |› dt        |j                  «      › �«       |j                  D �cg c]  }|j                  «       ‘Œ }}t        j                  j                  rOt        t        |j                  «      «      D �cg c]
  }|› d|› d�‘Œ }}| j                  |j                  ||«       y | j                  |j                  ||«       y c c}w c c}w )Nú = [None] * r}  r~  )rh   r¯   rº   r¯  r\   r­  r@   rf   r‰  rþ  r  r  r  )r  Úinvoke_subgraphr³   rR  r  r   r  s          rc   Úcodegen_invoke_subgraphz,PythonWrapperCodegen.codegen_invoke_subgraphq  sÚ   € Ø×'Ñ'Ó)ˆà�‰˜$˜˜|¬C°×0GÑ0GÓ,HÐ+IÐJÔKØ;J×;QÑ;QÖR°C˜×-Ñ-Õ/ÐRˆÐRä�7‰7×Òä(-¬c°/×2IÑ2IÓ.JÓ(KöØ#$�4�&˜˜!˜˜A’ðˆMð ð ×-Ñ-Ø×(Ñ(¨,¸õð ×!Ñ! /×":Ñ":¸LÈ$ÕOùò Sùòs   ÁC(ÂC-c                óð  — |j                  «       }|j                  D �cg c]  }|j                  «       ‘Œ }}|j                  j                  «       }t	        |j                  t
        j                  «      s|› d�}| j                  |› dt        |j                  «      › �«       | j                  d|› d�«       | j                  t        | |j                  j                  «      «       t        j                  j                  rOt        t        |j                  «      «      D �cg c]
  }|› d|› d�‘Œ }}| j!                  |j                  ||«       n| j#                  |j                  ||«       | j                  t%        | «      «       | j                  d«       | j                  t        | |j&                  j                  «      «       t        j                  j                  rOt        t        |j                  «      «      D �cg c]
  }|› d|› d�‘Œ }}| j!                  |j&                  ||«       n| j#                  |j&                  ||«       | j                  t%        | «      «       y c c}w c c}w c c}w )NrÁ  r3  rÆ   r‡  r}  r~  zelse:)rh   Úoperandsr­  Ú	predicaterw   r'   ÚShapeAsConstantBufferr¯   rº   r¯  rQ  Útrue_subgraphrf   r@   r‰  rþ  r  r  rw  Úfalse_subgraph)r  Úconditionalr³   rR  r  r8  r   r  s           rc   Úcodegen_conditionalz(PythonWrapperCodegen.codegen_conditional�  s  € Ø×#Ñ#Ó%ˆà;F×;OÑ;OÖP°C˜×-Ñ-Õ/ÐPˆÐPà×)Ñ)×;Ñ;Ó=ˆ	Ü˜+×/Ñ/´×1IÑ1IÔJà$˜+ WÐ-ˆIà�‰˜$˜˜|¬C°×0CÑ0CÓ,DÐ+EÐFÔGØ�‰˜˜Y˜K qÐ)Ô*Ø�‰Ô(¨¨{×/HÑ/H×/NÑ/NÓOÔPÜ�7‰7×ÒÜ5:¼3¸{×?RÑ?RÓ;SÓ5TÖU° ˜v Q q c¨š^ÐUˆMÐUØ×-Ñ-Ø×)Ñ)¨<¸õð ×!Ñ! +×";Ñ";¸\È4ÔPà�‰Ô'¨Ó-Ô.Ø�‰�wÔØ�‰Ô(¨¨{×/IÑ/I×/OÑ/OÓPÔQÜ�7‰7×ÒÜ5:¼3¸{×?RÑ?RÓ;SÓ5TÖU° ˜v Q q c¨š^ÐUˆMÐUØ×-Ñ-Ø×*Ñ*¨L¸-õð ×!Ñ! +×"<Ñ"<¸lÈDÔQØ�‰Ô'¨Ó-Õ.ùò9 Qùò Vùò Vs   ŸI)Ä"I.ÈI3c                óô  ‡ — ˆ fd„}|j                  «       }|j                  D �cg c]  }|j                  «       ‘Œ }}|j                  D �cg c]  }|j                  «       ‘Œ }}t	        |«      }‰ j                  |› dt	        |«      › �«       |r ‰ j                  |› dt	        |«      › d�«       t        |«      D ]  \  }	}
‰ j                  |› d|	› d|
› �«       Œ  g t        t	        |«      «      D �	cg c]
  }	|› d|	› d�‘Œ c}	¢|¢}|› d�g}t        |«      }|d	t	        |«       } ||j                  ||«       ‰ j                  d
|d   › �«       ‰ j                  d«       |rwt        |«      D ]h  \  }	}‰ j                  t        ‰ |j                  j                  «      «       ‰ j                  |› d|	› d|› d�«       ‰ j                  t        ‰ «      «       Œj nvt        |«      D ]h  \  }	}‰ j                  t        ‰ |j                  j                  «      «       ‰ j                  |› d|	› d|› d�«       ‰ j                  t        ‰ «      «       Œj ‰ j                  d«       ‰ j                  t        ‰ |j                  j                  «      «        ||j                  ||«       ‰ j                  t        ‰ «      «       |r‚‰ j                  t        ‰ |j                  j                  «      «       t        t	        |«      «      D ]"  }	‰ j                  |› d|	|z   › d|› d|	› d�«       Œ$ ‰ j                  t        ‰ «      «       ‰ j                  t        ‰ |j                  j                  «      «        ||j                  ||«       ‰ j                  t        ‰ «      «       ‰ j                  d|d   › �«       |r¯‰ j                  d«       t        t	        |«      «      D ]†  }	‰ j                  d|› d|	|z   › d�«       ‰ j                  t        ‰ |j                  j                  «      «       ‰ j                  |› d|	› d|› d|	|z   › d�«       ‰ j                  t        ‰ «      «       Œˆ y	y	c c}w c c}w c c}	w )z1while_loop is codegened as a host side while_loopc                óˆ   •— t         j                  j                  r‰j                  | ||«       y‰j	                  | ||«       y)z3Helper method to deduplicate subgraph codegen logicN)r@   rf   r‰  r  r0  )r  r  r  r  s      €rc   r  zAPythonWrapperCodegen.codegen_while_loop.<locals>.codegen_subgraph¥  s7   ø€ ä�w‰w×ÒØ×1Ñ1°(¸LÈ-ÕXà×<Ñ<Ø˜l¨Mõrb   r3  z.extend([[] for _ in range(z)])r}  z] = r~  Ú_cond_resultNzshould_loop = r   zif not should_loop:z.unsqueeze(0).clone()r¹  zwhile should_loop:z	].append(z])z    should_loop = z%# Stack outputs after loop completionzif len(z]) > 0:z] = torch.stack(z	], dim=0))rh   Úcarried_inputsr­  Úadditional_inputsrº   r¯   rð  rþ  r(  Úcond_subgraphrQ  Úbody_subgraphrf   rw  )r  Ú
while_loopÚstack_outputr  r³   rR  Úouter_carried_inputsÚouter_additional_inputsÚ
ckp_offsetr   ÚinpÚcond_outer_inputsÚcond_outer_outputsÚbody_outer_inputsÚbody_outer_outputsÚcarried_inputs   `               rc   Úcodegen_while_loopz'PythonWrapperCodegen.codegen_while_loop¢  s‡  ø€ ô	ð ×"Ñ"Ó$ˆà/9×/HÑ/Hö 
Ø(+ˆC×!Ñ!Õ#ð 
Ðð  
ð 0:×/KÑ/Kö#
Ø(+ˆC×!Ñ!Õ#ð#
Ðð #
ô Ð-Ó.ˆ
Ø�‰˜$˜˜|¬CÐ0DÓ,EÐ+FÐGÔHÙØ�N‰NØ�&Ð3´CÐ8LÓ4MÐ3NÈcÐRôô  Ð 4Ó5ò 	3‰FˆAˆsà�N‰N˜d˜V 1 Q C t¨C¨5Ð1Õ2ð	3ð
Ü&+¬CÐ0DÓ,EÓ&FÖG ���a˜�s˜!ŠnÒGð
à$ð
Ðð "&  lÐ3Ð4ÐÜ Øó
Ðð /Ð/J´Ð5IÓ1JÐKÐáØ×$Ñ$Ð&7Ð9Kô	
ð 	�‰˜Ð(:¸1Ñ(=Ð'>Ð?Ô@Ø�‰Ð,Ô-Ùä$-Ð.BÓ$Cò 7Ñ ��=Ø—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘ $  q¨¨¨4°¨Ð>SÐTÔUØ—‘Ô/°Ó5Õ6ñ7ô
 %.Ð.BÓ$Cò 7Ñ ��=Ø—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘ $  q¨¨¨4°¨¸hÐGÔHØ—‘Ô/°Ó5Õ6ð7ð
 	�‰Ð+Ô,à�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOÙØ×$Ñ$Ð&7Ð9Kô	
ð 	�‰Ô'¨Ó-Ô.ñ Ø�N‰NÔ,¨T°:×3KÑ3K×3QÑ3QÓRÔSÜœ3Ð3Ó4Ó5ò P�Ø—‘ $  q¨¨Z©Ð(8¸	À$ÀÀqÈÈÈ2ÐNÕOðPà�N‰NÔ+¨DÓ1Ô2ð 	�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOÙØ×$Ñ$Ð&7Ð9Kô	
ð 	�‰Ô'¨Ó-Ô.Ø�‰Ð+Ð,>¸qÑ,AÐ+BÐCÔDñ Ø�N‰NÐBÔCÜœ3Ð3Ó4Ó5ò 7�Ø—‘ ¨¨¨a°°J±Ð/?¸wÐGÔHØ—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘Ø�f˜A˜a˜SÐ 0°°°a¸¸J¹Ð7GÀyÐQôð —‘Ô/°Ó5Õ6ñ7ð ùòQ 
ùò#
ùò  Hs   ¥Q+ÁQ0Ã4Q5c                óÒ   — 	 t        | dd «      ry t        | t        «      r| S t        j                  j
                  j                  | «      }|€|S t        |«      S # t        $ r Y y w xY w)Nrg  )rð  rw   rY   r@   rf   Ú
_shape_envÚ_maybe_evaluate_staticr9  )r½   rò  s     rc   Ústatically_known_int_or_nonez1PythonWrapperCodegen.statically_known_int_or_none  sf   € ð	Ü�q˜.¨$Ô/ð Ü˜!œSÔ!Ø�Ü—'‘'×$Ñ$×;Ñ;¸AÓ>ˆCØˆ{Ø�
Ü�s“8ˆOøÜò 	Ùð	ús!   ‚A �A ¢,A Á
A Á	A&Á%A&c                ól   — g }| D ],  }t         j                  |«      }|€ y |j                  |«       Œ. |S re   )rR  rT  r�   )Úlstr£  r½   r)  s       rc   Ú%statically_known_list_of_ints_or_nonez:PythonWrapperCodegen.statically_known_list_of_ints_or_none  sA   € àˆØò 	ˆAÜ&×CÑCÀAÓFˆCØˆ{ÙØ�M‰M˜#Õð		ð
 ˆrb   c                ó0   — t         j                  | «      d uS re   )rR  rW  )rV  s    rc   Ú is_statically_known_list_of_intsz5PythonWrapperCodegen.is_statically_known_list_of_ints  s   € ô !×FÑFÀsÓKÐSWÐWð	
rb   c                óH   — t         j                  | j                  «       «      S re   )rR  rW  rM  ©r  s    rc   rD  z4PythonWrapperCodegen.static_shape_for_buffer_or_none!  s   € ä#×IÑIØ�O‰OÓó
ð 	
rb   c                ó0   — t         j                  | «      d uS re   )rR  rD  r[  s    rc   Ú!can_prove_buffer_has_static_shapez6PythonWrapperCodegen.can_prove_buffer_has_static_shape'  s   € ä#×CÑCÀFÓKÐSWÐWÐWrb   c                 ó   — y re   ra   )r  r°  Únode_schedules      rc   Úwrite_kernel_context_guardz/PythonWrapperCodegen.write_kernel_context_guard+  s   € ð
 	rb   c                 ó   — y)z<
        Mark the beginning of kernel context guard
        Nra   r  s    rc   Ú write_kernel_context_guard_beginz5PythonWrapperCodegen.write_kernel_context_guard_begin2  ó   € ð 	rb   c                 ó   — y)z6
        Mark the end of kernel context guard
        Nra   r  s    rc   Úwrite_kernel_context_guard_endz3PythonWrapperCodegen.write_kernel_context_guard_end:  rc  rb   re   )r  r  r  rË   r  úPythonWrapperCodegen | Noner  ú!ir.GraphPartitionSignature | Noner_  )r³   rW   r   rW   rÈ   r>  )rÑ  rW   )r>  ÚTritonMetaParamsrÈ   rW   ©rÈ   rÐ  )rÈ   z9dict[str, ir.TensorBox | ir.TorchBindObject | sympy.Expr]©rÈ   zlist[IRNode])rZ  rÐ  ©rÈ   rY   )rZ  zIterable[str]rÈ   r>  )
r³   rW   rˆ   rW   r‰   rW   r¹  rW   rÈ   r>  )r[  r7   r³   rW   rˆ   rW   r‰   rW   r¹  rW   rÈ   r>  )rv  rW   r¹  rW   rÈ   r>  )r[  r7   rv  rW   r¹  rW   rÈ   r>  )r�  rY   rÔ  rW   rÈ   rW   rô  )rA   N)r�  rY   r÷  rY   rù  zdict[int, int] | NonerÈ   r>  )rü  rY   rÈ   r  )rŸ  rÐ  rÈ   r>  ©r£  r7   rÈ   r>  )rp   zir.FallbackKernelrÈ   r>  )rp   r·  rÈ   r>  )rp   rš  )rp   r¤  rÈ   r>  )r  rW   rÉ  rW   rÊ  rË   rŸ  rÐ  r[   rW   rË  zOrderedSet[str] | NonerÈ   r>  )F)rp   rœ  )rp   rŒ  rÈ   r>  )rú  rW   r£  rW   rû  zCallable[[], Sequence[str]]rü  z6torch._ops.OpOverload | torch._ops.HigherOrderOperatorrÏ  ry  r¯  zSequence[ir.Buffer]rÈ   r>  )r  r7   r¯   úCallable[..., None]rÈ   zIterator[Callable[..., None]])r  r  )r³   rW   r  rk  ra  rl  )r_  r  )r½   r
   ro   r  rÈ   rW   )r½   r
   rÈ   rW   )r‚  rW   r³   rW   r6  rW   rÈ   rW   )rT  zSequence[Expr]rÈ   rW   )rÈ   ztuple[str, list[str]])r¯   rm  rÈ   rW   )r˜  z
bool | str)rX   rW   rÈ   r>  )rp   zir.MultiOutput)NTN)
r°  rW   rÝ  rW   rÞ  rË   rß  r  rà  rË   )r°  rW   rÝ  rW   rÞ  rË   )r³   rW   )rÖ   zlist[list[int | sympy.Expr]]r  ú$tuple[ir.ComputedBuffer, str] | None)r°  rW   rÓ  rË   )r§  r  rf   rP   rÈ   r>  )r\  rF   )r°  rW   )r  r/  )NF)rÆ  )r  rÄ  )rÌ  rÐ  )rÇ  rW   rÆ  rW   rÐ  rW   )rÓ  r‚  )r	  r/  r×  r/  rl  r  )r³   rW   rg  zir.ReinterpretViewrÈ   r>  ©r  rß  )ré  rß  rð  rß  )r®  rW   r¯  r   r±  r°  rÈ   r>  )r  rY   r  zir.GraphPartitionSignature)r   rY   )r°  rW   r_  z*Sequence[BaseSchedulerNode] | ExternKernel)§r]   r^   r_   rÁ  rÄ  r`   r%  rj  r  rå  rí  rè  r,  rê  r2   r6  r:  r<  rA  rE  rG  rJ  rC  rS  rU  rW  r\  ra  rc  ré  rj  rq  rp  r»  rw  rÂ  r~  r�  ró  r†  rX  r{  rT  ry  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Í   Úcontextmanagerr
  r  rÿ  r  rB  rN  r  rc  ri  rq  r  rv  ry  r{  r€  r§   rQ  r‚  r•  r™  r�  r   r©  r¼  r¾  rs  ræ  r  rë  ró  rã  r÷  Úclassmethodrû  rK  rR  r©  r^  r`  rX  rö  r  r  r  r  r  r  rq  r¡   r~  r‚  r×  r¯   rØ  r  rÐ  rK  r¼  r[  rÂ  rÈ  rÇ  rÍ  rÑ  rÔ  ro  rh  r°  rì  rë  ró  rõ  r÷  rù  r³  r  r  r!  r&  r*  r.  r0  r  r5  r=  rP  rT  rW  rY  rD  r]  r`  rb  re  r@  rA  s   @rc   rR  rR  é  sÙ  ø… ñð "Ð�dÓ!ôfBðP ð
 CGð	&Øð&à!ð&ð 4ð&ð @ò	&ó ð&ó'ó<óoBóbóð: òó ðó"	Bð ò+ó ð+ó!ð ò
ó ð
ó
ð$à	Bó$ó
%ó?ó0(ó
ó.óó8)ó:ó6-óKðPØðPØ"ðPØ,/ðPØ:=ðPà	óPðUàðUð ðUð ð	Uð
 ðUð ðUð 
óUóHð

Ø"ð

Ø14ð

Ø?Bð

à	ó

óó*Góò.ò1ò0òJò
/ó+ð Ø<@ð	*-àð*-ð ð*-ð %:ð	*-ð
 
ó*-óX5ðàðð 
$óó,4ó5ó$óóó:ó4ó,òð:8à ð8ð 
ó8ð 04ð;àð;ð ð;ð ð	;ð
 ð;ð ð;ð -ð;ð 
ó;ó ó(	óUòó
8òó(Bò<ð
	Vàð	Vð  ð	Vð .ð		Vð
 Lð	Vð  ð	Vð %ð	Vð 
ó	Vò0óð ×Ñð!àð!ð 'ð!ð 
'ò	!ó ð!ó2,òS
òj.Sò`:ò
'
óR%ð>Kàð>Kð ð>Kð -ó	>Kò@'?óRDòð @Dõ Wð CGõ +ó.ó&ó'ó6ðà	óð, ð<Pð 'ð<Pð 
ó<Pó|>ó5óSò
IòGò.6ò4ò*^ò@ð4  $ØØ%)ð
àð
ð ð
ð ð	
ð
 ð
ð #ó
ð& àCGð	Øð	Ø'*ð	Ø6@ò	ó ð	ð  $ØØ%)ð!àð!ð ð!ð ð	!ð
 ð!ð #ó!ó,>ð ñó ððuJð ,ðuJð >óuJôn	ð*AØ"ðAØ+8ðAà	óAó
%óN:ò.ò>ò8ò5òWòFò
ò
ò@
4ó=Gò~'ð ØØØØØØØ!ñ6
àó6
ðz ØØØØØØØØ!ØñP(àóP(òd ò!ò,ó%óP	
ð ñ8ó ð8ð TYó.ò`*ó_ó*óAóuð $(ðð !óó`ó 
ó53òn:ó(
ò 
óEò'ð
àð
ð ð
ð Eð	
ð
 
ó
ðKàðKð ðKð Eð	Kð
 
óKòZ+'ðZ<àð<ð 9ó<óBYò
ò$
ò"Kò*
òNòPó /òB]7ð~ ñó ðð ñó ðð ñ
ó ð
ð
 ñ
ó ð
ð
 ñXó ðXðàðð Bóòörb   rR  c                  óþ   ‡ — e Zd ZdZ	 d	 	 	 	 	 dˆ fd„Zdd„Zdd„Zd„ 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 ˆ fd„Z	 d	 	 	 	 	 d!d„Zd"d„Zdd„Zedd„«       Zedd„«       Zed#d„«       Zd„ Zˆ xZS )$r  a  
    A wrapper codegen that generates code for a subgraph. For most of the
    methods, we rely on the implementation in the PythonWrapperCodegen. But we
    override a few functions to produce cleaner code (like avoiding writing
    imports twice in the output code)
    c                óô   •— || _         || _        || _        t        ‰| �  «        | j                  «       }|j                  | _        |j                  | _        |j                  | _        |j                  | _	        y re   )
r  r  r  r$  r%  Úget_root_graphrÖ  r±   rÚ  râ  )r  r  r  r  Úrootr,  s        €rc   r%  z%SubgraphPythonWrapperCodegen.__init__K  sn   ø€ ð +ˆÔØ,ˆÔØ$8ˆÔ!ä‰ÑÔà×"Ñ"Ó$ˆà$(×$=Ñ$=ˆÔ!Ø%)×%?Ñ%?ˆÔ"à!×/Ñ/ˆÔà)-×)GÑ)GˆÕ&rb   c                ó&   — | j                   | _        y re   )r  rä  r  s    rc   rå  z1SubgraphPythonWrapperCodegen.set_launcher_fn_nameb  s   € ð !%× 2Ñ 2ˆÕrb   c                 ó   — y re   ra   r  s    rc   rè  z)SubgraphPythonWrapperCodegen.write_headerh  r-  rb   c                 ó   — y re   ra   ré  s     rc   r  z2SubgraphPythonWrapperCodegen.add_benchmark_harnessk  r-  rb   c                 ó   — y re   ra   ré  s     rc   ræ  z6SubgraphPythonWrapperCodegen.benchmark_compiled_modulen  r-  rb   c                 ó   — y re   ra   r  s    rc   rW  z5SubgraphPythonWrapperCodegen.write_async_compile_waitq  r-  rb   c                ó6   — | j                   j                  «       S re   )r  r�  r  s    rc   r�  z/SubgraphPythonWrapperCodegen.next_kernel_suffixt  s   € à×"Ñ"×5Ñ5Ó7Ð7rb   c                 ó   — y re   ra   r¢  s     rc   r¨  z2SubgraphPythonWrapperCodegen.generate_after_suffixx  rH  rb   c                ó\   — | j                   j                  d| j                  › d�«       d}|S )Nz
            def z(args):
            rA   )rÒ  rû   rä  r_  s     rc   ra  z>SubgraphPythonWrapperCodegen.write_launcher_fn_call_get_indent{  s<   € Ø�‰×ÑðØ×&Ñ&Ð'ð (ðô	
ð
 ˆØÐrb   c                 ó   — yr¸   ra   r  s    rc   r  z4SubgraphPythonWrapperCodegen.get_wrapper_call_indent„  s   € Ørb   c                óÆ   — | j                   x}r3|j                  |j                  D �ci c]  }t        |«      |“Œ c}z  }|S t        j
                  j                  }|S c c}w re   )r  Úinput_nodesr  rW   r@   rf   r  )r  r  rE  r\   s       rc   rJ  z-SubgraphPythonWrapperCodegen.get_graph_inputs‡  sh   € ð ×1Ñ1Ð1ˆ9Ð1Ø×*Ñ*Ø#,×#:Ñ#:ö.Ø”�A“˜‘	ò.ñ ˆFð
 ˆô —W‘W×)Ñ)ˆFØˆùò.s   ¨Ac                óô   — | j                   x}rJt        |j                  j                  «       «      |j                  D �cg c]  }|j
                  ‘Œ c}z   }|S t        j                  j                  }|S c c}w re   )	r  r(  r€  r  r  r³   r@   rf   rQ  )r  r  r  Únamess       rc   rc  z2SubgraphPythonWrapperCodegen.get_graph_input_names’  sr   € Ø×1Ñ1Ð1ˆ9Ð1Ü˜×.Ñ.×3Ñ3Ó5Ó6Ø6?×6MÑ6Mö:Ø&2�×!Ó!ò:ñ ˆEð
 ˆô —G‘G×-Ñ-ˆEØˆùò:s   ¿A5c                ór   — | j                   x}r|j                  }|S t        j                  j                  }|S re   )r  r  r@   rf   r  )r  r  r¯  s      rc   rC  z.SubgraphPythonWrapperCodegen.get_graph_outputs›  s;   € Ø×1Ñ1Ð1ˆ9Ð1Ø×,Ñ,ˆGð ˆô —g‘g×+Ñ+ˆGØˆrb   c                ó~   •— |j                  «       }| j                  x}r||j                  v ry t        ‰| �  |«       y re   )rh   r  r€  r$  r°  )r  r  r³   r  r,  s       €rc   r°  z/SubgraphPythonWrapperCodegen.codegen_allocation¢  s?   ø€ Ø�‰Ó ˆØ×2Ñ2Ð2ˆIÐ2¸À	×@UÑ@UÑ8Uð ä‰Ñ" 6Õ*rb   c                ó>  — | j                  «        t        j                  j                  j                  x}�D|› d�}| j                  |› dt        j                  j                  j                  |«      › �«       |S t        |«      }| j                  |› d|› d�«       |S )NÚ_rawrî   rƒ  rï   )	r6  r@   rf   rj   r—  r¯   rŠ  Ústream_handler0   )r  r�  rf   r—  r³   s        rc   Ú_write_get_raw_streamz2SubgraphPythonWrapperCodegen._write_get_raw_stream¬  sœ   € ð 	×%Ñ%Ô'ä#$§7¡7×#4Ñ#4×#HÑ#HÐHÐÐUØ)Ð*¨$Ð/ˆDØ�N‰NØ�&˜œAŸG™G×.Ñ.×<Ñ<Ð=PÓQÐRÐSôð ˆô ' zÓ2ˆDØ�N‰N˜d˜VÐ#5°j°\ÀÐCÔDØˆrb   c                ó.   — | j                  d|› d�«       y)z#Generate NVTX range push for graph.z"torch.cuda.nvtx.range_push('graph re  Nrt  )r  Úpost_grad_graph_ids     rc   Úcodegen_graph_nvtx_range_pushz:SubgraphPythonWrapperCodegen.codegen_graph_nvtx_range_push»  s   € à�‰Ð;Ð<NÐ;OÈrÐRÕSrb   c                ó&   — | j                  d«       y)z"Generate NVTX range pop for graph.ztorch.cuda.nvtx.range_pop()Nrt  r  s    rc   Úcodegen_graph_nvtx_range_popz9SubgraphPythonWrapperCodegen.codegen_graph_nvtx_range_pop¿  s   € à�‰Ð4Õ5rb   c                ó8   — | j                   j                  «        y re   )r  r6  r  s    rc   r6  z5SubgraphPythonWrapperCodegen.write_triton_header_onceÃ  s   € ð 	×Ñ×4Ñ4Õ6rb   c                ó8   — | j                   j                  «        y re   )r  r<  r  s    rc   r<  z=SubgraphPythonWrapperCodegen.write_get_raw_stream_header_onceÌ  s   € ð 	×Ñ×<Ñ<Õ>rb   c                óˆ   — | }t        |t        «      r|j                  }t        |t        «      rŒt        |t        «      sJ ‚|S re   )rw   r  r  rR  )r  ru  s     rc   rt  z+SubgraphPythonWrapperCodegen.get_root_graphÖ  sB   € àDHˆÜ˜Ô;Ô<Ø×&Ñ&ˆDô ˜Ô;Õ<ô ˜$Ô 4Ô5Ð5Ð5Øˆrb   c                 ó   — y re   ra   r  s    rc   r  z<SubgraphPythonWrapperCodegen.generate_and_run_autotune_blockß  s   € àrb   re   )r  rW   r  rR  r  rg  r_  rô  rl  rk  )rÈ   z@dict[str, ir.TensorBox | ir.TorchBindObject | sympy.Expr | None]ri  rj  ro  )r�  rY   rf   zGraphLowering | NonerÈ   rW   )rŠ  rY   rÈ   r>  )rÈ   rR  )r]   r^   r_   rÁ  r%  rå  rè  r  ræ  rW  r�  r¨  ra  r  rJ  rc  rC  r°  rˆ  r‹  r�  r2   r6  r<  rt  r  r@  rA  s   @rc   r  r  C  sí   ø„ ñð CGð	HàðHð -ðHð @õ	Hó.3óòòòó8óóóð	à	Ió	óóõ+ð >BðØðØ&:ðà	óóTó6ð ò7ó ð7ð ò?ó ð?ð òó ðörb   r  )rp   r/  rÈ   r<  )rp   r/  rÈ   r?  )r€   r/  r�   r/  )NN)r³   rW   rÕ   zlist[triton.Config]rÖ   zlist[TritonGrid]r£   rf  r×   rË   rÈ   ztuple[str, str]re   )r  rn  rÈ   rW   )ÂÚ
__future__r   r&  rÍ   rî  ró   rE  r9  r1  r  r  rR  rî  r3  Úcollections.abcr   Ú	itertoolsr   r   Útypingr   r   r	   r–   r
   r  Ú
torch._opsÚtorch.utils._pytreeÚutilsÚ_pytreerü  r   rv  Útorch._dynamo.utilsr   r   r   Ú#torch._inductor.codegen.debug_utilsr   Ú$torch._inductor.codegen.multi_kernelr   Ú%torch._inductor.runtime.runtime_utilsr   Útorch._library.opaque_objectr   r   Útorch._loggingr   Ú%torch.fx.experimental.symbolic_shapesr   r   r   r   r   Útorch.fx.noder   Ú torch.multiprocessing.reductionsr   Útorch.utils._ordered_setr   Útorch.utils._sympy.functionsr   r   Ú torch.utils._sympy.singleton_intr    rY  r"   r#   rÆ  r%   r&   r'   Ú	codecacher(   r)   r*   rþ   r+   Úruntime.hintsr,   Ústream_constantsr-   r.   r/   Ústream_utilsr0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   Úvirtualizedr@   r'  rB   rC   rD   rE   rF   rG   Ú	cpp_utilsrH   Úcustom_extern_kernel_codegenrI   Útriton_utilsrJ   rK   rL   rM   rN   rO   r¨   rf   rP   rQ   rj   rR   Úwrapper_fxirrS   Ú	getLoggerr]   ÚlogÚdoprintrx  r¦   r[   rW   r  rY   r<  r?  r†   r/  rN  r­  rV   rt   r|   r„   r“   r  rh  rÊ   rá   r  r  r!  rC  rT   rQ  rc  rl  rq  rw  r€  r“  r™  r£  r¶  rÃ  rÌ  rÜ  rç  rö  r  r  r
  r  r0  r=  rb  rC  rA  rv  r‹  r›  r¦  r­  r·  r¿  rî  ÚLinerR  r  ra   rb   rc   ú<module>r³     s™  ðå "ã Û Û Û 
Û Û Û Û Û 	Û Û 	Û Ý $ß "ß /Ñ /ã Ý ã Û ß $Ð $Ý &ß EÑ EÝ CÝ AÝ ;ß RÝ +÷õ õ .Ý ;Ý /ß 1Ý 9ß :ç (Ñ (Ý 'ß (Ý 'Ý ,ß WÑ Wß ?÷÷ ÷ ÷ õ  ÷÷ õ Ý Fß PÑ Pñ ß<Ñ<ãå%Ý!Ý-Ý)ð €g×Ò˜Ó!€á‹×Ò€ð �—’˜uŸ{™{¨C°°sÐ:Ñ;€Ø˜5Ÿ<š<¨¯©°cÐ;NÐPSÐSÑTÐ Ø�YŠY˜Ñ%€
Ø˜]˜O¨TÐ1Ñ2Ð ð ×Ò÷3ð 3ó ð3óó óò>QðF ˜˜S˜‘>Ð à	ˆ#�—
‘
Ñ
˜CÐ
Ñ  8Ð-=Ð,>ÀÀcÈ3ÀhÁÐ,OÑ#PÑPð ð ,0Ø'+ðk&Ø
ðk&à ðk&ð ðk&ð )ð	k&ð
 %ðk&ð ók&ð^ EIðq&ØAðq&àóq&ðh ×Ò÷ð ó ð÷%6ñ %6ôPQ�8ô Q÷Yñ Yð
 ×Òô2˜ó 2ó ð2ð ×Òô	/�kó 	/ó ð	/ð ×Òô+�+ó +ó ð+ð ×Òô	2˜ó 	2ó ð	2ð ×Òô1�{ó 1ó ð1ð ×Òô"@ Kó "@ó ð"@ôJ? ;ô ?ð ×Òô
7˜Kó 
7ó ð
7ð ×Òô5˜+ó 5ó ð5ð< ×ÒôP˜{ó Pó ðPð. ×Òô	(ˆ{ó 	(ó ð	(ð ×Òô /�[ó  /ó ð /ðF ×Òô5˜;ó 5ó ð5ð* ×Òô;˜ó ;ó ð;ð2 ×ÒôÐ2Oó ó ðð: ×ÒôÐ1Mó ó ðð ×Òô ó ó ðð ×Òô ó ó ð÷%
ñ %
ðP ×Òôh,Ð%ó h,ó ðh,ðV ×Òô+6Ð,ó +6ó ð+6ð\ ×Òô/Ð(ó /ó ð/ð& ×Òô)Ð"ó )ó ð)ô0(Ð!ô (ð
 ×Òô#0�kó #0ó ð#0ðL ×Òô6˜;ó 6ó ð6ð, ×Òô4˜+ó 4ó ð4ð8 ×Òô	5˜+ó 	5ó ð	5ð ×Òô8˜[ó 8ó ð8ð ×Òô6˜;ó 6ó ð6ð" ×ÒôS˜+ó Só ðSð  €
Ø˜KÑ'€ôW1˜7ô W1ôtb^Ð#7õ ^rb   