Ë
    g^(h#Â ã                  ó  — 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mZ d dlmZmZ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! d dl"m#Z# d d	l$m%Z% d d
l&m'Z' d dl(m)Z)m*Z*m+Z+m,Z,m-Z- d dl.m/Z/ d dl0m1Z1 d dl2m3Z3 d dl4m5Z5m6Z6 ddl7m8Z8m9Z9m:Z: ddl;m<Z< ddl:m=Z=m>Z> ddl?m@Z@ ddlAmBZB ddlmCZCmDZDmEZEmFZFmGZGmHZHmIZI ddlJmKZK ddlLmMZMmNZNmOZOmPZPmQZQmRZRmSZS ddlTmUZU ddlVmWZWmXZXmYZY erd dlZm[Z[m\Z\ d dl]Z]ddl^m_Z_  eQ«       jÀ                  ZaebejÆ                  ej:                  edf   Zeee:jÌ                  eRf   ZgdHd„ZhdId„ZidJd „ZjdKd!„ZkdLd"„Zlemedenf   Zoeebeenej,                  f   d#f   eeogebend#f   f   f   Zp	 dM	 	 	 	 	 	 	 	 	 dNd$„ZqdOd%„Zrejæ                   G d&„ d'«      «       Zt G d(„ d)«      Zu G d*„ d+«      Zvejæ                   G d,„ d-ev«      «       Zwejæ                   G d.„ d/ev«      «       Zxejæ                   G d0„ d1ev«      «       Zy G d2„ d3ev«      Zzejæ                   G d4„ d5ev«      «       Z{ejæ                   G d6„ d7e{«      «       Z|ejæ                   G d8„ d9e{«      «       Z}ejæ                   G d:„ d;e{«      «       Z~ G d<„ d=e{«      Zejæ                   G d>„ d?ev«      «       Z€ejæ                   G d@„ dAe€«      «       Z�ejæ                   G dB„ dCe€«      «       Z‚edZƒ G dD„ dEeN«      Z„ G dF„ dGe„«      Z…y)Pé    )ÚannotationsN)Úcount)ÚAnyÚCallableÚOptionalÚTYPE_CHECKINGÚUnion)ÚExpr)Údtype)ÚcountersÚdynamo_timed)ÚDebugPrinterManager)ÚMultiKernelState)Ú	cache_dir)ÚCallMethodKeyÚConvertIntKeyÚDivideByKeyÚresolve_unbacked_bindingsÚSymTypes)Ú_get_qualified_name)Ú
OrderedSet)ÚSingletonInt)Úsymbol_is_typeÚSymTé   )Úasync_compileÚconfigÚir)Úoutput_code_log)ÚIRNodeÚReinterpretView)Útriton_heuristics)ÚDeviceProperties)Úcache_on_selfÚget_benchmark_nameÚLineContextÚsympy_productÚ	sympy_strÚ
sympy_subsÚtriton_version_uses_attrs_dict)ÚVé   )ÚArgNameÚCodeGenÚDeferredLineÚIndentedBufferÚPythonPrinterÚWorkspaceArgÚWorkspaceZeroMode)Úcexpr)Ú	config_ofÚshould_unwrap_unspec_argÚsignature_to_meta)ÚIteratorÚSequence)ÚGraphLoweringc                óâ   — t         j                  j                  | «      }| j                  «       | j	                  «       t        t         j                  j                  j                  |«      «      fS ©N)r+   ÚgraphÚget_allocation_storage_sizeÚget_device_or_errorÚ	get_dtyper(   ÚsizevarsÚsimplify)ÚnodeÚstorage_sizes     ú]/var/www/skyplay_api_hub/venv/lib/python3.12/site-packages/torch/_inductor/codegen/wrapper.pyÚbuffer_reuse_keyrF   T   sU   € Ü—7‘7×6Ñ6°tÓ<€Là× Ñ Ó"Ø�‰Óô 	”!—'‘'×"Ñ"×+Ñ+¨LÓ9Ó:ðð ó    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)
r?   r@   r+   r=   rA   rB   r>   r(   Ústatically_known_geqÚstatically_known_leq)Ú	input_bufÚ
output_bufÚ
input_sizeÚoutput_sizes       rE   Úcan_match_buffer_sizerO   `   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ààrG   c                óœ  — ddl m}m} t        | j                  «      }|dk(  r.| j
                  �| j
                  j                  rd|› d�S d|› d�S ||v r||   }|S |j                  «       D ]Q  \  }}t        j                  |dz   |«      }t        |«      dk(  sŒ.|d   }||v sJ d	|› d
|› �«       ‚||   }	|› d|	› d�c S  t        d|› �«      ‚)Nr,   )ÚCONTAINER_PYTHON_TO_CPPÚPYTHON_TO_CPPÚTensorzat::ú&z const&z\[([a-zA-Z_]+)]r   zunsupported z type in convert_arg_type: ú<ú>zunsupport python_type: )ÚcpprQ   rR   ÚreprÚ	real_typeÚ
alias_infoÚis_writeÚitemsÚreÚfindallÚlenÚAssertionError)
ÚargrQ   rR   Úpython_typeÚcpp_typeÚpy_containerÚcpp_containerÚcontainer_matchÚcontained_typeÚcpp_contained_types
             rE   Úconvert_arg_typeri      s  € ß;ô �s—}‘}Ó%€Kà�hÒà�>‰>Ð%¨#¯.©.×*AÒ*AØ˜+˜ aÐ(Ð(à˜+˜ gÐ.Ð.à�mÑ#Ø  Ñ-ˆØˆð (?×'DÑ'DÓ'Fò <Ñ#ˆ�mÜŸ*™* \Ð4FÑ%FÈÓTˆÜˆÓ 1Ó$Ø,¨QÑ/ˆNØ! ]Ñ2ð Ø˜|˜nÐ,GÈÐGWÐXóÐ2ð "/¨~Ñ!>ÐØ#�_ AÐ&8Ð%9¸Ð;Ò;ð<ô Ð2°;°-Ð@Ó
AÐArG   c                ó¢   — t        | j                  «      }dddœ}|j                  |d «      }|€
J d|› �«       ‚|dk(  r| j                  �|dz  }|S )Nz
at::Tensorzstd::vector<at::Tensor>)rS   zList[Tensor]zNYI return type: rS   rT   )rX   rY   ÚgetrZ   )Úretrb   Úpython_to_cpprc   s       rE   Úconvert_return_typern   ž   sk   € ä�s—}‘}Ó%€KàØ1ñ€Mð
 × Ñ  ¨dÓ3€HØÐÐBÐ#4°[°MÐ!BÓBÐð �hÒ 3§>¡>Ð#=Ø�C‰ˆØ€OrG   c                ó¦  — | j                   j                  }| j                   j                  }t        |«      }|dkD  sJ d«       ‚|dk(  rt	        |d   «      }n3|dkD  r.dj                  |D �cg c]  }t	        |«      ‘Œ c}«      }d|› d�}|D �cg c]  }t        |«      › d|j                  › �‘Œ }}› ddj                  |«      › d	�S c c}w c c}w )
Nr   z#must have at least one return valuer,   ú, zstd::tuple<rV   ú ú(ú))Ú_schemaÚ	argumentsÚreturnsr_   rn   Újoinri   Úname)	ÚkernelÚargsrv   Únum_returnsÚcpp_return_valueÚrÚtuple_returnsra   Úcpp_arg_types	            rE   Úget_cpp_op_schemar€   °   sÚ   € Ø�>‰>×#Ñ#€DØ�n‰n×$Ñ$€Gä�g“,€KØ˜Š?ÐAÐAÓAˆ?à�aÒÜ.¨w°q©zÓ:ÑØ	�qŠØŸ	™	À7Ö"K¸aÔ#6°qÕ#9Ò"KÓLˆØ(¨¨°qÐ9ÐàEIÖJ¸cÔ'¨Ó,Ð-¨Q¨s¯x©x¨jÒ9ÐJ€LÐJØÐ˜q §¡¨<Ó!8Ð 9¸Ð;Ð;ùò	 #Lùò Ks   Á,C	Â!C.c                óÊ  ‡ ‡‡‡— t        «       Š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        |«      dk(  r ||d   «      \  }}	 |d	|› �d	|	› �«       nåt        |«      dkD  sJ ‚t        |«      t        |«      k(  sJ ‚t        t           «       }
t        t        ||«      d
„ d¬«      D ]�  \  }}|j                  rD|j                  j                  «       D � �cg c]  \  } }d| › d|› �‘Œ }} }dj                  |«      }nd} ||«      \  }}	d|› d|› �}||
v rŒp|
j!                  |«        ||d|› d|	› �«       Œ‘ d d d «       d d d «       |‰j#                  «       fS c c}} w # 1 sw Y   Œ)xY w# 1 sw Y   Œ-xY w)Nc                ód   — t        | t        j                  «      r| S t        j                  | «      S r<   )Ú
isinstanceÚsympyr
   ÚInteger)Úitems    rE   Ú_convert_to_sympy_exprz@user_defined_kernel_grid_fn_code.<locals>._convert_to_sympy_exprÐ   s#   € Ü! $¬¯
©
Ô3ˆtÐL¼¿¹ÀtÓ9LÐLrG   c                óð   •— ‰�t        | «      r| | fS t        ˆfd„| D «       «      }‰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r<   © )Ú.0Úgr‡   s     €rE   ú	<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r<   ©Úgenerate_example_arg_valueÚtype)r‹   rŒ   Úwrappers     €rE   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Ú
sympy_gridr‡   r’   s     €€rE   Údetermine_gridz8user_defined_kernel_grid_fn_code.<locals>.determine_gridÓ   s…   ø€ ð ˆ?œh tœnà˜�:ÐäÓC¸dÔCÓCˆ
à×.Ñ.¨zÓ:ô —=‘=×9Ò9ð ×2Ñ2Üó à!+ôó óð
ð 	
ð ð
ð 	
rG   c                óÀ   •— ‰j                  | «       ‰rJt        j                  j                  r/‰‰j                  vr ‰j
                  j                  |xs | «       y y y y r<   )Ú	writeliner   r—   r˜   Úkernel_autotune_namesÚkernel_autotune_calls)ÚlineÚexample_gridrx   Úoutputr’   s     €€€rE   r�   z3user_defined_kernel_grid_fn_code.<locals>.writelineï   sW   ø€ Ø×Ñ˜ÔáÜ—‘×6Ò6Ø˜G×9Ñ9Ñ9à×)Ñ)×3Ñ3°LÒ4HÀDÕIð :ð 7ð rG   Úgrid_wrapper_for_údef z(meta):r,   r   zreturn c                ó2   — t        | d   j                  «      S ©Nr,   ©r_   Úkwargs©Úxs    rE   ú<lambda>z2user_defined_kernel_grid_fn_code.<locals>.<lambda>  ó   € ´3°q¸±t·{±{Ó3C€ rG   T©ÚkeyÚreversezmeta['z'] == z and ÚTrueúif z	: return )r†   úUnion[int, sympy.Expr]Úreturnú
sympy.Expr)r™   Ú
TritonGridr<   )r    Ústrr¡   úOptional[str])r0   r   r—   r˜   rŸ   ÚindentÚ
contextlibÚnullcontextr_   r   r¶   ÚsortedÚzipr¨   r\   rw   ÚaddÚgetvalue)rx   ÚconfigsÚgridsr’   r›   r�   Úfn_nameÚkernel_autotune_calls_indentr™   r¡   ÚseenÚcÚvalÚguardsÚ	statementr‡   r¢   s   `  `           @@rE   Ú user_defined_kernel_grid_fn_coderÈ   È   s  û€ ô Ó€FóMð
Øö
÷8Jð Jð " $ Ð(€GÙ��W�I˜WÐ%Ô&ñ ”v—}‘}×=Ò=ð 	×%Ñ%×,Ñ,Ô.ä×#Ñ#Ó%ð !ð
 
�‰‹ñ LÐ6ñ LÜˆu‹:˜Š?Ù!/°°a±Ó!9ÑˆD�,Ù˜ ˜vÐ&¨'°,°Ð(@ÕAä�u“: ’>Ð!�>Ü�u“:¤ W£Ò-Ð-Ð-Üœc‘?Ó$ˆDô "Ü�E˜7Ó#Ñ)CÈTôò L‘��að —8’8àDEÇHÁHÇNÁNÓDT÷Ù7@°t¸S˜&   f¨S¨EÒ2ð�Fñ ð %Ÿ\™\¨&Ó1‘Fà#�FÙ%3°DÓ%9Ñ"��lØ! & ¨°4°&Ð9�	Ø Ñ$ØØ—‘˜Ô#Ù˜) s¨6¨(°)¸L¸>Ð%JÕKðL÷L÷ Lð< �F—O‘OÓ%Ð%Ð%ùó÷#Lð Lú÷ Lð Lús8   ÂGÂB/GÅGÅAGÆ%GÇGÇG	ÇGÇG"c                óà   ‡‡‡‡‡— t        «       Š‰j                  | j                  d¬«       ddlmŠ ddlmŠ t        | j                  g«      Šˆˆˆˆˆfd„Š ‰| «       ‰j                  «       S )zg
    Given a triton kernel function pointer collect the transitive closure of
    its dependencies
    T©Ústripr   )ÚJITFunction)Ú	constexprc           	     óÈ  •— 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        |t        t        t         ‰
f«      r·‰	j                  «        t        |‰
«      rd|j"                  ›d�}n|›}|j                  |«      x}rKt        |t$        «      rd|j&                  › d	|j(                  › �}nd|›�}‰	j                  |› |› d
|› �«       n‰	j                  |› d
|› �«       ‰j                  |«       �Œr||v s�Œx|dk7  s�Œt+        |d«      s�Œ�|j&                  j-                  d«      s�Œª‰	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     rE   r�   z^user_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse.<locals>.<genexpr>4  s(   è ø€ ò '
àØ�{‰{˜mÒ+ð �K�Kñ'
ùs   ‚%'Ú__annotations__z@triton.jitTrÊ   ztl.constexpr(rs   ú: ú.ú = ÚtlÚ
__module__r—   zfrom z import z as )r   ÚdisÚBytecodeÚfnÚ__globals__rk   Ú__code__Úco_namesrƒ   Únewliner�   ÚspliceÚsrcr½   Úintr¶   ÚboolÚvaluer‘   rÙ   Ú__name__ÚhasattrÚ
startswith)Ú
cur_kernelÚunqualified_loadsÚglobal_annotationsÚsymbol_nameÚsymbolÚ
symbol_strÚ
annotationÚannotation_coderÌ   Úcompile_wrapperrÍ   Úsymbols_includedÚtraverses           €€€€€rE   ró   zKuser_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse/  s?  ø€ ô
 'ñ '
äŸ™ Z§]¡]Ó3ô'
ó 
Ðð
 (Ÿ]™]×6Ñ6×:Ñ:Ð;LÈbÓQÐØ%Ÿ=™=×1Ñ1×:Ñ:ó -	6ˆKØÐ.Ñ.ØØ˜jŸm™m×7Ñ7Ò7Ø#Ÿ™×2Ñ2°;Ñ?�Ü˜f kÔ2Ø#×+Ñ+Ô-Ø#×-Ñ-¨mÔ<Ø#×*Ñ*¨6¯:©:¸TÐ*ÔBØ$×(Ñ(¨Ô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Ö>ð
 $×-Ñ-Ø × 1Ñ 1Ð2°(¸6¿?¹?Ð:KÈ4ÐP[È}Ð]ôð %×(Ñ(¨Ö5ñ[-	6rG   )
r0   rá   râ   r—   rÌ   Útriton.languagerÍ   r   ræ   r¾   )ry   rÌ   rñ   rÍ   rò   ró   s    @@@@@rE   Ú9user_defined_triton_kernel_transitive_closure_source_coderõ      sd   ü€ ô
 %Ó&€OØ×Ñ˜6Ÿ:™:¨TÐÔ2õ #Ý)ô " 6§?¡?Ð"3Ó4Ð÷86ð 86ñt ˆVÔØ×#Ñ#Ó%Ð%rG   c                  ó(   — e Zd ZU ded<   ded<   d„ Zy)ÚSymbolicCallArgr¶   Úinnerr´   Ú
inner_exprc                ó,   — t        | j                  «      S r<   )r¶   rø   ©Úselfs    rE   Ú__str__zSymbolicCallArg.__str__s  s   € Ü�4—:‘:‹ÐrG   N)ræ   rÙ   Ú__qualname__rÔ   rý   rŠ   rG   rE   r÷   r÷   m  s   … àƒJàÓórG   r÷   c                  ó6   ‡ — e Zd Zˆ fd„Zdd„Zdd„Zdd„Zˆ xZS )ÚMemoryPlanningStatec                ól   •— t         ‰| �  «        t        j                  t        «      | _        d| _        y ©Nr   )ÚsuperÚ__init__ÚcollectionsÚdefaultdictÚlistÚ
reuse_poolÚtotal_allocated_buffer_size)rü   Ú	__class__s    €rE   r  zMemoryPlanningState.__init__x  s-   ø€ Ü‰ÑÔä×#Ñ#¤DÓ)ð 	Œð 12ˆÕ(rG   c                óL   — t        | j                  j                  |d «      «      S r<   )rä   r  rk   )rü   r®   s     rE   Ú__contains__z MemoryPlanningState.__contains__  s   € Ü�D—O‘O×'Ñ'¨¨TÓ2Ó3Ð3rG   c                ó\   — | j                   |   j                  «       }|j                  rJ ‚|S r<   )r  ÚpopÚ	is_reused©rü   r®   r†   s      rE   r  zMemoryPlanningState.pop‚  s+   € Ø�‰˜sÑ#×'Ñ'Ó)ˆØ—>’>Ð!Ð!ØˆrG   c                ó\   — |j                   rJ ‚| j                  |   j                  |«       y r<   )r  r  Úappendr  s      rE   ÚpushzMemoryPlanningState.push‡  s&   € Ø—>’>Ð!Ð!Ø�‰˜Ñ×#Ñ# DÕ)rG   )r®   ÚReuseKeyr³   rä   )r®   r  r³   ÚFreeIfNotReusedLine)r®   r  r†   r  r³   ÚNone)ræ   rÙ   rþ   r  r  r  r  Ú__classcell__©r
  s   @rE   r   r   w  s   ø„ ô2ó4ó÷
*rG   r   c                  ó   — e Zd Zy)ÚWrapperLineN©ræ   rÙ   rþ   rŠ   rG   rE   r  r  Œ  ó   „ ØrG   r  c                  ó2   — e Zd ZU ded<   ded<   dd„Zd	d„Zy)
ÚEnterSubgraphLineÚPythonWrapperCodegenr’   r:   r=   c                ób   — | j                   j                  | j                   j                  «       y r<   )r’   Úpush_computed_sizesÚcomputed_sizesrû   s    rE   Ú__post_init__zEnterSubgraphLine.__post_init__•  s   € Ø�‰×(Ñ(¨¯©×)DÑ)DÕErG   c                ón   — | j                   j                  | j                  «       |j                  «        y r<   )r’   Úpush_codegened_graphr=   Ú	do_indent©rü   Úcodes     rE   ÚcodegenzEnterSubgraphLine.codegen˜  s"   € Ø�‰×)Ñ)¨$¯*©*Ô5Ø�‰ÕrG   N©r³   r  ©r(  r0   r³   r  ©ræ   rÙ   rþ   rÔ   r#  r)  rŠ   rG   rE   r  r  �  s   … à!Ó!ØÓóFôrG   r  c                  ó(   — e Zd ZU ded<   dd„Zdd„Zy)ÚExitSubgraphLiner  r’   c                óV   — | j                   j                  «       | j                   _        y r<   )r’   Úpop_computed_sizesr"  rû   s    rE   r#  zExitSubgraphLine.__post_init__¡  s   € Ø&*§l¡l×&EÑ&EÓ&Gˆ�‰Õ#rG   c                óX   — | j                   j                  «        |j                  «        y r<   )r’   Úpop_codegened_graphÚdo_unindentr'  s     rE   r)  zExitSubgraphLine.codegen¤  s   € Ø�‰×(Ñ(Ô*Ø×ÑÕrG   Nr*  r+  r,  rŠ   rG   rE   r.  r.  �  s   … à!Ó!óHôrG   r.  c                  ó*   — e Zd ZU ded<   ded<   dd„Zy)ÚEnterDeviceContextManagerLinerã   Ú
device_idxzOptional[int]Ú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+   r=   Úcpp_wrapperr�   Úaot_moder7  Ú
device_opsÚcpp_aoti_stream_guardr6  Úcpp_aoti_device_guardÚdevice_guardr&  Ú
set_devicer'  s     rE   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ÕJrG   Nr+  )ræ   rÙ   rþ   rÔ   r)  rŠ   rG   rE   r5  r5  ©  s   … àƒOØ"/Ó/ôKrG   r5  c                  ó   — e Zd Zdd„Zy)ÚExitDeviceContextManagerLinec                óZ   — t         j                  j                  s|j                  «        y y r<   )r+   r=   r<  r3  r'  s     rE   r)  z$ExitDeviceContextManagerLine.codegenÍ  s    € Ü�w‰w×"Ò"Ø×ÑÕð #rG   Nr+  ©ræ   rÙ   rþ   r)  rŠ   rG   rE   rD  rD  Ì  s   „ ôrG   rD  c                  ó0   — e Zd ZU ded<   dd„Zdd„Zd	d„Zy)
ÚMemoryPlanningLiner  r’   c                ó   — | S )zFirst pass to find reuserŠ   ©rü   Ústates     rE   ÚplanzMemoryPlanningLine.planÖ  s   € àˆrG   c                 ó   — y)zSecond pass to output codeNrŠ   r'  s     rE   r)  zMemoryPlanningLine.codegenÚ  s   � rG   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’   ú=rr   rp   rs   )ÚdataclassesÚfieldsrx   Úgetattrr  r‘   r   ÚBufferÚget_nameræ   rw   )rü   rz   ÚfieldrÅ   s       rE   rý   zMemoryPlanningLine.__str__Ý  s¡   € ð ˆÜ ×'Ñ'¨Ó-ò 	ˆEØ�z‰z˜YÒ&ØÜ˜$ §
¡
Ó+ˆCØ�K‰KØ—:‘:�,˜a°%·*±*ÄÇ	Á	Ñ2I §¡¤ÈsÐSÐTõð		ô �t“*×%Ñ%Ð& a¨¯	©	°$«Ð'8¸Ð:Ð:rG   N©rK  r   r³   rH  r+  ©r³   r¶   )ræ   rÙ   rþ   rÔ   rL  r)  rý   rŠ   rG   rE   rH  rH  Ò  s   … à!Ó!óó)ô;rG   rH  c                  ó(   — e Zd ZU ded<   dd„Zdd„Zy)ÚAllocateLineÚ
BufferLikerC   c           	     óŽ  — | j                   j                  «       t        j                  j                  v rt        | j                  «      S t        | j                   «      }t        j                  rG||v rC|j                  |«      }d|_        t        | j                  |j                   | j                   «      S | j                   j                  «       j                  dk(  rh| j                  j                  | j                   «      }|�A|xj                   t#        t%        j&                  t(        j*                  |d«      «      z  c_        | S )NTÚcpur,   )rC   rT  r+   r=   Úremoved_buffersÚNullLiner’   rF   r   Úallow_buffer_reuser  r  Ú	ReuseLiner?   r‘   Ústatic_shape_for_buffer_or_noner	  rã   Ú	functoolsÚreduceÚoperatorÚmul)rü   rK  r®   Ú	free_lineÚstatic_shapes        rE   rL  zAllocateLine.planð  sí   € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ñ:Ü˜DŸL™LÓ)Ð)ô ˜tŸy™yÓ)ˆÜ×$Ò$¨°©ØŸ	™	 #›ˆIØ"&ˆIÔÜ˜TŸ\™\¨9¯>©>¸4¿9¹9ÓEÐEà�9‰9×(Ñ(Ó*×/Ñ/°5Ò8ØŸ<™<×GÑGÈÏ	É	ÓRˆLØÐ'Ø×1Ò1´SÜ×$Ñ$¤X§\¡\°<ÀÓCó6ñ Õ1ð ˆrG   c                óÜ   — | j                   j                  «       t        j                  j                  vsJ ‚| j
                  j                  | j                   «      }|j                  |«       y r<   )rC   rT  r+   r=   r]  r’   Úmake_buffer_allocationr�   ©rü   r(  r    s      rE   r)  zAllocateLine.codegen  sK   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�|‰|×2Ñ2°4·9±9Ó=ˆØ�‰�tÕrG   NrV  r+  )ræ   rÙ   rþ   rÔ   rL  r)  rŠ   rG   rE   rY  rY  ì  s   … à
Óóô(rG   rY  c                  ó6   — e Zd ZU ded<   dZded<   d	d„Zd
d„Zy)r  rZ  rC   Frä   r  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!                  t#        | j                  «      | «       | S r  )r_   rC   Úget_inputs_that_alias_outputrƒ   Úlayoutr   ÚMultiOutputLayoutr  rT  r+   r=   r]  r^  r’   r   r_  r  rF   rJ  s     rE   rL  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Ó)Ð)Ü×$Ò$Ø�J‰JÔ'¨¯	©	Ó2°DÔ9ØˆrG   c                óò   — | j                   j                  «       t        j                  j                  vsJ ‚| j
                  s5|j                  | j                  j                  | j                   «      «       y y r<   )	rC   rT  r+   r=   r]  r  r�   r’   Úmake_buffer_freer'  s     rE   r)  zFreeIfNotReusedLine.codegen  sR   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�~Š~Ø�N‰N˜4Ÿ<™<×8Ñ8¸¿¹ÓCÕDð rG   NrV  r+  )ræ   rÙ   rþ   rÔ   r  rL  r)  rŠ   rG   rE   r  r  
  s   … à
ÓØ€IˆtÓó
ôErG   r  c                  ó@   — e Zd ZU ded<   ded<   dZded<   d
d„Zdd„Zy	)r`  rZ  rC   Ú	reused_asTrä   Ú
delete_oldc                ó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 r<   )rC   rT  r+   r=   r]  rs  r^  r’   rJ  s     rE   rL  zReuseLine.plan'  s   € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ñ:Ø—>‘>×*Ñ*Ó,´·±×0GÑ0GÑGÐGÐGÜ˜DŸL™LÓ)Ð)Ø�~‰~×&Ñ&Ó(´·±×0GÑ0GÑGÐGÐGØˆrG   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 r<   )
rC   rT  r+   r=   r]  rs  r�   r’   Úmake_buffer_reusert  r'  s     rE   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õ	
rG   NrV  r+  )ræ   rÙ   rþ   rÔ   rt  rL  r)  rŠ   rG   rE   r`  r`  !  s"   … à
ÓØÓØ€J�Óóô
rG   r`  c                  ó   — e Zd Zy)r^  Nr  rŠ   rG   rE   r^  r^  6  r  rG   r^  c                  óX   — e Zd ZU ded<   ded<   ed	d„«       Zed
d„«       Zedd„«       Zy)ÚCommBufferLiner  r’   ú	ir.BufferrC   c                óæ   — ddl m} | j                  j                  «       }| j                  j	                  «       } ||«      rt        d| j                  › �«      ‚t        |«      |j                  z  S )Nr   )Úis_symbolicz-The size of a comm buffer can't be symbolic: )Útorch._inductor.utilsr}  rC   Ú	get_numelr@   r`   rã   Úitemsize)rü   r}  Únumelr   s       rE   ÚsizezCommBufferLine.size?  sa   € å5à—	‘	×#Ñ#Ó%ˆØ—	‘	×#Ñ#Ó%ˆÙ�uÔÜ Ø?ÀÇ	Á	¸{ÐKóð ô �5‹z˜EŸN™NÑ*Ð*rG   c                ó†   — | j                   j                  «       }t        |t        j                  «      sJ ‚|j
                  S r<   )rC   Úget_output_specrƒ   r   ÚCommBufferLayoutÚcomm_buffer_type©rü   rn  s     rE   r†  zCommBufferLine.comm_buffer_typeK  s6   € à—‘×*Ñ*Ó,ˆÜ˜&¤"×"5Ñ"5Ô6Ð6Ð6Ø×&Ñ&Ð&rG   c                ó†   — | j                   j                  «       }t        |t        j                  «      sJ ‚|j
                  S r<   )rC   r„  rƒ   r   r…  Ú
