
    ^j                      d dl mZ d dlZd dlZd dlZd dlZd dlZd dlZd dlZd dl	Z	d dl
Z
d dlZd dlZd dlZd dlmZ d dlmZmZ d dlmZmZmZ d dlZd dlmZ d dlZd dlZd dlmc mZ d dlmZ  d dl!m"Z"m#Z#m$Z$ d d	l%m&Z& d d
l'm(Z( d dl)m*Z* d dl+m,Z,m-Z- d dl.m/Z/ d dl0m1Z1m2Z2m3Z3m4Z4m5Z5 d dl6m7Z7 d dl8m9Z9 d dl:m;Z; d dl<m=Z=m>Z> d dl?m@Z@ d dlAmBZBmCZC ddlDmEZEmFZFmGZG ddlHmIZI ddlGmJZJmKZK ddlLmMZM ddlNmOZO ddlPmQZQmRZRmSZS ddlTmUZUmVZV ddlmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZd ddlemfZf dd lgmhZhmiZimjZjmkZkmlZlmmZm dd!lnmoZo dd"lpmqZq dd#lrmsZsmtZtmuZu er&d d$lmvZvmwZwmxZx d dlyZydd%lzm{Z{ dd&lGm|Z| dd'l}m~Z~ dd(lmZ  ej                  e      Z ek       j                  Zeej                  ej>                  eeef   Zeej                  ej>                  ed)ef   ZeGj                  elz  Zed*gdf   Zej                    G d+ d,             Zd}d-Zd~d.Zdd/Zd0 Zeeef   Zeeej0                  z  d1f   eegeed1f   f   z  Z	 	 d	 	 	 	 	 	 	 	 	 	 	 dd2Z	 d	 	 	 dd3Zej                    G d4 d5             Z G d6 d7      Z G d8 d9e      Z G d: d*      Zej                    G d; d<e             Zej                    G d= d>e             Zej                    G d? d@e             Zej                    G dA dBe             Zej                    G dC dDe             Zej                    G dE dFe             Z G dG dHe      Zej                    G dI dJe             Zej                    G dK dLe             Zej                    G dM dNe             Zej                    G dO dPe             Zej                    G dQ dRe             Zej                    G dS dTe             Zej                    G dU dVe             Zej                    G dW dXe             Zej                    G dY dZe             Zej                    G d[ d\e             Zej                    G d] d^e             Z G d_ d`      Zej                    G da dbe             Zej                    G dc dde             Zej                    G de dfe             Zej                    G dg dhe             Z G di dje      Zej                    G dk dle             Zej                    G dm dne             Zej                    G do dpe             Zej                    G dq dre             Zej                    G ds dte             Zej                    G du dve             Zej                    G dw dxe             ZeZee_z  Z G dy dzei      Z G d{ d|e      Zy)    )annotationsN)Callable)chaincount)AnyProtocolTYPE_CHECKING)Exprdtype)countersdynamo_timedget_debug_dir)DebugPrinterManager)MultiKernelState)	cache_dir)get_opaque_obj_repris_opaque_value_type)trace_structured)CallMethodKeyConvertIntKeyDivideByKeyresolve_unbacked_bindingsSymTypes)_get_qualified_name)StorageWeakRef)
OrderedSet)MaxMin)SingletonIntsymbol_is_typeSymT   )async_compileconfigir)output_code_log)IRNodeReinterpretView)triton_heuristics)DeviceProperties)DEFAULT_STREAMDEFAULT_STREAM_IDXSTREAM_NAME_TEMPLATE)get_raw_stream_nameget_stream_name)cache_on_selfDeferredLineBaseDelayReplaceLineget_benchmark_nameget_dtype_sizeIndentedBuffer#is_codegen_graph_partition_subgraphis_using_cudagraph_partitionLineContextmake_codegen_buffersympy_product	sympy_str
sympy_substriton_version_uses_attrs_dict)V   )ArgNameCodeGenDeferredLinePythonPrinterWorkspaceArgWorkspaceZeroMode)cexpr)CUSTOM_EXTERN_KERNEL_CODEGEN)	config_ofshould_unwrap_unspec_argsignature_to_meta)IterableIteratorSequence)GraphLowering)ExternKernel)BaseSchedulerNode)FxConverterzir.CommBufferTypeWrapperLinec                  @    e Zd ZU ded<   ded<   ded<   ded<   ded	<   y
)BenchmarkStorageGroupstrbuffer_nameintnbytesr   devicer   z&dict[str, tuple[list[int], list[int]]]inputsN)__name__
__module____qualname____annotations__     j/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/codegen/wrapper.pyrV   rV   j   s    KKJ22rb   rV   c                   t         j                  j                  |       }| j                         t         j                  j                  v}t         j                  j
                  j                  | j                               }| j                         | j                         t        t         j                  j                  j                  |            ||fS N)r@   graphget_allocation_storage_sizeget_nameunaligned_buffers	schedulerget_buf_streamget_device_or_error	get_dtyper=   sizevarssimplify)nodestorage_size	alignmentstreams       rc   buffer_reuse_keyrt   s   s    7766t<Lqww'@'@@IWW--dmmo>F  " 	!''""++L9:	 	rb   c                f   t         j                  j                  |       }| j                         }t	        |t
        j                        sJ | j                         | j                         t        t         j                  j                  j                  |            |j                  |j                  fS re   )r@   rf   rg   get_output_spec
isinstancer'   CommBufferLayoutrl   rm   r=   rn   ro   comm_buffer_type
group_name)rp   rq   layouts      rc   comm_buffer_reuse_keyr|      s    7766t<L!!#Ffb11222  "!''""++L9: rb   c                   | j                         |j                         k7  ry| j                         |j                         k7  ryt        j                  j                  j                  t        j                  j                  |             }t        j                  j                  j                  t        j                  j                  |            }t        |      t        |      k(  sWt        j                  j                  j                  |d|z        r+t        j                  j                  j                  ||      ryy)NFgffffff?T)
rl   rm   r@   rf   rn   ro   rg   r=   statically_known_geqstatically_known_leq)	input_buf
output_buf
input_sizeoutput_sizes       rc   can_match_buffer_sizer      s     $$&**H*H*JJ
 4 4 66!!**	++I6J ''""++	++J7K 	*;!77 	
--k4*;LMGG11+zJrb   c                   t        | t        j                        r?| j                         }|j                  |j
                  |j                  |j                  dfS g }| }t        |t        j                  t        j                  t        j                  f      rj|j                         }|y|j                  |       |j                  }t        |t        j                  t        j                  t        j                  f      rjt        |t        j                        sy|D ]+  }|j                  |j                         j                  k7  s+ y |j                         }|j                  |j
                  |j                  |j                  dfS )a  
    Collapse a chain of ReinterpretView <- StorageBox
    <- ReinterpretView <- StorageBox.... <- buffer wrappers if every layer
    has the same offset as the innermost (base) buffer.

    Returns:
        (size, stride, offset, dtype, collapsible: bool)
    T)NNNNF)rw   r'   Buffer
get_layoutsizestrideoffsetr   	TensorBox
StorageBoxr*   appenddata)r   laylayoutscurbase_lays        rc   codegen_reinterpret_view_helperr      s!    $		"ooxxSZZD@@G
C
S2<<8J8JK
Lnn;0shh S2<<8J8JK
L c299%,  1::)00001 ~~H==(//8??HNNDPPrb   .c                    t               dd	 d	 	 	 dfd}dd fd}d  } |d| d       r4t        j                  j                  rj                  j                         nt        j                         }j                         5  |5  t        j                  j                  rV|rTt        j                  j                  r:|t        j                  j                  v rt        j                  j                  |   }	nd gt        |      z  }	t        |      dk(  r" ||d   |	d         \  }
} |d	|
 d	|        nt        |      dkD  sJ t        |      t        |      k(  sJ t               }t        t        |||	      d
 d      D ]  \  }
}}g }|j                  r:|j                  D ]+  }|dvs|j!                  d| d|j                  |           - |rdj#                  |      }nd} ||
|      \  }
}d| d|
 }||v r~|j%                  |        ||d| d|         d d d        d d d        |j'                         fS # 1 sw Y   #xY w# 1 sw Y   'xY w)Nc                d    t        | t        j                        r| S t        j                  |       S re   )rw   sympyr
   Integer)items    rc   _convert_to_sympy_exprz@user_defined_kernel_grid_fn_code.<locals>._convert_to_sympy_expr   s#    !$

3tLt9LLrb   c                    t        |       r| | fS t        fd| D              }|s|}j                  |      t        j                  j
                  r$j                  t        fd|D                    fS dfS )a'  
        This function return a tuple of two values: the first one is for the real grid
        which is used in the generated code; the second one is an example grid with
        concreate values which is used in the autotune block to run the generated
        kernels at compile time.
        Nc              3  .   K   | ]  } |        y wre   ra   ).0gr   s     rc   	<genexpr>zKuser_defined_kernel_grid_fn_code.<locals>.determine_grid.<locals>.<genexpr>   s     C1!4Cs   c              3  T   K   | ]  }j                  |t        |             ! y wre   generate_example_arg_valuetype)r   r   wrappers     rc   r   zKuser_defined_kernel_grid_fn_code.<locals>.determine_grid.<locals>.<genexpr>   s*        ::1d1gF   %()callabletuplecodegen_python_shape_tupler&   tritonautotune_at_compile_time)gridexample_grid
sympy_gridr   r   s      rc   determine_gridz8user_defined_kernel_grid_fn_code.<locals>.determine_grid   s     ?htn:CdCC
%L..z: ==99 22 !- 
 	
 
 	
rb   c                    j                  |        rJt        j                  j                  r/j                  vr j
                  j                  |xs |        y y y y re   )	writeliner&   r   r   kernel_autotune_nameskernel_autotune_calls)liner   nameoutputr   s     rc   r   z3user_defined_kernel_grid_fn_code.<locals>.writeline  sW    66G999))33L4HDI : 7 rb   grid_wrapper_for_def z(meta):rA   r   zreturn c                2    t        | d   j                        S NrA   lenkwargsxs    rc   <lambda>z2user_defined_kernel_grid_fn_code.<locals>.<lambda>,  s    c!A$++. rb   Tkeyreverse)matrix_instr_nonkdimwaves_per_eukpackzmeta['z'] == z and Trueif z	: return )r   int | sympy.Exprreturn
sympy.Exprre   )r   
TritonGridr   zTritonGrid | None)r   rW   r   
str | None)r7   r&   r   r   r   indent
contextlibnullcontextr@   rf   autotuning_gridsr   r   sortedzipr   r   joinaddgetvalue)r   configsgridsr   original_fxnode_namer   r   fn_namekernel_autotune_calls_indentexample_gridsr   r   seenc
guardslistkwargguards	statementr   r   s   `  `              @@rc    user_defined_kernel_grid_fn_coder      s    FM
 +/

'
>J J "$(GWIW%& v}}== 	%%,,.##% !
 
 .L6 .LMM22$(($(@(@@GG445IJM!FSZ/Mu:?!/a-:J!KD,v&',(@Au:>!>u:W---$.LD *0E7M2.* L%a
  
88!" W  ) 
 '--ugVAHHUOCT.UVW $\\*5F#F%3D,%G"l!&4&9	$#)s6()L>%JK1L-.L .L` FOO%%%a.L .L .L .Ls,   I!D%I<A7I3I!I	I!!I*c                  	 t               | j                  }|r|d   }j                  |d       ddlddlm ddlm ddlm	 t        | j                  g      	fd	 |        j                         S )
z
    Given a triton kernel function pointer collect the transitive closure of
    its dependencies

    epilogue_fusion: Optional[(fused epilogue node, modified kerel src code)]
    rA   Tstripr   N)JITFunction)	constexprr   c           	        t        d t        j                  | j                        D              }| j                  j                  j                  di       }| j                  j                  j                  D ]  }|v r	|| j                  j                  v s"| j                  j                  |   }t        |	      rX
j                          
j                  d       
j                  |j                  d       j                  |        |       t        d      rWt        |j                  j                   j"                        r,d|j                  v rdvr~|j                  j                  j                  d      }|rWt        |d      rK|j$                  j'                  d	      r0
j                  d
|j$                   d       j                  d       
j                          
j                  d       
j                  |j                  d       ||j                  j(                  k7  r*
j                  | d|j                  j(                          j                  |        |       t        |t*        t,        t.        f      r
j                          t        |      rd|j0                  d}n|}|j                  |      x}rKt        |t2              rd|j$                   d|j(                   }nd|}
j                  | | d|        n
j                  | d|        j                  |       ||v s|dk7  st        |d      s|j$                  j'                  d	      st        |      r!
j                  | d|j4                          n:t        |d      r.
j                  d
|j$                   d|j(                   d|        j                  |        y )Nc              3  R   K   | ]  }|j                   d k(  r|j                   ! yw)LOAD_GLOBALN)opnameargval)r   insts     rc   r   z^user_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse.<locals>.<genexpr>d  s(      '
{{m+ KK'
s   %'r`   z@triton.jitTr   constexpr_functionr   r^   r   zfrom z import dtype as dtypez@triton.constexpr_function = ztl.constexpr(): .tlz = tl.r]   z import z as )r   disBytecodefn__globals__get__code__co_namesrw   newliner   splicesrcr   hasattrruntimejitConstexprFunctionr^   
startswithr]   rY   rW   boolvaluer   r   )
cur_kernelunqualified_loadsglobal_annotationssymbol_namesymboldtype_symbol
symbol_str
annotationannotation_coder   compile_wrapperr   symbols_includedtraverser   triton_dtypes            rc   r  zKuser_defined_triton_kernel_transitive_closure_source_code.<locals>.traverse_  s   
 ' '
Z]]3'
 

 (]]66::;LbQ%==11:: K	6K..jmm777#22;?fk2#++-#--m<#**6::T*B$((5V$V%9:zNN&&88@
 &**,@P1P'-yy'<'<'@'@'I( 'l C , 7 7 B B8 L+55"'(?(?'@@V W -009#++-#--.JK#**6::T*B"fii&8&88'11*m3vyy/A/A.BC %((5V$c4(CD#++-!&)4'4V\\4DA%F
(.z
%7%;%;K%HHzH%j$7"$Z%:%:$;1Z=P=P<Q R , 13:..AO'11*mO+<C
|L (11[MZL2QR$((5#44#t+5 ))44X> "&,7'11[M}2UV 4'11#F$5$5#6hv>OtT_S`a %((5WK	6rb   )r7   r   r   r   r   triton.languager   triton.language.corer   r   r]   r   )
kernelepilogue_fusion
kernel_srcr   r  r   r  r  r   r  s
      @@@@@@@rc   9user_defined_triton_kernel_transitive_closure_source_coder  G  s~     %&OJ$Q'
:T2 "): "6??"34V6 V6p V##%%rb   c                  (    e Zd ZU ded<   ded<   d Zy)SymbolicCallArgsympy.Symbolinnerr   
inner_exprc                ,    t        | j                        S re   )rW   r  selfs    rc   __str__zSymbolicCallArg.__str__  s    4::rb   N)r]   r^   r_   r`   r  ra   rb   rc   r  r    s    rb   r  c                  Z     e Zd Z fdZddZd	dZd
dZddZddZ	 	 	 	 	 	 ddZ	 xZ
S )MemoryPlanningStatec                    t         |           t        j                  t              | _        t        j                  t              | _        d| _        y Nr   )super__init__collectionsdefaultdictlist
reuse_poolcomm_buffer_reuse_pooltotal_allocated_buffer_size)r  	__class__s    rc   r%  zMemoryPlanningState.__init__  sG     ##D) 	 ##D) 	# 12(rb   c                L    t        | j                  j                  |d             S re   )r  r)  r   r  r   s     rc   __contains__z MemoryPlanningState.__contains__  s    DOO''T233rb   c                \    | j                   |   j                         }|j                  rJ |S re   )r)  pop	is_reusedr  r   r   s      rc   r1  zMemoryPlanningState.pop  s+    s#'')>>!!rb   c                \    |j                   rJ | j                  |   j                  |       y re   )r2  r)  r   r3  s      rc   pushzMemoryPlanningState.push  s&    >>!!##D)rb   c                L    t        | j                  j                  |d             S re   )r  r*  r   r.  s     rc   comm_buffer_containsz(MemoryPlanningState.comm_buffer_contains  s     D//33C>??rb   c                \    | j                   |   j                         }|j                  rJ |S re   )r*  r1  r2  r3  s      rc   comm_buffer_popz#MemoryPlanningState.comm_buffer_pop  s-    **3/335>>!!rb   c                \    |j                   rJ | j                  |   j                  |       y re   )r2  r*  r   r3  s      rc   comm_buffer_pushz$MemoryPlanningState.comm_buffer_push  s*     >>!!##C(//5rb   )r   ReuseKeyr   r  )r   r<  r   FreeIfNotReusedLine)r   r<  r   r=  r   None)r   CommBufferReuseKeyr   r  )r   r?  r   r=  )r   r?  r   r=  r   r>  )r]   r^   r_   r%  r/  r1  r5  r7  r9  r;  __classcell__r,  s   @rc   r!  r!    s?    
24
*@
6%6-@6	6rb   r!  c                      e Zd ZddZy)HasWriteLinec                     y re   ra   r  r   s     rc   r   zHasWriteLine.writeline      rb   N)r   z$LineContext | DeferredLineBase | strr   r>  )r]   r^   r_   r   ra   rb   rc   rC  rC    s    Prb   rC  c                      e Zd ZddZy)rT   c                0    t        dt        |              )Nz&FX codegen not yet supported for type )NotImplementedErrorr   r  	converters     rc   
codegen_fxzWrapperLine.codegen_fx  s    !$J4PT:,"WXXrb   NrK  rS   r   FxConversionFuncr]   r^   r_   rL  ra   rb   rc   rT   rT     s    Yrb   c                  :    e Zd ZU ded<   ded<   d	dZd
dZddZy)EnterSubgraphLinePythonWrapperCodegenr   rP   rf   c                b    | j                   j                  | j                   j                         y re   )r   push_computed_sizescomputed_sizesr  s    rc   __post_init__zEnterSubgraphLine.__post_init__  s    (()D)DErb   c                n    | j                   j                  | j                         |j                          y re   )r   push_codegened_graphrf   	do_indentr  codes     rc   codegenzEnterSubgraphLine.codegen  s"    ))$**5rb   c                    |j                   S re   )_generate_enter_subgraphrJ  s     rc   rL  zEnterSubgraphLine.codegen_fx  s    111rb   Nr   r>  r[  r7   r   r>  rM  r]   r^   r_   r`   rV  r\  rL  ra   rb   rc   rQ  rQ    s    !!F2rb   rQ  c                  <    e Zd ZU ded<   ded<   ddZed	d       Zy)
ConditionalLinerR  r   zir.Conditionalrp   c                    t        d      )NzOnly supports FX codegenrI  rZ  s     rc   r\  zConditionalLine.codegen  s    !"<==rb   c                    | j                   S re   )_generate_conditionalrK  s    rc   rL  zConditionalLine.codegen_fx  s    ...rb   Nr`  rM  r]   r^   r_   r`   r\  staticmethodrL  ra   rb   rc   rc  rc    s'    !!
> / /rb   rc  c                  2    e Zd ZU ded<   ddZedd       Zy)CommentLiner:   r   c                :    |j                  | j                         y re   )r   r   rZ  s     rc   r\  zCommentLine.codegen  s    tyy!rb   c                    | j                   S re   )_generate_commentrh  s    rc   rL  zCommentLine.codegen_fx  s    ***rb   Nr`  rM  ri  ra   rb   rc   rl  rl    s!    
" + +rb   rl  c                  <    e Zd ZU ded<   ded<   ddZed	d       Zy)
DynamicScalarLinerR  r   zir.DynamicScalarrp   c                N    | j                   j                  | j                         y re   )r   _codegen_dynamic_scalarrp   rZ  s     rc   r\  zDynamicScalarLine.codegen$  s    ,,TYY7rb   c                    | j                   S re   )_generate_dynamic_scalarrh  s    rc   rL  zDynamicScalarLine.codegen_fx'  s    111rb   Nr`  rM  ri  ra   rb   rc   rq  rq    s'    !!
8 2 2rb   rq  c                  0    e Zd ZU ded<   ddZddZd	dZy)
ExitSubgraphLinerR  r   c                V    | j                   j                         | j                   _        y re   )r   pop_computed_sizesrU  r  s    rc   rV  zExitSubgraphLine.__post_init__0  s    &*ll&E&E&G#rb   c                X    | j                   j                          |j                          y re   )r   pop_codegened_graphdo_unindentrZ  s     rc   r\  zExitSubgraphLine.codegen3  s    ((*rb   c                    |j                   S re   )_generate_exit_subgraphrJ  s     rc   rL  zExitSubgraphLine.codegen_fx7  s    000rb   Nr_  r`  rM  ra  ra   rb   rc   rw  rw  ,  s    !!H1rb   rw  c                  2    e Zd ZU ded<   ded<   ddZd	dZy)
EnterDeviceContextManagerLinerY   
device_idx
int | Nonelast_seen_device_guard_indexc                x   t         j                  j                  r|j                  d       t         j                  j                  rg| j
                  ;|j                  t         j                  j                  j                          d       y | j
                  | j                  k(  s{J d       | j
                  H|j                  t         j                  j                  j                          d| j                   d       y |j                  d| j                   d       y y |j                  dt         j                  j                  j                  | j                         d       |j                          |j                  t         j                  j                  j                  | j                               y )	N
z) stream_guard(stream, this->device_idx_);z4AOTInductor only supports running on one CUDA devicez device_guard(z);zdevice_guard.set_index(with :)r@   rf   cpp_wrapperr   aot_moder  
device_opscpp_aoti_stream_guardr  cpp_aoti_device_guarddevice_guardrY  
set_devicerZ  s     rc   r\  z%EnterDeviceContextManagerLine.codegen@  sP   77NN4 ww 44<NN77--CCEFFop  <<O NO 44<NN77--CCEFnUYUdUdTeegh NN%<T__<MR#PQ P NNU177#5#5#B#B4??#S"TTUVWNNNN177--88IJrb   c                    |j                   S re   )&_generate_enter_device_context_managerrJ  s     rc   rL  z(EnterDeviceContextManagerLine.codegen_fx]  s    ???rb   Nr`  rM  r]   r^   r_   r`   r\  rL  ra   rb   rc   r  r  ;  s    O",,K:@rb   r  c                      e Zd ZddZddZy)ExitDeviceContextManagerLinec                Z    t         j                  j                  s|j                          y y re   r@   rf   r  r|  rZ  s     rc   r\  z$ExitDeviceContextManagerLine.codegenb  s     ww"" #rb   c                    |j                   S re   )%_generate_exit_device_context_managerrJ  s     rc   rL  z'ExitDeviceContextManagerLine.codegen_fxf  s    >>>rb   Nr`  rM  )r]   r^   r_   r\  rL  ra   rb   rc   r  r  a  s    ?rb   r  c                  2    e Zd ZU ded<   ded<   ddZd	dZy)
ExternKernelAllocLinerR  r   ir.ExternKernelAllocrp   c                    | j                   }g |j                         |j                         }| j                  j	                  | j                   |       y re   )rp   codegen_argscodegen_kwargsr   $_generate_extern_kernel_alloc_helper)r  r[  rp   argss       rc   r\  zExternKernelAllocLine.codegeno  sD    yy=""$=t':':'<=99$))TJrb   c                    |j                   S re   )_generate_extern_kernel_allocrJ  s     rc   rL  z ExternKernelAllocLine.codegen_fxt  s    666rb   Nr`  rM  r  ra   rb   rc   r  r  j  s    !!
K
7rb   r  c                  2    e Zd ZU ded<   ded<   ddZd	dZy)
ExternKernelOutLinerR  r   ir.ExternKernelOutrp   c           	     D   | j                   }g |j                         |j                  d      }|j                         }t        j
                  j                  r|j                  dk(  rd}n|j                         }|j                         x}r|j                  nt        j
                  j                  }| j                  j                  ||j                         |j                  r|j                  j                         nd ||| j                   j                                y )NT)skip_outztorch::inductor::_mm_plus_mmaoti_torch__mm_plus_mm_out)rp   r  r  get_kernel_namer@   rf   r  cpp_kernel_name
