tjO UddlmZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl m Z m Z mZmZmZddlmZmZddlmZddlmZddlmZddlmZmZmZmZmZmZmZm Z m!Z!m"Z"m#Z#dd l$m%Z%m&Z&m'Z'm(Z(m)Z)m*Z*dd l+m,Z,ddl-Z-dd l-m.Z.m/Z/m0Z0ddl1m2cm3cm4Z5ddl6m7cm8Z8ddl9m7cm:Z;ddlZ=ddl?m:cm@ZAdd lBmCZCdd lDmEZEddlFmGZGddlHmIZIddlJmKZKddlLmMZMddlNmOZOddlPmQZQmRZRddlSmTZTmUZUmVZVmWZWmXZXmYZYddlZm[Z[m\Z\m]Z]m^Z^m_Z_m`Z`maZambZbmcZcmdZdmeZeddlfmgZgddlhmiZiddljmkZkddllmmZmmnZnmoZompZpmqZqmrZrddlsmtZtddlumvZvmwZwddlxmyZymzZzm{Z{m|Z|m}Z}ddlwm~Z~mZmZmZmZmZddlmZdd lmZmZmZmZdd!lmZdd"lmZmZdd#l:mZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZdd$lmZmZmZe r+dd%lZmZdd&lfmZdd'lmZdd(lmZdd)lmZdd*lmZdd+l:mZneZd,ed-< ddlZejZd.Zn#e$rdZd/ZYnwxYwe(d0Ze"d1Ze"d2Ze"d3Zee.zZd,ed4<eeze.zZd,ed5<e=jje=jjzZd,ed6<ejeɦZeje jd78Ze=jjZ e#ed9eed9fd:d;ee#eeed9fd9d:d;fdzfZd,ed<<ddAZejd.BGdCdDZҐddHZӐddLZԐddPZՐddQZ֐ddWZgdXZgdYZ ddd^Zڐdd`Z dddaZeddddZe dddfZ dddiZݐddmZސddpZߐddrZddsZddwZdd|ZddZddZddZdZGdd;Zed/BGddZeGddeZddZeGddeZeGddeZedededededededdZded< dddZeGddeZde/dfddZe ee.ee.gefZd,ed<GddeZGddeZGddeZGddeZeGddeZeGddeZeGddeZddZddZddZ dddDŽZddȄZddʄZeGd˄deZeGd̈́dΐeZeGdτdАeZeGdфdҐeZeGdӄdԐeZeGdՄd֐eZeGdׄdؐeZeGdلdڐeZGdۄdܐeZeGd݄deZ eGd߄de Z eGdde Z ddZ ddZ GddZeGddeZGddeZGddeZGddeZGddeZGddeZeGddeZGddeZed/BGddeezZed/BGddeeZGddeZGddeZGddeZeGddeZeGddeZed/BGddeZejd.BGddZGd d eZ Gd d e Z!eeze"zeze#eezeze"zzZ$Gd dZ%Gdde%Z&Gdde!Z'Gdde Z(Gdde Z)Gdde Z*Gdde Z+ddZ,ed/BGdd eZ-Gd!d"e-Z.Gd#d$e.Z/ej0dd'Z1dd)Z2ed/BGd*d+e-Z3ed/BGd,d-e3Z4Gd.d/e4Z5Gd0d1e3Z6Gd2d3eZ7Gd4d5e3Z8Gd6d7e8Z9Gd8d9e8Z:Gd:d;e3Z;Gd<d=e3Z<Gd>d?e3Z=Gd@dAe3Z>GdBdCe3Z?GdDdEe?Z@GdFdGe6ZAGdHdIe3ZBGdJdKe3ZCGdLdMe4ZDGdNdOe3ZEGdPdQe3ZFGdRdSe3ZGGdTdUe3ZHed/BGdVdWZIGdXdYe6ZJed/BGdZd[eJZKGd\d]eJZLeGd^d_eZMGd`dae3ZNGdbdceNZOGdddeeNZPddhZQGdidjeJZRejGdkdleZSGdmd9eSZTGdndoeSZUed/BGdpdqeZVddsZWed/BGdtdue3ZXed/BGdvdwe3ZYdd{ZZed/BGd|d}e3Z[Gd~deJZ\GddeZ]eGdde]Z^eGdde]Z_eGdde]Z`eGdde]ZaGddeJZbGddebZcGddebZdGddebZeddZfddZgddZhdS() annotationsN)Callable GeneratorIterableIteratorSequence)AbstractContextManager nullcontext)Enum)partial)Number) AnycastClassVarLiteraloverload SupportsFloat SupportsInt TYPE_CHECKING TypeAliasTypeVarUnion) assert_neverNeveroverride ParamSpecSelfTypeIs)patch)ExprIntegerSymbol)enable_python_dispatcher)identity)GraphModuleSerializer)can_auto_functionalize)metricsget_free_symbols)FakeScriptObject)get_opaque_obj_repris_opaque_value)compute_required_storage_lengthis_boolean_dtype(is_contiguous_for_memory_format_or_falseis_float_dtypemake_channels_last_strides_for StrideType) &_remove_effect_token_unbacked_bindingscompute_unbacked_bindings free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNodehas_free_unbacked_symbols IterateExprsrebind_unbackedresolve_unbacked_bindingsShapeEnvSymTypes)Node OrderedSet)_disable_current_modes)CleanDivFloorDivMaxMinModModularIndexing)SymT)config dependencies)BackendFeature CodegenSymbolget_scheduling_for_deviceindex_prevent_reorderingKernel)Depextract_free_symbols#extract_input_node_reduction_rangesextract_read_writesSymbolUsageCollectorOpsHandler var_builder)LoopBody) OpCounterCSE OpCountResult ReductionType StoreMode) benchmarker)DeviceProperties ReductionHint)argsort argsort_sym cache_on_selfcache_on_self_and_argsceildivconvert_shape_to_inductorconvert_shape_to_symintdeveloper_warningdo_bench_using_profilingdtype_from_sizeget_dtype_sizeget_kernel_metadataGPU_ALIGN_BYTES ir_dataclass is_dynamicis_gpu sympy_dotsympy_index_symbolsympy_index_symbol_with_prefix sympy_product sympy_substensor_is_aligned)opsOpsValueV) SympyBoolean)Argument) IntLikeType)CUTLASSTemplate)PythonWrapperCodegen) GraphLowering)IndentedBufferrr{TF_P_T_U_V_IntLike_NumLike _OpOverloadsz prefix TensorBoxr"IRNode _NodeOrNodesxobjectreturnTypeIs[int | Integer]c:t|ttfSN) isinstanceintr!rs P/var/www/html/CA-Chatbot/venv/lib/python3.11/site-packages/torch/_inductor/ir.py _is_staticrs a#w ( (()frozencLeZdZUded<ded<ded<ded<d ed <d ed <d S)GraphPartitionSignatureOrderedSet[sympy.Symbol] symbol_inputsz*dict[str, IRNode | Expr | TorchBindObject] input_nodes list[IRNode] output_nodeszdict[str, bool]input_deallocationboolskip_cudagraph list[str]constant_namesN__name__ __module__ __qualname____annotations__rrrrsb,+++<;;;('''rr node_or_nodes_NodeOrNodes | NoneNonec*dfd |dS)Nnodesrrrc |dSt|ttfr|D] }|dSt|tr$|D] }|dSt|t t tttj j j tttt t"f sJdt%|ddS)NzFound zE, which is not a supported top level IR node. See [Note: Inductor IR])rlisttupledictvalues ExpandView DynamicScalar AssertScalarrsympylogicboolalgBooleanr rEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputtype)rnode_check_tensorboxs rrz%validate_ir.._check_tensorboxs = D e} - -  ' '  &&&& ' ' t $ $   ' '  &&&& ' '! K'/#)%   lekkk    r)rrrrr)rrs @r validate_irrs:>]#####rnamestrCallable[..., OpsValue]clttsJtdfd }|S)Nargsrkwargsrrvc:tt|i|Sr)getattrru)rrrs rfnzops_wrapper..fns"!wsD!!426222r)rrrrrrv)rrr)rrs` r ops_wrapperrsL dC ,,$t**,, 333333 Irorder Sequence[int]&Callable[[Sequence[_T]], Sequence[_T]]c tt|tt|dfd }|S)Nindex Sequence[_T]rcttksJfdttDS)Nc,g|]}|Srr).0ir inv_orders r z4inverse_reorder..reindex..'s"???il#???rlenrange)rrs`rreindexz inverse_reorder..reindex%sL5zzS^^++++?????U3u::->->????rrrrr)rziprr)rrrs @rinverse_reorderr"sSSc%jj 1 12233I@@@@@@ Nrcdfd }|S)NrrrcttksJfdttDS)Nc,g|]}|Srr)rrrrs rrz1same_reorder..reindex../s!;;;AeAh;;;rr)rrs`rrzsame_reorder..reindex-sL5zzSZZ'''';;;;;s5zz):):;;;;rrr)rrs` r same_reorderr,s(<<<<<< Nrreindex1&Callable[[Sequence[_U]], Sequence[_V]]reindex2&Callable[[Sequence[_T]], Sequence[_U]]&Callable[[Sequence[_T]], Sequence[_V]]cdfd }|S)Nrrr Sequence[_V]c,|Srr)rrrs rrz fuse_reindexing..reindex8sx(((r)rrrrr)rrrs`` rfuse_reindexingr4s.))))))) Nr)rrI)rrrrIseq#Sequence[int | torch.SymInt | Expr] shape_envShapeEnv | Nonec||td|Drt|}nt||}|S)z1 Convert strides to fill order (argsort) Nc3XK|]%}t|ttjfV&dSrrrrr!rss r z!get_fill_order..Hs3QQ 1sEM.B C CQQQQQQr)allr_r`)rr sorted_idxs rget_fill_orderrBsJ CQQSQQQQQ$+CLL !C00 rSequence[int | Integer]cdt|Dfdtt|D}|S)z Convert stride order to fill order For channel last format, stride order = [3, 0, 2, 1] and fill order = [1, 3, 2, 0] ci|]\}}|| Srrridxposs r z+stride_order2fill_order..Ws 8 8 883c3 8 8 8rc g|] }| Srrrrlookups rrz+stride_order2fill_order..Xs777&)777r) enumeraterr)r fill_orderrs @rstride_order2fill_orderrPsM9 8y'7'7 8 8 8F7777U3u::%6%6777J rct||}dtt|D}t|D] \}}|||< |S)z) Convert strides to stride order cg|]}dSrrr_s rrz$get_stride_order..cs & & &1 & & &r)rrrr)rrroutrelems rget_stride_orderr\s] !/sI > >J & &eCHHoo & & &CZ((4D Jrreplace_symbols_with_hintsrcdSrrrrs rir_node_to_tensorrisRURUr torch.TensorcdSrrrs rrrms 3r IRNode | Nonetorch.Tensor | Nonec|dS|rtjjjntfd|D}t |r&fd|jD}nt |}| }| }t|}t|}tjjj5t!j||||}dddn #1swxYwY|S)Nc&g|] }|Srrrrshape_fns rrz%ir_node_to_tensor..s! . . .AHHQKK . . .rc&g|] }|Srrrs rrz%ir_node_to_tensor..s!===!((1++===r)sizestridedtypedevice)rwgraphsizevarsoptimization_hintr$get_sizeis_storage_and_layout get_layoutrFlexibleLayoutcontiguous_strides get_dtype get_devicerersuppress_guardstorch empty_stridedzero_)rrrrr r!trs @rrrssm yt"7#5 . . . . . . .DQ9====q||~~'<===22488 KKMME \\^^F "4 ( (D $V , ,F   # 3 3 5 5  fE&   %''  Hs+D::D>D>valueSequence[_T] | NoneSequence[_T | None] | Nonec:t|tr|sdgS|Sr)rrr1s rmay_convert_to_optionalr6s+%uv Lr/IRNode | OutputSpec | torch.device | None | str str | NonecDt|ts||St|tjr|jSt|t t fr!t|Std|dt |j ddS)Nzget_device_type(: )) rrr-r!rr OutputSpecget_device_typer+rrrs rr=r=s!S/QY Au| $ $/v A + , ,/q||~~....s9 F FQ%(3q)$$a ( ( F F F F F FrrI)rrmaybe_get_strider6 get_strider%rAndOrrQLerwr"r#guard_or_false)rrMaligned_stridesaligned_last_dim is_aligneds ` ris_aligned_realized_tensorr\s  q& ! !     '  0 0 (  . . (ui F F F F!,,.."2E F F FOx #Q''!**,,r2BA)F)F?,<==J 7  * *: 6 66rstrides1Sequence[_IntLike]strides2shapeczt|t|kr t|t|ksJt|||D]f\}}}tjj|dr,tjjtj||sdSgdS)zP Returns true if the strides are equal, ignoring dimensions of size 1 . rIFT) rrrwr"r#statically_known_leqrXrrQ)r]r_r`dims1s2s rsignificant_strides_equalrfs u::X & &3x==CMM+I+I+I I5(H55 R 7  0 0a 8 8  w..uxB/?/?@@ 55  4rtensorstridesSequence[int | torch.SymInt]ct|s|Stdt||Dr|St |||s|St |\}}g|j}t|D]5\}}tj j |dr ||||<6t|j|j|j||j|j}t't)||S)a Tries to match the strides of the tensor to those in the meta_strides. Strides of insignificant dimensions - size 0 or 1 - will be updated. If there are real stride differences (NHWC vs NCHW), or the tensor is not realized, then the input will be returned c3fK|],\}}tjj||V-dSrrwr"r#statically_known_equalsrrdres rrz2try_match_insignificant_strides..sO   B 00R88      rrIdatalayout)r&rrrTrfr%as_storage_and_layoutrrrwr"r#rb FixedLayoutr!r roffset is_pinnedrReinterpretView)rgrhstorage old_layout new_striderr new_layouts rtry_match_insignificant_stridesr{sN ! ( (    '6#4#4#6#677    $Wf.?.?.A.A6??CTCT U U /77GZ%:$%J&//++,,''1 7  0 0A 6 6 '#AJJqM J _'*EEE F FFrgmtorch.fx.GraphModulec|jdd}dt|jD|jd<ddlm}||dS)Noutputoprcg|]\}}|Srr)rrr s rrz.gm_original_output_strides..s)444Q444ruser_visible_output_idxs)record_original_output_strides)r" find_nodesrrmetatorch._inductor.compile_fxr)r| output_noders rgm_original_output_stridesrs}(%%%2215K44#K$455444K/0JIIIII""2&&&&&rinputsSequence[IRNode] list[Expr]ct}|D]N}|t|dz}|t|dz}Ot |S)NF unbacked_only)r@r)r%rTr)rsym_varsinps rget_symbolic_inputsrsl!+HLL$S\\^^5IIII$S^^%5%5UKKKK >>rct|tr|j}t|tr|}t|t r|j}t|t r|ndSr)rrrpBaseView unwrap_view StorageBoxBufferget_namers r try_get_namer%st!Y F!X MMOO!Z   F%a00 :1::<<ddkZ?ddmZ@ dddoZAddqZBddsZCddtZDddvZE dddyZFddzZGdd|ZHdd}ZIdd~ZJddZKddZLddZMddZNeOre)ddZPdSdS)rzBase class for all intermediate representation (IR) nodes in TorchInductor. Note: This is an abstract base class. Most methods raise NotImplementedError and must be overridden by concrete subclasses. zClassVar[OrderedSet[Any]]_current_originsNzClassVar[int | None]_current_stream_idxF)initOrderedSet[Any]originslist[str] | None tracebacktorch.fx.Node | None origin_nodedict[str, Any]r int | None stream_idxOrderedSet[Node]rGenerator[None, None, None]c#Ktj}||zt_ dV|t_dS#|t_wxYwr)rr)rolds rcurrent_originszIRNode.current_originsDsL%"%- * EEE&)F # # #cF # ) ) ) )s1?c#~Ktj}|t_ dV|t_dS#|t_wxYwr)rr)rrs rcurrent_stream_idxzIRNode.current_stream_idxNsH (%/" - EEE),F & & &F & , , , ,s.<rrc^t|ttttt fSr)rComputedBuffer InputsKernel InputBufferrvTemplateBufferrs ris_realized_nodezIRNode.is_realized_nodeZs+      rc6t|Srrcreateselfs rwrap_for_loweringzIRNode.wrap_for_loweringgs%%%rattrrr1rrc>t|||dSr)r __setattr__)rrr1s r_post_init_setattrzIRNode._post_init_setattrjs" 4u-----rcRt|j}|d||dtjrt jnd|dd|di|d|jdS)Nrrrrr)r@rrrJdebug_ir_tracebackr format_stackr)rrs r __post_init__zIRNode.__post_init__psT233  7333  V5NX/111TX     t444  r222  d.FGGGGGrOrderedSet[str]cXtd|DS)Nc3$K|] }|jV dSrrrdeps rrz(IRNode.get_read_names..|$??s#(??????rr@ get_readsrs rget_read_nameszIRNode.get_read_names{)??dnn.>.>??????rc|jSr)rrs r get_tracebackzIRNode.get_traceback~s ~rc|jSrrrs rget_origin_nodezIRNode.get_origin_node rOperation | NonecdSrrrs rget_defining_opzIRNode.get_defining_optrlist[Subgraph]cgS)z'Return subgraphs contained in this noderrs r get_subgraphszIRNode.get_subgraphs rc\t}|j}t|tr+|}|jrt|g}|D]}t |dr"|jr||j4tj j j di |jg}t|ts|D]C}tj j j |d}|r||D|S)N stack_trace postToPre)r@rr ExternKernelrrhasattrraddr- _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrrrpre_grad_nodes node_namers rget_stack_traceszIRNode.get_stack_tracessH)3 , dL ) ) 4..00K 4$k]33 6 6Dt]++ 60@ 6  !12222O)JNN#c$)R(( ".$77!/66I-QUU%t #6$((5556rTshorten Sequence[str]c^dt|dd}|r t|dkr |ddd}|s|gSg}|D]D}|d||dz }|d E|g|zS) Nzorigins=r@=z...zstack_traces = { })rrrappendsplit)rrrstack_trace_strrs r common_reprzIRNode.common_reprs;WT9b99;;  +s7||b(( "***G$$&& 9 0022 ( (K  " "#5 6 6 6 {0066 6O  " "3 ' ' ' 'y?**rlinesSequence[object] multilinecZt|t||z}ttt|}|r>## #"   44   $$LayoutcBtdt|d)Nz#get_layout() is not implemented by !rrrs rr'zIRNode.get_layouts"!"UT "U"U"UVVVr Layout | NonecN |S#t$rYdSwxYwr)r'rrs rmaybe_get_layoutzIRNode.maybe_get_layout9 ??$$ $"   44 r r<c*|Sr)r'rs rget_output_speczIRNode.get_output_specs   rOutputSpec | NonecN |S#t$rYdSwxYwr)rrrs rmaybe_get_output_speczIRNode.maybe_get_output_specs; '')) )"   44 r cPt|tS)z4True for single tensor output (excludes MultiOutput))rrr rs rhas_tensor_outputzIRNode.has_tensor_outputs$4466???rr^cBtdt|d)Nz!get_size() is not implemented by r r rs rr%zIRNode.get_sizes"!"Sd4jj"S"S"STTTrSequence[_IntLike] | NonecN |S#t$rYdSwxYwr)r%rrs rmaybe_get_sizezIRNode.maybe_get_size7 ==?? ""   44 r )_IntLike | sympy.Rel | Sequence[_IntLike]c*|Srr%rs rr`z IRNode.shapes}}rr cDt|Sr)rrr%rs r get_numelzIRNode.get_numelsT]]__---rctjjt j|dSNrrwr"r#statically_known_truerrQr"rs ris_zero_elementszIRNode.is_zero_elements2w55eht~~?O?OQR6S6STTTrr8c@tdt|)a) If the IRNode refers to data which has not been materialized (e.g., it is a Pointwise/Reduction that could potentially have more compute fused into it), realize the IRNode into physical memory, ending the possibility of fusing into it, but allowing, e.g., multiple users to access the data without having to recompute. Check StorageBox.realize for a particularly notable implementation. TODO(ezyang): I think, in principle, every IRNode should have an implementation of this, and most of the time no-op is OK, but you really do have to audit each IRNode for this, so for now, raise an error if it's not implemented. Note that some code in graph.py will catch this thrown error and suppress it with a warning. zrealize NYI on r rs rrealizezIRNode.realizes! ""@DJJ"@"@AAArwriterIndentedBuffer | Nonec@tdt|)Nzcodegen_reference NYI on r rr+s rcodegen_referencezIRNode.codegen_reference s!"Jd4jj"J"JKKKrtorch.device | NonecdSrrrs rr+zIRNode.get_devicerr torch.devicec6|}|J|Srr+rr!s rget_device_or_errorzIRNode.get_device_or_errors"""!!! rcdSNFrrs rhas_exceeded_max_readszIRNode.has_exceeded_max_readsur$Callable[[Sequence[Expr]], OpsValue]cDtt|jrrrrrs r make_loaderzIRNode.make_loader!$t**"5666r Callable[[Sequence[Expr]], Expr]cDtt|jrr=rs r make_indexerzIRNode.make_indexerr?rcDtt|jrr=rs rrTzIRNode.get_strider?rcN |S#t$rYdSwxYwr)rTrrs rrSzIRNode.maybe_get_stride"rr cDtt|jrr=rs rrzIRNode.get_name(r?rcN |S#t$rYdSwxYwr)rrrs rmaybe_get_namezIRNode.maybe_get_name+rr cp |tjjvS#t$rYdSwxYwr8)rrwr" graph_inputsrrs ris_input_bufferzIRNode.is_input_buffer1sA ==??ag&:: :"   55 s $' 55 thresholdcdSr8rrrKs rhas_large_inner_fnzIRNode.has_large_inner_fn7r:r graph_reuseusersrrPcdS)al Hint that this node's value will be reused. `users` estimates the reuse count. When `graph_reuse` is true, this represents graph fanout and can trigger fanout-based realization heuristics. When false, the reuse comes from loop-level indexing, such as a broadcast expand, and should not be treated as graph fanout. NrrrQrPs r mark_reusezIRNode.mark_reuse:srcdSrrrs r realize_hintzIRNode.realize_hintD rcDtt|jrr=rs rrzIRNode.unwrap_viewGr?rcDtt|jrr=rs r freeze_layoutzIRNode.freeze_layoutJr?rrr allow_paddingcDtt|jrr=rrr[s rfreeze_layout_with_stride_orderz&IRNode.freeze_layout_with_stride_orderM"$t**"5666rcDtt|jrr=rrs rfreeze_layout_with_fill_orderz$IRNode.freeze_layout_with_fill_orderRr?rrcDtt|jrr=rrs rfreeze_layout_with_same_orderz$IRNode.freeze_layout_with_same_orderUr?r exact_stridescDtt|jrr=rrfr[s r freeze_layout_with_exact_stridesz'IRNode.freeze_layout_with_exact_stridesXr_rdependencies.ReadWritescDtt|jrr=rs rget_read_writeszIRNode.get_read_writes]r?rOrderedSet[Dep]c4|jSrrlreadsrs rrzIRNode.get_reads`##%%++rcDt|Sr)rrrs r num_readszIRNode.num_readscs4>>##$$$rrcDtt|jrr=rs rget_storage_numelzIRNode.get_storage_numelfr?rrrcDtt|jrr=rrs rget_free_symbol_useszIRNode.get_free_symbol_usesir_rcDtt|jrr=rs rget_reduction_typezIRNode.get_reduction_typenr?rSequence[Expr]cDtt|jrr=rs rget_reduction_sizezIRNode.get_reduction_sizeqr?rcdSr8rrs r is_externzIRNode.is_externtr:rcdSr8rrs ris_no_opzIRNode.is_no_opwr:rr!cDtt|jrr=r5s rconstant_to_devicezIRNode.constant_to_devicezr?rcDtt|jrr=rs rget_mutation_nameszIRNode.get_mutation_names}r?rcDtt|jrr=rs rget_operation_namezIRNode.get_operation_namer?rcDtt|jrr=rs rget_inputs_that_alias_outputz#IRNode.get_inputs_that_alias_outputr?rcdSrrrs rr z IRNode.dtypes(+r)rrrr)rrrrrrrrrr)rrr1rrrrrrr)rrrrrrrrT)rrrr)TT)rrrrrrrrrr)rrrr )rrrr<)rrrrrr^)rr)rrrr rr8rr+r,rrrr0rr2rr;rr@rrrKrrrrQrrPrrrFrrr[rrrrrrrrr^rrrfr^r[rrrrrjrrmrrrrrrrrrr{r!r2rrrr)Qrrr__doc__r@rrr dataclassesfieldrrrrr staticmethod contextlibcontextmanagerrrrrrrrrrrrrrrr*rr'rrrrr%rpropertyr`r"r'r*r/r+r6r9r>rBrTrSrrGrJrNrTrVrrZr^rbrerirlrrsrurxrzr}rrrrrrrr rrrrr/s3=*,,>>>>044444 1{0e<<.rrrrs rrzOperation.get_read_namesrrrmc4|jSrrors rrzOperation.get_readsrqr list[Buffer]ctrrrs r get_outputszOperation.get_outputsrrrctSrr?rs rget_unbacked_symbol_defsz"Operation.get_unbacked_symbol_defs ||rFrctS)a When unbacked_only=True: Returns the unbacked symbols which are required to be in scope in order to successfully perform codegen for this buffer. For example, a buffer that corresponds to an extern kernel call that takes i0 as an argument would return {i0} here. This is used to generate necessary dependencies that ensure we actually bind i0 in codegen before you try to use it. Note that this is NOT transitive; in particular, if this buffer takes in as input another buffer with dynamic shape (e.g., (i0,)), we will not report it here, because you will already have a dependency on that buffer, which will eventually have a dependency on i0 if necessary. When unbacked_only=False: Similar to `unbacked_only=True` but including all free symbols instead of only free unbacked symbols. r?rws rrxzOperation.get_free_symbol_usess,||rrcdS)z Gets extra global memory size needed by this buffer. Some algorithms (e.g. group gemm) may require extra global memory in the generated code. rrrs rget_workspace_sizezOperation.get_workspace_sizes qrNrrr)rr)rrr)rr)rrrrrr)rrrrrrrrrrrrr)rrrrr+rrrrrrrrrlrrrrrrxrrrrrrs2222""""    ####$$$$''''""""----@@@@,,,,""""%*0rrc`eZdZUdZded<ded<ded<ded <ed d=d>dZd?dZd@dZeZ dAdZ dBdZ dCdZ dCdZ edDdZeejfdEd#ZedFd%ZdGd'Zed@d(ZdHdId,Zd=dJd.ZdKd2ZdLd4ZdMd5ZdNd7ZdOd8ZdPd:ZdQd<Z d)S)RLoopsz:Base class for pointwise and reduction loop-body IR nodes.r2r!rr Callable[..., Any]inner_fnr^rangesFrrrrctjgfd|jD|RS)Nc38K|]}t|VdSrr(rers rrz-Loops.get_free_symbol_uses..s. F FQq-00 F F F F F Fr)r@unionrinner_fn_free_symbolsrws `rrxzLoops.get_free_symbol_usessY"z||! F F F F$+ F F F  & &} 5 5    rnamesrrcdjjdtjgfd|DzdjgzS)N'c:g|]}|dt|S=)r)rrrs rrz!Loops._to_str..s1AAA4$..t,,..AAAr origin_node=)rr!rrr  inner_fn_strr)rrs` r_to_strz Loops._to_strs'DK$'''DJ!!##  BAAA5AAA  B 3d.223  4   rc,|dSNrrrs r__str__z Loops.__str__||K(((rr0c|jSrr!rs rr+zLoops.get_device {rrc|jSrrrs rrzLoops.get_origin_node rrc|jSrrrs rr%zLoops.get_size rrc|jSrrrs rget_pointwise_sizezLoops.get_pointwise_sizerrrrrrc|dd}|dd}||i|}|d||d|p|jt|S)Nrr)poprrrr)clsrrrtbrs rrz Loops.createsjj55 ZZ T * * C   ]K888 ["*; <<<"""rrrHr{c:fdt|DS)Ncdg|],\}}|dkrtjjnt|-SrI)rSZerorq)rnrrs rrz Loops._index..!sH   1FFEGLL(Fvq(Q(Q   r)r)rrs `r_indexz Loops._indexs6    !&))    rrYc~ttj}tj|5t jt dd5|j|| cdddcdddS#1swxYwYddddS#1swxYwYdSNallow_indexingT) rXrw MockHandlerset_ops_handlerrrr(r inner_fn_argsgetvalue)r opcounters rinner_fn_opcountzLoops.inner_fn_opcount&sH 11 i ( ( ( ( L)94 @ @ ( ( DM4--// 0 0%%''  ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( ( (s5B20B B2B B2!B "B22B69B6Sequence[Sequence[_IntLike]]c8||jfSr)r rrs rrzLoops.inner_fn_args0s DK((**rc^tjj|jg|RSr)rwKernelFormatterHandler ir_to_stringrrrs rrzLoops.inner_fn_str3s8'4 M  ..00    rNrKrc|d}tj}|)t|r tj}n$tj}nt |t sJt||}|j |kSr$) rJrealize_opcount_thresholdrKrealize_cpu_opcount_threshold"_realize_opcount_threshold_defaultrrmaxrnum_ops)rrKrs rrNzLoops.has_large_inner_fn9s  I$*$D! $ ,d|| V,2,P)),2,U))7== = == #<== $$&&.::rOrderedSet[Symbol]cd||j}t|j||SNr)r rrRr)rrrs rrzLoops.inner_fn_free_symbolsGs, DK((#DM5 VVVVrsymbolr"rc||j}t|}tj|5||dddn #1swxYwY|jSr)r rrUrwrrusages)rr#rhandlers rcollect_inner_fn_symbol_usagez#Loops.collect_inner_fn_symbol_usageLs DK((088 g & & ! ! MM%  ! ! ! ! ! ! ! ! ! ! ! ! ! ! !~sA  A$'A$rmctjtdd5|rXt |||jcdddSt ||jcdddS#1swxYwYdSr ) rrr(rzrTr>r%r}rprs rrzLoops.get_readsUs  \.*:D A A  &&(( *$$&&MMOO++--        +$$&&MMOO                  sA C 8CCCcNt|jSr)r@r read_buffersrs rrzLoops.get_read_namescs$//11>???rrcNt|jSr)rrr*rs rrszLoops.num_readsfs4((**7888rcBtdt|d)Nz+get_reduction_size() is not implemented by r r rs rr}zLoops.get_reduction_sizei(! G$t** G G G   rr8cBtdt|d)Nz+get_reduction_type() is not implemented by r r rs rrzzLoops.get_reduction_typenr-rrcBtdt|d)Nz+constant_to_device() is not implemented by r r r5s rrzLoops.constant_to_devicesr-rrr)rrrrrrrr)rrrrrr)rr^rrHrr{)rrYrrrrrrrr )r#r"rrrrrrrr)!rrrrrrbrxrr__repr__r+rr%r classmethodrrrHINDEXr rarrrrNrr'rrrsr}rzrrrrrrsbDD    G$$$)    %$      ))))H     # # #[ #:>*    \ (((](++++   ] ; ; ; ; ;WWWWW     @@@@9999              rrrExpr | Sequence[Expr]r rrvc~|jr"tjtd|Stjd|S)Nnanr)is_floating_pointruconstantfloat)rr s r nop_loader_fnr;ys7 &|E%LL%000|Au%%%rcBeZdZddZddZeZddZdd ZddZddZ dS) Pointwiserr;cn|rtt|jS|jSNr)r'r r;r rrs rr>zPointwise.make_loaders3  " " <= ;;; ;}rrc,|dSrrrs rrzPointwise.__str__rrr^cgSrrrs rr}zPointwise.get_reduction_sizes rr8cdSrrrs rrzzPointwise.get_reduction_typerr output_nameindexer!Callable[[Sequence[Expr]], Never]varsr{rc~|}tj|pd||||SNunnamed)r>rustorerrCrDrFloaders r store_outputzPointwise.store_outputs> !!##y1 774==&&,,OOOrr!r2rc|}tjtd||}t ||j||jSFMove this to a given device. Requires that all reads are to constants.override_devicer!r rr)r>rrConstantBufferr=r rrr!rLs rrzPointwise.constant_to_devicesX!!##Hn.?HHPP*;     rNrrrr)rCr8rDrErFr{rrr) rrrr>rr2r}rzrMrrrrr=r=s))))HPPPP       rr=c8eZdZUded<dZded<dd ZddZdS)Scatterr@output_indexerNr[ scatter_moder!r2rrc|}tjtd||}t ||j||j|j|jS)rPrQ)r!r rrrWrX) r>rrrSrVr rrWrXrTs rrzScatter.constant_to_devicesd!!##Hn.?HHPP*;.*     rrCr8rDrErFr{rc|}|d}tj|||||||jS)NrI)mode)r>rurJrWrXrKs rrMzScatter.store_outputsh !!##  #Ky  GD''-- . . F4LL"     rr)rCr8rDrErFr{rr)rrrrrXrrMrrrrVrVs[4444"L""""           rrV logical_ormaximumminimummulr bitwise_xor)anyrminprodsumdotxor_sumz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNreduction_typearg_break_ties_leftCallable[..., object]ctvr tSdvr d fd }|Sdkrd d }|Std )Nargmaxargmin argmax_value argmin_valueargmax_with_valueargmin_with_valueatuple[object, object]brtuple[OpsValue, OpsValue]c|\}}|\}} dvrtj||}ntj||}tj||}t rztj||}tj||} tj|tj|| }tj|tj|| } rtj||ntj||} tj|tj|| }tj|||tj|||fS)N)rnrprr) rultgteqr0ner\ logical_andwhere)rsrua_valuea_indexb_valueb_indexmaskequala_isnanb_isnantierir rhs rargmax_combine_fnz3get_reduction_combine_fn..argmax_combine_fnsA ! GW GW!PPPvgw//vgw//F7G,,Ee$$ Q&'22&'22~dCF7G,D,DEEucogw.O.OPP'.w(((VGW--  >$s(C(CDDD $11 $11 rwelford_combine#tuple[OpsValue, OpsValue, OpsValue]c >|\}}}|\}}}tjtjtj|tj||tjdt j||z }||z} || z } ||| zz||z||z|z| zz| fS)Ng)rur}r|isinfrzr9r-float32) rsrua_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_ws rwelford_combine_fnz4get_reduction_combine_fn..welford_combine_fns&' "FD(%& "FD( I & 1 136&&3I3IJJ S%-00E "H,J :-I**t eemh6BB rzunknown reduction_type=)rsrtrurtrrv)rsrrurrr)rgr)rhr rirrs``` rget_reduction_combine_fnrs---#N33           :! , , ,    ."!""LN"L"LMMMrceZdZUded<ded<ded<ded<dUd ZeZeddVdWfd ZdXdZdYdZ dZdZ d[dZ d\d Z dVdWd!Z d]d%Ze d^d_d1Zed`d4Zeejd&fdad7Zedbd:Zedbd;Zedcd?ZedddAZe d^dedGZedfdNZedgdRZe d^dhdSZedidTZxZS)j Reductionr^reduction_rangesrZrhr src_dtyper^reduction_hintrrc,|dS)N)rrrhrrs rrzReduction.__str__,s||LMMMrFrrr cttjfd|jDzS)Nc38K|]}t|VdSrr(rs rrz1Reduction.get_free_symbol_uses..4s. P PQq-00 P P P P P Pr)superrxr@rrrr __class__s `rrxzReduction.get_free_symbol_uses1sNww++M::=OZ\\=O P P P P$:O P P P>   rc|jSr)rrs rr}zReduction.get_reduction_size7s $$rr8c|jSr)rhrs rrzzReduction.get_reduction_type: ""rrCrDrErFr{reduction_varsSequence[Symbol]rc tj|j|j|j|||}tj|pd|||dSrH)ru reductionr rrhrstore_reduction)rrCrDrFrr1s rrzReduction.store_reduction=s` J N   MM$ / /    K49ggdmmUKKKKKrrcTt|jt|jzSr)rrrrs r index_lengthzReduction.index_lengthLs"4;#d&;"<"<<rrrSrr rrrhrr^DEFAULTrTs rrzReduction.constant_to_device[so!!##Hn.?HHPP*;!2.n(0    rN dst_dtyperCallable[_P, OpsValue]rReductionType | Literal['scan']reduction_numelr input_nodertuple[ReductionHint, _IntLike]c X|t|g} tjj| st jdfStjj|} t|} tjj| } tj|tj o0t| p tj || |} d}|dkp| o||vo tj}|dkrt jdfSt!j|}|j}d}|rMt'jtj j|d}t'jtj j|d}ndd }|}| dkr|| | }|dkrt j|fS|t/|t0rt3jt6dd5t9|\}}dddn #1swxYwY|i|gtjjt||z}| |kr-t<d|||||t jdfSt j|fS| |ks | |dzdzkrt jdfStA||||||dkr|nd|t j}dd}||\}}|r||\}}tC|dkrt jdfStEj#|$|%\\}}}d} d}!|D]}"tjj&|"|}#tjj'|#|tQ|)}$tUd|$D}%|%r| dz } |!dz }!|!| krt j|| | fSt j+|| | fS)zd Choose the reduction hint and split count from shape and input stride information. rIrlscanre T)inner_reductionFreduction_numel_hintr numel_hintrcdSNrIr)rrs rinner_reduction_splitsz4Reduction.num_splits..inner_reduction_splitss qrNrzUse previous IRNode's range and reduction_ranges instead of split. current ranges: %s, current reduction ranges: %s, current split: %d, new ranges: %s, new reduction ranges: %srRrrdrrrtuple[Sequence[Expr], bool]c |}|Jtdt||||}|}|jJd|jD}tj|| \\}}}d|D}g}g}d} t|j dD]} | j j tfd|D} | otfd |D} | r|| j n| r|| j | s| rz| jt$jjvrbt$jj| j} t+| jd d}| t+| jd d|krd } |p|| fS) Nr!r rrrqrpcpg|]3}t|tt|tj1|4Srrr rr rvars rrzBReduction.num_splits..get_read_indices..sNc4((2.get_read_indices..sNc4((2z@Reduction.num_splits..get_read_indices..safrkeyc3 K|]}|vV dSrrrrr5s rrzAReduction.num_splits..get_read_indices.. s('R'R|(;'R'R'R'R'R'Rrc3 K|]}|vV dSrrrs rrzAReduction.num_splits..get_read_indices.. sANN,/C<'NNNNNNrrT)r+rr(r*r%rl range_varsrKindex_vars_squeezer}sortedrprr5rrarrrwr"name_to_bufferrrq decide_layout)rr!cb read_writesrr rindicesbroadcasted_reduction_indiceschangedmdis_full_size_readis_broadcasted_reduction_readbuforiginal_strider5s @rget_read_indicesz.Reduction.num_splits..get_read_indicesst\\^^F%%%%!++-- B,,..K)555&1J &2%D a2244&& " Q)N G,. )G[.4D4DEEE + +!x4 $''R'R'R'Rz'R'R'R$R$R!5F0E1#NNNN3ANNNKK-%CNN28,,,,2C188BBB$+(E+w!'"888g4RW=*1#*h*M*M))+++"3:x>>/QQ&*G;;WD Drrc3.K|]}|dkp|dkVdS)rrINrrs rrz'Reduction.num_splits..1s.99AQ!a%999999r)rrrrrr)rrrr),rrrwr"r#all_unbacked_explicitly_hintedr^rr$ has_featurerLREDUCE_TO_SINGLE_ELEMENTrIchoices should_use_cooperative_reductionrJsplit_reductionsr]rmulti_processor_count functoolsr reduction_split_factorINNERrrrrr(rS!replace_backed_symbols_with_hintslogrrrrKrr%r}simplify_with_ranges stride_hintsrkeysrOUTER)&r!rrrrrrhrrexprsrnumelrshould_reduce_to_single_elementarg_reduction_types should_splitpropsnum_smmin_elements_per_threadrouter_reduction_splitsr new_rangesnew_reduction_rangesextracted_numel_hintrrrrr rranges1 num_outer num_innerrjrhouters& r num_splitszReduction.num_splitsjs9"!-"7"78w>>uEE , (!+ + w/AA/RRf%%W%77>> +,'*=*= N;+ + + &!! ! y99 ( &/ / / (&99 ('  U " " (!+ + '//,"$  <@I@Q 0&$AAA "AJ@Q 0&%AAA " "      &< " ??**+?LLEzz$*E11%*Z*K*K%\.2BDIIHH>5(0 5 E5 E5 E5 En,+A..  -))!,,JGQ w<<1   (!+ +'3'F JJLL!..00( ( $NW    A 55aAAAg&33>4 #7#7G 9999999E Q Q y  &(>(>$j)) !&(>(>$j)) s/HHH.fn..NsCHUF++rc,g|]}t|Sr)rrrs rrz>Reduction._unroll_reduction_fn..fn..Qs===q%((===r)rreduce itertoolsproduct)r combine_fnrr s`rrz*Reduction._unroll_reduction_fn..fnKsa#"+"3==,<===# r)rnrmrprorrrqrrvcd|D}||}t|tr|S|tj|tjfS)Nc6g|]}tj|Srrexpandrrs rrzDReduction._unroll_reduction_fn..value_fn..gs :::a%,q//:::r)rrru index_exprr-int64)rrr1 flatten_indexrs rr z0Reduction._unroll_reduction_fn..value_fndsh;:6::: //eU++! LN==#8#8%+FFr)rrrq)rproc&|dSr$rrrs rrz0Reduction._unroll_reduction_fn..ssRRYYq\rc&|dSrrrs rrz0Reduction._unroll_reduction_fn..tsE1rN)rr^rr)rr^rr^rrv)rwr"r# guard_int_seqr_fixed_indexerr(r))rrrhrrrrr s`` @@@@r_unroll_reduction_fnzReduction._unroll_reduction_fn?s7+99:JKK-niHH             + 112BCCM       !KKK !AAA11111---- -HIrrrc Ntjjt dkrudfd } | d| d| d| ddvs Jd dfd } t ||| t|Sdkr0dvrdfd } ndfd } t || |Sttr~ttj krat |dkst|jr:dkr4t ||||S||||| \} }d fd }||}|t"jkr| }|dkro| Jt'jt*dd5t-| \}}dddn #1swxYwY|J|J||||||| S|dkr||||||| }d}tjjr$t|t6rd!d}||}|rct|jt:sJt|j|jjd|_|_ ||_!|_"|St6t;||||}|S)"zy Create a reduction node. May split the reduction to multiple layers to expose more parallelism. rvalrrbool | float | intc<tjkrt|Sjr;t|tsJt |t |St|tsJt |t|Sr) r-rr8rrrr:rr)r"rs rpy_cnstz!Reduction.create..py_cnsts **99$0$%c=99DD499DD9 ::%%c;77BBcBB7s88OrrI)rdrfrcra* not supported for zero-dimension tensors!rrrvc:tjSrrur9)rrrhrtypes_to_initss rconst_fnz"Reduction.create..const_fns|ON$CYOOOrrR)rnrmc.tjdSr$r()rrs rrzReduction.create..fns<9555rc4dD}||S)Nc0g|]}tjjSrrrr r s rrz0Reduction.create..fn..&N&N&Nuw|&N&N&Nrr)rreduction_indexrrs rrzReduction.create..fns)&N&N=M&N&N&NO#8E?;;;rrercjtr|S|dkrt|tjS|Sr)rrrJ min_num_split)rrs r_maybe_increase_splitz/Reduction.create.._maybe_increase_splits9/**  qyy5&"6777 rrRNrTcur_noderComputedBuffer | NonecL|}t|dkrdStt|}|tjjvrdStjj|}t|tsdS|j J|Sr) rrnextiterrwr"rrrrprz)r4 read_namesbufnamers r_find_split_reductionz/Reduction.create.._find_split_reductions"*!8!8!:!:J:!++#t"4 #3#344Gag&<<<#t'09C%c>::$#t86688DDDJrr)r"rrr#)rrrrv)rrrr)r4rrr5)#rwr"r#simplifyrrr=rrrr!rrJunroll_reductions_thresholdrnrr rr^rrrr(rS!create_multilayer_existing_rangescreate_multilayerrDmix_order_reductionrrprr _split_size_original_inner_fn_original_ranges_original_reduction_ranges)rr!rrrrrrhrrr%r*rhintrr3rrrsplit_reductionr;rr)s ` ` `` @@rrzReduction.createys$'*33MBR4S4STT a   $ $ $ $ $ $wqzz"71:: wqzz O"_444!MMM544 P P P P P P P P##!F|| $  a  !5556666666 <<<<<<<##YF$   0 0 O$$v'IIIv&&!++vfk/B/B+%'' ##11.  $ nn            e      &%e,, ]2 2 2!N B;;)))n.>EE  3V440 0               )))'33388 $   QYY''   C#O}0 =ZY5O5O =$#8"7"<"< N"/"6 BBO0113B/>.B.STU.V+5=2390=M:J !!1-#-       s$HH Hr _NumLike | Sequence[_NumLike]c |dvrHt|rtdSt|rdStj|jS|dvrHt|rtdSt|rdStj|jSt|rdnd}t|rdnd}||||||||f|||ftd|fd |S) N)rrmrorq-infF)rbrnrprrinfTrrI)rdrcrerfrawelford_reduceronline_softmax_reduce)r0r:r.r-iinforbr)rhr zeroones rdefault_accumulatorzReduction.default_accumulatorKs S S Se$$ .V}}$!%(( .u{5))-- S S Se$$ .U||#!%(( .t{5))--(//6uuQ&u--4dd1#T40 $dD1&+FmmT%:      rcH|dkrdSt||S)NrKr)rrPrhr s r default_valuezReduction.default_valueks+ - - -1,,^UCCCrrrrc|dkr|S|dkr"|dkr|tjkr tjS|dkr"|dkr|tjkr tjS|S)NrRi)r^r OUTER_TINY)rrrs r_multilayer_second_step_hintz&Reduction._multilayer_second_step_hintssm B;;! ! C<.wrapper_fn..bodys!vi'););<<.wrapper_fns"1 _*/ 'Y ?2_DG = = = = = = = = = -o>> vN7K88N?K@@z$g666tvv r)rrr0rrrv) rZViewdynamic_reshape_indexerrwr"r#r&rrQrF) rrLrrrr[r\r dense_indexrgrfrs ` ` `` @@r_multilayer_wrap_loaderz!Reduction._multilayer_wrap_loaders>> Z  .. /  (>> HS%00! 4 4              (r4Callable[[Sequence[Expr], Sequence[Expr]], OpsValue]original_rangesoriginal_reduction_rangesrSequence[Integer]rctdDs Jdt|t|t|zd fd }|S) Nc3"K|] }|dkV dS)rINrrrs rrzDReduction._multilayer_wrap_loader_existing_ranges..s&33a16333333rz8Only enabled for numel_hint == 1, found original_ranges= merged_indexr{new_reduction_indexrrvc |dt}|td}|t|t|zSr)rr)rsrt original_idxr`rLrmrs rrgzEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fnso((>#o*>*>(>?L$S%9%9%;%;.intermediate_fns!'&'A'A'ABB BrNr)rr^r0r^rrv)r-float16bfloat16r:rrr*r>rwr"r#r$rrrXrr)rr!rrrgrmrnrrrhrrintermediate_dtype intermediater}rr|s @rcreate_multilayer_helperz"Reduction.create_multilayer_helpers\0 ??? I  !''           *6688 C C C C C C W%77 o8V8VWW 99 :~  *-Cs?/C/C-C"DDDDD (&!+C,@,@,B,B!C-#-      rc t|} t| |dz z|} |||} |||| || | | }|||||||g||| g||| S)rzrI)rrrCrSrkr)rr!rrrrrrhrrrrr[r\rgs rr?zReduction.create_multilayer,s&((899o;UCC ##NI>>00         ++       f e  L      rc z||||||} ||||| ||g|||| d| S)rzrR)rwr) rr!rrrrmrnrrrhrrgs rr>z+Reduction.create_multilayer_existing_rangesZsm$@@   %    ++      % +o + +      rrrr1rr) rCr8rDrErFr{rrrrrrrrr)r!r2rrrrrrrr^rr^rhrrr rrrr) rrrr^rhrrrrr)r!r2rrrrrrrr^rr^rhrZrr^rrrrrhrr rrrG)rrrrrr^rr^)rrrrrr)rLrrr^rrrrr[rr\rGrrrrj) rLrlrmr{rnr{rrorrorrl)r!r2rrrrrgrrmr{rnr{rrrrxrhrZrrrr^rr)r!r2rrrrrrrr^rr^rhrZrrrr^rrrr)r!r2rrrrrrrmr^rnr^rrxrrxrhrZrr^rr) rrrrrr2rbrxr}rzrrrrrrrr r3r^rrrPrSrXrZrkrwrr?r> __classcell__rs@rrr$s((((!!!!!!!!NNNNHK((      )( %%%%#### L L L L====           %)RRRR\Rh777\7r)6(=$(NNNN[N`\>DDD\D\ [>%)(((([(T[8= = = [= ~%)+ + + + [+ Z$ $ $ [$ $ $ $ $ rrrrrrtr $Callable[[Sequence[_IntLike]], Expr]cdfd }|S)1A closure containing math to read a given elementrr^rr c t|tksJt|tksJ}t|D]\}}}|dkr|||zz}|Sr)rr)rresultrstszrtrrs rrDz_fixed_indexer..indexers!c%jjCKK&?&?&??5zzSYY&&&&ufd33 + +KCRQww#(* r)rr^rr r)rrrtrDs``` rrrs4 Nr INNER_FN_TYc4eZdZUded<dfd ZddZxZS)MultiOutputReductionr output_indexr!r2rr inner_fns#INNER_FN_TY | Sequence[INNER_FN_TY]rrorrhrZrrr^c trftdkr d} nd fd } t||| ||||| | |_dS) NrIrrr{ reduction_idxrtuple[OpsValue, ...]c@tfdDS)Nc30K|]}|VdSrr)rrrrs rrz@MultiOutputReduction.__init__..loader..s/HHRR]33HHHHHHrr)rrrs``rrLz-MultiOutputReduction.__init__..loaders-HHHHHiHHHHHHrr)rr{rr{rr)callablerr__init__r) rr!rrrrrhrrrrLrs ` rrzMultiOutputReduction.__init__s I   %" I y>>Q  q\FF I I I I I I -))  )rrCr8rDrErFr{rrrrc 8tj|j|j|j|||}t |ttfsJt|||j }tj |pd|||SrH) rurr rrhrrrrrrr)rrCrDrFrrr1s rrz$MultiOutputReduction.store_reductions J N   MM$ / /    &5$-00>>$v,,>>0t()";#;)WWT]]ERRRr)r!r2rrrrrrorrorhrZrrrr^rr) rCr8rDrErFr{rrrr)rrrrrrrrs@rrrsh#)#)#)#)#)#)JSSSSSSSSrrc:eZdZdZeejdfddZdS) ArgReductionzJMulti-output arg reduction that returns both the selected value and index.Nr!r2rrrrrrr{rrhrZrr^rrrSequence[TensorBox]c 4dvsJtjjt } | dkrt d| dkrkdfd } dd } t || t t tj | t fSt| trt| tjkrt dkst!jrYt$dfd t)fdt+|tj fDSt$|| | \} }|dks Jdt.jkr| t)fdt+|tj fD}|D]}||S)N)rqrrrr&rIrr{rrvcrdD}||}t|tr|dS|S)Nc0g|]}tjjSrr.r s rrz9ArgReduction.create..value_fn..s"J"J"JA57<"J"J"Jrrrr)rr0r1rrs rr z%ArgReduction.create..value_fnsH"J"J9I"J"J"J 88eU++$ 8O rc@tjdtjSr$)rur9r-rrs rindex_fnz%ArgReduction.create..index_fns|Au{333rrRrrcX|}t|tsJ||Srr)rrr unrolled_fns rprojectz$ArgReduction.create..project s3$U++!&%00000l++rc 3K|]A\}}t|t|tVBdS))rrRN)r=rr r)rrr r!rrs rrz&ArgReduction.create.. sq(L%   !$W<HHH<< !rz.arg reductions do not support split reductionsc3tK|]2\}}t| | V3dSrr) rrr rr!rrrrrhrs rrz&ArgReduction.create..+ sw  $ e   $""           rrr{rrv)rr{rrrrv)rwr"r#r<rrAssertionErrorr=rrr-rrr!rrJr=rnrrr rrrr^rr*)rr!rrrrrrhrrrr rrErresultsrrrs`` `````` @@rrzArgReduction.creates!KKKKK'*33MBR4S4STT a   !MMM  a          4 4 4 4  !#%<< !   !+%<< !    0 0 O$$v'IIIv&&!++vfk/B/B+#88*NIK , , , , , , ,5i5M+N+N  **            ezzzKzzz ]2 2 2!N           (1)U[1I'J'J        F NN    r)r!r2rrrrrrrr{rr{rhrZrr^rrrr)rrrrr3r^rrrrrrrsOTT)6(=$(gggg[gggrrc~eZdZdZeejdfddZeejdfddZe dddZ dS)OnlineSoftmaxReductionzAMulti-output reduction that computes softmax max/sum in one pass.Nr!r2rrrrrrr{r num_outputrrr^rrrrc tfdt|D} | D]} | | S)Nc3vK|]3}ttd| V4dS)rLN)rrr) r output_idxr!rrrrrrs rrz:OnlineSoftmaxReduction._create_no_split..P sq     $$+"          r)rrr*) rr!rrrrrrrrrr0s `````` ` r_create_no_splitz'OnlineSoftmaxReduction._create_no_splitC s          $J//        A IIKKKKrc Vtjjt |} t ||||||d| |  \} } |tjkr| }| dkr| |||||||| || S| ||||||||| S)NrL)rhrrrI) rwr"r#r<rrrrr^rr?r) rr!rrrrrrrrrrErs rrzOnlineSoftmaxReduction.created s'*33MBR4S4STT**      2+!+   e ]2 2 2!N 199((    ##            rrrc .t|t|dz z||| } t|g| t jjtj t|d d fd } |tj tjfvr|n tj} ||| | g||g|| }|D]}|t%d |Ddfd }t jjt|}|||| } |||| |||g|| S)NrIrrrr0rrvc |\}|^} |z|zd  fd }tjd} rt }tjtj|tj |}tj||t dtj||tjdfS||fS)Nrrvc0gSrr)rrr`rsrrazJOnlineSoftmaxReduction.create_multilayer..wrapper_fn..body s!x 77G9+=+=>>>rrIrIrrb)rur9rhrxrrcr:r})rr0rdrarOrerrr`r[rrfrrrs @@rrgz.wrapper_fn s"1 _*/ 'Y ?2_DG ? ? ? ? ? ? ? ? ?,q),,C #-o>> vN7K88N?K@@ JtT5==99IdCa)C)CDD tvvs{"rc3>K|]}|VdSr)r>rs rrz;OnlineSoftmaxReduction.create_multilayer.. s*$L$LQ]]__$L$L$L$L$L$Lrr^rc@tfdDS)Nc36K|]}|gVdSrr)rrLrr0s rrzTOnlineSoftmaxReduction.create_multilayer..intermediate_fn.. sI7=11122rr)rr0intermediate_loaderss``rr}zAOnlineSoftmaxReduction.create_multilayer..intermediate_fn sDAU r)rrr0rrrv)rr^r0r^rr)rrrCrZrhrirwr"r#r&rrQrFr-r~rr:rr*rr$rX)rr!rrrrrrrrrrjrgr intermediatesrr}rr[rrfrrs `` @@@@@rr?z(OnlineSoftmaxReduction.create_multilayer s ((899o;UCC >> Z  .. /  (>> HS%00! 4 4    # # # # # # # # # # #4 ??? I  ,,     f e  L       A IIKKKK$$L$Lm$L$L$LLL      W%77 f8M8MNN 99 :~  ##      G     r)r!r2rrrrrrrr{rr{rrrr^rrrrr)r!r2rrrrrrrr{rr{rrrrrr^rrrr) rrrrr3r^rrrr?rrrrr@ sKK)6(=$([@)6(=$(1 1 1 1 [1 f%)Z Z Z Z [Z Z Z rrcdeZdZeejfddZeddZeddZ dS)WelfordReductionr!r2r rrSequence[Callable[..., Any]]rrxrrhrZrr^rrc dvsJtjjr Jdtjjt}dfd } |dkr&| d} | d} | d} | | | fS|d krPdfd d kr&d| d| d fStfdDSt d|\} }tj kr| |d kr| |SfdtdD}|D]}||S)N)rKrzGwelford reduction usage is explicitly disabled, please check you configr"rrrcjdfd }t|tS)Nrr{rrvc.tjSrr()rr r"s rrz8WelfordReduction.create..const..inner_fn s|rrRrr{rrvr=rr)r"rr!r rs` rconstz&WelfordReduction.create..const sX        ##!F|| $ rrrIrLrlcjdfd }t|tS)Nrr{rrvc4dD}||S)Nc0g|]}tjjSrr.r s rrzKWelfordReduction.create..copy..inner_fn..# r/rr)rr0rLrs rrz7WelfordReduction.create..copy..inner_fn" s)&N&N=M&N&N&NO!6#777rrRrr)rLrr!r rrs` rcopyz%WelfordReduction.create..copy sX8888888!''!%<< (rrKc3.K|]}|VdSrr)rrrs rrz*WelfordReduction.create..0 s+::"TT"XX::::::r)rhrcng|]1}tt| 2Sr)rrr) rrr!r rrrrrhs rrz+WelfordReduction.create.._ sc       $""       rr)r"rrr)rLrlrr)rJmtiadisable_welford_reductionrwr"r#r<rrrrrr^rr?rr*)rr!r rrrrhrrrmeanm2weightrErrr0rs ``````` @rrzWelfordReduction.create s]!FFFFF;8  U  8'*33MBR4S4STT         a  588DqBU1XXFV# # a           !111tIaL))5588UU1XX==:::: ::::::& **    aL  )++   e ]2 2 2!N 199((              $Ahh      A IIKKKKrrrGcdS)N)rrrrrRs rrSzWelfordReduction.default_values s yrrrc Lttjjt jtd } | rP|dkrJdfd } ||dt| d t| d f|d| Std z zt |tfd|Dg|g||} | D]} | ddtjjt|} | |}t |tfd| D|gd|S)rzrrrr{rr1rrrvc.tj|Srr()rrr1r s rr9z4WelfordReduction.create_multilayer..constant s|E5111rr5rI)r!r rrrrhrrc 3NK|]}|dV dS)rr\N)rk)rrLr[rrrrs rrz5WelfordReduction.create_multilayer.. s`  ++$# ,      rrr0rLr;c |g||Srr)rr0rLs rintermediate_loader_fnzBWelfordReduction.create_multilayer..intermediate_loader_fn s 64E4O455 5rc3^K|]'}t|V(dS))rLN)r r>)rrrs rrz5WelfordReduction.create_multilayer.. sM.q}}GGGr)rr{rr{r1rrrv)rr{r0r{rLr;rrv)rrrwr"r#r&rrQrFr?r rCrrrr*r$rX)rr!r rrrrhrrrfr9rrrr[rrs` ` ` ` @@@rr?z"WelfordReduction.create_multilayery sS ((899(>> HS%00! 4 4     +<<< 2 2 2 2 2 2 ((aLHA...HA... !10-)   o;UCC (//           (      f e  L  #  &  A IIKKKK 6 6 6 6W%77 f8M8MNN 99 :~   &&   &    G     rN)r!r2r rrrrrxrrxrhrZrr^rrr)r!r2r rrrrrxrrxrhrZrrrr^rr) rrrr3r^rrrrSr?rrrrr s)6(=yyyy[yv\ \ \ \ [\ \ \ rrc(eZdZUded<ded<ded<ded<ded <d ed <d ed <ded<edd=d>fd Zd?fd Zd@d!ZdAd"ZdBd#Z dBd$Z dBd%Z dCd&Z dDd(Z d=d>d)Zeejfd*d+dEd3ZedFd<ZxZS)GScanrx scan_rangesr=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rr^rrrtuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]rFrrrr cttjfd|jDztjfd|jDzS)Nc38K|]}t|VdSrr(rs rrz,Scan.get_free_symbol_uses.. .OO"1m44OOOOOOrc38K|]}t|VdSrr(rs rrz,Scan.get_free_symbol_uses.. .HH"1m44HHHHHHr)rrxr@rrrrs `rrxzScan.get_free_symbol_uses s GG ( ( 7 7 jll OOOOd>NOOO !jll HHHHdiHHH   rrct|jt|jzt|jksJt dSr)rrrrrrrrs rrzScan.__post_init__ R4;#d&6"7"773ty>>IIII rrCr8rD%Callable[[Sequence[_IntLike]], Never]rFr{ scan_varsrrc|||tfd|jD}tj|j|j|}tj|pd|||jS)Nc3.K|]}|VdSrrrrrs rrz'Scan.store_reduction.. +DDxx}}DDDDDDrrI) rrrrurrrrJr)rrCrDrFrrrrs @rrzScan.store_reduction sll4++DDDDT^DDDDD$+t??y  $9ggcllF4;L4M   rcdS)Ncustomrrs rrzzScan.get_reduction_type sxrc|jSr)rrs rr}zScan.get_reduction_size rrc|jSrrrs rr%z Scan.get_size yrc|jSrrrs rrzScan.get_pointwise_size rrcTt|jt|jzSr)rrrrs rrzScan.index_length "4;#d&6"7"777rrc||j}||jtj}|||}|fSr)r rrrHrrrrrrs rrzScan.inner_fn_args G DK((T-t}==ll5&))v rc||j}||jtj}|||}t |j||Sr")r rrrHrrrRrrrrrrs rrzScan.inner_fn_free_symbols W DK((T-t}==ll5&))#DM3mTTTTrT)can_fallback_to_atenr!r2+tuple[Callable[[Sequence[Expr]], Any], ...]axisrrSequence[TensorBox | None]c  gddzdgtjtjsdgt zSt dkr=tjtjsdgt zStjj} | t} t t ksJ| tj | dr+fdtt DS|dd| \} t| dkrYt jjdupt&o t(dkot dk} | s|rdgt zSd} nt*d fd  f d tt D}|D]}||S)NrIcdg|],}t||-SrRr=rrrr!rrrs rrzScan.create..E V!   ! .&|4 !rr)r!r rrpointwise_rangesrr scan_numelz3.3.0rr{ scan_indexrrct|tksJt|tksJg|d||dSrr)rr rr rs rrzScan.create..reindexh gz??c+&6&66666u::%5!6!66666>U5D5\>J>tuu> >rc g|];}t d||   |d .m s   $#    ! .!&|4'+ +)##1!-     r)rr{r r{rr)rwr"rrLSCANrTUPLE_REDUCTIONr#r<rrr&rrWrrrr-versionhip has_tritontriton_version SplitScanr*)rr!rrrrrrrrr#r rsupports_splitrrr rrrs ``````` ` @@@@rrz Scan.create$ s=T%4%[<4q +;<Dzl w""6>+>?? (6CKK' ' v;;??17#6#6 N2$ $ ?6CKK' '7#&&}['A'ABB 6{{c)nn,,,,  ) )%(:q*A*A B B %*#f++$6$6 &)^^)q\-#!!&4 & & "  >> !T)Wj.V^w=V%v;;!# " &'# 6CKK//!"JJ%  ? ? ? ? ? ? ? ?               $!&c&kk 2 2%   *  F NN    rr rrr;r r r rc Xd fd } t|||| ||d|S) Nrr{rrrvcFg|d||dSrr)rrrrs rrgz#Scan.num_splits..wrapper_fn s38Fc%4%jF=F3tuu:FGG Grr)r!rrrrrrhr)rr{rr{rrv)rr) rr!r rrr rrr rgs `` rrzScan.num_splits se H H H H H H H###(!&$   rrr1r) rCr8rDrrFr{rrrrrrrr0)r!r2rrrrrrxrrrrrr^rrrrrr)r!r2r rrr;rrr rxrrxrrr r rr)rrrrrbrxrrrzr}r%rrrrr3r^rrrrrs@rrr sMMMMSSSS!!!!####----F##       $#                 8888 UUUUU )6(=`&*`````[`D   [     rrceZdZdS)rN)rrrrrrrr sDrrceZdZUded<ded<ded<ded<ded <d ed <d ed <ded<ded<edd2d3fd Zd4fd Zd5d Zd6d!Zd7d#Z d7d$Z d8d&Z d9d'Z d:d)Z d2d3d*Zeejfd;d1ZxZS)<Sortrx sort_rangesrz:Callable[[Sequence[Expr], Sequence[Expr]], Sequence[Expr]]rr^rrrrrrrrstable descendingFrrr cttjfd|jDztjfd|jDzS)Nc38K|]}t|VdSrr(rs rrz,Sort.get_free_symbol_uses.. rrc38K|]}t|VdSrr(rs rrz,Sort.get_free_symbol_uses.. rr)rrxr@rr rrs `rrxzSort.get_free_symbol_uses s GG ( ( 7 7 jll OOOOd>NOOO !jll HHHHdiHHH   rrct|jt|jzt|jksJt dSr)rrr rrrrs rrzSort.__post_init__ rrrCr8rDr@rFr{rrc|||tfd|jD}tj|j||j|j}tj|pd|||j S)Nc3.K|]}|VdSrrrs rrz'Sort.store_reduction.. rrrI) rrrrusortrr!r"rJr)rrCrDrFrrrrs @rrzSort.store_reduction sll400DDDDT^DDDDD$+vt{DOLLy  $9ggcllF4;L4M   rcdS)Nr)rrs rrzzSort.get_reduction_type svrroc|jSr)r rs rr}zSort.get_reduction_size rrc|jSrrrs rr%z Sort.get_size rrr^c|jSrrrs rrzSort.get_pointwise_size rrcTt|jt|jzSr)rrr rs rrzSort.index_length rrrc||j}||jtj}|||}|fSr)r rr rHrrrs rrzSort.inner_fn_args rrc||j}||jtj}|||}t |j||Sr")r rr rHrrrRrrs rrzSort.inner_fn_free_symbols rrr!r2'tuple[Callable[[list[Expr]], Any], ...]rrrc  gddzdgtjtjsdgt zStjj} | t} tj j rtj j } n;d} tj j o'| tj| | } | sdgt zSt t ksJ| tj| dr+fdt!t DSd fd  f d t!t D}|D]}||S) NrIrUcdg|],}t||-Srrr s rrzSort.create.. r rrr{ sort_indexrrct|tksJt|tksJg|d||dSrr)rr4rr r s rrzSort.create..reindex' rrc g|]@}ttd||  | d AS)) r!r rrrrrr rrrr!r"r)rrr) rrr"r!rrrr rrrr r!s rrzSort.create.., s   &%   ! .!&|4'+ +##1!-!)     r)rr{r4r{rr)rwr"rrLSORTrr#r<rrrJrDdecompose_sort_opspersistent_reductionsr&rrWrr*)rr!rrrrr!r"rrr# sort_numelis_persistent_kernel max_rblockrrr rr s ````````` @@@rrz Sort.create sJ=T%4%[<4q +;<Dzl w""6>+>?? (6CKK' '7#&&}['A'ABB = + #)=#F J 3U2258J 3S3STT !$ (6CKK' '6{{c)nn,,,,  ) )%(:q*A*A B B %*#f++$6$6  ? ? ? ? ? ? ? ?               &!&c&kk 2 2'   ,  F NN    rrr1r) rCr8rDr@rFr{rr{rrr)rrorrr)r!r2rrrr1rrxrrr!rr"rrr^rrrr)rrrrrbrxrrrzr}r%rrrrr3r^rrrrs@rrr sGGGG!!!!####----LLLF##       $#                 8888 UUUUU )6(=QQQQ[QQQQQrrcL t|ddS#t$rYdSwxYw)NFfreezeT)rrrrs rr&r&H sAa....t uus  ##c  t|d\}}|r}t|tsJt |t |j|j|j| |j |j|j|j |j }| S#t$rYdSwxYwNFr>)rrshould_pad_stridesrr(rrsr!r r _pad_stridesrrtru is_contiguousrr_bufferrqs r is_contiguous_storage_and_layoutrGP s/%@@@  $ $ & & fn55 C CtF|| C C5    ##FM6; MM  F##%%% uusB8B;; C C c t|d\}}|sdStjj||jt|j zS#t$rYdSwxYwrA) rrrDrwr"r#rm storage_sizertrrrrrEs r&is_dense_contiguous_storage_and_layoutrJd s/%@@@##%% 5w77    ! !6==3M3M#M    uus(A=AA== B  B r?want_contiguous stride_orderSequence[int | Integer] | Noner[rftuple[StorageBox, Layout]ct|trt|j|||||St|tr8t|j|||||\}}||jfSt|t r|r|r=||sJnH|| ||n.|| ||n| t ||fSt|tr"t|j|\}}||j fSt)z Try to simplify x into a StorageBox and a Layout. allow_padding only affect how we apply stride_order. When allow_padding is True, we have the freedom to add padding when applying the stride_order. r?rKrLr[rfNr[r>)rrrrrprr'rrZrDr^rirrvrqr) rr?rKrLr[rfr rqbuffers rrrrrp s!Y $ F+%''     !Z  &) F+%''    6!&##%%%%!V-  " "!!!||~~33555555)11 2*22!3!!!!}}allnn,,!_%% * F   qx rcx t|d\}}||S#t$rYdSwxYwrA)rris_stride_orderedr)rrLrFrqs r"is_stride_order_storage_and_layoutrU sS/%@@@'' 555 uus (+ 99rct|ttfrt|jSt|t rb|j}tjj |j t|j zt }t|jp|St|tr%|tjjvSdSr8)rrr is_unalignedrprvrqrwr"r#statically_known_multiple_ofrtrir rkrrunaligned_buffers)rrqhas_unaligned_layouts rrWrW s$J/00'DI&&&$((?#$7#3#P#P MN6<88 8/$ $  DI&&>*>>$<}}!'";;; 5rceZdZUded<edd3d4dZd5d Zd6d Zd7dZe d8dZ d9dZ d:dZ d;dZ dd!Zd?d"Zd@d$ZdAd%ZdBd'Zd?d(Zd?d)ZdCd+ZdDd-ZdEd.ZdFd1Zd2S)GrrrpFrrrr c6|j|Srrprxrws rrxzBaseView.get_free_symbol_uses sy--m<<.indexer 5&& &r)rr{rr )rprBr`)rrDrdrs @@rrBzBaseView.make_indexer sR &&((%%'' ' ' ' ' ' ' 'rr;cr|j|dfd }|S)Nrr{rrvc,|Srrrcs rrLz$BaseView.make_loader..loader rerr)rpr>r`)rrLrdrs @@rr>zBaseView.make_loader sR %%''%%'' ' ' ' ' ' ' ' rrc4|jSr)rpr*rs rr zBaseView.dtype sy""$$$rr c4|jSrrpr'rs rr'zBaseView.get_layout y##%%%rr0c4|jSrrpr+rs rr+zBaseView.get_device rkrrcdSrrrs rrzBaseView.get_origin_node rrrc4|jSrrprrs rrzBaseView.get_name y!!###rr{c*|Srr rs rrzBaseView.get_pointwise_size }}rTrOrQrrPrc:|j||SNrOrprTrSs rrTzBaseView.mark_reuse y##E{#CCCrc4|jSrrpr9rs rr9zBaseView.has_exceeded_max_reads y//111rr8c4|jSrrpr*rs rr*zBaseView.realize y  """rc8|jdSrrprVrs rrVzBaseView.realize_hint s      rrc4|jSrrprurs rruzBaseView.get_storage_numel y**,,,rc4|jSrrprrs rrzBaseView.is_extern y""$$$rct|jtsJt|j|jSr)rrprris_module_bufferrs rrzBaseView.is_module_buffer s>$)X..??TY??.y))+++rrc4|jSrrprrs rrzBaseView.get_read_names y'')))rrmctjtdd5t||jcdddS#1swxYwYdSr )rrr(rTr>r%rprs rrzBaseView.get_reads s \.*:D A A  &  ""                    s9A""A&)A&cl|}t|tr|j}t|t|Sr)rrrp)rrs rrzBaseView.unwrap_view s9H%% AH%% rr!r2c|}tjtd||}t ||||SrO)r>rrrSr=r*r%rTs rrzBaseView.constant_to_device sf!!##Hn.?HHPP..""==??     rNrr1rr^rrrrrrrrrrrrrrrrr)rrrrrbrxr`rBr>rr r'r+rrrrTr9r*rVrurrrrrrrrrrr sLLLJ''====('=CCCC%%%X%&&&&&&&&$$$$=ADDDDDD2222####!!!!----%%%%,,,,****       rrcZeZdZUded<eddZedd Zdd Zdd Z dS)rr{rrrnew_sizer^rctjj}d|D}|}dgt |t |z zt |z}t |t |ksJt t |D]}||dkr||J||||<$||*tjj||rWt|slt|s]||}||}|J|J||z }| |ddks%Jd|d|d||S) zReplace `-1` with correct sizesc6g|]}tj|Srrrs rrz.ExpandView._normalize_size..2 s 666ELOO666rNrRr)fallbackzBroadcast failed in ExpandView(, z) on dimension ) rwr"r#r%rrris_size_one_or_falser8r$)rrr#old_sizerv1v2diffs r_normalize_sizezExpandView._normalize_size. s7#66X666::<<6S]]S]]:;d8nnL8}}H ----s8}}%%  A{b  {...&qk !$(8(M(M ))$. /99 a[a[~~~~~~Bw..!"/  cajjllbbhbb_`bb rrcv|||}t|rt|\}}t|t|jz }|dksJt jjg|z}t|j |jD]O\}}| tj j |s|nt jjPt|j|jt%|||j|j} t+|| St-||S)Nrro)rpr)rr&rrrrrrr rrrrwr"r#rrsr!r rrtrurvr) rrrrwrxskipryrrrzs rrzExpandView.createV s3&&q(33  # # D"7":": GZx==3z#7#77D19999',$.J #J$5z G G   !!7+@@FF&FF %! X!$ J# CCC Cqx0000rc|jSrrrs rr%zExpandView.get_sizeq rrr^c|}|jt|tz dfd }|S)Nrr{rct|d}t|tksJttD]"}|dkrtjj||<#|Sr)rrrrrr )rractualrs rrz*ExpandView.make_reindexer..reindex{ sttuu&&Eu::V,,,,3v;;'' , ,!9>>$w|E!HLrrr{rr{)r%rpr)rtargetrrrs @@rr`zExpandView.make_reindexert sf##%%6{{S[[(       rN)rrrr^rr^)rrrr^rrrr) rrrrrrr3rr%r`rrrrr* s%%%\%N111[14rrcbeZdZUded<eddZedd Zdd Zdd ZddZ dS) PermuteViewrdimsrrrrrc||}t|ttt|ksJt |rct |\}t jjfd|Dfd|Dj j }t||St||S)Nc*g|]}j|Srrrrrxs rrz&PermuteView.create.. s 222#222rc*g|]}j|Srrrs rrz&PermuteView.create.. s!444!"1%444rro)rpr) _map_neg_dimsr@rrr&rrrsr!r rtrurvr)rrrrwrzrxs @rrzPermuteView.create s  &&$:eCII.>.>#?#?????  # # D"7":": GZ$! 2222T2224444t444!$ J# CCC C----r list[int]c fdDS)NcDg|]}|dkr|nt|zSr r)rrcrs rrz-PermuteView._map_neg_dims.. s/EEEsaxxSYY_EEErr)rrs `rrzPermuteView._map_neg_dims sEEEEEEEErr{ct||jttt |jksJ|jfd|jDS)Nc g|] }| Srr)rrrs rrz(PermuteView.get_size.. s+++AQ+++r)r@rrrrrpr%)rrs @rr%zPermuteView.get_size s$,,TY7788J #di.. ! != =     y!!##++++++++rct||jttt |jksJ|jfd|jDS)Nc g|] }| Srr)rrrs rrz*PermuteView.get_stride.. s---aq ---r)r@rrrrrprTrds @rrTzPermuteView.get_stride s$,,TY7788J #di.. ! != =     %%''----49----rr^c,dt|jDfdtt|jDt t tt|jksJdfd }|S)Nci|]\}}|| Srr)rrrs rrz.PermuteView.make_reindexer.. s5551q!555rc g|] }| Srr)rrinvs rrz.PermuteView.make_reindexer.. s555!s1v555rrr{rc"fdDS)Nc g|] }| Srr)rrrs rrz?PermuteView.make_reindexer..reindex.. s***E!H***rr)rrs`rrz+PermuteView.make_reindexer..reindex s +***c*** *rr)rrrrr@)rrrs @rr`zPermuteView.make_reindexer s65 $) 4 45555555uS^^44555#*U3ty>>-B-B"C"CCCCC + + + + + + rN)rrrrrr)rrrrrr) rrrrr3rrr%rTr`rrrrr s...[.$FFF[F,,,,....      rrcLeZdZeddddZedd ZddZdS) SqueezeViewN)rcrrrcrrc t|r_t|\}}g}g}LttsJt dkrt |jksJtt|j|j D]\}\}} Otj j |s*|||| Y|kr+|||| |dks Jdt|j|j|||j|j} t)|| S7t*|d|DS|dksJt*|fdt|DS)NrrIzexpected squeezed size to be 1rocZg|](}tjj|&|)Sr)rwr"r#rrs rrz&SqueezeView.create.. sA7+@@CCrc&g|] \}}|k |Srr)rrrrcs rrz&SqueezeView.create.. s#"U"U"UAAQTHH1HHHr)r&rrrrrrrrrrrwr"r#rrrsr!r rtrurvrhrr%) rrrcrwrxrryrrrrzs ` rrzSqueezeView.create s  # # D"7":": GZHJ!#s++66T#YY66+CxxC#jo*>*>$>$>$>>%.s:?JDU/V/V%W%W K K!>D&;7+@@FF2 ---"))&111Cxx ---"))&1111#qyyy*Jyyyy$! !$ J# CCC C ;;;ZZ\\ ::<<$))));;q"U"U"U"U1::<<1H1H"U"U"UVV Vrrr{>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]]cd|D}dt|Dt|dfd }||fS) Ncg|] }|dk| Srrrs rrz(SqueezeView.squeezer.. s...!qAvvAvvvrc$g|] \}}|dk |Srr)rrrs rrz(SqueezeView.squeezer.. s!;;;AAFF1FFFrrr{rtuple[Expr, ...]ct|tksJ|dtjjgz}t |D] \}}|||< t |S)N )rrrr rr)rr`rrlengthnot_ones rrz%SqueezeView.squeezer..reindex s|u::W---%/C/C'/C/C---%*W\NV$;Igu-- # #Q!" ### #r)rr{rr)rr)rrrrrs @@rsqueezerzSqueezeView.squeezer sm/.t...;;4;;;T $ $ $ $ $ $ $  rrprrc td)Nzuse SqueezeView.create())r)rrps rrzSqueezeView.__init__s7888r)rrrcrrr)rr{rr)rprrr)rrrr3rrrrrrrrr sz48+W+W+W+W+W[+WZ!!!\! 999999rrc`eZdZUded<ded<ddZddZdd ZeZeddZ ddZ dS) GenericViewr{rr^rrc|jSr)rrs rr`zGenericView.make_reindexer s |rrcdtt|jD}t||}ddt t|d|S)NcBg|]}ttj|Sr)rqrHr4rr s rrz+GenericView.reindex_str..s2   >? *4:q 9 9   rzlambda rr:)rrrrrrrr)r index_old index_news r reindex_strzGenericView.reindex_strsw  CHTYCXCX   i0011 F3sI#6#677FF9FFFrcv||jd|jd|gS)Nsize=zreindex=)rrprrrs rrzGenericView.__str__sD Y+ ++-L8H8H8J8J-L-L M   rrrrrc8||t||S)Nrprr)r)rrrrs rrzGenericView.creates!sX@@@@rc|jSrrrs rr%zGenericView.get_size'rrNrr)rrrr{rr^rrr) rrrrr`rrr2r3rr%rrrrrs7777 GGGG    HAAA[ArrceZdZdZeddZeedd ZeddZ e dddZ e dddZ dS)rhz This class handles tensor reshaping by computing appropriate index transformations to map the new shape back to the original storage layout. rr rrctj|}tj|}tjjjj}|tj|dr||z}|Sr$)rrrwr"r#r evaluate_exprLt)rrrs rhandle_negative_indexzView.handle_negative_index2s\l3|D!!(2@ =#q)) * * *C rrrrr{cfttsJt|\t jjr|Sttdkpttdk}t|}dd dfd }dvr!dfd }|t|S|r%|tSt|s ||St!|d\}}|j} t jj} t jj| } t jj} ddlm} | | | | |}|Ed|D}t+|j|j||j|j}t5||S||S)Nrrrrr{ryrrvc"t|t|}t |\}}t |j|j|||j |j }t||SNro) rrequire_exact_stridesr(r)r%rrrsr!r rtrurv)rrryrwrxrzs rcreate_reinterpret_viewz,View.create..create_reinterpret_viewKs44^66s||~~FFC#8"<"< GZ$! !$ J# CCC Crrc }|t|S#t$rBt|}|t cYSwxYw)a Handle the case where view is not possible with current strides. Try dynamic_reshape_indexer first; if it fails with unbacked symbols (guard_or_false can't resolve comparisons), fall back to making the tensor contiguous. r)rirr7rrequire_contiguousr(r))rrrrrrs r"handle_unbacked_or_dynamic_reshapez7View.create..handle_unbacked_or_dynamic_reshape\s 55hIIsXHHHH.   !33A66..x!B!B8!L!L  s04A B?Brrtuple[int, ...]cDtdgtzSr$)rr)rrs r fake_reindexz!View.create..fake_reindextsaS3x==0111rrFr>)_compute_stride)size_obliviouscng|]2}t|dr |jjntj|3Sr)rrexprrr!rs rrzView.create..sI 'q&11G u}Q7G7Grro)rrrr{ryr{rrvrrrr)rrrr)rrrresolve_negative_sizer%rwr"r#statically_known_list_equalsrr6rJrr(r)r&rrrto_symints_or_intstorch._subclasses.fake_implsrrsr!r rtrurv)rrrunbacked_symbols_in_sizesrDrrrwrx old_strideold_size_symintold_stride_symintnew_size_symintrnew_stride_symintryrzrrs` ` @@rrz View.create;s(H--==tH~~==- 66qzz||XNN( 7  8 88 L L H %h// 0 01 4 8(2233a7 "?qAA  D D D D"         , == 2 2 2 2 2 23ADNNLIII I **8^>>xHH  %Q'' 955a88 84AeDDD& '*==hGGG,?? KK'*==hGG@@@@@@,O   4      (*J %! !$ J# CCC C21!444rrtuple[list[Expr], list[Expr]]cd|D}d|D}t|}tt|D]Q}||dkrCtjj||<t t|t|||<nRtj j t|t|||fS)NcVg|]&}tjj|'Srrwr"r#r<rs rrz.View.resolve_negative_size..+CCCQAG$--a00CCCrcVg|]&}tjj|'Srrrs rrz.View.resolve_negative_size..rrrR) rrrrrOnerBrrrwr"r# check_equals)rrrs rrzView.resolve_negative_sizesDC(CCCCC(CCC>>s8}}%%  A{b  #gk &}X'>'> h@W@WXX ! %%mH&=&=}X?V?VWWW!!rNr^ dense_dimrrc  ||||}ni#tttf$rOt |g}|||}|||}t ||}YnwxYw|Sr)_dynamic_reshape_indexerrr7 IndexErrorrrr)rrrrrflatrrs rrizView.dynamic_reshape_indexers :228XyQQGG ;ZH : : :!(++,D33HdCCH33D(CCH%h99GGG  : sA#B?Br^cXtjjjdfd }dt t |Dt t|}t |}|duo(|t |d z kot |d k}|r.|J||}| |g|r$|r!|}|\} } |d kr= tj j | | | fn| d kr| |n|| |d kr | nj|| |d kr|| |d kr4|\} } | | z| z} | | z} || |d k4 | tjj | |n|| |d krtj j} |} t| | || |z} || |d krR|} t| | || |z} ||z}|| |d kRtjj | |nt |r|!|r_|}tjj |d  tj j |_|r>|\} } tjj | d |>|Rt |d kr?}||nt t |ksJdfd }|S)zG Perform a reshape entirely by modifying indexing math rsr rurrctj||rdStj||rdStj||rdStjj||rdStjj||rdSttj||)z Compare two symbolic sizes, returning -1 if a < b, 0 if a == b, 1 if a > b. For unbacked symbols, guard_or_false returns False, so we fall back to divisibility checks. rrRrI) rrQrGtrwr"r#rXr7)rsrurXs r compare_sizesz4View._dynamic_reshape_indexer..compare_sizess~ehq!nn-- q~ehq!nn-- r~ehq!nn-- qw<.s2   => *49a 8 8   rNrIrrr{ct|tks&Jt|tftt|tfdDS)Nc38K|]}t|VdSr)rs)rr replacementss rrzAView._dynamic_reshape_indexer..reindex..Hs-HHA|44HHHHHHr)rrrr)rr rF view_exprs @rrz.View._dynamic_reshape_indexer..reindexCsvu::T***SZZT,C***D% 0 011LHHHHiHHHHH Hr)rsr rur rrr)rwr"r#rXrrrrrrrrr rrrGrreverseinsert)rrrr stack_new stack_oldreordering_dense_dimold_dimsize_oldrsize_newvar2 size_new2divisormodulus dense_exprrrXrFr s @@@rrzView._dynamic_reshape_indexers)8 > > > > > >H  BGH BVBV   T8,,-- NN  T ! #S^^a// #H "   &(((mmI..G   W % % %  %I % }}H%MMOOMC1}}  ...  #x1111Q  ****x22a77  %%%%x22Q66#mHh77!;;&/mmooOD)/C/C')3H$mHh77!;;  %%% --hAAAAx22Q66'+"  gw!G!GHHH!G+#mHh77!;;'mmooG$$_S'7%K%KLLL%/G''1H $mHh77!;;  --hAAAA$$; %I %> + }}H G  ) )(A 6 6 6   UW\ * * * +  7%MMOOMC G  ) )(A 6 6 6 7  S]]a%7%7      "J   Y 3 3 3 3      9~~X.... I I I I I I Ir)rr rr rr )rrrr{rr)rr{rr{rrr)rr^rr^rrrr)rr{rr{rrrr^) rrrrrrr3rrrrirrrrrhrh+s \ n5n5n5X[n5`"""\"  !% [ !%zzzz\zzzrrhceZdZUdZded<d$fd Zd%dZeZd%d Zd&d Z d'd Z e d(dZ d)dZ d)dZd*dZd+dZd,dZd$dZed d-d.dZd/d0d!Zd1d#ZxZS)2rvz*Pretend our storage has a different layoutr rqrrctt|jtr5t |d|jdSdS)Nrp)rrrrprrrrrs rrzReinterpretView.__post_init__Ssf  di * * F   tVTY-B-B-D-D E E E E E F FrrcD||j|jgSr)rrprqrs rrzReinterpretView.__str__Xs'      rc4|jSrrprs rrzReinterpretView.get_namebrqrr0c|jjSr)rqr!rs rr+zReinterpretView.get_devicees {!!rrcdSrrrs rrzReinterpretView.get_origin_nodehrrrc|jjSr)rqr rs rr zReinterpretView.dtypeks {  rr{c4t|jjSr)rrqrrs rr%zReinterpretView.get_sizeosDK$%%%rc4t|jjSr)rrqrrs rrTzReinterpretView.get_stridersDK&'''rr;cdfd }|S)Nrr{rrvcj}tj||}jjjjkr%tj|jjjS|Sr)rqrBruloadrr rpto_dtype_bitcast)rrD tmp_loaderrs rrLz+ReinterpretView.make_loader..loadervslk..00G$--//775>>BBJ{ DIO33+J DIOTTT!!rrrrrLs` rr>zReinterpretView.make_loaderus( " " " " " " rr@c4|jSr)rqrBrs rrBzReinterpretView.make_indexer{'')))rc|jSrrqrs rr'zReinterpretView.get_layoutrrcdSrrrs rrZzReinterpretView.freeze_layoutrWrFrrrct|jj|t|jj|zt|jj|zSr)r)rqrrrtrws rrxz$ReinterpretView.get_free_symbol_usessK T[-} = =t{1=AA Bt{1=AA B rNr+r,ctjj|j|jj|jj|jj||j ntjjj |jj Sr?) rwr" wrapper_codecodegen_reinterpret_viewrprqrrrt writeliner r.s rr/z!ReinterpretView.codegen_referencescw#<< I K  K  K  & 2F  8L8V+# =   rrcdSrrrs rrszReinterpretView.num_readsqrrrrrrrrrrrrrrr)rrrrrrrr2rr+rrr r%rTr>rBr'rZrbrxr/rsrrs@rrvrvMs44NNNFFFFFF     H$$$$""""!!!X!&&&&((((    ****    -..$)    /.       rrvcjeZdZUdZded<edd Zdd ZeZe dd Z ddZ ddZ dS) DtypeViewz(Pretend our storage has a different typer target_dtyperr new_dtyperrct|rPt|\}}t|j||j|j|j|j}t||St||S)Nro)rpr7) r&rrrsr!rrrtrurvr6)rrr8rwrxrzs rrzDtypeView.creates|  # # D"7":": GZ$!!!$ J# CCC Cai8888rrcD||j|jgSr)rrpr7rs rrzDtypeView.__str__s 4+<=>>>rc|jSr)r7rs rr zDtypeView.dtypes   rr{c4|jSrrpr%rs rr%zDtypeView.get_sizerqrr;cJjdfd }|S)Nrr{rrvc`tj|jjjSr)rur&r7rpr )rrdrs rrLz%DtypeView.make_loader..loaders''c D4EtyWW Wrrrpr>)rrLrds` @rr>zDtypeView.make_loadersH %%'' X X X X X X X rN)rrr8rrrrrrr) rrrrrr3rrr2rr r%r>rrrr6r6s22 9 9 9[ 9????H !!!X!$$$$rr6cHeZdZdZedd Ze dddZdS) SliceViewzqView that represents a slice along a single dimension. Corresponds to tensor[..., start:end:step, ...]. rrrcrstartendrtuple[int, int]c tjj ||t d||fDrt t n9t d||fDrt t n j jd fd dfd }||dd}|||}||fS)zz Normalize start and end such that both are in the range [0, x.get_size()[dim]] and start <= end. c34K|]}t|VdSr)r6rs rrz0SliceView.normalize_start_end..s+HHA$Q''HHHHHHrc3K|]O}t|t|tjtjtt VPdSr)rr hasrrErDrs rrz0SliceView.normalize_start_end..s^  !T""  EE%)UYS 1 1      rrr lowerrupperrc||r|n ||}||r|n ||}|Sr)statically_known_geqrb)rrJrK clamped_lower clamped_fullmax_funcmin_funcr#s rclampz,SliceView.normalize_start_end..clampsj221e<<T((1eBTBT  00FF4 XmU33   rr"rr\ Expr | intcR||S|}|||Sr)r)r"rJrKr\rRrdim_sizes r clamp_wrapz1SliceView.normalize_start_end..clamp_wraps8{++C::C5eU++ +rr)rr rJrrKrrr ) r"rrJrrKrr\rSrrS) rwr"r#r%rarErD evaluate_min evaluate_max) rrrcrCrDrVrRrUrPrQr#s ` @@@@@rnormalize_start_endzSliceView.normalize_start_ends37#::<<$ HH%h1GHHH H H -HHH   S(+     - HHH,H,H  , , , , , , , , 5!Xq11jeXx88czrrITsteprRrc  tjttsdks J dkr|dkrdkr|Sn#t$rYnwxYwt | |r|||\}t|z dz z <t|r~t|\}}t |j } | z| <t|j |j | |j|j zz|j} t#|| Sd  fd } t%| | S) NrlrIrorr{rct|tksJd|dt|}|zz|<|S)Nz wrong ndim r)rr)rrcrrCrZs rrz!SliceView.create..reindex1sbu::X...0Pe0P0Ph0P0P...KKEsd*U2E#JLrrr)rrrr TypeErrorrr%rYrCr&rrrrsr!r rtrurvrB) rrrcrCrDrZrRrwrxryrzrrs `` ` @rrzSliceView.creates|D!!$%%741 zzcY..4199    D  %%  E00CDDJE3 uq!94@@  # # D"7":": GZj/00J(o4JsO$! !J$5c$:U$BB$ J# CCC C         ah@@@@sA A! A!N) rrrcrrCrrDrrrE)rIT)rrrcrrCrrDrrZrrRrrr)rrrrr3rYrrrrrBrBss ///[/b3A3A3A3A[3A3A3ArrBcDeZdZUded<ded<ddZdd Zdd Zdd ZdS) BaseConstantrr r2r!rr{cdSNrrrs rr%zBaseConstant.get_sizeBsrrr0c|jSrrrs rr+zBaseConstant.get_deviceErrrcdSrrrs rrzBaseConstant.get_origin_nodeHrrrmctSrr?rs rrzBaseConstant.get_readsKrrNrrrr)rrrrr%r+rrrrrr_r_=s~rr_cFeZdZUded<ded<ded<dd Zdd Zdd ZdS)Constantrr1rr r2r!rr;cdfd }|S)Nrr{rrvcBtjjjSr)rur9r1r rrs rrLz$Constant.make_loader..loaderVs< DJ77 7rrrr(s` rr>zConstant.make_loaderUs( 8 8 8 8 8 8 rr8cdSrrrs rr*zConstant.realize[rWrrc:t|j|j|S)N)r1r r!)rfr1r r5s rrzConstant.constant_to_device^sdj 6JJJJrNrrr)rrrrr>r*rrrrrfrfOsyJJJ     KKKKKKrrfc>eZdZUded<ded<ded<d d Zdd Zd S)IndexingConstantrrrr r2r!rr;cdfd }|S)Nrr{rrvcBtjjjSr)rurrr ris rrLz,IndexingConstant.make_loader..loaderis>$*dj99 9rrrr(s` rr>zIndexingConstant.make_loaderhs( : : : : : : rrc:t|j|j|S)N)rr r!)rmrr r5s rrz#IndexingConstant.constant_to_devicensdj 6RRRRrNrr)rrrrr>rrrrrmrmbseJJJ SSSSSSrrmcFd}d}ttt||D]q\}}|dkr tjj||s(tjj||sdS|td|z}||z}rdSNrIFT)reversedrrrwr"r#rmrD)rr`expected_strideexpected_stride_maxrys ris_contiguous_strides_for_shaperwrsOs5&112233  1 66 w77    '"::1>QRR 55s1ayy(1 4rc*tj|jzSr)rJpadding_alignment_bytesitemsizers rget_align_for_dtyper{s  )U^ ;;rc.eZdZdZd dZddZ ddd Zd S)r<zxAbstract base for Layout, MultiOutputLayout, NoneLayout. Represents the memory layout of the output of an Operation.rr0cDtt|jrr=rs rr+zOutputSpec.get_devicer?rrcDtt|jrr=rs rrIzOutputSpec.storage_sizer?rFrrrcDtt|jrr=rws rrxzOutputSpec.get_free_symbol_usesr_rNrrrr)rrrrr+rIrxrrrr<r<sfCC77777777%*7777777rr<ceZdZdZdeddfd=dZed>dZejd?dZed@dZ e jdAdZ edBdZ e jdCdZ dDdZ e Z dEdZ dFd ZdGd!ZedHd$ZdGd%ZedId)ZdJd,ZdGd-ZedKd/ZdLd0ZdGd1ZdMd3ZdNd5ZdOd8ZdBd9Zed dPdQd<ZdS)Rr zo Layout base class Carries tensor meta-information including offset and whether it is pinned. NrFr!r2r rrr^rSequence[Expr] | Nonertr rurrrcn|t|}||_||_t |t |ksJd|d|t d|DsJ||_||_||_||_ |j r|jj dks JddSdS)Nr , stride=c3@K|]}t|tVdSr)rrrs rrz"Layout.__init__..s,99q:a**999999rr@zOnly CPU tensors can be pinned) r(r)r!r rr_size_stride_offsetrur)rr!r rrrtrus rrzLayout.__init__s >#66t< <;4>;;MDzz" N Ndk&6 N8H N NTZ N NI N N(,  N5; N=J N N N rc|jSrrrs rr+zLayout.get_devicerrrctj5tjt |jt |j|j|j|j cdddS#1swxYwYdS)N)r r! pin_memory) rw fake_moder-r.rerrr r!rurs r get_examplezLayout.get_examples [  &' 22' 44j{>                   sA A&&A*-A*c6t|j|jSr)rwrrrs rrDzLayout.is_contiguouss.t{DIFFFrr`rhct|}|dvs |ddkrdSt|t||D]\}}}|dkr ||krdSdS)N)rrIFT)rrr1)r`rhndimleftrightrs ris_channels_last_contiguousz"Layout.is_channels_last_contiguouss|5zz v  qQ5!$  *5 1 1  " "    D%qyyTU]]uutrc t|jttt t|j|jD]\}}}|dkr ||krdSdSrr)rrrsr(r)rr)rrrrs r is_transposedzLayout.is_transposedsw!$ K ^66tHTY%> ? ? @rrrctjtksJdtjD}fd|D}fd|Dd d}|dgtz}t tD]}||||<t td z D]}||}||d z}t jjtj ||rM ||rd S ||sd Sd S) Ncbg|],\}}tjj|d*|-Srrl)rrrcs rrz,Layout.is_stride_ordered..#sG   37#;;CCC    rc*g|]}j|Srr)rrrs rrz,Layout.is_stride_ordered..)s888Q$+a.888rc g|] }| Srr)rrrs rrz,Layout.is_stride_ordered..*s@@@Qa@@@rarrrrc>t|fd|DS)Nc:g|]}|Srr)relement sorted_arrs rrzDLayout.is_stride_ordered..sorted_indices...s'AAA'J$$W--AAAr)r)rrs @rsorted_indicesz0Layout.is_stride_ordered..sorted_indices,s(JAAAASAAA ArrRrIFT)rrrr) rrrrrrwr"r#rXrrQr) rr non_1_indicesrrstride_orderedrrrs `` rrTzLayout.is_stride_ordereds4;3u::----  #DI..   9888-888@@@@-@@@ B B B B u%%E *s5zz"" 1 1A'-ayN58 $ $s5zzA~&&  A!!$D"1q5)Ew..uxe/D/DEE ,,T599 uu00== uu trc dgtttdt|jdz z}t|g|z}||SNrrI)rrsrrrrTras ris_channels_last_stride_orderedz&Layout.is_channels_last_stride_orderedCs]d8E!S-=-=-A$B$BCCDDDU u$%%e,,,r in_stridesct|}t|dkr|Stjst||r|St j}t|dr|j ddr|Stdtj ||D }tjs|r|Stt jdrt jjnddfd rt#fd |Dr|St%|}t'|}dt)t|D}d||d<d} t+|dddD]\} } || dz } || || z} t-| t.t0jfo| tjko| |zdkp t-| t6o tj}| || <|rt9| ||z|| <d} | s|St:xjdz c_|S)z The padding does not change stride order but makes sure all strides larger than the threshold are multiple of align. rrdislike_paddingFc3XK|]%}t|ttjfV&dSrrrs rrz&Layout._pad_strides..bsG   q3 . / /      r _shape_envNrrSrrczdSt|tsdStfd|jDS)NFc3BK|]}|VdSr)is_unbacked_symint)rrrs rrzILayout._pad_strides..contains_unbacked_symints..ps1RR1y33A66RRRRRRr)rr rar5)rrs rcontains_unbacked_symintsz6Layout._pad_strides..contains_unbacked_symintsksK udD)) uRRRR@QRRRRR Rrc3.K|]}|VdSrr)rrrs rrz&Layout._pad_strides..ss/NNa66q99NNNNNNrcg|]}dSr rr s rrz'Layout._pad_strides..ys999Qq999rrI)rCT)rrSrr)r{rrJpad_channels_lastr rrwget_current_noderrrrrchainpad_dynamic_shapesr"rrarrrrrrrr!padding_stride_thresholdr rcr'num_comprehensive_padding)rrr aligncurrent_fx_nodermrLr new_stridespaddedrankrprev_idxrrequire_paddingrrs @@rrCzLayout._pad_stridesIs$E** z??a   ' F,N,N *- -   ,.. ?F + + 0D0H0H u1 1     _Z66      ( Z  *1!'<*H*HRAG&&d  S S S S S S  NNNN:NNNNN  ' I>> ,\:: 99%J"8"8999 &' JqM"":abb>;;;  ID#!$(+H *T(^;F6C#788(V<<(UNa'JVT**Hv/H  &K  #*65#9#9E#A C    ))Q.))rct|tsJt||jJ||j|j|j|_dSr)rr(rrrCrr rs r pad_strideszLayout.pad_stridessW$//;;d;;/{&&&'' TY KK rcDtjot|tSr)rJcomprehensive_paddingrr(rs rrBzLayout.should_pad_stridess+P 40P0PPrrsct|tr|S|r|t|j|j|j|j|j|j Sr) rrsrBrr!r rrrtrurs ras_fixedzLayout.as_fixedso dK ( ( K  " " $ $        K J I K K N    rr@ctjs Jdt|jd|S)Nzconvert z to FixedLayout first)r(rrrrrBrs rrBzLayout.make_indexersR,  AtDzz* A A A  ,}}++---rotherrct|to_|j|jkoO|j|jko?|j|jko/|j|jko|j|jko|j|jkSr)rr r!r rrrtru)rrs r__eq__z Layout.__eq__s uf % % 2 u|+ 2 ek) 2 UZ' 2 u|+  2  u|+  2 %/1 rcBt|j|j|jSr)r-rrrtrs rrIzLayout.storage_sizes.ty$+t{SSSrrrct|j|t|j|zt|j|zSr)r)rrrtrws rrxzLayout.get_free_symbol_usess@ TY 6 6t{M:: ;t{M:: ; r)r!r2r rrr^rrrtr rurrrr)r1r^rrrr1r{rrrr1r rrrr)rrr)r`r^rhr^rr)rrrrrr)rrrr)rrrr{r rrrrrrsr)rrrrrr)rrrrr!rrrsetterrrtrr2r+rrDrrrrrTrrCrrBrrBrrIrbrxrrrr r s)-wqzz     2X [[X ]]X ]]     HGGGG\" @ @ @\ @""""H---- LLL\L\LLLL QQQQ     ....     TTTTH%%$)    &%   rr ceZdZdZddZdS)rsz A Tensor layout we cannot changerr@cBt|j|j|jS)r)rrrrtrs rrBzFixedLayout.make_indexersdidkBBBrNr)rrrrrBrrrrsrss4**CCCCCCrrsceZdZdZdZd4dZed5d Zed6d Zed7d Z ed8dZ ed9dZ e d:dZ e jd;dZ e d:dZejd;dZe dd?dZ d>d@d!ZdAd"ZdBd#ZdCd%ZdDd)Z dEdFfd3 ZxZS)Gr(z| A Tensor layout that we are allowed to change Assumption: layout change should NOT add or remove free symbols FrrscNtj|S)z Compute what the strides would be if this layout were frozen, without actually modifying the layout. This is used for speculative stride computation during Triton template code generation. )rdeepcopyrrs r!get_fixed_layout_without_freezingz0FlexibleLayout.get_fixed_layout_without_freezings }T""++---rsizesrrct|dkrgStjjg}t |ddD] }|||dz!t t |S)NrrIrR)rrrrrsrr)rreversed_stridesrs rr)z!FlexibleLayout.contiguous_stridess| u::??I!GK=U122Y'' A AD  # #D+;B+?$? @ @ @ @H-..///rrcttt|t|ks J||ftjj}dgt|z}|D]}|||<|||z}|S)z Create a stride based on the order the dimensions should be filled in. In this format, channels last would be: [1, 3, 2, 0] N)r@rrrrr)rr next_striderhrs r fill_orderedzFlexibleLayout.fill_ordereds%E ++,, 50A0AAAAE5>AAAgk &3u::% 1 1A$GAJ%a0KKrr{cttt|t|ksJt|}t||S)z Create a stride based on the sorted order of a permuted range. In this format, channels last would be: [3, 0, 2, 1] )r@rrrr(r)rrrs rrzFlexibleLayout.stride_orderedsV%E ++,, 50A0AAAAA,U33 **5*===r memory_formattorch.memory_formatcZ|tjkr t|tS|tjkr t|t S|tjkrt|St d|t)aq Create a stride based on a memory format. Memory format is translasted into a stride order, so channels_last is the same as: FlexibleLayout.stride_ordered(sizes, [3, 0, 2, 1]) This interface does not support memory_format `torch.preserve_format` which should be used to deduce a format from another source z>stride_ordered_for_memory_format, unsuppored memory_format: %s) r- channels_lastr(rNHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatr)rrr)rrs r stride_ordered_for_memory_formatz/FlexibleLayout.stride_ordered_for_memory_formats E/ / /!008IJJ J e4 4 4!008JKK K e5 5 5!44U;; ; IIP   & %rrr^c$t|t|ksJtjj|}t t t||j}t ||S)z Create a stride that has the same stride order as given stride For example, if given stride is [1000, 1, 100, 10], the fill order should be [1, 3, 2, 0] r) rrwr"r#guarding_hints_or_throwrr __getitem__r(r)rrrs r same_orderedzFlexibleLayout.same_ordered*sr5zzS[[((((!99&AAE#f++..F4FGGG **5*===rc|jSrrrs rrzFlexibleLayout.size9rrr1rc@|d|||_dS)Nr)!assert_free_symbol_uses_unchangedrrs rrzFlexibleLayout.size=s# ..vu=== rc|jSrrrs rrzFlexibleLayout.strideBrrc@|d|||_dS)Nr)rrrs rrzFlexibleLayout.strideF# ..x??? rr c|jSrrrs rrtzFlexibleLayout.offsetKrrc@|d|||_dS)Nrt)rrrs rrtzFlexibleLayout.offsetOrrr[rc||j|}|r#|r!|||j|j}t |j|j|j||j|jSr) rrrBrCr rsr!rtru)rrr[rys ras_stride_orderzFlexibleLayout.as_stride_orderTs((E::  " " $ $ N N**:ty$*MMJ K J I  K N    rrfc|}|r#|r!|||j|j}t |j|j|j||j|jSr)rBrCrr rsr!rtru)rrfr[rys ras_exact_strideszFlexibleLayout.as_exact_stridesdso#  " " $ $ N N**:ty$*MMJ K J I  K N    rc||j|}|r!|||j|j}t |j|j|j||j|jSr) rrrBrCr rsr!rtru)rrrys r as_fill_orderzFlexibleLayout.as_fill_ordertsw$($5$5di$G$G  " " $ $ N**:ty$*MMJ K J I  K N    rc||j|}|r!|||j|j}t |j|j|j||j|jSr) rrrBrCr rsr!rtru)rrrys r as_same_orderzFlexibleLayout.as_same_ordersw&&ty&99  " " $ $ N**:ty$*MMJ K J I  K N    r$dict[tuple[str, bool], sympy.Symbol]c i}dD]9}dD]4}||f}ttt|||||<5:|S)N)rrrtTF)r@r)r)rinitial_free_symbolsrrrs rget_initial_free_symbol_usesz+FlexibleLayout.get_initial_free_symbol_usessl!0  D!.   ]+,6$WT4%8%8-HH--$S))  $#rrrr9cdD]D}|j||f}tt||}||ksJd|d|EdS)Nr z)Expected free symbols unchanged, but got z vs )rr@r))rrr1rold_free_symbolsnew_free_symbolss rrz0FlexibleLayout.assert_free_symbol_uses_unchangedsy*  M#8$ 9NO )*:5-*P*PQQ #'7777d$11$ EEGG$77==G gKKK%)$E$E$G$G!!!rr)rrrr)rrrrrr)rrrrrr{)rrrrrr{)rrrr^rr{rrrrr)rrr[rrrs)rfr^r[rrrs)rrrrs)rr^rrs)rr )rrr1r9rrr8) r!r2r rrr{rLrMrurrr)rrrrrrrr)rrrrrrrrrtrrrr rrrrrs@rr(r(so N....000\0\  > > >\ >&&&\&4 > > >\ >X [[X ]]X ]] ;@     "HM               $ $ $ $8< HHHHHHHHHHHrr(cbeZdZdZdfd ZddZdd Zed dddZxZ S)NonOwningLayoutz,Is a view into the storage of another tensorviewBaseView | TensorBoxrrc|}t|j|j|j|j||_dSr)r'rrr!r rrr)rrrqrs rrzNonOwningLayout.__init__sO""  M L K M     rr@cN|Sr)rrBrs rrBzNonOwningLayout.make_indexers}}++---rrc|jj}|dkrdSddlm}t jj||S)NrTrI) ALIGNMENT) rr'rtutilsrrwr"r#rX)rrtrs rmaybe_guard_alignedz#NonOwningLayout.maybe_guard_alignedsS%%''. Q;;4$$$$$$w< srdefault_factoryrrcdgSr$rrrrrzNoneLayout. s1#rrrrcdSr$rrs rrIzNoneLayout.storage_size r4rr<c|Srrrs rrzNoneLayout.as_fixed rc|jSrrrs rr+zNoneLayout.get_devicerrNrrr) rrrrrrrrrIrr+rrrr/r/s 'k' DDDDDDDD) )++FFFFFFFFrr/ceZdZdfd ZeddZejdd Zdd ZddZd dZ e d!d"dZ d#dZ d$dZ xZS)%MutationLayoutSHOULDREMOVErrrrcTt|||d||_|}tj |dSr) rrr6r*r%r get_bufferrrwr"mark_buffer_mutated)rrrrs rrz#MutationLayoutSHOULDREMOVE.__init__s   & & ( (      OO          ))++ ##D)))))rr{c4|jSr) real_layoutrrs rrz!MutationLayoutSHOULDREMOVE.stride"s!!((rr1rcdSrrrs rrz!MutationLayoutSHOULDREMOVE.stride&s rr cN|Sr)r>rIrs rrIz'MutationLayoutSHOULDREMOVE.storage_size*s !!..000rrcdfd |j}t|tsJt||S)Nrrrct|tr|jSt|tr|St|t r|jS|Sr)rr9rrr MutableBoxrp)r unwrap_viewss rrDz;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_views.s&"<== 3#|FM222&(++ :#|F$6$6$8$8999&*-- 1#|FK000Mr)rrrr)rrrr)rrrDs @rr;z%MutationLayoutSHOULDREMOVE.get_buffer-s]      dk**&&))774<<77) rr cf|j}t|tsJ|Sr)r;rqrr )rrqs rr>z&MutationLayoutSHOULDREMOVE.real_layout;s/"")&&))))) rFsrcdst unsafe_aliasrc v|tj|t |t r|j}||st | | | dt||D}t |t t"fsJ|j}|t%|ds J|t |jjt(s!Jt+|jjt-||j_|jS)Nc^g|]*\}}tjj||+Srrwr"r#check_equals_and_simplifyrrsrus rrz;MutationLayoutSHOULDREMOVE.realize_into..ZsA1G$>>q!DDrrRrp)r*rwr"r<rrrrprVr=rr+r*r>rr%rrCrrqr(rr9)rrFrGrHrs r realize_intoz'MutationLayoutSHOULDREMOVE.realize_into@sp  ##CLLNN333 c9 % % (C  ##~~''mmoo** #CLLNNCLLNN C C $DdXz$:;; ; ;;)C sF##((S((##(/>::QQDr3rNrrBrrs@rr9r9s  * * * * * *)))X) ]   ] 1111     Z#xZ$S)Zrr8rr<rqrrctt|dddSr)rrrrs rrzBuffer.__post_init__ys5   t44444rr@cN|Sr)r'rBrs rrBzBuffer.make_indexer}s  --///rrc2|js J||jSrrrs rrzBuffer.get_namesy$yyrtorch.Tensor | torch.SymIntct|jtr|jSt t |jjr)rrqr rrrrrs rrzBuffer.get_examplesE dk6 * * -;**,, ,!$t{"3"3"<===rr0cN|Sr)rr+rs rr+zBuffer.get_devices ##%%00222rrcdSrrrs rrzBuffer.get_defining_oprrrc4|jSr)r'r rs rr z Buffer.dtypes  &&rr{c8g|jSr)r'rrs rr%zBuffer.get_sizes(""'((rrc8g|jSr)r'rrs rrTzBuffer.get_strides*"")**rr c4|jSr)r'rtrs r get_offsetzBuffer.get_offsets  ''rr ct|jtr|jStt |jjr)rrqr rrrrs rr'zBuffer.get_layouts9 dk6 * * ; !$t{"3"3"<===rc|jSrr,rs rrzBuffer.get_output_specrrrc*|Sr)r"rs rruzBuffer.get_storage_numels~~rrc4|jSr)r'rurs r get_is_pinnedzBuffer.get_is_pinneds  **rct|jtr:t|jts"|j|_dSdSdSr)rrqr rrrs rrZzBuffer.freeze_layouts] dk6 * * 1: K4 4  1+..00DKKK 1 1 1 1rFrrr[ct|jtsJt|j|j|||_dSNrQ)rrqr(rrr]s rr^z&Buffer.freeze_layout_with_stride_ordersN$+~66IIT[8I8III6k11%}1UU rct|jtsJt|j|j||_dSr)rrqr(rrras rrbz$Buffer.freeze_layout_with_fill_ordersG$+~66IIT[8I8III6k//66 rrct|jtsJt|j|j||_dSr)rrqr(rr rds rrez$Buffer.freeze_layout_with_same_ordersG$+~66IIT[8I8III6k//77 rrfct|jtsJt|j|j|||_dSre)rrqr(rrrhs rriz'Buffer.freeze_layout_with_exact_stridessU$+~66IIT[8I8III6k22 3   rctjjt j|dSr$r%rs rr'zBuffer.is_zero_elementsr(rr;cr(ttSdfd }|S)Nrrr{rrvcv}tjjpd||SrH)rBrur%rrrDrs rrLz"Buffer.make_loader..loaders5''))G8DI2GGENNCC Crr)r'r r;r*r(s` rr>zBuffer.make_loaders`  " " B=0@0@AAA A D D D D D D rNr+r,c*|Srrr.s rr/zBuffer.codegen_referencersrcdSrrrs rrzBuffer.decide_layoutrWrrcxt|jtr|jjgSdSra)rrqrrrrs rrz#Buffer.get_inputs_that_alias_outputs5 dk? 3 3 1K$--//0 0rrcxt|jtr|jjgSdSra)rrqr9rrrs rrzBuffer.get_mutation_namess6 dk#= > > 3K&//112 2rrrcFt|gSr)r@rrs rrzBuffer.get_read_namess4==??+,,,rrrctSrr?rws rrxzBuffer.get_free_symbol_uses||rctSrr?rs rrzBuffer.get_unbacked_symbol_defsrrcdSrrrs rr*zBuffer.realizerWrcdSr8rrs rshould_allocatezBuffer.should_allocatesurrrr)rrUrrrr)rrrrrrrrrr)rrrr)rfrr[rrrrrrrrrrr)%rrrrrrBrrr+rrr r%rTr]r'rrurbrZr^rbrerir'r>r/rrrrrbrxrr*rxrrs@rrros 5555550000>>>> 3333'''X'))))++++((((>>>>     ++++1111;@VVVVV 77778888 CH     UUUU          ----H%%$)&%     rrc4eZdZd dZd dZejZd dZdS) OperationBufferrrc|gSrrrs rrzOperationBuffer.get_outputss v rrc|Srrrs rrzOperationBuffer.get_defining_opr6rrcnt|t|dSr)rrrrs rrzOperationBuffer.__post_init__s0T"""%%%%%rNrrrr)rrrrrrrrrrrrzrzs[#5&&&&&&rrzceZdZddZdS)rrrcdSrrrs rrszInputBuffer.num_reads r4rNr)rrrrsrrrrr s(rrceZdZdZdS) DonatedBufferaY Represents a donated buffer which is a saved tensor that is not alias to any fwd inputs, fwd user outputs, and bwd outputs. We generally cannot inplace reuse the input tensor memory during backward since it might be used in another function. However, donated buffer can be inplace reused during backward to save memory. N)rrrrrrrrrsrrc.eZdZUdZded<d dZd d ZdS) rSNr0rQrr;cdfd }|S)Nrr{rrvc}tjtjj||Sr) r'rBrur%rwr" constant_namerrQrls rrLz*ConstantBuffer.make_loader..loaders]oo''4466G8%%dmmoot7KLL rrrr(s` rr>zConstantBuffer.make_loaders(       rr!r2rcttj|||jSNrrq)rSrwr"rrrqr5s rrz!ConstantBuffer.constant_to_device%s:&&t}}??     rrr)rrrrQrr>rrrrrSrSsO+/O////      rrSc^eZdZddZed dddZddd ZddZddZd S)NoneAsConstantBufferrrmctSrr?rs rrzNoneAsConstantBuffer.get_reads-rrFrrrctSrr?rws rrxz)NoneAsConstantBuffer.get_free_symbol_uses0rtrNr+r,rc.tjjjSr)rwr"r0none_strr.s rr/z&NoneAsConstantBuffer.codegen_reference6sw#,,rr<c"tdSNr)r/rs rrz$NoneAsConstantBuffer.get_output_spec9s&&&&rcdSr8rrs rrz&NoneAsConstantBuffer.has_tensor_output<r:rrrrrrrr) rrrrrbrxr/rrrrrrr+s233$)43 -----''''rrcZeZdZUded<ed dddZddd ZddZd S)rr rFrrrrc,t|j|Sr)r)rrws rrxz*ShapeAsConstantBuffer.get_free_symbol_usesDs  =999rNr+r,rcTtjj|jSr)rwr"r0codegen_sizevarrr.s rr/z'ShapeAsConstantBuffer.codegen_referenceJsw#33DI>>>rcdSr8rrs rrz'ShapeAsConstantBuffer.has_tensor_outputMr:rrrrrr)rrrrrbrxr/rrrrrr@sJJJ344$)::::54: ?????rrceZdZUdZded<dZded<dZded <dZd ed <dZd ed <dZ d ed<e j dEdZ e e j dEdZdFdZdGdZdHdZdIdZdJdZed dKdLd ZdMfd" ZdNd#ZdOd%ZdPd'ZdQd)ZedRd+Z dSdTd0Ze dUdVd:ZdWd<ZdWd=Z dFd>Z!dNd?Z"dNd@Z#dXdDZ$xZ%S)Yrzb Represents a buffer that is computed during kernel execution rather than being an input. rrpFzClassVar[bool]_force_realizeNrrACallable[..., Any] | NonerBrrCrDrIterator[None]c #2K|jJ|jJ|jJ|jJt |jt sJt|j|j}|j} t |j |j |j|j|j|j |j |j }||_t|j |j |j|_|j|dV||_||_dS#||_||_wxYw)Nr)rArBrCrDrrprrrqr!r rhrrrsget_default_sizes_body clear_cache)rold_datarxnew_datas rwith_original_inner_fnz%ComputedBuffer.with_original_inner_fn`s?+++&222$000.:::$)Y//EEDOO1EEE/9[  % n0,!%!@'6",'6   H!DI&%DK  ' 3 3D 9 9 9 EEE DI$DKKK!DI$DK $ $ $ $s -B DDc#~Ktj} dt_dV|t_dS#|t_wxYwNT)rr) old_values r force_realizezComputedBuffer.force_realizesH#1  6,0N ) EEE,5N ) ) )IN ) 5 5 5 5s.<r8cd|j|jSt|jdr |jjSdS)z Returns self.name if it exists, otherwise returns the name of the data node if that exists. If neither exist, returns None. Nr)rrrprs rget_computed_buffer_namez'ComputedBuffer.get_computed_buffer_names7 9 9  49f % % "9> !trrc4|jSrrprsrs rrszComputedBuffer.num_readsrrrmc4|jSrrprrs rrzComputedBuffer.get_readsrrrc4|jSrrrs rrzComputedBuffer.get_read_namesrrrjct|jtttt fs;t jtttStj tdd5|j r]t||j|jcdddSt||jcdddS#1swxYwYdS)Nrpwrites index_exprsrT)rrprrrr=rK ReadWritesr@rrr(rzrTget_store_functionrr}r%rs rrlzComputedBuffer.get_read_writesst$)itY%GHH * ll!||&LL  \.*:D A A  y++-- *++--I0022I0022        +++--I&&((                  sA*E;8EEErrrc|j||j|z}|r*|||z}|Sr)rqrxrphas_store_functionrl)rrrs rrxz#ComputedBuffer.get_free_symbol_usessv$11   I * *= 9 9:  " " $ $ Q d**,,AA-PP PF rr;c |sP|jtjjvr8|dkr |js|jStSr$) rzrrwr"mutated_buffersrsrrpr>rrs rr>zComputedBuffer.make_loadersu'')) + !888  A%%'&9((** *ww""$$$rc\t|jtttt fSr)rrprrrr=rs rrz!ComputedBuffer.has_store_functions$)itY%GHHHrCallable[..., None]c|}t|jt t tfr t|jj |j |St|jtsJt|jt|jj |j |Sr)r'rrBrrprrrr rrr=rrM)rrDs rrz!ComputedBuffer.get_store_functions//##,,..;;== di)T4!8 9 9 G494diII Idi33 D DT$)__ D D349149gFF Frlist[int] | Nonec:t|jtrtj|j|j\\}}|j }td|DsJfd|D}|rqt|jttfr|j |n|fd|D}ddlm}|||SdS)al If our layout is still flexible, try to determine the stride order based on stride orders of reads. TODO(jansel): A better algorithm here would look at downstream consumers of this value and try to do global graph-level layout optimization. This is also something just begging to be autotuned. c3bK|]*}t|tjtjfV+dSr)rrKStarDep MemoryDeprrs rrz0ComputedBuffer.get_fill_order..sK1|3\5KLMMrcg|];}t|tjt|jdDi|]}|dk|tjjSr r.rvs rrz..s*$W$W$WPQUVPVPVQ PVPVPVr)rrKrrsr)rrrs rrz1ComputedBuffer.get_fill_order..sWa!78817$W$Wn$W$W$WXXrcZg|]'}tjj|(Srrwr"r#r)rrrs rrz1ComputedBuffer.get_fill_order..s;"""EIAG$11$@@"""rrIpick_loop_orderN)rrqr(rKrrprr}rlrprrrr schedulerrr%)r index_varsr rpstride_lengthsrrrs @@rrzComputedBuffer.get_fill_ordersl dk> 2 2 H.:.M ,,.. 0L0L0N0N// + (Z!((**0E  E  Hdi$66)"i// NKKGG(G""""MR"""766666&~t}}GGGtrrct|jtrC|}|r||dS|dSdSr)rrqr(rrbrZras rrzComputedBuffer.decide_layout sk dk> 2 2 %''))E %22599999""$$$$$  % %rMtuple[tuple[list[Expr], list[Expr]], LoopBody, tuple[list[Expr], list[Expr]]]ctj||d\}}t jt d|5t| | r|n |dd|g|R}dddn #1swxYwYg}g}g}g}| D]t\}} ||dvr/|rJ| || | >||dvsJ| || | u||f|||ffS)NqrrQrIr) rKrrr}rrrSr+rWrrzitemsr) rr var_rangesrar reduce_vars index_size reduce_sizerrs rrz%ComputedBuffer.get_default_sizes_bodys(:  # # % %t'>'>'@'@   j\.*;T__=N=N O O  ''))0022@RaR D                !#   $$&& & &DAqDG||&&&!!!$$$!!!$$$$DG||||""1%%%""1%%%%K($[0IIIs,AB88B<?B<extra_indexing_constraints'tuple[dict[Any, Any], list[Any]] | Nonerecompute_sizes_body_func5tuple[tuple[list[Expr], list[Expr]], LoopBody | None]cl\\}}}\}}|r|||f|||f\\}}}\}}g|j|t|trt |dksJ|\}} t|t sJt|t| tsJt| td| DsJ|j } | |ks J| |ffd| D} | z g| tj tjs'|dfd } ||z} t't) p t*j } | || || \}}}| || || \}}}t/j||d\\}}}t3|||||g|||}||f|fS)an This is a main place where we do loop transformations in a backend-agnostic way. Here we: 1) Remove any 1 dimensions 2) Fuse contiguous dimensions together 3) Reorder dimensions based on stride orders Optional argument extra_indexing_constraints can be used to append additional indexing expressions to existing ones derived from buffer's body. This can be useful to fuse scheduler nodes with compatible ranges, e.g. (s0*s1*...,) and (s0, s1, s2, ...) on CPU by preventing indexing simplifications and obtaining index/reduce ranges for the scheduler node compatible with other nodes. Optional argument recompute_sizes_body_func can be used to recompute sizes and body on the default body. This can be useful to append additional loop transformations. Nrc3@K|]}t|tVdSr)rr )rfs rrz6ComputedBuffer.simplify_and_reorder..ds,HHqz!T**HHHHHHrcg|]}|v| Srr)rrindex_formulass rrz7ComputedBuffer.simplify_and_reorder..ls*###!>2I2I2I2I2Irx_varsSequence[sympy.Symbol] support_varsrrsimplify_loopsrrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]]c $ || \}}} dkr}tdkrjtt t}||ddkr,fd|D}t |}t |}||}|rJtjj ||t ||\}}} t||} n|} || |fS)Nrerrc g|] }| Srrrrrs rrzUComputedBuffer.simplify_and_reorder..simplify_and_reorder..s888Qa888r) _apply_loop_reorderingrzrrrrrrwr"r#_simplify_loopsrOr)rrrrnewsizesreindex0rrr_prunerr memory_addrsrs ` rsimplify_and_reorderzAComputedBuffer.simplify_and_reorder..simplify_and_reorderus.,0+F+F e\,, (Hh*&&((E11c%jjAooU3u::..//8E??1%**8888%888H+E22H.u55HXf%%F #-.W-=-M-M,^VXNN..*(F *(H=="Wh. .rpr) rrrrrrrrrr)rindexing_exprsrrrrrrrrrget_write_exprsrwr"rrLPREFER_STORE_LOOP_ORDERextendget_read_exprsrnr=rJloop_ordering_after_fusionrKindex_vars_no_squeezerW)rrrrrrarrextra_indexing_rangesextra_indexing_exprexpected_var_rangesrrshould_merge_loops iter_ranges iter_reindexr  reduce_rangesreduce_reindex iter_varsrrrs` @@rrz#ComputedBuffer.simplify_and_reorder6s4  ' ' ) )  %Z  %Z %  *)[)4*k1J  )[)[ 94.55778 % 15u== 233q8889:T 6 !#63T:: W WDAV.s@ --dJ MMrrz%Did not simplify complex index: %s %sc g|] }| Srrrs rrz9ComputedBuffer._apply_loop_reordering..s)))aq)))r)rrrrrs ExceptionrJrrwarningrrrrr)rrrrrrrhrs``` rrz%ComputedBuffer._apply_loop_reorderingsx /.....  L ,(Gw<<3|#4#444WQZCMM::://'5,"O"OPPQQEE , , ,|  =Z//00  s5zz**++EEE ,*)))5)))l5))?5+A+AAAsA>BA6DDr{c4|jSr)rprrs rrz!ComputedBuffer.get_pointwise_sizey++---rc4|jSrrpr}rs rr}z!ComputedBuffer.get_reduction_sizerrc4|jSrrprzrs rrzz!ComputedBuffer.get_reduction_typerrc4|jSr)rpr'rs rrzComputedBuffer.is_no_opsy))+++rcdSrrrs rrxzComputedBuffer.should_allocaterrr!r2rc6|j|S)rPrprr5s rrz!ComputedBuffer.constant_to_devicesy++F333r)rrrrrrrrrrr)rr)rrr)rrNN)rrrrrrr) rrrrrrrrrrrrrr)&rrrrrrrArBrCrDrrrrrrrsrrrlrbrxr>rrrrrarrrrr}rzrrxrrrs@rrrQsKKK%*N****#K""""488888266666<@@@@@ % % % %D666\6    %%%%%%%%*****,--$).-2 % % % % % %IIIIGGGG%%%%N%%%%JJJ]JDOS?CU2U2U2U2U2n *. %B%B%B%B\%BN............,,,,44444444rrc<eZdZUdZded<ded<ded<ded<dS) FinalizeCodegenResultzNStructured result from TemplateBuffer._finalize_codegen for external backends.rsourcerimports call_preamble call_argsN)rrrrrrrrrr sEXXKKKrrceZdZdZ d@dAfd ZedBdZdCdZdDdZdEdFdZ dGdZ dHdZ dId Z dJdKd%Z dId&ZdLd(ZdMd+ZdNd/ZedOd3Zedddd4dPd?ZxZS)Qrz Base class for template operators that support epilogue and prologue fusion. Subclasses: TritonTemplateBuffer (built-in Triton templates), HelionTemplateBuffer (Helion kernels), etc. Nrqr<rrmake_kernel_renderrmutated_inputsIterable[IRNode] | Noneallowed_prologue_inpsOrderedSet[str] | None named_inputsdict[str, IRNode] | Nonerrctd|t|_|_t j_ t j i_ i_ i_ |rt|ni_|_g_|ajd}t%|t&sJt)||fd|D_|p t-_d_d_dS)NrrcNg|]!}tt|"SrMutationOutputr/rrr!rs rrz+TemplateBuffer.__init__..?s@%%%z888#tDD%%%r)rrrunwrap_storagerrrwr"register_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsrmutation_outputsrrrr+r@rallow_epilogue_fusionallow_prologue_fusion) rrqrrrrr first_inputr!rs ` @rrzTemplateBuffer.__init__sb d6222"11&99 "4G++D11  ""4(((+-9;%?A# #/ 6D   B  -68  %+a.Kk622 E ED4E4E E E2 ++--F%%%%%)%%%D ! " 1Z\\ " 37"26"""rrct|jtrtd|jS)Nz1Multi-output templates do not have a single dtype)rrqMultiOutputLayoutrr'r rs rr zTemplateBuffer.dtypeLsA dk#4 5 5 %C   &&rrjc.|dS)NT normalize)rTrs rrlzTemplateBuffer.get_read_writesTs''$'777rr/rOrderedSet[dependencies.Dep]ct}|jD]}t|ttfsJt ||tjtsJt jjd fd }|tj | d|j z}|S) z(Build read dependencies from all inputs.r Sequence[Any]rrrct|dksJtj|Sr$)rrur%r)rrr inp_indexers rdummyz4TemplateBuffer._read_deps_from_inputs..dummy`s>6{{a''''x  E0B0BCCCrrr.rr2rr2rr) r@rrrvrrrqr rBrKrTr%rp)rr/rpinp_rawr5rr4s @@r_read_deps_from_inputsz%TemplateBuffer._read_deps_from_inputsWs.8ll{  Gg'@AA P P4== P PA,3Ccj&11 C C4 3C3C C C1*1133K D D D D D D D \5s||~~rY EE rFc^t|jtrtt j|tjdddg}t j | ||tddS|| d fd }t j ||d| }|xj| |zc_|S)aExtract read/write dependencies for this TemplateBuffer. When the layout is MultiOutputLayout (multi-output templates), the buffer itself has no data layout, so we cannot build an indexer. Instead, synthesize a trivial write dep and derive read deps from the named tensor inputs (``_named_inputs``). For single-output templates with a concrete layout, fall through to the standard path. rr) var_namesrN)rprrrrrr2rrrclt|dksJtj|dSNrfakerrurJrrrDrs rr5z1TemplateBuffer.extract_read_writes..dummy6v;;!####9T775>>6:: :rr.r6)rrqr,r@rKrrrr!rr8r'rBrTr%rp)rr/rr5depsrDrs @@rrTz"TemplateBuffer.extract_read_writesis8 dk#4 5 5 3= * q)9)9Rb44F *11)<<&LL  }}//##0022 ; ; ; ; ; ; ;/ 4==??B)    d11)<<<  rr{c$tjjSr)rrrrs rr}z!TemplateBuffer.get_reduction_sizes w{rr8cdSrrrs rrzz!TemplateBuffer.get_reduction_typerrcdSrrrs rrxzTemplateBuffer.should_allocaterrrrr9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]c2|gfdfSrr )rrrs rrz#TemplateBuffer.simplify_and_reorders&      rc6t|jtS)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrqr,rs ris_multi_outputs_templatez(TemplateBuffer.is_multi_outputs_templates$+'8999rrc|jSr)rrs rget_allowed_prologue_inpsz(TemplateBuffer.get_allowed_prologue_inpss ))rscheduler_noderc*|S)a3Return whether this template's aliasing/mutation blocks prologue fusion. The default preserves the scheduler's conservative behavior. External template subclasses may override this when they can prove a prologue producer only feeds independent, non-mutated template inputs. )has_aliasing_or_mutation)rrKs r,has_aliasing_or_mutation_for_prologue_fusionz;TemplateBuffer.has_aliasing_or_mutation_for_prologue_fusions66888r hook_outputsdict[str, str]FinalizeCodegenResult | NonecdS)aCalled after epilogue/prologue subgraph codegen with rendered hook outputs. ``hook_outputs`` maps placeholder keys (e.g. ````, ````) to the Triton code generated by Inductor for each fused subgraph. Return a ``FinalizeCodegenResult`` to provide custom source code and call metadata, or ``None`` to use the default codegen path. Nr)rrOs r_finalize_codegenz TemplateBuffer._finalize_codegens trrrrc:t|tr!t|jtr|jSt|}t|t r|j}t|jtr| |S)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrrp MultiOutputr realize_inputrrqr(rZ)rrrs rrealize_template_inputz%TemplateBuffer.realize_template_inputs b) $ $ BG[)I)I 7N++B// fj ) ) ! fm^ 4 4 #  " " " r)direct_alias_at_leafon_tensor_leafon_non_tensor_leaf template_buf structuredrrXdict[int, IRNode] | NonerYFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | NonerZCallable[[int], None] | Nonetuple[TensorBox, ...]citjdfd t|gS) zLWalk a structured output tree, creating MultiOutput nodes for tensor leaves.rrrlist[tuple[type, int]]rlist[TensorBox]c t|ttfrKg}t|D]7\}}||g|t ||f8|St }t|tjr r%| vr!t |gSt|}| vr |gStt||}|j|<  ||||t|}| |<|gS  |gSr)rrrrrrr7r-TensorrridrUFallbackKerneltensor_to_layoutr%r)rrrritemleaf_idxtidmorrX leaf_counterrZrY seen_outputsr[walks rroz0TemplateBuffer.build_multi_outputs..walks&4-00 +-(00NNGAtNN44.K.K$v,,9J.K#L#LMMMML))H&%,// 'NH8L,L,L%,,-A(-KLLMMjj,&&(-.. "33F;;\7FH 3BKKMMB!-"N2;;=="gxHHHr]]$& S!t !-""8,,,Ir)rrrrbrrc)rcountr) rr[r\rXrYrZrmrnros ` ```@@@rbuild_multi_outputsz"TemplateBuffer.build_multi_outputssu.0  ((             6TT*b))***r)NNN)rqr<rrrrrrrrrrrrrr)r/rrr0r)r/rrrjrrrr )rrrrrrEr)rKrrr)rOrPrrQ)rrrr) r[rr\rrXr]rYr^rZr_rr`)rrrrrrr rlr8rTr}rzrxrrHrJrNrSr3rWrqrrs@rrrs378<1507070707070707d'''X'8888$$$$$$L OS?C      ::::****9999       [  :>;?)+)+)+)+)+[)+)+)+)+)+rrcheZdZ ddfd Zed ddfd ZddZddZxZS)TritonTemplateBufferNrqr rrrCallable[_P, _T] | Nonerrrrrrct||||||jJ|j|ji|_d|_d|_dS)a NOTE:[TritonTemplates with multiple outputs] We want the ability for TritonTemplates to output multiple tensors. Triton kernels have no notion of outputs and this is done by creating tensors that are then mutated by the kernel. Currently our STORE_OUTPUT codegen doesn't support creating multinode outputs for triton templates. We work around this by creating an extra input buffer during the lowering and we mark them as mutated inputs. )rrN)rrrr$ subgraph_inps subgraph_outs)rrqrrrrrs rrzTritonTemplateBuffer.__init__sl"    )"7    y$$$)-DI(>%@D9=rFrrrct|}|jr|jng}|jr|jng}|D]}}t |t r$|t||;t |tr)|||y|J~|D]D}t |tr)|||@|JE|Sr) rrxrwrvrr updater)r)rrresrwrvrrrs rrxz)TritonTemplateBuffer.get_free_symbol_usessgg**=99.2.@H**b .2.@H**b   # #C#t$$ # +C??@@@@C(( # 333MBBCCCC{{{{  # #C#v&& # 333MBBCCCC{{{{ rrc|g|jSrr'rs rrz TritonTemplateBuffer.get_outputs3-t,--rrcd|jd}|S)NzTritonTemplateBuffer(layout=r;r,)rrs rrzTritonTemplateBuffer.__str__6s;T[;;; rr ) rqr rrrrtrrrrrrrrrr) rrrrrbrxrrrrs@rrsrss 378< >>>>>>><233$)43.....rrscleZdZdZdfd ZddZddZd dZddZddZ d!dZ d"dZ ddZ d#dZ xZS)$ ChoiceCallera1 Represents a possible choice used in autotune_process.py. During autotuning, self.benchmark() is first called to get benchmark result, and if this choice is selected, self.output_node() is called to get the output_node. Children classes: TritonTemplateCaller, CUTLASSTemplateCaller. rrrrrqr  descriptionrrct||_||_||_||_d|_d|_i|_d|_ d|_ i|_ i|_ dSr8) rrrrqrrfailed_benchmark_with_cudagraphsrr| decompositiondecomposition_kwargsconfig_patches)rrrrqrrs rrzChoiceCaller.__init__Gsy   &'! 05',./38<46!.0rrrrrr:c||jrtjfdStjrt fdStjd|idS)NcSrralgorsrrz(ChoiceCaller.benchmark..es TT4[rcSrrrsrrz(ChoiceCaller.benchmark..gs DD$Krrr) to_callablerr\benchmark_gpu_with_cuda_graphrJ/profile_bandwidth_with_do_bench_using_profilingrg benchmark)rrrrs `@rrzChoiceCaller.benchmarkbs!!  * R<=P=P=P=P=PQQ Q  A A+,?,?,?,?,?@@ @$T4%dKKKKrctrrrs r call_namezChoiceCaller.call_namejrrrctrrrs rrzChoiceCaller.to_callablemrrc*|S)z Hash key for the underlying kernel. By default, we assume there are no runtime params, so kernel hash key defaults to choice caller's hash key. )hash_keyrs rkernel_hash_keyzChoiceCaller.kernel_hash_keyps }}rctrrrs rrzChoiceCaller.hash_keywrrrctrrrs rrzChoiceCaller.output_nodezrr6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]ciS)zRInformation returned here is logged to the autotune log file when that is enabled.rrs r info_dictzChoiceCaller.info_dict}rrcdS)Nunsupported_choicerrs rautoheuristic_idzChoiceCaller.autoheuristic_ids##rcd|_dS)z Mark the choice as failed so that it can be removed later. Useful for when we decouple compilation and tuning. TN)rrs r mark_failedzChoiceCaller.mark_faileds  r) rrrrrqr rrrr)rrrrrr:r)rr)rr)rrr)rrrrrrrrrrrrrrrrs@rrr>s1111116LLLL""""""""""""""""$$$$rrceZdZddZdS)TritonTemplateCallerBaserrctrrrs rget_make_kernel_renderz/TritonTemplateCallerBase.get_make_kernel_renderrrN)rr)rrrrrrrrrs(""""""rrceZdZdZd!fd Zed"dZed#dZ d$d%dZe j d&dZ d'dZ d$d(dZ d)d ZxZS)*MultiTemplateBufferaG Represents a Buffer with multiple backing implementation choices. Choices can be TritonTemplates or ExternKernels. During scheduling if there is a potential epilogue we will benchmark each of the choices with the epilogue to determine an implementation. Otherwise, the fastest base choice will be chosen. rqr rrchoice_timings_fn1Callable[[int | None], dict[ChoiceCaller, float]]unfiltered_choiceslist[ChoiceCaller]rrrrct||d|||_i|_||_||_t d|D|_i|_dS)N)rqrrrc3K|]D}t|tp*t|tjjjo|jVEdSr)rrr-rselect_algorithmExternKernelCallerhas_out_variant)rchoices rrz/MultiTemplateBuffer.__init__..sh% %  v7 8 8 65?#C#VWW+* % % % % % % r) rr_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders)rrqrrrrrs rrzMultiTemplateBuffer.__init__s #"7    #4LN,> %!$% % - % % % " " <>!!!rrc|jS)z^ Are all possible choices TritonTemplates or Extern Kernels with out variants )rrs routput_plannablez$MultiTemplateBuffer.output_plannables %%rc|jSr)rrs rrzMultiTemplateBuffer.choicess }rN hint_overriderdict[ChoiceCaller, float]ch||jvr|||j|<|j|Sr)rr)rrs rchoice_timingsz"MultiTemplateBuffer.choice_timingss;  4 4 4262I2I-2X2XD  /#M22rcallerrrc#Kt|tjjjsJt ||j|jksJ|j}||_ dV||_dS#||_wxYwr) rr-rrTritonTemplateCallerrrqrr)rrrenders rswap_as_triton_callerz)MultiTemplateBuffer.swap_as_triton_callers EO4I     <<   {fm++++("("?"?"A"A - EEE&,D # # #fD # , , , ,s 1A>> Bc>t|tjjjsJt |||jjksJ| |jj ksJ| |_ dSr) rr-rrrrr%rqrrTrrr)rrs rfinalize_as_triton_callerz-MultiTemplateBuffer.finalize_as_triton_callers EO4I     <<   }}&-"44444  FM$88888"("?"?"A"Artuple[ChoiceCaller, float]cn||}t||j}|||fS)N)rr)rrbr)rrtimings min_choices rget_min_choicez"MultiTemplateBuffer.get_min_choices>%%M%BBgk222 GJ/00rcallers*dict[int | None, TritonTemplateCallerBase]c|D]!\}}||j|<"|jd|_dS)z;Finalize with multiple callers for different hint overridesN)rrrr)rrrrs rfinalize_as_triton_callersz.MultiTemplateBuffer.finalize_as_triton_callerssV&-]]__ W W !M67=7T7T7V7VD %m 4 4#'";D"Ar) rqr rrrrrrrrrrr)rrr)rrrr)rrrr)rrrr)rrrr)rrrr)rrrrrrrrrrrrrrrrrs@rrrs>>>>>>8&&&X& X+/33333 - - - -BBBB+/11111BBBBBBBBrrc0eZdZdfd ZddZddZxZS)CUTLASSTemplateBufferrqr rrrCallable[_P, _T]workspace_sizertemplater{supports_epilogue_fusionrrrcxt|||||_||_||_dSr)rrrrr)rrqrrrrrrs rrzCUTLASSTemplateBuffer.__init__s? );<<<,  (@%%%rc"|j|jndSr$rrs rrz(CUTLASSTemplateBuffer.get_workspace_sizes&*&9&Et""1Lrc|D]*}tj|dd+dSr)rrurJr)rrs remulate_store_fnz&CUTLASSTemplateBuffer.emulate_store_fnsH&&(( 5 5F Ifoo''t 4 4 4 4 5 5r)rqr rrrrrrrr{rrrrrr)rrrrrrrrs@rrrso A A A A A AMMMM55555555rrc,eZdZdfd Zdfd ZxZS)CppTemplateBufferrqr rrrrrr{rrrrcxt|||||_||_d|_dSr)rrrroutputs)rrqrrrrrs rrzCppTemplateBuffer.__init__ s; );<<<   ,0 rct|jtrt|jtsJt |j|jd}t|t sJt ||j}t|tsJt ||St Sr$) rrqr,rrrrr rr')r first_outputrqrs rr'zCppTemplateBuffer.get_layouts dk#4 5 5 (dlH55 I ItDL7I7I I I5<?LlF33 G GT,5G5G G G3!(Fff-- ; ;tF|| ; ;-M77%%'' 'r) rqr rrrrrr{rrrrr)rrrrr'rrs@rrr s[ 1 1 1 1 1 1 ( ( ( ( ( ( ( ( ( (rrc0eZdZdZ ddfd ZddZxZS)CuteDSLTemplateBufferz Buffer for CuteDSL (CUTLASS Python DSL) template kernels. Similar to other template buffers but specialized for CuteDSL operations. Nrqr rrrrrrrrrrct||||_|_g_|t jdts"Jtjdjd xjfd|Dz c_dSdS)NrcNg|]!}tt|"Srrr s rrz2CuteDSLTemplateBuffer.__init__..:s@z888#tDDr) rrrrrrrrrr+)rrqrrrrr!rs` @rrzCuteDSLTemplateBuffer.__init__*s );<<<  ,&*V  %dk!nf55 K KtDKN7K7K K K5[^..00F LL) LLLL & %rrc|jSrrrs rrz!CuteDSLTemplateBuffer.get_outputs? |rr) rqr rrrrrrrrrrr)rrrrrrrrs@rrr$se37 *rrcLeZdZdZ ddfd Zd dZd!dZ d"d#dZxZS)$NVUniversalGemmBufferz Buffer for NVIDIA Universal GEMM kernels. Unlike CuteDSL templates which use Jinja templates, this generates simpler Python code that directly calls the cutlass_api library. rNrqr rrkernelraccumulator_typevariantrr scale_type_a Any | None scale_type_bswizzle_type_aswizzle_type_brrc $t||d||_||_|g|_||_||_||_||_| |_ | |_ |j j |j j d|_|j|_dS)N)r) kernel_namemin_cc)rrrrrrrrrrrmetadatarrkernel_metadata_make_kernel_renderr) rrqrrrrrrrrrrs rrzNVUniversalGemmBuffer.__init__Ks DAAA 0&*V , ((,,"?6o,   #'":rc|jS)z#Return the workspace size in bytes.rrs rrz(NVUniversalGemmBuffer.get_workspace_sizels ""rrc|jSrrrs rrz!NVUniversalGemmBuffer.get_outputsprrout_noderrtuple[Any, Any]c  ddlm}ddlm}g}|jD]O}t |t r|j}t |tr|j}| |Pt|j }|||||j |j |j|j|j|j|j|j fd} |fS)z Create a kernel renderer for code generation. Returns (kernel, render) tuple where: - kernel: NVUniversalGemmKernel object with call_kernel() method - render: function that returns source code string r)NVUniversalGemmKernel) Placeholder) rrrrrrrrrrrc,Sr)r) render_kernelsrrz9NVUniversalGemmBuffer._make_kernel_render..renders '')) )r)Btorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelrtorch._inductor.utilsrrrrrprrr KERNEL_NAMErrrrrrrr) rrrrrrrrrrs @rrz)NVUniversalGemmBuffer._make_kernel_renderss       655555!# ; $ $C#y)) h#z** h   s # # # #+122 --##  0!2.L**..     * * * * *f$$r)rNNNN)rqr rrrrrrrrrrrrrrrrrrrrrrr)rrrrrr) rrrrrrrrrrs@rrrCs #'#'%)%);;;;;;;B####:>*%*%*%*%*%*%*%*%*%rrr#Sequence[IRNode | Sequence[IRNode]]TypeIs[Sequence[IRNode]]c4td|DS)Nc3@K|]}t|tVdSrrrrs rrz#is_node_sequence..s,44z!V$$444444r)r)rs ris_node_sequencer s! 44e444 4 44rceZdZUded<ddZdd Zdd ZeddZe ddZ ddZ ddZ e d d d!dZdS)"rr rrrrrcr|j|}t|tsJ|Srrrrr)rrinputs r input_namezInputsKernel.input_names3 A%(((((~~rrjc&ttj}tj|jD]}t |t r"|fd|D9t |trO| | ttjfd| D}tj ||tS)Nc3RK|]!}|V"dSrrn)rrrs rrz/InputsKernel.get_read_writes..s5BBqWWQZZ\\22BBBBBBrc3RK|]!}|V"dSrrn)rrrs rrz/InputsKernel.get_read_writes..sF. . (+GGCLLNN # #. . . . . . rr) r@rKrQrrrrryrrrrr)rrpr rrs @rrlzInputsKernel.get_read_writess<+,..&[ 5 5E%** 5 BBBBEBBBBBBBE#899 5 ''%.."2"2334444L,-. . . . /3/?/?/A/A. . .   &"     rrmc4|jSrrors rrzInputsKernel.get_readsrqrrrct|tr|j}t|tr|j}t|tr/t|t st |}t|tr||St|tr|St|tt fsJt||Sr) rrrprrrvrrVunwrap_storage_for_inputTorchBindObjectrrrrs rr z%InputsKernel.unwrap_storage_for_inputs a # # A a $ $ A a " " .:a+I+I .**1--A a # # 3 //22 2 a ) ) H!fo677@@a@@7rlist[IRNode | Sequence[IRNode]]cg}|D]S}t|tr d|D}nt|}||T|S)NcBg|]}t|Sr)rr rs rrz/InputsKernel.unwrap_storage..s&III!\::1==IIIr)rrrr r)r inputs_newrs rr!zInputsKernel.unwrap_storageso79  ! !A!X&& =IIqIII 99!<<   a rrcdSrrrs rrzInputsKernel.is_externrrcdSrrrs rrszInputsKernel.num_readsr4rFrrcttj}|jD]M}t |t r|||z}0|D]}|||z}N|Sr)r@rr"rrrrx)rrrr inner_inps rrxz!InputsKernel.get_free_symbol_usess u| $ & &; G GC#v&& GS--m<<<!$GGI77 FFFAAGrN)rrrrrrr)rr rr rrrr)rrrrr rlrr3r rr!rrsrbrxrrrrrs////        ,,,,,[$   \ N++$)    ,+   rrceZdZddZddZdS) NopKernelrrcdSrrrs rrzNopKernel.is_no_oprrrmctSrr?rs rrzNopKernel.get_readsrrNrr)rrrrrrrrr r s<rr ceZdZdZeddZe dddZed dddZeddZ ddZ d S) ConcatKernelzn There isn't actually a real kernel for concat, we just change the storage for the upstream data. rrrcrrrc  |d}|d}t|d}dg}||g}d|cxkrt |ksnJt dt |D]#}||} |||t | t |ksJ|||ksJ|||ksJt t |D]Q} | |kr|| | | z|| <tjj || | | || <R|||%t |} tjr't| ||dj} t t |D]x}||} t%| r_| } t)| t*r6t| j| jrt3|} nyt5d|D}tjjjd}|dur=t)|tr(t5d|Drt3|} t;d|D}|Jt=dt+|||| |g }t?|}g}tA|D]\}}t)|tBtDfsJtG||$|tJ&||||||d }t)|tNsJtG|t)|j(tsJtG|j(|j(|t)|j)tBr|j)*}n|j)}t)|t>rt|+r`|x}JtY|j#r6t[|s'||.t |dkrItj/|t`j1rtj2|tj3||_4|5|j(|_(tj6||S) z6 Create the concat kernel from inputs rrIc34K|]}t|VdSr)r&rs rrz&ConcatKernel.create..6s+-W-W1.CA.F.F-W-W-W-W-W-WrFc3K|]Y}d|jvoKt|jdtjp%t|jdtjVZdS)r"rN)rr/r-rrrargs rrz&ConcatKernel.create..=s! <&+&9  @&+&<rc3fK|],}t|o|jV-dSr)r&r'rurs rrz&ConcatKernel.create..QsL  FG !! $ $ A)A      rN)r!r rrrurrqrrR)7r+r*rr%rrrrwr"r#rLr(r)rJrr rCr r&r'rrsrrrr1ra current_noderrr" rrrrCrrNrBrrrrprrJrnrmrrrLFOREACHregister_operation_listr"rr!r#)rrrcr!r r offsets_start offsets_endr input_sizer output_striderrqany_input_is_storage_and_layout fx_node_argsru concat_kernelrop_namesrr!input_unwrappeddevs rrzConcatKernel.create s %%''q ##%%q **,,-- }o C''''#h--''''''q#f++&& . .A++--J  # / / /z??c(mm3333!9&&((E1111!9''))V33333x==))  88"*1+ 1 "=HQKK"#'"2"L"L  Z]##HQKK   x} - - - -'5'H'H'R'R  ' "//xM s6{{##  Aq A$Q'' K88fmTT%C8$L$LME*--W-WPV-W-W-W*W*W'w+03  ,u 4 4<.. 5( 5(;8DDM  KQ     !!!$$#     M**'' C CFAscHj#9:: E EDII E E:++  Cq!1;q>!L lF33 G GT,5G5G G G3m2D99 U U4 @T;U;U U U9  ' ' 5 5 5#(H-- +"%("6"6"8"8"%(?J77 C#3355 C NN,,,S938$$:"<00: ? ? A ABBB x==1  !4!4V^=S!T!T  G + +H 5 5 5W44]CC "11-2FGG  ""=111 rNrFrrGrrct|tr||j|St|tt fsJt |t|jtrt|jjtr |jj sdS|dSt| t| krdStdt| | DSt|jdo9t|jjt ot|jt" S)NFTc3fK|],\}}tjj||V-dSrrlrns rrz=ConcatKernel.can_realize_into_without_copy..sOB 88R@@rrq)rrcan_realize_into_without_copyrprrrrrqrsrrrTrrrr(ExternKernelAlloc)rrFrGs rr; z*ConcatKernel.can_realize_into_without_copysi c9 % % D44SXsCC C#*566AAS AA6 ch 3 4 4 sx << x0 u{t3>>##$$CNN,<,<(=(===u!#.."2"2CNN4D4DEE  CHh ' ' <38?N;; <sx):;;; rFrrc8t||Sr)r rxrws rrxz!ConcatKernel.get_free_symbol_usess--dMBBBrc t|ts2t|r#t|\}}t||}t|tsJt |t|t r||j|St|tra| t|jdsJ| ||r t||j_ |jSt|||dt'||D}|||S)Nrorqc^g|]*\}}tjj||+SrrKrMs rrz-ConcatKernel.realize_into..sAAq ::1a@@rrR)rrvr&rrrrrNrprr*rr; rrqr=rr+r*r>rr%)rrFrGrwrqpws rrNzConcatKernel.realize_intos #// C$S)) C"7"<"<%76BBB#//::c::/ c9 % % 3##CHc22 2 c: & & KKMMM38X.. . ..00c:: "1#"6"6x   >>##--//__&&  ??   C(((rcdSrrrs rrxzConcatKernel.should_allocaterr)rrrcrrrr)rFrrGrrrrr)rFrrGrrrr) rrrrr3rr; rbrxrNrxrrrr" r" s yyy[yv/3! ! ! ! [! FN++$)CCCC,+C )))[)Brr" r<rc#K|d} dV||dS#||wxYwr8)"_ignore_fresh_unbacked_symbols_set)rprevs r_track_fresh_unbacked_symbolsrE s[  7 7 > >D; 44T::::: 44T::::s 4A AbstractContextManager[None]ch|#|rt|StSr)"_ignore_fresh_unbacked_symbols_tlsrE r )rs r(_fallback_kernel_symbol_tracking_contextrI s3!M!M!O!O,Y777 ==rceZdZUdZdZded<ejeZ ded<dZ d ed <dZ d ed <dZ d ed <eje Zded<dZded<dZded<ejeZded<dZded<ejeZded<eje Zded< dkdlfd$ Zdmd&Zdnd(Zdofd* Zdpd+Zdpd,Z dqdrd0Zdsd1Zdqdtd2Zdud3Zdvd4Zdwd6Z e!dxd:Z"e#dyd@Z$e#dzdBZ%e#d{dCZ&e#d{dDZ'e# d|d}dLZ(e# d~ddNZ)e# d~ddPZ*e#d{dQZ+e#d{dRZ,e#d{dSZ-e#d{dTZ.dpdUZ/ddVZ0dqddZZ1dd[Z2dd]Z3d~dd_Z4dwd`Z5dsdaZ6dsdbZ7dsdcZ8ddeZ9ddgZ:e;d d~ddiZxZ?S)rz A class that represents Kernels which are not directly lowered to Inductor Loop Level IR, such as custom operators, or aten operators which we fallback to. rr2 constant_argsr1rrNReinterpretView | None output_viewr8python_kernel_namecpp_kernel_name Iterable[str]ordered_kwargs_for_cpp_kernel_OpOverloads | None op_overloadzlist[dict[str, Any]] | Nonearg_propertieszdict[str, dict[str, Any]]allarg_propertiesz dict[str, dict[str, Any]] | Nonekwarg_propertiesz"dict[sympy.Symbol, pytree.KeyPath]unbacked_bindingszlist[MutationOutput]r'rrqr<rr dict[str, Any] | Nonerrc pt|||||_|r|ni|_||_| |_||||| |_| i|_ g|_ tj j|_i|_dSNr* )rrrK rrM rS set_cpp_kernel_nameset_python_kernel_namerQ collect_arg_kwarg_propertiesrW r'rwr"r, fx_noder) rrrqrrK rrM rN rO rQ rS rs rrzExternKernel.__init__s     + &.ffB &&   111 ##$6777-J* ))+++!# "w+ +-rrc|g|jSrr|rs rrzExternKernel.get_outputs#r}rrctSrr?rs rrz%ExternKernel.get_unbacked_symbol_defs&rrrjctdfd }tj||j|jfdS) Nr1rrrct|trD|}|0jt j|dSdSdSr)rrrGrprrKr)r1rrs r add_ir_readz1ExternKernel.get_read_writes..add_ir_read,sg%(( F++--#%)),*>t*D*DEEEEE F F##rc,t|tSrr r5s rrz.ExternKernel.get_read_writes..5s*UF";";r)is_leaf)r1rrr)rrlpytree tree_map_rK r)rrc rrs @rrlzExternKernel.get_read_writes)swgg--//  F F F F F F     -;;    rcXt|jtjjrd|jjjDn*dtt|j D|_ t|jtjjrd|jjjDni|_ t|jtjjrI|j s d|jjjD|_ d|jjjD|_ dSg|_ dS)NcFg|]}|j |j|j|jdS))rrrS) kwarg_onlyr real_typerSrs rrz=ExternKernel.collect_arg_kwarg_properties..=sH    | FK%&_   rcg|]}iSrrrs rrz=ExternKernel.collect_arg_kwarg_properties..Gs666"666rc8i|]}|j|j|jdS))rrS)rrk rSrs rrz=ExternKernel.collect_arg_kwarg_properties..Js8   qOO   rc*g|]}|j |jSrrj rrs rrz=ExternKernel.collect_arg_kwarg_properties..Us1666 !,6F666rc g|] }|j | Sr)rj rs rrz=ExternKernel.collect_arg_kwarg_properties..Xs/"""""""r)rrS r-_ops OpOverload_schema argumentsrrrrT rU rQ schema_kwargsrs rr] z)ExternKernel.collect_arg_kwarg_properties9s[$*EJ,ABB 7   )1;     76eC $4$455666 $$*EJ,ABB    )1;       d& (= > > $5 66$($4$<$F6662""+3="""D   "$D   rct|jtr*||dSdSr)rrqr(apply_constraintrZrs rrzExternKernel.decide_layout^sJ dk> 2 2 !  ! ! # # #     ! !rwrapperr|rct||\}}|r|||s|}|r,ddlm}|||d}|||dSdS)NrI)'set_kernel_post_grad_provenance_tracingT)r)rj make_commenttry_get_kernel_namerrz write_provenance_debug_handle)rrx r origin_str_detailed_origin_strrz debug_handles rcodegen_commentzExternKernel.codegen_commentcs,?tW+M+M( (  -   , , , 52244K  M F F F F F FBBkTL  1 1+| L L L L L  M Mrctrrrrx s rcodegenzExternKernel.codegentrrc||_tjjr$t |jt jjsdS|j}|j|j dkr|j dkr |j ddn|j dd}ddlm}|d|j i}t#d|Dd }|d kr|d |}d |d |_dS|jj|_dSdS)Natenr\.rr inductor_fallback_opszaten.c3nK|]0}|dt|ddV1dS)rrIN) startswithrrs rrz3ExternKernel.set_cpp_kernel_name..sAKKAc9J9JKS122ZZKKKKKKrrIr_vz at::_ops::z::call)rO rwr" cpp_wrapperrrS r-rq rr namespace _overloadnamerrreplacetorchgen.aoti.fallback_opsr rrrs r)rrO ropnamer version_infolatest_versions rr[ z ExternKernel.set_cpp_kernel_namewsT.w" *  ej3+ +   F!   '6))+y88O))#..q1100c::MLLLLL4889R9R9RTVWW !$KKKKK""""A%% &::.::F'BF'B'B'B$$$'-~':$$$7 ( 'rc||_|dS|j}|dSt|tjjrd|j|_dS|jddd|j|_dS)Nztorch.ops.higher_order.._ops..ops.r ) rN rS rr-rq HigherOrderOperatorrrr )rrN rs rr\ z#ExternKernel.set_python_kernel_names"4  ) F! > D  > ? ? &Q&Q&QD # # #$,,Xw??SS&/SS  # # #rcddlm}|x}r|jntjj}tjjr|jStjj rxttjj |s&Jttjj |j dStjj |j |S|jS)NrI) CppWrapperCpu)codegen.cpp_wrapper_cpur r+rrwr" device_type fx_wrapperrN r rr0rO get_c_shim_func_name)rr dr!s rr| z ExternKernel.try_get_kernel_names::::::!%!2!22AL9L 7  +* * W +ag2MBB  D$EE  B#+t7'<<$f * *rrc6|}|J|Sr)r| rs rget_kernel_namezExternKernel.get_kernel_names$'')) rrrrc Ft||||||}||S)N)r!r rrrr) r=rr+r*r>r%rrr*)rr@ s r copy_inputzExternKernel.copy_inputsx   <<>>++--]]__::<<))++oo''     rrrrrftuple[Any, list[Any], list[Any], Callable[[Any, Any], Any], dict[sympy.Symbol, pytree.KeyPath] | None]cX ||d}tj|\}gg}g}g}|D]} | xtdxn\tjjj| d} d|| || }xtdx\d|| | j j } | j j dkr| J|tjj| xt"dxI\d|| || jit&dx.\d||  d|| || dfd } fd|D}|D]"} t)| rt+| d#g}|D]} t-| t.sb| tjjvr=|tjj| zt-| t.sb| tjjvr=|tjj| t-| t6r)|| /t-| t:r|| j`t-| tjj j r^| j j } | j j dkr| J|tjj| |tC| | ||\}}tj"j}tG5tI|5||i|}dddn #1swxYwYdddn #1swxYwYd}|rtj%j&'d}tQ}tj%j)tj*j+j,ur!|d}t[tj%}|5t]|tj%|dddn #1swxYwYt_|||}t-|t`tbfs|gn|}|D]u}t-|tj2rY|j3rRthj5sFd}tjj%j&'ddx}r|d|}|tj_6v|||| |fS)aPartition kernel args into tensor and non-tensor, realize tensor inputs, re-run fake tensor propagation with the realized strides, and return (example_output, tensor_args, non_tensor_args, unflatten_args, unbacked_bindings). unflatten_args(new_tensor_args, new_non_tensor_args) reconstructs the original (args, kwargs) tree from replacement lists. )rrrN)rEFrATnew_tensor_argsrnew_non_tensor_argsrtuple[list[_T], dict[str, _T]]c^g}t|}t|}D]I}|r#|t|'|t|Jtj|}|dg|difS)Nrr)r8rr7rf tree_unflattenr) r r r it_tensorsit_non_tensors is_tensorrargs_flat_is_tensor args_specs runflatten_argsz3ExternKernel.process_kernel..unflatten_argssFo..J!"566N0 8 8 8MM$z"2"23333MM$~"6"67777%fi88A55$$aeeHb&9&99 9rc:g|]}|SrrVrrrs rrz/ExternKernel.process_kernel..s'AAAs((++AAArr>r"rIzEsparsity not handled. Please file issue for sparse inference weights.rz Found from : )r rr rrr )7rf tree_flattenr rwr"r#rcreate_symintnoderGeneratorStater!rrr-rAdefault_generators clone_stateOpaqueObjectStater1rr&rrrrr constantstorchbind_constantsr get_valueropaque_example_valuerirrrr#rI r, rrr r_higher_order_opseffects with_effectsr3r:r4rrre is_sparserJgraph_partitiondisable_cudagraphs_reason)rrrr binded_args args_flat tensor_argsnon_tensor_argsreal_non_tensor_argsr( r device_indexr r example_argsnew_args new_kwargsrexample_outputrW node_meta_valctxexample_out_lir0msgrr r s` @@rprocess_kernelzExternKernel.process_kernels;" $v66 %2;?? 9*,$& (*   5 5CTVVVVVV7+5GGRVGWWD'..u555#**4000(//5555  &^%%%%%%'..u555#**3///#&:#3L:?f449Q9QQ(// 5lCOOQQ &)&(((((('..u555#**3///(// ::::) XXXXX'..t444&&s++++'..u555#**3///(//4444 : : : : : : :BAAA[AAA  6 6A$Q'' 6%a5555   : :Aa** :qzz||qw?P/P/P##AG$5ajjll$CDDDDq(++ :JJLLAG$???##AG$? $MNNNNA// :##AKKMM2222A011 :##A$:;;;;Au1@AA : x~ x}..<3K3KK##J1,?KKMM##$5a$8$89999-~l>JJJ K K K K K K K K K K K K K K K 9>=!!  ntUm<< ^     8 8A1el++ 8K 8. 8 ^"#'"6";"?"? t"T"TT;A @@;@@C471        sHR= R& R=&R* *R=-R* .R==SSU00U47U4rvc t|tsJt|t|tr|S|}t j|}|J| }|d|j vrt|tttfrt|j trt|j dt js&t|j dt jr5|t)|n|t/j|d\}}|d}||}t jj||}t jj||} t jj||} t=|| | z} || kr$t> d| | |tBt|j"tG|$|%|| | d  S) z In order to pass this to an extern kernel we need a ReinterpretView not a View. This allows us to avoid some unneeded copies. Nr"r& rrrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFr,ro)&rrrrvrrwr"r;rrrrrCrqr(r/r-rrrer1r%rZrKrrBr#r stride_vars offset_varrorrrrprsr6r*) rr x_unwrap_viewrx_unwrap_view_fx_node index_argsrrrrhrtexpecteds rconvert_to_reinterpret_viewz(ExternKernel.convert_to_reinterpret_viewus!X&&//Q//& a ) ) H  g  !7!7!9!9:: # 3 3 5 5 " -.333=?FJ*OPP4=/@@49).u5"'"5 4<).u5"'"84  7 7.}/E/E/G/GHH      ' ' ) ) )!-!@ JJLL" " "  J ]    ,, 55eZHH'"..ujAA!,,UJ??Z11F: H   IIR     & %,,..kkmmZZ\\     rc |tSt|ttjjjtfrt|St|trt5tj tj|j||cdddS#1swxYwYt|t&r|St|t(r||jSt|t.r;t/||j|St|t2r[|t7|r& ||S#t<$rYnwxYwt|t>r||St|t@ttBfr|S|"|S)N)r)r r!ro)#rrr rrrrrrrfrArwr"add_tensor_constantr-rgr1r*r+rSrrVrprvr'rr*r&rr rr NonTensorObjrr r s rrVzExternKernel.realize_inputsI 9')) ) a$ 3 ;SA B B 1(a000 0 a " " ())  w22L allnnUUU                 a ( ( H a # # -$$QV,, , a ) ) "&&qv..q||~~  a " "  IIKKK$Q]]__55 ::1===*D a $ $  IIKKKH a,(=?PQ R R H~~a   s%5ACC#&C# G G-,G-ct|rHt|dkr|S|D] }|dkr|cS ||Sr)r&rrTr )rrrs rrequire_stride1zExternKernel.require_stride1sr  # # 1<<>>""a'',,..  Q;;HHH~~a   rFrSequence[int] | Nonerfrr[rcL  |JdvrsStrttrV|r=t |}|rntdd||Stddd|Sd}|rOrMtt   }tttfry|r'|s<rNt!j rt%nS|xtttfrJt!|j rt%|Stx}t&rt|x}trt+dt|trC|r||s*r*t!|j rStt,rg|r'|s<r.~sj33M!4DaHH11!**,,q/1EE rrrI)/r"r&rr'r(rUrZrrrr rCr%r*rsrrTrfrr{r9r>rrrrprrvrrr< r require_stride_orderrrrwr"r#rrr-rloweringslice_r r6rTr)rrrrfr[use_current_stride_orderpadded_exact_stridesmutation_layoutr>r expanded_dims orig_sizercr#s ` ` @rrequire_strideszExternKernel.require_stridess? M$=$== ;;==F " "= "H ! # #X !,,...99# "0R500,0 )))).#',1).*7 H*#(-%)&3&3 H $(   '+'' qzz||Q[[]]SS(($!,,..;*HII) ) <<>>;;EBB) ") 2%q||~~'>  .!1<<>>#8!**,, H q) $ $ 168,, qv77 &Q]]__&DkEE    V,,   {/1BCC   88@@335 4#44= 5'    +/ JJLL  $w'Hs1::<<0011M% B BO,33AsAqAA NN1   !''       E)!**,,77 E*!,,..99 E 6a?? ? ?? ?  E(]-F-FF(//9==A21mDD Ds18Q*Q QQr^cF||d|D|S)Nc^g|]*}t|tjr |jjn|+Sr)rr-SymIntrrrs rrz6ExternKernel.require_exact_strides..s>FGz!U\::A r)rfr[r )rrrfr[s rrz"ExternKernel.require_exact_stridessE"" KX( #   rrc2||||S)N)rr[r )rrrr[s rr z!ExternKernel.require_stride_orders""1E"OOOrc8||tSr)r rr s rrequire_channels_lastz"ExternKernel.require_channels_lasts''+<===rc8||tSr)r rr s rrequire_channels_last_3dz%ExternKernel.require_channels_last_3ds''+=>>>rcdd}||r|S||t|S)Nrrrrc |}n#ttf$rYdSwxYw|tjjvotjj|jSr8)rAttributeErrorrrwr"r is_mkldnn)rrs ris_mkldnn_tensorz9ExternKernel.require_contiguous..is_mkldnn_tensorsb zz||"$78   uu 17,,R1B41H1R Rs ,,rrrrrr(r)r%)rrr s rrzExternKernel.require_contiguoussg S S S S  A   H,,>44QZZ\\BB rc||t|Srr r s rrequire_contiguous_stridesz'ExternKernel.require_contiguous_stridess7(( ~00>>   rcdSrrrs rrw zExternKernel.apply_constraintrWrct|tsJt|t|tst|}|js Jdt |}t |j}||kr|t d|j||z t||D]G}|j|d}| ||vr||n|j|dH|S)Nz/ExternKernel.arg_properties should not be emptyzv%s has %d unprovided positional arguments. Will check if they are in the keyword arguments or will use default values.rrS) rrrrrT rrrrS rr)rrrn_args n_pos_argsrarg_names rfill_non_provided_argsz#ExternKernel.fill_non_provided_argss&$))554::55)$%% ::D"UU$UUU"T,--  J   II^ V#     6:..  .q1&9 6))8$$,Q/@  rrrrctjjr4g}d}|rG|jr@t |jt |ks Jdd|jD}t |jD]\}}|9|J|||}|r|dnd}nXt |j|z}|jr8|t |jkr |j|dnd}| tjj |||Sd|jDS)NzDnames passed to codegen_const_args does not match self.constant_argsc:i|]}|d|Sr)rr' s rrz3ExternKernel.codegen_const_args.. s1***-0CGGFOOS***rrcVg|]&}tjj|'Srrwr"r0val_to_arg_str)rrss rrz3ExternKernel.codegen_const_args..s+WWWqAG(77::WWWr) rwr"r rT rrK rrrrr0r ) rrrname_to_arg_propertiesrrproptype_rs rcodegen_const_argszExternKernel.codegen_const_argss 7  XF &* " , 4-..#e**<<<Z=<<**484G***&"$"455 M M1)5 ,,,155eAh??D04>DHHV,,,$EEdk**Q.C ."36T=P9Q9Q3Q3Q+C044V<<<!  ag2AA!UKKLLLLMWWDDVWWW Wrctjjr3|j,|g|j|j|j}d}n |j}d}g}t|D]\}}tjjr|j r|t|j ks Jd|j | d}| tjj ||| tjj ||r'|||S)NFTz-Invalid access to ExternKernel.arg_propertiesr)rwr"r rS r rrK rrrT rrrr0r rr )rrneed_codegen_constant_argsrrrr s r codegen_argszExternKernel.codegen_argssV 7  .4#3#?003$+3 23T[F*/ & &[F)- &f%% D DDAqw" D*q3t7J3K3K/K/K/KC0L/KK+A.226:: AG0??5IIJJJJ AG0??BBCCCC % 3 KK//11 2 2 2 rr c ||vr||S||jvr|j|S|j|x}|dSt|d)zGiven an argument name, queries for values in (in order): 1. any provided kwargs for this function. 2. the class self.kwargs member. 3. any available default arguments in self.allarg_properties.NrSz not in self.allarg_properties)rrrU r)rr rr( s rget_kwargs_valuezExternKernel.get_kwargs_value8s v  ::h'' ' t{ " ";??8,, ,)--h77 7C D77?++ +HHHIIIrskip_outc:tjjr|jt |jdkrgSg}|jD]}|r|dkr ||}t|tr| |K|j J|j |i d}| tjj ||n#d|jD}|S)Nrrrcfg|].\}}|dtjj|/Srr )rkrs rrz/ExternKernel.codegen_kwargs..YsMAq??qw+::1==??r)rwr"r rS rru rQ r rr rrU rr0r rr)rr rr rr s rcodegen_kwargszExternKernel.codegen_kwargsEs3 7  +D4F0G0G10L0L F > Q QE 1 1))(33a&&QMM!$$$$1=== 266xDDHHPPEMM!'"6"E"Ea"O"OPPPP Q K--//F rc|jW|jj}t|dd}|dd}|ddd}|d|}nd}|S) Nrunknown_namespacer r r rIr unknown_op)r^ rrr rsplit)rr op_namespaceop_names r get_op_namezExternKernel.get_op_name_st < #\(F"6<9LMML'//'BBL'..sA66q9L%0000GG"GrctjsdSt|dkrdStjj|}tjj|}| }| }tjj rt|j tjjr^tjj|j jvrAt|jdt*sJ|jd }|||||dSr$)rJ size_assertsrrr%rwr"r0codegen_shape_tuplerTr rr rrS r-rq rr Tag inplace_viewtagsrrwrite_assert_size_stride)rrx rrr rs rcodegen_size_assertsz!ExternKernel.codegen_size_assertsjs$"  F  ) )Q . . Fw#77 HH%99$//:K:KLL""$$}} 7  5$*EJ,ABB 59)T-=-BBB%dk!nf=====;q>2244D((tVWEEEEErc Ftjrtjjs|}|tjjv}|}|r&|d|dtd|ddS|d|d|ddSdSdS)Nzassert_alignment(rr;z # buffer z (op: z) is assumed to be not aligned) rJalignment_assertsrwr"r rrYr r2rk)rrx ralignedr s rcodegen_alignment_assertsz&ExternKernel.codegen_alignment_asserts}s  # AG,? ==??D!'";;G&&((G !!OOOOO7OOO!!SSSGSSS    rctjjrtjjrdS||}|d|d|ddS)zc Track outputs of fallback operators if config.test_configs.track_memory_lifecycle Nz track_tensor(z, 'z')) rJ test_configstrack_memory_lifecyclerwr"r "write_memory_track_allocation_oncerr2)rrx rs rcodegen_memory_trackingz$ExternKernel.codegen_memory_trackingsp"9 QW=P  F22444}};$;;4;;;<<<<.s'BBBQ8--a00BBBrc2g|]}td|S)r )rprs rrz-ExternKernel.canonicalize..s'MMMa(Q11MMMrT)rr ci|]\}}|| Srrrs rrz-ExternKernel.canonicalize..sBBBxsC#sBBBrc g|] }| Srrrs rrz-ExternKernel.canonicalize..s777q777rc g|] }| Srr)rrrs rrz-ExternKernel.canonicalize..s333jm333rcc&g|] }|Srr)rradd_vars rrz-ExternKernel.canonicalize..s!3R3R3R1GGAJJ3R3R3Rr)rwr"r#r%rTrrrrrrBrrVrrrsrrr)rrrh index_orderrrDr new_sizesrrr  replacementr< rrr#s @@@@r canonicalizezExternKernel.canonicalizes 7# //##CBBB'BBBMM5U;L;LMMM U3w<<00g6ISWXXX BB9[+A+ABBB7777E#f++$6$67773333U333 ##%% ##%&W%5%E%E w& & " 7F !%% 73z773R3R3R3R 3R3R3R+S+STTUU 5<.. <<eI&&&&rrc|rtnt}t||}|jD]}|||z}|jD]}|||z}|Sr)maybe_free_unbacked_symbolsmaybe_free_symbolsrrxrK rr)rrmaybe_get_symbolsrr( s rrxz!ExternKernel.get_free_symbol_usess,9 P ' '>P   - -dM B B% ( (C ""3'' 'AA;%%'' ( (C ""3'' 'AArctdd}d|g}|fdtjDz }|dj|S)NrN zpython_kernel_name=cNg|]!}|jdt|j"Sr)rr)rrrs rrz(ExternKernel.__str__..sE   z 7 7GD%*55 7 7   rr)rrfieldsrrr)rrrs` rrzExternKernel.__str__sd$8$?? 1+ 1 1      $+D11      8D$488999u%%%rrNNNNrN)rr8rqr<rr rK r2rrX rM rL rN r8rO r8rQ rP rS rR rrrrrrr)rx r|rr8rrrx r|rrrO r8rr)rN r8rrrr)rrrr)rrrrrrrr )rrrrvr)NNF) rrrr rfrr[rrrr)rrrfr^r[rrr)rrrrr[rrr)rr2rrrr2)rrrrrr)r rrrrr)r rrr)rr0 )rr3 r)@rrrrrK rrrrrrM rN rO rrQ rS rT rU rV rW r'rrrrlr] rr r r[ r\ r| r rr r3r r rVr r rr r r rr rw r r r r r r r& r* r/ r2 r@ rbrxrr2rrs@rrrs $&M%%%%.[.tDDDFDDDD*.K....%)))))"&O&&&&4E;3D444!(,K++++26N66663D;3D444:>====t-T-T-TTTTT(*(,.2)-&*79+/.......@.... #$#$#$#$J!!!! HLMMMMM"""""#;#;#;#;#;J    ++++$    \ c c c [c JE E E [E N!!!!!![!!F!!![!'+37# tttt[tlQV     [  DIPPPP[P >>>[>???[?["   [     """"HXXXXXB4 J J J J4    FFFF&     = = = =$$$$''''BN++$)    ,+  & & & &HHHHHrrc@eZdZddZ ddfd ZddZxZS)ExternKernelOutrx r|rrc0||dSr)generate_extern_kernel_outr s rr zExternKernelOut.codegen**400000rrNrqr rrrK r2rrX rM rL rN r8rO rQ rS rR c h||} t| tsJt| t d|| ||pid|||| t j||_ t j |dSr) r!rrrrrrwr"r"rr#) rrqrrK rrM rN rO rQ rS unwrapped_inputsrs rrzExternKernelOut.__init__s ..v66*H55MMt ,      r)rprr!r2rrrrrrrrs@rrU rU s=          rrU cFeZdZddZ ddfd ZddZddZxZS)r< rx r|rrc0||dSr)generate_extern_kernel_allocr s rr zExternKernelAlloc.codegens,,T22222rrNrqr<rrrK r2rrX rN r8rO rQ rS rR c ||} td| DsJtd|t t t | ||pid|||| g|_tj ||_ tj |dS)Nc3@K|]}t|tVdSrr rs rrz-ExternKernelAlloc.__init__..)s,CCQ:a((CCCCCCr) r!rrrrrrrrwr"r"rr#) rrqrrK rrN rO rQ rS rR rs rrzExternKernelAlloc.__init__s ..v66CC2BCCCCCCCC    &!#3 4 4  Lb    )  ') G++D11  ""4(((((rrcdSr8rrs rrxz!ExternKernelAlloc.should_allocate=r:rctrrrs rrw z"ExternKernelAlloc.apply_constraint@rrrI )rNNNrN)rqr<rrrK r2rrX rN r8rO r8rQ r2rS rR rrrr)rrrr rrxrw rrs@rr< r< s3333(*(,)-&*79+/)))))))@""""""""rr< cDeZdZdZdfd Zdd Zdd ZddZddZxZ S)rzP An output buffer that represents the mutation of a pre-existing buffer rqr< mutated_noder mutating_noderrrctd||}tj||g|_||_tj||_ dSr) rrrrwr"r<mutation_namesrd r"r)rrqrc rd mutated_node_namers rrzMutationOutput.__init__Isy d6222(1133 ##$566601(5G++D11 rc|jSr)rd rs rrzMutationOutput.get_defining_opSs !!rrc|jSr)rf rs rrz!MutationOutput.get_mutation_namesVrrrcdSr8rrs rrxzMutationOutput.should_allocateYr:rrcV|}dd|DDS)Ncg|]}||Srr)rrs rrz7MutationOutput.get_mutation_buffers..^s%    rc3TK|]#}tj|V$dSr)rwr"try_get_buffer)rrs rrz6MutationOutput.get_mutation_buffers..`s2PP..t44PPPPPPr)r)rrf s rget_mutation_buffersz#MutationOutput.get_mutation_buffers\sB0022  PPPPP    r)rqr<rc rrd rrrr~rrrr) rrrrrrrrxro rrs@rrrDs222222""""####        rrcteZdZUdZiZded<edd Zedd Zdfd Z ddZ ddZ xZ S) TMADescriptorad An IR node representing a generic host-side TMA descriptor in the Triton API Mostly useful for user-defined Triton kernels relying on host-side TMA; but can, in principle, be used for Inductor's Triton templates, too. See TMADescriptorExperimental and TMADescriptorStable for the two implementations (the old API and the new API) zdict[Any, TMADescriptor]_CACHErgrtma_metatuple[str, tuple[Any, ...]]rct|dksJ|ddkrt|g|dRS|ddksJt|g|dRS)Nrr experimentalrIr!)rTMADescriptorExperimentalTMADescriptorStable)rrgrt s r _create_implzTMADescriptor._create_implsss8}}!!!! A;. ( (,VBhqkBBB BA;(****&v< <<< .sBBBG7<BBBrrFATEN) max_autotunemax_autotune_gemmmax_autotune_gemm_backends)rrr|r rwr"r"rr# make_subgraphsubgraphr rrrIgraph_input_namesr sym_inputstorch._inductor.configrrJset_graph_handlerrrun operationsrr)rrqrr|r r rr sym_inpinductor_config base_patchesmerged_patchesrrs rrzSubgraphBuffer.__init__s v{333,G++D11  ""4(((--dg~}UU  ,,,,,(55 ! A AG7>DM &w| 4 M + 2 27< @ @ @ @BBzBBB888888888   / / A A!&%*.4L .X -WAUSU-WN &&~66 8 8! !4#677 8 8 8 8 8 8 8 8 8 8 8 8 8 8 8  A-2AAB)).*=*=*?*?@@@@! A A A A A A A A A A A A A A A A A As6;$GF4 GF GF ;GGGrx r|rrcGdd}t|jsJd|jD}|||jg|j||jgdS)NceZdZddZdS),SubgraphBuffer.codegen..CodegenGraphr"r}c,||_|j|_dSr)r"r)rr"s rrz5SubgraphBuffer.codegen..CodegenGraph.__init__s" !J rN)r"r})rrrrrrr CodegenGraphr s( ' ' ' ' ' 'rr c6g|]}|Srr/rr0s rrz*SubgraphBuffer.codegen..s$CCC!++--CCCr)r r'codegen_subgraph_with_flattened_outputsr r r)rrx r outer_inputss rr zSubgraphBuffer.codegens ' ' ' ' ' ' ' '   ,,,,,CCt{CCC 77 L ' ' -do - - YK     rr) rqr rrr|r}r r r rrrX rI rrrrr rrs@rr r sc15,A,A,A,A,A,A,A\         rr ceZdZdZd!dZd"dZed#d Zd$d Zd%dZ e d d&d'fd Z d(dZ d)fd Z d*dZd+d ZxZS),UserDefinedTritonKernelz> A user-defined triton kernel (e.g. via @triton.jit). r(tuple[Kernel, Any, list[str], list[str]]cBddlm}ddlm}||jg}g}g}t |rtdr'|fdj Dn,tdsJ|j tdr0j D]'}| j j|(n,tdsJ|jj}j |||fS) Nr) Autotuner)kernel_side_table restore_idxc3<K|]}jj|VdSr)r arg_names)rrrs rrzBUserDefinedTritonKernel.get_kernel_and_metadata..0s>**/0FI'*******r restore_value reset_idx reset_to_zero)triton.runtime.autotunerr *torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrrrr r r rrr r configs)rr r r restore_value_argsreset_to_zero_argsrrs @rget_kernel_and_metadataz/UserDefinedTritonKernel.get_kernel_and_metadata#s666666PPPPPP"--do>>(*(* fi ( ( v}-- @"))****4:4F***v77777"))&*>???v{++ @)FFA&--fi.A!.DEEEEFv77777"))&*>???nGYFw 24FFFrrcptjsdS|jjsdSt |jjdkrdSt |jdksJt|jdtsdSt|jdj tsdSt|jdj j tsdSt|jdj j j tsdStd|jdj j j jDsdSdS)aE For kernels like @triton.jit def add_kernel(in_ptr0, in_ptr1, out_ptr, n_elements, BLOCK_SIZE: tl.constexpr): pid = tl.program_id(0) offs = pid * BLOCK_SIZE + tl.arange(0, BLOCK_SIZE) mask = offs < n_elements x = tl.load(in_ptr0 + offs, mask=mask) y = tl.load(in_ptr1 + offs, mask=mask) tl.store(out_ptr + offs, x + y, mask=mask) @torch.compile def fn(a, b): out = torch.empty_like(a) grid = (triton.cdiv(a.numel(), 1024),) add_kernel[grid](a, b, out, a.numel(), BLOCK_SIZE=1024) return out.relu() We can potentially fuse the relu epilogue into the add_kernel. We do this by pruning the `out` tensor allocation and directly writing the relu-output. FrIrc3"K|] }|dkV dSrPrrrs rrz.vs&NNa16NNNNNNrT)rJ*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrprrr=rrrs rr z)UserDefinedTritonKernel.can_fuse_epilogueCs00@ 5 2 5 t!( ) )Q . .54$%%****$+A. :: 5$+A.3Z@@ 5$+A.38.II 5$+A.38=yII 5NN4#4Q#7#<#A#F#MNNNNN 5trrx r|rc0||dS)N)epilogue_fusion_codegenr s rr zUserDefinedTritonKernel.codegen{s}}Wd};;;rr tuple[ComputedBuffer, str]c.|||S)zQ epilogue_fusion: (fused epilogue node, modified kerel src code) r )rrx r s rcodegen_with_epilogue_fusionz4UserDefinedTritonKernel.codegen_with_epilogue_fusions }}Wo666r!tuple[ComputedBuffer, str] | Nonec  ddlm}\}}}}j}i} |rjt jjjdksJttjjjj } |\} } ij| | i}| | i} | |||||j |\} }}}fdj D}|| d|jDd|jD}t!fd|D}g}g}g}g}t#j|t)t#jd|D]\}}||vr |r||||t/|t0rO||||t/|t6t8t:t<fr9|||t?|||vr1|d |t69|f |r1|d |t6w| | tCd t?|d |"|| |#| ||||||d $j%j  d S)YOverrides the parent member. See https://github.com/pytorch/pytorch/issues/151692r)triton_version_uses_attrs_dictrIc<i|]}||Srr )rr rs rrz4UserDefinedTritonKernel._codegen..s6   ,-At$$Q''   rcg|] }|j Srrrrs rrz4UserDefinedTritonKernel._codegen..s333QV333rc*g|]}|j |jSr) is_constexprnumr s rrz4UserDefinedTritonKernel._codegen..s!EEEanEaeEEErc3(K|] }|V dSrr)rrr s rrz3UserDefinedTritonKernel._codegen..s'$F$FaYq\$F$F$F$F$F$FrrrRNzUnsupported arg type: r:T) arg_typesraw_argsraw_keys triton_meta inductor_metarDr!original_fxnode_name)&r r r rrr rrr7r8r!define_user_defined_triton_kernelgridrQ ryparamsr@rrrrrepeatrrrr/r*rr:rr rrrr generate_kernel_callr+r^ )rrx r r rr r r kernel_kwargsepilogue_out_overridemutable_arg_nameepilogue_computed_bufferr new_namer r extra_launch_args named_args constexprsconstexpr_namesrr raw_keys_filteredraw_args_filteredrr( r s` @rr z UserDefinedTritonKernel._codegens IHHHHH  ( ( * *       02  Qt(4;<<AAAA#D):)F)M$N$NOOT *9 ' $aWt{W,<>VWWM%57O$P !  5 5      I           151S    /00033V]333 EEV]EEE $$F$F$F$F:$F$F$FFF! ')')"     I$4R$8$8:K L L  $ W$ WID#&&+I+I+K+K&  $ $T * * *  $ $S ) ) )#v&& W C1133444  1111C#udD!9:: W C     c++++(( B  %%%%2133,KKOOO$$S))))%))+++%))++++)*U499*U*UPS*U*UVVV Wh///$$  &&#'??$$!%!2 % rFrrcrt|t|j|zSr)rrxr)r rs rrxz,UserDefinedTritonKernel.get_free_symbol_usess8 ww++M::=M I}> >   rctSrr?rs rrz0UserDefinedTritonKernel.get_unbacked_symbol_defsrrr rr rtma_descriptor_metadatar kernel_argsc8g}i}g}D]\}} t| trmt| } ||vr!t | ||} || | ||<|| | ||<t|dksJ|d _ t|tsJt|tdt!j |t#|||_|_\} } } } t+| dsJfd| jD_ddlm}m}t| dkr | djni}ddl}| j_|j_ |j_!_"|| i||_#fdj#j$j%D_&fdj&D_'tPj)*dS)Nrrr cg|]}|v| Srr)rr( r s rrz4UserDefinedTritonKernel.__init__..( s*. . . sk/A/AC/A/A/Ar)identify_accessed_tensorsidentify_triton_storescg|]<}t|jt/|j=Sr)rrrr)rrr s rrz4UserDefinedTritonKernel.__init__..C sN   +//#(33Y??  !   rcXg|]&}ttj|'Sr)rr/r!)rrrs rrz4UserDefinedTritonKernel.__init__..I sB! ! !  :T[9993 E E! ! ! r)+rrrrr rVrr rrrr+r!rrrrr/rr r r rr rQ r r r rastrF kernel_srcparse kernel_astr r r rrr r'rwr"r#)rr r r r rrrK r rr0rr r r r autotuned_kwargsr rs` ` rrz UserDefinedTritonKernel.__init__s "$&&( %%''  DAq!Y''  99$:L:LQ:O:OPP///%,,Q0G0JKKA a   q $$Q'''q 6{{aQi**,, &(++99T&\\99+   dk * * *  -     % $ < < > >Av{+++++. . . . !+. . . *        14G q0@0@71:,,b !*))DO4433DODD&55  /{ /. / #      (4;    ! ! ! ! (! ! !  ""4(((((rrc*t|jSr)rr'rs rrz#UserDefinedTritonKernel.get_outputsO sD)***rr0c|jSrrrs rr+z"UserDefinedTritonKernel.get_deviceR rr)rr rrI )rx r|r r rr)rx r|r r rrrrr) r rr rr rr rrrrr)rrrrr r rr r r rbrxrrrr+rrs@rr r sDGGGG@6666p<<<X<7777i i i i V566$)      76 N)N)N)N)N)N)`++++rr cDeZdZdZddZddZdd Zdd Zdfd ZxZ S)InplaceBernoulliFallbackE This needs to be a custom class to handle mutation properly rx r|rrc td|jDsJd|jD\}tjjrd||d|ddtt|j d|j dS||d|ddtt|j d|j dS)Nc3@K|]}t|tVdSrr r s rrz3InplaceBernoulliFallback.codegen..\ s,>>Q:a((>>>>>>rc3dK|]+}tt|V,dSr)rrr/r s rrz3InplaceBernoulliFallback.codegen..] s6IIVQ1133IIIIIIrrrz, NULL)r;) rrrwr"r r2r rrreprrK ending)rrx rs rr z InplaceBernoulliFallback.codegen[ s'>>$+>>>>>>>>IIT[III 7     ''))rrArr3tTEW;X;X1Y1Yrrbibprr        ''))llAll3tTEW;X;X1Y1Yll\c\jll     rrcdSr8rrs rrxz(InplaceBernoulliFallback.should_allocatej r:rrc.|dgSr$r rs rrz+InplaceBernoulliFallback.get_mutation_namesm ""##rrctSrr?rs rrz1InplaceBernoulliFallback.get_unbacked_symbol_defsp rrrS rrrrK rctdt|||g||t j|t j ||_ t j |dS)NrrS ) rrr/r+r!rwr"r<rr"rr#)rrS rrK rs rrz!InplaceBernoulliFallback.__init__s s   allnn - - -    $ $ #     ##AJJLL111G++D11  ""4(((((rrI rrr)rS rrrrK rrr rrrrr rxrrrrrs@rr r V s    $$$$ ) ) ) ) ) ) ) ) ) )rr c`eZdZdZddZddZdd Zdd Zdfd Ze dd dZ xZ S)!InplaceCopyFallbackr rx r|rrcd|\}}}||||dSr)r codegen_device_copy)rrx rGrF non_blockings rr zInplaceCopyFallback.codegen s8#'#4#4#6#6 c<##Cl;;;;;rrcdSr8rrs rrxz#InplaceCopyFallback.should_allocate r:rrc.|dgSr$r rs rrz&InplaceCopyFallback.get_mutation_names r rrctSrr?rs rrz,InplaceCopyFallback.get_unbacked_symbol_defs rrrqr<rrrK r2cJtd|||ddtj|dtj||_tj|dS)Nz aten.copy_aoti_torch_copy_)rN rO r) rrrwr"r<rr"rr#)rrqrrK rs rrzInplaceCopyFallback.__init__ s     +.     ##F1I$6$6$8$8999G++D11  ""4(((((rFrGrrFr cfd||fD}|f}tt|||}|S)Nc:g|]}|Srr )rr0rs rrz.InplaceCopyFallback.create.. s';;;1###A&&;;;rr)r r/r+)rrGrFr rrK rs` rrzInplaceCopyFallback.create s]<;;;c ;;;% $ cnn.. / / /      rrI rrr)rqr<rrrK r2rrr)rGrrFrr rrr ) rrrrr rxrrrr3rrrs@rr r s<<<<$$$$))))))$K|]}|VdSrr r s rrz7MutatingFirstArgExternKernel.codegen.. s. 9 9a!!## 9 9 9 9 9 9rrrr;) r rrr rK r2r rr )rrx argrefss rr z$MutatingFirstArgExternKernel.codegen s ,,,,, 9 9T[ 9 9 9 t) * *  ##%% M M '(:(: M MW^ M M     rrcdSr8rrs rrxz,MutatingFirstArgExternKernel.should_allocate r:rrc.|dgSr$r rs rrz/MutatingFirstArgExternKernel.get_mutation_names r rrctSrr?rs rrz5MutatingFirstArgExternKernel.get_unbacked_symbol_defs rrcdSrrrs rhas_side_effectsz-MutatingFirstArgExternKernel.has_side_effects rrNrI rrr) rrrrr rxrrr* rrrr" r" s    $$$$rr" c eZdZdfd ZxZS) ResizeStorageBytesvariablerrrrrct|ts Jdtdt |||g|ftj | tj ||_ tj |d|_d|_t|t t"t$fsJt'|tjj|j dS)NzTODO: dynamic shapesr)rK z"inductor_ops.resize_storage_bytes_z&torch::inductor::resize_storage_bytes_)rrrrr/r+r!rwr"r<rr"rr#rN rO rrrrnever_reuse_buffersrrp)rr- rrs rrzResizeStorageBytes.__init__ s1(C((@@*@@@(   h1133 4 4 4    + +#+    ##H$5$5$7$7888G++D11  ""4((("FG(Xz9$EFFVVXVVF #'' (>(>(@(@AAAAAr)r- rrrrrrZ rs@rr, r, sGBBBBBBBBBBrr, c(eZdZd fd Zd dZxZS) SetSourceTensorKernel self_tensorrstorage_tensorrrcl||}|t|||gdt jjjj t|tttfsJt|tjj|jtjj|tjj||}t-t/|||t-t/|||g|_dS)Nz!torch.ops.aten.set_.source_Tensor)rN rS r)rVrZrrr'r-rur set_ source_Tensorrrrrrrwr"r/ rrprr+rr/r')rr2 r3 r!rs rrzSetSourceTensorKernel.__init__ ss++N;;$$&&&   % % ' ' . )B +9    +*i'HII  4 L L   I #'' (8(A(A(C(CDDD #''(?(?(A(ABBB #'' 888**,, :V444k4 H H :V444nd K K! rrcV|d|dgSrr rs rrz2SetSourceTensorKernel.get_inputs_that_alias_output s%""DOOA$6$677r)r2 rr3 rrrr)rrrrrrrs@rr1 r1 sQ      *88888888rr1 cLeZdZdZddZddZdd Zdd Zd ddd fdZxZ S)!ScatterFallbackz This needs to be a custom class to handle mutation properly. This class handles both aten.scatter_ and aten.scatter_reduce_. It also handle the case `src` being a scalar properly. rx r|rrc0||dSr)generate_scatter_fallbackr s rr zScatterFallback.codegen!s))$/////rrcdSr8rrs rrxzScatterFallback.should_allocate!r:rrct|jd}t|tsJ|gSr$r )rrs rrz"ScatterFallback.get_mutation_names !s4k!n#v&&&&& rrctSrr?rs rrz(ScatterFallback.get_unbacked_symbol_defs!rrNTr include_selfrS rrrrcrrrFrr8r@ c bt|t_jrfd|||fD}|f} nfd||fD}||f} tdt ||| ||dt|ddg|tj | tj _tj dS)Nc:g|]}|Srr rr0rs rrz,ScatterFallback.__init__..!!s'FFFt))!,,FFFrc:g|]}|Srr rC s rrz,ScatterFallback.__init__..$!s'AAAt))!,,AAArrr? rr@ )rN rQ rS )rr src_is_tensorrrr/r+r!rrwr"r<rr"rr#) rrS rrcrrFrr@ tensorsrK rs ` rrzScatterFallback.__init__!s5(Y77   'FFFFq%oFFFG FMMAAAAq%jAAAG #JM   allnn - - -    ( ( | < <";//+3^*D#   ##AJJLL111G++D11  ""4(((((rrI rrK r)rS rrrrcrrrrFrrr8r@ rrrr rs@rr9 r9 s 0000    "!!)!)!)!)!)!)!)!)!)!)!)!)rr9 cDeZdZdZddZddZdd Zdd Zdfd ZxZ S)IndexPutFallbackzQ This needs to be a custom class to handle mutation and indices properly rx r|rrc0||dSr)generate_index_put_fallbackr s rr zIndexPutFallback.codegen;!s++D11111rrcdSr8rrs rrxz IndexPutFallback.should_allocate>!r:rrc.|dgSr$r rs rrz#IndexPutFallback.get_mutation_namesA!r rrctSrr?rs rrz)IndexPutFallback.get_unbacked_symbol_defsD!rrrS torch._ops.OpOverloadrrrr rr2 accumulaterc |_d|D}fd||g|D}d}tdt|||fd||t j dt j _ t j dS)Ncg|]}||Srrrs rrz-IndexPutFallback.__init__..P!s===qq}}}}rc:g|]}|Srr )rrrs rrz-IndexPutFallback.__init__..R!s'NNNQ4%%a((NNNraoti_torch_index_put_outrzaten.index_put_)rN rO rS r) rrrr/r+r!rwr"r<r r"rr#) rrS rrrrO valid_indicesrF rO rs ` rrzIndexPutFallback.__init__G!s ==G=== NNNN1f2M}2MNNN4   allnn - - -    ( ( M0+#     ##DOOA$6$6777G++D11  ""4(((((rrI rrr) rS rN rrrr rr2rO rrrr rs@rrH rH 6!s2222$$$$))))))))))rrH c.eZdZeddZdd Zd S) DeviceCopyrrr!r2r rrc |}|J|st|tjjvrt d|Drutj j sdtjj r>tj |tj || |Stj |tj |td|f}t|}d}|r|}t'|jo |jdko|}|jdkot'|jo|}|r(t+|rd|_t1t3|||||||g|S)Nc3<K|]}|tjjvVdSr)rwr"r rrs rrz$DeviceCopy.create..k!s,GGqA**GGGGGGrzDeviceCopy in input programr@Tr)r+rrrwr"rrrrJ aot_inductoruse_runtime_constant_foldingr add_device_inforrfrrr%rTrnrr&r'rurV rsr*rV) rrr!r x_devicerK ris_destination_pinnedis_source_pinneds rrzDeviceCopy.createc!s<<>>###  0Qqw'>>>GGA4D4D4F4FGGGGG?'D?w" 2''///''111''// / ''' )))7888%  + +A . . ::<< $\\^^F 8= ! ! KfkU&: K|  MU " Kvfk':': K|   , 5a 8 8 ,'+ALLNN $   /      q ! ! "    rrx r|rc|}t|dksJ|jr;||d|j|dn5||d||dt |jtr5|jjr+| | dSdSdS)NrrrI) r rrM r r/rrqr ru sync_d2h_copyr)rrx rs rr zDeviceCopy.codegen!s  ""4yyA~~~~   T  ' 'Q);;==tAw      ' 'Q1G1G1I1I4PQ7 S S S dk6 * * 3t{/D 3  ! !$--// 2 2 2 2 2 3 3 3 3rN)rrr!r2r rrrrI )rrrr3rr rrrrV rV b!sG- - - [- ^ 3 3 3 3 3 3rrV creZdZdZddZddZdfd ZddZed dddZ ddZ xZ S) DynamicSelectStorageOffseta The result of computing a dynamic selection index is determined as follows: when the index in the select operation is unbacked, the actual index calculation is ambiguous for negative indices (index + size) versus non-negative indices (just index). To resolve this, we allocate an unbacked SymInt to represent the storage offset and decompose the select operation into a call to as_strided, computing the storage offset at runtime with this node. rrmctSrr?rs rrz$DynamicSelectStorageOffset.get_reads!rrrcdSr8rrs rrxz*DynamicSelectStorageOffset.should_allocate!r:runbacked_offset_symbol sympy.Symbolr base_offsetsympy.Symbol | intbase_dim_striderrRrctdttjdg||_||_||_||_||_ ||_ dSNr@r) rrr/r-r!re rrg ri rrR)rre rrg ri rrRrs rrz#DynamicSelectStorageOffset.__init__!si ze1D1DEEErJJJ'=# &.  rrc,t|jgSr)r@re rs rrz3DynamicSelectStorageOffset.get_unbacked_symbol_defs!s467888rFrct|j|t|j|t|j|t|j|Sr)r)rrrrg ri rws rrxz/DynamicSelectStorageOffset.get_free_symbol_uses!sh TZ 7 7 U#DI}== > > U#D$4mDD E E U#D$8-HH I I  rrx r|c>|||jdS)Nr+ )codegen_dynamic_select_indexrRr s rr z"DynamicSelectStorageOffset.codegen!s#,,T,DDDDDrrr)re rf rrf rg rh ri rh rrh rRrrrrrrrI rrrrrrxrrrbrxr rrs@rrb rb !s&9999899$)    :9 EEEEEEEErrb creZdZdZddZddZdfd ZddZed dddZ ddZ xZ S)DynamicSliceSizeac Computes the output size of a slice call, handling the correct semantics in codegen. We do this for flexible handling for unbacked indices (to not data-dependent error). Slicing has 4 semantics for indices, i.e. x[start:] could be: 1) start < -x.size(0) -> x[0:] # negative out-of-bounds 2) start in [-x.size(0), 0) -> x[x.size(0) + start:] # negative slicing 3) start in [0, x.size(0)) -> x[start:] # standard slicing 4) start >= x.size(0) -> empty slice # positive out-of-bounds If the appropriate semantics are known beforehand, the output size is computed based on the start & end indices. If not (with unbacked indices), a new unbacked symbol is created to represent the output size, and codegen handles computing the correct case. rrmctSrr?rs rrzDynamicSliceSize.get_reads!rrrcdSr8rrs rrxz DynamicSliceSize.should_allocate!r:runbacked_size_symbolrf rCrh rDrZrctdttjdg||_||_||_||_||_ dSrk ) rrr/r-r!ru rCrDrZr)rru rCrDrZrrs rrzDynamicSliceSize.__init__!s^ ze1D1DEEErJJJ$8!   rrc,t|jgSr)r@ru rs rrz)DynamicSliceSize.get_unbacked_symbol_defs!s445666rFrct|j|t|j|t|j|t|j|Sr)r)rCrrDrrZrws rrxz%DynamicSliceSize.get_free_symbol_uses!sf TZ 7 7 U#DHm<< = = U#DI}== > > U#DI}== > >  rrx r|rc0||dSr)codegen_dynamic_slice_sizer s rr zDynamicSliceSize.codegen"rP rrr) ru rf rCrh rDrh rZrh rrh rrrrI rp rs@rrr rr !s   7777.//$)    0/ 11111111rrr cDeZdZdZddZddZdfd ZddZddZxZ S)rz; The result of a call to aten._local_scalar_dense. rrmctSrr?rs rrzDynamicScalar.get_reads"rrrcdSr8rrs rrxzDynamicScalar.should_allocate"r:rsymrf keypathpytree.KeyPathrprrc|tdtt jd||g||_||_dSrk ) r*rrr/r-r!r!r~ r )rr~ r rprs rrzDynamicScalar.__init__"sn   *EL$7$7888$:M:Mtf:U:U    rrc,t|jgSr)r@r~ rs rrz&DynamicScalar.get_unbacked_symbol_defs!"s48*%%%rrx r|c0||dSr)codegen_dynamic_scalarr s rr zDynamicScalar.codegen$"s&&t,,,,,rrr)r~ rf r r rprrrrrI ) rrrrrrxrrr rrs@rrr "s&&&&--------rrcreZdZdZddZddZdfd Zdd Zed dddZ ddZ xZ S)rz5 The result of a call to aten._assert_scalar rrmctSrr?rs rrzAssertScalar.get_reads-"rrrcdSr8rrs rrxzAssertScalar.should_allocate0"r:rscalarrxr rrctdttjdg||_||_dSrk )rrr/r-r!r r )rr r rs rrzAssertScalar.__init__3"sQ   el511 2 2 2      rcdSrrrs rr* zAssertScalar.has_side_effects>"rrFrrc,t|j|Sr)r)r rws rrxz!AssertScalar.get_free_symbol_usesA"s  ];;;rrx r|c tjsdStt|d}t jjrdSt jjrWd|d}t jj |j d}| d|d|j d|d dSt jj |j d}| d |d | d t|j d| |d dS)NFrzstd::to_string(r;)r<zif (!(z()) { throw std::runtime_error("Expected z but received " + z); }zif not (z):z raise RuntimeError(z = None)rJscalar_assertsr7r8rxrwr"r r r0codegen_cpp_sizevarr r2r codegen_python_sizevarr r)rrx r# symbol_strsizevars rr zAssertScalar.codegenG"st$  F d44454IIJJKK 7  ; D W  ;46444Jg*>> e?G   xxx48xxgqxxx     g*AA eBG   4444 5 5 5   ITXIII J J J   999 : : : : :rrr)r rxr rrrrrrI ) rrrrrrxrr* rbrxr rrs@rrr("s      N++$)<<<<,+< ;;;;;;;;rrc$eZdZUded<ded<dS)ExternKernelNoderrzexport_schema.NoderNrrrrr r g"s' IIIrr ceZdZdZ d1ddd2fdZd3fd Zd4dZd5dZd6dZe d7d!Z d8d#Z d9d%Z d9d&Z d'Zed4d(Ze d:d,Zed;d/Zed.has_functionalize_impl"s]8AA 58/=L))HH(6"-G  rz'NYI: Can't generate FallbackKernel for infotorch._C.Argumentr( rrctjtjr3t|tt fsJt|t jjrt|t tfrJ|dSjdSdfd d fd }t j jr||D]}||dSdSt jjsJ||dS) Nr0rrrc j|jJjjrKjt t||dSdSr) alias_namesrr alias_infois_writer'rr/r+)r0r rs r add_aliaszPFallbackKernel.__init__..handle_aliasing_and_mutation..add_alias"s '' 555222?+)00&z'H'H'H!TRRrrc|dSt|ttgtRrdSt|tsJt ||dSr)rr r r=rr)r0r s radd_alias_if_graph_bufferz`FallbackKernel.__init__..handle_aliasing_and_mutation..add_alias_if_graph_buffer"sh9Fa&$!:!:!:;;F!!V,,55d1gg55, ! r)r0rrr)r0rrr) rrr-ListTyperr library_utilsis_tensor_like_typer is_tensorlist_like_type)r r( r optional_tensor_argr rs` @rhandle_aliasing_and_mutationz=FallbackKernel.__init__..handle_aliasing_and_mutation"sU$)U^44 A!#e}55@@tCyy@@50;; :&cE4=99999{&             4TY?? /?/2GG+112EFFFF#?GG%8CCCCC))#.....r)rrN rr)r r r( rrr)!rrruse_runtime_dispatchrW rr-rq rr r rrS r rrN rwr" warn_fallbackr rf rrs _libraryrmutates_and_returns_first_argrr is_mutabler&rrrK zip_schema)rrqrr r r rrW schemar rr r r( rs` rrzFallbackKernel.__init__t"s   +   . ! !    %*!!2!8b UZ*EJ,JK   X X W W W4<< W W W X X ","Nbb &222 d5666'))+ d& (F G G  F !1!6!6!8!8 8 8 F!) >  = =d>N O O    & &{1~'>'>'@'@ A A A F       *4+;<< +*4+;<<  &L$:JLL **4;8JKK f( /( /( /( /( /( /T-88vNN 4 4ID# ( (s 3 3 3 3 4 4rrjc2t}|jtjjjurX|jD]P}t|tr9| tj | }Q|Sr)rrlrS r-_prims rng_primsgraphsafe_run_with_rng_staterK rr with_readrKrr)rrr( rs rrlzFallbackKernel.get_read_writes"sgg--//  u|5R R R)  c>22"-"7"7$,S\\^^<<##Krrx r|c ||||jt|ddSNrW )(codegen_unbacked_symbol_defs_for_outputsrrrr s rcodegen_unbacked_symbol_defsz+FallbackKernel.codegen_unbacked_symbol_defs#s8?? MMOOT\749Ld+S+S   rrct|ddx}rIttjjj|}|Jt |St Sr rr;rwr"r#rr@rrrW resolveds rrz'FallbackKernel.get_unbacked_symbol_defs#h '.A4 H H H  0 *,=H'''hmmoo.. .<< rrc tjGddt|jsJfd|jD}|||j\}}t jjrdt|j tj j r@|||}dt|j jj|D}n d|D}|j||S)Nc"eZdZUded<ddZdS))FallbackKernel.codegen_args..Shimrrefrrc|jSr)r rs rr2z2FallbackKernel.codegen_args..Shim.__repr__#s xrNr)rrrrr2rrrShimr #s3 HHH      rr cJg|]}| Srr )rrr s rrz/FallbackKernel.codegen_args..#s/HHHqttA//1122HHHrchg|]/\}}tjj||j0Sr)rwr"r0r rk )rparamrs rrz/FallbackKernel.codegen_args..#sCE1$33AuGGrcVg|]&}tjj|'Srr rs rrz/FallbackKernel.codegen_args..##s+IIIqAG(77::IIIr)r dataclassr rr rK rwr"r rrS r-rq rr r rrs rt rry)rr rrr s @rr zFallbackKernel.codegen_args#s+              ,,,,,HHHHDKHHH **;8JKK f 7  J:d.> @U#V#V J..tV<..#s' J J J1:a+I+I JQ J J Jrc^g|]*}||+Srr4r' s rrz.FallbackKernel.find_device..4#s1SSSC#..BRBRSs~~''SSSrrr@c3LK|]}td|V dSr)rg find_devicers rrz-FallbackKernel.find_device..=#sH$$**433$$$$$$rcg|]}||Srr)rr!s rrz.FallbackKernel.find_device..C#sAAA&&AvAAArrI)rr-rer!r ScriptObjectr*r,rrr@rrnr)r r non_torch_bind_tensor_argsdevices device_setr!s rr zFallbackKernel.find_device)#s   J J J J J J # &   ;SS;SSSG1:  nel 3 3 )!( (  UX24DE   ' ^ , , '<&& & ntUm 4 4 #$$($$$J BAJAAAG7||q  qz! t! " "!&%,77777&+&&"!MMM"1: trrc.ddlm}||jS)Nr) is_impure)torch._library.utilsr rS )rr s rr* zFallbackKernel.has_side_effectsO#s)222222y)***rrct|jtjjtjjfs(Jd|jdt |jdt|jtjjsBd|jvr'|jjj rt|jrgS|j S)Nr r:r r ) rrS r-rq rr r rrrs r r&r rs rrz+FallbackKernel.get_inputs_that_alias_outputV#s  uz4ej6TU     6$2B 6 6D$%% 6 6 6   4+UZ-KLL $"$*:*?*?*A*AAA (3B&t'788BI# #rcDt|jdksJ|jSr)rrf rs rrz!FallbackKernel.get_mutation_namesk#s'4&''1,,,,""rc  tdjt t sJt jj \} | } fdj D}j}tj jsg||Stdg}||| }dd  t |t"jjjr(||d |d j}n |jj}t1|d kr0jrjnj}|d j} ||g} n! fd t9|jD} jJi} t jt"jjrt?j jd| d<tCtEj#j$|| | } tj%&| g||S)a ProxyExecutor Design Note We export the ExternFallbackNodes (for custom ops) into a serialized file and run it with a host side proxy executor to address the ABI problem This is currently only implemented for fbcode. Eventually, we will also make this work for OSS. Detailed design doc can be found at https://docs.google.com/document/d/1wC4DOZFaYym2t1Esz0X5yxlLI3RDnSiyRbUus3bkJ64/edit?usp=sharing z4Extern kernel node added for node %s with target %s.c,g|]}j|fiSrr )rrrrs rrz.#s>    "D !# 0 0 0 0   rN return_type1torch.TensorType | torch.ListType | torch.JitTyperIRNode | Sequence[IRNode]rexport_schema.Argumentct|tjtjfr|}t|tt frt |dksJ|d}t|tjr\t|tsJtj tj | S|Jtj dSt|tj rt|tjrVt|tsJt!|tj d|DSt|tjrt|tjr|>tj tj d St|tsJtj tj tj |  St|tjr tj | St)d t!|) NrIrr) as_tensorT)as_nonec\g|])}tj|*S)r) export_schemaTensorArgumentr)rrs rrzZFallbackKernel.export_extern_kernel_node..handle_single_output..#s>   &4#,,..III   r) as_tensors)as_optional_tensor)as_intzUnsupported return type )rr- TensorTypeNoneTyperrrrr ryrr rr getElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError)r rrs rhandle_single_outputzFFallbackKernel.export_extern_kernel_node..handle_single_output#s+(8%.'IJJ/ SftUm44$v;;!++++ )Ck5+;<<G%c622222(188"/">CLLNN"S"S"S9;;;(1888FFFK88! SZ**,,e.>>>! S"&(33AAT&\\AA3$-44  #)   5 K);<< S**,,e.>BB S>(188+8+O+V+V$(,W,,9 &ff55555(188+8+O+V+V&3&B%+__%6%6''',W,,9K77 S$-44F4CCC"#Qd;>O>O#Q#QRRRrrrIc8g|]\}}|j|Sr)rk )r return_schemarr s rrz.#sE    *M6 %$!+   r1torch_library_custom_op)rrrr)rr)r r rr rr )'rrrrS rrgrr rrK r rQ rwr"aot_moder%serialize_inputsr-r torchbind CallTorchBindr returnsrs rrr'rk rrq rr custom_ops_maybe_get_opdefr r r>rextern_kernel_nodesr)rrordered_kwargsr serializernamed_argumentsr rr output_argumentsrrr rs` @@rexport_extern_kernel_nodez(FallbackKernel.export_extern_kernel_nodeo#s B MMOO     $//;;d;;/**4;8JKK f**488     9   !w ,+T+N+ +*444 $55fdFKK3 S3 S3 S3 Sj fe5?M N N -mmDGT!W55=GGn,G w<<1  '+lMdll8MG!!*.K 4 4[' J JK       .1$,-G-G     +++ t')> ? ? 6+D,<==I25H. /#',,..&(!     $$T***''''rcj}|J|jdkrt|tjjsJt |tjj r9ddl m }t||vr"td|d_nk|jdkr7t|tjjsJt |n)tjj r|t jjv_tjj rt|tjjrjsdfd jj\}t-j|fd jD}t3fdt5||jjD_|jrs}jJjJ| !jfdj|j"rj"nj#nn|$tj%tLr?'|(|)|*|dS)r Nr rr zG%s is missing a c-shim implementation, using proxy executor as fallbackT _quantizedr0 torch.JitTyperrct|tjr|St|tjSr)rr-r r NumberType)r0 is_numbers rr z)FallbackKernel.codegen..is_number$sGa!3449$9Q%5%5%7%7888!!U%5666rc34K|]}j|fiVdSrr )rr rrs rrz)FallbackKernel.codegen..#$sL*D)!66v66rc3hK|],\}}t|to|jV-dSr)rcomplexrk )rrrsr s rrz)FallbackKernel.codegen..($sV,,Aq1g&&A99Q[+A+A,,,,,,rcXgSr)r r rsrrz(FallbackKernel.codegen..6$s*F$++--F0C0C0E0EFr)r0r rr)+rS r rr-rq rr rrwr"r r r rrrr rJrY custom_ops_to_c_shimsr rrK rrrQ rarrs rt r r rN ,generate_fallback_kernel_with_runtime_lookuprrr'generate_fallback_kernelrqr r& r* r/ r ) rrx rr r args_iter exported_argsr rs ` @@rr zFallbackKernel.codegen#s/!!!!  v % %fej&;<< J Jd6ll J J<w" 5LLLLLLv;;&;;;KKa15D-   - -fej&;<< J Jd6ll J J< J W  f1GG  % G  65:#899 -   7 7 7 7 7 7 ..t{D+CDD,,,))D % W%%%  $ 6 ::<t\fr|_/n2t|t`rt]|_/n|g_/|S)z9Create an instance of FallbackKernel from an _OpOverloadsNr)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variantrIr!r rr)rqrrK rrN rS c34K|]}t|VdSr)rWr' s rrz(FallbackKernel.create..$s*!K!K,s"3"3!K!K!K!K!K!Krr@rW has_unaligned_inputrrrW z"Not sure where to find device inforrrlist[tuple[Any, int]]rctttfr@tfdt t DStt r"fdDSttj rlt }tj ssts)tjj|j|Stt(rSttjr jjSttjjt4fst7r t9t;SJdtddS)Nc3bK|])}|t|fgzV*dSrr)rrgenerate_outputrrs rrzAFallbackKernel.create..generate_output..$sZ$$$OF1Iw4<<:K9L/LMM$$$$$$rc Vi|]%\}}||t|fgz&Srr9 )rrr"r: rrs rrzBFallbackKernel.create..generate_output..%sN Sg$v,,9L8M.MNNrrzFallbackKernel output type z is not supported)rrrrrrrrr-rerUrhrJ assume_unaligned_fallback_outputrtrwr"rYrrrr rrr r r*r,rr/)rrrrr!r: r4 packeds`` rr: z.FallbackKernel.create..generate_output$s &4-00( #tF||$$$$$$"3v;;//$$$FD))# $*LLNNFEL11 !((00 ;<*<-V44< G-11#(;;; FC((  FEL11 {''.0@A  (( )f--- ~~Q$v,,QQQ&~~tr)rrrr6 rr)1r *_fused_moving_avg_obs_fq_helper_functionalrr rwr"rr r rr-rq rr r+ r r is_builtinretorch._library._out_variantr- r. r/ r0 rs rrsr!r r`rrM r_make_out_variant_kernel_namerar r r r ru higher_orderprintr ExternKernelMultiOut try_creater/r,rrr)rrrrfake_incorrect_kernelscontextr r r r rW r- r. r/ r0 out_oprq out_resultrr!r: r4 r= s` @@@@rrzFallbackKernel.create$s#'"Q!S / / /1$79JKKGG!mmG  < <#"6;D;;;F;;  !  < < < < < < < < < < < < < < < fej3 4 4 F ,C ,V Ed E E Ef E E E vuz4 5 5 N(33F;; >5<88              F~fn-- 0'//~226::!c*;*;F*C*C&D&D&I&I$)0(.0>/05^22445  '! ,,"&"7"7!'DV'L'L & "!K!K{!K!K!KKKn== ) vu6@N O O )/555\%((F vuz4 5 5 "N(33F;; "G' "  " .88"3$79  J%!!  !S&)))"3FF ? ?? ? ?6S!000"3F) ) ) ) ) ) ) ) ) ) V"/."55 ge} - - '$FNN  & & '"7^^FNN%YFNsA66A:=A:rrqr<rrr rr r2r rrrX rW r rrrrI rrK )r r r r2rrrr)rrrrs)rrN rrrrrr)rrrrrrrrg)rrrrrrlr rr rr r* rrr rr rhr3r+ rrrs@rrgrgm"s)-@4HL@4@4@4@4@4@4@4@4D              0###\#J++++$$$$*####~(~(~(@S3S3S3XS3j    \  000[0deee[eeeeerrgc<eZdZdZddZddZddddfdZxZS) ComplexViewz9View a complex number as two dtyped numbers or vice versarrcdSr8rrs rrxzComplexView.should_allocate5%r:rrc.|dgSr$r rs rrz(ComplexView.get_inputs_that_alias_output8%s""##rNr5 rqr<rrr rr r2r rrrX rW r rc Xt|||||||dS)Nr5 )rr rrqrr r r rrW rs rrzComplexView.__init__<%sE      /      rrrrJ )rrrrrxrrrrs@rrL rL 1%s}CC$$$$)-GK            rrL ceZdZdZddZdS) MemoryCheckKernelz Custom kernel for memory checking that generates direct function calls TODO - the custom op was erroring with str inputs. should be able to custom op directly. rx r|rrc||j\}}}t|}t|}|r"|dd|d|d|d}n d|d|d}||dS)z.Override codegen to write direct function callzV# note: dont currently distinguish between buffers returned and dealloc'd in last stepzcheck_memory_step(allocated=z, freed=z, is_final_step=r;N)r. rK r r2)rrx alive_list dead_list is_final_step alive_repr dead_reprcalls rr zMemoryCheckKernel.codegenY%s 22444/3/A, I}*%% OO  S   h   r*qqiqqanqqqDDR*RRiRRRD$rNrI )rrrrr rrrrR rR R%s2       rrR c"eZdZUded<ddZdS)r,r2r!rr0c|jSrrrs rr+zMultiOutputLayout.get_deviceo%rrNr)rrrrr+rrrr,r,k%s6rr,creZdZddZ ddfd Zed dddZddZddZddZ xZ S)rUrx r|rrc|||js,||||dSdSr)codegen_multi_output!skip_size_stride_alignment_checksr& r* r s rr zMultiOutput.codegent%sY$$T***5 4  % %g . . .  * *7 3 3 3 3 3 4 4rFrqr<r rrlist[tuple[Any, ...]]r_ rctd||gdtj||_tj|||_||_dSra) rrrwr"r"rr#rr_ )rrqr rr_ rs rrzMultiOutput.__init__z%sg vw333G++D11  ""4((( 1R...rrrc|jd}t|ts J|||Sr$)rrrrx)rrrs rrxz MultiOutput.get_free_symbol_uses%sA[^ *f--99z99-..}===rcrt|jdkot|jdtS)NrIr)rrrrrs rrxzMultiOutput.should_allocate%s24;1$ t{1~'8 9 9 rrc$d|jDS)Ncg|]P}t|tr9t|dk<|QSr )rrgrrrrrs rrz.%s_   #~.. C446677!;; LLNN<;;r)rrs rrz(MultiOutput.get_inputs_that_alias_output%s%  {    rrjclt}|jD]P}t|tr9|t j|Q|| d fd }| }tj p|dupt|j }t j||d|}t j||jt S) Nrr2rrrclt|dksJtj|dSr<r>r?s rr5z*MultiOutput.get_read_writes..dummy%r@rrr.rr6)r@rrrrrKrrr'rBr+rJrrnrrTr%rr) rrprr5r!should_normalizewrite_rwrDrs @@rrlzMultiOutput.get_read_writes%s>/9ll; @ @C#v&& @ ,.s||~~>>??? }}//##0022 ; ; ; ; ; ; ;""1 1 '~ '&+&&&   3 4==??B2B   &?"     rrI r) rqr<r rrr` r_ rrrrrrr) rrrr rrbrxrxrrlrrs@rrUrUs%s444438 S S S S S S SM**$)>>>>+*>                 rrUcLeZdZdZdfd ZeddZddZddZxZ S)rzMultiOutput for opaque objects.rqr<r rrr` opaque_valuerrrc`t|||d||_dS)NTr_ )rrr )rrqr rrl rs rrzOpaqueMultiOutput.__init__%s5 SWXXX$0!!!rrc td)NzOpaqueMultiOutput has no dtype)r rs rr zOpaqueMultiOutput.dtype%s=>>>rc|Srrrs rrz#OpaqueMultiOutput.wrap_for_lowering%r6rrjc|t}|jD]P}t|tr9|t j|Qtt j|g}t j||tS)Nr) r@rrrrrKrrr)rrprrs rrlz!OpaqueMultiOutput.get_read_writes%s.8ll; @ @C#v&& @ ,.s||~~>>???/9  !$--// 2 2 30 0 &"     r) rqr<r rrr` rl rrr)rr)rrr) rrrrrrr rrlrrs@rrr%s))111111???X?         rrc"eZdZdZd dZd dZd S) AllocatingMultiOutputzMultiOutput with Inductor-controlled allocation for .out() variant ops. Overrides should_allocate()=True so Inductor allocates the output buffer, and skips tuple-indexing codegen since .out() writes directly into these buffers. rrcdSrrrs rrxz%AllocatingMultiOutput.should_allocate%rrrx r|rcl|js,||||dSdSr)r_ r& r* r s rr zAllocatingMultiOutput.codegen%sE5 4  % %g . . .  * *7 3 3 3 3 3 4 4rNrrI )rrrrrxr rrrrs rs %sF 444444rrs rH rN c~|j}|jjdd}|j}d|d|d|S)z8Build fully-qualified kernel name for an out-variant op.z::rIz torch.ops.r )r rs rrr )rH nsr rs rrA rA %sM  Bn!''--a0G#H 1 1 1W 1 1x 1 11rcbeZdZUdZded<ded<d%fd Zd&dZedddd'd$ZxZ S)(rD zMulti-output .out() variant lowering. Subclass of FallbackKernel that emits .out() calls with pre-allocated output buffers. Uses AllocatingMultiOutput child nodes for each output. r out_arg_nameszlist[AllocatingMultiOutput]out_variant_output_nodesrrrH rN rrrctj|i|||_g|_t ||_||_dSr)rrry rz rA rN rS )rrH ry rrrs rrzExternKernelMultiOut.__init__&sO $)&)))*(*%"?"G"G!rrx r|cZ||||dSr)r generate_extern_kernel_multi_outr s rr zExternKernelMultiOut.codegen&s0 W%%%0066666rNFr3 rr r!r2r rr r2r rrX rW r r4 r&Sequence[AllocatingMultiOutput] | Nonec ddlm} m} m} | |jsdSt |t tfsdS| |} | dS| | }td|DsdSt|t|krdS|t|||||||| | }g}t|D]\}}t|j |jg|jg|}t#||t%||fg}t&js| st+|s)t,jj|j||||_||_t |t rt |St|S) zGCreate an ExternKernelMultiOut if the op has a matching .out() variant.r)r- r. r0 Nc3JK|]}t|tjVdSr)rr-rer s rrz2ExternKernelMultiOut.try_create..5&s.GG1:a..GGGGGGrr)rrW rH ry r1 )rqr r)r@ r- r. r0 rs rrrrrr,rrsr!r r`rrs rrJr< rtrwr"rYrrrrz r)rrr r!r r r rrW r4 r- r. r0 rH ry r= rr tensor_outrq multi_outs rrE zExternKernelMultiOut.try_create&s5          ~fn-- 4.5$-88 4'' >4))&11 GGGGGGG 4 ~  #m"4"4 4 44 V , , ,    /'    02&~66 & &MAz !( &(z'(-**,,- F .~..23I 7 >& >)44 > )--in=== NN9 % % % %*1'  ne , , ">> !G}}r) rrrH rN ry rrrrrrI )rrN r rr!r2r rr r2r rrrX rW r r4 rrr~ ) rrrrrrr r3rE rrs@rrD rD %s 9999 " " " " " "7777HL$)JJJJJ[JJJJJrrD ceZdZUdZded<dRdZdSdZdTd ZdUd ZdVdZ dWdZ dXdYdZ dddZdZ d[dZ d\dZdRdZd[dZ d]d^d$Zd_d%Zd`d'Z d]dad)Zdbd+Zdcd-Zddd.Zded0Zdfd2Zdgd4ZdRd5ZdRd6Zdhd9Zdid;ZdWd<Zdid=Z dfd>Z!e"d d]djdAZ#dkdCZ$dldEZ%dXdmdHZ&e'dndJZ(dodLZ)dndMZ*dgdNZ+e'dpdPZ,dWdQZ-e-Z.dS)qrCzC TensorBox / StorageBox allow in-place mutation of Tensors rrprrc4|jSrryrs rr9z!MutableBox.has_exceeded_max_readsm&rzrr0c4|jSrrmrs rr+zMutableBox.get_devicep&rkrr;c4|jSrr@rs rr>zMutableBox.make_loaders&y$$&&&rr@c4|jSr)rprBrs rrBzMutableBox.make_indexerv&y%%'''rr^c4|jSr)rprTrs rrTzMutableBox.get_stridey&rkrrc4|jSrrprs rrzMutableBox.get_name|&rqrNrKrc6|j|Sr)rprNrMs rrNzMutableBox.has_large_inner_fn&sy++I666rTrOrQrrPrc:|j||SrurvrSs rrTzMutableBox.mark_reuse&rwrc4|jSrrrs rrVzMutableBox.realize_hint&r rc4|jSr)rprrs rrzMutableBox.unwrap_view&r rc4|jSr)rprJrs rrJzMutableBox.is_input_buffer&y((***rc4|jSr)rprZrs rrZzMutableBox.freeze_layout&sy&&(((rFrrr[c8|j||Sr)rpr^r]s rr^z*MutableBox.freeze_layout_with_stride_order&sy88 NNNrc6|j|Sr)rprbras rrbz(MutableBox.freeze_layout_with_fill_order&sy66u===rrc6|j|Sr)rprerds rrez(MutableBox.freeze_layout_with_same_order&sy66v>>>rrfc8|j||Sr)rprirhs rriz+MutableBox.freeze_layout_with_exact_strides&sy99-WWWrrjc4|jSr)rprlrs rrlzMutableBox.get_read_writes&r rrmc4|jSrrrs rrzMutableBox.get_reads&rrc4|jSrrrs rrszMutableBox.num_reads&rrrc4|jSrrrs rruzMutableBox.get_storage_numel&rrr8c4|jSrrrs rrzzMutableBox.get_reduction_type&rrr{c4|jSrrrs rr}zMutableBox.get_reduction_size&rrc4|jSrrrs rrzMutableBox.is_extern&rrc4|jSr)rprrs rrzMutableBox.is_no_op&rqrr!r2c6|j|Srr r5s rrzMutableBox.constant_to_device&sy++F333rrc4|jSr)rprrs rrzMutableBox.get_mutation_names&rrc4|jSr)rprrs rrzMutableBox.get_operation_name&rrc4|jSr)rprrs rrz'MutableBox.get_inputs_that_alias_output&sy55777rc4|jSrr|rs rr*zMutableBox.realize&r}rrrc6|j|Srr]rws rrxzMutableBox.get_free_symbol_uses&sy--m<<rBrTrrNrTrVrrJrZr^rbrerirlrrsrurzr}rrrrrrr*rbrxrrr/rrqr'rr%r rr2rrrrCrCe&sLLL2222&&&&''''((((&&&&$$$$77777=ADDDDDD((((''''++++))));@OOOOO >>>>????HMXXXXX ++++%%%%%%%%----........%%%%$$$$4444........8888####L))$)====*)= ****++++33333+++X+&&&&++++$$$$X    "HHHrrCcveZdZeed dZeed dZed dZdS) rrprrcdSrrrps rrzTensorBox.create&sFIcrrcdSrrr s rrzTensorBox.create's+.3rcht|tr|Stt|Sr)rrrrr s rrzTensorBox.create's0 d1 2 2 KD))***rN)rprrr)rprrr)rpr)rrrrrrrrrrr&sr III\XI ...\X.+++\+++rcfeZdZdZddZddZddZdd Zdd Zdd Z ddddZ ddddZ ddZ dS)rz7 StorageBox allow in-place mutation of Tensors rrct|jttfr*|jt jjvSdSr8)rrprrvrrwr"rIrs rrJzStorageBox.is_input_buffer's? di+!? @ @ @9%%''17+?? ?urct|jto)|jtjjvSr)rrprSrrwr"r rs rrzStorageBox.is_module_buffer's6 ty> 3 3 : ""$$(99 rr8c t|jr|jSt |jt t ttfsJt|j|j }|j }|j }|Jtdt||j|jd|j|_t"j|j|j_t"j|j|j|j_||j_||j_|jjj|j_|jjS)NF)r!r rrur)rrrprrr=rrrrrrr+rr(r*r%rwr"r"rr#rrrr)rrrr!s rr*zStorageBox.realize's  " "49 - - (9%%'' '$)iD$%GHH  $ IK K   Hi//11 I++-- %%''!!!"!i))++Y''))       00;;  ""49--- L  + ' #y~8 y~rrct|jttfr8|jjdkr|dSdSdS)zL Called on buffers we expect to be forced to realize later. rIN)rrpr=rrnontrivial_read_countr*rs rrVzStorageBox.realize_hint:'sZ ty9i"8 9 9  **,,BQFF LLNNNNN   FFrrKrcddlmfd|D}|sdSt|}t |}t |}||ko||z dko||kS)Nr)is_nonfreeable_bufferscdg|],}| tj|-Sr)rwr"get_dep_size_hint)rrr s rrzCStorageBox.has_accumulated_enough_reads_by_size..G'sJ   ))#.. G % %c * *   rFr)r r rrdrrb)rrK size_of_reads total_sizemax_sizemin_sizer s @r$has_accumulated_enough_reads_by_sizez/StorageBox.has_accumulated_enough_reads_by_sizeD's@@@@@@    ~~''    5'' }%%}%% ) # %X%* %H$ rctj}|.t|jr tj}n$tj}nt |tsJt |jtoX| |kp@| p,tj duo| tj Sr) rJrealize_acc_reads_thresholdrKrprealize_cpu_acc_reads_threshold$_realize_acc_reads_threshold_defaultrrr=rsrN realize_acc_reads_size_thresholdr )rr s rr9z!StorageBox.has_exceeded_max_readsW's&,&H# & .di   .4.T++?,+93?? ? ??$)Y//  NN  : : &&(( 7tC==; rTrOrQrPcf|dkrt|jttfrt |jr:|jgd}t fd|DrdS|rdS|o!|tj kSdS)zj A heuristic to decide if we should realize a tensor that is used multiple times. rI)exprlog10log1plog2sigmoidtanhc3*K|] }|jvVdSr)used_ops)rropcounts rrz5StorageBox.should_realize_on_reuse..'s+@@qG,,@@@@@@rTF) rrpr=rrKrrarNrsrJrealize_reads_threshold)rrQrP heavy_opsr s @rshould_realize_on_reusez"StorageBox.should_realize_on_reusem's 199DI 9/EFF9di  )4466 @@@@i@@@@@ 4&&(( tT4>>#3#3f6T#T Turc`|||r|dSdSru)r r*rSs rrTzStorageBox.mark_reuse's8  ' '; ' G G  LLNNNNN  rc4|jSrrrs rrszStorageBox.num_reads'rrNrrr)rKrrr)rQrrPrrrrr) rrrrrJrr*rVr r9r rTrsrrrrr 's     <    &    ,JN2=A%%%%%%rrc2eZdZUded<ded<dZded<dS)Subgraphrrr} graph_moduleNzGraphLowering | Noner")rrrrr"rrrr r 's8 III&&&&"&E&&&&&&rr bufferscd|D}ttd|Dt|kS)Ncdg|]-}t|tr|n|.SrrrvrrrRs rrz(_has_aliased_buffers..'sH !+6? C COrc34K|]}t|VdSr)rfr s rrz'_has_aliased_buffers..'s(;;"V**;;;;;;r)rr@)r s r_has_aliased_buffersr 'sRG z;;7;;;;; < .'s@@@qQVE]@@@rc:g|]}|Srr r s rrz)InvokeSubgraph.create..'s'!I!I!I1#"3"3A"6"6!I!I!Irr|r r r)r r rqrrindrr:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutputc t|ttfr|S|}|Jt t ||||| j | j t|fgdS)Nr,Trn ) rrrr+rUrsr*r%rTr'rtrur)rr r!invoke_subgraphs r create_outputz,InvokeSubgraph.create..create_output(s:N"OPP  **,,)))"%$..00#__..%0022%00229"("3"3"5"5"? $C[M6:    rc.g|]\}}||Srr)rrrr s rrz)InvokeSubgraph.create..(s)MMMYQ fa((MMMr)rrr rrr )r r rwr"r, rrrr-rurB r rr rrrr r rr r rr r graph_outputsr+r r,r)rr r r r, fake_operandsr rt fx_operands new_operandsroperandrr!outsr r s` @@rrzInvokeSubgraph.create's; 766666w+  +0445GHH H  AF"ei&<&III#(+uy/E/UUUUU,Q/8MMF"ei&<&III#(+uy/E/UUUUU'+FGG4K@@K@@@M "J!I!I!I!I!I!I%' %h//  LC/ARS ##G,,,,##,,WmC6HII   > !W22(,&m3HN $X^44 3 3""M22 3 3 3 3 3 3 3 3 3 3 3 3 3 3 3..  Gg'<==  ++-- !!!($F333          .NMMM)G:L:LMMM"& sG""G&)G&rx r|c0||dSr)codegen_invoke_subgraphr s rr zInvokeSubgraph.codegen (r r)r r r rrqr,rrr)r r r rrr rI )rrrrr rr rrrr3rr rrs@rr r 's!%H$$$$(,H,,,,'+G++++ ) ) ) ) ) )8888fff[fP........rr ceZdZUdZdZded<dZded<dZded<dZded <dZ d ed <d'fd Z d(dZ e d)dZ ed*d!Zd+d$Zd,d&ZxZS)- Conditionala IR node representing torch.cond Attributes: predicate: A boolean scalar tensor determining which branch to execute. operands: Input tensors passed to both true and false subgraphs. true_subgraph: Subgraph executed when predicate is True. false_subgraph: Subgraph executed when predicate is False. outputs: MultiOutput nodes representing the conditional's outputs. Nr predicater r r true_subgraphfalse_subgraphSequence[MultiOutput] | Nonerrrr rqr,rW r rrcL||_||_||_||_t |g|\}}t d||||||_tj ||_ tj |dSN)rrqrrK ) r r r r _split_by_sym_typerrrW rwr"r"rr#) rr r r r rqrW sym_argsr rs rrzConditional.__init__7(s#  *, 2I3I3I J J+ "     (%6D "G++D11  ""4(((((rrcg}|jr||j|jr||j|Sr)r rr r subgraphss rrzConditional.get_subgraphsS(sQ   1   T/ 0 0 0   2   T0 1 1 1rrint | torch.SymInt int | ExprcHt|tr|S|jjSr)rrrr)rs r _maybe_exprzConditional._maybe_expr[(s# a   Hv{rrtrue_fnfalse_fnrclist[MultiOutput]c j|}fdDtjjjd}t |t sJt|g}|D]M}t |tr!| |j d8| |Ntjjj d}dd } ||fD]} | jtj | j || j | _tj| j5| jj|| | jj|| j_d d d n #1swxYwY|jJ|jJ|jj} |jj} d | fd| ffD])\} }t#| rt%d| d|*t'| t'| ks J| | ft)t+| | D]\}\}}||ks J|||f||ks J|||f|j|jks J|||ft5d|gzDt7tjjjtjjj dd } Jdt?|||tA|fdt)t+| tjjj dD}|_!ddl"m#}||j \}}}}}||j \}}}}}tI|tI|z}fdtK|D_&|S)zNCreate a Sequence of IRNodes from a conditional statement (see .lowering.cond)c:g|]}|Srr r s rrz&Conditional.create..m(s';;;QC%%a((;;;rrRr"r r fake_tensorsSequence[torch.Tensor]rrc*g}t||D]\}}t|tr||0|tt ||d|SNFrQ)rrrrrrrr)r r retrr=s r_require_exact_stridesz2Conditional.create.._require_exact_strides}(sC #M< @ @   f&;<<JJv&&&&JJ$::%f--t{{}}E;JrNr r r zVOutput aliasing is currently not supported in compiled torch.cond. The outputs of the z% subgraph of torch.cond are aliased: c3hK|]-}t|t|V.dSr)rrr+)ros rrz%Conditional.create..(sO  a!677 LLNN      rrW zcannot determine devicer)r r r r rqrW c g|]\}\}}tt||n|d|Dd|D|j|jt|fgS)NcBg|]}t|Srr r rrs rrz1Conditional.create...(s&UUU"+11"55UUUrcBg|]}t|Srr! r" s rrz1Conditional.create...(s38: //33rr,) rUrsr+r*rrr'rtrur)rrr merged_output conditionalr!s rrz&Conditional.create..(s   (+*FM' ((**6",,... **,,UU @R@R@T@TUUU>K>R>R>T>T",,..5$//11;         rr)-check_input_alias_and_mutation_return_outputscTg|]$}t|j|%Sr)rrq)rrr% r s rrz&Conditional.create..(s@( ( (  8C=/# L L( ( ( r)r rr r rr)'rVrwr"r, rrrrr>rrr r rr r r r rrrrr+r*r'rtr7r;r#rrr r,rtorch._higher_order_ops.utilsr& r@rr')rr r r r r r fx_op fake_outputsr r true_outputs false_outputsrrrt_of_orW r& r true_mutated_inputsfalse_mutated_inputsmutated_operand_indicesr% r!s` ` @@rrzConditional.createa(sR%%i00 ;;;;(;;; ! 4 9" = +x00CC${2C2CCC0$&   , ,E%&& ,$$UZ%67777$$U++++w+07     $!(+  H~%!"!6!6,#0"*-"7"" (88&HN& 664J3I 4l44HN0 }(((~)))}2  4 (,7*m9TU  MD'#L11 $_*.__U\__ <  C $6$6666}8U666&s<'G'GHH U UMAzS>>##s~~'7'7777!S#777==??cmmoo5553}555>>##*cnn.>.>.EEEE3PS}EEEE    +      6 G  & G % ) )*=t D D  !!#<!!!!!#$F333/         (/8L!'"6";E"BCC//)   2&       : 9':N O O *Aq%q : 9(:O P P +Aq&#--@"A"AJ E E #  ( ( ( ( ( 566( ( ( $ s90E55E9 <E9 rx r|c |||||jt |didSr )codegen_conditionalr rrrr s rr zConditional.codegen(sV##D)))88 MMOOT\749Lb+Q+Q     rrct|ddx}rIttjjj|}|Jt |St Sr r r s rrz$Conditional.get_unbacked_symbol_defs)r r)r rr rr r r r rqr,rW r rrr)rr rr ) r rr r r r r rcrr rI r)rrrrr rr r r rrrrr r3rr rrrs@rr r $(s!   $I####(,H,,,,%)M))))&*N****,0G0000))))))8\ XXX[Xt            rr rr -tuple[list[ShapeAsConstantBuffer], list[Any]]cg}g}|D]G}t|tr||j2||H||fSr)rrrr)r non_sym_argsr r( s rr r )siLH%% c0 1 1 % OOCH % % % %    $ $ $ $ \ !!rceZdZUdZdZded<dZded<dZded<dZded<dZ d ed <d%fd Z d&dZ e d'dZ e d(dZed)dZd*d"Zd+d$ZxZS), WhileLoopzSThe IR node for while_loop and while_loop_stack_output. It supports input mutation.Nr carried_inputsadditional_inputsr cond_subgraph body_subgraphr rrr rqr,rW r stack_outputrrrc\||_||_||_||_t g||\}} t d|| ||||_||_tj ||_ tj |dSr )r: r; r< r= r rrrW r> rwr"r"rr#) rr: r; r< r= rqrW r> r r rs rrzWhileLoop.__init__%)s-!2** 2 1n 10 1! ! + "     (%6D "(G++D11  ""4(((((rrcg}|jr||j|jr||j|Sr)r< rr= r s rrzWhileLoop.get_subgraphsD)sQ   1   T/ 0 0 0   1   T/ 0 0 0rcxt|s|Sd|D}t}g}t||D]{\}}t||vr.|t |D|t|||||S)Ncdg|]-}t|tr|n|.Srr r s rrz3WhileLoop._clone_aliased_inputs..W)sH   %/v$G$G SF   V   r)r r@rrfrrr r)r: unwrapped_buffers seen_buffersroriginal_inputunwrapped_buffers r_clone_aliased_inputszWhileLoop._clone_aliased_inputsO)s#N33 "! !   (    )3 +-03NDU0V0V . . ,N,"##|33 l55nEEFFFF  $4!5!5666 n---- rrrc"t|tr|St|ttfrt|St|trt|St dt|)NzNYI unsupported output type: )rrrrvrUrr r)rs r_maybe_wrap_as_tensor_boxz#WhileLoop._maybe_wrap_as_tensor_boxi)s c9 % % LJ j/: ; ; LS>> ! [ ) ) L##C(( (JtCyyJJKK Krcond_fnbody_fnr c 4&ddlm}d'd }tjjjd }tjjjd } || z} d | D} d |D} d| D} fd|D}t |}||| }fd|D}||| }||z&||fD]}|jt| tsJt| tj |j | |j |_tj|j5|jj| ||urLt!|jjt!| ksJ||jj| |j_dddn #1swxYwY|jr|jsJ|jj}|jj}t%|rt'd|t!|dks J||d}t|t(s[|t,jks J|t!|dks J|t!&dks Jd&d}|Jt!|t!|ks J||ft5t7||D]\}\}}d(d}||||||||ksJ||||f||ks J|||f|Jt;tjjjtjjj !dd}t ||||tE|||}|j)t|jj#t,j$j%sJ||jj#| d}tM|}&fd|D}tO|}g}g|_(g|_)|rt!|dks Jd t5tjjj d!D]\} }!tUtW|!j,|!j-d"|!.Dd#|!/D$|t`| fg}"|j(1|"|1|"nDt5|D]3\} }!| |vrp| t!|ks Jd%te|}#|j)1tg|#j4|#||1|#ztUtW|!|!|!|!|!5j6&|t`| fg}"|j(1|"|1|"5t7||D]`\}$}%|$7tjj8vr6tjj9:|%7a|S))zcreate the while_loop IR node. stack_output controls whether it stack each iterations' output, which is necessary for training. r)check_input_alias_and_mutation tensor_boxesrr 'list[int | torch.SymInt | torch.Tensor]rrct|t|ksJg}t||D]\}}t|tjr]t |}|t || d||||Sr ) rrrr-rer9 rI rrrr)rN r r rfknew_tbs rr z0WhileLoop.create.._require_exact_strides)s|$$L(9(99999ClL99 # #Bb%,//#'@@DDFJJ$::"BIIKKu; JJrNNNNJrrRc(g|]}|jdSr r rs rrz$WhileLoop.create..)s@@@Q16%=@@@rc(g|]}|jdSr r rs rrz$WhileLoop.create..)sHHHqve}HHHrc(g|]}|jdSr r rs rrz$WhileLoop.create..)s!N!N!NA!&-!N!N!Nrc:g|]}|Srr r s rrz$WhileLoop.create..)s'HHHA3,,Q//HHHrc:g|]}|Srr r s rrz$WhileLoop.create..)s'NNNqc//22NNNrNr zOutput aliasing is currently not supported in compiled torch.while_loop. The outputs of the body_fn subgraph of torch.while_loop are aliased: rIz9torch.while_loop is assumed to have at least one operand. lhs_exprsSequence[int | Expr] rhs_exprsrct|t|ksJt||D]*\}}tjj||+dSr)rrrwr"r#r)rY r[ lhsrhss r_guard_list_equalsz,WhileLoop.create.._guard_list_equals)sf9~~Y7777 #Iy 9 9< rc g|] }| Srr)rr all_inputss rrz$WhileLoop.create..*sEEEc*S/EEErz-NYI: while_loop_stack_output input mutations.r"cBg|]}t|Srr! r" s rrz$WhileLoop.create..#*s&RRRbk55b99RRRrcBg|]}t|Srr! )rrs rrz$WhileLoop.create..$*s&VVV 7 7 ; ;VVVrr1 zonly carries can be mutated.)r!r rrrt)rN rr rO rr)rY rZ r[ rZ rr);r( rM rwr"r, rr9 rG rrrr r rr r rr r rrr*r-rr%r+rrrTr;r#rrrr,modulefx GraphModuler@r8rr'rUrsr!r rrrrr7rrqr'rtrrIr/ r)'rrJ rK r: r; r> rM r fx_carried_inputsfx_additional_inputs fx_all_inputsfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_r cond_outputs body_outputsrr!rrbor_ rW while_loop mutated_idxsmutated_idx_setrmutated_inputs_iter all_outputsrrr mutated_inputrrra s'` @rrzWhileLoop.createt)s~ QPPPPP    :G05b9 w38<),@@ @@-@@@HH6GHHH!N!N9M!N!N!NHHHHHHH#99/JJ00BUVVNNNN#?@@C/EE    8N7M$N8/884"}....}2 }2  - -  gXdgg  <  A%%%|%%% O!233 -;;==EJ......qzz||$$)))1))):""" G#""A))++!!!?##s<'8'8888  ; 888%S,%G%GHH A AKAxB < < < <  r{{}}bkkmm < < <  r}}  @ @ @==??bmmoo5552r67J555<<>>R\\^^333aR[3333!!!5 G  & G % ) )*=t D D   *0!!$F333/%    }(Z M %("6. . (( 65 M /   %\22EEEE_EEE#>22$&  &( # ) 2''1,,,?-,, ))=)B5)IJJ . . V'%}$lRRFKKMMRRRVVfmmooVVV  C[M   ")))444""9---- . )66 2 2 V/))^!4!44446T444$()<$=$=M/66&}';]JWW &&}5555 +##)#4#4#6#6"("2"2"4"4!'!2!2#)#4#4#6#6#)#4#4#6#6#= #  ! !I&--i888&&y1111NK88 @ @HC||~~!555 +// ???sA F00F4 7F4 rx r|c |||j|||jt |didSr )codegen_while_loopr> r rrrr s rr zWhileLoop.codegenN*s[""4):;;;88 MMOOT\749Lb+Q+Q     rrct|ddx}rIttjjj|}|Jt |St Sr r r s rrz"WhileLoop.get_unbacked_symbol_defsT*r r)r: rr; rr< r r= r rqr,rW r r> rrrr)r: rrr)rrrr) rJ r rK r r: rr; rr> rrr rI r)rrrrr: rr; r< r= rrrrrG rI r3rr rrrs@rr9 r9 )s@]].2N2222155555%)M))))%)M)))),0G0000))))))>\2LLL\LWWW[Wr            rr9 c>eZdZ ddddfdZdfd ZddZxZS)rNr rqr<rrr rr r2r rrrX rW r rrc  t|||||d|ddlm}||} | J| |_t jj| d|_ |t jj| <dS)Nr5 r) _get_effect) rrtorch._higher_order_ops.effectsr} effect_typerwr" effectful_opsrprev_effect_buffer) rrqrr r r rrW r} r rs rrzEffectfulKernel.__init__`*s      /     @?????!k&)) &&&&"#'"7";";K"N"N-1k***rrjct}|jC|jt j|j|Sr)rrlr rprrKrr)rrrs rrlzEffectfulKernel.get_read_writes}*sagg--//  " .   ! !$T%<%E%E%G%GHH   rrcdSrrrs rr* z EffectfulKernel.has_side_effects*rrrrJ rr)rrrrrlr* rrs@rrr_*s)-2HL22222222:rrcDeZdZed d d dZd dZd S) r FrrrrctSrr?rws rrxz!NonTensorObj.get_free_symbol_uses*rtrr8cdSrrrs rr*zNonTensorObj.realize*rrNrrr)rrrrbrxr*rrrr r *sZN++$),+ rr cNeZdZUded<ded<ddZddd Zdd Zdd ZddZdS)r rr%FakeScriptObject | torch.ScriptObjectr1rc|jSrrrs rrzTorchBindObject.get_name*rrNr+r,c|jSrrr.s rr/z!TorchBindObject.codegen_reference*rrc|jSrr5rs rr zTorchBindObject.get_value*rrtorch.ScriptObjectcft|jtjr|jS|jjSr)rr1r-r real_objrs r get_real_objzTorchBindObject.get_real_obj*s, dj%"4 5 5 ': :& &rrc>|}t|rdSt|dsJt|}t j|d}d|D}tjtj |dS)Nr__obj_flatten__cg|]E}t|tj||zFSr)rr-rer rrs rrz1TorchBindObject.get_buf_bytes..*sQ   !U\** NN  qwwyy (   r) r r,rrr rf r rroperatorr)rreal_script_obj flat_dict flat_elems flat_sizess r get_buf_byteszTorchBindObject.get_buf_bytes*s++-- ? + + 1(9:::::88::;; (33A6       j!<<.+s1. . .  . F. . . r) rrS r-rq rr rO rs rrt rQ )rrO rs rr[ z%_CollectiveKernel.set_cpp_kernel_name*sD$%%)>>>> :?>>!  D #)>#6D . . "N4. . . ***rrrrIRNode | list[IRNode]rrrc tjj5|j||g|Ri|\}}}}} dddn #1swxYwY| rJ|d| d |D]s} t | t r| tj|  | t Jd|d|t |||| tj |} j  fd| D jd| Dd|vrq j t#t |d  j|ddSdS)NrzIn-place collective zE requires at least one tensor argument; got only non-tensor IR nodes.rcNg|]!}tt|"Srr)rrr!r= s rrz4_CollectiveKernel.create_inplace..3+s0 T T T^Jf555sF C C T T Trc6g|]}|Srrnrf s rrz4_CollectiveKernel.create_inplace..7+s "B"B"Bc3<<>>"B"B"Brr)rwr"rr rr r*r<rr+r/rf tree_leavesr'rr rr)rrrrr_example_outputr r r rW tensor_arginpsr!r= s @@rcreate_inplacez _CollectiveKernel.create_inplace +sW  D D#"66CDCCCFCC  !  D D D D D D D D D D D D D D D%EE&E&E2C&E&EEE$% 1 1J*l33     G ' ' (;(;(=(= > > >~#..00!! 76 7 7 7"!! f % % %       !&))&& T T T T Tt T T T   !!"B"BT"B"B"BCCC F??  # * *z888&-PP      % %fUm&<&<&>&> ? ? ? ? ? ? 9==TensorBox | list[TensorBox]%list[MultiOutput] | _CollectiveKernelctjj5j||g|Ri|\}}}}} dddn #1swxYwY| rJ|d| |D]+} t | t s| ,t |tr||} | Jt| ||||fdt|D_ tj |D]I\} } tjst| s)tjj| jJj S|||||tjst|s)tjjjg_ S)Nrrcng|]1\}}t|t|fg2Sr)rUrhr)rrrgrr= s rrz9_CollectiveKernel.create_out_of_place..t+sV Av ((00AYKr)rwr"rr rr r*rr r,rrrrJr< rtrYrrrh)rrrrrr r r r rW r r!rrgr= s` @rcreate_out_of_placez%_CollectiveKernel.create_out_of_placeU+sW  D D#"66CDCCCFCC  !  D D D D D D D D D D D D D D D%FF&F&F3D&F&FFF$% % %Jj/:: %""$$$ nd + +% __[.AAF%%%S!000 F "+>!:!: FN #6>>BB < < V: !S$$^44 F6 ;>O?? ;)--fk:::$XFNMr rrrJ ) rrrr rrrrrr) rrrr rrrrrr ) rrrrxr* r[ r3r r rrrr r *s      (0@0@0@[0@P999[999rr c2eZdZ ddddfdZddZxZS)_AllReduce_KernelNr rqr<rrr rr r2r rrrX rW r rrc t|||||d||ddS)Nr5 +aoti_torch_cpu__c10d_functional_all_reduce_rrr[ rP s rrz_AllReduce_Kernel.__init__+X      /       !NOOOOOrrx r|c|d||t|jtr||dSdSNz+torch/csrc/inductor/aoti_torch/c/shim_cpu.hinclude_extra_headerr] rrqr r& r s rr z_AllReduce_Kernel.codegen+c$$%RSSS,,T222 dk6 * * /  % %g . . . . . / /rrrJ rI r rs@rr r +su)-PHLPPPPPPPP,////////rr c2eZdZ ddddfdZddZxZS)_AllReduceKernelNr rqr<rrr rr r2r rrrX rW r rrc t|||||d||ddS)Nr5 *aoti_torch_cpu__c10d_functional_all_reducer rP s rrz_AllReduceKernel.__init__+sX      /       !MNNNNNrrx r|c|d||t|jtr||dSdSr r r s rr z_AllReduceKernel.codegen+r rrrJ rI r rs@rr r +su)-OHLOOOOOOOO,////////rr c^eZdZ ddddfdZddZd dZed!dZd"fd ZxZ S)# _WaitKernelNr rqr<rrr rr r2r rrrX rW r rrc t|||||d||ddS)Nr5 +aoti_torch_cpu__c10d_functional_wait_tensorr rP s rrz_WaitKernel.__init__+r rrx r|c|d||t|jtr||dSdSr r r s rr z_WaitKernel.codegen+r rc|jd}t|tsJt|tr<|jd}t|tsJt ||gSt|t rB|jd}t|tr|jd\}}|j|gSgSgSr$)rrrr rrUr)rrrcollr rs rget_volatile_readsz_WaitKernel.get_volatile_reads+sk!n#v&&&&& c, - -  1 Aa(( 1 1$q'' 1 1(3J [ ) ) :a=D$ 122 *Q3 C())IIrrrctjj5|||\}}}}}dddn #1swxYwY|rJ|d||t |||||}|jtt |||dS)Nrr) rwr"rr r/r+r'rr) rrrr r r r rW r= s r create_waitz_WaitKernel.create_wait,s0 W  0 0""63//  !  0 0 0 0 0 0 0 0 0 0 0 0 0 0 0%EE&E&E2C&E&EEE$ cnn.. / / /        && :S^^-=-=>>>V L L     s ;??rjct}|}|D]@}|jt j|A|Sr)rrlr rprrKrr)rrvolatile_readsvrrs rrlz_WaitKernel.get_read_writes,smgg--// 0022  G GB   ! !,"6r{{}}"E"E F F F FrrrJ rI rp )rrrrrrr) rrrrr r r3r rlrrs@rr r +s)-PHLPPPPPPPP,////2   [ *rr rr clt|ttfrt|St|tt fr7t tj}|D]}|t|z}|St|tj rt|St Sr) rr=r r6rrr@rr"rB r-rerrr0s rrB rB ),s!h%&& $Q''' At} % %  u| $ & & 0 0A ,Q// /AA Au| $ $$Q'''||rclt|ttfrt|St|tt fr7t tj}|D]}|t|z}|St|tj rt|St Sr) rr=r r5rrr@rr"rC r-rer s rrC rC 9,s!h%&& A At} % %  u| $ & & ' 'A #A&& &AA Au| $ $A||rrrr  torch.fx.Nodec`t|trt|jtrtt|jjtr"|jjd|dSt|jjt r|jjd|t|jjtrKt|jjjtr'|jjjd|dSt|jjtrn|jjj s_t|jjj dt r7|jjj dd|dSdSdSdSdSdSdS)Nrr) rrrprrrrrrUrr)rr s rassign_origin_noder I,s&)$$TFK)L)LT fk& . . T K  / / q A A A A A  (& 1 1 T K  / / q A A A&+*N;; T   %uAA T  %88JJJJJ6;+[99 T (0 Tfk.5a8&AATK$+A.AA-QRSSSSSTTTT T T T T T TTTr)rrrr)rrrr)rrrr)rrrr)rrrrrrr)rrrrrr)rrrrr)rrrrrr)rrrrrr)rrrrrr)r1r2rr3)rr7rr8)rr>rr)rrLrMrrr)r]r^r_r^r`r^rr)rgrrhrirr)r|r}rr)rrrr)rr5r rrrvr)rhrr rrirrrj)rr^rrrtr rrr )TFNFN)rrr?rrKrrLrMr[rrfrMrrN)rrrLrrrr)rr^r`r^rr)r rrr)rr rr )rr<rr)rrrrF )rH rN rr)r rrr)rr rr5 )rrrr )rrr r rr(i __future__rrrrrrloggingr textwraprcollections.abcrrrrrr r enumr r numbersr typingrrrrrrrrrrrtyping_extensionsrrrrrr unittest.mockrrr r!r"torch._export.serde.schema_exportserder r torch._library.custom_opsr r r rr torch._loggingr-torch.fxtorch.utils._pytree_pytreerf torch._dispatch.pythonr#torch._dynamo.utilsr$torch._export.serde.serializer%*torch._higher_order_ops.auto_functionalizer&torch._inductorr'r r)"torch._library.fake_class_registryr*torch._library.opaque_objectr+r,torch._prims_commonr-r.r/r0r1r2%torch.fx.experimental.symbolic_shapesr3r4r5r6r7r8r9r:r;r<r= torch.fx.noder>torch.utils._ordered_setr@torch.utils._python_dispatchrAtorch.utils._sympy.functionsrBrCrDrErFrGtorch.utils._sympy.symbolrHrrJrKcodegen.commonrLrMrNrOrPrQrRrSrTrUrV loop_bodyrW ops_handlerrXrYrZr[runtime.benchmarkingr\ runtime.hintsr]r^r_r`rarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrt virtualizedrurvrwrxry torch.typesrzcodegen.cutlass.templater{codegen.wrapperr|r"r}r~rrrD __version__rr ImportErrorrrrrrrr:rrq rr r r getLoggerrrrr rrrrr rrrrrrrrrrrrr6r=rIrKr\rfr{rrrrrrr;r=rVrgrrrrrrrrrrrr&rGrJrrrUrWrrrrrrhrvr6rBr_rfrmrwr{r<r rsr(rr#r'r/r9rrzrrrSrrrrrrsrrPrimitiveInfoTyperrrrrrrr rr r" rrE rI rrM rU r< rrr rx ry r r r r r" r, r1 r9 rH rV rb rr rrr rgrL rR r,rUrrs rA rD rCrrr r r r r r9 rr r r r r r r r r rB rC r rrrr s$#""""""" MMMMMMMMMMMMMM::::::::                          UTTTTTTTTTTTTTTT ''''''''''222222222222.........,,,,,,,,,$$$$$$$$$;;;;;;((((((??????MMMMMM######222222??????MMMMMMMM                          //////??????+*****"""""""" NNNNNNNNNNNN------::::::::0*))))))))) (BBBBBB&&&&&&''''''999999555555$$$$$$%%%%%%%"(O''''MMM'NJJNJJJ Yt__ WT]] WT]] WT]]Dj EkD(((((*/%*2PP PPPPg!!  8?4 8 8 8 y~'T k     c4[();(JKdR      ))))d###$#$ $ $ $ $F!LL$__LP         LP      UUUU U 27  :?     D > > > >;;;;(''''7777*$$G$G$G$GN'''';;;Z,Z,Z,Z,Z,Z,Z,Z,z UUUUUUUUUpP P P P P FP P P f&&&&% % % % % % % % P     i   F ;| $ $ ;y ! ! ;y ! ! K   ;u   ;u  {=))<<JNJNJNJNJNJNZZ  Z  Z  Z  Z  Z  Z  Z  ~)-71::&"8D>8D>"BH"LM MMMM7S7S7S7S7S97S7S7Stkkkkk'kkk\s s s s s 1s s s l` ` ` ` ` +` ` ` FF F F F F 5F F F T        [[[[[5[[[|(    !3748 99999x$^ ^ ^ ^ ^ v^ ^ ^ B[[[[[[[[|44444(444nA9A9A9A9A9(A9A9A9H!!!!!(!!!H^^^^^;^^^B SSSSShSSSl&&&&&&&&RlAlAlAlAlAlAlAlA^6"KKKKK|KKK$ S S S S S| S S S(<<<< 7 7 7 7 7 7 7 7 p p p p p Zp p p f CCCCC&CCC[H[H[H[H[HV[H[H[H|!G!G!G!G!Gf!G!G!GHT%%%%%{%%%D.V*V*V*V*V*V*V*V*rUEEEEEV]EEEPU & & & & &fi & & & & K     [   &6(     F    Ut4t4t4t4t4_t4t4t4n d###$#f+f+f+f+f+_f+f+f+R<<<<<><<<~%K$&,tC#I4E4L/MMLLLLLLLL^"""""|""" \B\B\B\B\B.\B\B\B~55555N5550((((((((6N>Z%Z%Z%Z%Z%NZ%Z%Z%z5555 URRRRR?RRRj OOOOO9OOOd ;;;;Uttttt<tttnU"""""l"""J     /   (("("("("(" ("("("V     V   B=====L===@% % % % % % % % P     -   $< < < < < \< < < ~uuuuuluuup ))))))))))|))))))Z/////,///d<8BBBBB5BBB$88888-88846)6)6)6)6)l6)6)6)r))))))))))|))))))X;3;3;3;3;3;3;3;3|1E1E1E1E1E1E1E1Eh5151515151|515151p-----L---8<;<;<;<;<;<<;<;<;~U A A A A A &A A A HU     .   @        2 I I I I I ,I I I X         F 4 4 4 4 4K 4 4 4 2222fffff>fffV TTTTTTTTn + + + + + + + +%%%%%%%%DU'''''v''' LLLLUB.B.B.B.B.\B.B.B.JUe e e e e ,e e e P " " " "U@ @ @ @ @ @ @ @ F )))))n)))X6"="="="="=l"="="=Jl(\ "````````F/////)///>/////(///>SSSSS#SSSr         TTTTTTs: H HH