group_namer‡  s     rE   r‰  zCommBufferLine.group_nameQ  s6   € à—‘×*Ñ*Ó,ˆÜ˜&¤"×"5Ñ"5Ô6Ð6Ð6Ø× Ñ Ð rG   N©r³   rã   )r³   zir.CommBufferTyperW  )ræ   rÙ   rþ   rÔ   Úpropertyr‚  r†  r‰  rŠ   rG   rE   rz  rz  :  sG   … à!Ó!Ø
ƒOàò	+ó ð	+ð ò'ó ð'ð
 ò!ó ñ!rG   rz  c                  ó$   — e Zd Zdd„Zed„ «       Zy)ÚCommBufferAllocateLinec                ó"  — | j                   j                  «       t        j                  j                  vsJ ‚| j                   j                  «       }| j                   j                  «       }| j                   j                  «       }t        | j                   j                  «       «      }t        | j                   j                  «       «      }|j                  | j                  | j                  | j                  | j                  |||||«      «       y r<   )rC   rT  r+   r=   r]  Ú
get_devicer@   r•   Úget_sizeÚ
get_strider�   Úmake_allocation_liner†  r‰  r’   )rü   r(  rx   Údevicer   ÚshapeÚstrides          rE   r)  zCommBufferAllocateLine.codegenZ  sË   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÑBÐBÐBØ�y‰y×!Ñ!Ó#ˆØ—‘×%Ñ%Ó'ˆØ—	‘	×#Ñ#Ó%ˆÜ�d—i‘i×(Ñ(Ó*Ó+ˆÜ�t—y‘y×+Ñ+Ó-Ó.ˆØ�‰Ø×%Ñ%Ø×%Ñ%Ø—‘Ø—‘ØØØØØó	õ	
rG   c                óþ   — | t         j                  j                  k(  rS|› d|j                  |«      › d|j                  |«      › d|› d|j                  › d|› dt        j                  dd«      › d�S t        d	| › �«      ‚)
Nz = empty_strided_p2p(rp   z, torch.device("cuda:z"), group_name="z", alloc_id=r   l   ÿÿÿÿ rs   zUnsupported comm buffer type: )r   ÚCommBufferTypeÚSYMM_MEMÚcodegen_shape_tupleÚindexÚrandomÚrandintÚNotImplementedError)r†  r‰  r’   rx   r“  r   r”  r•  s           rE   r’  z+CommBufferAllocateLine.make_allocation_linen  s    € ð œr×0Ñ0×9Ñ9Ò9à�&Ð-Ø×.Ñ.¨uÓ5Ð6°bØ×.Ñ.¨vÓ6Ð7°rØ�'ð &Ø&,§l¡l ^ð 4Ø)˜lð +Ü"ŸN™N¨1¨iÓ8Ð9¸ð<ðô &Ø0Ð1AÐ0BÐCóð rG   Nr+  )ræ   rÙ   rþ   r)  Ústaticmethodr’  rŠ   rG   rE   r�  r�  X  s   „ ó
ð( ñó ñrG   r�  c                  ó   — e Zd Zdd„Zy)ÚCommBufferFreeLinec                ó¤   — | j                   j                  | j                  «      }|j                  |› d| j                  j
                  › d�«       y )Nz # z buffer free)r’   rq  rC   r�   r†  rå   rj  s      rE   r)  zCommBufferFreeLine.codegen„  s@   € Ø�|‰|×,Ñ,¨T¯Y©YÓ7ˆØ�‰˜$˜˜s 4×#8Ñ#8×#>Ñ#>Ð"?¸|ÐLÕMrG   Nr+  rF  rŠ   rG   rE   r   r   ‚  s   „ ôNrG   r   c                  ój  ‡ — e Zd ZdZˆ fd„Ze	 dz	 	 	 	 	 	 	 d{d„«       Zd|d„Zd}d„Zd|d„Z	d~d„Z
d|d	„Ze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dzd…d„Zd„ Z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&„ Z)d'„ Z*	 	 	 	 	 	 	 	 	 	 	 	 dŠd(„Z+d‹d)„Z,d*„ Z-d+„ Z.d,„ Z/	 	 	 dŒ	 	 	 	 	 	 	 	 	 d�d-„Z0d.„ Z1d„d/„Z2d0„ Z3d1„ Z4d2„ Z5d3„ Z6	 	 	 	 	 	 dŽd4„Z7d5„ Z8d�d6„Z9d7„ Z:d8d9œd�d:„Z;d8d9œd�d;„Z<d‘d<„Z=d’d=„Z>d“d>„Z?d“d?„Z@d†d@„ZA	 dz	 	 	 d”dA„ZBd•dB„ZCdC„ ZDdD„ ZEdE„ ZFdF„ ZG	 	 	 d–	 	 	 	 	 	 	 	 	 d—dG„ZHd˜dH„ZI	 	 d™dI„ZJdzdšdJ„ZKd›dK„ZLd›dL„ZMdM„ ZNdN„ ZOdO„ ZPdP„ ZQdQ„ ZRdR„ ZSdS„ ZTdT„ ZUdœdU„ZVdV„ ZWdd8ddddWœ	 d�dX„ZXdY„ ZYdZ„ ZZd[„ Z[dzd\„Z\džd]„Z]	 dzd^„Z^dŸd_„Z_d d`„Z`d¡da„Zad¢db„Zbd£dc„Zcd¤dd„Zdd¥de„Zedf„ Zfdzdg„Zgdh„ Zhd¦di„Zidj„ Zj	 	 	 	 	 	 	 	 d§dk„Zkdl„ Zldm„ Zm	 	 	 	 d¨dn„Znd©do„Zodp„ Zpdq„ Zqdr„ Zrds„ Zsdt„ Ztedu„ «       Zuedv„ «       Zvedw„ «       Zwedx„ «       Zxedy„ «       Zyˆ xZzS )ªr  zB
    Generate outer wrapper in Python that calls the kernels.
    c                ó†  •‡ — t         ‰‰ �  «        t        «       ‰ _        t	        «       ‰ _        t	        «       ‰ _        t	        «       ‰ _        t	        «       ‰ _        t	        «       ‰ _	        t	        «       ‰ _
        t	        «       ‰ _        t	        «       ‰ _        t	        «       ‰ _        t        t           «       ‰ _        i ‰ _        t        «       ‰ _        g ‰ _        d‰ _        d‰ _        d‰ _        d‰ _        d‰ _        t2        j4                  j6                  rdnd‰ _        t2        j4                  j6                  rdnd‰ _        d ‰ _        d‰ _        i ‰ _         t        t           «       ‰ _!        t        «       ‰ _"        d ‰ _#        ‰ jI                  «        g ‰ _%        g ‰ _&        ‰ jO                  «        ‰ jQ                  «        ‰ jS                  «        t2        j4                  jT                  sBt2        j4                  jV                  jY                  «       D ]  \  }}‰ j[                  ||«       Œ t        t\           «       ‰ _/        t        t\           «       ‰ _0        i ‰ _1         te        jf                  d «      ‰ jh                  «      ‰ _4        te        jf                  d «      d	ˆ fd„«       }|‰ _5        i ‰ _6        t        t           «       ‰ _7        tq        «       ‰ _9        t        t           «       ‰ _:        i ‰ _;        ty        tz        j|                  j~                  tz        j|                  j€                  ¬«      ‰ _A        g ‰ _B        y )
NÚ ú#r  z
std::move(rs   Tc                ó¨   •— ‰j                   j                  | «       t        j                  j                  r‰j
                  j                  | «       y y r<   )Úimportsr�   r   r—   r˜   rŸ   )r    rü   s    €rE   Úadd_import_oncez6PythonWrapperCodegen.__init__.<locals>.add_import_onceÍ  s;   ø€ à�L‰L×"Ñ" 4Ô(Ü�}‰}×5Ò5Ø×*Ñ*×4Ñ4°TÕ:ð 6rG   )Údebug_printer_levelÚuse_array_ref)r    r¶   r³   r  )Cr  r  r   Ú_names_iterr0   r§  ÚheaderÚprefixÚsuffixÚkernel_declarationsÚwrapper_callÚkernel_autotune_defsrŸ   Úsubgraph_definitionsr   r¶   rž   Úsrc_to_kernelÚkernel_numel_exprÚlinesÚdeclareÚdeclare_maybe_referenceÚendingÚcommentÚnone_strr+   r=   r<  Ú