get_devicer   device_typer   "_generate_extern_kernel_out_helpercodegen_referenceoutput_viewget_stack_traces)r  r[  rp   r  kernel_namedr[   s          rc   r\  zExternKernelOutLine.codegen}  s    yyJ""$Jt':':D':'IJ**,GG$$(FF 7K..0K!%!22A29L9L77""$484D4DD..0$II&&(	
rb   c                    |j                   S re   )_generate_extern_kernel_outrJ  s     rc   rL  zExternKernelOutLine.codegen_fx      444rb   Nr`  rM  r  ra   rb   rc   r  r  x  s    !!

,5rb   r  c                  .    e Zd ZU dZded<   ded<   ddZy)	ExternKernelMultiOutLinezCodegen line for multi-output .out() variant calls.

    Generates a kernel call with pre-allocated output buffers passed as
    keyword arguments. E.g. kernel(x, out0=buf0, out1=buf1).
    rR  r   ir.ExternKernelMultiOutrp   c           	        | j                   }|j                         }g |j                         |j                         }t	        |j
                  |j                        D ])  \  }}|j                  | d|j                                 + |j                  |j                          d| ddj                  |       d       y )N=r   (, r   )rp   r  r  r  r   out_arg_namesout_variant_output_nodesr   rh   r   r   )r  r[  rp   r  r  out_nameout_nodes          rc   r\  z ExternKernelMultiOutLine.codegen  s    yy**,=""$=t':':'<="% = =#
 	=Hh KK8*Ah&7&7&9%:;<	=
 	$--/*#k]!DIIdO;LANOrb   Nr`  r]   r^   r_   __doc__r`   r\  ra   rb   rc   r  r    s     "!
!!
Prb   r  c                  2    e Zd ZU ded<   ded<   ddZd	dZy)
FreeLinerR  r   BufferLike | ir.TorchBindObjectrp   c                    | j                   j                         t        j                  j                  vsJ |j                  | j                  j                  | j                                y re   )rp   rh   r@   rf   removed_buffersr   r   make_buffer_freerZ  s     rc   r\  zFreeLine.codegen  sF    yy!!#177+B+BBBBt||44TYY?@rb   c                    |j                   S re   )_generate_freerJ  s     rc   rL  zFreeLine.codegen_fx      '''rb   Nr`  rM  r  ra   rb   rc   r  r    s    !!
))A(rb   r  c                      e Zd ZU ded<   ded<   ded<   ded<   ded<   d	ed
<   ded<   ded<   ded<   ded<   ded<   ded<   dZded<   ddZddZy)KernelCallLinerR  r   rW   r  ztuple[Any, ...]	call_argsraw_keysraw_args	list[str]	arg_typesr  r   zdict[str, Any]triton_metazdict[str, Any] | Noneinductor_metaztorch.devicer[   
graph_namer   Nr  current_stream_idxc                B   | j                   j                  | j                  | j                  | j                  | j
                  | j                  | j                  | j                  | j                  | j                  | j                  | j                  | j                         y )N)
r   r  r  r  r  r  r[   r  r   r  )r   _generate_kernel_call_helperr  r  r   r  r  r  r  r  r[   r  r   r  rZ  s     rc   r\  zKernelCallLine.codegen  sw    11NN;;nn]]]]((,,;;!%!:!:#66 	2 	
rb   c                    |j                   S re   )_generate_kernel_callrJ  s     rc   rL  zKernelCallLine.codegen_fx      ...rb   r`  rM  )r]   r^   r_   r`   r  r\  rL  ra   rb   rc   r  r    s]    !!L((O%)
)
 /rb   r  c                  f    e Zd ZU ded<   ded<   ded<   dZded<   d	Zd
ed<   dZded<   ddZddZy)KernelDefinitionLinerR  r   rW   r  kernel_bodyNr   metadataTr  gpucpp_definitionc                    | j                   j                  | j                  | j                  | j                  | j
                  | j                         y N)r  r  r  )r   _define_kernel_helperr  r  r  r  r  rZ  s     rc   r\  zKernelDefinitionLine.codegen  sB    **]].. 	+ 	
rb   c                    |j                   S re   )_generate_kernel_definitionrJ  s     rc   rL  zKernelDefinitionLine.codegen_fx  r  rb   r`  rM  )	r]   r^   r_   r`   r  r  r  r\  rL  ra   rb   rc   r  r    s<    !!HjC!%NJ%
5rb   r  c                  0    e Zd ZU ded<   ddZddZd	dZy)
MemoryPlanningLinerR  r   c                    | S )zFirst pass to find reusera   r  states     rc   planzMemoryPlanningLine.plan  s    rb   c                     y)zSecond pass to output codeNra   rZ  s     rc   r\  zMemoryPlanningLine.codegen  rF  rb   c                r   g }t        j                  |       D ]t  }|j                  dk(  rt        | |j                        }|j	                  |j                   d|j
                  t        j                  u r|j                         n|        v t        |       j                   ddj                  |       dS )zF
        Emits a string representation that fits on one line.
        r   r  r  r  r   )dataclassesfieldsr   getattrr   r   r'   r   rh   r]   r   )r  r  fieldvals       rc   r  zMemoryPlanningLine.__str__  s      ''- 	EzzY&$

+CKK::,a%**		2IsST		 t*%%&a		$'8::rb   Nr  r!  r   r  r`  r   rW   )r]   r^   r_   r`   r  r\  r  ra   rb   rc   r  r    s    !!);rb   r  c                  d     e Zd ZU dZdZded<    ej                  e      Z	ded<   d	 fdZ
 xZS )
+EnterDeviceContextManagerWithStreamInfoLineu;  Enter a device context and retrieve user stream objects.

    Attributes:
        num_streams: Number of streams (determined by user annotations on nodes).
        stream_idx_to_user_obj_idx: Maps stream_idx → user_object_index for
            retrieving user stream objects via get_external_object_by_index.
    rA   rY   num_streams)default_factoryzdict[int, int]stream_idx_to_user_obj_idxc                   t         j                  j                  rt        |   |       yt        |   |       |j                  t         dt         j                  j                  j                                 | j                  dkD  rVt        d| j                        D ]<  }| j                  |   }|j                  t        j                  |       d| d       > yy)z5Generate context switching and stream retrieval code.r   rA   
stream_idxz  = get_external_object_by_index(r   N)r@   rf   r  r$  r\  r   r-   r  current_streamr  ranger  r/   format)r  r[  iuser_obj_idxr,  s       rc   r\  z3EnterDeviceContextManagerWithStreamInfoLine.codegen  s    77GOD!GOD!NNn-S1C1C1R1R1T0UVW!#q$"2"23 A#'#B#B1#ELNN/66!DE F::FqJ $rb   r`  )r]   r^   r_   r  r  r`   r  r  dictr  r\  r@  rA  s   @rc   r  r    s7     K1B1B1BSW1XX rb   r  c                  (    e Zd ZU dZdZded<   ddZy)*ExitDeviceContextManagerWithStreamInfoLinez{Exit a device context.

    Attributes:
        num_streams: Number of streams that were allocated (must match Enter).
    rA   rY   r  c                Z    t         j                  j                  s|j                          yy)zGenerate context exit code.Nr  rZ  s     rc   r\  z2ExitDeviceContextManagerWithStreamInfoLine.codegen8  s     ww"" #rb   Nr`  )r]   r^   r_   r  r  r`   r\  ra   rb   rc   r  r  .  s     Krb   r  c                  $    e Zd ZU dZded<   ddZy)EnterCudaStreamContextLinezEnter a context executed by the respective device stream.

    Attributes:
        stream_idx: The index number for the entering device stream context.
    rY   r  c                t    |j                  dt        | j                         d       |j                          y )Nr  r  )r   r1   r  rY  rZ  s     rc   r\  z"EnterCudaStreamContextLine.codegenH  s+    t?@BCrb   Nr`  r  ra   rb   rc   r  r  >  s     Orb   r  c                      e Zd ZdZddZy)ExitCudaStreamContextLinez1Generate code to exit the current stream context.c                $    |j                          y re   )r|  rZ  s     rc   r\  z!ExitCudaStreamContextLine.codegenQ  s    rb   Nr`  )r]   r^   r_   r  r\  ra   rb   rc   r
  r
  M  s
    ;rb   r
  c                  *    e Zd Zd ZddZddZddZy)EfficientPeakEstimatec                   ddl m}m} t        j                  j
                  j                  }t        t        j                  j                  j                               }t        t        j                  j                               } |||      } ||||      \  | _        }ddlm}  ||t        j                  t         d      | _        y )Nr$   )estimate_peak_memoryget_freeable_input_bufrA   )SegmentedTreer   )memoryr  r  r@   rf   rj   nodesr   graph_inputskeysget_output_namesoverall_peak_memorysegmented_treer  operatorr   max)	r  r  r  scheduler_nodesr  graph_outputsnames_to_freeable_bufspeak_by_scheduler_noder  s	            rc   r%  zEfficientPeakEstimate.__init__V  s    I''++11!!''"6"6";";"=>"177#;#;#=>!7!V;O"<
8 "8 	2+"HLL#q
rb   c                    t         j                  j                  j                  t         j                  j	                  |      d      t        |j                               z  S )Nr   fallback)r@   rf   rn   optimization_hintrg   r6   rm   r  rp   s     rc   	_get_sizezEfficientPeakEstimate._get_sizei  sL    ww11GG//5 2 
4>>+,- 	-rb   c                n    | j                   j                  |j                  dz   |j                  dz
        S r   )r  summarize_rangescheduler_node_indexr  line_aline_bs      rc   peak_betweenz"EfficientPeakEstimate.peak_betweenn  s6    ""22''!+V-H-H1-L
 	
rb   c                    |j                   dz   |j                   k(  ry | j                  j                  |j                   dz   |j                   dz
  | j                  |j                               y r   )r'  r  update_ranger$  rp   r(  s      rc   update_peak_betweenz)EfficientPeakEstimate.update_peak_betweens  s^    &&*f.I.II((''!+''!+NN6;;'	
rb   N)rp   
BufferLiker   rY   )r)  r=  r*  AllocateLine)r]   r^   r_   r%  r$  r+  r.  ra   rb   rc   r  r  U  s    
&-



rb   r  c                  X    e Zd ZU dZded<   dZded<   d ZddZdd	Zdd
Z	ddZ
ddZy)r0  z6Represents a buffer allocation during memory planning.r/  rp   Fr  comm_bufferc                   t         j                  j                  j                  J t         j                  j                  j                  j                  t         j                  j                  j                        | _        y re   r@   rf   rj   current_noder  indexr'  r  s    rc   rV  zAllocateLine.__post_init__  T    ww  --999$%GG$5$5$;$;$A$AGG**%
!rb   c                    | j                   ry|j                  dz   | j                  k(  ry| j                  j                  j                  }| j                  j                  j                  ||       }||z   }||k  S )NTrA   )r2  r'  r   estimate_peakr  r+  )r  	free_liner   r  peak_memory_in_rangenew_peak_memorys         rc   should_reuse_bufferz AllocateLine.should_reuse_buffer  su    ))A-1J1JJ"ll88LL#||99FFyRVW!55"555rb   c           	     (   | j                   j                         t        j                  j                  v rt        | j                        S | j                  r}t        | j                         }t        j                  rV|j                  |      rE|j                  |      }d|_        t        | j                  |j                   | j                   d      S | S t        | j                         }t        j                  r
||v r|j!                  |      }t        j                  j"                  j%                  t        j                  j'                  | j                         d      t)        | j                   j+                               z  }| j-                  ||      rXd|_        | j                  j.                  j1                  ||        t        | j                  |j                   | j                         S |j3                  ||       | S | j                   j5                         j6                  dk(  rh| j                  j9                  | j                         }|A|xj:                  t=        t?        j@                  tB        jD                  |d            z  c_        | S )NTr2  r   r   cpurA   )#rp   rh   r@   rf   r  NullLiner   r2  r|   r&   allow_buffer_reuser7  r9  r2  	ReuseLinert   r1  rn   r"  rg   r6   rm   r=  r9  r.  r5  rl   r   static_shape_for_buffer_or_noner+  rY   	functoolsreducer  mul)r  r  r   r:  r   static_shapes         rc   r  zAllocateLine.plan  s   99177#:#::DLL))'		2C((U-G-G-L!11#6	&*	# LL)..$))  K tyy)$$		#I77##5533DII> 6 tyy22456D ''	48&*	#**>>y$O y~~tyyII

3	*99((*//58<<GG		RL'11S$$X\\<C6 1 rb   c                   | j                   j                         t        j                  j                  vsJ | j
                  r| j                  |       y | j                  j                  | j                         }|j                  |       y re   )
rp   rh   r@   rf   r  r2  _codegen_comm_bufferr   make_buffer_allocationr   )r  r[  r   s      rc   r\  zAllocateLine.codegen  sb    yy!!#177+B+BBBB%%d+<<66tyyADNN4 rb   c                @   | j                   j                         }| j                   j                         }||j                  
J d|        | j                   j	                         }t        | j                   j                               }t        | j                   j                               }| j                   j                         }t        |t        j                        sJ |j                  }|j                  }	|t        j                  j                  k(  rh| d| j                   j#                  |       d| j                   j#                  |       d| d|j                   d|	 dt%        j&                  dd	       d
}
nt)        d|       |j+                  |
       y)z*Generate allocation code for comm buffers.Nz9Comm buffer requires a valid CUDA device with index, got z = empty_strided_p2p(r  z, torch.device("cuda:z"), group_name="z", alloc_id=r   l    r   zUnsupported comm buffer type: )rp   rh   r  r6  rm   r   get_size
get_striderv   rw   r'   rx   ry   rz   CommBufferTypeSYMM_MEMr   codegen_shape_tuplerandomrandintrI  r   )r  r[  r   r[   r   shaper   r{   ry   rz   r   s              rc   rJ  z!AllocateLine._codegen_comm_buffer  s   yy!!#%%'!fll&> 	
GxP	
> 		##%dii((*+tyy++-.**,&""5"5666!22&&
r00999&-<<33E:;2<<33F;<B' &&,ll^ 4)l +"NN1i89<  &01A0BC  	trb   c                J    | j                   r|j                  S |j                  S re   )r2  _generate_comm_buffer_allocate_generate_allocaterJ  s     rc   rL  zAllocateLine.codegen_fx  s#    ;;;+++rb   N)r:  r=  r   rY   r   r  r  r`  rM  )r]   r^   r_   r  r`   r2  rV  r=  r  r\  rJ  rL  ra   rb   rc   r0  r0  }  s4    @
K
6&P!>,rb   r0  c                  R    e Zd ZU ded<   dZded<   dZded<   d ZddZdd	Zdd
Z	y)r=  r/  rp   Fr  r2  r2  c                   t         j                  j                  j                  J t         j                  j                  j                  j                  t         j                  j                  j                        | _        y re   r4  r  s    rc   rV  z!FreeIfNotReusedLine.__post_init__  r7  rb   c                <   t        | j                  j                               dkD  r| S t        | j                  j                  t
        j                        r| S | j                  rJ | j                  j                         t        j                  j                  v rt        | j                        S t        j                  r\| j                   r)t#        | j                        }|j%                  ||        | S t'        | j                        }|j)                  ||        | S r#  )r   rp   get_inputs_that_alias_outputrw   r{   r'   MultiOutputLayoutr2  rh   r@   rf   r  rA  r   r&   rB  r2  r|   r;  rt   r5  )r  r  r   s      rc   r  zFreeIfNotReusedLine.plan  s    tyy55781<Kdii&&(<(<=K>>!!99177#:#::DLL))$$+DII6&&sD1  'tyy1

3%rb   c                   | j                   j                         t        j                  j                  vsJ | j
                  s| j                  j                  | j                         }| j                  rb| j                   j                         }t        |t        j                        sJ |j                  | d|j                  j                   d       y |j                  |       y y )Nz # z buffer free)rp   rh   r@   rf   r  r2  r   r  r2  rv   rw   r'   rx   r   ry   r  )r  r[  r   r{   s       rc   r\  zFreeIfNotReusedLine.codegen  s    yy!!#177+B+BBBB~~<<00;D224!&"*=*=>>>$s6+B+B+H+H*IVWt$ rb   c                J    | j                   r|j                  S |j                  S re   )r2  _generate_comm_buffer_free_generate_free_if_not_reusedrJ  s     rc   rL  zFreeIfNotReusedLine.codegen_fx  s#    777555rb   Nr  r`  rM  )
r]   r^   r_   r`   r2  r2  rV  r  r\  rL  ra   rb   rc   r=  r=    s0    
ItK
$	%6rb   r=  c                  D    e Zd ZU ded<   ded<   ded<   d
dZddZddZy	)ReinterpretLiner/  rp   	reused_asz	ir.Layoutr{   c                    | S re   ra   r  s     rc   r  zReinterpretLine.plan  s    rb   c                @   t        | j                  t        j                        sJ t        | j                  j                  t        j
                        sJ | j                  j                  | j                  j                         | j                  j                         y re   )
rw   r{   r'   NonOwningLayoutviewr*   r   codegen_deferred_allocationrc  rh   rZ  s     rc   r\  zReinterpretLine.codegen!  sj    $++r'9'9:::$++**B,>,>???00NN##%t{{'7'7	
rb   c                    |j                   S re   )_generate_reinterpretrJ  s     rc   rL  zReinterpretLine.codegen_fx(  r  rb   Nr  r`  rM  )r]   r^   r_   r`   r  r\  rL  ra   rb   rc   rb  rb    s#    

/rb   rb  c                  V    e Zd ZU ded<   ded<   dZded<   dZded<   dd	Zdd
ZddZy)rC  r/  rp   rc  Tr  
delete_oldFr2  c                p   | j                   j                         t        j                  j                  v rK| j
                  j                         t        j                  j                  v sJ t        | j                        S | j
                  j                         t        j                  j                  vsJ | S re   )rp   rh   r@   rf   r  rc  rA  r   r  s     rc   r  zReuseLine.plan3  s    99177#:#::>>**,0G0GGGGDLL))~~&&(0G0GGGGrb   c                p   | j                   j                         t        j                  j                  vsJ | j
                  j                         t        j                  j                  vsJ |j                  | j                  j                  | j                   | j
                  | j                               y re   )
rp   rh   r@   rf   r  rc  r   r   make_buffer_reuserl  rZ  s     rc   r\  zReuseLine.codegen:  sz    yy!!#177+B+BBBB~~&&(0G0GGGGLL**499dnndooV	
rb   c                    |j                   S re   )_generate_reuserJ  s     rc   rL  zReuseLine.codegen_fxA  s    (((rb   Nr  r`  rM  )	r]   r^   r_   r`   rl  r2  r  r\  rL  ra   rb   rc   rC  rC  ,  s1    
JK
)rb   rC  c                      e Zd ZddZy)rA  c                    |j                   S re   )_generate_nullrJ  s     rc   rL  zNullLine.codegen_fxF  r  rb   NrM  rO  ra   rb   rc   rA  rA  E  s    (rb   rA  c                  J    e Zd ZU dZded<   ded<   ded<   ded<   dd	Zdd
Zy)MultiOutputLinezU
    Given a MultiOutputLayout buffer, indexes actual buffer(s) from the result.
    rR  r   rW   result_namearg_nameSequence[Any]indicesc                      fd  j                    j                        }|j                   j                  j                    j
                   d|  j                  j                          y )Nc                l   t        |      dkD  r|d   \  }}t        |t              r |  d| d|dd        S t        |t              r<j                  j                  | j                  t        |            } ||dd        S t        |t              r |  d| d|dd        S t        d|      | S )Nr   []rA   z['z']znon supported index type: )
r   
issubclassr(  r   r   codegen_tuple_accessrw  rW   r  AssertionError)basenamerz  ityper   tuple_accesscodegen_list_tuple_accessr  s        rc   r  z:MultiOutputLine.codegen.<locals>.codegen_list_tuple_accessV  s    7|a"1:qeT*4z1#Q5GQRQSUUu-#'<<#D#D $"2"2CF$L 5\712;OOt,4zA3b5I7STSU;WW()EuMMrb   r   )rx  rz  r   r   declarerw  ending)r  r[  r  r  s   `  @rc   r\  zMultiOutputLine.codegenU  s]    	 $ *$--F||##$T%5%5$6c%ATAT@UV	
rb   c                    |j                   S re   )_generate_multi_outputrJ  s     rc   rL  zMultiOutputLine.codegen_fxm  s    ///rb   Nr`  rM  )r]   r^   r_   r  r`   r\  rL  ra   rb   rc   rv  rv  J  s*     "!M
00rb   rv  c                  <    e Zd ZU ded<   ded<   ded<   d
dZddZy	)IndexPutFallbackLinerR  r   ir.IndexPutFallbackrp   zlist[ir.IRNode | None]rz  c                   | j                   }t        j                  |j                        sJ d |j                  d d D        \  }}| j                  D cg c]*  }|r|j                         n| j                  j                  , }} | j                  j                  |j                         |||g|j                           y c c}w )Nc              3  <   K   | ]  }|j                           y wre   r  r   ts     rc   r   z/IndexPutFallbackLine.codegen.<locals>.<genexpr>z  s     Fq**,F   r$   )rp   r'   is_node_sequencer\   rz  r  r   none_str_generate_index_put_fallbackr  codegen_const_args)r  r[  rp   r   valuesidxrz  s          rc   r\  zIndexPutFallbackLine.codegenw  s    yy""4;;///Fdkk"1oFF ||
 (+C!!#0E0EE
 

 	211  "Aw	
9=9P9P9R	

s   /Cc                    |j                   S re   )r  rJ  s     rc   rL  zIndexPutFallbackLine.codegen_fx  s    555rb   Nr`  rM  r  ra   rb   rc   r  r  q  s    !!
##
6rb   r  c                  2    e Zd ZU ded<   ded<   ddZd	dZy)
ScatterFallbackLinerR  r   ir.ScatterFallbackrp   c           
     8   | j                   }t        j                  |j                        sJ |j                  rd |j                  D        \  }}}n%d |j                  D        \  }}|j
                  d   }|j                         x}r|j                  nt        j                  j                  }| j                  j                  |||j
                  d   ||g|j                  |j                  |j                  |j                  d   |j!                         |       y )Nc              3  <   K   | ]  }|j                           y wre   r  r  s     rc   r   z.ScatterFallbackLine.codegen.<locals>.<genexpr>  s     Jq224Jr  c              3  <   K   | ]  }|j                           y wre   r  r  s     rc   r   z.ScatterFallbackLine.codegen.<locals>.<genexpr>  s     EA!--/Er  rA   r   rF  )rp   r'   r  r\   src_is_tensorconstant_argsr  r   r@   rf   r  r   _generate_scatter_fallbackr  python_kernel_namer   r  )r  r[  rp   r   r6  r   r  r[   s           rc   r\  zScatterFallbackLine.codegen  s    yy""4;;///JdkkJOQsEEJQ$$Q'C!%!22A29L9L//""1%uc2  ##KK!!		
rb   c                    |j                   S re   )r  rJ  s     rc   rL  zScatterFallbackLine.codegen_fx  s    333rb   Nr`  rM  r  ra   rb   rc   r  r    s    !!

