Lj UddlmZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl m Z m Z mZmZmZddlmZmZddlmZddlmZddlmZmZmZmZmZmZmZmZmZm Z m!Z!ddl"m#Z#m$Z$m%Z%m&Z&m'Z'm(Z(dd l)m*Z*ddl+Z+dd l+m,Z,m-Z-m.Z.ddl/m0cm1cm2Z3ddl4m5cm6Z7ddl8Z9ddl:Z9ddl;m6cmm?Z?dd l@mAZAdd lBmCZCddlDmEZEddlFmGZGddlHmIZIddlJmKZKmLZLddlMmNZNmOZOmPZPmQZQmRZRmSZSddlTmUZUmVZVmWZWmXZXmYZYmZZZm[Z[m\Z\m]Z]m^Z^m_Z_ddl`maZaddlbmcZcddldmeZeddlfmgZgmhZhmiZimjZjddlkmlZlddlmmnZnmoZoddlpmqZqmrZrmsZsmtZtmuZuddlomvZvmwZwmxZxmyZymzZzm{Z{ddl|m}Z}ddl~mZmZmZmZddlmZdd lmZmZdd!l6mZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZdd"lmZmZmZer+dd#lTmZdd$l`mZdd%lmZdd&lmZdd'lmZdd(lmZdd)l6mZneZd*ed+< ddlZej`Zd,Ze&d.Ze d/Ze d0Ze d1Zee,zZd*ed2<eeze,zZd*ed3<e9jxjze9jxj|zZd*ed4<ejeZej*e jd56Ze9j>jZ e!ed7eed7fd8d9ee!eeed7fd7d8d9fdzfZd*ed:<d;d;Zejd,<Gd=d>ZʐddAZ͐d>dBZ d?dCZgdDZgdEZ d@ dAdFZҐdBdGZ d@ dAdHZedCdDdIZe dC dEdJZ dC dFdKZ dGdLZ dHdMZאdIdNZؐdIdOZِdJdPZ dKdQZ dLdRZܐdMdSZݐdNdTZdUZGdVd9Zed-<GdWdXZeGdYdZeZdOd[ZeGd\d]eZeGd^d_eZed`edaedbedceddeddededfZdgedh< dP dQdiZeGdjdkeZde-df dRdlZe ee,ee,gefZd*edm<GdndoeZGdpdqeZGdrdseZeGdtdueZeGdvdweZeGdxdyeZdSdzZdSd{Z dT dUd|Z dVd}ZdWd~ZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZeGddeZGddeZeGddeZeGddeZeGddeZ dXdZdYdZGddZeGddeZGddeZGddeZGddeZGddeZ GddeZ eGddeZ GddeZ ed-<GddeerZ ed-<Gdde eZGdde ZGddeZGddeZeGddeZeGddeZed-<GddeZejd,<GddZGddeZGddeZeezezezeeezezezzZGdd«ZGdÄdĐeZGdńdƐeZGdDŽdȐeZGdɄdʐeZGd˄d̐eZ Gd̈́dΐeZ! dZdτZ"ed-<GdЄdѐeZ#Gd҄dӐe#Z$GdԄdՐe$Z%ed-<Gdքdאe#Z&ed-<Gd؄dِe&Z'Gdڄdېe'Z(Gd܄dݐe&Z)Gdބdߐe Z*Gdde&Z+Gdde+Z,Gdde+Z-Gdde&Z.Gdde&Z/Gdde&Z0Gdde&Z1Gdde&Z2Gdde2Z3Gdde)Z4Gdde&Z5Gdde&Z6Gdde'Z7Gdde&Z8Gdde&Z9Gdde&Z:Gdde&Z;ed-<GddZ<Gdde)Z=ed-<Gdde=Z>Gdd e=Z?eGd d eZ@Gd d e&ZAGddeAZBGddeAZCd[dZDGdde=ZEejGddeZFGdd7eFZGGddeFZHed-<GddeZId\dZJed-<Gdde&ZKed-<Gdd e&ZL d]d!ZMed-<Gd"d#e&ZNGd$d%e=ZOGd&d'eZPeGd(d)ePZQeGd*d+ePZReGd,d-ePZSeGd.d/ePZTGd0d1e=ZUGd2d3eUZVGd4d5eUZWGd6d7eUZXd^d8ZYd^d9ZZd_d:Z[y#e$rdZd-ZY KwxYw(`) annotationsN)Callable GeneratorIterableIteratorSequence)AbstractContextManager nullcontext)Enum)partial) AnycastClassVarLiteraloverload SupportsFloat SupportsInt TYPE_CHECKING TypeAliasTypeVarUnion) assert_neverNeveroverride ParamSpecSelfTypeIs)patch)ExprIntegerSymbol)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)CleanDivFloorDivModModularIndexing)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)IndentedBufferrrwTF_P_T_U_V_IntLike_NumLike _OpOverloadsz prefix TensorBoxr!IRNode _NodeOrNodesc.t|ttfSN) isinstanceintr xs Q/var/www/html/ia-enhanced/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_inputsz0dict[str, IRNode | sympy.Expr | TorchBindObject] input_nodes list[IRNode] output_nodeszdict[str, bool]input_deallocationboolskip_cudagraph list[str]constant_namesN__name__ __module__ __qualname____annotations__rrrrs/,+BA('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 AssertScalarrsympylogicboolalgBooleanrrEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputtype)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)getattrrq)argskwargsnames rfnzops_wrapper..fns!wsD!42622r)robjectrrrrr)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..reindexs?5zS^+++-23u:->?il#???Ar Sequence[_T]rr)rziprr)orderrrs @rinverse_reorderrs*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..reindex.s((r)rrrz Sequence[_V]r)rrrs`` rfuse_reindexingr*s) Nr)rrE)rrrrEc`|td|Dr t|}|St||}|S)z1 Convert strides to fill order (argsort) c3\K|]$}t|ttjf&ywrrrrr .0ss r z!get_fill_order..>s Q 1sEMM.B CQ*,)allr[r\)seq shape_env sorted_idxs rget_fill_orderr8s< 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_orderrFsR(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_orderrRsV !/sI >JCHo &1 &C &Z(4D  J 's Acyrrrreplace_symbols_with_hintss rir_node_to_tensorr_sRUrcyrrrs rrrcsrc|y|r%tjjj}nt}|j Dcgc] }|| }}t |r.|jjDcgc] }|| }}ntj|}|j}|j}t|}t|}tjjjj5t!j"||||j%}ddd|Scc}wcc}w#1swYSxYw)N)sizestridedtypedevice)rsgraphsizevarsoptimization_hintr"get_sizeis_storage_and_layout get_layoutrFlexibleLayoutcontiguous_strides get_dtype get_devicerarsuppress_guardstorch empty_stridedzero_) rrshape_fnrrrrrts rrris 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_typer rrrs rrrsq!SQY Au|| $vv A + ,q||~..#A3ba)9)9(:!<=rct|}|dvrtt|ddk(ryy| t|x}yddlm}t |tsJt|t||S)N)cpucudaxpu_backendtritonTFrE)TritonScheduling) rrrFrJcodegen.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)rEqrB)rr alignments rrz-is_aligned_realized_tensor..s# FQ%((3q)$a ( Fs*-rE)rrmaybe_get_strider4 get_striderrAndOrr,Lersrrguard_or_false)rr-aligned_stridesaligned_last_dim is_aligneds ` ris_aligned_realized_tensorr8s 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 . rEFT) rrrsrrstatically_known_leqr4rr,)strides1strides2shapedims1s2s rsignificant_strides_equalrAs 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||3ywrrsrrstatically_known_equalsrr?r@s rrz2try_match_insignificant_strides..s3  B 00R8 79rEdatalayout)rrrr0rAras_storage_and_layoutrrrsrrr: FixedLayoutrrroffset is_pinnedrReinterpretView)tensorstridesstorage old_layout new_striderr new_layouts rtry_match_insignificant_stridesrVs ! (   '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_stridesrbsa((%%%215K#K$4$454Q4K/0J"2& 4s A$ct}|D]>}|t|jdz}|t|jdz}@t |S)NF unbacked_only)r>r'rr0r)inputssym_varsinps rget_symbolic_inputsrisW!+HL$S\\^5II$S^^%5UKKL >rct|tr |j}t|tr|j }t|t r |j}t|t r|jSdSr)rrrIBaseView unwrap_view StorageBoxBufferget_namers r try_get_namerpsW!Y FF!X MMO!Z FF%a01::<:d:rceZdZUdZeZded<dZded<ejdZ d ed <ejdZ d ed <ejdZ d ed<ejdZ ded<ejdZded<eej"dPdZeej" dQdZedRdZdSdZdTdZdUdZdVdZdWdZdXdZdYdZdZdZdVdZd[d\dZ d] d^d Zd_d!Z d`d"Z!dad#Z"dbd$Z#dcd%Z$ddd&Z%ded'Z&dfd(Z'dgd)Z(e)dhd*Z*did+Z+ded,Z,djd-Z-dkdld.Z.dmd/Z/dnd0Z0ded1Z1dod2Z2dpd3Z3dqd4Z4dgd5Z5drd6Z6djd7Z7ded8Z8dkdsd9Z9dtd:Z:dUd;Z;dSd<ZZ>dwd?Z?dxd@Z@ du dydAZAdzdBZBd{dCZCd|dDZDd}dEZE du d~dFZFdjdGZGdfdHZHdedIZIdedJZJddKZKddLZLdrdMZMddNZNeOr e)d_dOZPyy)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)rrr)rvolds rcurrent_originszIRNode.current_origins:s7%%"%- * &)F #cF #sA2 A ?Ac#Ktj}|t_ d|t_y#|t_wxYwwr)rrs)r}rs rcurrent_stream_idxzIRNode.current_stream_idxDs3 ((%/" - ),F &F &s?/ ? <?cLt|ttttt fSr)rComputedBuffer InputsKernel InputBufferrOTemplateBuffer)rs ris_realized_nodezIRNode.is_realized_nodePs&     rc,tj|Sr)rcreateselfs rwrap_for_loweringzIRNode.wrap_for_lowering]s%%rc2tj|||yr)r __setattr__)rattrrs r_post_init_setattrzIRNode._post_init_setattr`s 4u-rc>t|j}|jd||jdtjrt j nd|jdd|jdi|jd|jy)Nrvrxrzrr})r>rrrrFdebug_ir_tracebackrx format_stackrs)rrvs r __post_init__zIRNode.__post_init__fsT223  73  V5N5N//1TX   t4  r2  d.F.FGrcBtd|jDS)Nc34K|]}|jywrrrdeps rrz(IRNode.get_read_names..r?s#((?r> get_readsrs rget_read_nameszIRNode.get_read_namesq?dnn.>???rc|jSr)rxrs r get_tracebackzIRNode.get_tracebackts ~~rc|jSrrzrs rget_origin_nodezIRNode.get_origin_nodewrcyrrrs rget_defining_opzIRNode.get_defining_opzrcgS)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>rvr ExternKernelrrzhasattrraddr _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrvrzrpre_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,rcdt|dd}|rt|dkDr|ddd}|js|gSg}|jD]8}|jd||j dz }|jd :|g|zS) Nzorigins=rv@=z...zstack_traces = { })rrrappendsplit)rshortenrvstack_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 rr zIRNode.get_dtype zzrcB |jS#t$rYywxYwr)r NotImplementedErrorrs rmaybe_get_dtypezIRNode.maybe_get_dtypes& >># #"    c2tdt|d)Nz#get_layout() is not implemented by !rrrs rrzIRNode.get_layouts!$GT |ST"UVVrcB |jS#t$rYywxYwr)rrrs rmaybe_get_layoutzIRNode.maybe_get_layout& ??$ $"  rc"|jSr)rrs 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 rrzIRNode.get_sizes!$Ed4j\QR"STTrcB |jS#t$rYywxYwr)rrrs rmaybe_get_sizezIRNode.maybe_get_size% ==? ""  rc"|jSrrrs rr=z IRNode.shapes}}rc4t|jSr)rnrrs r get_numelzIRNode.get_numelsT]]_--rctjjjt j |j dSNrrsrrstatically_known_truerr,rrs 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_references!$=d4j\"JKKrcyrrrs rr zIRNode.get_devicerrc.|j}|J|Sr)r rrs rget_device_or_errorzIRNode.get_device_or_errors"!!! rcyNFrrs rhas_exceeded_max_readszIRNode.has_exceeded_max_reads rc>tt|jrrrrrs r make_loaderzIRNode.make_loader!$t*"5"566rc>tt|jrrrs r make_indexerzIRNode.make_indexerrrc>tt|jrrrs rr0zIRNode.get_striderrcB |jS#t$rYywxYwr)r0rrs rr/zIRNode.maybe_get_striderrc>tt|jrrrs rrozIRNode.get_namerrcB |jS#t$rYywxYwr)rorrs rmaybe_get_namezIRNode.maybe_get_name!rrcv |jtjjvS#t$rYywxYwr)rorsr graph_inputsrrs ris_input_bufferzIRNode.is_input_buffer's4 ==?agg&:&:: :"  s ), 88cyrrr thresholds rhas_large_inner_fnzIRNode.has_large_inner_fn-rrcyrrruserss r mark_reusezIRNode.mark_reuse0 rcyrrrs r realize_hintzIRNode.realize_hint3rrc>tt|jrrrs rrlzIRNode.unwrap_view6rrc>tt|jrrrs r freeze_layoutzIRNode.freeze_layout9rrc>tt|jrrrr allow_paddings rfreeze_layout_with_stride_orderz&IRNode.freeze_layout_with_stride_order<"$t*"5"566rc>tt|jrrrrs rfreeze_layout_with_fill_orderz$IRNode.freeze_layout_with_fill_orderArrc>tt|jrrrrs rfreeze_layout_with_same_orderz$IRNode.freeze_layout_with_same_orderDrrc>tt|jrrr exact_stridesrs r freeze_layout_with_exact_stridesz'IRNode.freeze_layout_with_exact_stridesGr rc>tt|jrrrs rget_read_writeszIRNode.get_read_writesLrrc6|jjSrr,readsrs rrzIRNode.get_readsO##%+++rc4t|jSr)rrrs r num_readszIRNode.num_readsRs4>>#$$rc>tt|jrrrs rget_storage_numelzIRNode.get_storage_numelUrrc>tt|jrrrres rget_free_symbol_useszIRNode.get_free_symbol_usesXr rc>tt|jrrrs rget_reduction_typezIRNode.get_reduction_type]rrc>tt|jrrrs rget_reduction_sizezIRNode.get_reduction_size`rrcyrrrs r is_externzIRNode.is_externcrrcyrrrs ris_no_opzIRNode.is_no_opfrrc>tt|jrrrs rconstant_to_devicezIRNode.constant_to_deviceirrc>tt|jrrrs rget_mutation_nameszIRNode.get_mutation_nameslrrc>tt|jrrrs rget_operation_namezIRNode.get_operation_nameorrc>tt|jrrrs rget_inputs_that_alias_outputz#IRNode.get_inputs_that_alias_outputrrrcyrrrs rrz IRNode.dtypews(+r)rvzOrderedSet[Node]rGenerator[None, None, None])r}r|rrIrrrrrr)rrrr rrrrrOrderedSet[str])rrwrryrzOperation | Nonerzlist[Subgraph]T)rrr Sequence[str])TT)rzSequence[object]rrrrrrr torch.dtype)rztorch.dtype | Nonerr)rz Layout | Nonerr)rzOutputSpec | NonerrrSequence[Expr])rSequence[_IntLike] | None)rz)_IntLike | sympy.Rel | Sequence[_IntLike]rrr str | NonerrzIndentedBuffer | Nonerrrtorch.device | Noner torch.devicer$Callable[[Sequence[Expr]], OpsValue]r Callable[[Sequence[Expr]], Expr]rSequence[_IntLike]rrrr|rrrrrrFr Sequence[int]rrrrrrorrrrirrr)rirrrrrdependencies.ReadWritesrzOrderedSet[Dep]rrrrrerrrrrcrrrrS)Qrrr__doc__r>rrrrs dataclassesfieldrvrxrzrr} staticmethod contextlibcontextmanagerrrrrrrrrrrrrrrr rrrrrrrrpropertyr=rrrrr rrrrr0r/ror r rrrrlrrr#r&r*r,rr2r4r7r9r;r=r?rArCrErGrrrrrrr%s3=,/>04-4 1{00e.rrrrs rrzOperation.get_read_namesrrc6|jjSrr.rs rrzOperation.get_readsr0rctrrrs 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=r6s rr7zOperation.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 rNrLr`rO)rru)rr|rj)rr{)rr{rrrXrs)rrrrrMrur list[Buffer]rrrmrxrv)rrrrr rrrrErrr=r?r,rrrrrr7rrrrrr{sr2" #$'"-@,"%*! !0rrcFeZdZUded<ded<ded<ded<ed d d d Zd!d Zd"d ZeZd#d Z d$d Z d%dZ d%dZ e d&dZeej"fd'dZed(dZd)dZed"dZd*d+dZdd,dZd-dZd.dZd/dZd0dZd%dZd1dZd2dZy)3LoopsrcrrUrCallable[..., Any]inner_fnrirangesctjgfd|jD|jS)Nc36K|]}t|ywrr&reres rrz-Loops.get_free_symbol_uses..s FQq-0 F)r>unionrinner_fn_free_symbolsr6s `rr7zLoops.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_strrrz)rnamesrs r_to_strz Loops._to_strsDKK$$%Q'DJJ!!#  r read_buffersrs rrzLoops.get_read_namesHs$//1>>??rcHt|jjSr)rrrrs rr2zLoops.num_readsKs4((*7788rc2tdt|d)Nz+get_reduction_size() is not implemented by rrrs rr;zLoops.get_reduction_sizeN!9$t*Q G  rc2tdt|d)Nz+get_reduction_type() is not implemented by rrrs rr9zLoops.get_reduction_typeSrrc2tdt|d)Nz+constant_to_device() is not implemented by rrrs rrAzLoops.constant_to_deviceXrrrmrx)rrSrrrjr`rOrY)rr rr rr)rrirrDrrZ)rrUrzSequence[Sequence[_IntLike]]rrkrerrOrderedSet[Symbol])rr!rrNrurMrvr]ry) rrrrr^r7rr__repr__r rrr classmethodrr~rDINDEXrr]rrrrrrrrr2r;r9rArrrrrs    G$$) ! ! %   )H  # #:>**  ((+  ; W  @9   rrc|jrtjtd|Stjd|S)Nnanr)is_floating_pointrqconstantfloat)rrs r nop_loader_fnr^s1 ||E%L%00||Au%%rcPeZdZddZd dZeZd dZd dZ d dZd dZ y) Pointwisecp|jrtt|jS|jSNr)rr rrrrs rrzPointwise.make_loadergs)  "= ; ;}}rc$|jdSrrrs rrzPointwise.__str__nrrcgSrrrs rr;zPointwise.get_reduction_sizess rcyrrrs rr9zPointwise.get_reduction_typevrrcp|j}tj|xsd||||SNunnamed)rrqstorer output_nameindexervarsloaders r store_outputzPointwise.store_outputys2 !!#yy1 74=&,OOrc|j}tjtd||}t ||j ||j SFMove this to a given device. Requires that all reads are to constants.override_devicerrrr)rrrConstantBufferrrrrrrs rrAzPointwise.constant_to_devicesK!!#Hn.?HP**;;   rNrdrj)rzSequence[sympy.Expr]r])rr^r!Callable[[Sequence[Expr]], Never]rrZrrry) rrrrrrr;r9rrArrrrresR)HPP3P P  P  rrcFeZdZUded<dZded<ddZ d dZy) Scatterrgoutput_indexerNrW scatter_modec|j}tjtd||}t ||j ||j |j|jS)rr)rrrrrr) rrrrrrrrrrs rrAzScatter.constant_to_devices]!!#Hn.?HP**;;..**   rc|j}|d}tj|||j||||jS)Nr )mode)rrqr rrr s rrzScatter.store_outputsT !!#  #Kyy  D''- . 4L""   rry)rr^rrrrZrr )rrrrrrArrrrrrsB44"L)"    3    rr logical_ormaximumminimummulr bitwise_xor)anyrminprodsumdotxor_sumz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNctvr tSdvr dfd }|Sdk(r dd}|Std)Nargmaxargminc|\}}|\}} dk(rtj||}ntj||}tj||}t rtj ||}tj ||} tj |tj|| }tj |tj|| } rtj||ntj||} tj |tj|| }tj|||tj|||fS)Nr1) rqltgteqr.ner" logical_andwhere)aba_valuea_indexb_valueb_indexmaskequala_isnanb_isnantiearg_break_ties_leftrreduction_types rargmax_combine_fnz3get_reduction_combine_fn..argmax_combine_fns ! GW GW)vvgw/vvgw/FF7G,Ee$&&'2&&'2~~dCFF7G,DEucoogw.OP'w(VVGW-  >>$s(CDD $1 $1 rwelford_combinecl|\}}}|\}}}||z }||z} || z } ||| zz||z||z|z| zz| fSrr) r9r:a_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_ws rwelford_combine_fnz4get_reduction_combine_fn..welford_combine_fnsm&' "FD(%& "FD(VOE!H,J :-I**t eemh6BB rzunknown reduction_type=)r9tuple[object, object]r:rSrtuple[OpsValue, OpsValue])r9#tuple[OpsValue, OpsValue, OpsValue]r:rUrrU)r-r)rErrDrFrRs``` rget_reduction_combine_fnrVs--#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 Reductionrireduction_rangesrVrErU src_dtyperZreduction_hintc$|jdS)N)rrYrErrs rrzReduction.__str__s||LMMrczt|tjfd|jDzS)Nc36K|]}t|ywrr&rs rrz1Reduction.get_free_symbol_uses.. s PQq-0 Pr)superr7r>rrYrre __class__s `rr7zReduction.get_free_symbol_usess:w+M:=OZ\=O=O P$:O:O P>   rc|jSr)rYrs rr;zReduction.get_reduction_sizes$$$rc|jSr)rErs rr9zReduction.get_reduction_type"""rc tj|j|j|j|j ||}tj |xsd|||yr )rq reductionrrZrErstore_reduction)rrrrreduction_varsrs rrgzReduction.store_reductionsT JJ NN    MM$ /   K49gdmUKrcXt|jt|jzSr)rrrYrs r index_lengthzReduction.index_length#s!4;;#d&;&;"<<.inner_reduction_splitsosrrzUse 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: %sr.rr*rpc 2 |j}|Jtdt||j|j |}|j }|j J|j Dcgc].}t|trt|tjs|0}}g}d}t|jdD] t fd|Ds|j j j t"j$j&vsZt"j$j& j }t)|j*dd}|j-t)|j*dd|k7sd}||fScc}w) NrrrrrJrIFc|jSrrrs rz@Reduction.num_splits..get_read_indices..s affrkeyc3NK|]}|jjvywr)rr3)rrmds rrzAReduction.num_splits..get_read_indices..s FaqBHH111Fs"%rT)r rr r rr, range_varsrrrNumbersortedr/rrrrrsrname_to_bufferrrJ decide_layout) rrcb read_writesrindiceschangedbuforiginal_striders @rget_read_indicesz.Reduction.num_splits..get_read_indicessh\\^F% %%%!++- B,,.K))5 55%//a&z!U\\/JJ GG[..4DE +F:FFNN288,ww!''"8"88gg44RWW=*1#**h*M))+"3::x>/Q&*G +G# #!s83Frc34K|]}|dk(xs|dkDyw)rrENrrs rrz'Reduction.num_splits..s9AQ!a%9r)rwrrxrrr)rrXrztuple[Sequence[Expr], bool])*rnrsrrall_unbacked_explicitly_hintedrZrqr has_featurerHREDUCE_TO_SINGLE_ELEMENTrFsplit_reductionsrYrmulti_processor_count functoolsr choicesreduction_split_factorINNERrrrrr rO!replace_backed_symbols_with_hintslogrrXrrGindex_vars_squeezerr;simplify_with_ranges stride_hintsrkeysrOUTER)#r dst_dtyperZrrrYrEreduction_numel input_nodeexprsrwrx should_splitpropsnum_smmin_elements_per_threadryouter_reduction_splitsr new_rangesnew_reduction_rangesextracted_numel_hintrrrrrrhranges1 num_outer num_innerrjrQouters# r num_splitszReduction.num_splitsAs#!-"78ww>>uE ((!+ + ww//AA/RWW%%77 f8MN %/ ##FN,S,ST T ( ( ''  U " ((!+ + ''/,,"$ @I@Q@Q 00&$A "AJ@Q@Q 00&%A "  &)    &< " ?*+?LEz$**E11%*Z*K\\.2BDIH5(00 ! $F,A. )!,JGQ w<1  ((!+ +'3'F'F JJL!..0( $NW   A  55aAAgg&&33>4 #7G 999EQ Q   y  &&(>$j) !&&(>$j) YHHs -P  Pctjjjt ||dfd |dvr1t t j dfd fdSS)z1Convert inner_fn from a reduction to an pointwisec tjfdtjDcgc] }t |c}DScc}w)Nc30K|] }|ywrr)rrmrvalue_fns rrz=Reduction._unroll_reduction_fn..fn..s UF+)rreduce itertoolsproductr)rr combine_fnrYrs` rrz*Reduction._unroll_reduction_fn..fnsH##"+"3"3,<=q%(=#  >sA r1r0c|Dcgc]}tj|}}||tj|tj fScc}wr)rexpandrq index_exprrint64)rrmr flatten_indexrs rrz0Reduction._unroll_reduction_fn..value_fn sP4::a%,,q/::UF+NN=#8%++F;sAc|dSrr)rrs rr~z0Reduction._unroll_reduction_fn..sE1r)rrirr )rrirmrirrT)rsrr guard_int_seqrV_fixed_indexerr r )rrYrErZrrrrs`` @@@@r_unroll_reduction_fnzReduction._unroll_reduction_fns77++99:JK-niH   1 1* 112BCM  ) 3E * . -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) rrrrrrrrr)valrs rpy_cnstz!Reduction.create..py_cnst2sg *9$00%c=9D49D9 :%%c;7BcB7s8OrrE)r*r,r)r'z* not supported for zero-dimension tensors!c6tjSrrqr)rrrErtypes_to_initss rconst_fnz"Reduction.create..const_fnHs||ON$CYOOrrrc0tjdSrr)rrs rrzReduction.create..fnVs<<955rcnDcgc]}tjj}}||Scc}wrrrr)rrreduction_indexrrYs rrzReduction.create..fn[s1=M&Nuww||&NO&N#E?;;'O!2r+c`tr|S|dkDrt|tjS|Sr)rrrF min_num_split)rrs r_maybe_increase_splitz/Reduction.create.._maybe_increase_splits//* qy5&"6"677 rr.NrTc:|j}t|dk7rytt|}|tj j vrytj j |}t|tsy|jjJ|Sr) rrnextiterrsrrrrrIr9)cur_node read_namesbufnamers r_find_split_reductionz/Reduction.create.._find_split_reductions"*!8!8!:J:!+#"4 #34Gagg&<&<<#''009C%c>:#88668DDDJrrp)rrrzbool | float | int)rrrrr)rrrr)rrrzComputedBuffer | None)#rsrrsimplifyrnrrrrr rrFunroll_reductions_thresholdrjrrrrZrqrrr rO!create_multilayer_existing_rangescreate_multilayerr!mix_order_reductionrrIrXrY _split_size_original_inner_fn_original_ranges_original_reduction_ranges)rrrrZrrrYrEr[rrrrhintrrrrrsplit_reductionrrrs ` ` `` @@rrzReduction.createsk$''**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)rr0z-infF)r(r1infTrrE)r*r)r+r,r'welford_reducerGonline_softmax_reduce)r.rr,riinfor(r)rErzeroones rdefault_accumulatorzReduction.default_accumulators . .e$V}$!%({{5)--- . .e$U|#!%({{5)---(/uQ&u-d1#T40 $dD1&+FmT%:     rc:|dk(rytj||S)Nrr)rXrrErs r default_valuezReduction.default_value s# - -,,^UCCrc|dk(r|S|dkr(|dkr#|tjk(rtjS|dkr(|dkr#|tjk(rtjS|S)Nr.i)rZr OUTER_TINY)rrxr[s r_multilayer_second_step_hintz&Reduction._multilayer_second_step_hintsg B;! ! C.wrapper_fn..bodyasi');<.wrapper_fnZs"1 _*/ 'Y ?2_DG = =-o> vvNN7K8NN?K@zz$g66v r)rSequence[Symbol]rrrrr) rViewdynamic_reshape_indexerrsrrrrr,rB) rrrYrrrrr dense_indexrrrs ` ` `` @@r_multilayer_wrap_loaderz!Reduction._multilayer_wrap_loaderEs>> Z .. / ((>> HHS%0! 4   # 6F    (rctdDs Jdtj|t|t|z dfd }|S)Nc3&K|] }|dk( yw)rENrrrs rrzDReduction._multilayer_wrap_loader_existing_ranges..ys3a163z8Only enabled for numel_hint == 1, found original_ranges=c |dt}|td}|t|t|zSr)rr) merged_indexnew_reduction_index original_idxrroriginal_rangesrs rrzEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fnsQ((>#o*>?L$S%9%;.intermediate_fns''A'A'AB BrNrp)rrirrirrr)rfloat16bfloat16rrXrrrrsrrrrnrrr)rrrrZrrrrrrErr[intermediate_dtype intermediaterrxrs @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)rrE)rnrArr r)rrrrZrrrYrErr[rrrrrs rrzReduction.create_multilayers&((89o;UC ##NI>00        ++       f e  L     rc j|j|||||} |j|||| ||g|||| d| S)rr.)rr) rrrrZrrrrrrEr[rs rrz+Reduction.create_multilayer_existing_rangessc$@@   %   ++      % +o + +     rrjrmrrYr]) rr^rrrrZrhrrrrvrzSequence[Sequence[Expr]]ryr)rrcrrUrZrUrzCallable[_P, OpsValue]rrirYrirEzReductionType | Literal['scan']rrr IRNode | Nonertuple[ReductionHint, _IntLike]) rzDD$/D &DD%(:G  &4A >%)('(-(" (  (  (/("( ((TD($2  &  0  J8= = =  = ' = ( = $2= = ,= &= = &=  = = ~%)+ + +  + % +  + )+ &+ + &+ "+  + + Z$ $ $  $ % $ ( $ $2$ "$ ,$ &$ &$  $ $ rrXcdfd }|S)1A closure containing math to read a given elementct|tk(sJt|tk(sJ}t|D]\}}}|dk7s |||zz}|Sr)rr)rresultrstszrMrrs rrz_fixed_indexer..indexer*ss!c%jCK&???5zSY&&&ufd3 +KCRQw#(* + r)rrorrr)rrrMrs``` rrr#s Nr INNER_FN_TYcleZdZUded< dfd Z ddZxZS)MultiOutputReductionr output_indexc trftdk(rd} n dfd } t | ||| |||||| |_y)NrErc2tfdDS)Nc30K|] }|ywrr)rrr reduction_idxs rrz@MultiOutputReduction.__init__..loader..SsHR]3Hr)r)rr; inner_fnss``rrz-MultiOutputReduction.__init__..loaderPsHiHHHrrp)rrZr;rZrztuple[OpsValue, ...])callablerr___init__r7) rrrr<rrYrErZr[r7rras ` 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 ) rqrfrrZrErrrrrr7rg)rrrrrhrrs rrgz$MultiOutputReduction.store_reductionas JJ NN    MM$ /   &5$-0>$v,>0t(()"";#;)WT]ERRr)rrcrrUr<z#INNER_FN_TY | Sequence[INNER_FN_TY]rr)rYr)rErVrZrUr[rZr7r) rr^rrrrZrhrrr )rrrrr>rgr,r-s@rr6r69s#)#)#)7 #) " #) , #)&#)#)&#)#)JSS3S S ) S Srr6c^eZdZeej df ddZy)OnlineSoftmaxReductionNc tfdt|D} | D]} | j| S)z> Create the reduction disregarding splitting. c3jK|]*}tjtd| ,yw)rN)rrr6) r output_idxrrrrr[rYrZs rrz0OnlineSoftmaxReduction.create..sI    $$+"   s03)rrr) rrrrZrrrY num_outputr[rresultsrs `````` ` rrzOnlineSoftmaxReduction.createtsG   $J/    A IIK r)rrcrrUrZrUrrrrZrYrZrErr[rZrr"rSequence[TensorBox])rrrrrZrqrrrrrArAss)6(=(=$(!!! ! % !  !)!!&!"! !!rrAceZdZeej 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)rrGzGwelford reduction usage is explicitly disabled, please check you configcXdfd }tj|tS)Nc0tjSrr)rrrs rrz8WelfordReduction.create..const..inner_fns||rrrrZrrrrrr)rrrrrs` rconstz&WelfordReduction.create..consts2  ##!F| $ rrrEcXdfd }tj|tS)NcnDcgc]}tjj}}||Scc}wrr)rrrrrYs rrz7WelfordReduction.create..copy..inner_fns1=M&Nuww||&NO&N!#77'OrrrMrN)rrrrrrYs` rcopyz%WelfordReduction.create..copys28!''!%< (rrc3.K|] }|ywrr)rrrRs rrz*WelfordReduction.create..s:"T"X:)rErr)rrrr)rr(rr)rFmtiadisable_welford_reductionrsrrrrnrrXrrZrqrrrrrIr)rrrr<rrYrEr[rrOmeanm2weightrrrDrFrrRs `` `` @rrzWelfordReduction.creates!FFFF;;88 U 8''**33MBR4ST  a 8DqB1XFV# # a  L   !11IaL)58U1X==: :::& **    aL  )++  e ]22 2!N 19((   2$Ah     $""     A IIK % s0-E8cy)N)rrrrrs rrzWelfordReduction.default_value 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) rrrGc0tj|Srr)rr;rrs rrz4WelfordReduction.create_multilayer..constant5 s||E511rrrE)rrr<rrYrErr[c 3LK|]}j|dyw)rrN)r )rrrrrrYrs rrz5WelfordReduction.create_multilayer..M s= ++$# , s!$c|g||Srr)rrrs rintermediate_loader_fnzBWelfordReduction.create_multilayer..intermediate_loader_fn` s 4E4O45 5rc3TK|]}t|j!yw))rN)r r)rrr`s rrz5WelfordReduction.create_multilayer..p s*.q}}GG%()rrZr;rZrrrrr)rrZrrZrrerrr)rnrsrrrrr,rBrr rArIrrrrr)rrrr<rrYrErr[rr intermediatesrrxrr`rs` ` ` ` @@@rrz"WelfordReduction.create_multilayer 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)rrcrrUr<Sequence[Callable[..., Any]]rr+rYr+rErVr[rZrrGr$)rrcrrUr<rdrr+rYr+rErVrrr[rZrrG) rrrrrZrqrr~rrrrrrIrIs)6(=(=yyy0 y  y ( y&y&y yyv$/ & \ \ \ 0 \  \ ( \ &\ \ &\  \ \ rrIcxeZdZUded<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)(Scanr+ scan_rangesr=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rrZr[rr7tuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]r<ct|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)r_r7r>rrgrr`s `rr7zScan.get_free_symbol_uses sa G ( 7 jl  Od>N>NO !jl  HdiiH   rct|jt|jzt|jk(sJt|yr)rrrgrr_rrras 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}DrTr ) rrr<rqrsrjrr r7)rrrr scan_varsrr1rs @rrgzScan.store_reduction soll4+DT^^DD$++t?yy  $9gclF4;L;L4M  rcy)Ncustomrrs rr9zScan.get_reduction_type src|jSr)rgrs rr;zScan.get_reduction_size rrc|jSrrrs rrz Scan.get_size yyrc|jSrrrs rrzScan.get_pointwise_size rrcXt|jt|jzSr)rrrgrs rrjzScan.index_length !4;;#d&6&6"777rc|j|j}|j|jtj}|j ||}|fSr)rrrgrDrlrrrrmrs rrzScan.inner_fn_args E DKK(T--t}}=ll5&)v rc|j|j}|j|jtj}|j ||}t |j||Sr)rrrgrDrlrrNrrrerrmrs 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) NrErr)rrraxispointwise_rangesrgr scan_numelz3.3.0ct|tk(sJt|tk(sJg|d||dSrr)r scan_indexrrrgs rrzScan.create..reindex Sz?c+&66 66u:%5!66 66>U5D\>J>tu> >r) rrrjrr<rrrgrrr[r7)rrZrrZrr*r)rsrrrHSCANrTUPLE_REDUCTIONrrrnrrr3rrrrrfrversionhip has_tritontriton_version SplitScanrr)rrrjr<rrrr[rrrrr7r scan_typesupports_splitrrFr1rrgs ` @@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)rr;rrs rrz#Scan.num_splits..wrapper_fn7 s*Fc%4jF=F3tu:FG Grrs)rrrZrrrYrEr)rrZr;rZrrr)rXr) rrrrrrrgrrrs `` rrzScan.num_splits* s; H###(!&$  rrmrrL) rr^rz%Callable[[Sequence[_IntLike]], Never]rrZrxrrr r]rYrvr)rrcrjrir<z+tuple[Callable[[Sequence[Expr]], Any], ...]rr+rrrrhr[rZrrrr rSequence[TensorBox | None])rrcrrUrrerrrr+rgr+rrhrrrr#)rrrrr^r7rrgr9r;rrrjrrrrZrqrrr,r-s@rrfrf| s MM SS!! ##--F#  $       7    $    8 U )6(=(=`&*``(`? `  `  `R`&`#`` $``D   7    (  # R   (  rrfc eZdZy)rN)rrrrrrrrG 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)%Sortr+ sort_rangesrz:Callable[[Sequence[Expr], Sequence[Expr]], Sequence[Expr]]rrZr[rr7rirjrkr<rstable descendingct|tjfd|jDztjfd|j DzS)Nc36K|]}t|ywrr&rs rrz,Sort.get_free_symbol_uses..b rnrc36K|]}t|ywrr&rs rrz,Sort.get_free_symbol_uses..e rpr)r_r7r>rrrr`s `rr7zSort.get_free_symbol_uses] sa G ( 7 jl  Od>N>NO !jl  HdiiH   rct|jt|jzt|jk(sJt|yr)rrrrr_rrrs rrzSort.__post_init__i rsrc,|j||tfd|jD}tj|j ||j |j}tj|xsd|||jS)Nc3.K|] }|ywrrrvs rrz'Sort.store_reduction..u rwrTr ) rrr<rqsortrjrrr r7)rrrrrhrr1rs @rrgzSort.store_reductionm sull40DT^^DD$++vt{{DOOLyy  $9gclF4;L;L4M  rcy)Nrrrs rr9zSort.get_reduction_type{ src|jSr)rrs rr;zSort.get_reduction_size~ rrc|jSrr}rs rrz Sort.get_size r~rc|jSrrrs rrzSort.get_pointwise_size rrcXt|jt|jzSr)rrrrs rrjzSort.index_length rrc|j|j}|j|jtj}|j ||}|fSr)rrrrDrlrrs rrzSort.inner_fn_args rrc|j|j}|j|jtj}|j ||}t |j||Sr)rrrrDrlrrNrrs 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)NrErrct|tk(sJt|tk(sJg|d||dSrr)r sort_indexrrrs rrzSort.create..reindex rr) rrrjrr<rrrrr[r7rr)rrZrrZrr*r)rsrrrHSORTrrrrnrFr!decompose_sort_opspersistent_reductionsrrr3rrrrrr)rrrjr<rrrrr[rr sort_numelis_persistent_kernel max_rblockr7rrFr1rrs ` @@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.9HrmrrL) rr^rrgrrZrhrZrr r]rYrvr!)rrcrjrir<z'tuple[Callable[[list[Expr]], Any], ...]rr+rrrrrrr[rZrr rr)rrrrr^r7rrgr9r;rrrjrrrrZrqrr,r-s@rrrL s>  GG!! ##-- LF#  $       2    '    8 U )6(=(=QQ(Q; Q  Q  QQQ&QQ $QQrrc> t|dy#t$rYywxYw)NFfreezeT)rKrrs rrr s&a. s  c t|d\}}|jr|j|jS#t$rYywxYwNFr)rKshould_pad_strides pad_strides is_contiguousr)r_bufferrJs r is_contiguous_storage_and_layoutr sR/%@  $ $ &    ##%% s?A A Act|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_orderrr)rr)rrrKrIrmrrnrrrr*rrOrJr) rrrrrr)rrJbuffers rrKrK sG!Y$ FF+%''   !Z ) FF+%''  6!&&##%%%!V !||~33555)11 2*22!3!!}alln,,!_%* FF qxx rcd t|d\}}|j|S#t$rYywxYwr)rKis_stride_orderedr)rrrrJs 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)rrrm is_unalignedrIrOrJrsrrstatically_known_multiple_ofrMrerrgrnrounaligned_buffers)rrJhas_unaligned_layouts rrrE s$J/0DII&&$(#$77#3#3#P#P MMN6<<8 8/$  DII&>*>>$}}!''";";;; rceZdZUded<eddddZddZddZddZe ddZ ddZ d d Z d!d Z d"d Zd#d Zd$d Zd%dZd&dZd'dZd(dZd%dZd%dZd)dZd*dZd+dZd,dZy)-rkrrIc8|jj|SrrIr7r6s rr7zBaseView.get_free_symbol_uses[ syy--m<.indexerf & &r)rrZrr)rIrr)rrrrs @@rrzBaseView.make_indexerb s/ &&(%%' 'rcl|jj|jdfd }|S)Nc |Srrrs rrz$BaseView.make_loader..loadero rrrM)rIrr)rrrrs @@rrzBaseView.make_loaderk s/ %%'%%' ' rc6|jjSr)rIr rs rrzBaseView.dtypet syy""$$rc6|jjSrrIrrs rrzBaseView.get_layoutx yy##%%rc6|jjSrrIr rs rr zBaseView.get_device{ rrcyrrrs rrzBaseView.get_origin_node~ rrc6|jjSrrIrors rrozBaseView.get_name yy!!##rc"|jSrrrs rrzBaseView.get_pointwise_size }}rc8|jj|SrrIrrs rrzBaseView.mark_reuse yy##E**rc6|jjSrrIrrs rrzBaseView.has_exceeded_max_reads yy//11rc6|jjSrrIrrs rrzBaseView.realize yy  ""rc8|jjyrrIrrs rrzBaseView.realize_hint s  rc6|jjSrrIr4rs rr4zBaseView.get_storage_numel yy**,,rc6|jjSrrIr=rs rr=zBaseView.is_extern yy""$$rct|jtsJt|j|jj Sr)rrIrkris_module_bufferrs rrzBaseView.is_module_buffer s6$))X.?TYY?.yy))++rc6|jjSrrIrrs rrzBaseView.get_read_names yy''))rctjtdd5t|j |j j cdddS#1swYyxYwr)rrr rPrrr/rs rrzBaseView.get_reads sL \\.*:D A &  " e    s 2AA!cd|}t|tr|j}t|tr|Sr)rrkrI)rrs rrlzBaseView.unwrap_view s+H%AH%rc|j}tjtd||}t ||j ||j Sr)rrrrrr rrs rrAzBaseView.constant_to_device sP!!#Hn.?HP.."==?   rNrmrr*Callable[[Sequence[Expr]], Sequence[Expr]]rfrdrTrVr`rOrjrYrlrXr]rLrwrMrurKry)rrrrr^r7rrrrrrr rrorrrrrr4r=rrrrlrArrrrkrkW s LJ'=(=C%%&&$+2#!-%,*  rrkcPeZdZUded<eddZed dZd dZ d dZ y) rrZrctjj}|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 sizesNr.rfallbackzBroadcast failed in ExpandView(, z) on dimension ) rsrrrrrrrris_size_one_or_falser6r) 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)NrrH)rIr)rrrKrrrrrrrrrsrrr rLrrrrMrNrOr) rrr rRrSskiprTrrrUs rrzExpandView.create s&&q(3  #"7": GZx=3z#77D19 9'',,$.J #J$5$5z G  !!77++@@F  %!!  X!!$$ J# C Cqx00rc|jSrr}rs rrzExpandView.get_size r~rc|j}|jjt|tz  dfd }|S)Nct|d}t|tk(sJttD](}|dk(s tjj ||<*|Sr)rrrrrr)rractualrs rrz*ExpandView.make_reindexer..reindex sftu&Eu:V, ,,3v;' ,!9>$ww||E!H ,LrrrZrrZ)rrIr)rtargetrrrs @@rrzExpandView.make_reindexer sP##%6{S[( !   rN)rrr rirri)rrr rirrkrYr) rrrrr~rrrrrrrrrr sA %%N114 3rrcPeZdZUded<eddZed dZd dZ d dZy) PermuteViewr*dimsc |j|}t|ttt|k(sJt |rt |\}}t |j|j|Dcgc]}|j|c}|Dcgc]}|j|c}|j|j}t||St||Scc}wcc}w)NrH)rIr) _map_neg_dimsr>rrrrKrLrrrrrMrNrOr)rrrrRrSrrUs 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)rrr>s rrzPermuteView._map_neg_dims/ s)@DEsaxSY_4EEEs$ct|j|jttt |jk(sJ|j j }|jDcgc]}|| c}Scc}wr)r>rrrrrIr)rrrs rrzPermuteView.get_size3 so$,,TYY78J #dii. !=    yy!!#!%+AQ+++s7 BcLt|jDcic]\}}|| c}}tt|jDcgc]}| c}t t tt|jk(sJ dfd }|Scc}}wcc}w)Nc4Dcgc]}|| c}Scc}wrr)rrinvs rrz+PermuteView.make_reindexer..reindexA s'**E!H* **s r)rrrrr>)rrrrr"s @rrzPermuteView.make_reindexer: s!*$)) 451q!t5$S^45!s1v5#*U3tyy>-B"CCCC +! +  + 65s B B!N)rrrrorrk)rror list[int]rYr) rrrrrrrrrrrrrr sB ..$FF,  3 rrcFeZdZeddddZe ddZddZy) SqueezeViewN)r>c 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)NrrEzexpected squeezed size to be 1rH)rrKrrrrrrrrrsrrr rrLrrrMrNrOrrr) rrr>rRrSr rTrrrrUrs rrzSqueezeView.createK 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)NrEct|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)rrZrztuple[Expr, ...])rr)rrr rrr*r+s @@rsqueezerzSqueezeView.squeezery s_ $.!qAvA..!*4;AAF1;T $  /;s AA AActd)Nzuse SqueezeView.create())AssertionError)rrIs rr>zSqueezeView.__init__ s788r)rrr>r|rr)rrZrz>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]])rIr rr)rrrrrr~r,r>rrrr%r%I sC48+W+WZ!! G!! 9rr%cleZdZUded<ded< d dZd dZd dZeZe d dZ dd Z y ) GenericViewrZrrrc|jSr)rrs rrzGenericView.make_reindexer s||rctt|jDcgc]}ttj |}}t |j|}ddjtt|d|Scc}w)Nzlambda r r) rrrrmrDrrrrrr)rr index_old index_news r reindex_strzGenericView.reindex_str ssCHTYYCX >? *4::q 9  i01 3sI#6789+FF  s!Bcz|j|jd|jd|jgS)Nsize=zreindex=)rrIrr5rs rrzGenericView.__str__ s= YY% {+x8H8H8J7K-L M  rc*||t||S)NrIrr)r)rrr rs rrzGenericView.create sX@@rc|jSrr}rs rrzGenericView.get_size r~rNrrj)rrr rZrrrrkrY) rrrrrr5rrrrrrrrr0r0 sp  77 3 G HA A!A< A  AArr0ceZdZdZeddZeed dZe d dZ e d d dZ e d d dZ y)rz 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)rrrsrrr evaluate_exprLt)rrr=s rhandle_negative_indexzView.handle_negative_index sZll3||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)Nrct|d\}}t|j|j|||j|j }t ||S)NT)rrH)rKrLrrrMrNrO)rhr rTrRrSrUs rcreate_reinterpret_viewz,View.create..create_reinterpret_view sX#8T"R 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. r9)rrr5rrequire_contiguousr r )rrrrBr r s r"handle_unbacked_or_dynamic_reshapez7View.create..handle_unbacked_or_dynamic_reshape sm 55hIXHH. !33A6.x!B!B8!L  s%);A'&A'c4tdgtzSr)rr)rr s r fake_reindexz!View.create..fake_reindex saS3x=011rr9Fr)_compute_stridesize_obliviousrrH)rhrr rZrTrZrrOrrrr)rr rztuple[int, ...]) rrrresolve_negative_sizerrsrrstatically_known_list_equalsrr4rrr r rrKrto_symints_or_intstorch._subclasses.fake_implsrHrrexprrr rLrrrMrNrO)rrr unbacked_symbols_in_sizesrrErGrRrS old_strideold_size_symintold_stride_symintnew_size_symintrHnew_stride_symintrrTrUrBr s` ` @@rrz View.create sb(H-=tH~=- 66qzz|XN( 77   8 88 LH %h/ 01 4 8(23a7 "9;  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)Nr.) rsrrrrrrrrOner@rn check_equals)r r rrs rrLzView.resolve_negative_size2 s;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_indexerr.r5 IndexErrorrnr)rr r  dense_dimrflatrrs rrzView.dynamic_reshape_indexerC s} :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. rr.rE) rr,r>Gtrsrrrr5)r9r:r4s r compare_sizesz4View._dynamic_reshape_indexer..compare_sizes_ sehhq!n-ehhq!n-ehhq!n-ww<.reindex.. sHA|4Hr)rrrr)rrer view_exprs @rrz.View._dynamic_reshape_indexer..reindex sOu:T* CSZT,C C*D% 01LHiHH Hr)r9rr:rrrr)rsrrr4rrrmrDVIEWrrrrrrrrYrXrCr.reverseinsert)r r r]rbr stack_new stack_oldreordering_dense_dimold_dimsize_oldvarsize_newvar2 size_new2divisormodulus dense_exprrr4rrfs @@@rr[zView._dynamic_reshape_indexerT s))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/)rrrrrr)rrr rZrr)r rZr rZrztuple[list[Expr], list[Expr]]r)r rir rir]r|r&Callable[[Sequence[_T]], Sequence[_V]])r rZr rZr]r|rr) rrrr{r~r?rrrrLrr[rrrrr s  k5k5Z" ",:" &""  !% $%  0  !%z z zz 4 zzrrceZdZUdZded<dfd ZddZeZddZddZ ddZ e dd Z dd Z dd Zdd Zdd ZddZddZed d ddZdddZd dZxZS)!rOz*Pretend our storage has a different layoutrrJct|t|jtr0t j |d|jjyy)NrI)r_rrrIrkrrrlrrs rrzReinterpretView.__post_init__ s@  dii *   tVTYY-B-B-D E +rcP|j|j|jgSr)rrIrJrs rrzReinterpretView.__str__ s&     rc6|jjSrrrs rrozReinterpretView.get_name rrc.|jjSr)rJrrs rr zReinterpretView.get_device s{{!!!rcyrrrs rrzReinterpretView.get_origin_node rrc.|jjSr)rJrrs rrzReinterpretView.dtype s{{   rc@t|jjSr)rrJrrs rrzReinterpretView.get_size sDKK$$%%rc@t|jjSr)rrJrrs rr0zReinterpretView.get_stride sDKK&&''rcdfd }|S)NcTjj}tjj ||}jj j j k7r5tj|j j j S|Sr)rJrrqloadrorrIto_dtype_bitcast)rr tmp_loaderrs rrz+ReinterpretView.make_loader..loader spkk..0G$--/75>BJ{{  DIIOO3++J DIIOOTT!!rrrZrrrrrrs` rrzReinterpretView.make_loader s " rc6|jjSr)rJrrs rrzReinterpretView.make_indexer{{''))rc|jSrrJrs rrzReinterpretView.get_layoutrrcyrrrs rrzReinterpretView.freeze_layout rrct|jj|t|jj|zt|jj|zSr)r'rJrrrMr6s rr7z$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) rsr wrapper_codecodegen_reinterpret_viewrIrJrrrM writelinerrs rrz!ReinterpretView.codegen_referencesww##<< II KK   KK   KK   & 2F  8L8L8V8V++## =  rcyrrrs rr2zReinterpretView.num_reads%rrLrjr`rOrTrYrdrfrVrmrxrr_rv)rrrr{rrrrror rrrrr0rrrrr^r7rr2r,r-s@rrOrO s4 NF  H$"!!&( * -.$) ! ! /   rrOc\eZdZUdZded<ed dZd dZeZe d dZ d dZ ddZ y ) DtypeViewz(Pretend our storage has a different typerU target_dtypect|r]t|\}}t|j||j|j |j |j}t||St||S)NrH)rIr) rrKrLrrrrMrNrOr)rr new_dtyperRrSrUs rrzDtypeView.create/sm  #"7": GZ$!!!!!!$$ J# C Cai88rcP|j|j|jgSr)rrIrrs rrzDtypeView.__str__>s  4+<+<=>>rc|jSr)rrs rrzDtypeView.dtypeCs   rc6|jjSrrIrrs rrzDtypeView.get_sizeGrrcLjjdfd }|S)Ncztj|jjjSr)rqrrrIr)rrrs rrz%DtypeView.make_loader..loaderMs*''c D4E4EtyyW WrrMrIr)rrrs` @rrzDtypeView.make_loaderJs" %%' X rN)rrrrUrrkrjrTrYrd) rrrr{rrrrrrrrrrrrrr)sE2 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 j t jnNt d||fDr!t j t jn 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)r4rrs rrz0SliceView.normalize_start_end..dsHA$Q'Hc3K|]B}t|tr0|jtjtj Dywr)rrhasrMinMaxrs rrz0SliceView.normalize_start_end..gs7 !T" EE%))UYY ' sAA cj||r|n||}j||r|}|S||}|Sr)statically_known_geqr:)rlowerupper clamped_lower clamped_fullmax_funcmin_funcrs rclampz,SliceView.normalize_start_end..clampss`221e<(1eBT  00F   mU3   rcD||Sj|}|||Sr)r?)rrrrrrdim_sizes r clamp_wrapz1SliceView.normalize_start_end..clamp_wrap~s0{++C:CeU+ +rr)rrrrrrrr) rr|rrrrr Expr | intrr) rsrrrr'rrr evaluate_min evaluate_max) rrr>startendrrrrrrs ` @@@@@rnormalize_start_endzSliceView.normalize_start_endYs77##::<$ H%h1GH HyyHyyH  S(+  yyHyyH,,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)NrlrErHct|tk(s Jd|dt|}|zz|<|S)Nz wrong ndim r))rr)rr>r rsteps rrz!SliceView.create..reindexsRu:X. P+eWAhZ0P P.KEsd*U2E#JLrr9r)rrrr TypeErrorrrrrArrKrrLrrrMrNrOr) rrr>rrrrrRrSrTrUrr s `` ` @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) rrr>rrrrrrztuple[int, int])rET)rrr>rrrrrrrrrrr)rrrr{rrrrrrrrSs // /),/36/ //b3A 3A3A 3A  3A  3A3A 3A3ArrcBeZdZUded<ded<d dZd dZd dZd dZy ) BaseConstantrUrrcrcyNrrrs rrzBaseConstant.get_sizesrc|jSrrrs rr zBaseConstant.get_devicerrcyrrrs rrzBaseConstant.get_origin_noderrctSrr=rs rrzBaseConstant.get_readsrrNrYr`rOru)rrrrrr rrrrrrrs"  rrcDeZdZUded<ded<ded<d dZd dZd d Zy )Constantr rrUrrcrcdfd }|S)NcXtjjjSr)rqrrrrrs rrz$Constant.make_loader..loaders<< DJJ7 7rrrrs` rrzConstant.make_loaders 8 rcyrrrs rrzConstant.realizerrcFt|j|j|S)N)rrr)rrrrs rrAzConstant.constant_to_devicesdjj 6JJrNrdr]ry)rrrrrrrArrrrrs# J    Krrc<eZdZUded<ded<ded<d dZd dZy ) IndexingConstantr rrUrrcrcdfd }|S)NcXtjjjSr)rqrrrrs rrz,IndexingConstant.make_loader..loaders>>$**djj9 9rrrrs` rrzIndexingConstant.make_loaders : rcFt|j|j|S)N)rrr)rrrrs rrAz#IndexingConstant.constant_to_devicesdjj 6RRrNrdry)rrrrrrArrrrrs J   SrrcLd}d}ttt||D]\}}|dk(r tjj j ||s,tjj j ||sy|tjd|z}||z}yNrEFT) reversedrrrsrrrErr)rr=expected_strideexpected_stride_maxrys ris_contiguous_strides_for_shapersOs5&123 1 6 ww77  ''""::1>QRuyyA.1  rc<tj|jzSr)rFpadding_alignment_bytesitemsizers rget_align_for_dtyper s  ) )U^^ ;;rc2eZdZdZddZddZ d d dZy) rzxAbstract base for Layout, MultiOutputLayout, NoneLayout. Represents the memory layout of the output of an Operation.c>tt|jrrrs rr zOutputSpec.get_devicerrc>tt|jrrrs r storage_sizezOutputSpec.storage_sizerrc>tt|jrrr6s rr7zOutputSpec.get_free_symbol_usesr rNr`rvrmrx)rrrr{r rr7rrrrrs,C77%*7!7 !7rrceZdZdZdeddf ddZeddZejddZeddZ e jdd Z edd Z e jd d Z d!d Z e Z d"d Z d#dZd$dZe d%dZd$dZd&dZd$dZe d'dZd(dZd$dZd)dZd*dZd+dZddZed d, d-dZy).rzo 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)Nr7 , stride=c3HK|]}t|ttfywr)rrrrs rrz"Layout.__init__..7szLayout.__init__(s >#66t>*4>>*:;MDz""#2dkk&6&6%78H7ITZZLYII;i }VH]O1 N rc|jSrrrs rr zLayout.get_deviceirrc tj5tjt |j t |j |j|j|jcdddS#1swYyxYw)N)rr pin_memory) rs fake_moderrrarrrrrNrs r get_examplezLayout.get_examplelsY [[ &&' 2' 4jj{{>>    s AA99BcBt|j|jSr)rrrrs rrzLayout.is_contiguousvs.t{{DIIFFrct|}|dvs|ddk(ryt|t||D]\}}}|dk7s ||k7syy)N)rrEFT)rrr/)r=rQndimleftrightrs ris_channels_last_contiguousz"Layout.is_channels_last_contiguousysb5z v qQ!$  *5 1  "   D%qyTU] rc t|jttj t t|j |j D]\}}}|dk7s ||k7syyr)rrrr r rr)rrrrs r is_transposedzLayout.is_transposedsc!$ KK ^66tHTYY.sorted_indicess*J=@A'J$$W-A AAs+r.TrIF)r rorro) rrrrrsrrrrrr _shape_envr=) rrrr> non_1_indicesrrstride_orderedrPs rrzLayout.is_stride_ordereds~4;;3u:--- $DII. 3ww11#1BaG   +88Q$++a.882?@Qa@@ B u%E *s5z" 1A'-ayN58 $ 1s5zA~& A!!$~a!e'<.s(  q3 . / rrNc|yt|tjsytfd|jDS)NFc3@K|]}j|ywr)is_unbacked_symint)rrrs rrzILayout._pad_strides..contains_unbacked_symints..sR1y33A6Rs)rrrr'r3)rPrs rcontains_unbacked_symintsz6Layout._pad_strides..contains_unbacked_symintss4 dEJJ/R@Q@QRR Rrc3.K|] }|ywrr)rrrs rrz&Layout._pad_strides..sNa6q9NrTrE)rT)rPzsympy.Expr | intrr)rrrFpad_channels_lastrrrsget_current_noderr^rrrchainpad_dynamic_shapesrrr'rrrrrrrr padding_stride_thresholdrr_r%num_comprehensive_padding) in_stridesrraligncurrent_fx_noderirrr new_stridespaddedrankrprev_idxrrequire_paddingrrs @@r _pad_strideszLayout._pad_stridess0$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'PVUZZ0NV5N5N  &K #*65#9E#A C    ))Q.);:s H(ct|tsJt||jJ|j |j|j |j |_yr)rr rrr+rrrs rrzLayout.pad_strides sM$/;d;/{{&&&'' TYY K rcFtjxrt|tSr)rFcomprehensive_paddingrr rs rrzLayout.should_pad_stridess++P 40PPrct|tr|S|jr|jt|j|j |j |j|j|jSr) rrLrrrrrrrMrNrs ras_fixedzLayout.as_fixeds` dK (K  " " $     KK JJ II KK KK NN   rctjsJdt|jd|j j S)Nzconvert z to FixedLayout first)r rrrr0rrs rrzLayout.make_indexer$sG,, tDz**++@ A ,}}++--rcft|txr|j|jk(xr|j|jk(xrj|j|jk(xrO|j |j k(xr4|j |j k(xr|j|jk(Sr)rrrrrrrMrN)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+rrrMrs rrzLayout.storage_size5s.tyy$++t{{SSrct|j|t|j|zt|j|zSr)r'rrrMr6s rr7zLayout.get_free_symbol_uses8s= TYY 6t{{M: ;t{{M: ; r)rrcrrUrrZrzSequence[Expr] | NonerMrrNrrrrYrrZrrr\rrrrrjrb)r torch.TensorrX)r=rirQrirr)rrorr)r#rorrZrrUrrorLrrLrf)r3rrrrmrx)rrrr{r r>rrsetterrrMrrr rrr~rrrrr+rrr0rr4rr^r7rrrrrs)-qz     &       2 [[ ]] ]]  HG!,> "!F- L!L)7L@KL LL\L Q  .  TH%$) ! ! & rrceZdZdZddZy)rLz A Tensor layout we cannot changecXt|j|j|jS)r/)rrrrMrs rrzFixedLayout.make_indexerFsdiidkkBBrNrf)rrrr{rrrrrLrLCs *CrrLceZdZdZdZddZeddZeddZeddZ 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. )rRdeepcopyr0rs r!get_fixed_layout_without_freezingz0FlexibleLayout.get_fixed_layout_without_freezingTs}}T"++--rct|dk(rgStjjg}t |ddD]}|j ||dzt t |S)NrrEr.)rrrrXrrr)sizesreversed_stridesrs rr z!FlexibleLayout.contiguous_strides^sh 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>rrrrrX)rCr next_striderQrs r fill_orderedzFlexibleLayout.fill_orderedgsx%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 rG)rCrrs rrzFlexibleLayout.stride_orderedxsB%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 rNHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatr rrr)rC 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) rrsrrguarding_hints_or_throwrr __getitem__r rG)rCrrs r same_orderedzFlexibleLayout.same_orderedsb5zS[(((!!99&AE#f+.F4F4FG **5*==rc|jSrrrs rrzFlexibleLayout.sizerrc6|jd|||_y)Nr)!assert_free_symbol_uses_unchangedrrs rrzFlexibleLayout.sizes ..vu= rc|jSrrrs rrzFlexibleLayout.striderrc6|jd|||_y)Nr)rWrrs rrzFlexibleLayout.stride ..x? rc|jSrrrs rrMzFlexibleLayout.offsetrrc6|jd|||_y)NrM)rWrrs rrMzFlexibleLayout.offsetrZrc0|j|j|}|jr)|r'|j||j|j}t |j |j|j||j|jSr) rrrr+rrLrrMrN)rrrrTs ras_stride_orderzFlexibleLayout.as_stride_ordersw((E:  " " $**:tyy$**MJ KK JJ II  KK NN   rc|}|jr)|r'|j||j|j}t |j |j|j||j |jSr)rr+rrrLrrMrN)rr)rrTs ras_exact_strideszFlexibleLayout.as_exact_stridessf#  " " $**:tyy$**MJ KK JJ II  KK NN   rc,|j|j|}|jr'|j||j|j}t |j |j|j||j|jSr) rGrrr+rrLrrMrN)rrrTs r as_fill_orderzFlexibleLayout.as_fill_orderss$($5$5dii$G  " " $**:tyy$**MJ KK JJ II  KK NN   rc,|j|j|}|jr'|j||j|j}t |j |j|j||j|jSr) rTrrr+rrLrrMrN)rrrTs 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)rrrMTF)r>r'r)rinitial_free_symbolsrrers 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)Nrfz)Expected free symbols unchanged, but got z vs )rgr>r')rrrreold_free_symbolsnew_free_symbolss rrWz0FlexibleLayout.assert_free_symbol_uses_unchangedsd* M#88$ 9NO )*:5-*PQ #'77 ;rhrg)rrrrrrNrQras rr>zFlexibleLayout.__init__sU $11$ EG$77=G gK%)$E$E$G!rr:)rCrorr*)rCrorrorr*)rCrorrorrZ)rCrorOztorch.memory_formatrrZ)rCrorrirrZrYr7r\r8rm)rrorrrrL)r)rirrrrL)rrorrL)rrirrL)rz$dict[tuple[str, bool], sympy.Symbol])rrrr7rrr) rrcrrUrrZrSequence[int | Integer] | NonerNrrr)rrrr{rrAr~r rGrrPrTrrr;rrMr^r`rbrdrhrWr>r,r-s@rr r Ks N.00  > >&&-@& &&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)rr_r>rrrrview)rrrrJras rr>zNonOwningLayout.__init__,sA"  MM LL KK MM   rc>|jjSr)r0rrs rrzNonOwningLayout.make_indexer6s}}++--rc|jjj}|dk(ryddlm}t j jj||S)NrTrE) ALIGNMENT) rrrrMutilsrursrrr)rrMrus rmaybe_guard_alignedz#NonOwningLayout.maybe_guard_aligned9sD%%'.. Q;$ww<rrwr^r7r,r-s@rrprp)sG6.P-.$)G!G !G/GrrpceZdZdZy)CommBufferTypesymm_memN)rrrSYMM_MEMrrrr|r|MsHrr|cFeZdZUdZded<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]. r|comm_buffer_typer group_namect|tr|jn|}t||j |j |j|j|j|j||_ ||_ y)NrrrrrMrN) rr r0r_r>rrrrrMrNrr)rrJrrfixedras rr>zCommBufferLayout.__init__`si &0%G!V <<++<<<<oo  !1$r)rJzFlexibleLayout | FixedLayoutrr|rr)rrrr{rr>r,r-s@rrrQs; %$O%,%)% %%rrceZdZUded<ej dZded<ej dZded<d d Zdd Z dd Z y ) NoneLayoutrarcdgSrrrrrr~zNoneLayout.~srdefault_factoryr#rcdgSrrrrrr~zNoneLayout.s1#rrcyrrrs rrzNoneLayout.storage_sizerrc|Srrrs rr0zNoneLayout.as_fixed rc|jSrrrs rr zNoneLayout.get_devicerrNrvrWr`) rrrrr|r}rrrr0r rrrrrssG '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) r_r>rr rr get_bufferrorsrmark_buffer_mutated)rrrras rr>z#MutationLayoutSHOULDREMOVE.__init__se   & & (     OO      ))+ ##D)rc6|jjSr) real_layoutrrs rrz!MutationLayoutSHOULDREMOVE.strides!(((rcyrrrs rrz!MutationLayoutSHOULDREMOVE.strides rc>|jjSr)rrrs rrz'MutationLayoutSHOULDREMOVE.storage_sizes!..00rcxdfd |j}t|tsJt||S)Nct|tr|jSt|tr|j St|t r|j S|Sr)rrrrkrl MutableBoxrI)r unwrap_viewss rrz;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_viewssY&"<=#FMM22&(+#F$6$6$899&*-#FKK00Mr)rr rr )rrrnr)rr1rs @rrz%MutationLayoutSHOULDREMOVE.get_buffers6 dkk*&&)74<7) rc^|jj}t|tsJ|Sr)rrJrr)rrJs rrz&MutationLayoutSHOULDREMOVE.real_layouts)"))&&))) 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)NrrI)rrsrrrorrrIrrrr r rrrrcheck_equals_and_simplifyrkrrrJr rr)rsrcdst unsafe_aliasr9r:rs r realize_intoz'MutationLayoutSHOULDREMOVE.realize_intosI  ##CLLN3 c9 %((C ##~~'mmo*!$CLLNCLLN C1GG$$>>q!D $DdXz$:; ;;))C sF#(S(##((//>:QDrrr;rrrrrr0rr,r-s@rrrs *)) ]]  1  rnr^rrrJcFt||jddyr)r_rrrrs rrzBuffer.__post_init__s   t4rc>|jjSr)rrrs rrzBuffer.make_indexers --//rc@|jsJ||jSrrrs rrozBuffer.get_namesyy$yyyrct|jtr|jjSt t |jj r)rrJrrrrrrs rrzBuffer.get_examples= dkk6 *;;**, ,!$t{{"3"<"<==rc>|jjSr)rr rs rr zBuffer.get_devices##%0022rcyrrrs rrzBuffer.get_defining_oprrc6|jjSr)rrrs rrz Buffer.dtypes &&&rc:g|jjSr)rrrs rrzBuffer.get_sizes("''((rc:g|jjSr)rrrs rr0zBuffer.get_stride s*"))**rc6|jjSr)rrMrs r get_offsetzBuffer.get_offsets '''rct|jtr |jStt |jj r)rrJrrrrrs rrzBuffer.get_layouts4 dkk6 *;; !$t{{"3"<"<==rc|jSrrrs rrzBuffer.get_output_specrrc"|jSr)rrs rr4zBuffer.get_storage_numels~~rc6|jjSr)rrNrs r get_is_pinnedzBuffer.get_is_pinneds ***rct|jtr;t|jts |jj |_yyyr)rrJrrpr0rs rrzBuffer.freeze_layouts> dkk6 *: KK4 ++..0DK4 *rct|jtsJt|j|jj |||_yNr)rrJr rr^rs rrz&Buffer.freeze_layout_with_stride_order%sA$++~6IT[[8II6kk11%}1U rct|jtsJt|j|jj ||_yr)rrJr rrbr"s rr#z$Buffer.freeze_layout_with_fill_order+s:$++~6IT[[8II6kk//6 rct|jtsJt|j|jj ||_yr)rrJr rrdr%s rr&z$Buffer.freeze_layout_with_same_order/s:$++~6IT[[8II6kk//7 rct|jtsJt|j|jj |||_yr)rrJr rr`r(s rr*z'Buffer.freeze_layout_with_exact_strides3sF$++~6IT[[8II6kk22 3  rctjjjt j |j dSrrrs rrzBuffer.is_zero_elements;rrcrjrttjSdfd }|S)Nrcxj}tjjxsd||Sr )rrqrrrrrs rrz"Buffer.make_loader..loaderCs/'')G88DII2GENC Crr)rr rr rs` rrzBuffer.make_loader>s0  "=0@A A D rc"|jSrrors rrzBuffer.codegen_referenceIrrcyrrrs rrzBuffer.decide_layoutLrrct|jtr%|jjj gSyr)rrJrprrrors rrGz#Buffer.get_inputs_that_alias_outputOs/ dkk? 3KK$$--/0 0rct|jtr%|jjj gSyr)rrJrrrors rrCzBuffer.get_mutation_namesTs0 dkk#= >KK&&//12 2rc6t|jgSr)r>rors rrzBuffer.get_read_namesYs4==?+,,rctSrr=r6s rr7zBuffer.get_free_symbol_uses\ |rctSrr=rs rrzBuffer.get_unbacked_symbol_defsbrrcyrrrs rrzBuffer.realizeerrcyrrrs rshould_allocatezBuffer.should_allocatehsrrLrfrj)rztorch.Tensor | torch.SymIntr`rPrTrY)rr*r\rVrWrvrXrmrnrp)rrorr)r)rorrrrrdrr_rzrMrxrr])%rrrrrrrorr rrrrr0rrrr4rrrr#r&r*rrrrrGrCrr^r7rrrr,r-s@rrnrns   50> 3'')+(>  +1;@V"V37V V 78 CH * ;?  U    -H%$)! !&  rrnc<eZdZddZddZej ZddZy)OperationBufferc|gSrrrs rrzOperationBuffer.get_outputsps v rc|Srrrs rrzOperationBuffer.get_defining_opsrrcXtj|tj|yr)rnrrrs rrzOperationBuffer.__post_init__ysT"%rNrrrrL)rrrrrrrErrrrrrms#55&rrceZdZddZy)rcyrrrs rr2zInputBuffer.num_readsrrNrv)rrrr2rrrrr~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)rrrr{rrrrrsrrc,eZdZUdZded<ddZddZy)rNrarcdfd }|S)Ncjj}tjtj j jj||Sr) rrrqrrsr constant_namerorrs rrz*ConstantBuffer.make_loader..loadersPoo'446G88%%dmmot7K7KL rrrrs` rrzConstantBuffer.make_loaders  rcttjj|j ||j SNrrJ)rrsrrrorJrs rrAz!ConstantBuffer.constant_to_devices/&&t}}?   rrdry)rrrrrrrArrrrrs+/O(/ rrcVeZdZddZed d d dZd d dZd dZd dZy)NoneAsConstantBufferctSrr=rs rrzNoneAsConstantBuffer.get_readsrrctSrr=r6s rr7z)NoneAsConstantBuffer.get_free_symbol_usesrrNcJtjjjSr)rsrrnone_strrs rrz&NoneAsConstantBuffer.codegen_referencesww##,,,rctdSNr)rrs rrz$NoneAsConstantBuffer.get_output_specs &&rcyrrrs rrz&NoneAsConstantBuffer.has_tensor_outputrrrurmrxrr_rWrX) rrrrr^r7rrrrrrrrsC23$)! !4 -'rrcReZdZUded<ed d ddZd d dZd dZy) rrrPc.t|j|Sr)r'rPr6s rr7z*ShapeAsConstantBuffer.get_free_symbol_usess  =99rNchtjjj|jSr)rsrrcodegen_sizevarrPrs rrz'ShapeAsConstantBuffer.codegen_references!ww##33DII>>rcyrrrs rrz'ShapeAsConstantBuffer.has_tensor_outputrrrmrxrr_rX)rrrrr^r7rrrrrrrs< 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. rrIFzClassVar[bool]_force_realizeNr|rCallable[..., Any] | Nonerr[rrc #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)Nrp)rrrrrrIrXrrJrrrErZr[rLget_default_sizes_body clear_cache)rold_datarSnew_datas rwith_original_inner_fnz%ComputedBuffer.with_original_inner_fns=+++&&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)rrrIrs rget_computed_buffer_namez'ComputedBuffer.get_computed_buffer_names7 99 99  499f %99>> !rc6|jjSrrIr2rs rr2zComputedBuffer.num_reads rrc6|jjSrrIrrs rrzComputedBuffer.get_readsrrc6|jjSrrrs rrzComputedBuffer.get_read_namesrrcXt|jtttt fs0t jtttStjtdd5|jjrTt|j|jj|jj!cdddSt|j|jj#cdddS#1swYyxYw)Nr/writes index_exprsrT)rrIrXrfrrrG ReadWritesr>rrr r9rPget_store_functionrr;rrs rr,zComputedBuffer.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)rJr7rIhas_store_functionr,)rrer1s rr7z#ComputedBuffer.get_free_symbol_uses+sc$11  II * *= 9:  " " $ d**,AA-P PF rc|js_|jtjjvr9|j dk(r&|j s|jjSt|!Sr) r9rrsrmutated_buffersr2rrIrr_rrs rrzComputedBuffer.make_loaderEs`'') !8!88 A%''99((* *w"$$rcVt|jtttt fSr)rrIrXrfrrrs rrz!ComputedBuffer.has_store_functionPs$))itY%GHHrc|jjj}t|jt t tfr+t|jj|j|St|jtsJt|jt|jj|j|Sr)rr0rrrIrXrfrr rgrrrr)rrs rrz!ComputedBuffer.get_store_functionSs//#,,.;;= 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)rrGStarDep MemoryDepr s rrz0ComputedBuffer.get_fill_order..js01|33\5K5KLMs46rrEpick_loop_orderN)rrJr rGrrIrr;r,r/rrrorrrrrfrrrsrrr schedulerrr) r index_varsrhrr/rvrrPstride_lengthsrs rrzComputedBuffer.get_fill_order[sj 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)rrJr rr#rr"s rrzComputedBuffer.decide_layouts@ 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)NqrrrEr) rGrrr;rrrr rSrr9itemsr) rr var_rangesrr reduce_vars index_size reduce_sizerrs rrz%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..sHqz!T*Hc j||| \}}}jdk(rat|dk(rStt t|}||ddk7r(|Dcgc]}|| }}t |}t |}||}|rGtjjj||t ||\}} } t|| } n|} || |fScc}w)Nr+rr) _apply_loop_reorderingr9rrrrrrsrr_simplify_loopsrKr)x_vars support_varsrCsimplify_loopsnewsizesreindex0rrrr_prunerindex_formulas memory_addrsrs rsimplify_and_reorderzAComputedBuffer.simplify_and_reorder..simplify_and_reorders,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&r/rCror'rrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]])rindexing_exprsrrrrrrrrrget_write_exprsrsrrrHPREFER_STORE_LOOP_ORDERextendget_read_exprsrjrrFloop_ordering_after_fusionrGindex_vars_no_squeezerS)rextra_indexing_constraintsrecompute_sizes_body_funcrrrrrextra_indexing_rangesextra_indexing_exprexpected_var_rangesrr-r&should_merge_loops iter_ranges iter_reindexr reduce_rangesreduce_reindex iter_varsrr+r,s` @@rr-z#ComputedBuffer.simplify_and_reorders4  ' ' )  %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 0runwrap_storagerfmake_kernel_renderrsrregister_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsmutated_inputsmutation_outputsrrrr MutationOutputrr>allowed_prologue_inpsallow_epilogue_fusionallow_prologue_fusion) rrJrfr^rdrg named_inputs first_inputrrras 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)rrJMultiOutputLayoutrrrrs rrzTemplateBuffer.dtypes6 dkk#4 5%C  &&&rc&|jdS)NT normalize)rPrs rr,zTemplateBuffer.get_read_writess''$'77rct}|jD]}t|ttfsJt ||tj tsJt j j jdfd }|tj|jd|jz}|S)z(Build read dependencies from all inputs.cxt|dk(sJtjj|Sr)rrqrro)rrmrh inp_indexers rdummyz4TemplateBuffer._read_deps_from_inputs..dummys16{a'''xx  E0BCCrrror Sequence[Any]rmrvrr ) r>rfrrOrnrrJrrrGrPrr/)rrpr/inp_rawrtrhrss @@r_read_deps_from_inputsz%TemplateBuffer._read_deps_from_inputss.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)r/rrrrc^t|dk(sJtj|dSNrfakerrqr rrmrrs rrtz1TemplateBuffer.extract_read_writes..dummy,v;!# ##99T75>6: :rroru)rrJrmr>rGrrorr rrxrrrPrr/)rrprrtdepsrrs @@rrPz"TemplateBuffer.extract_read_writess dkk#4 53= ** q)9Rb4F **11)<&L  }}//#002 ;// 4==?B)  d11)<<  rc6tjjSr)rrrXrs rr;z!TemplateBuffer.get_reduction_sizesww{{rcyrrrs rr9z!TemplateBuffer.get_reduction_typerrcyrrrs rrzTemplateBuffer.should_allocate rrc*|jgfdfSrr)rr8r9s rr-z#TemplateBuffer.simplify_and_reorder s$      rc6t|jtS)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrJrmrs ris_multi_outputs_templatez(TemplateBuffer.is_multi_outputs_templates$++'899rc|jSr)rgrs rget_allowed_prologue_inpsz(TemplateBuffer.get_allowed_prologue_inpss)))rcy)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_codegen!src(t|tr&t|jtr |jStj |}t|t r |j}t|jtr|j|S)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrrI MultiOutputr realize_inputrmrJr r)rrr1s rrealize_template_inputz%TemplateBuffer.realize_template_input/sf 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)rrrrr4rrrTensorrridrFallbackKerneltensor_to_layoutrbro)rXrrFritemleaf_idxtidmorr leaf_counterrr seen_outputs template_bufwalks rrz0TemplateBuffer.build_multi_outputs..walkJs5&4-0+-(0NGAtNN4.K.K$v,9J.K#LMNL)H&%,,/'H8L,L%,,-A(-KLMMj,&(-.. "33F;\7FH 33BKKMB!-"2;;="gxHr]$& S!t !-"8,Ir)rXrrzlist[tuple[type, int]]rlist[TensorBox])rcountr) rr structuredrrrrrrs ` ```@@@rbuild_multi_outputsz"TemplateBuffer.build_multi_outputs;s4.0  (   6T*b)**r)NNN)rJrrfSequence[IRNode]r^rrdIterable[IRNode] | NonergOrderedSet[str] | Nonerjzdict[str, IRNode] | NonerrrTrs)rprrzOrderedSet[dependencies.Dep]rm)rprrrtrYr]rXrS)r8rTr9rrz9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]rM)rzdict[str, str]rzFinalizeCodegenResult | None)rrrr) rrrrrzdict[int, IRNode] | NonerzFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | NonerzCallable[[int], None] | Nonerztuple[TensorBox, ...])rrrr{r>rrr,rxrPr;r9rr-rrrrrrr,r-s@rrrsH378<150707!076 07 0 07 6 07/07 07d''8$$L OS?C  $K  $=   C  :* *  %    :>;?)+$)+)+ 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. )rdrgN)r_r>rra subgraph_inps subgraph_outs)rrJrfr^rdrgras rr>zTritonTemplateBuffer.__init__is]"    )"7  yy$$$)-DII(>%FJ9=rct||}|jr |jng}|jr |jng}|D]m}t |t j r|jt||9t |tr!|j|j|j|mJ|D]7}t |tr!|j|j|4|7J|Sr) r_r7rrrrrupdater'r)rreresrrrhrras rr7z)TritonTemplateBuffer.get_free_symbol_usessg*=9.2.@.@**b .2.@.@**b   #C#uzz* +C?@C( 333MBC{"{  #! #C#v& 333MBC{"{  #  rc |g|jSrrers rrz TritonTemplateBuffer.get_outputs-t,,--rc&d|jd}|S)NzTritonTemplateBuffer(layout=rr)rrs rrzTritonTemplateBuffer.__str__s,T[[M; rrS) rJrrfrr^zCallable[_P, _T] | Nonerdrrgrrrrmrxrrj) rrrr>r^r7rrr,r-s@rrrhs 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) r_r>rrJr descriptionfailed_benchmark_with_cudagraphsrr` decompositiondecomposition_kwargsconfig_patches)rrrrJrras rr>zChoiceCaller.__init__sh   &'! 05',./38<46!.0rc|j|jrtjfdStj rt fdStjd|idS)NcSrralgorsrr~z(ChoiceCaller.benchmark..s T4[rcSrrrsrr~z(ChoiceCaller.benchmark..s D$Krrr) to_callablerrXbenchmark_gpu_with_cuda_graphrF/profile_bandwidth_with_do_bench_using_profilingrc benchmark)rrrrs `@rrzChoiceCaller.benchmarks\!  * *<<=PQ Q  A A+,?@ @$$T4%dKKrctrrrs r call_namezChoiceCaller.call_namerrctrrrs rrzChoiceCaller.to_callablerrc"|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_keys }}rctrrrs rrzChoiceCaller.hash_keyrrctrrrs rrazChoiceCaller.output_noderrciS)zRInformation returned here is logged to the autotune log file when that is enabled.rrs r info_dictzChoiceCaller.info_dictrrcy)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) rrrrrJrrrrr)rr rr9rrrj)rr)rr)rz6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]rL)rrrr{r>rrrrrrarrrr,r-s@rrrsk11"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)rJrfr^rgc3K|]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) r_r>_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders)rrJrfchoice_timings_fnunfiltered_choicesrgras rr>zMultiTemplateBuffer.__init__sj #"7  #4LN,> %!$% - % " <>!rc|jS)z^ Are all possible choices TritonTemplates or Extern Kernels with out variants )rrs routput_plannablez$MultiTemplateBuffer.output_plannable#s %%%rc|jSr)rrs rrzMultiTemplateBuffer.choices*s }}rcx||jvr|j||j|<|j|Sr)rr)r hint_overrides rchoice_timingsz"MultiTemplateBuffer.choice_timings.s>  4 4 4262I2I-2XD  /##M22rc#0Kt|tjjjsJt ||j |j k(sJ|j}|j|_ d||_y#||_wxYwwr) rrrrTritonTemplateCallerrrJr^r)rcallerrenders rswap_as_triton_callerz)MultiTemplateBuffer.swap_as_triton_caller5s EOO44II   <  {{fmm+++(("("?"?"A - &,D #fD #sA;B>B B BBcNt|tjjjsJt ||j |jjk(sJ|j|jjk(sJ|j|_ yr) rrrrrrrrJrr0rrr^)rrs rfinalize_as_triton_callerz-MultiTemplateBuffer.finalize_as_triton_callerCs EOO44II   <  }}&--"4"4444 FMM$8$8888"("?"?"Arcb|j|}t||j}|||fS)N)rr)rr(r)rrtimings min_choices rget_min_choicez"MultiTemplateBuffer.get_min_choiceKs7%%M%Bgkk2 GJ/00rc|jD]"\}}|j|j|<$|jd|_y)z;Finalize with multiple callers for different hint overridesN)rrrr^)rcallersrrs rfinalize_as_triton_callersz.MultiTemplateBuffer.finalize_as_triton_callersRsN&-]]_ W !M67=7T7T7VD % %m 4 W#'";";D"Ar) rJrrfrrz1Callable[[int | None], dict[ChoiceCaller, float]]rlist[ChoiceCaller]rgrNrrrX)rrr)rr|rzdict[ChoiceCaller, float])rrrrQ)rrrr)rr|rztuple[ChoiceCaller, float])rz*dict[int | None, TritonTemplateCallerBase]rr)rrrr{r>rrrrrrrrrrr,r-s@rrrs>>!>M > / > / > >8&& +/3'3 "3 - -B+/1'1 #1BAB BrrcLeZdZ dfd ZddZddZxZS)CUTLASSTemplateBuffercRt||||||_||_||_yr)r_r>workspace_sizetemplatesupports_epilogue_fusion)rrJrfr^rrrras rr>zCUTLASSTemplateBuffer.__init__^s. );<,  (@%rc6|j |jSdSrrrs rrz(CUTLASSTemplateBuffer.get_workspace_sizems&*&9&9&Et""L1Lrcx|jD]'}tj|jdd)yr)rrqr ro)rrXs remulate_store_fnz&CUTLASSTemplateBuffer.emulate_store_fnps1&&( 5F IIfoo't 4 5r)rJrrfrr^Callable[_P, _T]rrrrwrrrrrvrL)rrrr>rrr,r-s@rrr]s_ A A! A- A  A " A#' A  AM5rrcDeZdZ dfd Zdfd ZxZS)CppTemplateBuffercRt||||||_||_d|_yr)r_r>rroutputs)rrJrfr^rrras rr>zCppTemplateBuffer.__init__vs, );<   ,0 rcvt|jtrt|jtsJt |j|jd}t|t sJt ||j}t|tsJt ||St|%Sr) rrJrmr rrrnrr_r)r first_outputrJras rrzCppTemplateBuffer.get_layouts dkk#4 5dllH5 ItDLL7I I5<<?LlF3 GT,5G G3!((Fff- ;tF| ;-M7%' 'r) rJrrfrr^r rrwrr rrrV)rrrr>rr,r-s@rr r usL 1 1! 1- 1 " 1  1  1 ( (rr cFeZdZdZ 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) r_r>rrdr rrfrrr rfr) rrJrfr^rrdrrras rr>zCuteDSLTemplateBuffer.__init__s );<  ,&*V  %dkk!nf5 KtDKKN7K K5[[^..0F LL)z8#tD L &sB<c|jSrr rs rrz!CuteDSLTemplateBuffer.get_outputs ||rr) rJrrfrr^r rr rdrrrr)rrrr{r>rr,r-s@rrrsQ37 !-    0  *rrceZdZdZ 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)r_r>kernelaccumulator_typer rvariant scale_type_a scale_type_bswizzle_type_aswizzle_type_bmetadatarrkernel_metadata_make_kernel_renderr^) rrJrfrrrrrrr r!ras rr>zNVUniversalGemmBuffer.__init__s DA 0&*V , ((,,"??66oo,,  #'":":rc|jS)z#Return the workspace size in bytes.rrs rrz(NVUniversalGemmBuffer.get_workspace_sizes"""rc|jSrrrs rrz!NVUniversalGemmBuffer.get_outputsrrc  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) rrrar#rrrrrr r!c$jSr)r) render_kernelsrrz9NVUniversalGemmBuffer._make_kernel_render..renders '') )r)Btorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelr(torch._inductor.utilsr)rfrrrIrmrr KERNEL_NAMEr#rrrrrr r!) rout_noderr(r)rrhrrr+s @rr$z)NVUniversalGemmBuffer._make_kernel_renders  6!# ;; $C#y)hh#z*hh   s #  $+112 -##  00!22..LL****....   *f$$r)rNNNN)rJrrfrrr rr rr rrr Any | Nonerr0r r0r!r0rrrvrr)r/r rr|rztuple[Any, Any]) rrrr{r>rrr$r,r-s@rrrs #'#'%)%);;!; ;  ;  ;;!;!;#;#; ;B#:>*%*%,6*% *%rrc&td|DS)Nc3<K|]}t|tywrrr)rrs rrz#is_node_sequence..s4z!V$4r!)r)rs ris_node_sequencer4 s 4e4 44rceZdZUded<d dZd dZddZeddZe ddZ ddZ dd Z e d d dd Zy )r#Sequence[IRNode | Sequence[IRNode]]rfcd|j|}t|tsJ|jSrrfrrro)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)rrrs rrz/InputsKernel.get_read_writes.. sBqWQZZ\2B #c3JK|]}|jywrr)rrrs rrz/InputsKernel.get_read_writes..'s!. (+GCLLN #. r=r) r>rGrMrrfrrrrrrorr)rr/r9rrs @rr,zInputsKernel.get_read_writess<++,.&&[[ 5E%* BEBBE#89 '%.."234 5L,,-. /3/?/?/A.  &&"   rc6|jjSrr.rs rrzInputsKernel.get_reads1r0rct|tr |j}t|tr |j}t|tr%t|t st j|}t|tr|j|St|tr|St|tt fsJt||Sr) rrrIrmrkrOrrunwrap_storage_for_inputTorchBindObjectrnrrrs rrAz%InputsKernel.unwrap_storage_for_input4s 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)rrrrAr)rf inputs_newrrs rr]zInputsKernel.unwrap_storageGsl79  !A!X&GHI!\::1=II 99!<   a  !  JsA%cyrrrs rr=zInputsKernel.is_externTrrcyrrrs rr2zInputsKernel.num_readsWrrcttj}|jD]B}t |t r||j |z}(|D]}||j |z}D|Sr)r>rr!rfrrr7)rrerrh inner_inps rr7z!InputsKernel.get_free_symbol_usesZsv u|| $ &;; GC#v&S--m<<!$GI77 FFAG  G rN)rrrrrsrurK)rfr6rzlist[IRNode | Sequence[IRNode]]rXrvrmrx)rrrrr:r,rrrAr~r]r=r2r^r7rrrrrs //  ,,$ 3  (  N+$) !  ! , rrceZdZddZddZy) NopKernelcyrrrs rr?zNopKernel.is_no_opirrctSrr=rs rrzNopKernel.get_readslrrNrXru)rrrr?rrrrrKrKhs rrKceZdZdZeddZe 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 rrEc32K|]}t|ywr)rrs rrz&ConcatKernel.create..s-W1.CA.F-WrFc3K|]p}d|jvxr\|jdjtjxs-|jdjtjryw)rrON)r^rrrJrLrargs rrz&ConcatKernel.create..sw !HHUO11@S@S1Txx44&+&<&<5  sA6A8c3jK|]+}t|xr|jj-ywr)rrrNrs rrz&ConcatKernel.create..s/ FG !! $ A)A)A A s13N)rrrrrNrrJrfr)7r r rrrrrrsrrrr r rFr.rr+rrrrrLrrrr/r' current_noderrrOrmrrkrrrrrrnrfrIrlr rjrirErrHFOREACHregister_operation_listr_rr]r`)rrfr>rrr  offsets_start offsets_endr input_sizer output_striderrJany_input_is_storage_and_layout fx_node_argsrN concat_kernelrop_namesrhrzinput_unwrappeddevs rrzConcatKernel.createvs %%'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||3ywrrDrFs rrz=ConcatKernel.can_realize_into_without_copy..s3B  88R@rGrJ)rrcan_realize_into_without_copyrIrkrmrrrJrLrrr0rrrr ExternKernelAlloc)rrrs rrhz*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)rKr7r6s rr7z!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)NrHrJr)rrOrrKrrrrIrmrrrhrprJrrr r rrrrsrrr)rrrrRrJr9r:pws rrzConcatKernel.realize_intosL #/$S)"7"<%76B#/:c:/ c9 %##CHHc2 2 c: & KKM388X. ..00c:"1#"6xx   >>#--/__&   ?Aq  ::1a@  C(( s4F)cyrrrs rrzConcatKernel.should_allocate:rr)rfrr>rrrmr)rrrr"rrrmrx)rrrrrrrX) rrrr{rrrhr^r7rrrrrrOrOps uun/3! ! ,! ! ! FN+$)C!C !C,C ))BrrOceZdZUdZdZded<ejeZ ded<dZ d ed <dZ d ed <dZ d ed <eje Zded<dZded<dZded<ejeZded<dZded<ejeZded<eje Zded< dA dBfd ZdCdZdDdZdEdZdEd Z dF dGd!ZdHd"ZdFdId#ZdJd$ZdKd%ZdLd&Ze dMd'Z!e" dNd(Z#e"dOd)Z$e"dPd*Z%e"dPd+Z&e" dQ dRd,Z'e" dS dTd-Z(e" dS dUd.Z)e"dPd/Z*e"dPd0Z+e"dPd1Z,e"dPd2Z-dEd3Z. dVd4Z/dFdWd5Z0dXd6Z1dYd7Z2dSdZd8Z3dLd9Z4dHd:Z5dHd;Z6dHd<Z7d[d=Z8d\d>Z9e:d dS d]d?Z;dLd@ZS)^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. rrv constant_argsrr{rNReinterpretView | None output_viewr^python_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]rec Dt ||||||_|r|ni|_||_| |_|j ||j|| |_|ji|_ g|_ tjj|_i|_yNrW)r_r>rorrqrwset_cpp_kernel_nameset_python_kernel_namerucollect_arg_kwarg_propertiesr{rersrrYfx_noder) rrrJrfrorrqrrrsrurwras rr>zExternKernel.__init__Zs   + &fB &&   1 ##$67-J* ))+!# "ww++ +-rc |g|jSrrrs rrzExternKernel.get_outputszrrctSrr=rs rrz%ExternKernel.get_unbacked_symbol_defs}rrcNt|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)rrwr_ops OpOverload_schema arguments kwarg_onlyr real_typerrrrfrxryru schema_kwargs)rrrs rrz)ExternKernel.collect_arg_kwarg_propertiess$**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)rrJr apply_constraintrrs rrzExternKernel.decide_layouts- dkk> 2  ! ! #     3rct||\}}|r|j||s|j}|r$ddlm}|||d}|j ||yy)NrE)'set_kernel_post_grad_provenance_tracingT)r=)rf make_commenttry_get_kernel_namerrwrite_provenance_debug_handle)rwrapperr origin_str_detailed_origin_strr debug_handles rcodegen_commentzExternKernel.codegen_commentsf,?tW+M( (    ,224K  FBkTL  1 1+| L rctrrrrs rcodegenzExternKernel.codegenrrc^||_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)rrEN) startswithr)rrs rrz3ExternKernel.set_cpp_kernel_name..s%KAc9JS12ZKs,,rEr^_vz at::_ops::z::call)rsrsr cpp_wrapperrrwrrr namespace _overloadnamerrreplacetorchgen.aoti.fallback_opsrrrrr)rrsropnamer version_infolatest_versions rr~z ExternKernel.set_cpp_kernel_names.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) rrrwrrrHigherOrderOperatorrrr)rrrrs 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)NrE) CppWrapperCpu)codegen.cpp_wrapper_cpurr rrsr device_type fx_wrapperrrrrrrsget_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)rrrrrzrx) rrr r rrrrr)rrls r copy_inputzExternKernel.copy_input sa   <<>++-]]_::<))+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| | ||\}}||i|}d}tjFj x}rtjHjJjMd }tO}tjHjPtjRjTjVur|d }tYtjH}|5t[|tjH|dddt]|||}t/|t^t`fs|gn|}|D]}t/|tjbs|jds+tfjhr.unflatten_argsUsFo.J!"56N0 8 MM$z"23MM$~"67  8 %%fi8A55$aeeHb&99 9rrrrEzEsparsity not handled. Please file issue for sparse inference weights.rz Found from : )rrrrrztuple[list[_T], dict[str, _T]])6r tree_flattenrrsrrrcreate_symintnoderGeneratorStaterrrrrdefault_generators clone_stateOpaqueObjectStaterrrrrKrrkro constantstorchbind_constantsrB get_valueropaque_example_valuerirrrrYr^rr r_higher_order_opseffects with_effectsr1r8r2rrr is_sparserFgraph_partitiondisable_cudagraphs_reason)rrrr binded_args args_flat tensor_argsnon_tensor_argsreal_non_tensor_argsrUr device_indexrr example_argsnew_args new_kwargsexample_outputr{r node_meta_valctxexample_out_lirmsgrrrs @@rprocess_kernelzExternKernel.process_kernels" $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        YBj K Ks W8W==Xc *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. rrSrrrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFrrH)&rrkrrOrlrsrrrorr^rnrrJr r-rrJrLr&r/rrrGrrrr stride_vars offset_varrkrrrrIrLrr ) rr x_unwrap_viewrx_unwrap_view_fx_node index_argsrrrrQrMexpecteds rconvert_to_reinterpret_viewz(ExternKernel.convert_to_reinterpret_views$!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)rP)rrrH)#rrrrrrrrrrr?rsradd_tensor_constantrrPrr r rrrrIrOrrkrrrlrrrm NonTensorObjrrrCs 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 2 ||J|jdvr|s|St|rt|jtr|rt ||xr$t |jj }t|dd|rJttjjj|jjn|||St|ddd|||St|jttfrf|r|jj!|s5|rCt#||jj|j%r| t'||S|St|jx}t(rwt|j+x}tr t-dt|tr<|r|j!|s'|r't#||j|j%r|St|t.rX|r|jj!|s5|r5t#||jj|j%r|St|t0rt|j2t4rt|j2t6st|j9x}rtt;|drht|j2t<sN |j?|j2|_|r|jA|||S|r|jC|||S d} |j%} |tjj} tGtI|j%D cgc]<} | jK|| d r%| jM|j%| d r| >} } | D].} tNjPjRjU|| d d }0|jW|}t|dd||||rt ||sJ|S| r<| |JtNjPjRjY|| }t'||S|S#tD$rYHwxYwcc} w) N)rrETF)rrrrrzHthe MutationLayoutSHOULDREMOVE's real layout shouldn't be FlexibleLayoutrIrrrrE)-rrrrr rr4rrKrrsrrrRrLrprrArrVrrr.rrrIrkrOrlrrirrequire_stride_orderrequire_exact_stridesrrrrErrrloweringslice_rr)rrrr)ruse_current_stride_ordermutation_layoutrrl expanded_dims orig_sizerrr>s rrequire_strideszExternKernel.require_strides-sS M$=== ;;=F "=H ! #!,,..90R500K3ALLN4I4IJJ-*#(- 8 - ! 0 0 H H$%LLN$9$9!" "'&3H*#(-%)&3&3 HALLN[/,JK1<<>;;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  !''   5a? ?? (]-F FF((//9=A21mD DW'  s;5R1RAR RRc |j||Dcgc]4}t|tjr|jj n|6c}|Scc}w)N)r)r)rrrSymIntrrP)rrr)rrs rrz"ExternKernel.require_exact_stridessV"" KXFGz!U\\: A( #  s9A c*|j|||S)N)rr)r)rrrrs rrz!ExternKernel.require_stride_orders""1E"OOrc.|j|tSr)rrKrCs rrequire_channels_lastz"ExternKernel.require_channels_lasts''+<==rc.|j|tSr)rrMrCs rrequire_channels_last_3dz%ExternKernel.require_channels_last_3ds''+=>>rcdd}||r|S|j|tj|jS)Nc |j}|tjj vxr'tjj |j S#ttf$rYywxYwr)roAttributeErrorrrsrr is_mkldnn)rrs ris_mkldnn_tensorz9ExternKernel.require_contiguous..is_mkldnn_tensors] zz|177,,,R1B1B41H1R1R R#$78  sAA)(A)rrrrrr r r)rrr s rrDzExternKernel.require_contiguouss@ S A H,,>44QZZ\B rch|j|tj|jSrr rCs 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) rrrrrxrrrrwrr)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) rsrrrxrrorrrfrrval_to_arg_str) rrr1name_to_arg_propertiesrUrrproptype_rr9s rcodegen_const_argszExternKernel.codegen_const_args,s 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)rsrrrwr rfrorrrxrrrrr r4r )rrfneed_codegen_constant_argsrrrr s r codegen_argszExternKernel.codegen_argsMs5 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.rz not in self.allarg_properties)rrryr.)rr rrUs rget_kwargs_valuezExternKernel.get_kwargs_valuegsx 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)rsrrrwrrrur rrrryrrr rr)rskip_outrr rr ks rcodegen_kwargszExternKernel.codegen_kwargsts= 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_namespacerrrrEr unknown_op)rrrrrsplit)rr op_namespaceop_names r get_op_namezExternKernel.get_op_namesv << #\\((F"6<9LML'//'BL'..sA6q9L%ax0G#Grc tjrtjjst |j dk(rytjjj|j }tjjj|j}|j}|jd|jd|d|d|d yyy)Nrzassert_size_stride(r r) rF size_assertsrsrrrnrrcodegen_shape_tupler0r( rro)rrrrr' s rcodegen_size_assertsz!ExternKernel.codegen_size_assertss   qww':':T]]_-277'';;DMMOLDWW))==doo>OPF&&(G   %dmmo%6bb7+UVW (; rc Htjrtjjsw|j }|tjj v}|j}|r |jd|dtd|dy|jd|d|dyyy)Nzassert_alignment(r rz # buffer z (op: z) is assumed to be not aligned) rFalignment_assertsrsrrrorr( rrg)rrralignedr' s rcodegen_alignment_assertsz&ExternKernel.codegen_alignment_assertss  # #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')) rF test_configstrack_memory_lifecyclersrr"write_memory_track_allocation_onceror)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_symbolsrr7rorr)rremaybe_get_symbolsrrUs rr7z!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)Nrrzpython_kernel_name=rr)rr|fieldsrrrzr)rrrr}s rrzExternKernel.__str__sd$8$? !+ 1  $++D1 zzl!GD%**56 7    |D$4$4#789u%%  s,B rNNNNrN)rr^rJrrfr6rorvrdict[str, Any] | Nonerqrprrr^rsr^rurtrwrvrrrrrLr)rrxrr^rrrrxrrrsr^rr)rrr^rrr]rj)rrrr)rrrr rr rzftuple[Any, list[Any], list[Any], Callable[[Any, Any], Any], dict[sympy.Symbol, pytree.KeyPath] | None])rrrrOrK)NNF) rrrSequence[int] | Noner)r[rrrrrm)rrr)rirrrr)rrrrorrrr)rrvrr{rrv)rrwrrrr)r rrr rr )r rrr)rz'tuple[list[Sequence[Expr]], list[Expr]])rztuple[Expr, Sequence[Expr]]rx)?rrrr{rorr|r}rrrqrrrsrrurwrxryrzr{rer>rrrrrrr~rrrr~rrrrrrrrrr r rDr rr r r r r! r( r, r0 r5 r7 r> r^r7rrr,r-s@rrr>s $&M=%.[..tDFND*.K'.%) )"&OZ&4E;3D3D4!=(,K$+26N/63D;3D3D40:>6=->t-T*T(*(,.2)-&*79+/...4 . % . & .,.'.$.(5.). .@.#$J! HLM+M:DM M""#;J +$   ^ !^ *-^ 9<^  ^ ^ @E E N!!!!F!!'+37# a a$a1 a  a  aaFQV    '9  JN      DIPP,P=AP PP >>??"   "!"+9" "HXB4 J4    =$'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) r]rrrr_r>rsrr_rr`) rrJrfrorrqrrrsrurwunwrapped_inputsras rr>zExternKernelOut.__init__ s ..v6*H5Mtrr,r-s@rrL rL s1(*(,.2)-&*79+/))!)% ) & ) , )')$)(5))) ):rrL c eZdZdfd ZxZS) RandomSeedsc tjtj}t|t |tj|gg|j |j|ggddtjjy)Nr{zaten.randint.low_outzat::_ops::randint_low_out::call)rJrfrorrrsrw) rrrr_r>rLr(rrrandintlow_out)rrrlimitsras rr>zRandomSeeds.__init__.slU[[) kkW !::vzzE7;5> ,,  r)rrrrcrrrrrr>r,r-s@rrT rT -s   rrT cheZdZddZ d dfd ZddZd dZxZS) ric&|j|yr)generate_extern_kernel_allocrs rrzExternKernelAlloc.codegenBs,,T2rc F|j|} td| DsJt | d|t t t | ||xsid|||| g|_tjj||_ tjj|y)Nc3<K|]}t|tywrr3)rrs rrz-ExternKernelAlloc.__init__..QsCQ:a(Cr!) r]rr_r>rrrr rsrr_rr`) rrJrfrorrrrsrurwrQ ras rr>zExternKernelAlloc.__init__Es ..v6C2BCCCC    &!#3 4  Lb    )  ') GG++D1  ""4(rcyrrrs rrz!ExternKernelAlloc.should_allocateerrctrrrs rrz"ExternKernelAlloc.apply_constrainthrrrG )rNNNrN)rJrrfrrorvrrF rrr^rsr^rurvrwrvrrrXrL)rrrrr>rrr,r-s@rririAs3(*(,)-&*79+/))!)% ) & ) ' )$)(5))) )@"rricTeZdZdZ dfd ZddZd dZd dZd dZxZ S) rfzP An output buffer that represents the mutation of a pre-existing buffer ct|d||j}tjj ||g|_||_tjj||_ yr) r_r>rorsrrmutation_names mutating_noder_r)rrJ mutated_noderd mutated_node_nameras rr>zMutationOutput.__init__qsb d62(113 ##$5601(5GG++D1 rc|jSr)rd rs rrzMutationOutput.get_defining_op{s!!!rc|jSr)rc rs rrCz!MutationOutput.get_mutation_names~rdrcyrrrs rrzMutationOutput.should_allocaterrc^|j}d|DDcgc]}||c}Scc}w)Nc3ZK|]#}tjj|%ywr)rsrtry_get_buffer)rrs rrz6MutationOutput.get_mutation_buffers..sP..t4Ps)+)rC)rrc rs rget_mutation_buffersz#MutationOutput.get_mutation_bufferss=002QP     s *)rJrre rrd rrrrrzrXrr) rrrr{r>rrCrrm r,r-s@rrfrflsA2 2062GP2 2"# rrfceZdZUdZiZded<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 experimentalrEr)rTMADescriptorExperimentalTMADescriptorStable)rrPtma_metas r _create_implzTMADescriptor._create_impls\8}!!! A;. (,VBhqkB BA;(* **&v< < rprOrrrrnrrPrsrr_rr`)rrPrfroras rr>zTMADescriptor.__init__s   !,,.  &!6 * -   GG++D1  ""4(rc&|j|yr)generate_tma_descriptorrs rrzTMADescriptor.codegen''-rc|jSr)rPrs r get_tensorzTMADescriptor.get_tensorrr)rPrrv ztuple[str, tuple[Any, ...]]rrp )rPrrfrvrorvrrrG rK) rrrr{rq rrrw rr>rr~ r,r-s@rrp rp s(*F $)=='B= =='B ))&3)DQ) )..rrp c:eZdZdZ d dfd ZxZS)rt z the new host-side TMA Descriptor API: (the ones obtained via create_{1d,2d}_tma_descriptor calls). See also TMADescriptorStable for the new API. cbt|dvsJt|t|k(sJ||jj}||_||_||_t|j|_|g}g|j|j|j }t|!|||y)N)rErrPrfro) rr rr block_dims element_sizer(r_r>)rrPrr r rfroras rr>z"TMADescriptorExperimental.__init__s4yF"""4yC O+++  !++-66L $( N  YY __     '  rr) rPrrlist[int | torch.SymInt]r r r r|rrrrrr{r>r,r-s@rrt rt sG$(   ' -  !     rrt c$eZdZdZdfd ZxZS)ru z the new host-side TMA descriptor API (the ones obtained via TensorDescriptor.from_tensor). See also TMADescriptorExperimental for the old API. c:||_t| ||g|y)Nr ) block_shaper_r>)rrPr ras rr>zTMADescriptorStable.__init__s(& 8%  r)rPrr r r r-s@rru ru s  rru cBeZdZ d dfd ZddZxZS)SubgraphBufferc t|d||||_||_tj j ||_tj j|tj j|j|||_ t|jsJt|j}|D]T}||jj|j<|jjj!|jV|D cgc]} | jc} |_ddlmcm} t j*|j5dddd} i| |xsi} | j-| 5|jj.|jddd|r:|jj0D]!} | j3|j5#dddycc} w#1swYSxYw#1swYyxYw)NrFATEN) max_autotunemax_autotune_gemmmax_autotune_gemm_backends)r_r>r`example_inputsrsrr_rr` make_subgraphsubgraphr4rfrir graph_input_namesr sym_inputstorch._inductor.configrrFset_graph_handlerrrun operationsrrR)rrJrr`r subgraph_namerr sym_inpsym_varinductor_config base_patchesmerged_patchesrZras rr>zSubgraphBuffer.__init__s v{3,GG++D1  ""4(--dgg~}U  ,,,(5 ! AG7>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__8s" !JJ rN)rry)rrrr>rrr CodegenGraphr 7s 'rr )r4rfr'codegen_subgraph_with_flattened_outputsr r r)rrr r outer_inputss rrzSubgraphBuffer.codegen6sz ' '   ,,,7;{{C!++-C C77  ' -doo - - YYK DsBr) rJrrrr`torch.fx.GraphModuler list[Any]r rrrF rG rrrr>rr,r-s@rr r sT15,A,A",A ! ,A " ,A  ,A.,A\  rr ceZdZdZd dZddZeddZ ddZ ddZ e d d dfd Z ddZ dfd Z edfd Zdd 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)rrrs rrzBUserDefinedTritonKernel.get_kernel_and_metadata..Xs%*/0FII''**s#& restore_value reset_idx reset_to_zero)triton.runtime.autotunerr *torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrrr4r r r rrr r configs)rr r r restore_value_argsreset_to_zero_argsrrs @rget_kernel_and_metadataz/UserDefinedTritonKernel.get_kernel_and_metadataKs6P"--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. FrErc3&K|] }|dk( ywr+rr s rrz.sNa16Nr T)rF*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrIrmrrrrrs rr z)UserDefinedTritonKernel.can_fuse_epilogueks0@@  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.codegens}}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|||j|||j |\}} } } |j D cic]} | |j| } } |rzt|jjjdk(sJtt|jjjj}|| vsJ|\}}|| |<|jDcgc]}|jc}|jDcgc]}|j s|j"}}t%fd|D}g}g}g}g}t'j(| j+t-t'j.d| D]q\}}||vr|r|j1||j1|t3|t4r?|j1|j7|j1|j9t3|t:t<t>t@jBfr,|j1||j1tE|||vr(|j1d|j1t:|R |r(|j1d|j1t:8|jG|jGZtIdtE|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_dictrEc3(K|] }| ywrr)rrr s rrz3UserDefinedTritonKernel._codegen..s$FaYq\$Fsrr.NzUnsupported arg type: rT) arg_typesraw_argsraw_keys triton_meta inductor_metar!roriginal_fxnode_name))r-r r !define_user_defined_triton_kernelrgridrur rr rrrrrparams is_constexprnumr>rrrrrepeatrrrrr rrrrrrrrrgenerate_kernel_callr r)rrr r rr r r new_namer r extra_launch_argsr named_argsmutable_arg_nameepilogue_computed_bufferrr. constexprsconstexpr_namesrr raw_keys_filteredraw_args_filteredrrUr s @rr z UserDefinedTritonKernel._codegens+ I  ( ( *       5 5   KK   II      261S1S ,-At$$Q' '   t((44;;<A AA#D):):)F)F)M)M$NOTT #z1 11*9 ' $a+CJ' (%+]]3QVV3 %+]]EannaeeE E$$F:$FF! ')')"    I$4$4R$8:K L $ WID#&+I+K  $ $T *  $ $S )#v& C1134  1C#udEJJ!?@ C   c+( B  %23KKO$$S)%))+%))+),B49+RPSu*UVVI$ WL Wh/$$  &&#'??$!%!2!2 % u 4EsM(M-$M26M2cPt||t|j|zSr)r_r7r'r r`s rr7z,UserDefinedTritonKernel.get_free_symbol_usess- w+M:=M II}>   rctSrr=rs rrz0UserDefinedTritonKernel.get_unbacked_symbol_defs rrcg}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).rrrrrArrp rrrr rrrr_r>rrr r r rr rur r r rastr kernel_srcparse kernel_astr kernel_argsr rrrrr rfrersrr`)rr r tma_descriptor_metadatar rfrror rrrr rrUr r autotuned_kwargsr rrras 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 Ltjst| S|jj j Dcic]4}|j|j|jj6}}t|jj j}t||jDcic] \}}|j|j"}}}tjt!|jj j Dcgc]}|j#|c}t!|jj jDcgc]}|j#|c}t!}|Scc}wcc}}wcc}wcc}wNr)rFr r_r,r rr/rr rorrrrerGrr>rename) rformal_arg_dep read_renamesformal_arg_writes mut_output write_renamesrrras rr,z'UserDefinedTritonKernel.get_read_writesss@@7*, , #'"3"3"?"?"E"E    !1!1.2E2E!F!O!O!Q Q  !!2!2!>!>!E!EF/2!4#8#8/ *    !4!4!6 6  #-- $00<<BBJJ|,  $00<<CCJJ}- #  9   s9F?%FF "F!c,t|jSr)rrers rrz#UserDefinedTritonKernel.get_outputssD))**rc|jSrrrs rr z"UserDefinedTritonKernel.get_devicerr)rz(tuple[Kernel, Any, list[str], list[str]]rXrG )rrxr ztuple[ComputedBuffer, str]rr)rrxr z!tuple[ComputedBuffer, str] | Nonerrrmrxr) r rr r r r{r r{rrrsrr`)rrrr{r r rrr r r^r7rr>r,rr r,r-s@rr r FsG@6p<<7+7>X7 7e %e ;e   e N56$) ! ! 7 N)N) N) "0 N) $ N) N)`##J+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|tywrr3rrs rrz3InplaceBernoulliFallback.codegen..s>Q:a(>r!c3XK|]"}tt|j$ywr)rrrr s rrz3InplaceBernoulliFallback.codegen..sIVQ113Is(*rr z, NULL)r) rrfrsrrrrrrreprroending)rrrs rrz InplaceBernoulliFallback.codegens>$++>>>>IT[[I 77     '')*!A3b3tTEWEW;X1Y0ZZabibpbpaqr    '')*!A3b3tTEWEW;X1Y0ZZ[\c\j\j[kl rcyrrrs rrz(InplaceBernoulliFallback.should_allocaterrc&|jdgSrr:rs rrCz+InplaceBernoulliFallback.get_mutation_names"##rctSrr=rs rrz1InplaceBernoulliFallback.get_unbacked_symbol_defsrrc^t|dt|j|j |g||t j j|jt j j||_ t j j|y)Nrrw) r_r>rr r]rsrrror_rr`)rrwrroras rr>z!InplaceBernoulliFallback.__init__s   alln -    $ #   ##AJJL1GG++D1  ""4(rrG rXrzr)rwrrrror rr rrrr{rrrCrr>r,r-s@rr r sF $ )' ),2 )DG )  ) )rr cxeZdZdZddZd dZd dZd dZ d fd Ze d ddZ xZ S)InplaceCopyFallbackr cR|j\}}}|j|||yr)r codegen_device_copy)rrrr non_blockings rrzInplaceCopyFallback.codegens)#'#4#4#6 c<##Cl;rcyrrrs rrz#InplaceCopyFallback.should_allocaterrc&|jdgSrr rs rrCz&InplaceCopyFallback.get_mutation_namesr rctSrr=rs rrz,InplaceCopyFallback.get_unbacked_symbol_defsrrct|d|||ddtjj |dj tjj ||_tjj|y)Nz aten.copy_aoti_torch_copy_)rrrsr) r_r>rsrrror_rr`)rrJrfroras rr>zInplaceCopyFallback.__init__sr     +.   ##F1I$6$6$89GG++D1  ""4(rc||fDcgc]}|j|}}|f}tt|j||}|Scc}wr)rr rr )rrrr rrfror1s rrzInplaceCopyFallback.createsX25c ;1###A&;;% $ cnn. /     rrr,r-s@rr r s~<$))!)% )  )$.s 9a!!# 9r!rr r) r4rfrr rorrrr )rrargrefss rrz$MutatingFirstArgExternKernel.codegens~ ,,, 9T[[ 9 t)) *  ##%&a '(:';1W^^rr r]rsrrror_rr`rrrsrkrmrrnever_reuse_buffersrrI)rvariabler ras rr>zResizeStorageBytes.__init__s(C(@*@@(   h113 4    +#+  ##H$5$5$78GG++D1  ""4("FG(Xz9$EFVXVF ##'' (>(>(@Ar)r$ rr rrrrY r-s@rr! r! sBBrr! c(eZdZdfd ZddZxZS)SetSourceTensorKernelc|jt| |j||gdtj j jjt|tttfsJt|tjj j#|j$j'tjj j#|j'tjj j#|j'|j)}t+t-|||t+t-|||g|_y)Nz!torch.ops.aten.set_.source_Tensor)rrrwr)rr_r>rrrqrset_ source_Tensorrrkrmrrrsrr# rrIror rfrre)r self_tensorstorage_tensorrras rr>zSetSourceTensorKernel.__init__.s$$&   % % ' . )B ++99  +*i'HI 4 L  I ##'' (8(8(A(A(CD ##''(?(?(AB ##'' 8**, :V4k4 H :V4nd K! rcF|jd|jdgSrr rs rrGz2SetSourceTensorKernel.get_inputs_that_alias_outputBs"DOOA$677r)r* rr+ rrrrz)rrrr>rGr,r-s@rr& r& -s  (8rr& 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.codegenMs))$/rcyrrrs rrzScatterFallback.should_allocatePrrcf|jd}t|tsJ|jgSrr8rrhs rrCz"ScatterFallback.get_mutation_namesSs.kk!n#v&&& rctSrr=rs rrz(ScatterFallback.get_unbacked_symbol_defsXrrNTr include_selfc ft|t|_|jr%|||fDcgc]}|j|} }|f} n$||fDcgc]}|j|} }||f} t |dt |j|j| | ||dt|ddg|tjj|jtjj||_tjj!|ycc}wcc}w)Nrr5 rr6 )rrrurw)rr src_is_tensorrr_r>rr r]rrsrrror_rr`) rrwrr>rrrr6 rtensorsroras rr>zScatterFallback.__init__[s(Y7   78%oFt))!,FGF FM78%jAt))!,AGA #JM   alln -    ( | <";/+3^*D#   ##AJJL1GG++D1  ""4(%GBs D)D.rG rXrJ r)rwrrrr>rrrrrrr^r6 rrrr r-s@rr. r. Fs| 0 "!!)!!) !) !)  !)  !)!)!) !)!)rr. 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.codegens++D1rcyrrrs rrz IndexPutFallback.should_allocaterrc&|jdgSrr rs rrCz#IndexPutFallback.get_mutation_namesr rctSrr=rs rrz)IndexPutFallback.get_unbacked_symbol_defsrrc ||_|Dcgc]}|| }}||g|Dcgc]}|j|}}d} t | dt j |j ||fd| |tjj|jdtjj||_ tjj|ycc}wcc}w)Naoti_torch_index_put_outrzaten.index_put_)rrrsrwr)rrr_r>rr r]rsrrr:r_rr`) rrwrrr accumulater valid_indicesr9 rsras rr>zIndexPutFallback.__init__s $+=qq}= =34f2M}2MNQ4%%a(NN4   alln -    ( M0+#   ##DOOA$67GG++D1  ""4(>NsC/C/C4rG rXrzr) rwtorch._ops.OpOverloadrrrr rrvrC r rrr r-s@rr; r; 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)rsrrr s rrz$DeviceCopy.create..sGqA***Gs&(zDeviceCopy in input programrTrm)r r=rprsrr rrrF aot_inductoruse_runtime_constant_foldingradd_device_inforArbrrDrr0rjrrrrNrG rLr r) rrrr x_deviceroris_destination_pinnedis_source_pinneds rrzDeviceCopy.creates<<>### 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 |dy|j|d|j |dy)NrrrE)r rrqr r)rrrs rrzDeviceCopy.codegens{  "4yA~~     ' 'Q));;=tAw   ' 'Q1G1G1I4PQ7 SrN)rrrrcr rrrrG )rrrrrrrrrrG rG s- - ^TrrG 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_readsrrcyrrrs rrz*DynamicSelectStorageOffset.should_allocaterrct|dttjdg||_||_||_||_||_ ||_ yNrr) r_r>rrrunbacked_offset_symbolr base_offsetbase_dim_striderr)rrW rrX rY rrras rr>z#DynamicSelectStorageOffset.__init__sS ze1DErJ'=# &.  rc.t|jgSr)r>rW rs rrz3DynamicSelectStorageOffset.get_unbacked_symbol_defs s466788rc.t|j|Sr)r'rr6s rr7z/DynamicSelectStorageOffset.get_free_symbol_uses s  M::rc>|j||jy)NrX)codegen_dynamic_select_indexrrs rrz"DynamicSelectStorageOffset.codegen s,,T,DrrurX)rW sympy.Symbolrr^ rX sympy.Symbol | intrY r_ rr_ rrrrrrmrxrG rrrr{rrr>rr^r7rr,r-s@rrR rR s ,(  ,  !  &989$);!; !;:; ErrR 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||_||_||_||_||_ yrV ) r_r>rrrunbacked_size_symbolrrrr)rrf rrrrras rr>zDynamicSliceSize.__init__, sH ze1DErJ$8!   rc.t|jgSr)r>rf rs rrz)DynamicSliceSize.get_unbacked_symbol_defs< s444566rctt|j|jt|j|Sr)r'rrrr6s rr7z%DynamicSliceSize.get_free_symbol_uses? s0  M:@@ TXX} 5  rc&|j|yr)codegen_dynamic_slice_sizers rrzDynamicSliceSize.codegenG rO rrurX) rf r^ rr_ rr_ rr_ rr_ rrmrxrG r` r-s@rrb rb  s~ *"  !  !  7./$) ! ! 0 1rrb cTeZdZdZddZddZ d fd Zd dZd dZxZ S) rz; The result of a call to aten._local_scalar_dense. ctSrr=rs rrzDynamicScalar.get_readsP rrcyrrrs rrzDynamicScalar.should_allocateS rrc|jt| dtt j d|j |g||_||_yrV ) rr_r>rrrr]symkeypath)rro rp rIras rr>zDynamicScalar.__init__V sM   *ELL$78$:M:Mtf:U  rc.t|jgSr)r>ro rs rrz&DynamicScalar.get_unbacked_symbol_defs` s488*%%rc&|j|yr)codegen_dynamic_scalarrs rrzDynamicScalar.codegenc s&&t,rrurX)ro r^ rp zpytree.KeyPathrIrrrrrG ) rrrr{rrr>rrr,r-s@rrrK 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_readsl rrcyrrrs rrzAssertScalar.should_allocateo rrc~t|dttjdg||_||_yrV )r_r>rrrscalarr)rrx rras rr>zAssertScalar.__init__r s7   ell51 2    rcyrrrs rr zAssertScalar.has_side_effects} rrc.t|j|Sr)r'rx r6s rr7z!AssertScalar.get_free_symbol_uses 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)NFrdzstd::to_string(r)rzif (!(z()) { throw std::runtime_error("Expected z but received " + z); }zif not (z):z raise RuntimeError(z = None)rFscalar_assertsrrr7rsrrrrcodegen_cpp_sizevarrx rrcodegen_python_sizevarr ro)rrr symbol_strsizevars rrzAssertScalar.codegen s4$$  d44454IJK 77    WW *6(!4Jgg**>> e?G    !J488*Tfgqfrrwx gg**AA eBG    4 5    7TXX7GqI J    19 :rrurX)rx rtrrrrrmrxrG ) rrrr{rrr>r r^r7rr,r-s@rrrg sM N+$)<!< !<,< ;rrc"eZdZUded<ded<y)ExternKernelNoderrzexport_schema.NoderNrrrrr r s I rr ceZdZdZ ddd dfdZdfd ZddZddZddZe dd Z dd Z dd Z dd Z d ZeddZe 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. Nr{ct|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) Nr 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 rYs 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 }t jjr||D] }|| yyt jjsJ||y)Ncjj|jjJjjr?j jt t|j|yyr) alias_namesrro alias_infois_writererfrr )rinfors r add_aliaszPFallbackKernel.__init__..handle_aliasing_and_mutation..add_alias!sj  '' 5222??++))00&z'H!TR,r)rrrr) rrrListTyperr library_utilsis_tensor_like_typer is_tensorlist_like_type)r rUr optional_tensor_argrs` rhandle_aliasing_and_mutationz=FallbackKernel.__init__..handle_aliasing_and_mutation!s$))U^^4!#e}5@tCy@500;&cE4=999{& 44TYY??/27+!"567#%88CCC#r)rZrE rr)r ztorch._C.ArgumentrUr rr)!r_r>ruse_runtime_dispatchr{rrrrrrrwrrrrrsr warn_fallbackr rc rr_libraryrvmutates_and_returns_first_argrro is_mutabler$rrfro zip_schema)rrJrrnontensor_argsrrr{schemar rr r rUras` rr>zFallbackKernel.__init__ 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 >--88vN 4ID# (s 3 4rc,t|}|jtjj j urT|jD]E}t|ts|jtj|j}G|Sr)r_r,rwr_prims rng_primsgraphsafe_run_with_rng_staterorr with_readrGrro)rrrUras rr,zFallbackKernel.get_read_writes*!syg-/   u||55RR R)) c>2"-"7"7$,,S\\^<#K  rc n|j|j|jt|ddSNr{)(codegen_unbacked_symbol_defs_for_outputsror rrs rcodegen_unbacked_symbol_defsz+FallbackKernel.codegen_unbacked_symbol_defs6!s0?? MMOT\\749Ld+S  rct|ddx}rKttjjj |}|Jt |jSt Sr rr9rsrrrr>rrr{resolveds 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..Shimr refc|jSr)r rs rrz2FallbackKernel.codegen_args..Shim.__repr__J!s xxrNrj)rrrrrrrrShimr F!s H rr )r| dataclassr4rfrrrorsrrrrwrrrr rrrrr rrr)rr rrrrparams rr zFallbackKernel.codegen_argsE!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)Nrrc3HK|]}tjd|ywr)r find_devicers rrz-FallbackKernel.find_device..q!s'$**43$rrE)rrBr rrrr ScriptObjectr(r*rrr>rrjr)rrrnon_torch_bind_tensor_argsrUdevices 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 rw)rr s rr zFallbackKernel.has_side_effects!s2))**rct|jtjjtjj fs+Jd|jdt |jdt|jtjj sSd|jjvr7|jjjrt|jrgS|jS)Nr rr r ) rrwrrrrrrrr r$r rs rrGz+FallbackKernel.get_inputs_that_alias_output!s   uzz44ejj6T6TU  2$2B2B1C2D$$%&n 6  4++UZZ-K-KL"$*:*:*?*?*AA  ((33&t'7'78I## #rcNt|jdksJ|jSr)rrc rs rrCz!FallbackKernel.get_mutation_names!s'4&&'1,,,"""rc tjd|j|jt |t sJt ||j|j|j\}}|j||}|jDcgc]}|j|fi|}}|j}tjjsg||St!dg}|j#|||} dd}t |t$j&j(j*r#|j-|d|dj.} n|j0j.} t3| dk(r>|j4r |j4n |j6} | dj8} || | g} n9t;| |j4D cgc]\} }|| j8|} } }|jJt=|jt?j@|jjC|| i}tjDjG|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) NrErr) as_tensorT)as_none) as_tensors)as_optional_tensor)as_intzUnsupported return type )rr TensorTypeNoneTyperrrr export_schemarurTensorArgumentror getElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError) return_typerXrs 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!rrE)rrfr r")rr)r z1torch.TensorType | torch.ListType | torch.JitTyperXIRNode | Sequence[IRNode]rzexport_schema.Argument)$rrrorwrrrrrfror rur rsraot_moder#serialize_inputsrr torchbind CallTorchBindr returnsrrr rerrr r r<rextern_kernel_nodesr)rrrrordered_kwargsr serializernamed_argumentsr r r r output_arguments return_schemarXrs rexport_extern_kernel_nodez(FallbackKernel.export_extern_kernel_node!sj 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*''''K ` s I3I8c"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 rr z)FallbackKernel.codegen..is_numberC"s:a!3!34$Q%5%5%788!!U%5%566rc3DK|]}j|fiywr)r )rr rrs rrz)FallbackKernel.codegen..P"s+*D))!6v6s c3lK|]+\}}t|txr|j-ywr)rcomplexr)rrr9r s rrz)FallbackKernel.codegen..U"s5,Aq1g&A9Q[[+AA,s14cHgjjSr)r r! rsrr~z(FallbackKernel.codegen..c"s$F$++-F0C0C0EFr)rz torch.JitTyperr)+rwrrrrrrrsrrrrrrrEr rFrJ custom_ops_to_c_shimsrrfrorrrur'rrrrr rr,generate_fallback_kernel_with_runtime_lookupror regenerate_fallback_kernelrJrr, r0 r5 r ) rrrrr args_iter exported_argsr 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%  $ $ ::j@jBs&|tjDjFjHurtj,dt|tjjr>tj jJs$r"tLjO||| |  }||S|tQ || | n"sJd tS || | dfd |g}t|t6tTfr |_+|St|tXrtU|_+|S|g_+|S#1swYxYw)z9Create an instance of FallbackKernel from an _OpOverloadsNr)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variantrErrrr)rJrfrorrrrwc32K|]}t|ywr)rrTs rrz(FallbackKernel.create.."s!K,s"3!Krrr{has_unaligned_inputrrr{z"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_outputrrXs rrzAFallbackKernel.create..generate_output.."s5$$F1Iw4<:K9L/LM$rbrzFallbackKernel output type z is not supported)rrrrrrrrrrrrrF assume_unaligned_fallback_outputrprsrrrrrrrrPr r r(r*rr) rXrrrrrrr r packeds `` rr 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)rXr rzlist[tuple[Any, int]]rr )-r*_fused_moving_avg_obs_fq_helper_functionalrr rsrrr rrrrrrtorch._library._out_variantr r r r rrrLrrr=rrL r_make_out_variant_kernel_namer'r rr r rq higher_orderprintrExternKernelMultiOut try_createrrmrr r)rrrrfake_incorrect_kernelscontextrrrrr{r r r r out_oprJ out_resultr rr r r s` @@@@rrzFallbackKernel.create"s#'"Q"Q!S / /1$79J9JKG!mG  <#""6;D;F;  !  < fejj33 4 ELL:   Ffnn-'/~26:!c*;F*C&D&I$)00(..0>//05^2245  '! ,"&"7!'DV'L & "!K{!KKn= vu66@@NN O//555\\%(F vuzz44 5GG''-88"3$79 J%!!  !&)"3F ?? ?6!0"3F) ) V"."5 ge} -$FN   &"7^FN&YFNk < r,r rr r~r r rGrCr rrrrrr,r-s@rrr s)-u4HLu4u4u4& u4 & u4 + u4&u4Eu4 u4n   0#2#DQ# ##J+$*#w(rS3S3j    ]]rrcZeZdZdZddZddZddd d fdZxZS) ComplexViewz9View a complex number as two dtyped numbers or vice versacyrrrs rrzComplexView.should_allocate'#rrc&|jdgSrr rs rrGz(ComplexView.get_inputs_that_alias_output*#s"##rNr c 2t||||||||y)Nr )r_r> rrJrrr rrr{ras rr>zComplexView.__init__.#s,      /  rrXrzr )rrrr{rrGr>r,r-s@rr r ##sqC$)-GK   &  &  +  & E    rr ceZdZdZddZy)MemoryCheckKernelz Custom kernel for memory checking that generates direct function calls TODO - the custom op was erroring with str inputs. should be able to custom op directly. c|j|j\}}}t|}t|}|r|jdd|d|d|d}n d|d|d}|j|y)z.Override codegen to write direct function callzV# note: dont currently distinguish between buffers returned and dealloc'd in last stepzcheck_memory_step(allocated=z, freed=z, is_final_step=rN)r4 ror r)rr alive_list dead_list is_final_step alive_repr dead_reprcalls rrzMemoryCheckKernel.codegenK#s 224/3/A/A, I}*% O    h 2*Xi[P`an`oopqD1*Xi[PQRD$rNrG )rrrr{rrrrr r D#s   rr c eZdZUded<ddZy)rmrcrc|jSrrrs rr zMultiOutputLayout.get_devicea#rrNr`)rrrrr rrrrmrm]#s rrmc~eZdZddZ d d fd Zed d d dZd dZd dZd dZ xZ S)rc|j||js#|j||j|yyr)codegen_multi_output!skip_size_stride_alignment_checksr, r0 rs rrzMultiOutput.codegenf#s:$$T*55  % %g .  * *7 36rct|d||gdtjj ||_tjj |||_||_yr) r_r>rsrr_rr`rr% )rrJr9rr% ras rr>zMultiOutput.__init__l#sQ vw3GG++D1  ""4( 1R.rcp|jd}t|tsJ||j|Sr)rfrrr7)rrers rr7z MultiOutput.get_free_symbol_usesy#s7[[^ *f-9z9-..}==rcpt|jdk(xrt|jdtS)NrEr)rrfrr rs rrzMultiOutput.should_allocate#s04;;1$ t{{1~'8 9 rc|jDcgc]>}t|tr,t|j dkDr|j @c}Scc}wr)rfrrrrGror3 s rrGz(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|dSr|r~rs rrtz*MultiOutput.get_read_writes..dummy#rrrrorru)r>rfrrrrGrrorrr rFr6rjrrPrrr) rr/rhrtrshould_normalizewrite_rwrrs @@rr,zMultiOutput.get_read_writes#s/9l;; @C#v& ,..s||~>? @}}//#002 ;"11 1 '~ '&++&&   33 4==?B2B &&??"   rrG rm) rJrr9rrlist[tuple[Any, ...]]r% rrrrxrXrzrs) rrrrr>r^r7rrGr,r,r-s@rrre#s438 S S S' S ,0 S  SM*$)>!> !>+>    rrcZeZdZdZ dfd ZeddZddZd dZxZ S) rzMultiOutput for opaque objects.c:t||||d||_y)NTr% )r_r>r)rrJr9r opaque_valueras rr>zOpaqueMultiOutput.__init__#s$ SWX$0!rctd)NzOpaqueMultiOutput has no dtype)r rs rrzOpaqueMultiOutput.dtype#s=>>rc|Srrrs rrz#OpaqueMultiOutput.wrap_for_lowering#rrcXt}|jD]E}t|ts|j t j |jGtt j |jg}t j||tSr ) r>rfrrrrGrror)rr/rhrs rr,z!OpaqueMultiOutput.get_read_writes#s.8l;; @C#v& ,..s||~>? @0:  ! !$--/ 2 30 &&"   r) rJrr9rrr. r2 r rr)rr)rrrs) rrrr{r>rrrr,r,r-s@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. cyrrrs rrz%AllocatingMultiOutput.should_allocate#rrcb|js#|j||j|yyr)r% r, r0 rs rrzAllocatingMultiOutput.codegen#s,55  % %g .  * *7 36rNrXrG )rrrr{rrrrrr7 r7 #s 4rr7 c|j}|jjjdd}|j}d|d|d|S)z8Build fully-qualified kernel name for an out-variant op.z::rEz torch.ops.r)rrrrr)r nsr' rs rr r #sP   Bnn!!''-a0G##H t1WIQxj 11rceZdZUdZded<ded< d fd Zd dZedd d  dd ZxZ S)r 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)r_r>r= r> r rrrw)rr r= rrras 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)r r r Nc3PK|]}t|tj ywr)rrrr s rrz2ExternKernelMultiOut.try_create..'$sG1:a.Gs$&r)rr{r r= r )rJr9r)r r r r rrrrrrrmrrLrrr=rr7 rrFr rprsrrrrrr> r )rrrrrrrrr{r r r r r r= r r r tensor_outrJ multi_outs rr 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) rr r rE r= rrr rrrG )rrE rr rrcrrrrvrrrrF r{r r rrz&Sequence[AllocatingMultiOutput] | None) rrrr{rr>rrr r,r-s@rr r #s 99 " "& "! "  "  "7HL$)J%JJ J & J ' J+J&JEJ"J 0JJrr ceZdZUdZded<d,dZd-dZd.dZd/dZd0dZ d1d Z d2d3d Z d4d Z d5d Z d6dZd,dZd5dZ d7 d8dZd9dZd:dZ d7 d;dZddZd?dZd@dZdAdZd,dZd,dZdBdZdCdZd1dZdCd Z d@d!Z!e"d d7 dDd"Z#dEd#Z$dFd$Z%d2dGd%Z&e'dHd&Z(dId'Z)dHd(Z*dAd)Z+e'dJd*Z,d1d+Z-e-Z.y )KrzC TensorBox / StorageBox allow in-place mutation of Tensors rrIc6|jjSrrrs rrz!MutableBox.has_exceeded_max_reads_$rrc6|jjSrrrs rr zMutableBox.get_deviceb$rrc6|jjSrrrs rrzMutableBox.make_loadere$yy$$&&rc6|jjSr)rIrrs rrzMutableBox.make_indexerh$yy%%''rc6|jjSr)rIr0rs rr0zMutableBox.get_stridek$rrc6|jjSrrrs rrozMutableBox.get_namen$rrNc8|jj|Sr)rIrrs rrzMutableBox.has_large_inner_fnq$syy++I66rc8|jj|Srrrs rrzMutableBox.mark_reuset$rrc6|jjSrrrs rrzMutableBox.realize_hintw$rL rc6|jjSr)rIrlrs rrlzMutableBox.unwrap_viewz$rJ rc6|jjSr)rIr rs rr zMutableBox.is_input_buffer}$yy((**rc6|jjSr)rIrrs rrzMutableBox.freeze_layout$syy&&((rc:|jj||Sr)rIrrs rrz*MutableBox.freeze_layout_with_stride_order$syy88 NNrc8|jj|Sr)rIr#r"s rr#z(MutableBox.freeze_layout_with_fill_order$syy66u==rc8|jj|Sr)rIr&r%s rr&z(MutableBox.freeze_layout_with_same_order$syy66v>>rc:|jj||Sr)rIr*r(s rr*z+MutableBox.freeze_layout_with_exact_strides$syy99-WWrc6|jjSr)rIr,rs rr,zMutableBox.get_read_writes$rT rc6|jjSrrrs rrzMutableBox.get_reads$rrc6|jjSrrrs rr2zMutableBox.num_reads$rrc6|jjSrrrs rr4zMutableBox.get_storage_numel$rrc6|jjSrrLrs rr9zMutableBox.get_reduction_type$rHrc6|jjSrrJrs rr;zMutableBox.get_reduction_size$rHrc6|jjSrrrs rr=zMutableBox.is_extern$rrc6|jjSr)rIr?rs rr?zMutableBox.is_no_op$rrc8|jj|SrrPrs rrAzMutableBox.constant_to_device$syy++F33rc6|jjSr)rIrCrs rrCzMutableBox.get_mutation_names$rHrc6|jjSr)rIrErs rrEzMutableBox.get_operation_name$rHrc6|jjSr)rIrGrs rrGz'MutableBox.get_inputs_that_alias_output$syy5577rc6|jjSrrrs rrzMutableBox.realize$rrc8|jj|Srrr6s rr7zMutableBox.get_free_symbol_uses$syy--m<?HMX/X@DX X +%%-..%$4..8#L)$)=!= !=*= *+3++&+$ "HrrcVeZdZeeddZeeddZeddZy)rcyrrrIs rrzTensorBox.create$sFIrcyrrrv s rrzTensorBox.create$s+.rcNt|tr|Stt|Sr)rrrrmrv s rrzTensorBox.create$s" d1 2KD)**rN)rIrrr)rIrrr)rIr)rrrrr~rrrrrr$s@ II ..++rcXeZdZdZd dZd dZd dZddZddZd dZ ddZ dd Z dd Z y )rmz7 StorageBox allow in-place mutation of Tensors ct|jttfr4|jj t j jvSyr)rrIrrOrorsrr rs rr zStorageBox.is_input_buffer%s: dii+!? @99%%'177+?+?? ?rct|jtxr4|jjtj j vSr)rrIrrorsrrrs rrzStorageBox.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)rrrrNr|)rrrIrorrrXrfrrrrr rr r rrsrr_rr`rvrzrxr})rrzrxrs 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. rEN)rrIrrXrnontrivial_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) r-r rrsrget_dep_size_hintr*rr()rrr r size_of_reads total_sizemax_sizemin_sizes r$has_accumulated_enough_reads_by_sizez/StorageBox.has_accumulated_enough_reads_by_size6%s@~~' )#. GG % %c *  ' }%}% ) # %X%* %H$  s.Bc t|jtxrh|jtj kDxsE|j xs3tjduxr|jtjSr) rrIrr2rFrealize_acc_reads_thresholdr realize_acc_reads_size_thresholdr rs rrz!StorageBox.has_exceeded_max_readsI%sq$))Y/  NN vAA A &&( 77tC==;; rcF|dkDrt|jttfr{t |jr3|jj ddg}t fd|Dry|jtjkDxs|jSy)zj A heuristic to decide if we should realize a tensor that is used multiple times. rEexpsigmoidc3:K|]}|jvywr)used_ops)rropcounts rrz5StorageBox.should_realize_on_reuse.._%s@qG,,,@sTF) rrIrrXr(rr'r2rFrealize_reads_thresholdr)rr heavy_opsr s @rshould_realize_on_reusez"StorageBox.should_realize_on_reuseU%s 19DII 9/EFdii ))446"I. @i@@ 6#A#AA-**, rcH|j|r|jyyr)r rrs rrzStorageBox.mark_reuseg%s  ' ' . LLN /rc6|jjSrrrs rr2zStorageBox.num_readsk%rrNrXr]rL)rrrr)rrrrrlrv) rrrr{r rrrr rr rr2rrrrmrm$s4  < &  $%rrmc0eZdZUded<ded<dZded<y)Subgraphrrr graph_moduleNzGraphLowering | Noner)rrrrrrrrr r o%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)rrOrlrr>)buffersrs r_has_aliased_buffersr v%s^ !+6? COG z;7;; r rsrr_rr`)rr r rJras rr>zInvokeSubgraph.__init__%sQ   ! GG++D1  ""4(rc8|jr |jgSgSr)r rs rrzInvokeSubgraph.get_subgraphs%s"&-- 7R7rcddlm}tjj}d}|j j dx}rrd}|jtjjjur7|jdtjjjusJd}|d|d}nd}|jtjjjur7|jdtjjjusJd}|j|d}|D cgc]} | j d }} |D cgc]} |j| }} g} t|D]M\} } t!| t"t$t&fr| j)| 3| j)|| || O| }|j|tjj+|j,||j. |_tj0|j5|jj2|ddd|jj4} d}|D]$} t!| t"r| j7}n|Jt9||t;|   dfd }t| Dcgc]\}}|||}}}|_|Scc} wcc} w#1swYxYwcc}}w)zFor each operand, get a realized input, force it to have the same strides as the subgraph inputs, then use an InvokeSubgraphrE)constrain_to_fake_tensorNeager_input_valsrrrrrr`r r r)r r rJc Zt|ttfr|S|j}|Jt t ||j |j|j|jj|jjt|fgdS)NrTr1 ) rrrr rrLr rr0rrMrNr)rXindrinvoke_subgraphs r create_outputz,InvokeSubgraph.create..create_output%s:N"OP **,)))"%$..0#__.%002%00299"("3"3"5"?"? $C[M6:  r)rXrr rrz:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput)rr rsrrYr^rrrrqr rrr rrrrrrrr r rr r graph_outputsr r rmr )rr r r rY fake_operandsr rM fx_operandsr new_operandsroperandr rr rrXoutsr 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 rrJrmrrrQ)r r r rrz@list[ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput]rG )rrrr{r rr r r>rrrrr,r-s@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] | Noner c||_||_||_||_t |g|\}}t |d||||||_tjj||_ tjj|yN)rrJrfro) r r r r _split_by_sym_typer_r>r{rsrr_rr`) rr r r r rJr{sym_argsrras rr>zConditional.__init__&s#  *, 2I3I3I J+ "   (%6D "GG++D1  ""4(rcg}|jr|j|j|jr|j|j|Sr)r rr r subgraphss rrzConditional.get_subgraphs4&sG      T// 0      T00 1rcRt|tr|S|jjSr)rrrrP)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)r.rcg}t||D]e\}}t|tr|j|(|jtj t ||jdg|SNFr)rrrrrrrr)r fake_tensorsretrXr}s 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 rr{zcannot determine devicer)r r r r rJr{rr)-check_input_alias_and_mutation_return_outputs)r rr zSequence[torch.Tensor]rr)0rrsrrYrrrrr<rr^r r rr r r r r.rrrr r rrMrr9rrrr rmrrLrr rrNrr torch._higher_order_ops.utilsr r>rrfrJre)rr r r r rr r fx_op fake_outputsr r true_outputs false_outputsrr rt_of_orr{ conditionalrX merged_outputr3r rtrue_mutated_inputsfalse_mutated_inputsmutated_operand_indicesrs rrzConditional.createB&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 ror rrs 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 rJrmr{r rrrQ)rzint | torch.SymIntrzint | sympy.Expr) r rr r r r r rrzlist[MultiOutput]rG r)rrrr{r rr r r r r>rr~r rrrrr,r-s@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)rrrrP)r non_sym_argsr rUs 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 r c.||_||_||_||_t g||\}} t |d|| ||||_||_tjj||_ tjj|yr )r r r r r r_r>r{ stack_outputrsrr_rr`) rr r r r rJr{r r rras rr>zWhileLoop.__init__'s-!2** 2 1n 10 1! + "   (%6D "(GG++D1  ""4(rcg}|jr|j|j|jr|j|j|Sr)r rr r s rrzWhileLoop.get_subgraphs%'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 rrOrlr>rrrrrr)r runwrapped_buffers seen_buffersr1original_inputunwrapped_buffers r_clone_aliased_inputszWhileLoop._clone_aliased_inputs0'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: )rrrmrOrrr r)rs r_maybe_wrap_as_tensor_boxz#WhileLoop._maybe_wrap_as_tensor_boxJ'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 ) rrrrrr r rrrr) tensor_boxesr r rfknew_tbs rr z0WhileLoop.create.._require_exact_stridesc's|$L(99 99ClL9 #Bb%,,/'@@DFJJ$::"BIIKu; JJrN) #*Jrr.rNr zOutput aliasing is currently not supported in compiled torch.while_loop. The outputs of the body_fn subgraph of torch.while_loop are aliased: rEz9torch.while_loop is assumed to have at least one operand.ct|t|k(sJt||D]/\}}tjjj ||1yr)rrrsrrrY) lhs_exprs rhs_exprslhsrhss r_guard_list_equalsz,WhileLoop.create.._guard_list_equals'sN9~Y777 #Iy 9r r rsrrYrr^rr r rrrr r rr r rr r r.rr rrrr rrr0r9rrrrmmodulefx GraphModuler>rr rerrLrrrr r rrrrrfrJrrMror r# r)*rcond_fnbody_fnr r r r r fx_carried_inputsfx_additional_inputs fx_all_inputsrfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_ all_inputsr cond_outputs body_outputsr.rrrZbor r{ while_loop mutated_idxsmutated_idx_setrrdmutated_inputs_iter all_outputsrXr3r2rE mutated_inputrhrs* rrzWhileLoop.createU'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 <5 <5 < < 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_loopr r ror rrs rrzWhileLoop.codegen/(sA""4):):;88 MMOT\\749Lb+Q rct|ddx}rKttjjj |}|Jt |jSt Sr r r s rrz"WhileLoop.get_unbacked_symbol_defs5(r r)r rr rr r r r rJrmr{r r rrrrQ)r rrr)rrrr) r r r r r rr rr rrr rG r)rrrr{r rr r r r r>rr~r r rrrrr,r-s@rr r &s ].2N+215.5%)M?)%)M?),0G )0)(),) )  ) " )E)) )>2LLWWW) W , W  W #WWr  rr c\eZdZ ddd dfdZdfd Zd dZxZS) rNr c t ||||||d|ddlm}||} | J| |_t j jj| d|_ |t j j| <y)Nr r) _get_effect) r_r>torch._higher_order_ops.effectsr effect_typersr effectful_opsrprev_effect_buffer) rrJrrr rrr{r r ras rr>zEffectfulKernel.__init__A(s      /   @!&) &&&&"#''"7"7";";K"N-1k*rct|}|jF|jj t j |jj|Sr)r_r,r r/rrGrro)rrras rr,zEffectfulKernel.get_read_writes^(sUg-/  " " .    ! !$$T%<%<%E%E%GH rcyrrrs rr z EffectfulKernel.has_side_effectsh(rrrr rsrX)rrrr>r,r r,r-s@rrr@(sr)-2HL222& 2 & 2 + 2&2E2 2:rrc4eZdZed d ddZy)rctSrr=r6s rr7z!NonTensorObj.get_free_symbol_usesm(rrNrmrx)rrrr^r7rrrrrl(s,N+$)! !,rrcLeZdZUded<ded<d dZd d dZddZdd Zdd Zy)rBrr%FakeScriptObject | torch.ScriptObjectrc|jSrrrs rrozTorchBindObject.get_namey(r~rNc|jSrrrs rrz!TorchBindObject.codegen_reference|(r~rc|jSrrrs rrzTorchBindObject.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+ rrrrrr numelrroperatorr)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 rjrr_)rr# )rztorch.ScriptObjectrv) rrrrrorrr) r2 rrrrBrBt(s& I 00' =rrBc.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. r rc,t|jSr)r rrs rroz OpaqueValueTypeConstant.get_name(sDJJrNct|j\}}|jD]"\}}|tjj |<$|Sr)r)rrrsropaque_value_type_classes)rrobj_repr opaque_typesrrs rrz)OpaqueValueTypeConstant.codegen_reference(sJ!4TZZ!@, &&( 5DAq34AGG - -a 0 5rrjrr_rrrr{rrorrrrr4 r4 (s J rr4 c4eZdZUded<ded<ddZd d dZy) rrrrcrc|jSrrrs rrozGeneratorState.get_name(r~rNc|jSrrrs rrz GeneratorState.codegen_reference(r~rrjrr_)rrrrrorrrrrr(s I rrc8eZdZUdZded<ded<d dZd d dZy) rz 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. rrr rc|jSrrrs rrozOpaqueObjectState.get_name(r~rNc|jSrrrs rrz#OpaqueObjectState.codegen_reference(r~rrjrr_r: rrrrr(s I JrrcreZdZddZddZdd dZe d dZe d dZy) _CollectiveKernelcyrrrs rrz!_CollectiveKernel.should_allocate(rrcyrrrs rr 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) rrwrrrrsrrrrru)rrsrrs rr~z%_CollectiveKernel.set_cpp_kernel_name(sD$$%)>)>> : >!!  D #)>>#6#6D #NN44.  AFF. *. s B)B)c tjj5|j||g|i|\}}}}} ddd r J|d| D]?} | j tjj | j A|dj} |t| ||} tj|} | jj| Dcgc]}tt| || c}| jj| Dcgc]}|j c}d|vr`| jjtt| |d| | jj|dj yy#1swYtxYwcc}wcc}w)Nr)rrr)rsrrrrrror rr tree_leavesrer4rfr r)rrrfrr_example_outputrrrr{ tensor_argrr inpsrrhs rcreate_inplacez _CollectiveKernel.create_inplace(sWW   D#""66CDCFC  !  D%E2C1D&EE$% ?J    GG ' ' (;(;(= > ?Q**, f %      !!&)&&OS T^Jf5sF C T !!T"Bc3<<>"BC F?  # # * *z8&-P     % %fUm&<&<&> ? ; D D0 U#CsF+*F8(F=+F5c 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)Nr r)rsrrrrrBrrr rmrrrr rrFr rprrr)rrrfrrrrrrr{rI rr rrPrs rcreate_out_of_placez%_CollectiveKernel.create_out_of_place,)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-rXrrH ) rrrfzIRNode | list[IRNode]rr rr rr) rrrfzTensorBox | list[TensorBox]rr rr rz%list[MultiOutput] | _CollectiveKernel) rrrrr r~rrK rM rrrrB rB (s  ()@)@&)@ )@  )@  )@)@B99,9 9  9 / 99rrB cPeZdZ ddd dfdZddZxZS)_AllReduce_KernelNr c Tt||||||d||jdy)Nr +aoti_torch_cpu__c10d_functional_all_reduce_r_r>r~r s rr>z_AllReduce_Kernel.__init__j)=      /     !NOrc|jd|j|t|jtr|j |yyNz+torch/csrc/inductor/aoti_torch/c/shim_cpu.hinclude_extra_headerr\ rrJrr, rs rrz_AllReduce_Kernel.codegen)@$$%RS,,T2 dkk6 *  % %g . +rrr rG r r-s@rrO rO i)sw)-PHLPPP& P & P + P&PEP P,/rrO cPeZdZ ddd dfdZddZxZS)_AllReduceKernelNr c Tt||||||d||jdy)Nr *aoti_torch_cpu__c10d_functional_all_reducerR r s rr>z_AllReduceKernel.__init__)s=      /     !MNrc|jd|j|t|jtr|j |yyrU rV rs rrz_AllReduceKernel.codegen)rX rrr rG r r-s@rrZ rZ )sw)-OHLOOO& O & O + O&OEO O,/rrZ cveZdZ ddd d fdZd dZd dZed dZd fd ZxZ S) _WaitKernelNr c Tt||||||d||jdy)Nr +aoti_torch_cpu__c10d_functional_wait_tensorrR r s rr>z_WaitKernel.__init__)rS rc|jd|j|t|jtr|j |yyrU rV rs rrz_WaitKernel.codegen)rX rcr|jd}t|tsJt|tr2|jd}t|tsJt ||gSt|t rC|jd}t|tr"|j d\}}|j|gSgSgSr)rfrrrB rrr)rrhrcollrrs 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)Nr)r) rsrrrrr rerrf) rrrhrH rrrr{r 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)r_r,re r/rrGrro)rrvolatile_readsvrras rr,z_WaitKernel.get_read_writes)sZg-/ 002  GB    ! !,"6"6r{{}"E F Grrr rG rn )rrrhrrrrs) rrrr>rre rrg r,r,r-s@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;rr4rrr>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;rr3rrr>rr!rA rrrl s rrA rA *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)Nrzr) rrrIrmrrrnrrrrf)r1rs rassign_origin_nodero *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)rrrr&)rrorz&Callable[[Sequence[_T]], Sequence[_T]])rz&Callable[[Sequence[_U]], Sequence[_V]]rz&Callable[[Sequence[_T]], Sequence[_U]]rrvr)rz#Sequence[int | torch.SymInt | Expr]rzShapeEnv | Nonerro)rSequence[int | Integer]rrorm)rrrrrr)rrrrrr9)rr"rrrztorch.Tensor | None)rzSequence[_T] | NonerzSequence[_T | None] | None)rz/IRNode | OutputSpec | torch.device | None | strrr^)rz"IRNode | torch.device | None | strrr)rzBuffer | TensorBoxr-rrr)r;rir<rir=rirr)rPrrQzSequence[int | torch.SymInt]rr)r`r rr)rfrrr*)rzExpr | Sequence[Expr]rrUrrrrR)rErrrUrDrrr')rrorrI rMrrrgr )TFNFN)rrrrrrrrnrrr)rnrztuple[StorageBox, Layout])rrrrp rrrJ)rrir=rirr)rrUrr)rr6rzTypeIs[Sequence[IRNode]])r rE rr)r rrr)rr rz-tuple[list[ShapeAsConstantBuffer], list[Any]])rrrr)r1r rz torch.fx.Noderr(\ __future__rrrRr|rrloggingr- textwraprxcollections.abcrrrrrr r enumr r typingr rrrrrrrrrrtyping_extensionsrrrrrr unittest.mockrrrr r!torch._export.serde.schema_exportserder r r r rvr torch._loggingrtorch.fxtorch.utils._pytree_pytreertorch._dynamo.utilsr"torch._export.serde.serializer#*torch._higher_order_ops.auto_functionalizer$torch._inductorr%r-r'"torch._library.fake_class_registryr(torch._library.opaque_objectr)r*torch._prims_commonr+r,r-r.r/r0%torch.fx.experimental.symbolic_shapesr1r2r3r4r5r6r7r8r9r:r; torch.fx.noder<torch.utils._ordered_setr>torch.utils._python_dispatchr?torch.utils._sympy.functionsr@rArBrCtorch.utils._sympy.symbolrDrrFrGcodegen.commonrHrIrJrKrLrMrNrOrPrQrR loop_bodyrS ops_handlerrTrUrVrWruntime.benchmarkingrX runtime.hintsrYrZr[r\r]r^r_r`rarbrcrdrerfrgrhrirjrkrlrmrnrorp virtualizedrqrrrsrtru torch.typesrvcodegen.cutlass.templaterwcodegen.wrapperrxrryrzrrr! __version__rr ImportErrorr{r|r}r~rrrrrrrr getLoggerrrrrrrrrr rrrrrrrKrMrrrrrrr&r(r8rArVrbrirprrrrrrr-rVrXrr4r6rArIrfrrrrrKrrrkrrr%r0rrOrrrrrrrrrrLr rpr|rrrrnrrrrrrrrVrrrrPrimitiveInfoTyperrrrr rrr4rrKrOrrL rT rirfrp rt ru r r r r r r! r& r. r; rG rR rb rrr rr r rmrrr7 r r rrrmr r r r r r rrrBr4 rrrB rO rZ r_ r@ rA ro rrrr sO" MM:    UT ''22,,$$(?M#2?M    /?QQ*" 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';S,S,l UUUUpF FF F R&% % % P i  F | $ y ! y !   u  u =)<8JN<N<N +<NBF<N<N~d  d  d  R$(1:   & &"8D>8D>"BH"LM YM7S97St#1#L` +` FF 5F F T   [5[[| !3748 9 9 991 9  9 2 99x 4 $^ v^ ^ B[[[|-(--`A9(A9A9H!(!!H[;[[|ShSSl&&&RlAlA^6"K|KK$ S| S S '9 (< 7 7 ` Z` ` F C&C[HV[H|!Gf!GHT%{%D.V*V*rUEV]EEPU &fi & & & K [ &6( F   Ut4_t4t4n d#$]+_]+@<><~%K$&,tC#I4E4L/MMLL^"|" \B.\B~5N50((6N>Z%NZ%z5 .55 UR?RRj K9K\UE<EEPU"l""J / ((" ("V V B=L=@% % P - $< \< ~WlWt ))|))Z/,/d<8B5B$8-826)l6)r))|))X9T9Tx,E,E^21|21j-L-8<;<<;~U t &t nU .  @  2 I ,I X   F 4K 4 2f>fV TTTn + +n%n%bU'v'' LUB.\B.B.JUe ,e e P "  "2 "U@ @ @ F )n)X6"=l"="=Jl(\ "YYx/)/>/(/>S#Sr   TiLNJs!l33 m?m