move_beginÚmove_endr7  Úsupports_intermediate_hooksÚuser_defined_kernel_cacheÚunbacked_symbol_declsr"  Úlauncher_fn_nameÚset_launcher_fn_nameÚcodegened_graph_stackÚcomputed_sizes_stackÚwrite_headerÚwrite_prefixÚ!write_kernel_autotune_defs_headerr=  Úconstant_reprsr\   Úwrite_constantÚ
BufferNameÚ	allocatedÚfreedÚreusesrb  Ú	lru_cacheÚwrite_get_raw_streamr¨  Ú_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)rü   rx   Úhashedr¨  r
  s   `   €rE   r  zPythonWrapperCodegen.__init__‘  s�  ù€ Ü‰ÑÔÜ*/«'ˆÔÜ%Ó'ˆŒÜ$Ó&ˆŒÜ$Ó&ˆŒÜ$Ó&ˆŒÜ#1Ó#3ˆÔ Ü*Ó,ˆÔÜ$2Ó$4ˆÔ!Ü%3Ó%5ˆÔ"Ü$2Ó$4ˆÔ!Ü%/´¡_Ó%6ˆÔ"ð .0ˆÔÜHRËˆÔØCEˆŒ
ØˆŒØ')ˆÔ$ØˆŒØˆŒØˆŒÜ*+¯'©'×*=Ò*=™,À2ˆŒÜ Ÿw™w×2Ò2™¸ˆŒØ;?ˆÔ)Ø+/ˆÔ(ØQSˆÔ&Ü%/´¡_Ó%6ˆÔ"Ü8B»ˆÔØ $ˆÔà×!Ñ!Ô#ð &(ˆÔ"Ø$&ˆÔ!à×ÑÔØ×ÑÔØ×.Ñ.Ô0ä�w‰w×ÒÜ !§¡× 6Ñ 6× <Ñ <Ó >ò 2‘��fà×#Ñ# D¨&Õ1ð2ô $¤JÑ/Ó1ˆŒÜ¤
Ñ+Ó-ˆŒ
ð 57ˆŒà$=¤I×$7Ñ$7¸Ó$=Ø×%Ñ%ó%
ˆÔ!ô 
×	Ñ	˜TÓ	"ô	;ó 
#ð	;ð
  /ˆÔØ&(ˆŒÜ$¤S™/Ó+ˆŒÜ"2Ó"4ˆÔÜ+5´c©?Ó+<ˆÔ(Ø46ˆÔ!ô 1Ü &× 3Ñ 3× TÑ TÜ ×-Ñ-×DÑDô
ˆÔð !#ˆÕrG   Nc                óD   — | r|€J ‚|€J ‚t        |||«      S t        «       S r<   )ÚSubgraphPythonWrapperCodegenr  )Úis_subgraphÚsubgraph_nameÚparent_wrapperÚpartition_signaturess       rE   ÚcreatezPythonWrapperCodegen.createã  s?   € ñ Ø Ð,Ð,Ð,Ø!Ð-Ð-Ð-Ü/Ø˜~Ð/Cóð ô $Ó%Ð%rG   c                ó   — d| _         y )NÚcall)rÀ  rû   s    rE   rÁ  z)PythonWrapperCodegen.set_launcher_fn_nameò  s
   € Ø &ˆÕrG   c                óD   — | j                   j                  |› d|› �«       y )Nz = None  # )r¬  r�   )rü   rx   rÙ  s      rE   rÈ  z#PythonWrapperCodegen.write_constantõ  s   € Ø�‰×Ñ   k°&°Ð:Õ;rG   c           	     óT  — t         j                  j                  j                  «       }d}|�|j                  �d|j                  › �}d}t        t        j                  j                  «      dkD  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«       y y # t         t"        f$ r Y Œ>w xY w)Nr¤  z
# AOT ID: r   zRfrom torch._inductor.codegen.debug_utils import _print_debugging_tensor_value_infoz
                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 zÂ import AsyncCompile
                from torch._inductor.select_algorithm import extern_kernels
                from torch._inductor.codegen.multi_kernel import MultiKernelCall
                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
                empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
                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 nvtx)ÚtorchÚ_guardsÚTracingContextÚtry_getÚaot_graph_namerã   r   rÔ  rÕ  r§  rá   r   ræ   r¬  Útorch._C._distributed_c10drå  ÚAttributeErrorÚImportErrorÚannotate_trainingr�   )rü   ÚcontextÚaot_config_commentÚaot_inductor_debug_utilsrå  s        rE   rÄ  z!PythonWrapperCodegen.write_headerø  sF  € Ü—-‘-×.Ñ.×6Ñ6Ó8ˆØÐØÐ 7×#9Ñ#9Ð#EØ#-¨g×.DÑ.DÐ-EÐ!FÐØ#%Ð ÜŒv×"Ñ"×CÑCÓDÀqÒHØ'{Ð$Ø�‰×ÑðØ#Ð$ð %ô $×,Ñ,Ð-ð .ð *Ð*ð +ð#ð& ð) 	ô 	
ð, 	�‰×Ñðð ð 	ô 	
ð	õ Dà�K‰K×Ñðð ð	 ô ô ×#Ò#Ø�K‰K×!Ñ!Ð"?Õ@ð $øô ¤Ð,ò 	Ùð	ús   Ã#D ÄD'Ä&D'c                 ó   — y r<   rŠ   )rü   r¬  s     rE   Úinclude_extra_headerz)PythonWrapperCodegen.include_extra_header5  ó   € ØrG   c                ó^   — | j                   j                  dt        j                  › d�«       y )Na	  
                import torch
                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±  rá   r   ræ   rû   s    rE   rÆ  z6PythonWrapperCodegen.write_kernel_autotune_defs_header8  s3   € Ø×!Ñ!×(Ñ(ðô
 $×,Ñ,Ð-ð .ðõ	
rG   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+   r=   r>  Úimport_get_raw_stream_asr<  r§  )rü   Ú
import_strs     rE   Úwrite_triton_header_oncez-PythonWrapperCodegen.write_triton_header_onceH  sÇ   € ðô $×,Ñ,Ð-ð .ðˆ
ô
 �=‰=×1Ò1Ø×&Ñ&×-Ñ-¨jÔ9Ø×&Ñ&×0Ñ0Ü—‘×"Ñ"×;Ñ;Ð<LÓMôô �w‰w×"Ò"Ø�L‰L×Ñ 
°$ÐÔ7Ø�L‰L×"Ñ"Ü—‘×"Ñ"×;Ñ;Ð<LÓMõð #rG   c                óv  — t         j                  j                  rB| j                  j	                  t
        j                  j                  j                  d«      «       t
        j                  j                  sC| j                  j	                  t
        j                  j                  j                  d«      «       y y )Nr÷  )r   r—   r˜   rŸ   r�   r+   r=   r>  rø  r<  r§  rû   s    rE   Ú write_get_raw_stream_header_oncez5PythonWrapperCodegen.write_get_raw_stream_header_onceZ  s{   € ä�=‰=×1Ò1Ø×&Ñ&×0Ñ0Ü—‘×"Ñ"×;Ñ;Ð<LÓMôô �w‰w×"Ò"Ø�L‰L×"Ñ"Ü—‘×"Ñ"×;Ñ;Ð<LÓMõð #rG   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×   )rX   rÏ  r_   r¬  r�   r   r—   r˜   rŸ   rÐ  r½   )rü   rþ  Úvars      rE   Úadd_meta_oncez"PythonWrapperCodegen.add_meta_oncee  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Ñ Ð rG   c                óz   — | j                  «       D �cg c]  }|j                  | j                  «      ‘Œ c}S c c}w r<   )Úget_graph_outputsÚcodegen_referencer°  ©rü   rª   s     rE   Úget_output_refsz$PythonWrapperCodegen.get_output_refsp  s<   € ð =A×<RÑ<RÓ<Tö
Ø78ˆA×Ñ × 1Ñ 1Õ2ò
ð 	
ùò 
s   “"8c                 ó   — y r<   rŠ   rû   s    rE   Úmark_output_typez%PythonWrapperCodegen.mark_output_typev  ó   € ØrG   c                ó6   — t         j                  j                  S r<   )r+   r=   Úgraph_inputsrû   s    rE   Úget_graph_inputsz%PythonWrapperCodegen.get_graph_inputsy  s   € ô �w‰w×#Ñ#Ð#rG   c                ó6   — t         j                  j                  S r<   )r+   r=   Úgraph_outputsrû   s    rE   r  z&PythonWrapperCodegen.get_graph_outputs~  s   € Ü�w‰w×$Ñ$Ð$rG   c           
     ó  — | j                  «       j                  «       D ]ç  \  }}t        |t        j                  t
        j                  f«      rŒ1|t        j                  j                  vst        |t
        j                  «      rŒht        |j                  «       «      dk(  rŒ…| j                  |j                  «       «      }| j                  |j                  «       «      }| j                  j!                  d|› d|› d|› d�«       Œé y )Nr   zassert_size_stride(rp   rs   )r  r\   rƒ   r„   r
   r   ÚTorchBindObjectr+   r=   Úgraph_input_namesÚGeneratorStater'   r�  r–   r‘  r­  r�   )rü   rx   Úbufr‚  r•  s        rE   Úcodegen_input_size_assertsz/PythonWrapperCodegen.codegen_input_size_asserts�  sÝ   € Ø×.Ñ.Ó0×6Ñ6Ó8ò 	S‰IˆD�#Ü˜#¤§
¡
¬B×,>Ñ,>Ð?Ô@Øð œ1Ÿ7™7×4Ñ4Ñ4¼
Ø”R×&Ñ&ô9ð ô ˜SŸ\™\›^Ó,°Ò1ØØ×2Ñ2°3·<±<³>ÓBˆDØ×4Ñ4°S·^±^Ó5EÓFˆFØ�K‰K×!Ñ!Ð$7¸°v¸RÀ¸vÀRÈÀxÈqÐ"QÕRñ	SrG   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�   r  r\   rƒ   r„   r
   r   r  )rü   rx   r  r    s       rE   Úcodegen_input_nan_assertsz.PythonWrapperCodegen.codegen_input_nan_asserts“  s˜   € Ø�‰×ÑÐHÔIØ×.Ñ.Ó0×6Ñ6Ó8ò 	(‰IˆD�#Ü˜#¤§
¡
¬B×,>Ñ,>Ð?Ô@Øà   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Ô'Ø   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Õ'ñ	(rG   c                ó:   — | j                   j                  d«       y )NzV

            async_compile.wait(globals())
            del async_compile
            )r­  rá   rû   s    rE   Úwrite_async_compile_waitz-PythonWrapperCodegen.write_async_compile_waitž  s   € Ø�‰×Ñðõ	
rG   c                ó¾   — dj                  |«      }t        |«      dk(  r|dz  }| j                  j                  |› d�«       | j                  j                  d«       y )Nrp   r,   ú,z = argszargs.clear())rw   r_   r­  r�   )rü   Úinput_namesÚlhss      rE   Ú
write_argszPythonWrapperCodegen.write_args§  sP   € Ø�i‰i˜Ó$ˆÜˆ{Ó˜qÒ Ø�3‰JˆCØ�‰×Ñ   W˜oÔ.Ø�‰×Ñ˜nÕ-rG   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):
                r,   )r   Úgraph_partitionr­  rá   rÀ  ©rü   Úprefix_indents     rE   Ú!write_launcher_fn_call_get_indentz6PythonWrapperCodegen.write_launcher_fn_call_get_indent®  sm   € Ü×!Ò!Ø�K‰K×Ñðôð ˆMð Ðð �K‰K×ÑðØ×*Ñ*Ð+ð ,ðôð
 ˆMàÐrG   c                ó6   — t         j                  j                  S r<   )r+   r=   r  rû   s    rE   Úget_graph_input_namesz*PythonWrapperCodegen.get_graph_input_namesÊ  s   € Ü�w‰w×(Ñ(Ð(rG   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#                  «        | j%                  «        d d d «       y # 1 sw Y   y xY w)Nz0training_annotation = nvtx._device_range_start('ú'))rÀ  r  r!  r­  r¸   r   r—   Údebug_sync_graphr�   r+   r=   r>  ÚsynchronizeÚget_training_phaserî  r#  r  Úcodegen_inputsÚ"codegen_input_size_and_nan_asserts)rü   r   Úphaser  s       rE   rÅ  z!PythonWrapperCodegen.write_prefixÍ  sü   € Ø×$Ñ$Ð0Ð0Ð0Ø×%Ñ%Ô'Ø×>Ñ>Ó@ˆà�[‰[×Ñ Ó.ñ 	6Ü�}‰}×-Ò-Ø—‘×%Ñ%¤a§g¡g×&8Ñ&8×&DÑ&DÓ&FÔGÜ—G‘G×.Ñ.Ó0ˆEÜ×'Ò'Ø—‘×%Ñ%ØFÀuÀgÈRÐPôð %)×$>Ñ$>Ó$@Ð@Ð Ð@Ø—‘Ð 1Ô2à×ÑÔ!Ø×3Ñ3Ô5÷	6÷ 	6ñ 	6ús   Á
C,D?Ä?Ec                ó†   — t         j                  r| j                  «        t         j                  r| j	                  «        y y r<   )r   Úsize_assertsr  Únan_assertsr  rû   s    rE   r*  z7PythonWrapperCodegen.codegen_input_size_and_nan_assertsá  s1   € Ü×ÒØ×+Ñ+Ô-Ü×ÒØ×*Ñ*Õ,ð rG   c                ó  — | j                  «        d|› �}t        j                  j                  r=| j                  j                  |› d|› d�«       t        j                  j                  r|S | j                  |› d|› d�«       |S )NÚstreamú = get_raw_stream(rs   )	rü  r   r—   r˜   rŸ   r�   r+   r=   r<  )rü   r6  r=   rx   s       rE   rÎ  z)PythonWrapperCodegen.write_get_raw_streamê  s   € Ø×-Ñ-Ô/Ø˜
�|Ð$ˆÜ�=‰=×1Ò1Ø×&Ñ&×0Ñ0Ø�&Ð*¨:¨,°aÐ8ôô �w‰w×"Ò"à�Ø�‰˜$˜Ð1°*°¸QÐ?Ô@ØˆrG   c                ó    — | j                   d   S )Néÿÿÿÿ)rÂ  rû   s    rE   Úget_codegened_graphz(PythonWrapperCodegen.get_codegened_graph÷  s   € Ø×)Ñ)¨"Ñ-Ð-rG   c                ó:   — | j                   j                  |«       y r<   )rÂ  r  )rü   r=   s     rE   r%  z)PythonWrapperCodegen.push_codegened_graphú  s   € Ø×"Ñ"×)Ñ)¨%Õ0rG   c                ó6   — | j                   j                  «       S r<   )rÂ  r  rû   s    rE   r2  z(PythonWrapperCodegen.pop_codegened_graphý  s   € Ø×)Ñ)×-Ñ-Ó/Ð/rG   c                óP   — ddl m} | j                  j                   ||«      «      S )Nr   )Údeepcopy)Úcopyr8  rÃ  r  )rü   r"  r8  s      rE   r!  z(PythonWrapperCodegen.push_computed_sizes   s!   € Ý!à×(Ñ(×/Ñ/±¸Ó0HÓIÐIrG   c                ó6   — | j                   j                  «       S r<   )rÃ  r  rû   s    rE   r0  z'PythonWrapperCodegen.pop_computed_sizes  s   € Ø×(Ñ(×,Ñ,Ó.Ð.rG   c                ó.   — t        | j                  «      › S r<   )Únextr«  rû   s    rE   Únext_kernel_suffixz'PythonWrapperCodegen.next_kernel_suffix  s   € Ü�t×'Ñ'Ó(Ð)Ð*rG   c                ó8  — | j                  t        || j                  «      «       t        j                  j
                  rÔ| j                  «        | j                  j                  dt        j                  j                  j                  |«      › d�«       | j                  j                  «        | j                  j                  t        j                  j                  j                  |«      «       | j                  j                  d|› d|› d�«       || _        y )Nr:  r;  r0  r1  rs   )r�   r5  r7  r   r—   r˜   rú  rŸ   r+   r=   r>  rA  r&  rB  )rü   r6  s     rE   Úcodegen_device_guard_enterz/PythonWrapperCodegen.codegen_device_guard_enter  sâ   € Ø�‰Ü)¨*°d×6WÑ6WÓXô	
ô �=‰=×1Ò1à×)Ñ)Ô+Ø×&Ñ&×0Ñ0ØœŸ™×*Ñ*×7Ñ7¸
ÓCÐDÀAÐFôð ×&Ñ&×0Ñ0Ô2Ø×&Ñ&×0Ñ0Ü—‘×"Ñ"×-Ñ-¨jÓ9ôð ×&Ñ&×0Ñ0Ø˜˜Ð$6°z°lÀ!ÐDôð -7ˆÕ)rG   c                ó    — | j                  t        «       «       t        j                  j                  r| j
                  j                  «        y y r<   )r�   rD  r   r—   r˜   rŸ   r3  rû   s    rE   Úcodegen_device_guard_exitz.PythonWrapperCodegen.codegen_device_guard_exit  s6   € Ø�‰Ô3Ó5Ô6Ü�=‰=×1Ò1Ø×&Ñ&×2Ñ2Õ4ð 2rG   c                ó    — |r1| j                   j                  ddj                  |«      z   dz   «       y | j                   j                  d«       y )Nzreturn (rp   ú, )z	return ())r°  r�   rw   )rü   Úoutput_refss     rE   Úgenerate_returnz$PythonWrapperCodegen.generate_return#  s@   € ÙØ×Ñ×'Ñ'¨
°T·Y±Y¸{Ó5KÑ(KÈeÑ(SÕTà×Ñ×'Ñ'¨Õ4rG   c                 ó   — y r<   rŠ   ©rü   Úresults     rE   Úgenerate_before_suffixz+PythonWrapperCodegen.generate_before_suffix)  r  rG   c                óÀ   — t         j                  rNdj                  | j                  «      t	        | j                  «      dk(  rdndz   }|j                  d|› d�«       y y )Nrp   r,   r  r¤  z-
                runner = Runner(partitions=[z{])
                call = runner.call
                recursively_apply_fns = runner.recursively_apply_fns
                )r   r  rw   Úall_partition_namesr_   rá   )rü   rH  Úall_partition_name_lists      rE   Úgenerate_after_suffixz*PythonWrapperCodegen.generate_after_suffix,  se   € Ü×!Ò!Ø&*§i¡i°×0HÑ0HÓ&IÜ˜4×3Ñ3Ó4¸Ò9‘¸rñ'Ð#ð �M‰Mð-Ø-DÐ,Eð Fðõð "rG   c                 ó   — y r<   rŠ   rG  s     rE   Úgenerate_endz!PythonWrapperCodegen.generate_end:  r  rG   c                ó(   — | j                  ||«       y r<   )Úgenerate_extern_kernel_alloc)rü   Úfallback_kernelrz   s      rE   Úgenerate_fallback_kernelz-PythonWrapperCodegen.generate_fallback_kernel=  s   € Ø×)Ñ)¨/¸4Õ@rG   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_complexz.clone()rr   rp   rs   r×   ÚinductorÚintermediate_hooksr,   zrun_intermediate_hooks()rƒ   rn  r   Ú
NoneLayoutrT  Úget_origin_nodeÚget_kernel_namer¸  r   Úmemory_planningr�   r¶  rw   r½  Úgenerate_intermediate_hooksr   rx   )rü   Úextern_kernelrz   Ú	no_returnÚoutput_nameÚorigin_nodeÚkernel_namer¸  s           rE   rQ  z1PythonWrapperCodegen.generate_extern_kernel_alloc@  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ð 1rG   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=rr   rp   rs   )r+   r=   Úwrapper_coder×  Úset_printer_argsr  r�   rw   )rü   ry   ÚoutÚout_viewrz   r“  Údebug_printer_managers          rE   Úgenerate_extern_kernel_outz/PythonWrapperCodegen.generate_extern_kernel_out]  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                  }|r$t        d„ |D «       «      }t        d„ |D «       «      }|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 )Nc              3  ón   K  — | ]-  }t         j                  j                  j                  |«      –— Œ/ y ­wr<   ©r+   r=   rA   Úatomically_apply_size_hint©r‹   Úds     rE   r�   zEPythonWrapperCodegen._generate_tma_descriptor_call.<locals>.<genexpr>p  s%   è ø€ ÒVÈAœŸ™×)Ñ)×DÑDÀQ×GÑVùó   ‚35c              3  ón   K  — | ]-  }t         j                  j                  j                  |«      –— Œ/ y ­wr<   rl  rn  s     rE   r�   zEPythonWrapperCodegen._generate_tma_descriptor_call.<locals>.<genexpr>q  s*   è ø€ ò ØCD”—‘× Ñ ×;Ñ;¸A×>ñùrp  z.data_ptr()rp   c              3  óJ   •K  — | ]  }t         j                  ‰|«      –— Œ y ­wr<   ©r  Úval_to_arg_str©r‹   Údimrü   s     €rE   r�   zEPythonWrapperCodegen._generate_tma_descriptor_call.<locals>.<genexpr>w  s   øè ø€ ÒXÈCÔ-×<Ñ<¸TÀ3×GÑXùó   ƒ #c              3  óJ   •K  — | ]  }t         j                  ‰|«      –— Œ y ­wr<   rs  ru  s     €rE   r�   zEPythonWrapperCodegen._generate_tma_descriptor_call.<locals>.<genexpr>x  s$   øè ø€ ò 
Ø?BÔ ×/Ñ/°°c×:ñ
ùrw  z$triton.tools.experimental_descriptorz.create_Úd_tma_descriptorrr   rs   )
ÚdimsÚ
block_dimsr•   Útensorr  rw   r  rt  Úelement_sizeÚrank)rü   ÚdescÚapply_size_hintsrz  r{  Úptrr}  r­  rÜ   rz   râ  s   `          rE   Ú_generate_tma_descriptor_callz2PythonWrapperCodegen._generate_tma_descriptor_calll  sô   ø€ Ø�y‰yˆØ—_‘_ˆ
ÙÜÑVÐQUÔVÓVˆDÜñ ØHRôó ˆ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ˆˆØˆrG   c                ó‚   — | j                  |«      }|j                  › d|› | j                  › �}| j                  |«       y ©Nr×   )r‚  rx   r¸  r�   )rü   r  râ  r    s       rE   Úgenerate_tma_descriptorz,PythonWrapperCodegen.generate_tma_descriptor‚  s:   € Ø×1Ñ1°$Ó7ˆØ—)‘)�˜C ˜v d§k¡k ]Ð3ˆØ�‰�tÕrG   c                óò   — |› ddj                  t        t        |«      «      › �}|j                  d«      r|dj                  dg|z   «      z  }n|r|dt	        |«      › �z  }|dz  }| j                  |«       y )Nrr   r  zaten.scatter_reducerp   r¤  z	, reduce=rs   )rw   Úmapr¶   rè   rX   r�   )	rü   r¢   ÚinputsÚcpp_kernel_nameÚpython_kernel_nameÚsrc_is_tensorrc  r¨   r    s	            rE   Úgenerate_scatter_fallbackz.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ÕrG   c                ó|   — ddj                  |«      › d�}||||g}| j                  | j                  ||«      «       y )Nú[rp   ú])rw   r�   Úwrap_kernel_call)rü   ry   rª   ÚindicesÚvaluesÚ
accumulateÚindices_strrz   s           rE   Úgenerate_index_put_fallbackz0PythonWrapperCodegen.generate_index_put_fallbackš  sA   € Ø˜$Ÿ)™) GÓ,Ð-¨QÐ/ˆØ�; ¨
Ð3ˆØ�‰�t×,Ñ,¨V°TÓ:Õ;rG   c           	     óV   — | j                  |› d|› ddj                  |«      › d�«       y )Nr×   rr   rp   rs   )r�   rw   )rü   Úbuf_namerŠ  r‰  Úcodegen_argsÚop_overloadÚraw_argsÚoutputss           rE   Ú,generate_fallback_kernel_with_runtime_lookupzAPythonWrapperCodegen.generate_fallback_kernel_with_runtime_lookupŸ  s0   € ð 	�‰˜(˜ 3Ð'9Ð&:¸!¸D¿I¹IÀlÓ<SÐ;TÐTUÐVÕWrG   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     rE   ÚgeneratezPythonWrapperCodegen.generate«  s,   € ÜÐ9Ó:ñ 	0Ø—>‘> ,Ó/÷	0÷ 	0ò 	0ús   Œ'§0c                ó&   — t         j                  ryy)Nr   r,   )r   r  rû   s    rE   Úget_wrapper_call_indentz,PythonWrapperCodegen.get_wrapper_call_indent¯  s   € Ü×!Ò!ØàrG   c                óD	  — t         j                  r| j                  «        t        «       }|j	                  | j
                  «       |j                  d«       |j	                  | j                  «       t        j                  j                  r>t        j                  j                  r$t        j                  j                  r