(4rb   r  c                  <    e Zd ZU ded<   ded<   ded<   d
dZddZy	)SymbolicCallArgLinerR  r   r  argrP   rf   c                d    | j                   j                  | j                  | j                         y re   )r   "_generate_symbolic_call_arg_helperr  rf   rZ  s     rc   r\  zSymbolicCallArgLine.codegen  s    77$**Mrb   c                    |j                   S re   )_generate_symbolic_call_argrJ  s     rc   rL  zSymbolicCallArgLine.codegen_fx  r  rb   Nr`  rM  r  ra   rb   rc   r  r    s    !!	N5rb   r  c                  F    e Zd ZU ded<   ded<   ded<   ded<   dd	Zdd
Zy)UnbackedSymbolDefsLinerR  r   rW   output_namer   outputs)dict[sympy.Symbol, pytree.KeyPath] | Noneunbacked_bindingsc                z    | j                   j                  | j                  | j                  | j                         y re   )r   )_codegen_unbacked_symbol_defs_for_outputsr  r  r  rZ  s     rc   r\  zUnbackedSymbolDefsLine.codegen  s+    >>dllD,B,B	
rb   c                    |j                   S re   )_generate_unbacked_symbol_defsrJ  s     rc   rL  z!UnbackedSymbolDefsLine.codegen_fx  s    777rb   Nr`  rM  r  ra   rb   rc   r  r    s#    !!L@@

8rb   r  c                  ^    e Zd ZU ded<   ded<   ded<   ded<   dZded<   dd	Zedd
       Zy)AssertSizeStrideLinerR  r   rW   r   r   r   inputop_namec                    | j                   j                  || j                  | j                  | j                  | j
                         y re   )r   _codegen_assert_size_strider   r   r   r  rZ  s     rc   r\  zAssertSizeStrideLine.codegen  s/    00$))TYYT\\	
rb   c                    | j                   S re   )_generate_assert_size_striderh  s    rc   rL  zAssertSizeStrideLine.codegen_fx  s    555rb   Nr`  rM  )r]   r^   r_   r`   r  r\  rj  rL  ra   rb   rc   r  r    s:    !!
I
IKGS

 6 6rb   r  c                  8    e Zd ZU dZded<   ded<   ded<   d	dZy)
AssertDivByZeroLinea  Deferred AOTI runtime check that a sizevar divisor is non-zero.

    Queued from CppWrapperCpu.codegen_cpp_sizevar during the build phase
    and replayed at the position of use so the check lands after the
    declarations of any unbacked symbols it references.
    rR  r   rW   divisorr  c                f    | j                   j                  || j                  | j                         y re   )r   _codegen_assert_div_by_zeror  r  rZ  s     rc   r\  zAssertDivByZeroLine.codegen  s     00t||T\\Rrb   Nr`  r  ra   rb   rc   r  r    s     "!LLSrb   r  c                      e Zd ZU dZdZded<    fdZe	 d	 	 	 	 	 	 	 dd       ZddZ	dd	Z
dd
ZddZddZedd       ZddZedd       ZddZedd       ZddZ	 	 ddZddZddZddZddZddZddZddZddZddZddZ	 	 	 	 	 	 	 	 	 	 ddZ 	 	 	 	 	 	 	 	 	 	 	 	 ddZ!dd Z"	 	 	 	 	 	 	 	 dd!Z#dd"Z$dd#Z%dd$Z&d% Z'd& Z(d' Z)d( Z*d) Z+dd*Z,	 	 d	 	 	 	 	 	 	 dd+Z-dd,Z.	 	 	 	 dd-Z/dd.Z0dd/Z1dd0Z2dd1Z3dd2Z4dd3Z5dd4Z6dd5Z7d6 Z8	 	 	 	 dd7Z9	 d	 	 	 	 	 	 	 	 	 	 	 	 	 dd8Z:dd9Z;dd:Z<dd;Z=d< Z>dd=Z?d> Z@dd?ZAd@ ZB	 	 	 	 	 	 	 	 	 	 	 	 	 	 ddAZCdB ZDddCZEeFj                  	 	 	 	 	 	 ddD       ZHddEZIdF ZJdG ZKdH ZLdI ZMddJZN	 	 	 	 	 	 ddKZOdL ZPddMZQdN ZRddOddPZSddOddQZTddRZUddSZVddTZWddUZX	 	 ddVZY	 d	 	 	 ddWZZddXZ[ddYZ\ddZZ]d[ Z^d\ Z_d] Z`d^ Zad_ Zbd` Zc	 	 	 d	 	 	 	 	 	 	 	 	 ddaZde	 d	 	 	 	 	 ddb       Ze	 	 	 d	 	 	 	 	 	 	 	 	 ddcZfdddZgehde        Zi	 	 	 	 ddfZjdddgZk	 	 	 	 	 	 ddhZlddiZmddjZndk Zodl Zpdm Zqdn Zrdo Zsdp Ztdq Zudr Zvds ZwddtZxdu Zydddddddddv	 ddwZzddddddddxdddy
	 ddzZ{d{ Z|d| Z}d} Z~dd~ZddZed        Z	 ddZd ZddZddZddZddZ	 d	 ddZddZddZddZd ZddZd ZddZd Z	 	 	 	 	 	 	 	 ddZ	 	 	 	 	 	 	 	 ddZd Z	 	 	 	 ddZddZd Zd Zd Zd Zd Zd ZddZd Zed        Zed        Zed        Zed        Zed        Z	 	 	 	 ddZd Zd Z xZS )rR  zB
    Generate outer wrapper in Python that calls the kernels.
    Tr  supports_cachingc                    t                    i  _        t                _        t                _        i  _        t                _	        t                _        t                _        t                _        t                _        t                _        t                _        t                _        t                _        t                _        i  _        d _        i  _        t                _        g  _        d _        d _        d _        d _        d _        t<        j>                  j@                  rdnd _!        t<        j>                  j@                  rdnd _"        d  _#        d _$        i  _%        t                _&        t                _'        d  _(         jS                          g  _*        g  _+         jY                          t[               s j]                           j_                          t<        j>                  j`                  sBt<        j>                  jb                  je                         D ]  \  }} jg                  ||        t        th                   _5        t        th                   _6        i  _7         tq        jr                  d        jt                         _:        tp        jv                  d
 fd       }| _<        i  _=        t                _>        t                _@        t                _A        i  _B        t        t        j                  j                  t        j                  j                  	       _H        g  _I        t                _J        y )Nr    #r>  z
std::move(r   Tc                    j                   j                  |        t        j                  j                  rj
                  j                  |        y y re   )importsr   r&   r   r   r   )r   r  s    rc   add_import_oncez6PythonWrapperCodegen.__init__.<locals>.add_import_once=  s;    LL""4(}}55**44T: 6rb   )debug_printer_leveluse_array_ref)r   rW   r   r>  )Kr$  r%  _pending_input_assertsr   _pending_alignment_copiesr   _names_iterargs_to_buffersr7   r  r;   headerprefixsuffixkernel_declarationswrapper_callkernel_autotune_defsr   subgraph_definitionsr   kernel_autotune_example_argskernel_autotune_tmp_arg_idxsrc_to_kernelkernel_numel_exprlinesr  declare_maybe_referencer  commentr  r@   rf   r  
move_beginmove_endr  supports_intermediate_hooksuser_defined_kernel_cacheunbacked_symbol_declsrU  launcher_fn_nameset_launcher_fn_namecodegened_graph_stackcomputed_sizes_stackwrite_headerr8   write_prefix!write_kernel_autotune_defs_headerr  constant_reprsitemswrite_constant
BufferName	allocatedfreedreusesrE  	lru_cachewrite_get_raw_streamcacher  _metas
_meta_varsr   multi_kernel_statealready_codegened_subgraphsallocated_workspacesr   r&   aot_inductor debug_intermediate_value_printerallow_stack_allocationdebug_printeradditional_filesexternal_kernel_libs)r  r   hashedr  r,  s   `   rc   r%  zPythonWrapperCodegen.__init__  s   BD#:D,&*/'  	 &')+)+$&#1#3 /1$2$4!%3%5"$2$4!6@l" IK)01( .0HR!#
')$*+''*=*=,2 ww228<)+/(  	& L 	" 9C $!!# &("$&!248..0ww ! 6 6 < < > 2f##D&12 $J/1
+-
 57$=I$7$7$=%%%
! 
	; 
	;
  /&(+5<"2"4<FL(46! 1 & 3 3 T T --DD
 !#
 6@\!rb   Nc                D    | r|J |J t        |||      S t               S re   )SubgraphPythonWrapperCodegenrR  )is_subgraphsubgraph_nameparent_wrapperpartition_signaturess       rc   createzPythonWrapperCodegen.createX  s?      ,,,!---/~/C  $%%rb   c                    d| _         y )Ncall)r  r  s    rc   r  z)PythonWrapperCodegen.set_launcher_fn_nameg  s
     &rb   c                D    | j                   j                  | d|        y )Nz = None  # )r  r   )r  r   r   s      rc   r  z#PythonWrapperCodegen.write_constantk  s    k&:;rb   c                   t         j                  j                  j                         }d}||j                  d|j                   }d}t        t        j                  j                        dkD  rd}n0t         j                  j                  j                  j                  rd}| j                  j                  d| dt        j                   d	| d
d       | j                   j                  dd       	 ddlm} | j                   j                  dd       t        j*                  r| j                   j-                  d       t        j.                  j0                  r| j                   j-                  d       | j                   j-                  d       | j                   j-                  d       | j                   j-                  d       | j                   j-                  d       | j                   j-                  d       | j                   j-                  d       t        j.                  j2                  xs( t4        j6                  j9                  t;               d      }| j                   j-                  d| d       d| d}d| d}t        j.                  j<                  }t        j.                  j>                  }	t        j.                  j@                  }
| j                   j-                  d       | j                   j                  d| d | d!|	 d"|
 d#	       | j                   j-                  d$| d%       | j                   j-                  d&       yy# t&        t(        f$ r Y Qw xY w)'z>Write the header section of the generated Python wrapper code.r  Nz
# AOT ID: r   zRfrom torch._inductor.codegen.debug_utils import _print_debugging_tensor_value_infozFfrom torch._inductor.runtime.debug_utils import tracked_empty_strided
z
                aH  
                from ctypes import c_void_p, c_long, c_int
                import torch
                import math
                import random
                import os
                import tempfile
                from math import inf, nan
                from cmath import nanj
                from torch._inductor.hooks import run_intermediate_hooks
                from torch._inductor.utils import maybe_profile
                from torch._inductor.codegen.memory_planning import _align as align
                from torch import device, empty_strided
                from zq import AsyncCompile
                from torch._inductor.select_algorithm import extern_kernels
                z
            Tr   a  
                aten = torch.ops.aten
                inductor_ops = torch.ops.inductor
                _quantized = torch.ops._quantized
                assert_size_stride = torch._C._dynamo.guards.assert_size_stride
                assert_alignment = torch._C._dynamo.guards.assert_alignment
                empty_strided_cpu = torch._C._dynamo.guards._empty_strided_cpu
                empty_strided_cpu_pinned = torch._C._dynamo.guards._empty_strided_cpu_pinned
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
                empty_strided_mtia = torch._C._dynamo.guards._empty_strided_mtia
                reinterpret_tensor = torch._C._dynamo.guards._reinterpret_tensor
                alloc_from_pool = torch.ops.inductor._alloc_from_pool
                async_compile = AsyncCompile()
            )_SymmetricMemoryzs
                empty_strided_p2p = torch._C._distributed_c10d._SymmetricMemory.empty_strided_p2p
                zfrom torch.cuda import nvtxz import triton.profiler as protonz%import triton.profiler.language as plzCfrom triton.profiler.hooks import HookManager as _ProtonHookManagerzimport tritonzimport atexitz	import oszetriton.set_allocator(lambda size, align, stream: torch.empty(size, dtype=torch.uint8, device='cuda'))protonzos.makedirs("z", exist_ok=True)zos.path.join("z", "inductor")z", "inductor.chrome_trace")z^from torch._inductor.runtime.proton_utils import process_proton_trace as _proton_process_tracez
                def _proton_finalize_and_postprocess():
                    proton.finalize()
                    _trace_path = z
                    if os.path.exists(_trace_path):
                        _proton_process_trace(
                            _trace_path,
                            group_by_sm=z0,
                            split_invocations=z0,
                            per_cta_occupancy=z,,
                        )
                z5if not _ProtonHookManager.active_hooks: proton.start(z], backend="instrumentation", data="trace"); atexit.register(_proton_finalize_and_postprocess)pl.enable_semantic("triton"))!torch_guardsTracingContexttry_getaot_graph_namerY   r&   r  r  	_inductortest_configstrack_memory_lifecycler  r   r%   r]   r  torch._C._distributed_c10dr  AttributeErrorImportErrorannotate_trainingr   r   proton_profilingproton_output_dirospathr   r   proton_group_by_smproton_split_invocationsproton_per_cta_occupancy)r  contextaot_config_commentinductor_debug_utilsr  
output_dirproton_name
trace_pathgroup_by_smsplit_invocationsper_cta_occupancys              rc   r  z!PythonWrapperCodegen.write_headern  s4   --..6687#9#9#E#-g.D.D-E!F!v""CCDqH#w __##00GG#l #$ % $,,- .%& '!$ ' 	 	
* 	 ! 	 	
$	 DKK 	   ##KK!!"?@==))KK!!"DEKK!!"IJKK!!U KK!!/2KK!!/2KK!!+.KK!!G  88 BGGLL=J KK!!M*=N"OP*:,nEK)*5PQJ --::K & F F & F FKK!!p KK# $., /) *5 6//@.A B//@.A B KK!!  +} -DD
 KK!!"@A[ *	 , 		s   6#M M+*M+c                     y re   ra   )r  r  s     rc   include_extra_headerz)PythonWrapperCodegen.include_extra_header      rb   c                    | j                   j                  dt        j                   d       	 ddlm} | j                   j                  dd       y # t        t        f$ r Y y w xY w)Na3  
                import torch
                from math import inf, nan
                from torch._dynamo.testing import rand_strided
                from torch._dynamo.utils import preserve_rng_state
                from torch._inductor.select_algorithm import AlgorithmSelectorCache
                from aH   import AsyncCompile

                async_compile = AsyncCompile()
                generate_example_value = AlgorithmSelectorCache.generate_example_value
                empty_strided_cuda = torch._C._dynamo.guards._empty_strided_cuda
                empty_strided_xpu = torch._C._dynamo.guards._empty_strided_xpu
            r   )_cuda_getCurrentRawStreamzU
                get_raw_stream = torch._C._cuda_getCurrentRawStream
                Tr   )r  r   r%   r]   torch._Cr/  r  r  )r  r/  s     rc   r  z6PythonWrapperCodegen.write_kernel_autotune_defs_header  sw    !!(( $,,- .	
 
	:%%,, 	 -  ^, 		s   #A A%$A%c                   dt         j                   d}t        j                  j                  r]| j
                  j                  |       | j
                  j                  t        j                  j                  j                  d             t        j                  j                  s`| j                  j                  |d       | j                  j                  t        j                  j                  j                  d             y y )NzU
            import triton
            import triton.language as tl
            from z+ import start_graph, end_graph
            get_raw_streamTr   )r+   r]   r&   r   r   r   r   r   r@   rf   r  import_get_raw_stream_asr  r  r  
import_strs     rc   write_triton_header_oncez-PythonWrapperCodegen.write_triton_header_once  s     $,,- .

 ==11&&--j9&&00"";;<LM ww""LL
