^jO NUddlmZddlZddlZddlZddlZddlZddlZddlZddl 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-< ddlZejpZd.Ze(d0Ze"d1Ze"d2Ze"d3Zee.zZd,ed4<eeze.zZd,ed5<e=jje=jjzZd,ed6<ejeɫZej*e jd78Ze=jNjZ e#ed9eed9fd:d;ee#eeed9fd9d:d;fdzfZd,ed<<dBd=Zejd.>Gd?d@ZҐdCdAZӐdDdBZԐdEdCZՐdEdDZ dFdEZgdFZgdGZ dG dHdHZڐdIdIZ dG dHdJZedJdKdKZe dJ dLdLZ dJ dMdMZ dNdNZ dOdOZߐdPdPZdPdQZdQdRZ dRdSZ dSdTZdTdUZdUdVZdWZGdXd;Zed/>GdYdZZeGd[d\eZdVd]ZeGd^d_eZeGd`daeZedbedceddedeedfedfedgdhZdiedj< dW dXdkZeGdldmeZde/df dYdnZe ee.ee.gefZd,edo<GdpdqeZGdrdseZGdtdueZGdvdweZeGdxdyeZeGdzd{eZeGd|d}eZdZd~ZdZdZdZdZ d[ d\dZ d]dZd^dZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZGddeZeGddeZ eGdde Z eGdde Z d_dZ d`dZ GddZeGddeZGddeZGddeZGddeZGddeZGddeZeGddeZGddeZed/>GddeezZed/>GddeeZGddeZGddeZGddeZeGddeZeGddeZed/>GddeZejd.>GddZGd„dÐeZ GdĄdŐe Z!eeze"zeze#eezeze"zzZ$GdƄdǫZ%GdȄdɐe%Z&Gdʄdːe!Z'Gd̄d͐e Z(Gd΄dϐe Z)GdЄdѐe Z*Gd҄dӐe Z+ dadԄZ,ed/>GdՄd֐eZ-Gdׄdؐe-Z.Gdلdڐe.Z/ej`dbdۄZ1 dcd܄Z2ed/>Gd݄dސe-Z3ed/>Gd߄de3Z4Gdde4Z5Gdde3Z6GddeZ7Gdde3Z8Gdde8Z9Gdde8Z:Gdde3Z;Gdde3Z<Gdde3Z=Gdde3Z>Gdde3Z?Gdde?Z@Gdde6ZAGdde3ZBGdde3ZCGdde4ZDGdde3ZEGdde3ZFGdde3ZGGdde3ZHed/>Gd d ZIGd d e6ZJed/>Gd deJZKGddeJZLeGddeZMGdde3ZNGddeNZOGddeNZPdddZQGddeJZRejGddeZSGdd9eSZTGdd eSZUed/>Gd!d"eZVded#ZWed/>Gd$d%e3ZXed/>Gd&d'e3ZY dfd(ZZed/>Gd)d*e3Z[Gd+d,eJZ\Gd-d.eZ]eGd/d0e]Z^eGd1d2e]Z_eGd3d4e]Z`eGd5d6e]ZaGd7d8eJZbGd9d:ebZcGd;d<ebZdGd=d>ebZedgd?Zfdgd@ZgdhdAZhy#e$rdZd/ZY wxYw(i) 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 _NodeOrNodesc.t|ttfSN) isinstanceintr!xs ]/var/www/ramen.bs-engineer-server.com/venv/lib/python3.12/site-packages/torch/_inductor/ir.py _is_staticrs a#w (()frozencJeZdZUded<ded<ded<ded<d ed <d ed <y )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__rrrrs.,+<;('rrc"dfd |y)Nc |yt|ttfr|D] }| yt|tr|j D] }| yt|t t tttjjjtttt t"f sJdt%|dy)NzFound zE, which is not a supported top level IR node. See [Note: Inductor IR])rlisttupledictvalues ExpandView DynamicScalar AssertScalarrsympylogicboolalgBooleanr rEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputtype)nodesnode_check_tensorboxs rrz%validate_ir.._check_tensorboxs =  e} - ' & ' t $  ' & '! KK''//#)%  e %jk r)r_NodeOrNodes | NonereturnNoner) node_or_nodesrs @r validate_irrs>]#rcTttsJtdfd }|S)Nc0tt|i|Sr)getattrru)argskwargsnames rfnzops_wrapper..fns!wsD!42622r)robjectrrrrv)rstrr)rrs` r ops_wrapperrs( dC ,$t*, 3 Irc btt|tt|dfd }|S)Nct|tk(sJtt|Dcgc] }|| c}Scc}wrlenrange)indexi inv_orders rreindexz inverse_reorder..reindex%s?5zS^+++-23u:->?il#???Ar Sequence[_T]rr)rziprr)orderrrs @rinverse_reorderr"s*Sc%j 123I@ Nrcdfd }|S)Nct|tk(sJtt|Dcgc] }|| c}Scc}wrr)rrrs rrzsame_reorder..reindex-s>5zSZ''').s5z):;AeAh;;;rrr)rrs` r same_reorderr,s< Nrcdfd }|S)Nc |Srr)rreindex1reindex2s rrz fuse_reindexing..reindex8s((r)rrrz Sequence[_V]r)rrrs`` rfuse_reindexingr4s) Nr)rrI)rrrrIc`|td|Dr t|}|St||}|S)z1 Convert strides to fill order (argsort) c3\K|]$}t|ttjf&ywrrrrr!.0ss r z!get_fill_order..Hs Q 1sEMM.B CQ*,)allr_r`)seq shape_env sorted_idxs rget_fill_orderrBs< CQSQQ$+CL  !C0 rct|Dcic]\}}|| }}}tt|Dcgc]}|| }}|Scc}}wcc}w)z Convert stride order to fill order For channel last format, stride order = [3, 0, 2, 1] and fill order = [1, 3, 2, 0] ) enumeraterr)ridxposlookupr fill_orders rstride_order2fill_orderrPsR(1'7 883c3h 8F 8%*3u:%67&)7J7 97s A A ct||}tt|Dcgc]}d}}t|D] \}}|||< |Scc}w)z) Convert strides to stride order r)rrrr)rrr_outrelems rget_stride_orderr\sV !/sI >JCHo &1 &C &Z(4D  J 's Acyrrrreplace_symbols_with_hintss rir_node_to_tensorrisRUrcyrrrs rrrmsrc|y|r%tjjj}nt}|j Dcgc] }|| }}t |r.|jjDcgc] }|| }}ntj|}|j}|j}t|}t|}tjjjj5t!j"||||j%}ddd|Scc}wcc}w#1swYSxYw)N)sizestridedtypedevice)rwgraphsizevarsoptimization_hintr$get_sizeis_storage_and_layout get_layoutrFlexibleLayoutcontiguous_strides get_dtype get_devicerersuppress_guardstorch empty_stridedzero_) rrshape_fnrrrrrts rrrss y"77##55!" .AHQK .D .Q'(||~'<'<=!(1+==2248 KKME \\^F "4 (D $V ,F    # # 3 3 5   fE& %'  H /> HsD<;E (EEc0t|tr|sdgS|Sr)rrvalues rmay_convert_to_optionalrs%uv Lrc t|ts||St|tjr |jSt|t t frt|jStd|dt |jdy)Nzget_device_type(: )) rrrrrr OutputSpecget_device_typerrrrs rrrsq!SQY Au|| $vv A + ,q||~..#A3ba)9)9(:!<=rct|}|dvrtt|ddk(ryy| t|x}yddlm}t |tsJt|t||S)N)cpucudaxpu_backendtritonTFrI)TritonScheduling) rrrJrNcodegen.tritonr&rr issubclass)rrdevice_schedulingr&s r is_tritonr*sz Q F'' 6fXX. /8 ;!:6!B B  K0 ' .G5F0GG . ')9 ::rct|dk(S)Nr!)rrs ris_cpur,s 1  &&rc@t|trB|j2t|j st|j ryt jfd|j ddD}t jt j|j ddt j|j dd}t j||}tjjj|S)NFc3^K|]$}tjt|d&ywrN)rEqrF)rr alignments rrz-is_aligned_realized_tensor..s# FQ%((3q)$a ( Fs*-rI)rrmaybe_get_strider6 get_strider rAndOrr0Lerwrrguard_or_false)rr1aligned_stridesaligned_last_dim is_aligneds ` ris_aligned_realized_tensorr<s q& !    '  0  .ii F!,,."2E FOxx #Q'!**,r2BA)F?,<=J 77   * *: 66rcdt|t|k(rt|t|k(sJt|||D]p\}}}tjjj |dr2tjjj tj||rpyy)zP Returns true if the strides are equal, ignoring dimensions of size 1 . rIFT) rrrwrrstatically_known_leqr8rr0)strides1strides2shapedims1s2s rsignificant_strides_equalrEs u:X &3x=CM+II I5(H5 R 77   0 0a 8 ww..uxxB/?@  rcZt|s|Stdt||jDr|St ||j|j s|St |\}}g|j}t|j D]8\}}tjjj|ds1||||<:t|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 c3vK|]1\}}tjjj||3ywrrwrrstatically_known_equalsrrCrDs rrz2try_match_insignificant_strides..s3  B 00R8 79rIdatalayout)r rrr4rEr as_storage_and_layoutrrrwrrr> FixedLayoutrrroffset is_pinnedrReinterpretView)tensorstridesstorage old_layout new_striderr new_layouts rtry_match_insignificant_stridesrZs ! (   '6#4#4#67  $Wf.?.?.A6??CT U /7GZ%:$$%J&//+,'1 77   0 0A 6#AJJqM' J _'*E FFrc|jjdd}t|jDcgc]\}}| c}}|jd<ddlm}||ycc}}w)Noutputopruser_visible_output_idxs)record_original_output_strides)r find_nodesrrmetatorch._inductor.compile_fxr`)gm output_noderrr`s rgm_original_output_stridesrfsa((%%%215K#K$4$454Q4K/0J"2& 4s A$ct}|D]>}|t|jdz}|t|jdz}@t |S)NF unbacked_only)r@r)r r4r)inputssym_varsinps rget_symbolic_inputsrmsW!+HL$S\\^5II$S^^%5UKKL >rct|tr |j}t|tr|j }t|t r |j}t|t r|jSdSr)rrrMBaseView unwrap_view StorageBoxBufferget_namers r try_get_namert%sW!Y FF!X MMO!Z FF%a01::<:d:rceZdZUdZeZded<dZded<ejdZ d ed <ejdZ d ed <ejdZ d ed<ejdZ ded<ejdZded<eej"dRdZeej" dSdZedTdZdUdZdVdZdWdZdXdZdYdZdZdZd[dZd\dZdXdZd]d^d Z d_ d`d!Zdad"Z dbd#Z!dcd$Z"ddd%Z#ded&Z$dfd'Z%dgd(Z&dhd)Z'did*Z(e)djd+Z*dkd,Z+dgd-Z,dld.Z-dmdnd/Z.dod0Z/dpd1Z0dgd2Z1dqd3Z2drd4Z3dhd5Z4did6Z5dsd7Z6dld8Z7dgd9Z8dmdtd:Z9dd;dud<Z:dWd=Z;dUd>ZdxdAZ?dydBZ@ dv dzdCZAd{dDZBd|dEZCd}dFZDd~dGZE dv ddHZFdldIZGddJZHdgdKZIdgdLZJddMZKddNZLdsdOZMddPZNeOr e)dadQZPyy)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_idxc#Ktj}||zt_ d|t_y#|t_wxYwwr)rrv)rzolds rcurrent_originszIRNode.current_originsDs7%%"%- * &)F #cF #sA2 A ?Ac#Ktj}|t_ d|t_y#|t_wxYwwr)rrw)rrs rcurrent_stream_idxzIRNode.current_stream_idxNs3 ((%/" - ),F &F &s?/ ? <?cLt|ttttt fSr)rComputedBuffer InputsKernel InputBufferrSTemplateBuffer)rs ris_realized_nodezIRNode.is_realized_nodeZs&     rc,tj|Srrcreateselfs rwrap_for_loweringzIRNode.wrap_for_loweringgs%%rc2tj|||yr)r __setattr__)rattrrs r_post_init_setattrzIRNode._post_init_setattrjs 4u-rc>t|j}|jd||jdtjrt j nd|jdd|jdi|jd|jy)Nrzr|r~rr)r@rvrrJdebug_ir_tracebackr| format_stackrw)rrzs r __post_init__zIRNode.__post_init__psT223  73  V5N5N//1TX   t4  r2  d.F.FGrcBtd|jDS)Nc34K|]}|jywrrrdeps rrz(IRNode.get_read_names..|?s#((?r@ get_readsrs rget_read_nameszIRNode.get_read_names{?dnn.>???rc|jSr)r|rs r get_tracebackzIRNode.get_traceback~s ~~rc|jSrr~rs rget_origin_nodezIRNode.get_origin_nodercyrrrs rget_defining_opzIRNode.get_defining_oprcgS)z'Return subgraphs contained in this noderrs r get_subgraphszIRNode.get_subgraphs rctt}|j}t|tr(|j }|j r t|g}|D]}t |dr(|jr|j|j7tjjjjdij|jg}t|ts|D]J}tjjj j|d}|s:|j|L|S)N stack_trace postToPre)r@rzr ExternKernelrr~hasattrraddr _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrzr~rpre_grad_nodes node_namers rget_stack_traceszIRNode.get_stack_tracess)3 ,, dL )..0K$k]3 6Dt]+0@0@  !1!12OO))JJNN#c$))R( ".$7!/6I--QQUU%t #$((56 6,rTcdt|dd}|rt|dkDr|ddd}|js|gSg}|jD]8}|jd||j dz }|jd :|g|zS) Nzorigins=rz@=z...zstack_traces = { })rrrappendsplit)rshortenrzstack_trace_strrs r common_reprzIRNode.common_reprsWT9b9:; s7|b( "c*G$$&9 002 (K  " "#5 6 {006 6O  " "3 ' (y?**rc$t|t|j|z}ttt|}|r5t dj |}t |jd|dSt |jd|dS)Nz, z( z )(r)rrmaprindentjoinrr)rlinesr multiline new_liness r str_helperzIRNode.str_helpersU d4#3#3G#<==Se_% uzz%01I4j))*#i[< <4j))*!E7!4 4rc|jSrrrs rrzIRNode.get_dtype zzrcB |jS#t$rYywxYwr)rNotImplementedErrorrs rmaybe_get_dtypezIRNode.maybe_get_dtypes& >># #"    c2tdt|d)Nz#get_layout() is not implemented by !rrrs rr zIRNode.get_layouts!$GT |ST"UVVrcB |jS#t$rYywxYwr)r rrs rmaybe_get_layoutzIRNode.maybe_get_layout& ??$ $"  rc"|jSr)r rs rget_output_speczIRNode.get_output_specs  rcB |jS#t$rYywxYwr)rrrs rmaybe_get_output_speczIRNode.maybe_get_output_specs( '') )"  rc>t|jtS)z4True for single tensor output (excludes MultiOutput))rrLayoutrs rhas_tensor_outputzIRNode.has_tensor_outputs$446??rc2tdt|d)Nz!get_size() is not implemented by rrrs rr zIRNode.get_sizes!$Ed4j\QR"STTrcB |jS#t$rYywxYwr)r rrs rmaybe_get_sizezIRNode.maybe_get_size% ==? ""  rc"|jSrr rs rrAz IRNode.shapes}}rc4t|jSr)rrr rs r get_numelzIRNode.get_numelsT]]_--rctjjjt j |j dSNrrwrrstatically_known_truerr0rrs ris_zero_elementszIRNode.is_zero_elements0ww55ehht~~?OQR6STTrc0tdt|)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 rrs rrealizezIRNode.realizes "ODJ<"@AArc0tdt|)Nzcodegen_reference NYI on rrwriters rcodegen_referencezIRNode.codegen_reference s!$=d4j\"JKKrcyrrrs rrzIRNode.get_devicerrc.|j}|J|Sr)rrrs rget_device_or_errorzIRNode.get_device_or_errors"!!! rcyNFrrs rhas_exceeded_max_readszIRNode.has_exceeded_max_readsrc>tt|jrrrrrs r make_loaderzIRNode.make_loader!$t*"5"566rc>tt|jrrrs r make_indexerzIRNode.make_indexerrrc>tt|jrrrs rr4zIRNode.get_striderrcB |jS#t$rYywxYwr)r4rrs rr3zIRNode.maybe_get_stride"rrc>tt|jrrrs rrszIRNode.get_name(rrcB |jS#t$rYywxYwr)rsrrs rmaybe_get_namezIRNode.maybe_get_name+rrcv |jtjjvS#t$rYywxYwr)rsrwr graph_inputsrrs ris_input_bufferzIRNode.is_input_buffer1s4 ==?agg&:&:: :"  s ), 88cyrrr thresholds rhas_large_inner_fnzIRNode.has_large_inner_fn7rr graph_reusecy)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. Nrrusersrs r mark_reusezIRNode.mark_reuse:srcyrrrs r realize_hintzIRNode.realize_hintD rc>tt|jrrrs rrpzIRNode.unwrap_viewGrrc>tt|jrrrs r freeze_layoutzIRNode.freeze_layoutJrrc>tt|jrrrr allow_paddings rfreeze_layout_with_stride_orderz&IRNode.freeze_layout_with_stride_orderM"$t*"5"566rc>tt|jrrrrs rfreeze_layout_with_fill_orderz$IRNode.freeze_layout_with_fill_orderRrrc>tt|jrrrrs rfreeze_layout_with_same_orderz$IRNode.freeze_layout_with_same_orderUrrc>tt|jrrr exact_stridesr%s r freeze_layout_with_exact_stridesz'IRNode.freeze_layout_with_exact_stridesXr'rc>tt|jrrrs rget_read_writeszIRNode.get_read_writes]rrc6|jjSrr3readsrs rrzIRNode.get_reads`##%+++rc4t|jSr)rrrs r num_readszIRNode.num_readscs4>>#$$rc>tt|jrrrs rget_storage_numelzIRNode.get_storage_numelfrrc>tt|jrrrris rget_free_symbol_useszIRNode.get_free_symbol_usesir'rc>tt|jrrrs rget_reduction_typezIRNode.get_reduction_typenrrc>tt|jrrrs rget_reduction_sizezIRNode.get_reduction_sizeqrrcyrrrs r is_externzIRNode.is_externtrrcyrrrs ris_no_opzIRNode.is_no_opwrrc>tt|jrrrs rconstant_to_devicezIRNode.constant_to_devicezrrc>tt|jrrrs rget_mutation_nameszIRNode.get_mutation_names}rrc>tt|jrrrs rget_operation_namezIRNode.get_operation_namerrc>tt|jrrrs rget_inputs_that_alias_outputz#IRNode.get_inputs_that_alias_outputrrcyrrrs rrz IRNode.dtypes(+r)rzzOrderedSet[Node]rGenerator[None, None, None])rrrrPrrrrrr)rrrrrrrrrOrderedSet[str])rr{rr}rzOperation | Nonerzlist[Subgraph]T)rrr Sequence[str])TT)rzSequence[object]rrrrrrr torch.dtype)rztorch.dtype | Nonerr)rz Layout | Nonerr)rzOutputSpec | NonerrrSequence[_IntLike])rSequence[_IntLike] | None)rz)_IntLike | sympy.Rel | Sequence[_IntLike]rr r str | NonerrzIndentedBuffer | Nonerrrtorch.device | Noner torch.devicer$Callable[[Sequence[Expr]], OpsValue]r Callable[[Sequence[Expr]], Expr]rrrrrrrrrrrrFr Sequence[int]r%rrrrrtrrrrarrr0rar%rrrrdependencies.ReadWritesrzOrderedSet[Dep]rrrrrirrrrSequence[Expr]rrjrrrrZ)Qrrr__doc__r@rvrrw dataclassesfieldrzr|r~rr staticmethod contextlibcontextmanagerrrrrrrrrrrrrrrrrr rrrrr rpropertyrArrrrrrrrr r4r3rsrrrrrrpr"r&r*r-r1r3rr9r;r>r@rBrDrFrHrJrLrNrrrrrrr/s3=,/>04-4 1{00e.rrrrs rrzOperation.get_read_namesrrc6|jjSrr5rs rrzOperation.get_readsr7rctrrrs r get_outputszOperation.get_outputsrrctSrr?rs rget_unbacked_symbol_defsz"Operation.get_unbacked_symbol_defs |rctS)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?r=s rr>zOperation.get_free_symbol_usess ,|rcy)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 rNrSrgrV)rry)rrro)rr)rrrrr_rx)rrrrrTrzr list[Buffer]rrrrr}r{)rrrrrrrrrLrrrDrFr3rrrrrr>rrrrrrsr2" #$'"-@,"%*! !0rrcJeZdZUdZded<ded<ded<ded <ed d d!d Zd"d Zd#d ZeZ d$d Z d%dZ d&dZ d&dZ ed'dZeej$fd(dZed)dZd*dZed#dZd+d,dZd d-dZd.dZd/dZd0dZd1dZd2dZd3dZd4dZ y)5Loopsz:Base class for pointwise and reduction loop-body IR nodes.rjrr\rCallable[..., Any]inner_fnrarangesctjgfd|jD|jS)Nc36K|]}t|ywrr(reris rrz-Loops.get_free_symbol_uses..s FQq-0 F)r@unionrinner_fn_free_symbolsr=s `rr>zLoops.get_free_symbol_usess@"z|!! F$++ F  & &} 5  rc |jd|jjdt|j|j g|Dcgc]}|dt ||c}zd|jgzScc}w)N'= origin_node=)rrrrr inner_fn_strrr~)rnamesrs r_to_strz Loops._to_strsDKK$$%Q'DJJ!!#  >??rcHt|jjSr)rrrrs rr9zLoops.num_readsfs4((*7788rc2tdt|d)Nz+get_reduction_size() is not implemented by rrrs rrBzLoops.get_reduction_sizei!9$t*Q G  rc2tdt|d)Nz+get_reduction_type() is not implemented by rrrs rr@zLoops.get_reduction_typenrrc2tdt|d)Nz+constant_to_device() is not implemented by rrrs rrHzLoops.constant_to_devicesrrrrr})rrZrrrorgrVr`)rrrrrr)rrarrHrr)rrYrzSequence[Sequence[_IntLike]]rrprirrOrderedSet[Symbol])rr"rrUrzrTr{r~rdr)!rrrrrrbr>rr__repr__rrr r classmethodrrrHINDEXrrarrrrrrrrr9rBr@rHrrrrrsD    G$$) ! ! %   )H  # #:>**  ((+  ;W  @9   rrc|jrtjtd|Stjd|S)Nnanr)is_floating_pointruconstantfloat)rrs r nop_loader_fnr ys1 ||E%L%00||Au%%rcPeZdZddZd dZeZd dZd dZ d dZd dZ y) Pointwisecp|jrtt|jS|jSNr)rr r rrrs rrzPointwise.make_loaders)  "= ; ;}}rc$|jdSrrrs rrzPointwise.__str__rrcgSrrrs rrBzPointwise.get_reduction_sizes rcyrrrs rr@zPointwise.get_reduction_typerrcp|j}tj|xsd||||SNunnamed)rrustorer output_nameindexervarsloaders r store_outputzPointwise.store_outputs2 !!#yy1 74=&,OOrc|j}tjtd||}t ||j ||j SFMove this to a given device. Requires that all reads are to constants.override_devicerrrr)rrrConstantBufferr rrrrrs rrHzPointwise.constant_to_devicesK!!#Hn.?HP**;;   rNrkror`rd)rrer!Callable[[Sequence[Expr]], Never]rrrrr) rrrrrrrBr@rrHrrrr r sR)HPP3P P  P  rr cFeZdZUded<dZded<ddZ d dZy) Scatterrnoutput_indexerNr[ scatter_modec|j}tjtd||}t ||j ||j |j|jS)rr)rrrrr&r') rrrr!r%rrr&r'r"s rrHzScatter.constant_to_devices]!!#Hn.?HP**;;..**   rc|j}|d}tj|||j||||jS)Nr)mode)rrurr&r'rs rrzScatter.store_outputsT !!#  #Kyy  D''- . 4L""   rr)rrerr#rrrr)rrrrr'rHrrrrr%r%sB44"L)"    3    rr% logical_ormaximumminimummulr bitwise_xor)anyrminprodsumdotxor_sumz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNctvr tSdvr dfd }|Sdk(r dd}|Std)Nargmaxargmin argmax_value argmin_valueargmax_with_valueargmin_with_valuec|\}}|\}} dvrtj||}ntj||}tj||}t rtj ||}tj ||} tj |tj|| }tj |tj|| } rtj||ntj||} tj |tj|| }tj|||tj|||fS)N)r:r<r>) rultgteqr0ner+ logical_andwhere)aba_valuea_indexb_valueb_indexmaskequala_isnanb_isnantiearg_break_ties_leftrreduction_types rargmax_combine_fnz3get_reduction_combine_fn..argmax_combine_fns ! GW GW!PPvvgw/vvgw/FF7G,Ee$&&'2&&'2~~dCFF7G,DEucoogw.OP'w(VVGW-  >>$s(CDD $1 $1 rwelford_combinec P|\}}}|\}}}tjtjtj|tj||tj dt j||z }||z} || z } ||| zz||z||z|z| zz| fS)Ng)rurErDisinfrBrrfloat32) rFrGa_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_ws rwelford_combine_fnz4get_reduction_combine_fn..welford_combine_fns&' "FD(%& "FD( II & 1366&&3IJ S%--0E "H,J :-I**t eemh6BB rzunknown reduction_type=)rFtuple[object, object]rGrbrtuple[OpsValue, OpsValue])rF#tuple[OpsValue, OpsValue, OpsValue]rGrdrrd)r6r)rRrrQrSras``` rget_reduction_combine_fnres--#N33   $ )>  & :! , , 2 2 1 ."!"$;N;K"LMMrc eZdZUded<ded<ded<ded<dd ZeZedd d!fd Zd"d Zd#d Z d$d Z d%dZ d&dZ d d!dZ d'dZe d( d)dZe d*dZeej(df d+dZe d,dZe d,dZe d-dZe d.dZe d( d/dZe d0dZe d1dZe d( d2dZe d3dZxZS)4 Reductionrareduction_rangesrZrRr\ src_dtyper^reduction_hintc$|jdS)N)rrhrRrrs rrzReduction.__str__,s||LMMrczt|tjfd|jDzS)Nc36K|]}t|ywrr(rs rrz1Reduction.get_free_symbol_uses..4s PQq-0 Pr)superr>r@rrhrri __class__s `rr>zReduction.get_free_symbol_uses1s:w+M:=OZ\=O=O P$:O:O P>   rc|jSr)rhrs rrBzReduction.get_reduction_size7s$$$rc|jSr)rRrs rr@zReduction.get_reduction_type:"""rc tj|j|j|j|j ||}tj |xsd|||yr)ru reductionrrirRrstore_reduction)rrrrreduction_varsrs rrvzReduction.store_reduction=sT JJ NN    MM$ /   K49gdmUKrcXt|jt|jzSr)rrrhrs r index_lengthzReduction.index_lengthLs!4;;#d&;&;"<<.inner_reduction_splitssrNrzUse 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: %sr2rr3rc |j}|Jtdt||j|j |}|j }|j J|j Dcgc].}t|trt|tjs|0}}tj|j |j\\}}}|Dcgc].}t|trt|tjs|0}}g}g} d} t|jdD]} | j j"t%fd|D} | xrt'fd|D} | r|j)| j n| r| j)| j | s| s| j*t,j.j0vst,j.j0| j*}t3|j4dd}|j7t3|j4dd|k7sd } |xs| | fScc}wcc}w) NrrrrrNrMFc|jSrrrs rz@Reduction.num_splits..get_read_indices..s affrkeyc3&K|]}|v ywrrrvarr5s rrzAReduction.num_splits..get_read_indices.. s'R|(;'Rc3&K|]}|v ywrrrs rrzAReduction.num_splits..get_read_indices.. sN,/C<'NrrT)rrr rr r3 range_varsrr rr rKindex_vars_squeezerBsortedr6rr5rr0rrrwrname_to_bufferrrN decide_layout)rrcb read_writesrrrrwindicesbroadcasted_reduction_indiceschangedmdis_full_size_readis_broadcasted_reduction_readbuforiginal_strider5s @rget_read_indicesz.Reduction.num_splits..get_read_indicess:\\^F% %%%!++- B,,.K))5 55'11c4(C1NJ &2%D%D a224& " Q *c4(C1NN G,. )G[..4DE +!xx44 $''Rz'R$R!5F0E1#N3ANK-%NN288,2188B$(Eww!''"8"88gg44RWW=*1#**h*M))+"3::x>/Q&*G' +(;;WD DIs 83I!)3I&rc34K|]}|dk(xs|dkDyw)rrINrrs rrz'Reduction.num_splits..1s9AQ!a%9r)rrrrrr)rrgrztuple[Sequence[Expr], bool]),rrrwrrall_unbacked_explicitly_hintedr^rr  has_featurerLREDUCE_TO_SINGLE_ELEMENTr*choices should_use_cooperative_reductionrJsplit_reductionsr]rmulti_processor_count functoolsr reduction_split_factorINNERrrrrr rS!replace_backed_symbols_with_hintslogrrgrrKrr rBsimplify_with_ranges stride_hintsrkeysrOUTER)&r dst_dtyperirrrhrRreduction_numel input_nodeexprsrnumelrshould_reduce_to_single_elementarg_reduction_types should_splitpropsnum_smmin_elements_per_threadrouter_reduction_splitsr new_rangesnew_reduction_rangesextracted_numel_hintrrrrrrwranges1 num_outer num_innerrjrUouters& r num_splitszReduction.num_splitsjsj"!-"78ww>>uE ((!+ + ww//AA/Rf%WW%%77> +,''*=*= N;;+ + &! ! yy99 ( &/ / / (&99 (''  U " ((!+ + ''/,,"$ @I@Q@Q 00&$A "AJ@Q@Q 00&%A "  &)    &< " ?*+?LEz$**E11%*Z*K\\.2BDIH5(00 5 En,A. )!,JGQ w<1  ((!+ +'3'F'F JJL!..0( $NW   A  55aAAgg&&33>4 #7G 999EQ Q   y  &&(>$j) !&&(>$j) AHHs $QQ ctjjjt ||dfd |dvr@t t j d fd |dvrS|dvrfdSfdSS) z1Convert inner_fn from a reduction to an pointwisec tjfdtjDcgc] }t |c}DScc}w)Nc30K|] }|ywrr)rr|rvalue_fns rrz=Reduction._unroll_reduction_fn..fn..Ns UF+)rreduce itertoolsproductr)rr combine_fnrhrs` rrz*Reduction._unroll_reduction_fn..fnKsH##"+"3"3,<=q%(=#  >sA )r:r9r<r;r>r=c|Dcgc]}tj|}}||}t|tr|S|t j |t jfScc}wr)rexpandrrru index_exprrint64)rr|rr flatten_indexrs rrz0Reduction._unroll_reduction_fn..value_fndsd4::a%,,q/:: /eU+ LNN=#8%++F ;sA+)r>r=)r<r;c|dSrrrrs rrz0Reduction._unroll_reduction_fn..ssRYq\rc|dSrrrs rrz0Reduction._unroll_reduction_fn..tsE1r)rrarr)rrar|rarrc)rwrr guard_int_seqre_fixed_indexerr r)rrhrRrirrrrs`` @@@@r_unroll_reduction_fnzReduction._unroll_reduction_fn?s77++99:JK-niH     + 112BCM ) 3E * !KK !AA11- -HIrc tjjjt dk(r]dfd } | d| d| d| ddvs Jddfd } t j ||| t|Sdk(r+dvrdfd } ndfd } t j || |SttrottjkrSt |dk7st|jr0d k7r+t j ||j||S|j!|||| \} }dfd }||}|t"j$k(r| }|d k(rX| Jt'j(t*dd5t-| \}}dddJJ|j/|||||| S|dkDr|j1|||||| }d}tj2j4r t|t6r dd}||}|rht|j8t:sJt|j8|j8j<d|_|_ ||_!|_"|St6j t;||||}|S#1swYxYw)zy Create a reduction node. May split the reduction to multiple layers to expose more parallelism. rctjk(r t|Sjr+t|tsJt |t |St|tsJt |t|Sr) rrrrrrr rr)valrs rpy_cnstz!Reduction.create..py_cnstsg *9$00%c=9D49D9 :%%c;7BcB7s8OrrI)r3r5r2r0* not supported for zero-dimension tensors!c6tjSrrur)rrrRrtypes_to_initss rconst_fnz"Reduction.create..const_fns||ON$CYOOrr )r:r9c0tjdSrr)rrs rrzReduction.create..fns<<955rcnDcgc]}tjj}}||Scc}wrrrr)rrreduction_indexrrhs rrzReduction.create..fns1=M&Nuww||&NO&N#E?;;'O!2r4c`tr|S|dkDrt|tjS|Sr)rrrJ min_num_split)rrs r_maybe_increase_splitz/Reduction.create.._maybe_increase_splits//* qy5&"6"677 rr2NrTc:|j}t|dk7rytt|}|tj j vrytj j |}t|tsy|jjJ|Sr) rrnextiterrwrrrrrMr@)cur_node read_namesbufnamers r_find_split_reductionz/Reduction.create.._find_split_reductions"*!8!8!:J:!+#"4 #34Gagg&<&<<#''009C%c>:#88668DDDJrr)rrrzbool | float | int)rrrrv)rrrr)rrrzComputedBuffer | None)#rwrrsimplifyrrr rrrr!rrJunroll_reductions_thresholdrnrrrr^rrrr rS!create_multilayer_existing_rangescreate_multilayerr%mix_order_reductionrrMrgrh _split_size_original_inner_fn_original_ranges_original_reduction_ranges)rrrrirrrhrRrjrrrrhintrrrrrsplit_reductionrrrs ` ` `` @@rrzReduction.createysk$''**33MBR4ST a  $qz"1: qz O"_4 !""LM 4 P##!F| $  a !556 <##YF$   0O$v'I'IIv&!+vfkk/B%' ##11.  $ nn           e &e, ]22 2!N B;) ))n.>E 3V40 0 ) ))'3 3388 $  QY''  C#O}}00ZY5O'*$#8"<"/"6"6 BO0012B/>.B.B.S.STU.V+5=2390=M:J !!1-#-    k  s -KKc |dvrAt|r tdSt|rytj|j S|dvrAt|r tdSt|rytj|j St|rdnd}t|rdnd}||||||||f|||ftd|fd |S) N)rr9r;r=-infF)r1r:r<r>infTrrI)r3r2r4r5r0welford_reducerTonline_softmax_reduce)r0r r.riinfor1r)rRrzeroones rdefault_accumulatorzReduction.default_accumulatorKs S Se$V}$!%({{5)--- S Se$U|#!%({{5)---(/uQ&u-d1#T40 $dD1&+FmT%:     rc:|dk(rytj||S)Nrr)rgr rRrs r default_valuezReduction.default_valueks# - -,,^UCCrc|dk(r|S|dkr(|dkr#|tjk(rtjS|dkr(|dkr#|tjk(rtjS|S)Nr2i)r^r OUTER_TINY)rrrjs r_multilayer_second_step_hintz&Reduction._multilayer_second_step_hintssg B;! ! C.wrapper_fn..bodysi');<.wrapper_fns"1 _*/ 'Y ?2_DG = =-o> vvNN7K8NN?K@zz$g66v r)rSequence[Symbol]rr"rrv) rViewdynamic_reshape_indexerrwrrrrr0rF) rrrhrrrrr dense_indexr!r rs ` ` `` @@r_multilayer_wrap_loaderz!Reduction._multilayer_wrap_loaders>> Z .. / ((>> HHS%0! 4   # 6F    (rctdDs Jdtj|t|t|z dfd }|S)Nc3&K|] }|dk( yw)rINrrrs rrzDReduction._multilayer_wrap_loader_existing_ranges..s3a163z8Only enabled for numel_hint == 1, found original_ranges=c |dt}|td}|t|t|zSr)rr) merged_indexnew_reduction_index original_idxrroriginal_rangesrs rr!zEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fnsQ((>#o*>?L$S%9%;.intermediate_fns''A'A'AB BrNr)rrarrarrv)rfloat16bfloat16r rgrrrrwrrr rrrrr)rrrrir!r/r0rrrRrrjintermediate_dtype intermediater6rr5s @rcreate_multilayer_helperz"Reduction.create_multilayer_helpers(0 ??   !''          *668 C% C8J C  C WW%%77 o8VW 99 :~ *-Cs?/C"DDDD (&!+C,@,B!C-#-   rc t|} t| |dz z|} |j||} |j||| || | | }|j ||||||g||| g||| S)r3rI)rrrCr r&r;)rrrrirrrhrRrrjrrrrr!s rrzReduction.create_multilayer,s&((89o;UC ##NI>00        ++       f e  L     rc j|j|||||} |j|||| ||g|||| d| S)r3r2)r1r;) rrrrirr/r0rrrRrjr!s rrz+Reduction.create_multilayer_existing_rangesZsc$@@   %   ++      % +o + +     rrorrrr`rd) rrerr#rrrwr"rrr{rzSequence[Sequence[Expr]]rr)rrjrr\rir\rzCallable[_P, OpsValue]rrarhrarRzReductionType | Literal['scan']rr r IRNode | Nonertuple[ReductionHint, _IntLike]) rzrBr@rvryrrrHrrrrr^rrr r rrr&r1r;rr __classcell__rps@rrgrg$s((!!!!NHK( ) %# L L3 L L ) L  L=    %)RRRR) R # R - R8RR"R (RRh7N7,77 7 2 77r)6(=(=$(NNN N % N # N-N&N&N"N NN`$/ &>DD$/D &DD%(:G  &4A >%)('(-(" (  (  (/("( ((TD($2  &  0  >8= = =  = ' = ( = $2= = ,= &= = &=  = = ~%)+ + +  + % + # + -+ &+ + &+ "+  + + Z$ $ $  $ % $ , $ $6$ "$ ,$ &$ &$  $ $ rrgcdfd }|S)1A closure containing math to read a given elementct|tk(sJt|tk(sJ}t|D]\}}}|dk7s |||zz}|Sr)rr)rresultrstszrQrrs rrz_fixed_indexer..indexerss!c%jCK&???5zSY&&&ufd3 +KCRQw#(* + r)rrarr r)rrrQrs``` rrrs Nr INNER_FN_TYcleZdZUded< dfd Z ddZxZS)MultiOutputReductionr output_indexc trftdk(rd} n dfd } t | ||| |||||| |_y)NrIrc2tfdDS)Nc30K|] }|ywrr)rrr reduction_idxs rrz@MultiOutputReduction.__init__..loader..sHR]3Hrr)rrX inner_fnss``rrz-MultiOutputReduction.__init__..loadersHiHHHrr)rrrXrrtuple[OpsValue, ...])callablerrn__init__rT) rrrrZrrhrRrirjrTrrps ` rr]zMultiOutputReduction.__init__s I " I y>Q q\F I# I4B I% I -))  )rc :tj|j|j|j|j ||}t |ttfsJt|||j}tj|xsd|||Sr) rururrirRrrrrrrTrv)rrrrrwrrs rrvz$MultiOutputReduction.store_reductions JJ NN    MM$ /   &5$-0>$v,>0t(()"";#;)WT]ERRr)rrjrr\rZz#INNER_FN_TY | Sequence[INNER_FN_TY]rrFrhrFrRrZrir\rjr^rTr) rrerr#rrrwr"rr)rrrrr]rvrIrJs@rrSrSs#)#)#)7 #) " #) , #)&#)#)&#)#)JSS3S S ) S SrrScbeZdZdZeej df ddZy) ArgReductionzJMulti-output arg reduction that returns both the selected value and index.Nc dvsJtjjjt } | dk(rt d| dk(r]d fd } d d} t j|| tt jtj| tfSt| trt| tjkrpt dk7st!j"rMt$j'd fd t)fd t+|tjfDSt$j-|| | \} }|dk(sJd t.j0k(r| t)fd t+|tjfD}|D]}|j3|S)N)r=r>rrrIcDcgc]}tjj}}||}t|tr|dS|Scc}wr)rrrrr)rrrrrrhs rrz%ArgReduction.create..value_fnsH9I"JA577<<"J"J 8eU+ 8O #Ks!A cJtjdtjSr)rurrr)rs rindex_fnz%ArgReduction.create..index_fns||Au{{33rr cB|}t|tsJ||Sr)rr)rrTrN unrolled_fns rprojectz$ArgReduction.create..project s($U+!&%000l++rc 3~K|]4\}}tj|t|t6yw))rTr N)r rr r)rrTrrrgrs rrz&ArgReduction.create.. sF(L%   !$W<H< !s:=z.arg reductions do not support split reductionsc3jK|]*\}}tj| | ,ywrr) rrTrrrrrrjrhrRris rrz&ArgReduction.create..+ sM $ e   $""    03rrrrv)rrrTrrrv)rwrrrrrAssertionErrorr rrrrrr!rrJrrnrrgrrrrr^rr)rrrrirrrhrRrjrrrrdrrresultsrNrgrfs`` `````` @@rrzArgReduction.creates!KKKK''**33MBR4ST a  !""LM  a   4  !#%< !   !++%< !    0O$v'I'IIv&!+vfkk/B#88*NIK , ,5i5M+N  **           ezKKKz ]22 2!N  (1)U[[1I'J    F NN  r)rrjrr\rir\rrrrrhrrRrZrjr^rr?rSequence[TensorBox])rrrrrr^rrrrrr`r`sT)6(=(=$(ggg g % g  g)g&g&g"g ggrr`ceZdZdZeej df ddZeej df ddZe d ddZ y) OnlineSoftmaxReductionzAMulti-output reduction that computes softmax max/sum in one pass.Nc tfdt|D} | D]} | j| S)Nc3jK|]*}tjtd| ,yw)rN)rrrS) r output_idxrrrrrjrhris rrz:OnlineSoftmaxReduction._create_no_split..P sI    $$+"   rj)rrr) rrrrirrrh num_outputrjrrmrs `````` ` r_create_no_splitz'OnlineSoftmaxReduction._create_no_splitC sG  $J/    A IIK rc Btjjjt |} t j ||||||d| |  \} } |tjk(r| }| dkDr|j|||||||| || S|j||||||||| S)Nr)rRrrrI) rwrrrrrrgrr^rrru) rrrrirrrhrtrjrrrrs rrzOnlineSoftmaxReduction.created s''**33MBR4ST**      2+!+  e ]22 2!N 19((   ##           rc t|t|dz z||j| } tj |g| t j jjtjt|d  dfd } |tjtjfvr|ntj} |j!|| | g||g|| }|D]}|j#t%d|D dfd }t j jj't|}|j)||| } |j!||| |||g|| S)NrIrc |\}|^} |z|zd  fd }tjd} rt }tjtj|tj |}tj ||t dtj||tjdfS||fS)Nc$gSrr)rrrrsrrzJOnlineSoftmaxReduction.create_multilayer..wrapper_fn..body s 7G9+=>>rrIrrr)rurrhr@rrr rE)rrrrr rrLrrrrr rrris @@rr!z.wrapper_fn s"1 _*/ 'Y ?2_DG ? ?,,q),C-o> vvNN7K8NN?K@ JJtT5=9IIdCa)CD vs{"rc3<K|]}|jywr)rrrs rrz;OnlineSoftmaxReduction.create_multilayer.. s$LQ]]_$Lc2tfdDS)Nc36K|]}|gywrr)rrrrs rrzTOnlineSoftmaxReduction.create_multilayer..intermediate_fn.. s&7=1112rrY)rrintermediate_loaderss``rr6zAOnlineSoftmaxReduction.create_multilayer..intermediate_fn sAU r)rr"rr"rrc)rrarrarr[)rrrCrr#r$rwrrrrr0rFrr7r8r rurrr r)rrrrirrrhrtrrjrr%r!r9 intermediatesrr6rrrr rrs `` @@@@@rrz(OnlineSoftmaxReduction.create_multilayer s((89o;UC >> Z .. / ((>> HHS%0! 4   ## #6F # & # #4 ??   ,,     f e  L     A IIK  %$Lm$LL % 8J  ! WW%%77 f8MN 99 :~ ##      G    r)rrjrr\rir\rrrrrhrrtrrjr^rr?rrnr)rrjrr\rir\rrrrrhrrtrrrrjr^rr?rrn) rrrrrr^rrurrrrrrprp@ sK)6(=(=$(  %   )&" @)6(=(=$(1 1 1  1 % 1  1 )1 1 &1 "1  1 1 f%)Z Z Z  Z % Z  Z )Z Z Z &Z "Z  Z Z rrpceZdZeej f ddZe ddZe ddZ y)WelfordReductionc|dvsJtjjrJdtjj j t}d fd } |dk(r| d} | d} | d} | | | fS|dk(r@ d fd |dk(r|d| d| dfStfd|DStj|d|| \} }|tjk(r| }|dkDr|j||||Std Dcgc](}tj!t#|||| *}}|D]}|j%|Scc}w) N)rrTzGwelford reduction usage is explicitly disabled, please check you configcXdfd }tj|tS)Nc0tjSrr)rrrs rrz8WelfordReduction.create..const..inner_fn s||rr rrrrvr rr)rrrrrs` rconstz&WelfordReduction.create..const s2  ##!F| $ rrrIcXdfd }tj|tS)NcnDcgc]}tjj}}||Scc}wrr)rrrrrhs rrz7WelfordReduction.create..copy..inner_fn" s1=M&Nuww||&NO&N!#77'Orr rr)rrrrrrhs` rcopyz%WelfordReduction.create..copy s28!''!%< (rrc3.K|] }|ywrr)rrrs rrz*WelfordReduction.create..0 s:"T"X:)rRrr)rrrr)rrErr)rJmtiadisable_welford_reductionrwrrrrrrrgrr^rrrrrrr)rrrrZrrhrRrjrrmeanm2weightrrrsrmrrs `` `` @rrzWelfordReduction.create s!FFFF;;88 U 8''**33MBR4ST  a 8DqB1XFV# # a  L   !11IaL)58U1X==: :::& **    aL  )++  e ]22 2!N 19((   2$Ah     $""     A IIK % s0-E8cy)N)rrrrr s rr zWelfordReduction.default_values src ttjjj t j td } | rH|dk7rC d fd } j||dt| dt| df|d|Stdz ztj|tfd|Dg|g||} | D]} | j d dtjjjt|} j!| |}tj|tfd | D|gd|S) r3rrTc0tj|Srr)rrXrrs rrz4WelfordReduction.create_multilayer..constant s||E511rrrI)rrrZrrhrRrrjc 3LK|]}j|dyw)rrN)r&)rrrrrrhrs rrz5WelfordReduction.create_multilayer.. s= ++$# , s!$c|g||Srr)rrrs rintermediate_loader_fnzBWelfordReduction.create_multilayer..intermediate_loader_fn s 4E4O45 5rc3TK|]}t|j!yw))rN)r r)rrrs rrz5WelfordReduction.create_multilayer.. s*.q}}GG%()rrrXrrrrrv)rrrrrrlrrv)rrrwrrrrr0rFrr rCrrrrr r)rrrrZrrhrRrrjr rrrrrrrs` ` ` ` @@@rrz"WelfordReduction.create_multilayery s ((89((>> HHS%0! 4   +<< 2# 24B 2KN 2 2 ((aLHA.HA. !10-)  o;UC (//    (   f e  L  # & A IIK  6! 6+ 69 6  6WW%%77 f8MN 99 :~  &&   &   G    rN)rrjrr\rZSequence[Callable[..., Any]]rrHrhrHrRrZrjr^rrnrA)rrjrr\rZrrrHrhrHrRrZrrrjr^rrn) rrrrr^rrrr rrrrrr s)6(=(=yyy0 y  y ( y&y&y yyv$/ & \ \ \ 0 \  \ ( \ &\ \ &\  \ \ rrcxeZdZUded<ded<ded<ded<ded <d ed <d ed <ded<edddfd Zd fd Z d!dZd"dZd#dZ d#dZ d#dZ d$dZ d%dZ dddZeej"fdd d&dZe d'dZxZS)(ScanrH scan_rangesr=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rr^rjrrTtuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]rZct|tjfd|jDztjfd|j DzS)Nc36K|]}t|ywrr(rs rrz,Scan.get_free_symbol_uses.. O"1m4Orc36K|]}t|ywrr(rs rrz,Scan.get_free_symbol_uses.. H"1m4Hr)rnr>r@rrrros `rr>zScan.get_free_symbol_uses sa G ( 7 jl  Od>N>NO !jl  HdiiH   rct|jt|jzt|jk(sJt|yr)rrrrrnrrrps rrzScan.__post_init__ =4;;#d&6&6"773tyy>III rc|j||tfd|jD}tj|j |j |}tj|xsd|||jS)Nc3.K|] }|ywrrrrrs rrz'Scan.store_reduction.. Dx}Drr) rrrZrurrrrrT)rrrr scan_varsrrNrs @rrvzScan.store_reduction soll4+DT^^DD$++t?yy  $9gclF4;L;L4M  rcy)Ncustomrrs rr@zScan.get_reduction_type src|jSr)rrs rrBzScan.get_reduction_size rrc|jSrrrs rr z Scan.get_size yyrc|jSrrrs rrzScan.get_pointwise_size rrcXt|jt|jzSr)rrrrs rryzScan.index_length !4;;#d&6&6"777rc|j|j}|j|jtj}|j ||}|fSr)rrrrHr{rrrr|rs rrzScan.inner_fn_args E DKK(T--t}}=ll5&)v rc|j|j}|j|jtj}|j ||}t |j||Sr)rrrrHr{rrRrrrirr|rs rrzScan.inner_fn_free_symbols Q DKK(T--t}}=ll5&)#DMM3mTTrT)can_fallback_to_atenc g|d|dzd|gtjj|tjsdgt |zSt |dkDr=tjj|tj sdgt |zStjj} | jt} t |t |k(sJ| jtj| dr?tt |D cgc]!} tj||| || |#c} S|j!||d|d|| \}} t"}| dkDr[t$j&j(duxst*xr t,dk\xrt |dk(}|s|rdgt |zSd} nt.}dfd }tt |D cgc]0} t0j|d ||| ||| |||||| d | 2}} |D]}|j3|Scc} wcc} w) NrIr r)rrraxispointwise_rangesrr scan_numelz3.3.0ct|tk(sJt|tk(sJg|d||dSrr)r scan_indexrrrs rrzScan.create..reindexh Sz?c+&66 66u:%5!66 66>U5D\>J>tu> >r) rrrrrZrrrrrrjrT)rrrrrrGr)rwrrrLSCANrTUPLE_REDUCTIONrrrrrrr7rr rrrrversionhip has_tritontriton_version SplitScanrr)rrrrZrrrrjrrrrrTr scan_typesupports_splitrrmrNrrs ` @@rrz Scan.create$ s=T%4[<4q +;<Dzl ww""6>+>+>?6CK' ' v;?177#6#6 N22$ 6CK' '77##&&}['AB 6{c)n,,,  ) )%((:q*A B%*#f+$6!   ! .&|4 ! &)^^)q\-#!!&4 & "  > !!T)Wj.V^w=V%v;!# "' 6CK//!"J%  ?.!&c&k 2% $#   ! .!&|4'+ +)##1!-   * F NN  AP s )&I85I c Ndfd } tj|||| ||d|S)Nc,g|d||dSrr)rrXrrs rr!z#Scan.num_splits..wrapper_fn s*Fc%4jF=F3tu:FG Grr)rrrirrrhrRr)rrrXrrrv)rgr) rrrrrrrrrr!s `` rrzScan.num_splits s; H###(!&$  rrrrrS) rrerz%Callable[[Sequence[_IntLike]], Never]rrrr"rrrdr~r{r)rrjrrrZz+tuple[Callable[[Sequence[Expr]], Any], ...]rrHrrrrrjr^rrrrrSequence[TensorBox | None])rrjrr\rrlrrrrHrrHrrrr rr@)rrrrrbr>rrvr@rBr rryrrrr^rrrrIrJs@rrr s MM SS!! ##--F#  $       7    $    8 U )6(=(=`&*``(`? `  `  `R`&`#`` $``D   7    (  # R   (  rrc eZdZy)rN)rrrrrrrr srrcFeZdZUded<ded<ded<ded<ded <d ed <d ed <ded<ded<edddfd Zdfd Z ddZd dZd!dZ d!dZ d"dZ d#dZ d$dZ dddZeej"f d%dZxZS)&SortrH sort_rangesrz:Callable[[Sequence[Expr], Sequence[Expr]], Sequence[Expr]]rr^rjrrTrrrrZrstable descendingct|tjfd|jDztjfd|j DzS)Nc36K|]}t|ywrr(rs rrz,Sort.get_free_symbol_uses.. rrc36K|]}t|ywrr(rs rrz,Sort.get_free_symbol_uses.. rr)rnr>r@rrrros `rr>zSort.get_free_symbol_uses sa G ( 7 jl  Od>N>NO !jl  HdiiH   rct|jt|jzt|jk(sJt|yr)rrrrrnrrs rrzSort.__post_init__ rrc,|j||tfd|jD}tj|j ||j |j}tj|xsd|||jS)Nc3.K|] }|ywrrrs rrz'Sort.store_reduction.. rrr) rrrZrusortrrrrrT)rrrrrwrrNrs @rrvzSort.store_reduction sull40DT^^DD$++vt{{DOOLyy  $9gclF4;L;L4M  rcy)Nrrrs rr@zSort.get_reduction_type src|jSr)rrs rrBzSort.get_reduction_size rrc|jSrrrs rr z Sort.get_size rrc|jSrrrs rrzSort.get_pointwise_size rrcXt|jt|jzSr)rrrrs rryzSort.index_length rrc|j|j}|j|jtj}|j ||}|fSr)rrrrHr{rrs rrzSort.inner_fn_args rrc|j|j}|j|jtj}|j ||}t |j||Sr)rrrrHr{rrRrrs rrzSort.inner_fn_free_symbols rrc  g|d|dzd|gtjj|tjsdgt |zStjj } | jt} tjjrtjj} nCd} tjjxr%| jtj| | } | sdgt |zSt |t |k(sJ| jtj| dr?t!t |Dcgc]!}t"j%||||||#c}Sdfd }t!t |Dcgc]4}t&j%t)d|||||||||||||d | 6}}|D]}|j+|Scc}wcc}w)NrIrr ct|tk(sJt|tk(sJg|d||dSrr)r sort_indexrrrs rrzSort.create..reindex' rr) rrrrrZrrrrrjrTrr)rrrrrrGr)rwrrrLSORTrrrrrrJr%decompose_sort_opspersistent_reductionsrrr7rr rrrr)rrrrZrrrrrjrr sort_numelis_persistent_kernel max_rblockrTrrmrNrrs ` @@rrz Sort.create s=T%4[<4q +;<Dzl ww""6>+>+>?6CK' '77##&&}['AB == + +#)==#F#F J 33U22588J 3ST !$6CK' '6{c)n,,,  ) )%((:q*A B%*#f+$6!   ! .&|4 !  ?0!&c&k 2' &%   ! .!&|4'+ +##1!-!)   , F NN  Q s '&H.9HrrrrS) rrerrnrrrwrrrrd)rrFr`r{r>)rrjrrrZz'tuple[Callable[[list[Expr]], Any], ...]rrHrrrrrrrjr^rrrr)rrrrrbr>rrvr@rBr rryrrrr^rrrIrJs@rrr s>  GG!! ##-- LF#  $       2    '    8 U )6(=(=QQ(Q; Q  Q  QQQ&QQ $QQrrc> t|dy#t$rYywxYw)NFfreezeT)rOrrs rr r H s&a. s  c  t|d\}}|jrt|tsJt |t |j |j|j|j|j|j|j|j|j}|jS#t$rYywxYwNFr)rOshould_pad_stridesrr rrPrrr _pad_stridesrrQrR is_contiguousrr_bufferrNs r is_contiguous_storage_and_layoutrP s/%@  $ $ &fn5 CtF| C5    ##FMM6;; M   F##%% sCC CCc t|d\}}|jsytjjj |j |jt|jzS#t$rYywxYwr) rOrrwrrrI storage_sizerQrrrrrs r&is_dense_contiguous_storage_and_layoutrd sw/%@##%ww77    !6===3M#M   s A;AA;; BBct|trt|j|||||St|tr:t|j|||||\}}||jj fSt|t r|rn|r0|j|j js>J||j||n&||j||n|jt ||j fSt|tr(t|j|\}}||jfSt)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. rwant_contiguous stride_orderr%r0r%r)rrrOrMrqr rrr"rr&r1rrSrNr) rrrrr%r0rrNbuffers rrOrOp sG!Y$ FF+%''   !Z ) FF+%''  6!&&##%%%!V !||~33555)11 2*22!3!!}alln,,!_%* FF qxx rcd t|d\}}|j|S#t$rYywxYwr)rOis_stride_orderedr)rrrrNs r"is_stride_order_storage_and_layoutr s:/%@'' 55 s # //ct|ttfrt|jSt|t rt|j }tjjj|jt|jzt }t|jxs|St|tr*|j!tjj"vSyr)rrrq is_unalignedrMrSrNrwrrstatically_known_multiple_ofrQrirrkrrrsunaligned_buffers)rrNhas_unaligned_layouts rr r  s$J/0DII&&$(#$77#3#3#P#P MMN6<<8 8/$  DII&>*>>$}}!''";";;; rceZdZUded<eddddZddZddZddZe d dZ d!dZ d"d Z d#d Z d$d Zd%d Zd dd&dZd'dZd(dZd)dZd*dZd'dZd'dZd+dZd,dZd-dZd.dZy)/rorrMc8|jj|SrrMr>r=s rr>zBaseView.get_free_symbol_uses syy--m<.indexer & &r)rrrr )rMr r)rrrrs @@rr zBaseView.make_indexer s/ &&(%%' 'rcl|jj|jdfd }|S)Nc |Srrrs rrz$BaseView.make_loader..loader rrr)rMrr)rrrrs @@rrzBaseView.make_loader s/ %%'%%' ' rc6|jjSr)rMrrs rrzBaseView.dtype syy""$$rc6|jjSrrMr rs rr zBaseView.get_layout yy##%%rc6|jjSrrMrrs rrzBaseView.get_device rrcyrrrs rrzBaseView.get_origin_node rrc6|jjSrrMrsrs rrszBaseView.get_name yy!!##rc"|jSrrrs rrzBaseView.get_pointwise_size }}rTrc<|jj||SNrrMrrs rrzBaseView.mark_reuse yy##E{#CCrc6|jjSrrMrrs rrzBaseView.has_exceeded_max_reads yy//11rc6|jjSrrMrrs rrzBaseView.realize yy  ""rc8|jjyrrMrrs rrzBaseView.realize_hint s  rc6|jjSrrMr;rs rr;zBaseView.get_storage_numel yy**,,rc6|jjSrrMrDrs rrDzBaseView.is_extern yy""$$rct|jtsJt|j|jj Sr)rrMroris_module_bufferrs rr9zBaseView.is_module_buffer s6$))X.?TYY?.yy))++rc6|jjSrrMrrs rrzBaseView.get_read_names yy''))rctjtdd5t|j |j j cdddS#1swYyxYwr)rrr rTrr r6rs rrzBaseView.get_reads sL \\.*:D A &  " e    s 2AA!cd|}t|tr|j}t|tr|Sr)rrorM)rrs rrpzBaseView.unwrap_view s+H%AH%rc|j}tjtd||}t ||j ||j Sr)rrrr!r rr r"s rrHzBaseView.constant_to_device sP!!#Hn.?HP.."==?   rNrrrr*Callable[[Sequence[Expr]], Sequence[Expr]]rmrkr[r]rgrVror~rqr_rdrSr|rTrzrRr)rrrrrbr>rr rrrr rrrsrrrrrr;rDr9rrrprHrrrroro s LJ'=(=C%%&&$=AD2#!-%,*  rrocPeZdZUded<eddZed dZd dZ d dZ y) rrrctjj}|Dcgc]}tj|}}|j }dgt |t |z zt|z}t |t |k(sJtt |D]}||dk(r||J||||<||,tjjj||rMt|rYt|re||}||}|J|J||z }|j|ddk(rJd|j d|d||Scc}w)zReplace `-1` with correct sizesNr2r)fallbackzBroadcast failed in ExpandView(, z) on dimension ) rwrrrrr rrris_size_one_or_falser8r ) rnew_sizerrold_sizerv1v2diffs r_normalize_sizezExpandView._normalize_size. s77##-56ELLO66::<6S]S]:;d8nL8}H ---s8}% A{b {...&qk !$(8(8(M(M )./9 a[a[~%~~%~Bw..!"/ 6ajjl^2hZ_`^ab- <E7sEc|j||}t|rt|\}}t|t|jz }|dk\sJt j jg|z}t|j|jD]Y\}}|jtjjj|s|nt j j[t|j |j"t%|||j&|j(} t+|| St-||S)NrrL)rMr)rLr rOrrrrrrrrrwrrrFrPrrrrQrRrSr) rrrGrVrWskiprXrrrYs rrzExpandView.createV s&&q(3  #"7": GZx=3z#77D19 9'',,$.J #J$5$5z G  !!77++@@F  %!!  X!!$$ J# C Cqx00rc|jSrrrs rr zExpandView.get_sizeq rrc|j}|jjt|tz  dfd }|S)Nct|d}t|tk(sJttD](}|dk(s tjj ||<*|Sr)rrrrrr)rractualrNs rrz*ExpandView.make_reindexer..reindex{ sftu&Eu:V, ,,3v;' ,!9>$ww||E!H ,Lrrrrr)r rMr)rtargetrrRrNs @@rrzExpandView.make_reindexert sP##%6{S[( !   rN)rrrGrarra)rrrGrarror~r@) rrrrrrLrrr rrrrrr* sA %%N114 3rrcXeZdZUded<ed dZed dZd dZd dZ d dZ y) PermuteViewrGdimsc |j|}t|ttt|k(sJt |rt |\}}t |j|j|Dcgc]}|j|c}|Dcgc]}|j|c}|j|j}t||St||Scc}wcc}w)NrL)rMrW) _map_neg_dimsr@rrr rOrPrrrrrQrRrSrV)rrrWrVrWrrYs rrzPermuteView.create s  &$:eCI.>#????  #"7": GZ$!!  -12#2/34!""1%4!!$$ J# C C--34s 5C C#cR|Dcgc]}|dk\r|n t||zc}Scc}wrr)rrWrBs rrYzPermuteView._map_neg_dims s)@DEsaxSY_4EEEs$ct|j|jttt |jk(sJ|j j }|jDcgc]}|| c}Scc}wr)r@rYrWrrrMr )rrrs rr zPermuteView.get_size so$,,TYY78J #dii. !=    yy!!#!%+AQ+++7 Bct|j|jttt |jk(sJ|j j }|jDcgc]}|| c}Scc}wr)r@rYrWrrrMr4)rrrs rr4zPermuteView.get_stride so$,,TYY78J #dii. !=    %%'#'99-aq ---r\cLt|jDcic]\}}|| c}}tt|jDcgc]}| c}t t tt|jk(sJ dfd }|Scc}}wcc}w)Nc4Dcgc]}|| c}Scc}wrr)rrinvs rrz+PermuteView.make_reindexer..reindex s'**E!H* **s rS)rrWrrr@)rrrrr`s @rrzPermuteView.make_reindexer s!*$)) 451q!t5$S^45!s1v5#*U3tyy>-B"CCCC +! +  + 65s B B!N)rrrWrtrro)rWrtr list[int]r~r@) rrrrrrrYr r4rrrrrVrV sG ..$FF,.  3 rrVcFeZdZeddddZe ddZddZy) SqueezeViewN)rBc Dt|rOt|\}}g}g}|?t|tsJt |d|kr|t |j ksJtt|j |jD]\}\}} |Mtjjj|r5|j||j| X||k7r#|j||j| |dk(rJdt|j |j"|||j$|j&} t)|| S|[t*j-||j/D cgc]-} tjjj| s| /c} S|j/|dk(sJt*j-|t|j/D cgc] \}} ||k7s | c} }Scc} wcc} }w)NrrIzexpected squeezed size to be 1rL)r rOrrrrrrrrrwrrrFrrPrrrQrRrSr#rr ) rrrBrVrWrGrXrrrrYrs rrzSqueezeView.create s  #"7": GZHJ!#s+6T#Y6+CxC#joo*>$>>>%.s:??JDUDU/V%W K!>D&;77++@@F -"))&1Cx -"))&1#qyJ*JJy K%!!  !!$$ J# C C ;;;ZZ\77++@@C ::<$) ));;q1::<1H"UAAQTH1"UV V#Vs2H = H H c|Dcgc] }|dk7s | }}t|Dcgc] \}}|dk7s |c}}t|dfd }||fScc}wcc}}w)NrIct|tk(s J|dtjjgz}t |D] \}}|||< t |S)N )rrrrrr)rrrrlengthnot_ones rrz%SqueezeView.squeezer..reindex sku:W- C%'/C C-%*WW\\NV$;Igu- #Q!" # ## #r)rrrztuple[Expr, ...])rr)rrrGrrrhris @@rsqueezerzSqueezeView.squeezer s_ $.!qAvA..!*4;AAF1;T $  /;s AA AActd)Nzuse SqueezeView.create())rl)rrMs rr]zSqueezeView.__init__s788r)rrrBrrr)rrrz>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]])rMrrr)rrrrrrrjr]rrrrcrc sC48+W+WZ!! G!! 9rrccleZdZUded<ded< d dZd dZd dZeZe d dZ dd Z y ) GenericViewrrrArc|jSr)rrs rrzGenericView.make_reindexer s||rctt|jDcgc]}ttj |}}t |j|}ddjtt|d|Scc}w)Nzlambda rEr) rrrrqrHrrrrrr)rr index_old index_news r reindex_strzGenericView.reindex_strssCHTYYCX >? *4::q 9  i01 3sI#6789+FF  s!Bcz|j|jd|jd|jgS)Nsize=zreindex=)rrMrrrrs rrzGenericView.__str__s= YY% {+x8H8H8J7K-L M  rc*||t||S)NrMrr)r)rrrGrs rrzGenericView.createsX@@rc|jSrrrs rr zGenericView.get_size'rrNr@ro)rrrGrrrArror~) rrrrrrrrrrrr rrrrmrmsp  77 3 G HA A!A< A  AArrmceZdZdZeddZeed dZe d dZ e d d dZ e d d dZ y)r#z This class handles tensor reshaping by computing appropriate index transformations to map the new shape back to the original storage layout. ctj|}tj|}tjjj j }|tj|dr||z}|Sr)rrrwrrr evaluate_exprLt)rrrzs rhandle_negative_indexzView.handle_negative_index2sZll3||D!((22@@ #q) **C rcttsJtj|j \t j jjr|SttdkDxsttdkD}t|} d d d fd }dvrdfd }|t|S|r|tjSt|s||St!|d\}}|j"} t j jj%} t j jj%| } t j jj%} ddlm} | | | | | }||Dcgc]9}t+|d r|j,j.nt1j2|;}}t5|j6|j8||j:|j<}t?|| S||Scc}w)Nrctj|tj|j }t |\}}t |j|j|||j|j}t||SNrL) rrequire_exact_stridesr rr rOrPrrrQrRrS)rlrGrXrVrWrYs rcreate_reinterpret_viewz,View.create..create_reinterpret_viewKs~44^66s||~FC#8"< GZ$!!  !!$$ J# C Crc j}|t|S#t$r5tj |}|t j 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. rv)r$rr7rrequire_contiguousr r)rrrrrGrHs r"handle_unbacked_or_dynamic_reshapez7View.create..handle_unbacked_or_dynamic_reshape\sm 55hIXHH. !33A6.x!B!B8!L  s%);A'&A'c4tdgtzSr)rr)rrHs r fake_reindexz!View.create..fake_reindextsaS3x=011rrvFr)_compute_stride)size_obliviousrrL)rlrrGrrXrrrSrrrr)rrrztuple[int, ...]) rrrresolve_negative_sizer rwrrstatically_known_list_equalsrr6rrr rr rOrto_symints_or_intstorch._subclasses.fake_implsrrrexprrr!rPrrrQrRrS)rrrGunbacked_symbols_in_sizesrrrrVrW old_strideold_size_symintold_stride_symintnew_size_symintrnew_stride_symintrrXrYrrHs` ` @@rrz View.create;sb(H-=tH~=- 66qzz|XN( 77   8 88 LH %h/ 01 4 8(23a7 "?qA  D D#1 D?M D  D"     , = 2ADNLI I*8^>>xH  %Q'5a8 84AeD&& ''**==hGGG,,?? K''**==hG@,   4    (+ 'q&1 u}}Q7GGJ %!!  !!$$ J# C C2!44!s>IcF|Dcgc]+}tjjj|-}}|Dcgc]+}tjjj|-}}t |}t t |D]J}||dk(s tjj||<tt|t|||<ntjjjt|t|||fScc}wcc}w)Nr2) rwrrrrrrrrOnerBrr check_equals)rHrGrrs rrzView.resolve_negative_sizes;CCQAGG$$--a0CC:BCQAGG$$--a0CC>s8}% A{b #ggkk &}X'> h@WX    %%mH&=}X?VW!!DCs 0D0DNc |j|||}|S#tttf$r@t |g}|j||}|j||}t ||}Y|SwxYwr)_dynamic_reshape_indexerrlr7 IndexErrorrrr)rrHrG dense_dimrflatrrs rr$zView.dynamic_reshape_indexers} :228XyQG  ;ZH :!(+,D33HdCH33D(CH%h9G  :sAA+*A+ctjjjdfd }t t |Dcgc]}t tj|c}tt|}t|}|duxr!|t |dz k7xrt |dk(}|r&|J|j|}|j|g|r|r|j} |j\} } | dk(r>jtjj|j| | fn| dk(r|j| ni|| | dk(rj| nJ|| | dkru|| | dkr-|j\} } | | z| z} | | z} || | dkr-j| tjjj!| | n|| | dkDrtjj"}| }jt%| ||||z}|| | dkDrC|j}jt%| ||||z}| |z} || | dkDrCtjjj!| | nt&|r|r|rf|j} tjjj!| djtjj|rf|r@|j\} } tjjj!| d|r@|At |dk(r3j)j}j+||nj)t t |k(sJ dfd }|Scc}w)zG Perform a reshape entirely by modifying indexing math ctj||rytj||rytj||rytj j j||rytj j j||ryttj||)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. rr2rI) rr0r{Gtrwrrr r7)rFrGr8s r compare_sizesz4View._dynamic_reshape_indexer..compare_sizessehhq!n-ehhq!n-ehhq!n-ww<.reindex..HsHA|4Hr)rrrr)rrr view_exprs @rrz.View._dynamic_reshape_indexer..reindexCsOu:T* CSZT,C C*D% 01LHiHH Hr)rFr rGr rrrS)rwrrr8rrrqrHVIEWrrrrrrrrrrGrlreverseinsert)rHrGrrr stack_new stack_oldreordering_dense_dimold_dimsize_oldrsize_newvar2 size_new2divisormodulus dense_exprrr8rrs @@@rrzView._dynamic_reshape_indexers))88 >JCHH BV => *499a 8 T8,- N  T ! #S^a// #H "  ( ((mmI.G   W % I }}H%MMOMC1}  .  #x1Q  *x2a7  %x2Q6#Hh7!;&/mmoOD)/C/C')3H$Hh7!;  %  --hAx2Q6''++"  gw!GH!G+#Hh7!;'mmoG$$_S'7%KL%/G''1H $Hh7!;   --hA$$;I> }}H GG   ) )(A 6   UWW\\ * %MMOMC GG   ) )(A 6  S]a%7    "J   Y 3    9~X... I! I  IY s!O/)rr rr rr )rrrGrrr)rHrrGrrztuple[list[Expr], list[Expr]]r)rHrarGrarrr&Callable[[Sequence[_T]], Sequence[_V]])rHrrGrrrrrA) rrrrrr|rrrrr$rrrrr#r#+s  n5n5`" ",:" &""  !% $%  0  !%z z zz 4 zzrr#ceZdZUdZded<dfd ZddZeZddZddZ ddZ e dd Z dd Z dd Zdd Zdd ZddZddZed d ddZdddZd dZxZS)!rSz*Pretend our storage has a different layoutrrNct|t|jtr0t j |d|jjyy)NrM)rnrrrMrorrrprs rrzReinterpretView.__post_init__Ss@  dii *   tVTYY-B-B-D E +rcP|j|j|jgSr)rrMrNrs rrzReinterpretView.__str__Xs&     rc6|jjSrr"rs rrszReinterpretView.get_namebr#rc.|jjSr)rNrrs rrzReinterpretView.get_devicees{{!!!rcyrrrs rrzReinterpretView.get_origin_nodehrrc.|jjSr)rNrrs rrzReinterpretView.dtypeks{{   rc@t|jjSr)rrNrrs rr zReinterpretView.get_sizeosDKK$$%%rc@t|jjSr)rrNrrs rr4zReinterpretView.get_stridersDKK&&''rcdfd }|S)NcTjj}tjj ||}jj j j k7r5tj|j j j S|Sr)rNr ruloadrsrrMto_dtype_bitcast)rr tmp_loaderrs rrz+ReinterpretView.make_loader..loadervspkk..0G$--/75>BJ{{  DIIOO3++J DIIOOTT!!rrkrrrs` rrzReinterpretView.make_loaderus " rc6|jjSr)rNr rs rr zReinterpretView.make_indexer{{''))rc|jSrrNrs rr zReinterpretView.get_layoutrrcyrrrs rr"zReinterpretView.freeze_layoutrrct|jj|t|jj|zt|jj|zSr)r)rNrrrQr=s rr>z$ReinterpretView.get_free_symbol_usessQ T[[--} =t{{11=A Bt{{11=A B rcttjjj|j|j j |j j|j j| |jn#tjjj|j jSr) rwr wrapper_codecodegen_reinterpret_viewrMrNrrrQ writelinerrs rrz!ReinterpretView.codegen_referencesww##<< II KK   KK   KK   & 2F  8L8L8V8V++## =  rcyrrrs rr9zReinterpretView.num_readsrrSrorgrVr[r~rkrmr]rrr}rrfr{)rrrrrrrrrsrrrrr r4rr r r"rbr>rr9rIrJs@rrSrSMs4 NF  H$"!!&( * -.$) ! ! /   rrSc\eZdZUdZded<ed dZd dZeZe d dZ d dZ ddZ y ) DtypeViewz(Pretend our storage has a different typer\ target_dtypect|r]t|\}}t|j||j|j |j |j}t||St||S)NrL)rMr) r rOrPrrrrQrRrSr)rr new_dtyperVrWrYs rrzDtypeView.createsm  #"7": GZ$!!!!!!$$ J# C Cai88rcP|j|j|jgSr)rrMrrs rrzDtypeView.__str__s  4+<+<=>>rc|jSr)rrs rrzDtypeView.dtypes   rc6|jjSrrMr rs rr zDtypeView.get_sizer#rcLjjdfd }|S)Ncztj|jjjSr)rurrrMr)rrrs rrz%DtypeView.make_loader..loaders*''c D4E4EtyyW WrrrMr)rrrs` @rrzDtypeView.make_loaders" %%' X rN)rrrr\rroror[r~rk) rrrrrrrrrrrr rrrrrrsE2 9 9?H !!$rrcheZdZdZe ddZe d ddZy) SliceViewzqView that represents a slice along a single dimension. Corresponds to tensor[..., start:end:step, ...]. c tjj |j|t d||fDr t t n:t d||fDr t 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. c32K|]}t|ywr)r6rrs rrz0SliceView.normalize_start_end..sHA$Q'Hc3K|]L}t|tr:|jtjtj tt Nywr)rr hasrrErDrs rrz0SliceView.normalize_start_end..s; !T" EE%))UYYS 1 sAAcj||r|n||}j||r|}|S||}|Sr)statically_known_geqr>)rlowerupper clamped_lower clamped_fullmax_funcmin_funcrs rclampz,SliceView.normalize_start_end..clamps`221e<(1eBT  00F   mU3   rcD||Sj|}|||Sr)r|)rrrrrrdim_sizes r clamp_wrapz1SliceView.normalize_start_end..clamp_wraps0{++C:CeU+ +rr)rr rrrrrr ) rrrrrrr Expr | intrr) rwrrr r0rErD evaluate_min evaluate_max) rrrBstartendrrrrrrs ` @@@@@rnormalize_start_endzSliceView.normalize_start_ends77##::<$ H%h1GH HHH  S(+  HH,,H,,H  , ,$' ,03 ,>H ,  ,5!Xq1eXx8czrc  tjtts dkDsJ dk(r |dk\rdk(r|St |j  |r|j||\}t|z dz z <t|rt|\}}t |j} | z| <t|j|j | |j|jzz|j } t#|| S d fd } t%| | S#t$rYwxYw)NrlrIrLct|tk(s Jd|dt|}|zz|<|S)Nz wrong ndim rg)rr)rrBrGrsteps rrz!SliceView.create..reindex1sRu:X. P+eWAhZ0P P.KEsd*U2E#JLrrvrS)rrrr TypeErrorrr rrCr rOrrPrrrQrRrSr) rrrBrrrrrVrWrXrYrrGs `` ` @rrzSliceView.createsk||D!$%7471 zcY.419 % 00CDJE3 uq!94@  #"7": GZj//0J(o4JsO$!!  !!J$5$5c$:U$BB$$ J# C C !    ah@@K   sD>> E  E N) rrrBrrrrrrztuple[int, int])rIT)rrrBrrrrrrrrrrr)rrrrrrrrrrrrs // /),/36/ //b3A 3A3A 3A  3A  3A3A 3A3ArrcBeZdZUded<ded<d dZd dZd dZd dZy ) BaseConstantr\rrjrcyNrrrs rr zBaseConstant.get_sizeBsrc|jSrrrs rrzBaseConstant.get_deviceErrcyrrrs rrzBaseConstant.get_origin_nodeHrrctSrr?rs rrzBaseConstant.get_readsKrrNr~rgrVrz)rrrrr rrrrrrrr=s"  rrcDeZdZUded<ded<ded<d dZd dZd d Zy )Constantrrr\rrjrcdfd }|S)NcXtjjjSr)rurrrrrs rrz$Constant.make_loader..loaderVs<< DJJ7 7rrkrrs` rrzConstant.make_loaderUs 8 rcyrrrs rrzConstant.realize[rrcFt|j|j|S)N)rrr)rrrrs rrHzConstant.constant_to_device^sdjj 6JJrNrkrdr)rrrrrrrHrrrrrOs# J    Krrc<eZdZUded<ded<ded<d dZd dZy ) IndexingConstantrrr\rrjrcdfd }|S)NcXtjjjSr)rurrrrs rrz,IndexingConstant.make_loader..loaderis>>$**djj9 9rrkrrs` rrzIndexingConstant.make_loaderhs : rcFt|j|j|S)N)rrr)rrrrs rrHz#IndexingConstant.constant_to_devicensdjj 6RRrNrkr)rrrrrrHrrrrrbs J   Srrc8d}d}ttt||D]u\}}|dk(r tjj j ||s,tjj j ||sy|td|z}||z}wyNrIFT)reversedrrrwrrrIrD)rrAexpected_strideexpected_stride_maxrys ris_contiguous_strides_for_shaperrsOs5&123 1 6 ww77  ''""::1>QRs1ay(1  rc<tj|jzSr)rJpadding_alignment_bytesitemsizers rget_align_for_dtypers  ) )U^^ ;;rc2eZdZdZddZddZ d d dZy) rzxAbstract base for Layout, MultiOutputLayout, NoneLayout. Represents the memory layout of the output of an Operation.c>tt|jrrrs rrzOutputSpec.get_devicerrc>tt|jrrrs rrzOutputSpec.storage_sizerrc>tt|jrrr=s rr>zOutputSpec.get_free_symbol_usesr'rNrgr{rrr})rrrrrrr>rrrrrs,C77%*7!7 !7rrceZdZdZdeddf ddZeddZejddZed dZ e jd!d Z ed"d Z e jd#d Z d$d Z e Z d%d Z d&dZd'dZe d(dZd'dZed)dZd*dZd'dZe d+dZd,dZd'dZd-dZd.dZd/dZd"dZed d0 d1dZy)2rzo Layout base class Carries tensor meta-information including offset and whether it is pinned. NrFcP|tj|}||_||_t |t |k(s Jd|d|t d|DsJ||_||_||_||_ |jr!|jjdk(sJdyy)Nrt , stride=c3<K|]}t|tywr)rrrs rrz"Layout.__init__..s9q:a*9r|r!zOnly CPU tensors can be pinned) r rrrrr_size_stride_offsetrRr)rrrrrrQrRs rr]zLayout.__init__s >#66t>*4>>*:;MDz""#2dkk&6&6%78H7ITZZLYII;i }VH]O1 N rc|jSrrrs rrzLayout.get_devicerrc tj5tjt |j t |j |j|j|jcdddS#1swYyxYw)N)rr pin_memory) rw fake_moderrrerrrrrRrs r get_examplezLayout.get_examplesY [[ &&' 2' 4jj{{>>    s AA99BcBt|j|jSr)rrrrs rrzLayout.is_contiguouss.t{{DIIFFrct|}|dvs|ddk(ryt|t||D]\}}}|dk7s ||k7syy)N)rrIFT)rrr1)rArUndimleftrightrs ris_channels_last_contiguousz"Layout.is_channels_last_contiguoussb5z v qQ!$  *5 1  "   D%qyTU] rc t|jttj t t|j |j D]\}}}|dk7s ||k7syyr )rrrr rrr)rr<r=rs r is_transposedzLayout.is_transposedsc!$ KK ^66tHTYY ? @rct|jt|k(sJt|jDcgc]1\}}tj j j|ds|3}}}|Dcgc]}|j|}}|Dcgc]}|| }}dd}||}dgt|z}tt|D] }|||||<tt|dz D]u}||}||dz} tj j jtj|| rN|j|| ry|j| |ruyycc}}wcc}wcc}w)NrIc`t|}|Dcgc]}|j|c}Scc}wr)rr)arr sorted_arrelements rsorted_indicesz0Layout.is_stride_ordered..sorted_indices,s*J=@A'J$$W-A AAs+r2FT)rFrtrrt) rrrrrwrrrIrr8rr0rC) rrrrB non_1_indicesrrIstride_orderedr<r=s rrzLayout.is_stride_orderedsy4;;3u:--- $DII. 377##;;CC   +88Q$++a.882?@Qa@@ B u%E *s5z" 1A'-ayN58 $ 1s5zA~& A!!$D"1q5)Eww..uxxe/DE,,T5900= = 9@s6E29E8 E=c dgtttdt|jdz z}t|g|z}|j |SNrrI)rrrrrrr)s ris_channels_last_stride_orderedz&Layout.is_channels_last_stride_orderedCsNd8E!S-=-A$BCDDU u$%%e,,rcFt|}t|dk(r|Stjstj ||r|St j}t|dr|jjddr|Stdtj||D }tjs|r|Stt jdrt jj ndd fd rt#fd |Dr|St%|}t'|}t)t|Dcgc]}d} }d | |d<d} t+|d dd D]\} } || d z } | | || z}t-|t.t0j2fxr|tj4kDxr||zdk7xs"t-|t6xrtj}|| | <|st9|||z| | <d } | s|St:xj<d z c_| Scc}w)z The padding does not change stride order but makes sure all strides larger than the threshold are multiple of align. rrbdislike_paddingFc3\K|]$}t|ttjf&ywrrrs rrz&Layout._pad_strides..bs(  q3 . / r _shape_envNchyt|tsytfd|jDS)NFc3@K|]}j|ywr)is_unbacked_symint)rrrs rrzILayout._pad_strides..contains_unbacked_symints..psR1y33A6Rs)rr r0r5)rrs rcontains_unbacked_symintsz6Layout._pad_strides..contains_unbacked_symintsks0 dD)R@Q@QRR Rrc3.K|] }|ywrr)rrrVs rrz&Layout._pad_strides..ssNa6q9NrrI)rT)rrrr)rrrJpad_channels_lastrr>rwget_current_noderrbrrrchainpad_dynamic_shapesrrRr0rrrrrrrr!padding_stride_thresholdr rcr'num_comprehensive_padding) in_stridesrraligncurrent_fx_nodermrrr new_stridespaddedrankrprev_idxrrequire_paddingrVrs @@rrzLayout._pad_stridesIs,$E* z?a  ''F,N,N *-  ,,. ?F +0D0D0H0H u1   __Z6   ((Z *1!''<*HAGG&&d  S N:NN ' I> ,\: "'J"89Qq9 9&' JqM"":ab>; ID#!$(+H *T(^;F6C#78(V<<<(UNa'JVT*Hv/H/H  &K #*65#9E#A C    ))Q.);:s Hct|tsJt||jJ|j |j|j |j |_yr)rr rrrrrrs r pad_strideszLayout.pad_stridessM$/;d;/{{&&&'' TYY K rcFtjxrt|tSr)rJcomprehensive_paddingrr rs rrzLayout.should_pad_stridess++P 40PPrct|tr|S|jr|jt|j|j |j |j|j|jSr) rrPrrgrrrrrQrRrs ras_fixedzLayout.as_fixeds` dK (K  " " $     KK JJ II KK KK NN   rctjsJdt|jd|j j S)Nzconvert z to FixedLayout first)r rrrrkr rs rr zLayout.make_indexersG,, tDz**++@ A ,}}++--rcft|txr|j|jk(xr|j|jk(xrj|j|jk(xrO|j |j k(xr4|j |j k(xr|j|jk(Sr)rrrrrrrQrR)rothers r__eq__z Layout.__eq__s uf % 2 u||+ 2 ekk) 2 UZZ' 2 u||+  2  u||+  2 %//1 rcXt|j|j|jSr)r-rrrQrs rrzLayout.storage_sizes.tyy$++t{{SSrct|j|t|j|zt|j|zSr)r)rrrQr=s rr>zLayout.get_free_symbol_usess= TYY 6t{{M: ;t{{M: ; r)rrjrr\rrarzSequence[Expr] | NonerQr rRrrrr`)rrarrr~rrrrrcrr rrrori)r torch.Tensorr_)rArarUrarr)r<rr=rrr)rrtrr)r^rtrrrr\rrtrSrrPrm)rnrrrrrr})rrrrr!r]rrsetterrrQrrrr7rrr>r@rCrrNrrgrrkr rorrbr>rrrrrs)-qz   !  &       2 [[ ]] ]]  HG!,> " @ @"H- L!L)7L@KL LL\L Q  .  TH%$) ! ! & rrceZdZdZddZy)rPz A Tensor layout we cannot changecXt|j|j|jS)rL)rrrrQrs rr zFixedLayout.make_indexersdiidkkBBrNrm)rrrrr rrrrPrPs *CrrPceZdZdZdZddZeddZeddZeddZ e ddZ e ddZ e dd Z e jdd Z e dd Zejdd Ze dd ZejddZ d d!dZ d d"dZd#dZd$dZd%dZd&dZ d' d(fd ZxZS))r z| A Tensor layout that we are allowed to change Assumption: layout change should NOT add or remove free symbols FcHtj|jS)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. )rdeepcopyrkrs r!get_fixed_layout_without_freezingz0FlexibleLayout.get_fixed_layout_without_freezings}}T"++--rct|dk(rgStjjg}t |ddD]}|j ||dzt t |S)NrrIr2)rrrrrrr)sizesreversed_stridesrs rrz!FlexibleLayout.contiguous_stridessh u:?I!GGKK=U12Y' AD  # #D+;B+?$? @ AH-.//rcttt|t|k(s 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)r~r next_striderUrs r fill_orderedzFlexibleLayout.fill_orderedsx%E +, 50AAQE5>QAggkk &3u:% 1A$GAJ%a0K 1rcttt|t|k(sJt|}tj ||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)r~rrs rrKzFlexibleLayout.stride_orderedsB%E +, 50AAAA,U3 **5*==rc>|tjk(rtj|tS|tj k(rtj|t S|tjk(rtj|Stjd|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 rKNHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatrrrr)r~ memory_formats r stride_ordered_for_memory_formatz/FlexibleLayout.stride_ordered_for_memory_formats E// /!008IJ J e44 4!008JK K e55 5!44U; ; IIP & %rct|t|k(sJtjjj |}t t t||j}tj||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) rrwrrguarding_hints_or_throwrr __getitem__r r)r~rrs r same_orderedzFlexibleLayout.same_ordered*sb5zS[(((!!99&AE#f+.F4F4FG **5*==rc|jSrr#rs rrzFlexibleLayout.size9r$rc6|jd|||_y)Nr)!assert_free_symbol_uses_unchangedrr&s rrzFlexibleLayout.size=s ..vu= rc|jSrr(rs rrzFlexibleLayout.strideBr)rc6|jd|||_y)Nr)rr r&s rrzFlexibleLayout.strideF ..x? rc|jSrr-rs rrQzFlexibleLayout.offsetKr)rc6|jd|||_y)NrQ)rr!r&s rrQzFlexibleLayout.offsetOrrc0|j|j|}|jr)|r'|j||j|j}t |j |j|j||j|jSr) rKrrrrrPrrQrR)rrr%rXs ras_stride_orderzFlexibleLayout.as_stride_orderTsw((E:  " " $**:tyy$**MJ KK JJ II  KK NN   rc|}|jr)|r'|j||j|j}t |j |j|j||j |jSr)rrrrrPrrQrR)rr0r%rXs ras_exact_strideszFlexibleLayout.as_exact_stridesdsf#  " " $**:tyy$**MJ KK JJ II  KK NN   rc,|j|j|}|jr'|j||j|j}t |j |j|j||j|jSr) rrrrrrPrrQrR)rrrXs r as_fill_orderzFlexibleLayout.as_fill_ordertss$($5$5dii$G  " " $**:tyy$**MJ KK JJ II  KK NN   rc,|j|j|}|jr'|j||j|j}t |j |j|j||j|jSr) rrrrrrPrrQrR)rrrXs r as_same_orderzFlexibleLayout.as_same_orderss&&tyy&9  " " $**:tyy$**MJ KK JJ II  KK NN   rc ri}dD]/}dD](}||f}ttt|||||<*1|S)N)rrrQTF)r@r)r)rinitial_free_symbolsrrirs rget_initial_free_symbol_usesz+FlexibleLayout.get_initial_free_symbol_usessY!0 D!.  ]+,6$WT4%8-H-$S)  $#rcdD]:}|j||f}tt||}||k(r/Jd|d|y)Nrz)Expected free symbols unchanged, but got z vs )rr@r))rrrriold_free_symbolsnew_free_symbolss rrz0FlexibleLayout.assert_free_symbol_uses_unchangedsd* M#88$ 9NO )*:5-*PQ #'77 ; >&&-@& &&4 > >&8 >  > > [[ ]] ]] ;@ " 37  "HM / @D      $8< HHH H 5 H  H HHrr c\eZdZdZdfd ZddZddZed d d dZxZ S) NonOwningLayoutz,Is a view into the storage of another tensorc|j}t| |j|j|j |j ||_yr)r rnr]rrrrview)rrrNrps rr]zNonOwningLayout.__init__sA"  MM LL KK MM   rc>|jjSr)rkr rs rr zNonOwningLayout.make_indexers}}++--rc|jjj}|dk(ryddlm}t j jj||S)NrTrI) ALIGNMENT) rr rQutilsrrwrrr )rrQrs rmaybe_guard_alignedz#NonOwningLayout.maybe_guard_alignedsD%%'.. Q;$ww<)rribox input_buffers rr>z$NonOwningLayout.get_free_symbol_usessw$))_555iinn#z*5DI5*xx ,/:c:/""77 FFr)rzBaseView | TensorBoxrrrmr_rrr}) rrrrr]r rrbr>rIrJs@rrrsG6.P-.$)G!G !G/GrrceZdZdZy)CommBufferTypesymm_memN)rrrSYMM_MEMrrrrrsHrrcFeZdZUdZded<ded< dfd ZxZS)CommBufferLayoutax A layout that signifies the buffer is a comm buffer. In terms of striding, the layout is identical to `FixedLayout`. Buffers with this layout do not participate in in-place reuse - it can be neither the source nor the target for in-place reuse. For detailed motivation and usage of this layout, see NOTE [lowering-time collective optimization]. rcomm_buffer_typer group_namect|tr|jn|}t||j |j |j|j|j|j||_ ||_ y)NrrrrrQrR) rr rkrnr]rrrrrQrRrr)rrNrrfixedrps rr]zCommBufferLayout.__init__si &0%G!V <<++<<<<oo  !1$r)rNzFlexibleLayout | FixedLayoutrrrr)rrrrrr]rIrJs@rrrs; %$O%,%)% %%rrceZdZUded<ej dZded<ej dZded<d d Zdd Z dd Z y ) NoneLayoutrhrcdgSrrrrrrzNoneLayout. srdefault_factoryrarcdgSrrrrrrzNoneLayout. s1#rrcyrrrs rrzNoneLayout.storage_size rrc|Srrrs rrkzNoneLayout.as_fixed rc|jSrrrs rrzNoneLayout.get_devicerrNr{r^rg) rrrrrrrrrrkrrrrrrsG 'k'' DD)D) ))+FFIFrrceZdZd fd Zed dZej d dZd dZddZddZ e d ddZ ddZ dd Z xZS)MutationLayoutSHOULDREMOVEct||j|j|j d||_|j j}tjj|yr) rnr]rrr rT get_bufferrsrwrmark_buffer_mutated)rrTrrps rr]z#MutationLayoutSHOULDREMOVE.__init__se   & & (     OO      ))+ ##D)rc6|jjSr) real_layoutrrs rrz!MutationLayoutSHOULDREMOVE.stride"s!(((rcyrrr&s rrz!MutationLayoutSHOULDREMOVE.stride&s rc>|jjSr)rrrs rrz'MutationLayoutSHOULDREMOVE.storage_size*s!..00rcxdfd |j}t|tsJt||S)Nct|tr|jSt|tr|j St|t r|j S|Sr)rrrTrorp MutableBoxrM)rT unwrap_viewss rrz;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_views.sY&"<=#FMM22&(+#F$6$6$899&*-#FKK00Mr)rTrrr)rTrrrr)rrNrs @rrz%MutationLayoutSHOULDREMOVE.get_buffer-s6 dkk*&&)74<7) rc^|jj}t|tsJ|Sr)rrNrr)rrNs rrz&MutationLayoutSHOULDREMOVE.real_layout;s)"))&&))) rc |jtjj|j t |t r |j}|j|stj|j|j|jt|j|jDcgc]/\}}tjj j#||1c}}}t |t$t&fsJ|j}|jt)|dsJ|t |jj*t,s$Jt/|jj*t1||j_|jScc}}w)Nr rM)rrwrrrsrrrMrr rrrrrr rcheck_equals_and_simplifyrorrrNr rr)rsrcdst unsafe_aliasrFrGrs r realize_intoz'MutationLayoutSHOULDREMOVE.realize_into@sI  ##CLLN3 c9 %((C ##~~'mmo*!$CLLNCLLN C1GG$$>>q!D $DdXz$:; ;;))C sF#(S(##((//>:QDrrrerrrNcFt||jddyr)rnrrrs rrzBuffer.__post_init__ys   t4rc>|jjSr)r r rs rr zBuffer.make_indexer}s --//rc@|jsJ||jSrrrs rrszBuffer.get_namesyy$yyyrct|jtr|jjSt t |jj r)rrNrr7rrrrs rr7zBuffer.get_examples= dkk6 *;;**, ,!$t{{"3"<"<==rc>|jjSr)rrrs rrzBuffer.get_devices##%0022rcyrrrs rrzBuffer.get_defining_oprrc6|jjSr)r rrs rrz Buffer.dtypes &&&rc:g|jjSr)r rrs rr zBuffer.get_sizes("''((rc:g|jjSr)r rrs rr4zBuffer.get_strides*"))**rc6|jjSr)r rQrs r get_offsetzBuffer.get_offsets '''rct|jtr |jStt |jj r)rrNrrrrrs rr zBuffer.get_layouts4 dkk6 *;; !$t{{"3"<"<==rc|jSrrrs rrzBuffer.get_output_specrrc"|jSr)rrs rr;zBuffer.get_storage_numels~~rc6|jjSr)r rRrs r get_is_pinnedzBuffer.get_is_pinneds ***rct|jtr;t|jts |jj |_yyyr)rrNrrrkrs rr"zBuffer.freeze_layouts> dkk6 *: KK4 ++..0DK4 *rct|jtsJt|j|jj |||_yNr)rrNr rrr$s rr&z&Buffer.freeze_layout_with_stride_ordersA$++~6IT[[8II6kk11%}1U rct|jtsJt|j|jj ||_yr)rrNr rrr)s rr*z$Buffer.freeze_layout_with_fill_orders:$++~6IT[[8II6kk//6 rct|jtsJt|j|jj ||_yr)rrNr rrr,s rr-z$Buffer.freeze_layout_with_same_orders:$++~6IT[[8II6kk//7 rct|jtsJt|j|jj |||_yr)rrNr rrr/s rr1z'Buffer.freeze_layout_with_exact_stridessF$++~6IT[[8II6kk22 3  rctjjjt j |j dSrrrs rrzBuffer.is_zero_elementsrrcrjrttjSdfd }|S)Nrcxj}tjjxsd||Sr)r rurrrrrs rrz"Buffer.make_loader..loaders/'')G88DII2GENC Crrk)rr r rrs` rrzBuffer.make_loaders0  "=0@A A D rc"|jSrrsrs rrzBuffer.codegen_referencer%rcyrrrs rrzBuffer.decide_layoutrrct|jtr%|jjj gSyr)rrNrrrsrs rrNz#Buffer.get_inputs_that_alias_outputs/ dkk? 3KK$$--/0 0rct|jtr%|jjj gSyr)rrNrrTrsrs rrJzBuffer.get_mutation_namess0 dkk#= >KK&&//12 2rc6t|jgSr)r@rsrs rrzBuffer.get_read_namess4==?+,,rctSrr?r=s rr>zBuffer.get_free_symbol_uses |rctSrr?rs rrzBuffer.get_unbacked_symbol_defsrrcyrrrs rrzBuffer.realizerrcyrrrs rshould_allocatezBuffer.should_allocatesrrSrmro)rztorch.Tensor | torch.SymIntrgrWr[r~)rrGrcr]r^r{r_rrrsru)rrtrr)r0rtr%rrrrkrrfrrTr}rrd)%rrrrrr rsr7rrrrr r4rr rr;rr"r&r*r-r1rrrrrNrJrrbr>rrrrIrJs@rrrrros   50> 3'')+(>  +1;@V"V37V V 78 CH * ;?  U    -H%$)! !&  rrrc<eZdZddZddZej ZddZy)OperationBufferc|gSrrrs rrzOperationBuffer.get_outputss v rc|Srrrs rrzOperationBuffer.get_defining_oprrcXtj|tj|yr)rrrrrs rrzOperationBuffer.__post_init__sT"%rNrrrrS)rrrrrrrLrrrrr r s#55&rr ceZdZddZy)rcyrrrs rr9zInputBuffer.num_reads rrNr{)rrrr9rrrrr srrceZdZdZy) 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<ddZddZy)r!Nrhrcdfd }|S)Ncjj}tjtj j jj||Sr) r r rurrwr constant_namersrrs rrz*ConstantBuffer.make_loader..loadersPoo'446G88%%dmmot7K7KL rrkrrs` rrzConstantBuffer.make_loaders  rcttjj|j ||j SNrrN)r!rwrrrsrNrs rrHz!ConstantBuffer.constant_to_device%s/&&t}}?   rrkr)rrrrrrrHrrrr!r!s+/O(/ rr!cVeZdZddZed d d dZd d dZd dZd dZy)NoneAsConstantBufferctSrr?rs rrzNoneAsConstantBuffer.get_reads-rrctSrr?r=s rr>z)NoneAsConstantBuffer.get_free_symbol_uses0rrNcJtjjjSr)rwrrnone_strrs rrz&NoneAsConstantBuffer.codegen_reference6sww##,,,rctdSNr)rrs rrz$NoneAsConstantBuffer.get_output_spec9s &&rcyrrrs rrz&NoneAsConstantBuffer.has_tensor_output<rrrzrrr}rrfr^r_) rrrrrbr>rrrrrrrr+sC23$)! !4 -'rrcReZdZUded<ed d ddZd d dZd dZy) rr rc.t|j|Sr)r)rr=s rr>z*ShapeAsConstantBuffer.get_free_symbol_usesDs  =99rNchtjjj|jSr)rwrrcodegen_sizevarrrs rrz'ShapeAsConstantBuffer.codegen_referenceJs!ww##33DII>>rcyrrrs rrz'ShapeAsConstantBuffer.has_tensor_outputMrrrrr}rrfr_)rrrrrbr>rrrrrrr@s< J34$):!: !:5: ?rrceZdZUdZded<dZded<dZded <dZd ed <dZd ed <dZ d ed<e jd%dZ e e jd%dZd&dZd'dZd(dZd)dZd*dZed d+ d,dZd-fd Zd.dZd/dZd0dZd1dZe d2dZ d3 d4dZe d5 d6dZd7dZd7d Z d&d!Z!d.d"Z"d.d#Z#d8d$Z$xZ%S)9rzb Represents a buffer that is computed during kernel execution rather than being an input. rrMFzClassVar[bool]_force_realizeNrrCallable[..., Any] | Nonerrbrrc #K|jJ|jJ|jJ|jJt |j t sJt|j |j }|j} t |j|j|j|j|j|j|j|j}||_t|j|j|j|_|jj!|d||_||_y#||_||_wxYww)Nr)rrrrrrMrgrrNrrrRrirjrPget_default_sizes_body clear_cache)rold_datarWnew_datas rwith_original_inner_fnz%ComputedBuffer.with_original_inner_fn`s=+++&&222$$000..:::$))Y/EDO3DE/99[[  % nn00,,!%!@!@'66",,'66 H!DI&%%DK  ' ' 3 3D 9  DI$DK!DI$DKsBE$ B9EE$E!!E$c#Ktj} dt_d|t_y#|t_wxYwwNT)rr() old_values r force_realizezComputedBuffer.force_realizes3#11  6,0N ) ,5N )IN )s?/ ? <?c|j |jSt|jdr|jjSy)z Returns self.name if it exists, otherwise returns the name of the data node if that exists. If neither exist, returns None. Nr)rrrMrs rget_computed_buffer_namez'ComputedBuffer.get_computed_buffer_names7 99 99  499f %99>> !rc6|jjSrrMr9rs rr9zComputedBuffer.num_readsr7rc6|jjSrrMrrs rrzComputedBuffer.get_readsr7rc6|jjSrr;rs rrzComputedBuffer.get_read_namesr<rcXt|jtttt fs0t jtttStjtdd5|jjrTt|j|jj|jj!cdddSt|j|jj#cdddS#1swYyxYw)Nr6writes index_exprsrT)rrMrgrrr rK ReadWritesr@rrr r@rTget_store_functionrrBr rs rr3zComputedBuffer.get_read_writess$))itY%GH** l!|&L  \\.*:D A yy++-*++-II002II002  +++-II&&(   s6A%D %1D  D)c|jj||jj|z}|jr"||j j|z}|Sr)rNr>rMhas_store_functionr3)rrirNs rr>z#ComputedBuffer.get_free_symbol_usessc$11  II * *= 9:  " " $ d**,AA-P PF rc|js_|jtjjvr9|j dk(r&|j s|jjSt|!Sr) r@rrwrmutated_buffersr9r(rMrrnrs rrzComputedBuffer.make_loaders`'') !8!88 A%''99((* *w"$$rcVt|jtttt fSr)rrMrgrrr rs rrBz!ComputedBuffer.has_store_functions$))itY%GHHrc|jjj}t|jt t tfr+t|jj|j|St|jtsJt|jt|jj|j|Sr)r rkr rrMrgrrr rvrr rr)rrs rr@z!ComputedBuffer.get_store_functions//#,,.;;= dii)T4!8 949944diiI Idii3 DT$))_ D349911499gF Frc Pt|jtr{tj|j j |j j\\}}}|jj}td|DsJ|Dcgc]_}t|tjrCt|j|Dcic]#}|dk7s |tjj %c}a}}}|rt|j t"t$fr|j j'||}n|}|Dcgc],}t(j*j,j/||.} }ddlm} | | |j5Sycc}wcc}}wcc}w)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. c3pK|].}t|tjtjf0ywr)rrKStarDep MemoryDepr)s rrz0ComputedBuffer.get_fill_order..s01|33\5K5KLMs46rrIpick_loop_orderN)rrNr rKrrMrrBr3r6rrJrsrrrrrrrrwrrr schedulerrLr ) r index_varsrwrr6rvrrstride_lengthsrLs rrzComputedBuffer.get_fill_ordersj dkk> 2.:.M.M ,,. 0L0L0N/ + (Z!((*00E  a!7!78177n$WPQUVPVQ _$WXE dii$6"ii// NKG(GMR"EIAGG$$11$@""7&~t}}GG#%X"s$3F FF6 F1F#Fct|jtr5|j}|r|j |y|j yyr)rrNr rr*r"r)s rrzComputedBuffer.decide_layout s@ dkk> 2'')E2259""$ 3rcRtj|j|jd\}}t j t d|j5t|j|jr|n|dd|g|}dddg}g}g}g}|jD]^\}} ||dvr'|rJ|j||j| 4||dvsJ|j||j| `||f||ffS#1swYxYw)NqrrrIr) rKrrrBrrr!rrWr@r@itemsr) rr var_rangesrrN reduce_vars index_size reduce_sizerOrs rr+z%ComputedBuffer.get_default_sizes_bodys@(::  # # %t'>'>'@ j\\.*;T__=N O '')002Ra D  !#   $$& &DAqDG|&&!!!$!!!$DG|#|""1%""1% &K($[0III)  s !2DD&cj\\}}}\}}|r|||f|||f\\}}}\}}g|jj|t|trt |dk(sJ|\}} t|t sJt|t| tsJt| td| DsJ|j} | |k(s J| |f| D cgc] } | vs|  } } | z g|jtjjtj sj#|j% dfd } ||z} t't) xst*j, }| || ||\}}}| || ||\}}}t/j0||d\\}}}t3|||||g|||}||f|fScc} w)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. rc3<K|]}t|tywr)rr )rfs rrz6ComputedBuffer.simplify_and_reorder..dsHqz!T*Hr|c j||| \}}}jdk(rat|dk(rStt t|}||ddk7r(|Dcgc]}|| }}t |}t |}||}|rGtjjj||t ||\}} } t|| } n|} || |fScc}w)Nr4rr) _apply_loop_reorderingr@rrrrrrwrr_simplify_loopsrOr)x_vars support_varsr~simplify_loopsnewsizesreindex0rrrr_prunerindex_formulas memory_addrsrs rsimplify_and_reorderzAComputedBuffer.simplify_and_reorder..simplify_and_reorderus,0+F+F e\, (Hh*&&(E1c%jAoU3u:./E?1%*278Qa8H8+E2H.u5Hf%F-.WW-=-=-M-M,^VXN.*(F *(H="Wh. .! 9s* C&pr) r_Sequence[sympy.Symbol]r`rir~rtrarrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]])r+indexing_exprsrrrrrrrrrUget_write_exprsrwrrrLPREFER_STORE_LOOP_ORDERextendget_read_exprsrnrrJloop_ordering_after_fusionrKindex_vars_no_squeezerW)rextra_indexing_constraintsrecompute_sizes_body_funcrWrXrrNrVextra_indexing_rangesextra_indexing_exprexpected_var_rangesrrgr`should_merge_loops iter_ranges iter_reindexr reduce_rangesreduce_reindex iter_varsrUrerfs` @@rrgz#ComputedBuffer.simplify_and_reorder6s4  ' ' )  %Z  %Z % *[)4*k1J  )[)[ 94..5578 % 15u=23q8 9:T 6 !#63T: WDAV2I# # 1 1N0--/0 ww""4)O)OP    3 3 5 64 /*4 /04 /!4 /! 4 /  4 /l"K/ t,- - VV5V5V1V (<     ( $ \1,@ {4F, ( ~q 0rrBr@rrrar+rgr]rrBr@rFrrHrIrJs@rrrQs K%*NN*#K"481826/6<@ 9@ % %D66 %%**,-$)! !.2 %IG%N%J JJDOS?CU2$KU2$=U2 ? U2n *. %B*%B,%B%B! %B ' %B   %B%BN...,4rrc:eZdZUdZded<ded<ded<ded<y) FinalizeCodegenResultzNStructured result from TemplateBuffer._finalize_codegen for external backends.rsourcerimports call_preamble call_argsN)rrrrrrrrrr sX K rrceZdZdZ d dfd ZeddZddZddZdddZ ddZ dd Z dd Z d dd Z dd Zdd Zd dZ d!dZed"dZedddd d#dZxZS)$rz Base class for template operators that support epilogue and prologue fusion. Subclasses: TritonTemplateBuffer (built-in Triton templates), HelionTemplateBuffer (Helion kernels), etc. Nct |d|tj||_||_t jj||_ t jj|i|_ i|_ i|_ |r t|ni|_||_g|_|i|jd}t%|t&sJt)||j+}|D cgc]} t-t/|| |c} |_|xs t1|_d|_d|_ycc} w)Nrrr)rnr]runwrap_storagerjmake_kernel_renderrwrregister_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsmutated_inputsmutation_outputsrrrrMutationOutputrr@allowed_prologue_inpsallow_epilogue_fusionallow_prologue_fusion) rrNrjrrr named_inputs first_inputrrrps rr]zTemplateBuffer.__init__s/ d62"11&9 "4GG++D1  ""4(+-9;%?A# #/D B  -68  %++a.Kk62 ED4E E2 ++-F*%z8#tD%D ! " 1Z\ " 37"26"%s3D;ct|jtr td|j j S)Nz1Multi-output templates do not have a single dtype)rrNMultiOutputLayoutrr rrs rrzTemplateBuffer.dtypeLs6 dkk#4 5%C  &&&rc&|jdS)NT normalize)rTrs rr3zTemplateBuffer.get_read_writesTs''$'77rct}|jD]}t|ttfsJt ||tj tsJt j j jdfd }|tj|jd|jz}|S)z(Build read dependencies from all inputs.cxt|dk(sJtjj|Sr)rrurrs)rr|rl inp_indexers rdummyz4TemplateBuffer._read_deps_from_inputs..dummy`s16{a'''xx  E0BCCrrrr Sequence[Any]r|rrr) r@rjrrSrrrrNrr rKrTr r6)rrr6inp_rawrrlrs @@r_read_deps_from_inputsz%TemplateBuffer._read_deps_from_inputsWs.8l{{ Gg'@A P4= PA,3Ccjj&1 C4 3C C1**113K D \55s||~rYe E  rc(t|jtrutt j |j tjdddg}t j|j||tddS|j |jjdfd }t j||jd|}|xj|j|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)r6r=r>rrUc^t|dk(sJtj|dSNrfakerrurrr|rrs rrz1TemplateBuffer.extract_read_writes..dummy,v;!# ##99T75>6: :rrr)rrNrr@rKrJrsrr!r?rr r rTr r6)rrr=rdepsrrs @@rrTz"TemplateBuffer.extract_read_writesis dkk#4 53= ** q)9Rb4F **11)<&L  }}//#002 ;// 4==?B)  d11)<<  rc6tjjSr)rrrrs rrBz!TemplateBuffer.get_reduction_sizesww{{rcyrrrs rr@z!TemplateBuffer.get_reduction_typerrcyr1rrs rrzTemplateBuffer.should_allocaterrc*|jgfdfSrr)rrrrss rrgz#TemplateBuffer.simplify_and_reorders$      rc6t|jtS)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrNrrs ris_multi_outputs_templatez(TemplateBuffer.is_multi_outputs_templates$++'899rc|jSr)rrs rget_allowed_prologue_inpsz(TemplateBuffer.get_allowed_prologue_inpss)))rc"|jS)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)rscheduler_nodes r,has_aliasing_or_mutation_for_prologue_fusionz;TemplateBuffer.has_aliasing_or_mutation_for_prologue_fusions6688rcy)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)r hook_outputss r_finalize_codegenz TemplateBuffer._finalize_codegensrc(t|tr&t|jtr |jStj |}t|t r |j}t|jtr|j|S)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrrM MultiOutputr realize_inputrqrNr r")rrrNs rrealize_template_inputz%TemplateBuffer.realize_template_inputsf b) $BGG[)I77N++B/ fj )[[F fmm^ 4  " r)direct_alias_at_leafon_tensor_leafon_non_tensor_leafcxitjdfd t|gS)zLWalk a structured output tree, creating MultiOutput nodes for tensor leaves.c Lt|ttfr>g}t|D],\}}|j |g|t ||f.|St }t|tjr r| vrtj |gSt|}| vr |gSttj||}|j|j!<  |j!|||t|}| |<|gS  |gSr)rrrrrnrrrTensorrridrFallbackKerneltensor_to_layoutrrs)r\rrmritemleaf_idxtidmorr leaf_counterrr seen_outputs template_bufwalks rrz0TemplateBuffer.build_multi_outputs..walks5&4-0+-(0NGAtNN4.K.K$v,9J.K#LMNL)H&%,,/'H8L,L%,,-A(-KLMMj,&(-.. "33F;\7FH 33BKKMB!-"2;;="gxHr]$& S!t !-"8,Ir)r\rrzlist[tuple[type, int]]rlist[TensorBox])rcountr) rr structuredrrrrrrs ` ```@@@rbuild_multi_outputsz"TemplateBuffer.build_multi_outputss4.0  (   6T*b)**r)NNN)rNrrjSequence[IRNode]rr)rIterable[IRNode] | NonerOrderedSet[str] | Nonerzdict[str, IRNode] | Nonerrr[rx)rrrzOrderedSet[dependencies.Dep]rr)rrrryr~rdr_r)rrrrsr)rz9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]rT)rrrr)rzdict[str, str]rzFinalizeCodegenResult | None)rrrr) rrrrrzdict[int, IRNode] | NonerzFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | NonerzCallable[[int], None] | Nonerztuple[TensorBox, ...])rrrrr]rrr3rrTrBr@rrgrrrrrrrrIrJs@rrrsM378<150707!076 07 0 07 6 07/07 07d''8$$L OS?C  $K  $=   C  :*9 *  %    :>;?)+$)+)+ 7 )+  )+9)+ )+)+rrcxeZdZ d dfd Zed d dfd Zd dZd dZxZS) TritonTemplateBufferct|||||||jJ|j|ji|_d|_d|_y)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)rnr]rr subgraph_inps subgraph_outs)rrNrjrrrrps rr]zTritonTemplateBuffer.__init__s]"    )"7  yy$$$)-DII(>%@D9=rct||}|jr |jng}|jr |jng}|D]c}t |t r|j t||/t |tr!|j |j|`|cJ|D]7}t |tr!|j |j|4|7J|Sr) rnr>rrrr updater)r)rriresrrrlrrps rr>z)TritonTemplateBuffer.get_free_symbol_usessg*=9.2.@.@**b .2.@.@**b   #C#t$ +C?@C( 333MBC{"{  #! #C#v& 333MBC{"{  #  rc |g|jSrrrs rrz TritonTemplateBuffer.get_outputs3-t,,--rc&d|jd}|S)NzTritonTemplateBuffer(layout=rr)rrs rrzTritonTemplateBuffer.__str__6s,T[[M; rr) rNrrjrrzCallable[_P, _T] | Nonerrrrrrrrr}rro) rrrr]rbr>rrrIrJs@rrrs 378< >>!>4 > 0 > 6 > ><23$)! !4..rrceZdZdZ d fd Zd dZddZddZddZddZ ddZ dd Z dd Z dd Z 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. ct|||_||_||_||_d|_d|_i|_d|_ d|_ i|_ i|_ yr) rnr]rrNr descriptionfailed_benchmark_with_cudagraphsrrd decompositiondecomposition_kwargsconfig_patches)rrrrNrrps rr]zChoiceCaller.__init__Gsh   &'! 05',./38<46!.0rc|j|jrtjfdStj rt fdStjd|idS)NcSrralgorsrrz(ChoiceCaller.benchmark..es T4[rcSrrrsrrz(ChoiceCaller.benchmark..gs D$Krrr) to_callablerr\benchmark_gpu_with_cuda_graphrJ/profile_bandwidth_with_do_bench_using_profilingrg benchmark)rrrrs `@rrzChoiceCaller.benchmarkbs\!  * *<<=PQ Q  A A+,?@ @$$T4%dKKrctrrrs r call_namezChoiceCaller.call_namejrrctrrrs rrzChoiceCaller.to_callablemrrc"|jS)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 rr zChoiceCaller.hash_keywrrctrrrs rrezChoiceCaller.output_nodezrrciS)zRInformation returned here is logged to the autotune log file when that is enabled.rrs r info_dictzChoiceCaller.info_dict}rrcy)Nunsupported_choicerrs rautoheuristic_idzChoiceCaller.autoheuristic_ids#rcd|_y)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) rrrrrNrrrrr)rrrrtrr ro)rr)rr)rz6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]rS)rrrrr]rrrr r rerrrrIrJs@rrr>sk11"1 1  1  16L""""$rrceZdZddZy)TritonTemplateCallerBasectrrrs rget_make_kernel_renderz/TritonTemplateCallerBase.get_make_kernel_renderrrN)rr)rrrrrrrrrs"rrceZdZdZ d fd Zed dZed dZ d ddZe jddZ ddZ d ddZ dd 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. ct|||d|||_i|_||_||_t d|D|_i|_y)N)rNrjrrc3K|]R}t|txs<t|tjjj xr |j Tywr)rrrrselect_algorithmExternKernelCallerhas_out_variant)rchoices rrz/MultiTemplateBuffer.__init__..sT%  v7 8 65??#C#C#V#VW+** % sAA) rnr]_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders)rrNrjchoice_timings_fnunfiltered_choicesrrps rr]zMultiTemplateBuffer.__init__sj #"7  #4LN,> %!$% - % " <>!rc|jS)z^ Are all possible choices TritonTemplates or Extern Kernels with out variants )r$rs routput_plannablez$MultiTemplateBuffer.output_plannables %%%rc|jSr)r"rs rrzMultiTemplateBuffer.choicess }}rcx||jvr|j||j|<|j|Sr)r!r )r hint_overrides rchoice_timingsz"MultiTemplateBuffer.choice_timingss>  4 4 4262I2I-2XD  /##M22rc#0Kt|tjjjsJt ||j |j k(sJ|j}|j|_ d||_y#||_wxYwwr) rrrrTritonTemplateCallerrrNrr)rcallerrenders rswap_as_triton_callerz)MultiTemplateBuffer.swap_as_triton_callers EOO44II   <  {{fmm+++(("("?"?"A - &,D #fD #sA;B>B B BBcNt|tjjjsJt ||j |jjk(sJ|j|jjk(sJ|j|_ yr) rrrrr/rr rNrr4rrr)rr0s rfinalize_as_triton_callerz-MultiTemplateBuffer.finalize_as_triton_callers EOO44II   <  }}&--"4"4444 FMM$8$8888"("?"?"Arcb|j|}t||j}|||fS)N)r,r)r-r1r)rr,timings min_choices rget_min_choicez"MultiTemplateBuffer.get_min_choices7%%M%Bgkk2 GJ/00rc|jD]"\}}|j|j|<$|jd|_y)z;Finalize with multiple callers for different hint overridesN)rTrr%r)rcallersr,r0s rfinalize_as_triton_callersz.MultiTemplateBuffer.finalize_as_triton_callerssN&-]]_ W !M67=7T7T7VD % %m 4 W#'";";D"Ar) rNrrjrr&z1Callable[[int | None], dict[ChoiceCaller, float]]r'list[ChoiceCaller]rrUrrr_)rr<r)r,rrzdict[ChoiceCaller, float])r0rrr)r0rrr)r,rrztuple[ChoiceCaller, float])r:z*dict[int | None, TritonTemplateCallerBase]rr)rrrrr]rr)rr-rrr2r4r8r;rIrJs@rrrs>>!>M > / > / > >8&& +/3'3 "3 - -B+/1'1 #1BAB BrrcLeZdZ dfd ZddZddZxZS)CUTLASSTemplateBuffercRt||||||_||_||_yr)rnr]workspace_sizetemplatesupports_epilogue_fusion)rrNrjrr@rArBrps rr]zCUTLASSTemplateBuffer.__init__s. );<,  (@%rc6|j |jSdSrr@rs rrz(CUTLASSTemplateBuffer.get_workspace_sizes&*&9&9&Et""L1Lrcx|jD]'}tj|jdd)yr)rrurrs)rr\s remulate_store_fnz&CUTLASSTemplateBuffer.emulate_store_fns1&&( 5F IIfoo't 4 5r)rNrrjrrCallable[_P, _T]r@rrAr{rBrrrr{rS)rrrr]rrFrIrJs@rr>r>s_ A A! A- A  A " A#' A  AM5rr>cDeZdZ dfd Zdfd ZxZS)CppTemplateBuffercRt||||||_||_d|_yr)rnr]rAroutputs)rrNrjrrArrps rr]zCppTemplateBuffer.__init__ s, );<   ,0 rcvt|jtrt|jtsJt |j|jd}t|t sJt ||j}t|tsJt ||St|%Sr) rrNrrKrrrrrrnr )r first_outputrNrps rr zCppTemplateBuffer.get_layouts dkk#4 5dllH5 ItDLL7I I5<<?LlF3 GT,5G G3!((Fff- ;tF| ;-M7%' 'r) rNrrjrrrGrAr{rrrrr])rrrr]r rIrJs@rrIrI sL 1 1! 1- 1 " 1  1  1 ( (rrIcFeZdZdZ d dfd ZddZxZS)CuteDSLTemplateBufferz Buffer for CuteDSL (CUTLASS Python DSL) template kernels. Similar to other template buffers but specialized for CuteDSL operations. c t||||||_||_|g|_|t |j dtsJt|j d|j dj}|xj|Dcgc]}tt|||c}z c_yycc}w)Nrr) rnr]rArrKrrjrrrrr) rrNrjrrArrrrps rr]zCuteDSLTemplateBuffer.__init__*s );<  ,&*V  %dkk!nf5 KtDKKN7K K5[[^..0F LL)z8#tD L &sB<c|jSrrKrs rrz!CuteDSLTemplateBuffer.get_outputs? ||rr) rNrrjrrrGrArrrrrr)rrrrr]rrIrJs@rrOrO$sQ37 !-    0  *rrOceZdZdZ d dfd ZddZd 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. c 0t |||d||_||_|g|_||_||_||_||_| |_ | |_ |jj|jjd|_|j|_y)N)r) kernel_namemin_cc)rnr]kernelaccumulator_typerKr@variant scale_type_a scale_type_bswizzle_type_aswizzle_type_bmetadatarWrXkernel_metadata_make_kernel_renderr) rrNrjrYrZr[r@r\r]r^r_rps rr]zNVUniversalGemmBuffer.__init__Ks DA 0&*V , ((,,"??66oo,,  #'":":rc|jS)z#Return the workspace size in bytes.rDrs rrz(NVUniversalGemmBuffer.get_workspace_sizels"""rc|jSrrRrs rrz!NVUniversalGemmBuffer.get_outputsprSrc  ddlm}ddlm}g}|jD]K}t |t r |j}t |tr |j}|j|Mt|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) rWrrerarZr@r[r\r]r^r_c$jSr)r1) render_kernelsrr1z9NVUniversalGemmBuffer._make_kernel_render..renders '') )r)Btorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelrftorch._inductor.utilsrgrjrrrMrqrr KERNEL_NAMErarZr@r[r\r]r^r_) rout_noder,rfrgrrlrWr1ris @rrbz)NVUniversalGemmBuffer._make_kernel_renderss  6!# ;; $C#y)hh#z*hh   s #  $+112 -##  00!22..LL****....   *f$$r)rNNNN)rNrrjrrYrrZrr[rr@rr\ Any | Noner]rnr^rnr_rnrrr{rr)rmrr,rrztuple[Any, Any]) rrrrr]rrrbrIrJs@rrUrUCs #'#'%)%);;!; ;  ;  ;;!;!;#;#; ;B#:>*%*%,6*% *%rrUc&td|DS)Nc3<K|]}t|tywrrr)rrs rrz#is_node_sequence..s4z!V$4r|)r)rs ris_node_sequencerrs 4e4 44rceZdZUded<d dZd dZddZeddZe ddZ ddZ dd Z e d d dd Zy )r#Sequence[IRNode | Sequence[IRNode]]rjcd|j|}t|tsJ|jSrrjrrrs)rrinputs r input_namezInputsKernel.input_names, A%(((~~rcttj}tj|jD]c}t |t r|jfd|D.t |tr?|j|jettjfd|jD}tj||tS)Nc3JK|]}|jywrr)rrrIs rrz/InputsKernel.get_read_writes..sBqWQZZ\2B #c3JK|]}|jywrr)rrrIs rrz/InputsKernel.get_read_writes..s!. (+GCLLN #. r{r<) r@rKrQrIrjrrrrrrsrr?)rr6rwr=rIs @rr3zInputsKernel.get_read_writess<++,.&&[[ 5E%* BEBBE#89 '%.."234 5L,,-. /3/?/?/A.  &&"   rc6|jjSrr5rs rrzInputsKernel.get_readsr7rct|tr |j}t|tr |j}t|tr%t|t st j|}t|tr|j|St|tr|St|tt fsJt||Sr) rrrMrqrorSrrunwrap_storage_for_inputTorchBindObjectrrrrrs rrz%InputsKernel.unwrap_storage_for_inputs a #A a $A a ":a+I**1-A a # //2 2 a )H!fo67@a@7rcg}|D][}t|tr#|Dcgc]}tj|}}ntj|}|j |]|Scc}wr)rrrrr)rj inputs_newrrs rrzInputsKernel.unwrap_storagesl79  !A!X&GHI!\::1=II 99!<   a  !  JsA%cyr1rrs rrDzInputsKernel.is_externrrcyrrrs rr9zInputsKernel.num_readsrrcttj}|jD]B}t |t r||j |z}(|D]}||j |z}D|Sr)r@rr"rjrrr>)rrirrl inner_inps rr>z!InputsKernel.get_free_symbol_usessv u|| $ &;; GC#v&S--m<<!$GI77 FFAG  G rN)rrrrrxrzr)rjrtrzlist[IRNode | Sequence[IRNode]]r_r{rrr})rrrrrxr3rrrrrrDr9rbr>rrrrrs //  ,,$ 3  (  N+$) !  ! , rrceZdZddZddZy) NopKernelcyr1rrs rrFzNopKernel.is_no_oprrctSrr?rs rrzNopKernel.get_readsrrNr_rz)rrrrFrrrrrrs rrceZdZdZeddZe d d dZed d d dZed dZ ddZ y) ConcatKernelzn There isn't actually a real kernel for concat, we just change the storage for the upstream data. c Z |dj}|dj}t|dj}dg}||g}d|cxkrt |ksJJt dt |D]}||j} |j ||t | t |k(sJ||j|k(sJ||j|k(sJt t |D]I} | |k(r|| | | z|| <tjjj|| | | || <K|j ||tj|} tjr$tj!| ||dj"} t t |D]k}||} t%| s| j'} t)| t*s5tj-| j.| j0s`t3|} nt5d|D}tjj6j8d}|dur-t)|trt5d|Dr t3|} t;d|D}|Jt=dt+|||| |g }t?|}g}tA|D]\}}t)|tBtDfsJtG||jI|tJjM||||||d }t)|tNsJtG|t)|jPtsJtG|jP|jPj |t)|jRtBr|jRjU}n |jR}t)|t>s|jWs1|jx}EtY|jFs\t[|ri|j |j]t |dkDrMtjj_|t`jbrtjje|tjjg||_4|jk|jP|_(tjjm||S) z6 Create the concat kernel from inputs rrIc32K|]}t|ywr)r rs rrz&ConcatKernel.create..6s-W1.CA.F-WrFc3K|]f}d|jvxrRt|jdtjxs(t|jdtjhyw)rrN)rbr/rrrrargs rrz&ConcatKernel.create..=sm! <&+&9&9  @&+&<&< sA,A.c3jK|]+}t|xr|jj-ywr)r r rRrs rrz&ConcatKernel.create..Qs/ FG !! $ A)A)A A s13N)rrrrrRrrNrjr)7rrrr rrrrwrrrr rrJrirrrr r rrPr>rrr1r0 current_noderrrrqrrorrrrrrrrjrMrprrnrmrLrrLFOREACHregister_operation_listrrrr)rrjrBrrrG offsets_start offsets_endr input_sizer output_striderrNany_input_is_storage_and_layout fx_node_argsrR concat_kernelrYop_namesrlrinput_unwrappeddevs rrzConcatKernel.create s %%'q ##%q **,- }o C'#h-'''''q#f+& .A++-J  # /z?c(m3 33!9&&(E1 11!9'')V3 333x=) 8"*1+ 1 "=HQK"#''"2"2"L"L  Z]#HQK     x} - .(6'H'H'R  ' '"//xM s6{# Aq A$Q'K88fmmT$B8$LM +.-WPV-W*W'ww++003  ,u 4<.($;8DM KQ  !!!$$#   M*' CFAscHj#9: EDI E:++  Cq!1;q>!L lF3 GT,5G G3m22D9 U4 @T@T;U U9  ' ' 5#((H-"%(("6"6"8"%((?J7#335NN,,S9388$"<0 ? ? AB1 C4 x=1 !4!4V^=S=S!T GG + +H 5WW44]C "11-2F2FG  ""=1 rNct|tr|j|j|St|tt fsJt |t|jtrt|jjtr|jjsy|yt|jt|jk7rytdt|j|jDSt|jdxrAt|jjt xrt|jt" S)NFTc3vK|]1\}}tjjj||3ywrrHrJs rrz=ConcatKernel.can_realize_into_without_copy..s3B  88R@rKrN)rrcan_realize_into_without_copyrMrorqrrrNrPr)rr4rrrr ExternKernelAlloc)rrrs rrz*ConcatKernel.can_realize_into_without_copys c9 %44SXXsC C#*56AS A6 chh 3 4sxx <xx00{3>>#$CNN,<(==!#.."2CNN4DE  CHHh ' <388??N; <sxx):;; rc.tj||Sr)rr>r=s rr>z!ConcatKernel.get_free_symbol_usess--dMBBrc ^t|ts&t|rt|\}}t||}t|tsJt |t|t r|j |j|St|tr`|jt|jdsJ|j||r&t||j_ |jStj|j!|j#|j%t'|j)|j)Dcgc]/\}}t*j,j.j1||1c}}}|j ||Scc}}w)NrLrNr )rrSr rOrrrrMrqrrrrrNr rrrrrr rwrrr)rrrrVrNrFrGpws rrzConcatKernel.realize_intosL #/$S)"7"<%76B#/:c:/ c9 %##CHHc2 2 c: & KKM388X. ..00c:"1#"6xx   >>#--/__&   ?Aq  ::1a@  C(( s4F)cyr1rrs rrzConcatKernel.should_allocaterr)rjrrBrrrqr)rrrr?rrrrr})rrrrrrr_) rrrrrrrrbr>rrrrrrrs yyv/3! ! ,! ! ! FN+$)C!C !C,C ))Brrc#K|jd} d|j|y#|j|wxYwwr)"_ignore_fresh_unbacked_symbols_set)rprevs r_track_fresh_unbacked_symbolsrs:  7 7 >D; 44T: 44T:sA+A>AcP||jr t|StSr)"_ignore_fresh_unbacked_symbols_tlsrr )rs r(_fallback_kernel_symbol_tracking_contextrs(!M!M!O,Y77 =rceZdZUdZdZded<ejeZ ded<dZ d ed <dZ d ed <dZ d ed <eje Zded<dZded<dZded<ejeZded<dZded<ejeZded<eje Zded< dB dCfd ZdDdZdEdZdFfd ZdGd ZdGd!Z dH dId"ZdJd#ZdHdKd$ZdLd%ZdMd&ZdNd'Z e!dOd(Z"e# dPd)Z$e#dQd*Z%e#dRd+Z&e#dRd,Z'e# dS dTd-Z(e# dU dVd.Z)e# dU dWd/Z*e#dRd0Z+e#dRd1Z,e#dRd2Z-e#dRd3Z.dGd4Z/ dXd5Z0dHdYd6Z1dZd7Z2d[d8Z3dUd\d9Z4dNd:Z5dJd;Z6dJd<Z7dJd=Z8d]d>Z9d^d?Z:e;d dU d_d@ZxZ?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. rr constant_argsrrrNReinterpretView | None output_viewrepython_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]rc Dt ||||||_|r|ni|_||_| |_|j ||j|| |_|ji|_ g|_ tjj|_i|_yNr)rnr]rrrrset_cpp_kernel_nameset_python_kernel_namercollect_arg_kwarg_propertiesrrrwrrfx_noder) rrrNrjrrrrrrrrps rr]zExternKernel.__init__s   + &fB &&   1 ##$67-J* ))+!# "ww++ +-rc |g|jSrrrs rrzExternKernel.get_outputs#rrctSrr?rs rrz%ExternKernel.get_unbacked_symbol_defs&rrct|dfd }tj||j|j fdS)Nct|trB|j}|/jj t j |yyyr)rrrr6rrKrI)rrrs r add_ir_readz1ExternKernel.get_read_writes..add_ir_read,sJ%(++-#%%)),*>*>t*DE$)rc"t|tSrrqrs rrz.ExternKernel.get_read_writes..5s *UF";r)is_leaf)rrrr)rnr3pytree tree_map_rr)rrrrps @rr3zExternKernel.get_read_writes)sFg-/  F      -; rcNt|jtjjr\|jj j Dcgc]2}|js$|j|j|jd4c}n+tt|jDcgc]}ic}|_t|jtjjrP|jj j Dcic]&}|j|j|jd(c}ni|_t|jtjjr|j sJ|jj j Dcgc]}|js|jc}|_|jj j Dcgc]}|js|c}|_yg|_ycc}wcc}wcc}wcc}wcc}w)N)rrr )rr )rrr_ops OpOverload_schema arguments kwarg_onlyr real_typer rrrjrrr schema_kwargs)rrrs rrz)ExternKernel.collect_arg_kwarg_properties9s$**EJJ,A,AB))11;;   || FFKK%&__ $C $456"6 $$**EJJ,A,AB))11;; qOO    d&& (=(= >55$($4$4$<$<$F$F6 !,,AFF62 ++33==""D "$D ? 7 6"s*7H+ H +HH/H'H"9H"czt|jtr!|j|j yyr)rrNr apply_constraintr"rs rrzExternKernel.decide_layout^s- dkk> 2  ! ! #     3rct||\}}|r|j||s|j}|r$ddlm}|||d}|j ||yy)NrI)'set_kernel_post_grad_provenance_tracingT)rD)rj make_commenttry_get_kernel_namerrwrite_provenance_debug_handle)rwrapperrW origin_str_detailed_origin_strr debug_handles rcodegen_commentzExternKernel.codegen_commentcsf,?tW+M( (    ,224K  FBkTL  1 1+| L rctrrrrs rcodegenzExternKernel.codegentrrc^||_tjjr.t |j t jjsy|j }|j|jdk(r|jdk(r|jjddn|jjdd}ddlm}|j!d|ji}t#d|Dd }|d kDr|d |}d |d |_y|j$j&|_yy)Natenr.rrinductor_fallback_opszaten.c3\K|]$}|jdst|dd&yw)rOrIN) startswithr)rrOs rrz3ExternKernel.set_cpp_kernel_name..s%KAc9JS12ZKs,,rIr_vz at::_ops::z::call)rrwr cpp_wrapperrrrrr namespace _overloadnamerrreplacetorchgen.aoti.fallback_opsrrrrr)rrrYopnamer version_infolatest_versions rrz ExternKernel.set_cpp_kernel_namews.ww""*   ejj33+  !!    '6)++y8OO))#.q100c:M4885@Q9RTVW !$KK""A% &xr.)9:F)3F86'B$'-~~':':$7 (rc||_|y|j}|yt|tjj rd|j |_y|jjddd|j |_y)Nztorch.ops.higher_order.._ops..ops.r) rrrrrHigherOrderOperatorrrr)rrrYs rrz#ExternKernel.set_python_kernel_names"4  ) !! >   > > ?(??P&QD #$$,,Xw?@&//ARS  #rc0ddlm}|jx}r |jntj j }tj jr |jStj jrttj j|s(Jttj j|jytj jj|j|S|jS)NrI) CppWrapperCpu)codegen.cpp_wrapper_cpurrrrwr device_type fx_wrapperrrrrrget_c_shim_func_name)rrdrs rrz ExternKernel.try_get_kernel_names:!%!22A29L9L 77  ** * WW agg22MB D$$E B##+77''<<$$f ** *rc.|j}|J|Sr)rrs rget_kernel_namezExternKernel.get_kernel_names!'') rc tj|j|j|j |j |j |j}|j|S)N)rrrrr~r|) r rrrrr rrr)rrs r copy_inputzExternKernel.copy_inputsa   <<>++-]]_::<))+oo'    rc ||d}tj|\}gg}g}g}|D]} | xtdxl\tjj j j| d} jd|j| |j| {xtdx\jd|j| | jj} | jjdk(r| J|jtjj| j!xt"dxB\jd|j| |j| j$gt&dx&\jd|j|  jd|j| |j|  dfd } |D cgc]} |j)| }} |D]} t+| st-| d g}|D]} t/| t0se| j3tjj4vr;|jtjj4| j3yt/| t0se| j3tjj6vr;|jtjj6| j3t/| t8r!|j| j;t/| t<r|j| j>Lt/| tj@jBjro| jj} | jjdk(r| J|jtjj| j!|jtE| | ||\}}tjFj }tI5tK|5||i|}ddddddd}|rtjLjNjQd }tS}tjLjTtjVjXjZur|d }t]tjL}|5t_|tjLdddta||}t/tbtdfs|gn|}|D]}t/|tjfs|jhs+tjjlr.unflatten_argssFo.J!"56N0 8 MM$z"23MM$~"67  8 %%fi8A55$aeeHb&99 9rrrrIzEsparsity not handled. Please file issue for sparse inference weights.rz Found from : )r rr rrztuple[list[_T], dict[str, _T]])8r tree_flattenr rwrrrcreate_symintnoderGeneratorStaterrrrr"default_generators clone_stateOpaqueObjectStaterrrr rOrrors constantstorchbind_constantsr get_valueropaque_example_valuerirrr6r#rrrbrr rT_higher_order_opseffects with_effectsr3r:r4rrr is_sparserJgraph_partitiondisable_cudagraphs_reason)rrYrr binded_args args_flat tensor_argsnon_tensor_argsreal_non_tensor_argsrr device_indexr r example_argsnew_args new_kwargsrexample_outputr node_meta_valctxexample_out_lirmsgrr r s @@rprocess_kernelzExternKernel.process_kernelsV" $v6 %22;? 9*,$& (*   5CTV77++55GGRVGWD'..u5#**40(//5  &^%'..u5#**3/#&::#3#3L::??f49QQQ(// 55lCOOQ &)&('..u5#**3/(// :) X'..t4&&s+'..u5#**3/(//4; 5> :) :@L : + :6AAs((+A A 6A$Q'%a5 6   :Aa*qzz|qww?P?P/P##AGG$5$5ajjl$CDq(+JJLAGG$?$??##AGG$?$? $MNA/##AKKM2A01##A$:$:;Au11@@A xx~~ xx}}.<3KKK##JJ11,?KKM##$5a$89+ :. .l>>J K 9>=!  ntUm<    8A1ell+KK..]"#''"6"6";";"?"? t"TT;T E!2;-@C471 8        cBV = = = = K Ks6 X> X1 X$X18X>$X. )X11X;>Yc *t|tsJt|t|tr|S|j }t j j|j}|J|j}|d|jvrt|tttfrt|jtryt|jdt j"s(t|jdt j$r)|j't)|j+n|j-t/j0|j+d\}}|d}|j3|}t j j4j7||}t j j4j9||} t j j4j;||} t=|| | z} || k7rt>jAd| | |tBt|jDtG|jI|jK|j+| | 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. rrrrrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFrrL)&rrorrSrprwrrrsrrbrrrrNr r/rrrr-r1r r"rKrr rr stride_vars offset_varrorrrrMrPrr) rr x_unwrap_viewrx_unwrap_view_fx_node index_argsrUrrrUrQexpecteds rconvert_to_reinterpret_viewz(ExternKernel.convert_to_reinterpret_viewus$!X&/Q/& a )H  gg  !7!7!9: # 3 3 5 " -.333=?FJ*OP=//@8)..u5"'"5"5<)..u5"'"8"8  7 7.}/E/E/GH   ' ' )!-!@!@ JJL"  J ]   ,  55eZH''""..ujA!!,,UJ?Z1F: H  IIR   & %,,.kkmZZ\   rc | tSt|ttjj j tfr t|St|trot5tjjtj|j |j#|j%cdddSt|t&r|St|t(r|j+|j,St|t.r4t/|j+|j,|j1St|t2r;|j5t7|j9r |j;|St|t>r|j5|St|t@ttBfr|S|jE|S#1swY&xYw#t<$rYhwxYw)N)r)rrrL)#rrr rrrrrrrrArwradd_tensor_constantrrTrrrr!rrrMrSr rorr rpr8 rrq NonTensorObjrr rs rrzExternKernel.realize_inputs 9') ) a$ 3 3 ; ;SA B(a0 0 a "() ww22LL allnU   a (H a #$$QVV, , a )"&&qvv.q||~  a " IIK$Q]]_5::1== a $ IIKH a,(=?PQ RH~~a  3  "+s,AG- G:-G7: HHct|r"a',,. Q;H ~~a  rc( ||J|jdvr|s|St|r-t|jtrG|r2t ||}|r|j |St|dd|||St|ddd|||Sd}|r>|r} } | D].}tRjTjVjY||d d }0|j[|}t|dd||||rBt]|js)t]|j_st ||sJ|S| r<| |JtRjTjVja|| }t%||S|S#tB$rYzwxYwcc} w) zUEnsure x has the requested stride order or exact strides, inserting a copy if needed.N)rrITF)rrrr%rzHthe MutationLayoutSHOULDREMOVE's real layout shouldn't be FlexibleLayoutrMrrrrI)1rr rr r rr"rOrrrr rrPrrrErrZrrrlrrrMrorSrprrr8 require_stride_orderrrrwrrrrrIrrrloweringslice_r r6r4r)rrrr0r%use_current_stride_orderpadded_exact_stridesmutation_layoutrrp expanded_dims orig_sizerrrBs rrequire_strideszExternKernel.require_stridess M$=== ;;=F "=H ! #!,,..90R50,0)H.#',1).*7 H*#(-%)&3&3 H $( '+'' qzz|Q[[]S($!,,.;*HI1<<>;;EB!1%q||~'<'-!1<<>#8#8!**, H q) $1668,qvv7%Q]]_&DkE V,{//1BC 88@335 4#44= 5#+/ JJL  $ww''Hs1::<0133M!4DaH11!**,q/1EM% BOO,,33AsAqA B NN1  !''   )!**,7)!,,.95a? ?? (]-F FF((//9=A21mD D_'  s5S?:S?#AT? T  T c |j||Dcgc]4}t|tjr|jj n|6c}|Scc}w)N)r0r%)rG rrSymIntrr)rrr0r%rs rrz"ExternKernel.require_exact_stridessV"" KXFGz!U\\: A( #  s9A c*|j|||S)N)rr%)rG )rrrr%s rr? z!ExternKernel.require_stride_orders""1E"OOrc.|j|tSr)r? rrs rrequire_channels_lastz"ExternKernel.require_channels_lasts''+<==rc.|j|tSr)r? rrs rrequire_channels_last_3dz%ExternKernel.require_channels_last_3ds''+=>>rcdd}||r|S|j|tj|jS)Nc |j}|tjj vxr'tjj |j S#ttf$rYywxYwr)rsAttributeErrorrrwrr is_mkldnn)rrs ris_mkldnn_tensorz9ExternKernel.require_contiguous..is_mkldnn_tensors] zz|177,,,R1B1B41H1R1R R#$78  sAA)(A)rrrrrr rr )rrrS s rrzExternKernel.require_contiguouss@ S A H,,>44QZZ\B rch|j|tj|jSrrU rs rrequire_contiguous_stridesz'ExternKernel.require_contiguous_stridess-(( ~00>  rcyrrrs rrzExternKernel.apply_constraintrrct|tsJt|t|ts t|}|jsJdt |}t |j}||krqt jd|j||z t||D]>}|j|d}|j||vr||n|j|d@|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.rr ) rrrrrrrrrrr)rrrn_args n_pos_argsrarg_names rfill_non_provided_argsz#ExternKernel.fill_non_provided_argss$)54:5)$%:D""U$UU"T,,-  J  II^  V#   6:. ..q1&9 6)8$,,Q/@  rc\tjjrGg}d}|r]|jrQt |j t |k(sJd|jDcic]}|j d|}}t|j D]\}}|.|J|j ||}|r|j dnd}n\t |j|z} |jr6| t |jkr|j| j dnd}|jtjjj|||S|j D cgc]+} tjjj| -c} Scc}wcc} w)NzDnames passed to codegen_const_args does not match self.constant_argsrr) rwrrrrrrrrjrrval_to_arg_str) rrrNname_to_arg_propertiesrrrproptype_rrFs rcodegen_const_argszExternKernel.codegen_const_argss 77  F &* ",,4--.#e*<Z<594G4G*-0CGGFOS(*&*"$"4"45 M1)5 ,,,155eAh?D04DHHV,$Edkk*Q.C ..3T=P=P9Q3Q++C044V<!  agg22AA!UKL MMDHDVDVWqAGG((77:W W'*&Xs $F$10F)ctjjrC|j7|j g|j |j |j}d}n|j }d}g}t|D]\}}tjjr|jr|t|jksJd|j|jd}|jtjjj|||jtjjj||r|j|j!|S)NFTz-Invalid access to ExternKernel.arg_propertiesr)rwrrrr] rjrrrrrrrrr_ rnrc )rrjneed_codegen_constant_argsrrrrb s r codegen_argszExternKernel.codegen_argss5 77  4#3#3#?003$++3 2 23T[[F*/ &[[F)- &f% DDAqww""**q3t7J7J3K/KCK++A.226: AGG00??5IJ AGG00??BC D & KK//1 2 rc ||vr|j|S||jvr|jj|S|jj|x}|jdSt|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.r z not in self.allarg_properties)rrrrl)rr\ rrs rget_kwargs_valuezExternKernel.get_kwargs_value8sx v ::h' ' t{{ ";;??8, ,))--h7 7C D77?+ +z)GHIIrc tjjr|jt |j dk(rgSg}|j D]}|r|dk(r |j|}t|tr|j|>|jJ|jj|ijd}|jtjjj|||S|jj!Dcgc]3\}}|dtjjj|5}}}|Scc}}w)Nrrrr)rwrrrrrrrh rr rrrrr_ rrT)rskip_outrr\ rOrb ks rcodegen_kwargszExternKernel.codegen_kwargsEs= 77  +D4F4F0G10L F >> QE 1))(3a&MM!$11=== 2266xDHHPEMM!''"6"6"E"Ea"OP Q" !KK--/Aq#Qqww++::1=>?F s 8Ec|jS|jj}t|dd}|jdd}|j ddd}|d|}|Sd}|S) Nrunknown_namespacerrrrIr unknown_op)rrTrrrsplit)rrT op_namespaceop_names r get_op_namezExternKernel.get_op_name_sv << #\\((F"6<9LML'//'BL'..sA6q9L%ax0G#Grc tjsyt|jdk(rytj j j|j}tj j j|j}|j}|j}tj jrt|jtjj rltj"j$|jj&vr224D((tVWErc Htjrtjjsw|j }|tjj v}|j}|r |jd|dtd|dy|jd|d|dyyy)Nzassert_alignment(rErz # buffer z (op: z) is assumed to be not aligned) rJalignment_assertsrwrrrsr rs rrk)rrralignedrr s rcodegen_alignment_assertsz&ExternKernel.codegen_alignment_asserts}s  # #AGG,?,?==?D!''";";;G&&(G!!'vR/@7+QO!!vVG94RS-@ #rctjjrtjj ry|j |j}|jd|d|dy)zc Track outputs of fallback operators if config.test_configs.track_memory_lifecycle Nz track_tensor(z, 'z')) rJ test_configstrack_memory_lifecyclerwrr"write_memory_track_allocation_oncersr)rrrs rcodegen_memory_trackingz$ExternKernel.codegen_memory_trackingsV""99QWW=P=P 224}}M$s4&; F?$ G6 G +Gc|rtnt}tj||}|jD] }|||z}|j j D] }|||z}|Sr)maybe_free_unbacked_symbolsmaybe_free_symbolsrr>rrr)rrimaybe_get_symbolsrrs rr>z!ExternKernel.get_free_symbol_usess|,9 '>P   - -dM B%% (C "3' 'A (;;%%' (C "3' 'A (rc "t|dd}d|g}|tj|Dcgc]'}|jdt||j)c}z }|j d|j |j |Scc}w)Nrzpython_kernel_name=rr)rrfieldsrrr~r)rrWrrs rrzExternKernel.__str__sd$8$? !+ 1  $++D1 zzl!GD%**56 7    |D$4$4#789u%%  s,B rNNNNrN)rrerNrrjrtrrrdict[str, Any] | NonerrrrerrerrrrrrrrrxrSr)rr|rWrerrrr|rrrrerr)rrerrrdro)rrrr)rYrrrrrrzftuple[Any, list[Any], list[Any], Callable[[Any, Any], Any], dict[sympy.Symbol, pytree.KeyPath] | None])rrrrSr)NNF) rrrzSequence[int] | Noner0rbr%rrrrr)rrr0rar%rrr)rrrrtr%rrr)rrrrrr)rr{rrrr)r\ rrrrr)rj rrr)rz'tuple[list[Sequence[Expr]], list[Expr]])rztuple[Expr, Sequence[Expr]]r})@rrrrrrrrrrrrrrrrrrrrrr]rrr3rrrrrrrr rr rr0 r8 rr= rG rr? rL rN rrW rr] rc rf rh rl rs r{ r r r r rbr>rrrIrJs@rrrs  $&M=%.[..tDFND*.K'.%) )"&OZ&4E;3D3D4!=(,K$+26N/63D;3D3D40:>6=->t-T*T(*(,.2)-&*79+/...4 . % . & .,.'.$.(5.). .@. #$J! HLM+M:DM M""#;J +$   c !c *-c 9<c  c c JE E N!!!!F!!'+37# t t$t1 t  t  ttlQV    '9  JN      DIPP,P=AP PP >>??"   "!"+9" "HXB4 J4 F&  =$'BN+$) !  ! ,  &HrrcfeZdZddZ d dfd ZddZxZS)ExternKernelOutc&|j|yr)generate_extern_kernel_outrs rrzExternKernelOut.codegen**40rc &|j|} t| tsJt| t |d|| ||xsid|||| t jj||_ t jj|yr) rrrrrnr]rwrrrr) rrNrjrrrrrrrunwrapped_inputsrps rr]zExternKernelOut.__init__s ..v6*H5Mt ,,  r)rrrrjrrrrrr]rIrJs@rr r s   rr cheZdZddZ d dfd ZddZd dZxZS) rc&|j|yr)generate_extern_kernel_allocrs rrzExternKernelAlloc.codegens,,T2rc F|j|} td| DsJt | d|t t t | ||xsid|||| g|_tjj||_ tjj|y)Nc3<K|]}t|tywrrqr{s rrz-ExternKernelAlloc.__init__..)sCQ:a(Cr|) rrrnr]rrrrKrwrrrr) rrNrjrrrrrrr rps rr]zExternKernelAlloc.__init__s ..v6C2BCCCC    &!#3 4  Lb    )  ') GG++D1  ""4(rcyrrrs rrz!ExternKernelAlloc.should_allocate=rrctrrrs rrz"ExternKernelAlloc.apply_constraint@rrr )rNNNrN)rNrrjrrrrr rrerrerrrrrrr_rS)rrrrr]rrrIrJs@rrrs3(*(,)-&*79+/))!)% ) & ) ' )$)(5))) )@"rrcTeZdZdZ dfd ZddZd dZd dZd dZxZ S) rzP An output buffer that represents the mutation of a pre-existing buffer ct|d||j}tjj ||g|_||_tjj||_ yr) rnr]rsrwrrmutation_names mutating_noderr)rrN mutated_noder mutated_node_namerps rr]zMutationOutput.__init__Isb d62(113 ##$5601(5GG++D1 rc|jSr)r rs rrzMutationOutput.get_defining_opSs!!!rc|jSr)r rs rrJz!MutationOutput.get_mutation_namesVrsrcyrrrs rrzMutationOutput.should_allocateYrrc^|j}d|DDcgc]}||c}Scc}w)Nc3ZK|]#}tjj|%ywr)rwrtry_get_buffer)rrs rrz6MutationOutput.get_mutation_buffers..`sP..t4Ps)+)rJ)rr rs rget_mutation_buffersz#MutationOutput.get_mutation_buffers\s=002QP     s *)rNrr rr rrrr rr_rr) rrrrr]rrJrr rIrJs@rrrDsA2 2062GP2 2"# rrceZdZUdZiZded<e d dZe d dZ d fd Z d dZ d dZ 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]_CACHEct|dk(sJ|ddk(rt|g|dS|ddk(sJt|g|dS)Nrr experimentalrIr)rTMADescriptorExperimentalTMADescriptorStable)rrTtma_metas r _create_implzTMADescriptor._create_implss\8}!!! A;. (,VBhqkB BA;(* **&v< < DMM & &w|| 4 MM + + 2 27<< @ A8BBG7<<B88   / A!&%*.4L .X -WAUSU-WN &&~6 8! !!4#6#67 8 --22AB)).*=*=*?@A A A C 8 8 A As+G3 "H$G8&AH8H =HH cGdd}t|jsJ|jDcgc]}|j}}|j||jg|j ||j gycc}w)NceZdZddZy),SubgraphBuffer.codegen..CodegenGraphc4||_|j|_yr)rr)rrs rr]z5SubgraphBuffer.codegen..CodegenGraph.__init__s" !JJ rN)rr})rrrr]rrr CodegenGraphr s 'rr )rrrjr'codegen_subgraph_with_flattened_outputsr r r)rrr r outer_inputss rrzSubgraphBuffer.codegensz ' '   ,,,7;{{C!++-C C77  ' -doo - - YYK DsBr) rNrrrrdtorch.fx.GraphModuler list[Any]r rrr r rrrr]rrIrJs@rr r sT15,A,A",A ! ,A " ,A  ,A.,A\  rr ceZdZdZd dZd dZeddZ ddZ ddZ e d d dfd Z ddZ dfd Z dd Zdd ZxZS)UserDefinedTritonKernelz> A user-defined triton kernel (e.g. via @triton.jit). cDddlm}ddlm}|j |j g}g}g}t |rtdr%|jfdjDn)tdsJ|jjtdr:jD]*}|jjj|,n)tdsJ|jjj }j|||fS) Nr) Autotuner)kernel_side_table restore_idxc3PK|]}jj|ywr)r arg_names)rrrYs rrzBUserDefinedTritonKernel.get_kernel_and_metadata..0s%*/0FII''**s#& restore_value reset_idx reset_to_zero)triton.runtime.autotunerr *torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrrrnr r r rrr r configs)rr r r restore_value_argsreset_to_zero_argsrrYs @rget_kernel_and_metadataz/UserDefinedTritonKernel.get_kernel_and_metadata#s6P"--doo>(*(* fi (v}-"))*4:4F4F*v777"))&*>*>?v{+))FA&--fii.A.A!.DEFv777"))&*>*>?nnGYYFw 24FFFrctjsy|jjsyt |j j dk7ryt |jdk(sJt|jdtsyt|jdjtsyt|jdjjtsyt|jdjjjtsytd|jdjjjjDsyy)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|] }|dk( ywr/rr)s rrz.vsNa16Nr*T)rJ*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrMrqrr rrrs rr z)UserDefinedTritonKernel.can_fuse_epilogueCs0@@  22 t!!(( )Q .4$$%***$++A. :$++A.33Z@$++A.3388.I$++A.3388==yIN4#4#4Q#7#<#<#A#A#F#F#M#MNNrc(|j|dS)N)epilogue_fusion_codegenrs rrzUserDefinedTritonKernel.codegen{s}}Wd};;rc&|j||S)zQ epilogue_fusion: (fused epilogue node, modified kerel src code) r )rrr s rcodegen_with_epilogue_fusionz4UserDefinedTritonKernel.codegen_with_epilogue_fusions }}Wo66rc ddlm}|j\}}}}|j}i} |rt |j j jdk(sJtt|j j jj} |\} } i|j| | i}| | i} |j||||||j|\} }}}|jDcic]}||j|}}|j| |j Dcgc]}|jc}|j Dcgc]}|j"s|j$}}t'fd|D}g}g}g}g}t)j*|j-t/t)j0d|D]f\}}||vr|r|j3||j3|t5|t6r?|j3|j9|j3|j;t5|t<t>t@tBfr,|j3||j3tE|||vr'|j3d|j3t<|R |r(|j3d|j3t<-|jG|jGOtIdtE|d ||jK|| |jM| ||||||d |jO|jPj ycc}wcc}wcc}w) YOverrides the parent member. See https://github.com/pytorch/pytorch/issues/151692r)triton_version_uses_attrs_dictrIc3(K|] }| ywrr)rrr s rrz3UserDefinedTritonKernel._codegen..s$FaYq\$Fsrr2NzUnsupported arg type: rT) arg_typesraw_argsraw_keys triton_meta inductor_metar%roriginal_fxnode_name))rkr r rrr rr=rrr!define_user_defined_triton_kernelgridrrh rparams is_constexprnumr@rrZrTrrepeatrrrrrrr rr rrrrgenerate_kernel_callrr)rrr r rYr r r kernel_kwargsepilogue_out_overridemutable_arg_nameepilogue_computed_bufferrnew_namer r extra_launch_argsrk named_argsrh constexprsconstexpr_namesrr raw_keys_filteredraw_args_filteredrrr s @rr z UserDefinedTritonKernel._codegensM I  ( ( *       02 t((44;;<A AA#D):):)F)F)M)M$NOTT *9 ' $aWt{{W,<>VWM%57O$P !  5 5      II      261S1S ,-At$$Q' '   /0%+]]3QVV3 %+]]EannaeeE E$$F:$FF! ')')"    I$4$4R$8:K L $ WID#&+I+K  $ $T *  $ $S )#v& C1134  1C#udD!9: C   c+( B  %23KKO$$S)%))+%))+),B49+RPSu*UVVI$ WL Wh/$$  &&#'??$!%!2!2 % i 4Es&M= NNNcPt||t|j|zSr)rnr>r)r# ros rr>z,UserDefinedTritonKernel.get_free_symbol_usess- w+M:=M II}>   rctSrr?rs rrz0UserDefinedTritonKernel.get_unbacked_symbol_defsrrcg}i}g}|jD]\}} t| trXtj |j | } ||vrt j| ||} |j| | ||<n|j| | ||<t|dk7sJ|dj|_ t|tsJt|t|=dt!|j|t#||||_||_|j)\} } } } t+| dsJ| j,Dcgc] }||vs| c}|_ddlm}m}t| dkDr| dj6ni}ddl}| j:|_|j?|j<|_ ||j<|_!||_"|| i||||_#|jFjHjJDcgc]:}t|jM|jNtr||jN<c}|_(|jPDcgc]#}tSt!|j||%c}|_*tVjXj[|ycc}wcc}wcc}w)Nrrr )identify_accessed_tensorsidentify_triton_stores).rTrrrrrr rrrrrrrrnr]rrr r# r rr rr r7 r8 rastr kernel_srcparse kernel_astr kernel_argsr rr=rrr rrrwrr)rr r# tma_descriptor_metadatar= rjrrrk rOrrYr rrr7 r8 autotuned_kwargsr9 rrrps rr]z UserDefinedTritonKernel.__init__s "$&&( %%' DAq!Y' 99$:L:LQ:OP//%,,Q0G0JKA a q $$Q'q  6{aQi**, &(+9T&\9+   dkk *  -   % $ < < >Av{+++!++. sk/AC. * 14G q0@71:,,b!**))DOO43DOOD&5  /{ /. / # ((44;; +//#((3Y?  ! ((!  :T[[93 E!  ""4(K. 6 ! s J?J??K/(K c,t|jSr)rrrs rrz#UserDefinedTritonKernel.get_outputsO sD))**rc|jSrrrs rrz"UserDefinedTritonKernel.get_deviceR rr)rz(tuple[Kernel, Any, list[str], list[str]]r_r )rr|r ztuple[ComputedBuffer, str]rr)rr|r z!tuple[ComputedBuffer, str] | Nonerrrrr}r) r rr# rr> rr= rrrrrg)rrrrr r rrr r rbr>rr]rrrIrJs@rr r sG@6p<<7+7>X7 7i %i ;i   i V56$) ! ! 7 N)N) N) "0 N) $ N) N)`+rr cTeZdZdZddZddZd dZd dZ d fd ZxZ S) InplaceBernoulliFallbackE This needs to be a custom class to handle mutation properly c  td|jDsJd|jD\}tjjr\|j |j d|ddjtt|jd|jy|j |j d|ddjtt|jd|jy)Nc3<K|]}t|tywrrqrrs rrz3InplaceBernoulliFallback.codegen..\ s>Q:a(>r|c3XK|]"}tt|j$ywr)rrrrG s rrz3InplaceBernoulliFallback.codegen..] sIVQ113Is(*rrEz, NULL)r) rrjrwrrrr rrreprrending)rrrs rrz InplaceBernoulliFallback.codegen[ s>$++>>>>IT[[I 77     '')*!A3b3tTEWEW;X1Y0ZZabibpbpaqr    '')*!A3b3tTEWEW;X1Y0ZZ[\c\j\j[kl rcyrrrs rrz(InplaceBernoulliFallback.should_allocatej rrc&|jdgSrrxrs rrJz+InplaceBernoulliFallback.get_mutation_namesm "##rctSrr?rs rrz1InplaceBernoulliFallback.get_unbacked_symbol_defsp rrc^t|dt|j|j |g||t j j|jt j j||_ t j j|y)Nrr) rnr]rrrrwrrrsrrr)rrrrrps rr]z!InplaceBernoulliFallback.__init__s s   alln -    $ #   ##AJJL1GG++D1  ""4(rr r_rr)rrrrrrrr rrrrrrrJrr]rIrJs@rrC rC V sF $ )' ),2 )DG )  ) )rrC cxeZdZdZddZd dZd dZd dZ d fd Ze d ddZ xZ S)InplaceCopyFallbackrD cR|j\}}}|j|||yr)rf codegen_device_copy)rrrr non_blockings rrzInplaceCopyFallback.codegen s)#'#4#4#6 c<##Cl;rcyrrrs rrz#InplaceCopyFallback.should_allocate rrc&|jdgSrrM rs rrJz&InplaceCopyFallback.get_mutation_names rN rctSrr?rs rrz,InplaceCopyFallback.get_unbacked_symbol_defs rrct|d|||ddtjj |dj tjj ||_tjj|y)Nz aten.copy_aoti_torch_copy_)rrr) rnr]rwrrrsrrr)rrNrjrrps rr]zInplaceCopyFallback.__init__ sr     +.   ##F1I$6$6$89GG++D1  ""4(rc||fDcgc]}|j|}}|f}tt|j||}|Scc}wr )rrT rr)rrrrW rrjrrNs rrzInplaceCopyFallback.create sX25c ;1###A&;;% $ cnn. /     . s 9a!!# 9r|rrEr) rrrjrrI rrr rrJ )rrargrefss rrz$MutatingFirstArgExternKernel.codegen s~ ,,, 9T[[ 9 t)) *  ##%&a '(:';1W^^(>(@Ar)rl rrGrrrr rJs@rri ri sBBrri c(eZdZdfd ZddZxZS)SetSourceTensorKernelc |j|}|jt| |j ||gdt j jjjt|tttfsJt|tj j"j%|j&j)tj j"j%|j)tj j"j%|j)|j+}t-t/|||t-t/|||g|_y)Nz!torch.ops.aten.set_.source_Tensor)rrr)rr"rnr]r rrurset_ source_Tensorrrorqrrrwrrk rrMrsrrrr)r self_tensorstorage_tensorrrps rr]zSetSourceTensorKernel.__init__ s,++N;$$&   % % ' . )B ++99  +*i'HI 4 L  I ##'' (8(8(A(A(CD ##''(?(?(AB ##'' 8**, :V4k4 H :V4nd K! rcF|jd|jdgSrMrM rs rrNz2SetSourceTensorKernel.get_inputs_that_alias_output s"DOOA$677r)rr rrs rrrr)rrrr]rNrIrJs@rrn rn s  *8rrn cjeZdZdZd dZd dZd dZd dZddd dfd ZxZ 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. c&|j|yr)generate_scatter_fallbackrs rrzScatterFallback.codegen!s))$/rcyrrrs rrzScatterFallback.should_allocate!rrcf|jd}t|tsJ|jgSrrvrrls rrJz"ScatterFallback.get_mutation_names !s.kk!n#v&&& rctSrr?rs rrz(ScatterFallback.get_unbacked_symbol_defs!rrNTr include_selfc ft|t|_|jr%|||fDcgc]}|j|} }|f} n$||fDcgc]}|j|} }||f} t |dt |j|j| | ||dt|ddg|tjj|jtjj||_tjj!|ycc}wcc}w)Nrr} rr~ )rrr)rr src_is_tensorrrnr]rrrrrwrrrsrrr) rrrrBrrrr~ rtensorsrrps rr]zScatterFallback.__init__!s(Y7   78%oFt))!,FGF FM78%jAt))!,AGA #JM   alln -    ( | <";/+3^*D#   ##AJJL1GG++D1  ""4(%GBs D)D.r r_r r)rrrrrBrrrrrrrer~ rrrrR rJs@rrv rv s| 0 "!!)!!) !) !)  !)  !)!)!) !)!)rrv c\eZdZdZddZddZd dZd dZ d fd ZxZ S) IndexPutFallbackzQ This needs to be a custom class to handle mutation and indices properly c&|j|yr)generate_index_put_fallbackrs rrzIndexPutFallback.codegen;!s++D1rcyrrrs rrz IndexPutFallback.should_allocate>!rrc&|jdgSrrM rs rrJz#IndexPutFallback.get_mutation_namesA!rN rctSrr?rs rrz)IndexPutFallback.get_unbacked_symbol_defsD!rrc ||_|Dcgc]}|| }}||g|Dcgc]}|j|}}d} t | dt j |j ||fd| |tjj|jdtjj||_ tjj|ycc}wcc}w)Naoti_torch_index_put_outrzaten.index_put_)rrrr)rrrnr]rrrrwrrrxrrr) rrrrr accumulater valid_indicesr rrps rr]zIndexPutFallback.__init__G!s $+=qq}= =34f2M}2MNQ4%%a(NN4   alln -    ( M0+#   ##DOOA$67GG++D1  ""4(>NsC/C/C4r r_rr) rtorch._ops.OpOverloadrrrr rrr rrrrR rJs@rr r 6!s`2$)*) ) )  )  ) ))rr c&eZdZeddZddZy) DeviceCopyc Z|j}|J|jst|tjj vrt d|jDrtjjsitjjr>tjj|tjj||j|Stjj|tjj|td|f}tj!|}d}|j#r|j%}t'|j(xr|j(dk(xr|}|j(dk(xrt'|j(xr|}|r t+|rd|j-_t1t3||j5|j#|||j7|g|S)Nc3TK|] }|tjjv"ywr)rwrr r)s rrz$DeviceCopy.create..k!sGqA***Gs&(zDeviceCopy in input programr!Tr)rrDrtrwrrDrrrJ aot_inductoruse_runtime_constant_foldingradd_device_inforHrfrrr r4rnrr r rRr rPrr) rrrrW x_devicerris_destination_pinnedis_source_pinneds rrzDeviceCopy.createc!s<<>### Qqww'>'>>GA4D4D4FGG''DDww""''/''1''/ / ' )78%  + +A . ::<\\^F 8== ! KfkkU&: K|  MMU " Kvfkk': K|   5a 8'+ALLN $   /    q ! "   rc|j}t|dk(sJ|jr2|j|d|jj |dn'|j|d|j |dt |j tr7|j jr |j|jyyy)NrrrI) rf rrrV rrrNrrR sync_d2h_copyrs)rrrs rrzDeviceCopy.codegen!s  "4yA~~     ' 'Q));;=tAw   ' 'Q1G1G1I4PQ7 S dkk6 *t{{/D/D  ! !$--/ 20E *rN)rrrrjrW rrrr )rrrrrrrrrr r b!s- - ^ 3rr ceZdZdZddZd dZ d fd Zd dZed d d dZ 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. ctSrr?rs rrz$DynamicSelectStorageOffset.get_reads!rrcyrrrs rrz*DynamicSelectStorageOffset.should_allocate!rrct|dttjdg||_||_||_||_||_ ||_ yNr!r) rnr]rrrunbacked_offset_symbolr base_offsetbase_dim_striderr)rr rr r rrrps rr]z#DynamicSelectStorageOffset.__init__!sS ze1DErJ'=# &.  rc.t|jgSr)r@r rs rrz3DynamicSelectStorageOffset.get_unbacked_symbol_defs!s466788rct|j|jt|j|jt|j|jt|j |Sr)r)rrrr r r=s rr>z/DynamicSelectStorageOffset.get_free_symbol_uses!s] TZZ 7 U#DII}= > U#D$4$4mD E U#D$8$8-H I  rc>|j||jy)Nr)codegen_dynamic_select_indexrrs rrz"DynamicSelectStorageOffset.codegen!s,,T,Drrzr_)r sympy.Symbolrr r sympy.Symbol | intr r rr rrrrrrrr}r rrrrrrr]rrbr>rrIrJs@rr r !s ,(  ,  !  &989$) ! ! : Err ceZdZdZddZd dZ d fd Zd dZed d d dZ 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. ctSrr?rs rrzDynamicSliceSize.get_reads!rrcyrrrs rrz DynamicSliceSize.should_allocate!rrct|dttjdg||_||_||_||_||_ yr ) rnr]rrrunbacked_size_symbolrrrr)rr rrrrrps rr]zDynamicSliceSize.__init__!sH ze1DErJ$8!   rc.t|jgSr)r@r rs rrz)DynamicSliceSize.get_unbacked_symbol_defs!s444566rct|j|jt|j|jt|j|jt|j |Sr)r)rrrrrr=s rr>z%DynamicSliceSize.get_free_symbol_uses!sY TZZ 7 U#DHHm< = U#DII}= > U#DII}= >  rc&|j|yr)codegen_dynamic_slice_sizers rrzDynamicSliceSize.codegen"r rrzr_) r r rr rr rr rr rrrr}r r rJs@rr r !s~ *"  !  !  7./$) ! ! 0 1rr cTeZdZdZddZddZ d fd Zd dZd dZxZ S) rz; The result of a call to aten._local_scalar_dense. ctSrr?rs rrzDynamicScalar.get_reads"rrcyrrrs rrzDynamicScalar.should_allocate"rrc|jt| dtt j d|j |g||_||_yr ) rrnr]rrrrsymkeypath)rr r rMrps rr]zDynamicScalar.__init__"sM   *ELL$78$:M:Mtf:U  rc.t|jgSr)r@r rs rrz&DynamicScalar.get_unbacked_symbol_defs!"s488*%%rc&|j|yr)codegen_dynamic_scalarrs rrzDynamicScalar.codegen$"s&&t,rrzr_)r r r zpytree.KeyPathrMrrrrr ) rrrrrrr]rrrIrJs@rrr "sA*8@F &-rrcleZdZdZddZd dZd fd Zd dZed d d dZ d dZ xZ S)rz5 The result of a call to aten._assert_scalar ctSrr?rs rrzAssertScalar.get_reads-"rrcyrrrs rrzAssertScalar.should_allocate0"rrc~t|dttjdg||_||_yr )rnr]rrrscalarr/ )rr r/ rps rr]zAssertScalar.__init__3"s7   ell51 2    rcyr1rrs rrg zAssertScalar.has_side_effects>"rrc.t|j|Sr)r)r r=s rr>z!AssertScalar.get_free_symbol_usesA"s  ];;rc tjsytt|j d}t j jryt j jrad|d}t j jj|jd}|jd|d|jd|d yt j jj|jd}|jd |d |jd t|jd|j|j!d y)NFrhzstd::to_string(r)rzif (!(z()) { throw std::runtime_error("Expected z but received " + z); }zif not (z):z raise RuntimeError(z = None)rJscalar_assertsrrr>rwrrrrcodegen_cpp_sizevarr rr/ codegen_python_sizevarrI rs)rrr symbol_strsizevars rrzAssertScalar.codegenG"s4$$  d44454IJK 77    WW *6(!4Jgg**>> e?G    !J488*Tfgqfrrwx gg**AA eBG    4 5    7TXX7GqI J    19 :rrzr_)r rxr/ rrrrrr}r ) rrrrrrr]rg rbr>rrIrJs@rrr("sM N+$)<!< !<,< ;rrc"eZdZUded<ded<y)ExternKernelNoderrzexport_schema.NoderNrrrrr r g"s I rr ceZdZdZ ddd dfdZdfd ZddZddZddZe dd Z dd Z dd Z dd Z d ZeddZe ddZe ddZeddZxZS)rz A class that represents a fallback kernel for handling operators that are not directly support by inductor. It currently supports functional ops, view ops, inplace aten ops, and mutating ops that are auto-functionalizable. Nrct|t|t||d_|xsi_t |t jjt jjfsJd|dt|d|_ |_ |in|_ jJtj j#jg_g_t jt jjrydjj)vryjj*}t j,j.j1jr-j&j3|dj5yd d} |j6r?t9js*| jst;d jjj<j>\} }d fd } t j,j.jA|| |D]\} } | | | y) NrQ F#Fails to create FallbackKernel for r not supported_c10d_functionalrc"tjj|jtjjj xs>t |dxr0tjjj |jvS)N py_kernels)r_C%_dispatch_has_kernel_for_dispatch_keyr DispatchKey Functionalizerr r]s rhas_functionalize_implz7FallbackKernel.__init__..has_functionalize_impl"sg88AA 588//==L)HHH((66"--G  rz'NYI: Can't generate FallbackKernel for ctjtjr&t|tt fsJt|t jjrt|t tfrJ|yjydfd dfd }t jjr||D] }|| yyt jjsJ||y)Ncjj|jjJjjr?j jt t|j|yyr ) alias_namesrrs alias_infois_writerrrr)rinfors r add_aliaszPFallbackKernel.__init__..handle_aliasing_and_mutation..add_alias"sj  '' 5222??++))00&z'H!TR,rc|yt|ttgtryt|tsJt ||yr)rr r r=rr)rr s radd_alias_if_graph_bufferz`FallbackKernel.__init__..handle_aliasing_and_mutation..add_alias_if_graph_buffer"sD9a&$!:!:;!!V,5d1g5,! r)rrrr)rrrr) rrrListTyperr library_utilsis_tensor_like_typer is_tensorlist_like_type)r rr optional_tensor_argr rs` @rhandle_aliasing_and_mutationz=FallbackKernel.__init__..handle_aliasing_and_mutation"s$))U^^4!#e}5@tCy@500;&cE4=999{&  44TYY??/2G+12EFG#%88CCC)#.r)r^r rr)r ztorch._C.Argumentrrrr)!rnr]ruse_runtime_dispatchrrrrrrrrr rrrwr warn_fallbackr r rr_libraryrmutates_and_returns_first_argrrs is_mutabler&rrjr zip_schema)rrNrYr$ nontensor_argsr rrschemar rr r rrps` rr]zFallbackKernel.__init__t"s&   +  . !  %*!!2!8b UZZ**EJJ,J,JK  X 04<. W X ","Nb &&222 d556'))+ d&& (F(F G  !1!1!6!6!8 8 !!)) >>   = =d>N>N O    & &{1~'>'>'@ A     *4+;+;<*4+;+;<%9$:J:J9KL **4;;8J8JK f( /T--88vN 4ID# (s 3 4rc,t|}|jtjj j urT|jD]E}t|ts|jtj|j}G|Sr)rnr3rr_prims rng_primsgraphsafe_run_with_rng_staterrr with_readrKrIrs)rrrrps rr3zFallbackKernel.get_read_writes"syg-/   u||55RR R)) c>2"-"7"7$,,S\\^<#K  rc n|j|j|jt|ddSNr)(codegen_unbacked_symbol_defs_for_outputsrsrKrrs rcodegen_unbacked_symbol_defsz+FallbackKernel.codegen_unbacked_symbol_defs#s0?? MMOT\\749Ld+S  rct|ddx}rKttjjj |}|Jt |jSt Sr rr;rwrrrr@rrrresolveds rrz'FallbackKernel.get_unbacked_symbol_defs#_ '.A4 H H  H0  **,=H' ''hmmo. .< rcZtjGdd}t|jsJ|jDcgc]}||j }}|j ||j \}}tjjrt|jtjjr|j||}t!|jj"j$|Dcgc]9\}}tjj&j)||j*;}}}n6|Dcgc]+}tjj&j)|-}}|j,j/||Scc}wcc}}wcc}w)Nc eZdZUded<ddZy))FallbackKernel.codegen_args..Shimrrefc|jSr)r rs rrz2FallbackKernel.codegen_args..Shim.__repr__#s xxrNro)rrrrrrrrShimr #s H rr )r dataclassrrrjrr rrwrrrrrrrr] rrrrr_ rrr)rr rr$ rrparams rrf zFallbackKernel.codegen_args#sY         ,,,<@KKHqtA//12H H**;8J8JK f 77  :d.>.> @U@U#V..tVF"0F(c|r!|Dcgc]}t|tr|c}nd}|r7|sJ|Dcgc]#}|js|j%}}|dSt|tjr |j St|tj jtfs t|rtj dSt|ttfr{td|D}|Dcgc]}|s| }}t|dk(r|dS|sy|D]7}t|tj sJt|js5|cS|dSycc}wcc}wcc}w)Nrr!c3HK|]}tjd|ywr)r find_devicers rrz-FallbackKernel.find_device..=#s'$**43$s "rI)rrrrrrr ScriptObjectr*r,rrr@rrnr)r$ r+ rnon_torch_bind_tensor_argsrdevices device_setrs rr zFallbackKernel.find_device)#s_ $ J1:a+IQ J # & ;3>SC#..BRs~~'SGS1:  nell 3!(( (  UXX224DE ^ ,<<& & ntUm 4#$($J -7A&&vAGA7|q qz!! "!&%,,777&++&!M "1: ? K TBs!E!E!E&E&;E+E+c2ddlm}||jS)Nr) is_impure)torch._library.utilsr r)rr s rrg zFallbackKernel.has_side_effectsO#s2))**rct|jtjjtjj fs+Jd|jdt |jdt|jtjj sSd|jjvr7|jjjrt|jrgS|jS)Nr rr r ) rrrrrrrrrr r&r rs rrNz+FallbackKernel.get_inputs_that_alias_outputV#s   uzz44ejj6T6TU  2$2B2B1C2D$$%&n 6  4++UZZ-K-KL"$*:*:*?*?*AA  ((33&t'7'78I## #rcNt|jdksJ|jSr)rr rs rrJz!FallbackKernel.get_mutation_namesk#s'4&&'1,,,"""rc tjd|j|jt |t sJt ||j|j|j\}}|j||}|jDcgc]}|j|fi|}}|j}tjjsg||St!dg}|j#|||} d d}t |t$j&j(j*r#|j-|d|dj.} n|j0j.} t3| dk(r>|j4r |j4n |j6} | dj8} || | g} n9t;| |j4D cgc]\} }|| j8|} } }|jJi}t |jt$j<j>r$tAjB|jd|d<tE|jtGjH|jjK|| | }tjLjO|g||Scc}wcc}} w) 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.Nc t|tjtjfr|}t|tt frt |dk(sJ|d}t|tjrTt|tsJtjjtj|jS|JtjjdSt|tjrt|jtjrpt|tsJt!|tjj|Dcgc]&}tj|j(c}St|tj"rt|jtjr|>tjjtj$jdSt|tsJtjjtj$jtj|jSt|tj&r tjj| St)d t!|cc}w) NrIrr) as_tensorT)as_none) as_tensors)as_optional_tensor)as_intzUnsupported return type )rr TensorTypeNoneTyperrrr export_schemaryrTensorArgumentrsr getElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError) return_typer\rs rhandle_single_outputzFFallbackKernel.export_extern_kernel_node..handle_single_output#sh+(8(8%..'IJftUm4v;!+++ )Ck5+;+;<%c6222(1188"/">">CLLN"S9;&;(11888FFK8Z**,e.>.>>"&(3AT&\A3$--44$* &44#,,.I 5 K););<**,e.>.>B>(1188+8+O+O+V+V$(,W,9 &ff555(1188+8+O+O+V+V&3&B&B%+__%6',W,9K7$--44F4CC"%=d;>O=P#QRR7 s)+K!rrI1torch_library_custom_op)rTrjrKr`)rr)r" z1torch.TensorType | torch.ListType | torch.JitTyper\IRNode | Sequence[IRNode]rzexport_schema.Argument)(rrrsrrrrr rjrr] rrh rwraot_moder%serialize_inputsrr torchbind CallTorchBindr returnsrrrKrrrrr custom_ops_maybe_get_opdefr r r>rextern_kernel_nodesr)rrrrordered_kwargsrT serializernamed_argumentsr# r+ rKr" output_arguments return_schemar\r`rs rexport_extern_kernel_nodez(FallbackKernel.export_extern_kernel_nodeo#s B MMO    $/;d;/**4;;8J8JK f**4899  "D ! !# 0 0  !!ww+T+N+ +*44 $55fdFK3 SJ3 S-3 S$3 Sj fe55??MM NmmDGT!W5==Gnn,,G w<1 '+lldll8M8MG!!*..K 4[' JK .1$,,-G   *M6 %!++   +++ t'')>)> ?++D,<,<=I25H. /##'',,.&(!   $$T*''''Y ` s KK c"j}|J|jdk(rt|tjj sJt |tjjrddl m }t||vrtjd|d_n~|jdk(r4t|tjj sKJt |tjjr!|t j"j$v_tjjrt|tjj rjsd fd j'j(j*\}t-j.|fd j0D}t3fd t5||j6j8D_j;|jrj=}j>JjJ|jAjCj>fd j|jDr jDn jFn^|jItjJtLr3jO|jQ|jS|jU|y) r NrrrzG%s is missing a c-shim implementation, using proxy executor as fallbackT _quantizedct|tjr|jSt|tjSr)rrr r NumberType)r is_numbers rr9 z)FallbackKernel.codegen..is_number$s:a!3!34$Q%5%5%788!!U%5%566rc3DK|]}j|fiywr)rh )rrk rrs rrz)FallbackKernel.codegen..#$s+*D))!6v6s c3lK|]+\}}t|txr|j-ywr)rcomplexr)rrOrFr9 s rrz)FallbackKernel.codegen..($s5,Aq1g&A9Q[[+AA,s14cHgjjSr)rf rl rsrrz(FallbackKernel.codegen..6$s$F$++-F0C0C0EFr)rz torch.JitTyperr)+rrrrrrrrwrrrrrrrr rJr custom_ops_to_c_shimsr rjrrrZrr0rrrrr4 r,generate_fallback_kernel_with_runtime_lookuprsrKrgenerate_fallback_kernelrNrr{ r r r ) rrrYrr args_iter exported_argsr9 rs ` @@rrzFallbackKernel.codegen#so!!!!!   v %fejj&;&;< Jd6l J<ww""Lv;&;;KKa15D-    -fejj&;&;< Jd6l J< WW f11GGG  % GG  65::#8#89-- 7 ..t{{D>+C+CD,)D % W%  $ $ ::tj jRs$r"tTjW||| |  }||S|tY || | n"sJd t[ || | dfd |g}t|t>t\fr |_/|St|t`rt]|_/|S|g_/|S#1swYxYw)z9Create an instance of FallbackKernel from an _OpOverloadsNr)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variantrIrrrr)rNrjrrrrc32K|]}t|ywr)r rs rrz(FallbackKernel.create..$s!K,s"3!Krr!rhas_unaligned_inputrrrz"Not sure where to find device infoc tttfr.tfdt t DStt r:jDcic]\}}||t|fgz c}}Sttjrntj }tjs s ts3tj j"j%|j&|Stt(rSttj*rj,j.Sttj0j2t4fs t7rt9t; SJdtdycc}}w)Nc3TK|]}|t|fgz!ywr)r)rrgenerate_outputrr\s rrzAFallbackKernel.create..generate_output..$s5$$F1Iw4<:K9L/LM$rrzFallbackKernel output type z is not supported)rrrrrrrrTrrrrrJ assume_unaligned_fallback_outputrtrwrr rrrrI rrr r r*r,rr) r\rrrrrrra r] packeds `` rra z.FallbackKernel.create..generate_output$s&4-0#tF|$"3v;/$FD)%+LLN Sg$v,9L8M.MNNFELL1!((0 ;;*,V4GG--11#((; FC( FELL1{{'''..0@A ((f- ~1$v,?PQ~Es+#G)r\rrzlist[tuple[Any, int]]rr)1r*_fused_moving_avg_obs_fq_helper_functionalrr rwrr6r r0 rrrrrT r r is_builtinrtorch._library._out_variantrV rW rX rY rrrPrrrArr r_make_out_variant_kernel_namer0r r r) r* ru higher_orderprintrExternKernelMultiOut try_createrrrrKr)rrYrrfake_incorrect_kernelscontextr+ r$ r% r rrV rW rX rY out_oprN out_resultrKrra r] rc s` @@@@rrzFallbackKernel.create$sU#'"Q"Q!S / /1$79J9JKG!mG  <#""6;D;F;  !  < fejj33 4 ,C , ,V Ed Ef E vuzz44 5NN((33F;>5<<8  Ffnn-'/~26:!c*;F*C&D&I$)00(..0>//05^2245  '! ,"&"7!'DV'L & "!K{!KKn= vu66@@NN O//555\\%(F vuzz44 5NN((33F;GG''-88"3$79 J%!!  !&)"3F ?? ?6!0"3F) ) V"."5 ge} -$FN   &"7^FN&YFN{ < )rrirs rr>z MultiOutput.get_free_symbol_uses%s7[[^ *f-9z9-..}==rcpt|jdk(xrt|jdtS)NrIr)rrjrrIrs rrzMultiOutput.should_allocate%s04;;1$ t{{1~'8 9 rc|jDcgc]>}t|tr,t|j dkDr|j @c}Scc}wr)rjrrrrNrsr{ s rrNz(MultiOutput.get_inputs_that_alias_output%sN{{ #~.C4467!; LLN   sAAcLt}|jD]E}t|ts|j t j |jG|j|jjdfd }|j}tj xs|duxst|j }t j||j!d|}t j"||j$tS)Nc^t|dk(sJtj|dSrrrs rrz*MultiOutput.get_read_writes..dummy%rrrrr<r)r@rjrrrrKrIrsr r rrJrprnrrTr r?r=) rr6rlrrshould_normalizewrite_rwrrs @@rr3zMultiOutput.get_read_writes%s/9l;; @C#v& ,..s||~>? @}}//#002 ;"11 1 '~ '&++&&   33 4==?B2B &&??"   rr rr) rNrrwrrlist[tuple[Any, ...]]r rrrr}r_rrx) rrrrr]rbr>rrNr3rIrJs@rrrs%s438 S S S' S ,0 S  SM*$)>!> !>+>    rrcZeZdZdZ dfd ZeddZddZd dZxZ S) rzMultiOutput for opaque objects.c:t||||d||_y)NTr )rnr]r )rrNrwr opaque_valuerps rr]zOpaqueMultiOutput.__init__%s$ SWX$0!rctd)NzOpaqueMultiOutput has no dtype)rQ rs rrzOpaqueMultiOutput.dtype%s=>>rc|Srrrs rrz#OpaqueMultiOutput.wrap_for_lowering%rrcXt}|jD]E}t|ts|j t j |jGtt j |jg}t j||tS)Nr<) r@rjrrrrKrIrsr?)rr6rlr=s rr3z!OpaqueMultiOutput.get_read_writes%s.8l;; @C#v& ,..s||~>? @0:  ! !$--/ 2 30 &&"   r) rNrrwrrr r rrr)rr)rrrx) rrrrr]rrrr3rIrJs@rrr%sY)111' 1  1  1??  rrc eZdZdZddZddZy)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. cyr1rrs rrz%AllocatingMultiOutput.should_allocate%rrcb|js#|j||j|yyr)r r{ r rs rrzAllocatingMultiOutput.codegen%s,55  % %g .  * *7 36rNr_r )rrrrrrrrrr r %s 4rr c|j}|jjjdd}|j}d|d|d|S)z8Build fully-qualified kernel name for an out-variant op.z::rIz torch.ops.r)rrrrr)rn nsrr rs rrg rg %sP   Bnn!!''-a0G##H t1WIQxj 11rceZdZUdZded<ded< d fd Zd dZedd d  dd ZxZ S)rj 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_nodescnt||i|||_g|_t ||_||_yr)rnr]r r rg rr)rrn r rrrps rr]zExternKernelMultiOut.__init__&s< $)&)*(*%"?"G!rcH|j||j|yr)r generate_extern_kernel_multi_outrs rrzExternKernelMultiOut.codegen&s W%006rNFr\ c "ddlm} m} m} | |jsyt |t tfsy| |} | y| | }td|Dsyt|t|k7ry|t|||||||| | }g}t|D]\}}t|j|jg|jg|j!}t#||t%||fg}t&j(s | s t+|s3t,j.j0j3|j4|j7|||_||_t |t r t |St|S) zGCreate an ExternKernelMultiOut if the op has a matching .out() variant.r)rV rW rY Nc3PK|]}t|tj ywr)rrrrG s rrz2ExternKernelMultiOut.try_create..5&sG1:a.Gs$&r)rrrn r rZ )rNrwr)rf rV rW rY rrrrrrrrrPrrrArr rrJrb rtrwrr rrrr rK)rrYr+ rr$ r% r rrr] rV rW rY rn r rc rKr tensor_outrN multi_outs rrk zExternKernelMultiOut.try_create&s  fnn-.5$-8' >)&1 GGG ~ #m"4 4 V ,    /'  02&~6 &MAz !(( &&(z''(-**,- F .~.23I 77&(4))--inn= NN9 %% &(+2'  ne ,> !G}r) rrrn r r rrrrrr )rYr r+ rrrjr$ rr% rr rrr rrq r] rrz&Sequence[AllocatingMultiOutput] | None) rrrrrr]rrrk rIrJs@rrj rj %s 99 " "& "! "  "  "7HL$)J%JJ J & J ' J+J&JEJ"J 0JJrrj ceZdZUdZded<d.dZd/dZd0dZd1dZd2dZ d3d Z d4d5d Z d d d6dZ d7dZ d8dZd.dZd7dZ d9 d:dZd;dZddZd?dZd@dZdAdZdBdZdCdZd.dZd.dZdDdZdEd Zd3d!ZdEd"Z dBd#Z!e"d d9 dFd$Z#dGd%Z$dHd&Z%d4dId'Z&e'dJd(Z(dKd)Z)dJd*Z*dCd+Z+e'dLd,Z,d3d-Z-e-Z.y )MrzC TensorBox / StorageBox allow in-place mutation of Tensors rrMc6|jjSrr+rs rrz!MutableBox.has_exceeded_max_readsm&r,rc6|jjSrrrs rrzMutableBox.get_devicep&rrc6|jjSrrrs rrzMutableBox.make_loaders&yy$$&&rc6|jjSr)rMr rs rr zMutableBox.make_indexerv&yy%%''rc6|jjSr)rMr4rs rr4zMutableBox.get_stridey&rrc6|jjSrr"rs rrszMutableBox.get_name|&r#rNc8|jj|Sr)rMrrs rrzMutableBox.has_large_inner_fn&syy++I66rTrc<|jj||Sr'r(rs rrzMutableBox.mark_reuse&r)rc6|jjSrr1rs rrzMutableBox.realize_hint&r rc6|jjSr)rMrprs rrpzMutableBox.unwrap_view&r rc6|jjSr)rMrrs rrzMutableBox.is_input_buffer&yy((**rc6|jjSr)rMr"rs rr"zMutableBox.freeze_layout&syy&&((rc:|jj||Sr)rMr&r$s rr&z*MutableBox.freeze_layout_with_stride_order&syy88 NNrc8|jj|Sr)rMr*r)s rr*z(MutableBox.freeze_layout_with_fill_order&syy66u==rc8|jj|Sr)rMr-r,s rr-z(MutableBox.freeze_layout_with_same_order&syy66v>>rc:|jj||Sr)rMr1r/s rr1z+MutableBox.freeze_layout_with_exact_strides&syy99-WWrc6|jjSr)rMr3rs rr3zMutableBox.get_read_writes&r rc6|jjSrr9rs rrzMutableBox.get_reads&r7rc6|jjSrr7rs rr9zMutableBox.num_reads&r7rc6|jjSrr3rs rr;zMutableBox.get_storage_numel&r4rc6|jjSrrrs rr@zMutableBox.get_reduction_type&rrc6|jjSrrrs rrBzMutableBox.get_reduction_size&rrc6|jjSrr6rs rrDzMutableBox.is_extern&r7rc6|jjSr)rMrFrs rrFzMutableBox.is_no_op&r#rc8|jj|Srrrs rrHzMutableBox.constant_to_device&syy++F33rc6|jjSr)rMrJrs rrJzMutableBox.get_mutation_names&rrc6|jjSr)rMrLrs rrLzMutableBox.get_operation_name&rrc6|jjSr)rMrNrs rrNz'MutableBox.get_inputs_that_alias_output&syy5577rc6|jjSrr.rs rrzMutableBox.realize&r/rc8|jj|Srrr=s rr>zMutableBox.get_free_symbol_uses&syy--m<rrrrrNr rr rrrrrrrre&si L2&'(&$7=AD('+);@O"O37O O >?HMX/X@DX X +%%-..%$4..8#L)$)=!= !=*= *+3++&+$ "HrrcVeZdZeeddZeeddZeddZy)rcyrrrMs rrzTensorBox.create&sFIrcyrrr s rrzTensorBox.create's+.rcNt|tr|Stt|Sr)rrrrqr s rrzTensorBox.create's" d1 2KD)**rN)rMrrr)rMrrr)rMr)rrrrrrrrrrr&s@ II ..++rcdeZdZdZddZddZddZddZddZddZ dd dd Z dd dd Z dd Z y )rqz7 StorageBox allow in-place mutation of Tensors ct|jttfr4|jj t j jvSyr)rrMrrSrsrwrrrs rrzStorageBox.is_input_buffer's: dii+!? @99%%'177+?+?? ?rct|jtxr4|jjtj j vSr)rrMr!rsrwrr rs rr9zStorageBox.is_module_buffer's9 tyy> 3 : ""$(9(99 rc tj|jr|jjSt |jt t ttfsJt|j|jj}|jj}|jj}|Jtdt||jj|jj!d|j|_t"j$j'|j|j_t"j$j+|j|j,|j_||j_||j_|jjj2|j_|jj(S)NF)rrrrRr)rrrMrsrr rgrrrrrrrr rr rwrrrrrzr~r|r)rr~r|rs rrzStorageBox.realize'sj  " "499 -99%%' '$))iD$%GH $ IIK  Hii//1 II++- %%'!!!"!ii))+YY'')     00;  ""499- LL  + ' #yy~~88 yy~~rct|jttfr9|jj j dkDr|j yyy)zL Called on buffers we expect to be forced to realize later. rIN)rrMr rgrnontrivial_read_countrrs rrzStorageBox.realize_hint:'sF tyy9i"8 9 **,BBQF LLNG :rcddlm}|jDcgc])}||stjj |+}}|syt |}t|}t|}||k\xr||z dk\xr||k(Scc}w)Nr)is_nonfreeable_buffersFr) rkr rrwrget_dep_size_hintr3rr1)rrr r size_of_reads total_sizemax_sizemin_sizes r$has_accumulated_enough_reads_by_sizez/StorageBox.has_accumulated_enough_reads_by_sizeD's@~~' )#. GG % %c *  ' }%}% ) # %X%* %H$  s.Bctj}|7t|jrtj}n#tj }nt |tsJt |jtxrZ|j|kDxsE|jxs3tjduxr|jtjSr) rJrealize_acc_reads_thresholdr,rMrealize_cpu_acc_reads_threshold$_realize_acc_reads_threshold_defaultrrr r9r realize_acc_reads_size_thresholdr )rr s rrz!StorageBox.has_exceeded_max_readsW's&,&H&H# & .dii .4.T.T+??,93? ??$))Y/  NN : : &&( 77tC==;; rTrcL|dkDrt|jttfr~t |jr3|jj gd}t fd|Dry|jry|xr!|jtjkDSy)zj A heuristic to decide if we should realize a tensor that is used multiple times. rI)exprlog10log1plog2sigmoidtanhc3:K|]}|jvywr)used_ops)rropcounts rrz5StorageBox.should_realize_on_reuse..'s@qG,,,@sTF) rrMr rgr,rr0rr9rJrealize_reads_threshold)rrr heavy_opsr s @rshould_realize_on_reusez"StorageBox.should_realize_on_reusem's 19DII 9/EFdii ))446 @i@@&&(T4>>#3f6T6T#T TrcL|j||r|jyyr')r rrs rrzStorageBox.mark_reuse's$  ' '; ' G LLN Hrc6|jjSrr7rs rr9zStorageBox.num_reads'r7rNr_rdrS)rrrr)rrrrrrrqr{) rrrrrr9rrr rr rr9rrrrqrq 'sA  < & ,JN2=A%rrqc0eZdZUded<ded<dZded<y)Subgraphrrr graph_moduleNzGraphLowering | Noner)rrrrrrrrr r 's I&&"&E &rr c|Dcgc]$}t|tr|jn|&}}tt d|Dt|kScc}w)Nc32K|]}t|ywr)r)rrs rrz'_has_aliased_buffers..'s;"V*;r)rrSrprr@)buffersrs r_has_aliased_buffersr 's^ !+6? COG z;7;; .create_output(s:N"OP **,)))"%$..0#__.%002%00299"("3"3"5"?"? $C[M6:  r)r\rr rrz:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput)r@ r rwrrrbrrTrrurh r rr rrrrr r rr r rr r graph_outputsrr rrK)rr r r r fake_operandsr rQ fx_operandsr new_operandsroperandrKrr rr\outsr s @rrzInvokeSubgraph.create's 7ww++  +00445GH H  HF""eii&<&<&I&II#((+uyy/E/E/U/UUUU,Q/8MF""eii&<&<&I&II#((+uyy/E/E/U/UUUU'++FG4K4?@qQVVE]@M@ AI!I1#"3"3A"6!I!I%' %h/ LC/ARS##G,##,WmC6HI    >> !WW22((,&mm3HN $$X^^4 3"""M2 3.... Gg'<= ++- !!!($F3    !$  G .;DG:LMYQ fa(MM"& SA "J. 3 3TNs/K" K'*K,K8,K5c&|j|yr)codegen_invoke_subgraphrs rrzInvokeSubgraph.codegen (r r)r r r rrNrrrrX)r r r rrz@list[ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput]r )rrrrr rr rKr]rrrrrIrJs@rr r 's!%Ho$(,H%,'+G $+ )  ),< )FW )  )8ff,2f IffP.rr ceZdZUdZdZded<dZded<dZded<dZded <dZ d ed < dfd Z dd Z e ddZ e ddZddZddZxZS) 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] | NonerKc||_||_||_||_t |g|\}}t |d||||||_tjj||_ tjj|yN)rrNrjr) r r r r _split_by_sym_typernr]rrwrrrr) rr r r r rNrsym_argsr$ rps rr]zConditional.__init__7(s#  *, 2I3I3I J+ "   (%6D "GG++D1  ""4(rcg}|jr|j|j|jr|j|j|Sr)r rr r subgraphss rrzConditional.get_subgraphsS(sG      T// 0      T00 1rcRt|tr|S|jjSr)rrrr)rs r _maybe_exprzConditional._maybe_expr[(s a Hvv{{rc |j|}|Dcgc]}|j|}}tjjjd}t |t sJt|g}|D]B}t |tr|j|jd2|j|Dtjjjd} dd} ||fD]} | jtjj| j|| j| _tj| j5| jj|| | jj | | j_ddd|jJ|jJ|jj } |jj } d| fd| ffD]!\}}t#| st%d|d |t'| t'| k(s J| | ft)t+| | D]\}\}}|j-|j-k(s J|||f|j/|j/k(s J|||f|j1j2|j1j2k(rJ|||ft5d ||gzD}t7tjj8j:tjjjj=d d}|Jd t?||||tA| |}t)t+| tjjjdDcgc]\}\}}tCtE|j-|j-n||j/|jGDcgc]}t>jI|c}|jKDcgc]}t>jI|c}|j1j2|j1jL|tN|fg}}}}}||_(ddl)m*}||j\}}}}}||j\}}}}}tW|tW|z}tY|Dcgc]}t[||j\|||!c}|_/|Scc}w#1swYkxYwcc}wcc}wcc}}}}wcc}w)zNCreate a Sequence of IRNodes from a conditional statement (see .lowering.cond)r2rcg}t||D]e\}}t|tr|j|(|jtj t ||jdg|SNFr)rrrrrrrr)r fake_tensorsretr\rs r_require_exact_stridesz2Conditional.create.._require_exact_strides}(srC #M< @  f&;<JJv&JJ$::%f-t{{}E; JrNr true_fnfalse_fnzVOutput aliasing is currently not supported in compiled torch.cond. The outputs of the z% subgraph of torch.cond are aliased: c3\K|]$}t|ts|j&ywr)rrr)ros rrz%Conditional.create..(s) a!67 LLN rrzcannot determine devicer)r r r r rNrrr)-check_input_alias_and_mutation_return_outputs)r rr, zSequence[torch.Tensor]rr)0rrwrrrrrrr>rrbr r rr r r r rlrrrrrr rQrr;rrrr rrrPrr( rrRrrKtorch._higher_order_ops.utilsr3 r@rrrNr)rr r/ r0 r rr r fx_op fake_outputsr. r true_outputs false_outputsrrKrt_of_orr conditionalr\ merged_outputrPr3 rtrue_mutated_inputsfalse_mutated_inputsmutated_operand_indicesrs rrzConditional.createa(s`%%i0 2:;QC%%a(;; ! 4 4 9 9" = +x0C${2CC0$&   ,E%&$$UZZ%67$$U+  , ww++007  + 0  $!(+ H~~%!"!6!6,,#0"*--"7" ((8&HNN&& 64J 44l4HNN0  "}}(((~~)))}}22  44 (,7*m9TU MD'#L1$**./TU\T]_ < C $66U}8UU6&s<'GH UMAzS>>#s~~'77 F!S# F7==?cmmo5 D3} D5>>#**cnn.>.E.EE T3PS} TE U  +   6 GG   & & GG % % ) )*=t D !<#<<!!!#$F3/  :/8L!''"6"6";";E"BC/)  (+*FM' ((*6",,. **,@M@R@R@TU"+11"5U>K>R>R>T8: //3",,.55$//1;;     2&  :':N:N O *Aq%q :(:O:O P +Aq&#--@"AJ E # 56(  8C=//# L( $ [UAUAU- (U#U- U(4A U- #$U5U # U- c |j||j|j|jt |diyr )codegen_conditionalr rsrKrrs rrzConditional.codegen(s9##D)88 MMOT\\749Lb+Q rct|ddx}rKttjjj |}|Jt |jSt Sr r r s rrz$Conditional.get_unbacked_symbol_defs)r r)r rr rr r r r rNrrrq rrrX)rzint | torch.SymIntrz int | Expr) r rr/ r r0 r r rrzlist[MultiOutput]r r)rrrrr rr r r rKr]rrr( rrrrrIrJs@rr r $(s  $I}#(,H%,%)M?)&*NO*,0G )0))#) ) ! ) " )E) )8 XXX X " X  XXt  rr cg}g}|D]?}t|tr|j|j/|j|A||fSr)rrrr)r non_sym_argsr# rs rr" r" )sSLH% c0 1 OOCHH %    $ % \ !!rceZdZUdZdZded<dZded<dZded<dZded<dZ d ed < dfd Z dd Z e dd Z e ddZe ddZddZddZxZS) WhileLoopzSThe IR node for while_loop and while_loop_stack_output. It supports input mutation.Nr carried_inputsadditional_inputsr cond_subgraph body_subgraphr rKc.||_||_||_||_t g||\}} t |d|| ||||_||_tjj||_ tjj|yr! )rG rH rI rJ r" rnr]r stack_outputrwrrrr) rrG rH rI rJ rNrrL r# r$ rps rr]zWhileLoop.__init__%)s-!2** 2 1n 10 1! + "   (%6D "(GG++D1  ""4(rcg}|jr|j|j|jr|j|j|Sr)rI rrJ r% s rrzWhileLoop.get_subgraphsD)sG      T// 0      T// 0rct|s|S|Dcgc]$}t|tr|jn|&}}t }g}t ||D]b\}}t ||vr%|jtj|8|jt ||j|d|Scc}wr) r rrSrpr@rrrrr r)rG runwrapped_buffers seen_buffersrNoriginal_inputunwrapped_buffers r_clone_aliased_inputszWhileLoop._clone_aliased_inputsO)s#N3! !) %/v$GF   V S  )3 +-03NDU0V . ,N,"#|3 l55nEF  $4!56 n-  . ! s)B<ct|tr|St|ttfr t|St|trtj |St dt|)NzNYI unsupported output type: )rrrqrSrrr! r)rs r_maybe_wrap_as_tensor_boxz#WhileLoop._maybe_wrap_as_tensor_boxi)s\ c9 %J j/: ;S> ! [ )##C( (!>tCykJK Krc"ddlm} dd}tjjj d}tjjj d} || z} | D cgc]} | j d} } |D cgc]} | j d} } | D cgc]} | j d}} |D cgc]} |j| }} tj|}||| }|D cgc]} |j| }} |||}||z}||fD] }|jt| tsJt| tjj|j| |j|_tj |j5|jj"| ||urYt%|jj&t%| k(sJ||jj&| |j_ddd |jr |jsJ|jj&}|jj&}t)|rt+d |t%|d k(sJ||d}t|t,sK|j/t0j2k(sJ|t%|j5dk(sJ|t%|dkDsJd |dj7}|Jt%|t%|k(s J||ft9t;||D]\}\}} dd }||j5|j5||j=|j=|j7|j7k(s J||||f|j/|j/k(rJ|||f|Jt?tjj@jBtjjj jEd d}t||||tG|||}|j8t|jjHt0jJjLsJ||jjH| d}tO|}|Dcgc]}|| } }tQ| }!g}"g|_)g|_*|rt%|dk(sJdt9tjjj dD]\}}#tWtY|#jZ|#j\|#j_D$cgc]}$t`jc|$c}$|#jeD%cgc]}%t`jc|%c}%|tf|fg}&|jRji|&|"ji|&nt9|D] \}}#||vrb|t%|ksJdtk|!}'|jTjitm|'jn|'||"ji|'mtWtY|#j7|#j/|#j5|#j=|#jqjr|tf|fg}&|jRji|&|"ji|& t;||"D]g\}(})|(jutjjvvs1tjjxj{|)jui|"Scc} wcc} wcc} wcc} wcc} w#1swYxYwcc}wcc}$wcc}%w)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_mutationcPt|t|k(sJg}t||D]{\}}t|tjrKt j |}|jtj||jdk|j|}|Sr+ ) rrrrrrF rU rrrr) tensor_boxesr, r- rfknew_tbs rr. z0WhileLoop.create.._require_exact_strides)s|$L(99 99ClL9 #Bb%,,/'@@DFJJ$::"BIIKu; JJrN) #*Jrr2rNr 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.ct|t|k(sJt||D]/\}}tjjj ||1yr)rrrwrrr) lhs_exprs rhs_exprslhsrhss r_guard_list_equalsz,WhileLoop.create.._guard_list_equals)sN9~Y777 #Iy 9r4 rW rwrrrrbrrF rS rrrr r rr r rr r rlrrrrr rrrr4r;rrrrmodulefx GraphModuler@rrKrrrPrrrr r( rrrrrrNr rQrsrrk r)*rcond_fnbody_fnrG rH rL rW r. fx_carried_inputsfx_additional_inputs fx_all_inputsrfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_ all_inputsr cond_outputs body_outputsrhrrr^borb r while_loop mutated_idxsmutated_idx_setrrmutated_inputs_iter all_outputsr\rPrOr mutated_inputrlrs* rrzWhileLoop.createt)s Q * A  :GG0055b9 ww3388<),@@ 2?@Q166%=@@6GHqvve}HH9M!NA!&&-!N!N9GHA3,,Q/HH#99/J0BUV>#?#?@C/E  8N$NN88/84 4}}..}}22 }}22  - XXdWeg  < A%3|3% O!23;;=EJJ. 1 1.qzz|$) ,1 ,):" G "A))+!!!?#s<'88  ;  8%S,%GH AKAxB </ </ < < r{{}bkkm < r}}  @==?bmmo5 J2r67J J5<<>R\\^3 @aR[ @3 A"!!!5 GG   & & GG % % ) )*=t D  *0!!$F3/%  }}(Z MM %(("6"6.   6 MM /  %\25DEc*S/EE#>2$&  &( # '1, ? , ))=)=)B)B5)IJ . V'%}}$llDJKKMRbk55b9RFLmmoV 7 7 ;V  C[M  "")))4""9- . )6 2 V/)^!44T6TT4$()<$=M//66&}';';]JW &&}5 +##)#4#4#6"("2"2"4!'!2#)#4#4#6#)#4#4#6#=#= #  !I&&--i8&&y1- 20NK8 @HC||~!5!55 ++// ? @UAH!NHOtF"SVs=(__! _&<_+8_0/A7_5 `:`*` 5_? c |j||j|j|j|jt |diyr )codegen_while_looprL r rsrKrrs rrzWhileLoop.codegenN*sA""4):):;88 MMOT\\749Lb+Q rct|ddx}rKttjjj |}|Jt |jSt Sr r r s rrz"WhileLoop.get_unbacked_symbol_defsT*r r)rG rrH rrI r rJ r rNrrrq rL rrrrX)rG rrr)rrrr) rg r rh r rG rrH rrL rrr& r r)rrrrrG rrH rI rJ rKr]rrrS rU rrrrrIrJs@rrF rF )s ].2N+215.5%)M?)%)M?),0G )0)(),) )  ) " )E)) )>2LLWWW) W , W  W #WWr  rrF c\eZdZ ddd dfdZdfd Zd dZxZS) rNr c t ||||||d|ddlm}||} | J| |_t j jj| d|_ |t j j| <y)Nr^ r) _get_effect) rnr]torch._higher_order_ops.effectsr effect_typerwr effectful_opsrprev_effect_buffer) rrNrYr$ r r rrr r rps rr]zEffectfulKernel.__init__`*s      /   @!&) &&&&"#''"7"7";";K"N-1k*rct|}|jF|jj t j |jj|Sr)rnr3r r6rrKrIrs)rrrps rr3zEffectfulKernel.get_read_writes}*sUg-/  " " .    ! !$$T%<%<%E%E%GH rcyr1rrs rrg z EffectfulKernel.has_side_effects*rrrrp rxr_)rrrr]r3rg rIrJs@rrr_*sr)-2HL222& 2 & 2 + 2&2E2 2:rrc<eZdZed d ddZddZy)r; ctSrr?r=s rr>z!NonTensorObj.get_free_symbol_uses*rrcyrrrs rrzNonTensorObj.realize*rrNrrr}rd)rrrrbr>rrrrr; r; *s1N+$)! !, rr; cLeZdZUded<ded<d dZd d dZddZdd Zdd Zy)rrr%FakeScriptObject | torch.ScriptObjectrc|jSrrrs rrszTorchBindObject.get_name*rrNc|jSrrrs rrz!TorchBindObject.codegen_reference*rrc|jSrrrs rr zTorchBindObject.get_value*rrct|jtjr |jS|jjSr)rrrr real_objrs r get_real_objzTorchBindObject.get_real_obj*s0 djj%"4"4 5:: ::&& &rc|j}t|ryt|dsJt|j }t j |d}|Dcgc]=}t|tjr!|j|jz?}}tjtj|dScc}w)Nr__obj_flatten__)r r,rrr rr rrrr rrroperatorr)rreal_script_obj flat_dict flat_elemsr flat_sizess r get_buf_byteszTorchBindObject.get_buf_bytes*s++- ? +(9:::88:; ((3A6   !U\\* NN qwwy (   j!<<  s AC rorrf)rr )rztorch.ScriptObjectr{) rrrrrsrr r r rrrrr*s& I 00' =rrc.eZdZUdZded<ddZdd dZy) OpaqueValueTypeConstantaIR node for opaque value type constants that appear directly in graph outputs. Unlike TorchBindObject (which references named constants loaded at runtime), this inlines the value's repr into the generated code since value types are reconstructed from their repr. rrc,t|jSr)rI rrs rrsz OpaqueValueTypeConstant.get_name*sDJJrNct|j\}}|jD]"\}}|tjj |<$|Sr)r+rrTrwropaque_value_type_classes)rrobj_repr opaque_typesrrs rrz)OpaqueValueTypeConstant.codegen_reference*sJ!4TZZ!@, &&( 5DAq34AGG - -a 0 5rrorrfrrrrrrsrrrrr r *s J rr c4eZdZUded<ded<ddZd d dZy) r rrrjrc|jSrrrs rrszGeneratorState.get_name*rrNc|jSrrrs rrz GeneratorState.codegen_reference*rrrorrf)rrrrrsrrrrr r *s I rr c8eZdZUdZded<ded<d dZd d dZy) r z Represents an opaque object (e.g., ProcessGroup) that is passed through as a graph input. Similar to GeneratorState, this wraps the object with its placeholder name so codegen can reference it properly. rrrrc|jSrrrs rrszOpaqueObjectState.get_name*rrNc|jSrrrs rrz#OpaqueObjectState.codegen_reference*rrrorrfr rrrr r *s I Jrr creZdZddZddZdd dZe d dZe d dZy) _CollectiveKernelcyrrrs rrz!_CollectiveKernel.should_allocate*rrcyr1rrs rrg z"_CollectiveKernel.has_side_effects*rrNc\t|jtjjusJd|j}|||_n|j j|_|j jDcgc]}|js|jc}|_ ycc}w)Nz,Setting cpp kernel needs a valid op_overload) rrrrrrrrrrr)rrrYrs rrz%_CollectiveKernel.set_cpp_kernel_name*sD$$%)>)>> : >!!  D #)>>#6#6D #NN44.  AFF. *. s B)B)c tjj5|j||g|i|\}}}}} ddd r J|d| d} D]c} t | t r| j tjj| j| T| j} e| Jd|d|t| ||} tj|} | jj| Dcgc]}tt| || c}| j j| Dcgc]}|jc}d|vr`| jj#tt| |d| | j j#|djyy#1swYxYwcc}wcc}w)NrgzIn-place collective zE requires at least one tensor argument; got only non-tensor IR nodes.rr)rwrr6r0 rr; rrrsrrr tree_leavesrrnrr r)rrYrjrr_example_outputr$ r% r rr tensor_argrc inpsrrls rcreate_inplacez _CollectiveKernel.create_inplace +sWW   D#""66CDCFC  !  D%E2C1D&EE$% 1J*l3    GG ' ' (;(;(= >~#..0  1! "6(+6 7 ! f %      !!&)&&OS T^Jf5sF C T !!T"Bc3<<>"BC F?  # # * *z8&-P     % %fUm&<&<&> ? I D D> U#CsG  GG Gc tjj5|j||g|i|\}}}}} ddd r J|d| D]#} t | t r| j %t tr|j||} | J|t| ||} t|D cgc](\} }t|j|| t| fg*c}} | _ t| j|D]T\}}tj s t#|r"tjj$j'|j(V| jS||j|||} tj s t#|s3tjj$j'| j(| g| _ | S#1swYxYwcc}} w)NrEr)rwrr6r0 rrrrr rrrrrKrrJrb rtr rr)rrYrjrrr+ r$ r% r rr rrc rrTrs rcreate_out_of_placez%_CollectiveKernel.create_out_of_placeU+sWW   D#""66CDCFC  !  D%F3D2E&FF$% %Jj/:""$ % nd +__[.AF% %%!0 F"+>!:  Av ((0AYKFN #6>>>B < V::BSCGG--11#((;  < >> !$$^4 F66>O?))--fkk:$XFNMe D D.sG#-G0#G-r_rr ) rYrrjzIRNode | list[IRNode]rrrrrr) rYrrjzTensorBox | list[TensorBox]rrrrrz%list[MultiOutput] | _CollectiveKernel) rrrrrg rrr r rrrr r *s  (0@0@&0@ 0@  0@  0@0@P99,9 9  9 / 99rr cPeZdZ ddd dfdZddZxZS)_AllReduce_KernelNr c Tt||||||d||jdy)Nr^ +aoti_torch_cpu__c10d_functional_all_reduce_rnr]rrw s rr]z_AllReduce_Kernel.__init__+=      /     !NOrc|jd|j|t|jtr|j |yyNz+torch/csrc/inductor/aoti_torch/c/shim_cpu.hinclude_extra_headerr rrNrr{ rs rrz_AllReduce_Kernel.codegen+@$$%RS,,T2 dkk6 *  % %g . +rrrp r r rJs@rr r +sw)-PHLPPP& P & P + P&PEP P,/rr cPeZdZ ddd dfdZddZxZS)_AllReduceKernelNr c Tt||||||d||jdy)Nr^ *aoti_torch_cpu__c10d_functional_all_reducer rw s rr]z_AllReduceKernel.__init__+s=      /     !MNrc|jd|j|t|jtr|j |yyr r rs rrz_AllReduceKernel.codegen+r rrrp r r rJs@rr r +sw)-OHLOOO& O & O + O&OEO O,/rr cveZdZ ddd d fdZd dZd dZed dZd fd ZxZ S) _WaitKernelNr c Tt||||||d||jdy)Nr^ +aoti_torch_cpu__c10d_functional_wait_tensorr rw s rr]z_WaitKernel.__init__+r rc|jd|j|t|jtr|j |yyr r rs rrz_WaitKernel.codegen+r rcr|jd}t|tsJt|tr2|jd}t|tsJt ||gSt|t rC|jd}t|tr"|j d\}}|j|gSgSgSr)rjrrr rrr)rrlrcollrrs rget_volatile_readsz_WaitKernel.get_volatile_reads+skk!n#v&&& c, - 1 Aa( 1$q' 1(3J [ )::a=D$ 12Q3 C())IIrcrtjj5|j||\}}}}}dddr J|d||t |j |}|j jtt |j ||y#1swYzxYw)Nrgr) rwrr6r0 rrrrr) rrYrlr r$ r% r rrc s r create_waitz_WaitKernel.create_wait,s WW   0""63/  !  0%E2C1D&EE$ cnn. /       && :S^^-=>V L ! 0 0s B--B6ct|}|j}|D]>}|jj t j |j@|Sr)rnr3r r6rrKrIrs)rrvolatile_readsvrrps rr3z_WaitKernel.get_read_writes,sZg-/ 002  GB    ! !,"6"6r{{}"E F Grrrp r r )rYrrlrrrrx) rrrr]rr rr r3rIrJs@rr r +s)-PHLPPP& P & P + P&PEP P,/2  *rr c2t|ttfr t|St|tt fr2t tj}|D]}|t|z}|St|tjr t|St Sr) rr=r r6rrr@rr"r rrrrrs rr r ),s!h%&$Q'' At} % u|| $ & 0A ,Q/ /A 0 Au|| $$Q''|rc2t|ttfr t|St|tt fr2t tj}|D]}|t|z}|St|tjr t|St Sr) rr=r r5rrr@rr"r rrr s rr r 9,s~!h%&A At} % u|| $ & 'A #A& &A ' Au|| $A|rct|trt|jtrt|jjtr'|jjj d|yt|jjt rU|jjj d|t|jjtr_t|jjjtr1|jjjj d|yt|jjtr|jjjsft|jjjdt r4|jjjdj d|yyyyyyy)Nr~r) rrrMrqrrrrrrrrj)rNrs rassign_origin_noder I,sV&)$FKK)L fkk&& . KK   / / q A  ((& 1 KK   / / q A&++**N;    %%uA   %%88J6;;++[9 ((00fkk..55a8&AKK$$++A.AA-QRSB1:2*M$r)rrrzTypeIs[int | Integer])rrrr)rrrrC)rrtrz&Callable[[Sequence[_T]], Sequence[_T]])rz&Callable[[Sequence[_U]], Sequence[_V]]rz&Callable[[Sequence[_T]], Sequence[_U]]rrr)rz#Sequence[int | torch.SymInt | Expr]rShapeEnv | Nonerrt)rSequence[int | Integer]rrtrr)rrrrrr)rrrrrrt)rr?rrrztorch.Tensor | None)rzSequence[_T] | NonerzSequence[_T | None] | None)rz/IRNode | OutputSpec | torch.device | None | strrre)rz"IRNode | torch.device | None | strrr)rzBuffer | TensorBoxr1rrr)r?rar@rarArarr)rTrrUzSequence[int | torch.SymInt]rr)rdr rr)rjrrrG)rzExpr | Sequence[Expr]rr\rrvrY)rRrrr\rQrrrD)rrarrbrQr rz$Callable[[Sequence[_IntLike]], Expr]rT )TFNFN)rrrrrrrrr%rr0rrztuple[StorageBox, Layout])rrrr rrrQ)rrarArarr)rr\rr)rrtrzTypeIs[Sequence[IRNode]])rr<rr)rr rzAbstractContextManager[None])rn r rr)r rrr)rr rz-tuple[list[ShapeAsConstantBuffer], list[Any]])rrrr)rNrrz torch.fx.Noderr(i __future__rrrrrrloggingr textwrapr|collections.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._loggingrtorch.fxtorch.utils._pytree_pytreertorch._dispatch.pythonr#torch._dynamo.utilsr$torch._export.serde.serializer%*torch._higher_order_ops.auto_functionalizer&torch._inductorr'rkr)"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|rr}r~rrr% __version__rr ImportErrorrrrrrrr rrrrr getLoggerrrrrrrrrr rrrrrrrrrrrrrrr*r,r<rErZrfrmrtrrrr r r%r6rergrrQrSr`rprrrrr rrrOrr rorrVrcrmr#rSrrrrrrrrrrPr rrrrrrrr rrr!rrrrrrrrPrimitiveInfoTyperrrr>rIrOrUrrrrrrrrrr r rrr r r r r rC rT r_ ri rn rv r r r r rrr rrs ry rrrr rg rj rrrqr r r r r" rF rr; rr r r r r r r r r r rrrr s" MM:    UT ''22..,,$$;(?M#2?M    /?+" NN-:0*)B&'95$%"(OY'''NJ t_ T] T] T]Dj) EkD()(**//%**2P2PP iPg!   8??4 8 yy~~'T k     c4[();(JKdR   i )d#$$ $F44,!$LP  , 9H   LP  , 9H   U U 27 +/  :?  26  D  >6 > >;('7*    $$G $G )$G $GN';Z,Z,z UUUUpP FP P f&% % % P i  F | $ y ! y !   u  u =)<8JNJNJN +JNBFJNJNZZ  Z  Z  ~)-1:  % * &"8D>8D>"BH"LM YM7S97Stk'k\s 1s l` +` FF 5F F T   [5[[|( !3748 9 9 991 9  9 2 99x 4 $^ v^ ^ B[[[|4(44nA9(A9A9H!(!!H^;^^B ShSSl&&&RlAlA^6"K|KK$ S| S S '9 (< 7 7 p Zp p f C&C[HV[H|!Gf!GHT%{%D.V*V*rUEV]EEPU &fi & & & K [ &6( F   Ut4_t4t4n d#$f+_f+R<><~%K$&,tC#I4E4L/MMLL^"|" \B.\B~5N50((6N>Z%NZ%z5 .55 UR?RRj O9Od ;;!Ut<ttnU"l""J / ((" ("V V B=L=@% % P - $< \< ~ulup ))|))Z/,/d<8B5B$8-846)l6)r))|))X;3;3|1E1Eh51|51p-L-8<;<<;~U A &A HU .  @  2 I ,I X   F 4K 4 2f>fV TTTn + +%%DU'v'' LUB.\B.B.JUe ,e e P "  "2 "U@ @ @ F )n)X6"=l"="=Jl(\ "``F/)/>/(/>S#Sr   Tg]NJs:n n$#n$