t        «       }|j	                  | j                  «       t        j                  «       5 }|j!                  | j"                  j%                  «       «       t         j&                  r| j)                  |«       t         j                  r| j+                  «        |r!t         j,                  r| j/                  «        n| j1                  «        t         j2                  j4                  r*t         j2                  j6                  s| j9                  «        | j:                  D ]I  }t=        |t>        «      r|jA                  | j"                  «       Œ/| j"                  j                  |«       ŒK | jC                  «       }| jE                  «        t         j2                  jF                  rA| j"                  j                  t        j                  jH                  jK                  «       «       t         j                  r| jM                  «        t         j2                  j4                  r*t         j2                  j6                  s| jO                  «        t         j2                  j6                  r| jQ                  «        t         jR                  r+t         j                  s| j"                  j                  d«       | jU                  |«       d d d «       | jW                  «        |j	                  | jX                  «       | j[                  «       }|j%                  |«      5  |j	                  | j"                  «       d d d «       | j]                  |«       |j	                  | j^                  «       | ja                  |«       | jc                  |«       | je                  |«       |jg                  «       | jh                  jg                  «       fS # 1 sw Y   �ŒxY w# 1 sw Y   ŒŸxY w)Nr¤  z+nvtx._device_range_end(training_annotation))5r   Úprofile_bandwidthrú  r0   rá   r§  r�   r¬  r+   r=   r=  r<  Úis_const_graphr²  r¹   Ú	ExitStackÚenter_contextr°  r¸   Úprofiler_mark_wrapper_callÚ#generate_profiler_mark_wrapper_callÚgenerate_start_graphr[  Úmemory_planÚmemory_plan_reuser—   Ústore_cubinr˜   Ú!generate_reset_kernel_saved_flagsrµ  rƒ   r  r)  r  r  r&  r>  r'  Úgenerate_end_graphÚ generate_save_uncompiled_kernelsÚgenerate_and_run_autotune_blockrî  rE  Úfinalize_prefixr­  r¢  rI  r®  rM  rO  Úadd_benchmark_harnessÚgetvaluewithlinemapr¯  )rü   rŸ  rH  Ústackr    rD  Úwrapper_call_indents          rE   rž  zPythonWrapperCodegen._generateµ  s?  € Ü×#Ò#Ø×)Ñ)Ô+ÜÓ!ˆØ�‰�d—l‘lÔ#Ø×Ñ˜ÔØ�‰�d—k‘kÔ"ô �7‰7×Ò¤§¡× 3Ò 3¼¿¹×8NÒ8NÜ#Ó%ˆFð 	�‰�d×/Ñ/Ô0ä×!Ñ!Ó#ð )	. uØ×Ñ × 1Ñ 1× 8Ñ 8Ó :Ô;Ü×0Ò0Ø×8Ñ8¸Ô?Ü×'Ò'Ø×)Ñ)Ô+ñ ¤× 6Ò 6Ø× Ñ Õ"à×&Ñ&Ô(ä�}‰}×(Ò(´·±×1WÒ1WØ×6Ñ6Ô8àŸ
™
ò 6�Ü˜d¤KÔ0Ø—L‘L ×!2Ñ!2Õ3à×%Ñ%×/Ñ/°Õ5ð	6ð ×.Ñ.Ó0ˆKØ×!Ñ!Ô#Ü�}‰}×-Ò-Ø×!Ñ!×+Ñ+¬A¯G©G×,>Ñ,>×,JÑ,JÓ,LÔMä×'Ò'Ø×'Ñ'Ô)ä�}‰}×(Ò(´·±×1WÒ1WØ×5Ñ5Ô7ä�}‰}×5Ò5Ø×4Ñ4Ô6ô ×'Ò'´×0BÒ0BØ×!Ñ!×+Ñ+ØAôð × Ñ  Ô-÷S)	.ðV 	×ÑÔØ�‰�d—k‘kÔ"à"×:Ñ:Ó<Ðà�]‰]Ð.Ó/ñ 	-Ø�M‰M˜$×+Ñ+Ô,÷	-ð 	×#Ñ# FÔ+Ø�‰�d—k‘kÔ"Ø×"Ñ" 6Ô*à×Ñ˜&Ô!à×"Ñ" 6Ô*ð ×&Ñ&Ó(Ø×$Ñ$×8Ñ8Ó:ð
ð 	
÷w)	.ñ )	.ú÷`	-ð 	-ús   Ã9JR	ÏRÒ	RÒRc                ó8  — | j                   j                  d«       i }| 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        ||«       y# 1 sw Y   Œ-xY w# t         $ r}t#        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
        r9  z.pyF)Údirr®  Údeletezutf-8NzAuto-tuning code written to %sz%Failed to run autotuning code block: )r±  rá   r¾   rŸ   r   ÚlevelÚloggingÚDEBUGÚtempfileÚNamedTemporaryFiler   ÚwriteÚencoderx   ÚdebugÚexecÚ	ExceptionÚRuntimeError)rü   ÚscopeÚtuning_codeÚfÚ	file_pathÚes         rE   r±  z4PythonWrapperCodegen.generate_and_run_autotune_block  s  € ð
 	×!Ñ!×(Ñ(ðô	
ð ˆà×%Ñ%×.Ñ.Ó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$   Â-C0Ã#C< Ã0C9Ã<	DÄDÄDc                ó\   — ddl m}  || «      j                  | j                  «      | _        y )Nr,   )ÚMemoryPlanner)r[  rË  rL  rµ  )rü   rË  s     rE   r«  z PythonWrapperCodegen.memory_plan'  s    € Ý2á" 4Ó(×-Ñ-¨d¯j©jÓ9ˆ�
rG   c                óÆ  — 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 }t        t        | j                  «      «      D ]�  }| j                  |   }t	        |t
        «      r"|j                  |d   «      | j                  |<   ŒDt	        |t        «      r|j                  t        «       «       Œnt	        |t        «      sŒ|j                  |j                  «       «       ŒŸ |j                  |j                  «       «       t        |«      dk(  sJ ‚t!        d„ |D «       «      }y )Nr3  r   c              3  ó4   K  — | ]  }|j                   –— Œ y ­wr<   )r	  )r‹   Úss     rE   r�   z9PythonWrapperCodegen.memory_plan_reuse.<locals>.<genexpr>I  s   è ø€ ò +
Ø./ˆA×)Õ)ñ+
ùs   ‚)r+   r=   Úget_output_namesrµ  rƒ   rH  rC   rx   r  r   Úranger_   rL  r  r  r.  Úsum)rü   Ú	out_namesÚplanning_statesÚpast_planning_statesÚir    Ú_total_allocated_buffer_sizes          rE   r¬  z&PythonWrapperCodegen.memory_plan_reuse,  s}  € Ü—G‘G×,Ñ,Ó.ˆ	ð �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Ø×&Ñ&Ô':Ó'<Õ=Ü˜DÔ"2Õ3Ø$×+Ñ+¨O×,?Ñ,?Ó,AÕBð	Cð 	×#Ñ# O×$7Ñ$7Ó$9Ô:Ü�?Ó# qÒ(Ð(Ð(ô
 (+ñ +
Ø3Gô+
ó (
Ñ$rG   c           	     ó4  ‡	— | j                   Š	t        j                  d «      ˆ	fd„«       }t        j                  d «      ˆ	fd„«       }t        |t        j
                  «      rGt        |t        j                  «      r||v ry ‰	j                  |› d|› �«       |j                  |«       y t        |t        j                  «      råt        |j                  «       «      D ]V  \  }}t        |t        j                  «      sŒ!||vsŒ&‰	j                  |› d ||«      › d|› d�«       |j                  |«       ŒX t        |j                  «       «      D ]V  \  }}t        |t        j                  «      sŒ!||vsŒ&‰	j                  |› d ||«      › d|› d�«       |j                  |«       ŒX y t        |t        j                  «      ry t        |t        j                  «      ry t         j"                  j$                  j&                  ry t)        dt+        |«      › �«      ‚)Nc                ó<   •— ‰j                  | › d| › d�«       | › d�S )Nz_size = z.size()Ú_size©r�   ©rx   r(  s    €rE   ÚsizeofzDPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.sizeofU  s(   ø€ à�N‰N˜d˜V 8¨D¨6°Ð9Ô:Ø�V˜5�>Ð!rG   c                ó<   •— ‰j                  | › d| › d�«       | › d�S )Nz
_stride = z	.stride()Ú_striderÚ  rÛ  s    €rE   ÚstrideofzFPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.strideofZ  s)   ø€ à�N‰N˜d˜V :¨d¨V°9Ð=Ô>Ø�V˜7Ð#Ð#rG   r×   rŽ  r�  zUnknown value type: )r­  rb  rÍ  rƒ   r„   r
   ÚSymbolr�   r½   r   Ú	TensorBoxÚ	enumerater�  r‘  r  r  ræ  Ú	_inductorr   r  r`   r‘   )
rü   rx   rå   Ú
bound_varsrÜ  rß  rv  r‚  r•  r(  s
            @rE   Úcodegen_input_symbol_assignmentz4PythonWrapperCodegen.codegen_input_symbol_assignmentM  sÃ  ø€ ð �{‰{ˆä	×	Ñ	˜TÓ	"ó	"ó 
#ð	"ô 
×	Ñ	˜TÓ	"ó	$ó 
#ð	$ô �eœUŸZ™ZÔ(Ü˜e¤U§\¡\Ô2°e¸zÑ6IØØ�N‰N˜e˜W C¨ vÐ.Ô/Ø�N‰N˜5Õ!Ü˜œrŸ|™|Ô,Ü& u§~¡~Ó'7Ó8ò )‘	��TÜ˜d¤E§L¡LÕ1°dÀ*Ò6LØ—N‘N d V¨3©v°d«|¨n¸A¸c¸UÀ!Ð#DÔEØ—N‘N 4Õ(ð)ô  )¨×)9Ñ)9Ó);Ó<ò +‘��VÜ˜f¤e§l¡lÕ3¸ÀjÒ8PØ—N‘N f X¨S±¸$³Ð0@ÀÀ#ÀÀaÐ#HÔIØ—N‘N 6Õ*ñ+ô ˜œr×1Ñ1Ô2ØÜ˜œr×0Ñ0Ô1Øä�‰×%Ñ%×5Ò5Øä$Ð';¼DÀ»K¸=Ð%IÓJÐJrG   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                  |||«       Œ yc c}}w c c}}w )z$Assign all symbolic shapes to localsN)r   r„   rà  r  r\   rƒ   rå  )rü   rä  r
  ÚkÚvrˆ  rx   rå   s           rE   r)  z#PythonWrapperCodegen.codegen_inputsw  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ùó
ùãXs   ¿"CÁ"CÁ>"CÂ!Cc                óP  — t        |t        j                  «      rŒt        |t        j
                  «      rq|| j                  v ry | j                  j                  |«       t        j                  j                  j                  |   }| j                  |› dt        |«      › �«       y y y r„  )rƒ   r„   rà  r   r   ÚPRECOMPUTED_SIZEr"  r½   r+   r=   rA   Úinv_precomputed_replacementsr�   Úpexpr)rü   ÚsymÚexprs      rE   Úensure_size_computedz)PythonWrapperCodegen.ensure_size_computedˆ  s‚   € Ü�cœ5Ÿ<™<Ô(¬^¸CÄ×AVÑAVÔ-WØ�d×)Ñ)Ñ)ØØ×Ñ×#Ñ# CÔ(Ü—7‘7×#Ñ#×@Ñ@ÀÑEˆDØ�N‰N˜c˜U #¤e¨D£k ]Ð3Õ4ð .XÐ(rG   c                 ó   — y r<   rŠ   rû   s    rE   r²  z$PythonWrapperCodegen.finalize_prefix�  rô  rG   T©rB   c               ó   — t        d«      ‚)Nz8codegen_cpp_sizevar is only implemented for cpp_wrapper!)rÄ  ©rü   rª   rB   s      rE   Úcodegen_cpp_sizevarz(PythonWrapperCodegen.codegen_cpp_sizevar“  s   € ÜÐUÓVÐVrG   c               ó   — t        ||¬«      S )Nrñ  )rì  ró  s      rE   Úcodegen_python_sizevarz+PythonWrapperCodegen.codegen_python_sizevar–  s   € Ü�Q Ô*Ð*rG   c                ó$   — | j                  |«      S r<   )rö  r  s     rE   Úcodegen_sizevarz$PythonWrapperCodegen.codegen_sizevar™  s   € Ø×*Ñ*¨1Ó-Ð-rG   c                ó   — |› d|› d�S )NrŽ  r�  rŠ   )rü   Úbasenamerx   rš  s       rE   Úcodegen_tuple_accessz)PythonWrapperCodegen.codegen_tuple_accessœ  s   € Ø�˜1˜U˜G 1Ð%Ð%rG   c                ó¨   — g t        | j                  |«      ¢}t        |«      dk(  ryt        |«      dk(  r	d|d   › d�S ddj                  |«      › d�S )Nr   ú()r,   rr   rC  rp   rs   )r‡  rö  r_   rw   )rü   r”  Úpartss      rE   r–   z/PythonWrapperCodegen.codegen_python_shape_tupleŸ  s^   € Ø:”#�d×1Ñ1°5Ó9Ð:ˆÜˆu‹:˜Š?ØÜˆu‹:˜Š?Ø�u˜Q‘x�j Ð$Ð$Ø�4—9‘9˜UÓ#Ð$ AÐ&Ð&rG   c                ó$   — | j                  |«      S r<   )r–   )rü   r”  s     rE   r™  z(PythonWrapperCodegen.codegen_shape_tuple§  s   € Ø×.Ñ.¨uÓ5Ð5rG   c                ó¬   — dj                  dj                  |t        |«      t        |«      | j	                  |«      | j	                  |«      g«      «      S )Nzalloc_from_pool({})rp   )Úformatrw   rì  r¶   r–   )rü   rx   Úoffsetr   r”  r•  s         rE   Úcodegen_alloc_from_poolz,PythonWrapperCodegen.codegen_alloc_from_poolª  sS   € Ø$×+Ñ+Ø�I‰IàÜ˜&“MÜ˜“JØ×3Ñ3°EÓ:Ø×3Ñ3°FÓ;ðóó