$7LL"""";;<LM #rb   c                   t         j                  j                  j                  d      }t        j
                  j                  r6| j                  j                  |      s| j                  j                  |       t         j                  j                  s8| j                  j                  |      s| j                  j                  |       y y y )Nr2  )r@   rf   r  r3  r&   r   r   r   containsr   r  r  )r  import_get_raw_stream_strs     rc   write_get_raw_stream_headerz0PythonWrapperCodegen.write_get_raw_stream_header  s    $%GG$6$6$O$O%
! ==11--667PQ**445NOww""<<(()BC&&'@A D #rb   c                $    | j                          y re   )r:  r  s    rc    write_get_raw_stream_header_oncez5PythonWrapperCodegen.write_get_raw_stream_header_once  s    ((*rb   c                   t        |      }|| j                  vrdt        | j                         }|| j                  |<   | j                  j	                  | d|        t
        j                  j                  r;| j                  j	                  | d|        | j                  j                  |       | j                  |   S )Nmetar   )reprr  r   r  r   r&   r   r   r   r  r   )r  r>  vars      rc   add_meta_oncez"PythonWrapperCodegen.add_meta_once   s    Dzt{{"T[[)*+C #DKKKK!!SETF"34}}55**44uCv5FG##C({{4  rb   c                z    | j                         D cg c]  }|j                  | j                         c}S c c}w re   )get_graph_outputsr  r  r  r   s     rc   get_output_refsz$PythonWrapperCodegen.get_output_refs.  s<     =A<R<R<T
78A 1 12
 	
 
s   "8c                     y re   ra   r  s    rc   mark_output_typez%PythonWrapperCodegen.mark_output_type4      rb   c                6    t         j                  j                  S re   )r@   rf   r  r  s    rc   get_graph_inputsz%PythonWrapperCodegen.get_graph_inputs7  s     ww###rb   c                6    t         j                  j                  S re   )r@   rf   r  r  s    rc   rC  z&PythonWrapperCodegen.get_graph_outputs<  s    ww$$$rb   c                   | j                         j                         D ]  \  }}t        |t        j                  t
        j                  t
        j                  t
        j                  f      rO|t        j                  j                  vrlt        |j                               dk(  r| j                  |j                               }| j                  |j                               }||f| j                   |<    y r#  )rJ  r  rw   r   Basicr'   TorchBindObjectGeneratorStateOpaqueObjectStater@   rf   graph_input_namesr<   rM  rQ  rN  r  )r  r   bufr   r   s        rc   codegen_input_size_assertsz/PythonWrapperCodegen.codegen_input_size_asserts?  s    ..0668 	?ID#KK&&%%((	  177444 S\\^,1++CLLN;D--cnn.>?F15vD''-+	?rb   c                `   | j                   j                  d       | j                         j                         D ]r  \  }}t	        |t
        j                  t        j                  f      r1d| d}| j                   j                  |       d| d}| j                   j                  |       t y )Nz(# make sure graph inputs are not nan/infzassert not z.isnan().any().item()z.isinf().any().item())	r  r   rJ  r  rw   r   rM  r'   rN  )r  r   rR  r   s       rc   codegen_input_nan_assertsz.PythonWrapperCodegen.codegen_input_nan_assertsW  s    HI..0668 	(ID##R-?-?@A &;<DKK!!$' &;<DKK!!$'	(rb   c                :    | j                   j                  d       y )NzV

            async_compile.wait(globals())
            del async_compile
            )r  r   r  s    rc   write_async_compile_waitz-PythonWrapperCodegen.write_async_compile_waita  s    	
rb   c                    dj                  |      }t        |      dk(  r|dz  }| j                  j                  | d       | j                  j                  d       y )Nr  rA   ,z = argszargs.clear())r   r   r  r   )r  input_nameslhss      rc   
write_argszPythonWrapperCodegen.write_argsj  sP    ii${q 3JCWo.n-rb   c                    t         j                  r| j                  j                  d       d}|S | j                  j                  d| j                   d       d}|S )Na  
                class Runner:
                    def __init__(self, partitions):
                        self.partitions = partitions

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

                    def call(self, args):
                r$   z
                def z(args):
                rA   )r&   graph_partitionr  r   r  r  prefix_indents     rc   !write_launcher_fn_call_get_indentz6PythonWrapperCodegen.write_launcher_fn_call_get_indentq  sm    !!KK M  KK**+ ,
 Mrb   c                6    t         j                  j                  S re   )r@   rf   rQ  r  s    rc   get_graph_input_namesz*PythonWrapperCodegen.get_graph_input_names  s    ww(((rb   c                   | j                   J | j                          | j                         }| j                  j	                  |      5  t
        j                  j                  rA| j                  j                  t        j                  j                  j                                t        j                  j                         }t
        j                  r| j                  j                  d| d       | j                         x}r| j!                  |       | j#                          t%               rt'        |       r| j)                          d d d        y # 1 sw Y   y xY w)Nz0training_annotation = nvtx._device_range_start(''))r  rW  ra  r  r   r&   r   debug_sync_graphr   r@   rf   r  synchronizeget_training_phaser  rc  r\  codegen_inputsr9   r8   "codegen_input_size_and_nan_asserts)r  r`  phaserQ  s       rc   r  z!PythonWrapperCodegen.write_prefix  s   $$000%%'>>@[[. 	:}}--%%agg&8&8&D&D&FGGG..0E''%%FugRP %)$>$>$@@ @ 12!
 -.<TB779)	: 	: 	:s   
DEEc                    t         j                  r| j                          t         j                  r| j	                          y y re   )r&   size_assertsrS  nan_assertsrU  r  s    rc   rj  z7PythonWrapperCodegen.codegen_input_size_and_nan_asserts  s1    ++-**, rb   c                    |D ]C  }|| j                   v s| j                   j                  |      \  }}| j                  |||d       E y )Nr  )r  r1  write_assert_size_stride)r  rZ  r   r   r   s        rc   codegen_deferred_input_assertsz3PythonWrapperCodegen.codegen_deferred_input_asserts  sO     	KDt222#::>>tDf--dD&'J	Krb   c           	     @    | j                  t        | ||||             y)z7Queue an assert_size_stride for emission during replay.N)r   r  )r  r   r   r   r  s        rc   rp  z-PythonWrapperCodegen.write_assert_size_stride  s     	+D$fgNOrb   c                @    |j                  d| d| d| d|d	       y)zEmit one assert_size_stride line to `code` (replay-phase target).

        Subclasses override to change the emitted form (e.g., C++ assert with
        an AOTI runtime env guard).
        zassert_size_stride(r  r   Nr   )r  r[  r   r   r   r  s         rc   r  z0PythonWrapperCodegen._codegen_assert_size_stride  s-     	,TF"TF"VHBwkQRSTrb   c                <    | j                  t        | ||             y)zQueue a div-by-zero AOTI check for emission during replay.

        codegen_cpp_sizevar is called in both build and replay phases;
        set_writeline routes the WrapperLine appropriately in either case.
        N)r   r  )r  divisor_strr  s      rc   write_assert_div_by_zeroz-PythonWrapperCodegen.write_assert_div_by_zero  s     	*4gFGrb   c                    t        d      )zEmit one div-by-zero AOTI check to `code` (replay-phase target).

        Only emitted by C++ wrappers; the Python wrapper relies on Python's
        native ZeroDivisionError instead of pre-checking divisors.
        z6AOTI div-by-zero check is only emitted by C++ wrappersre  )r  r[  rv  r  s       rc   r  z0PythonWrapperCodegen._codegen_assert_div_by_zero  s     "D
 	
rb   c                   t         j                  j                  ryt         j                  j                  }|syt	        t         j                  j
                        }|D ]?  }||vst         j                  j                  |   }| j                  j                  |       A | j                  r1dt         j                  _	        | j                  j                  d       yy)zPopulate pending alignment copies for non-mutated inputs.
        Called from the scheduler after mutated_input_idxs is computed.NTz6from torch._C._dynamo.guards import copy_if_misaligned)r@   rf   r  inputs_to_checkr   mutated_input_idxsrQ  r  r   _defers_input_alignmentr  r   )r  rz  mutated_idxsr  r   s        rc   register_alignment_check_inputsz4PythonWrapperCodegen.register_alignment_check_inputs  s     77''11 "!''"<"<=" 	9C,&ww005..2248	9 )).2AGG+LL""H *rb   c                    t         j                  j                  ry|D ]C  }|| j                  v s| j                  j	                  |       | j                  | d| d       E y)z~Emit alignment check + clone just before the first kernel
        that reads each input, hiding the cost behind GPU execution.Nz = copy_if_misaligned(r   )r@   rf   r  r  discardr   )r  rZ  r   s      rc   !codegen_deferred_alignment_copiesz6PythonWrapperCodegen.codegen_deferred_alignment_copies  sa     77 	GDt555..66t<$'=dV1EF	Grb   c                   | j                          t        |      }t        j                  j                  r=| j
                  j                  | d| d       t        j                  j                  r|S | j                  | d| d       |S )N = get_raw_stream(r   )
r:  r0   r&   r   r   r   r   r@   rf   r  )r  r  r  r   s       rc   r  z)PythonWrapperCodegen.write_get_raw_stream  s}    ((*":.==11&&00&*:,a8 ww""$1*Q?@rb   c                     | j                   d   S )N)r  r  s    rc   get_codegened_graphz(PythonWrapperCodegen.get_codegened_graph  s    ))"--rb   c                :    | j                   j                  |       y re   )r  r   )r  rf   s     rc   rX  z)PythonWrapperCodegen.push_codegened_graph  s    ""))%0rb   c                6    | j                   j                         S re   )r  r1  r  s    rc   r{  z(PythonWrapperCodegen.pop_codegened_graph  s    ))--//rb   c                P    ddl m} | j                  j                   ||            S )Nr   )deepcopy)copyr  r  r   )r  rU  r  s      rc   rT  z(PythonWrapperCodegen.push_computed_sizes  s!    !((//0HIIrb   c                6    | j                   j                         S re   )r  r1  r  s    rc   ry  z'PythonWrapperCodegen.pop_computed_sizes!  s    ((,,..rb   c                .    t        | j                         S re   )nextr  r  s    rc   next_kernel_suffixz'PythonWrapperCodegen.next_kernel_suffix$  s    t''()*rb   c                   |dkD  rd|J d}| j                   j                  |      s| j                   j                  |       | j                  t        || j                  ||             n%| j                  t        || j                               t        j                  j                  r| j                          | j                  j                  dt        j                  j                  j                  |       d       | j                  j                          t!        |       r| j#                          | j                  j                  t%        |       d| d       || _        || _        y )NrA   zLfrom torch._dynamo.graph_bytecode_inputs import get_external_object_by_indexr  r  r  r   )r  r8  r   r  r  r  r&   r   r   r6  r   r@   rf   r  r  rY  r8   r:  r0   _num_streams)r  r  r  r  import_lines        rc   codegen_device_guard_enterz/PythonWrapperCodegen.codegen_device_guard_enter'  sE    ?-999/  <<((5&&{3NN;55.	 NN- A A
 ==11))+&&00**77
CDAF &&002248002&&00&z233Ej\QRS -7)!,rb   c                    t        | d      r4| j                  dkD  r%| j                  t        | j                               n| j                  t	                      t
        j                  j                  r| j                  j                          y y )Nr  rA   )
r   r  r   r  r  r&   r   r   r   r|  r  s    rc   codegen_device_guard_exitz.PythonWrapperCodegen.codegen_device_guard_exitS  si    4(T->->-BNN:4;L;LM NN79:==11&&224 2rb   c                    t         j                  j                  j                  x}t	        d| dt        |             t        |      }| j                  |       |S )a  Generate data structure for entering a CUDA Stream context.

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

        Raises:
            ValueError: If this function is called while the previous stream context isn't exited.
        z!Nested stream context switching: z -> r  )r@   rf   rj   current_stream_name
ValueErrorr1   r  r   )r  r  r  ctx_entrances       rc   codegen_cuda_stream_enterz.PythonWrapperCodegen.codegen_cuda_stream_enter]  sf     $%77#4#4#H#HHU34G3H":./1  2ZH|$rb   c                6    | j                  t                      y)z:Generate data structure for exiting a CUDA Stream context.N)r   r
  r  s    rc   codegen_cuda_stream_exitz-PythonWrapperCodegen.codegen_cuda_stream_exits  s    023rb   c                   |r,t         j                  r| j                  j                  ddj	                  |      z   dz          | j                  j                  d       | j                  j                          | j                  j                  d       | j                  j                          | j                  j                  d       | j                  j                  d       | j                  j                  d       | j                  j                  d	dj	                  |      z   dz          y | j                  j                  d
       y )Nzreturn_vars = (r  , )zfor var in return_vars:z!if isinstance(var, torch.Tensor):z#assert not var.isnan().any().item()z#assert not var.isinf().any().item()r$   zreturn (z	return ())r&   rn  r  r   r   rY  r|  )r  output_refss     rc   generate_returnz$PythonWrapperCodegen.generate_returnw  s   !!!!++%		+(>>F !!++,EF!!++-!!++,OP!!++-!!++,QR!!++,QR!!--a0''
TYY{5K(Ke(ST''4rb   c                     y re   ra   r  results     rc   generate_before_suffixz+PythonWrapperCodegen.generate_before_suffix  rH  rb   c                    t         j                  rNdj                  | j                        t	        | j                        dk(  rdndz   }|j                  d| d       y y )Nr  rA   rY  r  z-
                runner = Runner(partitions=[z{])
                call = runner.call
                recursively_apply_fns = runner.recursively_apply_fns
                )r&   r^  r   all_partition_namesr   r   )r  r  all_partition_name_lists      rc   generate_after_suffixz*PythonWrapperCodegen.generate_after_suffix  se    !!&*ii0H0H&I43349r'# MM--D,E F "rb   c                     y re   ra   r  s     rc   generate_endz!PythonWrapperCodegen.generate_end  rH  rb   c                    |j                   }|1|t        v r)t        |   j                  }| ||| j                         y | j                  t	        | |             y re   )r  rI   pythonr   r  )r  rp   r  custom_codegens       rc   generate_fallback_kernelz-PythonWrapperCodegen.generate_fallback_kernel  sV    ))7.J#J9'BIIN)tT^^4,T489rb   c                   |j                   D ]  }| j                  |        | j                  t        | |             |j                   D ]8  }t	        |j
                  t        j                        s(|j                  |        : y)z7Generate .out() call with pre-allocated output buffers.N)	r  codegen_allocationr   r  rw   r{   r'   Layoutcodegen_size_asserts)r  rp   r  s      rc    generate_extern_kernel_multi_outz5PythonWrapperCodegen.generate_extern_kernel_multi_out  sn    55 	.H##H-	./d;<55 	4H(//2995--d3	4rb   c                    |j                  |        | j                  t        | |             t        |j                  t
        j                        r|j                  |        y y re   )codegen_commentr   r  rw   r{   r'   r  r  r#  s     rc   generate_extern_kernel_allocz1PythonWrapperCodegen.generate_extern_kernel_alloc  sI    T",T489dkk299-%%d+ .rb   c           
        t        |j                  t        j                        }|j	                         }|j                         }|j                         }| j                  }t        j                  r	d|v rd| }|r5| j                  | j                   | ddj                  |       d|        y | j                  | j                   | d| ddj                  |       d|        | j                  rKt        j                  r:|7t        d   dxx   d	z  cc<   | j                  d
|j                   d| d       y y y y )Nview_as_complex.clone()r  r  r   r   inductorintermediate_hooksrA   zrun_intermediate_hooks()rw   r{   r'   
NoneLayoutrh   get_origin_noder  r  r&   memory_planningr   r  r   r  generate_intermediate_hooksr   r   )r  extern_kernelr  	no_returnr  origin_noder  r  s           rc   r  z9PythonWrapperCodegen._generate_extern_kernel_alloc_helper  s;    }33R]]C	#,,.#335#335!!&7;&F  x(FNNdll^K=$))D/9J!F8TUNN<<.SQtyy>OqQWPXY 0066+$%9:a?:-k.>.>-AK=PQR , 7 1rb   c                \    |j                  |        | j                  t        | |             y re   )r  r   r  r#  s     rc   generate_extern_kernel_outz/PythonWrapperCodegen.generate_extern_kernel_out  s&     	T"*467rb   c                    t         j                  j                  j                  }|j	                  ||d d d       |j                  d|r|n|        |5  | j                  | ddj                  |       d       d d d        y # 1 sw Y   y xY w)Nexternzout=r  r  r   )r@   rf   wrapper_coder  set_printer_argsr   r   r   )r  r  outout_viewr  r[   stack_tracesdebug_printer_managers           rc   r  z7PythonWrapperCodegen._generate_extern_kernel_out_helper  s     !" 4 4 B B..tVT4Rdx8S9:;" 	;NNfXQtyy&7q9:	; 	; 	;s   'BBc                    |j                   }|j                  }|rRt        j                  j                  j                  |      }t        j                  j                  j                  |      }|j                  j                          d}dj                   fd|D              }dj                   fd|D              }t        j                   |j                        }d}| d|j                   d}| d| d| d| }	| d|	 d	}
|
S )
Nz.data_ptr()r  c              3  J   K   | ]  }t         j                  |        y wre   rR  val_to_arg_strr   dimr  s     rc   r   zRPythonWrapperCodegen._generate_tma_descriptor_call_experimental.<locals>.<genexpr>  s     XC-<<T3GX    #c              3  J   K   | ]  }t         j                  |        y wre   r  r  s     rc   r   zRPythonWrapperCodegen._generate_tma_descriptor_call_experimental.<locals>.<genexpr>  s$      
?B //c:
r  z$triton.tools.experimental_descriptorz.create_d_tma_descriptorr  r   )dims
block_dimsr@   rf   rn   r"  optimization_hintstensorr  r   rR  r  element_sizerank)r  descapply_size_hintsr  r  ptrr  r  r   r  r	  s   `          rc   *_generate_tma_descriptor_call_experimentalz?PythonWrapperCodegen._generate_tma_descriptor_call_experimental  s    yy__
77##55d;D))<<ZHJ..01=yyXSWXXYY 
FP
 

 ,::4ARARS7xx		{*:;bbB|n=QtfArb   c                    |j                   }|r)t        j                  j                  j	                  |      }d}| d}|j
                  j                          d| }| d| d}|S )Nz/triton.tools.tensor_descriptor.TensorDescriptorz.from_tensorr  r  r   )block_shaper@   rf   rn   r  r  r  )r  r  r  r  r  r   r  r	  s           rc   $_generate_tma_descriptor_call_stablez9PythonWrapperCodegen._generate_tma_descriptor_call_stable  so    &&''**==kJKBx|$++//12"[MBQtfArb   c                    t        |t        j                        r| j                  ||      S t        |t        j                        sJ | j                  ||      S re   )rw   r'   TMADescriptorExperimentalr  TMADescriptorStabler  )r  r  r  s      rc   _generate_tma_descriptor_callz2PythonWrapperCodegen._generate_tma_descriptor_call	  sW    dB889BB&  dB$:$:;;;<<TCSTTrb   c                    | j                  |      }|j                   d| | j                   }| j                  |       y Nr   )r  r   r  r   )r  r  r	  r   s       rc   generate_tma_descriptorz,PythonWrapperCodegen.generate_tma_descriptor  s:    11$7))Cvdkk]3trb   c                :    | j                  t        | |             y re   )r   r  r#  s     rc   generate_scatter_fallbackz.PythonWrapperCodegen.generate_scatter_fallback  s    *467rb   c	                    | ddj                  t        t        |             }	|j                  d      r|	dj                  dg|z         z  }	n|r|	dt	        |       z  }	|	dz  }	| j                  |	       y )Nr  rY  zaten.scatter_reducer  r  z	, reduce=r   )r   maprW   r  r?  r   )
r  r   r\   r  r  r  rF  r   r[   r   s
             rc   r  z/PythonWrapperCodegen._generate_scatter_fallback  s{     %%QsxxC0@'A&BC(()>?DIIrdVm,,D)DL>22trb   c                `   g }|j                   dd  }t        |      }t        |j                        D ]^  \  }}|j                  |   9t	        |      }t        |t        j                        sJ |j                  |       N|j                  d        ` | j                  t        | ||             y )Nr$   )r\   iter	enumeraterz  r  rw   r'   r)   r   r   r  )r  rp   rz  valid_indicesiter_valid_indicesr   _r6  s           rc   generate_index_put_fallbackz0PythonWrapperCodegen.generate_index_put_fallback.  s    *,AB!-0dll+ 	%DAq||A*/0!%333u%t$	% 	+D$@Arb   c                |    ddj                  |       d}||||g}| j                  | j                  ||             y )Nr}  r  r~  )r   r   wrap_kernel_call)r  r  r   rz  r  
accumulateindices_strr  s           rc   r  z1PythonWrapperCodegen._generate_index_put_fallback=  sA    $))G,-Q/;
3t,,VT:;rb   c           
     `    | j                  | d| ddj                   |              d       y )Nr   r  r  r   )r   r   )r  buf_namer  get_argsop_overloadr  r  s          rc   ,generate_fallback_kernel_with_runtime_lookupzAPythonWrapperCodegen.generate_fallback_kernel_with_runtime_lookupB  s2     	(3'9&:!DIIhj<Q;RRSTUrb   c                f    t        d      5  | j                  |      cd d d        S # 1 sw Y   y xY w)NzPythonWrapperCodegen.generate)r   	_generater  is_inferences     rc   generatezPythonWrapperCodegen.generateM  s,    9: 	0>>,/	0 	0 	0s   '0c                &    t         j                  ryy)Nr$   rA   )r&   r^  r  s    rc   get_wrapper_call_indentz,PythonWrapperCodegen.get_wrapper_call_indentQ  s    !!rb   c              #  t   K   | j                   }dfd}	 || _         | || _         y# || _         w xY ww)zTemporarily redirect self.writeline to write into `target`.

        Plain strings go through `writeline`.
        WrapperLines are codegen'd into `target` directly.
        c                Z    t        | t              r| j                         y  |        y re   )rw   rT   r\  )r   targetr   s    rc   dispatchz4PythonWrapperCodegen.set_writeline.<locals>.dispatchd  s     $,V$$rb   N)r   z2LineContext | DeferredLineBase | WrapperLine | strr   r>  rt  )r  r  r   oldr  s    ``  rc   set_writelinez"PythonWrapperCodegen.set_writelineW  s6      nn	 	!%DNN DNSDNs   8, 8	58c                    | j                   j                  }t        j                  j                  r| j
                  j                  |       y | j                  j                  |       y re   )r  kernel_defsr&   r   r   r  r   r  )r  r  s     rc   _write_multi_kernel_defsz-PythonWrapperCodegen._write_multi_kernel_defsq  sF    --99==11%%,,[9KK{+rb   c                
   t         j                  r| j                          t        j                         5 }|j                  | j                  j                                t         j                  r| j                  |       t         j                  r| j                          | j                  |       t         j                  j                  r*t         j                  j                  s| j                          | j!                  | j                  | j                  j"                        5  | j$                  D ]I  }t'        |t(              r|j+                  | j                         /| j                  j#                  |       K 	 d d d        | j-                          | j/                         }| j1                          t         j                  j2                  rA| j                  j#                  t4        j6                  j8                  j;                                t         j                  r| j=                          t         j                  j>                  r| jA                          t         j                  j                  r*t         j                  j                  s| jC                          t         j                  j                  r| jE                          t         jF                  r+t         jH                  s| j                  j#                  d       | jK                  |       d d d        tM               }|jO                  | jP                         |j#                  d       |jO                  | jR                         t4        j6                  jT                  r>t4        j6                  jH                  r$t4        j6                  jV                  r
tM               }|jO                  | jX                         | j[                          |jO                  | j\                         | j_                         }|j                  |      5  |jO                  | j                         d d d        | ja                  |       |jO                  | jb                         | je                  |       | jg                  |       | ji                  |       |jk                         | jl                  jk                         fS # 1 sw Y   ^xY w# 1 sw Y   xY w# 1 sw Y   xY w)Nz+nvtx._device_range_end(training_annotation)r  )7r&   profile_bandwidthr6  r   	ExitStackenter_contextr  r   profiler_mark_wrapper_call#generate_profiler_mark_wrapper_callgenerate_start_graphrun_wrapper_ir_passesr   store_cubinr   !generate_reset_kernel_saved_flagsr
  r   r  rw   rT   r\  r  rE  rG  rf  r@   rf   r  rg  generate_end_graphr  generate_proton_finalize generate_save_uncompiled_kernelsgenerate_and_run_autotune_blockr  r  r  r7   r   r  r  r  is_const_graphr  finalize_prefixr  r  r  r  r  r  add_benchmark_harnessgetvaluewithlinemapr  )r  r  stackr   r  r  wrapper_call_indents          rc   r  zPythonWrapperCodegen._generatex  s   ##))+!!# .	.u 1 1 8 8 :;0088?''))+&&|4}}((1W1W668 ##D$5$5t7H7H7R7RS : JJ :D!$4T%6%67))33D9:: ))+..0K!!#}}--!!++AGG,>,>,J,J,LM'''')}}----/}}((1W1W557}}55446 ''0B0B!!++A   -].	.b  !dll#dkk" 77 3 38N8N#%F 	d//0dkk""::<]]./ 	-MM$++,	- 	##F+dkk"""6*&!""6* &&($$88:
 	
}: :.	. .	.B	- 	-s2   C0S(%AS?FS(.S5S%	 S((S25S>c                \   | j                   j                  d       i }t        j                  j                  r_t
        j                  j                  rEt        t
        j                  j                        D ci c]  \  }}| j                  |      | }}}| j                   j                         dz   | j                  j                         z   t        j                  t        j                  k(  rkt!        j"                  t%               dd      5 }|j'                  j)                  d             |j*                  }ddd       t        j,                  d       t/        d	d
 fd       	 t1        |       yc c}}w # 1 sw Y   ExY w# t2        $ r}t5        d|       |d}~ww xY w)z
        Compose self.kernel_autotune_defs and self.kernel_autotune_calls into a single block of
        code and execute it to trigger Triton kernel compilation and auto-tuning
        zQ
            async_compile.wait(globals())
            del async_compile
        r  z.pyF)dirr  deletezutf-8NzAuto-tuning code written to %sartifactc                     dddS )N&inductor_autotune_at_compile_time_codestring)r   encodingra   ra   rb   rc   r   zFPythonWrapperCodegen.generate_and_run_autotune_block.<locals>.<lambda>  s    @$! rb   c                      S re   ra   )tuning_codes   rc   r   zFPythonWrapperCodegen.generate_and_run_autotune_block.<locals>.<lambda>  s    { rb   )metadata_fn
