ó
    EñiGx ã                  óž	  • S SK Jr  S SKrS SKrS SKrS SKrS SKrS SKrS SKrS SK	r	S SK
r
S SKrS SKrS SKrS SKJr  S SKJrJr  S SKJrJrJrJr  S SKrS SKJr  S SKrS SKrS SKJs  Jr  S SKJ r!  S SK"J#r#J$r$J%r%  S S	K&J'r'  S S
K(J)r)  S SK*J+r+  S SK,J-r-J.r.  S SK/J0r0  S SK1J2r2J3r3J4r4J5r5J6r6  S SK7J8r8  S SK9J:r:  S SK;J<r<  S SK=J>r>J?r?  SSK@JArAJBrBJCrC  SSKDJErE  SSKCJFrFJGrG  SSKHJIrI  SSKJJKrK  SSKJLrLJMrMJNrNJOrOJPrPJQrQJRrRJSrSJTrTJUrUJVrVJWrW  SSKXJYrY  SSKZJ[r[J\r\J]r]J^r^J_r_J`r`  SSKaJbrb  SSKcJdrd  SSKeJfrfJgrgJhrh  \(       a$  S S KJiriJjrj  S SKkrkSS!KlJmrm  SS"KCJnrn  SS#KoJprp  SS$KqJrrr  \Ræ                  " \t5      ru\^" 5       Rì                  rw\x\Rò                  \R@                  \z\{4   r|\x\Rò                  \R@                  \zS%\z4   r}\\CRü                  \_4   r\S&/S4   r€SgS' jr�ShS( jr‚SiS) jrƒS* r„\…\z\†4   r‡\\x\\†\R2                  4   S+4   \\‡/\x\†S+4   4   4   rˆ  Sj           SkS, jjr‰SlS- jrŠ\GR                   " S. S/5      5       rŒ " S0 S15      r� " S2 S&5      rŽ\GR                   " S3 S4\Ž5      5       r�\GR                   " S5 S6\Ž5      5       r�\GR                   " S7 S8\Ž5      5       r‘\GR                   " S9 S:\Ž5      5       r’\GR                   " S; S<\Ž5      5       r“\GR                   " S= S>\Ž5      5       r” " S? S@\Ž5      r•\GR                   " SA SB\Ž5      5       r–\GR                   " SC SD\Ž5      5       r—\GR                   " SE SF\Ž5      5       r˜\GR                   " SG SH\Ž5      5       r™\GR                   " SI SJ\Ž5      5       rš\GR                   " SK SL\Ž5      5       r› " SM SN5      rœ\GR                   " SO SP\›5      5       r�\GR                   " SQ SR\›5      5       rž\GR                   " SS ST\›5      5       rŸ\GR                   " SU SV\›5      5       r  " SW SX\›5      r¡\GR                   " SY SZ\Ž5      5       r¢\GR                   " S[ S\\Ž5      5       r£\GR                   " S] S^\Ž5      5       r¤\GR                   " S_ S`\Ž5      5       r¥\GR                   " Sa Sb\Ž5      5       r¦\zr§\\›\S4   r¨ " Sc Sd\\5      r© " Se Sf\©5      rªg)mé    )ÚannotationsN)ÚCallable)ÚchainÚcount)ÚAnyÚOptionalÚTYPE_CHECKINGÚUnion)ÚExpr)Údtype)ÚcountersÚdynamo_timedÚget_debug_dir)ÚDebugPrinterManager)ÚMultiKernelState)Ú	cache_dir)Úget_opaque_obj_reprÚis_opaque_value_type)Útrace_structured)ÚCallMethodKeyÚConvertIntKeyÚDivideByKeyÚresolve_unbacked_bindingsÚSymTypes)Ú_get_qualified_name)Ú
OrderedSet)ÚSingletonInt)Úsymbol_is_typeÚSymTé   )Úasync_compileÚconfigÚir)Úoutput_code_log)ÚIRNodeÚReinterpretView)Útriton_heuristics)ÚDeviceProperties)Úcache_on_selfÚDelayReplaceLineÚget_benchmark_nameÚget_dtype_sizeÚIndentedBufferÚ#is_codegen_graph_partition_subgraphÚis_using_cudagraph_partitionÚLineContextÚsympy_productÚ	sympy_strÚ
sympy_subsÚtriton_version_uses_attrs_dict)ÚVé   )ÚArgNameÚCodeGenÚDeferredLineÚPythonPrinterÚWorkspaceArgÚWorkspaceZeroMode)Úcexpr)ÚCUSTOM_EXTERN_KERNEL_CODEGEN)Ú	config_ofÚshould_unwrap_unspec_argÚsignature_to_meta)ÚIteratorÚSequence)ÚGraphLowering)ÚExternKernel)ÚBaseSchedulerNode)ÚFxConverterzir.CommBufferTypeÚWrapperLinec                ó:  • [         R                  R                  U 5      nU R                  5       [         R                  R                  ;  nU R                  5       U R                  5       [        [         R                  R                  R                  U5      5      U4$ ©N)
r5   ÚgraphÚget_allocation_storage_sizeÚget_nameÚunaligned_buffersÚget_device_or_errorÚ	get_dtyper2   ÚsizevarsÚsimplify)ÚnodeÚstorage_sizeÚ	alignments      Ú\/home/mande/repo/quber/.venv/lib/python3.13/site-packages/torch/_inductor/codegen/wrapper.pyÚbuffer_reuse_keyrW   d   sr   € Ü—7‘7×6Ñ6°tÓ<€LØ—‘“¤q§w¡w×'@Ñ'@Ñ@€Ià× Ñ Ó"Ø�‰Óô 	”!—'‘'×"Ñ"×+Ñ+¨LÓ9Ó:Øðð ó    c                óp  • [         R                  R                  U 5      nU R                  5       n[	        U[
        R                  5      (       d   eU R                  5       U R                  5       [        [         R                  R                  R                  U5      5      UR                  UR                  4$ rJ   )r5   rK   rL   Úget_output_specÚ
isinstancer#   ÚCommBufferLayoutrO   rP   r2   rQ   rR   Úcomm_buffer_typeÚ
group_name)rS   rT   Úlayouts      rV   Úcomm_buffer_reuse_keyr`   r   sŒ   € ä—7‘7×6Ñ6°tÓ<€LØ×!Ñ!Ó#€FÜ�fœb×1Ñ1×2Ñ2Ð2Ð2à× Ñ Ó"Ø�‰ÓÜ”!—'‘'×"Ñ"×+Ñ+¨LÓ9Ó:Ø×ÑØ×Ñðð rX   c                óš  • U R                  5       UR                  5       :w  a  gU R                  5       UR                  5       :w  a  g[        R                  R                  R                  [        R                  R                  U 5      5      n[        R                  R                  R                  [        R                  R                  U5      5      n[        U5      [        U5      :X  d`  [        R                  R                  R                  USU-  5      (       a/  [        R                  R                  R                  X25      (       a  gg)NFgffffffî?T)
rO   rP   r5   rK   rQ   rR   rL   r2   Ústatically_known_geqÚstatically_known_leq)Ú	input_bufÚ
output_bufÚ
input_sizeÚoutput_sizes       rV   Úcan_match_buffer_sizerh   €   sñ   € ð ×$Ñ$Ó&¨*×*HÑ*HÓ*JÓJØà×ÑÓ 
× 4Ñ 4Ó 6Ó6Øä—‘×!Ñ!×*Ñ*Ü	�‰×+Ñ+¨IÓ6ó€Jô —'‘'×"Ñ"×+Ñ+Ü	�‰×+Ñ+¨JÓ7ó€Kô 	�*Ó¤¨;Ó!7Ó7ô 	
�‰×Ñ×-Ñ-¨k¸4À*Ñ;L×MÑMÜ�G‰G×Ñ×1Ñ1°+×JÑJààrX   c                óH  • [        U [        R                  5      (       a?  U R                  5       nUR                  UR
                  UR                  UR                  S4$ / nU n[        U[        R                  [        R                  [        R                  45      (       aq  UR                  5       nUc  gUR                  U5        UR                  n[        U[        R                  [        R                  [        R                  45      (       a  Mq  [        U[        R                  5      (       d  gU H-  nUR                  UR                  5       R                  :w  d  M-    g   UR                  5       nUR                  UR
                  UR                  UR                  S4$ )zí
Collapse a chain of ReinterpretView <- StorageBox
<- ReinterpretView <- StorageBox.... <- buffer wrappers if every layer
has the same offset as the innermost (base) buffer.

Returns:
    (size, stride, offset, dtype, collapsible: bool)
T)NNNNF)r[   r#   ÚBufferÚ
get_layoutÚsizeÚstrideÚoffsetr   Ú	TensorBoxÚ
StorageBoxr&   ÚappendÚdata)rr   ÚlayÚlayoutsÚcurÚbase_lays        rV   Úcodegen_reinterpret_view_helperrw   Ÿ   s(  € ô �$œŸ	™	×"Ñ"Ø�o‰oÓˆØ�x‰x˜Ÿ™ S§Z¡Z°·±¸DÐ@Ð@à€GØ
€CÜ
�Sœ2Ÿ<™<¬¯©¼×8JÑ8JÐK×
LÑ
LØ�n‰nÓˆØ‰;Ø0Ø�‰�sÔØ�h‰hˆô �Sœ2Ÿ<™<¬¯©¼×8JÑ8JÐK×
LÓ
Lô �cœ2Ÿ9™9×%Ñ%Ø,ó ˆØ�:‰:˜Ÿ™Ó)×0Ñ0Õ0Ù0ñ ð �~‰~Ó€HØ�=‰=˜(Ÿ/™/¨8¯?©?¸H¿N¹NÈDÐPÐPrX   .c                óR  ^ ^^^• [        5       mSS jm S   SUU4S jjjnSSU UU4S jjjnST  3nU" SU S35        T(       a9  [        R                  R                  (       a  TR                  R                  5       O[        R                  " 5       nTR                  5          U   [        R                  R                  (       ab  U(       a[  [        R                  R                  (       a<  U[        R                  R                  ;   a  [        R                  R                  U   n	OS /[        U5      -  n	[        U5      S:X  a!  U" US   U	S   5      u  p«U" S	U
 3S	U 35        Oü[        U5      S:”  d   e[        U5      [        U5      :X  d   e[        5       n[        [        X!U	5      S
 SS9 H­  u  p­n/ nUR                  (       a?  UR                   H/  nUS;  d  M  UR!                  SU SUR                  U    35        M1     U(       a  SR#                  U5      nOSnU" X«5      u  p«SU SU
 3nUU;   a  M�  UR%                  U5        U" USU SU 35        M¯     S S S 5        S S S 5        UTR'                  5       4$ ! , (       d  f       N(= f! , (       d  f       N1= f)Nc                óp   • [        U [        R                  5      (       a  U $ [        R                  " U 5      $ rJ   )r[   Úsympyr   ÚInteger)Úitems    rV   Ú_convert_to_sympy_exprÚ@user_defined_kernel_grid_fn_code.<locals>._convert_to_sympy_exprÑ   s&   € Ü! $¬¯
©
×3Ñ3ˆtÐL¼¿ºÀtÓ9LÐLrX   c                ó  >• Tb  [        U 5      (       a  X 4$ [        U4S jU  5       5      nU(       d  UnTR                  U5      [        R                  R
                  (       a%  TR                  [        U4S jU 5       5      5      4$ S4$ )zÿ
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  ó4   >#   • U  H  nT" U5      v •  M     g 7frJ   © )Ú.0Úgr}   s     €rV   Ú	<genexpr>ÚKuser_defined_kernel_grid_fn_code.<locals>.determine_grid.<locals>.<genexpr>â   s   øé € ÐCºd¸Ñ1°!×4Ð4ºdùs   ƒc              3  óZ   >#   • U  H   nTR                  U[        U5      5      v •  M"     g 7frJ   ©Úgenerate_example_arg_valueÚtype)r‚   rƒ   Úwrappers     €rV   r„   r…   é   s,   øé € ð â!-˜Að  ×:Ñ:¸1¼dÀ1»g×FÐFÚ!-ùó   ƒ(+)ÚcallableÚtupleÚcodegen_python_shape_tupler"   ÚtritonÚautotune_at_compile_time)ÚgridÚexample_gridÚ
sympy_gridr}   rŠ   s      €€rV   Údetermine_gridÚ8user_defined_kernel_grid_fn_code.<locals>.determine_gridÔ   s�   ø€ ð ‰?œh tŸn™nà�:ÐäÔC¹dÓCÓCˆ
ÞØ%ˆLà×.Ñ.¨zÓ:ô —=‘=×9×9ð ×2Ñ2Üô á!-óó óð
ð 	
ð ð
ð 	
rX   c                óâ   >• TR                  U 5        T(       aV  [        R                  R                  (       a6  TTR                  ;  a%  TR
                  R                  U=(       d    U 5        g g g g rJ   )Ú	writeliner"   r�   r�   Úkernel_autotune_namesÚkernel_autotune_calls)Úliner’   ÚnameÚoutputrŠ   s     €€€rV   r—   Ú3user_defined_kernel_grid_fn_code.<locals>.writelineó   sW   ø€ Ø×Ñ˜ÔæÜ—‘×6×6Ø˜G×9Ñ9Ó9à×)Ñ)×3Ñ3°L×4HÀDÕIð :ð 7ð rX   Úgrid_wrapper_for_údef z(meta):r6   r   zreturn c                ó2   • [        U S   R                  5      $ ©Nr6   ©ÚlenÚkwargs©Úxs    rV   Ú<lambda>Ú2user_defined_kernel_grid_fn_code.<locals>.<lambda>  s   € œc ! A¡$§+¡+Ô.rX   T©ÚkeyÚreverse)Úmatrix_instr_nonkdimÚwaves_per_euÚkpackzmeta['z'] == z and ÚTrueúif z	: return )r|   úUnion[int, sympy.Expr]Úreturnú
sympy.ExprrJ   )r‘   Ú
TritonGridr’   zOptional[TritonGrid])rš   Ústrr’   úOptional[str])r-   r"   r�   r�   r™   ÚindentÚ
contextlibÚnullcontextr5   rK   Úautotuning_gridsr£   r   ÚsortedÚzipr¤   rq   ÚjoinÚaddÚgetvalue)r›   ÚconfigsÚgridsrŠ   Úoriginal_fxnode_namer”   r—   Úfn_nameÚkernel_autotune_calls_indentÚexample_gridsr‘   r’   ÚseenÚcÚ
guardslistÚkwargÚguardsÚ	statementr}   rœ   s   `  `              @@rV   Ú user_defined_kernel_grid_fn_coderÌ   È   sb  û€ ô Ó€FôMð
 .2ð
Øð
à*÷
ñ 
÷>Jó Jð " $ Ð(€GÙ��W�I˜WÐ%Ô&ö ”v—}‘}×=×=ð 	×%Ñ%×,Ñ,Ô.ä×#Ò#Ó%ð !ð
 
�‰�Ò6ä�M‰M×2×2Þ$Ü—‘×(×(Ø$¬¯©×(@Ñ(@Ó@äŸG™G×4Ñ4Ð5IÑJ‰Mà!˜F¤S¨£ZÑ/ˆMÜˆu‹:˜‹?Ù!/°°a±¸-ÈÑ:JÓ!KÑˆDÙ˜ ˜vÐ&¨'°,°Ð(@ÕAä�u“: “>Ð!�>Ü�u“:¤ W£Ó-Ð-Ð-Ü$.£LˆDô *0Ü�E MÓ2Ù.Øô*Ñ%�˜ð
  �
Ø—8—8à!"§¤˜Ø ð )õ ð
 '×-Ñ-°°u°g¸VÀAÇHÁHÈUÁOÐCTÐ.UÖVñ "*ö Ø$Ÿ\™\¨*Ó5‘Fà#�FÙ%3°DÓ%GÑ"�Ø! & ¨°4°&Ð9�	Ø Ó$ÙØ—‘˜Ô#Ù˜) s¨6¨(°)¸L¸>Ð%JÖKñ1*÷- 7�ð` �F—O‘OÓ%Ð%Ð%÷a 7Õ6ú��ús,   Â%JÂ(D8JÇ$BJÉ%JÊ
J	ÊJÊ
J&c                óè   ^^^^^^• [        5       mTR                  U R                  SS9  SSKmSSKJm  SSKJm  [        U R                  /5      mUUUUUU4S jmT" U 5        TR                  5       $ )z[
Given a triton kernel function pointer collect the transitive closure of
its dependencies
T©Ústripr   N)ÚJITFunction)Ú	constexprc           	     ó\  >• [        S [        R                  " U R                  5       5       5      nU R                  R                  R                  S0 5      nU R                  R                  R                   GH°  nUT;   a  M  X0R                  R                  ;   d  M'  U R                  R                  U   n[        UT5      (       aV  T	R                  5         T	R                  S5        T	R                  UR                  SS9  TR                  U5        T" U5        M§  [        TS5      (       a†  [        UTR                  R                   R"                  5      (       aW  T	R                  5         T	R                  S5        T	R                  UR                  SS9  TR                  U5        T" U5        GM>  [        U[$        [&        [(        T
45      (       aÊ  T	R                  5         [        UT
5      (       a  SUR*                  < S	3nOU< nUR                  U5      =n(       aQ  [        U[,        5      (       a  S
UR.                   SUR0                   3nOS
U< 3nT	R                  U U SU 35        OT	R                  U SU 35        TR                  U5        GM)  X1;   d  GM1  US:w  d  GM:  [        US5      (       d  GMN  UR.                  R3                  S5      (       d  GMq  T	R                  SUR.                   SUR0                   SU 35        TR                  U5        GM³     g )Nc              3  ó^   #   • U  H#  nUR                   S :X  d  M  UR                  v •  M%     g7f)ÚLOAD_GLOBALN)ÚopnameÚargval)r‚   Úinsts     rV   r„   Ú^user_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse.<locals>.<genexpr>K  s)   é € ð '
â3�Ø�{‰{˜mÑ+ó ˆD�KŽKÚ3ùs   ‚-š-Ú__annotations__z@triton.jitTrÎ   Úconstexpr_functionz@triton.constexpr_functionztl.constexpr(Ú)ú: Ú.ú = ÚtlÚ
__module__r�   zfrom z import z as )r   ÚdisÚBytecodeÚfnÚ__globals__ÚgetÚ__code__Úco_namesr[   Únewliner—   ÚspliceÚsrcr¾   ÚhasattrÚruntimeÚjitÚConstexprFunctionÚintrµ   ÚboolÚvaluer‰   rà   Ú__name__Ú
startswith)Ú
cur_kernelÚunqualified_loadsÚglobal_annotationsÚsymbol_nameÚsymbolÚ
symbol_strÚ
annotationÚannotation_coderÐ   Úcompile_wrapperrÑ   Úsymbols_includedÚtraverser�   s           €€€€€€rV   rþ   ÚKuser_defined_triton_kernel_transitive_closure_source_code.<locals>.traverseF  sÃ  ø€ ô
 'ñ '
äŸš Z§]¡]Ô3ó'
ó 
Ðð
 (Ÿ]™]×6Ñ6×:Ñ:Ð;LÈbÓQÐØ%Ÿ=™=×1Ñ1×:Õ:ˆKØÐ.Ó.ÙØŸm™m×7Ñ7Õ7Ø#Ÿ™×2Ñ2°;Ñ?�Ü˜f k×2Ñ2Ø#×+Ñ+Ô-Ø#×-Ñ-¨mÔ<à#×*Ñ*¨6¯:©:¸TÐ*ÑBØ$×(Ñ(¨Ô5Ù˜VÖ$Ü˜VÐ%9×:Ñ:¼zØØ—N‘N×&Ñ&×8Ñ8÷@ñ @ð $×+Ñ+Ô-Ø#×-Ñ-Ð.JÔKØ#×*Ñ*¨6¯:©:¸TÐ*ÑBØ$×(Ñ(¨Ô5Ù˜V×$Ü ¬¬c´4¸Ð(C×DÑDØ#×+Ñ+Ô-Ü! &¨)×4Ñ4Ø'4°V·\±\Ñ4DÀAÐ%F™
à(.¡z˜
Ø%7×%;Ñ%;¸KÓ%HÐH�zÕHÜ% j´$×7Ñ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Ø# tÖ+Ü ¨×5Ô5ð ×)Ñ)×4Ñ4°X×>Ô>ð
 $×-Ñ-Ø × 1Ñ 1Ð2°(¸6¿?¹?Ð:KÈ4ÐP[È}Ð]ôð %×(Ñ(¨×5òo ;rX   )
r-   ré   rê   r�   rÐ   Útriton.languagerÑ   r   rò   r¿   )ÚkernelrÐ   rü   rÑ   rý   rþ   r�   s    @@@@@@rV   Ú9user_defined_triton_kernel_transitive_closure_source_coder  6  si   ý€ ô
 %Ó&€OØ×Ñ˜6Ÿ:™:¨TÐÑ2ó Ý"Ý)ô " 6§?¡?Ð"3Ó4Ð÷B6ò B6ñH ˆVÔØ×#Ñ#Ó%Ð%rX   c                  ó0   • \ rS rSr% S\S'   S\S'   S rSrg)	ÚSymbolicCallArgiŽ  úsympy.SymbolÚinnerr³   Ú
inner_exprc                ó,   • [        U R                  5      $ rJ   )rµ   r  ©Úselfs    rV   Ú__str__ÚSymbolicCallArg.__str__”  s   € Ü�4—:‘:‹ÐrX   r�   N)rò   rà   Ú__qualname__Ú__firstlineno__rÙ   r  Ú__static_attributes__r�   rX   rV   r  r  Ž  s   ‡ àÓàÓõrX   r  c                  óp   ^ • \ rS rSrU 4S jrS
S jrSS jrSS jrSS jrSS jr	      SS jr
S	rU =r$ )ÚMemoryPlanningStatei˜  c                ó¬   >• [         TU ]  5         [        R                  " [        5      U l        [        R                  " [        5      U l        SU l        g ©Nr   )ÚsuperÚ__init__ÚcollectionsÚdefaultdictÚlistÚ
reuse_poolÚcomm_buffer_reuse_poolÚtotal_allocated_buffer_size)r
  Ú	__class__s    €rV   r  ÚMemoryPlanningState.__init__™  sG   ø€ Ü‰ÑÔô ×#Ò#¤DÓ)ð 	Œô ×#Ò#¤DÓ)ð 	Ô#ð 12ˆÕ(rX   c                óL   • [        U R                  R                  US 5      5      $ rJ   )rð   r  rå   ©r
  rª   s     rV   Ú__contains__Ú MemoryPlanningState.__contains__¥  s   € Ü�D—O‘O×'Ñ'¨¨TÓ2Ó3Ð3rX   c                óf   • U R                   U   R                  5       nUR                  (       a   eU$ rJ   )r  ÚpopÚ	is_reused©r
  rª   r|   s      rV   r#  ÚMemoryPlanningState.pop¨  s+   € Ø�‰˜sÑ#×'Ñ'Ó)ˆØ—>—>Ð!Ð!ØˆrX   c                óf   • UR                   (       a   eU R                  U   R                  U5        g rJ   )r$  r  rq   r%  s      rV   ÚpushÚMemoryPlanningState.push­  s&   € Ø—>—>Ð!Ð!Ø�‰˜Ñ×#Ñ# DÕ)rX   c                óL   • [        U R                  R                  US 5      5      $ rJ   )rð   r  rå   r  s     rV   Úcomm_buffer_containsÚ(MemoryPlanningState.comm_buffer_contains±  s    € Ü�D×/Ñ/×3Ñ3°C¸Ó>Ó?Ð?rX   c                óf   • U R                   U   R                  5       nUR                  (       a   eU$ rJ   )r  r#  r$  r%  s      rV   Úcomm_buffer_popÚ#MemoryPlanningState.comm_buffer_pop´  s-   € Ø×*Ñ*¨3Ñ/×3Ñ3Ó5ˆØ—>—>Ð!Ð!ØˆrX   c                óf   • UR                   (       a   eU R                  U   R                  U5        g rJ   )r$  r  rq   r%  s      rV   Úcomm_buffer_pushÚ$MemoryPlanningState.comm_buffer_push¹  s*   € ð —>—>Ð!Ð!Ø×#Ñ# CÑ(×/Ñ/°Õ5rX   )r  r  r  )rª   ÚReuseKeyr²   rð   )rª   r3  r²   ÚFreeIfNotReusedLine)rª   r3  r|   r4  r²   ÚNone)rª   ÚCommBufferReuseKeyr²   rð   )rª   r6  r²   r4  )rª   r6  r|   r4  r²   r5  )rò   rà   r  r  r  r   r#  r(  r+  r.  r1  r  Ú__classcell__©r  s   @rV   r  r  ˜  sD   ø† õ
2ô4ôô
*ô@ôð
6Ø%ð6Ø-@ð6à	÷6ò 6rX   r  c                  ó   • \ rS rSrSS jrSrg)rH   iÀ  c                ó0   • [        S[        U 5       35      e)Nz&FX codegen not yet supported for type )ÚNotImplementedErrorr‰   ©r
  Ú	converters     rV   Ú
codegen_fxÚWrapperLine.codegen_fxÁ  s   € Ü!Ð$JÌ4ÐPTË:È,Ð"WÓXÐXrX   r�   N©r=  rG   r²   ÚFxConversionFunc©rò   rà   r  r  r>  r  r�   rX   rV   rH   rH   À  s   † ÷YrX   c                  óH   • \ rS rSr% S\S'   S\S'   SS jrSS jrSS jrS	rg
)ÚEnterSubgraphLineiÅ  ÚPythonWrapperCodegenrŠ   rD   rK   c                ób   • U R                   R                  U R                   R                  5        g rJ   )rŠ   Úpush_computed_sizesÚcomputed_sizesr	  s    rV   Ú__post_init__ÚEnterSubgraphLine.__post_init__Ê  s   € Ø�‰×(Ñ(¨¯©×)DÑ)DÕErX   c                ón   • U R                   R                  U R                  5        UR                  5         g rJ   )rŠ   Úpush_codegened_graphrK   Ú	do_indent©r
  Úcodes     rV   ÚcodegenÚEnterSubgraphLine.codegenÍ  s"   € Ø�‰×)Ñ)¨$¯*©*Ô5Ø�‰ÕrX   c                ó   • UR                   $ rJ   )Ú_generate_enter_subgraphr<  s     rV   r>  ÚEnterSubgraphLine.codegen_fxÑ  s   € Ø×1Ñ1Ð1rX   r�   N©r²   r5  ©rO  r-   r²   r5  r@  ©	rò   rà   r  r  rÙ   rI  rP  r>  r  r�   rX   rV   rD  rD  Å  s   ‡ à!Ó!ØÓôFô÷2rX   rD  c                  óH   • \ rS rSr% S\S'   S\S'   S
S jr\SS j5       rSrg	)ÚConditionalLineiÕ  rE  rŠ   zir.ConditionalrS   c                ó   • [        S5      e)NzOnly supports FX codegen)r;  rN  s     rV   rP  ÚConditionalLine.codegenÚ  s   € Ü!Ð"<Ó=Ð=rX   c                ó   • U R                   $ rJ   )Ú_generate_conditional©r=  s    rV   r>  ÚConditionalLine.codegen_fxÝ  s   € à×.Ñ.Ð.rX   r�   NrV  r@  ©	rò   rà   r  r  rÙ   rP  Ústaticmethodr>  r  r�   rX   rV   rY  rY  Õ  s'   ‡ à!Ó!Ø
Óô>ð ó/ó ó/rX   rY  c                  ó>   • \ rS rSr% S\S'   SS jr\S	S j5       rSrg)
ÚCommentLineiâ  r0   rš   c                ó:   • UR                  U R                  5        g rJ   )r—   rš   rN  s     rV   rP  ÚCommentLine.codegenæ  s   € Ø�‰�t—y‘yÕ!rX   c                ó   • U R                   $ rJ   )Ú_generate_commentr^  s    rV   r>  ÚCommentLine.codegen_fxé  s   € à×*Ñ*Ð*rX   r�   NrV  r@  r`  r�   rX   rV   rc  rc  â  s!   ‡ à
Óô"ð ó+ó ó+rX   rc  c                  óH   • \ rS rSr% S\S'   S\S'   S
S jr\SS j5       rSrg	)ÚDynamicScalarLineiî  rE  rŠ   zir.DynamicScalarrS   c                óN   • U R                   R                  U R                  5        g rJ   )rŠ   Ú_codegen_dynamic_scalarrS   rN  s     rV   rP  ÚDynamicScalarLine.codegenó  s   € Ø�‰×,Ñ,¨T¯Y©YÕ7rX   c                ó   • U R                   $ rJ   )Ú_generate_dynamic_scalarr^  s    rV   r>  ÚDynamicScalarLine.codegen_fxö  s   € à×1Ñ1Ð1rX   r�   NrV  r@  r`  r�   rX   rV   rj  rj  î  s'   ‡ à!Ó!Ø
Óô8ð ó2ó ó2rX   rj  c                  ó>   • \ rS rSr% S\S'   S	S jrS
S jrSS jrSrg)ÚExitSubgraphLineiû  rE  rŠ   c                óV   • U R                   R                  5       U R                   l        g rJ   )rŠ   Úpop_computed_sizesrH  r	  s    rV   rI  ÚExitSubgraphLine.__post_init__ÿ  s   € Ø&*§l¡l×&EÑ&EÓ&Gˆ�‰Õ#rX   c                óX   • U R                   R                  5         UR                  5         g rJ   )rŠ   Úpop_codegened_graphÚdo_unindentrN  s     rV   rP  ÚExitSubgraphLine.codegen  s   € Ø�‰×(Ñ(Ô*Ø×ÑÕrX   c                ó   • UR                   $ rJ   )Ú_generate_exit_subgraphr<  s     rV   r>  ÚExitSubgraphLine.codegen_fx  s   € Ø×0Ñ0Ð0rX   r�   NrU  rV  r@  rW  r�   rX   rV   rr  rr  û  s   ‡ à!Ó!ôHô÷1rX   rr  c                  ó>   • \ rS rSr% S\S'   S\S'   S
S jrSS jrSrg	)ÚEnterDeviceContextManagerLinei
  rï   Ú
device_idxúOptional[int]Úlast_seen_device_guard_indexc                ó’  • [         R                  R                  (       Ga  UR                  S5        [         R                  R                  (       aj  U R
                  c;  UR                  [         R                  R                  R                  5        S35        g U R
                  U R                  :X  d   S5       eg U R
                  cH  UR                  [         R                  R                  R                  5        SU R                   S35        g UR                  SU R                   S35        g UR                  S[         R                  R                  R                  U R                  5       S35        UR                  5         UR                  [         R                  R                  R                  U R                  5      5        g )	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 Ú:)r5   rK   Úcpp_wrapperr—   Úaot_moder�  Ú
device_opsÚcpp_aoti_stream_guardr  Úcpp_aoti_device_guardÚdevice_guardrM  Ú
set_devicerN  s     rV   rP  Ú%EnterDeviceContextManagerLine.codegen  sM  € Ü�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ð �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ÕJrX   c                ó   • UR                   $ rJ   )Ú&_generate_enter_device_context_managerr<  s     rV   r>  Ú(EnterDeviceContextManagerLine.codegen_fx,  s   € Ø×?Ñ?Ð?rX   r�   NrV  r@  ©rò   rà   r  r  rÙ   rP  r>  r  r�   rX   rV   r~  r~  
  s   ‡ àƒOØ"/Ó/ôK÷:@rX   r~  c                  ó(   • \ rS rSrSS jrSS jrSrg)ÚExitDeviceContextManagerLinei0  c                ód   • [         R                  R                  (       d  UR                  5         g g rJ   )r5   rK   r†  rx  rN  s     rV   rP  Ú$ExitDeviceContextManagerLine.codegen1  s    € Ü�w‰w×"×"Ø×ÑÕð #rX   c                ó   • UR                   $ rJ   )Ú%_generate_exit_device_context_managerr<  s     rV   r>  Ú'ExitDeviceContextManagerLine.codegen_fx5  s   € Ø×>Ñ>Ð>rX   r�   NrV  r@  )rò   rà   r  r  rP  r>  r  r�   rX   rV   r“  r“  0  s   † ô÷?rX   r“  c                  ó>   • \ rS rSr% S\S'   S\S'   S
S jrSS jrSrg	)ÚExternKernelAllocLinei9  rE  rŠ   úir.ExternKernelAllocrS   c                ó¬   • U R                   n/ UR                  5       QUR                  5       QnU R                  R	                  U R                   U5        g rJ   )rS   Úcodegen_argsÚcodegen_kwargsrŠ   Ú$_generate_extern_kernel_alloc_helper)r
  rO  rS   Úargss       rV   rP  ÚExternKernelAllocLine.codegen>  sD   € Ø�y‰yˆØ=�×"Ñ"Ó$Ð= t×':Ñ':Ó'<Ð=ˆØ�‰×9Ñ9¸$¿)¹)ÀTÕJrX   c                ó   • UR                   $ rJ   )Ú_generate_extern_kernel_allocr<  s     rV   r>  Ú ExternKernelAllocLine.codegen_fxC  s   € Ø×6Ñ6Ð6rX   r�   NrV  r@  r‘  r�   rX   rV   rš  rš  9  s   ‡ à!Ó!Ø
ÓôK÷
7rX   rš  c                  ó>   • \ rS rSr% S\S'   S\S'   S
S jrSS jrSrg	)ÚExternKernelOutLineiG  rE  rŠ   úir.ExternKernelOutrS   c           	     ó^  • U R                   n/ UR                  5       QUR                  SS9QnUR                  5       n[        R
                  R                  (       a  UR                  S:X  a  SnOUR                  5       nUR                  5       =n(       a  UR                  O[        R
                  R                  nU R                  R                  UUR                  5       UR                  (       a  UR                  R                  5       OS UUU R                   R                  5       5        g )NT)Úskip_outztorch::inductor::_mm_plus_mmÚaoti_torch__mm_plus_mm_out)rS   r�  rž  Úget_kernel_namer5   rK   r†  Úcpp_kernel_nameÚ
get_devicer‰   Údevice_typerŠ   Ú"_generate_extern_kernel_out_helperÚcodegen_referenceÚoutput_viewÚget_stack_traces)r
  rO  rS   r   Úkernel_nameÚdÚdevices          rV   rP  ÚExternKernelOutLine.codegenL  sé   € Ø�y‰yˆØJ�×"Ñ"Ó$ÐJ t×':Ñ':ÀDÐ':Ð'IÐJˆØ×*Ñ*Ó,ˆä�G‰G××Ø×$Ñ$Ð(FÓFð 7‰Kà×.Ñ.Ó0ˆKØ!%§¡Ó!2Ð2˜AÕ2�—’¼¿¹×9LÑ9LˆØ�‰×7Ñ7ØØ×"Ñ"Ó$Ø48×4D×4DˆD×Ñ×.Ñ.Ô0È$ØØØ�I‰I×&Ñ&Ó(õ	
rX   c                ó   • UR                   $ rJ   )Ú_generate_extern_kernel_outr<  s     rV   r>  ÚExternKernelOutLine.codegen_fxb  ó   € Ø×4Ñ4Ð4rX   r�   NrV  r@  r‘  r�   rX   rV   r¦  r¦  G  s   ‡ à!Ó!Ø
Óô
÷,5rX   r¦  c                  ó>   • \ rS rSr% S\S'   S\S'   S
S jrSS jrSrg	)ÚFreeLineif  rE  rŠ   ú%Union[BufferLike, ir.TorchBindObject]rS   c                óÜ   • U R                   R                  5       [        R                  R                  ;  d   eUR                  U R                  R                  U R                   5      5        g rJ   )rS   rM   r5   rK   Úremoved_buffersr—   rŠ   Úmake_buffer_freerN  s     rV   rP  ÚFreeLine.codegenk  sF   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÓBÐBÐBØ�‰�t—|‘|×4Ñ4°T·Y±YÓ?Õ@rX   c                ó   • UR                   $ rJ   )Ú_generate_freer<  s     rV   r>  ÚFreeLine.codegen_fxo  ó   € Ø×'Ñ'Ð'rX   r�   NrV  r@  r‘  r�   rX   rV   r¼  r¼  f  s   ‡ à!Ó!Ø
/Ó/ôA÷(rX   r¼  c                  ó¢   • \ rS rSr% S\S'   S\S'   S\S'   S\S'   S\S	'   S
\S'   S\S'   S\S'   S\S'   S\S'   S\S'   S\S'   SS jrSS jrSrg)ÚKernelCallLineis  rE  rŠ   rµ   r³  ztuple[Any, ...]Ú	call_argsÚraw_keysÚraw_argsú	list[str]Ú	arg_typesrð   r�   zdict[str, Any]Útriton_metazOptional[dict[str, Any]]Úinductor_metaztorch.devicerµ  Ú
graph_namerÂ   c                ó&  • U R                   R                  U R                  U R                  U R                  U R
                  U R                  U R                  U R                  U R                  U R                  U R                  U R                  S9  g )N)	r�   rÌ  rÉ  rÊ  rÍ  rÎ  rµ  rÏ  rÂ   )rŠ   Ú_generate_kernel_call_helperr³  rÈ  r�   rÌ  rÉ  rÊ  rÍ  rÎ  rµ  rÏ  rÂ   rN  s     rV   rP  ÚKernelCallLine.codegen‚  sn   € Ø�‰×1Ñ1Ø×ÑØ�N‰NØ—;‘;Ø—n‘nØ—]‘]Ø—]‘]Ø×(Ñ(Ø×,Ñ,Ø—;‘;Ø—‘Ø!%×!:Ñ!:ð 	2ò 	
rX   c                ó   • UR                   $ rJ   )Ú_generate_kernel_callr<  s     rV   r>  ÚKernelCallLine.codegen_fx‘  ó   € Ø×.Ñ.Ð.rX   r�   NrV  r@  r‘  r�   rX   rV   rÇ  rÇ  s  sR   ‡ à!Ó!ØÓØÓØÓØÓØÓØƒLØÓØ+Ó+ØÓØƒOØÓô
÷/rX   rÇ  c                  ór   • \ rS rSr% S\S'   S\S'   S\S'   SrS\S	'   S
rS\S'   SrS\S'   SS jrSS jr	Sr
g)ÚKernelDefinitionLinei•  rE  rŠ   rµ   r³  Úkernel_bodyNr¶   ÚmetadataTrð   ÚgpuÚcpp_definitionc                ó¢   • U R                   R                  U R                  U R                  U R                  U R
                  U R                  S9  g ©N)rÚ  rÛ  rÜ  )rŠ   Ú_define_kernel_helperr³  rÙ  rÚ  rÛ  rÜ  rN  s     rV   rP  ÚKernelDefinitionLine.codegenž  sB   € Ø�‰×*Ñ*Ø×ÑØ×ÑØ—]‘]Ø—‘Ø×.Ñ.ð 	+ò 	
rX   c                ó   • UR                   $ rJ   )Ú_generate_kernel_definitionr<  s     rV   r>  ÚKernelDefinitionLine.codegen_fx§  rº  rX   r�   rV  r@  )rò   rà   r  r  rÙ   rÚ  rÛ  rÜ  rP  r>  r  r�   rX   rV   rØ  rØ  •  s<   ‡ à!Ó!ØÓØÓØ"€HˆmÓ"Ø€CˆÓØ$(€N�MÓ(ô
÷5rX   rØ  c                  ó>   • \ rS rSr% S\S'   S	S jrS
S jrSS jrSrg)ÚMemoryPlanningLinei«  rE  rŠ   c                ó   • U $ )zFirst pass to find reuser�   ©r
  Ústates     rV   ÚplanÚMemoryPlanningLine.plan¯  s   € àˆrX   c                ó   • g)zSecond pass to output codeNr�   rN  s     rV   rP  ÚMemoryPlanningLine.codegen³  s   � rX   c                ó|  • / n[         R                  " U 5       Hw  nUR                  S:X  a  M  [        XR                  5      nUR	                  UR                   SUR
                  [        R                  L a  UR                  5       OU 35        My     [        U 5      R                   SSR                  U5       S3$ )z6
Emits a string representation that fits on one line.
rŠ   Ú=Ú(ú, rÛ   )ÚdataclassesÚfieldsr›   Úgetattrrq   r‰   r#   rj   rM   rò   r½   )r
  r   ÚfieldÚvals       rV   r  ÚMemoryPlanningLine.__str__¶  sš   € ð ˆÜ ×'Ò'¨Ö-ˆEØ�z‰z˜YÓ&ÙÜ˜$§
¡
Ó+ˆCØ�K‰KØ—:‘:�,˜a°%·*±*ÄÇ	Á	Ò2I §¡¤ÈsÐSÐTöñ	 .ô �t“*×%Ñ%Ð& a¨¯	©	°$«Ð'8¸Ð:Ð:rX   r�   N©rè  r  r²   rå  rV  ©r²   rµ   )	rò   rà   r  r  rÙ   ré  rP  r  r  r�   rX   rV   rå  rå  «  s   ‡ à!Ó!ôô)÷;rX   rå  c                  ó8   • \ rS rSrS rSS jrS	S jrS	S jrSrg)
ÚEfficientPeakEstimateiÅ  c                ó�  • SSK JnJn  [        R                  R
                  R                  n[        [        R                  R                  R                  5       5      n[        [        R                  R                  5       5      nU" X45      nU" UUU5      u  U l        nSSKJn  U" U[        R                  [         S5      U l        g )Nr    )Úestimate_peak_memoryÚget_freeable_input_bufr6   )ÚSegmentedTreer   )Úmemoryrü  rý  r5   rK   Ú	schedulerÚnodesr   Úgraph_inputsÚkeysÚget_output_namesÚoverall_peak_memoryÚsegmented_treerþ  Úoperatorr¾   Úmax)	r
  rü  rý  Úscheduler_nodesr  Úgraph_outputsÚnames_to_freeable_bufsÚpeak_by_scheduler_noderþ  s	            rV   r  ÚEfficientPeakEstimate.__init__Æ  s”   € ßIäŸ'™'×+Ñ+×1Ñ1ˆÜ!¤!§'¡'×"6Ñ"6×";Ñ";Ó"=Ó>ˆÜ"¤1§7¡7×#;Ñ#;Ó#=Ó>ˆÙ!7¸Ó!VÐÙ;OØØ"Øó<
Ñ8ˆÔ Ð"8õ 	2á+Ø"¤H§L¡L´#°qó
ˆÕrX   c                óÀ   • [         R                  R                  R                  [         R                  R	                  U5      SS9[        UR                  5       5      -  $ )Nr   ©Úfallback)r5   rK   rQ   Ú	size_hintrL   r,   rP   ©r
  rS   s     rV   Ú	_get_sizeÚEfficientPeakEstimate._get_sizeÙ  sL   € Ü�w‰w×Ñ×)Ñ)Ü�G‰G×/Ñ/°Ó5Àð *ð 
ä˜4Ÿ>™>Ó+Ó,ñ-ð 	-rX   c                ón   • U R                   R                  UR                  S-   UR                  S-
  5      $ r¡   )r  Úsummarize_rangeÚscheduler_node_index©r
  Úline_aÚline_bs      rV   Úpeak_betweenÚ"EfficientPeakEstimate.peak_betweenÞ  s6   € Ø×"Ñ"×2Ñ2Ø×'Ñ'¨!Ñ+¨V×-HÑ-HÈ1Ñ-Ló
ð 	
rX   c                óà   • UR                   S-   UR                   :X  a  g U R                  R                  UR                   S-   UR                   S-
  U R                  UR                  5      5        g r¡   )r  r  Úupdate_ranger  rS   r  s      rV   Úupdate_peak_betweenÚ)EfficientPeakEstimate.update_peak_betweenã  s^   € Ø×&Ñ&¨Ñ*¨f×.IÑ.IÓIØØ×Ñ×(Ñ(Ø×'Ñ'¨!Ñ+Ø×'Ñ'¨!Ñ+Ø�N‰N˜6Ÿ;™;Ó'õ	
rX   )r  r  N)rS   Ú
BufferLiker²   rï   )r  r4  r  ÚAllocateLine)	rò   rà   r  r  r  r  r  r  r  r�   rX   rV   rú  rú  Å  s   † ò
ô&-ô

÷

rX   rú  c                  ój   • \ rS rSr% SrS\S'   SrS\S'   S rSS	 jrSS
 jr	SS jr
SS jrSS jrSrg)r"  ií  z6Represents a buffer allocation during memory planning.r!  rS   Frð   Úcomm_bufferc                ó  • [         R                  R                  R                  c   e[         R                  R                  R                  R                  [         R                  R                  R                  5      U l        g rJ   ©r5   rK   r   Úcurrent_noder  Úindexr  r	  s    rV   rI  ÚAllocateLine.__post_init__ô  óT   € Ü�w‰w× Ñ ×-Ñ-Ñ9Ð9Ð9Ü$%§G¡G×$5Ñ$5×$;Ñ$;×$AÑ$AÜ�G‰G×Ñ×*Ñ*ó%
ˆÕ!rX   c                óü   • U R                   (       a  gUR                  S-   U R                  :X  a  gU R                  R                  R                  nU R                  R                  R                  X5      nX$-   nXS:*  $ )NTr6   )r$  r  rŠ   Úestimate_peakr  r  )r
  Ú	free_linerl   r  Úpeak_memory_in_rangeÚnew_peak_memorys         rV   Úshould_reuse_bufferÚ AllocateLine.should_reuse_bufferú  sl   € Ø××ØØ×)Ñ)¨AÑ-°×1JÑ1JÓJØØ"Ÿl™l×8Ñ8×LÑLÐØ#Ÿ|™|×9Ñ9×FÑFÀyÓWÐØÑ5ˆØÑ5Ð5rX   c           	     óT  • U R                   R                  5       [        R                  R                  ;   a  [        U R                  5      $ U R                  (       a„  [        U R                   5      n[        R                  (       aX  UR                  U5      (       aB  UR                  U5      nSUl        [        U R                  UR                   U R                   SS9$ U $ [        U R                   5      n[        R                  (       Ga
  X!;   Ga  UR!                  U5      n[        R                  R"                  R%                  [        R                  R'                  U R                   5      SS9[)        U R                   R+                  5       5      -  nU R-                  X45      (       aW  SUl        U R                  R.                  R1                  X05        [        U R                  UR                   U R                   5      $ UR3                  X#5        U $ U R                   R5                  5       R6                  S:X  aj  U R                  R9                  U R                   5      nUbB  U=R:                  [=        [>        R@                  " [B        RD                  US5      5      -  sl        U $ )NT©r$  r   r  Úcpur6   )#rS   rM   r5   rK   r¿  ÚNullLinerŠ   r$  r`   r"   Úallow_buffer_reuser+  r.  r$  Ú	ReuseLinerW   r#  rQ   r  rL   r,   rP   r0  r,  r  r(  rO   r‰   Ústatic_shape_for_buffer_or_noner  rï   Ú	functoolsÚreducer  Úmul)r
  rè  rª   r-  rl   Ústatic_shapes         rV   ré  ÚAllocateLine.plan  sì  € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ó:Ü˜DŸL™LÓ)Ð)à××ä'¨¯	©	Ó2ˆCÜ×(×(¨U×-GÑ-GÈ×-LÑ-LØ!×1Ñ1°#Ó6�	Ø&*�	Ô#Ü Ø—L‘L )§.¡.°$·)±)Èñð ð ˆKô ˜tŸy™yÓ)ˆÜ×$×$Ð$¨¬ØŸ	™	 #›ˆIÜ—7‘7×#Ñ#×-Ñ-Ü—‘×3Ñ3°D·I±IÓ>Èð .ð ä˜tŸy™y×2Ñ2Ó4Ó5ñ6ˆDð ×'Ñ'¨	×8Ñ8Ø&*�	Ô#Ø—‘×*Ñ*×>Ñ>¸yÔOÜ  §¡¨y¯~©~¸t¿y¹yÓIÐIà—
‘
˜3Ô*Ø�à�9‰9×(Ñ(Ó*×/Ñ/°5Ó8ØŸ<™<×GÑGÈÏ	É	ÓRˆLØÑ'Ø×1Ò1´SÜ×$Ò$¤X§\¡\°<ÀÓCó6ñ Õ1ð ˆrX   c                ó&  • U R                   R                  5       [        R                  R                  ;  d   eU R
                  (       a  U R                  U5        g U R                  R                  U R                   5      nUR                  U5        g rJ   )
rS   rM   r5   rK   r¿  r$  Ú_codegen_comm_bufferrŠ   Úmake_buffer_allocationr—   )r
  rO  rš   s      rV   rP  ÚAllocateLine.codegen+  sb   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÓBÐBÐBØ××Ø×%Ñ% dÕ+à—<‘<×6Ñ6°t·y±yÓAˆDØ�N‰N˜4Õ rX   c                óR  • U R                   R                  5       nU R                   R                  5       nUb  UR                  c
   SU 35       eU R                   R	                  5       n[        U R                   R                  5       5      n[        U R                   R                  5       5      nU R                   R                  5       n[        U[        R                  5      (       d   eUR                  nUR                  n	U[        R                  R                  :X  ai  U SU R                   R#                  U5       SU R                   R#                  U5       SU SUR                   SU	 S[$        R&                  " SS	5       S
3n
O[)        SU 35      eUR+                  U
5        g)z*Generate allocation code for comm buffers.Nz9Comm buffer requires a valid CUDA device with index, got z = empty_strided_p2p(rð  z, torch.device("cuda:z"), group_name="z", alloc_id=r   l   ÿÿÿÿ rÛ   zUnsupported comm buffer type: )rS   rM   r­  r(  rP   r�   Úget_sizeÚ
get_striderZ   r[   r#   r\   r]   r^   ÚCommBufferTypeÚSYMM_MEMrŠ   Úcodegen_shape_tupleÚrandomÚrandintr;  r—   )r
  rO  r›   rµ  r   Úshaperm   r_   r]   r^   rš   s              rV   r?  Ú!AllocateLine._codegen_comm_buffer3  s‚  € à�y‰y×!Ñ!Ó#ˆØ—‘×%Ñ%Ó'ˆØÑ! f§l¡lÑ&>ð 	
ØGÈÀxÐPó	
Ð>ð —	‘	×#Ñ#Ó%ˆÜ�d—i‘i×(Ñ(Ó*Ó+ˆÜ�t—y‘y×+Ñ+Ó-Ó.ˆØ—‘×*Ñ*Ó,ˆÜ˜&¤"×"5Ñ"5×6Ñ6Ð6Ð6Ø!×2Ñ2ÐØ×&Ñ&ˆ
àœr×0Ñ0×9Ñ9Ó9à�&Ð-Ø—<‘<×3Ñ3°EÓ:Ð;¸2Ø—<‘<×3Ñ3°FÓ;Ð<¸BØ�'ð &Ø&,§l¡l ^ð 4Ø)˜lð +Ü"ŸNšN¨1¨iÓ8Ð9¸ð<ñ ô &Ø0Ð1AÐ0BÐCóð ð 	�‰�tÕrX   c                óT   • U R                   (       a  UR                  $ UR                  $ rJ   )r$  Ú_generate_comm_buffer_allocateÚ_generate_allocater<  s     rV   r>  ÚAllocateLine.codegen_fxR  s#   € Ø××Ø×;Ñ;Ð;Ø×+Ñ+Ð+rX   ©r  N)r-  r4  rl   rï   r²   rð   r÷  rV  r@  )rò   rà   r  r  Ú__doc__rÙ   r$  rI  r0  ré  rP  r?  r>  r  r�   rX   rV   r"  r"  í  s4   ‡ á@à
ÓØ€K�Óò
ô6ô%ôN!ô÷>,rX   r"  c                  ó`   • \ rS rSr% S\S'   SrS\S'   SrS\S'   S rSS	 jrSS
 jr	SS jr
Srg)r4  iX  r!  rS   Frð   r$  r$  c                ó  • [         R                  R                  R                  c   e[         R                  R                  R                  R                  [         R                  R                  R                  5      U l        g rJ   r&  r	  s    rV   rI  Ú!FreeIfNotReusedLine.__post_init__^  r*  rX   c                óf  • [        U R                  R                  5       5      S:”  a  U $ [        U R                  R                  [
        R                  5      (       a  U $ U R                  (       a   eU R                  R                  5       [        R                  R                  ;   a  [        U R                  5      $ [        R                  (       a_  U R                   (       a(  [#        U R                  5      nUR%                  X 5        U $ ['        U R                  5      nUR)                  X 5        U $ r  )r£   rS   Úget_inputs_that_alias_outputr[   r_   r#   ÚMultiOutputLayoutr$  rM   r5   rK   r¿  r5  rŠ   r"   r6  r$  r`   r1  rW   r(  )r
  rè  rª   s      rV   ré  ÚFreeIfNotReusedLine.pland  sÔ   € Üˆt�y‰y×5Ñ5Ó7Ó8¸1Ó<ØˆKÜ�d—i‘i×&Ñ&¬×(<Ñ(<×=Ñ=ØˆKØ—>—>Ð!Ð!Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ó:Ü˜DŸL™LÓ)Ð)Ü×$×$Ø××ä+¨D¯I©IÓ6�Ø×&Ñ& sÔ1ð ˆô ' t§y¡yÓ1�Ø—
‘
˜3Ô%ØˆrX   c                óô  • U R                   R                  5       [        R                  R                  ;  d   eU R
                  (       d¯  U R                  R                  U R                   5      nU R                  (       ag  U R                   R                  5       n[        U[        R                  5      (       d   eUR                  U SUR                  R                   S35        g UR                  U5        g g )Nz # z buffer free)rS   rM   r5   rK   r¿  r$  rŠ   rÀ  r$  rZ   r[   r#   r\   r—   r]   rñ   )r
  rO  rš   r_   s       rV   rP  ÚFreeIfNotReusedLine.codegenv  s±   € Ø�y‰y×!Ñ!Ó#¬1¯7©7×+BÑ+BÓBÐBÐBØ�~�~Ø—<‘<×0Ñ0°·±Ó;ˆDØ××ØŸ™×2Ñ2Ó4�Ü! &¬"×*=Ñ*=×>Ñ>Ð>Ð>Ø—‘ $  s¨6×+BÑ+B×+HÑ+HÐ*IÈÐVÕWà—‘˜tÕ$ð rX   c                óT   • U R                   (       a  UR                  $ UR                  $ rJ   )r$  Ú_generate_comm_buffer_freeÚ_generate_free_if_not_reusedr<  s     rV   r>  ÚFreeIfNotReusedLine.codegen_fx�  s#   € Ø××Ø×7Ñ7Ð7Ø×5Ñ5Ð5rX   rP  Nr÷  rV  r@  )rò   rà   r  r  rÙ   r$  r$  rI  ré  rP  r>  r  r�   rX   rV   r4  r4  X  s0   ‡ à
ÓØ€IˆtÓØ€K�Óò
ôô$	%÷6rX   r4  c                  óR   • \ rS rSr% S\S'   S\S'   S\S'   SS jrSS jrSS	 jrS
rg)ÚReinterpretLinei‡  r!  rS   Ú	reused_asz	ir.Layoutr_   c                ó   • U $ rJ   r�   rç  s     rV   ré  ÚReinterpretLine.plan�  s   € ØˆrX   c                óT  • [        U R                  [        R                  5      (       d   e[        U R                  R                  [        R
                  5      (       d   eU R                  R                  U R                  R                  5       U R                  R                  5        g rJ   )
r[   r_   r#   ÚNonOwningLayoutÚviewr&   rŠ   Úcodegen_deferred_allocationra  rM   rN  s     rV   rP  ÚReinterpretLine.codegen�  sp   € Ü˜$Ÿ+™+¤r×'9Ñ'9×:Ñ:Ð:Ð:Ü˜$Ÿ+™+×*Ñ*¬B×,>Ñ,>×?Ñ?Ð?Ð?Ø�‰×0Ñ0Ø�N‰N×#Ñ#Ó% t§{¡{×'7Ñ'7õ	
rX   c                ó   • UR                   $ rJ   )Ú_generate_reinterpretr<  s     rV   r>  ÚReinterpretLine.codegen_fx—  rÖ  rX   r�   Nr÷  rV  r@  )	rò   rà   r  r  rÙ   ré  rP  r>  r  r�   rX   rV   r`  r`  ‡  s#   ‡ à
ÓØÓØÓôô
÷/rX   r`  c                  ód   • \ rS rSr% S\S'   S\S'   SrS\S'   SrS\S	'   SS
 jrSS jrSS jr	Sr
g)r7  i›  r!  rS   ra  Trð   Ú
delete_oldFr$  c                ó|  • U R                   R                  5       [        R                  R                  ;   aM  U R
                  R                  5       [        R                  R                  ;   d   e[        U R                  5      $ U R
                  R                  5       [        R                  R                  ;  d   eU $ rJ   )rS   rM   r5   rK   r¿  ra  r5  rŠ   rç  s     rV   ré  ÚReuseLine.plan¢  s   € Ø�9‰9×ÑÓ¤1§7¡7×#:Ñ#:Ó:Ø—>‘>×*Ñ*Ó,´·±×0GÑ0GÓGÐGÐGÜ˜DŸL™LÓ)Ð)Ø�~‰~×&Ñ&Ó(´·±×0GÑ0GÓGÐGÐGØˆrX   c                óx  • U R                   R                  5       [        R                  R                  ;  d   eU R
                  R                  5       [        R                  R                  ;  d   eUR                  U R                  R                  U R                   U R
                  U R                  5      5        g rJ   )
rS   rM   r5   rK   r¿  ra  r—   rŠ   Úmake_buffer_reuserm  rN  s     rV   rP  Ú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õ	
rX   c                ó   • UR                   $ rJ   )Ú_generate_reuser<  s     rV   r>  ÚReuseLine.codegen_fx°  s   € Ø×(Ñ(Ð(rX   r�   Nr÷  rV  r@  )rò   rà   r  r  rÙ   rm  r$  ré  rP  r>  r  r�   rX   rV   r7  r7  ›  s1   ‡ à
ÓØÓØ€J�ÓØ€K�Óôô
÷)rX   r7  c                  ó   • \ rS rSrSS jrSrg)r5  i´  c                ó   • UR                   $ rJ   )Ú_generate_nullr<  s     rV   r>  ÚNullLine.codegen_fxµ  rÅ  rX   r�   Nr@  rB  r�   rX   rV   r5  r5  ´  s   † ÷(rX   r5  c                  óV   • \ rS rSr% SrS\S'   S\S'   S\S'   S\S	'   SS
 jrSS jrSrg)ÚMultiOutputLinei¹  zM
Given a MultiOutputLayout buffer, indexes actual buffer(s) from the result.
rE  rŠ   rµ   Úresult_nameÚarg_nameúSequence[Any]Úindicesc                óè   ^ ^• UU 4S jmT" T R                   T R                  5      nUR                  T R                  R                   T R
                   SU T R                  R                   35        g )Nc                óˆ  >• [        U5      S:”  a±  US   u  p#[        U[        5      (       a  T" U  SU S3USS  5      $ [        U[        5      (       a;  TR                  R                  U TR                  [        U5      5      nT" XASS  5      $ [        U[        5      (       a  T" U  SU S3USS  5      $ [        SU5      eU $ )Nr   Ú[Ú]r6   z['z']znon supported index type: )
r£   Ú
issubclassr  r�   rŠ   Úcodegen_tuple_accessr|  rµ   ÚdictÚAssertionError)Úbasenamer  ÚitypeÚiÚtuple_accessÚcodegen_list_tuple_accessr
  s        €€rV   rŒ  Ú:MultiOutputLine.codegen.<locals>.codegen_list_tuple_accessÅ  sÐ   ø€ Ü�7‹|˜aÓØ" 1™:‘�Ü˜e¤T×*Ñ*Ù4¸°zÀÀ1À#ÀQÐ5GÈÐQRÐQSÈÓUÐUÜ ¤u×-Ñ-à#'§<¡<×#DÑ#DØ  $×"2Ñ"2´C¸³Fó$�Lñ 5°\È1È2À;ÓOÐOÜ ¤t×,Ñ,Ù4¸°zÀÀAÀ3ÀbÐ5IÈ7ÐSTÐSUÈ;ÓWÐWä(Ð)EÀuÓMÐMà�rX   rÞ   )r}  r  r—   rŠ   Údeclarer|  Úending)r
  rO  rñ   rŒ  s   `  @rV   rP  ÚMultiOutputLine.codegenÄ  s]   ù€ ö	 ñ$ *¨$¯-©-¸¿¹ÓFˆØ�‰Ø�|‰|×#Ñ#Ð$ T×%5Ñ%5Ð$6°c¸%¸ÀÇÁ×ATÑATÐ@UÐVõ	
rX   c                ó   • UR                   $ rJ   )Ú_generate_multi_outputr<  s     rV   r>  ÚMultiOutputLine.codegen_fxÜ  s   € Ø×/Ñ/Ð/rX   r�   NrV  r@  )	rò   rà   r  r  rQ  rÙ   rP  r>  r  r�   rX   rV   r{  r{  ¹  s*   ‡ ñð "Ó!ØÓØƒMØÓô
÷00rX   r{  c                  óH   • \ rS rSr% S\S'   S\S'   S\S'   SS jrSS	 jrS
rg)ÚIndexPutFallbackLineià  rE  rŠ   úir.IndexPutFallbackrS   zlist[Optional[ir.IRNode]]r  c                ó¨  • U R                   n[        R                  " UR                  5      (       d   eS UR                  S S  5       u  p4U R                   Vs/ s H0  nU(       a  UR                  5       OU R                  R                  PM2     nnU R                  R                  " UR                  5       X6U/UR                  5       Q76   g s  snf )Nc              3  ó@   #   • U  H  oR                  5       v •  M     g 7frJ   ©r°  ©r‚   Úts     rV   r„   Ú/IndexPutFallbackLine.codegen.<locals>.<genexpr>é  s   é € ÐF²o°×*Ñ*×,Ð,²oùó   ‚r    )rS   r#   Úis_node_sequenceÚinputsr  r°  rŠ   Únone_strÚ_generate_index_put_fallbackr«  Úcodegen_const_args)r
  rO  rS   r¦   ÚvaluesÚidxr  s          rV   rP  ÚIndexPutFallbackLine.codegenæ  s¶   € Ø�y‰yˆÜ×"Ò" 4§;¡;×/Ñ/Ð/Ð/ÙF°d·k±kÀ"À1±oÓF‰ˆð —|’|ó
â#�ö (+ˆC×!Ñ!Ô#°·±×0EÑ0EÒEÙ#ð 	ð 
ð
 	�‰×1Ò1Ø× Ñ Ó" A°ð	
Ø9=×9PÑ9PÓ9Rô	
ùò
s   Á7Cc                ó   • UR                   $ rJ   )r¡  r<  s     rV   r>  ÚIndexPutFallbackLine.codegen_fxó  s   € Ø×5Ñ5Ð5rX   r�   NrV  r@  r‘  r�   rX   rV   r•  r•  à  s   ‡ à!Ó!Ø
ÓØ&Ó&ô
÷6rX   r•  c                  ó>   • \ rS rSr% S\S'   S\S'   S
S jrSS jrSrg	)ÚScatterFallbackLinei÷  rE  rŠ   úir.ScatterFallbackrS   c           
     óP  • U R                   n[        R                  " UR                  5      (       d   eUR                  (       a  S UR                   5       u  p4nO$S UR                   5       u  p4UR
                  S   nUR                  5       =n(       a  UR                  O[        R                  R                  nU R                  R                  UX2R
                  S   XE/UR                  UR                  UR                  UR                  S   UR!                  5       U5        g )Nc              3  ó@   #   • U  H  oR                  5       v •  M     g 7frJ   r™  rš  s     rV   r„   Ú.ScatterFallbackLine.codegen.<locals>.<genexpr>   s   é € ÐJºk¸×2Ñ2×4Ð4ºkùr�  c              3  ó@   #   • U  H  oR                  5       v •  M     g 7frJ   r™  rš  s     rV   r„   r­    s   é € ÐEº°A×-Ñ-×/Ð/ºùr�  r6   r   r:  )rS   r#   rž  rŸ  Úsrc_is_tensorÚconstant_argsr­  r‰   r5   rK   r®  rŠ   Ú_generate_scatter_fallbackr¬  Úpython_kernel_namer¤   rž  )r
  rO  rS   r¦   r(  rê   r´  rµ  s           rV   rP  ÚScatterFallbackLine.codegenü  så   € Ø�y‰yˆÜ×"Ò" 4§;¡;×/Ñ/Ð/Ð/Ø××ÙJ¸d¿kºkÓJ‰OˆQ‘sáE¸¿ºÓE‰JˆQØ×$Ñ$ QÑ'ˆCØ!%§¡Ó!2Ð2˜AÕ2�—’¼¿¹×9LÑ9LˆØ�‰×/Ñ/ØØ×"Ñ" 1Ñ% uÐ2Ø× Ñ Ø×#Ñ#Ø×ÑØ�K‰K˜Ñ!Ø×ÑÓ!Øõ		
rX   c                ó   • UR                   $ rJ   )r±  r<  s     rV   r>  ÚScatterFallbackLine.codegen_fx  s   € Ø×3Ñ3Ð3rX   r�   NrV  r@  r‘  r�   rX   rV   r©  r©  ÷  s   ‡ à!Ó!Ø
Óô
÷(4rX   r©  c                  óH   • \ rS rSr% S\S'   S\S'   S\S'   SS jrSS	 jrS
rg)ÚSymbolicCallArgLinei  rE  rŠ   r  ÚargrD   rK   c                ód   • U R                   R                  U R                  U R                  5        g rJ   )rŠ   Ú"_generate_symbolic_call_arg_helperr¸  rK   rN  s     rV   rP  ÚSymbolicCallArgLine.codegen  s   € Ø�‰×7Ñ7¸¿¹À$Ç*Á*ÕMrX   c                ó   • UR                   $ rJ   )Ú_generate_symbolic_call_argr<  s     rV   r>  ÚSymbolicCallArgLine.codegen_fx  rº  rX   r�   NrV  r@  r‘  r�   rX   rV   r·  r·    s   ‡ à!Ó!Ø	ÓØÓôN÷5rX   r·  c                  óR   • \ rS rSr% S\S'   S\S'   S\S'   S\S	'   SS
 jrSS jrSrg)ÚUnbackedSymbolDefsLinei!  rE  rŠ   rµ   Úoutput_namer   Úoutputsú,Optional[dict[sympy.Symbol, pytree.KeyPath]]Úunbacked_bindingsc                óz   • U R                   R                  U R                  U R                  U R                  5        g rJ   )rŠ   Ú)_codegen_unbacked_symbol_defs_for_outputsrÁ  rÂ  rÄ  rN  s     rV   rP  ÚUnbackedSymbolDefsLine.codegen(  s+   € Ø�‰×>Ñ>Ø×Ñ˜dŸl™l¨D×,BÑ,Bõ	
rX   c                ó   • UR                   $ rJ   )Ú_generate_unbacked_symbol_defsr<  s     rV   r>  Ú!UnbackedSymbolDefsLine.codegen_fx-  s   € Ø×7Ñ7Ð7rX   r�   NrV  r@  r‘  r�   rX   rV   rÀ  rÀ  !  s#   ‡ à!Ó!ØÓØƒLØCÓCô
÷
8rX   rÀ  c            
      óØ  ^ • \ rS rSr% SrSrS\S'   U 4S jr\ Sœ       S�S jj5       r	SžS	 jr
SŸS
 jrSžS jrS S jrSžS jr\SžS j5       rSžS jr\SžS j5       rS¡S jr\S¢S j5       rSžS jr  S£S jrS¤S jrSžS jrSžS jrSžS jrS¥S jrS¦S jrS¢S jrSžS jrSžS jrS§S jr S r!S  r"S! r#S" r$S# r%S¨S$ jr&S©S% jr'SžS& jr(SªS' jr)S«S( jr*S«S) jr+S«S* jr,S¬S+ jr-S­S, jr.S- r/    S®S. jr0 Sœ             S¯S/ jjr1S°S0 jr2S°S1 jr3S°S2 jr4S3 r5S±S4 jr6S5 r7S²S6 jr8S7 r9              S³S8 jr:S9 r;S¦S: jr<\=R|                  S´S; j5       r?SžS< jr@S= rAS> rBS? rCS@ rDSµSA jrE      S¶SB jrFSC rGS·SD jrHSE rISSF.S¸SG jjrJSSF.S¸SH jjrKS¹SI jrLSºSJ jrMS»SK jrNS»SL jrO  S¼SM jrP Sœ   S½SN jjrQS¾SO jrRS¿SP jrSSQ rTSR rUSS rVST rWSU rXSV rY   SÀ         SÁSW jjrZ\ Sœ     SÂSX jj5       r[   SÀ         SÁSY jjr\SÃSZ jr]\^S[ 5       r_  SÄS\ jr`SœSÅS] jjra      SÆS^ jrbSÇS_ jrcSÇS` jrdSa reSb rfSc rgSd rhSe riSf rjSg rkSh rlSi rmSœSj jrnSk roSSSSSSSSSl. SÈSm jjrpSSSSSSSSnSSo.	 SÈSp jjrqSq rrSr rsSs rtSœSt jruSÉSu jrv\Sv 5       rw SÊSw jrxSx rySËSy jrzSÌSz jr{SÍS{ jr|SÎS| jr} Sœ SÏS} jjr~SÐS~ jrSÑS jr€SÒS€ jr�S� r‚SœS‚ jrƒSƒ r„SÓS„ jr…S… r†        SÔS† jr‡        SÔS‡ jrˆSˆ r‰    SÕS‰ jrŠSÖSŠ jr‹S‹ rŒSŒ r�S� rŽSŽ r�S� r�S� r‘SžS‘ jr’S’ r“\S“ 5       r”\S” 5       r•\S• 5       r–\S– 5       r—\S— 5       r˜    S×S˜ jr™S™ ršSš r›S›rœU =r�$ )ØrE  i5  z:
Generate outer wrapper in Python that calls the kernels.
Trð   Úsupports_cachingc                ó°  >^ • [         TT ]  5         [        5       T l        0 T l        [        5       T l        [        5       T l        [        5       T l        [        5       T l	        [        5       T l
        [        5       T l        [        5       T l        [        5       T l        [        5       T l        [        5       T l        0 T l        ST l        0 T l        [        5       T l        / T l        ST l        ST l        ST l        ST l        ST l        [6        R8                  R:                  (       a  SOST l        [6        R8                  R:                  (       a  SOST l        S T l         ST l!        0 T l"        [        5       T l#        [        5       T l$        S T l%        T RM                  5         / T l'        / T l(        T RS                  5         [U        T 5      (       d  T RW                  5         T RY                  5         [6        R8                  RZ                  (       dB  [6        R8                  R\                  R_                  5        H  u  pT Ra                  X5        M     [        [b           " 5       T l2        [        [b           " 5       T l3        0 T l4        [j        Rl                  " S 5      " T Rn                  5      T l7        [j        Rp                  S
U 4S jj5       nUT l9        0 T l:        [        5       T l;        [y        5       T l=        [        5       T l>        0 T l?        [�        [‚        R„                  R†                  [‚        R„                  Rˆ                  S	9T lE        / T lF        g )Nr   Ú Ú#r5  z
std::move(rÛ   Tc                ó²   >• TR                   R                  U 5        [        R                  R                  (       a  TR
                  R                  U 5        g g rJ   )Úimportsr—   r"   r�   r�   r™   )rš   r
  s    €rV   Úadd_import_onceÚ6PythonWrapperCodegen.__init__.<locals>.add_import_once‡  s;   ø€ à�L‰L×"Ñ" 4Ô(Ü�}‰}×5×5Ø×*Ñ*×4Ñ4°TÕ:ð 6rX   )Údebug_printer_levelÚuse_array_ref)rš   rµ   r²   r5  )Gr  r  r   Ú_names_iterÚargs_to_buffersr-   rÑ  ÚheaderÚprefixÚsuffixÚkernel_declarationsÚwrapper_callÚkernel_autotune_defsr™   Úsubgraph_definitionsr   r˜   Úkernel_autotune_example_argsÚkernel_autotune_tmp_arg_idxÚsrc_to_kernelÚkernel_numel_exprÚlinesrŽ  Údeclare_maybe_referencer�  Úcommentr   r5   rK   r†  Ú
move_beginÚmove_endr�  Úsupports_intermediate_hooksÚuser_defined_kernel_cacheÚunbacked_symbol_declsrH  Úlauncher_fn_nameÚset_launcher_fn_nameÚcodegened_graph_stackÚcomputed_sizes_stackÚwrite_headerr.   Úwrite_prefixÚ!write_kernel_autotune_defs_headerr‡  Úconstant_reprsÚitemsÚwrite_constantÚ
BufferNameÚ	allocatedÚfreedÚreusesr9  Ú	lru_cacheÚwrite_get_raw_streamÚcacherÒ  Ú_metasÚ
_meta_varsr   Úmulti_kernel_stateÚalready_codegened_subgraphsÚallocated_workspacesr   r"   Úaot_inductorÚ debug_intermediate_value_printerÚallow_stack_allocationÚdebug_printerÚadditional_files)r
  r›   ÚhashedrÒ  r  s   `   €rV   r  ÚPythonWrapperCodegen.__init__<  s¥  ù€ Ü‰ÑÔÜ*/«'ˆÔð ð 	Ôô &Ó'ˆŒÜ$Ó&ˆŒÜ$Ó&ˆŒÜ$Ó&ˆŒÜ#1Ó#3ˆÔ Ü*Ó,ˆÔÜ$2Ó$4ˆÔ!Ü%3Ó%5ˆÔ"Ü$2Ó$4ˆÔ!Ü6@³lˆÔ"ð IKˆÔ)Ø01ˆÔ(ð .0ˆÔÜHRËˆÔØ!#ˆŒ
ØˆŒØ')ˆÔ$ØˆŒØˆŒØˆŒÜ*+¯'©'×*=×*=™,À2ˆŒÜ Ÿw™w×2×2™¸ˆŒØ;?ˆÔ)Ø+/ˆÔ(ð ð 	Ô&ô ‹Lð 	Ô"ô 9C»ˆÔØ $ˆÔà×!Ñ!Ô#ð &(ˆÔ"Ø$&ˆÔ!à×ÑÔä2°4×8Ñ8à×ÑÔà×.Ñ.Ô0ä�w‰w××Ü !§¡× 6Ñ 6× <Ñ <Ö >‘�à×#Ñ# DÖ1ñ !?ô $¤JÒ/Ó1ˆŒÜ¤
Ò+Ó-ˆŒ
ð 57ˆŒä$-×$7Ò$7¸Ô$=Ø×%Ñ%ó%
ˆÔ!ô 
�‰ö	;ó 
ð	;ð
  /ˆÔØ&(ˆŒÜ+5«<ˆŒÜ"2Ó"4ˆÔÜ<F»LˆÔ(Ø46ˆÔ!ô 1Ü &× 3Ñ 3× TÑ TÜ ×-Ñ-×DÑDñ
ˆÔð !#ˆÕrX   Nc                óP   • U (       a  Uc   eUc   e[        XU5      $ [        5       $ rJ   )ÚSubgraphPythonWrapperCodegenrE  )Úis_subgraphÚsubgraph_nameÚparent_wrapperÚpartition_signaturess       rV   ÚcreateÚPythonWrapperCodegen.create�  s=   € ö Ø Ñ,Ð,Ð,Ø!Ñ-Ð-Ð-Ü/ØÐ/Cóð ô $Ó%Ð%rX   c                ó   • SU l         g )NÚcall)rë  r	  s    rV   rì  Ú)PythonWrapperCodegen.set_launcher_fn_name¬  s
   € à &ˆÕrX   c                óD   • U R                   R                  U SU 35        g )Nz = None  # )rØ  r—   )r
  r›   r  s      rV   rô  Ú#PythonWrapperCodegen.write_constant°  s   € Ø�‰×Ñ   k°&°Ð:Õ;rX   c                óú  • [         R                  R                  R                  5       nSnUb  UR                  b  SUR                   3nSn[        [        R                  R                  5      S:”  a  SnO5[         R                  R                  R                  R                  (       a  SnU R                  R                  SU S[        R                   S	U S
3SS9  U R                   R                  SSS9   SSKJn  U R                   R                  SSS9  [        R*                  (       a  U R                   R-                  S5        [        R.                  R0                  (       Gaý  U R                   R-                  S5        U R                   R-                  S5        U R                   R-                  S5        U R                   R-                  S5        U R                   R-                  S5        U R                   R-                  S5        U R                   R-                  S5        [        R.                  R2                  =(       d(    [4        R6                  R9                  [;        5       S5      nU R                   R-                  SU S35        SU S3nSU S3n[        R.                  R<                  n[        R.                  R>                  n	[        R.                  R@                  n
U R                   R-                  S5        U R                   R                  SU S U S!U	 S"U
 S#3	5        U R                   R-                  S$U S%35        U R                   R-                  S&5        gg! [&        [(        4 a     GNaf = f)'z>Write the header section of the generated Python wrapper code.rÎ  Nz
# AOT ID: r   zRfrom torch._inductor.codegen.debug_utils import _print_debugging_tensor_value_infozFfrom torch._inductor.runtime.debug_utils import tracked_empty_strided
z
                aH  
                from ctypes import c_void_p, c_long, c_int
                import torch
                import math
                import random
                import os
                import tempfile
                from math import inf, nan
                from cmath import nanj
                from torch._inductor.hooks import run_intermediate_hooks
                from torch._inductor.utils import maybe_profile
                from torch._inductor.codegen.memory_planning import _align as align
                from torch import device, empty_strided
                from zq import AsyncCompile
                from torch._inductor.select_algorithm import extern_kernels
                z
            TrÎ   a–  
                aten = torch.ops.aten
                inductor_ops = torch.ops.inductor
                _quantized = torch.ops._quantized
                assert_size_stride = torch._C._dynamo.guards.assert_size_stride
                assert_alignment = torch._C._dynamo.guards.assert_alignment
                empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
                empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
                empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
                reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
                alloc_from_pool = torch.ops.inductor._alloc_from_pool
                async_compile = AsyncCompile()
            )Ú_SymmetricMemoryzs
                empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
                zfrom torch.cuda import nvtxz import triton.profiler as protonz%import triton.profiler.language as plzCfrom triton.profiler.hooks import HookManager as _ProtonHookManagerzimport tritonzimport atexitz	import oszetriton.set_allocator(lambda size, align, stream: torch.empty(size, dtype=torch.uint8, device='cuda'))Úprotonzos.makedirs("z", exist_ok=True)zos.path.join("z", "inductor")z", "inductor.chrome_trace")z^from torch._inductor.runtime.proton_utils import process_proton_trace as _proton_process_tracez�
                def _proton_finalize_and_postprocess():
                    proton.finalize()
                    _trace_path = zµ
                    if os.path.exists(_trace_path):
                        _proton_process_trace(
                            _trace_path,
                            group_by_sm=z0,
                            split_invocations=z0,
                            per_cta_occupancy=z,,
                        )
                z5if not _ProtonHookManager.active_hooks: proton.start(z], backend="instrumentation", data="trace"); atexit.register(_proton_finalize_and_postprocess)úpl.enable_semantic("triton"))!ÚtorchÚ_guardsÚTracingContextÚtry_getÚaot_graph_namerï   r"   r  r  Ú	_inductorÚtest_configsÚtrack_memory_lifecyclerÑ  ré   r!   rò   rØ  Útorch._C._distributed_c10dr  ÚAttributeErrorÚImportErrorÚannotate_trainingr—   r�   Úproton_profilingÚproton_output_dirÚosÚpathr½   r   Úproton_group_by_smÚproton_split_invocationsÚproton_per_cta_occupancy)r
  ÚcontextÚaot_config_commentÚinductor_debug_utilsr  Ú
output_dirÚproton_nameÚ
trace_pathÚgroup_by_smÚsplit_invocationsÚper_cta_occupancys              rV   rï  Ú!PythonWrapperCodegen.write_header³  s7  € ä—-‘-×.Ñ.×6Ñ6Ó8ˆØÐØÑ 7×#9Ñ#9Ñ#EØ#-¨g×.DÑ.DÐ-EÐ!FÐØ!ÐÜŒv×"Ñ"×CÑCÓDÀqÓHØ#wÑ Ü�_‰_×#Ñ#×0Ñ0×G×GØ#lÐ à�‰×ÑðØ#Ð$ð %ô $×,Ñ,Ð-ð .à%Ð&ð 'ð!ð$ ð' 	ñ 	
ð* 	�‰×Ñðð ð! 	ñ 	
ð$	õ Dà�K‰K×Ñðð ð	 ñ ô ×#×#Ø�K‰K×!Ñ!Ð"?Ô@Ü�=‰=×)×)Ð)Ø�K‰K×!Ñ!Ð"DÔEØ�K‰K×!Ñ!Ð"IÔJØ�K‰K×!Ñ!ØUôð �K‰K×!Ñ! /Ô2Ø�K‰K×!Ñ! /Ô2Ø�K‰K×!Ñ! +Ô.Ø�K‰K×!Ñ!ðGôô  Ÿ™×8Ñ8÷ ¼B¿G¹G¿L¹LÜ“ ó=ˆJð �K‰K×!Ñ! M°*°Ð=NÐ"OÔPØ*¨:¨,°nÐEˆKØ)¨*¨Ð5PÐQˆJÜ Ÿ-™-×:Ñ:ˆKÜ &§¡× FÑ FÐÜ &§¡× FÑ FÐØ�K‰K×!Ñ!Øpôð �K‰K×Ñð#ð $. ,ð /)ð *5¨ð 6/Ø/@Ð.Að B/Ø/@Ð.Að Bðôð �K‰K×!Ñ!ð Ø +˜}ð -DðDôð
 �K‰K×!Ñ!Ð"@ÕAð[ *øô	 ¤Ð,ó 	Úð	ús   Ã8 M& Í&M:Í9M:c                ó   • g rJ   r�   )r
  rØ  s     rV   Úinclude_extra_headerÚ)PythonWrapperCodegen.include_extra_header$  ó   € ØrX   c                óÌ   • U R                   R                  S[        R                   S35         SSKJn  U R                   R                  SSS9  g ! [        [        4 a     g f = f)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   )Ú_cuda_getCurrentRawStreamzU
                get_raw_stream = torch._C._cuda_getCurrentRawStream
                TrÎ   )rÝ  ré   r!   rò   Útorch._Cr;  r#  r"  )r
  r;  s     rV   rñ  Ú6PythonWrapperCodegen.write_kernel_autotune_defs_header'  sw   € Ø×!Ñ!×(Ñ(ðô
 $×,Ñ,Ð-ð .ðô	
ð
	Ý:à×%Ñ%×,Ñ,ðð ð	 -ò øô œ^Ð,ó 	Ùð	ús   ¯ A ÁA#Á"A#c                ó  • S[         R                   S3n[        R                  R                  (       a]  U R
                  R                  U5        U R
                  R                  [        R                  R                  R                  S5      5        [        R                  R                  (       d]  U R                  R                  USS9  U R                  R                  [        R                  R                  R                  S5      5        g g )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—   r5   rK   rˆ  Úimport_get_raw_stream_asr†  rÑ  ©r
  Ú
import_strs     rV   Úwrite_triton_header_onceÚ-PythonWrapperCodegen.write_triton_header_onceC  sÇ   € ðô $×,Ñ,Ð-ð .ðˆ
ô
 �=‰=×1×1Ø×&Ñ&×-Ñ-¨jÔ9Ø×&Ñ&×0Ñ0Ü—‘×"Ñ"×;Ñ;Ð<LÓMôô �w‰w×"×"Ø�L‰L×Ñ 
°$ÐÑ7Ø�L‰L×"Ñ"Ü—‘×"Ñ"×;Ñ;Ð<LÓMõð #rX   c                óÂ  • [         R                  R                  R                  S5      n[        R
                  R                  (       a;  U R                  R                  U5      (       d  U R                  R                  U5        [         R                  R                  (       d=  U R                  R                  U5      (       d  U R                  R                  U5        g g g )Nr?  )r5   rK   rˆ  r@  r"   r�   r�   r™   Úcontainsr—   r†  rÑ  )r
  Úimport_get_raw_stream_strs     rV   Úwrite_get_raw_stream_headerÚ0PythonWrapperCodegen.write_get_raw_stream_headerU  s�   € Ü$%§G¡G×$6Ñ$6×$OÑ$OØó%
Ð!ô �=‰=×1×1Ø×-Ñ-×6Ñ6Ð7P×QÑQØ×*Ñ*×4Ñ4Ð5NÔOÜ�w‰w×"×"Ø—<‘<×(Ñ(Ð)B×CÑCØ—‘×&Ñ&Ð'@ÕAð Dð #rX   c                ó$   • U R                  5         g rJ   )rH  r	  s    rV   Ú write_get_raw_stream_header_onceÚ5PythonWrapperCodegen.write_get_raw_stream_header_once`  s   € à×(Ñ(Õ*rX   c                ó”  • [        U5      nXR                  ;  a   S[        U R                  5       3nX R                  U'   U R                  R	                  U SU 35        [
        R                  R                  (       a;  U R                  R	                  U SU 35        U R                  R                  U5        U R                  U   $ )NÚmetarÞ   )Úreprrü  r£   rØ  r—   r"   r�   r�   r™   rý  r¾   )r
  rN  Úvars      rV   Úadd_meta_onceÚ"PythonWrapperCodegen.add_meta_onced  s¡   € ä�D‹zˆØ—{‘{Ó"Øœ˜TŸ[™[Ó)Ð*Ð+ˆCà #�K‰K˜ÑØ�K‰K×!Ñ! S E¨¨T¨FÐ"3Ô4Ü�}‰}×5×5Ø×*Ñ*×4Ñ4¸°u¸CÀ¸vÐ5FÔGØ—‘×#Ñ# CÔ(à�{‰{˜4Ñ Ð rX   c                ó~   • U R                  5        Vs/ s H  oR                  U R                  5      PM     sn$ s  snf rJ   )Úget_graph_outputsr°  rÜ  ©r
  r¦   s     rV   Úget_output_refsÚ$PythonWrapperCodegen.get_output_refsr  s?   € ð =A×<RÑ<RÔ<Tó
Ú<T°q×Ñ × 1Ñ 1Ö2Ñ<Tñ
ð 	
ùò 
s   “$:c                ó   • g rJ   r�   r	  s    rV   Úmark_output_typeÚ%PythonWrapperCodegen.mark_output_typex  ó   € ØrX   c                ó6   • [         R                  R                  $ rJ   )r5   rK   r  r	  s    rV   Úget_graph_inputsÚ%PythonWrapperCodegen.get_graph_inputs{  s   € ô �w‰w×#Ñ#Ð#rX   c                ó6   • [         R                  R                  $ rJ   )r5   rK   r
  r	  s    rV   rT  Ú&PythonWrapperCodegen.get_graph_outputs€  s   € Ü�w‰w×$Ñ$Ð$rX   c           
     ó6  • U R                  5       R                  5        H÷  u  p[        U[        R                  [
        R                  45      (       a  M6  U[        R                  R                  ;  d  [        U[
        R                  5      (       a  Mu  [        UR                  5       5      S:X  a  M”  U R                  UR                  5       5      nU R                  UR                  5       5      nU R                  R!                  SU SU SU S35        Mù     g )Nr   zassert_size_stride(rð  rÛ   )r]  ró  r[   rz   r   r#   ÚTorchBindObjectr5   rK   Úgraph_input_namesÚGeneratorStater1   rC  rŽ   rD  rÙ  r—   )r
  r›   Úbufrl   rm   s        rV   Úcodegen_input_size_assertsÚ/PythonWrapperCodegen.codegen_input_size_assertsƒ  sÜ   € Ø×.Ñ.Ó0×6Ñ6Ö8‰IˆDÜ˜#¤§
¡
¬B×,>Ñ,>Ð?×@Ñ@Ùð œ1Ÿ7™7×4Ñ4Ó4¼
Ø”R×&Ñ&÷9ñ 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ò 9rX   c                ón  • U R                   R                  S5        U R                  5       R                  5        Hx  u  p[	        U[
        R                  [        R                  45      (       a  M6  SU S3nU R                   R                  U5        SU S3nU R                   R                  U5        Mz     g )Nz(# make sure graph inputs are not nan/infzassert not z.isnan().any().item()z.isinf().any().item())	rÙ  r—   r]  ró  r[   rz   r   r#   rb  )r
  r›   re  rš   s       rV   Úcodegen_input_nan_assertsÚ.PythonWrapperCodegen.codegen_input_nan_asserts•  s”   € Ø�‰×ÑÐHÔIØ×.Ñ.Ó0×6Ñ6Ö8‰IˆDÜ˜#¤§
¡
¬B×,>Ñ,>Ð?×@Ñ@Ùà   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Ô'Ø   Ð&;Ð<ˆDØ�K‰K×!Ñ! $Ö'ò 9rX   c                ó:   • U R                   R                  S5        g )NzV

            async_compile.wait(globals())
            del async_compile
            )rÙ  ré   r	  s    rV   Úwrite_async_compile_waitÚ-PythonWrapperCodegen.write_async_compile_wait   s   € Ø�‰×Ñðõ	
rX   c                óÀ   • SR                  U5      n[        U5      S:X  a  US-  nU R                  R                  U S35        U R                  R                  S5        g )Nrð  r6   Ú,z = argszargs.clear())r½   r£   rÙ  r—   )r
  Úinput_namesÚlhss      rV   Ú
write_argsÚPythonWrapperCodegen.write_args©  sP   € Ø�i‰i˜Ó$ˆÜˆ{Ó˜qÓ Ø�3‰JˆCØ�‰×Ñ   W˜oÔ.Ø�‰×Ñ˜nÕ-rX   c                óÄ   • [         R                  (       a  U R                  R                  S5        SnU$ U R                  R                  SU R                   S35        SnU$ )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):
                r6   )r"   Úgraph_partitionrÙ  ré   rë  ©r
  Úprefix_indents     rV   Ú!write_launcher_fn_call_get_indentÚ6PythonWrapperCodegen.write_launcher_fn_call_get_indent°  sm   € Ü×!×!Ø�K‰K×Ñðôð ˆMð Ðð �K‰K×ÑðØ×*Ñ*Ð+ð ,ðôð
 ˆMàÐrX   c                ó6   • [         R                  R                  $ rJ   )r5   rK   rc  r	  s    rV   Úget_graph_input_namesÚ*PythonWrapperCodegen.get_graph_input_namesÌ  s   € Ü�w‰w×(Ñ(Ð(rX   c                óþ  • U R                   c   eU R                  5         U R                  5       nU R                  R	                  U5         [
        R                  R                  (       aA  U R                  R                  [        R                  R                  R                  5       5        [        R                  R                  5       n[
        R                  (       a  U R                  R                  SU S35        U R                  5       =n(       a  U R!                  U5        U R#                  5         [%        5       (       a  ['        U 5      (       a  U R)                  5         S S S 5        g ! , (       d  f       g = f)Nz0training_annotation = nvtx._device_range_start('ú'))rë  rl  rx  rÙ  r·   r"   r�   Údebug_sync_graphr—   r5   rK   rˆ  ÚsynchronizeÚget_training_phaser$  r{  rr  Úcodegen_inputsr/   r.   Ú"codegen_input_size_and_nan_asserts)r
  rw  Úphaserc  s       rV   rð  Ú!PythonWrapperCodegen.write_prefixÏ  s	  € Ø×$Ñ$Ñ0Ð0Ð0Ø×%Ñ%Ô'Ø×>Ñ>Ó@ˆà�[‰[×Ñ Õ.Ü�}‰}×-×-Ø—‘×%Ñ%¤a§g¡g×&8Ñ&8×&DÑ&DÓ&FÔGÜ—G‘G×.Ñ.Ó0ˆEÜ×'×'Ø—‘×%Ñ%ØFÀuÀgÈRÐPôð %)×$>Ñ$>Ó$@Ð@Ð Õ@Ø—‘Ð 1Ô2à×ÑÔ!ô
 -×.Ñ.Ü<¸T×BÑBà×7Ñ7Ô9÷) /×.Ö.ús   ÁDE.Å.
E<c                óš   • [         R                  (       a  U R                  5         [         R                  (       a  U R	                  5         g g rJ   )r"   Úsize_assertsrf  Únan_assertsri  r	  s    rV   rƒ  Ú7PythonWrapperCodegen.codegen_input_size_and_nan_assertsê  s1   € Ü××Ø×+Ñ+Ô-Ü××Ø×*Ñ*Õ,ð rX   c                ó   • U R                  5         SU 3n[        R                  R                  (       aB  U R                  R                  U SU S35        [        R                  R                  (       a  U$ U R                  U SU S35        U$ )NÚstreamú = get_raw_stream(rÛ   )	rH  r"   r�   r�   r™   r—   r5   rK   r†  )r
  r  rÏ  r›   s       rV   rú  Ú)PythonWrapperCodegen.write_get_raw_streamó  s   € Ø×(Ñ(Ô*Ø˜
�|Ð$ˆÜ�=‰=×1×1Ø×&Ñ&×0Ñ0Ø�&Ð*¨:¨,°aÐ8ôô �w‰w×"×"à�Ø�‰˜$˜Ð1°*°¸QÐ?Ô@ØˆrX   c                ó    • U R                   S   $ )Néÿÿÿÿ)rí  r	  s    rV   Úget_codegened_graphÚ(PythonWrapperCodegen.get_codegened_graph   s   € Ø×)Ñ)¨"Ñ-Ð-rX   c                ó:   • U R                   R                  U5        g rJ   )rí  rq   )r
  rK   s     rV   rL  Ú)PythonWrapperCodegen.push_codegened_graph  s   € Ø×"Ñ"×)Ñ)¨%Õ0rX   c                ó6   • U R                   R                  5       $ rJ   )rí  r#  r	  s    rV   rw  Ú(PythonWrapperCodegen.pop_codegened_graph  s   € Ø×)Ñ)×-Ñ-Ó/Ð/rX   c                óP   • SSK Jn  U R                  R                  U" U5      5      $ )Nr   )Údeepcopy)Úcopyr—  rî  rq   )r
  rH  r—  s      rV   rG  Ú(PythonWrapperCodegen.push_computed_sizes	  s!   € Ý!à×(Ñ(×/Ñ/±¸Ó0HÓIÐIrX   c                ó6   • U R                   R                  5       $ rJ   )rî  r#  r	  s    rV   rt  Ú'PythonWrapperCodegen.pop_computed_sizes  s   € Ø×(Ñ(×,Ñ,Ó.Ð.rX   c                ó.   • [        U R                  5       $ rJ   )ÚnextrÖ  r	  s    rV   Únext_kernel_suffixÚ'PythonWrapperCodegen.next_kernel_suffix  s   € Ü�t×'Ñ'Ó(Ð)Ð*rX   c                óú  • U R                  [        XR                  5      5        [        R                  R
                  (       a²  U R                  5         U R                  R                  S[        R                  R                  R                  U5       S35        U R                  R                  5         [        U 5      (       a  U R                  5         U R                  R                  SU SU S35        Xl        g )Nr„  r…  r‹  rŒ  rÛ   )r—   r~  r�  r"   r�   r�   rC  r™   r5   rK   rˆ  r‹  rM  r.   rH  )r
  r  s     rV   Úcodegen_device_guard_enterÚ/PythonWrapperCodegen.codegen_device_guard_enter  sÇ   € Ø�‰Ü)¨*×6WÑ6WÓXô	
ô �=‰=×1×1à×)Ñ)Ô+Ø×&Ñ&×0Ñ0ØœŸ™×*Ñ*×7Ñ7¸
ÓCÐDÀAÐFôð ×&Ñ&×0Ñ0Ô2Ü2°4×8Ñ8à×0Ñ0Ô2Ø×&Ñ&×0Ñ0Ø˜˜Ð$6°z°lÀ!ÐDôð -7Õ)rX   c                óª   • U R                  [        5       5        [        R                  R                  (       a  U R
                  R                  5         g g rJ   )r—   r“  r"   r�   r�   r™   rx  r	  s    rV   Úcodegen_device_guard_exitÚ.PythonWrapperCodegen.codegen_device_guard_exit'  s6   € Ø�‰Ô3Ó5Ô6Ü�=‰=×1×1Ø×&Ñ&×2Ñ2Õ4ð 2rX   c                ó¬  • U(       Ga1  [         R                  (       aë  U R                  R                  SSR	                  U5      -   S-   5        U R                  R                  S5        U R                  R                  5         U R                  R                  S5        U R                  R                  5         U R                  R                  S5        U R                  R                  S5        U R                  R                  S5        U R                  R                  S	SR	                  U5      -   S-   5        g U R                  R                  S
5        g )Nzreturn_vars = (rð  ú, )zfor var in return_vars:z!if isinstance(var, torch.Tensor):z#assert not var.isnan().any().item()z#assert not var.isinf().any().item()r    zreturn (z	return ())r"   rˆ  rÜ  r—   r½   rM  rx  )r
  Úoutput_refss     rV   Úgenerate_returnÚ$PythonWrapperCodegen.generate_return,  s  € ßÜ×!×!Ø×!Ñ!×+Ñ+Ø%¨¯	©	°+Ó(>Ñ>ÀÑFôð ×!Ñ!×+Ñ+Ð,EÔFØ×!Ñ!×+Ñ+Ô-Ø×!Ñ!×+Ñ+Ð,OÔPØ×!Ñ!×+Ñ+Ô-Ø×!Ñ!×+Ñ+Ð,QÔRØ×!Ñ!×+Ñ+Ð,QÔRØ×!Ñ!×-Ñ-¨aÔ0à×Ñ×'Ñ'¨
°T·Y±Y¸{Ó5KÑ(KÈeÑ(SÕTà×Ñ×'Ñ'¨Õ4rX   c                ó   • g rJ   r�   ©r
  Úresults     rV   Úgenerate_before_suffixÚ+PythonWrapperCodegen.generate_before_suffix>  r[  rX   c                óÌ   • [         R                  (       aO  SR                  U R                  5      [	        U R                  5      S:X  a  SOS-   nUR                  SU S35        g g )Nrð  r6   ro  rÎ  z-
                runner = Runner(partitions=[z{])
                call = runner.call
                recursively_apply_fns = runner.recursively_apply_fns
                )r"   ru  r½   Úall_partition_namesr£   ré   )r
  r­  Úall_partition_name_lists      rV   Úgenerate_after_suffixÚ*PythonWrapperCodegen.generate_after_suffixA  se   € Ü×!×!Ø&*§i¡i°×0HÑ0HÓ&IÜ˜4×3Ñ3Ó4¸Ó9‘¸rñ'Ð#ð �M‰Mð-Ø-DÐ,Eð Fðõð "rX   c                ó   • g rJ   r�   r¬  s     rV   Úgenerate_endÚ!PythonWrapperCodegen.generate_endO  r[  rX   c                ó¼   • UR                   nUb3  U[        ;   a)  [        U   R                  nUb  U" XR                  5        g U R                  [	        X5      5        g rJ   )r²  r>   Úpythonr—   rš  )r
  rS   Úop_nameÚcustom_codegens       rV   Úgenerate_fallback_kernelÚ-PythonWrapperCodegen.generate_fallback_kernelR  sR   € à×)Ñ)ˆØÑ 7Ô.JÓ#JÜ9¸'ÑB×IÑIˆNØÑ)Ù˜t§^¡^Ô4ØØ�‰Ô,¨TÓ8Õ9rX   c                óÐ   • UR                  U 5        U R                  [        X5      5        [        UR                  [
        R                  5      (       a  UR                  U 5        g g rJ   )Úcodegen_commentr—   rš  r[   r_   r#   ÚLayoutÚcodegen_size_assertsr  s     rV   Úgenerate_extern_kernel_allocÚ1PythonWrapperCodegen.generate_extern_kernel_alloc\  sJ   € Ø×Ñ˜TÔ"Ø�‰Ô,¨TÓ8Ô9Ü�d—k‘k¤2§9¡9×-Ñ-Ø×%Ñ% dÕ+ð .rX   c           
     ó°  • [        UR                  [        R                  5      nUR	                  5       nUR                  5       nUR                  5       nU R                  n[        R                  (       a  SU;   a  SU 3nU(       a5  U R                  U R                   U SSR                  U5       SU 35        g U R                  U R                   U SU SSR                  U5       SU 35        U R                  (       aR  [        R                  (       a<  Ub8  [        S   S==   S	-  ss'   U R                  S
UR                   < SU S35        g g g g )NÚview_as_complexú.clone()rï  rð  rÛ   rÞ   ÚinductorÚintermediate_hooksr6   zrun_intermediate_hooks()r[   r_   r#   Ú
NoneLayoutrM   Úget_origin_noder«  r�  r"   Úmemory_planningr—   rŽ  r½   rè  Úgenerate_intermediate_hooksr   r›   )r
  Úextern_kernelr   Ú	no_returnrÁ  Úorigin_noder³  r�  s           rV   rŸ  Ú9PythonWrapperCodegen._generate_extern_kernel_alloc_helperb  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ð 1rX   c                óZ   • UR                  U 5        U R                  [        X5      5        g rJ   )r¿  r—   r¦  r  s     rV   Úgenerate_extern_kernel_outÚ/PythonWrapperCodegen.generate_extern_kernel_out  s$   € ð 	×Ñ˜TÔ"Ø�‰Ô*¨4Ó6Õ7rX   c                ó2  • [         R                  R                  R                  nUR	                  XAS S S5        UR                  SU(       a  UOU 35        U   U R                  U SSR                  U5       S35        S S S 5        g ! , (       d  f       g = f)NÚexternzout=rï  rð  rÛ   )r5   rK   Úwrapper_coder  Úset_printer_argsrq   r—   r½   )r
  r  ÚoutÚout_viewr   rµ  Ústack_tracesÚdebug_printer_managers           rV   r¯  Ú7PythonWrapperCodegen._generate_extern_kernel_out_helper†  sw   € ô !"§¡× 4Ñ 4× BÑ BÐØ×.Ñ.¨t¸TÀ4ÈÔRØ�‰�d¦x™8°SÐ9Ð:Ô;Ú"Ø�N‰N˜f˜X Q t§y¡y°£Ð&7°qÐ9Ô:÷ #×"Ö"ús   Á'BÂ
Bc                ó"  ^ • UR                   nUR                  nU(       aR  [        R                  R                  R                  U5      n[        R                  R                  R                  U5      nUR                  R                  5        S3nSR                  U 4S jU 5       5      nSR                  U 4S jU 5       5      n[        R                  T UR                  5      nSnU SUR                   S3nU SU SU SU 3n	U SU	 S	3n
U
$ )
Nz.data_ptr()rð  c              3  óP   >#   • U  H  n[         R                  TU5      v •  M     g 7frJ   ©rE  Úval_to_arg_str©r‚   Údimr
  s     €rV   r„   ÚRPythonWrapperCodegen._generate_tma_descriptor_call_experimental.<locals>.<genexpr>Ÿ  s$   øé € ÐXÒSWÈCÔ-×<Ñ<¸TÀ3×GÐGÒSWùó   ƒ#&c              3  óP   >#   • U  H  n[         R                  TU5      v •  M     g 7frJ   rß  rá  s     €rV   r„   rã     s%   øé € ð 
ÚFP¸sÔ ×/Ñ/°°c×:Ð:Âjùrä  z$triton.tools.experimental_descriptorz.create_Úd_tma_descriptorrï  rÛ   )ÚdimsÚ
block_dimsr5   rK   rQ   Úoptimization_hintÚoptimization_hintsÚtensorr°  r½   rE  rà  Úelement_sizeÚrank)r
  ÚdescÚapply_size_hintsrç  rè  Úptrrì  rÙ  rã   r   r  s   `          rV   Ú*_generate_tma_descriptor_call_experimentalÚ?PythonWrapperCodegen._generate_tma_descriptor_call_experimental–  sþ   ø€ Ø�y‰yˆØ—_‘_ˆ
ÞÜ—7‘7×#Ñ#×5Ñ5°dÓ;ˆDÜŸ™×)Ñ)×<Ñ<¸ZÓHˆJà—‘×.Ñ.Ó0Ð1°Ð=ˆà�y‰yÔXÑSWÓXÓXˆØ—Y‘Yô 
ÙFPó
ó 
ˆ
ô ,×:Ñ:¸4À×ARÑARÓSˆØ7ˆØˆx�x §	¡	˜{Ð*:Ð;ˆØ��b˜˜˜b  ¨B¨|¨nÐ=ˆØ��Q�t�f˜AˆˆØˆrX   c                óÚ   • UR                   nU(       a)  [        R                  R                  R	                  U5      nSnU S3nUR
                  R                  5        SU 3nU SU S3nU$ )Nz/triton.tools.tensor_descriptor.TensorDescriptorz.from_tensorrð  rï  rÛ   )Úblock_shaper5   rK   rQ   rê  rë  r°  )r
  rî  rï  rô  rÙ  rã   r   r  s           rV   Ú$_generate_tma_descriptor_call_stableÚ9PythonWrapperCodegen._generate_tma_descriptor_call_stableª  so   € Ø×&Ñ&ˆÞÜŸ'™'×*Ñ*×=Ñ=¸kÓJˆKàBˆØˆx�|Ð$ˆØ—+‘+×/Ñ/Ó1Ð2°"°[°MÐBˆØ��Q�t�f˜AˆˆØˆrX   c                óÆ   • [        U[        R                  5      (       a  U R                  X5      $ [        U[        R                  5      (       d   eU R                  X5      $ rJ   )r[   r#   ÚTMADescriptorExperimentalrñ  ÚTMADescriptorStablerõ  )r
  rî  rï  s      rV   Ú_generate_tma_descriptor_callÚ2PythonWrapperCodegen._generate_tma_descriptor_callµ  sW   € Ü�dœB×8Ñ8×9Ñ9Ø×BÑBØóð ô ˜d¤B×$:Ñ$:×;Ñ;Ð;Ð;Ø×<Ñ<¸TÓTÐTrX   c                ó‚   • U R                  U5      nUR                   SU U R                   3nU R                  U5        g ©NrÞ   )rú  r›   r�  r—   )r
  rî  r  rš   s       rV   Úgenerate_tma_descriptorÚ,PythonWrapperCodegen.generate_tma_descriptor¾  s:   € Ø×1Ñ1°$Ó7ˆØ—)‘)�˜C ˜v d§k¡k ]Ð3ˆØ�‰�tÕrX   c                ó8   • U R                  [        X5      5        g rJ   )r—   r©  r  s     rV   Úgenerate_scatter_fallbackÚ.PythonWrapperCodegen.generate_scatter_fallbackÃ  s   € Ø�‰Ô*¨4Ó6Õ7rX   c	                ó  • U SSR                  [        [        U5      5       3n	UR                  S5      (       a  U	SR                  S/U-   5      -  n	OU(       a  U	S[	        U5       3-  n	U	S-  n	U R                  U	5        g )Nrï  ro  zaten.scatter_reducerð  rÎ  z	, reduce=rÛ   )r½   Úmaprµ   ró   rO  r—   )
r
  rœ   rŸ  r¬  r²  r¯  r:  r¤   rµ  rš   s
             rV   r±  Ú/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ÕrX   c                ón  • / nUR                   SS  n[        U5      n[        UR                  5       He  u  pVUR                  U   b?  [	        U5      n[        U[        R                  5      (       d   eUR                  U5        MT  UR                  S 5        Mg     U R                  [        XU5      5        g )Nr    )rŸ  ÚiterÚ	enumerater  r�  r[   r#   r%   rq   r—   r•  )r
  rS   r  Úvalid_indicesÚiter_valid_indicesrŠ  Ú_r(  s           rV   Úgenerate_index_put_fallbackÚ0PythonWrapperCodegen.generate_index_put_fallbackÚ  s”   € à-/ˆØŸ™ A B˜ˆÜ! -Ó0ÐÜ˜dŸl™lÖ+‰DˆAØ�|‰|˜A‰Ñ*ÜÐ/Ó0�Ü! %¬¯©×3Ñ3Ð3Ð3Ø—‘˜uÖ%à—‘˜tÖ$ñ ,ð 	�‰Ô+¨D¸Ó@ÕArX   c                óv   • SSR                  U5       S3nX&XE/nU R                  U R                  X5      5        g )Nr‚  rð  rƒ  )r½   r—   Úwrap_kernel_call)r
  r  r¦   r  r£  Ú
accumulateÚindices_strr   s           rV   r¡  Ú1PythonWrapperCodegen._generate_index_put_fallbacké  s;   € Ø˜$Ÿ)™) GÓ,Ð-¨QÐ/ˆØ Ð3ˆØ�‰�t×,Ñ,¨VÓ:Õ;rX   c           
     ó`   • U R                  U SU SSR                  U" 5       5       S35        g )NrÞ   rï  rð  rÛ   )r—   r½   )r
  Úbuf_namer²  Úget_argsÚop_overloadrÊ  rÂ  s          rV   Ú,generate_fallback_kernel_with_runtime_lookupÚAPythonWrapperCodegen.generate_fallback_kernel_with_runtime_lookupî  s2   € ð 	�‰˜(˜ 3Ð'9Ð&:¸!¸D¿I¹IÁhÃjÓ<QÐ;RÐRSÐTÕUrX   c                óp   • [        S5         U R                  U5      sS S S 5        $ ! , (       d  f       g = f)NzPythonWrapperCodegen.generate)r   Ú	_generate©r
  Úis_inferences     rV   ÚgenerateÚPythonWrapperCodegen.generateù  s#   € ÜÐ9Õ:Ø—>‘> ,Ó/÷ ;×:×:ús   Œ'§
5c                ó0   • [         R                  (       a  gg)Nr    r6   )r"   ru  r	  s    rV   Úget_wrapper_call_indentÚ,PythonWrapperCodegen.get_wrapper_call_indentý  s   € Ü×!×!ØàrX   c              #  ó\   #   • U R                   n Xl         Uv •  X l         g ! X l         f = f7frJ   ©r—   )r
  ÚnewÚolds      rV   Úset_writelineÚ"PythonWrapperCodegen.set_writeline  s'   é € à�n‰nˆð	!Ø ŒNØŠIà �Nø˜S�Nüs   ‚,�
! š,¡)©,c                óÜ   • U R                   R                  n[        R                  R                  (       a  U R
                  R                  U5        g U R                  R                  U5        g rJ   )rþ  Úkernel_defsr"   r�   r�   rÝ  ré   rØ  )r
  r)  s     rV   Ú_write_multi_kernel_defsÚ-PythonWrapperCodegen._write_multi_kernel_defs  sF   € Ø×-Ñ-×9Ñ9ˆÜ�=‰=×1×1Ø×%Ñ%×,Ñ,¨[Õ9à�K‰K×Ñ˜{Õ+rX   c                óº
  • [         R                  (       a  U R                  5         [        R                  " 5        nUR                  U R                  R                  5       5        [         R                  (       a  U R                  U5        [         R                  (       a  U R                  5         U R                  U5        [         R                  R                  (       a/  [         R                  R                  (       d  U R                  5         U R!                  U R                  R"                  5         U R$                   HP  n['        U[(        5      (       a  UR+                  U R                  5        M5  U R                  R#                  U5        MR     S S S 5        U R-                  5         U R/                  5       nU R1                  5         [         R                  R2                  (       aA  U R                  R#                  [4        R6                  R8                  R;                  5       5        [         R                  (       a  U R=                  5         [         R                  R>                  (       a  U RA                  5         [         R                  R                  (       a/  [         R                  R                  (       d  U RC                  5         [         R                  R                  (       a  U RE                  5         [         RF                  (       a0  [         RH                  (       d  U R                  R#                  S5        U RK                  U5        S S S 5        [M        5       nURO                  U RP                  5        UR#                  S5        URO                  U RR                  5        [4        R6                  RT                  (       aH  [4        R6                  RH                  (       a)  [4        R6                  RV                  (       a
  [M        5       nURO                  U RX                  5        U R[                  5         URO                  U R\                  5        U R_                  5       nUR                  U5         URO                  U R                  5        S S S 5        U Ra                  U5        URO                  U Rb                  5        U Re                  U5        U Rg                  U5        U Ri                  U5        URk                  5       U Rl                  Rk                  5       4$ ! , (       d  f       GNš= f! , (       d  f       GNí= f! , (       d  f       N»= f)Nz+nvtx._device_range_end(training_annotation)rÎ  )7r"   Úprofile_bandwidthrC  r¸   Ú	ExitStackÚenter_contextrÜ  r·   Úprofiler_mark_wrapper_callÚ#generate_profiler_mark_wrapper_callÚgenerate_start_graphÚrun_wrapper_ir_passesr�   Ústore_cubinr�   Ú!generate_reset_kernel_saved_flagsr&  r—   rã  r[   rH   rP  r*  rV  rY  r  r5   rK   rˆ  r€  Úgenerate_end_graphr%  Úgenerate_proton_finalizeÚ generate_save_uncompiled_kernelsÚgenerate_and_run_autotune_blockr$  r†  r©  r-   ré   rÑ  rØ  r‡  Úis_const_graphrÞ  Úfinalize_prefixrÙ  r   r®  rÚ  r³  r¶  Úadd_benchmark_harnessÚgetvaluewithlinemaprÛ  )r
  r  Ústackrš   r¨  r­  Úwrapper_call_indents          rV   r  ÚPythonWrapperCodegen._generate  si  € Ü×#×#Ø×)Ñ)Ô+ä×!Ò!Ô# uØ×Ñ × 1Ñ 1× 8Ñ 8Ó :Ô;Ü×0×0Ø×8Ñ8¸Ô?Ü×'×'Ø×)Ñ)Ô+à×&Ñ& |Ô4ä�}‰}×(×(´·±×1W×1WØ×6Ñ6Ô8ð ×#Ñ# D×$5Ñ$5×$?Ñ$?Õ@Ø ŸJœJ�DÜ! $¬×4Ñ4àŸ™ T×%6Ñ%6Ö7à×)Ñ)×3Ñ3°DÖ9ñ '÷ Að ×)Ñ)Ô+à×.Ñ.Ó0ˆKØ×!Ñ!Ô#Ü�}‰}×-×-Ø×!Ñ!×+Ñ+¬A¯G©G×,>Ñ,>×,JÑ,JÓ,LÔMä×'×'Ø×'Ñ'Ô)ä�}‰}×-×-Ø×-Ñ-Ô/ä�}‰}×(×(´·±×1W×1WØ×5Ñ5Ô7ä�}‰}×5×5Ø×4Ñ4Ô6ô ×'×'´×0B×0BØ×!Ñ!×+Ñ+ØAôð × Ñ  Ô-÷] $ôb  Ó!ˆØ�‰�d—l‘lÔ#Ø×Ñ˜ÔØ�‰�d—k‘kÔ"ô �7‰7××¤§¡× 3× 3¼¿¹×8N×8NÜ#Ó%ˆFð 	�‰�d×/Ñ/Ô0Ø×ÑÔØ�‰�d—k‘kÔ"à"×:Ñ:Ó<Ðà�]‰]Ð.Õ/Ø�M‰M˜$×+Ñ+Ô,÷ 0ð 	×#Ñ# FÔ+Ø�‰�d—k‘kÔ"Ø×"Ñ" 6Ô*à×Ñ˜&Ô!à×"Ñ" 6Ô*ð ×&Ñ&Ó(Ø×$Ñ$×8Ñ8Ó:ð
ð 	
÷} AÖ@ú÷ $Ö#ú÷B 0Õ/ús2   »C9T:Ä4A!T(ÆF?T:Ñ;UÔ(
T7	Ô2T:Ô:
U	Õ
Uc                ó|  ^• U R                   R                  S5        0 n[        R                  R                  (       af  [
        R                  R                  (       aG  [        [
        R                  R                  5       VVs0 s H  u  p#U R                  U5      U_M     nnnU R                   R                  5       S-   U R                  R                  5       -   m[        R                  [        R                  :X  aj  [         R"                  " [%        5       SSS9 nUR'                  TR)                  S5      5        UR*                  nSSS5        [        R,                  " SW5        [/        S	S
 U4S jS9   [1        TU5        gs  snnf ! , (       d  f       NI= f! [2         a  n[5        SU 35      UeSnAff = f)z¢
Compose self.kernel_autotune_defs and self.kernel_autotune_calls into a single block of
code and execute it to trigger Triton kernel compilation and auto-tuning
zQ
            async_compile.wait(globals())
            del async_compile
        rƒ  z.pyF)ÚdirrÚ  Údeletezutf-8NzAuto-tuning code written to %sÚartifactc                 ó   • SSS.$ )NÚ&inductor_autotune_at_compile_time_codeÚstring)r›   Úencodingr�   r�   rX   rV   r§   ÚFPythonWrapperCodegen.generate_and_run_autotune_block.<locals>.<lambda>Œ  s   € Ø@Ø$ò!rX   c                 ó   >• T $ rJ   r�   )Útuning_codes   €rV   r§   rI  �  s   ø€ ™{rX   )Úmetadata_fnÚ
payload_fnz%Failed to run autotuning code block: )rÝ  ré   r"   r�   r�   r5   rK   Úautotuning_inputsr  Úget_autotuning_input_namer¿   r™   r$   ÚlevelÚloggingÚDEBUGÚtempfileÚNamedTemporaryFiler   ÚwriteÚencoder›   Údebugr   ÚexecÚ	ExceptionÚRuntimeError)r
  Úscoper¤  ÚvÚfÚ	file_pathÚerK  s          @rV   r9  Ú4PythonWrapperCodegen.generate_and_run_autotune_blockh  s„  ø€ ð
 	×!Ñ!×(Ñ(ðô	
ð ˆÜ�=‰=×1×1´a·g±g×6O×6Oô (¬¯©×(AÑ(AÔBôâB‘F�Cð ×.Ñ.¨sÓ3°QÒ6ÙBð ñ ð
 ×%Ñ%×.Ñ.Ó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*   ÂFÄ-FÅ9F Æ
FÆ
F;Æ'F6Æ6F;c                ó\   • SSK Jn  U" U 5      R                  U R                  5      U l        g )Nr6   )ÚMemoryPlanner)rË  rb  ré  rã  )r
  rb  s     rV   Úmemory_planÚ PythonWrapperCodegen.memory_plan˜  s    € Ý2á" 4Ó(×-Ñ-¨d¯j©jÓ9ˆ�
rX   c                óD  • U R                  5       n[        R                  R                  U5      nU R                  (       a¿  [        U R                  S   [        5      (       a�  U R                  S   R                  R                  U;  av  U R                  R                  5         U R                  (       aK  [        U R                  S   [        5      (       a)  U R                  S   R                  R                  U;  a  Mv  [        5       /n/ n[        [        U R                  5      5       H°  nU R                  U   n[        U[        5      (       a#  UR                  US   5      U R                  U'   MJ  [        U[        5      (       a  UR                  [        5       5        Mz  [        U[         5      (       d  M‘  UR                  UR                  5       5        M²     UR                  UR                  5       5        [        U5      S:X  d   e[#        S U 5       5      ng )Nr�  r   c              3  ó8   #   • U  H  oR                   v •  M     g 7frJ   )r  )r‚   Úss     rV   r„   Ú9PythonWrapperCodegen.memory_plan_reuse.<locals>.<genexpr>»  s   é € ð +
Ú3G¨a×)Ö)Ò3Gùs   ‚)rT  r5   rK   Ú_get_output_namesrã  r[   rå  rS   r›   r#  r  Úranger£   ré  rD  rq   rr  Úsum)r
  rÂ  Ú	out_namesÚplanning_statesÚpast_planning_statesrŠ  rš   Ú_total_allocated_buffer_sizes           rV   Úmemory_plan_reuseÚ&PythonWrapperCodegen.memory_plan_reuse�  s•  € Ø×(Ñ(Ó*ˆÜ—G‘G×-Ñ-¨gÓ6ˆ	ð �J�JÜ˜4Ÿ:™: b™>Ô+=×>Ñ>à—
‘
˜2‘×#Ñ#×(Ñ(°	Ó9ð �J‰J�N‰NÔð �J�JÜ˜4Ÿ:™: b™>Ô+=×>Ñ>à—
‘
˜2‘×#Ñ#×(Ñ(°	Õ9ô /Ó0Ð1ˆØ!ÐÜ”s˜4Ÿ:™:“Ö'ˆAØ—:‘:˜a‘=ˆDÜ˜$Ô 2×3Ñ3Ø $§	¡	¨/¸"Ñ*=Ó >�—
‘
˜1“Ü˜DÔ"3×4Ñ4Ø×&Ñ&Ô':Ó'<Ö=Ü˜DÔ"2×3Ó3Ø$×+Ñ+¨O×,?Ñ,?Ó,AÖBñ (ð 	×#Ñ# O×$7Ñ$7Ó$9Ô:Ü�?Ó# qÓ(Ð(Ð(ô
 (+ñ +
Ù3Gó+
ó (
Ñ$rX   c                óÆ   • U(       a&  [         R                  (       a  U R                  5         g [         R                  (       a  [	        5       U l        U R                  5         g rJ   )r"   rË  rc  r6  rú  r,  rp  r  s     rV   r3  Ú*PythonWrapperCodegen.run_wrapper_ir_passes¿  s9   € æœF×2×2Ø×ÑÕä×(×(Ü%:Ó%<�Ô"Ø×"Ñ"Õ$rX   c           	     ó†  ^	• U R                   m	[        R                  U	4S j5       n[        R                  U	4S j5       n[        U[        R
                  5      (       aM  [        U[        R                  5      (       a  X#;   a  g T	R                  U SU 35        UR                  U5        g [        U[        R                  5      (       a÷  [        UR                  5       5       H^  u  pg[        U[        R                  5      (       d  M&  Xs;  d  M-  T	R                  U SU" U5       SU S35        UR                  U5        M`     [        UR                  5       5       H^  u  ph[        U[        R                  5      (       d  M&  Xƒ;  d  M-  T	R                  U SU" U5       SU S35        UR                  U5        M`     g [        U[        R                  5      (       a  g [        U[        R                  5      (       a  g [         R"                  R$                  R&                  (       a  g [)        S[+        U5       35      e)Nc                ó<   >• TR                  U  SU  S35        U  S3$ )Nz_size = z.size()Ú_sizer#  ©r›   rO  s    €rV   ÚsizeofÚDPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.sizeofÐ  s(   ø€ à�N‰N˜d˜V 8¨D¨6°Ð9Ô:Ø�V˜5�>Ð!rX   c                ó<   >• TR                  U  SU  S35        U  S3$ )Nz
_stride = z	.stride()Ú_strider#  rw  s    €rV   ÚstrideofÚFPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.strideofÕ  s)   ø€ à�N‰N˜d˜V :¨d¨V°9Ð=Ô>Ø�V˜7Ð#Ð#rX   rÞ   r‚  rƒ  zUnknown value type: )rÙ  r9  rû  r[   rz   r   ÚSymbolr—   r¾   r#   ro   r  rC  rD  rb  rd  r  r  r"   ru  r‡  r‰   )
r
  r›   rñ   Ú
bound_varsrx  r|  râ  rl   rm   rO  s
            @rV   Úcodegen_input_symbol_assignmentÚ4PythonWrapperCodegen.codegen_input_symbol_assignmentÈ  s¶  ø€ ð �{‰{ˆä	�‰ô	"ó 
ð	"ô 
�‰ô	$ó 
ð	$ô �eœUŸZ™Z×(Ñ(Ü˜e¤U§\¡\×2Ñ2°eÓ6IØØ�N‰N˜e˜W C¨ vÐ.Ô/Ø�N‰N˜5Õ!Ü˜œrŸ|™|×,Ñ,Ü& u§~¡~Ó'7Ö8‘	�Ü˜d¤E§L¡L×1Ó1°dÕ6LØ—N‘N d V¨3©v°d«|¨n¸A¸c¸UÀ!Ð#DÔEØ—N‘N 4Ö(ñ 9ô  )¨×)9Ñ)9Ó);Ö<‘�Ü˜f¤e§l¡l×3Ó3¸Õ8PØ—N‘N f X¨S±¸$³Ð0@ÀÀ#ÀÀaÐ#HÔIØ—N‘N 6Ö*ò  =ô ˜œr×1Ñ1×2Ñ2ØÜ˜œr×0Ñ0×1Ñ1Øä�‰×%Ñ%×5×5Øä$Ð';¼DÀ»K¸=Ð%IÓJÐJrX   c           	     óF  • [         [        R                     " 5       nU R                  5       nUR	                  5        VVs/ s H)  u  p4[        U[        R                  5      (       d  M&  X44PM+     snnUR	                  5        VVs/ s H)  u  p4[        U[        R                  5      (       a  M&  X44PM+     snn-   nU H  u  pgU R                  XgU5        M         SS jnU H.  u  p—[        U[        R                  5      (       d  M&  U" Xq5        M0     gs  snnf s  snnf )z$Assign all symbolic shapes to localsc                óx  • [         R                  " U R                  5       U R                  5       /5       H}  n[	        U[
        5      (       a  [	        U[        R                  5      (       a  M9  UR                   Vs/ s H  o3U;  d  M
  UPM     nn[        U5      S:”  d  Mm  [        SU SU S35      e   g s  snf )Nr   zFor z, expected z to have been codegen-ed.)r   Úfrom_iterablerC  rD  r[   r   rz   r~  Úfree_symbolsr£   r‡  )rñ   r  ÚexprÚsymÚundefined_symbolss        rV   Ú_verify_input_symbol_assignmentÚLPythonWrapperCodegen.codegen_inputs.<locals>._verify_input_symbol_assignment  s°   € ô ×+Ò+¨U¯^©^Ó-=¸u×?OÑ?OÓ?QÐ,RÖS�Ü! $¬×-Ñ-´¸DÄ%Ç,Á,×1OÑ1OÙð $(×#4Ò#4ó%Ú#4˜CÀ:Ñ8M—CÑ#4ð "ð %ô Ð(Ó)¨AÕ-Ü(Ø˜t˜f KÐ0AÐ/BÐB[Ð\óð ò Tùò%s   Á<	B7Â	B7N)rñ   úir.TensorBoxr  úOrderedSet[sympy.Symbol])	r   rz   r~  r]  ró  r[   r€  r#   ro   )
r
  r  r  Úkr\  rŸ  r›   rñ   r‰  r  s
             rV   r‚  Ú#PythonWrapperCodegen.codegen_inputsò  sô   € ä¤§¡Ò-Ó/ˆ
ð ×,Ñ,Ó.ˆà+×1Ñ1Ô3ô
Ú3‘t�q´zÀ!ÄUÇ\Á\×7R‹FˆQ‹FÑ3ò
à ,× 2Ñ 2Ô 4ÔXÒ 4™˜¼JÀqÌ%Ï,É,×<W‹Vˆa‹VÑ 4ÒXñYˆó "‰KˆDØ×0Ñ0°¸jÖIñ "ð	Øð	à0ô	ó& ‰HˆAÜ˜e¤R§\¡\×2Ñ2ÙÙ+¨EÖ>ò ùó3
ùãXs   Á %DÁ)DÂ%DÂ/Dc                ó�  • [        U[        R                  5      (       a§  [        U[        R
                  5      (       a‡  XR                  ;   a  g U R                  R                  U5        [        R                  R                  R                  U   n[        X5      nU R                  [        X[        R                  5      5        g g g rJ   )r[   rz   r~  r   r   ÚPRECOMPUTED_SIZErH  r¾   r5   rK   rQ   Úinv_precomputed_replacementsr  r—   r·  )r
  r‡  r†  r¸  s       rV   Úensure_size_computedÚ)PythonWrapperCodegen.ensure_size_computed  s�   € Ü�cœ5Ÿ<™<×(Ñ(¬^¸CÄ×AVÑAV×-WÑ-WØ×)Ñ)Ó)ØØ×Ñ×#Ñ# CÔ(Ü—7‘7×#Ñ#×@Ñ@ÀÑEˆDÜ! #Ó,ˆCØ�N‰NÔ.¨t¼!¿'¹'ÓBÕCð .XÐ(rX   c                ó   • g rJ   r�   r	  s    rV   r;  Ú$PythonWrapperCodegen.finalize_prefix$  r9  rX   ©rR   c               ó   • [        S5      e)Nz8codegen_cpp_sizevar is only implemented for cpp_wrapper!)rZ  ©r
  r¦   rR   s      rV   Úcodegen_cpp_sizevarÚ(PythonWrapperCodegen.codegen_cpp_sizevar'  s   € ÜÐUÓVÐVrX   c               ó   • [        XS9$ )Nr–  )Úpexprr˜  s      rV   Úcodegen_python_sizevarÚ+PythonWrapperCodegen.codegen_python_sizevar*  s   € Ü�QÑ*Ð*rX   c                ó$   • U R                  U5      $ rJ   )r�  rU  s     rV   Úcodegen_sizevarÚ$PythonWrapperCodegen.codegen_sizevar-  s   € Ø×*Ñ*¨1Ó-Ð-rX   c                ó   • U SU S3$ )Nr‚  rƒ  r�   )r
  rˆ  r›   r(  s       rV   r…  Ú)PythonWrapperCodegen.codegen_tuple_access0  s   € Ø�˜1˜U˜G 1Ð%Ð%rX   c                ó¬   • / [        U R                  U5      Qn[        U5      S:X  a  g[        U5      S:X  a	  SUS    S3$ SSR                  U5       S3$ )Nr   ú()r6   rï  r§  rð  rÛ   )r  r�  r£   r½   )r
  rJ  Úpartss      rV   rŽ   Ú/PythonWrapperCodegen.codegen_python_shape_tuple3  s^   € Ø:”#�d×1Ñ1°5Ó9Ð:ˆÜˆu‹:˜‹?ØÜˆu‹:˜‹?Ø�u˜Q‘x�j Ð$Ð$Ø�4—9‘9˜UÓ#Ð$ AÐ&Ð&rX   c                ó$   • U R                  U5      $ rJ   )rŽ   )r
  rJ  s     rV   rG  Ú(PythonWrapperCodegen.codegen_shape_tuple;  s   € Ø×.Ñ.¨uÓ5Ð5rX   c                ó°   • SR                  SR                  U[        U5      [        U5      U R	                  U5      U R	                  U5      /5      5      / 4$ )Nzalloc_from_pool({})rð  )Úformatr½   rœ  rµ   rŽ   )r
  r›   rn   r   rJ  rm   s         rV   Úcodegen_alloc_from_poolÚ,PythonWrapperCodegen.codegen_alloc_from_pool>  s_   € ð %×+Ñ+Ø�I‰IàÜ˜&“MÜ˜“JØ×3Ñ3°EÓ:Ø×3Ñ3°FÓ;ðóó

ð ð
ð 
	rX   c                ó¬  ^ • [        U5      u  pxpšnU 4S jnUR                  5       nU=(       a    XI:H  nU(       a  X':H  =(       a    X8:H  nU
nObX!R                  R                  :H  =(       a7    X1R                  R                  :H  =(       a    XAR                  R
                  :H  nUR                  nU(       a  Ub  UU:w  a	  SU SU S3$ U $ U" XÒX4UU5      $ )Nc           	     ó®   >• T
R                  U5      nT
R                  U5      nT
R                  U5      nSU  SU SU SU S3	n	Ub  XE:w  a	  SU	 SU S3$ U	$ )Nzreinterpret_tensor(rð  rÛ   úaten.view.dtype()rŽ   r   )r›   Útgt_sizeÚ
tgt_strideÚ
tgt_offsetÚ
cast_dtypeÚ
base_dtyperg  ÚstÚoffr†  r
  s             €rV   Úapply_reinterpretÚHPythonWrapperCodegen.codegen_reinterpret_view.<locals>.apply_reinterpretl  sz   ø€ ð ×/Ñ/°Ó9ˆAØ×0Ñ0°Ó<ˆBØ×&Ñ& zÓ2ˆCØ(¨¨¨b°°°2°b°T¸¸C¸5ÀÐBˆDØÑ%¨*Ó*BØ)¨$¨¨r°*°¸QÐ?Ð?ØˆKrX   r°  rð  rÛ   )rw   rM   r_   rl   rm   rn   r   )r
  rr   rl   rm   rn   r—   r   Úd_sizeÚd_strideÚd_offsetÚd_dtypeÚcollapsibler¸  r›   Ú	collapsedÚsame_layoutrµ  s   `                rV   Úcodegen_reinterpret_viewÚ-PythonWrapperCodegen.codegen_reinterpret_viewM  sØ   ø€ ô8 ,¨DÓ1ñ 	9ˆ˜(¨[õ		ð �}‰}‹ˆØ×6 FÑ$6ˆ	ÞØ™.×?¨VÑ-?ˆKØ ‰Jð Ÿ™×(Ñ(Ñ(÷ 1ØŸk™k×0Ñ0Ñ0÷1àŸk™k×0Ñ0Ñ0ð ð
 Ÿ™ˆJæØÑ  U¨jÓ%8Ø)¨$¨¨r°%°¸Ð:Ð:Ø�VÐá  ¨V¸UÀJÓOÐOrX   c                ó8   • U R                  U SU SU S35        g )Nz.copy_(rð  rÛ   r#  )r
  rê   ÚdstÚnon_blockings       rV   Úcodegen_device_copyÚ(PythonWrapperCodegen.codegen_device_copy‹  s!   € Ø�‰˜#˜˜g c U¨"¨\¨N¸!Ð<Õ=rX   c                ó�   • UR                  5       nUR                  S5      nU R                  [        XX1R                  5      5        g r  )rM   Ú
input_namer—   r{  r  )r
  rS   r|  r}  s       rV   Úcodegen_multi_outputÚ)PythonWrapperCodegen.codegen_multi_outputŽ  s2   € Ø—m‘m“oˆØ—?‘? 1Ó%ˆØ�‰” t¸(ÇLÁLÓQÕRrX   c           
     ót  • UR                    SUR                   SUR                    SUR                    3nU(       a  SUR                   SU S3nU R                  UR                   SUR                   SUR
                   SU S	35        U R                  R                  [        UR                  5      5        g )
Nú + ú if z
 < 0 else zmax(0, min(rð  z))rÞ   z * (rÛ   )	r(  rl   r—   Úunbacked_offset_symbolÚbase_offsetÚbase_dim_striderê  r¾   rµ   )r
  rS   ÚclampÚ	index_strs       rV   Úcodegen_dynamic_select_indexÚ1PythonWrapperCodegen.codegen_dynamic_select_index“  s«   € Ø—z‘z�l # d§i¡i [°°T·Z±Z°LÀ
È4Ï:É:È,ÐWˆ	ÞØ% d§i¡i [°°9°+¸RÐ@ˆIØ�‰Ø×*Ñ*Ð+¨3¨t×/?Ñ/?Ð.@ÀÀD×DXÑDXÐCYÐY]Ð^gÐ]hÐhiÐjô	
ð 	×"Ñ"×&Ñ&¤s¨4×+FÑ+FÓ'GÕHrX   c                ó˜  ^ ^• UU 4S jnU 4S jnTR                   nU" TR                  5      nU" TR                  5      nT R                  U SU 35        T R                  U SU 35        U" U S3U S3TR                  5      nT R                  U SU S35        T R
                  R                  [        TR                   5      5        g )	Nc           	     óh  >• TR                  [        R                  " S[        R                  " U TR                  5      5      5      nTR                  [        R                  " S[        R                  " U TR                  -   TR                  5      5      5      nTR                  U 5      nU SU SU 3$ )Nr   rÎ  z >= 0 else )r   rz   ÚMaxÚMinrl   )r¦   ÚposÚnegÚx_condrS   r
  s       €€rV   Úclamp_indexÚDPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.clamp_indexž  s‹   ø€ Ø×&Ñ&¤u§y¢y°´E·I²I¸aÀÇÁÓ4KÓ'LÓMˆCØ×&Ñ&Ü—	’	˜!œUŸYšY q¨4¯9©9¡}°d·i±iÓ@ÓAóˆCð ×)Ñ)¨!Ó,ˆFØ�U˜$˜v˜h k°#°Ð7Ð7rX   c                ó\   >• US:X  a  U SU  3$ TR                  U5      nSU SU  SU SU 3$ )Nr6   z - rï  rÍ  z	 - 1) // )r   )Ú	start_varÚend_varÚstepÚstep_r
  s       €rV   Úcodegen_with_stepÚJPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.codegen_with_step¦  sJ   ø€ Ø�q‹yØ!˜ # i [Ð1Ð1Ø×(Ñ(¨Ó.ˆEØ�w�i˜s 9 +¨S°°°yÀÀÐHÐHrX   z	_start = z_end = Ú_startÚ_endz
 = max(0, rÛ   )Úunbacked_size_symbolÚstartÚendr—   râ  rê  r¾   rµ   )r
  rS   rÝ  rä  r‡  ré  rê  Ú	with_steps   ``      rV   Úcodegen_dynamic_slice_sizeÚ/PythonWrapperCodegen.codegen_dynamic_slice_size�  s»   ù€ ö	8õ	Ið ×'Ñ'ˆÙ˜DŸJ™JÓ'ˆÙ˜$Ÿ(™(Ó#ˆØ�‰˜#˜˜i¨ wÐ/Ô0Ø�‰˜#˜˜g c UÐ+Ô,Ù%¨¨¨V n¸¸¸T°lÀDÇIÁIÓNˆ	Ø�‰˜#˜˜j¨¨°1Ð5Ô6Ø×"Ñ"×&Ñ&¤s¨4×+DÑ+DÓ'EÕFrX   c                ó8   • U R                  [        X5      5        g rJ   )r—   rj  r  s     rV   Úcodegen_dynamic_scalarÚ+PythonWrapperCodegen.codegen_dynamic_scalar¶  s   € Ø�‰Ô(¨Ó4Õ5rX   c                óØ  • S UR                    5       u  n[        UR                  5      S:X  a#  U R                  UR                   SU S35        GOw[        UR                  5      S:X  aE  [        UR                  S   [        5      (       a#  U R                  UR                   SU S35        GO[        UR                  5      S:X  aè  [        UR                  S   [        5      (       aÆ  U R                  UR                   SU S35        U R                  S	UR                   S
UR                  S   R                   SUR                   SUR                  S   R                   S3	5        U R                  UR                   SUR                   SUR                  S   R                   35        O[        SUR                   35      eU R                  UR                  5        S35        g )Nc              3  ó@   #   • U  H  oR                  5       v •  M     g 7frJ   r™  rš  s     rV   r„   Ú?PythonWrapperCodegen._codegen_dynamic_scalar.<locals>.<genexpr>º  s   é € Ð>²+¨Q×&Ñ&×(Ð(²+ùr�  r   rÞ   ú.item()r6   z = 1 if z.item() else 0z_undivided = zassert z_undivided % z
 == 0, f'{z_undivided} not divisible by Ú'z_undivided // úunrecognized keypath ú = None)rŸ  r£   Úkeypathr—   r‡  r[   r   r   Údivisorr‡  rM   )r
  rS   rr   s      rV   rl  Ú,PythonWrapperCodegen._codegen_dynamic_scalar¹  sŒ  € Ù>°$·+²+Ó>‰ˆÜˆt�|‰|Ó Ó!Ø�N‰N˜dŸh™h˜Z s¨4¨&°Ð8Ö9Ü�—‘Ó !Ó#¬
°4·<±<À±?ÄM×(RÑ(RØ�N‰N˜dŸh™h˜Z x°¨v°^ÐDÖEÜ�—‘Ó !Ó#¬
°4·<±<À±?ÄK×(PÑ(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Õ3rX   c           
     óà  ^ ^• UU 4S jnU4S jnU4S jnTR                  / SQ5        TR                  5          TR                  SSS9  [        R                  R
                  R                  5        HT  u  pVTR                  SU 35        U" XVR                  5       UR                  5       UR                  UR                  5        MV     [        [        R                  R                  5      S	:”  a^  TR                  S
5        [        R                  R                  R                  5        H!  u  pWTR                  SU 35        U" XW5        M#     [        R                  R                  R                  5        GH9  u  pV[        U[         R"                  5      (       aI  [        [        R                  R$                  R&                  R)                  US5      [*        5      (       a  Mn  [        U[,        R.                  5      (       ad  [        [        R                  R                  5      S	:X  a  TR                  S
5        TR                  SU 35        U" XVR1                  5       5        Mñ  [        U[         R2                  5      (       a2  U" U[        R                  R$                  R5                  USS95        GMB  [        U[,        R6                  5      (       a$  U" USUR                  R8                   S35        GM…  UR;                  5        Vs/ s H+  n[        R                  R$                  R5                  USS9PM-     n	nUR=                  5        Vs/ s H+  n[        R                  R$                  R?                  USS9PM-     n
nU" UU	U
URA                  5       URC                  5       5        GM<     TR                  SSRE                  [        R                  R                  RG                  5       5       S35        SSS5        TR                  / SQ5        TR                  5          TR                  SSS9  SSS5        gs  snf s  snf ! , (       d  f       NU= f! , (       d  f       g= f)z2Write out codegen for benchmarking the output codec                ó‚   >• TR                  U  STR                  U5       STR                  U5       SU SU S3
5        g )Nz = rand_strided(rð  ú
, device='ú	', dtype=rÛ   )r—   rŽ   )r›   rJ  rm   rµ  r   rœ   r
  s        €€rV   Úadd_fake_inputÚFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_fake_inputÑ  sT   ø€ Ø×ÑØ�&Ð(Ø×2Ñ2°5Ó9Ð:¸"Ø×2Ñ2°6Ó:Ð;ð <Ø!˜( )¨E¨7°!ð5õrX   c                ó2   >• TR                  U  SU 35        g rý  r#  )r›   rõ  rœ   s     €rV   Úadd_expr_inputÚFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_expr_inputÙ  s   ø€ Ø×Ñ ˜v S¨¨Ð.Õ/rX   c                ó8  >• Uc  TR                  U  S35        g SS Kn TR                  U  SUR                  U5      < S35        g ! [        [        UR
                  4 a7  nTR                  S[        U5       SU  S[        U5       S35         S nAg S nAff = f)	Nr÷  r   z = pickle.loads(rÛ   z.raise TypeError("Failed to pickle opaque type z for variable rÜ   z"))r—   ÚpickleÚdumpsÚ	TypeErrorr"  ÚPicklingErrorr‰   rµ   )r›   rñ   r  r_  rœ   s       €rV   Úadd_torchbind_inputÚKPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_torchbind_inputÜ  s¨   ø€ Ø‰}Ø× Ñ  D 6¨Ð!1Ô2ØãðØ× Ñ  D 6Ð)9¸&¿,¹,ÀuÓ:MÑ9PÐPQÐ!RÕSøÜœ~¨v×/CÑ/CÐDó Ø× Ñ ØDÄTÈ%Ã[ÀMÐQ_Ð`dÐ_eÐegÔhkÐlmÓhnÐgoÐoqÐr÷ñ ûðús   Ÿ'A ÁBÁ"-BÂB)rÎ  rÎ  zdef get_args():zP
                from torch._dynamo.testing import rand_strided
                TrÎ   zglobal r   zimport pickleNé*   r  ztorch.cuda.default_generators[z].graphsafe_get_state()zreturn [rð  rƒ  )rÎ  rÎ  z9def benchmark_compiled_module(args, times=10, repeat=10):zÌ
                from torch._inductor.utils import print_performance
                fn = lambda: call(list(args))
                return print_performance(fn, times=times, repeat=repeat)
                )$Ú
writelinesr·   ré   r5   rK   Ú	constantsró  r—   rl   rm   rµ  r   r£   Útorchbind_constantsr  r[   rz   r~  rQ   Úbacked_var_to_valrå   r   r#   rb  Úget_real_objr   ré  rd  r(  rC  rD  r  r­  rP   r½   r  )r
  rœ   rÿ  r  r	  r›   rñ   Útorchbind_objr¦   rJ  rm   s   ``         rV   Úbenchmark_compiled_moduleÚ.PythonWrapperCodegen.benchmark_compiled_moduleÎ  s�  ù€ ö	õ	0õ	ð 	×ÑÒ5Ô6Ø�]‰]�_Ø�M‰Mðð ð	 ñ ô  !Ÿw™w×0Ñ0×6Ñ6Ö8‘�ð × Ñ  7¨4¨&Ð!1Ô2ÙØŸ*™*›,¨¯©«¸¿¹ÀeÇkÁköñ	  9ô ”1—7‘7×.Ñ.Ó/°!Ó3Ø× Ñ  Ô1Ü+,¯7©7×+FÑ+F×+LÑ+LÖ+NÑ'�Dð ×$Ñ$ w¨t¨fÐ%5Ô6Ù'¨Ö<ñ	 ,Oô  !Ÿw™w×3Ñ3×9Ñ9×;‘�Ü˜e¤U§\¡\×2Ñ2´zÜ—G‘G×$Ñ$×6Ñ6×:Ñ:¸5À$ÓGÌ÷8ñ 8ñ Ü˜e¤R×%7Ñ%7×8Ñ8Üœ1Ÿ7™7×6Ñ6Ó7¸1Ó<à×(Ñ(¨Ô9Ø×$Ñ$ w¨t¨fÐ%5Ô6Ù'¨×.@Ñ.@Ó.BÖCÜ ¤u§z¡z×2Ñ2ñ
 #ØœaŸg™g×.Ñ.×@Ñ@ÀÐQSÐ@ÐT÷ô   ¤r×'8Ñ'8×9Ñ9Ù"ØØ8¸¿¹×9KÑ9KÐ8LÐLcÐd÷ð "'§¡Ô!1óâ!1˜Aô Ÿ™×(Ñ(×:Ñ:¸1ÀrÐ:ÓJÙ!1ð ð ð "'×!1Ñ!1Ô!3óâ!3˜Aô Ÿ™×(Ñ(×2Ñ2°1¸rÐ2ÓBÙ!3ð ð ñ #ØØØØ×(Ñ(Ó*ØŸ™Ó)÷ñI  <ðX ×Ñ˜x¨¯	©	´!·'±'×2FÑ2F×2KÑ2KÓ2MÓ(NÐ'OÈqÐQÔR÷I ðN 	×ÑÚQô	
ð �]‰]�_Ø�M‰Mðð
 ð ñ ÷ ˆ_ùò-ùò÷q �_ú÷T �_ús7   ¹KQÌ2QÍ QÍ2Q	ÎA9QÐ*QÑ
QÑ
QÑ
Q-c                ó  • [         R                  (       d  gU R                  U5        UR                  / SQ5        UR	                  5          UR                  SSS[        5        S3/5        SSS5        g! , (       d  f       g= f)z<
Append a benchmark harness to generated code for debugging
N)rÎ  rÎ  zif __name__ == "__main__":zBfrom torch._inductor.wrapper_benchmark import compiled_module_mainzargs = get_args()zcompiled_module_main('zU', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat)))r"   Úbenchmark_harnessr  r  r·   r+   ©r
  rœ   s     rV   r<  Ú*PythonWrapperCodegen.add_benchmark_harness@	  sr   € ô ×'×'Øà×&Ñ& vÔ.à×ÑÒ@ÔAØ�]‰]�_Ø×ÑàXØ'à0Ô1CÓ1EÐ0Fð Gmð mð	ô	÷ �_Ž_ús   Á!A5Á5
Bc                ó>   • U R                  [        U UUUUUS95        g rÞ  )r—   rØ  )r
  r³  rÙ  rÚ  rÛ  rÜ  s         rV   Údefine_kernelÚ"PythonWrapperCodegen.define_kernelV	  s*   € ð 	�‰Ü ØØØØ!ØØ-ñõ		
rX   c                óÌ   • [         R                  R                  (       a,  U(       a%  [        R                  " SSU[        R
                  S9nU(       a  U S3OSnSU U  SU 3nU$ )Nz^// z# )Úflagsrƒ  rÎ  z

rÞ   )r"   r�   r�   ÚreÚsubÚ	MULTILINE)r³  rÙ  rÚ  Úmetadata_commentÚbodys        rV   Ú_format_kernel_definitionÚ.PythonWrapperCodegen._format_kernel_definitioni	  sZ   € ô �=‰=×1×1¶hô —v’v˜g t¨X¼R¿\¹\ÑJˆHÞ.6˜h˜Z r™?¸BÐØÐ&Ð'¨ }°C¸°}ÐEˆØˆrX   c                ó<  • [         R                  R                  (       aR  U(       aK  U R                  XUS9nU R                  R                  U5        [        R                  R                  (       a  g U R                  XUS9nU R                  R                  U5        g )N)rÚ  )
r"   r�   r�   r"  rÝ  ré   r5   rK   r†  rØ  )r
  r³  rÙ  rÚ  rÛ  rÜ  r!  s          rV   rß  Ú*PythonWrapperCodegen._define_kernel_helperu	  s�   € ô �=‰=×1×1¶cØ×1Ñ1Ø°8ð 2ð ˆDð ×%Ñ%×,Ñ,¨TÔ2Ü�w‰w×"×"àà×-Ñ-Ø¨xð .ð 
ˆð 	�‰×Ñ˜4Õ rX   c                óN   • U R                   R                  UR                  5        g rJ   )rÞ  ré   rñ   )r
  r›   Úsubgraph_codes      rV   Údefine_subgraph_launcher_fnÚ0PythonWrapperCodegen.define_subgraph_launcher_fn‹	  s   € Ø×!Ñ!×(Ñ(¨×)<Ñ)<Õ=rX   c                ó   • SSK Jn  U$ )Nr6   )ÚTritonKernel)r�   r+  )Úclsr+  s     rV   Ú_get_triton_info_kernel_clsÚ0PythonWrapperCodegen._get_triton_info_kernel_clsŽ	  s   € õ 	)àÐrX   c                óX  ^^/^0^1^2^3^4• SSK JnJnJn	  SSKJm/Jn
JnJnJ	n  UR                  n/ m40 m2/ m1/ nU1U44S jm0S-U/U0U2U4S jjnUR                   Vs/ s H  nUR                  PM     nnUR                   Vs/ s H"  nUR                  (       d  M  UR                  PM$     nn[        U5       GHÒ  u  nnUU;   a  U" UT/" US9SS	9  M  UT;  a  M#  TU   nTU   c  U" UT/" US9SS
9  M<  [!        U["        R$                  5      (       a^  [!        U["        R&                  5      (       a'  SUR(                  UR*                  R-                  5       4OSu  nnnU" UU" UUUUS95        M¹  [!        U["        R.                  5      (       a.  U" UU" UUR1                  5       UR-                  5       S95        GM  [!        U["        R2                  5      (       aM  U" UU" UUR4                  R1                  5       UR-                  5       UR6                  R8                  S95        GMr  [!        U[:        [<        R>                  45      =(       a*    [@        RB                  RD                  RG                  US5      nU" UU" UU5      US9  GMÕ     [I        T4S T1URJ                   Vs/ s H  n[M        U5      PM     snS9nU[N        RP                  " [@        RB                  RS                  5       5      0 T2E[T        RW                  US5      E[Y        T4T1S9/S.nU(       a  [[        U5      US'   U(       a  [[        U5      US'   []        U5      S:X  a0  UR_                  5       n/ [a        [<        Rb                  US   5      QnOöS.U34S jjn 0 m3U V!s/ s H  n!/ [a        U U!5      QPM     nn!U(       a  []        U5      []        U5      :X  d   e/ n"[e        [g        Xb5      S SS9 HR  u  n!n#U"Ri                  U" U#5      / [a        [j        U!5      Q/ [a        [l        U!5      Q/ [a        [j        U!5      QS.5        MT     U	R                  U"/ [a        [n        T3Rq                  5       5      QS.n/ T3Rs                  5       Qn[u        URv                  5      /n$[]        U5      S:”  aY  TRq                  5        HE  n[!        U["        R.                  ["        R2                  45      (       a  M4  U$Ri                  U5        MG     U$Ri                  [o        U5      5        U$Ry                  [o        U5      5        [[        U$5      n$U$U Rz                  ;   a  U Rz                  U$   u  n%nn&U%UU&U4$ U S[]        U Rz                  5       3n%[}        5       n'[~        R€                  R‚                  (       a  U'R…                  SU%< S35        OU'R…                  SU< S35        U%US'   U R‡                  5       n(UR‰                  U(R‹                  5       5        U'R�                  U(R�                  5       5        [~        R€                  R�                  (       a  U'R…                  S 5        U'R�                  S!/ [a        Xr5      Q< S"U< S#U< S$35        [“        U5      n)[~        R€                  R‚                  (       a  U)R•                  S%U S&3S%U% S&35      n)U)R•                  S'S(5      n)U'R�                  U)5        [@        RB                  RS                  5       n*U'R…                  S)U*R–                   S*35        [˜        Rš                  " URv                  5      u  n+n,[˜        Rœ                  " URv                  5      n-S+U- S,U, 3n.U RŸ                  U%U'R¡                  5       U.5        U%UU4U Rz                  U$'   U%UUU4$ s  snf s  snf s  snf s  sn!f )/Nr    )Úconfig_to_dictÚ	FixedGridÚPrecomputedGridr6   )ÚConstexprArgÚKernelArgTypeÚSizeArgÚ	TensorArgÚTMADescriptorArgc                óJ   >• TR                  U5        TR                  U 5        g rJ   )rq   )r¤  r¸  Úarg_indicesÚ	signatures     €€rV   Úadd_to_signatureÚPPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_to_signature´	  s   ø€ Ø×Ñ˜SÔ!Ø×Ñ˜sÕ#rX   c                óØ  >• U(       aE  [        5       (       a  T" X5        UR                  T;   a  TUR                     TUR                  '   g g UR                  T;   d   eU(       a?  [        5       (       a  T" U T" UR                  S95        OT" X5        STUR                  '   g U(       a6  [        5       (       a  T" U T" UR                  S95        S TUR                  '   g T" X5        g )N©r›   r6   )r4   r›   )	r¤  r¸  Úis_constexprÚequals_1Úequals_noner3  r;  r  r¤   s	        €€€€rV   Úadd_argÚGPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_arg¸	  sÅ   ø€ ÞÜ1×3Ñ3ñ % SÔ.à—8‘8˜vÓ%ð +1°·±Ñ*:�I˜cŸh™hÒ'ð &ð —x‘x 6Ó)Ð)Ð)æÜ5×7Ñ7ñ
 )¨©lÀÇÁÑ.IÕJá(¨Ô2Ø*+�I˜cŸh™hÒ'Þ Ü5×7Ñ7ñ )¨©lÀÇÁÑ.IÔJØ*.�I˜cŸh™hÒ'á$ SÕ.rX   r>  T)r?  )rA  Ústable)ÚexperimentalNN)r›   Úapi_typerô  r   )r›   Úbufferr   )r›   rG  r   rn   )r@  )Ú
size_dtyper  Úargdefs)r  )r:  rµ  r  rÀ   Úrestore_valueÚreset_to_zeror   c                ót  >• [        U [        R                  5      (       al  / U R                  QnU(       d  U $ UR	                  [
        S9  U H0  nUT;   a  M  [        R                  " S[        T5       35      TU'   M2     [        U T5      $ [        U [        5      (       d   e[        R                  " U 5      $ )N)rª   Ú_launcher_s)r[   rz   r   r…  Úsortrµ   r~  r£   r3   rï   r{   )r†  Úsymbolsr‡  Úextra_launcher_argss      €rV   Úrename_sizes_for_launcherÚYPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.rename_sizes_for_launcherA
  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Ó*Ð*rX   c                ó2   • [        U S   R                  5      $ r¡   r¢   r¥   s    rV   r§   ÚHPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.<lambda>W
  s   € ´3°q¸±t·{±{Ô3CrX   r©   )r"   r¹  ÚcppÚpython_slow)Ú	grid_typeÚprecomputed_gridsrP  r  zasync_compile.triton(z, '''r³  r  zG
            @triton_heuristics.user_autotune(
                configs=z ,
                inductor_meta=z,
                triton_meta=z{,
                filename=__file__,
                custom_kernel=True,
            )
            @triton.jit
            rŸ   rï  z'''z\'\'\'z''', device_str='r~  z# Original path: r…  )FFF)r†  r±   r²   r³   )QÚruntime.triton_heuristicsr0  r1  r2  Úcommonr3  r4  r5  r6  r7  rò   Úparamsr›   r?  Únumr  r[   r#   ÚTMADescriptorrù  rô  rë  rP   rj   rM   r&   rr   r_   rn   rï   rz   r{   r5   rK   rQ   Ústatically_known_equalsrA   Ú	arg_namesr7   r(   r  Úget_current_device_or_throwr†  Úfromkeysr?   r�   r£   Úsetup_grid_as_argsr  Úsympifyr»   r¼   rq   rœ  r=   rµ   r£  r  Úidrã   Úextendré  r-   r"   r�   Úunique_user_kernel_namesr—   r-  ÚupdateÚinductor_meta_commonré   Úgen_common_triton_importsr%  r  Úreplacer‰   ÚinspectÚgetsourcelinesÚgetsourcefiler  r¿   )5r
  r  rÀ   r¤   Úrestore_value_argsÚreset_to_zero_argsrÁ   r0  r1  r2  r4  r5  r6  r7  Úoriginal_nameÚequal_to_1_argsrB  Úpr_  Ú
constexprsr¤  rª   r¸  rF  rô  r   r@  r¦   Útriton_signaturerÍ  rÎ  Úextra_launcher_call_argsrQ  r‘   rX  ÚcfgÚ	cache_keyr›   Úcached_inductor_metarü   Útriton_info_kernel_clsÚ
kernel_srcÚcurrent_devicer  ÚlinenoÚsrcfilerÚ  r3  r;  r9  r  rP  r:  s5      `                                           @@@@@@rV   Ú!define_user_defined_triton_kernelÚ6PythonWrapperCodegen.define_user_defined_triton_kernel˜	  s*  þ€ ÷	
ñ 	
÷
	
õ 	
ð Ÿ™ˆØ)+ˆ	Ø$&ˆ	Ø!#ˆØ%'ˆö	$÷"	/ò "	/ðH &,§]¢]Ó3¢] �Q—V”V¡]ˆ	Ð3Ø%+§]¢]ÓE¢] °a·nµn“e�a—e”e¡]ˆ
ÐEÜ! )×,‰HˆC�Ø�jÓ Ù˜™\¨sÑ3À$ÒGÙà˜&Ó Ùà˜‘+ˆCà�c‰{Ñ"Ù˜™\¨sÑ3ÀÔFä˜c¤2×#3Ñ#3×4Ñ4ô & c¬2×+AÑ+A×BÑBð " 3§?¡?°C·J±J×4HÑ4HÓ4JÑKà9ñ 1�H˜k¨5ñ
 ØÙ(Ø!$Ø%-Ø(3Ø"'ñ	öô   ¤R§Y¡Y×/Ñ/ÙØÙ!Ø!$Ø#&§<¡<£>Ø"%§-¡-£/ñ÷ô   ¤R×%7Ñ%7×8Ñ8ñ ØÙ!Ø!$Ø#&§8¡8×#4Ñ#4Ó#6Ø"%§-¡-£/Ø#&§:¡:×#4Ñ#4ñ	÷ô  *Øœc¤5§=¡=Ð1ó ÷  äŸ'™'×*Ñ*×BÑBØØóð ñ ˜C¡¨¨cÓ!2¸XÕFñs -ôv -ØØØØ)/×)9Ò)9Ó:Ò)9 A”W˜Q–ZÑ)9Ñ:ñ	
Ðð *Ü&×-Ò-¬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Ó;Ó<ÉˆEÐOæœS ›Z¬3¨w«<Ó7Ð7Ð7Ø "ÐÜ#Ü�EÓ#Ñ)CÈTô‘	��cð "×(Ñ(á"0°Ó"5Ø"5¤C¬¨tÓ$4Ð"5Ø2¤¤U¨DÓ!1Ð2Ø':¬¬U°DÓ)9Ð':ñ	öñð -×5Ñ5Ø%6Ø'P¬¬SÐ2E×2LÑ2LÓ2NÓ)OÐ'PñˆMð
 (EÐ)<×)AÑ)AÓ)CÐ'DÐ$ô ˜VŸY™Y›-˜ˆ	Üˆw‹<˜!ÓØ—}‘}–�ä! #¬¯	©	´2×3EÑ3EÐ'F×GÓGØ×$Ñ$ SÖ)ñ 'ð 	×Ñœ˜[Ó)Ô*Ø×Ñœ˜]Ó+Ô,Ü˜)Ó$ˆ	Ø˜×6Ñ6Ó6Ø6:×6TÑ6TØñ7Ñ3ˆD�+Ð3ð ØØ$Ø(ð	ð ð  � ¤# d×&DÑ&DÓ"EÐ!FÐGˆä(Ó*ˆÜ�=‰=×1×1Ø×%Ñ%Ð(=¸d¹XÀUÐ&KÕLà×%Ñ%Ð(=¸mÑ=NÈeÐ&TÔUà'+ˆ�mÑ$Ø!%×!AÑ!AÓ!CÐØ×ÑÐ3×HÑHÓJÔKà×ÑÐ5×OÑOÓQÔRÜ�=‰=×)×)Ø×%Ñ%Ð&DÔEà×Ñðà8œ3˜~Ó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:¸;ÈÐ4Vˆ×&Ñ& yÑ1Ø�[ -Ð1IÐIÐIùòa 4ùÚEùò@ ;ùòj Ps   Á^Á:^Â^Ê^"Î^'c                ó   • U SUR                    S3nUb  USU 3-  n[        R                  " USSS9n[        XRR                  5      nUS:H  nU(       d)  U R                  [        X[        R                  5      5        U$ )Nr  ÚnumelT)Ú
is_integerÚis_positiverÎ  )	rÙ  rz   r~  r  r�  r—   r·  r5   rK   )r
  r³  ÚtreerÚ  Úsym_namer‡  r¸  Úis_benchmark_kernels           rV   Úgenerate_numel_exprÚ(PythonWrapperCodegen.generate_numel_expr®
  s{   € Ø!�] ! D§K¡K =°Ð6ˆØÑØ˜!˜F˜8˜Ñ$ˆHÜ�lŠl˜8°À$ÑGˆô ˜c§:¡:Ó.ˆà)¨RÑ/ÐÞ"Ø�N‰NÔ.¨t¼!¿'¹'ÓBÔCàˆ
rX   c                ój   • U R                  UR                   S[        UR                  5       35        g rý  )r—   r  rœ  r  )r
  r¸  rK   s      rV   rº  Ú7PythonWrapperCodegen._generate_symbolic_call_arg_helperÃ
  s)   € ð 	�‰˜#Ÿ)™)˜ C¬¨c¯n©nÓ(=Ð'>Ð?Õ@rX   c                óè  • UR                  5       n[        X5      nUR                  [        R                  :X  a  U R                  U5        GOBUR                  [        R                  :X  a2  U R                  U5        U R                  U R                  U5      5        OòUR                  [        R                  :X  a¿  U R                  R                  U5      nU(       a]  [        U[        5      (       a  [        UR                  [        5      (       d   e[        R                  " UR                  U5      Ul        OUU R                  U5        U R                  U R                  U5      5        X0R                  U'   O[        UR                  5      e[         R"                  R$                  (       aÅ  U R&                  R                  [(        R+                  U UUR,                  UR.                  [0        R2                  R4                  R7                  UR8                  5      4SS95        UR                  [        R                  :w  a/  U R&                  R                  [(        R                  X5      5        g g g )N)r6   )rJ  rm   )rM   r"  Ú	zero_moder<   ÚUNINITIALIZEDr—   ÚZERO_ON_CALLÚmake_zero_bufferÚZERO_PER_GRAPHr   rå   r[   rS   r;   Úmaximumr‡  r"   r�   r�   r™   rE  Úmake_allocationrµ  r   r5   rK   rQ   r  r   )r
  Úwsr›   rš   Úpriors        rV   Úgenerate_workspace_allocationÚ2PythonWrapperCodegen.generate_workspace_allocationÈ
  sÉ  € Ø�{‰{‹}ˆÜ˜DÓ%ˆØ�<‰<Ô,×:Ñ:Ó:Ø�N‰N˜4Ö Ø�\‰\Ô.×;Ñ;Ó;Ø�N‰N˜4Ô Ø�N‰N˜4×0Ñ0°Ó6Õ7Ø�\‰\Ô.×=Ñ=Ó=Ø×-Ñ-×1Ñ1°$Ó7ˆEÞÜ! %¬×6Ñ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õð ?ð 2rX   c                óv   • UR                   [        R                  :w  a  U R                  [	        X5      5        g g rJ   )rŒ  r<   r�  r—   r4  )r
  r“  s     rV   Úgenerate_workspace_deallocationÚ4PythonWrapperCodegen.generate_workspace_deallocationï
  s,   € Ø�<‰<Ô,×;Ñ;Ó;Ø�N‰NÔ.¨tÓ8Õ9ð <rX   c                ó$   • U SU R                    3$ )Nz.zero_())r�  )r
  r›   s     rV   r�  Ú%PythonWrapperCodegen.make_zero_bufferó
  s   € Ø��x §¡˜}Ð-Ð-rX   c                óH   • U SSR                  U5       SU R                   3$ )Nrï  rð  rÛ   )r½   r�  )r
  r›   rÈ  s      rV   r  Ú%PythonWrapperCodegen.wrap_kernel_callö
  s'   € Ø��q˜Ÿ™ 9Ó-Ð.¨a°·±¨}Ð=Ð=rX   c                óú   • U R                   R                  S5        U R                   R                  S[        R                  R                   S35        UR                  U R                   R                  5       5        g )Nz*from torch.profiler import record_functionzwith record_function('graph_z_inductor_wrapper_call'):)rÜ  r—   r5   rK   Úgraph_idr/  r·   )r
  r>  s     rV   r1  Ú8PythonWrapperCodegen.generate_profiler_mark_wrapper_callù
  sb   € Ø×Ñ×#Ñ#Ð$PÔQØ×Ñ×#Ñ#Ø*¬1¯7©7×+;Ñ+;Ð*<Ð<UÐVô	
ð 	×Ñ˜D×-Ñ-×4Ñ4Ó6Õ7rX   c                ó:   • U R                   R                  S5        g )Nzstart_graph())rÜ  r—   r	  s    rV   r2  Ú)PythonWrapperCodegen.generate_start_graph   s   € Ø×Ñ×#Ñ# OÕ4rX   c                ó`   • U R                   R                  S[        R                  < S35        g )Nz
end_graph(rÛ   )rÜ  r—   r"   Úprofile_bandwidth_outputr	  s    rV   r6  Ú'PythonWrapperCodegen.generate_end_graph  s'   € Ø×Ñ×#Ñ# j´×1PÑ1PÑ0SÐSTÐ$UÕVrX   c                ó†   • U R                   R                  [        R                  R                  R                  5       5        g)z<Synchronize GPU to ensure proton captures all kernel events.N)rÜ  r—   r5   rK   rˆ  r€  r	  s    rV   r7  Ú-PythonWrapperCodegen.generate_proton_finalize  s)   € à×Ñ×#Ñ#¤A§G¡G×$6Ñ$6×$BÑ$BÓ$DÕErX   c                ó^   • U R                   R                  S[        R                   S35        g )NúU
            for kernel in globals().values():
                if isinstance(kernel, zU.CachingAutotuner):
                    kernel.cuda_kernel_saved = False
            ©rÜ  ré   r'   rò   r	  s    rV   r5  Ú6PythonWrapperCodegen.generate_reset_kernel_saved_flags
  s2   € Ø×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð Cðõ	
rX   c                ó^   • U R                   R                  S[        R                   S35        g)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(
                            stream="stream",  # use dummy stream
                            launcher=kernel.launchers[0],
                        )
            Nrª  r	  s    rV   r8  Ú5PythonWrapperCodegen.generate_save_uncompiled_kernels  s4   € ð 	×Ñ× Ñ ð'ä'8×'AÑ'AÐ&Bð Cð
õ	
rX   c                óB   • S nU Vs/ s H
  o2" U5      PM     sn$ s  snf )Nc                ó  • [        U [        5      (       a  [        U 5      (       a  U S-   $ U $ [        U [        [        [
        [        45      (       a  [        U 5      $ [        [        R                  R                  R                  U 5      5      $ )Nrô  )r[   rµ   r@   rï   Úfloatrð   r  rœ  r5   rK   rQ   rR   )r¸  s    rV   Úwrap_argÚAPythonWrapperCodegen.prepare_triton_kernel_call.<locals>.wrap_arg-  sg   € Ü˜#œs×#Ñ#ä*BÀ3×*GÑ*G�s˜Y‘ÐPÈSÐPÜ˜C¤#¤u¬d´OÐ!D×EÑEÜ˜3“x�äœQŸW™W×-Ñ-×6Ñ6°sÓ;Ó<Ð<rX   r�   )r
  rÈ  r±  r¸  s       rV   Úprepare_triton_kernel_callÚ/PythonWrapperCodegen.prepare_triton_kernel_call,  s%   € ò	=ñ *3Ó3ª #�˜–©Ñ3Ð3ùÒ3s   ˆc                óÞ  ^ • [        U[        5      (       GaY  [        U[        R                  5      (       a.  UR	                  5       R                  5       nT R                  U   nObT R                  R                  U5      (       a  UnT R                  U   nO0Uc   S5       eST R                   3nUnT =R                  S-  sl        Uc
   SU 35       e[        R                  R                  R                  UR                  5       5      n[        R                  R                  R                  [        R                  R                  U5      5      n[        R                  R                  R                  UR                  5       5      nUR!                  5       n	UR#                  5       n
[        R                  R                  R%                  UR'                  5       R(                  5      nSU SU SU	 SU
 SU SU S	3nT R*                  R-                  U S
U 35        [        U[        R                  5      (       a2  T R/                  USS9nUnT R*                  R-                  U S
U 35        U$ [1        U[2        R4                  5      (       d  [        U[6        5      (       aÏ  [        U[8        5      (       a  UT R:                  ;   a  U$ Uc  gUn[        U[6        5      (       a  UR<                  nU[        R                  R                  R>                  ;   a'  [        R                  R                  R>                  U   n[9        [        R                  R                  R%                  U5      5      $ [        U[8        [@        [B        [D        45      (       a  [9        U5      $ [        U[F        5      (       a  SSRI                  U 4S jU 5       5       S3$ [K        S[M        U5       35      e)NzBV.graph.get_buffer(arg) and raw_arg can't be None at the same timeÚtmp_arg_r6   z Failed to find a buffer for arg zgenerate_example_value(rð  z, 'z', rÛ   rÞ   T)rî  rï  r5  r‚  c              3  óZ   >#   • U  H   nTR                  U[        U5      5      v •  M"     g 7frJ   r‡   )r‚   Úar
  s     €rV   r„   ÚBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>s  s(   øé € Ð ZÒVYÐQR ×!@Ñ!@ÀÄDÈÃG×!LÐ!LÒVYùr‹   rƒ  zUnsupported type )'r[   Útorch_dtyper#   r]  Ú
get_tensorrM   r×  rå   rà  r5   rK   rQ   rê  rC  Úget_allocation_sizerD  r­  rP   ré  rk   rn   r™   r—   rú  r„  rz   ÚBasicr  rµ   rý  r  r‘  rï   r°  rð   r  r½   r;  r‰   )r
  r¸  Úarg_typeÚraw_argr  re  rl   Úallocation_sizerm   rµ  r   rn   rñ   s   `            rV   rˆ   Ú/PythonWrapperCodegen.generate_example_arg_value8  s<  ø€ Ü�h¤×,Ò,Ü˜'¤2×#3Ñ#3×4Ñ4à"×-Ñ-Ó/×8Ñ8Ó:�Ø×*Ñ*¨3Ñ/‘Ø×%Ñ%×)Ñ)¨#×.Ñ.Ø�Ø×*Ñ*¨3Ñ/‘àÑ*ð ØXóÐ*ð & d×&FÑ&FÐ%GÐH�Ø�Ø×0Ò0°AÑ5Õ0à‘?ÐLÐ&FÀsÀeÐ$LÓL�?Ü—7‘7×#Ñ#×6Ñ6°s·|±|³~ÓFˆDÜŸg™g×.Ñ.×AÑAÜ—‘×+Ñ+¨CÓ0óˆOô —W‘W×%Ñ%×8Ñ8¸¿¹Ó9IÓJˆFà—^‘^Ó%ˆFØ—M‘M“OˆEÜ—W‘W×%Ñ%×7Ñ7¸¿¹Ó8H×8OÑ8OÓPˆFØ-¨d¨V°2°f°X¸SÀÀÈÈEÈ7ÐRTÐU[ÐT\Ð\^Ð_nÐ^oÐopÐqˆEØ×&Ñ&×0Ñ0°H°:¸SÀÀÐ1HÔIä˜'¤2×#3Ñ#3×4Ñ4ð ×:Ñ:Ø Ø%)ð ;ð �ð �Ø×*Ñ*×4Ñ4¸°zÀÀUÀGÐ5LÔMàˆOÜ˜¤%§+¡+×.Ñ.´*¸SÄ/×2RÑ2Rä˜#œs×#Ñ#Ø˜$Ÿ/™/Ó)Ø�JØ‘?Ø!Ø�Ü˜#œ×/Ñ/Ø—n‘n�Ø”a—g‘g×&Ñ&×CÑCÓCÜ—g‘g×&Ñ&×CÑCÀCÑH�ä”q—w‘w×'Ñ'×9Ñ9¸#Ó>Ó?Ð?ä˜œc¤3¬¬tÐ4×5Ñ5Ü�s“8ˆOÜ˜œT×"Ñ"Ø�t—y‘yÔ ZÑVYÓ ZÓZÐ[Ð[\Ð]Ð]ä%Ð(9¼$¸s»)¸Ð&EÓFÐFrX   c                ó†   ^ • [        U[        5      (       a!  SSR                  U 4S jU 5       5      -   S-   $ [        U5      $ )Nr‚  rð  c              3  óF   >#   • U  H  nTR                  U5      v •  M     g 7frJ   )Ú_grid_dim_str)r‚   r|   r
  s     €rV   r„   Ú5PythonWrapperCodegen._grid_dim_str.<locals>.<genexpr>z  s   øé € ÐRÂ\¸T × 2Ñ 2°4× 8Ð 8Â\ùs   ƒ!rƒ  )r[   r  r½   rœ  )r
  Úgrid_per_dims   ` rV   rÄ  Ú"PythonWrapperCodegen._grid_dim_strw  s?   ø€ Ü�l¤D×)Ñ)à�d—i‘iÔRÁ\ÓRÓRÑRÐUXÑXðô ˜Ó&Ð&rX   )rµ  r�   rÌ  rÉ  rÊ  rÍ  rÎ  rÂ   c               ó”  • U R                   R                  U Vs0 s H:  n[        U[        5      (       d  M  U[        R
                  R                  U5      _M<     sn5        U=(       d    [        R
                  R                  5       nU R                  [        U UUUUUUUU	U[        R
                  R                  U
S95        gs  snf )z®
Generates kernel call code.

triton: Defines whether the backend uses Triton for codegen. Otherwise it uses the CUDA language when gpu=True,
        and C++ when gpu=False.
)r³  rÈ  rÉ  rÊ  rÌ  r�   rÍ  rÎ  rµ  rÏ  rÂ   N)r×  rg  r[   rµ   r5   rK   Útry_get_bufferr`  r—   rÇ  r›   )r
  r³  rÈ  rµ  r�   rÌ  rÉ  rÊ  rÍ  rÎ  rÂ   r¸  s               rV   Úgenerate_kernel_callÚ)PythonWrapperCodegen.generate_kernel_call  s³   € ð. 	×Ñ×#Ñ#ñ %óâ$�CÜ˜c¤3×'ó 1�”Q—W‘W×+Ñ+¨CÓ0Ò0Ù$ñô	
ð ×@œ1Ÿ7™7×>Ñ>Ó@ˆØ�‰ÜØØ'Ø#à!à!à#Øà'Ø+ØÜŸ7™7Ÿ<™<à%9ñ#õ	
ùòs
   šC·$CrÎ  )	rµ  r�   rÌ  rÉ  rÊ  rÍ  rÎ  rÏ  rÂ   c       	   
     óJ  ^ ^^• U=(       d    [         R                  R                  5       nU(       d�  UR                  S:w  a€  UR                  S:X  a"  T R	                  T R                  TU5      5        g UR                  S:X  a%  T R	                  T R                  T S3U5      5        g [        SUR                   S35      eT R                  U5      nSR                  U5      n[        R                  T UR                  U
5      nU(       d$  SU S	3nT R	                  T S
T SU SU S	35        g T R                  5         [        R                  R                  (       GaÆ  TT R                   ;  Gaµ  Ub  [#        U5      [#        U5      :X  d   S5       eS mU(       aI  [         R                  R$                  (       a*  [         R                  R$                  R'                  US 5      mSUU 4S jjnUU 4S jn/ nUc)  Ub   S5       eS /[#        U5      -  nS /[#        U5      -  nO[#        U5      [#        U5      :X  d   S5       e0 n[)        [+        X%Xg5      5       GH°  u  nu  nnnnS n[-        U[.        5      (       a#  S[/        U5      ;   a  UR1                  S5      u  nnS nT(       a  UT;   a  T R3                  TU   5      nU(       aQ  Un[-        U[4        5      (       d9  [7        U[8        R:                  5      (       d  [-        U[<        5      (       a  UUU'   OªUS:X  a  U" XgUU5      (       a  UU   nO�[-        U[4        5      (       ag  [>        R@                  " SU5      (       a  UnO6UT RB                  ;  a  T RE                  UUU5      nOT RB                  U   S   nUT4T RB                  U'   OT RE                  UUU5      n[-        U[.        5      (       a  [G        U5      (       a  US-  nURI                  Uc  UOU SU 35        GM³     T RJ                  R	                  S[         R                  RL                  RO                  UR                  5       S35        T RJ                  RQ                  5         T RJ                  R	                  T SSR                  U5       SU S	35        T RJ                  RS                  5         T RJ                  R	                  [U        SUS5      5        T R                   RW                  T5        [         R                  RX                  (       a  g [         R                  RZ                  R\                  nUR_                  UTUS 5        U   T R	                  T SU SU S	35        S S S 5        T R                  5         g ! , (       d  f       N= f)NÚcudar4  Úmpsz.generated_kernelzdevice z nyirð  z	c_void_p(rÛ   rÝ   rï  z$call_args and arg_types do not matchc                 ó¶   >• TR                   R                  5        V Vs/ s H  u  pUT:X  d  M  U PM     nn nU(       a  SSR                  U5       S3$ gs  snn f )zÖAfter all the autotune kernel calls have been written (i.e.
self.kernel_autotune_example_args is complete), returns a deletion call
for all autotune example tensors that are unnecessary after kernel_name
is called.údel rð  rƒ  rÎ  )rß  r£  r½   )rë  ÚknÚtensors_to_deleter³  r
  s      €€rV   Úget_autotune_deletion_callÚUPythonWrapperCodegen._generate_kernel_call_helper.<locals>.get_autotune_deletion_callî  se   ø€ ð '+×&GÑ&G×&NÑ&NÔ&Pô%â&P™
˜Ø˜[Ñ(÷ Ù&Pð "ñ %ö
 %Ø! $§)¡)Ð,=Ó">Ð!?¸rÐBÐBØùó%s
   ŸA¯Ac                óŒ  >• X   nXC;   a  g[        [        X5      5       H�  u  nu  pgXR:X  d  [        U[        5      (       d  M$  SnT(       a  UT;   a  TR	                  TU   5      nUS:X  a  MO   UR                  5       n	[        U	R                  5       H  u  p«X´:X  d  M  U SU
 S3X4'       g   M’     g! [         a     M¢  f = f)z­We try to infer raw_arg (i.e. raw_args[idx]) from remaining raw_args.
This is particularly useful for jagged cases, where the dimension is often
being passed in as an input.TrÎ  z.shape[rƒ  F)r  r¼   r[   r%   rO  rk   rl   r;  )rÉ  rÊ  r¤  Úreused_argsÚ
target_argrŠ  Úraw_keyr¿  Útriton_inputr_   râ  rg  Úautotune_argsr
  s               €€rV   Úinfer_arg_by_inputsÚNPythonWrapperCodegen._generate_kernel_call_helper.<locals>.infer_arg_by_inputsü  sÝ   ø€ ð
 &™]�
ØÓ,Øä-6´s¸8Ó7NÖ-OÑ)�AÑ)˜Ø“x¤z°'¼6×'BÑ'BÙ à#%�LÞ$¨°MÓ)AØ'+×'EÑ'EØ)¨'Ñ2ó(˜ð $ rÓ)Ù ð	!Ø!(×!3Ñ!3Ó!5˜Ü&/°·±Ö&<™F˜CØ �Ø=I¸NÈ'ÐRUÐQVÐVWÐ:X Ñ 7Ú'+ó '=ñ .Pð, øô	 /ó !ò !ð!ús   Á0-B5Â!B5Â.B5Â5
CÃCzkeys are not None but args arez#call_args and raw_args do not matchrî  rÎ  z^(workspace|semaphore)r   rô  r„  r…  z.run(z	, stream=z
<del_call>rø  )0r5   rK   r`  r‰   r—   r  rZ  r³  r½   rE  rú  r(  rC  r"   r�   r�   r˜   r£   Úautotuning_mappingrå   r  r¼   r[   rµ   ÚsplitrO  rº  r„  rz   r½  r  r  Úmatchrß  rˆ   r@   rq   r™   rˆ  r‹  rM  rx  r*   r¾   r†  rÖ  r  r×  )r
  r³  rÈ  rµ  r�   rÌ  rÉ  rÊ  rÍ  rÎ  rÏ  rÂ   Úcall_args_strÚstream_nameÚ
stream_ptrrÓ  rÛ  Úall_argsrÖ  rŠ  r¸  r¾  rØ  r¿  rª   rÙ  Úarg_strrÛ  rÚ  s   ``                          @rV   rÑ  Ú1PythonWrapperCodegen._generate_kernel_call_helperµ  s  ú€ ð ×@œ1Ÿ7™7×>Ñ>Ó@ˆÞ˜&Ÿ+™+¨Ó/Ø�{‰{˜eÓ#Ø—‘˜t×4Ñ4°[À)ÓLÔMð ð —‘ Ó%à—‘Ø×)Ñ)¨[¨MÐ9JÐ*KÈYÓWôð
 ô # W¨V¯[©[¨M¸Ð#>Ó?Ð?ð ×7Ñ7¸	ÓBˆØŸ	™	 -Ó0ˆÜ*×?Ñ?Ø�&—,‘, 
ó
ˆö Ø$ [ M°Ð3ˆJØ�N‰NØ�-˜q  ¨Q¨}¨o¸RÀ
¸|È1ÐMôð à×%Ñ%Ô'ô �M‰M×2×2Ð2Ø 4×#=Ñ#=Ô=ð Ñ(¬S°«^¼sÀ9»~Ó-Mð Ø6óÐMð !ˆMÞ#¬¯©×(B×(BÜ !§¡× :Ñ :× >Ñ >Ø(¨$ó!�÷ð öðB ˆHØÑàÑ'ÐIÐ)IÓIÐ'Ø ˜6¤C¨	£NÑ2�Ø ˜6¤C¨	£NÑ2‘ä˜8“}¬¨I«Ó6ð Ø9óÐ6ð ˆKÜ8Aä�I¨(Ó=÷9Ñ4�Ñ4�C˜ 7¨Gð �Ü˜c¤3×'Ñ'¨C´3°s³8«Oà"Ÿy™y¨›~‘H�C˜à.2�Þ  W°Ó%=Ø#'×#AÑ#AØ% gÑ.ó$�Lö  Ø*�GÜ% h´×<Ñ<Ü" 8¬U¯[©[×9Ñ9Ü% c¬?×;Ñ;à/6˜ GÑ,øØ “]Ñ':Ø¨¨;÷(ñ (ð
 *¨'Ñ2‘GÜ ¬+×6Ñ6ô —x’xÐ 9¸3×?Ñ?Ø"%™Ø D×$EÑ$EÓEØ"&×"AÑ"AØ ¨7ó#™ð #'×"CÑ"CÀCÑ"HÈÑ"K˜Ø>EÀ{Ð=S�D×5Ñ5°cÒ:à"×=Ñ=¸cÀ8ÈWÓU�GÜ˜c¤3×'Ñ'Ô,DÀS×,IÑ,IØ˜yÑ(�GØ—‘¨3©;¡¸s¸eÀ1ÀWÀIÐ<N×OñY9ð^ ×&Ñ&×0Ñ0ØœŸ™×*Ñ*×7Ñ7¸¿¹ÓEÐFÀaÐHôð ×&Ñ&×0Ñ0Ô2Ø×&Ñ&×0Ñ0Ø�-˜u T§Y¡Y¨xÓ%8Ð$9¸À;À-ÈqÐQôð ×&Ñ&×2Ñ2Ô4à×&Ñ&×0Ñ0Ü  Ð/IÈ<ÓXôð ×&Ñ&×*Ñ*¨;Ô7Ü�w‰w×"×"àô !"§¡× 4Ñ 4× BÑ BÐØ×.Ñ.¨y¸+ÀyÐRVÔWÚ"Ø�N‰N˜k˜]¨%°¨¸iÈÀ}ÐTUÐVÔW÷ #à×%Ñ%Õ'÷ #Õ"ús   Õ VÖ
V"c                ó:   • U R                   R                  U5        g rJ   )rã  rq   ©r
  rš   s     rV   r—   ÚPythonWrapperCodegen.writelinep  s   € Ø�
‰
×Ñ˜$ÕrX   c                ó8   • U H  nU R                  U5        M     g rJ   r#  )r
  rã  rš   s      rV   r  ÚPythonWrapperCodegen.writeliness  s   € ÛˆDØ�N‰N˜4Ö ò rX   c                óL   • U R                   R                  [        U5      5        g rJ   )rã  rq   r0   )r
  Úctxs     rV   r/  Ú"PythonWrapperCodegen.enter_contextw  s   € Ø�
‰
×Ñœ+ cÓ*Õ+rX   c                ó@  ^ ^	• SSK Jn  U" 5       (       a  SS Kn[        U[        5      (       a  [        UR                  R                  5      $ [        U[        R                  5      (       a  [        U5      $ [        U[        [        45      (       aB  [        R                   " S S5      5       m	[        [        U5      " U	U 4S jU 5       5      5      $ [        U[         R"                  R$                  5      (       a  ['        U5      $ [        U[(        R*                  [(        R,                  [.        45      (       a  UR1                  5       $ U" 5       (       a0  [        UWR2                  R4                  5      (       a  [        U5      $ [        U[(        R6                  5      (       a  UR1                  5       $ [9        [        U5      5      (       aE  [;        U5      u  pVUR=                  5        H"  u  pxU[>        R@                  RB                  U'   M$     U$ [        U5      $ )Nr   )Úhas_triton_packagec                  ó&   • \ rS rSr% S\S'   S rSrg)Ú1PythonWrapperCodegen.val_to_arg_str.<locals>.Shimi†  r   Úrefc                ó   • U R                   $ rJ   )rò  r	  s    rV   Ú__repr__Ú:PythonWrapperCodegen.val_to_arg_str.<locals>.Shim.__repr__Š  s   € ØŸ8™8�OrX   r�   N)rò   rà   r  r  rÙ   rô  r  r�   rX   rV   ÚShimrñ  †  s   ‡ à“õ$rX   rö  c              3  ó\   >#   • U  H!  nT" [         R                  TU5      5      v •  M#     g 7frJ   rß  )r‚   r¸  rö  r
  s     €€rV   r„   Ú6PythonWrapperCodegen.val_to_arg_str.<locals>.<genexpr>�  s)   øé € ÐVÒTUÈq™Ô1×@Ñ@ÀÀqÓI×JÐJÒTUùs   ƒ),)"Útorch.utils._tritonrï  r�   r[   r   rœ  rS   r†  rz   r   r�   r  rñ  Ú	dataclassrO  r‰   r  Ú_opsÚ
OpOverloadr   r#   rj   Ú
MutableBoxr&   r°  Úlanguager   rd  r   r   ró  r5   rK   Úopaque_value_type_classes)
r
  rg  Útype_rï  r�   Úobj_reprÚopaque_typesÚnr›  rö  s
   `        @rV   rà  Ú#PythonWrapperCodegen.val_to_arg_strz  sˆ  ù€ Ý:á×ÑÛä�aœ×"Ñ"Ü˜Ÿ™Ÿ™Ó%Ð%Ü˜œ5Ÿ:™:×&Ñ&Ü˜“8ˆOÜ˜œE¤4˜=×)Ñ)ä×"Ñ"÷$ð $ó #ð$ô Ü�Q”ÕVÑTUÓVÓVóð ô ˜œ5Ÿ:™:×0Ñ0×1Ñ1Ü& qÓ)Ð)Ü˜œBŸI™I¤r§}¡}´oÐF×GÑGØ×&Ñ&Ó(Ð(Ù×!Ñ!¤j°°F·O±O×4IÑ4I×&JÑ&JÜ˜“7ˆNÜ˜œ2×,Ñ,×-Ñ-Ø×&Ñ&Ó(Ð(Ü!¤$ q£'×*Ñ*Ü%8¸Ó%;Ñ"ˆHØ$×*Ñ*Ö,‘�Ø78”—‘×1Ñ1°!Ó4ñ -àˆOä˜“7ˆNrX   c           	     óZ  • UR                  5       nUR                  5       n[        UR                  5       5      n[        [        R
                  R                  U5      5      n[        UR                  5       5      nUR                  5       nU R                  UR                  5       X#XFXW5      $ rJ   )r­  rP   r�   rC  r5   rK   r¼  rD  Úget_is_pinnedr’  rM   )r
  rG  rµ  r   rJ  Úallocation_shaperm   Ú	is_pinneds           rV   r@  Ú+PythonWrapperCodegen.make_buffer_allocation¢  sŽ   € Ø×"Ñ"Ó$ˆØ× Ñ Ó"ˆÜ�f—o‘oÓ'Ó(ˆÜ ¤§¡×!<Ñ!<¸VÓ!DÓEÐÜ�v×(Ñ(Ó*Ó+ˆØ×(Ñ(Ó*ˆ	Ø×#Ñ#Ø�O‰OÓ˜v¨eÐ=Mó
ð 	
rX   c                ó|   • Sn[         R                  R                  (       d  U R                  R	                  USS9  g g )Nzi
            from torch._inductor.runtime.debug_utils import check_memory_step, track_tensor
            TrÎ   )r5   rK   r†  rÑ  ré   rA  s     rV   Ú"write_memory_track_allocation_onceÚ7PythonWrapperCodegen.write_memory_track_allocation_once­  s4   € ðˆ
ô �w‰w×"×"Ø�L‰L×Ñ 
°$ÐÒ7ð #rX   c                ó  • Uc  UnU R                  U5      nU R                  U5      n	U R                  U5      n
[        R                  R                  R                  R
                  (       a  U SU	 SU
 SU SUR                   SU S3nOmUR                  S:X  a  U(       a  U SU	 SU
 SU S	3nOGUR                  S
;   a  U SUR                   SU	 SU
 SU S	3
nOU SU	 SU
 SUR                   SU S	3
nX‰:w  a  USU SU
 S	3-   nU$ )Nz = tracked_empty_strided(rð  z, dtype=rý  z	', name='r~  r4  z = empty_strided_cpu_pinned(rÛ   )r4  rÍ  ÚxpuÚmtiaz = empty_strided_rï  z = empty_strided(rþ  z.as_strided()rŽ   r  r  r"   r  r   r‰   )r
  r›   rµ  r   rJ  rm   r  r  rG  Úcodegen_allocation_shape_tupleÚcodegen_stride_tuplerØ  s               rV   r’  Ú$PythonWrapperCodegen.make_allocationµ  s…  € ð Ñ#Ø$Ðà"×=Ñ=¸eÓDÐØ)-×)HÑ)HØó*
Ð&ð  $×>Ñ>¸vÓFÐÜ�?‰?×!Ñ!×.Ñ.×E×Eà�&Ð1Ø1Ð2°"Ø'Ð(ð )Ø˜ð  Ø!Ÿ;™;˜-ð (Ø˜˜bð"ñ ð �[‰[˜EÓ!¦ià�&Ð4Ø1Ð2°"Ø'Ð(¨Ø�'˜ðñ ð �[‰[Ð:Ó:ð �&Ð)¨&¯+©+¨°aØ1Ð2°"Ø'Ð(¨Ø�'˜ðñ ð �&Ð)Ø1Ð2°"Ø'Ð(ð )Ø!Ÿ;™;˜- y°°°qð:ð ð Ó@à˜,Ð':Ð&;¸2Ð>RÐ=SÐSTÐUÑUˆCØˆ
rX   c                ó8   • U R                  [        U5      5        g rJ   )r—   rc  rç  s     rV   Úmake_commentÚ!PythonWrapperCodegen.make_commentå  s   € Ø�‰”{ 4Ó(Õ)rX   c           	     ó`   • U R                    U SU U R                   SU R                   SU 3	$ )NrÞ   ú  Ú )rŽ  r�  rå  )r
  Únew_nameÚold_namerå  s       rV   Úmake_tensor_aliasÚ&PythonWrapperCodegen.make_tensor_aliasè  s6   € Ø—,‘,� ˜z¨¨X¨J°t·{±{°mÀ2ÀdÇlÁlÀ^ÐSTÐU\ÐT]Ð^Ð^rX   c                ó(   • SUR                  5        3$ )NrÐ  )rM   )r
  rG  s     rV   rÀ  Ú%PythonWrapperCodegen.make_buffer_freeë  s   € Ø�f—o‘oÓ'Ð(Ð)Ð)rX   c                ó8   • SSR                  S U 5       5       3$ )NrÐ  rð  c              3  ó$   #   • U  H  ov •  M     g 7frJ   r�   )r‚   r›   s     rV   r„   Ú:PythonWrapperCodegen.make_free_by_names.<locals>.<genexpr>ï  s   é € Ð>²¨¤²ùs   ‚)r½   )r
  Únames_to_dels     rV   Úmake_free_by_namesÚ'PythonWrapperCodegen.make_free_by_namesî  s   € Ø�d—i‘iÑ>±Ó>Ó>Ð?Ð@Ð@rX   c           	     ó`   • U R                    U SU U U R                   SU R                   S3	$ )NrÞ   r  ú reuse)rä  r�  rå  )r
  r  r  Údel_lines       rV   Úcodegen_exact_buffer_reuseÚ/PythonWrapperCodegen.codegen_exact_buffer_reuseñ  s@   € Ø×.Ñ.Ð/°¨z¸¸X¸JÀxÀjÐQU×Q\ÑQ\ÐP]Ð]_Ð`d×`lÑ`lÐ_mÐmsÐtÐtrX   c                óR   • Ub$  U R                  U R                   SU SU 35        g g )Nz [Provenance debug handles] r…  )r—   rå  )r
  r³  Údebug_handles      rV   Úwrite_provenance_debug_handleÚ2PythonWrapperCodegen.write_provenance_debug_handleô  s4   € ð
 Ñ#Ø�N‰NØ—<‘<�.Ð <¸[¸MÈÈ<È.ÐYõð $rX   c                ó‚  • UR                  5       UR                  5       :X  d   eUR                  5       nUR                  5       nSnU[        R                  R	                  5       ;  a  U(       a  SU R                  U5       3nUR                  5       UR                  5       :X  a4  UR                  5       UR                  5       :X  a  U R                  XEU5      $ U R                  XR                  5       UR                  5       SU R                  R                  5      nU R                   U SU U SU R                   S3$ )NÚ;z; r   rÞ   r  r&  )rP   rM   r5   rK   r  rÀ  rC  rD  r(  rÁ  rÜ  r—   rŽ  rå  )r
  r%  r$  rm  r  r  r'  Úreinterpret_views           rV   rq  Ú&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ÓPÐPà×8Ñ8Ø—‘“ §¡Ó!1°1°d×6GÑ6G×6QÑ6Qó
Ðð —,‘,� ˜z¨Ð-=Ð,>¸x¸jÈÈ4Ï<É<È.ÐX^Ð_Ð_rX   c                ó¬   • U R                  [        UU R                   U SUR                  5        U R                   SU R
                   S35      5        g )NrÞ   r  z alias)r—   r9   rŽ  r°  r�  rå  )r
  r›   rf  s      rV   rg  Ú0PythonWrapperCodegen.codegen_deferred_allocation  sS   € Ø�‰ÜØØ—<‘<�.   c¨$×*@Ñ*@Ó*BÐ)CÀDÇKÁKÀ=ÐPRÐSW×S_ÑS_ÐR`Ð`fÐgóõ	
rX   c                óÞ  ^• UR                  5       nU[        R                  R                  ;   dM  X R                  ;   d>  [        U[        R                  [        R                  [        R                  45      (       a  g U R                  R                  U5        [        UR                  5       [        R                  [        R                  45      (       a  UR                  5       (       d  g UR                  5       n[        U[        R                   5      (       a  g [        U[        R"                  5      (       a  g [        U[        R$                  5      (       Ga?  [        UR&                  [        R(                  5      (       d*   S[+        UR&                  5       SUR&                   35       eUR&                  R,                  n[        U[        R.                  5      (       d   [+        U5      5       eUR,                  n[        U[        R0                  [        R(                  45      (       d   [+        U5      5       e[        U[        R(                  5      (       a  SU4S jjmT" U5      nU R3                  U5        U R5                  [7        XX5      5        g [        U[        R8                  5      (       a  U R5                  [;        XSS95        g U R5                  [;        X5      5        g )Nzunexpected rÜ   c                ó2  >• [        U [        R                  5      (       a  T" U R                  5       5      $ [        U [        R                  5      (       a  T" U R
                  5      $ [        U [        R                  5      (       d   [        U 5      5       eU $ rJ   )r[   r#   ÚBaseViewÚunwrap_viewrý  rr   rj   r‰   )ÚtargetÚunwrap_viewss    €rV   r9  Ú=PythonWrapperCodegen.codegen_allocation.<locals>.unwrap_views:  sm   ø€ Ü! &¬"¯+©+×6Ñ6Ù+¨F×,>Ñ,>Ó,@ÓAÐAÜ! &¬"¯-©-×8Ñ8Ù+¨F¯K©KÓ8Ð8Ü% f¬b¯i©i×8Ñ8ÐF¼$¸v»,ÓFÐ8Ø!�MrX   Tr3  )r²   ú	ir.Buffer)rM   r5   rK   r¿  rö  r[   r#   ÚDonatedBufferÚSubgraphBufferÚInputBufferr¾   Úget_defining_opÚExternKernelAllocÚMultiOutputÚshould_allocaterZ   ÚMutationLayoutSHOULDREMOVErÉ  re  rf  r&   r‰   rr   rp   rj   Úcodegen_allocationr—   r`  r\   r"  )r
  rG  r›   r_   ÚboxÚinput_bufferr9  s         @rV   rD  Ú'PythonWrapperCodegen.codegen_allocation  s$  ø€ Ø�‰Ó ˆð ”A—G‘G×+Ñ+Ó+Ø—~‘~Ó%Ü˜&¤2×#3Ñ#3´R×5FÑ5FÌÏÉÐ"W×XÑXàØ�‰×Ñ˜4Ô äØ×&Ñ&Ó(Ü×%Ñ%¤r§~¡~Ð6÷ñ ð ×*Ñ*×,Ñ,àà×'Ñ'Ó)ˆÜ�fœb×;Ñ;×<Ñ<ØÜ�fœbŸm™m×,Ñ,ØÜ�fœb×0Ñ0×1Ò1Ü˜fŸk™k¬2×+=Ñ+=×>Ñ>ð Øœd 6§;¡;Ó/Ð0°°6·;±;°-Ð@óÐ>ð —+‘+×"Ñ"ˆCÜ˜c¤2§=¡=×1Ñ1Ð<´4¸³9Ó<Ð1ØŸ8™8ˆLÜ˜l¬R¯Y©Y¼×8JÑ8JÐ,K×LÑLð ÌdØóOó ÐLô ˜,¬×(:Ñ(:×;Ñ;÷"ñ  ,¨LÓ9�Ø×#Ñ# LÔ1Ø�N‰Nœ?¨4¸vÓNÔOØä�fœb×1Ñ1×2Ñ2Ø�N‰Nœ<¨À$ÑGÔHØà�‰”| DÓ1Õ2rX   c                óÞ  • UR                  5       n[        U[        R                  [        R                  45      (       a  U R                  [        X5      5        g [        UR                  5       [        R                  5      (       a  U R                  [        XSS95        g U R                  U5      (       d  g U R                  R                  U5        U R                  [        X5      5        g )NTr3  )rM   r[   r#   r>  rb  r—   r¼  rZ   r\   r4  Ú	can_reuser÷  r¾   )r
  rG  r›   s      rV   Úcodegen_freeÚ!PythonWrapperCodegen.codegen_freeM  s©   € Ø�‰Ó ˆô �fœrŸ~™~¬r×/AÑ/AÐB×CÑCØ�N‰Nœ8 DÓ1Ô2Øä�f×,Ñ,Ó.´×0CÑ0C×DÑDð �N‰NÔ.¨tÈÑNÔOØà�~‰~˜f×%Ñ%ØØ�
‰
�‰�tÔà�‰Ô*¨4Ó8Õ9rX   c                ó2  • UR                  5       nU[        R                  R                  ;   =(       dß    U[        R                  R                  ;   =(       a:    [        [        R                  R                  U   [        R                  5      (       + =(       dz    U[        R                  R                  ;   =(       dV    U[        R                  R                  ;   =(       d2    U[        R                  R                  ;   =(       d    X0R                  ;   (       + $ rJ   )rM   r5   rK   r¿  r  r[   Úgraph_inputs_originalr#   r<  r  r  Únever_reuse_buffersr÷  )r
  rF  Úoutput_bufferr›   s       rV   rI  ÚPythonWrapperCodegen.can_reusea  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÷
"ð —z‘zÑ!ô
ð 	
rX   c                ó¦   • UR                  5       U R                  ;   =(       a.    U R                  UR                  5          UR                  5       :H  $ rJ   )rM   rø  )r
  rG  Úreused_buffers      rV   Ú	did_reuseÚPythonWrapperCodegen.did_reuseq  sC   € ð �O‰OÓ §¡Ñ,÷ KØ—‘˜FŸO™OÓ-Ñ.°-×2HÑ2HÓ2JÑJð	
rX   c                óz  • [        X5      (       d   eU R                  U5        U R                  R                  UR	                  5       5        U R
                  R                  UR	                  5       5        UR	                  5       U R                  UR	                  5       '   U R                  [        XU5      5        g rJ   )	rh   rD  r÷  r¾   rM   rö  rø  r—   r7  )r
  rF  rO  s      rV   Úcodegen_inplace_reuseÚ*PythonWrapperCodegen.codegen_inplace_reusey  sˆ   € Ü$ \×AÑAÐAÐAØ×Ñ Ô-Ø�
‰
�‰�|×,Ñ,Ó.Ô/Ø�‰×Ñ˜=×1Ñ1Ó3Ô4Ø0<×0EÑ0EÓ0Gˆ�‰�M×*Ñ*Ó,Ñ-Ø�‰”y °]ÓCÕDrX   c                óŽ   • [        U5      nX R                  ;   a  U$ U R                  R                  U5        U R                  U-   $ rJ   )rµ   rê  r¾   rŽ  )r
  rø   r›   s      rV   Úcodegen_unbacked_symbol_declÚ1PythonWrapperCodegen.codegen_unbacked_symbol_decl�  sA   € Ü�6‹{ˆØ×-Ñ-Ó-ØˆKð ×&Ñ&×*Ñ*¨4Ô0Ø—<‘< $Ñ&Ð&rX   c                ó–   • [        [        R                  R                  R                  U5      nU R                  [        XX#5      5        g rJ   )r   r5   rK   rQ   Ú	shape_envr—   rÀ  )r
  rÁ  rÂ  rÄ  s       rV   Ú(codegen_unbacked_symbol_defs_for_outputsÚ=PythonWrapperCodegen.codegen_unbacked_symbol_defs_for_outputsŠ  s=   € ô 6Ü�G‰G×Ñ×&Ñ&Ð(9ó
Ðð 	�‰Ü" 4°gÓQõ	
rX   c                óÞ   ^^^^• U(       d  g UR                  5        HM  u  nmSU4S jjmUUUU4S jnU R                  U R                  U5       SU" 5        U R                   35        MO     g )Nc                óü  >• US:X  a  U $ [        U5      S:¼  ai  [        US   [        5      (       aQ  [        US   [        R                  5      (       a/  T" U  SUS   R
                   SUS   R                   S3USS  5      $ [        US   [        5      (       a  T" U  SUS   R
                   S3USS  5      $ [        US   [        R                  5      (       a^  [        R                  R                  (       a   T" S	US   R                   S
U  S3USS  5      $ T" U  SUS   R                   S3USS  5      $ [        US   [        5      (       a  T" U  SUS   R                   S3USS  5      $ [        SU 35      e)Nr�   r    r   r6   rÝ   rï  rÛ   r¥  z	std::get<z>(r‚  rƒ  z.__floordiv__(rö  )r£   r[   r   ÚpytreeÚSequenceKeyr›   r¤  r5   rK   r†  r   rù  r‡  )r†  rø  Úgos     €rV   rc  ÚJPythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs.<locals>.goª  s•  ø€ Ø˜b“=Ø�Kô ˜“L AÓ%Ü" 7¨1¡:¬}×=Ñ=Ü" 7¨1¡:¬v×/AÑ/A×BÑ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Ñ8ñ    n°W¸Q±Z×5GÑ5GÐ4HÈÐJÈGÐTUÐTVÈKÓXÐXä(Ð+@ÀÀ	Ð)JÓKÐKrX   c                 óÎ  >• [         R                  R                  (       a½  [        T5      S:X  a`  TS   n T" TS   R	                  5       [        U [        R                  5      (       a"  [        U R                  5      S:w  a	  TSS  5      $ T5      $ [        TS   [        R                  5      (       d   eT" TTS   R                     R	                  5       TSS  5      $ T" TT5      $ )Nr6   r   )r5   rK   r†  r£   rM   r[   r#   rA  r  ra  rb  r¤  )rØ  rc  rø  rÁ  rÂ  s    €€€€rV   Úgo_outerÚPPythonWrapperCodegen._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ÐIÙ! '¨'°!©*¯.©.Ñ"9×"BÑ"BÓ"DÀgÈaÈbÀkÓRÐRá˜k¨7Ó3Ð3rX   rÞ   )r†  rµ   rø  zpytree.KeyPath)ró  r—   rY  r�  )r
  rÁ  rÂ  rÄ  rg  rf  rc  rø  s    ``   @@rV   rÆ  Ú>PythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs—  sg   û€ ö !Øð ,×1Ñ1Ö3‰JˆAˆw÷
L÷<4ð 4ð. �N‰NØ×4Ñ4°QÓ7Ð8¸¹H»J¸<ÈÏÉÀ}ÐUöòu 4rX   c                óä  ^ ^^^• UU U4S jnUU U4S jn T R                  TR                  5        T R                  T R                   STR                   35        U" 5         [
        R                  n[
        R                  " TR                  5         TR                  R                  US9  S S S 5        U" 5         T R                  5         g ! , (       d  f       N&= f! T R                  5         f = f)Nc                 ó  >• [        TR                  R                  5      [        T5      :X  d   e[        TR                  R                  T5       H3  u  pTR	                  TR
                   U  SU TR                   35        M5     g rý  )r£   rK   r  r¼   r—   rŽ  r�  )Úinner_inputÚouter_inputÚouter_inputsr
  Úsubgraphs     €€€rV   Ú_codegen_subgraph_prefixÚSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_prefixñ  sr   ø€ Ü�x—~‘~×2Ñ2Ó3´s¸<Ó7HÓHÐHÐHÜ,/Ø—‘×+Ñ+¨\ö-Ñ(�ð —‘Ø—|‘|�n [ M°°[°MÀ$Ç+Á+ÀÐOöò-rX   c                 ó  >• [        TR                  R                  5      [        T5      :X  d   e[        TR                  R                  T5       H5  u  pTR	                  U SU R                  5        TR                   35        M7     g rý  )r£   rK   r
  r¼   r—   r°  r�  )Úinner_outputÚouter_outputÚouter_outputsr
  rn  s     €€€rV   Ú_codegen_subgraph_suffixÚSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_suffixú  st   ø€ Ü�x—~‘~×3Ñ3Ó4¼¸MÓ8JÓJÐJÐJÜ.1Ø—‘×,Ñ,¨mö/Ñ*�ð —‘Ø#�n C¨×(FÑ(FÓ(HÐ'IÈ$Ï+É+ÈÐWöò/rX   ú subgraph: )Úparent_graph)	rL  rK   r—   rå  r›   r5   Úset_graph_handlerÚcodegen_subgraphrw  )r
  rn  rm  rt  ro  ru  rx  s   ````   rV   Úcodegen_subgraph_by_inliningÚ1PythonWrapperCodegen.codegen_subgraph_by_inliningã  s¯   û€ ÷	÷	ð	'Ø×%Ñ% h§n¡nÔ5Ø�N‰N˜dŸl™l˜^¨;°x·}±}°oÐFÔGÙ$Ô&ÜŸ7™7ˆLÜ×$Ò$ X§^¡^Õ4Ø—‘×/Ñ/Ø!-ð 0ñ ÷ 5ñ %Ô&à×$Ñ$Õ&÷ 5Õ4ûð ×$Ñ$Õ&ús$   –A<C ÂCÂ,C Ã
CÃC ÃC/c           	     óä  • UR                   nUR                  n[        UR                  5       5      UR                   Vs/ s H  oUR
                  PM     sn-   nSR                  U5      [        U5      S:X  a  SOS-   nU Vs/ s H  oˆR                  5       PM     n	nSR                  U	5      [        U5      S:X  a  SOS-   n
U R                  SU SU S35        UR                  5        VVs/ s H  u  p¼U(       d  M  UPM     nnnU(       a#  U R                  SSR                  U5       35        U R                  S	U
 S
U SU S35        U R                  SU S35        gs  snf s  snf s  snnf )z'Generate code to call a graph partitionrð  r6   ro  rÎ  Ú	partitionú	_args = [rƒ  rÐ  rï  z) = self.partitions[z](partitionú_args)zdel partitionÚ_argsN)Úinput_deallocationÚoutput_nodesr  r  Úsymbol_inputsr›   r½   r£   rM   r—   ró  )r
  Úpartition_idr  r‚  rƒ  Úsymbol_inputrp  rŸ  rS   Úoutput_namesrÂ  r›   Ú
deallocater"  s                 rV   Úcodegen_partition_callÚ+PythonWrapperCodegen.codegen_partition_call  sn  € ð 2×DÑDÐØ+×8Ñ8ˆäÐ-×2Ñ2Ó4Ó5Ø2F×2TÒ2Tó9
Ú2T ,×ÔÑ2Tñ9
ñ 
ˆð —‘˜;Ó'´#°kÓ2BÀaÓ2G©3ÈRÑPˆá4@ÓA²L¨DŸ™ž±LˆÐAØ—)‘)˜LÓ)´C¸Ó4EÈÓ4J©SÐPRÑSˆð 	�‰˜ < .°	¸&¸ÀÐCÔDð *<×)AÑ)AÔ)Cô
Ú)CÑ%�TÄz�DÑ)Cð 	ñ 
ö Ø�N‰N˜T $§)¡)¨LÓ"9Ð!:Ð;Ô<ð 	�‰Ø�ˆyÐ,¨\¨N¸+ÀlÀ^ÐSYÐZô	
ð 	�‰˜ | n°EÐ:Õ;ùò-9
ùò Bùó
s   ¿E"ÂE'Ã-E,Ã>E,c                óV   • [        U5       Vs/ s H  nSU 3PM
     snU l        g s  snf )NÚ
partition_)rj  r±  )r
  Únum_partitionsr¤  s      rV   Úset_all_partition_namesÚ,PythonWrapperCodegen.set_all_partition_names1  s*   € ÜBGÈÔBWÓ#XÒBW¸3 j°°Ó$6ÑBWÑ#XˆÕ ùÒ#Xs   Ž&c           	     ót  • SR                  U5      [        U5      S:X  a  SOS-   nSR                  U5      [        U5      S:X  a  SOS-   nU R                  UR                  R                   SU S35        U R                  SU SUR                  R                   SUR                  R                   S	35        g )
Nrð  r6   ro  rÎ  r  rƒ  rï  z) = r€  )r½   r£   r—   rK   r›   )r
  rn  rm  Úouter_flattened_outputsÚouter_output_namesÚouter_input_namess         rV   Ú,codegen_subgraph_call_with_flattened_outputsÚAPythonWrapperCodegen.codegen_subgraph_call_with_flattened_outputs4  s¹   € ð "ŸY™YÐ'>Ó?ÜÐ.Ó/°1Ó4‰C¸"ñ
Ðð !ŸI™I lÓ3Ü�|Ó$¨Ó)‰C¨rñ
Ðð 	�‰˜(Ÿ.™.×-Ñ-Ð.¨iÐ8IÐ7JÈ!ÐLÔMð 	�‰ØÐ"Ð# 4¨¯©×(;Ñ(;Ð'<¸A¸h¿n¹n×>QÑ>QÐ=RÐRXÐYõ	
rX   c                óx  • SR                  U5      [        U5      S:X  a  SOS-   nU R                  UR                  R                   SU S35        [
        R                  R                  R                  5         U R                  U SUR                  R                   SUR                  R                   S	35        g )
Nrð  r6   ro  rÎ  r  rƒ  rÞ   rï  r€  )r½   r£   r—   rK   r›   r5   r   Úfree_buffers)r
  rn  rm  Úouter_buffer_namer“  s        rV   Úcodegen_subgraph_callÚ*PythonWrapperCodegen.codegen_subgraph_callF  s¢   € à ŸI™I lÓ3Ü�|Ó$¨Ó)‰C¨rñ
Ðð 	�‰˜(Ÿ.™.×-Ñ-Ð.¨iÐ8IÐ7JÈ!ÐLÔMô 	
�‰×Ñ×&Ñ&Ô(ð 	�‰Ø Ð!  X§^¡^×%8Ñ%8Ð$9¸¸8¿>¹>×;NÑ;NÐ:OÈvÐVõ	
rX   c                ó  • U R                  UR                  5        U R                  S5        U R                  U R                   SUR                   35        [
        R                  nUR                  UR                  l        UR                  UR                  l        UR                  R                  U R                  ;  a©  [
        R                  " UR                  5         [        R                  " SS5         UR                  R                  5       u  p4S S S 5        S S S 5        UR                  R                  nU R                  R                  U5        U R                  UW5        g g ! , (       d  f       N[= f! , (       d  f       Nd= f)NrÎ  rw  ru  F)rL  rK   r  rå  r›   r5   r†  Ú
fx_wrapperrÿ  ry  r"   ÚpatchrP  r¾   r(  )r
  rn  rx  r'  r  r  s         rV   Úcodegen_subgraph_commonÚ,PythonWrapperCodegen.codegen_subgraph_commonW  s  € Ø×!Ñ! (§.¡.Ô1Ø×Ñ˜"ÔØ×Ñ˜TŸ\™\˜N¨+°h·m±m°_ÐEÔFä—w‘wˆØ%1×%=Ñ%=ˆ�‰Ô"Ø$0×$;Ñ$;ˆ�‰Ô!à�>‰>×Ñ d×&FÑ&FÓFô ×$Ò$ X§^¡^Õ4ä—\’\Ð"3°UÕ;à'/§~¡~×'=Ñ'=Ó'?Ñ$�M÷ <÷ 5ð
 %ŸN™N×/Ñ/ˆMØ×,Ñ,×0Ñ0°Ô?Ø×,Ñ,¨]¸MÕJð G÷
 <Õ;ú÷ 5Õ4ús$   Ã!E<Ã9E+ÄE<Å+
E9	Å5E<Å<
F
c                óJ   • U R                  U5        U R                  XU5        g rJ   )rž  r”  )r
  rn  rm  r‘  s       rV   Ú'codegen_subgraph_with_flattened_outputsÚ<PythonWrapperCodegen.codegen_subgraph_with_flattened_outputsl  s&   € ð 	×$Ñ$ XÔ.Ø×9Ñ9ØÐ$;õ	
rX   c                óJ   • U R                  U5        U R                  XU5        g rJ   )rž  r™  )r
  rn  rm  r˜  s       rV   rz  Ú%PythonWrapperCodegen.codegen_subgrapht  s#   € ð 	×$Ñ$ XÔ.Ø×"Ñ" 8Ð;LÕMrX   c                óò  • UR                  5       nU R                  U S[        UR                  5       35        UR                   Vs/ s H  o3R                  5       PM     nn[        R                  R                  (       aP  [        [        UR                  5      5       Vs/ s H
  oR SU S3PM     nnU R                  UR                  XF5        g U R                  UR                  XB5        g s  snf s  snf )Nú = [None] * r‚  rƒ  )rM   r—   r£   rÂ  rŸ  r°  r5   rK   r‡  rj  r{  rn  rz  )r
  Úinvoke_subgraphr›   re  rm  rŠ  rt  s          rV   Úcodegen_invoke_subgraphÚ,PythonWrapperCodegen.codegen_invoke_subgraphz  sØ   € Ø×'Ñ'Ó)ˆà�‰˜$˜˜|¬C°×0GÑ0GÓ,HÐ+IÐJÔKØ;J×;QÒ;QÓRÒ;Q°C×-Ñ-Ö/Ñ;QˆÐRä�7‰7××ä(-¬c°/×2IÑ2IÓ.JÔ(KóÚ(K 1�&˜˜!˜˜A“Ñ(Kð ð ð ×-Ñ-Ø×(Ñ(¨,õð ×!Ñ! /×":Ñ":¸LÕOùò Sùòs   ÁC/Â"C4c                ó  • UR                  5       nUR                   Vs/ s H  o3R                  5       PM     nnUR                  R                  5       n[	        UR                  [
        R                  5      (       d  U S3nU R                  U S[        UR                  5       35        U R                  SU S35        U R                  [        XR                  R                  5      5        [        R                  R                  (       aP  [        [        UR                  5      5       Vs/ s H
  ob SU S3PM     nnU R!                  UR                  XG5        OU R#                  UR                  XB5        U R                  [%        U 5      5        U R                  S5        U R                  [        XR&                  R                  5      5        [        R                  R                  (       aP  [        [        UR                  5      5       Vs/ s H
  ob SU S3PM     nnU R!                  UR&                  XG5        OU R#                  UR&                  XB5        U R                  [%        U 5      5        g s  snf s  snf s  snf )Nrô  r¦  r°   r…  r‚  rƒ  zelse:)rM   Úoperandsr°  Ú	predicater[   r#   ÚShapeAsConstantBufferr—   r£   rÂ  rD  Útrue_subgraphrK   r5   r‡  rj  r{  rz  rr  Úfalse_subgraph)r
  Úconditionalr›   re  rm  r¬  rŠ  rt  s           rV   Úcodegen_conditionalÚ(PythonWrapperCodegen.codegen_conditionalŠ  sþ  € Ø×#Ñ#Ó%ˆà;F×;OÒ;OÓPÒ;O°C×-Ñ-Ö/Ñ;OˆÐPà×)Ñ)×;Ñ;Ó=ˆ	Ü˜+×/Ñ/´×1IÑ1I×JÑJà$˜+ WÐ-ˆIà�‰˜$˜˜|¬C°×0CÑ0CÓ,DÐ+EÐFÔGØ�‰˜˜Y˜K qÐ)Ô*Ø�‰Ô(¨×/HÑ/H×/NÑ/NÓOÔPÜ�7‰7××Ü5:¼3¸{×?RÑ?RÓ;SÔ5TÓUÒ5T°˜v Q q c¨›^Ñ5TˆMÐUØ×-Ñ-Ø×)Ñ)¨<õð ×!Ñ! +×";Ñ";¸\ÔPà�‰Ô'¨Ó-Ô.Ø�‰�wÔØ�‰Ô(¨×/IÑ/I×/OÑ/OÓPÔQÜ�7‰7××Ü5:¼3¸{×?RÑ?RÓ;SÔ5TÓUÒ5T°˜v Q q c¨›^Ñ5TˆMÐUØ×-Ñ-Ø×*Ñ*¨Lõð ×!Ñ! +×"<Ñ"<¸lÔQØ�‰Ô'¨Ó-Õ.ùò9 Qùò Vùò Vs   ŸI8Ä-I=ÈJc                ó,	  ^ • U 4S jnUR                  5       nUR                   Vs/ s H  oUR                  5       PM     nnUR                   Vs/ s H  oUR                  5       PM     nn[	        U5      nT R                  U S[	        U5       35        U(       a   T R                  U S[	        U5       S35        [        U5       H  u  pšT R                  U SU	 SU
 35        M      / [        [	        U5      5       V	s/ s H
  o” SU	 S3PM     sn	QUQnU S3/n[        U5      nUS	[	        U5       nU" UR                  X¼5        T R                  S
US    35        T R                  S5        U(       ax  [        U5       Hh  u  pŸT R                  [        T UR                  R                  5      5        T R                  U SU	 SU S35        T R                  [        T 5      5        Mj     Ow[        U5       Hh  u  pŸT R                  [        T UR                  R                  5      5        T R                  U SU	 SU S35        T R                  [        T 5      5        Mj     T R                  S5        T R                  [        T UR                  R                  5      5        U" UR                  XÞ5        T R                  [        T 5      5        U(       aƒ  T R                  [        T UR                  R                  5      5        [        [	        U5      5       H"  n	T R                  U SX˜-    SU SU	 S35        M$     T R                  [        T 5      5        T R                  [        T UR                  R                  5      5        U" UR                  X¼5        T R                  [        T 5      5        T R                  SUS    35        U(       a¯  T R                  S5        [        [	        U5      5       H…  n	T R                  SU SX˜-    S35        T R                  [        T UR                  R                  5      5        T R                  U SU	 SU SX˜-    S35        T R                  [        T 5      5        M‡     g	g	s  snf s  snf s  sn	f )z1while_loop is codegened as a host side while_loopc                óŽ   >• [         R                  R                  (       a  TR                  XU5        gTR	                  XU5        g)z3Helper method to deduplicate subgraph codegen logicN)r5   rK   r‡  r{  r¡  )rn  rm  rt  r
  s      €rV   rz  ÚAPythonWrapperCodegen.codegen_while_loop.<locals>.codegen_subgraph®  s3   ø€ ä�w‰w××Ø×1Ñ1°(È-ÕXà×<Ñ<Ø¨MõrX   r¦  z.extend([[] for _ in range(z)])r‚  z] = rƒ  Ú_cond_resultNzshould_loop = r   zif not should_loop:z.unsqueeze(0).clone()rÆ  zwhile should_loop:z	].append(z])z    should_loop = z%# Stack outputs after loop completionzif len(z]) > 0:z] = torch.stack(z	], dim=0))rM   Úcarried_inputsr°  Úadditional_inputsr£   r—   r  rj  r  Úcond_subgraphrD  Úbody_subgraphrK   rr  )r
  Ú
while_loopÚstack_outputrz  r›   re  Úouter_carried_inputsÚouter_additional_inputsÚ
ckp_offsetrŠ  ÚinpÚcond_outer_inputsÚcond_outer_outputsÚbody_outer_inputsÚbody_outer_outputsÚcarried_inputs   `               rV   Úcodegen_while_loopÚ'PythonWrapperCodegen.codegen_while_loop«  s]  ø€ õ	ð ×"Ñ"Ó$ˆà/9×/HÒ/Hó 
Ú/H¨×!Ñ!Ö#Ñ/Hð 	ð  
ð 0:×/KÒ/Kó#
Ú/K¨×!Ñ!Ö#Ñ/Kð 	 ð #
ô Ð-Ó.ˆ
Ø�‰˜$˜˜|¬CÐ0DÓ,EÐ+FÐGÔHÞØ�N‰NØ�&Ð3´CÐ8LÓ4MÐ3NÈcÐRôô  Ð 4Ö5‰FˆAà�N‰N˜d˜V 1 Q C t¨C¨5Ð1Ö2ñ 6ð
Ü&+¬CÐ0DÓ,EÔ&FÓGÒ&F ��a˜�s˜!‹nÑ&FÑGð
à$ð
Ðð "&  lÐ3Ð4ÐÜ Øó
Ðð /Ð/J´Ð5IÓ1JÐKÐáØ×$Ñ$Ð&7ô	
ð 	�‰˜Ð(:¸1Ñ(=Ð'>Ð?Ô@Ø�‰Ð,Ô-Þä$-Ð.BÖ$CÑ �Ø—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘ $  q¨¨¨4°¨Ð>SÐTÔUØ—‘Ô/°Ó5Ö6ò %Dô
 %.Ð.BÖ$CÑ �Ø—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘ $  q¨¨¨4°¨¸hÐGÔHØ—‘Ô/°Ó5Ö6ñ %Dð
 	�‰Ð+Ô,à�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOÙØ×$Ñ$Ð&7ô	
ð 	�‰Ô'¨Ó-Ô.ö Ø�N‰NÔ,¨T°:×3KÑ3K×3QÑ3QÓRÔSÜœ3Ð3Ó4Ö5�Ø—‘ $  q¨©Ð(8¸	À$ÀÀqÈÈÈ2ÐNÖOñ 6à�N‰NÔ+¨DÓ1Ô2ð 	�‰Ô(¨¨z×/GÑ/G×/MÑ/MÓNÔOÙØ×$Ñ$Ð&7ô	
ð 	�‰Ô'¨Ó-Ô.Ø�‰Ð+Ð,>¸qÑ,AÐ+BÐCÔDö Ø�N‰NÐBÔCÜœ3Ð3Ó4Ö5�Ø—‘ ¨¨¨a°±Ð/?¸wÐGÔHØ—‘Ô0°°z×7OÑ7O×7UÑ7UÓVÔWØ—‘Ø�f˜A˜a˜SÐ 0°°°a¸¹Ð7GÀyÐQôð —‘Ô/°Ó5Ö6ò 6ð ùòQ 
ùò#
ùò  Hs   ¦RÁRÃ?Rc                óê   •  [        U SS 5      (       a  g [        U [        5      (       a  U $ [        R                  R
                  R                  U 5      nUc  U$ [        U5      $ ! [         a     g f = f)Nr…  )ró  r[   rï   r5   rK   Ú
_shape_envÚ_maybe_evaluate_staticrY  )r¦   rõ  s     rV   Ústatically_known_int_or_noneÚ1PythonWrapperCodegen.statically_known_int_or_none
  sl   € ð	Ü�q˜.¨$×/Ñ/ð Ü˜!œS×!Ñ!Ø�Ü—'‘'×$Ñ$×;Ñ;¸AÓ>ˆCØ‰{Ø�
Ü�s“8ˆOøÜó 	Ùð	ús!   ‚A% •A% ¬-A% Á
A% Á%
A2Á1A2c                ór   • / nU  H.  n[         R                  U5      nUc    g UR                  U5        M0     U$ rJ   )rE  rË  rq   )Úlstr­  r¦   r\  s       rV   Ú%statically_known_list_of_ints_or_noneÚ:PythonWrapperCodegen.statically_known_list_of_ints_or_none  s<   € àˆÛˆAÜ&×CÑCÀAÓFˆCØ‰{ÙØ�M‰M˜#Öñ	 ð
 ˆrX   c                ó0   • [         R                  U 5      S L$ rJ   )rE  rÏ  )rÎ  s    rV   Ú is_statically_known_list_of_intsÚ5PythonWrapperCodegen.is_statically_known_list_of_ints$  s   € ô !×FÑFÀsÓKÐSWÐWð	
rX   c                óH   • [         R                  U R                  5       5      $ rJ   )rE  rÏ  rC  ©rG  s    rV   r8  Ú4PythonWrapperCodegen.static_shape_for_buffer_or_none*  s   € ä#×IÑIØ�O‰OÓó
ð 	
rX   c                ó0   • [         R                  U 5      S L$ rJ   )rE  r8  rÕ  s    rV   Ú!can_prove_buffer_has_static_shapeÚ6PythonWrapperCodegen.can_prove_buffer_has_static_shape0  s   € ä#×CÑCÀFÓKÐSWÐWÐWrX   c                ó   • g rJ   r�   )r
  r³  Únode_schedules      rV   Úwrite_kernel_context_guardÚ/PythonWrapperCodegen.write_kernel_context_guard4  s   € ð
 	rX   c                ó   • g)z,
Mark the beginning of kernel context guard
Nr�   r	  s    rV   Ú write_kernel_context_guard_beginÚ5PythonWrapperCodegen.write_kernel_context_guard_begin;  ó   € ð 	rX   c                ó   • g)z&
Mark the end of kernel context guard
Nr�   r	  s    rV   Úwrite_kernel_context_guard_endÚ3PythonWrapperCodegen.write_kernel_context_guard_endC  rá  rX   )/rý  rü  rÖ  rÒ  r  r±  rö  r   rÿ  r×  rí  rå  rH  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Ú  rè  rê  ré  rÜ  rú  r—   rJ   )r
  rð   r  r¶   r  úOptional[PythonWrapperCodegen]r  ú$Optional[ir.GraphPartitionSignature]rU  )r›   rµ   r  rµ   r²   r5  )rØ  rµ   )rN  ÚTritonMetaParamsr²   rµ   ©r²   rË  )r²   z>dict[str, Union[ir.TensorBox, ir.TorchBindObject, sympy.Expr]]©r²   zlist[IRNode])rp  rË  ©r²   rï   )r  rï   rÏ  rµ   r²   rµ   rø  )r  rï   r²   r5  )r¨  rË  r²   r5  ©r­  r-   r²   r5  )rS   zir.FallbackKernelr²   r5  )rS   r›  )rS   r§  r²   r5  )r  rµ   rØ  rµ   rÙ  r¶   r   rË  rµ  rµ   rÚ  zOptional[OrderedSet[str]]r²   r5  )F)rS   rª  )rS   r–  r²   r5  )r  rµ   r²  rµ   r  zCallable[[], Sequence[str]]r  z<Union[torch._ops.OpOverload, torch._ops.HigherOrderOperator]rÊ  r~  rÂ  zSequence[ir.Buffer]r²   r5  )r$  úCallable[..., None]r²   zIterator[Callable[..., None]])r  rð   )r›   rµ   rñ   r‹  r  rŒ  )r‡  r  )r¦   r   rR   rð   r²   rµ   )r¦   r   r²   rµ   )rˆ  rµ   r›   rµ   r(  rµ   r²   rµ   )rJ  zSequence[Expr]r²   rµ   )r²   ztuple[str, list[str]])r—   rì  r²   rµ   )rÅ  zUnion[bool, str])rS   zir.MultiOutput)NTN)
r³  rµ   rÙ  rµ   rÚ  r¶   rÛ  rð   rÜ  r¶   )r³  rµ   rÙ  rµ   rÚ  r¶   )r›   rµ   )rÁ   z"list[list[Union[int, sympy.Expr]]])r³  rµ   rÚ  r¶   )r¸  r  rK   rD   r²   r5  )r“  r;   )r³  rµ   )rG  r!  )NF)rÎ  )rG  r½  )r"  rË  )r  rµ   r  rµ   r'  rµ   )r+  r€  )r%  r!  r$  r!  rm  rð   )r›   rµ   rf  zir.ReinterpretViewr²   r5  ©rG  r;  )rF  r;  rO  r;  )rÁ  rµ   rÂ  r   rÄ  rÃ  r²   r5  )r…  rï   r  zir.GraphPartitionSignature)r�  rï   )r³  rµ   rÛ  z0Union[Sequence[BaseSchedulerNode], ExternKernel])žrò   rà   r  r  rQ  rÌ  rÙ   r  ra  r  rì  rô  rï  r7  rñ  r)   rC  rH  rK  rQ  rV  rY  r]  rT  rf  ri  rl  rr  rx  r{  rð  rƒ  rú  r�  rL  rw  rG  rt  rž  r¡  r¤  r©  r®  r³  r¶  r¼  rÂ  rŸ  rÒ  r¯  rñ  rõ  rú  rþ  r  r±  r  r¡  r  r  r   r¸   Úcontextmanagerr&  r*  r  r9  rc  rp  r3  r€  r‚  r’  r;  r™  r�  r   r…  rŽ   rG  r¬  rÁ  rÆ  rÊ  rÔ  rì  rï  rl  r  r<  r  r"  rß  r(  Úclassmethodr-  r~  r‡  rº  r•  r˜  r�  r  r1  r2  r6  r7  r5  r8  r³  rˆ   rÄ  rÊ  rÑ  r—   r  r/  rà  r@  r  r’  r  r  rÀ  r#  r(  r,  rq  rg  rD  rJ  rI  rS  rV  rY  r]  rÆ  r{  r‰  rŽ  r”  r™  rž  r¡  rz  r¨  r±  rÆ  rË  rÏ  rÒ  r8  rØ  rÜ  rß  rã  r  r7  r8  s   @rV   rE  rE  5  s¨  ø‡ ñð "Ð�dÓ!õ_#ðB ð
 FJð	&Øð&à$ð&ð 7ð&ð Cô	&ó ð&ô'ô<ôoBôbôð8 óó ðô"	Bð ó+ó ð+ô!ð ó
ó ð
ô
ð$à	Gô$ô
%ôSô$	(ô
ô.ôô8)ô:ô6-ôò.ò1ò0òJò
/ô+ô7ô&5ô
5ô$ôôô:ô,òð:8à ð8ð 
ô8ð 37ð;àð;ð ð;ð  ð	;ð
 ð;ð ð;ð 0ð;ð 
õ;ô ô(	ôUòô
8òô(Bò<ð
	Vàð	Vð  ð	Vð .ð		Vð
 Rð	Vð  ð	Vð %ð	Vð 
ô	Vò0ôð ×Ñó!ó ð!ô,òS
òj.Sò`:ò
 
ôD%ð(Kàð(Kð ð(Kð -ô	(KòT'?ôRDòð @D÷ Wð CG÷ +ô.ô&ô'ô6ðà	ôð, ð<Pð 'ð<Pð 
õ<Pô|>ôSò
IòGò26ò4ò*pòdð4 #'ØØ(,ð
àð
ð ð
ð  ð	
ð
 ð
ð &õ
ð& àFJð	Øð	Ø'*ð	Ø6Cô	ó ð	ð #'ØØ(,ð!àð!ð ð!ð  ð	!ð
 ð!ð &õ!ô,>ð ñó ððTJð 2ôTJölð*AØ"ðAØ+8ðAà	ôAô
%ôN:ò.ò>ò8ò5òWòFò
ò
ò2
4ô=Gò~'ð ØØØØØØØ!ñ4
àõ4
ðv ØØØØØØØØ!ñy(àõy(òv ò!ò,ô%ôP	
ð ñ8ó ð8ð TYô.ò`*ô_ô*ôAôuð '+ðð $õô`ô 
ô53òn:ô(
ò 
ôEò'ð
àð
ð ð
ð Hð	
ð
 
ô
ðJàðJð ðJð Hð	Jð
 
ôJòX+'ðZ<àð<ð 9ô<ôBYò
ò$
ò"Kò*
òNòPô /òB]7ð~ ñó ðð ñó ðð ñ
ó ð
ð
 ñ
ó ð
ð
 ñXó ðXðàðð Hôò÷ð rX   rE  c                  ó  ^ • \ rS rSrSr S     SU 4S jjjrSS jrSS jrS rS r	S r
SS	 jrSS
 jrSS jrSS jr  SS jrSS jrSS jrSU 4S jjr\SS j5       r\SS j5       r\S S j5       rS rSrU =r$ )!r	  iL  zü
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                óî   >• Xl         X l        X0l        [        TU ]  5         U R                  5       nUR                  U l        UR                  U l        UR                  U l        UR                  U l	        g rJ   )
r  r  r  r  r  Úget_root_graphrÝ  r™   rá  ré  )r
  r  r  r  Úrootr  s        €rV   r  Ú%SubgraphPythonWrapperCodegen.__init__T  sh   ø€ ð +ÔØ,ÔØ$8Ô!ä‰ÑÔà×"Ñ"Ó$ˆà$(×$=Ñ$=ˆÔ!Ø%)×%?Ñ%?ˆÔ"à!×/Ñ/ˆÔà)-×)GÑ)GˆÕ&rX   c                ó&   • U R                   U l        g rJ   )r  rë  r	  s    rV   rì  Ú1SubgraphPythonWrapperCodegen.set_launcher_fn_namek  s   € ð !%× 2Ñ 2ˆÕrX   c                ó   • g rJ   r�   r	  s    rV   rï  Ú)SubgraphPythonWrapperCodegen.write_headerq  r9  rX   c                ó   • g rJ   r�   r  s     rV   r<  Ú2SubgraphPythonWrapperCodegen.add_benchmark_harnesst  r9  rX   c                ó   • g rJ   r�   r  s     rV   r  Ú6SubgraphPythonWrapperCodegen.benchmark_compiled_modulew  r9  rX   c                ó   • g rJ   r�   r	  s    rV   rl  Ú5SubgraphPythonWrapperCodegen.write_async_compile_waitz  r9  rX   c                ó6   • U R                   R                  5       $ rJ   )r  rž  r	  s    rV   rž  Ú/SubgraphPythonWrapperCodegen.next_kernel_suffix}  s   € à×"Ñ"×5Ñ5Ó7Ð7rX   c                ó   • g rJ   r�   r¬  s     rV   r³  Ú2SubgraphPythonWrapperCodegen.generate_after_suffix�  r[  rX   c                ó\   • U R                   R                  SU R                   S35        SnU$ )Nz
            def z(args):
            r6   )rÙ  ré   rë  rv  s     rV   rx  Ú>SubgraphPythonWrapperCodegen.write_launcher_fn_call_get_indent„  s<   € Ø�‰×ÑðØ×&Ñ&Ð'ð (ðô	
ð
 ˆØÐrX   c                ó   • gr¡   r�   r	  s    rV   r   Ú4SubgraphPythonWrapperCodegen.get_wrapper_call_indent�  s   € ØrX   c                óÖ   • U R                   =n(       a6  UR                  UR                   Vs0 s H  n[        U5      U_M     sn-  nU$ [        R
                  R                  nU$ s  snf rJ   )r  Úinput_nodesr„  rµ   r5   rK   r  )r
  r:  rg  rŸ  s       rV   r]  Ú-SubgraphPythonWrapperCodegen.get_graph_inputs�  sm   € ð ×1Ñ1Ð1ˆ9Õ1Ø×*Ñ*Ø#,×#:Ò#:ó.Ú#:˜a”�A“˜’	Ñ#:ñ.ñ ˆFð
 ˆô —W‘W×)Ñ)ˆFØˆùò.s   ­A&c                ó  • U R                   =n(       aL  [        UR                  R                  5       5      UR                   Vs/ s H  o"R
                  PM     sn-   nU$ [        R                  R                  nU$ s  snf rJ   )	r  r  r  r  r„  r›   r5   rK   rc  )r
  r:  r†  Únamess       rV   r{  Ú2SubgraphPythonWrapperCodegen.get_graph_input_names›  su   € Ø×1Ñ1Ð1ˆ9Õ1Ü˜×.Ñ.×3Ñ3Ó5Ó6Ø6?×6MÒ6Mó:Ú6M l×!Ô!Ñ6Mñ:ñ ˆEð
 ˆô —G‘G×-Ñ-ˆEØˆùò:s   ÁA<c                ó|   • U R                   =n(       a  UR                  nU$ [        R                  R                  nU$ rJ   )r  rƒ  r5   rK   r
  )r
  r:  rÂ  s      rV   rT  Ú.SubgraphPythonWrapperCodegen.get_graph_outputs¤  s;   € Ø×1Ñ1Ð1ˆ9Õ1Ø×,Ñ,ˆGð ˆô —g‘g×+Ñ+ˆGØˆrX   c                óŠ   >• UR                  5       nU R                  =n(       a  X#R                  ;   a  g [        TU ]  U5        g rJ   )rM   r  r  r  rD  )r
  rG  r›   r:  r  s       €rV   rD  Ú/SubgraphPythonWrapperCodegen.codegen_allocation«  s=   ø€ Ø�‰Ó ˆØ×2Ñ2Ð2ˆIÕ2¸×@UÑ@UÓ8Uð ä‰Ñ" 6Õ*rX   c                ó8   • U R                   R                  5         g rJ   )r  rC  r	  s    rV   rC  Ú5SubgraphPythonWrapperCodegen.write_triton_header_onceµ  s   € ð 	×Ñ×4Ñ4Õ6rX   c                ó8   • U R                   R                  5         g rJ   )r  rK  r	  s    rV   rK  Ú=SubgraphPythonWrapperCodegen.write_get_raw_stream_header_once¾  s   € ð 	×Ñ×<Ñ<Õ>rX   c                ó¨   • U n[        U[        5      (       a#  UR                  n[        U[        5      (       a  M#  [        U[        5      (       d   eU$ rJ   )r[   r	  r  rE  )r
  ró  s     rV   rò  Ú+SubgraphPythonWrapperCodegen.get_root_graphÈ  sK   € àDHˆÜ˜Ô;×<Ñ<Ø×&Ñ&ˆDô ˜Ô;×<Ó<ô ˜$Ô 4×5Ñ5Ð5Ð5ØˆrX   c                ó   • g rJ   r�   r	  s    rV   r9  Ú<SubgraphPythonWrapperCodegen.generate_and_run_autotune_blockÑ  s   € àrX   )r™   rÝ  rë  r  r  rá  r  ré  rJ   )r  rµ   r  rE  r  ræ  rU  rø  rë  rê  )r²   zDdict[str, Union[ir.TensorBox, ir.TorchBindObject, sympy.Expr, None]]rè  ré  rí  )r²   rE  )rò   rà   r  r  rQ  r  rì  rï  r<  r  rl  rž  r³  rx  r   r]  r{  rT  rD  r)   rC  rK  rò  r9  r  r7  r8  s   @rV   r	  r	  L  sÊ   ø† ñð FJð	HàðHð -ðHð C÷	Hð Hô.3ôòòòô8ôôôð	à	Mô	ôô÷+ð ó7ó ð7ð ó?ó ð?ð óó ð÷ð rX   r	  )rS   r!  r²   r3  )rS   r!  r²   r6  )rd   r!  re   r!  )NN)r›   rµ   rÀ   zlist[triton.Config]rÁ   zlist[TritonGrid]rŠ   rå  rÂ   r¶   r²   ztuple[str, str]rø  )«Ú
__future__r   r  r¸   rñ  rá   r9  rk  rQ  r  r'  rH  r  rS  Úcollections.abcr   Ú	itertoolsr   r   Útypingr   r   r	   r
   rz   r   r  Ú
torch._opsÚtorch.utils._pytreeÚutilsÚ_pytreera  r   rº  Útorch._dynamo.utilsr   r   r   Ú#torch._inductor.codegen.debug_utilsr   Ú$torch._inductor.codegen.multi_kernelr   Ú%torch._inductor.runtime.runtime_utilsr   Útorch._library.opaque_objectr   r   Útorch._loggingr   Ú%torch.fx.experimental.symbolic_shapesr   r   r   r   r   Útorch.fx.noder   Útorch.utils._ordered_setr   Ú torch.utils._sympy.singleton_intr   Útorch.utils._sympy.symbolr   r   rÎ  r!   r"   r#   Ú	codecacher$   r%   r&   rì   r'   Úruntime.hintsr(   r)   r*   r+   r,   r-   r.   r/   r0   r1   r2   r3   r4   Úvirtualizedr5   rZ  r7   r8   r9   r:   r;   r<   Ú	cpp_utilsr=   Úcustom_extern_kernel_codegenr>   Útriton_utilsr?   r@   rA   rB   rC   r�   rK   rD   rE   r   rF   Úwrapper_fxirrG   Ú	getLoggerrò   ÚlogÚdoprintrœ  r�   rµ  rµ   rð   r3  r6  rj   r!  rA  rW   r`   rh   rw   r†  rï   rç  r´   rÌ   r  rú  r  r  rH   rD  rY  rc  rj  rr  r~  r“  rš  r¦  r¼  rÇ  rØ  rå  rú  r"  r4  r`  r7  r5  r{  r•  r©  r·  rÀ  rõ  ÚLinerE  r	  r�   rX   rV   Ú<module>r7     sc  ðå "ã Û Û Û 
Û Û Û Û Û 	Û Û 	Û Ý $ß "ß 6Ó 6ã Ý ã Û ß $Ð $Ý &ß EÑ EÝ CÝ AÝ ;ß RÝ +÷õ õ .Ý /Ý 9ß :ç (Ñ (Ý 'ß (Ý 'Ý ,÷÷ ÷ ó õ ÷÷ õ Ý Fß PÑ Pö ß2ãå%Ý!Ý-Ý)ð ×Ò˜Ó!€á‹×Ñ€ð �—‘˜uŸ{™{¨C°Ð5Ñ6€Ø˜5Ÿ<™<¨¯©°cÐ;NÐPSÐSÑTÐ Ø�2—9‘9˜lÐ*Ñ+€
Ø˜]˜O¨TÐ1Ñ2Ð ôôôò>QðF ˜˜S˜‘>Ð ØØ	ˆ%��U—Z‘Z�Ñ
  #Ð
%Ñ&¨Ð2BÐ1CÀUÈ3ÐPSÈ8Á_Ð1TÑ(UÐUñ€
ð /3Ø*.ðk&Ø
ðk&à ðk&ð ðk&ð ,ð	k&ð
 (ðk&ð õk&ô\U&ðp ×Ò÷ð ó ð÷%6ñ %6÷PYñ Yð
 ×Òô2˜ó 2ó ð2ð ×Òô	/�kó 	/ó ð	/ð ×Òô+�+ó +ó ð+ð ×Òô	2˜ó 	2ó ð	2ð ×Òô1�{ó 1ó ð1ð ×Òô"@ Kó "@ó ð"@ôJ? ;ô ?ð ×Òô
7˜Kó 
7ó ð
7ð ×Òô5˜+ó 5ó ð5ð< ×Òô	(ˆ{ó 	(ó ð	(ð ×Òô/�[ó /ó ð/ðB ×Òô5˜;ó 5ó ð5ð* ×Òô;˜ó ;ó ð;÷2%
ñ %
ðP ×Òôg,Ð%ó g,ó ðg,ðT ×Òô+6Ð,ó +6ó ð+6ð\ ×Òô/Ð(ó /ó ð/ð& ×Òô)Ð"ó )ó ð)ô0(Ð!ô (ð
 ×Òô#0�kó #0ó ð#0ðL ×Òô6˜;ó 6ó ð6ð, ×Òô4˜+ó 4ó ð4ð8 ×Òô	5˜+ó 	5ó ð	5ð ×Òô8˜[ó 8ó ð8ð €
ØÐ Ð,Ñ-€ôT,˜7ô T,ônXGÐ#7õ GrX   