ð 
	
rG   c                ó  — ||j                   j                  k(  rk||j                   j                  k(  rR||j                   j                  k(  r9|�&||j                  k7  rd|j                  «       › d|› d�S |j                  «       › S | j                  |«      }| j                  |«      }| j                  |«      }|�/||j                  k7  r d|j                  «       › d|› d|› d|› d|› d�S d|j                  «       › d|› d|› d|› d�	S )Nzaten.view.dtype(rp   rs   z#aten.view.dtype(reinterpret_tensor(z), zreinterpret_tensor()rn  r‚  r•  r  r   rT  r–   rø  )rü   Údatar‚  r•  r  r�   r   s          rE   Úcodegen_reinterpret_viewz-PythonWrapperCodegen.codegen_reinterpret_view·  s  € ð �D—K‘K×$Ñ$Ò$Ø˜$Ÿ+™+×,Ñ,Ò,Ø˜$Ÿ+™+×,Ñ,Ò,àÐ  U¨d¯j©jÒ%8Ø)¨$¯-©-«/Ð):¸"¸U¸GÀ1ÐEÐEàŸ-™-›/Ð*Ð+à×2Ñ2°4Ó8ˆDØ×4Ñ4°VÓ<ˆFØ×)Ñ)¨&Ó1ˆFØÐ  U¨d¯j©jÒ%8Ø<¸T¿]¹]»_Ð<MÈRÐPTÈvÐUWÐX^ÐW_Ð_aÐbhÐaiÐilÐmrÐlsÐstÐuÐuð *¨$¯-©-«/Ð):¸"¸T¸FÀ"ÀVÀHÈBÈvÈhÐVWÐXðrG   c                ó8   — | j                  |› d|› d|› d�«       y )Nz.copy_(rp   rs   rÚ  )rü   râ   ÚdstÚnon_blockings       rE   Úcodegen_device_copyz(PythonWrapperCodegen.codegen_device_copyÔ  s!   € Ø�‰˜#˜˜g c U¨"¨\¨N¸!Ð<Õ=rG   c                ó`   — | j                  | j                  › |› d|› | j                  › �«       y r„  )r�   r¶  r¸  )rü   rx   rå   s      rE   Úcodegen_multi_outputz)PythonWrapperCodegen.codegen_multi_output×  s)   € Ø�‰˜$Ÿ,™,˜¨ v¨S°°¸¿¹°}ÐEÕFrG   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r<   )r  )r‹   Úts     rE   r�   z>PythonWrapperCodegen.codegen_dynamic_scalar.<locals>.<genexpr>Û  s   è ø€ Ò>¨Q�1×&Ñ&×(Ñ>ùs   ‚r   r×   ú.item()r,   z = 1 if z.item() else 0z_undivided = zassert z_undivided % z
 == 0, f'{z_undivided} not divisible by ú'z_undivided // úunrecognized keypath z = None)rˆ  r_   Úkeypathr�   rí  rƒ   r   r   Údivisorr`   rT  )rü   rC   r  s      rE   Úcodegen_dynamic_scalarz+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Õ3rG   c           
     ó0  ‡ ‡— ˆˆ fd„}ˆfd„}ˆfd„}‰j                  g d¢«       ‰j                  «       5  ‰j                  dd¬«       t        j                  j
                  j                  «       D ]U  \  }}‰j                  d|› �«        |||j                  «       |j                  «       |j                  |j                  «       ŒW t        t        j                  j                  «      d	kD  r^‰j                  d
«       t        j                  j                  j                  «       D ]"  \  }}‰j                  d|› �«        |||«       Œ$ t        j                  j                  j                  «       D �]  \  }}t        |t         j"                  «      rCt        t        j                  j$                  j&                  j)                  |d «      t*        «      rŒdt        |t,        j.                  «      rct        t        j                  j                  «      d	k(  r‰j                  d
«       ‰j                  d|› �«        |||j1                  «       «       Œát        |t         j2                  «      r4 ||t        j                  j$                  j5                  |d¬«      «       �Œ/t        |t,        j6                  «      r# ||d|j                  j8                  › d�«       �Œl|j;                  «       D �cg c]-  }t        j                  j$                  j5                  |d¬«      ‘Œ/ }	}|j=                  «       D �cg c]-  }t        j                  j$                  j5                  |d¬«      ‘Œ/ }
} |||	|
|j?                  «       |jA                  «       «       �Œ" ddjC                  t        j                  j                  jE                  «       «      › d�}‰j                  d|› �«       ‰j                  d«       d d d «       y c c}w c c}w # 1 sw Y   y xY w)Nc                ó‚   •— ‰j                  | › d‰j                  |«      › d‰j                  |«      › d|› d|› d�
«       y )Nz = rand_strided(rp   ú
, device='ú	', dtype=rs   )r�   r–   )rx   r”  r•  r“  r   r¢   rü   s        €€rE   Úadd_fake_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_fake_inputð  sT   ø€ Ø×ÑØ�&Ð(Ø×2Ñ2°5Ó9Ð:¸"Ø×2Ñ2°6Ó:Ð;ð <Ø!˜( )¨E¨7°!ð5õrG   c                ó2   •— ‰j                  | › d|› �«       y r„  rÚ  )rx   rÅ   r¢   s     €rE   Úadd_expr_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_expr_inputø  s   ø€ Ø×Ñ ˜v S¨¨Ð.Õ/rG   c                ó’   •— dd l }t        |t        j                  «      sJ ‚‰j	                  | › d|j                  |«      ›d�«       y )Nr   z = pickle.loads(rs   )Úpicklerƒ   ræ  ÚScriptObjectr�   Údumps)rx   rå   r  r¢   s      €rE   Úadd_torchbind_inputzKPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_torchbind_inputû  sB   ø€ Ûä˜e¤U×%7Ñ%7Ô8Ð8Ð8à×Ñ ˜vÐ%5°f·l±lÀ5Ó6IÐ5LÈAÐNÕOrG   )r¤  r¤  z3def benchmark_compiled_module(times=10, repeat=10):z”
                from torch._dynamo.testing import rand_strided
                from torch._inductor.utils import print_performance
                TrÊ   zglobal r   zimport pickleé*   ©Úfallbackztorch.cuda.default_generators[z].graphsafe_get_state()zcall([rp   z])zfn = lambda: z8return print_performance(fn, times=times, repeat=repeat))#Ú
writelinesr¸   rá   r+   r=   Ú	constantsr\   r�   r‚  r•  r“  r   r_   Útorchbind_constantsr
  rƒ   r„   rà  rA   Ú
var_to_valrk   r   r   r  Úget_real_objr
   Ú	size_hintr  rš  r�  r‘  r�  r@   rw   Úkeys)rü   r¢   r  r  r!  rx   rå   Útorchbind_objrª   r”  r•  Úcall_strs   ``          rE   Úbenchmark_compiled_modulez.PythonWrapperCodegen.benchmark_compiled_moduleï  sy  ù€ õ	ô	0ô	Pð 	×ÑÚKô	
ð �]‰]‹_ñ E	YØ�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Õ<ð	=ô  !Ÿw™w×3Ñ3×9Ñ9Ó;ó (‘��eÜ˜e¤U§\¡\Ô2´zÜ—G‘G×$Ñ$×/Ñ/×3Ñ3°E¸4Ó@Ä,ô8ð Ü˜e¤R×%7Ñ%7Ô8Üœ1Ÿ7™7×6Ñ6Ó7¸1Ò<à×(Ñ(¨Ô9Ø×$Ñ$ w¨t¨fÐ%5Ô6Ù'¨¨e×.@Ñ.@Ó.BÕCÜ ¤u§z¡zÔ2ñ
 # 4¬¯©×)9Ñ)9×)CÑ)CÀEÐTVÐ)CÓ)WÖXÜ ¤r×'8Ñ'8Ô9Ù"ØØ8¸¿¹×9KÑ9KÐ8LÐLcÐdöð "'§¡Ó!1öàô Ÿ™×(Ñ(×2Ñ2°1¸rÐ2ÕBð�Eð ð "'×!1Ñ!1Ó!3öàô Ÿ™×(Ñ(×2Ñ2°1¸rÐ2ÕBð�Fð ñ #ØØØØ×(Ñ(Ó*ØŸ™Ó)öðE(ðT   §	¡	¬!¯'©'×*>Ñ*>×*CÑ*CÓ*EÓ FÐGÀrÐJˆHØ×Ñ˜}¨X¨JÐ7Ô8Ø×ÑÐWÔX÷KE	Yð E	Yùòfùò÷oE	Yð E	Yús+   ¶J>PË42PÌ&PÌ:2PÍ,BPÐ
PÐPc                óö   — t         j                  sy| j                  |«       |j                  g d¢«       |j	                  «       5  |j                  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compiled_module_main('z', benchmark_compiled_module))r   Úbenchmark_harnessr.  r%  r¸   r%   ©rü   r¢   s     rE   r³  z*PythonWrapperCodegen.add_benchmark_harnessL  ss   € ô ×'Ò'Øà×&Ñ& vÔ.à×ÑÒ@ÔAØ�]‰]‹_ñ 	Ø×ÑàXØ,Ô-?Ó-AÐ,BÐB_Ð`ðô÷	÷ 	ñ 	ús   Á A/Á/A8c                ó  — t         j                  j                  r>d|› d|› �}| j                  j	                  |«       t
        j                  j                  ry |r|› d�nd}d|› |› d|› �}| j                  j	                  |«       y )Nz

r×   r9  r¤  )	r   r—   r˜   r±  rá   r+   r=   r<  r¬  )rü   ra  Úkernel_bodyÚmetadataÚgpuÚcpp_definitionÚbodyÚmetadata_comments           rE   Údefine_kernelz"PythonWrapperCodegen.define_kernel^  s�   € ô �=‰=×1Ò1à˜+˜ c¨+¨Ð7ˆDØ×%Ñ%×,Ñ,¨TÔ2Ü�w‰w×"Ò"àÙ.6˜h˜Z r™?¸BÐØÐ&Ð'¨ }°C¸°}ÐEˆØ�‰×Ñ˜4Õ rG   c                ó:   — | j                   j                  |«       y r<   )r²  rá   )rü   Úfn_codes     rE   Údefine_subgraph_launcher_fnz0PythonWrapperCodegen.define_subgraph_launcher_fnq  s   € Ø×!Ñ!×(Ñ(¨Õ1rG   c                óà  ‡‡*‡+‡,‡-‡.‡/— ddl m} ddlm}m}	m}
 ddlmŠ*m}m	}m
}m} ddlm}m}  |«        |j                  }g Š/i Š-g Š,g }ˆ,ˆ/fd„Š+d)ˆ*ˆ+ˆ-ˆfd	„	}t!        |j"                  «      D �]„  \  }}||j$                  v r || ‰*|¬
«      d¬«       Œ(|‰vrŒ-‰|   }‰|   € || ‰*|¬
«      d¬«       ŒJt'        |t(        j*                  «      r || ||¬
«      «       Œut'        |t(        j,                  «      r/ || |||j/                  «       |j1                  «       ¬«      «       Œ¾t'        |t(        j2                  «      rO || |||j4                  j/                  «       |j1                  «       |j6                  j8                  ¬«      «       �Œ't'        |t:        t<        j>                  f«      xr* t@        jB                  jD                  jG                  |d«      } || |||«      |¬«       �Œ‡ tI        ‰/d ‰,|j"                  D �cg c]  }tK        |«      ‘Œ c}¬«      }|tM        jN                  t@        jB                  jQ                  «       «      i ‰-¥tR        jU                  |d«      ¥tW        ‰/‰,¬«      gdœ}|rtY        |«      |d<   |rtY        |«      |d<   t[        |«      dk(  r0|	j]                  «       }g t_        t<        j`                  |d   «      ¢}nÜd*ˆ.fd„}i Š.|D �cg c]  }g t_        ||«      ¢‘Œ }}|rt[        |«      t[        |«      k(  sJ ‚g }tc        te        ||«      d„ d¬«      D ]@  \  }} |jg                   || «      g t_        th        |«      ¢g t_        tj        |«      ¢dœ«       ŒB |
j                  |g t_        tl        ‰.jo                  «       «      ¢dœ}g ‰.jq                  «       ¢}ts        |jt                  «      g}!t[        |«      dkD  rQ‰jo                  «       D ]>  }t'        |t(        j,                  t(        j2                  f«      rŒ.|!jg                  |«       Œ@ |!jg                  tm        |«      «       |!jw                  tm        |«      «       tY        |!«      }!|!| jx                  v rg | jx                  |!   ¢|‘­S |› dt[        | jx                  «      › �}"t{        «       }#t|        j                  j~                  r|#j�                  d|"›d�«       n|#j�                  d|›d�«       |"|d<   |jƒ                  |j…                  «       «       |#j‡                   |«       «       |#j‡                  dg t_        ||«      ¢›d |›d!|›d"�«       t‰        |«      }$t|        j                  j~                  r|$j‹                  d#|› d$�d#|"› d$�«      }$|#j‡                  |$«       t@        jB                  jQ                  «       }%|#j�                  d%|%jŒ                  › d&�«       t�        j�                  |jt                  «      \  }&}'t�        j’                  |jt                  «      }(d'|(› d(|'› �})| j•                  |"|#j—                  «       |)«       |"|f| jx                  |!<   |"||fS c c}w c c}w )+Nr   )Úpatch_triton_dtype_reprr   )Úconfig_to_dictÚ	FixedGridÚPrecomputedGridr,   )ÚConstexprArgÚKernelArgTypeÚSizeArgÚ	TensorArgÚTMADescriptorArg)Úgen_common_triton_importsÚTritonKernelc                óJ   •— ‰j                  |«       ‰j                  | «       y r<   )r  )Úidxra   Úarg_indicesÚ	signatures     €€rE   Úadd_to_signaturezPPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_to_signature”  s   ø€ Ø×Ñ˜SÔ!Ø×Ñ˜sÕ#rG   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©rx   r,   )r*   rx   )	rJ  ra   Úis_constexprÚequals_1Úequals_nonerB  rM  r&  r¨   s	        €€€€rE   Úadd_argzGPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_arg˜  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¨#Õ.rG   rO  T)rP  )rR  )rx   Úbufferr   )rx   rT  r   r  )rQ  )Ú
size_dtyper‘  Úargdefs)r‘  )rL  r“  r&  r¿   Úrestore_valueÚreset_to_zeroc                ó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)rƒ   r„   r
   Úfree_symbolsÚsortr¶   rà  r_   r)   rã   r…   )rî  Úsymbolsrí  Úextra_launcher_argss      €rE   Ú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Ó*Ð*rG   c                ó2   — t        | d   j                  «      S r¦   r§   r©   s    rE   r«   zHPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.<lambda>-  r¬   rG   r­   )r   ÚpythonrW   )Ú	grid_typeÚprecomputed_gridsr^  Ú_zasync_compile.triton(z, '''ra  zG
            @triton_heuristics.user_autotune(
                configs=z ,
                inductor_meta=z,
                triton_meta=z{,
                filename=__file__,
                custom_kernel=True,
            )
            @triton.jit
            r¤   rr   z''', device_str='r%  z# Original path: r;  )FFF)rî  r²   r³   r´   )LÚtorch.utils._tritonr>  Úruntime.triton_heuristicsr?  r@  rA  ÚcommonrB  rC  rD  rE  rF  r—   rG  rH  ræ   râ  Ú	arg_namesÚ
constexprsrƒ   r   ÚTMADescriptorrS  rT  r@   r!   r  rn  r  rã   r„   r…   r+   r=   rA   Ústatically_known_equalsr7   r-   r#   rà  Úget_current_device_or_throwÚdictÚfromkeysr5   r•   r_   Úsetup_grid_as_argsr‡  Úsympifyr»   r¼   r  rì  r4   r¶   r’  r+  ÚidrÜ   Úextendr¾  r0   r   Úunique_user_kernel_namesr�   ÚupdateÚinductor_meta_commonrá   rõ   Úreplacer‘   ÚinspectÚgetsourcelinesÚgetsourcefiler9  r¾   )0rü   ry   r¿   r¨   Úrestore_value_argsÚreset_to_zero_argsrÀ   r>  r?  r@  rA  rC  rD  rE  rF  rG  rH  Úoriginal_nameÚequal_to_1_argsrS  rJ  r®   ra   rQ  rª   Útriton_signatureÚtriton_metaÚinductor_metaÚextra_launcher_call_argsr_  r™   rc  ÚcfgÚ	cache_keyrx   rñ   Ú
kernel_srcÚcurrent_devicerd  ÚlinenoÚsrcfiler4  rB  rM  rK  r&  r^  rL  s0      `                                      @@@@@@rE   Ú!define_user_defined_triton_kernelz6PythonWrapperCodegen.define_user_defined_triton_kernelt  sO  þ€ õ 	@÷	
ñ 	
÷
	
õ 	
÷ 	DáÔ!ØŸ™ˆØ)+ˆ	Ø$&ˆ	Ø!#ˆØ%'ˆõ	$÷"	/ð "	/ôH " &×"2Ñ"2Ó3ó 1	G‰HˆC�Ø�f×'Ñ'Ñ'Ù˜™\¨sÔ3À$ÕGØà˜&Ñ Øà˜‘+ˆCà�c‰{Ð"Ù˜™\¨sÔ3ÀÖFä˜c¤2×#3Ñ#3Ô4ÙØÙ(Ø!$ôõô   ¤R§Y¡YÔ/ÙØÙ!Ø!$Ø#&§<¡<£>Ø"%§-¡-£/ôõô   ¤R×%7Ñ%7Ô8ñ ØÙ!Ø!$Ø#&§8¡8×#4Ñ#4Ó#6Ø"%§-¡-£/Ø#&§:¡:×#4Ñ#4ô	öô  *Øœc¤5§=¡=Ð1ó ò  äŸ'™'×*Ñ*×BÑBØØóð ñ ˜C¡¨¨cÓ!2¸X×Fðc1	Gôf -ØØØØ)/×)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ñõð	ð -×5Ñ5Ø%6Ø'P¬¬SÐ2E×2LÑ2LÓ2NÓ)OÐ'PñˆMð
 (EÐ)<×)AÑ)AÓ)CÐ'DÐ$ô ˜VŸY™Y›-˜ˆ	Üˆw‹<˜!ÒØ—}‘}“ò *�ä! #¬¯	©	´2×3EÑ3EÐ'FÕGØ×$Ñ$ SÕ)ð*ð 	×Ñœ˜[Ó)Ô*Ø×Ñœ˜]Ó+Ô,Ü˜)Ó$ˆ	Ø˜×6Ñ6Ñ6ðØ×/Ñ/°	Ñ:ðà(ñð ð
  � ¤# d×&DÑ&DÓ"EÐ!FÐGˆä(Ó*ˆÜ�=‰=×1Ò1Ø×%Ñ%Ð(=¸d¸XÀUÐ&KÕLà×%Ñ%Ð(=¸mÐ=NÈeÐ&TÔUà'+ˆ�mÑ$Ø×Ñ˜\×>Ñ>Ó@ÔAà×ÑÑ8Ó:Ô;Ø×Ñðà8œ3˜~¨wÓ7Ð8Ð;ð <Ø,Ð/ð 0Ø(˜Oð ,ð		ô	
ô OÈvÓVˆ
Ü�=‰=×1Ò1à#×+Ñ+¨d°=°/ÀÐ,CÀtÈDÈ6ÐQRÀ^ÓTˆJØ×Ñ˜zÔ*äŸ™×<Ñ<Ó>ˆØ×!Ñ!Ð$5°n×6IÑ6IÐ5JÈ"Ð"MÔNÜ×*Ñ*¨6¯9©9Ó5‰	ˆˆ6Ü×'Ñ'¨¯	©	Ó2ˆØ& w i¨q°°Ð9ˆØ×ÑØØ×$Ñ$Ó&Øô	
ð 6:¸;Ð4Gˆ×&Ñ& yÑ1Ø�[Ð":Ð:Ð:ùòI ;ùòj Ps   ÈY&Ë)Y+c                ó¸   — |› d|j                   › d�}|�|d|› �z  }| j                  |› dt        |j                  «      › �«       t	        ||j                  «      S )Nrd  r�  r×   )r­  r�   rì  r�  r÷   )rü   ra  Útreer®  rî  s        rE   Úgenerate_numel_exprz(PythonWrapperCodegen.generate_numel_expry  sb   € Ø�˜a §¡˜}¨EÐ2ˆØÐØ�a˜�x�LÑ ˆDØ�‰˜$˜˜s¤5¨¯©Ó#4Ð"5Ð6Ô7ô ˜t T§Z¡ZÓ0Ð0rG   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)r,   )r”  r•  )rT  rY  Ú	zero_moder3   ÚUNINITIALIZEDr�   ÚZERO_ON_CALLÚmake_zero_bufferÚZERO_PER_GRAPHrÓ  rk   rƒ   rC   r2   Úmaximumr`   r   r—   r˜   rŸ   r  Úmake_allocationr“  r   r+   r=   rA   r*  r   )rü   Úwsrx   r    Úpriors        rE   Ú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×+Ñ+×5Ñ5°b·h±hÓ?ÐAØð 5ó ô	ð �|‰|Ô0×>Ñ>Ò>Ø×*Ñ*×4Ñ4Ü(×9Ñ9¸$ÀÓEõð ?ð 2rG   c                óv   — |j                   t        j                  k7  r| j                  t	        | |«      «       y y r<   )r�  r3   r‘  r�   r  )rü   r”  s     rE   Úgenerate_workspace_deallocationz4PythonWrapperCodegen.generate_workspace_deallocation®  s.   € Ø�<‰<Ô,×;Ñ;Ò;Ø�N‰NÔ.¨t°RÓ8Õ9ð <rG   c                ó$   — |› d| j                   › �S )Nz.zero_())r¸  )rü   rx   s     rE   r�  z%PythonWrapperCodegen.make_zero_buffer²  s   € Ø��x §¡˜}Ð-Ð-rG   c                óH   — |› ddj                  |«      › d| j                  › �S )Nrr   rp   rs   )rw   r¸  )rü   rx   Ú	call_argss      rE   r�  z%PythonWrapperCodegen.wrap_kernel_callµ  s'   € Ø��q˜Ÿ™ 9Ó-Ð.¨a°·±¨}Ð=Ð=rG   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+   r=   Úgraph_idr§  r¸   )rü   rµ  s     rE   r©  z8PythonWrapperCodegen.generate_profiler_mark_wrapper_call¸  sb   € Ø×Ñ×#Ñ#Ð$PÔQØ×Ñ×#Ñ#Ø*¬1¯7©7×+;Ñ+;Ð*<Ð<UÐVô	
ð 	×Ñ˜D×-Ñ-×4Ñ4Ó6Õ7rG   c                ó:   — | j                   j                  d«       y )Nzstart_graph())r°  r�   rû   s    rE   rª  z)PythonWrapperCodegen.generate_start_graph¿  s   € Ø×Ñ×#Ñ# OÕ4rG   c                ó^   — | j                   j                  dt        j                  ›d�«       y )Nz
end_graph(rs   )r°  r�   r   Úprofile_bandwidth_outputrû   s    rE   r¯  z'PythonWrapperCodegen.generate_end_graphÂ  s'   € Ø×Ñ×#Ñ# j´×1PÑ1PÐ0SÐSTÐ$UÕVrG   c                ó^   — | j                   j                  dt        j                  › d�«       y )NúU
            for kernel in globals().values():
                if isinstance(kernel, zU.CachingAutotuner):
                    kernel.cuda_kernel_saved = False
            ©r°  rá   r"   ræ   rû   s    rE   r®  z6PythonWrapperCodegen.generate_reset_kernel_saved_flagsÅ  s2   € Ø×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð Cðõ	