payload_fnz%Failed to run autotuning code block: )r  r   r&   r   r   r@   rf   autotuning_inputsr  get_autotuning_input_namer   r   r(   levelloggingDEBUGtempfileNamedTemporaryFiler   writeencoder   debugr   exec	ExceptionRuntimeError)r  scoper  vf	file_pather+  s          @rc   r  z4PythonWrapperCodegen.generate_and_run_autotune_block  s   
 	!!((	
 ==11agg6O6O ((A(ABC ..s3Q6E 
 %%..0((1134 	
   GMM1 ,,Ke #**734FF		#
 !!0 	 +	
	Se$?# #(  	S!FqcJKQRR	Ss*   9E<-F/F F	F+F&&F+c                \    ddl m}  ||       j                  | j                        | _        y )NrA   )MemoryPlanner)r  rA  r  r  )r  rA  s     rc   memory_planz PythonWrapperCodegen.memory_plan  s     2"4(--djj9
rb   c                   | j                         }t        j                  j                  |      }| j                  rt        | j                  d   t              r| j                  d   j                  j                  |vri| j                  j                          | j                  rCt        | j                  d   t              r&| j                  d   j                  j                  |vrit               g}g }g }t        t        | j                              D ](  }| j                  |   }t        |t              r"|j                  |d         | j                  |<   Et        |t              r}|j                  t                      t         j"                  s|j                  | j$                         t        j&                  |j                        5  t)               | _        d d d        t        |t*              s|j                  |j                                t         j"                  s|j                         | _        + |j                  |j                                t        |      dk(  sJ t-        d |D              }y # 1 sw Y   xxY w)Nr  r   c              3  4   K   | ]  }|j                     y wre   )r+  )r   ss     rc   r   z9PythonWrapperCodegen.memory_plan_reuse.<locals>.<genexpr>'	  s      +
./A))+
s   )rC  r@   rf   _get_output_namesr  rw   r  rp   r   r1  r!  r  r   r  rQ  r   r&   rB  r9  set_graph_handlerr  rw  sum)	r  r  	out_namesplanning_statespast_planning_statespeak_estimate_stackr   r   _total_allocated_buffer_sizes	            rc   memory_plan_reusez&PythonWrapperCodegen.memory_plan_reuse	  s   ((*GG--g6	 JJ4::b>+=>

2##((	9 JJNN JJ4::b>+=>

2##((	9 /01!;=s4::' 	CA::a=D$ 23 $		/"*= >

1D"34&&':'<=,,'..t/A/AB,,TZZ8 E-B-D*E ED"23$++O,?,?,AB,,)<)@)@)BD&	C 	##O$7$7$9:?#q(((
 (+ +
3G+
 (
$E Es   JJ	c                    |r!t         j                  r| j                          y t         j                  rt	               | _        | j                          y re   )r&   r  rB  rB  r  r9  rN  r   s     rc   r  z*PythonWrapperCodegen.run_wrapper_ir_passes+	  s9    F22((%:%<"""$rb   c           	     X  
 | j                   
t        j                  
fd       }t        j                  
fd       }d
fd}t        |t        j
                        rOt        |t        j                        r|v ry 
j                  | d|        j                  |        ||       y t        |t        j                        rt        |j                               D ]^  \  }}t        |t        j                        s!|vs&
j                  | d ||       d| d       j                  |        ||       ` t        |j                               D ]V  \  }}	t        |	t        j                        s!|	vs&
j                  |	 d ||       d| d       j                  |	       X y t        |t        j                  t        j                  t        j                   f      ry t"        j$                  j&                  j(                  ry t+        dt-        |             )	Nc                <    j                  |  d|  d       |  dS )Nz_size = z.size()_sizert  r   r[  s    rc   sizeofzDPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.sizeof<	  s(    NNdV8D69:V5>!rb   c                <    j                  |  d|  d       |  dS )Nz
_stride = z	.stride()_stridert  rS  s    rc   strideofzFPythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.strideofA	  s)    NNdV:dV9=>V7##rb   c                x   ddl m}m} t        j                  j
                  j                  j                  j                         D ]r  \  }}|| k(  st        |t        j                        s'|vs, |||j                  |j                  f      rLj                  | d|         j                  |       t y )Nr   r!   r   )torch.utils._sympy.symbolr"   r#   r@   rf   rn   	shape_envreplacementsr  rw   r   SymbolUNBACKED_INTUNBACKED_FLOATr   r   )symr"   r#   r   tgt
bound_varsr[  s        rc   maybe_emit_replacement_aliasesz\PythonWrapperCodegen.codegen_input_symbol_assignment.<locals>.maybe_emit_replacement_aliasesF	  s     GGG,,66CCIIK 
(S3J"35:-*d//1D1DE NNcU#cU#34NN3'
(rb   r   r}  r~  zUnknown value type: )r_  r  r   r>  )r  rE  r  rw   r   r
   r\  r   r   r'   r   r  rM  rN  rN  rO  rP  r  r  r&   r^  r  r   )r  r   r  ra  rT  rW  rb  r  r   r   r[  s      `      @rc   codegen_input_symbol_assignmentz4PythonWrapperCodegen.codegen_input_symbol_assignment4	  s    {{		" 
	" 
	$ 
	$	(( eUZZ(eU\\2ez6INNeWCv./NN5!*51r||,&u~~'78 9	TdELL1d*6LNNdV3vd|nAcU!#DENN4(248	9
  ))9)9);< +Vfell3j8PNNfXS$0@#a#HINN6*+ B&&(9(92;O;OP
 %%55$';DK=%IJJrb   c           	        t        t        j                            }| j                         }|j	                         D cg c]$  \  }}t        |t        j                        s!||f& c}}|j	                         D cg c]$  \  }}t        |t        j                        r!||f& c}}z   }|D ]  \  }}| j                  |||        	 	 	 	 dd}|D ])  \  }	}t        |t        j                        s! |||       + yc c}}w c c}}w )z$Assign all symbolic shapes to localsc                P   t        j                  | j                         | j                         g      D ]k  }t	        |t
              rt	        |t        j                        r.|j                  D cg c]	  }||vs| }}t        |      dkD  s[t        d| d| d       y c c}w )Nr   zFor z, expected z to have been codegen-ed.)r   from_iterablerM  rN  rw   r
   r   r\  free_symbolsr   r  )r  ra  exprr_  undefined_symbolss        rc   _verify_input_symbol_assignmentzLPythonWrapperCodegen.codegen_inputs.<locals>._verify_input_symbol_assignment	  s     ++U^^-=u?O?O?Q,RS 
!$-D%,,1O $(#4#4%:8MC%! % ()A-(tfK0A/BB[\ 
%s   0	B#:B#N)r  ir.TensorBoxra  OrderedSet[sympy.Symbol])	r   r   r\  rJ  r  rw   rc  r'   r   )
r  ra  r  kr<  r\   r   r  rj  r  s
             rc   ri  z#PythonWrapperCodegen.codegen_inputst	  s    -/
 ,,.+113
q!z!U\\7RQF
 , 2 2 4X1Jq%,,<WaVXY " 	JKD%00ujI	J		0	&  	?HAueR\\2+E:>	?3
Xs   "D"D>"D!Dc                ~   t        |t        j                        rt        |t        j
                        r|| j                  v ry | j                  j                  |       t        j                  j                  j                  |   }t        ||      }| j                  t        | |t        j                               y y y re   )rw   r   r\  r"   r#   PRECOMPUTED_SIZErU  r   r@   rf   rn   inv_precomputed_replacementsr  r   r  )r  r_  rh  r  s       rc   ensure_size_computedz)PythonWrapperCodegen.ensure_size_computed	  s    c5<<(^CAVAV-Wd)))##C(77##@@ED!#t,CNN.tS!''BC .X(rb   c                     y re   ra   r  s    rc   r  z$PythonWrapperCodegen.finalize_prefix	  r-  rb   ro   c                   t        d      )Nz8codegen_cpp_sizevar is only implemented for cpp_wrapper!)r:  r  r   ro   s      rc   codegen_cpp_sizevarz(PythonWrapperCodegen.codegen_cpp_sizevar	  s    UVVrb   c                   t        ||      S )Nrs  )pexprru  s      rc   codegen_python_sizevarz+PythonWrapperCodegen.codegen_python_sizevar	  s    Q**rb   c                $    | j                  |      S re   )ry  rD  s     rc   codegen_sizevarz$PythonWrapperCodegen.codegen_sizevar	  s    **1--rb   c                    | d| dS )Nr}  r~  ra   )r  r  r   r6  s       rc   r  z)PythonWrapperCodegen.codegen_tuple_access	  s    1UG1%%rb   c                    g t        | j                  |      }t        |      dk(  ryt        |      dk(  r	d|d    dS ddj                  |       dS )Nr   ()rA   r  r  r  r   )r  ry  r   r   )r  rT  partss      rc   r   z/PythonWrapperCodegen.codegen_python_shape_tuple	  s^    :#d1159:u:?u:?uQxj$$499U#$A&&rb   c                $    | j                  |      S re   )r   )r  rT  s     rc   rQ  z(PythonWrapperCodegen.codegen_shape_tuple	  s    ..u55rb   c                    dj                  dj                  |t        |      t        |      | j	                  |      | j	                  |      g            g fS )Nzalloc_from_pool({})r  )r  r   rx  rW   r   )r  r   r   r   rT  r   s         rc   codegen_alloc_from_poolz,PythonWrapperCodegen.codegen_alloc_from_pool	  s_     %++II&MJ33E:33F;

 
 
	rb   c                ~    t        |      \  }}}	}
} fd}|j                         }|xr ||	k(  }|r||k(  xr ||k(  }|
}n[||j                  j                  k(  xr4 ||j                  j                  k(  xr ||j                  j
                  k(  }|j                  }|r|||k7  r	d| d| dS | S  |||||||      S )Nc           	         
j                  |      }
j                  |      }
j                  |      }d|  d| d| d| d	}	|||k7  r	d|	 d| dS |	S )Nzreinterpret_tensor(r  r   aten.view.dtype()r   r{  )r   tgt_size
tgt_stride
tgt_offset
cast_dtype
base_dtyperE  stoffrh  r  s             rc   apply_reinterpretzHPythonWrapperCodegen.codegen_reinterpret_view.<locals>.apply_reinterpret	  s|     //9A00<B&&z2C(b2bTC5BD%*
*B)$r*Q??Krb   r  r  r   )r   rh   r{   r   r   r   r   )r  r   r   r   r   r   r   d_sized_strided_offsetd_dtypecollapsibler  r   	collapsedsame_layoutr  s   `                rc   codegen_reinterpret_viewz-PythonWrapperCodegen.codegen_reinterpret_view	  s    8 ,D1 	9(G[		 }}6Fh$6	&.?Vx-?K J ((( 1dkk0001dkk000 
 J Uj%8)$r%::V tVVUJOOrb   c                8    | j                  | d| d| d       y )Nz.copy_(r  r   rt  )r  r   dstnon_blockings       rc   codegen_device_copyz(PythonWrapperCodegen.codegen_device_copy
  s!    #gcU"\N!<=rb   c                    d| }| j                  | d       | j                  | d       | j                  | d       y )N_d2h_event_z = torch.Event()z	.record()z.synchronize()rt  )r  rX   	event_vars      rc   sync_d2h_copyz"PythonWrapperCodegen.sync_d2h_copy
  sH    !+/	)$456)I./)N34rb   c                    |j                         }|j                  d      }| j                  t        | |||j                               y r#  )rh   
input_namer   rv  rz  )r  rp   rw  rx  s       rc   codegen_multi_outputz)PythonWrapperCodegen.codegen_multi_output
  s6    mmo??1%t[(DLLQRrb   c           
     j   |j                    d|j                   d|j                    d|j                    }|rd|j                   d| d}| j                  |j                   d|j                   d|j
                   d| d	       | j                  j                  t        |j                               y )
N +  if z
 < 0 else zmax(0, min(r  z))r   z * (r   )	r6  r   r   unbacked_offset_symbolbase_offsetbase_dim_strider  r   rW   )r  rp   clamp	index_strs       rc   codegen_dynamic_select_indexz1PythonWrapperCodegen.codegen_dynamic_select_index
  s    zzl#dii[TZZL
4::,W	%dii[9+R@I**+3t/?/?.@DDXDXCYY]^g]hhij	
 	""&&s4+F+F'GHrb   c                     fd} fd}j                   } |j                        } |j                        } j                  | d|         j                  | d|         || d| dj                        } j                  | d| d        j
                  j                  t        j                                y )	Nc           	        j                  t        dt        | j                                    }j                  t        dt        | j                  z   j                                    }j                  |       }| d| d| S )Nr   r  z >= 0 else )r{  r   r   r   )r   posnegx_condrp   r  s       rc   clamp_indexzDPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.clamp_index&
  sv    &&s1c!TYY.?'@AC&&s1c!dii-.K'LMC))!,FU$vhk#77rb   c                Z    |dk(  r| d|  S j                  |      }d| d|  d| d| S )NrA   z - r  r  z	 - 1) // )r{  )	start_varend_varstepstep_r  s       rc   codegen_with_stepzJPythonWrapperCodegen.codegen_dynamic_slice_size.<locals>.codegen_with_step,
  sJ    qy!#i[11((.Ewis9+SyHHrb   z	_start = z_end = _start_endz
 = max(0, r   )unbacked_size_symbolstartendr   r  r  r   rW   )r  rp   r  r  r_  r  r  	with_steps   ``      rc   codegen_dynamic_slice_sizez/PythonWrapperCodegen.codegen_dynamic_slice_size%
  s    	8	I ''DJJ'$((##iw/0#gcU+,%VnTlDIIN	#j156""&&s4+D+D'EFrb   c                :    | j                  t        | |             y re   )r   rq  r#  s     rc   codegen_dynamic_scalarz+PythonWrapperCodegen.codegen_dynamic_scalar<
  s    (t45rb   c                   d |j                   D        \  }t        |j                        dk(  r#| j                  |j                   d| d       nkt        |j                        dk(  r@t        |j                  d   t              r#| j                  |j                   d| d       nt        |j                        dk(  rt        |j                  d   t              r| j                  |j                   d| d       | j                  d	|j                   d
|j                  d   j                   d|j                   d|j                  d   j                   d	       | j                  |j                   d|j                   d|j                  d   j                          nt        d|j                         | j                  |j                          d       y )Nc              3  <   K   | ]  }|j                           y wre   r  r  s     rc   r   z?PythonWrapperCodegen._codegen_dynamic_scalar.<locals>.<genexpr>@
  s     >Q1&&(>r  r   r   .item()rA   z = 1 if z.item() else 0z_undivided = zassert z_undivided % z
 == 0, f'{z_undivided} not divisible by 'z_undivided // unrecognized keypath  = None)r\   r   keypathr   r_  rw   r   r   r  r  rh   )r  rp   r   s      rc   rs  z,PythonWrapperCodegen._codegen_dynamic_scalar?
  s   >$++>t||!NNdhhZs4&89!#
4<<?M(RNNdhhZxv^DE!#
4<<?K(PNNdhhZ}TF'BCNN$((=a1H1H0I Jxxj >t||A?V?V>WWXZ NN88*CzQ8O8O7PQ !#8!GHH 	$--/*'23rb   c                     fd}fd}fd}i }i }t         j                  j                  }|t        t         j                  j                  j                         |      D ]  \  }}	t        |	t        j                        s!|	j                         }
|
j                         }t        |
      }|j                  |      }|5t        dt        |       ||	j                  |	j                   i       }|||<   t#        |	j$                        t#        |	j'                               f|j(                  |<    j+                  g d       j-                         5  j/                  dd	
       t         j                  j0                  j3                         D ]U  \  }}j5                  d|         |||j7                         |j'                         |j                  |j                          W t        t         j                  j8                        dkD  r^j5                  d       t         j                  j8                  j3                         D ]"  \  }}j5                  d|         |||       $ |j;                         D ]  }t        |j(                        dk  r|j                  t        j<                  j?                  |j                         z  }j5                  |j@                   d| d|j                   d|j                    d       |j(                  j3                         D ]  \  }\  }}|j@                  ||f||<     t         j                  j                  j3                         D ]U  \  }}t        |tB        jD                        rCt        t         j                  jF                  jH                  j                  |d      tJ              rdt        |tL        jN                        rj5                  | d       t        |tB        jP                        r3 ||t         j                  jF                  jS                  |d             t        |tB        jT                        r ||d       t        |tL        jV                        r# ||d|j                  jX                   d       Bt        |tL        jZ                        rj5                  | d       r||v rF||   \  }}}j5                  | d| d j]                  |       d j]                  |       d       t         j                  jF                  j_                  |ja                         d      }t         j                  jF                  j_                  |jc                         d      } |||||je                         |jg                                X j5                  ddji                  t         j                  j                  j                                d       ddd       j+                  g d       j-                         5  j/                  dt         j                  jj                   dd	
       ddd       y# 1 sw Y   exY w# 1 sw Y   yxY w) z2Write out codegen for benchmarking the output codec                    j                  |  dj                  |       dj                  |       d| d| d
       y )Nz = rand_strided(r  
, device='	', dtype=r   )r   r   )r   rT  r   r[   r   r   r  s        rc   add_fake_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_fake_inputW
  sT    &(2259:"226:; <!()E7!5rb   c                2    j                  |  d|        y r  rt  )r   r  r   s     rc   add_expr_inputzFPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_expr_input_
  s    vS./rb   c                2   |j                  |  d       y dd l}	 j                  |  d|j                  |      d       y # t        t        |j
                  f$ r7}j                  dt        |       d|  dt        |       d       Y d }~y d }~ww xY w)	Nr  r   z = pickle.loads(r   z.raise TypeError("Failed to pickle opaque type z for variable r   z"))r   pickledumps	TypeErrorr  PicklingErrorr   rW   )r   r  r  r?  r   s       rc   add_torchbind_inputzKPythonWrapperCodegen.benchmark_compiled_module.<locals>.add_torchbind_inputb
  s    }  D6!12  D6)9&,,u:M9PPQ!RS~v/C/CD   DT%[MQ_`d_eeghklmhngooqr s   &A B-BBN_shared_storage_)rX   rZ   r[   r   r\   )r  r  zdef get_args():zP
                from torch._dynamo.testing import rand_strided
                Tr   zglobal r   zimport pickler$   z = rand_strided((z,), (1,), device='r  r   r  *   r   Fztorch.cuda.default_generators[z].graphsafe_get_state()z = torch.as_strided(r  zreturn [r~  )r  r  z9def benchmark_compiled_module(args, times=10, repeat=10):z
                from torch._inductor.utils import print_performance
                fn = lambda: call(list(args))
                return print_performance(fn, times=times, repeat=repeat, device='z')
                )6r@   rf   example_inputsr   r  r  rw   r  Tensoruntyped_storagerZ   r   r   rV   r   r[   r   r(  rT  r   r\   
writelinesr   r   	constantsr  r   r   torchbind_constantsr  _utils_element_sizerX   r   r\  rn   backed_var_to_valr    r'   rN  r
   r"  rM  rO  r6  rP  r   r  rM  rN  r  rm   r   r  )r  r   r  r  r  storage_groupsaliased_input_specsr  r   exstoragerZ   storage_keygroupr  torchbind_objnumelrT  r   r  s   ``                  rc   benchmark_compiled_modulez.PythonWrapperCodegen.benchmark_compiled_moduleT
  s   		0	  GI?A//% 4 4 9 9 ;^L Ib!"ell3,,. ),W5&**;7=1&6s>7J6K$L%!yy hh!E 38N;/&*288nd299;6G%HT"#I( 	56]]_ X	SMM 	    !ww00668 e   74&!12%**,ekk	 177../!3  1+,77+F+F+L+L+N ='D- $$wtf%56'm<	= (..0 Su||$q((B(B5;;(OO  (()):5'ASTYT`T`Saajkpkvkvjwwxy .3\\-?-?-A S)D/5&161B1BE60R'-SS  !ww3399; 3eeU\\2zGG$$66::5$G8 eR%7%78$$vW%56uzz2
 #agg..@@QS@T  u{{3 #4/r'8'89"89K9K8LLcd  r';';<$$vW%5600.A$.G+HeV$$& 4XJb::5AB"::6BC1F GG,,??(2 @ E WW--@@((*R A F #((*)[3j x		!''2F2F2K2K2M(N'OqQRqX	Sv 	Q	
 ]]_ 	MMR STRYRYReReQf g
   	 	}X	S X	S|	 	s   R9Y7>0Z7Z Zc                    t         j                  sy| j                  |       |j                  g d       |j	                         5  |j                  dddt                dg       ddd       y# 1 sw Y   yxY w)zL
        Append a benchmark harness to generated code for debugging
        N)r  r  zif __name__ == "__main__":zBfrom torch._inductor.wrapper_benchmark import compiled_module_mainzargs = get_args()zcompiled_module_main('zU', lambda times, repeat: benchmark_compiled_module(args, times=times, repeat=repeat)))r&   benchmark_harnessr  r  r   r5   r  r   s     rc   r  z*PythonWrapperCodegen.add_benchmark_harness
  s}     ''&&v.@A]]_ 
	X'01C1E0F Gm m		
	 
	 
	s   !A00A9c           
     D    | j                  t        | |||||             y r  )r   r  )r  r  r  r  r  r  s         rc   define_kernelz"PythonWrapperCodegen.define_kernel
  s*     	 !-		
rb   c                    t         j                  j                  r)|r't        j                  dd|t        j
                        }|r| dnd}d| |  d| }|S )Nz^// z# )flagsr  r  z

r   )r&   r   r   resub	MULTILINE)r  r  r  metadata_commentbodys        rc   _format_kernel_definitionz.PythonWrapperCodegen._format_kernel_definition  sZ     ==11h vvgtXR\\JH.6hZr?B&'}C}Erb   c                .   t         j                  j                  rL|rJ| j                  |||      }| j                  j                  |       t        j                  j                  ry | j                  |||      }| j                  j                  |       y )N)r  )
r&   r   r   r  r  r   r@   rf   r  r  )r  r  r  r  r  r  r  s          rc   r  z*PythonWrapperCodegen._define_kernel_helper)  s     ==11c11[8 2 D %%,,T2ww""--x . 
 	4 rb   c                N    | j                   j                  |j                         y re   )r  r   r  )r  r   subgraph_codes      rc   define_subgraph_launcher_fnz0PythonWrapperCodegen.define_subgraph_launcher_fn?  s    !!(()<)<=rb   c                    ddl m} |S )NrA   )TritonKernel)r   r  )clsr  s     rc   _get_triton_info_kernel_clsz0PythonWrapperCodegen._get_triton_info_kernel_clsB  s     	)rb   c                  123456 ddl m}m}	m}
 ddlm1m}m}m}m	} |j                  }g 6i 4g 3g }36fd2d5124fd	}|j                  D cg c]  }|j                   }}|j                  D cg c]  }|j                  s|j                   }}t        |      D ]  \  }}||v r || 1|      d	       |vr#|   }|    || 1|      d       @t!        |t"        j$                        r[t!        |t"        j&                        r'd|j(                  |j*                  j-                         fnd\  }}} || |||||             t!        |t"        j.                        r/ || |||j1                         |j-                                      t!        |t"        j2                        rO || |||j4                  j1                         |j-                         |j6                  j8                               gt!        |t:        t<        j>                  f      xr* t@        jB                  jD                  jG                  |d      } || |||      |        tI        6d
3|jJ                  D cg c]  }tM        |       c}      }|tO        jP                  t@        jB                  jS                               i 4tT        jW                  |d      tY        63d      gd}|rt[        |      |d<   |rt[        |      |d<   t]        |      dk(  r0|	j_                         }g ta        t<        jb                  |d         } nd65fd}!i 5|D "cg c]  }"g ta        |!|"       }}"|rt]        |      t]        |      k(  sJ g }#te        tg        ||      d d      D ]Q  \  }"}$|#ji                   ||$      g ta        tj        |"      g ta        tl        |"      g ta        tj        |"      d       S |
j                  |#g ta        tn        5jq                               d}g 5js                         } |r|D %cg c]  }%||%   	 c}%|d<   tu        |jv                        g}&t]        |      dkD  rQjq                         D ]>  }t!        |t"        j.                  t"        j2                  f      r.|&ji                  |       @ |&ji                  to        |             |&jy                  to        |             |'|&ji                  |d   j1                         |d   f       t[        |&      }&|&| jz                  v r| jz                  |&   \  }'}}(|'||(| fS | dt]        | jz                         }'|'j}                  d       r|'j                  d       s|'dd
 }'t               })t        j                  j                  r|)j                  d!|'d"       n|)j                  d!|d"       |'|d#<   | j                         }*|j                  |*j                                |)j                  |*j                                t        j                  j                  r|)j                  d$       |)j                  d%g ta        ||      d&|d'|d(       t        ||      }+t        j                  j                  r|+j                  d)| d*d)|' d*      }+|+j                  d+d,      }+t        j                  r|+j                  d-d.      }+|+j                  d/d0      }+|)j                  |+       t@        jB                  jS                         },|)j                  d1|,j                   d2       t        j                  |jv                        \  }-}.t        j                  |jv                        }/d3|/ d4|. }0| j                  |'|)j                         |0       |'||f| jz                  |&<   |'||| fS c c}w c c}w c c}w c c}"w c c}%w )7a  Codegen a user-defined Triton kernel and return its cache entry.

        Emits the ``async_compile.triton(...)`` wrapper, assigns a graph-unique
        name (with a leading dunder stripped to avoid Python class-based name
        mangling at the call site), and records the kernel in
        ``user_defined_kernel_cache``. Returns ``(name, triton_meta,
        inductor_meta, extra_launcher_call_args)``; subsequent calls with the
        same ``cache_key`` reuse the previously assigned name.
        r$   )config_to_dict	FixedGridPrecomputedGridrA   )ConstexprArgKernelArgTypeSizeArg	TensorArgTMADescriptorArgc                J    j                  |       j                  |        y re   )r   )r  r  arg_indices	signatures     rc   add_to_signaturezPPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_to_signaturer  s    S!s#rb   c                   |r?t               r	 | |       |j                  v r|j                     |j                  <   y y |j                  v sJ |r>t               r |  |j                               n	 | |       d|j                  <   y |r4t               r |  |j                               d |j                  <   y  | |       y )Nr   rA   )r?   r   )	r  r  is_constexprequals_1equals_noner   r  r  r   s	        rc   add_argzGPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.add_argv  s    13 %S#.88v% +1*:Ichh' & xx6)))57
 )l.IJ(c2*+Ichh' 57 )l.IJ*.Ichh'$S#.rb   r
  T)r  N)r  stable)experimentalNN)r   api_typer  r   )r   bufferr   )r   r  r   r   )r  )
size_dtyperz  argdefsra   )rz  pointer_range_override)r  r[   r  r   restore_valuereset_to_zeror   c                N   t        | t        j                        rdg | j                  }|s| S |j	                  t
               |D ]+  }|v rt        j                  dt                     |<   - t        |       S t        | t              sJ t        j                  |       S )N)r   _launcher_s)rw   r   r
   rg  sortrW   r\  r   r>   rY   r   )rh  symbolsr_  extra_launcher_argss      rc   rename_sizes_for_launcherzYPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.rename_sizes_for_launcher   s    dEJJ/2 1 12G"#LLSL)& "55$38<<)#.A*B)CD4+C0 &d,?@@!$,,,}}T**rb   c                2    t        | d   j                        S r   r   r   s    rc   r   zHPythonWrapperCodegen.define_user_defined_triton_kernel.<locals>.<lambda>  s    3qt{{3C rb   r   )r&   r  cpppython_slow)	grid_typeprecomputed_gridsr  declared_constexpr_namesr  __zasync_compile.triton(z, '''r  r  zG
            @triton_heuristics.user_autotune(
                configs=z ,
                inductor_meta=z,
                triton_meta=z{,
                filename=__file__,
                custom_kernel=True,
            )
            @triton.jit
            r   r  \z\\z"""z\"\"\"z'''z\'\'\'z''', device_str='re  z# Original path: r  )FFF)rh  r   r   r   )Truntime.triton_heuristicsr  r  r  commonr   r  r  r  r  r]   paramsr   r  numr  rw   r'   TMADescriptorr  r  r  rm   r   rh   r*   r   r{   r   rY   r   r   r@   rf   rn   statically_known_equalsrL   	arg_namesrB   r,   r  get_current_device_or_throwr  fromkeysrJ   r   r   setup_grid_as_argsr  sympifyr   r   r   rx  rH   rW   r  r  idr   extendr  r  endswithr7   r&   r   unique_user_kernel_namesr   r  updateinductor_meta_commonr   gen_common_triton_importsr  r  replacer  r   inspectgetsourcelinesgetsourcefiler  r   )7r  r  r   r   restore_value_argsreset_to_zero_argsr   r  r  r  r  r  r  r  r  original_nameequal_to_1_argsr  pr,  
constexprsr  r   r  r  r  r   r  r   triton_signaturer  r  extra_launcher_call_argsr  r   r"  cfgr   	cache_keyr   cached_inductor_metar  triton_info_kernel_clsr  current_devicer  linenosrcfiler  r   r  r  r  r  r  s7      `                                             @@@@@@rc   !define_user_defined_triton_kernelz6PythonWrapperCodegen.define_user_defined_triton_kernelL  s   &	
 	

	
 	
 )+	$&	!#%'	$"	/ "	/H &,]]3QVV3	3%+]]EannaeeE
E!), 9	GHCj \s3$G& +Cc{"\s3Fc2#3#34 &c2+A+AB "3??CJJ4H4H4JK9 1Hk5
 (!$%-(3"'	  RYY/!!$#&<<>"%--/  R%7%78 !!$#&88#4#4#6"%--/#&::#4#4	  *c5==1   ''**BB  Cc!2XFs9	Gv -)/)9)9:AWQZ:	
 *&--agg.Q.Q.ST--3
 '+-'
. +01C+DK(+01C+DK(u:?,5,H,H,JM'FU]]E!H)E'F$+  EGINO<s4d;<OEOSZ3w<777 "#E7#)CT 
	c "(("0"5"5Ct$4"52UD!12':UD)9':	
 -55%6'PS2E2L2L2N)O'PM
 (E)<)A)A)C'D$&09!"	!9M45
 VYY-	w<!}} *!#		23E3E'FG$$S)* 	[)*]+,&oa099;_Q=OPQ)$	6666:6T6T73D+3 $(	   #d&D&D"E!FG ??4 t)<8D(*==11%%(=dXU&KL%%(=m=Ne&TU'+m$!%!A!A!C3HHJK5OOQR==))%%&DE83~w78; <,/ 0(O ,			
 OO

 ==11#++d=/,CtD6QR^TJ''f5
 $++E;?J''{;
z*<<>!!$5n6I6I5J""MN**6995	6''		2&wiq9$$&	
 6:;4V&&y1[-1IIIO 4E@ ;l P.9s$   _&4_+_+?_0_5_:c                    | d|j                    d}||d| z  }t        j                  |dd      }t        ||j                        }|dk(  }|s*| j                  t        | |t        j                               |S )Nr  r  T)
is_integeris_positiver  )	r  r   r\  r  r  r   r  r@   rf   )r  r  treer  sym_namer_  r  is_benchmark_kernels           rc   generate_numel_exprz(PythonWrapperCodegen.generate_numel_expr  s    !]!DKK=6!F8$Hll8$G c4::.)R/"NN.tS!''BC
rb   c                j    | j                  |j                   dt        |j                                y r  )r   r  rx  r  )r  r  rf   s      rc   r  z7PythonWrapperCodegen._generate_symbolic_call_arg_helper  s)     	#))Ccnn(='>?@rb   c                   |j                         }t        | |      }|j                  t        j                  k(  r| j                  |       n1|j                  t        j                  k(  r2| j                  |       | j                  | j                  |             n|j                  t        j                  k(  r| j                  j                  |      }|rRt        |t              rt        |j                  t              sJ t        j                  |j                  |      |_        nV| j                  |       | j                  | j                  |             || j                  |<   nt        |j                        t         j"                  j$                  r| j&                  j                  t(        j+                  | ||j,                  |j.                  t0        j2                  j4                  j7                  |j8                        fd             |j                  t        j                  k7  r0| j&                  j                  t(        j                  | |             y y y )N)rA   )rT  r   )rh   r0  	zero_moderG   UNINITIALIZEDr   ZERO_ON_CALLmake_zero_bufferZERO_PER_GRAPHr  r   rw   rp   rF   maximumr  r&   r   r   r   rR  make_allocationr[   r   r@   rf   rn   r"  r   )r  wsr   r   priors        rc   generate_workspace_allocationz2PythonWrapperCodegen.generate_workspace_allocation  s   {{}D"%<<,:::NN4 \\.;;;NN4 NN40067\\.===--11$7E!%6:JJ<   *11%**bA
t$t44T:;26))$/ ..==11&&00$44IIHH77++==bhhGI 5 	 ||0>>>**44(99$E ? 2rb   c                v    |j                   t        j                  k7  r| j                  t	        | |             y y re   )rU  rG   rY  r   r=  )r  r\  s     rc   generate_workspace_deallocationz4PythonWrapperCodegen.generate_workspace_deallocation  s.    <<,;;;NN.tR89 <rb   c                $    | d| j                    S )Nz.zero_())r  )r  r   s     rc   rX  z%PythonWrapperCodegen.make_zero_buffer  s    x}--rb   c                H    | ddj                  |       d| j                   S )Nr  r  r   )r   r  )r  r   r  s      rc   r  z%PythonWrapperCodegen.wrap_kernel_call  s'    q9-.a}==rb   c                    | j                   j                  d       | j                   j                  dt        j                  j                   d       |j                  | j                   j                                y )Nz*from torch.profiler import record_functionzwith record_function('graph_z_inductor_wrapper_call'):)r  r   r@   rf   graph_idr  r   )r  r   s     rc   r  z8PythonWrapperCodegen.generate_profiler_mark_wrapper_call  sb    ##$PQ##*177+;+;*<<UV	
 	D--4467rb   c                :    | j                   j                  d       y )Nzstart_graph())r  r   r  s    rc   r  z)PythonWrapperCodegen.generate_start_graph  s    ##O4rb   c                ^    | j                   j                  dt        j                  d       y )Nz
end_graph(r   )r  r   r&   profile_bandwidth_outputr  s    rc   r  z'PythonWrapperCodegen.generate_end_graph  s'    ##j1P1P0SST$UVrb   c                    | j                   j                  t        j                  j                  j                                y)z<Synchronize GPU to ensure proton captures all kernel events.N)r  r   r@   rf   r  rg  r  s    rc   r  z-PythonWrapperCodegen.generate_proton_finalize  s)    ##AGG$6$6$B$B$DErb   c                ^    | j                   j                  dt        j                   d       y )NU
            for kernel in globals().values():
                if isinstance(kernel, z.CachingAutotuner):
                    kernel.cuda_kernel_saved = False
                    kernel.cpu_kernel_saved = False
            r  r   r+   r]   r  s    rc   r  z6PythonWrapperCodegen.generate_reset_kernel_saved_flags  s2      ''8'A'A&B C	
rb   c                ^    | j                   j                  dt        j                   d       y)a  
        Precompile and save the per-config kernel artifacts (CUBINs on GPU,
        ``.so`` + launcher ``.so`` on CPU) for Triton kernels that haven't
        been precompiled and saved as a side effect of running the generated
        JIT model (Python wrapper). This can happen when the model contains
        control flow: only one pass through the control flow operators covers
        the kernels that are saved, the remaining kernels are not launched,
        hence not saved. The main purpose of this codegen is to compile and
        save the Triton kernels outside the active control flow path for
        subsequent AOTInductor code generation and compilation.
        rj  a  .CachingAutotuner):
                    if kernel.device_props.type == "cpu":
                        if not kernel.cpu_kernel_saved:
                            if len(kernel.launchers) == 0:
                                kernel.precompile()
                            kernel.save_cpu_kernel(launcher=kernel.launchers[0])
                    else:
                        if not kernel.cuda_kernel_saved:
                            if len(kernel.launchers) == 0:
                                kernel.precompile()
                            kernel.save_gpu_kernel(
                                stream="stream",  # use dummy stream
                                launcher=kernel.launchers[0],
                            )
            Nrk  r  s    rc   r  z5PythonWrapperCodegen.generate_save_uncompiled_kernels  s4     	  ''8'A'A&B C	
rb   c                >    d }|D cg c]
  } ||       c}S c c}w )Nc                    t        | t              rt        |       r| dz   S | S t        | t        t        t
        t        f      rt        |       S t        t        j                  j                  j                  |             S )Nr  )rw   rW   rK   rY   floatr  r  rx  r@   rf   rn   ro   )r  s    rc   wrap_argzAPythonWrapperCodegen.prepare_triton_kernel_call.<locals>.wrap_arg
  s^    #s#*B3*GsYPSPC#udO!DE3xQWW--66s;<<rb   ra   )r  r  rp  r  s       rc   prepare_triton_kernel_callz/PythonWrapperCodegen.prepare_triton_kernel_call	  s!    	= *33#333s   c                p    t        |t              rKt        |t        j                        r.|j	                         j                         } j                  |   }n\ j                  j                  |      r|} j                  |   }n/|J d       d j                   }|} xj                  dz  c_        |
J d|        t        j                  j                  j                  |j                               }t        j                  j                  j                  t        j                  j                  |            }t        j                  j                  j                  |j                               }|j!                         }	|j#                         }
t        j                  j                  j%                  |j'                         j(                        }d| d| d|	 d|
 d| d| d	} j*                  j-                  | d
|        t        |t        j                        r5 j/                  |d      }|} j*                  j-                  | d
|        |S t1        |t2        j4                        st        |t6              rt        |t8              r| j:                  v r|S |y|}t        |t6              r|j<                  }|t        j                  j                  j>                  v r't        j                  j                  j>                  |   }t9        t        j                  j                  j%                  |            S t        |t8        t@        tB        tD        f      rt9        |      S t        |tF              rddjI                   fd|D               dS tK        dtM        |             )NzBV.graph.get_buffer(arg) and raw_arg can't be None at the same timetmp_arg_rA   z Failed to find a buffer for arg zgenerate_example_value(r  z, 'z', r   r   T)r  r  r>  r}  c              3  T   K   | ]  }j                  |t        |             ! y wre   r   )r   ar  s     rc   r   zBPythonWrapperCodegen.generate_example_arg_value.<locals>.<genexpr>P  s#      ZQR!@!@DG!L Zr   r~  zUnsupported type )'rw   torch_dtyper'   r*  
get_tensorrh   r  r   r  r@   rf   rn   r  rM  get_allocation_sizerN  r  rm   r"  r   r   r   r   r  r  r   rM  r  rW   r  r  rp  rY   ro  r  r(  r   rI  r   )r  r  arg_typeraw_argr  rR  r   allocation_sizer   r[   r   r   r  s   `            rc   r   z/PythonWrapperCodegen.generate_example_arg_value  s   h,'2#3#34"--/88:**3/%%))#.**3/* X* &d&F&F%GH00A50?L&Fse$LL?77##66s||~FDgg..AA++C0O WW%%889IJF^^%FMMOEWW%%778H8O8OPF-dV2fXSE7RTU[T\\^_n^oopqE&&00H:S1HI'2#3#34 :: %) ;  **44zUG5LMO%++.*S/2R#s#$//)J?!#/nnagg&&CCCgg&&CCCHqww''99#>??c3t45s8OT"tyy ZVY ZZ[[\]]%(9$s)&EFFrb   c                z     t        |t              r ddj                   fd|D              z   dz   S t        |      S )Nr}  r  c              3  @   K   | ]  }j                  |        y wre   )_grid_dim_str)r   r   r  s     rc   r   z5PythonWrapperCodegen._grid_dim_str.<locals>.<genexpr>W  s     RT 2 24 8Rs   r~  )rw   r(  r   rx  )r  grid_per_dims   ` rc   r~  z"PythonWrapperCodegen._grid_dim_strT  s<    lD)diiR\RRRUXX &&rb   )r[   r   r  r  r  r  r  r   c                  | j                   j                  |D ci c]2  }t        |t              r |t        j
                  j                  |      4 c}       |xs t        j
                  j                         }t        j
                  j                  j                  }| j                  t        | ||||||||	|t        j
                  j                  |
|             yc c}w )z
        Generates kernel call code.

        triton: Defines whether the backend uses Triton for codegen. Otherwise it uses the CUDA language when gpu=True,
                and C++ when gpu=False.
        )r  r  r  r  r  r   r  r  r[   r  r   r  N)r  r5  rw   rW   r@   rf   try_get_bufferr-  rj   r  r   r  r   )r  r  r  r[   r   r  r  r  r  r  r   r  r  s                rc   generate_kernel_callz)PythonWrapperCodegen.generate_kernel_call\  s    . 	## %c3' QWW++C00	
 @177>>@WW..AA'#!!#'+77<<%9#5%	
s   7Cr  )
r[   r   r  r  r  r  r  r  r   r  c       
   
         |xs t         j                  j                         }|s|j                  dvr{|j                  dk(  r" j	                   j                  |             y |j                  dk(  r" j	                   j                  |             y t        d|j                   d       j                  |      }dj                  |      }|=|t        k7  r4 j                          d} j	                  | d|j                   d	       n!t        j                   |j                  |
      }|s$d