rG   c                ó^   — | j                   j                  dt        j                  › d�«       y)a[  
        Precompile and save the CUBINs of the 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.
        r¢  aÀ  .CachingAutotuner):
                    if not kernel.cuda_kernel_saved:
                        if len(kernel.launchers) == 0:
                            kernel.precompile()
                        kernel.save_gpu_kernel(
                            grid=(0, 0, 0),   # use dummy grid
                            stream="stream",  # use dummy stream
                            launcher=kernel.launchers[0],
                        )
            Nr£  rû   s    rE   r°  z5PythonWrapperCodegen.generate_save_uncompiled_kernelsÎ  s4   € ð 	×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð 	Cðõ	
rG   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  )rƒ   r¶   r6   rã   Úfloaträ   r÷   rì  r+   r=   rA   rB   )ra   s    rE   Ú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Ó;Ó<Ð<rG   rŠ   )rü   r›  r¨  ra   s       rE   Úprepare_triton_kernel_callz/PythonWrapperCodegen.prepare_triton_kernel_callè  s!   € ò	=ð *3Ö3 #‘˜•Ò3Ð3ùÒ3s   ˆc                ó  ‡ — t        |t        «      �rt        |t        j                  «      r:|j                  j                  «       }t        j                  j                  |«      }nQt        j                  j                  |«      �"|}t        j                  j                  |«      }n|€J d«       ‚d|› �}|}t        d„ |j                  «       D «       «      }t        d„ t        j                  j                  |«      D «       «      }t        d„ |j                  «       D «       «      }	|j                  «       }
|j                  «       }t        j                  j                   j#                  |j%                  «       j&                  t(        j*                  ¬«      }d|› d|	› d	|
› d
|› d|› d|› d�}‰ j,                  j/                  |› d|› �«       t        |t        j                  «      r5‰ j1                  |d¬«      }|}‰ j,                  j/                  |› d|› �«       |S t3        |t4        j6                  «      st        |t8        «      rÐt        |t:        «      r|‰ j<                  v r|S |€y|}t        |t8        «      r|j>                  }|t        j                  j                   j@                  v r't        j                  j                   j@                  |   }t;        t        j                  j                   jC                  |t(        j*                  ¬«      «      S t        |t:        tD        tF        tH        f«      rt;        |«      S t        |tJ        «      rddjM                  ˆ fd„|D «       «      › d�S tO        dtQ        |«      › �«      ‚)NzBV.graph.get_buffer(arg) and raw_arg can't be None at the same timeÚtmp_arg_c              3  óŽ   K  — | ]=  }t         j                  j                  j                  |t        j
                  ¬ «      –— Œ? y­w©r#  N©r+   r=   rA   rm  r   Úunbacked_symint_fallback©r‹   rÉ  s     rE   r�   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>  s@   è ø€ ò ð
 ô	 —‘× Ñ ×;Ñ;ØÜ#×<Ñ<ð <÷ ñùó   ‚AAc              3  óŽ   K  — | ]=  }t         j                  j                  j                  |t        j
                  ¬ «      –— Œ? y­wr­  r®  r°  s     rE   r�   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>  s@   è ø€ ò $ð
 ô	 —‘× Ñ ×;Ñ;ØÜ#×<Ñ<ð <÷ ñ$ùr±  c              3  óŽ   K  — | ]=  }t         j                  j                  j                  |t        j
                  ¬ «      –— Œ? y­wr­  r®  r°  s     rE   r�   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>  s@   è ø€ ò ð
 ô	 —‘× Ñ ×;Ñ;ØÜ#×<Ñ<ð <÷ ñùr±  r#  zgenerate_example_value(rp   z, 'z', rs   r×   T)r  r€  r  rŽ  c              3  óT   •K  — | ]  }‰j                  |t        |«      «      –— Œ! y ­wr<   r�   )r‹   Úarü   s     €rE   r�   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>C  s#   øè ø€ Ò ZÐQR ×!@Ñ!@ÀÄDÈÃG×!LÑ Zùr“   r�  zUnsupported type ))rƒ   Útorch_dtyper   rj  r|  rT  r+   r=   Ú
get_bufferÚtry_get_bufferr•   r�  Úget_allocation_sizer‘  r�  r@   rA   r*  Ú
get_layoutr  r   r¯  rŸ   r�   r‚  Ú
issubclassr„   ÚBasicr÷   r¶   rÐ  rù   rë  rm  rã   r§  rä   r  rw   r�  r‘   )rü   ra   Úarg_typeÚraw_argrš  r—  r  r‚  Úallocation_sizer•  r“  r   r  rå   s   `             rE   r�   z/PythonWrapperCodegen.generate_example_arg_valueô  s  ø€ Ü�h¤Õ,Ü˜'¤2×#3Ñ#3Ô4à"Ÿ>™>×2Ñ2Ó4�Ü—g‘g×(Ñ(¨Ó2‘Ü—‘×'Ñ'¨Ó,Ð8Ø�Ü—g‘g×(Ñ(¨Ó-‘àÐ*ð ØXóÐ*ð & e WÐ-�Ø�äñ ð
 Ÿ™›ôó ˆDô $ñ $ô
 Ÿ™×4Ñ4°SÓ9ô$ó ˆOô ñ ð
 Ÿ™Ó)ôó ˆFð —^‘^Ó%ˆFØ—M‘M“OˆEÜ—W‘W×%Ñ%×/Ñ/Ø—‘Ó ×'Ñ'Ü×8Ñ8ð 0ó ˆ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�äÜ—‘× Ñ ×;Ñ;Ø¤&×"AÑ"Að <ó óð ô ˜œc¤3¬¬tÐ4Ô5Ü�s“8ˆOÜ˜œTÔ"Ø�t—y‘yÓ ZÐVYÔ ZÓZÐ[Ð[\Ð]Ð]ä%Ð(9¼$¸s»)¸Ð&EÓFÐFrG   c                óz   ‡ — t        |t        «      r ddj                  ˆ fd„|D «       «      z   dz   S t        |«      S )NrŽ  rp   c              3  ó@   •K  — | ]  }‰j                  |«      –— Œ y ­wr<   )Ú_grid_dim_str)r‹   r†   rü   s     €rE   r�   z5PythonWrapperCodegen._grid_dim_str.<locals>.<genexpr>J  s   øè ø€ ÒR¸T × 2Ñ 2°4× 8ÑRùs   ƒr�  )rƒ   r  rw   rì  )rü   Úgrid_per_dims   ` rE   rÂ  z"PythonWrapperCodegen._grid_dim_strG  s<   ø€ Ü�l¤DÔ)à�d—i‘iÓRÀ\ÔRÓRÑRÐUXÑXðô ˜Ó&Ð&rG   )r“  r—   Ú	arg_typesrš  r  c          
     óŒ  — |xs t         j                  j                  «       }|s1|j                  dk7  s"| j	                  | j                  ||«      «       y| j                  |«      }dj                  |«      }t        j                  | |j                  t         j                  «      }	|s$d|	› d�}
| j	                  |› d|› d|› d|
› d�«       y| j                  «        t        j                  j                  �rä|| j                  v�rÕ|�t!        |«      t!        |«      k(  sJ d«       ‚i }g }|€dgt!        |«      z  }nt!        |«      t!        |«      k(  sJ d	«       ‚t#        t%        |||«      «      D ]¼  \  }\  }}}d}t'        |t(        «      r!d
t)        |«      v r|j+                  d
«      \  }}t'        |t,        «      rBt/        j0                  d|«      r|}|||<   n8||vr| j3                  ||||«      }|||<   n||   }n| j3                  ||||«      }|j5                  |€|n|› d
|› �«       Œ¾ | j6                  j	                  |› ddj                  |«      › d|	› d�«       | j6                  j	                  ddj                  d„ |j9                  «       D «       «      › d�«       | j                  j;                  |«       t         j                  j<                  ryt         j                  j>                  j@                  }|jC                  |||d«       |5  | j	                  |› d|› d|	› d�«       ddd«       y# 1 sw Y   yxY 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\  Nrp   z	c_void_p(rs   rÖ   rr   z$call_args and arg_types do not matchz#call_args and raw_args do not matchrO  z^(workspace|semaphore)z.run(z	, stream=údel c              3  ó    K  — | ]  }|–— Œ y ­wr<   rŠ   )r‹   ra   s     rE   r�   z<PythonWrapperCodegen.generate_kernel_call.<locals>.<genexpr>£  s   è ø€ Ò E¨¤Ñ Eùó   ‚r9  )"r+   r=   rl  r‘   r�   r�  r©  rw   r  rÎ  rš  rú  r   r—   r˜   rž   r_   râ  r¼   rƒ   r¶   Úsplitr¶  r]   Úmatchr�   r  rŸ   r’  r½   r<  rd  r×  re  )rü   ra  r›  r“  r—   rÄ  rš  r  Úcall_args_strÚstream_nameÚ
stream_ptrÚtensor_argsÚall_argsrÕ  ra   r½  r¾  r®   Úarg_strrh  s                       rE   Úgenerate_kernel_callz)PythonWrapperCodegen.generate_kernel_callO  sE  € ð" Ò@œ1Ÿ7™7×>Ñ>Ó@ˆÙ˜&Ÿ+™+¨Ò.Ø�N‰N˜4×0Ñ0°¸iÓHÔIØà×7Ñ7¸	ÓBˆØŸ	™	 -Ó0ˆÜ*×?Ñ?Ø�&—,‘,¤§¡ó
ˆñ Ø$ [ M°Ð3ˆJØ�N‰NØ�-˜q  ¨Q¨}¨o¸RÀ
¸|È1ÐMôð à×%Ñ%Ô'ô �M‰M×2Ó2Ø 4×#=Ñ#=Ò=ð Ð(¬S°«^¼sÀ9»~Ò-Mð Ø6óÐMð ˆKØˆHØÐà ˜6¤C¨	£NÑ2‘ä˜8“}¬¨I«Ò6ð Ø9óÐ6ô 09Ü�I˜y¨(Ó3ó0ò PÑ+�Ñ+�C˜ 7ð �Ü˜c¤3Ô'¨C´3°s³8©Oà"Ÿy™y¨›~‘H�C˜ä˜h¬Ô4ô —x‘xÐ 9¸3Ô?Ø"%˜Ø+2˜ CÒ(Ø KÑ/Ø"&×"AÑ"AØ ¨7°Aó#˜ð ,3˜ CÒ(à"-¨cÑ"2™à"×=Ñ=¸cÀ8ÈWÐVWÓX�GØ—‘¨3¨;¡¸s¸eÀ1ÀWÀIÐ<NÕOð/Pð2 ×&Ñ&×0Ñ0Ø�-˜u T§Y¡Y¨xÓ%8Ð$9¸À;À-ÈqÐQôð ×&Ñ&×0Ñ0Ø�t—y‘yÑ E°×0BÑ0BÓ0DÔ EÓEÐFÀbÐIôð ×&Ñ&×*Ñ*¨;Ô7Ü�w‰w×"Ò"àô !"§¡× 4Ñ 4× BÑ BÐØ×.Ñ.¨y¸+ÀyÐRVÔWØ"ñ 	XØ�N‰N˜k˜]¨%°¨¸iÈÀ}ÐTUÐVÔW÷	X÷ 	Xñ 	Xús   ÌL:Ì:Mc                ó:   — | j                   j                  |«       y r<   )rµ  r  )rü   r    s     rE   r�   zPythonWrapperCodegen.writeline°  s   € Ø�
‰
×Ñ˜$ÕrG   c                ó4   — |D ]  }| j                  |«       Œ y r<   rÚ  )rü   rµ  r    s      rE   r%  zPythonWrapperCodegen.writelines³  s   € Øò 	!ˆDØ�N‰N˜4Õ ñ	!rG   c                óL   — | j                   j                  t        |«      «       y r<   )rµ  r  r&   )rü   Úctxs     rE   r§  z"PythonWrapperCodegen.enter_context·  s   € Ø�
‰
×Ñœ+ cÓ*Õ+rG   c                ó&  ‡ ‡— ddl m}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.                  t0        f«      r|j3                  «       S  |«       r(t	        |j4                  j6                  «      r ||«      S t	        |t*        j8                  «      r|j3                  «       S t        |«      S )Nr   )Údtype_to_stringÚhas_triton_packagec                  ó   — e Zd ZU ded<   d„ Zy)ú1PythonWrapperCodegen.val_to_arg_str.<locals>.Shimr   Úrefc                ó   — | j                   S r<   )rÛ  rû   s    rE   Ú__repr__z:PythonWrapperCodegen.val_to_arg_str.<locals>.Shim.__repr__Ê  s   € ØŸ8™8�OrG   N)ræ   rÙ   rþ   rÔ   rÝ  rŠ   rG   rE   ÚShimrÚ  Æ  s   … à“ó$rG   rÞ  c              3  óV   •K  — | ]   } ‰t         j                  ‰|«      «      –— Œ" y ­wr<   rs  )r‹   rµ  rÞ  rü   s     €€rE   r�   z6PythonWrapperCodegen.val_to_arg_str.<locals>.<genexpr>Ï  s$   øè ø€ ÒVÈq™Ô1×@Ñ@ÀÀqÓI×JÑVùs   ƒ&))re  r×  rØ  r—   rƒ   r   rì  rC   rî  r„   r
   r•   r  rP  Ú	dataclassrX   r‘   ræ  Ú_opsÚ