| d	} j	                   d d| d| d	       y  j                          t        j                   j"                  r j$                  vr|t'        |      t'        |      k(  sJ d       d |rDt         j                  j(                  r*t         j                  j(                  j+                  |d       t        j,                  j.                  }d fd} fd}g }g }|(|J d       d gt'        |      z  }d gt'        |      z  }nt'        |      t'        |      k(  sJ d       i }t1        t3        ||||            D ]  \  }\  }}}}d }t5        |t6              r!dt7        |      v r|j9                  d      \  }}d }r|v r j;                  |         }|rB|}t5        |t<              st?        |t@        jB                        st5        |tD              r|||<   n|dk(  r |||||      r||   }nt5        |t<              rtG        jH                  d|      r|}n[|r% jK                  |||      }|jM                  |       n4| jN                  vr jK                  |||      }n jN                  |   d   }|s%|f jN                  |<   n jK                  |||      }t5        |t6              rtQ        |      r|dz  }|jM                  ||n| d|          jR                  j	                  dt         j                  jT                  jW                  |j                         d        jR                  jY                           jR                  j	                   ddj                  |       d| d	        jR                  j[                          |r0|r. jR                  j	                  ddj                  |              n(|s& jR                  j	                  t]        d|d              j$                  j_                         t         j                  j`                  ry t         j                  jb                  jd                  }|jg                  ||d        |5   j	                   d| d| d	       d d d         j                          y # 1 sw Y   xY w)N)cudaxpur@  mpszdevice z nyir  
raw_streamr  r   z	c_void_p(r   r  z$call_args and arg_types do not matchc                     j                   j                         D  cg c]  \  } }|k(  r|  }} }|rddj                  |       dS yc c}} w )z;Returns del for tensors whose last consumer is this kernel.del r  r  r  )r  r  r   )r  kntensors_to_deleter  r  s      rc   get_autotune_deletion_callzUPythonWrapperCodegen._generate_kernel_call_helper.<locals>.get_autotune_deletion_call  se     '+&G&G&N&N&P%"[( %! %
 %!$)),=">!?rBB%s   Ac                j   ||   }||v ryt        t        | |            D ]  \  }\  }}||k(  st        |t              sd}r|v rj	                  |         }|dk(  rA	 |j                         }	t        |	j                        D ]  \  }
}||k(  s| d|
 d||<     y  y# t        $ r Y w xY w)zWe try to infer raw_arg (i.e. raw_args[idx]) from remaining raw_args.
                This is particularly useful for jagged cases, where the dimension is often
                being passed in as an input.Tr  z.shape[r~  F)r  r   rw   r)   r/  r   r   rI  )r  r  r  reused_args
target_argr   raw_keyrz  triton_inputr{   r  rE  autotune_argsr  s               rc   infer_arg_by_inputszNPythonWrapperCodegen._generate_kernel_call_helper.<locals>.infer_arg_by_inputs  s    
 &c]
,-6s8X7N-O !)A)Cxz'6'B #%L$M)A'+'E'E)'2( $r) 	!!(!3!3!5&/&< ,FC J=IN'RUQVVW:XJ 7'+,!, 	 / ! !!s   #0B&B&"B&&	B21B2zkeys are not None but args arez#call_args and raw_args do not matchr  r  z^(workspace|semaphore)r   r  r  r  z.run(z	, stream=r  z
<del_call>r  )4r@   rf   r-  r   r   r  r:  rq  r   r.   r:  r6  rR  r  r6  r&   r   r   r   r   autotuning_mappingr   r  autotune_per_kernel_allocr  r   rw   rW   splitr/  rv  r  r   rM  r  r  matchr   r   r  rK   r   r  r  rY  r|  r4   r   r  r  r  r  ) r  r  r  r[   r   r  r  r  r  r  r  r   r  call_args_strstream_name
stream_ptr_per_kernelr  r  all_argstensor_arg_strsr  r   r  ry  r  rz  r   r  arg_strr  r  s    ``                             @rc   r  z1PythonWrapperCodegen._generate_kernel_call_helper  s     @177>>@&++_<{{e#t44[)LM  %t44[)LM  #WV[[M#>?? 77	B		-0).@DV.V ,,.&KNNk]*<V\\N!LM.CCfllJK $[M3JNN-qQ}oR
|1M %%' MM224#=#== (S^s9~-M 6M !M#(B(B ! : : > >($! !--GGK	B H O'I)II' 6C	N2 6C	N28}I6 96 K8AIy(H=9 3P44C7G c3'C3s8O"yy~HC+/ W%=#'#A#A%g.$L  *G%h<"8U[[9%c?;/6G,]':h;(
 *'2G+6 xx 93?"%$"&"A"A7# (..w7D$E$EE"&"A"A7# #'"C"CC"H"K&BI;AW99#>"==c8WUGc3',DS,Iy(G3;se1WI<NOg3Pl &&00**77EFaH &&002&&00-uTYYx%8$9;-qQ &&224**44499_567 !**44$$2$ &&**;7ww"" !" 4 4 B B..y+yRVW" 	XNNk]%i}TUVW	X%%'	X 	Xs   %WW"c                :    | j                   j                  |       y re   )r  r   rE  s     rc   r   zPythonWrapperCodegen.writelinef  s    

$rb   c                4    |D ]  }| j                  |        y re   rt  )r  r  r   s      rc   r  zPythonWrapperCodegen.writelinesi  s     	!DNN4 	!rb   c                L    | j                   j                  t        |             y re   )r  r   r:   )r  ctxs     rc   r  z"PythonWrapperCodegen.enter_contextm  s    

+c*+rb   c                   	 ddl m}  |       rdd l}t        |t              rt        |j                  j                        S t        |t        j                        rt        |      S t        |t        t        f      rAt        j                   G d d             	t         t        |      	 fd|D                    S t        |t         j"                  j$                        rt'        |      S t        |t(        j*                  t(        j,                  t.        f      r|j1                         S  |       r+t        |j2                  j4                        rt        |      S t        |t(        j6                  t(        j8                  f      r|j1                         S t;        t        |            rEt=        |      \  }}|j?                         D ]"  \  }}|t@        jB                  jD                  |<   $ |S t        |      S )Nr   )has_triton_packagec                      e Zd ZU ded<   d Zy)1PythonWrapperCodegen.val_to_arg_str.<locals>.Shimr   refc                    | j                   S re   )r  r  s    rc   __repr__z:PythonWrapperCodegen.val_to_arg_str.<locals>.Shim.__repr__  s    88Orb   N)r]   r^   r_   r`   r  ra   rb   rc   Shimr  |  s    $rb   r  c              3  V   K   | ]   } t         j                  |             " y wre   r  )r   ru  r  r  s     rc   r   z6PythonWrapperCodegen.val_to_arg_str.<locals>.<genexpr>  s$     Vq1@@qIJVs   &))#torch.utils._tritonr  r   rw   r   rx  rp   rh  r   r
   r   r(  r  	dataclassr?  r   r  _ops
OpOverloadr   r'   r   
MutableBoxr*   r  languager   rO  rP  r   r   r  r@   rf   opaque_value_type_classes)
r  rE  type_r  r   obj_repropaque_typesnr  r  s
   `        @rc   r  z#PythonWrapperCodegen.val_to_arg_strp  s   :a"%%5::&8OE4=)""$ $ #$ QVTUVV  5::001&q))BIIr}}oFG&&((!jFOO4I4I&J7NB--r/C/CDE&&((!$q'*%8%;"Hl$**, 917811!49O7Nrb   c           	     `   |j                         }|j                         }t        |j                               }t        t        j
                  j                  |            }t        |j                               }|j                         }| j                  |j                         ||||||      S re   )r  rm   r   rM  r@   rf   rx  rN  get_is_pinnedr[  rh   )r  r  r[   r   rT  allocation_shaper   	is_pinneds           rc   rK  z+PythonWrapperCodegen.make_buffer_allocation  s    ""$  "foo'( !<!<V!DEv((*+((*	##OOvueV=My
 	
rb   c                x    d}t         j                  j                  s| j                  j	                  |d       y y )Nzi
            from torch._inductor.runtime.debug_utils import check_memory_step, track_tensor
            Tr   )r@   rf   r  r  r   r4  s     rc   "write_memory_track_allocation_oncez7PythonWrapperCodegen.write_memory_track_allocation_once  s4    
 ww""LL
$7 #rb   c                   ||}| j                  |      }| j                  |      }	| j                  |      }
t        j                  j                  j                  j
                  r| d|	 d|
 d| d|j                   d| d}ne|j                  dk(  r|r| d|	 d|
 d| d	}nE|j                  d
v r| d|j                   d|	 d|
 d| d	
}n| d|	 d|
 d|j                   d| d	
}||	k7  r|d| d|
 d	z   }|S )Nz = tracked_empty_strided(r  z, dtype=r  z	', name='re  r@  z = empty_strided_cpu_pinned(r   )r@  r  r  mtiaz = empty_strided_r  z = empty_strided(r  z.as_strided()r   r  r  r&   r  r  r   )r  r   r[   r   rT  r   r  r  rQ  codegen_allocation_shape_tuplecodegen_stride_tupler  s               rc   r[  z$PythonWrapperCodegen.make_allocation  s    #$"==eD)-)H)H*
&  $>>vF??!!..EE&112"'( )  !;;- (b"  [[E!i&412"'('  [[:: &)&++a12"'('  &)12"'( )!;;-yq:  "@@,':&;2>R=SSTUUC
rb   c                8    | j                  t        |             y re   )r   rl  rE  s     rc   make_commentz!PythonWrapperCodegen.make_comment  s    {4()rb   c           	     `    | j                    | d| | j                   d| j                   d| 	S )Nr      )r  r  r  )r  new_nameold_namer  s       rc   make_tensor_aliasz&PythonWrapperCodegen.make_tensor_alias  s6    ,,zXJt{{m2dll^STU\T]^^rb   c                (    d|j                          S )Nr  )rh   )r  r  s     rc   r  z%PythonWrapperCodegen.make_buffer_free  s    foo'())rb   c                8    ddj                  d |D               S )Nr  r  c              3      K   | ]  }|  y wre   ra   )r   r   s     rc   r   z:PythonWrapperCodegen.make_free_by_names.<locals>.<genexpr>  s     >>s   )r   )r  names_to_dels     rc   make_free_by_namesz'PythonWrapperCodegen.make_free_by_names  s    dii>>>?@@rb   c           	     `    | j                    | d| | | j                   d| j                   d	S )Nr   r   reuse)r  r  r  )r  r  r  del_lines       rc   codegen_exact_buffer_reusez/PythonWrapperCodegen.codegen_exact_buffer_reuse  s@    ../zXJxjQUQ\Q\P]]_`d`l`l_mmsttrb   c                P    |$| j                  | j                   d| d|        y y )Nz [Provenance debug handles] r  )r   r  )r  r  debug_handles      rc   write_provenance_debug_handlez2PythonWrapperCodegen.write_provenance_debug_handle  s4    
 #NN<<. <[M<.Y $rb   c                r   |j                         |j                         k(  sJ |j                         }|j                         }d}|t        j                  j	                         vr|rd| j                  |       }|j                         |j                         k(  r4|j                         |j                         k(  r| j                  |||      S | j                  ||j                         |j                         d| j                  j                        }| j                   | d| | d| j                   dS )N;z; r   r   r  r  )rm   rh   r@   rf   r  r  rM  rN  r  r  r  r   r  r  )r  r	  newrl  r  r  r  reinterpret_views           rc   ro  z&PythonWrapperCodegen.make_buffer_reuse  s   }}#--/111<<><<>1773355*D11#678H<<>S\\^+0@CNNDT0T228XxPP88!11d6G6G6Q6Q
 ,,z-=,>xj4<<.X^__rb   c                    | j                  t        || j                   | d|j                          | j                   d| j
                   d             y )Nr   r  z alias)r   rD   r  r  r  r  )r  r   rg  s      rc   rh  z0PythonWrapperCodegen.codegen_deferred_allocation  sS    <<.c$*@*@*B)CDKK=PRSWS_S_R``fg	
rb   c                t   |j                         }|t        j                  j                  v sG|| j                  v s9t        |t        j                  t        j                  t        j                  f      ry | j                  j                  |       t        |j                         t        j                  t        j                  f      r|j                         sy |j                         }t        |t        j                         ry t        |t        j"                        ry t        |t        j$                        r+t        |j&                  t        j(                        s*J dt+        |j&                         d|j&                          |j&                  j,                  }t        |t        j.                        sJ t+        |             |j,                  }t        |t        j0                  t        j(                  f      sJ t+        |             t        |t        j(                        rdfd |      }| j3                  |       | j5                  t7        | |||             y t        |t        j8                        r| j5                  t;        | |d             y | j5                  t;        | |             y )Nzunexpected r   c                   t        | t        j                        r | j                               S t        | t        j                        r | j
                        S t        | t        j                        sJ t        |              | S re   )rw   r'   BaseViewunwrap_viewr  r   r   r   )r  unwrap_viewss    rc   r  z=PythonWrapperCodegen.codegen_allocation.<locals>.unwrap_views0  sd    !&"++6+F,>,>,@AA!&"--8+FKK88%fbii8F$v,F8!Mrb   Tr?  )r   	ir.Buffer)rh   r@   rf   r  r  rw   r'   DonatedBufferSubgraphBufferInputBufferr   get_defining_opExternKernelAllocMultiOutputshould_allocaterv   MutationLayoutSHOULDREMOVEr  rf  rg  r*   r   r   r   r   r  r   rb  rx   r0  )r  r  r   r{   boxinput_bufferr  s         @rc   r  z'PythonWrapperCodegen.codegen_allocation  s     AGG+++t~~%&2#3#3R5F5F"WX4 &&(%%r~~6 **,'')fb;;<fbmm,fb001fkk2+=+=> d6;;/06;;-@> ++""Cc2==1<49<188LlRYY8J8J,KL dO L ,(:(:;"  ,L9##L1NN?4vvNOfb112NN<f$GH|D&12rb   c                   |j                         }t        |t        j                  t        j                  f      r| j                  t        | |             y t        |j                         t        j                        r| j                  t        | |d             y | j                  |      sy | j                  j                  |       | j                  t        | |             y )NTr?  )rh   rw   r'   r  rN  r   r  rv   rx   r=  	can_reuser  r   )r  r  r   s      rc   codegen_freez!PythonWrapperCodegen.codegen_freeC  s      fr~~r/A/ABCNN8D&12f,,.0C0CD NN.tVNO~~f%

t*489rb   c                   |j                         }|t        j                  j                  v xs |t        j                  j                  v xr6 t        t        j                  j                  |   t        j                         xsh |t        j                  j                  v xsJ |t        j                  j                  v xs, |t        j                  j                  v xs || j                  v  S re   )rh   r@   rf   r  r  rw   graph_inputs_originalr'   r  r  r  never_reuse_buffersr  )r  r  output_bufferr   s       rc   r  zPythonWrapperCodegen.can_reuseW  s    $$&AGG+++ 
",,, "GG11$79I9I 
" qww(((
" qww222
" qww222
" tzz!
 	
rb   c                    |j                         | j                  v xr. | j                  |j                            |j                         k(  S re   )rh   r  )r  r  reused_buffers      rc   	did_reusezPythonWrapperCodegen.did_reuseg  sC     OO, KFOO-.-2H2H2JJ	
rb   c                t   t        ||      sJ | j                  |       | j                  j                  |j	                                | j
                  j                  |j	                                |j	                         | j                  |j	                         <   | j                  t        | ||             y re   )	r   r  r  r   rh   r  r  r   rC  )r  r  r  s      rc   codegen_inplace_reusez*PythonWrapperCodegen.codegen_inplace_reuseo  s    $\=AAA-

|,,./=11340<0E0E0GM**,-y|]CDrb   c                    t        |      }|| j                  v r|S | j                  j                  |       | j                  |z   S re   )rW   r  r   r  )r  r  r   s      rc   codegen_unbacked_symbol_declz1PythonWrapperCodegen.codegen_unbacked_symbol_declw  sC    6{4---K &&**40<<$&&rb   c                    t        t        j                  j                  j                  |      }| j                  t        | |||             y re   )r   r@   rf   rn   rZ  r   r  )r  r  r  r  s       rc   (codegen_unbacked_symbol_defs_for_outputsz=PythonWrapperCodegen.codegen_unbacked_symbol_defs_for_outputs  sB     6GG&&(9
 	"4g?PQ	
rb   c                    |sy |j                         D ]I  \  }dfdfd}| j                  | j                  |       d |        | j                          K y )Nc                0   |dk(  r| S t        |      dk\  r_t        |d   t              rLt        |d   t        j                        r/ |  d|d   j
                   d|d   j                   d|dd        S t        |d   t              r |  d|d   j
                   d|dd        S t        |d   t        j                        rYt        j                  j                  r  d	|d   j                   d
|  d|dd        S  |  d|d   j                   d|dd        S t        |d   t              rYt        j                  j                  r  d|  d|d   j                   d|dd        S  |  d|d   j                   d|dd        S t        d|       )Nra   r$   r   rA   r   r  r   r~  z	std::get<z>(r}  r~  z / z.__floordiv__(r  )r   rw   r   pytreeSequenceKeyr   r  r@   rf   r  r   r  r  )rh  r  gos     rc   r  zJPythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs.<locals>.go  s   b=K LA%"71:}="71:v/A/AB&'!*//!2!GAJNN3C1Ewqr{   
M:a
'8;WQR[II
F,>,>? 77.. Ywqz~~&6ba@'!"+N  4&'!*..)9 ;WQR[I
  
K8ww**!AdV3wqz/A/A.B!"DgabkRRnWQZ5G5G4HJGTUTVKXX(+@	)JKKrb   c                    t         j                  j                  rt              dk(  rZd   }  d   j	                         t        | t        j                        r!t        | j                        dk7  r	dd        S       S t        d   t        j                        sJ  d   j                     j	                         dd        S        S )NrA   r   )r@   rf   r  r   rh   rw   r'   r  rz  r  r  r  )r  r  r  r  r  s    rc   go_outerzPPythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs.<locals>.go_outer  s    77&&
 7|q(%aj  "#AJ//1)#r~~>3s{{CSWXCX $ABK   ")	    *'!*f6H6HIII!''!*.."9"B"B"DgabkRRk733rb   r   )rh  rW   r  zpytree.KeyPath)r  r   r  r  )r  r  r  r  rE  r   r  r  s    ``   @@rc   r  z>PythonWrapperCodegen._codegen_unbacked_symbol_defs_for_outputs  sg     ! ,113 =	JAw
L>4. NN44Q78HJ<}Uw=	rb   c                     fd} fd}	  j                  j                          j                   j                   dj                           |        t
        j                  }t        j                  j                        5  j                  j                  |       d d d         |         j                          y # 1 sw Y   !xY w#  j                          w xY w)Nc                    t        j                  j                        t              k(  sJ t        j                  j                        D ]3  \  } }j	                  j
                   |  d| j                          5 y r  )r   rf   r  r   r   r  r  )inner_inputouter_inputouter_inputsr  subgraphs     rc   _codegen_subgraph_prefixzSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_prefix  sy    x~~223s<7HHHH,/++\- ([ ||n[M[M$++Orb   c                    t        j                  j                        t              k(  sJ t        j                  j                        D ]5  \  } }j	                  | d| j                          j                          7 y r  )r   rf   r  r   r   r  r  )inner_outputouter_outputouter_outputsr  r  s     rc   _codegen_subgraph_suffixzSPythonWrapperCodegen.codegen_subgraph_by_inlining.<locals>._codegen_subgraph_suffix  s{    x~~334M8JJJJ.1,,m/ *l #nC(F(F(H'I$++Wrb    subgraph: )parent_graph)	rX  rf   r   r  r   r@   rG  codegen_subgraphr{  )r  r  r  r  r  r  r  s   ````   rc   codegen_subgraph_by_inliningz1PythonWrapperCodegen.codegen_subgraph_by_inlining  s    			'%%hnn5NNdll^;x}}oFG$&77L$$X^^4 //!- 0  %&$$&  $$&s$   A;C C,C CC C*c           	        |j                   }|j                  }t        |j                               |j                  D cg c]  }|j
                   c}z   }dj                  |      t        |      dk(  rdndz   }|D cg c]  }|j                          }	}dj                  |	      t        |      dk(  rdndz   }
| j                  d| d| d       |j                         D cg c]
  \  }}|s	| }}}|r#| j                  ddj                  |              | j                  d	|
 d
| d| d       | j                  d| d       yc c}w c c}w c c}}w )z'Generate code to call a graph partitionr  rA   rY  r  	partition	_args = [r~  r  r  z) = self.partitions[z](partition_args)zdel partition_argsN)input_deallocationoutput_nodesr(  r  symbol_inputsr   r   r   rh   r   r  )r  partition_idr  r  r  symbol_inputrZ  r\   rp   output_namesr  r   
deallocater  s                 rc   codegen_partition_callz+PythonWrapperCodegen.codegen_partition_call  sj    2DD+88-22452F2T2T9
".L9
 
 ;'#k2Ba2G3RP4@ADAA))L)C4E4JSPRS 	<.	&CD *<)A)A)C
%T:zD
 
 NNT$))L"9!:;< 	y,\N+l^SYZ	
 	|nE:;-9
 B
s   E?E'
E2Ec                P    t        |      D cg c]  }d| 	 c}| _        y c c}w )N
partition_)r  r  )r  num_partitionsr  s      rc   set_all_partition_namesz,PythonWrapperCodegen.set_all_partition_names(  s$    BGBW#X3j$6#X #Xs   #c           	     p   dj                  |      t        |      dk(  rdndz   }dj                  |      t        |      dk(  rdndz   }| j                  |j                  j                   d| d       | j                  d| d|j                  j                   d|j                  j                   d	       y )
Nr  rA   rY  r  r  r~  r  z) = r  )r   r   r   rf   r   )r  r  r  outer_flattened_outputsouter_output_namesouter_input_namess         rc   ,codegen_subgraph_call_with_flattened_outputszAPythonWrapperCodegen.codegen_subgraph_call_with_flattened_outputs+  s     "YY'>?./14C"
 !IIl3|$)Cr
 	(..--.i8I7J!LM 	"#4(;(;'<Ahnn>Q>Q=RRXY	
rb   c                v   dj                  |      t        |      dk(  rdndz   }| j                  |j                  j                   d| d       t
        j                  j                  j                          | j                  | d|j                  j                   d|j                  j                   d	       y )
Nr  rA   rY  r  r  r~  r   r  r  )r   r   r   rf   r   r@   rj   free_buffers)r  r  r  outer_buffer_namer%  s        rc   codegen_subgraph_callz*PythonWrapperCodegen.codegen_subgraph_call=  s     IIl3|$)Cr
 	(..--.i8I7J!LM 	
&&( 	 !X^^%8%8$98>>;N;N:OvV	
rb   c                    | j                  |j                         | j                  d       | j                  | j                   d|j                          t
        j                  }|j                  |j                  _        |j                  |j                  _        |j                  j                  | j                  vrt        j                  |j                        5  t        j                  dd      5  |j                  j                         \  }}d d d        d d d        |j                  j                  }| j                  j                  |       | j                  |       y y # 1 sw Y   VxY w# 1 sw Y   ZxY w)Nr  r  r^  F)rX  rf   r  r  r   r@   r  
fx_wrapperr  rG  r&   patchr\  r   r  )r  r  r  r  r  r  s         rc   codegen_subgraph_commonz,PythonWrapperCodegen.codegen_subgraph_commonN  s/   !!(..1"T\\N+hmm_EFww%1%=%="$0$;$;!>>d&F&FF $$X^^4 @\\"3U; @'/~~'='='?$M1@@
 %NN//M,,00?,,]MJ G
@ @@ @s$   E45E(E4(E1	-E44E=c                L    | j                  |       | j                  |||       y re   )r.  r&  )r  r  r  r#  s       rc   'codegen_subgraph_with_flattened_outputsz<PythonWrapperCodegen.codegen_subgraph_with_flattened_outputsc  s(     	$$X.99l$;	
rb   c                L    | j                  |       | j                  |||       y re   )r.  r*  )r  r  r  r)  s       rc   r  z%PythonWrapperCodegen.codegen_subgraphk  s%     	$$X.""8\;LMrb   c                   |j                         }| j                  | dt        |j                                |j                  D cg c]  }|j                          }}t        j                  j                  rOt        t        |j                              D cg c]
  }| d| d }}| j                  |j                  ||       y | j                  |j                  ||       y c c}w c c}w )N = [None] * r}  r~  )rh   r   r   r  r\   r  r@   rf   r  r  r  r  r  )r  invoke_subgraphr   rR  r  r   r  s          rc   codegen_invoke_subgraphz,PythonWrapperCodegen.codegen_invoke_subgraphq  s    '')$|C0G0G,H+IJK;J;Q;QRC--/RR77(-c/2I2I.J(K#$4&!AM  --((, !!/":":L$O Ss   C(C-c                   |j                         }|j                  D cg c]  }|j                          }}|j                  j                         }t	        |j                  t
        j                        s| d}| j                  | dt        |j                                | j                  d| d       | j                  t        | |j                  j                               t        j                  j                  rOt        t        |j                              D cg c]
  }| d| d }}| j!                  |j                  ||       n| j#                  |j                  ||       | j                  t%        |              | j                  d       | j                  t        | |j&                  j                               t        j                  j                  rOt        t        |j                              D cg c]
  }| d| d }}| j!                  |j&                  ||       n| j#                  |j&                  ||       | j                  t%        |              y c c}w c c}w c c}w )Nr  r3  r   r  r}  r~  zelse:)rh   operandsr  	predicaterw   r'   ShapeAsConstantBufferr   r   r  rQ  true_subgraphrf   r@   r  r  r  r  rw  false_subgraph)r  conditionalr   rR  r  r8  r   r  s           rc   codegen_conditionalz(PythonWrapperCodegen.codegen_conditional  s   ##%;F;O;OPC--/PP));;=	+//1I1IJ$+W-I$|C0C0C,D+EFGYKq)*({/H/H/N/NOP775:3{?R?R;S5TUvQqc^UMU--))< !!+";";\4P'-.w({/I/I/O/OPQ775:3{?R?R;S5TUvQqc^UMU--**L- !!+"<"<lDQ'-.9 Q V Vs   I)"I.I3c                     fd}|j                         }|j                  D cg c]  }|j                          }}|j                  D cg c]  }|j                          }}t	        |      } j                  | dt	        |              |r  j                  | dt	        |       d       t        |      D ]  \  }	}
 j                  | d|	 d|
          g t        t	        |            D 	cg c]
  }	| d|	 d c}	|}| dg}t        |      }|d	t	        |       } ||j                  ||        j                  d
|d            j                  d       |rwt        |      D ]h  \  }	} j                  t         |j                  j                                j                  | d|	 d| d        j                  t                      j nvt        |      D ]h  \  }	} j                  t         |j                  j                                j                  | d|	 d| d        j                  t                      j  j                  d        j                  t         |j                  j                                ||j                  ||        j                  t                      |r j                  t         |j                  j                               t        t	        |            D ]"  }	 j                  | d|	|z    d| d|	 d       $  j                  t                       j                  t         |j                  j                                ||j                  ||        j                  t                       j                  d|d           |r j                  d       t        t	        |            D ]  }	 j                  d| d|	|z    d        j                  t         |j                  j                                j                  | d|	 d| d|	|z    d        j                  t                       y	y	c c}w c c}w c c}	w )z1while_loop is codegened as a host side while_loopc                    t         j                  j                  rj                  | ||       yj	                  | ||       y)z3Helper method to deduplicate subgraph codegen logicN)r@   rf   r  r  r0  )r  r  r  r  s      rc   r  zAPythonWrapperCodegen.codegen_while_loop.<locals>.codegen_subgraph  s7    ww11(L-X<<lMrb   r3  z.extend([[] for _ in range(z)])r}  z] = r~  _cond_resultNzshould_loop = r   zif not should_loop:z.unsqueeze(0).clone()r  zwhile should_loop:z	].append(z])z    should_loop = z%# Stack outputs after loop completionzif len(z]) > 0:z] = torch.stack(z	], dim=0))rh   carried_inputsr  additional_inputsr   r   r  r  r(  cond_subgraphrQ  body_subgraphrf   rw  )r  
while_loopstack_outputr  r   rR  outer_carried_inputsouter_additional_inputs
ckp_offsetr   inpcond_outer_inputscond_outer_outputsbody_outer_inputsbody_outer_outputscarried_inputs   `               rc   codegen_while_loopz'PythonWrapperCodegen.codegen_while_loop  s   	 ""$/9/H/H 
(+C!!# 
  
 0:/K/K#
(+C!!##
 #
 -.
$|C0D,E+FGHNN&3C8L4M3NcR   45 	3FAsNNdV1QCtC512	3
&+C0D,E&FGas!nG
$
 "&l34 
 //J5I1JK$$&79K	
 	(:1(='>?@,-$-.B$C 7 =0z7O7O7U7UVW$q4>STU/567
 %..B$C 7 =0z7O7O7U7UVW$q4hGH/567
 	+,(z/G/G/M/MNO$$&79K	
 	'-. NN,T:3K3K3Q3QRS3345 P$qZ(8	$q2NOPNN+D12 	(z/G/G/M/MNO$$&79K	
 	'-.+,>q,A+BCD NNBC3345 7aJ/?wGH0z7O7O7U7UVWfAaS 0aJ7GyQ /567 Q 
#
  Hs   Q+Q04Q5c                    	 t        | dd       ry t        | t              r| S t        j                  j
                  j                  |       }||S t        |      S # t        $ r Y y w xY w)Nrg  )r  rw   rY   r@   rf   
_shape_env_maybe_evaluate_staticr9  )r   r  s     rc   statically_known_int_or_nonez1PythonWrapperCodegen.statically_known_int_or_none  sf    	q.$/ !S!''$$;;A>C{
s8O 		s!   A A ,A 
A 	A&%A&c                l    g }| D ],  }t         j                  |      }| y |j                  |       . |S re   )rR  rT  r   )lstr  r   r)  s       rc   %statically_known_list_of_ints_or_nonez:PythonWrapperCodegen.statically_known_list_of_ints_or_none  sA     	A&CCAFC{MM#		
 rb   c                0    t         j                  |       d uS re   )rR  rW  )rV  s    rc    is_statically_known_list_of_intsz5PythonWrapperCodegen.is_statically_known_list_of_ints  s     !FFsKSWW	
rb   c                H    t         j                  | j                               S re   )rR  rW  rM  r  s    rc   rD  z4PythonWrapperCodegen.static_shape_for_buffer_or_none!  s    #IIOO
 	
rb   c                0    t         j                  |       d uS re   )rR  rD  r[  s    rc   !can_prove_buffer_has_static_shapez6PythonWrapperCodegen.can_prove_buffer_has_static_shape'  s    #CCFKSWWWrb   c                     y re   ra   )r  r  node_schedules      rc   write_kernel_context_guardz/PythonWrapperCodegen.write_kernel_context_guard+  s    
 	rb   c                     y)z<
        Mark the beginning of kernel context guard
        Nra   r  s    rc    write_kernel_context_guard_beginz5PythonWrapperCodegen.write_kernel_context_guard_begin2       	rb   c                     y)z6
        Mark the end of kernel context guard
        Nra   r  s    rc   write_kernel_context_guard_endz3PythonWrapperCodegen.write_kernel_context_guard_end:  rc  rb   re   )r  r  r  r   r  PythonWrapperCodegen | Noner  !ir.GraphPartitionSignature | Noner_  )r   rW   r   rW   r   r>  )r  rW   )r>  TritonMetaParamsr   rW   r   r  )r   z9dict[str, ir.TensorBox | ir.TorchBindObject | sympy.Expr]r   zlist[IRNode])rZ  r  r   rY   )rZ  zIterable[str]r   r>  )
r   rW   r   rW   r   rW   r  rW   r   r>  )r[  r7   r   rW   r   rW   r   rW   r  rW   r   r>  )rv  rW   r  rW   r   r>  )r[  r7   rv  rW   r  rW   r   r>  )r  rY   r  rW   r   rW   r  )rA   N)r  rY   r  rY   r  zdict[int, int] | Noner   r>  )r  rY   r   r  )r  r  r   r>  r  r7   r   r>  )rp   zir.FallbackKernelr   r>  )rp   r  r   r>  )rp   r  )rp   r  r   r>  )r  rW   r  rW   r  r   r  r  r[   rW   r  zOrderedSet[str] | Noner   r>  )F)rp   r  )rp   r  r   r>  )r  rW   r  rW   r  zCallable[[], Sequence[str]]r  z6torch._ops.OpOverload | torch._ops.HigherOrderOperatorr  ry  r  zSequence[ir.Buffer]r   r>  )r  r7   r   Callable[..., None]r   zIterator[Callable[..., None]])r  r  )r   rW   r  rk  ra  rl  )r_  r  )r   r
   ro   r  r   rW   )r   r
   r   rW   )r  rW   r   rW   r6  rW   r   rW   )rT  zSequence[Expr]r   rW   )r   ztuple[str, list[str]])r   rm  r   rW   )r  z
bool | str)rX   rW   r   r>  )rp   zir.MultiOutput)NTN)
r  rW   r  rW   r  r   r  r  r  r   )r  rW   r  rW   r  r   )r   rW   )r   zlist[list[int | sympy.Expr]]r  $tuple[ir.ComputedBuffer, str] | None)r  rW   r  r   )r  r  rf   rP   r   r>  )r\  rF   )r  rW   )r  r/  )NF)r  )r  r  )r  r  )r  rW   r  rW   r  rW   )r  r  )r	  r/  r  r/  rl  r  )r   rW   rg  zir.ReinterpretViewr   r>  r  r  )r  r  r  r  )r  rW   r  r   r  r  r   r>  )r  rY   r  zir.GraphPartitionSignature)r   rY   )r  rW   r_  z*Sequence[BaseSchedulerNode] | ExternKernel)r]   r^   r_   r  r  r`   r%  rj  r  r  r  r  r,  r  r2   r6  r:  r<  rA  rE  rG  rJ  rC  rS  rU  rW  r\  ra  rc  r  rj  rq  rp  r  rw  r  r~  r  r  r  rX  r{  rT  ry  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r  r   contextmanagerr
  r  r  r  rB  rN  r  rc  ri  rq  r  rv  ry  r{  r  r   rQ  r  r  r  r  r  r  r  r  rs  r  r  r  r  r  r  classmethodr  rK  rR  r  r^  r`  rX  r  r  r  r  r  r  r  rq  r   r~  r  r  r   r  r  r  rK  r  r[  r  r  r  r  r  r  ro  rh  r  r  r  r  r  r  r  r  r  r  r!  r&  r*  r.  r0  r  r5  r=  rP  rT  rW  rY  rD  r]  r`  rb  re  r@  rA  s   @rc   rR  rR    s    "d!fBP 
 CG	&&!& 4& @	& &'<oBb:  "	B + +! 
 

$	B$
%?0(
.8):6-KPP"P,/P:=P	PUU U 	U
 U U 
UH

"

14

?B

	

*G.10J
/+ <@	*-*- *- %:	*-
 
*-X5 
$,45$:4,:8 8 
8 04;; ; 	;
 ; ; -; 
; (	U
8(B<
	V	V  	V .		V
 L	V  	V %	V 
	V0 !! '! 
'	! !2,S
j.S`:
'
R%>K>K >K -	>K@'?RD @D W CG +.&'6	, <P '<P 
<P|>5S
IG.64*^@4  $%)

 
 	

 
 #
& CG		'*	6@	 	  $%)!! ! 	!
 ! #!,>  uJ ,uJ >uJn	*A"A+8A	A
%N:.>85WF

@
4=G~' !6
6
z !P(P(d !,%P	
 8 8 TY.`*_*Au $( !` 
53n:(
 
E'

 
 E	

 

KK K E	K
 
KZ+'Z<< 9<BY
$
"K*
NP /B]7~     
 

 
 

 X X Brb   rR  c                       e Zd ZdZ	 d	 	 	 	 	 d fdZddZddZd Zd Zd Z	ddZ
dd	Zdd
ZddZ	 	 ddZddZddZd  fdZ	 d	 	 	 	 	 d!dZd"dZddZedd       Zedd       Zed#d       Zd Z xZS )$r  a  
    A wrapper codegen that generates code for a subgraph. For most of the
    methods, we rely on the implementation in the PythonWrapperCodegen. But we
    override a few functions to produce cleaner code (like avoiding writing
    imports twice in the output code)
    c                    || _         || _        || _        t        |           | j                         }|j                  | _        |j                  | _        |j                  | _        |j                  | _	        y re   )
r  r  r  r$  r%  get_root_graphr  r   r  r  )r  r  r  r  rootr,  s        rc   r%  z%SubgraphPythonWrapperCodegen.__init__K  sn     +,$8!""$$($=$=!%)%?%?"!//)-)G)G&rb   c                &    | j                   | _        y re   )r  r  r  s    rc   r  z1SubgraphPythonWrapperCodegen.set_launcher_fn_nameb  s     !% 2 2rb   c                     y re   ra   r  s    rc   r  z)SubgraphPythonWrapperCodegen.write_headerh  r-  rb   c                     y re   ra   r  s     rc   r  z2SubgraphPythonWrapperCodegen.add_benchmark_harnessk  r-  rb   c                     y re   ra   r  s     rc   r  z6SubgraphPythonWrapperCodegen.benchmark_compiled_modulen  r-  rb   c                     y re   ra   r  s    rc   rW  z5SubgraphPythonWrapperCodegen.write_async_compile_waitq  r-  rb   c                6    | j                   j                         S re   )r  r  r  s    rc   r  z/SubgraphPythonWrapperCodegen.next_kernel_suffixt  s    ""5577rb   c                     y re   ra   r  s     rc   r  z2SubgraphPythonWrapperCodegen.generate_after_suffixx  rH  rb   c                \    | j                   j                  d| j                   d       d}|S )Nz
            def z(args):
            rA   )r  r   r  r_  s     rc   ra  z>SubgraphPythonWrapperCodegen.write_launcher_fn_call_get_indent{  s<    &&' (	

 rb   c                     yr   ra   r  s    rc   r  z4SubgraphPythonWrapperCodegen.get_wrapper_call_indent  s    rb   c                    | j                   x}r3|j                  |j                  D ci c]  }t        |      | c}z  }|S t        j
                  j                  }|S c c}w re   )r  input_nodesr  rW   r@   rf   r  )r  r  rE  r\   s       rc   rJ  z-SubgraphPythonWrapperCodegen.get_graph_inputs  sh     11191**#,#:#:.A	. F
  WW))F.s   Ac                    | j                   x}rJt        |j                  j                               |j                  D cg c]  }|j
                   c}z   }|S t        j                  j                  }|S c c}w re   )	r  r(  r  r  r  r   r@   rf   rQ  )r  r  r  namess       rc   rc  z2SubgraphPythonWrapperCodegen.get_graph_input_names  sr    11191..33566?6M6M:&2!!: E
  GG--E:s   A5c                r    | j                   x}r|j                  }|S t        j                  j                  }|S re   )r  r  r@   rf   r  )r  r  r  s      rc   rC  z.SubgraphPythonWrapperCodegen.get_graph_outputs  s;    11191,,G  gg++Grb   c                ~    |j                         }| j                  x}r||j                  v ry t        |   |       y re   )rh   r  r  r$  r  )r  r  r   r  r,  s       rc   r  z/SubgraphPythonWrapperCodegen.codegen_allocation  s?     222I2	@U@U8U "6*rb   c                >   | j                          t        j                  j                  j                  x}D| d}| j                  | dt        j                  j                  j                  |              |S t        |      }| j                  | d| d       |S )N_rawr   r  r   )	r6  r@   rf   rj   r  r   r  stream_handler0   )r  r  rf   r  r   s        rc   _write_get_raw_streamz2SubgraphPythonWrapperCodegen._write_get_raw_stream  s     	%%'#$77#4#4#H#HHU)*$/DNN&AGG..<<=PQRS  'z2DNNdV#5j\CDrb   c                .    | j                  d| d       y)z#Generate NVTX range push for graph.z"torch.cuda.nvtx.range_push('graph re  Nrt  )r  post_grad_graph_ids     rc   codegen_graph_nvtx_range_pushz:SubgraphPythonWrapperCodegen.codegen_graph_nvtx_range_push  s    ;<N;OrRSrb   c                &    | j                  d       y)z"Generate NVTX range pop for graph.ztorch.cuda.nvtx.range_pop()Nrt  r  s    rc   codegen_graph_nvtx_range_popz9SubgraphPythonWrapperCodegen.codegen_graph_nvtx_range_pop  s    45rb   c                8    | j                   j                          y re   )r  r6  r  s    rc   r6  z5SubgraphPythonWrapperCodegen.write_triton_header_once  s     	446rb   c                8    | j                   j                          y re   )r  r<  r  s    rc   r<  z=SubgraphPythonWrapperCodegen.write_get_raw_stream_header_once  s     	<<>rb   c                    | }t        |t              r|j                  }t        |t              rt        |t              sJ |S re   )rw   r  r  rR  )r  ru  s     rc   rt  z+SubgraphPythonWrapperCodegen.get_root_graph  sB    DH;<&&D ;< $ 4555rb   c                     y re   ra   r  s    rc   r  z<SubgraphPythonWrapperCodegen.generate_and_run_autotune_block  s    rb   re   )r  rW   r  rR  r  rg  r_  r  rl  rk  )r   z@dict[str, ir.TensorBox | ir.TorchBindObject | sympy.Expr | None]ri  rj  ro  )r  rY   rf   zGraphLowering | Noner   rW   )r  rY   r   r>  )r   rR  )r]   r^   r_   r  r%  r  r  r  r  rW  r  r  ra  r  rJ  rc  rC  r  r  r  r  r2   r6  r<  rt  r  r@  rA  s   @rc   r  r  C  s     CG	HH -H @	H.38		I	+ >B&:	T6 7 7 ? ?  rb   r  )rp   r/  r   r<  )rp   r/  r   r?  )r   r/  r   r/  )NN)r   rW   r   zlist[triton.Config]r   zlist[TritonGrid]r   rf  r   r   r   ztuple[str, str]re   )r  rn  r   rW   )
__future__r   r&  r   r  r   rE  r9  r1  r  r  rR  r  r3  collections.abcr   	itertoolsr   r   typingr   r   r	   r   r
   r  
torch._opstorch.utils._pytreeutils_pytreer  r   rv  torch._dynamo.utilsr   r   r   #torch._inductor.codegen.debug_utilsr   $torch._inductor.codegen.multi_kernelr   %torch._inductor.runtime.runtime_utilsr   torch._library.opaque_objectr   r   torch._loggingr   %torch.fx.experimental.symbolic_shapesr   r   r   r   r   torch.fx.noder    torch.multiprocessing.reductionsr   torch.utils._ordered_setr   torch.utils._sympy.functionsr   r    torch.utils._sympy.singleton_intr    rY  r"   r#   r  r%   r&   r'   	codecacher(   r)   r*   r   r+   runtime.hintsr,   stream_constantsr-   r.   r/   stream_utilsr0   r1   r2   r3   r4   r5   r6   r7   r8   r9   r:   r;   r<   r=   r>   r?   virtualizedr@   r'  rB   rC   rD   rE   rF   rG   	cpp_utilsrH   custom_extern_kernel_codegenrI   triton_utilsrJ   rK   rL   rM   rN   rO   r   rf   rP   rQ   rj   rR   wrapper_fxirrS   	getLoggerr]   logdoprintrx  r   r[   rW   r  rY   r<  r?  r   r/  rN  r  rV   rt   r|   r   r   r  rh  r   r   r  r  r!  rC  rT   rQ  rc  rl  rq  rw  r  r  r  r  r  r  r  r  r  r  r  r  r
  r  r0  r=  rb  rC  rA  rv  r  r  r  r  r  r  r  LinerR  r  ra   rb   rc   <module>r     s   "    
     	  	  $ " / /     $ $ & E E C A ; R +  . ; / 1 9 : ( ( ' ( ' , W W ?        F P P <<%!-) g! u{{Cs:;5<<c;NPSST YY%
]OT12  3 3 3 >QF S> 	#


C
 8-=,>c3h,O#PP  ,0'+k&
k& k& k& )	k&
 %k& k&^ EIq&Aq&q&h   %6 %6PQ8 QY Y
 2 2 2 	/k 	/ 	/ ++ + + 	2 	2 	2 1{ 1 1 "@K "@ "@J?; ? 
7K 
7 
7 5+ 5 5< P{ P P. 	({ 	( 	(  /[  /  /F 5; 5 5* ; ; ;2 2O  : 1M        %
 %
P h,% h, h,V +6, +6 +6\ /( / /& )" ) )0(! (
 #0k #0 #0L 6; 6 6, 4+ 4 48 	5+ 	5 	5 8[ 8 8 6; 6 6" S+ S S  
K'W17 W1tb^#7 ^rb   