OpOverloadr   r   rS  Ú
MutableBoxr!   r  Úlanguager   r  )rü   rÎ  Útype_r×  rØ  r—   rÞ  s   `     @rE   rt  z#PythonWrapperCodegen.val_to_arg_strº  s'  ù€ ßKáÔÛä�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Ù" 1Ó%Ð%Ü˜œ2×,Ñ,Ô-Ø×&Ñ&Ó(Ð(ä˜“7ˆNrG   c                ó>  — |j                  «       }|j                  «       }t        |j                  «       «      }t        t        j
                  j                  |«      «      }t        |j                  «       «      }| j                  |j                  «       |||||«      S r<   )
r�  r@   r•   r�  r+   r=   r¹  r‘  r“  rT  )rü   rT  r“  r   r”  Úallocation_shaper•  s          rE   ri  z+PythonWrapperCodegen.make_buffer_allocationÝ  s„   € Ø×"Ñ"Ó$ˆØ× Ñ Ó"ˆÜ�f—o‘oÓ'Ó(ˆÜ ¤§¡×!<Ñ!<¸VÓ!DÓEÐÜ�v×(Ñ(Ó*Ó+ˆØ×#Ñ#Ø�O‰OÓ˜v u¨e°VÐ=Mó
ð 	
rG   c           
     ó   — |€|}| j                  |«      }| j                  |«      }| j                  |«      }	|j                  dv r|› d|j                  › d|› d|	› d|› d�
}
n|› d|› d|	› d|j                  › d|› d�
}
||k7  r|
d	|› d|	› d�z   }
|
S )
N)r\  ÚcudaÚxpuz = empty_strided_rr   rp   rs   z = empty_strided(r  r  z.as_strided()r–   r‘   )rü   rx   r“  r   r”  r•  rç  r™  Úcodegen_allocation_shape_tupleÚcodegen_stride_tuplerf  s              rE   r“  z$PythonWrapperCodegen.make_allocationç  sð   € ð Ð#Ø$Ðà"×=Ñ=¸eÓDÐØ)-×)HÑ)HØó*
Ð&ð  $×>Ñ>¸vÓFÐØ�;‰;Ð0Ñ0ð �&Ð)¨&¯+©+¨°aØ1Ð2°"Ø'Ð(¨Ø�'˜ðñ ð �&Ð)Ø1Ð2°"Ø'Ð(ð )Ø!Ÿ;™;˜- y°°°qð:ð ð Ð"@Ò@à˜,Ð':Ð&;¸2Ð>RÐ=SÐSTÐUÑUˆCØˆ
rG   c           	     ó`   — | j                   › |› d|› | j                  › d| j                  › d|› �	S )Nr×   ú  rq   )r¶  r¸  r¹  )rü   Únew_nameÚold_namer¹  s       rE   Úmake_tensor_aliasz&PythonWrapperCodegen.make_tensor_alias	  s6   € Ø—,‘,� ˜z¨¨X¨J°t·{±{°mÀ2ÀdÇlÁlÀ^ÐSTÐU\ÐT]Ð^Ð^rG   c                ó(   — d|j                  «       › �S )NrÆ  )rT  )rü   rT  s     rE   rq  z%PythonWrapperCodegen.make_buffer_free
	  s   € Ø�f—o‘oÓ'Ð(Ð)Ð)rG   c                ó8   — ddj                  d„ |D «       «      › �S )NrÆ  rp   c              3  ó    K  — | ]  }|–— Œ y ­wr<   rŠ   )r‹   rx   s     rE   r�   z:PythonWrapperCodegen.make_free_by_names.<locals>.<genexpr>	  s   è ø€ Ò>¨¤Ñ>ùrÈ  )rw   )rü   Únames_to_dels     rE   Úmake_free_by_namesz'PythonWrapperCodegen.make_free_by_names	  s   € Ø�d—i‘iÑ>°Ô>Ó>Ð?Ð@Ð@rG   c           	     ó`   — | j                   › |› d|› |› | j                  › d| j                  › d�	S )Nr×   rî  ú reuse)r·  r¸  r¹  )rü   rð  rï  Údel_lines       rE   Ú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ÐtrG   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ø  )r@   rT  r+   r=   rÏ  rq  r�  r‘  rú  r  r°  r�   r¶  r¹  )rü   ÚoldÚnewrt  rð  rï  rù  Úreinterpret_views           rE   rw  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^Ð_Ð_rG   c                ó¬   — | j                  t        || j                  › |› d|j                  «       › | j                  › d| j
                  › d�«      «       y )Nr×   rî  z alias)r�   r/   r¶  r  r¸  r¹  )rü   rx   Úviews      rE   Úcodegen_deferred_allocationz0PythonWrapperCodegen.codegen_deferred_allocation#	  sS   € Ø�‰ÜØØ—<‘<�.   c¨$×*@Ñ*@Ó*BÐ)CÀDÇKÁKÀ=ÐPRÐSW×S_ÑS_ÐR`Ð`fÐgóõ	
rG   c                ój  — |j                  «       }|t        j                  j                  v s(|| j                  v st        |t        j                  «      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                   «      �rHt        |j"                  t        j$                  «      s*J dt'        |j"                  «      › d|j"                  › �«       ‚t        |j"                  j(                  t        j*                  «      s$J t'        |j"                  j(                  «      «       ‚t        |j"                  j(                  j(                  t        j,                  «      s$J t'        |j"                  j(                  «      «       ‚| j/                  |j"                  j(                  j(                  «       | j1                  ||j"                  «       y t        |t        j2                  «      r| j5                  t7        | |«      «       y | j5                  t9        | |«      «       y )Nzunexpected rÕ   )rT  r+   r=   r]  rÊ  rƒ   r   ÚDonatedBufferr½   Úget_defining_opÚExternKernelAllocÚMultiOutputÚshould_allocater„  ÚMutationLayoutSHOULDREMOVErX  ÚNonOwningLayoutr  r!   r‘   r  Ú
StorageBoxrS  Úcodegen_allocationr  r…  r�   r�  rY  )rü   rT  rx   rn  s       rE   r  z'PythonWrapperCodegen.codegen_allocation+	  sæ  € Ø�‰Ó ˆð ”A—G‘G×+Ñ+Ñ+Ø�t—~‘~Ñ%Ü˜&¤"×"2Ñ"2Ô3àØ�‰×Ñ˜4Ô äØ×&Ñ&Ó(Ü×%Ñ%¤r§~¡~Ð6ôð ×*Ñ*Ô,àà×'Ñ'Ó)ˆÜ�fœb×;Ñ;Ô<ØÜ�fœbŸm™mÔ,ØÜ�fœb×0Ñ0Õ1Ü˜fŸk™k¬2×+=Ñ+=Ô>ð Øœd 6§;¡;Ó/Ð0°°6·;±;°-Ð@óÐ>ô ˜fŸk™k×.Ñ.´·±Ô>ÐVÄÀVÇ[Á[×EUÑEUÓ@VÓVÐ>Ü˜fŸk™k×.Ñ.×3Ñ3´R·Y±YÔ?ÐWÄÀfÇkÁk×FVÑFVÓAWÓWÐ?Ø×#Ñ# F§K¡K×$4Ñ$4×$9Ñ$9Ô:Ø×,Ñ,¨T°6·;±;Ô?Øä�fœb×1Ñ1Ô2Ø�N‰NÔ1°$¸Ó?Ô@Øà�‰”| D¨&Ó1Õ2rG   c                óÒ  — |j                  «       }t        |t        j                  t        j                  f«      r!| j                  | j                  |«      «       y t        |j                  «       t        j                  «      r| j                  t        | |«      «       y | j                  |«      sy | j                  j                  |«       | j                  t        | |«      «       y r<   )rT  rƒ   r   ÚInputBufferr  r�   rq  r„  r…  r   Ú	can_reuserË  r½   r  )rü   rT  rx   s      rE   Úcodegen_freez!PythonWrapperCodegen.codegen_freeS	  s¨   € Ø�‰Ó ˆô �fœrŸ~™~¬r×/AÑ/AÐBÔCØ�N‰N˜4×0Ñ0°Ó8Ô9Øä�f×,Ñ,Ó.´×0CÑ0CÔDð �N‰NÔ-¨d°FÓ;Ô<Øà�~‰~˜fÔ%ØØ�
‰
�‰�tÔà�‰Ô*¨4°Ó8Õ9rG   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 r<   )rT  r+   r=   r]  r
  rƒ   Úgraph_inputs_originalr   r  r&  r'  Únever_reuse_buffersrË  )rü   Úinput_bufferÚoutput_bufferrx   s       rE   r  zPythonWrapperCodegen.can_reuseg	  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Ð!ð
ð 	
rG   c                óš   — |j                  «       | j                  v xr. | j                  |j                  «          |j                  «       k(  S r<   )rT  rÌ  )rü   rT  Úreused_buffers      rE   Ú	did_reusezPythonWrapperCodegen.did_reusew	  sC   € ð �O‰OÓ §¡Ð,ò KØ—‘˜FŸO™OÓ-Ñ.°-×2HÑ2HÓ2JÑJð	
rG   c                ót  — t        ||«      sJ ‚| j                  |«       | j                  j                  |j	                  «       «       | j
                  j                  |j	                  «       «       |j	                  «       | j                  |j	                  «       <   | j                  t        | ||«      «       y r<   )	rO   r  rË  r½   rT  rÊ  rÌ  r�   r`  )rü   r  r  s      rE   Úcodegen_inplace_reusez*PythonWrapperCodegen.codegen_inplace_reuse	  s‰   € Ü$ \°=ÔAÐAÐAØ×Ñ Ô-Ø�
‰
�‰�|×,Ñ,Ó.Ô/Ø�‰×Ñ˜=×1Ñ1Ó3Ô4Ø0<×0EÑ0EÓ0Gˆ�‰�M×*Ñ*Ó,Ñ-Ø�‰”y  |°]ÓCÕDrG   c                óŒ   — t        |«      }|| j                  v r|S | j                  j                  |«       | j                  |z   S r<   )r¶   r¿  r½   r¶  )rü   rí   rx   s      rE   Úcodegen_unbacked_symbol_declz1PythonWrapperCodegen.codegen_unbacked_symbol_decl‡	  sC   € Ü�6‹{ˆØ�4×-Ñ-Ñ-ØˆKð ×&Ñ&×*Ñ*¨4Ô0Ø—<‘< $Ñ&Ð&rG   c                ó&  ‡‡‡‡— t        t        j                  j                  j                  |«      }|sy |j                  «       D ]I  \  }Šdˆfd„Šˆˆˆˆfd„}| j                  | j                  |«      › d |«       › | j                  › �«       ŒK y )Nc                ó¼  •— |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        «      r ‰| › d|d   j                  › d�|dd  «      S t        d|› �«      ‚)NrŠ   r   r   r,   rÖ   rr   rs   rý  z	std::get<z>(rŽ  r�  z.__floordiv__(r  )r_   rƒ   r   ÚpytreeÚSequenceKeyrx   rJ  r+   r=   r<  r   r  r`   )rî  r  Úgos     €rE   r!  zIPythonWrapperCodegen.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ñ    n°W¸Q±Z×5GÑ5GÐ4HÈÐJÈGÐTUÐTVÈKÓXÐXä(Ð+@ÀÀ	Ð)JÓKÐKrG   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 )Nr,   r   )r+   r=   r<  r_   rT  rƒ   r   r  r‘  r  r   rJ  )rf  r!  r  r_  r›  s    €€€€rE   Úgo_outerzOPythonWrapperCodegen.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Ð3rG   r×   )rî  r¶   r  zpytree.KeyPath)	r   r+   r=   rA   Ú	shape_envr\   r�   r  r¸  )rü   r_  r›  Úunbacked_bindingsrÎ  r#  r!  r  s    ``   @@rE   Ú(codegen_unbacked_symbol_defs_for_outputsz=PythonWrapperCodegen.codegen_unbacked_symbol_defs_for_outputs�	  sŠ   û€ ô 6Ü�G‰G×Ñ×&Ñ&Ð(9ó
Ðñ !Øð ,×1Ñ1Ó3ò <	‰JˆAˆwõ
L÷<4ð. �N‰NØ×4Ñ4°QÓ7Ð8¸¹H»J¸<ÈÏÉÀ}ÐUõñu<	rG   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_   r=   r
  r¼   r�   r¶  r¸  )Úinner_inputÚouter_inputÚouter_inputsrü   Úsubgraphs     €€€rE   Ú_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õñrG   c                 ó  •— t        ‰j                  j                  «      t        ‰«      k(  sJ ‚t        ‰j                  j                  ‰«      D ]5  \  } }‰j	                  |› d| j                  «       › ‰j                  › �«       Œ7 y r„  )r_   r=   r  r¼   r�   r  r¸  )Úinner_outputÚouter_outputÚouter_outputsrü   r,  s     €€€rE   Ú_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õñrG   ú subgraph: )Úparent_graph)	r%  r=   r�   r¹  rx   r+   Úset_graph_handlerÚcodegen_subgraphr2  )rü   r,  r+  r1  r-  r2  r4  s   ````   rE   Ú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                óH  — t        |«      t        |j                  j                  «      k(  s!J d|j                  j                  › d|› �«       ‚t        |j                  j                  |«      D ]3  \  }}| j	                  | j
                  › |› d|› | j                  › �«       Œ5 y )Nzgraph_input_names:z, outer_inputs: r×   )r_   r=   r  r¼   r�   r¶  r¸  )rü   r,  r+  r1  r)  r*  s         rE   Úcodegen_subgraph_prefixz,PythonWrapperCodegen.codegen_subgraph_prefix
  s    € ô �<Ó ¤C¨¯©×(HÑ(HÓ$IÒIð 	
Ø  §¡×!AÑ!AÐ BÐBRÐS_ÐR`Ðaó	
ÐIô ),Ø�N‰N×,Ñ,¨ló)
ò 	XÑ$ˆK˜ð �N‰N˜dŸl™l˜^¨K¨=¸¸K¸=ÈÏÉÈÐVÕWñ	XrG   c           	     óT  — |j                   }|j                  }dj                  |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 )z'Generate code to call a graph partitionrp   r,   r  r¤  Ú	partitionú	_args = [r�  rÆ  rr   z) = self.partitions[z](partitionú_args)zdel partitionÚ_argsN)Úinput_deallocationÚoutput_nodesrw   r+  r_   rT  r�   r\   )rü   Úpartition_idrß  r?  r@  rˆ  rC   Úoutput_namesr›  rx   Ú
deallocaterõ  s               rE   Úcodegen_partition_callz+PythonWrapperCodegen.codegen_partition_call
  s<  € ð 2×DÑDÐØ+×8Ñ8ˆà—‘Ð-×2Ñ2Ó4Ó5ÜÐ)Ó*¨aÒ/‰C°Rñ
ˆð 5AÖ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Ð:Õ;ùò! Bùó
s   ÁDÂ7
D$ÃD$c                óP   — t        |«      D �cg c]  }d|› �‘Œ	 c}| _        y c c}w )NÚ
partition_)rÐ  rK  )rü   Únum_partitionsrJ  s      rE   Úset_all_partition_namesz,PythonWrapperCodegen.set_all_partition_names7
  s$   € ÜBGÈÓBWÖ#X¸3 j°°Ò$6Ò#XˆÕ ùÒ#Xs   Ž#c           	     óê  — |j                   j                  }dj                  |«      }t        |«      dk(  r|dz  }dj                  |«      t        |«      dk(  rdndz   }| j	                  |j                   j
                  › d|› d�«       |d t        |«       D ]  }| j	                  d|› �«       Œ | j	                  d|› d	|j                   j
                  › d|j                   j
                  › d
�«       y )Nrp   r,   r  r¤  r<  r�  rÆ  rr   z) = r=  )r=   r  rw   r_   r�   rx   )rü   r,  r+  r1  r  Úinner_inputsÚouter_output_namesr)  s           rE   Úcodegen_subgraph_callz*PythonWrapperCodegen.codegen_subgraph_call:
  sï   € à—n‘n×6Ñ6ˆØ—y‘y Ó-ˆÜˆ{Ó˜qÒ Ø˜CÑˆLà!ŸY™Y }Ó5Ü�}Ó%¨Ò*‰C°ñ
Ðð
 	�‰˜(Ÿ.™.×-Ñ-Ð.¨i¸°~ÀQÐGÔHØ&Ð':¬¨\Ó):Ð;ò 	1ˆKØ�N‰N˜T + Ð/Õ0ð	1ð 	�‰ØÐ"Ð# 4¨¯©×(;Ñ(;Ð'<¸A¸h¿n¹n×>QÑ>QÐ=RÐRXÐYõ	
rG   c                ó~  — t         j                  j                  r| j                  |||«       y | j	                  |j                  «       | j                  d«       | j                  | j                  › d|j                  › �«       | j                  |||«       t         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!                  j"                  «       | j%                  |||«       y # 1 sw Y   ŒoxY w# 1 sw Y   ŒsxY w)Nr¤  r3  r  F)r+   r=   r=  r7  r%  r�   r¹  rx   r9  r<  rÒ  r5  r   Úpatchr)  r½   r<  rå   rL  )rü   r,  r+  r1  r4  Úsubgraph_coderd  s          rE   r6  z%PythonWrapperCodegen.codegen_subgraphO
  sZ  € ô �7‰7×ÒØ×-Ñ-¨h¸ÀmÔTØà×!Ñ! (§.¡.Ô1Ø�‰�rÔØ�‰˜$Ÿ,™,˜ {°8·=±=°/ÐBÔCØ×$Ñ$ X¨|¸]ÔKä—w‘wˆØ%1×%=Ñ%=ˆ�‰Ô"à�>‰>×Ñ d×&FÑ&FÑFô ×$Ñ$ X§^¡^Ó4ñ @ä—\‘\Ð"3°UÓ;ñ @à'/§~¡~×'=Ñ'=Ó'?Ñ$�M 1÷@÷@ð
 ×,Ñ,×0Ñ0°·±×1DÑ1DÔEØ×,Ñ,¨]×-@Ñ-@ÔAà×"Ñ" 8¨\¸=ÕI÷@ð @ú÷@ð @ús$   ÄF3ÄF'Ä9F3Æ'F0	Æ,F3Æ3F<c                ót  — |j                  «       }| j                  |› dt        |j                  «      › �«       |j                  D �cg c]  }|j                  «       ‘Œ }}t        t        |j                  «      «      D �cg c]
  }|› d|› d�‘Œ }}| j                  |j                  ||«       y c c}w c c}w )Nú = [None] * rŽ  r�  )	rT  r�   r_   r›  rˆ  r  rÐ  r6  r,  )rü   Úinvoke_subgraphrx   r  r+  rÕ  r1  s          rE   Úcodegen_invoke_subgraphz,PythonWrapperCodegen.codegen_invoke_subgraphk
  s¥   € Ø×'Ñ'Ó)ˆà�‰˜$˜˜|¬C°×0GÑ0GÓ,HÐ+IÐJÔKØ;J×;QÑ;QÖR°C˜×-Ñ-Õ/ÐRˆÐRÜ16´s¸?×;RÑ;RÓ7SÓ1TÖU¨A˜D˜6  1 # QšÐUˆÐUØ×Ñ˜o×6Ñ6¸ÀmÕTùò SùÚUs   ÁB0ÂB5c                ó¤  — |j                  «       }|j                  D �cg c]  }|j                  «       ‘Œ }}t        t	        |j
                  «      «      D �cg c]
  }|› d|› d�‘Œ }}|j                  j                  «       }t        |j                  t        j                  «      s|› d�}| j                  |› dt	        |j
                  «      › �«       | j                  d|› d�«       | j                  t        | |j                  j                  «      «       | j                  |j                  ||«       | j                  t        | «      «       | j                  d«       | j                  t        | |j                   j                  «      «       | j                  |j                   ||«       | j                  t        | «      «       y c c}w c c}w )NrŽ  r�  r  rQ  r±   r;  zelse:)rT  Úoperandsr  rÐ  r_   r›  Ú	predicaterƒ   r   ÚShapeAsConstantBufferr�   r  Útrue_subgraphr=   r6  r.  Úfalse_subgraph)rü   Úconditionalrx   r  r+  rÕ  r1  rV  s           rE   Úcodegen_conditionalz(PythonWrapperCodegen.codegen_conditionals
  sx  € Ø×#Ñ#Ó%ˆà;F×;OÑ;OÖP°C˜×-Ñ-Õ/ÐPˆÐPÜ16´s¸;×;NÑ;NÓ7OÓ1PÖQ¨A˜D˜6  1 # QšÐQˆÐQà×)Ñ)×;Ñ;Ó=ˆ	Ü˜+×/Ñ/´×1IÑ1IÔJà$˜+ WÐ-ˆIà�‰˜$˜˜|¬C°×0CÑ0CÓ,DÐ+EÐFÔGØ�‰˜˜Y˜K qÐ)Ô*Ø�‰Ô(¨¨{×/HÑ/H×/NÑ/NÓOÔPØ×Ñ˜k×7Ñ7¸À}ÔUØ�‰Ô'¨Ó-Ô.Ø�‰�wÔØ�‰Ô(¨¨{×/IÑ/I×/OÑ/OÓPÔQØ×Ñ˜k×8Ñ8¸,ÈÔVØ�‰Ô'¨Ó-Õ.ùò! QùÚQs   ŸGÁGc                óô  — |j                  «       }|j                  D �cg c]  }|j                  «       ‘Œ }}|j                  D �cg c]  }|j                  «       ‘Œ }}| j	                  |› dt        |«      › �«       t        |«      D ]  \  }}| j	                  |› d|› d|› �«       Œ  g t        t        |«      «      D �cg c]
  }|› d|› d�‘Œ c}¢|¢}|› d�g}	t        |«      }
|
d t        |«       }| j	                  d«       | j	                  t        | |j                  j                  «      «       | j                  |j                  ||	«       | j	                  d|	d   › d	�«       | j	                  t        | «      «       | j	                  t        | |j                  j                  «      «       | j                  |j                  |
|«       | j	                  t        | «      «       y c c}w c c}w c c}w )
NrQ  rŽ  z] = r�  Ú_cond_resultzwhile True:zif not r   z: break)rT  Úcarried_inputsr  Úadditional_inputsr�   r_   râ  rÐ  r  r  Úcond_subgraphr=   r6  r.  Úbody_subgraph)rü   Ú
while_looprx   r  Úouter_carried_inputsÚouter_additional_inputsrÕ  ÚinpÚcond_outer_inputsÚcond_outer_outputsÚbody_outer_inputsÚbody_outer_outputss               rE   Úcodegen_while_loopz'PythonWrapperCodegen.codegen_while_loopˆ
  s  € Ø×"Ñ"Ó$ˆà/9×/HÑ/Hö 
Ø(+ˆC×!Ñ!Õ#ð 
Ðð  
ð 0:×/KÑ/Kö#
Ø(+ˆC×!Ñ!Õ#ð#
Ðð #
ð 	�‰˜$˜˜|¬CÐ0DÓ,EÐ+FÐGÔHÜÐ 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Ðà�‰�}Ô%Ø�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOØ×ÑØ×$Ñ$Ð&7Ð9Kô	
ð 	�‰ØÐ(¨Ñ+Ð,¨GÐ4ô	
ð 	�‰Ô'¨Ó-Ô.Ø�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOØ×ÑØ×$Ñ$Ð&7Ð9Kô	
ð 	�‰Ô'¨Ó-Õ.ùòM 
ùò#
ùò Hs   ŸG+ÁG0ÃG5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)Nr[  )rR  rƒ   rã   r+   r=   Ú
_shape_envÚ_maybe_evaluate_staticrÃ  )rª   rÅ   s     rE   Ú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 r<   )r  rn  r  )ÚlstrH  rª   Únums       rE   Ú%statically_known_list_of_ints_or_nonez:PythonWrapperCodegen.statically_known_list_of_ints_or_noneÂ
  sA   € àˆØò 	ˆAÜ&×CÑCÀAÓFˆCØˆ{ÙØ�M‰M˜#Õð		ð
 ˆrG   c                ó0   — t         j                  | «      d uS r<   )r  rr  )rp  s    rE   Ú is_statically_known_list_of_intsz5PythonWrapperCodegen.is_statically_known_list_of_intsÌ
  s   € ô !×FÑFÀsÓKÐSWÐWð	
rG   c                óH   — t         j                  | j                  «       «      S r<   )r  rr  r�  ©rT  s    rE   ra  z4PythonWrapperCodegen.static_shape_for_buffer_or_noneÒ
  s   € ä#×IÑIØ�O‰OÓó
ð 	
rG   c                ó0   — t         j                  | «      d uS r<   )r  ra  rv  s    rE   Ú!can_prove_buffer_has_static_shapez6PythonWrapperCodegen.can_prove_buffer_has_static_shapeØ
  s   € ä#×CÑCÀFÓKÐSWÐWÐWrG   r<   )rÜ  rä   rÝ  r·   rÞ  úOptional[PythonWrapperCodegen]rß  ú$Optional[ir.GraphPartitionSignature]r*  )rx   r¶   rÙ  r¶   r³   r  )r¬  r¶   )rþ  ÚTritonMetaParamsr³   r¶   ©r³   ú	list[str]©r³   z>dict[str, Union[ir.TensorBox, ir.TorchBindObject, sympy.Expr]]©r³   zlist[IRNode])r  r}  rŠ  )r6  rã   r³   r¶   rW  )r6  rã   r³   r  )rD  r}  r³   r  ©rH  r0   r³   r  )ry   r¶   rf  r¶   rg  r·   rz   r}  r“  r¶   r³   r  )F)NNN)
r—  r¶   rŠ  r¶   r‰  r¶   r˜  r}  r™  zOptional[torch._ops.OpOverload])rx   r¶   rå   zir.TensorBoxrä  zOrderedSet[sympy.Symbol])rí  zsympy.Symbol)rª   r
   rB   rä   r³   r¶   )rª   r
   r³   r¶   )rú  r¶   rx   r¶   rš  r¶   r³   r¶   )r”  zSequence[Expr]r³   r¶   )r�   zCallable[..., None]r³   r¶   )r	  rä   )NTN)
ra  r¶   r3  r¶   r4  r·   r5  rä   r6  r·   )r;  r¶   )rÀ   z"list[list[Union[int, sympy.Expr]]])ra  r¶   r®  r·   )r”  r2   )NN)ra  r¶   )rT  rZ  )r¤  )rT  z%Union[BufferLike, ir.TorchBindObject])rõ  r}  )rð  r¶   rï  r¶   rù  r¶   )rý  rZ  rþ  rZ  rt  rä   )rx   r¶   r  zir.ReinterpretViewr³   r  ©rT  r{  )r  r{  r  r{  )r_  r¶   r›  r   r%  z,Optional[dict[sympy.Symbol, pytree.KeyPath]]r³   r  )rA  rã   rß  zir.GraphPartitionSignature)rG  rã   ){ræ   rÙ   rþ   Ú__doc__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Î  r4  r%  r2  r!  r0  r=  r?  rA  rE  rI  rM  rO  rS  rQ  ri  r‚  r…  rŒ  r•  rœ  r   r¢  rž  r±  r«  r¬  rå  r)  rï  r²  rô  rö  rø  rû  r–   r™  r  r  r
  r  r  r.  r³  r9  r<  rˆ  r‹  r–  r˜  r�  r�  r©  rª  r¯  r®  r°  r©  r�   rÂ  rÑ  r�   r%  r§  rt  ri  r“  rñ  rq  rö  rú  rw  r  r  r  r  r  r  r  r&  r7  r9  rD  rH  rL  r6  rS  r[  rj  rn  rr  rt  ra  rx  r  r  s   @rE   r  r  Œ  s�  ø„ ñôP#ðd ð
 FJð	&Øð&à$ð&ð 7ð&ð Cò	&ó ð&ó'ó<ó;Aózó
ð  òó ðð" òó ðó	!ð ò
ó ð
ó
ð$à	Gó$ó
%óSó$	(ó
ó.óó8)ó6ó(-ôò.ò1ò0òJò
/ó+ó7ó&5ó
5óóóòAòð:;àð;ð ð;ð  ð	;ð
 ð;ð ð;ð 
ó;óò,ò
ò&<ð 8<ØØð
Xàð
Xð  ð
Xð ð	
Xð
  ð
Xð 5ó
Xò0óòM
ò^!SòF:ò

ðB(Kàð(Kð ð(Kð -ó	(KòTJó"5òð @Dõ Wð CGõ +ó.ó&ó'ó6ó
ð( ðð 'ðð 
óó:>òGò4ò*[Yòzð, #'ØØ(,ð!àð!ð ð!ð  ð	!ð
 ð!ð &ó!ó&2ðC;ð 2óC;ôJ1ó%óN:ò.ò>ò8ò5òWò
ò
ò4
4óQGòf'ð ØØØØñ_Xàó_XòB ò!ò,ó óF
ð DHóó@_ó*óAóuó`ó 
ó&3òP:ó(
ò 
óEò'ðNàðNð ðNð Hð	Nð
 
óNò`+'òZ	Xð<àð<ð 9ó<ó>Yò
ò*Jò8Uò/ò*(/ðT ñó ðð ñó ðð ñ
ó ð
ð
 ñ
ó ð
ð
 ñXó ôXrG   r  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edd„«       Zed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                óL   •— || _         || _        || _        t        ‰| �  «        y r<   )rÝ  rÞ  rß  r  r  )rü   rÝ  rÞ  rß  r
  s       €rE   r  z%SubgraphPythonWrapperCodegen.__init__å
  s(   ø€ ð +ˆÔØ,ˆÔØ$8ˆÔ!ä‰ÑÕrG   c                ó&   — | j                   | _        y r<   )rÝ  rÀ  rû   s    rE   rÁ  z1SubgraphPythonWrapperCodegen.set_launcher_fn_nameó
  s   € ð !%× 2Ñ 2ˆÕrG   c                 ó   — y r<   rŠ   rû   s    rE   rÄ  z)SubgraphPythonWrapperCodegen.write_headerø
  rô  rG   c                 ó   — y r<   rŠ   r1  s     rE   r³  z2SubgraphPythonWrapperCodegen.add_benchmark_harnessû
  rô  rG   c                 ó   — y r<   rŠ   r1  s     rE   r.  z6SubgraphPythonWrapperCodegen.benchmark_compiled_moduleþ
  rô  rG   c                 ó   — y r<   rŠ   rû   s    rE   r  z5SubgraphPythonWrapperCodegen.write_async_compile_wait  rô  rG   c                ó6   — | j                   j                  «       S r<   )rÞ  r=  rû   s    rE   r=  z/SubgraphPythonWrapperCodegen.next_kernel_suffix  s   € à×"Ñ"×5Ñ5Ó7Ð7rG   c                 ó   — y r<   rŠ   rG  s     rE   rM  z2SubgraphPythonWrapperCodegen.generate_after_suffix  r  rG   c                ó\   — | j                   j                  d| j                  › d�«       d}|S )Nz
            def z(args):
            r,   )r­  rá   rÀ  r  s     rE   r!  z>SubgraphPythonWrapperCodegen.write_launcher_fn_call_get_indent  s<   € Ø�‰×ÑðØ×&Ñ&Ð'ð (ðô	
ð
 ˆØÐrG   c                 ó   — yr¦   rŠ   rû   s    rE   r¢  z4SubgraphPythonWrapperCodegen.get_wrapper_call_indent  s   € ØrG   c                ór   — | j                   x}r|j                  }|S t        j                  j                  }|S r<   )rß  Úinput_nodesr+   r=   r
  )rü   rL  rˆ  s      rE   r  z-SubgraphPythonWrapperCodegen.get_graph_inputs  s=   € ð ×1Ñ1Ð1ˆ9Ð1Ø×*Ñ*ˆFð ˆô —W‘W×)Ñ)ˆFØˆrG   c                ó    — | j                   x}r%t        |j                  j                  «       «      }|S t        j
                  j                  }|S r<   )rß  r  r�  r+  r+   r=   r  )rü   rL  Únamess      rE   r#  z2SubgraphPythonWrapperCodegen.get_graph_input_names   sI   € Ø×1Ñ1Ð1ˆ9Ð1Ü˜×.Ñ.×3Ñ3Ó5Ó6ˆEð ˆô —G‘G×-Ñ-ˆEØˆrG   c                ór   — | j                   x}r|j                  }|S t        j                  j                  }|S r<   )rß  r@  r+   r=   r  )rü   rL  r›  s      rE   r  z.SubgraphPythonWrapperCodegen.get_graph_outputs'  s;   € Ø×1Ñ1Ð1ˆ9Ð1Ø×,Ñ,ˆGð ˆô —g‘g×+Ñ+ˆGØˆrG   c                ó~   •— |j                  «       }| j                  x}r||j                  v ry t        ‰| �  |«       y r<   )rT  rß  r�  r  r  )rü   rT  rx   rL  r
  s       €rE   r  z/SubgraphPythonWrapperCodegen.codegen_allocation.  s?   ø€ Ø�‰Ó ˆØ×2Ñ2Ð2ˆIÐ2¸À	×@UÑ@UÑ8Uð ä‰Ñ" 6Õ*rG   c                ó8   — | j                   j                  «        y r<   )rÞ  rú  rû   s    rE   rú  z5SubgraphPythonWrapperCodegen.write_triton_header_once8  s   € ð 	×Ñ×4Ñ4Õ6rG   c                ó8   — | j                   j                  «        y r<   )rÞ  rü  rû   s    rE   rü  z=SubgraphPythonWrapperCodegen.write_get_raw_stream_header_onceA  s   € ð 	×Ñ×<Ñ<Õ>rG   r<   )rÝ  r¶   rÞ  r  rß  rz  r*  rW  r€  rŠ  r~  r|  r  r�  )ræ   rÙ   rþ   r‚  r  rÁ  rÄ  r³  r.  r  r=  rM  r!  r¢  r  r#  r  r  r$   rú  rü  r  r  s   @rE   rÛ  rÛ  Ý
  s¢   ø„ ñð FJð	àðð -ðð Cõ	ó3ó
òòòó8óóóðà	Góóóõ+ð ò7ó ð7ð ò?ó ô?rG   rÛ  )rC   rZ  r³   r  )rK   rZ  rL   rZ  )ra   útorch.Argumentr³   r¶   )rl   r–  r³   r¶   )ry   ztorch._ops.OpOverloadr³   r¶   r<   )
rx   r¶   r¿   zlist[triton.Config]rÀ   zlist[TritonGrid]r’   ry  r³   ztuple[str, str]rW  )†Ú
__future__r   r  r¹   rP  rÚ   rb  rw  r»  rd  r›  r]   r½  Ú	itertoolsr   Útypingr   r   r   r   r	   r„   r
   ræ  Ú
torch._opsÚtorch.utils._pytreeÚutilsÚ_pytreer  r   r¶  Útorch._dynamo.utilsr   r   Ú#torch._inductor.codegen.debug_utilsr   Ú$torch._inductor.codegen.multi_kernelr   Ú%torch._inductor.runtime.runtime_utilsr   Ú%torch.fx.experimental.symbolic_shapesr   r   r   r   r   Útorch.fx.noder   Útorch.utils._ordered_setr   Ú torch.utils._sympy.singleton_intr   Útorch.utils._sympy.symbolr   r   r¤  r   r   r   Ú	codecacher   r    r!   Úruntimer"   Úruntime.hintsr#   r$   r%   r&   r'   r(   r)   r*   Úvirtualizedr+   rg  r-   r.   r/   r0   r1   r2   r3   Ú	cpp_utilsr4   Útriton_utilsr5   r6   r7   Úcollections.abcr8   r9   r—   r=   r:   Údoprintrì  r•   r“  r¶   r  rS  rZ  rF   rO   ri   rn   r€   rm  rã   r{  rµ   rÈ   rõ   rà  r÷   r   r  r  r.  r5  rD  rH  rY  r  r`  r^  rz  r�  r   rÉ  r  rÛ  rŠ   rG   rE   ú<module>r¯     sr  ðå "ã Û Û Û 
Û Û Û Û Û Û 	Û Ý ß @Õ @ã Ý ã Û ß $Ð $Ý &ß 6Ý CÝ AÝ ;÷õ õ .Ý /Ý 9ß :ç (Ñ (Ý 'ß (Ý 'Ý ,÷÷ ñ õ ÷÷ ñ õ ß PÑ Pñ ß2ãå%ñ 	‹×Ñ€ð �—‘˜uŸ{™{¨CÐ/Ñ0€Ø�2—9‘9˜lÐ*Ñ+€
ó	óó>Bó>ó$<ð$ ˜˜S˜‘>Ð ØØ	ˆ%��U—Z‘Z�Ñ
  #Ð
%Ñ&¨Ð2BÐ1CÀUÈ3ÐPSÈ8Á_Ð1TÑ(UÐUñ€
ð /3ð	U&Ø
ðU&à ðU&ð ðU&ð ,ð	U&ð
 óU&ópJ&ðZ ×Ñ÷ð ó ð÷*ñ *÷*	ñ 	ð ×Ñô	˜ó 	ó ð	ð ×Ñô�{ó ó ðð ×ÑôK Kó Kó ðKôD ;ô ð ×Ñô;˜ó ;ó ð;ð2 ×ÑôÐ%ó ó ðð: ×ÑôEÐ,ó Eó ðEð, ×Ñô
Ð"ó 
ó ð
ô(	Ð!ô 	ð ×Ñô!�[ó !ó ð!ð: ×Ñô&˜^ó &ó ð&ðR ×ÑôN˜ó Nó ðNð €
ôN!X˜7ô N!XôbBl?Ð#7õ l?rG   