+ jya0t$^RIHt^RIt^RIt^RIt^RIt^RIt^RIt^RI t ^RI t ^RI t ^RI H t HtHtHtHt^RIHtHt^RIHt^RIHt^RIHt^RIHtHtHtHtHtHtH t H!t!H"t"H#t#H$t$H%t%^R I&H't'H(t(H)t)H*t*H+t+H,t,H-t-^R I.H/t/^RI0t0^R I0H1t1H2t2H3t3^RI4H5uH6uH7t8^RI9H:uH;t;^RI^RI?t@^RIAt@^RIBH=uHCtD^R IEHFtF^R IGHHtH^RIIHJtJ^RIKHLtL^RIMHNtN^RIOHPtP^RIQHRtRHStS^RITHUtUHVtV^RIWHXtX^RIYHZtZH[t[H\t\H]t]H^t^H_t_^RI`HataHbtbHctcHdtdHeteHftfHgtgHhthHitiHjtjHktk^RIlHmtm^RInHoto^RIpHqtq^RIrHstsHtttHutuHvtvHwtwHxtx^RIyHztz^RI{H|t|H}t}^RI~HtHtHtHtHt^RI}HtHtHtHtHtHt^RIHt^R IHtHtHtHt^R!IHt^R"IHtHt^R#I=HtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHt^R$IHtHtHt]"'d2^R%I`Ht^R&IlHt^R'IHt^R(IHt^R)IHt^R*IHt^R+IHt^R,I=HtM]tR-]R.&^RIt]EPtR/t]*!R14t]$!R24t]$!R34t]$!R44t]]1,tR-]R5&]],]1,tR-]R6&]@EPEP]@EPEP,tR-]R7&!R8R9],R0R:7t]EP!]4t]P,!] EPR;R<7t]@EP\EPt]%]R=]]R=3,R>R?]]%]]]R=3,R=R>R?3,R,,3,tR-]R@&RARBlt]EP!R/RC7!RDRE44tRFRGltRHRIltRJRKltRLRMltRNROlt.EROt.EROtERRPRQlltRRRSltERRTRUllt]ERRVRWll4t]ERRXRYll4tERRZR[lltR\R]ltR^R_ltR`RaltRbRcltRdReltRfRgltRhRiltRjRkltRlRmltRnt!RoR?4t]!R0RC7!RpRq44t]!RrRs]44tRtRult]!RvRw]44t]!RxRy]44tRz]!R{4R|]!R}4R~]!R4R]!R4R]!R4R]!R4R]!R4R]!R4/tR]R&ERRRllt]!RR]44tR]2!^43RRllt] ]]1,]]1,.]3,tR-]R&!RR]4t!RR]4t!RR]4t!RR]4Et]!RR]44Et]!RRE]44Et]!RR]44EtRRlEtRRlEtRRlEtERRRllEtRRlEtRRlEt ]!RR]44Et ]!RRE] 44Et ]!RRE] 44Et ]!RRE] 44Et ]!RRE] 44Et]!RRE]44Et]!RRE] 44Et]!RRE] 44Et!RRE]4Et]!RR]44Et]!RRE]44Et]!RRE]44EtRRlEtRRlEt!RR4Et]!RRE]44Et!RRE]4Et!RRE]4Et!RRE]4Et!RR]4Et!RRE]4Et]!RRE]44Et!RRE]4Et ]!R0RC7!RR]]44Et!]!R0RC7!RRE]!]44Et"!RRE]!4Et#!RRE]#4Et$!RRE]#4Et%]!RR]44Et&]!RR]44Et']EP!R/RC7!RR44Et(]!R0RC7!RRE]"44Et)]EP!R/RC7!RR44Et*!RR]4Et+!RRE]"4Et,!RRE],4Et-]],E].,],E]/]],],E].,,,Et0!RR4Et1!RRE]14Et2REs3R]R&RRlEt4!RRE]-4Et5!RRE],4Et6!ERERE],4Et7!ERERE],4Et8!ERERE],4Et9ERERlEt:]!R0RC7!ERER E]"44Et;!ER ER E];4Et<!ER ER E]<4Et=]EP|ERERl4Et?ERERlEt@]EP!R/RC7!ERER44EtA]!R0RC7!ERERE];44EtB]!R0RC7!ERERE]B44EtC!ERERE]C4EtD!ERERE]B4EtE!ERERE]!4EtF!ERERE]<4EtG!ER ER!E]B4EtH!ER"ER#E]H4EtI!ER$ER%E]H4EtJ!ER&ER'E]B4EtK!ER(ER)E]B4EtL!ER*ER+E]B4EtM!ER,ER-E]B4EtN!ER.ER/E]B4EtO!ER0ER1E]O4EtP!ER2ER3E]E4EtQ!ER4ER5E]E4EtR!ER6ER7E]B4EtS!ER8ER9E]B4EtT!ER:ER;E]C4EtU!ER<ER=E]B4EtV!ER>ER?E]B4EtW!ER@ERAE]B4EtX!ERBERCE]B4EtY]!R0RC7!ERDERE44EtZ!ERFERGE]E4Et[]!R0RC7!ERHERIE][44Et\!ERJERKE][4Et]]!ERLERME]44Et^!ERNEROE]B4Et_!ERPERQE]_4Et`!ERRERSE]_4EtaERTERUlEtb!ERVERWE][4Etc]EP!ERXERY]44Etd!ERZR=E]d4Ete!ER[ER\E]d4Etf]!R0RC7!ER]ER^]44EtgER_ER`lEth]!R0RC7!ERaERbE]B44EtiERcERdlEtj]!R0RC7!EReERfE]B44EtkERgERhlEtl]!R0RC7!ERiERjE]B44Etm!ERkERlE][4Etn!ERmERn]4Eto]!ERoERpE]o44Etp]!ERqERrE]o44Etq]!ERsERtE]o44Etr]!ERuERvE]o44Ets!ERwERxE][4Ett!ERyERzE]t4Etu!ER{ER|E]t4Etv!ER}ER~E]t4EtwERERlEtxERERlEtyERERlEtzR# ]d RtR0tE Li;i() annotationsN)Callable GeneratorIterableIteratorSequence)AbstractContextManager nullcontext)Enum)partial)Number) AnycastClassVarLiteraloverloadProtocol SupportsFloat SupportsInt TYPE_CHECKING TypeAliasTypeVarUnion) assert_neverNeveroverride ParamSpecSelf TypedDictTypeIs)patch)ExprIntegerSymbol)enable_python_dispatcher)identity)GraphModuleSerializer)can_auto_functionalize)metricsget_free_symbols)FakeScriptObjectmaybe_to_fake_obj)get_opaque_obj_repris_custom_class_obj)inner_tree_plan)compute_required_storage_lengthis_boolean_dtype(is_contiguous_for_memory_format_or_falseis_float_dtypemake_channels_last_strides_for StrideType) &_remove_effect_token_unbacked_bindingscompute_unbacked_bindings free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNodehas_free_unbacked_symbols IterateExprsrebind_unbackedresolve_unbacked_bindingsShapeEnvSymTypes)Node OrderedSet)_disable_current_modes)CleanDivFloorDivMaxMinModModularIndexing)SymT)config dependencies)BackendFeature CodegenSymbolget_scheduling_for_deviceindex_prevent_reorderingKernel)Depextract_free_symbols#extract_input_node_reduction_rangesextract_read_writesSymbolUsageCollectorOpsHandler var_builder)LoopBody) OpCounterCSE OpCountResult ReductionType StoreMode) benchmarker)DeviceProperties ReductionHint)argsort argsort_sym cache_on_selfcache_on_self_and_argsceildivconvert_shape_to_inductorconvert_shape_to_symintdeveloper_warningdo_bench_using_profilingdtype_from_sizeget_dtype_sizeget_kernel_metadataGPU_ALIGN_BYTES ir_dataclass is_dynamicis_gpu sympy_dotsympy_index_symbolsympy_index_symbol_with_prefix sympy_product sympy_substensor_is_aligned VarRanges)opsOpsValueV) SympyBoolean)Argument) IntLikeType)CUTLASSTemplate)PythonWrapperCodegen) GraphLowering)GemmReductionPlan)IndentedBufferrrTF_P_T_U_V_IntLike_NumLike _OpOverloadsc6]tRt^t$R]R&R]R&R]R&RtR#) ArgPropertystrname torch.JitTypetypeobject default_valueN__name__ __module__ __qualname____firstlineno____annotations____static_attributes__rS/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/ir.pyrrs I rr)totalz prefix TensorBoxr$IRNode _NodeOrNodesc V^8dQhRRRR/#)xrreturnzTypeIs[int | Integer]r)formats"r __annotate__rs))&)2)rc.\V\\34#N) isinstanceintr#rs&r _is_staticrs a#w ((r)frozencT]tRt^t$R]R&R]R&R]R&R]R&R ]R &R ]R &R tR#)GraphPartitionSignatureOrderedSet[sympy.Symbol] symbol_inputsz*dict[str, IRNode | Expr | TorchBindObject] input_nodes list[IRNode] output_nodeszdict[str, bool]input_deallocationboolskip_cudagraph list[str]constant_namesrNrrrrrrs.,+<;('rrc V^8dQhRRRR/#)r node_or_nodes_NodeOrNodes | NonerNoner)rs"rrr s!$!$2!$t!$rc*aRV3RlloS!V4R#)c V^8dQhRRRR/#)rnodesrrrr)rs"rr!validate_ir..__annotate__ s 3rc <VfR#\V\\34'dVF pS!V4K R#\V\4'd"VP 4F pS!V4K R#\V\ \ \\\PPP\\\\ \"3 4'g\%R\'V4 R24hR#)NzFound zE, which is not a supported top level IR node. See [Note: Inductor IR])rlisttupledictvalues ExpandView DynamicScalar AssertScalarrsympylogicboolalgBooleanr"rEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputAssertionErrorr)rnode_check_tensorboxs& rr%validate_ir.._check_tensorbox s =  e} - - & t $ $  &'! KK''//#)% %T%[M)norNr)rrs&@r validate_irr s@]#rc V^8dQhRRRR/#)rrrrCallable[..., OpsValue]r)rs"rrr-sc5rcna\S\4'g\\S44hRV3RllpV#)c$V^8dQhRRRRRR/#)rargsrkwargsrrzr)rs"rr!ops_wrapper..__annotate__1s!33&3F3x3rc0<\\S4!V/VB#r)getattrry)rrrs*,rfnops_wrapper..fn1ssD!42622r)rrrr)rrsf r ops_wrapperr-s0 dC T$Z((33 Irc V^8dQhRRRR/#rorder Sequence[int]rz&Callable[[Sequence[_T]], Sequence[_T]]r)rs"rrr7s=-Src ha\\V\\V4444oRV3RllpV#)c V^8dQhRRRR/#rindex Sequence[_T]rr)rs"rr%inverse_reorder..__annotate__:s@@|@ @rc<\V4\S48wd \R4h\\V44Uu.uFqSV,,NK up#uupi)z%Expected len(index) == len(inv_order)lenrrange)ri inv_orders& rreindex inverse_reorder..reindex:sL u:Y ' !HI I-23u:->?->il##->???A)rziprr)rrrs& @rinverse_reorderr7s0Sc%j 123I@@ Nrc V^8dQhRRRR/#rr)rs"rrrBs *PrcaRV3RllpV#)c V^8dQhRRRR/#rr)rs"rr"same_reorder..__annotate__Cs<<|< .reindexCsJ u:U # !DE E).s5z):;):AeAh):;;;rr)rrsf r same_reorderrBs<< Nrc$V^8dQhRRRRRR/#)rreindex1z&Callable[[Sequence[_U]], Sequence[_V]]reindex2z&Callable[[Sequence[_T]], Sequence[_U]]r&Callable[[Sequence[_T]], Sequence[_V]]r)rs"rrrKs&44,rcaaRVV3RllpV#)c V^8dQhRRRR/#)rrrrz Sequence[_V]r)rs"rr%fuse_reindexing..__annotate__Os))|) )rc <S!S!V44#rr)rrrs&rr fuse_reindexing..reindexOs((rr)rrrsff rfuse_reindexingr Ks)) Nrc$V^8dQhRRRRRR/#rseqz#Sequence[int | torch.SymInt | Expr] shape_envShapeEnv | Nonerrr)rs"rrrYs$   , 9H  rcVe;\;QJdRV4F 'dK RM RM !RV44'd\V4pV#\W4pV#)z) Convert strides to fill order (argsort) c3b"TF%p\V\\P34xK' R#5irrrrr#.0ss& r !get_fill_order.._s#QS 1sEMM.B C CS-/FT)allrbrc)r r  sorted_idxs&& rget_fill_orderrYsM CCQSQCCCQSQQQ$+CL  !0 rc V^8dQhRRRR/#)rrSequence[int | Integer]rrr)rs"rrrgs  #: } rc\V4UUu/uFwrW!bK ppp\\V44Uu.uF qCV,NK ppV#uuppiuupi)zx 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_orderr#gsT(1'7 8'783ch'7F 8%*3u:%67%6))%6J7 97s A Ac$V^8dQhRRRRRR/#r r)rs"rrrss$   , 9H  rc\W4p\\V44Uu.uFp^NK pp\V4F wrVWTV&K V#uupi)z! Convert strides to stride order )rrrr)r r r_outrelems&& rget_stride_orderr)ssL !/s >JCHo &o1oC &Z(D ) J 's A c$V^8dQhRRRRRR/#)rrrreplace_symbols_with_hintsrrr)rs"rrrsUUU4UDUrcR#rrrr+s&&rir_node_to_tensorr.sRUrc$V^8dQhRRRRRR/#)rrrr+rr torch.Tensorr)rs"rrrs$ +/rcR#rrr-s&&rr.r.src$V^8dQhRRRRRR/#)rr IRNode | Noner+rrztorch.Tensor | Noner)rs"rrrs$   26  rc VfR#V'd&\PPPpM\pVP 4Uu.uF q2!V4NK pp\ V4'd0VP4PUu.uF q2!V4NK ppM\PV4pVP4pVP4p\V4p\V4p\PPPP4;_uu_4\ P"!WEWgR7P%4pRRR4V#uupiuupi +'giX#;i)N)sizestridedtypedevice)r{graphsizevarsoptimization_hintr&get_sizeis_storage_and_layout get_layoutr6FlexibleLayoutcontiguous_strides get_dtype get_devicerhr suppress_guardstorch empty_stridedzero_) rr+shape_fnrr5r6r7r8ts && rr.r.s y"77##55!" .AHQKD .Q'(||~'<'<=' 6 5 Hs E E#'E E/ c V^8dQhRRRR/#)rvaluezSequence[_T] | NonerzSequence[_T | None] | Noner)rs"rrrs rcH\V\4'd V'gR.#V#r)rrrJs&rmay_convert_to_optionalrMs!%uv Lrc V^8dQhRRRR/#)rrz/IRNode | OutputSpec | torch.device | None | strr str | Noner)rs"rrrs > >6 > >rcJ\V\4'gVfV#\V\P4'd VP#\V\ \ 34'd\VP44#\RV R\ V4P R24R#)Nzget_device_type(: )) rrrDr8rr OutputSpecget_device_typerBrrrs&rrTrTsz!SQY Au|| $ $vv A + , ,q||~..#A3ba)9)9(:!<=rc V^8dQhRRRR/#rrz"IRNode | torch.device | None | strrrr)rs"rrrs;;3;;rc\V4pVR9d\\V R24R8XdR#R#Ve\V4;pfR#^RIHp\ V\4'g\\V44h\W#4#)cpu_backendtritonTF)TritonScheduling)rXcudaxpu) rTrrMrQcodegen.tritonr[rrr issubclass)rr8device_schedulingr[s& r is_tritonrasz Q F'' 6fXX. /8 ;!:6!B B  K0 ' . .T"3455 ' ::rc V^8dQhRRRR/#rVr)rs"rrrs''0'T'rc\V4R8H#)rX)rTrs&ris_cpurds 1  &&rc$V^8dQhRRRRRR/#)rrzBuffer | TensorBox alignmentrrrr)rs"rrrs"77"4777rca\V\4'dQVP4e?\VP 44'g \VP 44'dR#\ P!V3RlVP 4RR4!p\ P!\ P!VP 4R,^4\ P!VP 4R,^44p\ P!W#4p\PPPV4#)NFc3f<"TF&p\P!\VS4^4xK( R#5irN)rEqrJ)rrrfs& rr-is_aligned_realized_tensor..s( F2EQ%((3q)$a ( (2Es.1)rrmaybe_get_strider: get_strider<rAndOrrjLer{r9r:guard_or_false)rrfaligned_stridesaligned_last_dim is_aligneds&f ris_aligned_realized_tensorrvs q& ! !    '  0 0  . .ii F!,,."2E FOxx #Q'!**,r2BA)F?=J 77   * *: 66rc(V^8dQhRRRRRRRR/#)rstrides1Sequence[_IntLike]strides2shaperrr)rs"rrrs0    rc\V4\V48Xd\V4\V48Xg \R4h\;QJdRV4F 'gK RM RM !RV44'dR#\W V4F~wr4p\P P PV^4'dK8\P P P\P!WE44'dK}R# R#)z Returns true if the strides are equal, ignoring dimensions of size 0 or 1. If any dimension is size 0, all strides are insignificant since the tensor is empty. zGExpected len(shape) == len(strides1) and len(strides1) == len(strides2)c3v"TF/p\PPPV^4xK1 R#5irir{r9r:statically_known_equals)rdims& rr,significant_strides_equal..s* Mu177   3 3C ; ;u79TF) rranyrr{r9r:statically_known_leqrrrrj)rxrzr{rs1s2s&&& rsignificant_strides_equalrs J#h- 'CMS],J U   s Mu Msss Mu MMM5H5  77   0 0a 8 8 ww..uxx/?@@ 6 rc$V^8dQhRRRRRR/#)rtensorrstrideszSequence[int | torch.SymInt]rr)rs"rrr s+%G%G %G )%G %Grc,\V4'gV#\;QJd3R\WP444F 'dK RM( RM$!R\WP4444'dV#\ WP4VP 44'gV#\ V4wr#.VPOpVP4p\VP 44FIwrgV'g3\PPPV^4'gK?W,WF&KK \VPVP VP"VVP$VP&4p\)\+W(R74#)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 c3x"TF0wr\PPPW4xK2 R#5irr~rrrs& rr2try_match_insignificant_strides..s1 7FB 008878:FTdatalayout)r=rrrnrr<as_storage_and_layoutr6is_zero_elementsrr{r9r:r FixedLayoutr8r7r5offset is_pinnedrReinterpretView) rrstorage old_layout new_strideis_emptyrr new_layouts && rtry_match_insignificant_stridesr s? ! ( (  s '#4#4#67 sss '#4#4#67  $W.?.?.A6??CT U U /7G%:$$%J&&(H&//+, qww''<rc&\V\4'd VPp\V\4'dVP 4p\V\ 4'd VPp\V\ 4'dVP4#R#r)rrrBaseView unwrap_view StorageBoxBufferget_namers&r try_get_namerEsc!Y FF!X MMO!Z  FF%a001::<:d:rc*]tRtRt$Rt]!4tR]R&RtR]R&Rt R]R &] P!R R 7t R ]R &] P!R R 7t R]R&] P!R R 7tR]R&] P!R R 7tR]R&] P!R R 7tR]R&] P!R R 7tR]R&]]P(RRl44t]]P(RRl44t]]P(RRl44t]RRl4tR R!ltR"R#ltR$R%ltR&R'ltR(R)ltR*R+ltR,R-ltR.R/lt R0R1lt!RR3R4llt"RR5R6llt#R7R8lt$R9R:lt%R;R<lt&R=R>lt'R?R@lt(RARBlt)RCRDlt*RERFlt+RGRHlt,]-RIRJl4t.RKRLlt/RMRNlt0RORPlt1RRQRRllt2RSRTlt3RURVlt4RWRXlt5RYRZlt6R[R\lt7R]R^lt8R_R`lt9RaRblt:RcRdlt;ReRfltRlRmlt?RnRolt@RpRqltARRrRslltBRtRultCRvRwltDRRxRylltERzR{ltFR|R}ltGR~RltHRRltIRRRlltJRRltKRRltLRRltMRRltNRRltORRltPRRltQRRltR]S'd]-RRl4tTRtUR#RtUR#)riOzBase 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_idxz ClassVar[tuple[int, int] | None]_current_mempoolF)initOrderedSet[Any]originslist[str] | None tracebacktorch.fx.Node | None origin_nodezdict[str, object]r int | None stream_idxtuple[int, int] | Nonemempoolc V^8dQhRRRR/#)rrzOrderedSet[Node]rGenerator[None, None, None]r)rs"rrIRNode.__annotate__is**!1*6Q*rc# "\PpW,\nRxV\nR# T\ni;i5ir)rr)rolds& rcurrent_originsIRNode.current_originsgs4%%"%- * &)F #cF #s"A6 A AAc V^8dQhRRRR/#)rrrrrr)rs"rrrss--- $-rc# "\PpV\nRxV\nR# T\ni;i5ir)rr)rrs& rcurrent_stream_idxIRNode.current_stream_idxqs2 ((%/" - ),F &F &A0 A =Ac V^8dQhRRRR/#)rrrrrr)rs"rrrs**'* $*rc# "\PpV\nRxV\nR# T\ni;i5ir)rr)rrs& rcurrent_mempoolIRNode.current_mempool}s2 %%") * &)F #cF #rc V^8dQhRRRR/#rrrrrr)rs"rrrs   v  $  rc L\V\\\\\ 34#r)rComputedBuffer InputsKernel InputBufferrTemplateBuffer)rs&ris_realized_nodeIRNode.is_realized_nodes&     rcV^8dQhRR/#rrrr)rs"rrrs&&6&rc ,\PV4#rrcreateselfs&rwrap_for_loweringIRNode.wrap_for_lowerings%%rc$V^8dQhRRRRRR/#)rattrrrJrrrr)rs"rrrs!..s.3.4.rc 2\PWV4R#r)r __setattr__)rrrJs&&&r_post_init_setattrIRNode._post_init_setattrs 4u-rcV^8dQhRR/#rrrr)rs"rrrs B Bt Brc \VP4pVPRV4TPR\P'd\ P !4MR4VPRR4VPR/4VPRVP4VPRVP4R#)rrNrrrr) rDrrrMdebug_ir_tracebackr format_stackrr)rrs& r __post_init__IRNode.__post_init__sT223  73  V5N5N5N//1TX   t4  r2  d.F.FG  4+@+@ArcV^8dQhRR/#rrOrderedSet[str]r)rs"rrr@@@rc B\RVP444#)c38"TFqPxK R#5irrrdeps& rr(IRNode.get_read_names..?.>s((.>rD get_readsrs&rget_read_namesIRNode.get_read_names?dnn.>???rcV^8dQhRR/#)rrrr)rs"rrrs/rc VP#r)rrs&r get_tracebackIRNode.get_tracebacks ~~rcV^8dQhRR/#rrrr)rs"rrr  !5 rc VP#rrrs&rget_origin_nodeIRNode.get_origin_nodercV^8dQhRR/#rrzOperation | Noner)rs"rrr!1rc R#rrrs&rget_defining_opIRNode.get_defining_oprcV^8dQhRR/#rrlist[Subgraph]r)rs"rrrs~rc .#)z'Return subgraphs contained in this noderrs&r get_subgraphsIRNode.get_subgraphs rcV^8dQhRR/#rr)rs"rrrs/rc \4pVPp\V\4'd/VP 4pVP 'd \V.4pVEFp\ VR4'd0VP'dVPVP4KE\PPPPR/4PVP.4p\V\4'gKVFRp\PPP PVR4pV'gKAVPV4KT EK V#) stack_trace postToPreN)rDrr ExternKernelrrhasattrr)addrD _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrrrpre_grad_nodes node_namer)s& rget_stack_tracesIRNode.get_stack_tracess)3 ,, dL ) )..0K$k]3Dt]++0@0@0@  !1!12OO))JJNN#c$))R( ".$77!/I--QQUU%t #{$((5"0,rTc V^8dQhRRRR/#)rshortenrr Sequence[str]r)rs"rrrs + +4 += +rc TR\VRR4 2pV'd\V4^@8d VR, R2pVP4'gV.#.pVP4F<pVPR4W4P R4, pVPR4K> V.V,#)zorigins=r:N=Nz...zstack_traces = { })rrr6appendsplit)rr9rstack_trace_strr)s&& r common_reprIRNode.common_reprsWT9b9:; s7|b( c*G$$&&9 002K  " "#5 6 006 6O  " "3 '3y?**rc(V^8dQhRRRRRRRR/#)rlineszSequence[object]r9r multilinerrr)rs"rrrs, 5 5% 504 5HL 5  5rc 8\V4\VPV44,p\\\V44pV'd6\ RP V44p\ V4P RV R2#\ V4P RV R2#)z, z( z )(rR)rrCmaprindentjoinrr)rrFr9rG new_liness&&&& r str_helperIRNode.str_helpersU d4#3#3G#<==Se_% uzz%01I4j))*#i[< <4j))*!E7!4 4rcV^8dQhRR/#rr torch.dtyper)rs"rrrs;rc VP#rr7rs&rrAIRNode.get_dtype zzrcV^8dQhRR/#)rrztorch.dtype | Noner)rs"rrrs!3rc HVP4# \dR#i;ir)rANotImplementedErrorrs&rmaybe_get_dtypeIRNode.maybe_get_dtypes& >># #"    !!cV^8dQhRR/#rrLayoutr)rs"rrrsWWFWrc 2\R\V4 R24h)z#get_layout() is not implemented by !rYrrs&rr>IRNode.get_layouts!$GT |ST"UVVrcV^8dQhRR/#)rrz Layout | Noner)rs"rrrs-rc HVP4# \dR#i;ir)r>rYrs&rmaybe_get_layoutIRNode.maybe_get_layout& ??$ $"  r\cV^8dQhRR/#rrrSr)rs"rrr s!!!rc "VP4#r)r>rs&rget_output_specIRNode.get_output_spec s  rcV^8dQhRR/#)rrzOutputSpec | Noner)rs"rrr s'8rc HVP4# \dR#i;ir)rlrYrs&rmaybe_get_output_specIRNode.maybe_get_output_spec s( '') )"  r\cV^8dQhRR/#rrrr)rs"rrrs@@4@rc >\VP4\4#)z4True for single tensor output (excludes MultiOutput))rrpr_rs&rhas_tensor_outputIRNode.has_tensor_outputs$446??rcV^8dQhRR/#rrryr)rs"rrrsUU,Urc 2\R\V4 R24h)z!get_size() is not implemented by rarbrs&rr<IRNode.get_sizes!$Ed4j\QR"STTrcV^8dQhRR/#rrSequence[_IntLike] | Noner)rs"rrrs 9rc HVP4# \dR#i;ir)r<rYrs&rmaybe_get_sizeIRNode.maybe_get_size% ==? ""  r\cV^8dQhRR/#)rrz)_IntLike | sympy.Rel | Sequence[_IntLike]r)rs"rrr s@rc "VP4#rr<rs&rr{ IRNode.shapes}}rcV^8dQhRR/#rrr"r)rs"rrr#s..4.rc 4\VP44#r)rur<rs&r get_numelIRNode.get_numel#sT]]_--rcV^8dQhRR/#rsr)rs"rrr&UU$Urc \PPP\P !VP 4^44#rr{r9r:statically_known_truerrjrrs&rrIRNode.is_zero_elements&0ww55ehht~~?OQR6STTrcV^8dQhRR/#rrrOr)rs"rrr)sBBBrc 0\R\V4 24h)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 rbrs&rrealizeIRNode.realize)s "ODJ<"@AArc V^8dQhRRRR/#rwriterzIndentedBuffer | Nonerrr)rs"rrr;sLL(=LLrc 0\R\V4 24h)zcodegen_reference NYI on rbrrs&&rcodegen_referenceIRNode.codegen_reference;s!$=d4j\"JKKrcV^8dQhRR/#rrtorch.device | Noner)rs"rrr>s/rc R#rrrs&rrBIRNode.get_device>rrcV^8dQhRR/#rr torch.devicer)rs"rrrAs\rc DVP4pVf \R4hV#)NExpected device is not None)rBrrr8s& rget_device_or_errorIRNode.get_device_or_errorAs$" > !>? ? rcV^8dQhRR/#rsr)rs"rrrGsrc R#Frrs&rhas_exceeded_max_readsIRNode.has_exceeded_max_readsGrcV^8dQhRR/#rr$Callable[[Sequence[Expr]], OpsValue]r)rs"rrrJs77A7rc >\\V4P4hrrYrrrs&r make_loaderIRNode.make_loaderJ!$t*"5"566rcV^8dQhRR/#rr Callable[[Sequence[Expr]], Expr]r)rs"rrrMs77>7rc >\\V4P4hrrrs&r make_indexerIRNode.make_indexerMrrcV^8dQhRR/#rxr)rs"rrrPs77.7rc >\\V4P4hrrrs&rrnIRNode.get_stridePrrcV^8dQhRR/#r|r)rs"rrrSs";rc HVP4# \dR#i;ir)rnrYrs&rrmIRNode.maybe_get_strideSrhr\cV^8dQhRR/#rrrr)rs"rrrYs77#7rc >\\V4P4hrrrs&rrIRNode.get_nameYrrcV^8dQhRR/#rr)rs"rrr\s rc HVP4# \dR#i;ir)rrYrs&rmaybe_get_nameIRNode.maybe_get_name\rr\cV^8dQhRR/#rsr)rs"rrrbsrc ~VP4\PP9# \dR#i;ir)rr{r9 graph_inputsrYrs&ris_input_bufferIRNode.is_input_bufferbs4 ==?agg&:&:: :"  s *- <<c V^8dQhRRRR/#r thresholdrrrr)rs"rrrhsJ$rc R#rrrrs&&rhas_large_inner_fnIRNode.has_large_inner_fnhrr graph_reusec$V^8dQhRRRRRR/#rusersrrrrrr)rs"rrrks!   T T rc R#)a< Hint that this node's value will be reused. `users` estimates the reuse count. When `graph_reuse` is true, this represents graph fanout and can trigger fanout-based realization heuristics. When false, the reuse comes from loop-level indexing, such as a broadcast expand, and should not be treated as graph fanout. Nrrrrs&&$r mark_reuseIRNode.mark_reuseksrcV^8dQhRR/#rr)rs"rrrus  d rc R#rrrs&r realize_hintIRNode.realize_hintu rcV^8dQhRR/#rr)rs"rrrxs77V7rc >\\V4P4hrrrs&rrIRNode.unwrap_viewxrrcV^8dQhRR/#rr)rs"rrr{s77t7rc >\\V4P4hrrrs&r freeze_layoutIRNode.freeze_layout{rrc$V^8dQhRRRRRR/#rrr allow_paddingrrrr)rs"rrr~s$77"7377 7rc >\\V4P4hrrrrrs&&&rfreeze_layout_with_stride_order&IRNode.freeze_layout_with_stride_order~"$t*"5"566rc V^8dQhRRRR/#rrrrrr)rs"rrrs77=7T7rc >\\V4P4hrrrrs&&rfreeze_layout_with_fill_order$IRNode.freeze_layout_with_fill_orderrrc V^8dQhRRRR/#rr6ryrrr)rs"rrrs774F747rc >\\V4P4hrrrr6s&&rfreeze_layout_with_same_order$IRNode.freeze_layout_with_same_orderrrc$V^8dQhRRRRRR/#r exact_stridesryrrrrr)rs"rrrs$77/7@D7 7rc >\\V4P4hrrrrrs&&&r freeze_layout_with_exact_strides'IRNode.freeze_layout_with_exact_stridesrrcV^8dQhRR/#rrdependencies.ReadWritesr)rs"rrrs77!87rc >\\V4P4hrrrs&rget_read_writesIRNode.get_read_writesrrcV^8dQhRR/#rrzOrderedSet[Dep]r)rs"rrr,,?,rc 6VP4P#rr readsrs&rr IRNode.get_reads##%+++rcV^8dQhRR/#rrrr)rs"rrr%%3%rc 4\VP44#r)rr rs&r num_readsIRNode.num_readss4>>#$$rcV^8dQhRR/#rrrr)rs"rrrs7787rc >\\V4P4hrrrs&rget_storage_numelIRNode.get_storage_numelrrc V^8dQhRRRR/#rrrrrr)rs"rrr77!7 !7rc >\\V4P4hrrrrs&&rget_free_symbol_usesIRNode.get_free_symbol_usesrrcV^8dQhRR/#rr)rs"rrrs77J7rc >\\V4P4hrrrs&rget_reduction_typeIRNode.get_reduction_typerrcV^8dQhRR/#rrSequence[Expr]r)rs"rrrs77N7rc >\\V4P4hrrrs&rget_reduction_sizeIRNode.get_reduction_sizerrcV^8dQhRR/#rsr)rs"rrr4rc R#rrrs&r is_externIRNode.is_externrrcV^8dQhRR/#rsr)rs"rrr$rc R#rrrs&ris_no_opIRNode.is_no_oprrc V^8dQhRRRR/#rr8rrrr)rs"rrrs777&7rc >\\V4P4hrrrs&&rconstant_to_deviceIRNode.constant_to_devicerrcV^8dQhRR/#rrr:r)rs"rrrs77M7rc >\\V4P4hrrrs&rget_mutation_namesIRNode.get_mutation_namesrrcV^8dQhRR/#rr)rs"rrrs77C7rc >\\V4P4hrrrs&rget_operation_nameIRNode.get_operation_namerrcV^8dQhRR/#rCr)rs"rrrs77m7rc >\\V4P4hrrrs&rget_inputs_that_alias_output#IRNode.get_inputs_that_alias_outputrrcV^8dQhRR/#rQr)rs"rrrs ++;+rc R#rrrs&rr7 IRNode.dtypes(+rrT)TTrr)Vrrrr__doc__rDrrrr dataclassesfieldrrrrrr staticmethod contextlibcontextmanagerrrrrrrrr rrrr$r6rCrNrArZr>rfrlrprur<rpropertyr{rrrrrBrrrrrnrmrrrrrrrrrrrrr r rr r'r+r1r6r;r@rErIrMrr7rrrrrrOs3=,/>04-49=6= +00e.rrrrs&rr Operation.get_read_namesr rcV^8dQhRR/#rr)rs"rrr]rrc 6VP4P#rrrs&rr Operation.get_readsrrcV^8dQhRR/#rr list[Buffer]r)rs"rrr]s""\"rc \hrrers&r get_outputsOperation.get_outputsrgrcV^8dQhRR/#rrrr)rs"rrr]*Brc \4#rrCrs&rget_unbacked_symbol_defs"Operation.get_unbacked_symbol_defs |rc V^8dQhRRRR/#r#r)rs"rrr]s! !rc \4#)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. rCr&s&&rr'Operation.get_free_symbol_usess ,|rcV^8dQhRR/#rr)rs"rrr]sCrc ^#)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. rrs&rget_workspace_sizeOperation.get_workspace_sizes r)r`r_Nr)rrrrrrBrrmrrrvrIr}rrr6r;r rr r rrr'rrrrrr[r[sl2"    # $'"-@,"0rr[c]tRtRt$RtR]R&R]R&R]R&R ]R &]!R4R9R R ll4tR RltRRlt ] t RRlt RRlt RRlt RRlt]RRl4t]]P&3RRll4t]RRl4tRR lt]R!R"l4tR:R$R%lltR:R&R'lltR9R(R)lltR*R+ltR,R-ltR.R/ltR0R1ltR2R3lt R4R5lt!R6R7lt"R8t#R##);Loopsi"z:Base class for pointwise and reduction loop-body IR nodes.rr8rRr7Callable[..., Any]inner_fnryrangesc V^8dQhRRRR/#r#r)rs"rrLoops.__annotate__,s  ! ! rc a\4P!.V3RlVP4OVPS4N5!#)c3<<"TFp\VS4xK R#5irr*rers& rr-Loops.get_free_symbol_uses..0s F+Qq-00+)rDunionrinner_fn_free_symbolsr&s&frr'Loops.get_free_symbol_uses+s>|!! F$++ F  & &} 5  rc V^8dQhRRRR/#)rnamesr:rrr)rs"rrr4s   ]  s  rc  TPRVPP R2\VP4VP 4.VUu.uFq" R\ W4 2NK up,RVP: 2.,4#uupi)'= origin_node=)rNr8rrr7 inner_fn_strrr)rrrs&& r_to_str Loops._to_str4sDKK$$%Q'DJJ!!#  \PpM-\ V\ 4'g\R\V4 24h\W4#)Nz,expected int realize_opcount_threshold, got ) rMrealize_opcount_thresholdrdrealize_cpu_opcount_threshold"_realize_opcount_threshold_defaultrrrrmax)rrr s&& rget_realize_opcount_threshold#Loops.get_realize_opcount_thresholdvsy  I$*$D$D! $ ,d||,2,P,P),2,U,U)7==$B4HaCbBcd988rc V^8dQhRRRR/#rr)rs"rrrs  J $ rc ZVP4PVPV48#r)rnum_opsrrs&&rrLoops.has_large_inner_fns.$$&..1S1S 2   rc V^8dQhRRRR/#rrrrOrderedSet[Symbol]r)rs"rrrsWW4WDVWrc fVPVP4p\VPW!R7#r)rrrUr)rrrs&& rrLoops.inner_fn_free_symbolss% DKK(#DMM5VVrc V^8dQhRRRR/#)rsymbolr$rrr)rs"rrrsFrc VPVP4p\V4p\P!V4;_uu_4VP V4RRR4VP # +'giTP #;ir)rrrXr{rrusages)rrrhandlers&& rcollect_inner_fn_symbol_usage#Loops.collect_inner_fn_symbol_usagesa DKK(08  g & & MM%  '~~ ' &~~s A++ B cV^8dQhRR/#rr)rs"rrrs  ? rc \P!\RR4;_uu_4VP4'dL\ VP 4VP 4VP44PuuRRR4#\ VP 4VP 44PuuRRR4# +'giR#;ir) r!rr?r+rWrr<r1rrs&rr Loops.get_readss \\.*:D A A&&((*$$&MMO++-% B A+$$&MMO%B A A AsAC1C C cV^8dQhRR/#rr)rs"rrrrrc H\VP4P4#r)rDr read_buffersrs&rr Loops.get_read_namess$//1>>??rcV^8dQhRR/#rr)rs"rrrs9939rc H\VP4P4#r)rrr)rs&rrLoops.num_readss4((*7788rcV^8dQhRR/#r.r)rs"rrrs  N rc 2\R\V4 R24h)z+get_reduction_size() is not implemented by rarbrs&rr1Loops.get_reduction_size!9$t*Q G  rcV^8dQhRR/#rr)rs"rrrs  J rc 2\R\V4 R24h)z+get_reduction_type() is not implemented by rarbrs&rr+Loops.get_reduction_typer1rc V^8dQhRRRR/#r>r)rs"rrr   & rc 2\R\V4 R24h)z+constant_to_device() is not implemented by rarbrs&&rr@Loops.constant_to_devicer1rrrr)$rrrrrSrrer'rr__repr__rBrr<r classmethodrrVrLINDEXrrdrrrrrrr"r r rr1r+r@rrrrrr"sD    G$ %   )H  # #:>**  ((+  9  W  @9    rrc$V^8dQhRRRRRR/#)rrzExpr | Sequence[Expr]r7rRrrzr)rs"rrrs"&&,& &&rcVP'd!\P!\R4V4#\P!^V4#)nan)is_floating_pointryconstantfloat)rr7s&$r nop_loader_fnrBs4 ||E%L%00||Au%%rcb]tRtRtRRltRRlt]tRRltRR ltR R lt R R lt Rt R#) PointwiseicV^8dQhRR/#rr)rs"rrPointwise.__annotate__sArc |VP4'd\\VPR7#VP#)rT)rr rBr7rrs&rrPointwise.make_loaders,  " "= ; ;}}rcV^8dQhRR/#rr)rs"rrrFrrc $VPR4#rrrs&rrPointwise.__str__rrcV^8dQhRR/#rxr)rs"rrrFs$6rc .#rrrs&rr1Pointwise.get_reduction_size rcV^8dQhRR/#rr)rs"rrrFJrc R#rrrs&rr+Pointwise.get_reduction_typerrc(V^8dQhRRRRRRRR/#) r output_namerOindexer!Callable[[Sequence[Expr]], Never]varsr/rrr)rs"rrrFs8PPP3P P  Prc ~VP4p\P!T;'gRV!V4V!V44#unnamed)rrystorerrUrVrXloaders&&&& r store_outputPointwise.store_outputs5 !!#yy11 74=&,OOrc V^8dQhRRRR/#r>r)rs"rrrF     &  rc VP4p\P!\RV4!V4p\ VVP VVP R7#FMove this to a given device. Requires that all reads are to constants.override_devicer8r7rr)rr!rConstantBufferrDr7rrr8r^s&& rr@Pointwise.constant_to_devicesI!!#n.?HP**;;   rrN) rrrrrrr9r1r+r_r@rrrrrDrDs-)HP   rrDcH]tRtRt$R]R&RtR]R&RRltR R ltR tR#) Scatteriroutput_indexerNr^ scatter_modec V^8dQhRRRR/#r>r)rs"rrScatter.__annotate__s     &  rc  VP4p\P!\RV4!V4p\ VVP VVP VPVPR7#)rerf)r8r7rrrmrn) rr!rrhrlr7rrmrnris&& rr@Scatter.constant_to_devices[!!#n.?HP**;;..**   rc(V^8dQhRRRRRRRR/#) rrUrOrVrWrXr/rrr)rs"rrrps2   3    rc VP4pVfRp\P!VV!VPV44V!V4VPR7#)Nr[)mode)rryr\rmrnr]s&&&& rr_Scatter.store_outputsT !!#  #Kyy  D''- . 4L""   rr) rrrrrrnr@r_rrrrrlrls!44"L)"    rrlr logical_orrmaximumfmaxfmaximumminminimumprodmulsumr-dotxor_sum bitwise_xorz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNc(V^8dQhRRRRRRRR/#) rreduction_typerr7rRarg_break_ties_leftrrCallable[..., object]r)rs"rrrs8XNXNXN +XNBFXNXNrcaaaaS\9dI\S,o\P'd$SR 9d\S4'd RV3RllpV#S#SR 9dRVVV3RllpV#SR8Xd RRlpV#\ RS 24h) rc$V^8dQhRRRRRR/#)rarbrrzr)rs"rr.get_reduction_combine_fn..__annotate__(s! 9 9 9F 9x 9rcj<S!W4p\P!\P!W4W4#r)rywhereeq)rrrJ combine_fns&& rstrict_signed_zero_combine_fn?get_reduction_combine_fn..strict_signed_zero_combine_fn(s&"1(yyq88rc$V^8dQhRRRRRR/#)rrztuple[object, object]rrtuple[OpsValue, OpsValue]r)rs"rrr9s$  $ )>  & rc<Vwr#VwrES R9d\P!W$4pM\P!W$4p\P!W$4p\ S 4'd\P !W"4p\P !WD4p \P !V\P!W44p\P !V\P!W44pS 'd\P!W54M\P!W54p \P !V\P!Wz44p\P!WbV4\P!WcV43#)argmin)r argmin_valueargmin_with_value) ryltgtrr4nerw logical_andr)rra_valuea_indexb_valueb_indexmaskequala_isnanb_isnantierr7rs&& rargmax_combine_fn3get_reduction_combine_fn..argmax_combine_fn9s ! G G!PPvvg/vvg/FF7,Ee$$&&2&&2~~dCFF7,DEucoog.OP'w(VVG-  >>$(CDD $1 $1 rwelford_combinec$V^8dQhRRRRRR/#)rrz#tuple[OpsValue, OpsValue, OpsValue]rrr)rs"rrrZs&  2 2 1 rc Vwr#pVwrVp\P!\P!\P!V4\P!W%44\P !R\ P4WR, 4pWG,p Wy, p W(V ,,W6,W,V,V ,,V 3#)g)ryrrisinfrr@rDfloat32) rra_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_ws && rwelford_combine_fn4get_reduction_combine_fn..welford_combine_fnZs&' "F(%& "F( II & 1366&3IJ S%--0E ",J -I** emh6BB rzunknown reduction_type=)rr{argmaxr argmax_valuerargmax_with_valuer)rrMstrict_signed_zeror4rY)rr7rrrrrsfff @rget_reduction_combine_fnrs--).9   % % %.0u%% 9 91 0    :! , , ."!"$;N;K"LMMrc>a]tRtRt$RtR]R&R]R&R]R&R ]R &R tR ]R &RtR]R&RRlt]t ] !R4R=RV3Rlll4t RRlt RRlt RRltRRltRRltR=RR lltR!R"lt]R>R#R$ll4t]R%R&l4t]]P0R3R'R /R(R)lll4t]R*R+l4t]R,R-l4t]R.R/l4t]R0R1l4t]R>R2R3ll4t]R4R5l4t]R=R6R7ll4t ]R>R'R /R8R9lll4t!]R:R;l4t"R.s P:OQq-00:Or)superr'rDrrrr __class__s&frr'Reduction.get_free_symbol_usess7w+M:Z\=O=O P$:O:O P>   rcV^8dQhRR/#rxr)rs"rrrs%%$6%rc VP#r)rrs&rr1Reduction.get_reduction_sizes$$$rcV^8dQhRR/#rr)rs"rrrs##J#rc VP#r)rrs&rr+Reduction.get_reduction_type"""rc ,V^8dQhRRRRRRRRR R /#) rrUrOrVrWrXr/reduction_varsSequence[Symbol]rrr)rs"rrrsC L L L3 L L ) L  Lrc  \P!VPVPVPVP W444p\P !T;'gRV!V4V4R#)r[N)ry reductionr7rrrstore_reduction)rrUrVrXrrJs&&&&& rrReduction.store_reductionsU JJ NN    MM$ /   K449gdmUKrcV^8dQhRR/#rr)rs"rrrs==c=rc `\VP4\VP4,#r)rrrrs&r index_lengthReduction.index_lengths!4;;#d&;&;"<<r)rs"rrrs   & rc  8VP4p\P!\RV4!V4p\ VVP VVP VPVPVP\PVPVPR7 #)rerf r8r7rrrrrrrr)rr!rrhrr7rrrrraDEFAULTrrris&& rr@Reduction.constant_to_devices}!!#n.?HP**;;!22..nn(00 $ 8 8"44  rc@V^8dQhRRRRRRRRRR R R R R R RRRRR/ #)rr8r dst_dtyperRrrzCallable[_P, OpsValue]rryrrzReductionType | Literal['scan']reduction_numelr" input_noder3rtuple[ReductionHint, _IntLike]r)rs"rrrsxTTTTT) T # T - T8TT"T (Trc  V\V4.p \PPP V 4'g\ P ^3#\PPPV4p \V4p \PPPV 4p \PPV\P4;'d9\V4'*;'g!\PPW V4p RpVR8H;'g,V '*;'dWn9;'d\PpVR8Xd\ P ^3#\ P"!V4pVP$p^ pV'dd\&P(!\PP*VRR7p\&P(!\PP*VRR7pMRRlpTpV ^8XdV!W4pV^8Xd\ P,V3#Ve\/V\04'd\2P4!\6R R4;_uu_4\9V4wppRRR4XeqXem\PPP;\VV,44pV V8Xd-\<P?R VVVVV4\ P,R3#\ P,V3#V V8:gV V^,^ ,8d\ P ^3#\ATTTTTVR8wdTMR V\ P R 7pR RlpV!V4wppV'd V!V4wpp\CV4^8Xd\ P ^3#\DPF!VPI4VPK44wwppp^p ^p!VFp"\PPPMV"V4p#\PPPOV#V\QVPS444p$\T;QJdRV$4F 'dK RM RM !RV$44p%V%'d V ^, p KV!^, p!K V!V 8d\ P,V!W43#\ PVV!W43# +'giELn;i)zT Choose the reduction hint and split count from shape and input stride information. scanrT)inner_reductionFc$V^8dQhRRRRRR/#)rreduction_numel_hintr numel_hintrr)rs"rr*Reduction.num_splits..__annotate__s&  &)   rc^#rr)rrs&&rinner_reduction_splits4Reduction.num_splits..inner_reduction_splitssrNrzUse previous IRNode's range and reduction_ranges instead of split. current ranges: %s, current reduction ranges: %s, current split: %d, new ranges: %s, new reduction ranges: %srr8r7rrrrrrc V^8dQhRRRR/#)rrrrztuple[Sequence[Expr], bool]r)rs"rrr@s7 E7 E 7 E.I7 Erc PaVP4pVf \R4h\R\VVP 4VP 4R7VR7pVP 4pVPf \R4hVPUu.uF?p\V\4'gK\V\P4'dK=VNKA pp\P!VP 4VP44wwrgpVUu.uF?p\V\4'gK\V\P4'dK=VNKA pp.p.p Rp \VP RR7EFp V P"P$o\&;QJdV3RlV4F 'dK RM R M!V3RlV44p V '*;'d;\(;QJdV3R lV4F 'gK R M RM!V3R lV44p V 'dVP+V P"4M#V 'dV P+V P"4V 'g V 'gKV P,\.P0P29gEK"\.P0P2V P,,p\5VP6R R4pVP94\5VP6R R4V8wgEKR p EK T;'gT V 3#uupiuupi) Nrr8r7r5rrrz+Expected read_writes.range_vars is not NoneFcVP#rrrs&r@Reduction.num_splits..get_read_indices..csaffrkeyc3,<"TF qS9xK R#5irrrvarr9s& rrAReduction.num_splits..get_read_indices..es'Rz|(;zTc3,<"TF qS9xK R#5irrr s& rrr isN3AC<'>rr6)rBrrr?rAr<r  range_varsrr"rr rNindex_vars_squeezer1sortedrrr9rrr@rr{r9name_to_bufferrr decide_layout)rr8cb read_writesr rr&rindicesbroadcasted_reduction_indiceschangedmdis_full_size_readis_broadcasted_reduction_readbuforiginal_strider9s& @rget_read_indices.Reduction.num_splits..get_read_indices@s\\^F~$%BCC%!++- B,,.K%%-$%RSS'111Cc4(1;C1N1  &2%D%D a224& " Q *)Cc4(1;C1N)  G,. )G[..4DEE!xx44 $'C'Rz'RCCC'Rz'R$R!5F0E11##N3AN###N3ANK-%NN288,2188B$(E(Eww!''"8"88gg44RWW=*1#**h*M))+"3::x>/Q&*G'F(;;;WD DIs$L+L LL#.L#L#c3D"TFq^8H;'gV^8xK R#5irirrs& rr'Reduction.num_splits..s9AQ!a%s  rrl),rur{r9r:all_unbacked_explicitly_hintedrarr; has_featurerOREDUCE_TO_SINGLE_ELEMENTrachoices should_use_cooperative_reductionrMsplit_reductionsr`rmulti_processor_count functoolsr reduction_split_factorINNERrrr!rr?rV!replace_backed_symbols_with_hintslogr/rrrNrr<r1simplify_with_ranges stride_hintsrkeysrOUTER)&r8rrrrrrrrexprsrnumelrshould_reduce_to_single_elementarg_reduction_types should_splitpropsnum_smmin_elements_per_threadrouter_reduction_splitsrA new_rangesnew_reduction_rangesextracted_numel_hintrrrrr&rranges1 num_outer num_innerrjrouters&&&&&&&&&& r num_splitsReduction.num_splitssl"!-"78ww>>uEE ((!+ + ww//AA/Rf%WW%%77> +,''*=*= N;;+ + + &! !  yy99 ( &/  / / ( (9 ( (''  U " ((!+ + ''/,,"$ @I@Q@Q 00&$A "AJ@Q@Q 00&%A "   &< " ?*+?LEz$**E11%*Z*K*K\\.2BDII5(00 7 Er,A. )!,JGQ w<1  ((!+ +'3'F'F JJL!..0( $NW  A  55aAAgg&&33>4 #7G C99CCC999EQ Q  y  &&(>$) !&&(>$) EJIIs *S S/ c ,V^8dQhRRRRRRRRR R /#) rrz.__annotate__s  ( S rc <a\P!SVV3Rl\P!SUu.uFp\ V4NK up!44#uupi)c36<"TFpS!SV4xK R#5irr)rrrvalue_fns& rr=Reduction._unroll_reduction_fn..fn..s&#UF++#)r*reduce itertoolsproductr)rrrrrLsf rr*Reduction._unroll_reduction_fn..fnsN##"+"3"3,<=,sAc$V^8dQhRRRRRR/#)rrryrrrr)rs"rrrIs$  ) 3E * rc<VUu.uFp\P!V4NK ppS!W4p\V\4'dV#V\P !S!V4\ P43#uupir)rexpandrrry index_exprrDint64)rrrrJ flatten_indexrs&& rrL0Reduction._unroll_reduction_fn..value_fnsf4::6a%,,q/6: /eU++ LNN=#8%++F ;sA4c"<S!V4^,#rrrrs&rr0Reduction._unroll_reduction_fn..s RYq\rc"<S!V4^,#rrr[s&rrr\s E1r)rrrrrr)rr)rr)r{r9r: guard_int_seqr_fixed_indexerr?r@)rrrrrrXrrLsff&&@@@@r_unroll_reduction_fnReduction._unroll_reduction_fns77++99:JK-nH      + 112BCM  !KK !AA11- -HIr strict_sumcDV^8dQhRRRRRRRRRR R R R R R RRRRRRR/ #)rr8rrrRrrrrryrrr]rrarr3rbrrrr)rs"rrrsjjjj j % j # j-j&j&j"jj jrc  aaaaa"a#a$a%a&\PPP\ S44o%S%^8XdnRV3Rllp RV !^4RV !^4RV !^4RV !^4/o&SS&9d\ S R24hRVVV&3R llp \ PVVV \V4R 7#R p R pR pV 'EdV\P\P39Ed\ V4p\PPPS%V.4'Ed\PPPS%4p\P!VP 4p\#V^ V,4pR"pVV8*;'d\PPP%\&P(!\&P*!V^4\&P,!VV4\&P,!VV,V44R R 7pV'd^p VV\P.,8*pV'd5\#VV4p^V^, P14,pVV,pM]\P2!V^V4pVP4pVP6\P88dVV,^8Xd VP6p V R Jp V'd{Vf \ R4hS%o"So$VP;SS4o#RV"V#V$3Rllp\=\>R\@3,V4o\&PB!V4.o\&PB!V4o%S%^8Xd=V 'g5SR#9d RV3RllpM RVV3Rllp\ PVSVVR 7#\ES%\B4'd\GS%4\HPJ8df\ V4^8wg\MVPN4'd;SR8wd4V 'g,\ PVSVPQSSSV4VR 7#V fVPSVSVSVSSS%V 4 wppM\PT ppRV%3RllpV f V!V4pV\P8XdTpVR$8XdV f \ R4h\TPV!\XRR4;_uu_4\[V 4wppR R R 4Xf \ R4hXf \ R4hVP]VSVSVSVVSV4 #V^8dVP_VSVSVSSVVV V R7 pR p \HP`Pb'd%\EV\d4'dRR lp!V!!V4p V 'dw\EV Pf\h4'g \ \OV Pf4 4hV PfPj^,V n6SV n7VV n8SV n9V#\dP\iVSSVSSVVVVR!7 4pV# +'giEL\;i)%za Create a reduction node. May split the reduction to multiple layers to expose more parallelism. c V^8dQhRRRR/#)rvalrrzbool | float | intr)rs"rr&Reduction.create..__annotate__s $ $V $(: $rc<<S\P8Xd \V4#SP'd6\V\4'g\ \ V44h\V4#\V\4'g\ \ V44h\V4#r) rDrr?rrrrrArr)rfrs&rpy_cnst!Reduction.create..py_cnstsp *9$000%c=99,T#Y77 :%%c;77,T#Y77s8Orrrr}r* not supported for zero-dimension tensors!c V^8dQhRRRR/#rrrrrzr)rs"rrrgs P P P Prc@<\P!SS,S4#rryr@)rrrrtypes_to_initss&rconst_fn"Reduction.create..const_fns||ON$CYOOrrgNF)fallback_valuez.strict multirow sum requires a reduction blockc$V^8dQhRRRRRR/#rrrreduction_indexrrzr)rs"rrrgKs$ 7 7' 7:J 7 7rc<Vwp\S4p\P!\P!W#4\P!SV44p\P !SW4p\P !WES4#r)rkryrrVr*r masked) rrvr index_dtyperbodyactual_reduction_numeldefaultoriginal_inner_fns && rpadded_inner_fn)Reduction.create..padded_inner_fnKsf, -.DE vvNN67NN#9;G!(():ESzz$g66r.c V^8dQhRRRR/#rmr)rs"rrrg`s66c6h6rc2<\P!^S4#rro)rrs&rrReduction.create..fn`s<<955rc V^8dQhRRRR/#rmr)rs"rrrges<._maybe_increase_splits2/** qy5&"6"677 rzExpected input_node is not NonerTzExpected new_ranges is not Nonez)Expected new_reduction_ranges is not Nonerbc V^8dQhRRRR/#)rcur_noderrzComputedBuffer | Noner)rs"rrrgs'*rctVP4p\V4^8wdR#\\V44pV\P P 9dR#\P P V,p\V\4'gR#VPP4f \R4hV#)rNz2Expected buf.data.get_reduction_type() is not None) r rnextiterr{r9rrrrr+r)r read_namesbufnamers& r_find_split_reduction/Reduction.create.._find_split_reductions"*!8!8!:J:!+#"4 #34Gagg&<&<<#''009C%c>::#xx224<,PJrri)rrrl):r{r9r:simplifyrurrDrrrarr,r# guard_intr0vec_sizeitemsizerf evaluate_exprrroGtrq_K_MULTIROW_MAX_LOADS bit_lengthcompute_inner_tree_paramsbatch_total_elements num_batches_K_TWO_KERNEL_THRESHOLDrrrrr#rrrMunroll_reductions_thresholdrqrr`rDr!rr?rV!create_multilayer_existing_rangescreate_multilayerrZmix_order_reductionrrrr _split_size_original_inner_fn_original_ranges_original_reduction_ranges)'rr8rrrrrrrrrbrirq strict_splitrr num_outputsguarded_reduction_numel vector_sizenum_batches_ub int32_max indexing_safe num_loadsparamsr~rhintrArr<r=r'split_reductionrr{r|r}rrps'&&f&f&ff&&$ @@@@@rrReduction.creates(''**33MBR4ST a  $ $wqz71: wqz O_4$%&&PQ P P##!F| $  #  :.  ! !   -  (/Kww>> +.+,''*:*:*D*D_*U'-66y7I7IJ !()@"{BR!S% +y8((66 !HH[!4!HH[)<!HH[>%A9M (- 7!#$L/&)N)NNO(+$+,C[$Q $%)a-)C)C)E$E ,5 ,C)!0!J!J3'" -3,G,G)"..1X1XX 7:K Kq P+1+=+=L!-  ($%UVV%4 " ( '' BG 7 7HS#X.@H % .? @A #mm,=>O a  !5566 <<##YF$   0 0O$v'I'IIv&!+vfkk/B/B%' ##11.  $   ..  KD%(--|%D    )%0E ]22 2!N B;!$%FGGn.>EE3V40 0F!$%FGG#+$%PQQ88 $  QY'' %( C#O}}000ZY5O5O*#8"<"/"6"6 BB(D1E1E,F+GII.=.B.B.S.STU.V+5=2390=M:J !!1-#-$7"3    yFEEs X.. X? c$V^8dQhRRRRRR/#rrrr7rRr_NumLike | Sequence[_NumLike]r)rs"rrrs$###$/# &#rc VR 9dP\V4'd \R4#\V4'dR#\P!V4P #VR9dP\V4'd \R4#\V4'dR#\P!V4P #\V4'dRM^p\V4'dRM^pRVRVRVRVR VR W"V3R W"V3R \R4V3/V,#)r-infFinfTrr}rrrwelford_reduceronline_softmax_reduce)rryrrr)r{rrr)r4rAr2rDiinfor{r)rr7zeroones&& rdefault_accumulatorReduction.default_accumulators   e$$V}$!%(({{5)--- S Se$$U|#!%(({{5)---(//uQ&u--d1 4 C 4 t 4 t40 D1 #eFmT%:     rc$V^8dQhRRRRRR/#rr)rs"rrr*s)DDD$/D &Drc >VR8Xd^#\PW4#)r)rrrr7s&&rrReduction.default_value)s! - -,,^CCrc(V^8dQhRRRRRRRR/#)rrArrrrrarr)rs"rrr2s,%(:G rc VR8XdV#VR8:d-VR8:d&V\P8Xd\P#VR8:d-VR8:d&V\P8Xd\P#V#)rirl)rar2 OUTER_TINY)rArrs&&&r_multilayer_second_step_hint&Reduction._multilayer_second_step_hint1sg B;! ! C.__annotate__ws$  # 6F   rc6<aaVwpVEvopSV,V,oRVV VV 3RllpS 'dd\S 4p\P!\P!SV4\P!S V44p\P!WSS 4#V!4#)cV^8dQhRR/#rrrzr)rs"rrKReduction._multilayer_wrap_loader..wrapper_fn..__annotate__~s = =( =rc$<S!SS!S.44#rr)rr^ new_indexrsrrzCReduction._multilayer_wrap_loader..wrapper_fn..body~si');<.wrapper_fnws"1 _*/ 'Y ?2_DG = =-o> vvNN7K8NN?K@zz$g66v r) rViewdynamic_reshape_indexerr{r9r:rrrjrJ) rr^rrrArr|r dense_indexrrrs &f&f&ff& @@r_multilayer_wrap_loader!Reduction._multilayer_wrap_loaderbs~>> Z .. / ((>> HHS%0! 4    (rc 0V^8dQhRRRRRRRRRRR R/#) rr^4Callable[[Sequence[Expr], Sequence[Expr]], OpsValue]original_rangesr/original_reduction_rangesr<Sequence[Integer]r=rr)rs"rrrsFD($2  &  0  >rc aaa\;QJdRS4F 'dK RM RM !RS44'g\RS: 24h\PV\ V4\ V4,4oRVVV3RllpV#)c3*"TF q^8HxK R#5i)rNrrrs& rrDReduction._multilayer_wrap_loader_existing_ranges..s3?a6?FTz8Only enabled for numel_hint == 1, found original_ranges=c$V^8dQhRRRRRR/#)r merged_indexr/new_reduction_indexrrzr)rs"rrGReduction._multilayer_wrap_loader_existing_ranges..__annotate__s&  ( !/  rc <VR\S4pV\S4RpS!VS!\V4\V4,44#r)rr)rr original_idxrr^rrs&& rrEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fnsQ((>#o*>?L$S%9%;.__annotate__s) C C% C8J C  Crc<S!.VOVO4#rr)rrvintermediate_loaders&&rintermediate_fn;Reduction.create_multilayer_helper..intermediate_fns''A'A'AB BrNz>Expected original_ranges == new_ranges[: len(original_ranges)]) r8r7rrrrrrr)rDfloat16bfloat16rArrrrr{r9r:r;rurrrr)rr8rrrrrr<r=rrArrbintermediate_dtype intermediaterrrs&&&&&&&&&&&&& @rcreate_multilayer_helper"Reduction.create_multilayer_helpers>2 ??   !''       !(   *668 C C WW%%77 o8VW 99 :~  )?3+?@ @ P  (&!+C,@,B!C-#-'1!   8<   rcHV^8dQhRRRRRRRRRR R R R R R RRRRRRRRR/ #)rr8rrrRrrrrryrrr]rArrrarr3rbrrrr)rs"rrrs. . . .  . % . # . -. &. . &. ". .  . rc  \V4p \W^, ,V4p VPWr4pVPVVV VV VV 4pVP VVVVVV.VOVNV .VVV V 4 #)r)rurGrrr)rr8rrrrrrrArrrbrrr|rs&&&&&&&&&&&$ rrReduction.create_multilayers*((89o;UC ##N>00        ++       f e  L      rcDV^8dQhRRRRRRRRRR R R R R R R RRRRRR/ #)rr8rrrRrrrrryrr<rr=rr]rrarrr)rs"rrr# sx$ $ $ $  $ % $ , $ $6$ "$ ,$ &$ &$  $ rc  jVPVVVVV4p VPVVVV VV.VOVOVV RV 4 #)rrl)rr) rr8rrrrrr<r=rrrs &&&&&&&&&&& rr+Reduction.create_multilayer_existing_ranges" sc$@@   %   ++      % +o + +     rrrr)%rrrrrSrrrrr9rer'r1r+rrrrr@rVrDr`r:rarrrrrrrrrrrr __classcell__rs@rrrwsR((!!!! %%$(z(NHK(  ) %# L=   "TTl77r)6(=(=$(j!jjX ##JDD >((T:C C J. !. . `$ $ rrc(V^8dQhRRRRRRRR/#) rr5ryr6r}rr"rz$Callable[[Sequence[_IntLike]], Expr]r)rs"rrrJ s0  % * rc"aaaRVVV3RllpV#)1A closure containing math to read a given elementc V^8dQhRRRR/#)rrryrr"r)rs"rr$_fixed_indexer..__annotate__Q s  ) d rc<Se\V4\S48Xg \R4h\V4\S48wd \R4hSp\VSS4Fwr#pV^8wgKWV,,pK V#)Nz9Expected stride is not None and len(index) == len(stride)z Expected len(index) == len(size))rrr)rresultrstszrr5r6s& rrV_fixed_indexer..indexerQ s|"s5zS['@ K  u:T " !CD Dufd3KCRQw(*4 rr)r5r6rrVsfff rr_r_J s   Nr INNER_FN_TYcFa]tRtRt$R]R&RV3RlltRRltRtV;t#) MultiOutputReductionid r output_indexc<V^8dQhRRRRRRRRR RR R R RR RRR/ #)rr8rrrR inner_fnsz#INNER_FN_TY | Sequence[INNER_FN_TY]rrrrr]rrrarrr)rs"rr!MultiOutputReduction.__annotate__g sd#)#)#)#)7 #) " #) , #)&#)#)&#)#)rc  <a\S4'dS3o\S4^8Xd S^,p M RV3Rllp \S V` VVV VVVVVR7WnR#)rc$V^8dQhRRRRRR/#)rrr/ reduction_idxrtuple[OpsValue, ...]r)rs"rr3MultiOutputReduction.__init__..__annotate__{ s) I I# I4B I% Ircp<aa\;QJd.VV3RlS4FNK 5#!VV3RlS44#)c34<"TF q!SS4xK R#5irr)rrrr%s& rr@MultiOutputReduction.__init__..loader..~ sHiR]33ir)rr%r!sffrr^-MultiOutputReduction.__init__..loader{ s'uHiHuHuHiHHHrrN)callablerr__init__r) rr8rr!rrrrrrr^rs &&&f&&&&&& rr/MultiOutputReduction.__init__g sn I  " I y>Q q\F I I -))  )rc ,V^8dQhRRRRRRRRR R /#) rrUrOrVrWrXr/rrrrr)rs"rrr" sCSSS3S S ) S Src  b\P!VPVPVPVP W444p\ V\\34'g\\V44hWPP,p\P!T;'gRV!V4V4#rZ) ryrr7rrrrrrrrrr)rrUrVrXrrrJs&&&&& rr$MultiOutputReduction.store_reduction s JJ NN    MM$ /   &5$-00 f. .(()"";#;#;)WT]ERRrr) rrrrrr/rrrrs@rrrd s #)#)JSSrrcL]tRtRtRt]]PR3RRll4tRt R#) ArgReductioni zJMulti-output arg reduction that returns both the selected value and index.Nc@V^8dQhRRRRRRRRRR R R R R R RRRRR/ #)rr8rrrRrrrrr/rrr]rrarr3rSequence[TensorBox]r)rs"rrArgReduction.__annotate__ szjjjj j % j  j)j&j&j"j jrc  aaaaaaaaaaSR 9d\RS: 24h\PPP \ S44p V ^8Xd\S R24hV ^8XdcRVV3Rllp RRlp \ PSVV \S4R7\ PS\PV \S4R73#\V \4'd\V 4\P8d\ S4^8wg\!SP"4'd\$P'SSSS4oRV3R llo\(;QJd2.VVV3R l\+V\P344FNK 5#!VVV3R l\+V\P3444#\$P-SVSSSSSV V 4 wrV^8wd \R 4hS\.P08XdV o\(;QJd7.VVVVVVVV3R l\+V\P344FNK 5M0!VVVVVVVV3R l\+V\P3444pVFpVP34K V#)rzunexpected reduction_type rkc V^8dQhRRRR/#rrr/rrzr)rs"rr)ArgReduction.create..__annotate__ s   8 rc<SUu.uFp\PPNK ppS!W4p\V\4'd V^,#V#uupir)rrrrr)rr&rvrJrrs& rrL%ArgReduction.create..value_fn sL9I"J9IA577<<9I"J 8eU++ 8O #Ks#Ac V^8dQhRRRR/#r<r)rs"rrr= s 4 4 48 4rcL\P!^\P4#r)ryr@rDrWrs&rindex_fn%ArgReduction.create..index_fn s||Au{{33rrgc$V^8dQhRRRRRR/#)rrr/rrrrzr)rs"rrr= s! , ,~ ,S ,X ,rc~<S!V4p\V\4'g\R\V4 24hW!,#)zexpected tuple result, got )rrrr)rrr unrolled_fns&& rproject$ArgReduction.create..project s:$U+!&%00(+FtF|n)UVV++rc 3<"TF4wr\PSV\SVR7\S4R7xK6 R#5i)r4rgN)rDrr r)rrr7r8rHrs& rr&ArgReduction.create.. sM,O'L   !$W<H< ! ,Os. sO 0    $$+"   0s36)rrr) rr8rrrrrrSrrrMrHs &ffffff&f& r_create_no_split'OnlineSoftmaxReduction._create_no_split sa%  $J/ %%  $J/   A IIKrc@V^8dQhRRRRRRRRRR R R R R R RRRRR/ #rRr)rs"rrrT5 sn1 1 1 1  1 % 1  1 )1 1 &1 "1  1 rc  H\PPP\ V44p \ P VVVVVVRV V R7 wrV\P8XdT pV ^8dVPVVVVVVVV VV 4 #VPVVVVVVVVV 4 #)r)rrr) r{r9r:rrurrDrarrrY) rr8rrrrrrSrrrrrAs &&&&&&&&&& rrOnlineSoftmaxReduction.create4 s''**33MBR4ST**      2+!+   ]22 2!N 19((   ##           rcDV^8dQhRRRRRRRRRR R R R R R RRRRRRR/ #)rr8rrrRrrrrr/rrSrrArrrarr3rr8r)rs"rrrTi sZ Z Z Z  Z % Z  Z )Z Z Z &Z "Z  Z rc  daaaaaaa\V4o\SV^, ,V4oVPSV 4p \P VS.V 4o\ P PP\P!\SV4^44'*oRVVVVVV3Rllp V\P\P39dTM\Pp VP!VV SV .VOVNS.VV 4pVFpVP#4K \$;QJd.RV4FNK 5M !RV44oRV3Rllp\ P PP'\V44pVP)VVV 4p VP!VVV VVV.VV 4#)rc$V^8dQhRRRRRR/#)rrrrvrrr)rs"rr>OnlineSoftmaxReduction.create_multilayer..__annotate__ s$ # ## #6F # & #rc <aaVwpVEvopS V,V,oRVV VV 3Rllp\P!^S4pS 'd\S 4p\P!\P!SV4\P!S V44p\P !Wc\ R44\P!Wd\P!^S443#V!4V3#)cV^8dQhRR/#rr)rs"rrROnlineSoftmaxReduction.create_multilayer..wrapper_fn..__annotate__ s ? ?( ?rc$<S!SS!S.44#rr)rrrrsrrzJOnlineSoftmaxReduction.create_multilayer..wrapper_fn..body s 7G9+=>>rr)ryr@rkrrVrxrAr)rrvrrzrryrrrrrrrrrs&& @@rr.wrapper_fn s"1 _*/ 'Y ?2_DG ? ?,,q),C-o> vvNN7K8NN?K@ JJt5=9IIda)CD vs{"rc3@"TFqP4xK R#5ir)rrrs& rr;OnlineSoftmaxReduction.create_multilayer.. s$Lm]]__mc$V^8dQhRRRRRR/#)rrryrvrr&r)rs"rrra s$  % 8J  ! rcp<aa\;QJd.VV3RlS4FNK 5#!VV3RlS44#)c3:<"TFq!.SOSO4xK R#5irr)rr^rrvs& rrTOnlineSoftmaxReduction.create_multilayer..intermediate_fn.. s(AUv11122AUsr,)rrvintermediate_loaderssffrrAOnlineSoftmaxReduction.create_multilayer..intermediate_fn s75AU5 5AU r)rurGrrrr{r9r:rrrjrJrDrrrArYrrr;r)rr8rrrrrrSrArrrrr intermediatesrrrrrprrrs&&&ff&&&&&& @@@@@rr(OnlineSoftmaxReduction.create_multilayerh s((89o;UC >> Z .. / ((>> HHS%0! 4   # #4 ??   ,,     f e  L    A IIK %u$Lm$Luu$Lm$LL  WW%%77 f8MN 99 :~ ##      G    rrr) rrrrrSr:rarrYrrrrrrrPrP s_K)6(=(=$(@)6(=(=$(1 1 fZ Z rrPcr]tRtRt]]P 3RRll4t]RRl4t ]RRl4t Rt R #) WelfordReductioni c8V^8dQhRRRRRRRRR RR R R R RR/#)rr8rr7rRr!Sequence[Callable[..., Any]]rrrrr]rrarr8r)rs"rrWelfordReduction.__annotate__ sd}}}}0 }  } ( }&}&} }rc aaaaaVR 9d \R4h\PP'd \R4h\P P P\S44pRVVV3Rllp V^8XdV !^4p V !^4p V !^4p WV 3#V^8XddRVVVV3RlloVR8XdS!V^,4V !^4V !^43#\;QJd.V3RlV4FNK 5#!V3RlV44#\PSSSV^,SSVVR7wrV\P8XdT pV^8dVPSSVSSVVV4#\^4Uu.uF)p\ P#\%SSVSSVSVV4 4NK+ ppVFpVP'4K V#uupi) rz@Expected reduction_type in ("welford_reduce", "welford_combine")zGwelford reduction usage is explicitly disabled, please check you configc V^8dQhRRRR/#)rrfrrrr)rs"rr-WelfordReduction.create..__annotate__ s  s y rc^<aRVV3Rllp\PSSV\S4R7#)c V^8dQhRRRR/#rrr/rrzr)rs"rr.const..__annotate__ s  n  rc2<\P!SS4#rro)rr7rfs&rr8WelfordReduction.create..const..inner_fn s||rrgrDrr)rfrr8r7rsf rconst&WelfordReduction.create..const s7   ##!F| $ rc V^8dQhRRRR/#)rr^rrrr)rs"rrr{ s  L  rc^<aRVV3Rllp\PSSV\S4R7#)c V^8dQhRRRR/#r~r)rs"rr;WelfordReduction.create..copy..__annotate__ s88.8X8rcp<SUu.uFp\PPNK ppS!W4#uupirr)rr&rvr^rs& rr7WelfordReduction.create..copy..inner_fn s2=M&N=Muww||=MO&N!#77'Orrgr)r^rr8r7rrsf rcopy%WelfordReduction.create..copy s788!''!%< (rc34<"TF pS!V4xK R#5irr)rrrs& rr*WelfordReduction.create.. s: "T"XX r+)rr)rr)rrMmtiadisable_welford_reductionr{r9r:rrurrrDrarrrrrrur)rr8r7r!rrrrrrmeanm2weightrrArWrMrHrs&ff&ff&& @rrWelfordReduction.create s !F F R  ;; 0 0 0 Y ''**33MBR4ST   a 8DqB1XFV# # a   !11IaL)58U1X==u: :u:u: :::& **    aL  )++   ]22 2!N 19((   2$Ah '     $""  '  A IIK% s6/Gc$V^8dQhRRRRRR/#rr)rs"rrrxH s$$/ &rc R#)r)rrrrrs&&rrWelfordReduction.default_valueG s rc<V^8dQhRRRRRRRRR RR R R R RRRR/ #)rr8rr7rRr!rwrrrrr]rArrrarr8r)rs"rrrxN so\ \ \ \ 0 \  \ ( \ &\ \ &\  \ rc  aaaaaaa\S4o\PPP \ P !\SS4^44'*p V 'dJVR8wdCRV3Rllp SPVSV^,\V ^R7\V ^R73VSRSVR7#\SS^, ,S4o\PTS\;QJd.VVVVV3RlV4FNK 5M!VVVVV3RlV44.VOSNS.VV4p V Fp V P4K RRlo\PPP\V44p SP!SW4p\PTS\;QJd.V3R lV 4FNK 5M!V3R lV 44VS.RV4#) rrc(V^8dQhRRRRRRRR/#)rrr/r%rJrrrzr)rs"rr8WelfordReduction.create_multilayer..__annotate__f s, 2 2# 24B 2KN 2 2rc2<\P!VS4#rro)rr%rJr7s&&&rr@4WelfordReduction.create_multilayer..constantf s||E511rrL)r8r7r!rrrrArc 3R<"TFpSPVSSSS^R7xK R#5i)rr|N)r)rr^rrrrrAs& rr5WelfordReduction.create_multilayer..~ sC (F++$# ,(s$'c(V^8dQhRRRRRRRR/#)rrr/rvr^rrrzr)rs"rrr s0 6 6! 6+ 69 6  6rcV!.VOVO4#rr)rrvr^s&&&rintermediate_loader_fnBWelfordReduction.create_multilayer..intermediate_loader_fn s 4E4O45 5rc3Z<"TF p\SVP4R7xK" R#5i))r^N)r r)rrrs& rrr s)&A.q}}GG&s(+)rur{r9r:rrrjrJrr rGrurrrr;r)rr8r7r!rrrrArrr@rrrrrrrsf&f&&f&f& @@@rr"WelfordReduction.create_multilayerM s ((89((>> HHS%0! 4   +<< 2 2 ((aLHA.HA. !10-)  o;UC (//   E  ( EE  (   f e  L  # &A IIK 6WW%%77 f8MN 99 :  &&   E&EE&   G    rrN) rrrrr:rarrrVrrrrrrruru sL)6(=(=}}~ \ \ rruc^a]tRtRt$R]R&R]R&R]R&R]R&R ]R &R ]R &R ]R&R]R&]!R4R,RV3Rlll4tRV3RlltRRltRRlt RRlt RRlt RRlt RR lt R!R"ltR,R#R$llt]]P$3R%R&/R'R(lll4t]R)R*l4tR+tV;t#)-Scani r scan_rangesr5=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rrarrrtuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]r!c V^8dQhRRRR/#rr)rs"rrScan.__annotate__ s   $  CU  rc <a\SV`S4\4P!V3RlVP4!,\4P!V3RlVP 4!,#)c3<<"TFp\VS4xK R#5irr*rs& rr,Scan.get_free_symbol_uses.. O>N"1m44>Nrc3<<"TFp\VS4xK R#5irr*rs& rrr Hi"1m44ir)rr'rDrrr5rs&frr'Scan.get_free_symbol_uses s] G ( 7l  Od>N>NO l  HdiiH   rcV^8dQhRR/#rr)rs"rrr   t rc <\VP4\VP4,\VP48wd \ R4h\ SV`4R#)zCExpected len(self.ranges) + len(self.scan_ranges) == len(self.size)N)rrrr5rrrrrs&rrScan.__post_init__ I t{{ c$"2"23 3s499~ E U  rc ,V^8dQhRRRRRRRRR R /#) rrUrOrVz%Callable[[Sequence[_IntLike]], Never]rXr/ scan_varsrrrr)rs"rrr s<     7    $   rc zaVPW44o\;QJd!.V3RlVP4FNK 5M!V3RlVP44p\P!VP VP V4p\P!T;'gRV!S4W`P,4#)c32<"TF q!S4xK R#5irrrrrs& rr'Scan.store_reduction.. D^x}}^r[) rrr!ryrrrr\r)rrUrVrXrrrrs&&&&& @rrScan.store_reduction sll4+DT^^DDT^^DD$++t?yy  $ $9gclF;L;L4M  rcV^8dQhRR/#rr)rs"rrr sJrc R#)customrrs&rr+Scan.get_reduction_type srcV^8dQhRR/#r.r)rs"rrr s  N rc VP#r)rrs&rr1Scan.get_reduction_size rrcV^8dQhRR/#r.r)rs"rrr .rc VP#rr5rs&rr< Scan.get_size yyrcV^8dQhRR/#r.r)rs"rrr Nrc VP#rrrs&rrScan.get_pointwise_size rrcV^8dQhRR/#rr)rs"rrr 88c8rc `\VP4\VP4,#r)rrrrs&rrScan.index_length !4;;#d&6&6"777rcV^8dQhRR/#rr)rs"rrr s;rc VPVP4pVPVP\P4pVP W4pV3#r)rrrrLrrrrrrs& rrScan.inner_fn_args C DKK(T--t}}=ll5)v rc V^8dQhRRRR/#rr)rs"rrr UU4UDVUrc VPVP4pVPVP\P4pVP W#4p\ VPWAR7#r)rrrrLrrrUrrrrrrs&& rrScan.inner_fn_free_symbols M DKK(T--t}}=ll5)#DMM3TTrcan_fallback_to_atenTc@V^8dQhRRRRRRRRR R R R R RRRRRRR/ #)rr8rrrr!z+tuple[Callable[[Sequence[Expr]], Any], ...]r5raxisrrrrrarrrrrSequence[TensorBox | None]r)rs"rrr s{hhh(h? h  h  hRh&h#hh $hrc aaa.VRSOVS^,ROoVS,.o\PPV\P4'gR.\ V4,#\ V4^8dH\PPV\P 4'gR.\ V4,#\PPp V P\S44p \ V4\ V48wd \R4hV P\P!V ^44'dH\\ V44U u.uF(p \P!VW,,W<,VR7NK* up #VP#VV^,V^,SSSVV R7wr}\$pV ^8d\&P(P*RJ;'g\,;'d \.R8pT;'d1\ V4^8H;'d\&P0!4'*pV'gV'dR.\ V4,#^p M\2pRVVV3Rllp\\ V44U u.uFBp \4P!V!RRVRW,,R VR W<,R VR VR SRSRVRVRVRV / V B4NKD pp VFpVP74K V#uup iuup i)N&Expected len(dtypes) == len(inner_fns)rg)r8r7rrpointwise_rangesrr scan_numelz3.3.0c$V^8dQhRRRRRR/#)rrr/ scan_indexrrr)rs"rr!Scan.create..__annotate__E ! ? ?> ?~ ?* ?rc<\V4\S48wd \R4h\V4\S48wd \R4h.VRSOVOVSRO#)z,Expected len(scan_index) == len(scan_ranges),Expected len(index) == len(pointwise_ranges)Nrr)rrrrrs&&rrScan.create..reindexE _:#k"22$%STT5zS!122$%STT>U5D\>J>tu> >rr8r7rrr!r5rrrrrrr)r{r9r$rOSCANrTUPLE_REDUCTIONr:rrurrrrqrrDrrDrrDversionhip has_tritontriton_version$are_deterministic_algorithms_enabled SplitScanrr)rr8rr!r5rrrrrr:rrrD scan_typetriton_supports_splitsupports_splitrrMrrrs&&&&&f&&$, @@rr Scan.create s=T%4[<4q +;<Dzl ww""6>+>+>??6CK' ' v;?177#6#6 N22$ $ 6CK' '77##&&}['AB v;#i. ( !IJ J  ) )%((:q*A B B%*#f+$6%7L   ! .&4 ! %7 &)^^)q\-#!!&4 & " > !!T)WWj.V.V^w=V " &EEK1$EEBBDD  "' 6CK//!"J%  ? ?2!&c&k 2% $!3 #   ! ."'4  (   ,!, *$$2". "!3%  *F NN O^ s .K'AK,c<V^8dQhRRRRRRRRR R R R R R RRRR/ #)rr8rr7rRrrrrrrrrrrr"rrr)rs"rrrg se    7    (  # R   ( rc  TaaRVV3Rllp \PVVVV VVRVR7#)c$V^8dQhRRRRRR/#)rrr/r%rrzr)rs"rr%Scan.num_splits..__annotate__s s& H HN H> Hh Hrc,<S!.VRSOVOVSRO4#rr)rr%rrs&&rr#Scan.num_splits..wrapper_fns s*Fc%4jF=F3tu:FG Grr)r8rrrrrrr)rrD) rr8r7rrrrrrrs &&&ff&&&& rrDScan.num_splitsf sA H H###(!&$  rrr)rrrrrrer'rrr+r1r<rrrrr:rarrrDrrrs@rrr s MM SS!! ##--F#   $       8 U )6(=(=h&*hhT  rrc]tRtRtRtR#)ri rN)rrrrrrrrrr srrcJa]tRtRt$R]R&R]R&R]R&R]R&R ]R &R ]R &R ]R&R]R&R]R&]!R4R)RV3Rlll4tRV3RlltRRltRRlt RRlt RRlt RRlt R R!lt R"R#ltR)R$R%llt]]P$3R&R'll4tR(tV;t#)*Sorti r sort_rangesr5z:Callable[[Sequence[Expr], Sequence[Expr]], Sequence[Expr]]rrarrrrrrr!rstable descendingc V^8dQhRRRR/#rr)rs"rrSort.__annotate__ s   $  CU  rc <a\SV`S4\4P!V3RlVP4!,\4P!V3RlVP 4!,#)c3<<"TFp\VS4xK R#5irr*rs& rr,Sort.get_free_symbol_uses.. rrc3<<"TFp\VS4xK R#5irr*rs& rrr rr)rr'rDrrr5rs&frr'Sort.get_free_symbol_uses s] G ( 7l  Od>N>NO l  HdiiH   rcV^8dQhRR/#rr)rs"rrr rrc <\VP4\VP4,\VP48wd \ R4h\ SV`4R#)zCExpected len(self.ranges) + len(self.sort_ranges) == len(self.size)N)rrrr5rrrrs&rrSort.__post_init__ rrc ,V^8dQhRRRRRRRRRR /#) rrUrOrVrrXr/rrrr)rs"rrr s<     2    '   rc aVPW44o\;QJd!.V3RlVP4FNK 5M!V3RlVP44p\P!VP WPP VP4p\P!T;'gRV!S4W`P,4#)c32<"TF q!S4xK R#5irrrs& rr'Sort.store_reduction.. rrr[) rrr!rysortrrrr\r)rrUrVrXrrrrs&&&&& @rrSort.store_reduction sll40DT^^DDT^^DD$++v{{DOOLyy  $ $9gclF;L;L4M  rcV^8dQhRR/#rr)rs"rrr sJrc R#)rrrs&rr+Sort.get_reduction_type srcV^8dQhRR/#rrrr)rs"rrr s  $5 rc VP#r)rrs&rr1Sort.get_reduction_size rrcV^8dQhRR/#r%r)rs"rrr s+rc VP#rrrs&rr< Sort.get_size rrcV^8dQhRR/#rxr)rs"rrr rrc VP#rrrs&rrSort.get_pointwise_size rrcV^8dQhRR/#rr)rs"rrr rrc `\VP4\VP4,#r)rrrrs&rrSort.index_length rrcV^8dQhRR/#rr)rs"rrr s7rc VPVP4pVPVP\P4pVP W4pV3#r)rrrrLrrrs& rrSort.inner_fn_args rrc V^8dQhRRRR/#rr)rs"rrr rrc VPVP4pVPVP\P4pVP W#4p\ VPWAR7#r)rrrrLrrrUrrs&& rrSort.inner_fn_free_symbols rrc@V^8dQhRRRRRRRRR R R R R R RRRRRR/ #)rr8rrrr!z'tuple[Callable[[list[Expr]], Any], ...]r5rrrrrrrrarrrrr)rs"rrr szTTT(T; T  T  TTT&TT $Trc " aaa.VRSOVS^,ROoVS,.o\PPV\P4'gR.\ V4,#\PP p V P\S44p \PP'd\PPp MIRp \PP;'d&V P\P!W44p V 'gR.\ V4,#\ V4\ V48wd \!R4hV P\P!V ^44'dH\#\ V44Uu.uF(p\$P'VW.,W>,VR7NK* up#RVVV3Rllp\#\ V44Uu.uFGp\(P'\+RRVRW.,RVR W>,R VR VR SR SRVRVRVRVRV/ V B4NKI ppVFpVP-4K V#uupiuupi)Nrrrgc$V^8dQhRRRRRR/#)rrr/ sort_indexrrr)rs"rr!Sort.create..__annotate__ rrc<\V4\S48wd \R4h\V4\S48wd \R4h.VRSOVOVSRO#)z,Expected len(sort_index) == len(sort_ranges)rNr)rr:rrrs&&rrSort.create..reindex rrr8r7rrr!r5rrrrrrrr)r{r9r$rOSORTrr:rrurMrZdecompose_sort_opspersistent_reductionsrrrqrrrDrrr r)rr8rr!r5rrrrrr: sort_numelis_persistent_kernel max_rblockrrrMrrrs&&&&&f&&&, @@rr Sort.create s=T%4[<4q +;<Dzl ww""6>+>+>??6CK' '77##&&}['AB == + + +#)==#F#F J 33UU22588J3ST !$6CK' ' v;#i. ( !IJ J  ) )%((:q*A B B%*#f+$6%7L   ! .&4 ! %7  ? ?4!&c&k 2' &!3 %   ! ."'4  (   ,!,$$2"." * $!3'  ,F NN U" s !.I3A I"rr)rrrrrrer'rrr+r1r<rrrrr:rarrrrrs@rr r  s  GG!! ##-- LF#   $       8 U )6(=(=TTrr c V^8dQhRRRR/#rrrrrr)rs"rrr- sVrcF\VRR7R# \dR#i;i)FfreezeT)rrYrs&rr=r=- s&a. s   c V^8dQhRRRR/#rFr)rs"rrr5 s4rc \VRR7wrVP4'd\V\4'g\ \ V44h\ VPVPVPVPVPVPVP4VPVP4pVP4# \dR#i;iFrH)rshould_pad_stridesrr?rrrr8r7r5 _pad_stridesr6rr is_contiguousrYr_bufferrs& r is_contiguous_storage_and_layoutrR5 s/%@  $ $ & &fn55$T&\22    ##FMM6;; M   F##%% sCC C! C!c V^8dQhRRRR/#rFr)rs"rrrJ s  f  rc.\VRR7wrVP4'gR#\PPP VP 4VP\VP4,4# \dR#i;irL) rrOr{r9r:r storage_sizerrur5rYrPs& r&is_dense_contiguous_storage_and_layoutrVJ sx/%@##%%ww77    !6===3M#M   s%BAB BBc4V^8dQhRRRRRRRRRRR RR R /#) rrrrIrwant_contiguous stride_orderSequence[int | Integer] | Nonerrrztuple[StorageBox, Layout]r)rs"rrrV sN:: : ::1 :  : 2 ::rc \V\4'd\VPVVVVVR7#\V\4'd9\VPVVVVVR7wrgWPP 43#\V\ 4'dV'dV'dAVP4VP 4P4'g \R4hM>VeVPW4R7M'VeVPWTR7MVP4\ V4VP 43#\V\4'd'\VPVR7wrWP3#\h)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. rIrXrYrrz'Expected x.get_layout().is_contiguous()rrH)rrrrrr>rrrOrrrrrrrY) rrIrXrYrrr&rbuffers &&&&&& rrrV sU!Y$ FF+%''   !Z  ) FF+%''  &&##%%%!V !||~3355()RSS6)11 2*22!3!!}alln,,!_%%* FF xx rc$V^8dQhRRRRRR/#)rrrrYrrrr)rs"rrr s$ 4 rch\VRR7wr#VPV4# \dR#i;irL)ris_stride_orderedrY)rrYrQrs&& r"is_stride_order_storage_and_layoutrb s8/%@'' 55 s " 11c V^8dQhRRRR/#rr)rs"rrr sv$rc\V\\34'd\VP4#\V\ 4'dVP p\PPPVP\VP4,\4'*p\VP4;'gT#\V\4'd,VP!4\PP"9#R#r)rrr is_unalignedrrrr{r9r:statically_known_multiple_ofrrlr7rnrrunaligned_buffers)rrhas_unaligned_layouts& rrere s$J/00DII&&$((#$77#3#3#P#P MMN6<<8 8/$  DII&>>*>>$}}!''";";;; rcJ]tRtRt$R]R&]!R4R2RRll4tRRltRR ltR R lt ] R R l4t RRlt RRlt RRltRRltRRltRR/RRlltRRltRRltR R!ltR"R#ltR$R%ltR&R'ltR(R)ltR*R+ltR,R-ltR.R/ltR0tR1#)3ri rrc V^8dQhRRRR/#rr)rs"rrBaseView.__annotate__ s==$=CU=rc 8VPPV4#rrr'r&s&&rr'BaseView.get_free_symbol_uses syy--m<rc raaVPP4oVP4oRVV3RllpV#)c V^8dQhRRRR/#)rrr/rr"r)rs"rr+BaseView.make_indexer..__annotate__ s ' ' 'D 'rc <S!S!V44#rrrinnerrs&rrV&BaseView.make_indexer..indexer & &r)rrrs)rrVr{rs& @@rrBaseView.make_indexer s4 &&(%%' ' 'rcV^8dQhRR/#rr)rs"rrrk sArc raaVPP4oVP4oRVV3RllpV#)c V^8dQhRRRR/#r~r)rs"rr*BaseView.make_loader..__annotate__ s ' ' '8 'rc <S!S!V44#rrrzs&rr^$BaseView.make_loader..loader r}r)rrrs)rr^r{rs& @@rrBaseView.make_loader s4 %%'%%' ' ' rcV^8dQhRR/#rQr)rs"rrrk s%%{%rc 6VPP4#r)rrArs&rr7BaseView.dtype syy""$$rcV^8dQhRR/#r^r)rs"rrrk &&F&rc 6VPP4#rrr>rs&rr>BaseView.get_layout yy##%%rcV^8dQhRR/#rr)rs"rrrk &&/&rc 6VPP4#rrrBrs&rrBBaseView.get_device rrcV^8dQhRR/#rr)rs"rrrk !5rc R#rrrs&rrBaseView.get_origin_node rrcV^8dQhRR/#rr)rs"rrrk $$#$rc 6VPP4#rrrrs&rrBaseView.get_name yy!!##rcV^8dQhRR/#r.r)rs"rrrk sNrc "VP4#rrrs&rrBaseView.get_pointwise_size }}rrTc$V^8dQhRRRRRR/#rr)rs"rrrk &DDDTDTDrc :VPPWR7#rrrrs&&$rrBaseView.mark_reuse yy##E#CCrcV^8dQhRR/#rsr)rs"rrrk 222rc 6VPP4#rrrrs&rrBaseView.has_exceeded_max_reads yy//11rcV^8dQhRR/#rr)rs"rrrk ###rc 6VPP4#rrrrs&rrBaseView.realize yy  ""rcV^8dQhRR/#rr)rs"rrrk s!!d!rc :VPP4R#rrrrs&rrBaseView.realize_hint s  rcV^8dQhRR/#rr)rs"rrrk --8-rc 6VPP4#rrr rs&rr BaseView.get_storage_numel yy**,,rcV^8dQhRR/#rsr)rs"rrrk %%4%rc 6VPP4#rrr6rs&rr6BaseView.is_extern yy""$$rcV^8dQhRR/#rsr)rs"rrrk s,,$,rc \VP\4'g\\ VP44hVPP 4#r)rrrrris_module_bufferrs&rrBaseView.is_module_buffer s9$))X.. dii1 1yy))++rcV^8dQhRR/#rr)rs"rrrk ***rc 6VPP4#rrr rs&rr BaseView.get_read_names yy''))rcV^8dQhRR/#rr)rs"rrrk s?rc \P!\RR4;_uu_4\VP 4VP 44P uuRRR4# +'giR#;ir)r!rr?rWrr<rrs&rr BaseView.get_reads sJ \\.*:D A A&  " e B A A As 2A"" A3 cV^8dQhRR/#rr)rs"rrrksVrc RTp\V\4'dVPpK$V#r)rrr)rrs& rrBaseView.unwrap_views"H%%Arc V^8dQhRRRR/#r>r)rs"rrrkrbrc VP4p\P!\RV4!V4p\ VVP 4VVP 4R7#rd)rr!rrhrDrAr<ris&& rr@BaseView.constant_to_devicesN!!#n.?HP.."==?   rrNr)rrrrrrer'rsrrrYr7r>rBrrrrrrrr r6rr r rr@rrrrrr s LJ'=(=C%%&&$DDD2#!-%, *   rrcf]tRtRt$R]R&]RRl4t]RRl4tRR lt R R lt R t R #)rir/r5c$V^8dQhRRRRRR/#)rrrnew_sizeryrr)rs"rrExpandView.__annotate__s#**6*-?*DV*rc  \PPpVUu.uFp\P!V4NK ppVP 4pR.\ V4\ V4, ,\V4,p\ V4\ V48wd \R4h\\ V44EFpW,R 8Xd"WE,f \R4hWE,W&K2WE,e6\PPPWE,4'dKs\V4'dK\V4'dKW,pWE,pVf \R4hVf \R4hWg, pVPV^R7^8wgK\RVP 4 RV R V 24h V#uupi) zReplace `-1` with correct sizesNz'Expected len(new_size) == len(old_size)z Expected old_size[i] is not NonezExpected v1 is not NonezExpected v2 is not None)fallbackzBroadcast failed in ExpandView(, z) on dimension rl) r{r9r:rrUr<rrrris_size_one_or_falser<r;) rrr:rold_sizerv1v2diffs && r_normalize_sizeExpandView._normalize_sizes77##-56XELLOX6::<6S]S]:;d8nL x=CM ) !JK Ks8}%A{b ;&()KLL&k $(8(8(M(M ))./99 [[:()BCC:()BCCw..!"/  )9!**,r(Sbcdbef?&DO7sG c$V^8dQhRRRRRR/#)rrrrryrrr)rs"rrrDs"11v1);11rc VPW4p\V4'Ed7\V4wr4\V4\VP4, pV^8d \ R4h\ PP.V,p\VPVP4F_wrxTP\PPPV4'gTM\ PP4Ka \!VP"VP$\'V4VVP(VP*4p \-W9R7#\/WR7#)rzExpected skip >= 0r)rr5)rr=rrr5rrrrrr6r@r{r9r:rrr8r7rrrrr) rrrrrskiprr6r5rs &&& rrExpandView.createCs&&q3  # #"7": Gx=3z#77Dax$%9::'',,$.J #J$5$5z G !!77++@@FF!H %!!  X!!$$ J#C Cq00rcV^8dQhRR/#r.r)rs"rrr_rrc VP#rrrs&rr<ExpandView.get_size_rrcV^8dQhRR/#rpr)rs"rrrbs 3rc aaVP4pVPP4o\V4\S4, oRVV3RllpV#)c V^8dQhRRRR/#rrr/rr)rs"rr/ExpandView.make_reindexer..__annotate__is  !   rc<\VSR4p\V4\S48wd \R4h\\S44F/pSV,^8XgK\P P W&K1 V#)Nz"Expected len(index) == len(actual))rrrrrrr)rractualrs& rr*ExpandView.make_reindexer..reindexisetu&E5zS[($%IJJ3v;'!9>$ww||EH(Lr)r<rr)rtargetrrrs& @@rrsExpandView.make_reindexerbsE##%6{S[(  rrN) rrrrrrVrr:rr<rsrrrrrrs> **X116rrcr]tRtRt$R]R&]RRl4t]RRl4tRR ltR R lt R R lt Rt R#) PermuteViewixrdimsc$V^8dQhRRRRRR/#)rrrrrrrr)rs"rrPermuteView.__annotate__}s!..v.].x.rc  VPV4p\V4\\\V4448wd \ R4h\ V4'd\ V4wr4\VPVPVUu.uFqTPV,NK upVUu.uFqTPV,NK upVPVP4p\W6R7#\WR7#uupiuupi)z9Expected OrderedSet(dims) == OrderedSet(range(len(dims)))r)rr) _map_neg_dimsrDrrrr=rrr8r7r5r6rrrr)rrrrrrrs&&& rrPermuteView.create|s  & d z%D *:; ; K  ! # #"7": G$!!  -12T##T2/34t!""1%%t4!!$$ J#C C--34s C6 &C;c V^8dQhRRRR/#)rrrr list[int]r)rs"rrrsFFF9Frc `VUu.uFq"^8dTM\V4V,NK up#uupir)r)rrrs&& rrPermuteView._map_neg_dimss+@DEaxSY_4EEEs#+cV^8dQhRR/#r.r)rs"rrrs,,.,rc 6\VPVP44\\\ VP4448wd \ R4hVP P4pVPUu.uF q!V,NK up#uupi)zYExpected OrderedSet(self._map_neg_dims(self.dims)) == OrderedSet( range(len(self.dims)) ))rDrrrrrrr<)rr5rs& rr<PermuteView.get_sizes{ d((3 4 #dii. !9  !k yy!!#!%+AQ+++BcV^8dQhRR/#r.r)rs"rrrs..N.rc 6\VPVP44\\\ VP4448Xg \ R4hVP P4pVPUu.uF q!V,NK up#uupi)z.dims must be a permutation of range(len(dims)))rDrrrrrrrn)rr6rs& rrnPermuteView.get_stridess t))$))4 5%DII/0 1!!QR R%%'#'99-9aq 9---rcV^8dQhRR/#rpr)rs"rrrs 3rc la\VP4UUu/uFwrW!bK uppo\\VP44Uu.uF pSV,NK upo\ S4\ \\VP4448wd \ R4hRV3RllpV#uuppiuupi)z=Expected OrderedSet(inv) == OrderedSet(range(len(self.dims)))c V^8dQhRRRR/#rr)rs"rr0PermuteView.make_reindexer..__annotate__s + +! +  +rc><SUu.uF qV,NK up#uupirr)rrinvs& rr+PermuteView.make_reindexer..reindexs!'**c!HHc* **s)rrrrrDr)rrrBrrs& @rrsPermuteView.make_reindexers!*$)) 45 4qt 45$S^454!s1vv45 c?js499~)>? ? O  + + 65s B+ B1rN) rrrrrr:rrr<rnrsrrrrrrxsD ..*FF,.rrcV]tRtRt]RR/RRll4t]RRl4tRR ltR t R#) SqueezeViewirNc$V^8dQhRRRRRR/#)rrrrrrr)rs"rrSqueezeView.__annotate__s&1W1Wv1Wz1WV1Wrc  \V4'Edz\V4wr4.p.pVeW\V\4'g\ \ V44h^V8:dV\ VP48g \ R4h\\VPVP44FwpwrVfV\PPPV4'g%VPV4VPV 4K_KaWr8wd%VPV4VPV 4KV^8wgK\ R4h \!VP"VP$VVVP&VP(4p \+W:R7#Vff\,P/TVP14U u.uF6p \PPPV 4'dK4V NK8 up 4#VP14V,^8wd \ R4h\,P/T\VP144UU u.uFwr{Wr8wgK V NK up p4#uup iuup pi)Nz0Expected 0 <= dim and dim < len(old_layout.size)zexpected squeezed size to be 1rzExpected x.get_size()[dim] == 1)r=rrrrrrr5rrr6r{r9r:rr@rr8r7rrrrrr<) rrrrrrrrr5r6rrs &&$ rrSqueezeView.creates  # #"7": GHJ!#s++(c33SS3z+?%?(J&/s:??JDUDU/V%W!>D;77++@@FF -"))&1Gx -"))&119"01Q"RR&X%!!  !!$$ J#C C ;;;ZZ\)77++@@CA) zz|C A%$%FGG;;q1::<1H"U1HAH111H"UV V#Vs21I% (I%  I* I* c V^8dQhRRRR/#)rr5r/rz>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]]r)rs"rrrs!!! G!rc aaVUu.uF q^8wgK VNK pp\V4UUu.uFwr1V^8wgKVNK uppo\V4oRVV3RllpW$3#uupiuuppi)rc V^8dQhRRRR/#)rrr/rztuple[Expr, ...]r)rs"rr*SqueezeView.squeezer..__annotate__s $ $> $.> $rc<\V4\S48wd\V RS 24h\PP.S,p\ SV4F wr#W1V&K \ V4#) )rrrrrrr)rrrrlengthnot_ones& rr%SqueezeView.squeezer..reindexsb5zS\)$way%9::%*WW\\NV$;Igu-!"#.# #r)rr)r5rrrrrr s& @@rsqueezerSqueezeView.squeezersg $.t!AvAAt.!*4;AF11;T $ $  /;sAA A!A!c V^8dQhRRRR/#)rrrrrr)rs"rrrs99S9T9rc \R4h)zuse SqueezeView.create())r)rrs&&rr/SqueezeView.__init__s788rr) rrrrr:rrVr"r/rrrrrrs=1WD1W1Wf!!"99rrcv]tRtRt$R]R&R]R&RRltRR ltR R lt]t] R R l4t RRlt Rt R#) GenericViewir/r5rqrcV^8dQhRR/#rpr)rs"rrGenericView.__annotate__ s 3rc VP#r)rrs&rrsGenericView.make_reindexer s||rcV^8dQhRR/#rr)rs"rrr*sGGSGrc \\VP44Uu.uFp\\P V4NK pp\ VPV44pRRP\\V44 RV 2#uupi)zlambda rrQ) rrr5rtrLr;rrrLrJr)rr index_old index_news& r reindex_strGenericView.reindex_strsvCHTYYCX CXa *4::q 9CX  i01 3sI#6789+FF  s#BcV^8dQhRR/#rr)rs"rrr*s   rc zVPVPRVP 2RVP4 2.4#)size=zreindex=)rNrr5r1rs&rrGenericView.__str__s= YY% {+x8H8H8J7K-L M  rc(V^8dQhRRRRRRRR/#) rrrrr/rrqrrr)rs"rrr* s8AA A!A< A  Arc *V!V\V4VR7#)rr5r)r)rrrrs&&&&rrGenericView.createsX@@rcV^8dQhRR/#r.r)rs"rrr*(rrc VP#rrrs&rr<GenericView.get_size(rrrN) rrrrrrsr1rr9r:rr<rrrrr(r(sC  77 G HAArr(c]tRtRtRt]RRl4t]]RRl44t ]RRl4t ]RR R ll4t ]RR R ll4t Rt R #)ri,z This class handles tensor reshaping by computing appropriate index transformations to map the new shape back to the original storage layout. c$V^8dQhRRRRRR/#)rrr"r5rr)rs"rrView.__annotate__4s!4trc \P!V4p\P!V4p\PPP P pV!\P!V^44'd W,pV#r)rrUr{r9r:r rLt)rr5rs&& rhandle_negative_indexView.handle_negative_index3s[ll3||D!((22@@ #q) * **C rc$V^8dQhRRRRRR/#)rrrrr/rr)rs"rrr@>s&o5o5vo5o5Fo5rc aaaa\S\4'g\\S44hSP VP 4S4woo\ PPPSS4'dV#\\S44^8;'g\\S44^8p\V4pRRloRVVVV3Rllp^S9dRV3RllpS!V\S4VR7#V'dS!VS\PS44#\!V4'g V!V4#\#VRR 7wrxVP$p \ PPP'S4p \ PPP'V 4p \ PPP'S4p ^R IHp V !V V V VR 7pVeVUu.uFAp\-VR 4'dVP.P0M\2P4!V4NKC pp\7VP8VP:SVVP<VP>4p\AVVR 7#V!V4#uupi)rc(V^8dQhRRRRRRRR/#)rrrrr/rrrr)rs"rr!View.create..__annotate__Ms2 D D D#1 D?M D  Drc\PV\PVP 444p\ V4wr4\ VPVPVVVPVP4p\W5R7#)r) r+require_exact_stridesr?r@r<rrr8r7rrr)rrrrrrs&&& rcreate_reinterpret_view,View.create..create_reinterpret_viewMsz44^66s||~FC#8"< G$!!  !!$$ J#C Crc V^8dQhRRRR/#rrrrr)rs"rrrH^s     rc<SPSS4pS!V\S4VR7# \d6\P T4pS!TS\ P S44u#i;i)z 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)rrr;r+require_contiguousr?r@)rrrrKrrs& r"handle_unbacked_or_dynamic_reshape7View.create..handle_unbacked_or_dynamic_reshape^sm 55hIXHH. !33A6.x!B!B8!L  s%)=A)(A)c V^8dQhRRRR/#)rrrrztuple[int, ...]r)rs"rrrHvs 2 2C 2O 2rc<<\^.\S4,4#r)rr)rrs&r fake_reindex!View.create..fake_reindexvsaS3x=011rr9FrH)_compute_stride)size_obliviousrr)!rrrrresolve_negative_sizer<r{r9r:statically_known_list_equalsrr:rVrr?r@r=rr6to_symints_or_intstorch._subclasses.fake_implsrWr,rexprrr#rr8r7rrr)rrrunbacked_symbols_in_sizesrOrQrUrr old_strideold_size_symintold_stride_symintnew_size_symintrWnew_stride_symintrrrrKrsf&f @@rr View.create<sE(H-- h0 0 66qzz|XN( 77   8 88 L LH %h/ 01 4 8 8(23a7 "?qA  D"  , = 2 2ADNLI I*8^>>xH  %Q''5a8 84AeD&& ''**==hGGG,,?? K''**==hG@,   4    (+*A 'q&11 u}}Q7GG*  %!!  !!$$ J# C C2!44!s)AI?c$V^8dQhRRRRRR/#)rrr/rrztuple[list[Expr], list[Expr]]r)rs"rrr@s$"" ",:" &"rc VVUu.uF,p\PPPV4NK. ppVUu.uF,p\PPPV4NK. pp\ V4p\ \ V44FNpW,R8XgK\PPW&\\V4\V44W&M \PPP\V4\V44W3#uupiuupi)rrl) r{r9r:rrrrrrOnerFru check_equals)rrrrs&& rrYView.resolve_negative_sizes;CC(QAGG$$--a0(C:BC(QAGG$$--a0(C>s8}%A{b #ggkk &}X'> h@WX  & %%mH&=}X?VW!!DCs 2D!2D&Nc(V^8dQhRRRRRRRR/#)rrryr dense_dimrrrr)rs"rrr@s2$%  0 rc VPWV4pV# \\\3d>\ T4.pTPY4pTPYR4p\ Yg4pT#i;ir)_dynamic_reshape_indexerrr; IndexErrorrur )rrrrkrflatrrs&&&& rrView.dynamic_reshape_indexersu :228yQG  ;ZH :!(+,D33HCH33DCH%h9G  :sAA)(A)c(V^8dQhRRRRRRRR/#)rrr/rrkrrrqr)rs"rrr@s6}} } }} 4 }rc taaa\PPPoRV3Rllp\ \ V44Uu.uFp\ \PV4NK upo\\SV44p\V4pVRJ;'d,V\ V4^, 8g;'d\ V4^8HpV'd2Vf \R4hVPV4pVPV4.oV'EdV'EdVP4p VP4wrV ^8Xd>SP\PP 4VPW34KxV ^8XdVPV 4KV!W4^8XdSPV 4KV!W4^8dvV!W4^8d,VP4wrW,V ,p W,p K8SPV 4\PPP#W4EK4V!W4^8d\PP$pT pSP\'WV44W,pV!W4^8d>VP4pSP\'WV44W,pW,p KJ\PPP#W4EK\hV'dfVP4p \PPP#V ^4SP\PP 4KmV'd?VP4wr\PPP#V ^4KFVeD\ V4^8Xd4SP)4SP4pSP+VV4MSP)4\ S4\ V48wd \R4hRVV3RllpV#uupi)z7 Perform a reshape entirely by modifying indexing math c$V^8dQhRRRRRR/#)rrr"rrrr)rs"rr3View._dynamic_reshape_indexer..__annotate__s! > >T >d >s >rc<S!\P!W44'd^#S!\P!W44'dR#S!\P!W44'd^#\P P PW4'dR#\P P PW4'd^#\\P!W44h)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. rl) rrjrBrr{r9r:rfr;)rrrrs&&r compare_sizes4View._dynamic_reshape_indexer..compare_sizessehhqn--ehhqn-- ehhqn--ww<.reindex..MsHiA|44ir)rrrrr)rr{rX view_exprs&@rr.View._dynamic_reshape_indexer..reindexGs`5zSY&$c%j#d)%<==D% 01L5HiH5 H5HiHH Hr)r{r9r:rrrrrtrLVIEWrrrrr@rrrrhrgrKreverseinsert)rrrkrvr stack_new stack_oldreordering_dense_dimold_dimsize_oldr size_newvar2 size_new2divisormodulus dense_exprrrrrXr}s&&& @@@rrmView._dynamic_reshape_indexersC))88 > >JCHH BV BVQ *499a 8BV T8,- N  T ! # #S^a// # #H "   $%EFFmmI.G   W % iII }}H%MMOMC1}  .  #1Q  *x2a7  %x2Q6#H7!;&/mmoOD/C/C'3H  %  --hAx2Q6''++"  w!GH!+#H7!;'mmoG$$_S7%KL%/G'1H  --hA$$ }}H GG   ) )(A 6   UWW\\ *%MMOMC GG   ) )(A 6  S]a%7    "J   Y 3     y>S] * !KL L I I_ s#P5rr)rrrrrSrVrCr:rrrYrrmrrrrrr,sx  o5o5b""  }}rrca]tRtRt$RtR]R&RV3RlltRRlt]tR R lt R R lt R Rlt ] RRl4t RRltRRltRRltRRltRRltRRlt]!R4R$RRll4tR%RR lltR!R"ltR#tV;t#)&riRz*Pretend our storage has a different layoutr_rcV^8dQhRR/#rr)rs"rrReinterpretView.__annotate__XsFFtFrc <\SV`4\VP\4'd2\ P VRVPP44R#R#)rN)rrrrrrrrrs&rrReinterpretView.__post_init__XsC  dii * *   tVTYY-B-B-D E +rcV^8dQhRR/#rr)rs"rrr]s   rc PVPVPVP.4#r)rNrrrs&rrReinterpretView.__str__]s&     rcV^8dQhRR/#rr)rs"rrrgrrc 6VPP4#rrrs&rrReinterpretView.get_namegrrcV^8dQhRR/#rr)rs"rrrjrcrc .VPP#r)rr8rs&rrBReinterpretView.get_devicejs{{!!!rcV^8dQhRR/#rr)rs"rrrmrrc R#rrrs&rrReinterpretView.get_origin_nodemrrcV^8dQhRR/#rQr)rs"rrrq!!{!rc .VPP#r)rr7rs&rr7ReinterpretView.dtypeps{{   rcV^8dQhRR/#r.r)rs"rrrts&&.&rc @\VPP4#r)rrr5rs&rr<ReinterpretView.get_sizetsDKK$$%%rcV^8dQhRR/#r.r)rs"rrrws((N(rc @\VPP4#r)rrr6rs&rrnReinterpretView.get_stridewsDKK&&''rcV^8dQhRR/#rr)rs"rrrz  A rc aRV3RllpV#)c V^8dQhRRRR/#r<r)rs"rr1ReinterpretView.make_loader..__annotate__{s " ". "X "rc\<SPP4p\P!SP 4V!V44pSPP SP P 8wd7\P!VSP SP P 4#V#r)rrryloadrr7rto_dtype_bitcast)rrV tmp_loaderrs& rr^+ReinterpretView.make_loader..loader{spkk..0G$--/75>BJ{{  DIIOO3++J DIIOOTT!!rrrr^sf rrReinterpretView.make_loaderzs " " rcV^8dQhRR/#rr)rs"rrr**>*rc 6VPP4#r)rrrs&rrReinterpretView.make_indexer{{''))rcV^8dQhRR/#r^r)rs"rrrFrc VP#rrrs&rr>ReinterpretView.get_layoutrrcV^8dQhRR/#rr)rs"rrr  t rc R#rrrs&rrReinterpretView.freeze_layoutrrc V^8dQhRRRR/#r#r)rs"rrr  ! ! rc \VPPV4\VPPV4,\VPPV4,#r)r+rr5r6rr&s&&rr'$ReinterpretView.get_free_symbol_usessQ T[[--} =t{{11=A Bt{{11=A B rc V^8dQhRRRR/#rr)rs"rrrs   (=    rc  x\PPPVPVP P VP PVP PVe VPM#\PPPVP PR7#)NrT) r{r9 wrapper_codecodegen_reinterpret_viewrrr5r6r writeliner7rs&&rr!ReinterpretView.codegen_referencesww##<< II KK   KK   KK   & 2F  8L8L8V8V++## =  rcV^8dQhRR/#rr)rs"rrr3rc ^#rrrs&rrReinterpretView.num_readsrrrr)rrrrrSrrrr9rrBrrYr7r<rnrrr>rrer'rrrrrs@rrrRs4 NFF  H$"!!&( * -. /   rrc]tRtRt$RtR]R&]RRl4tRRlt]t ] R R l4t R R lt R Rlt RRltRtR#) DtypeViewiz(Pretend our storage has a different typerR target_dtypec$V^8dQhRRRRRR/#)rrr new_dtyperRrrr)rs"rrDtypeView.__annotate__s! 9 9v 9+ 9( 9rc \V4'd\\V4wr4\VPVVPVP VP VP4p\W5R7#\WR7#)r)rr) r=rrr8r5r6rrrr)rrrrrrs&&& rrDtypeView.createsj  # #"7": G$!!!!!!$$ J#C Ca88rcV^8dQhRR/#rr)rs"rrrs???rc PVPVPVP.4#r)rNrrrs&rrDtypeView.__str__s  4+<+<=>>rcV^8dQhRR/#rQr)rs"rrrrrc VP#r)rrs&rr7DtypeView.dtypes   rcV^8dQhRR/#r.r)rs"rrr$$.$rc 6VPP4#rrr<rs&rr<DtypeView.get_sizerrcV^8dQhRR/#rpr)rs"rrrs Jrc RRlpV#)c V^8dQhRRRR/#rr)rs"rr.DtypeView.make_reindexer..__annotate__s  > n rcV#rrrBs&rr)DtypeView.make_reindexer..reindexsLrr)rrs& rrsDtypeView.make_reindexers rcV^8dQhRR/#rr)rs"rrrsArc RaaSPP4oRVV3RllpV#)c V^8dQhRRRR/#r~r)rs"rr+DtypeView.make_loader..__annotate__s X X X8 Xrc|<\P!S!V4SPSPP4#r)ryrrrr7)rr{rs&rr^%DtypeView.make_loader..loaders*''c D4E4EtyyW Wrrr)rr^r{sf @rrDtypeView.make_loaders( %%' X X rrN)rrrrrSrr:rrr9rYr7r<rsrrrrrrrsO2 9 9?H !!$ rrc`]tRtRtRt]RRl4t]RRl4t]R RRll4tR t R #) SliceViewiziView that represents a slice along a single dimension. Corresponds to tensor[..., start:end:step, ...]. c 0V^8dQhRRRRRRRRRRR R/#) rrrrrstartr"r5steprr)rs"rrSliceView.__annotate__sF;; ;; ;  ;  ; ;rc taaaa\VP44oVSS&RVVVV3RllpV!VSVR7#)c V^8dQhRRRR/#rr)rs"rr0SliceView.create_with_size..__annotate__s  !   rc<\V4\S48wd\RV RS 24h\V4pVS,S,S,VS&V#)z wrong ndim r)rrr)rrrrrs&rr+SliceView.create_with_size..reindexsR5zS]*${5'8*%EFFKEsd*U2E#JLrr9)rr<)rrrrr5rrrs&&ff&f @rcreate_with_sizeSliceView.create_with_sizes9 %   '::rc ,V^8dQhRRRRRRRRRR/#) rrrrrrendrztuple[int, int]r)rs"rrrs4/// /),/36/ /rc (aaaaa a \PPo VP4V,o\;QJdRW4S34F 'gK RM RM!RW4S344'd\ o \ oMc\;QJdRW4S34F 'gK RM RM!RW4S344'd\ o \ oMS Po S PoRVV V 3RlloRVVV3RllpV!V^S^4pV!WCSS4pW43#) zb Normalize start and end such that both are in the range [0, x.get_size()[dim]] and start <= end. c38"TFp\V4xK R#5ir)r:rrs& rr0SliceView.normalize_start_end..sH1GA$Q''1GrTFc3"TFUp\V\4'gKVP\P\P \\ 4xKW R#5ir)rr"hasrrIrHr s& rrr s? ,!T" 2AEE%))UYYS 1 1+s AAAc(V^8dQhRRRRRRRR/#)rrr"lowerrupperrr)rs"rr3SliceView.normalize_start_end..__annotate__s( T # c d rc<SPW4'dTMS!W4pSPW24'dTpV#S!W24pV#r)statically_known_geqr)rr r clamped_lower clamped_fullmax_funcmin_funcr:s&&& rclamp,SliceView.normalize_start_end..clamps^221<<(1BT  00FF   m3   rc ,V^8dQhRRRRRRRRRR/#) rrfrr rr r| Expr | intrr)rs"rrr s4 , , ,$' ,03 ,>H ,  ,rcF<VfV#SPVS4pS!WV4#r)rC)rfr r r|r rdim_sizes&&&&r clamp_wrap1SliceView.normalize_start_end..clamp_wraps.{++C:CU+ +r) r{r9r:r<rrIrH evaluate_min evaluate_max) rrrrr r r r r r r:s f&&&& @@@@@rnormalize_start_endSliceView.normalize_start_ends77##::<$ 3H%h1GH333H%h1GH H HHH S (+ SSS (+   HH,,H,,H  , ,5!Xq1Xx8zrc4V^8dQhRRRRRRRRRRRR R R/#) rrrrrrr rr rrr)rs"rrr)sY,H,H ,H,H ,H  ,H  ,H,H ,Hrc  \P!V4p\V\4'gV^8g \ V4hV^8XdVR8d V^8XdV#\ TP44pT'dTPYY44wr4\YC, T^, ,T4Yr&\T4'd\T4wr\ T P4p Y,T,Y&\T PT PTT T P T PT,T,,T P"4p \%YR7#TP'YY7T,T4# \ dEL&i;i)rrl)rrUrr"r TypeErrorrr<r rGr=rr6rr8r7rrrr) rrrrr rr rrrrrs &&&&&&& rrSliceView.create(sM||D!4&&$( & & zcY.419 % 00DJE q!94@  # #"7": Gj//0J(o4JO$!!  !!J$5$5c$:U$BB$$ J#C C##AEC=$GG;   sE// E>=E>rN)rT) rrrrrSr:rr rrrrrrrsF ;;,//b,H,Hrrc\]tRtRt$R]R&R]R&RRltRR ltR R ltR R ltRt R#) BaseConstantiXrRr7rr8cV^8dQhRR/#r.r)rs"rrBaseConstant.__annotate__]s.rc R#Nrrrs&rr<BaseConstant.get_size]rOrcV^8dQhRR/#rr)rs"rrr& `rrc VP#rrrs&rrBBaseConstant.get_device`rrcV^8dQhRR/#rr)rs"rrr& crrc R#rrrs&rrBaseConstant.get_origin_nodecrrcV^8dQhRR/#rr)rs"rrr& f?rc \4#rrCrs&rr BaseConstant.get_readsfrrrN) rrrrrr<rBrr rrrrr$ r$ Xs'  rr$ cZ]tRtRt$R]R&R]R&R]R&RR ltR R ltR R ltRtR#)ConstantijrrJrRr7rr8cV^8dQhRR/#rr)rs"rrConstant.__annotate__pArc aRV3RllpV#)c V^8dQhRRRR/#r<r)rs"rr*Constant.make_loader..__annotate__qs 8 8. 8X 8rcZ<\P!SPSP4#r)ryr@rJr7rrs&rr^$Constant.make_loader..loaderqs<< DJJ7 7rrrsf rrConstant.make_loaderps 8 8 rcV^8dQhRR/#rr)rs"rrr7 v   rc R#rrrs&rrConstant.realizevrrc V^8dQhRRRR/#r>r)rs"rrr7 ysKKK&Krc F\VPVPVR7#))rJr7r8)r5 rJr7rs&&rr@Constant.constant_to_deviceysdjj 6JJrrN) rrrrrrrr@rrrrr5 r5 js) J    KKrr5 cN]tRtRt$R]R&R]R&R]R&RR ltR R ltR tR #)IndexingConstanti}rrrRr7rr8cV^8dQhRR/#rr)rs"rrIndexingConstant.__annotate__r8 rc aRV3RllpV#)c V^8dQhRRRR/#r<r)rs"rr2IndexingConstant.make_loader..__annotate__s : :. :X :rcZ<\P!SPSP4#r)ryrVrr7r= s&rr^,IndexingConstant.make_loader..loaders>>$**djj9 9rrrsf rrIndexingConstant.make_loaders : : rc V^8dQhRRRR/#r>r)rs"rrrJ sSSS&Src F\VPVPVR7#))rr7r8)rH rr7rs&&rr@#IndexingConstant.constant_to_devicesdjj 6RRrrN)rrrrrrr@rrrrrH rH }s$ J   SSrrH c$V^8dQhRRRRRR/#)rr6ryr{rrr)rs"rrrs$ '9 rcd^p^p\\\W444FwrEV^8XdK\PP P WR4'g3\PP P WS4'gR#V\^V4,pW$,pK R#rFT)reversedrrr{r9r:rrH)r6r{expected_strideexpected_stride_maxrys&& ris_contiguous_strides_for_shaper[ sOs5123 6 ww77   ''""::1RRs1ay(4 rc V^8dQhRRRR/#)rr7rRrrr)rs"rrrs<<{\\V4P4hrrrs&rrBOutputSpec.get_devicerrcV^8dQhRR/#rr)rs"rrrb s77c7rc >\\V4P4hrrrs&rrUOutputSpec.storage_sizerrc V^8dQhRRRR/#r#r)rs"rrrb r$rc >\\V4P4hrrr&s&&rr'OutputSpec.get_free_symbol_usesrrrNr) rrrrrSrBrUr'rrrrrSrSsC7777rrSc]tRtRtRtR]!^4R3RRllt]RRl4t]PR R l4t]R R l4t ] PR Rl4t ]RRl4t ] PRRl4t RRlt ] t RRltRRltRRlt]RRl4tRRlt]RR l4tR!R"ltR#R$lt]R%R&l4tR'R(ltR)R*ltR+R,ltR-R.ltR/R0ltR1R2lt]!R4R6R3R4ll4tR5t R#)7r_iz_ Layout base class Carries tensor meta-information including offset and whether it is pinned. NFc4V^8dQhRRRRRRRRR R R R R R/#)rr8rr7rRr5ryr6zSequence[Expr] | Nonerr"rrrrr)rs"rrLayout.__annotate__sYCCCC! C & C  CC Crc Vf\PV4pWnW n\ V4\ V48wd\ RV RV 24h\ ;QJdRV4F 'dK RM RM !RV44'g \ R4hW0nW@nWPn W`n VP'd)VPPR8Xg \ R4hR#R#) Nr5 , stride=c3B"TFp\V\4xK R#5ir)rrrs& rr"Layout.__init__..s9Dq:a**DFTz3Expected all(isinstance(s, _IntLike) for s in size)rXzOnly CPU tensors can be pinned) r?r@r8r7rrr_size_stride_offsetrr)rr8r7r5r6rrs&&&&&&&rr/Layout.__init__s >#66t>>*4>>*:;MDz""#2dkk&6&6%78H7ITZZLYII;i }VH]O1 N rcV^8dQhRR/#rr)rs"rrrm sLrc VP#rrrs&rrBLayout.get_devicerrcV^8dQhRR/#)rrr0r)rs"rrrm s\rc  .\P;_uu_4\P!\ VP 4\ VP 4VPVPVPR7uuRRR4# +'giR#;i))r7r8 pin_memoryN) r{ fake_moderDrErhr5r6r7r8rrs&r get_exampleLayout.get_examplesR [[[&&' 2' 4jj{{>> [[[s AB B cV^8dQhRR/#rsr)rs"rrrm sGGtGrc B\VPVP4#r)r[ r6r5rs&rrOLayout.is_contiguous s.t{{DIIFFrc$V^8dQhRRRRRR/#)rr{ryrrrr)rs"rrrm s$!,> rc \V4pVR9gV^,^8XdR#\V\V4V4Fwr4pV^8wgKW48wgKR# R#)FT)r )rrr5)r{rndimleftrightr5s&& ris_channels_last_contiguous"Layout.is_channels_last_contiguoussY5z v qQ!$  *5 1  " DqyT]" rcV^8dQhRR/#rsr)rs"rrrm "strc  \VP\\P \ \VP 4444VP 4FwrpV^8wgKW8wgKR# R#rV )rr6rW r?r@rr5)rr r r5s& r is_transposedLayout.is_transposed"sZ!$ KK ^66tHTYY.__annotate__Ms B B  B- Brcb\V4pVUu.uFq!PV4NK up#uupir)rr)r sorted_arrelements& rsorted_indices0Layout.is_stride_ordered..sorted_indicesMs,J=@AS'$$W-SA AAs,FTrl)rr6rrr5r{r9r:rrrrrrjr ) rrrr non_1_indicesr6r stride_orderedr r s && rraLayout.is_stride_ordered?su t{{ s5z ) !JK K $DII. .77##;;CC A.  +88-Q++a..-82?@-Qa-@ B u%E *s5z"A'-yN8 $#s5zA~&A!$D"q5)Eww..uxx/DEE,,T9900=='= 9@s4GG G.G cV^8dQhRR/#rsr)rs"rrrm ds---rc  ^.\\\^\VP4^, 444,p\V4.V,pVP V4#r)rrW rrr6rars& ris_channels_last_stride_ordered&Layout.is_channels_last_stride_ordereddsNd8E!S-=-A$BCDDU u$%%e,,rc(V^8dQhRRRRRRRR/#)r in_stridesrr5r/r7rRrr)rs"rrrm ks2LL!L)7L@KL Lrc aa\V4p\V4^8XdV#\P'g\P W4'dV#\ P!4p\VR4'd%VPPRR4'dV#\;QJd0R\P!W44F 'dK RM% RM!!R\P!W444'*p\P'g V'dV#\\ PR4'd\ PP MRoRV3R lloS'dC\";QJdV3R lV4F 'gK RM RM!V3R lV44'dV#\%VS4p\'V4p\)\V44Uu.uFp^NK p p^W^,&Rp \+VR ,^R 7FwrW{^, ,p W,W,,p\-V\.\0P234;'d/V\P48;'dW,^8g;'g)\-V\64;'d\PpWV &V'gK\9W4V,W&Rp K V 'gV#\:;P<^, unV #uupi) zv The padding does not change stride order but makes sure all strides larger than the threshold are multiple of align. rdislike_paddingFc3b"TF%p\V\\P34xK' R#5irrrs& rr&Layout._pad_strides..s* 6 q3 . / /6rT _shape_envNc V^8dQhRRRR/#)rr]r rrr)rs"rr)Layout._pad_strides..__annotate__s S SJ S4 Src<SfR#\V\4'gR#\;QJd)V3RlVP4F 'gK R# R#!V3RlVP44#)NFc3F<"TFpSPV4xK R#5ir)is_unbacked_symint)rrr s& rrILayout._pad_strides..contains_unbacked_symints..s!R@Q1y33A66@Qs!T)rr"rr9)r]r s&rcontains_unbacked_symints6Layout._pad_strides..contains_unbacked_symintssQ dD))3R@Q@QR33 R3 R3R@Q@QRR Rrc34<"TF pS!V4xK R#5irr)rrr s& rrr sN:a6q99:r+rNN)r)r_ rrMpad_channels_lastr_r r{get_current_noder,rr1rrPchainpad_dynamic_shapesr9r rr)r#rrrrrr#padding_stride_thresholdr"rfr)num_comprehensive_padding)r r5r7aligncurrent_fx_noderprYr"r& new_stridespaddedrankrprev_idxr6require_paddingr r s&&& @@rrNLayout._pad_stridesjst$E* z?a  '''F,N,N - -  ,,. ?F + +0D0D0H0H u1 1   __Z6  __Z6   (((Z *1!''<*H*HAGG&&d  S S N:NN:NNN ' I> ,\: "'J"89"8Qq"8 9&' qM"":b>;ID!(+H *T^;F6C#78((V<<<((Na'JJVT*HHv/H/H  & #*6#9E#A  <  ))Q.);:s LcV^8dQhRR/#rr)rs"rrrm sLLTLrc \V\4'g\\V44hVPf \R4hVP VPVP VP4VnR#)Nz Expected self.stride is not None)rr?rrr6rNr5r7rs&r pad_stridesLayout.pad_stridessV$// d, , ;;  !CD D'' TYY K rcV^8dQhRR/#rsr)rs"rrrm sQQDQrc R\P;'d\V\4#r)rMcomprehensive_paddingrr?rs&rrMLayout.should_pad_stridess++PP 40PPrcV^8dQhRR/#rrrr)rs"rrrm s   +  rc \V\4'dV#VP4'dVP4\VPVP VP VPVPVP4#r) rrrMr r8r7r5r6rrrs&ras_fixedLayout.as_fixedsf dK ( (K  " " $ $     KK JJ II KK KK NN   rcV^8dQhRR/#rr)rs"rrrm s..>.rc \P'g#\R\V4P R24hVP 4P 4#)zconvert z to FixedLayout first)r?rrrrr rrs&rrLayout.make_indexersC,,, 8DJ,?,?+@@U!VW W}}++--rc V^8dQhRRRR/#)rotherrrrr)rs"rrrm s   F  t  rc \V\4;'dVPVP8H;'dVPVP8H;'d}VPVP8H;'d\VP VP 8H;'d;VP VP 8H;'dVPVP8H#r)rr_r8r7r5r6rr)rr s&&r__eq__ Layout.__eq__s uf % 2 2 u||+ 2 2 ekk) 2 2 UZZ' 2 2 u||+  2 2  u||+  2 2 %//1 rcV^8dQhRR/#rr)rs"rrrm sTTdTrc X\VPVPVP4#r)r1r5r6rrs&rrULayout.storage_sizes.tyy$++t{{SSrc V^8dQhRRRR/#r#r)rs"rrrm rrc \VPV4\VPV4,\VPV4,#r)r+r5r6rr&s&&rr'Layout.get_free_symbol_usess= TYY 6t{{M: ;t{{M: ; r)ru rs rt r8r7rr6r)!rrrrrSr#r/rYr5setterr6rrr9rBr rOrVr r r rar rNr rMr rr rUrer'rrrrr_r_sE)-qzC4 [[ ]] ]]  HG"  $#J- LL\LQ  .  TH% & rr_c&]tRtRtRtRRltRtR#)riz A Tensor layout we cannot changecV^8dQhRR/#rr)rs"rrFixedLayout.__annotate__sCC>Crc X\VPVPVP4#)r)r_r5r6rrs&rrFixedLayout.make_indexersdiidkkBBrrN)rrrrrSrrrrrrrs*CCrrca]tRtRtRtRtRRlt]RRl4t]RR l4t ]R R l4t ]R R l4t ]RRl4t ] RRl4t]PRRl4t] RRl4t]PRRl4t] RRl4t]PRRl4tR+RRlltR+RRlltR R!ltR"R#ltR$R%ltR&R'ltR,R(V3R)llltR*tV;t#)-r?izp A Tensor layout that we are allowed to change Assumption: layout change should NOT add or remove free symbols FcV^8dQhRR/#r r)rs"rrFlexibleLayout.__annotate__s..;.rc J\P!V4P4#)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. )rdeepcopyr rs&r!get_fixed_layout_without_freezing0FlexibleLayout.get_fixed_layout_without_freezings}}T"++--rc V^8dQhRRRR/#)rsizesrrrr)rs"rrr s00-0J0rc \V4^8Xd.#\PP.p\ VR,4F!pVP W!R,,4K# \ \ V44#)rr rl)rrrrgrW r@r)r reversed_stridesr5s& rr@!FlexibleLayout.contiguous_strides s\ u:?I!GGKK=U2Y'D  # #DB+?$? @(H-.//rc$V^8dQhRRRRRR/#)rr rrrrr)rs"rrr s!M-Jrc \\\V444\V48wd \W34h\P P pR.\V4,pVFpW#V&W V,,pK V#)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)rDrrrrrrg)r r next_striderrs&& r fill_orderedFlexibleLayout.fill_orderedsm eCJ' (Ju,= = %0 0ggkk &3u:%A$AJ%a0Krc$V^8dQhRRRRRR/#)rr rrrr/r)rs"rrr 's! > >m >M >n >rc \\\V444\V48wd \R4h\ V4p\ P W4#)zz Create a stride based on the sorted order of a permuted range. In this format, channels last would be: [3, 0, 2, 1] z;Expected OrderedSet(range(len(sizes))) == OrderedSet(order))rDrrrr#r?r )r rr"s&& rr FlexibleLayout.stride_ordered&sL eCJ' (Ju,= = M -U3 **5==rc$V^8dQhRRRRRR/#)rr r memory_formatztorch.memory_formatrr/r)rs"rrr 6s$&&&-@& &rc JV\P8Xd\PV\4#V\P 8Xd\PV\ 4#V\P8Xd\PV4#\PRV4\h)a9 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) rD channels_lastr?r NHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatr@r.r/rY)r r s&&r stride_ordered_for_memory_format/FlexibleLayout.stride_ordered_for_memory_format5s E// /!008IJ J e44 4!008JK K e55 5!44U; ; IIP & %rc$V^8dQhRRRRRR/#)rr rr6ryrr/r)rs"rrr Qs$ > > >&8 >  >rc \V4\V48wd \R4h\PPP V4p\ \\V44VPR7p\PW4#)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] z"Expected len(sizes) == len(stride)r) rrr{r9r:guarding_hints_or_throwrr __getitem__r?r )r r6r"s&& r same_orderedFlexibleLayout.same_orderedPsf u:V $ !EF F!!99&AE#f+.F4F4FG **5==rcV^8dQhRR/#r.r)rs"rrr asnrc VP#rry rs&rr5FlexibleLayout.size`r{ rc V^8dQhRRRR/#r r)rs"rrr es.Trc 6VPRV4WnR#)r5N)!assert_free_symbol_uses_unchangedrs r~ s&&rr5r, ds ..vu= rcV^8dQhRR/#r.r)rs"rrr jr rc VP#rr rs&rr6FlexibleLayout.strideir rc V^8dQhRRRR/#r r)rs"rrr nsNtrc 6VPRV4WnR#)r6N)r/ rt r~ s&&rr6r2 m ..x? rcV^8dQhRR/#rr)rs"rrr sr rc VP#rr rs&rrFlexibleLayout.offsetrr rc V^8dQhRRRR/#r r)rs"rrr wsDTrc 6VPRV4WnR#)rN)r/ ru r~ s&&rrr8 vr5 rc$V^8dQhRRRRRR/#)rrrrrrrr)rs"rrr {s$  " 37  rc FVPVPV4pVP4'd/V'd'VPW0PVP4p\ VP VPVPVVPVP4#r) r r5rMrNr7rr8rr)rrrrs&&& ras_stride_orderFlexibleLayout.as_stride_order{sx((E:  " " $ $**:yy$**MJ KK JJ II  KK NN   rc$V^8dQhRRRRRR/#)rrryrrrrr)rs"rrr s$  / @D  rc TpVP4'd/V'd'VPW0PVP4p\ VP VPVPVVP VP4#r)rMrNr5r7rr8rr)rrrrs&&& ras_exact_stridesFlexibleLayout.as_exact_stridessg#  " " $ $**:yy$**MJ KK JJ II  KK NN   rc V^8dQhRRRR/#)rrrrrr)rs"rrr s   =  [  rc 6VPVPV4pVP4'd'VPW PVP4p\ VP VPVPVVPVP4#r) r r5rMrNr7rr8rr)rrrs&& r as_fill_orderFlexibleLayout.as_fill_orderst$($5$5dii$G  " " $ $**:yy$**MJ KK JJ II  KK NN   rc V^8dQhRRRR/#)rr6ryrrr)rs"rrr s   $6  ;  rc 6VPVPV4pVP4'd'VPW PVP4p\ VP VPVPVVPVP4#r) r( r5rMrNr7rr8rr)rr6rs&& r as_same_orderFlexibleLayout.as_same_orderst&&tyy&9  " " $ $**:yy$**MJ KK JJ II  KK NN   rcV^8dQhRR/#)rrz$dict[tuple[str, bool], sympy.Symbol]r)rs"rrr s $ $.R $rc  t/pRF/pRF&pW#3p\\\W4V44W&K( K1 V#)r5)r5r6rTF)rDr+r)rinitial_free_symbolsrrr s& rget_initial_free_symbol_uses+FlexibleLayout.get_initial_free_symbol_usessI!0D!. +,6$WT%8-H-$)"/1$#rc$V^8dQhRRRRRR/#)rrrrJr=rrr)rs"rrr s"c,SWrc RFBpVPW3,p\\W#44pWT8wgK3\RV RV 24h R#)Tz)Expected free symbols unchanged, but got z vs NrM )rN rDr+r)rrrJrold_free_symbolsnew_free_symbolss&&& rr/ 0FlexibleLayout.assert_free_symbol_uses_unchangedsY*M#88$9NO )*:5*PQ 3$?@P?QQUVfUgh +rc 0V^8dQhRRRRRRRRR R R R /#) rr8rr7rRr5r/rYrZrrrrr)rs"rrr sNHHHH H 5 H  H Hrc <V'd\PW44pM\PV4p\SV`WW6VR7VP 4VnR#)rN)r?r r@rr/rO rN )rr8r7r5rYrrrs&&&&&& rr/FlexibleLayout.__init__sO $11$EG$77=G K%)$E$E$G!r)ru rs rt rN r)NF)rrrrrSrr rVr@r r r" r( rYr5r r6rr= rA rE rI rO r/ r/rrrs@rr?r?s, N.00" > >&&4 > > [[ ]] ]]      $HHHrr?cpa]tRtRtRtRV3RlltRRltRRlt]!R4R R R ll4t R t V;t #) NonOwningLayoutiz,Is a view into the storage of another tensorc V^8dQhRRRR/#)rviewzBaseView | TensorBoxrrr)rs"rrNonOwningLayout.__annotate__s1drc <VP4p\SV` VPVPVP VP 4WnR#r)r>rr/r8r7r5r6r] )rr] rrs&& rr/NonOwningLayout.__init__s?"  MM LL KK MM   rcV^8dQhRR/#rr)rs"rrr^ s..>.rc >VP4P4#r)r rrs&rrNonOwningLayout.make_indexers}}++--rcV^8dQhRR/#rsr)rs"rrr^ sPPTPrc VPP4PpV^8XdR#^RIHp\ P PPW4#)rT) ALIGNMENT) r] r>rutilsrf r{r9r:rf)rrrf s& rmaybe_guard_aligned#NonOwningLayout.maybe_guard_alignedsB%%'.. Q;$ww<4srdefault_factoryrr5c ^.#rrrrrrr 5s1#rr6cV^8dQhRR/#rr)rs"rrNoneLayout.__annotate__7scrc ^#rrrs&rrUNoneLayout.storage_size7rrcV^8dQhRR/#rjr)rs"rrr :s*rc V#rrrs&rr NoneLayout.as_fixed: rcV^8dQhRR/#rr)rs"rrr =rrc VP#rrrs&rrBNoneLayout.get_device=rrrN) rrrrrrTrUr5r6rUr rBrrrrr~ r~ )sH !'' DD)D#))+FFIFrr~ ca]tRtRtRV3Rllt]RRl4t]PRRl4tRR ltR R lt R R lt ] RRRll4t RRlt RRltRtV;t#)MutationLayoutSHOULDREMOVEiAc V^8dQhRRRR/#)rrrrrr)rs"rr'MutationLayoutSHOULDREMOVE.__annotate__Bs * *v *$ *rc <\SV`VP4VP4VP 4R4WnVP 4P4p\PPV4R#r) rr/rrAr<r get_bufferrr{r9mark_buffer_mutated)rrrrs&& rr/#MutationLayoutSHOULDREMOVE.__init__Bsc   & & (     OO      ))+ ##D)rcV^8dQhRR/#r.r)rs"rrr Ns)))rc 6VP4P#r) real_layoutr6rs&rr6!MutationLayoutSHOULDREMOVE.strideMs!(((rc V^8dQhRRRR/#)rrJrrrr)rs"rrr Rs  E d rc R#rrr~ s&&rr6r Qs rcV^8dQhRR/#rr)rs"rrr Us11d1rc >VP4P4#r)r rUrs&rrU'MutationLayoutSHOULDREMOVE.storage_sizeUs!..00rcV^8dQhRR/#)rrrr)rs"rrr Xs  F rc aRV3RlloS!VP4p\V\4'g\\ V44hV#)c V^8dQhRRRR/#)rrrrr)rs"rr;MutationLayoutSHOULDREMOVE.get_buffer..__annotate__Ys    rc<\V\4'dS!VP4#\V\4'dS!VP 44#\V\ 4'dS!VP 4#V#r)rr rrr MutableBoxr)r unwrap_viewss&rr ;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_viewsYsb&"<==#FMM22&(++#F$6$6$899&*--#FKK00Mr)rrrrr)rrr s& @rr %MutationLayoutSHOULDREMOVE.get_bufferXs>  dkk*&&)) f. . rcV^8dQhRR/#r^r)rs"rrr gsVrc |VP4Pp\V\4'g \ R4hV#)z#Expected isinstance(layout, Layout))r rrr_r)rrs& rr &MutationLayoutSHOULDREMOVE.real_layoutgs2"))&&)) !FG G rc(V^8dQhRRRRRRRR/#)rsrcrdst unsafe_aliasrrr)rs"rrr ns,***%*59* *rc  VP4\PPVP 44\ V\ 4'd VPpVP4V'g\PVP4VP4VP4\VP4VP44UUu.uF.wrE\PP P#WE4NK0 uppR7p\ V\$\&34'g \)R4hVPpVP4\+VR4'g \)V4h\ VPP,\.4'g)\)\1VPP,44h\3V4VPnVP#uuppi)rgz1Expected isinstance(node, (BaseView, MutableBox))r)rr{r9r rrrrrrDrrBrArrr<r:check_equals_and_simplifyrr rr,rr?rr )rr r r rrrs&&&& r realize_into'MutationLayoutSHOULDREMOVE.realize_intomsb  ##CLLN3 c9 % %((C ##~~'mmo*!$CLLNCLLN C CGG$$>>qD C $DdXz$:;;$G))C sF## % %#((//>:: chhoo!67 74S9xx#s4G4cV^8dQhRR/#)rrrr)rs"rrr $rc V#rrrs&rr #MutationLayoutSHOULDREMOVE.as_fixedr rcV^8dQhRR/#rr)rs"rrr rrc 6VPP4#r)rrrs&rr'MutationLayoutSHOULDREMOVE.make_indexerrr)rr)rrrrr/rYr6r rUr r r:r r rrrrs@rr r Asl * *)) ]]  1  **X**rr ca]tRtRt$R]R&R]R&RtR]R&R V3R lltR R ltR RltRRlt RRlt RRlt RRlt ] RRl4tRRltRRltRRltRR ltR!R"ltR#R$ltR%R&ltR'R(ltRHR)R*lltR+R,ltR-R.ltRHR/R0lltR1R2ltR3R4ltRIR5R6lltR7R8ltR9R:ltR;R<lt R=R>lt!]"!R4RHR?R@ll4t#RARBlt$RCRDlt%RERFlt&RGt'V;t(#)JrirOrrSrFClassVar[bool] ordering_onlycV^8dQhRR/#rr)rs"rrBuffer.__annotate__s55t5rc H<\SV`4VPRR4R#rN)rrrrs&rrBuffer.__post_init__s   t4rcV^8dQhRR/#rr)rs"rrr s00>0rc >VP4P4#r)r>rrs&rrBuffer.make_indexers --//rcV^8dQhRR/#rr)rs"rrr s#rc TVP'g \V4hVP#r)rrrs&rrBuffer.get_names yyy & &yyrcV^8dQhRR/#rr)rs"rrr src VP#rrrs&rr}Buffer.get_buffer_namerrcV^8dQhRR/#)rrztorch.Tensor | torch.SymIntr)rs"rrr s>>8>rc \VP\4'dVPP4#\ \ VP4P 4hr)rrr_r rYrrrs&rr Buffer.get_examples@ dkk6 * *;;**, ,!$t{{"3"<"<==rcV^8dQhRR/#rr)rs"rrr s33/3rc >VP4P4#r)rlrBrs&rrBBuffer.get_devices##%0022rcV^8dQhRR/#rr)rs"rrr rrc R#rrrs&rrBuffer.get_defining_oprrcV^8dQhRR/#rQr)rs"rrr s''{'rc 6VP4P#r)r>r7rs&rr7 Buffer.dtypes &&&rcV^8dQhRR/#r.r)rs"rrr s)).)rc :.VP4PO#r)r>r5rs&rr<Buffer.get_sizes("''((rcV^8dQhRR/#)rrrr)rs"rrr s++J+rc :.VP4PO#r)r>r6rs&rrnBuffer.get_strides*"))**rcV^8dQhRR/#rr)rs"rrr s((D(rc 6VP4P#r)r>rrs&r get_offsetBuffer.get_offsets '''rcV^8dQhRR/#r^r)rs"rrr s>>F>rc \VP\4'd VP#\\ VP4P 4hr)rrr_rYrrrs&rr>Buffer.get_layouts7 dkk6 * *;; !$t{{"3"<"<==rcV^8dQhRR/#rjr)rs"rrr src VP#rrrs&rrlBuffer.get_output_specrrcV^8dQhRR/#rr)rs"rrr s  3 rc "VP4#r)rrs&rr Buffer.get_storage_numels~~rcV^8dQhRR/#rsr)rs"rrr s++t+rc 6VP4P#r)r>rrs&r get_is_pinnedBuffer.get_is_pinneds ***rcV^8dQhRR/#rr)rs"rrr s11t1rc \VP\4'dD\VP\4'g"VPP 4VnR#R#R#r)rrr_r[ r rs&rrBuffer.freeze_layoutsF dkk6 * *: KK4 4 ++..0DK4 *rc$V^8dQhRRRRRR/#rr)rs"rrr s)VV"V37V Vrc \VP\4'g\\ VP44hVPP WR7VnR#r]N)rrr?rrr= rs&&&rr&Buffer.freeze_layout_with_stride_ordersB$++~66 dkk!23 3kk11%1U rc V^8dQhRRRR/#rr)rs"rrr s77=7T7rc \VP\4'g\\ VP44hVPP V4VnR#r)rrr?rrrE rs&&rr$Buffer.freeze_layout_with_fill_orders=$++~66 dkk!23 3kk//6 rc V^8dQhRRRR/#)rr6rrrr)rs"rrr s88M8d8rc \VP\4'g\\ VP44hVPP V4VnR#r)rrr?rrrI rs&&rr$Buffer.freeze_layout_with_same_orders=$++~66 dkk!23 3kk//7 rc$V^8dQhRRRRRR/#)rrrrrrrr)rs"rrr s$  * ;?  rc \VP\4'g\\ VP44hVPP WR7VnR#r )rrr?rrrA rs&&&rr'Buffer.freeze_layout_with_exact_stridessG$++~66 dkk!23 3kk22 3  rcV^8dQhRR/#rsr)rs"rrr rrc \PPP\P !VP 4^44#rrrs&rrBuffer.is_zero_elementsrrcV^8dQhRR/#rr)rs"rrr rrc aSP4'd \\SP4R7#RV3RllpV#)rTc V^8dQhRRRR/#r<r)rs"rr(Buffer.make_loader..__annotate__ s D D. DX Drc<SP4p\P!SP;'gRV!V44#rZ)rryrrrrVrs& rr^"Buffer.make_loader..loader s2'')G88DII22GENC Cr)rr rBrArsf rrBuffer.make_loaders9  " "=0@A A D D rc V^8dQhRRRR/#rr)rs"rrr s(=rc "VP4#rrrs&&rrBuffer.codegen_referencerrcV^8dQhRR/#rr)rs"rrr rrc R#rrrs&rrBuffer.decide_layoutrrcV^8dQhRR/#rCr)rs"rrr smrc \VP\4'd&VPPP 4.#R#r( )rrr[ r] rrs&rrM#Buffer.get_inputs_that_alias_outputs4 dkk? 3 3KK$$--/0 0 rcV^8dQhRR/#rCr)rs"rrr sMrc \VP\4'd&VPPP 4.#R#r( )rrr rrrs&rrEBuffer.get_mutation_namess5 dkk#= > >KK&&//12 2 rcV^8dQhRR/#rr)rs"rrr s---rc 6\VP4.4#r)rDrrs&rr Buffer.get_read_namess4==?+,,rc V^8dQhRRRR/#r#r)rs"rrr #! !rc \4#rrCr&s&&rr'Buffer.get_free_symbol_uses" |rcV^8dQhRR/#rr)rs"rrr (rrc \4#rrCrs&rrBuffer.get_unbacked_symbol_defs(rrcV^8dQhRR/#rr)rs"rrr +rA rc R#rrrs&rrBuffer.realize+rrcV^8dQhRR/#rsr)rs"rrr .src R#rrrs&rshould_allocateBuffer.should_allocate.srrrr))rrrrrr rrrr}r rBrrYr7r<rnr r>rlr r rrrrrrrrrrMrEr rer'rrr' rrrs@rrrs  $)M>) 550 > 3'')+(>  +1 V7 8  U    -H%&  rrcR]tRtRtRRltRRlt]PtRRltRt R #) OperationBufferi3cV^8dQhRR/#rr)rs"rrOperationBuffer.__annotate__6s\rc V.#rrrs&rrOperationBuffer.get_outputs6s v rcV^8dQhRR/#rrr[r)rs"rrr, 9src V#rrrs&rrOperationBuffer.get_defining_op9r rcV^8dQhRR/#rr)rs"rrr, ?s&&t&rc Z\PV4\PV4R#r)rrr[rs&rrOperationBuffer.__post_init__?sT"%rrN) rrrrrrr[rIrrrrrr* r* 3s$#55&&rr* c"]tRtRtRRltRtR#)riDcV^8dQhRR/#rr)rs"rrInputBuffer.__annotate__Errc ^#rrrs&rrInputBuffer.num_readsErrrN)rrrrrrrrrrrDs rrc]tRtRtRtRtR#) DonatedBufferiIaA 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. rN)rrrrrSrrrrr< r< Isrr< c>]tRtRt$RtR]R&RRltRRltR tR#) rhiSNrrfcV^8dQhRR/#rr)rs"rrConstantBuffer.__annotate__VsArc aRV3RllpV#)c V^8dQhRRRR/#r<r)rs"rr0ConstantBuffer.make_loader..__annotate__Ws  . X rc<SP4P4p\P!\P P SP4SP4V!V44#r) r>rryrr{r9 constant_namerrfr s& rr^*ConstantBuffer.make_loader..loaderWsPoo'446G88%%dmmot7K7KL rrrsf rrConstantBuffer.make_loaderVs   rc V^8dQhRRRR/#r>r)rs"rrr? `r6rc \\PPVP 4V4VP R7#)rr)rhr{r9rD rrrs&&rr@!ConstantBuffer.constant_to_device`s/&&t}}?   rr) rrrrrfrrr@rrrrrhrhSs+/O(/  rrhcp]tRtRtRRlt]!R4RRRll4tRRRlltR R ltR R lt R t R#)NoneAsConstantBufferifcV^8dQhRR/#rr)rs"rr!NoneAsConstantBuffer.__annotate__hr1 rc \4#rrCrs&rr NoneAsConstantBuffer.get_readshrrc V^8dQhRRRR/#r#r)rs"rrrN lr rc \4#rrCr&s&&rr')NoneAsConstantBuffer.get_free_symbol_useskr rNc V^8dQhRRRR/#rr)rs"rrrN qs--(=--rc J\PPP#r)r{r9rnone_strrs&&rr&NoneAsConstantBuffer.codegen_referenceqsww##,,,rcV^8dQhRR/#rjr)rs"rrrN ts'''rc \RR7#Nr)r~ rs&rrl$NoneAsConstantBuffer.get_output_spects &&rcV^8dQhRR/#rsr)rs"rrrN w4rc R#rrrs&rru&NoneAsConstantBuffer.has_tensor_outputwrrrrr) rrrrr rer'rrlrurrrrrL rL fs5234 -'rrL cd]tRtRt$R]R&]!R4R RRll4tR RRlltR R ltR t R#)ri{r"r]c V^8dQhRRRR/#r#r)rs"rr"ShapeAsConstantBuffer.__annotate__s::!: !:rc .\VPV4#r)r+r]r&s&&rr'*ShapeAsConstantBuffer.get_free_symbol_usess  =99rNc V^8dQhRRRR/#rr)rs"rrrb s??(=??rc h\PPPVP4#r)r{r9rcodegen_sizevarr]rs&&rr'ShapeAsConstantBuffer.codegen_references!ww##33DII>>rcV^8dQhRR/#rsr)rs"rrrb r] rc R#rrrs&rru'ShapeAsConstantBuffer.has_tensor_outputrrrrr) rrrrrrer'rrurrrrrr{s. J34:5: ?rrc0]tRtRt$RtR]R&R]R&RtR#) ExtraIndexingConstraintsiaExtra indexing constraints appended during simplify_and_reorder. Produced by the cpp fusion path to force compatible index/reduce ranges across scheduler nodes, then threaded through simplify_and_reorder and recompute_size_and_body. Replaces a bare tuple[VarRanges, list[Expr]]. rxrzlist[sympy.Expr]r3rNrrrrrSrrrrrrm rm s  rrm ca]tRtRt$RtR]R&RtR]R&RtR ]R &RtR ]R &Rt R ]R&Rt R ]R&] PRRl4t ]] PRRl44tRRltRRltRRltRRltRRlt]!R4R=RRll4tR V3R!lltR"R#ltR$R%ltR&R'ltR(R)lt]R*R+l4tR>R,R-llt]R?R.R/ll4tR0R1lt R2R3lt!R4R5lt"R6R7lt#R8R9lt$R:R;lt%R6rc# "\PpR\nRxV\nR# T\ni;i5i)TN)rrp ) old_values r force_realizeComputedBuffer.force_realizes2#11  6,0N ) ,5N )IN )sA0 A =AcV^8dQhRR/#rr)rs"rrru s  * rc VPe VP#\VPR4'dVPP#R#)z} Returns self.name if it exists, otherwise returns the name of the data node if that exists. If neither exist, returns None. Nr)rr,rrs&rget_computed_buffer_name'ComputedBuffer.get_computed_buffer_names: 99 99  499f % %99>> !rcV^8dQhRR/#rr)rs"rrru rrc 6VPP4#rrrrs&rrComputedBuffer.num_readsrrcV^8dQhRR/#rr)rs"rrru %%?%rc 6VPP4#rrr rs&rr ComputedBuffer.get_readsrrcV^8dQhRR/#rr)rs"rrru rrc 6VPP4#rrrs&rr ComputedBuffer.get_read_namesrrcV^8dQhRR/#r r)rs"rrru s!8rc \VP\\\\ 34'g2\ P!\4\4\4R7#\P!\RR4;_uu_4VPP4'dV\VP4VPP4VPP!44uuRRR4#\VP4VPP#44uuRRR4# +'giR#;i)rwrites index_exprsrTN)rrrrr rDrN ReadWritesrDr!rr?r+rWget_store_functionrr1r<rs&rr ComputedBuffer.get_read_writess$))itY%GHH** l!|&L  \\.*:D A Ayy++--*++-II002II002B A+++-II&&(B A A AsA+D8<1D88 E c V^8dQhRRRR/#r#r)rs"rrru s! !rc VPPV4VPPV4,pVP4'd&W P 4PV4,pV#r)rr'rhas_store_functionr )rrrs&& rr'#ComputedBuffer.get_free_symbol_usessd$11  II * *= 9:  " " $ $ **,AA-P PF rcV^8dQhRR/#rr)rs"rrru s % %A %rc <VP4'gkVP\PP9dBVP 4^8Xd-VP 'gVPP4#\SV`!4#r) r+rr{r9mutated_buffersrrp rrrrs&rrComputedBuffer.make_loadersf'')) !8!88 A%'''99((* *w"$$rcV^8dQhRR/#rsr)rs"rrru *sIIDIrc V\VP\\\\ 34#r)rrrrr rDrs&rr !ComputedBuffer.has_store_function*s$))itY%GHHrcV^8dQhRR/#)rrzCallable[..., None]r)rs"rrru -sGG$7Grc VP4P4P4p\VP\ \ \34'd,\VPPVPV4#\VP\4'g\\VP44h\VPPVPV4#r)r>r rrrrrr r rrrDrrr_)rrVs& rr !ComputedBuffer.get_store_function-s//#,,.;;= dii)T4!8 9 949944diiI Idii33$T$))_5549911499gF FrcV^8dQhRR/#)rrlist[int] | Noner)rs"rrru 6s(( 0(rc  \VP\4'Ed\P!VP P 4VP P44wwrpVP4Pp\;QJdRV4F 'dK RM RM !RV44'g \R4hVUUu.uFlp\V\P4'gK%\VPVUu/uF&qf^8wgK V\P P"bK( up4NKn pppV'd\VP \$\&34'dVP P)W4pMTpVUu.uF,p\*P,P.P1W4NK. p p^RIHp V !WP744#R#uupiuuppiuupi)aD 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. c3v"TF/p\V\P\P34xK1 R#5ir)rrNStarDep MemoryDeprs& rr0ComputedBuffer.get_fill_order..Es2A1|33\5K5KLMMrFTz\Expected all( isinstance(r, (dependencies.StarDep, dependencies.MemoryDep)) for r in reads )pick_loop_orderN)rrr?rNrrrr1r rrrr rvrrrrrr rr{r9r:r0 schedulerr r<) r index_varsrr&rrvrr]stride_lengthsr s & rrComputedBuffer.get_fill_order6s dkk> 2 2.:.M.M ,,. 0L0L0N/ + (Z!((*00E3333%r Aa!7!78Y 177n$WnUVPV_Q _n$WX  dii$66"ii// KG(GMR"MRTAGG$$11$@U"7&~}}GG#%X"s*"G.9G. G)G)9 G.2G4)G.cV^8dQhRR/#rr)rs"rrru `s%%t%rc \VP\4'd>VP4pV'dVP V4R#VP 4R#R#r)rrr?rrrrs& rrComputedBuffer.decide_layout`sC dkk> 2 2'')E2259""$ 3rcV^8dQhRR/#)rrzMtuple[tuple[list[Expr], list[Expr]], LoopBody, tuple[list[Expr], list[Expr]]]r)rs"rrru is J J  Jrc \P!VP4VP4RR7wr\P !\ RVP44;_uu_4\VP4VP4'dTMVR,V.VO5!pRRR4.p.p.p.pVP4FwrW^,9d8V'd \R4hVPV4VPV 4KIW^,9d \R4hVPV4VPV 4K Wg3XWE33# +'giL;i)qrrf:NrNNzExpected not reduce_varszExpected v in args[1])rNrrr1r!rrhrBrZr r+itemsrr@) rr var_rangesrzr reduce_vars index_size reduce_sizer rs & rrw %ComputedBuffer.get_default_sizes_bodyhs=(::  # # %t'>'>'@ \\.*;T__=N O O'')0022R DP !#   $$&DAG|()CDD!!!$!!!$G#()@AA""1%""1%'($0III-P Os *=E E% c$V^8dQhRRRRRR/#)rextra_indexing_constraintsExtraIndexingConstraints | Nonerecompute_sizes_body_funcrq rz5tuple[tuple[list[Expr], list[Expr]], LoopBody | None]r)rs"rrru s-O2O2$CO2$=O2 ? O2rc aaaSP4wwr4pwrgV'dV!W43WVV34wwr4pwrg.VPP4OoVe`VPpWP8wd\ VVP34hVP U u.uF qS9gK V NK p p SV , o.VP4Oo\PPS\P4'g SPVP44RVVV3Rllp Wg,p \\!S44'*;'g\"P$'*p V !VV VV 4wrpV !W|WM4wppp\&P(!VVRR7wwppp\+VV!V4V!V4.VVV4pVV3V3#uup i)a 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. c ,V^8dQhRRRRRRRRRR /#) rx_varsSequence[sympy.Symbol] support_varsr rsimplify_loopsrrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]]r)rs"rr9ComputedBuffer.simplify_and_reorder..__annotate__s:4 /4 /*4 /04 /!4 /! 4 /  4 /rc <SPWVS 4wrEpSP4R8Xdo\V4^8Xd_\\ \V444pV!V4^,^8wd.VUu.uF qV,NK pp\ V4p\ V4pV!V4pV'dE\PPPVV\S W44wrIp \Wi4p MTp WKV3#uupi)r) _apply_loop_reorderingr+rrrrrr{r9r:_simplify_loopsrRr )r r r r newsizesreindex0rrrr_prunerindex_formulas memory_addrsrs&&&& rsimplify_and_reorderAComputedBuffer.simplify_and_reorder..simplify_and_reorders,0+F+Fe\, (H*&&(E1c%jAoU3u:./E?1%*278%Qa%H8+E2H.u5Hf%F-.WW-=-=-M-M,^VN.*F *(="h. .! 9s2C5pr)rw indexing_exprsrr rrr3get_write_exprsr{r9r$rOPREFER_STORE_LOOP_ORDERextendget_read_exprsrqrTrMloop_ordering_after_fusionrNindex_vars_no_squeezerZ)rr r r r rzr r expected_var_rangesrextra_indexing_exprr r should_merge_loops iter_ranges iter_reindexr& reduce_rangesreduce_reindex iter_varsr r r sf&& @@rr #ComputedBuffer.simplify_and_reorders4  ' ' )  %Z  %Z % *)4k1J  )) 94..5578 % 1"&// "&G&GG$+2996;;#;a?V; # 1 1N0--/0 ww""4)O)OPP    3 3 5 64 /4 /l"/ t,- - V VV5V5V1V (<     ( $ 1,@ {, ( ~q 0r)rs"rrru Vs444&4rc 8VPPV4#)rerr@rs&&rr@!ComputedBuffer.constant_to_deviceVsyy++F33rrrNNr)(rrrrrSrrp rrrrrWrXr{ rVr r rr r r rer'rr r rrrdrw r r rr1r+r;r' r@rrrs@rrrs0 K%*NN*#K"481826/6<@ 9@'%'%R66 %%**,-.2 % %IG(T% J JDO2b(B(BT...,44rrcD]tRtRt$RtR]R&R]R&R]R&R]R&R tR #) FinalizeCodegenResulti[zNStructured result from TemplateBuffer._finalize_codegen for external backends.rsourcerimports call_preamble call_argsrNrn rrrr r [sX K rr c&]tRtRtRtRRltRtR#)_HasAliasingOrMutationiezMinimal view of scheduler.BaseSchedulerNode used by prologue fusion. ir.py cannot import scheduler (circular), so this documents the single method consumed by has_aliasing_or_mutation_for_prologue_fusion instead of typing the argument as Any. cV^8dQhRR/#rsr)rs"rr#_HasAliasingOrMutation.__annotate__ms 33$3rc R#rrrs&rhas_aliasing_or_mutation/_HasAliasingOrMutation.has_aliasing_or_mutationmsrrN)rrrrrSr rrrrr r es43rr ca]tRtRtRtR&RV3Rlllt]RRl4tRR ltR R lt R'R R llt RRlt RRlt RRlt R(RRlltRRltRRltRRltRRlt]RRl4t]R RR!RR"R/R#R$ll4tR%tV;t#))ripz Base class for template operators that support epilogue and prologue fusion. Subclasses: TritonTemplateBuffer (built-in Triton templates), HelionTemplateBuffer (Helion kernels), etc. Nc4V^8dQhRRRRRRRRR R R R R R/#)rrrSrrmake_kernel_renderrq mutated_inputsIterable[IRNode] | Noneallowed_prologue_inpsOrderedSet[str] | None named_inputszdict[str, IRNode] | Nonerrr)rs"rrTemplateBuffer.__annotate__wsP171717!176 17 0 17 6 17/17 17rc  <\S V`RVR7\PV4VnW0n\ PPV4Vn \ PPV4/Vn /Vn /Vn V'd \V4M/VnW@n.VnVeyVP^,p\%V\&4'g\)\+V44hVP-4pVU u.uFp \/\1VR7W4NK up VnT;'g \34VnRVnRVnR#uup i)NrI r)rr/runwrap_storagerr r{r9register_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsr mutation_outputsrrrrrBMutationOutputr~ rDr allow_epilogue_fusionallow_prologue_fusion) rrrr r r r first_inputr8rrs &&&&&&& rr/TemplateBuffer.__init__ws4 d62"11&9 "4GG++D1  ""4(.09;%?A# #/D B  -68  %++a.Kk622$T+%677 ++-F*%)Cz8#D)%D ! " 1 1Z\ " 37"26"%sEcV^8dQhRR/#rQr)rs"rrr s''{'rc \VP\4'd \R4hVP 4P #)z1Multi-output templates do not have a single dtype)rrMultiOutputLayoutrYr>r7rs&rr7TemplateBuffer.dtypes9 dkk#4 5 5%C  &&&rcV^8dQhRR/#r r)rs"rrr s88!88rc &VPRR7#)T normalize)rWrs&rr TemplateBuffer.get_read_writess''$'77rc V^8dQhRRRR/#)rr3 rrzOrderedSet[dependencies.Dep]r)rs"rrr s9Urc  aa\4pVPFp\V\\34'g\ \ V44hVo\SP\4'g\ \ SP44hSPP4oRVV3RllpV\P!VSP4RVR7P,pK V#)z(Build read dependencies from all inputs.c$V^8dQhRRRRRR/#rr Sequence[Any]rrrr)rs"rr;TemplateBuffer._read_deps_from_inputs..__annotate__s& D D] DM Dc Drc<\V4^8wd \R4h\P!SP 4S!V44#)rExpected len(rindex) == 0)rrryrr)rrr inp_indexers&&rdummy4TemplateBuffer._read_deps_from_inputs..dummys7v;!#()DEExx  E0BCCrr2 r)rDrrrrrrrr_rrNrWr<r)rr3 rinp_rawr> rr= s&& @@r_read_deps_from_inputs%TemplateBuffer._read_deps_from_inputss.8l{{Gg'@AA$T']33,3Ccjj&11$T#**%566**113K D D \55s||~rYe E#  rc V^8dQhRRRR/#)rr3 rrr r)rs"rrr s%%T%>U%rc  Haa\VP\4'dy\\P !VP 4\P!^4RRR7.4p\P!VPV4V\4RRR7#VP 4oVP4P4oRVV3Rllp\P!W0P4RVR7pV;PVPV4,unV#)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. ) var_namesr5N)rr r rr c$V^8dQhRRRRRR/#r8 r)rs"rr8TemplateBuffer.extract_read_writes..__annotate__! ; ; ;  ;# ;rcv<\V4^8wd \R4h\P!SS!V4R4#rr< fakerrryr\rrrVrs&&rr> 1TemplateBuffer.extract_read_writes..dummy26{a$%@AA99T75>6: :rr2 r)rrr. rDrNr rrr#r rA r>rrWr<r)rr3 r r> depsrVrs&& @@rrW"TemplateBuffer.extract_read_writess dkk#4 5 53= ** q)9Rb4F **11)<&L  }}//#002 ; ; // ==?B)  d11)<<  rcV^8dQhRR/#r.r)rs"rrr rrc 6\PP#r)rrrgrs&rr1!TemplateBuffer.get_reduction_sizesww{{rcV^8dQhRR/#rr)rs"rrr rQrc R#rrrs&rr+!TemplateBuffer.get_reduction_typerrcV^8dQhRR/#rsr)rs"rrr r rc R#rRrrs&rr' TemplateBuffer.should_allocaterrc$V^8dQhRRRRRR/#)rr r r rq rz9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]r)rs"rrr s)   $C  $=   C  rc *VP4.3R3#rr)rr r s&&&rr #TemplateBuffer.simplify_and_reorders$      rcV^8dQhRR/#rsr)rs"rrr s::4:rc 6\VP\4#)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrr. rs&ris_multi_outputs_template(TemplateBuffer.is_multi_outputs_templates$++'899rcV^8dQhRR/#rr)rs"rrr s**?*rc VP#r)r rs&rget_allowed_prologue_inps(TemplateBuffer.get_allowed_prologue_inps s)))rc V^8dQhRRRR/#)rscheduler_noder rrr)rs"rrr s 9 94 9  9rc "VP4#)aReturn whether this template's aliasing/mutation blocks prologue fusion. The default preserves the scheduler's conservative behavior. External template subclasses may override this when they can prove a prologue producer only feeds independent, non-mutated template inputs. )r )rrg s&&r,has_aliasing_or_mutation_for_prologue_fusion;TemplateBuffer.has_aliasing_or_mutation_for_prologue_fusions6688rc V^8dQhRRRR/#)r hook_outputszdict[str, str]rzFinalizeCodegenResult | Noner)rs"rrr s  *  % rc R#)ahCalled 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)rrl s&&r_finalize_codegen TemplateBuffer._finalize_codegensrc V^8dQhRRRR/#)rrrrrr)rs"rrr (s   f rc X\V\4'd-\VP\4'd VP#\P V4p\V\ 4'd VPp\VP\4'dVP4V#)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrr MultiOutputr+ realize_inputrrr?r)rrrs&& rrealize_template_input%TemplateBuffer.realize_template_input'sr b) $ $BGG[)I)I77N++B/ fj ) )[[F fmm^ 4 4  " rdirect_alias_at_leafon_tensor_leafon_non_tensor_leafc 0V^8dQhRRRRRRRRR R R R /#) r template_bufr structuredrrv zdict[int, IRNode] | Nonerw zFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | Nonerx zCallable[[int], None] | Nonerztuple[TensorBox, ...]r)rs"rrr 4sF)+)+$)+)+ 7 )+  )+9)+ )+rc aaaaaaa/o\P!4oRVVVVVVV3Rllo\S!V.44#)zLWalk a structured output tree, creating MultiOutput nodes for tensor leaves.c$V^8dQhRRRRRR/#)rrrrzlist[tuple[type, int]]rlist[TensorBox]r)rs"rr8TemplateBuffer.build_multi_outputs..__annotate__Bs"   *@ _ rc <\V\\34'd@.p\V4F,wr4VP S!V.VO\ V4V3N44K. V#\ S 4p\V\P4'dS 'd%VS 9d\PS V,4.#\V4pVS 9d S V,.#\\PV4SV4pVSPVP!4&S eS !VP!4WqV4\V4pVS V&V.#S e S !V4.#r)rrrrr rrrDTensorrridrr FallbackKerneltensor_to_layoutr$ r)rrrMritemleaf_idxtidmorrv leaf_counterrx rw seen_outputsrz walks&& rr 0TemplateBuffer.build_multi_outputs..walkBs0&4-00+-(0GANN4.K.K$v,9J.K#LM 1L)H&%,,//'H8L,L%,,-A(-KLMMj,&(-.. "33F;\7FH 33BKKMB!-"2;;="xHr]$& S!t !-"8,Ir)rPcountr) rrz r{ rv rw rx r r r s &f&ddd@@@rbuild_multi_outputs"TemplateBuffer.build_multi_outputs3s4.0  (   6T*b)**r) r$ r% r( r) r rr# rr r r& r)NNNrr )rrrrrSr/rYr7r rA rWr1r+r' r r` rd ri rn r:rt r rrrs@rrrps 1717f''8*%N  :* 9   )+ :> )+)+<@)+)+rrcva]tRtRtR RV3Rlllt]!R4R RV3Rlll4tRRltRR ltR t V;t #) TritonTemplateBufferi`c 0V^8dQhRRRRRRRRR R R R /#) rrr_rrr zCallable[_P, _T] | Noner r r r rrr)rs"rr!TritonTemplateBuffer.__annotate__asF>>>!>4 > 0 > 6 > >rc <\SV`VVVVVR7VPf \R4hVPVP/VnRVnRVnR#)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. )r r NzExpected self.name is not None)rr/rrr# subgraph_inps subgraph_outs)rrrr r r rs&&&&&&rr/TritonTemplateBuffer.__init__asc"    )"7   99  !AB B)-DII(>%@D9=rc V^8dQhRRRR/#r#r)rs"rrr s! !rc J<\SV`V4pVP'd VPM.pVP'd VPM.pVF|p\ V\ 4'dVP \WQ44K5\ V\4'd#VP VPV44KmVfKs\R4h VFJp\ V\4'd#VP VPV44K;VfKA\R4h V#)NzExpected inp is NoneExpected out is None) rr'r r rr"updater+rr)rrresr r rr'rs&& rr')TritonTemplateBuffer.get_free_symbol_usessg*=9.2.@.@.@**b .2.@.@.@**b  C#t$$ +C?@C(( 333MBC?()?@@!!C#v&& 333MBC?()?@@ ! rcV^8dQhRR/#rr)rs"rrr ..\.rc V.VPO#rr& rs&rr TritonTemplateBuffer.get_outputs-t,,--rcV^8dQhRR/#rr)rs"rrr src &RVP R2pV#)zTritonTemplateBuffer(layout=rRr)rr's& rrTritonTemplateBuffer.__str__s,T[[M; r)r# r r r r) rrrrr/rer'rrrrrs@rr r `s;>>>2342.rr ca]tRtRtRtRV3RlltRRltRRltR R ltR R lt R Rlt RRlt RRlt RRlt RRltRtV;t#) ChoiceCalleria 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. c ,V^8dQhRRRRRRRRRR /#) rrrrrrr_ descriptionrrr)rs"rrChoiceCaller.__annotate__s<111"1 1  1  1rc <\SV`4WnW0nW nW@nRVnRVn/VnRVn RVn /Vn /Vn R#FN) rr/rrrr failed_benchmark_with_cudagraphsrr decompositiondecomposition_kwargsconfig_patches)rrrrr rs&&&&&rr/ChoiceCaller.__init__s`   &'! 05'/1/38<79!.0rc$V^8dQhRRRRRR/#)rrrr'r0rrAr)rs"rrr s&LLsLL%Lrc aaVP4oVP'd\P!VV3Rl4#\P 'd\ VV3Rl4#\P!SSRV/RR7#)c<S!S!#rralgorsrr(ChoiceCaller.benchmark..s T4[rc<S!S!#rrr srrr s D$Krr'Nr) to_callabler r_benchmark_gpu_with_cuda_graphrM/profile_bandwidth_with_do_bench_using_profilingrj benchmark)rr'rr s&$j@rr ChoiceCaller.benchmarksb!  * * *<<=PQ Q  A A A+,?@ @$$T4%dKKrcV^8dQhRR/#rr)rs"rrr s""3"rc \hrrers&r call_nameChoiceCaller.call_namergrcV^8dQhRR/#)rrrr)rs"rrr s""/"rc \hrrers&rr ChoiceCaller.to_callablergrcV^8dQhRR/#rr)rs"rrr src "VP4#)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_keyChoiceCaller.kernel_hash_keys }}rcV^8dQhRR/#rr)rs"rrr s""#"rc \hrrers&rr ChoiceCaller.hash_keyrgrcV^8dQhRR/#)rrrr)rs"rrr s""Y"rc \hrrers&rrChoiceCaller.output_nodergrcV^8dQhRR/#)rrz6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]r)rs"rrr sQrc /#)zRInformation returned here is logged to the autotune log file when that is enabled.rrs&r info_dictChoiceCaller.info_dictr&rcV^8dQhRR/#rr)rs"rrr s$$#$rc R#)unsupported_choicerrs&rautoheuristic_idChoiceCaller.autoheuristic_ids#rcV^8dQhRR/#rr)rs"rrr sTrc RVnR#)zp Mark the choice as failed so that it can be removed later. Useful for when we decouple compilation and tuning. TN)r rs&r mark_failedChoiceCaller.mark_faileds  r) r rr r r r r rrrr)rrrrrSr/r r r r r rr r r rrrs@rr r sE116L""""$rr c"]tRtRtRRltRtR#)TritonTemplateCallerBaseicV^8dQhRR/#)rrrr)rs"rr%TritonTemplateCallerBase.__annotate__s"""rc \hrrers&rget_make_kernel_render/TritonTemplateCallerBase.get_make_kernel_renderrgrrN)rrrrr rrrrr r s ""rr z type | None_NVUniversalGemmCallerClassc V^8dQhRRRR/#)rchoicer rrr)rs"rrrs ; ;| ; ;rch\f ^RIHpVs\ V\4# \dR#i;i)NNVUniversalGemmCallerF)r )torch._inductor.codegen.nv_universal_gemmr ImportErrorr)r r s& r"_is_output_plannable_nvgemm_choicer s:"*  W*? ' f9 ::  s " 11ca]tRtRtRtRV3Rllt]RRl4t]RRl4tRR R llt ] PR R l4t R Rlt ] PRRl4tRRltRRRlltRRltRtV;t#)MultiTemplateBufferia3 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. c 0V^8dQhRRRRRRRRR R R R /#) rrr_rrchoice_timings_fnz1Callable[[int | None], dict[ChoiceCaller, float]]unfiltered_choiceslist[ChoiceCaller]r rrrr)rs"rr MultiTemplateBuffer.__annotate__sG888!8M 8 / 8 / 8 8rc <\SV`VVRVR7W0n/VnW@nW n\ ;QJdRV4F 'dK RM RM !RV44Vn/VnRVn RVn R#)N)rrr r c3"TFrp\V\4;'gV\V4;'gC\V\PP P 4;'d VPxKt R#5ir)rr r rDr.select_algorithmExternKernelCallerhas_out_variant)rr s& rr/MultiTemplateBuffer.__init__..#ss% - v7 8  1&9  65??#C#C#V#VW++**   -sA<A<5A<'A<FT) rr/_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders _render_kind_render_caller)rrrr r r rs&&&&&&rr/MultiTemplateBuffer.__init__s #"7  #4LN,> %!$% -% % -% " <>!(,48rcV^8dQhRR/#rsr)rs"rrr 3s&&$&rc VP#)zgTrue when all choices are TritonTemplates, NVUniversalGemmCallers, or Extern Kernels with out variants.)r rs&routput_plannable$MultiTemplateBuffer.output_plannable2s%%%rcV^8dQhRR/#)rrr r)rs"rrr 8s+rc VP#r)r rs&rr&MultiTemplateBuffer.choices7s }}rc V^8dQhRRRR/#)r hint_overriderrzdict[ChoiceCaller, float]r)rs"rrr ;s33'3 "3rc WP9dVPV4VPV&VPV,#r)r r )rr s&&rchoice_timings"MultiTemplateBuffer.choice_timings;s<  4 4 4262I2I-2XD  /##M22rc V^8dQhRRRR/#)rcallerr rrt r)rs"rrr Cs..,D..rc# "\V\PPP4'g\ \ V44hVPVP8wd \ R4hVPpVPpVPpVP4VnRVn Wn RxW nW0n W@n R# Y nY0n Y@n i;i5i)%Expected self.layout == caller.layoutrZN) rrDr.r TritonTemplateCallerrrrr r r r )rr render prev_kind prev_callers&& rswap_as_triton_caller)MultiTemplateBuffer.swap_as_triton_callerBs EOO44II  !f. . ;;&-- ' !HI I((%% )) "("?"?"A$$ . &, # ) "- '- # ) "- sB5C'8C<C'C$$C'c V^8dQhRRRR/#)rr r rrr)rs"rrr Xs % %0H %T %rc \V\PPP4'g\ \ V44hVP4VPP8wd \ R4hVP4VPP8wd \ R4hVP4Vn RVnWnR#).Expected self.get_size() == caller.layout.size2Expected self.get_stride() == caller.layout.striderZN)rrDr.r r rrr<rr5rnr6r r r r )rr s&&rfinalize_as_triton_caller-MultiTemplateBuffer.finalize_as_triton_callerXs EOO44II  !f. . ==?fmm00 0 !QR R ??  4 4 4 !UV V"("?"?"A$$rc V^8dQhRRRR/#)rr r rrt r)rs"rrr fs..L.^.rc# "^RIHp\W4'g\R\ V4 24hVP VP 8wd \R4hVP pVPpVPpVP4VnRVnWnRxW0nW@nWPnR# Y0nY@nYPni;i5i)rr $expected NVUniversalGemmCaller, got r nvgemmN) r r rrrrr r r r )rr r r r r s&& rswap_as_nvgemm_caller)MultiTemplateBuffer.swap_as_nvgemm_calleresS&88 #GV ~!VW W ;;&-- ' !HI I((%% )) "("?"?"A$$ . &, # ) "- '- # ) "- sBC B6"C 6C  C c V^8dQhRRRR/#)rr r rrr)rs"rrr {s % % % %rc v^RIHp\W4'g\R\ V4 24hVP 4VP P8wd \R4hVP4VP P8wd \R4hVP4Vn RVn Wn R#)rr r# r r r$ N)r r rrrr<rr5rnr6r r r r )rr r s&& rfinalize_as_nvgemm_caller-MultiTemplateBuffer.finalize_as_nvgemm_caller{sS&88 #GV ~!VW W ==?fmm00 0 !QR R ??  4 4 4 !UV V"("?"?"A$$rc V^8dQhRRRR/#)rr rrztuple[ChoiceCaller, float]r)rs"rrr s11'1 #1rc fVPVR7p\W"PR7pW2V,3#))r r)r r{r1)rr timings min_choices&& rget_min_choice"MultiTemplateBuffer.get_min_choices3%%M%Bkk2 J/00rc V^8dQhRRRR/#)rcallersz*dict[int | None, TritonTemplateCallerBase]rrr)rs"rrr s , ,A ,  ,rc VP4F"wr#VP4VPV&K$ VPR,VnRVnVR,VnR#)z;Finalize with multiple callers for different hint overridesNrZ)r r r r r r )rr2 r r s&& rfinalize_as_triton_callers.MultiTemplateBuffer.finalize_as_triton_callerssY&-]]_ !M7=7T7T7VD % %m 4&5#'";";D"A$%dmr) r r r r r r r r r r)rrrrrSr/rYr r&r rWrXr r r% r) r/ r4 rrrs@rr r s88B&&3..* %..* %1 , ,rr cFa]tRtRtRV3RlltRRltRRltRtV;t#) CUTLASSTemplateBufferic4V^8dQhRRRRRRRRR R R R R R/#)rrr_rrr Callable[_P, _T]workspace_sizertemplatersupports_epilogue_fusionrrrr)rs"rr"CUTLASSTemplateBuffer.__annotate__sY A A A! A- A  A " A#' A  Arc L<\SV`WV4W@nWPnW`nR#r)rr/r: r; r< )rrrr r: r; r< rs&&&&&&&rr/CUTLASSTemplateBuffer.__init__s& );<, (@%rcV^8dQhRR/#rr)rs"rrr= sMMCMrc :VPe VP#^#rr: rs&rr(CUTLASSTemplateBuffer.get_workspace_sizes&*&9&9&Et""L1LrcV^8dQhRR/#rr)rs"rrr= s55$5rc VP4F)p\P!VP4RR4K+ R#r)rryr\r)rrs& remulate_store_fn&CUTLASSTemplateBuffer.emulate_store_fns,&&(F IIfoo't 4)r)r< r; r: ) rrrrr/rrF rrrs@rr7 r7 s A AM55rr7 c@a]tRtRtRV3RlltRV3RlltRtV;t#)CppTemplateBufferic 0V^8dQhRRRRRRRRR R R R /#) rrr_rrr r9 r; rr rrrr)rs"rrCppTemplateBuffer.__annotate__sF 1 1 1! 1- 1 " 1  1  1rc N<\SV`WV4W@nWPnRVnR#r)rr/r; r outputs)rrrr r; r rs&&&&&&rr/CppTemplateBuffer.__init__s& );<  ,0 rcV^8dQhRR/#r^r)rs"rrrK s ( (F (rc <\VP\4'd\VP\4'g\ \ VP44hVP^,p\V\4'g\ \ V44hVPp\V\4'g\ \ V44hV#\SV`)4#r) rrr. rM rrrrr_rr>)r first_outputrrs& rr>CppTemplateBuffer.get_layouts dkk#4 5 5dllH55$T$,,%788<<?LlF33$T,%788!((Fff--$T&\22M7%' 'r)r rM r; )rrrrr/r>rrrs@rrI rI s 1 1 ( (rrI cBa]tRtRtRtRRV3RllltRRltRtV;t#) CuteDSLTemplateBufferiz Buffer for CuteDSL (CUTLASS Python DSL) template kernels. Similar to other template buffers but specialized for CuteDSL operations. c 0V^8dQhRRRRRRRRR R R R /#) rrr_rrr r9 r; rr r rrr)rs"rr"CuteDSLTemplateBuffer.__annotate__sF!-    0  rc  <\SV`WV4W@nWPnV.VnVe\ VP ^,\4'g&\\VP ^,44hVP ^,P4pT;PVUu.uFp\\VR7Wp4NK up, unR#R#uupirZ ) rr/r; r rM rrrrrrBr' r~ ) rrrr r; r r8rrs &&&&&& rr/CuteDSLTemplateBuffer.__init__s );< ,&*V  %dkk!nf55$T$++a.%9::[[^..0F LL))Cz8#D) L &s'CcV^8dQhRR/#rr)rs"rrrV \rc VP#rrM rs&rr!CuteDSLTemplateBuffer.get_outputs ||r)r rM r; r) rrrrrSr/rrrrs@rrT rT s ,rrT c^a]tRtRtRtR RV3RllltRRltRRltR R R lltR t V;t #)NVUniversalGemmBufferiz Buffer for NVIDIA Universal GEMM kernels. Unlike CuteDSL templates which use Jinja templates, this generates simpler Python code that directly calls the cutlass.operators library. cPV^8dQhRRRRRRRRRRR R R R R R RR RR RRRRRRRR/#)rrr_rrkernelraccumulator_typevariantr: r scale_type_az Any | None scale_type_bswizzle_type_aswizzle_type_br< rswap_ab bias_nodez Buffer | Nonerrr)rs"rr"NVUniversalGemmBuffer.__annotate__s';';';!'; ';  ';  ';';!';!';#';#';#'';';!'; ';rc 2<\SV`WRR7W0nW@nV.VnW`nWPnWpnWnWn Wn Wn Wn Wn RVPPRVP /VnVP$VnR#)N)r kernel_namemin_cc)rr/rb rc rM r: rd re rf rg rh r< ri rj metadata operator_namedesigned_for_min_cckernel_metadata_make_kernel_renderr )rrrrb rc rd r: re rf rg rh r< ri rj rs&&&&&&&&&&&&&&rr/NVUniversalGemmBuffer.__init__s" DA 0&*V , ((,,(@% # 6??88 f00  #'":":rcV^8dQhRR/#rr)rs"rrrk #s##C#rc VP#)z#Return the workspace size in bytes.rB rs&rr(NVUniversalGemmBuffer.get_workspace_size#s"""rcV^8dQhRR/#rr)rs"rrrk 'rZ rc VP#rr\ rs&rr!NVUniversalGemmBuffer.get_outputs'r^ rc8V^8dQhRRRRRRRRR RR R R R RR/#)rout_noderr repilogue_fn_coderOepilogue_readsrepilogue_writesepilogue_var_renamesdict[str, Any] | None local_reducezGemmReductionPlan | Nonerztuple[Any, Any]r)rs"rrrk *sdC%C%C%"C%% C% ) C% * C%4C%/C% C%rc aVeVf \R4h^RIHp^RIHp .p VP FXp \ V \4'd V Pp \ V \4'd V Pp V PV 4KZ Rp VPeV R,p V RRp \V P4p V!R/RV bRV bRVbRVPbR VPbR VP bR VP"bR VP$bR VP&bRVP(bRVP*bRVbRVbRVbRVbRVbRVP,bRV boV3RlpSV3#)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 Nz.epilogue_fn_code requires epilogue_var_renames)NVUniversalGemmKernel) Placeholderrm rrrr rc r: rd re rf rg rh r} r~ r r r ri rj c$<SP4#r)r ) render_kernelsrr 9NVUniversalGemmBuffer._make_kernel_render..renderjs '') )rrlr)rBtorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelr torch._inductor.utilsr rrrrrr@rj r KERNEL_NAMErr rc r: rd re rf rg rh ri )rr| r r} r~ r r r r r rrrj rm r r s&&&&&&&& @rrs )NVUniversalGemmBuffer._make_kernel_render*s"  '#+$%UVV  6!# ;;C#y))hh#z**hh   s #  >> %#BI%cr*K+112 - # # ! !00  "22   ..  LL ** **  ..  .. . * , "6 &! "LL# $ % * *f$$r)rc rj rb rr r rM re rf r< ri rg rh rd r: )rNNNNFFN)NNNNNN) rrrrrSr/rrrs rrrs@rr` r` s)';';R#C%C%rr` c V^8dQhRRRR/#)rr#Sequence[IRNode | Sequence[IRNode]]rzTypeIs[Sequence[IRNode]]r)rs"rrrps55 .55rcj\;QJdRV4F 'dK R# R#!RV44#)c3B"TFp\V\4xK R#5irrr)rrs& rr#is_node_sequence..ss4ez!V$$err FT)r)rs&ris_node_sequencer ps- 34e43343434e4 44rc]tRtRt$R]R&RRltRRltRR lt]R R l4t ] R R l4t RRlt RRlt ]!R4RRRll4tRtR#)rivr rc V^8dQhRRRR/#)rrrrrr)rs"rrInputsKernel.__annotate__zs  C C rc VPV,p\V\4'g \R4hVP 4#)z"Expected isinstance(input, IRNode)rrrrr)rrinputs&& r input_nameInputsKernel.input_namezs5 A%(( !EF F~~rcV^8dQhRR/#r r)rs"rrr s  !8 rc a\\P,!4p\PoVPFsp\ V\ 4'dVPV3RlV44K6\ V\4'dKNVPS!VP444Ku \\P,!V3RlVP444p\P!VV\4R7#)c3P<"TFpS!VP44xK R#5irr )rrr s& rr/InputsKernel.get_read_writes..sBEqWQZZ\22E#&c3P<"TFpS!VP44xK R#5irr )rrr s& rrr s#. /AGCLLN # #/Ar r ) rDrNrTr rrrr rr-rrr )rrr r r s& @rr InputsKernel.get_read_writess<++,.&&[[E%** BEBBE#899 '%.."234!L,,-. /3/?/?/A.  &&"   rcV^8dQhRR/#rr)rs"rrr rrc 6VP4P#rrrs&rr InputsKernel.get_readsrrc V^8dQhRRRR/#rNr)rs"rrr sFrc \V\4'd VPp\V\4'd VPp\V\4'd,\V\ 4'g\ PV4p\V\4'dVPV4#\V\4'dV#\V\\ 34'g\\V44hV#r) rrrrrrr+rs unwrap_storage_for_inputTorchBindObjectrrrrrs&&rr %InputsKernel.unwrap_storage_for_inputs a # #A a $ $A a " ":a+I+I**1-A a # # //2 2 a ) )H!fo677 a) )rc V^8dQhRRRR/#)rrr rzlist[IRNode | Sequence[IRNode]]r)rs"rrr s  3  ( rc .pVFdp\V\4'd&VUu.uFp\PV4NK ppM\PV4pVP V4Kf V#uupir)rrrr r@)r inputs_newrrs& rr InputsKernel.unwrap_storagesm79 A!X&&GHIq!\::1=qI 99!<   a   JsA/cV^8dQhRR/#rsr)rs"rrr s4rc R#rRrrs&rr6InputsKernel.is_externrrcV^8dQhRR/#rr)rs"rrr rrc ^#rrrs&rrInputsKernel.num_readsrrc V^8dQhRRRR/#r#r)rs"rrr s  !  ! rc \\P,!4pVPFRp\ V\ 4'dW#P V4,pK2VFpW$P V4,pK KT V#r)rDrr$rrrr')rrrr inner_inps&& rr'!InputsKernel.get_free_symbol_usessg u|| $ &;;C#v&&--m<<!$I77 FFA"%  rrNr)rrrrrr r r r:r rVr r6rrer'rrrrrrvsb //  ,,&  N+ , rrc.]tRtRtRRltRRltRtR#) NopKernelicV^8dQhRR/#rsr)rs"rrNopKernel.__annotate__r rc R#rRrrs&rr;NopKernel.is_no_oprrcV^8dQhRR/#rr)rs"rrr r1 rc \4#rrCrs&rr NopKernel.get_readsrrrN)rrrrr;r rrrrr r srr c]tRtRtRt]RRl4t]RRRll4t]!R4RRR ll4t ]R R l4t R R lt Rt R#) ConcatKernelizb There isn't actually a real kernel for concat, we just change the storage for the upstream data. c$V^8dQhRRRRRR/#)rrrrrrrr)rs"rrConcatKernel.__annotate__s'AA,A3A:Arc  V^,P4pV^,P4p\V^,P44p^.pWR,.p^Tu;8:d\ V48gM\ R4h\ ^\ V44EF*pW,P4p VPWR,4\ V 4\ V48wd \ R4hW,P4V8wd \ R4hW,P4V8wd \ R4h\ \ V44FZp W8XdWZ,W,,WZ&K"\PPPWZ,W,4WZ&K\ VPWR,4EK- \PV4p \P'd(\ P#WV^,P$4p \ \ V44FpW,p \'V 4'gKV P)4p \+V \,4'gKF\ P/V P0V P24'gKx\5V4p M \6;QJdRV4F 'gK RM RM !RV44p\PP8P:^,pVRJd\\+V\4'dF\6;QJdRV4F 'gK RM RM !RV44'd \5V4p \<;QJdR V4F 'dK RM RM !R V44pVf \ R 4h\?R \-VVVV VR 7.R 7p\AV4p.p\CV4EFwpp\+V\D\F34'g\ \IV44hVPKV\LPOVW&V,Wx,RR74p\+V\P4'g\ \IV44h\+VPR\4'g\ \IVPR44hVPRPV4\+VPT\D4'dVPTPW4pM VPTp\+V\@4'gEKPVPY4'gEKiVP4;pfEK\[VPH4'gEK\]V4'dEKVPVP_44EK \ V4^8dT\PPaV\bPd4'd \PPgV4\PPiV4Vn5VPmVPR4Vn)\PPoV4V#)z& Create the concat kernel from inputs z!Expected 0 <= dim < len(new_size)z)Expected len(input_size) == len(new_size)z'Expected inputs[i].get_dtype() == dtypez)Expected inputs[i].get_device() == devicec38"TFp\V4xK R#5ir)r=r s& rr&ConcatKernel.create.. s-WPV1.CA.F.FPVrTFc3"TF|pRVP9;'da\VPR,\PR7;'g-\VPR,\PR7xK~ R#5i)rfr N)rr3rDr r rargs& rrr sx(C!  <&+&9&9   @&+&<&< (sB3B5Bc3|"TF2p\V4;'dVP4PxK4 R#5ir)r=r>rr s& rrr 's0 KQa !! $ A A)A)A A6s<#<Nr)r8r7r5r6rrrrr )8rBrArr<rrrr@r{r9r:r r?r@rMr r_rNr7r=r>rrr r5r6r5r current_noderrr rrrr rr rrrrrrrrqrprIr$rOFOREACHregister_operation_listr! rr r" )rrrr8r7r offsets_start offsets_endr input_sizerB output_striderrany_input_is_storage_and_layout fx_node_argsr concat_kernelrb op_namesrrm input_unwrappeddevs&&& rrConcatKernel.creates %%'q ##%q **,- }o S(3x=( !DE Eq#f+&A++-J   /:#h-/$%PQQy""$-$%NOOy##%/$%PQQ3x=)8"*+ "=HK"#''"2"2"L"L Z]#HK *   x} -!'$(6'H'H'R  ' ' '"//M s6{#A A$Q''K88fmmTT$B8$LM$+.#-WPV-W###-WPV-W*W'ww++003  ,u 4<..(($;8DMC KQ CCC KQ   > !>? ?$$#   M*'FAscHj#9::$T#Y//++  Cq!1;>!L lF33$T,%788m22D99$T-*>*>%?@@  ' ' 5#((H--"%(("6"6"8"%((?J77#3355NN,,S9388$$"<00 ? ? AB7(: x=1 !4!4V^=S=S!T!T GG + +H 5WW44]C "11-2F2FG  ""=1 rNc$V^8dQhRRRRRR/#)rr rr r3rrr)rs"rrr as$" " " ," " rc \V\4'dVPVPV4#\V\\ 34'g\ \V44h\VP\4'Ed\VPP\4'dVPP'gR#VfR#\VP44\VP448wdR#\;QJdBR\VP4VP444F 'dK R# R#!R\VP4VP4444#\!VPR4;'dL\VPP\"4;'d \VP\$4'*#)FTc3x"TF0wr\PPPW4xK2 R#5irr~rs& rr=ConcatKernel.can_realize_into_without_copy..zs1EFB  88@@Err)rrcan_realize_into_without_copyrrrrrr rrr rrnrrr,r?ExternKernelAlloc)rr r s&&&rr *ConcatKernel.can_realize_into_without_copy`sd c9 % %44SXXsC C#*566 c+ + chh 3 4 4sxx <<xx000{3>>#$CNN,<(==3!#.."2CNN4DE33 3 3!#.."2CNN4DE  CHHh ' < <388??N; < <sxx):;; rc V^8dQhRRRR/#r#r)rs"rrr s CC!C !Crc ,\PW4#r)r r'r&s&&rr'!ConcatKernel.get_free_symbol_usess--dBBrc$V^8dQhRRRRRR/#)rr rr rr)rs"rrr s!!)!)v!)F!)v!)rc  \V\4'g+\V4'd\V4wr4\W4R7p\V\4'g\ \ V44h\V\ 4'dVPVPV4#\V\4'duVP4\VPR4'g \ R4hVPW4'd'\V4VPnVP#\P!VP#4VP%4VP'4\)VP+4VP+44UUu.uF.wrV\,P.P0P3WV4NK0 uppR7pVPWr4#uuppi)rrz$Expected hasattr(src.data, "layout")rg)rrr=rrrrr rrrr,r r[ rrDrrBrArrr<r{r9r:r )rr r rrrrpws&&& rr ConcatKernel.realize_intos` #//$S))"7"<%7B#// c+ + c9 % %##CHHc2 2 c: & & KKM388X..$%KLL00::"1#"6xx   >>#--/__&   ??DA  ::1@?  (( s4GcV^8dQhRR/#rsr)rs"rrr r rc R#rRrrs&rr' ConcatKernel.should_allocaterrrrr) rrrrrSr:rr rer'r r' rrrrr r sm AAF" " HN+C,C !)!)Frr c V^8dQhRRRR/#)rr r@rrt r)rs"rrrs;;X;.;rc#"VPR4pRxVPV4R# TPT4i;i5ir )"_ignore_fresh_unbacked_symbols_set)r prevs& r_track_fresh_unbacked_symbolsr s9  7 7 >D; 44T: 44T:sA,A?Ac V^8dQhRRRR/#)rr rrzAbstractContextManager[None]r)rs"rrrs!rc`Ve"VP4'd \V4#\4#r)"_ignore_fresh_unbacked_symbols_tlsr r )r s&r(_fallback_kernel_symbol_tracking_contextr s+!M!M!O!O,Y77 =rcN]tRtRt$RtR]R&R]R&R]R&R ]R &R ]R &R tR#)ProcessKernelResultizStructured result of ExternKernel.process_kernel. Replaces the positional 5-tuple that was unpacked at every call site. unflatten_args(new_tensor_args, new_non_tensor_args) reconstructs the original (args, kwargs) tree from replacement lists. rexample_outputr tensor_argsz list[object]non_tensor_argszCallable[[Any, Any], Any]unflatten_args)dict[sympy.Symbol, pytree.KeyPath] | Noneunbacked_bindingsrNrn rrrr r s'!!--@@rr ca]tRtRt$RtRmtR]R&]P!] R7t R]R&Rt R ]R &Rt R ]R &Rt R ]R &]P!]R7tR]R&RtR]R&RtR]R&]P!] R7tR]R&RtR]R&]P!] R7tR]R&]P!]R7tR]R&RnRV3RllltRRltR R!ltR"V3R#lltR$R%ltR&R'ltRoR(R)lltR*R+ltRoR,R-lltR.R/ltR0R1lt R2R3lt!]"R4R5l4t#]$R6R7l4t%]$R8R9l4t&]$R:R;l4t']$R<R=l4t(]$RpR>R?ll4t)]$RqR@RAll4t*]$RqRBRCll4t+]$RDREl4t,]$RFRGl4t-]$RHRIl4t.]$RJRKl4t/RLRMlt0RNROlt1RoRPRQllt2RRRSlt3RTRUlt4RqRVRWllt5RXRYlt6RZR[lt7R\R]lt8R^R_lt9R`Ralt:RbRclt;RdRelt!R4RqRhRill4t?RjRklt@]@tARltBV;tC#)rr+iz 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. r9 constant_argsr rrNReinterpretView | None output_viewrOpython_kernel_namecpp_kernel_name Iterable[str]ordered_kwargs_for_cpp_kernel_OpOverloads | None op_overloadzlist[ArgProperty] | Nonearg_propertieszdict[str, ArgProperty]allarg_propertieszdict[str, ArgProperty] | Nonekwarg_propertiesz"dict[sympy.Symbol, pytree.KeyPath]rzlist[MutationOutput]r& cDV^8dQhRRRRRRRRR R R R R RRRRRRRRR/ #)rrrOrrSrr rr9 rr rrrrrrr r rrr)rs"rrExternKernel.__annotate__sx11114 1 % 1 & 1,1'1$1(51)1 1rc  J<\S V`VVVR7W@nV'dTM/VnW`nWnVP V4VPV4WnVP4/Vn .Vn \PPVn/VnR#)r N)rr/rrrr set_cpp_kernel_nameset_python_kernel_namercollect_arg_kwarg_propertiesrr& r{r9r fx_noder) rrrrrrrrrrr rs &&&&&&&&&&&rr/ExternKernel.__init__s   + &fB &&   1 ##$67-J* ))+!# "ww++ .0rcV^8dQhRR/#rr)rs"rrrr rc V.VPO#rr rs&rrExternKernel.get_outputsr rcV^8dQhRR/#rr)rs"rrrrrc \4#rrCrs&rr%ExternKernel.get_unbacked_symbol_defsrrcV^8dQhRR/#r r)rs"rrrs!8rc <a\SV`4oRV3Rllp\P!VVPVP 3RR7S#)c V^8dQhRRRR/#)rrJrrrr)rs"rr2ExternKernel.get_read_writes..__annotate__s F Fs Ft Frc<\V\4'dHVP4pVe2SPP \ P !V44R#R#R#r)rrrrr-rNr )rJrrs& r add_ir_read1ExternKernel.get_read_writes..add_ir_readsM%((++-#%%)),*>*>t*DE$)rc"\V\4#rr rLs&rr.ExternKernel.get_read_writes.. s *UF";r)is_leaf)rr pytree tree_map_rr)rr!rrs& @rr ExternKernel.get_read_writessLg-/  F F      -; rcV^8dQhRR/#rr)rs"rrr$s#$#$d#$rc  \VP\PP4'diVPP P Uu.uF=pVP'dKRVPRVPRVP/NK? upM-\\VP44Uu.uFp/NK upVn\VP\PP4'dTVPP P Uu/uF(pVPRVPRVP/bK* upM/Vn\VP\PP4'dVP 'gSVPP P Uu.uF"qP'gKVPNK$ upVnVPP P Uu.uFqP'gKVNK upVnR#.VnR#uupiuupiuupiuupiuupi)rrrN)rr rD_ops OpOverload_schema arguments kwarg_onlyr real_typerrrrr r r schema_kwargs)rrrs& rr)ExternKernel.collect_arg_kwarg_properties$s$**EJJ,A,ABB))11;;    >555$($4$4$<$<$F$F6$Fq,,FAFF$F62 ++33=="=a="D "$D ? 7 6"s0I 0*I = I&.IIII0IcV^8dQhRR/#rr)rs"rrrIs!!t!rc \VP\4'd#VP4VP 4R#R#r)rrr?apply_constraintrrs&rrExternKernel.decide_layoutIs0 dkk> 2 2  ! ! #     3rc$V^8dQhRRRRRR/#)rwrapperrrm rOrrr)rs"rrrNs)MM+M:DM Mrc \W4wr4V'dVPV4V'gVP4pV'd$^RIHpV!WRR7pVP W&4R#R#)r)'set_kernel_post_grad_provenance_tracingT)r6N)rm make_commenttry_get_kernel_namer/r:write_provenance_debug_handle)rr8rm origin_str_detailed_origin_strr: debug_handles&&& rcodegen_commentExternKernel.codegen_commentNs],?t+M(    ,224K  FBTL  1 1+ L rc V^8dQhRRRR/#rr8rrrr)rs"rrr_s""3""rc \hrrerr8s&&rcodegenExternKernel.codegen_rgrc V^8dQhRRRR/#rrrOrrr)rs"rrrbs#;#;:#;#;rc Wn\PP'd5\ VP \ PP4'gR#VP pVPfVPR8XdVPR8Xd#VPPR4^,MVPPRR4p^RIHpVP!RVP 2/4p\#RV4^R7pV^8dV R V 2pR V R 2VnR#VP$P&VnR#R#) Natenr|.r&inductor_fallback_opszaten.c3v"TF/qPR4'gK\VR,4xK1 R#5i)r r N) startswithr)rr s& rr3ExternKernel.set_cpp_kernel_name..~s%KAc9JZS2ZZs99r_vz at::_ops::z::call)rr{r9 cpp_wrapperrr rDr+r, namespace _overloadnamerrAreplacetorchgen.aoti.fallback_opsrOr1rr-r)rrrb opnamerO version_infolatest_versions&& rr ExternKernel.set_cpp_kernel_namebs".ww"""*   ejj33+ +  !!    '6)++y8OO))#.q100c:M4885@Q9RTVW !$KK""A% &xr.)9:F)3F86'B$'-~~':':$7 (rc V^8dQhRRRR/#)rrrOrrr)rs"rrrs    rc "WnVeR#VPpVfR#\V\PP 4'dRVP 2VnR#VPPRR4 RVP 2VnR#)Nztorch.ops.higher_order.._ops..ops.rM) rr rrDr+HigherOrderOperatorrrrW)rrrb s&& rr#ExternKernel.set_python_kernel_names"4  ) !! >   > > ? ?(??P&QD #$$,,Xw?@&//ARS  #rcV^8dQhRR/#rr)rs"rrrs++Z+rc n^RIHpVP4;p'd VPM\P P p\P P'd VP#\P P'd\\P PV4'g-\\\P P44hVPfR#\P PPVPV4#VP#)r) CppWrapperCpuN)codegen.cpp_wrapper_cpurerBrr{r9 device_type fx_wrapperrrTrrrrget_c_shim_func_name)rredr8s& rr< ExternKernel.try_get_kernel_names:!%!22A29L9L 77   ** * WW agg22MBB$T!''*>*>%?@@##+77''<<$$f ** *rcV^8dQhRR/#rr)rs"rrrsrc DVP4pVf \R4hV#)NzExpected name is not None)r<rrs& rget_kernel_nameExternKernel.get_kernel_names&'') < !<= = rc V^8dQhRRRR/#)rrrrrr)rs"rrrs  f  rc  \PVP4VP4VP 4VP 4VP 4VP4R7pVP4V#))r8r7rrrr) rDrrBrArr<rrr)rr s& r copy_inputExternKernel.copy_inputsa   <<>++-]]_::<))+oo'    rc(V^8dQhRRRRRRRR/#)rrb rrrrrr r)rs"rrrs2v v !v *-v 9<v v rc  JaaRVRV/p\P!V4wpo.o.p.p.pVEFp T ;\Rc;enw\PP P PV RR7p SPR4VPV 4VPV 4K~;\Rc;ewSPR4VPV 4V PPp V PPR9dV f \R4hVP\PP!V 4P#44EK/;\$Rc;eDwSPR4VPV 4VPV P&4EK\(Rc;e(wSPR4VPV 4EKSPR4VPV 4VPV 4EK R VV3R llp VU u.uFqP+V 4NK pp VF#p \-V 4'gK\/V RR 7K% .p\1V\P2P4P64;'g?\P8P:\P2P<P>JpVEFp \1V \@4'gnV PC4\PPD9dAVP\PPDV PC4,4K\1V \@4'goV PC4\PPF9dBVP\PPFV PC4,4EK \1V \H4'dV'd\PPJpV PC4V9d:\M\PNV PQ44VV PC4&VPVV PC4,4EKVPV PS44EK\1V \T4'dVPV PV4EK\1V \PXPZP4'dV PPp V PPR9dV f \R 4hVP\PP!V 4P#44EKVP\]V 44EK V !W4wpp\PNP p\_4;_uu_4\aV4;_uu_4V!V/VBpRRR4RRR4RpV'd\P8PbPeR 4p\g4p\P8P:\P2P<P>Jd#V^,p\i\P84pV;_uu_4\kV\P8X4RRR4\mVXV4p\1X\n\p34'gV.MTpVFp\1V\Pr4'gK%VPt'gK9\vPx'dKQRp\PP8PbPeRR4;p'dV RV 2pV\Pn=K \}VVVV VR7#uup i +'giEL;i +'giEL;i +'giEL-;i)zPartition kernel args into tensor and non-tensor, realize tensor inputs, re-run fake tensor propagation with the realized strides, and return a ProcessKernelResult (see that class for field semantics). rrrN)rFzHExpected arg.device.type in ["cuda", "xpu"] and device_index is not NoneTc$V^8dQhRRRRRR/#)rnew_tensor_argsrnew_non_tensor_argsrztuple[list[_T], dict[str, _T]]r)rs"rr1ExternKernel.process_kernel..__annotate__s$ : :) :@L : + :rc8<.p\V4p\V4pSFApV'dVP\V44K'VP\V44KC \P!VS4pVP R.4VP R/43#)rr)rr@rr&tree_unflattenr1) rwrxr it_tensorsit_non_tensors is_tensorrargs_flat_is_tensor args_specs && rr 3ExternKernel.process_kernel..unflatten_argssFo.J!"56N0 MM$z"23MM$~"67 1 %%fi8A55$aeeHb&99 9rrHzFExpected x.device.type in ["cuda", "xpu"] and device_index is not NonerfzEsparsity not handled. Please file issue for sparse inference weights.r)z Found from : )r r r r r)r\r])?r& tree_flattenr"r{r9r:r create_symintnoder@GeneratorStater8rrrrD_C _accelerator_getDefaultGenerator clone_stateOpaqueObjectStaterJrrs r=rr_higher_order_ops torchbind CallTorchBindr reffects with_effectsrr constantstorchbind_constantsr torchbind_replay_objectsr-r get_real_obj get_valueropaque_example_valuer.irr.r%r rr1r r7r>r8rrr is_sparserMgraph_partitiondisable_cudagraphs_reasonr )rrb rr binded_args args_flatr r real_non_tensor_argsr r device_indexr r example_argsreplay_torchbindreplay_objectsnew_args new_kwargsr r r node_meta_valctxexample_out_lirHmsgr)rrs&&*, @@rprocess_kernelExternKernel.process_kernelstXv6 %22;? 9*,$& (*  CTV77++55GGRVGWD'..u5#**40(//5  &^%'..u5#**3/#&::#3#3L ?:|?W,f)//AA(%+-& )&('..u5#**3/(// :) X'..t4&&s+'..u5#**3/(//4IL : :6AA[((+[ AA$Q''%a5   vu66@@NN O U U~~$$(?(?(G(G(T(TT Aa**qzz|qww?P?P/P##AGG$5$5ajjl$CDq(++JJLAGG$?$??##AGG$?$? $MNA//#%&WW%E%ENzz|>97HKK)98qzz|4!''qzz|(DE '' 6A011##A$:$:;Au11@@AA xx~~  8\=U(`##HH==$!km ##$5a$89EH .lQ*KK)) $ & & 4Y ? ?#X<>>J 9>=!  ntUm<<    A1ell++KKK...]"#''"6"6";";"?"? t"TT;T E!2;-@C471 #)#+)/   MBB @ ? ? ' & &s6&_$;_= _) _=`) _: 4 _== `  `" c V^8dQhRRRR/#)rrrrrr)rs"rrrtsG G FG G rc  \V\4'g\\V44h\V\4'dV#VP 4p\ PPVP44pVf \R4hVP4pVeRVP9d\V\\\34'd\VP\4'd\!VPR,\"P$R7'g3\!VPR,\"P&R7'd*VP)\+VP-444MVP/4\0P2!VP-4RR7wrVV^,pVP54!V4p\ PP6P9W4p\ PP6P;W4p \ PP6P=W4p \?Wy4V ,p W8wd\@PCRV V V4\Dh\ VPF\IVPK4VPM4VP-4V V RR7R 7#) 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. zExpected buf is not Nonerfr rrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFrz r)'rrrrrrr{r9r rrrrr rr?r3rDr r rr5r<rrNrrr:r/ stride_vars offset_varrrr.r/rYrrrrA) rr x_unwrap_viewrx_unwrap_view_fx_node index_argsr rrrrexpecteds && rconvert_to_reinterpret_view(ExternKernel.convert_to_reinterpret_viewss3!X&& a) ) a ) )H  gg  !7!7!9: ; !;< < # 3 3 5 " -.333=?FJ*OPP=//@@8)..u5"'"5"5<)..u5"'"8"8  7 7.}/E/E/GH   ' ' )!-!@!@ JJL"   ]  ,  55eH''""..uA!!,,U?Z1F:   IIR   & %,,.kkmZZ\   rc V^8dQhRRRR/#rNr)rs"rrrs!!!!f!!!!rc  Vf \4#\V\\PP P \34'd \VR7#\V\4'dz\4;_uu_4\PP\P!VP VP#4VP%4R74uuRRR4#\V\&4'dV#\V\(4'dVP+VP,4#\V\.4'd5\/VP+VP,4VP14R7#\V\24'dBVP54\7VP944'dVP;V4#\V\>4'dVP54V#\V\@\\B34'dV#VPEV4# +'giELW;i \<dL}i;i)N)r]r7r8r)#rL rr"rrrrrrr5 rEr{r9add_tensor_constantrDrrJrArBrhrrs rrr>rrr=rrrYr NonTensorObjrrrr s&&rrs ExternKernel.realize_inputs 9') ) a$ 3 3 ; ;SA B B(a0 0 a " "())ww22LL allnU*) a ( (H a # #$$QVV, , a ) )"&&qvv.q||~  a " " IIK$Q]]_55::1== a $ $ IIKH a,(=?PQ R RH~~a  3*))"+sAH/I/ I  IIc V^8dQhRRRR/#rNr)rs"rrrs!!!6!rc \V4'dC\VP44^8XdV#VP4FpV^8XgK Vu# VPV4#r)r=rrnrr)rrr6s&& rrequire_stride1ExternKernel.require_stride1sR  # #1<<>"a',,.Q;H)~~a  rc ,V^8dQhRRRRRRRRR R/#) rrrrzSequence[int] | Nonerr}rrrr)rs"rrrsCVV V$V1 V  V  Vrc  VfVf \R4hVP4R9d V'gV#\V4'Ed\VP 4\ 4'dSV'd8\ W4pV'dVP4V#\VRRVVR7V#\VRRRVVR7V#RpV'dDV'd<\\PW1P4VP444p\VP 4\\34'dV'd&VP 4P!V4'gAV'dJ\#W1P 4P$VP44'dVe \'W4#T#Ven\VP 4\\34'dD\#WaP 4P$VP44'd \'W4#\VP 4;p\(4'd\VP+4;p\ 4'd \R4h\V\4'dTV'dVP!V4'g3V'd-\#W8P$VP44'dV#\V\,4'dpV'd&VP 4P!V4'gAV'd;\#W1P 4P$VP44'dV#\V\.4'd\VP0\24'd\VP0\44'g\VP74;p 4'd\9V R4'dy\V P0\:4'gYVP=VP04VnV'dVP?WVR 7#V'dVPAWVR 7#Rp VP4p Ve\DPFPHp \K\MVP444U u.uFTp V POW=,^4'gK#V PQVP4V ,^4'gKRV NKV p p V F.p\RPTPVPYW^^4pK0 VP4PZ'EdVeI\ VP]4VP4VP44P_W4R 7pMVVf \R 4h\ VP]4VP4VP44PaW$R 7pTp\RPTPVPcVP4VP$VP4VP]4R 7p\dPgVV4M"VPiV4p\VRRVVVR7V'd^\kVP44'g?\kVPm44'g \ W4'g \R 4hV#V 'dHV eVf \R 4h\RPTPVPoW4p\'W4#V# \BdELi;iuup i)zUEnsure x has the requested stride order or exact strides, inserting a copy if needed.Nz7Expected order is not None or exact_strides is not NoneTF)rIrXrYrr\zHthe MutationLayoutSHOULDREMOVE's real layout shouldn't be FlexibleLayoutrr]z5Expected order is not None when exact_strides is Nonerz5Expected is_stride_order_storage_and_layout(x, order)zr?rbrrrr_rNr<rArr[ rarr6rr r rrrrrrr,r rrequire_stride_orderrJrYr{r9r:rrrr rDr.loweringslice_ is_complexrrA r= rEInplaceCopyFallbackrrrr:rnrU)rrrrruse_current_stride_orderpadded_exact_stridesmutation_layoutr r expanded_dims orig_sizer:rrrr s&&&&& rrequire_stridesExternKernel.require_stridess!]%> I  ;;=F "=H ! # #!,,..990R0,0)H.#',1).*7 H*#(-%)&3&3 H $( '+'' zz|Q[[]S($!,,.;*HII1<<>;;EBB!1%||~'<'>-!<<>#8#8!**, H q) $ $1668,,qvv77%Q]]_&DkEE V,,{//1BCC 88@33 4#44 5#+/ JJL  $ww''Hs1::<011A33M4DaH11!**,q/1E1 %OO,,33AAqA% ;;= # # #('))+Q[[]AJJL""="N=(O())+Q[[]AJJL!/%/EC((66  kkm,,. 7A & &q# .q!A ! %"++   )!**,77)!,,.995a??$K@)m.G$R((//=A21D D]'  s6'] 8] ] ] ; ]2 *]2]2 ]/.]/c(V^8dQhRRRRRRRR/#)rrrrryrrrr)rs"rrrs,     '9  JN    rc  TPTVUu.uF;p\V\P4'dVPP MTNK= upVR7#uupi))rr)rrrDSymIntrr])rrrrrs&&&& rrJ"ExternKernel.require_exact_stridess]"" KXKXaz!U\\:: A=( #  sAA c(V^8dQhRRRRRRRR/#)rrrrrrrrr)rs"rrrs2PPP,P=AP Prc (VPWVR7#))rr)r)rrrrs&&&&rr!ExternKernel.require_stride_orders""1"OOrc V^8dQhRRRR/#rNr)rs"rrrs>>f>>rc .VPV\4#r)rr r s&&rrequire_channels_last"ExternKernel.require_channels_lasts''+<==rc V^8dQhRRRR/#rNr)rs"rrrs???F?rc .VPV\4#r)rr r s&&rrequire_channels_last_3d%ExternKernel.require_channels_last_3ds''+=>>rc V^8dQhRRRR/#rNr)rs"rrrs6frc RRlpV!V4'dV#VPV\PVP444#)c V^8dQhRRRR/#rFr)rs"rr5ExternKernel.require_contiguous..__annotate__s S S S4 SrcVP4pT\PP 9;'d,\PP T,P # \\3dR#i;ir)rAttributeErrorrYr{r9r is_mkldnn)rrs& ris_mkldnn_tensor9ExternKernel.require_contiguous..is_mkldnn_tensors` zz|177,,,RR1B1B41H1R1R R#$78  sA""A76A7rJr?r@r<)rrrs&& rrPExternKernel.require_contiguoussC S A  H,,>44QZZ\B rc V^8dQhRRRR/#rNr)rs"rrrs  6 f rc hVPV\PVP444#rrr s&&rrequire_contiguous_strides'ExternKernel.require_contiguous_stridess-(( ~00>  rcV^8dQhRR/#rr)rs"rrrs  $ rc R#rrrs&rr5ExternKernel.apply_constraintrrc$V^8dQhRRRRRR/#)rrr9 rrrr)rs"rrrs$$$!$+9$ $rc 6\V\4'g\\V44h\V\4'g \ V4pVP 'g \R4h\ V4p\ VP 4pW48d\PRVPWC, 4\W44FTpVP V,R,pTPWb9d W&,MVP V,R,4KV V#)z/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) rrrrrr rr.r/r rr@)rrrn_args n_pos_argsrarg_names&&& rfill_non_provided_args#ExternKernel.fill_non_provided_argss$)) d, ,$%%:D""" !RS ST,,-    II^  #   6...q1&9 )$,,Q/@/ rc V^8dQhRRRR/#)rrrrrr)rs"rrr!s#X#X(8#XI#Xrc \PP'Ed.pRpV'dsVP'da\ VP 4\ V48wd \ R4hVPUu/uFpVPR4;pfKWTbK pp\VP 4FwrgVeCVf \ R4hVPW,4pV'dVPR4MRp Ml\ VP4V,p VP'd=V \ VP48d#VPV ,PR4MRp VP\PPPWy44K V#VP U u.uF,p \PPPV 4NK. up #uupiuup i)NzDnames passed to codegen_const_args does not match self.constant_argsrzExpected names is not Noner) r{r9rTr rrrr1rrr@rval_to_arg_str) rrrname_to_arg_propertiesr rrrproptype_rrs && rcodegen_const_argsExternKernel.codegen_const_args!s 77   F EI ",,,t))*c%j8(^ $22*2 #/DI2'* "$"4"45)5},-IJJ155eh?D04DHHV,$Edkk*Q.C ...3T=P=P9Q3Q++C044V<!  agg22AA!KL6MDHDVDVWDVqAGG((77:DVW W-*,Xs<G-G-82G2cV^8dQhRR/#rrrr)rs"rrrFsirc <\PP'dFVPe8VP .VP OVP OVP4pRpMVP pRp.p\V4FwrE\PP'dVP'dV\VP48g \R4hVPV,PR4pVP\PPPWV44KVP\PPPV44K V'd VP!VP#44V#)NFTz-Invalid access to ExternKernel.arg_propertiesr)r{r9rTr rrrrrr rrr1r@rrr r)rrneed_codegen_constant_argsrrrrs& r codegen_argsExternKernel.codegen_argsFs3 77   4#3#3#?003$++3 2 23T[[F*/ &[[F)- &f%DAww"""+++C8K8K4L0L(G++A.226: AGG00??IJ AGG00??BC& & KK//1 2 rc$V^8dQhRRRRRR/#)rrrrrrr)rs"rrras& J J J J Jrc W9dVPV4#WP9dVPPV4#VPPV4;peVPR4#\V R24h)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)r1rr r)rrrr s&&, rget_kwargs_valueExternKernel.get_kwargs_valueast  ::h' ' {{ ";;??8, ,))--h7 7C D77?+ +z)GHIIrc V^8dQhRRRR/#)rskip_outrrrr)rs"rrrnst rc  \PP'Ed VPe\ VP 4^8Xd.#.pVP FpV'd VR8XdKVPV4p\V\4'dVPV4KNVPf \R4hVPPV/4PR4pVP\PPPWE44K V#VP P#4UUu.uF3wrdV R\PPPV4 2NK5 pppV#uuppi)Nr'z+Expected self.allarg_properties is not Nonerr)r{r9rTr rr1rrrr"r@r rr1rrrr )rrrrr rks&& rcodegen_kwargsExternKernel.codegen_kwargsnsG 77   +D4F4F0G10L F >>E 1))(3a&&MM!$--5,I!2266xDHHPEMM!''"6"6"E"Ea"OP?( !KK--//DA#Qqww++::1=>?/  s9FcV^8dQhRR/#rr)rs"rrrs  S rc VPeXVPPp\VRR4pVPRR4pVP R^4^,pV RV 2pV#RpV#)Nrunknown_namespacer_r`rM unknown_op)rrrrWrsplit)rr op_namespaceop_names& r get_op_nameExternKernel.get_op_namesv << #\\((F"6<9LML'//'BL'..sA6q9L%ax0G#GrcV^8dQhRR/#rsr)rs"rrrs   rc \VP\PP4;'d2\P P VPP9#r)rr rDr+r,Tag inplace_viewtagsrs&ris_inplace_viewExternKernel.is_inplace_viewsK t'')>)> ? @ @ &&$*:*:*?*?? rc V^8dQhRRRR/#)rrrrrr)rs"rrrs   3  4  rc JVP4'*;'dVPR4'*;'djVP\PP P P\PP P P39#)z#torch.ops.aten.quantize_per_tensor.) rrQr rDryrLquantize_per_tensorr|tensor_qparams)rrs&&rshould_assert_dtype ExternKernel.should_assert_dtypes$$& &  &&'LMM    IINN66>>IINN66EE rc V^8dQhRRRR/#rDr)rs"rrrsMM,@MTMrc L\P'gR#VP4'd#\PP 'gR#VP 4pVP4p\PP 'djVP4'dT\VP^,\4'g \R4hVP^,P4p\PPPVP44p\PPPVP44pVP!V4'dVP#4MRpVP%W4WRV4R#)Nz+Expected isinstance(self.inputs[0], IRNode))rM size_assertsrr{r9rTrrrrrrrcodegen_shape_tupler<rnrrAwrite_assert_size_stride)rr8rrr5r6r7s&& rcodegen_size_asserts!ExternKernel.codegen_size_assertss"""     ! !!''*=*=*= ""$}} 77   ##%%!$++a.&99()VWW{{1~..0ww##77 H%%99$//:KL$($<$G5Gc V^8dQhRRRR/#r#r)rs"rrr s  !  ! rc V'd\M\p\PW4pVPFpW2!V4,pK VP P 4FpW2!V4,pK V#r)maybe_free_unbacked_symbolsmaybe_free_symbolsrr'rrr)rrmaybe_get_symbolsrr s&& rr'!ExternKernel.get_free_symbol_uses sp,9 '>P   - -d B%%C "3' 'A&;;%%'C "3' 'A(rcV^8dQhRR/#rr)rs"rrr s & & &rc  2\VRR4pRV: 2.pT\P!V4Uu.uF'pVP R\WP4 2NK) up, pVP RVP : 24VP V4#uupi)rNzpython_kernel_name=rr)rrTfieldsrr@rrN)rrm rFrUs& rrExternKernel.__str__ sd$8$? !+ 1  $++D1 1zzl!GD**56 71    |D$4$4#789u%%  s-B)r rr rrrrr& r rrrr1rrrNNNNrNr)NNFr)DrrrrrSrrrTrUrrrrrrrr r r r rr& r/rrr rrrArGrrr<rnrVrrr:rrrs rrrJrrrrPrr5rrrrrrrrr$r*r1r5r>rer'rr9rrrs@rr+r+s $&M=%(..tDFND*.K'.%) )"&OZ&4?3D3D4!=(,K$+/3N,30;0A0ARV0W-W6:3:->t-T*T11@. #$J! M""#;J +"   v v pG G R!!!!F!!VVp    PP >>??"   $L#XJ6 J:   M&  =$'BN+ ,  &Hrr+cJa]tRtRtRRltR RV3RllltRRltRtV;t#) ExternKernelOuti c V^8dQhRRRR/#rDr)rs"rrExternKernelOut.__annotate__" 11311rc (VPV4R#r)generate_extern_kernel_outrFs&&rrGExternKernelOut.codegen" **40rc@V^8dQhRRRRRRRRR R R R R R RRRRRR/ #)rrr_rrrr9 rr rrrrOrrr r rrr)rs"rrrN% sn)))!)% ) & ) , )')$)(5))) )rc  H<VPV4p \V \4'g\\ V 44h\ S T`RTT TT;'g/RVVVV 4 \PPV4Vn \PPV4R#r) r rrrrrr/r{r9r! rr" ) rrrrrrrrrr unwrapped_inputsrs &&&&&&&&&& rr/ExternKernelOut.__init__% s ..v6*H55 &6!78 8      LLb    )  GG++D1  ""4(rcV^8dQhRR/#rsr)rs"rrrNC r rc R#rRrrs&rr' ExternKernelOut.should_allocateC rrrrJ) rrrrrGr/r' rrrs@rrLrL s1))<rrLc.a]tRtRtRV3RlltRtV;t#) RandomSeedsiG c$V^8dQhRRRRRR/#)rr rr8rrrr)rs"rrRandomSeeds.__annotate__H s!  c < D rc  <\P!\P4p\SV`\ V\PV.R7.VP VPV..RR\PPR7R#)rzaten.randint.low_outzat::_ops::randint_low_out::call)rrrrrr N) rDrrWrr/rr{rrLrandintlow_out)rr r8limitsrs&&& rr/RandomSeeds.__init__H slU[[) kkW !::vzzE7;5> ,,  rrrrrrr/rrrs@rr\r\G s   rr\cVa]tRtRtRRltR RV3RllltRRltRR ltR tV;t #) r i[ c V^8dQhRRRR/#rDr)rs"rrExternKernelAlloc.__annotate__\ s33333rc (VPV4R#r)generate_extern_kernel_allocrFs&&rrGExternKernelAlloc.codegen\ s,,T2rc<V^8dQhRRRRRRRRR R R R R RR RRR/ #)rrrSrrrr9 rr rrOrrr r rrr)rs"rrrg_ sd!)!)!)!!)% !) & !) ' !)$!)(5!))!) !)rc  <VPV4p \;QJdRV 4F 'dK RM RM !RV 44'g \R4h\S T`RT\ \ \,V 4TT;'g/RVVVV4 .Vn\PPV4Vn \PPV4R#)c3B"TFp\V\4xK R#5irr ris& rr-ExternKernelAlloc.__init__..k sC2BQ:a((2Brr FTz=Expected all(isinstance(i, IRNode) for i in unwrapped_inputs)N)r rrrr/rrrrM r{r9r! rr" ) rrrrrrrrr rVrs &&&&&&&&& rr/ExternKernelAlloc.__init___ s ..v6sC2BCsssC2BCCC O     &!#3 4  LLb    )  ') GG++D1  ""4(rcV^8dQhRR/#rsr)rs"rrrg rc R#rrrs&rr' !ExternKernelAlloc.should_allocate rrcV^8dQhRR/#rr)rs"rrrg s""$"rc \hrrers&rr5"ExternKernelAlloc.apply_constraint rgr)rrM )rNNNrN) rrrrrGr/r' r5rrrs@rr r [ s"3!)!)F""rr cba]tRtRtRtRV3RlltRRltRRltR R ltR R lt R t V;t #)r' i zH An output buffer that represents the mutation of a pre-existing buffer c(V^8dQhRRRRRRRR/#) rrrS mutated_noder mutating_noder[rrr)rs"rrMutationOutput.__annotate__ s,22 2062GP2 2rc <\SV`RVR7VP4p\PP V4V.VnW0n\PPV4Vn R#)NrI ) rr/rr{r9r mutation_namesrzr! r)rrryrzmutated_node_namers&&&& rr/MutationOutput.__init__ s` d62(113 ##$5601(5GG++D1 rcV^8dQhRR/#r0 r)rs"rrr{ s"""rc VP#r)rzrs&rrMutationOutput.get_defining_op s!!!rcV^8dQhRR/#rCr)rs"rrr{ ##M#rc VP#rr}rs&rrE!MutationOutput.get_mutation_names rrcV^8dQhRR/#rsr)rs"rrr{ rqrc R#rrrs&rr' MutationOutput.should_allocate rrcV^8dQhRR/#rrrr)rs"rrr{ s  &6 rc jVP4pRV4Uu.uF pVfK VNK up#uupi)c3`"TF$p\PPV4xK& R#5ir)r{r9try_get_buffer)rrs& rr6MutationOutput.get_mutation_buffers.. s"P..t44s,.)rE)rr}rs& rget_mutation_buffers#MutationOutput.get_mutation_buffers sB002QP P CP   s00)rzr}r) rrrrrSr/rrEr' rrrrs@rr' r' s+22"#  rr' cNa]tRtRtRtRtRV3RlltRRltRR ltR t V;t #) OrderingBarrieri aA no-op that creates a rename chain for scheduling order. Takes a source buffer and produces a new buffer name. The scheduler sees this as a node with a rename (source -> self), so future readers of the source are redirected through this barrier. Uses WeakDep(is_fake=True) instead of StarDep via the ``ordering_only`` flag, avoiding lifetime extension and mark_buffer_mutated side effects. No kernel is generated and no allocation occurs. Tc V^8dQhRRRR/#)r source_noderrrr)rs"rrOrderingBarrier.__annotate__ s ) )F )t )rc 0<VP4\SV` R\VP 4R7V.R7VP 4.Vn\PPV4Vn \PPV4R#)Nrr ) rrr/r~ rBrr}r{r9r! rr" )rrrs&&rr/OrderingBarrier.__init__ sx [%;%;%=>=  +3356GG++D1  ""4(rcV^8dQhRR/#rCr)rs"rrr rrc VP#rrrs&rrE"OrderingBarrier.get_mutation_names rrcV^8dQhRR/#rsr)rs"rrr rqrc R#rrrs&rr' OrderingBarrier.should_allocate rr)r}r) rrrrrSr r/rEr' rrrs@rrr s( M ) )#rrca]tRtRt$Rt/tR]R&]RRl4t]RRl4t R V3R llt R R lt R Rlt Rt V;t#) TMADescriptori aL 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]_CACHEc$V^8dQhRRRRRR/#rrrtma_metaztuple[str, tuple[Any, ...]]rrr)rs"rrTMADescriptor.__annotate__ s$ = = ='B =  =rc \V4^8wd \R4hV^,R8Xd\V.V^,O5!#V^,R8wd \R4h\V.V^,O5!#)rzExpected len(tma_meta) == 2 experimentalrz Expected tma_meta[0] == "stable")rrTMADescriptorExperimentalTMADescriptorStable)rrrs&&&r _create_implTMADescriptor._create_impl sh x=A  !>? ? A;. (,VBhqkB B{h&$%GHH&v< < rrrrrr{r9r! rr" )rrrrrs&&&&rr/TMADescriptor.__init__ s   !,,.  &!6 * -   GG++D1  ""4(rc V^8dQhRRRR/#rDr)rs"rrr!..3..rc (VPV4R#r)generate_tma_descriptorrFs&&rrGTMADescriptor.codegen!''-rcV^8dQhRR/#rr)rs"rrr!rrc VP#r)rrs&r get_tensorTMADescriptor.get_tensor!rr)rr)rrrrrSrrr:rrr/rGrrrrs@rrr sT(*F $) = =))..rrc6a]tRtRtRtRRV3RllltRtV;t#)ri !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. c ,V^8dQhRRRRRRRRRR /#) rrrrlist[int | torch.SymInt] block_dims element_sizerrrr)rs"rr&TMADescriptorExperimental.__annotate__!s<   ' -  !    rc <\V4R9d \R4h\V4\V48wd \R4hVfVP4PpW nW0nW@n\VP4VnV.p.VPOVP OVP Np\SV`%VVVR7R#)rzExpected len(dims) in (1, 2)z%Expected len(dims) == len(block_dims)Nrrr)rr) rrrArrrrr rr/)rrrrrrrrs&&&&& rr/"TMADescriptorExperimental.__init__!s t9F " !?@ @ t9J ' !HI I  !++-66L $( N  YY __     '  r)rrrr rrrrrrSr/rrrs@rrr !s   rrc2a]tRtRtRtRV3RlltRtV;t#)ri6!z the new host-side TMA descriptor API (the ones obtained via TensorDescriptor.from_tensor). See also TMADescriptorExperimental for the old API. c V^8dQhRRRR/#)rrr block_shaperr)rs"rr TMADescriptorStable.__annotate__>!s  v 4L rc :<W n\SV` VV.VR7R#)rN)rrr/)rrrrs&&&rr/TMADescriptorStable.__init__>!s&& 8%  r)rrrs@rrr6!s  rrc>a]tRtRtRRV3RllltRRltRtV;t#)SubgraphBufferiH!c 0V^8dQhRRRRRRRRR R R R /#) rrr_rrrrexample_inputs list[Any] subgraph_namerr r r)rs"rrSubgraphBuffer.__annotate__I!sN-A-A-A"-A ! -A " -A  -A.-Arc  <\SV`RW4W0nW@n\P P V4Vn\P PV4\P PVPWE4Vn \VP4'g \R4h\VP4pVFTpWPPVP&VPP P#VP4KV VU u.uFqPNK up Vn^RIHuHp \P,!VP4;_uu_4RRRRRR/p /V CT;'g/Cp V P/V 4;_uu_4VPP0!VP!RRR4V'd=VPP2F"p V P5VP744K$ RRR4R#uup i +'giLc;i +'giR#;i)N&Expected is_node_sequence(self.inputs) max_autotuneFmax_autotune_gemmmax_autotune_gemm_backendsATEN)rr/rrr{r9r! rr" make_subgraphsubgraphr rrrrgraph_input_namesr@ sym_inputstorch._inductor.configr.rMset_graph_handlerr!run operationsrr)rrrrrrr rsym_inpsym_varinductor_config base_patchesmerged_patchesrrs&&&&&&& rr/SubgraphBuffer.__init__I!s v3,GG++D1  ""4(--dgg~U  ,, !IJ J(5 !G7>MM & &w|| 4 MM + + 2 27<< @"8BBzG<.CodegenGraphiy!cV^8dQhRR/#)rr9rr)rs"rr9SubgraphBuffer.codegen..CodegenGraph.__annotate__z!s ' 'm 'rc 4WnVPVnR#rr9r)rr9s&&rr/5SubgraphBuffer.codegen..CodegenGraph.__init__z!s" !JJ rrN)rrrrr/rrrr CodegenGraphry!s  ' 'rrrN)r rrr'codegen_subgraph_with_flattened_outputsrrr)rr8rrH outer_inputss&& rrGSubgraphBuffer.codegenx!s ' '   ,, !IJ J7;{{C{!++-{ C77  ' -doo - - YYK DsB)rrrrrrrrrrr/rGrrrs@rrrH!s-A-A^  rrca]tRtRtRtRRltRRlt]RRl4tR R lt R R lt ] !R4RR V3Rlll4t RRlt RV3RlltRRltRRltRtV;t#)UserDefinedTritonKerneli!z6 A user-defined triton kernel (e.g. via @triton.jit). cV^8dQhRR/#)rrz(tuple[Kernel, Any, list[str], list[str]]r)rs"rr$UserDefinedTritonKernel.__annotate__!s G G)Q Grc a^RIHp^RIHpVP VP 4o.p.p.p\ SV4'Ed\SR4'd'VPV3RlSP44M8\SR4'g \R4hVPSP4\SR4'dASPF/pVPSPPV,4K1 M8\SR4'g \R 4hVPSP 4SP"pSPoSW4V3#) r) Autotuner)kernel_side_table restore_idxc3^<"TF"pSPPV,xK$ R#5ir)r arg_names)rrrb s& rrBUserDefinedTritonKernel.get_kernel_and_metadata..!s'*4FqFII''**4Fs*- restore_valuez)Expected hasattr(kernel, "restore_value") reset_idx reset_to_zeroz)Expected hasattr(kernel, "reset_to_zero"))triton.runtime.autotunerr*torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrr,r rrrrr@rrrconfigs)rrrrrestore_value_argsreset_to_zero_argsrrb s& @rget_kernel_and_metadata/UserDefinedTritonKernel.get_kernel_and_metadata!s#6P"--doo>(*(* fi ( (v}--"))*4:4F4F*v77()TUU"))&*>*>?v{++))A&--fii.A.A!.DE*v77()TUU"))&*>*>?nnGYYFw4FFFrcV^8dQhRR/#rsr)rs"rrr!s7747rc $\P'gR#VPP'gR#\ VP P 4^8wdR#\ VP4^8wd \R4h\VP^,\4'gR#\VP^,P\4'gR#\VP^,PP\4'gR#\VP^,PPP\4'gR#\;QJdURVP^,PPPP 4F 'dK RMJ RMF!RVP^,PPPP 44'gR#R#)a 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. Fz$Expected len(self.mutable_args) == 1c3*"TF q^8HxK R#5irirrs& rr.!sN#Ma6#MrT)rM*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrrrrrDrrrs&rr)UserDefinedTritonKernel.can_fuse_epilogue!sd0@@@  222 t!!(( )Q . t  !Q & !GH H$++A. ::$++A.33Z@@$++A.3388.II$++A.3388==yIIsN4#4#4Q#7#<#<#A#A#F#F#M#MNsssN4#4#4Q#7#<#<#A#A#F#F#M#MNNNrc V^8dQhRRRR/#rDr)rs"rrr!s<<3<X7 7rc $VPW4#)zA epilogue_fusion: (fused epilogue node, modified kerel src code) r)rr8rs&&&rcodegen_with_epilogue_fusion4UserDefinedTritonKernel.codegen_with_epilogue_fusion!s }}W66rc$V^8dQhRRRRRR/#)rr8rrz!tuple[ComputedBuffer, str] | Nonerrr)rs"rrr!s-m m %m ;m   m rc  a^RIHpVP4wppppVPp/p V'd\ VP P P4^8wd6\R\ VP P P4 24h\\VP P P44Pp Vwr/VPCW/CpW/p VPVVVVVVPVVP4wp pppVPUu/uFpVVP!V4bK ppVP#V 4VP$Uu.uFpVPNK upoVP$Uu.uF#pVP&'gKVP(NK% pp\+V3RlV44p.p.p.p.p\,P.!VP14\3\,P4!R4V44EFwppVV9dV!4'dKVP7V4VP7V4\9V\:4'dAVP7VP=44VP7VP?44K\9V\@\B\D\F34'd.VP7V4VP7\IV44KVV9d*VP7R 4VP7\@4EKVf[V!4'd*VP7R 4VP7\@4EKSVPK4VPK4EKv\MR\IV4 RV 24h VPOW4VPQV VVVVVVRVPS4VPTPR 7 R#uupiuupiuupi) QOverrides the parent member. See https://github.com/pytorch/pytorch/issues/151692)triton_version_uses_attrs_dictzexpected one write, got c36<"TFpSV,xK R#5irr)rrrs& rr3UserDefinedTritonKernel._codegen..*"s$F:aYq\\:rNr<NzUnsupported arg type: rQT) arg_typesraw_argsraw_keys triton_meta inductor_metarZr8original_fxnode_namerl)+r r%r rrrrr rrrr!define_user_defined_triton_kernelgrid launch_kwargsrrr r is_constexprnumrDrPr r rrepeatr@rrrrArrArr"rrrYrAgenerate_kernel_callrBr)rr8rr%rb rr r  kernel_kwargsepilogue_out_overridemutable_arg_nameepilogue_computed_bufferr&new_namer+r,extra_launch_argsr named_argsr constexprsconstexpr_namesrr(raw_keys_filteredraw_args_filteredrr rs&&& @rr UserDefinedTritonKernel._codegen!s I  ( ( *       02 4$$00778A=$.s43D3D3P3P3W3W/X.YZ $D):):)F)F)M)M$NOTT *9 ' $Wt{{W,<WM%5$P !  5 5      II         261S1S 1SAAt$$Q' '1S   /0%+]]3]QVV]3 %+]]E]anneaee] E$$F:$FF! ')')"    I$4$4R$8:K L ID#&+I+K+K  $ $T *  $ $S )#v&& C1134  1C#udD!9:: C   c+( B  %233KKO$$S)%))+%))+),B49+RPSu*UVVI L W/$$  &&#'??$!%!2!2 % i 4Es)O8%O= P$Pc V^8dQhRRRR/#r#r)rs"rrre"rrc X<\SV`V4\VPV4,#r)rr'r+r/rs&&rr',UserDefinedTritonKernel.get_free_symbol_usesd"s- w+M:=M II}>   rcV^8dQhRR/#rr)rs"rrrn"rrc \4#rrCrs&rr0UserDefinedTritonKernel.get_unbacked_symbol_defsn"rrc 0V^8dQhRRRRRRRRRR R R /#) rrrr/rtma_descriptor_metadatar kernel_argsr0ztuple[str, ...]rrr)rs"rrrq"sPS)S)S) S) "0 S) $ S)'S) S)rc  ><.p/p.pVP4Fwr\V \4'd^\P VP V 44p W9d\ PWV ,4p VPV 4WV &KxVPV 4WV &K \V4^8Xd \R4hV^,P4Vn \V\4'g\\V44h\SV`AR\#VPR7V\%V4V4WnW nWPnVP-4wrp\/V R4'g \R4hV P0Uu.uF qV9gK VNK upVn^RIHpHp\V 4^8dV ^,P:M/p^RIpV P>Vn VPCVP@4Vn"V!VP@4Vn#W@n$V!V /VCVCV4Vn%VPJPLPNUu.uFGp\VPQVPR4\4'gK4VVPR,NKI upVn*VPTUu.uF$p\W\#VPR7VV4NK& upVn,\ZP\P_V4R#uupiuupiuupi)rzExpected len(inputs) != 0Nrrz%Expected hasattr(kernel, "arg_names"))identify_accessed_tensorsidentify_triton_stores)0r rrrr rs rrr@rrrBr8rrrr/r~ rrr/r0r r,rrrrKrLrastr kernel_srcparse kernel_astrrIrrr r1rrr' r& r{r9r" )rrr/rHrIr0rrrrr rHrb rr&r rKrLautotuned_kwargsrMr rrs&$$$$$ rr/ UserDefinedTritonKernel.__init__q"s "$&&( %%'DA!Y'' 99$:L:LQ:OP/%,,Q0JKA a q $$Q'q ( v;!  !<= =Qi**, &(++ f. .   dkk *  -   % * $ < < >Av{++ !HI I!++. +Ck/ACC+. * 14G q0@71:,,b!**))DOO43DOOD&5  /{ /. / # ((44;; ;+//#((3Y? "K ! !; ((! ( :T[[93 E(!  ""4(K. 6 ! s LL1LL=*LcV^8dQhRR/#rr)rs"rrr"s++\+rc ,\VP4#r)rr& rs&rr#UserDefinedTritonKernel.get_outputs"sD))**rcV^8dQhRR/#rr)rs"rrr"rrc VP#rrrs&rrB"UserDefinedTritonKernel.get_device"rr) rr8r/rIrPrrNrr0rr& rr)rrrrrSr rrrGr rrer'rr/rrBrrrs@rrr!su GD7r<<7m ^56  7 S)S)j+rrcba]tRtRtRtRRltRRltRRltR R ltR V3R llt R t V;t #)InplaceBernoulliFallbacki"= This needs to be a custom class to handle mutation properly c V^8dQhRRRR/#rDr)rs"rr%InplaceBernoulliFallback.__annotate__"s3rc  \;QJd&RVP4F 'dK RM RM!RVP44'g \R4hRVP4wp\PP 'd^VP VP4 RV RRP\\VP44 RVP 24R #VP VP4 RV RRP\\VP44 RVP 24R #) c3B"TFp\V\4xK R#5irr rrHs& rr3InplaceBernoulliFallback.codegen.."s>+Q:a((+rr FTz8Expected all(isinstance(t, IRNode) for t in self.inputs)c3^"TF#p\\V4P4xK% R#5ir)rrrr`s& rrra"s"I[VQ1133[s+-rIrz, NULL)rRN) rrrr{r9rTrrnrLrJreprrending)rr8rs&& rrG InplaceBernoulliFallback.codegen"ss>$++>sss>$++>>> J JT[[I 77      '')*!A3b3tTEWEW;X1Y0ZZabibpbpaqr    '')*!A3b3tTEWEW;X1Y0ZZ[\c\j\j[kl rcV^8dQhRR/#rsr)rs"rrr]"rqrc R#rrrs&rr' (InplaceBernoulliFallback.should_allocate"rrcV^8dQhRR/#rCr)rs"rrr]"$$M$rc &VP^4.#rr rs&rrE+InplaceBernoulliFallback.get_mutation_names""##rcV^8dQhRR/#rr)rs"rrr]"rrc \4#rrCrs&rr1InplaceBernoulliFallback.get_unbacked_symbol_defs"rrc(V^8dQhRRRRRRRR/#) rr rrrrrrrr)rs"rrr]"s, ) )' ),2 )DG )  )rc `<\SV`R\VP4R7VP V.4VVR7\ P PVP44\ P PV4Vn \ P PV4R#)Nrr ) rr/r~ rBr r{r9r rr! rr" )rr rrrs&&&*rr/!InplaceBernoulliFallback.__init__"s   alln -    $ #   ##AJJL1GG++D1  ""4(rr rrrrrSrGr' rErr/rrrs@rrZrZ"s&$$ ) )rrZc|a]tRtRtRtRRltRRltRRltR R ltR V3R llt ] RR Rll4t Rt V;t #)ri"r[c V^8dQhRRRR/#rDr)rs"rr InplaceCopyFallback.__annotate__#s<<3<.8#s 9[!!##[rkrIrrRN) r rrrJrcrrrnrLrd)rr8argrefss&& rrG$MutatingFirstArgExternKernel.codegen4#s ,, !IJ J 9T[[ 9 t)) *  ##%&a '(:';1W^^(>(@Ar)rrrrdrs@rrrL#sBBrrc:a]tRtRtRV3RlltRRltRtV;t#)SetSourceTensorKerneli`#c$V^8dQhRRRRRR/#r self_tensorrstorage_tensorrrr)rs"rr"SetSourceTensorKernel.__annotate__a#!  F F t rc 0<VPV4pVP4\SV` VP 4W.R\ P PPPR7\V\\\34'g\\V44h\ P"P$P'VP(P+44\ P"P$P'VP+44\ P"P$P'VP+44VP-4p\/\1VR7W4\/\1VR7W 4.VnR#)z!torch.ops.aten.set_.source_Tensorrr rNrs rrr/r>rDryrLset_ source_Tensorrrrrrrr{r9rr-rrrBr' r~ r& rrrr8rs&&& rr/SetSourceTensorKernel.__init__a#++N;$$&   % % '  )B ++99  +*i'HII k!23 3 ##'' (8(8(A(A(CD ##''(?(?(AB ##'' 8**, :V4k H :V4n K! rcV^8dQhRR/#rCr)rs"rrru#88m8rc FVP^4VP^4.#rrlrs&rrM2SetSourceTensorKernel.get_inputs_that_alias_outputu#"DOOA$677rr rrrrr/rMrrrs@rrr`#  (88rrc:a]tRtRtRV3RlltRRltRtV;t#)ShallowCopyDataKerneliy#c$V^8dQhRRRRRR/#rr)rs"rr"ShallowCopyDataKernel.__annotate__z#rrc 0<VPV4pVP4\SV` VP 4W.R\ P PPPR7\V\\\34'g\\V44h\ P"P$P'VP(P+44\ P"P$P'VP+44\ P"P$P'VP+44VP-4p\/\1VR7W4\/\1VR7W 4.VnR#)z!torch.ops.aten.shallow_copy_data_rrNrrs&&& rr/ShallowCopyDataKernel.__init__z#rrcV^8dQhRR/#rCr)rs"rrr#rrc FVP^4VP^4.#rrlrs&rrM2ShallowCopyDataKernel.get_inputs_that_alias_output#rrr rrs@rrry#rrrcna]tRtRtRtRRltRRltRRltR R ltR R R R/RV3Rlllt Rt V;t #)ScatterFallbacki#z 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 V^8dQhRRRR/#rDr)rs"rrScatterFallback.__annotate__#s00300rc (VPV4R#r)generate_scatter_fallbackrFs&&rrGScatterFallback.codegen#s))$/rcV^8dQhRR/#rsr)rs"rrr#rqrc R#rrrs&rr' ScatterFallback.should_allocate#rrcV^8dQhRR/#rr)rs"rrr#s  I rc VP^,p\V\4'g \R4hVP 4.#rz Expected isinstance(inp, IRNode)r rrs& rrE"ScatterFallback.get_mutation_names#s7kk!n#v&& !CD D rcV^8dQhRR/#rr)rs"rrr#rrc \4#rrCrs&rr(ScatterFallback.get_unbacked_symbol_defs#rrrON include_selfTc8V^8dQhRRRRRRRRRRR R R R R R/#)rr rrrrrrr rOrOrrrrr)rs"rrr#sZ!)!)!!) !) !)  !)  !)!)!) !)rc  t<\V\4VnVP'd&W$V3Uu.uFqPV4NK p pV3p M#W$3Uu.uFqPV4NK p pW53p \S V`R\ VP4R7VPV 4V RVRV/\V4RR.VR7\PPVP44\PPV4Vn\PP!V4R#uupiuupi)NrrOr)rrr )rr src_is_tensorrs rr/r~ rBr rr{r9r rr! rr" ) rr rrrr rOrrHtensorsrrs &&&&&&$$ rr/ScatterFallback.__init__#s(Y7    78oFo))!,oGF FM78jAj))!,jGA JM   alln -    (  v~| <";/+3^*D#   ##AJJL1GG++D1  ""4(%GBs D0D5)rrrvrs@rrr#s? 0 !)"!)"!)!)!)rrcba]tRtRtRtRRltRRltRRltR R ltR V3R llt R t V;t #)IndexPutFallbacki#zI This needs to be a custom class to handle mutation and indices properly c V^8dQhRRRR/#rDr)rs"rrIndexPutFallback.__annotate__#s22322rc (VPV4R#r)generate_index_put_fallbackrFs&&rrGIndexPutFallback.codegen#s++D1rcV^8dQhRR/#rsr)rs"rrr#rqrc R#rrrs&rr' IndexPutFallback.should_allocate#rrcV^8dQhRR/#rCr)rs"rrr#rjrc &VP^4.#rrlrs&rrE#IndexPutFallback.get_mutation_names#rnrcV^8dQhRR/#rr)rs"rrr#rrc \4#rrCrs&rr)IndexPutFallback.get_unbacked_symbol_defs#rrc 0V^8dQhRRRRRRRRR R R R /#) rr torch._ops.OpOverloadrrrrrr9 accumulaterrrr)rs"rrr#sF))*) ) )  )  ) )rc  <W0nVUu.uF qffKVNK ppW$.VOUu.uFq PV4NK ppRp \S V` R\ XP 4R7VP V4V3RV VR7\PPVP^44\PPV4Vn \PPV4R#uupiuupi)Naoti_torch_index_put_outrzaten.index_put_)rrr )rrs rr/r~ rBr r{r9r r r! rr" ) rr rrrrr valid_indicesrrrs &&&&&& rr/IndexPutFallback.__init__#s $+=GqG =342M}2MN2MQ%%a(2MN4   alln -    ( M0+#   ##DOOA$67GG++D1  ""4(>NsC3C3C8)rrrvrs@rrr#s&2$))rrc8]tRtRt]RRl4tRRltRtR#) DeviceCopyi#c(V^8dQhRRRRRRRR/#)rrrr8rr|rrr)rs"rrDeviceCopy.__annotate__#s(. . v. |. 4. F. rc  \VP4pVf \R4hVP4'Eg\V4\P P 9d\;QJd*RVP44F 'dK RM RM!RVP444'd\PP'gp\P P'd?\P PV4\P PV4VPV4#\P PV4\P PV4\R4V3p\ P#V4pRpVP%4'dVP'4p\)VP*4;'dVP*R8H;'dTpVP*R8H;'d \)VP*4;'dTpV'd'\-V4'dRVP/4n\3\5VVP74VP%4VVR7VP9V4.V4#)NzExpected x_device is not Nonec3Z"TF!q\PP9xK# R#5ir)r{r9rrs& rr$DeviceCopy.create..$sG4Fq***4Fs)+FTzDeviceCopy in input programrXrX )rBrr6rr{r9r rr rM aot_inductoruse_runtime_constant_foldingrTadd_device_infor@rir+rPr<rnrqrr=r>rrrrArs ) rrr8r|x_devicerr6is_destination_pinnedis_source_pinneds &&&& rrDeviceCopy.create#s<<>   !@A A Qqww'>'>>GA4D4D4FGGA4D4D4FGGG''DDDww"""''/''1''/ / ' )78%  + +A . ::<<\\^F 8== ! K KfkkU&: K K|  MMU " K Kvfkk': K K|   5a 8 8'+ALLN $   /    q ! "   rc V^8dQhRRRR/#rDr)rs"rrr*$s 3 33 3 3rc VP4p\V4^8wd \R4hVP'd;VP V^,VPP 4V^,4M/VP V^,VP 4V^,4\ VP\4'd@VPP'd"VPVP44R#R#R#)rzExpected len(args) == 2N) rrrrr{rrrr_r sync_d2h_copyr)rr8rs&& rrGDeviceCopy.codegen*$s  " t9> !:; ;      ' 'Q));;=tAw   ' 'Q1G1G1I4PQ7 S dkk6 * *t{{/D/D/D  ! !$--/ 20E *rrN)rrrrr:rrGrrrrrr#s. . ` 3 3rrca]tRtRtRtRRltRRltRV3RlltR R lt] !R4RR R ll4t R Rlt Rt V;t #)DynamicSelectStorageOffseti8$a 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. cV^8dQhRR/#rr)rs"rr'DynamicSelectStorageOffset.__annotate__A$r1 rc \4#rrCrs&rr $DynamicSelectStorageOffset.get_readsA$rrcV^8dQhRR/#rsr)rs"rrrD$rqrc R#rrrs&rr' *DynamicSelectStorageOffset.should_allocateD$rrc4V^8dQhRRRRRRRRRRRR R R /#) runbacked_offset_symbol sympy.Symbolr base_offsetsympy.Symbol | intbase_dim_strider5r rrrr)rs"rrrG$sP ,(  ,  !  rc <\SV`R\\P!R4R7.4WnW nW0nW@nWPn W`n R#NrXr) rr/r~ rDr8rrrrr5r )rrrrrr5r rs&&&&&&&rr/#DynamicSelectStorageOffset.__init__G$sG ze1DErJ'=# &.  rcV^8dQhRR/#rr)rs"rrrZ$s99*B9rc .\VP.4#r)rDrrs&rr3DynamicSelectStorageOffset.get_unbacked_symbol_defsZ$s466788rc V^8dQhRRRR/#r#r)rs"rrr^$  ! ! rc \VPV4P\VPV44P\VPV44P\VP V44#r)r+rrr5rrr&s&&rr'/DynamicSelectStorageOffset.get_free_symbol_uses]$s] TZZ 7 U#DII}= > U#D$4$4mD E U#D$8$8-H I  rc V^8dQhRRRR/#rDr)rs"rrrh$sEE3EErc >VPWPR7R#)r N)codegen_dynamic_select_indexr rFs&&rrG"DynamicSelectStorageOffset.codegenh$s,,T,Dr)rrr rr5rrrrrrrSr r' r/rrer'rGrrrs@rrr8$sG&989 : EErrca]tRtRtRtRRltRRltRV3RlltR R lt] !R4RR R ll4t R Rlt Rt V;t #)DynamicSliceSizeil$a7 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. cV^8dQhRR/#rr)rs"rrDynamicSliceSize.__annotate__|$r1 rc \4#rrCrs&rr DynamicSliceSize.get_reads|$rrcV^8dQhRR/#rsr)rs"rrr/$rqrc R#rrrs&rr' DynamicSliceSize.should_allocate$rrc ,V^8dQhRRRRRRRRRR/#)runbacked_size_symbolrrrr rr5r)rs"rrr/$s<*"  !  ! rc <\SV`R\\P!R4R7.4WnW nW0nW@nWPn R#r) rr/r~ rDr8r6rr rr5)rr6rr rr5rs&&&&&&rr/DynamicSliceSize.__init__$s> ze1DErJ$8!   rcV^8dQhRR/#rr)rs"rrr/$s77*B7rc .\VP.4#r)rDr6rs&rr)DynamicSliceSize.get_unbacked_symbol_defs$s444566rc V^8dQhRRRR/#r#r)rs"rrr/$r$rc \VPV4P\VPV44P\VPV44P\VP V44#r)r+rrr r5rr&s&&rr'%DynamicSliceSize.get_free_symbol_uses$sY TZZ 7 U#DHHm< = U#DII}= > U#DII}= >  rc V^8dQhRRRR/#rDr)rs"rrr/$rOrc (VPV4R#r)codegen_dynamic_slice_sizerFs&&rrGDynamicSliceSize.codegen$rSr)r r5rrr6rr+rs@rr-r-l$sE  7./ 0 11rr-cba]tRtRtRtRRltRRltRV3RlltR R ltR R lt R t V;t #)ri$z3 The result of a call to aten._local_scalar_dense. cV^8dQhRR/#rr)rs"rrDynamicScalar.__annotate__$r1 rc \4#rrCrs&rr DynamicScalar.get_reads$rrcV^8dQhRR/#rsr)rs"rrrE$rqrc R#rrrs&rr' DynamicScalar.should_allocate$rrc(V^8dQhRRRRRRRR/#) rsymrkeypathzpytree.KeyPathrrrrr)rs"rrrE$s,*8@F rc <VP4\SV` R\\P !R4R7VP V.44WnW nR#r) rrr/r~ rDr8r rLrM)rrLrMrrs&&&&rr/DynamicScalar.__init__$sI   *ELL$78$:M:Mtf:U  rcV^8dQhRR/#rr)rs"rrrE$s&&*B&rc .\VP.4#r)rDrLrs&rr&DynamicScalar.get_unbacked_symbol_defs$s488*%%rc V^8dQhRRRR/#rDr)rs"rrrE$s--3--rc (VPV4R#r)codegen_dynamic_scalarrFs&&rrGDynamicScalar.codegen$s&&t,r)rMrL) rrrrrSr r' r/rrGrrrs@rrr$s+&--rrca]tRtRtRtRRltRRltRV3RlltR R lt] !R4RR R ll4t R Rlt Rt V;t #)ri$z- The result of a call to aten._assert_scalar cV^8dQhRR/#rr)rs"rrAssertScalar.__annotate__$r1 rc \4#rrCrs&rr AssertScalar.get_reads$rrcV^8dQhRR/#rsr)rs"rrrY$rqrc R#rrrs&rr' AssertScalar.should_allocate$rrc$V^8dQhRRRRRR/#)rscalarr|rrrrr)rs"rrrY$s!  | # $ rc ~<\SV`R\\P!R4R7.4WnW nR#r)rr/r~ rDr8r`r)rr`rrs&&&rr/AssertScalar.__init__$s3   ell51 2    rcV^8dQhRR/#rsr)rs"rrrY$rrc R#rRrrs&rrAssertScalar.has_side_effects$rrc V^8dQhRRRR/#r#r)rs"rrrY$s<<!< !> e?G    !J488*Tfgqfrrwx gg**AA eBG    4 5    7TXX7GqI J    19 :r)rr`r)rrrrrSr r' r/rrer'rGrrrs@rrr$sD  N+<,< ;;rrc,]tRtRt$R]R&R]R&RtR#)ExternKernelNodei$rrzexport_schema.NoderrNrrrrrrrr$s I rrrca]tRtRtRtR#RR/RV3RlllltRV3RlltR R ltR R ltR Rlt RRlt ] RRl4t RRlt RRltRRltRt]RRl4t] RRl4t]RRl4t]R R!l4tR"tV;t#)$r i%z 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. Nrc8V^8dQhRRRRRRRRR R R R R RRR/#rrrSrb rr rnontensor_argsr9 r rrr rr rrr)rs"rrFallbackKernel.__annotate__ %seN4N4N4N4& N4 & N4 + N4&N4EN4 N4rc <a\SS`V\V4\V4VR7RSnT;'g/Sn\ V\ PP\ PP34'g\RV R\V4 R24hVSn VSn Vf/MTSnSPf \R4h\ P"P%SP4.Sn.Sn\ SP\ PP4'dR#RSPP+49dR#SPP,p\ P.P0P3SP4'dcSP(P5V^,P744SP&P5V^,P744R#RR lp VP8'dL\;SP4'g1V !SP4'g\=R SP 24hSPSP>SP@4wrR V3R llp \ P.P0PCWV4F wrV !W4K R#) rtF#Fails to create FallbackKernel for rQ not supportedN,Expected self.python_kernel_name is not None_c10d_functionalc V^8dQhRRRR/#)rrrrrr)rs"rr-FallbackKernel.__init__..__annotate__U%s  '<  rc<\PPVP4\PPP 4;'gF\ VR4;'d2\PPP VP9#) py_kernels)rDr%_dispatch_has_kernel_for_dispatch_keyr DispatchKey Functionalizer,rrs&rhas_functionalize_impl7FallbackKernel.__init__..has_functionalize_implU%sr88AA 588//==L)HHHH((66"--G  rz'NYI: Can't generate FallbackKernel for c$V^8dQhRRRRRR/#)rinfoztorch._C.Argumentr rrrr)rs"rrr~h%s"/ // //@/ /s/ /t/ /rc<aa\SP\P4'd1\V\\ 34'g\ \V44h\P!SP4'd(\V\ \34'd \ R4hVfR#SPfR#RVV3RlloRV3Rllp\P!SP4'dVeVF pV!V4K R#R#\P!SP4'g \ R4hV!V4R#)z+Expected not isinstance(arg, (tuple, list))Nc V^8dQhRRRR/#)rrHrrrr)rs"rrSFallbackKernel.__init__..handle_aliasing_and_mutation..__annotate__x%s  V  rcD<SPPVP44SPf \ R4hSPP 'dASP P\\VP4R7VS44R#R#)Nz$Expected info.alias_info is not Noner) alias_namesr@r alias_inforis_writer& r' r~ rB)rHrrs&r add_aliasPFallbackKernel.__init__..handle_aliasing_and_mutation..add_aliasx%ss  '' 5??*()OPP??+++))00&z'H!TR,rc V^8dQhRRRR/#)rrHrrrr)rs"rrr%s  S T rc<VfR#\V\\.\O54'dR#\V\4'g\ \ V44hS!V4R#r)rr r"rArrr)rHrs&radd_alias_if_graph_buffer`FallbackKernel.__init__..handle_aliasing_and_mutation..add_alias_if_graph_buffer%sJ9a&$!:!:;;!!V,,(a11! rz5Expected library_utils.is_tensor_like_type(info.type)) rrrDListTyperrr library_utilsis_tensor_like_typeris_tensorlist_like_type)rr roptional_tensor_argrrsf& @rhandle_aliasing_and_mutation=FallbackKernel.__init__..handle_aliasing_and_mutationh%s$))U^^44!#e}55(c3300;;cE4=11()VWW{&    44TYY???/2+12EF03#%88CC(O*#.r)"rr/ruse_runtime_dispatchrrrDr+r,rarrr r rrr{r9 warn_fallbackrr}rr-_libraryrg mutates_and_returns_first_argr@r is_mutabler(rYrr zip_schema)rrrb r rvr rrschemarrrrr rsf&&&&&&$ rr/FallbackKernel.__init__ %s]   +  . !  %*!!2!8!8b UZZ**EJJ,J,JK  !5fXRV ~^\ ","Nb  " " * !OP P d556'))+ d&& (F(F G G  !1!1!6!6!8 8 !!)) >>   = =d>N>N O O    & &{1~'>'>'@ A    # #KN$;$;$= >      *4+;+;<<*4+;+;<<%9$:J:J9KL **4;;8J8JK / // /b--88vNID ( 3OrcV^8dQhRR/#r r)rs"rrrw%s  !8 rc D<\SV`4pVP\PP P Jd_VPFNp\V\4'gKVP\P!VP444pKP V#r)rr r rD_prims rng_primsgraphsafe_run_with_rng_staterrr with_readrNr r)rrr rs& rr FallbackKernel.get_read_writes%swg-/   u||55RR R))c>22"-"7"7$,,S\\^<#K* rc V^8dQhRRRR/#rDr)rs"rrrw%s  4H T rc  vVPVP4VP4\VRR44#rN)(codegen_unbacked_symbol_defs_for_outputsrcodegen_outputsrrFs&&rcodegen_unbacked_symbol_defs+FallbackKernel.codegen_unbacked_symbol_defs%s7?? MMO  " D-t 4  rcV^8dQhRR/#)rrr9 r)rs"rrrw%s%%%rc VP'd VP#\VP\4'dV.#VP#r)rM rrr_r& rs&rrFallbackKernel.codegen_outputs%s9 <<<<<  dkk6 * *6M$$$rcV^8dQhRR/#rr)rs"rrrw%   *B  rc \VRR4;p'dW\\PPP V4pVf \ R4h\VP44#\4#rNzExpected resolved is not None rr?r{r9r:r rrDr1rrresolveds& rr'FallbackKernel.get_unbacked_symbol_defs%e '.A4 H H  H0  **,=H$%DEEhmmo. .< rcV^8dQhRR/#rr)rs"rrrw%sirc \P!RR44p\VP4'g \ R4hVPUu.uFq!!VP 44NK ppVP W0P4wrE\PP'd\VP\PP4'dVP!WE4p\#VPP$P&V4UUu.uF8wrb\PP(P+W&P,4NK: pppM8VUu.uF,p\PP(P+V4NK. ppVP.P1V4V#uupiuuppiuupi)c.]tRtRt$R]R&RRltRtR#))FallbackKernel.codegen_args..Shimi%rrefcV^8dQhRR/#rr)rs"rr6FallbackKernel.codegen_args..Shim.__annotate__%s  # rc VP#r)rrs&rr92FallbackKernel.codegen_args..Shim.__repr__%s xxrrN)rrrrrr9rrrrShimr%sH  rrr)rT dataclassr rrrr rr{r9rTrr rDr+r,rrr-r.rrr0rr )rrrr rrparams& rrFallbackKernel.codegen_args%se         ,, !IJ J<@KKHKqtA//12K H**;8J8JK  77   :d.>.> @U@U#V#V..tF='2Gc$V^8dQhRRRRRR/#)rr Sequence[IRNode] | Noner r9 rrr)rs"rrrw%s$%%,%>K% %rc V'd+VUu.uFp\V\4'dKVNK upMRpV'dSV'g \R4hVUu.uF*qDP4'gKVP4NK, ppV^,#\V\P 4'd VP #\V\PP\34'g\V4'd\P !R4#\V\\34'd\RV44pVUu.uFqw'gK VNK pp\V4^8Xd V^,#V'gR#VFMp\V\P 4'g \R4h\VP 4'gKKVu# V^,#R#uupiuupiuupi)NzExpected tensor_argsrXc3N"TFp\PRV4xK R#5ir)r find_devicer s& rr-FallbackKernel.find_device..%s)$(A**433's#%z)Expected isinstance(device, torch.device))rr rrBrDr r8r ScriptObjectr,r/rrrDrrqr)r r rHnon_torch_bind_tensor_argsr devices device_setr8s&& rrFallbackKernel.find_device%s $ J 1:a+IQQ J # &$%;<<3>S;C..BR's~~';GS1:  nell 3 3!(( (  UXX224DE   0 0<<& & ntUm 4 4#$($J -7AJ&&vvJGA7|q qz!!!&%,,77()TUU&++&&!M " 1: C KTBs"GGG 0G  G%G%cV^8dQhRR/#rsr)rs"rrrw&s++$+rc 2^RIHpV!VP4#)r) is_impure)torch._library.utilsrr )rrs& rrFallbackKernel.has_side_effects&s2))**rcV^8dQhRR/#rCr)rs"rrrw &s$$m$rc @\VP\PP\PP 34'g0\ RVP R\VP4 R24h\VP\PP 4'gcRVPP49dDVPPP'd\VP4'd.#VP#)ryrQrzr|) rr rDr+r,rarrrr-rr(rrs&rrM+FallbackKernel.get_inputs_that_alias_output &s   uzz44ejj6T6TU  !5d6F6F5Gr(()*.: 4++UZZ-K-KLL"$*:*:*?*?*AA  ((333&t'7'788I## #rcV^8dQhRR/#rCr)rs"rrrw&s##M#rc d\VP4^8d \R4hVP#)rz&Expected len(self.mutation_names) <= 1)rr}rrs&rrE!FallbackKernel.get_mutation_names&s- t"" #a ' !IJ J"""rc  \PRVP4VP4\ V\ 4'g\ \V44hVPVPVP4wrVPW4pVPUu.uFpVP!V3/VBNK ppVPp\PP 'g.VOVO#\#R.4pVP%WQV4pRRlp\ V\&P(P*P,4'd,VP/V^,V^,4P0p MVP2P0p \5V 4^8Xd.VP74p V ^,P8p V!W4.p M9\;WP<4U Uu.uFwrV!V P8V4NK p p pVPf \ R4h/p\ VP\&P>P@4'd(\BPD!VP4eRVR&\GVP4\HPJ!VPPM4VV VR7R 7p\PNPQV4.VOVO#uupiuupp i) 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$V^8dQhRRRRRR/#)r return_typez1torch.TensorType | torch.ListType | torch.JitTyperIRNode | Sequence[IRNode]rzexport_schema.Argumentr)rs"rr>FallbackKernel.export_extern_kernel_node..__annotate__E&s+8 S8 SJ8 S-8 S$8 Src \V\P\P34'dTp\V\\ 34'd%\ V4^8wd \R4hV^,p\V\P4'de\V\4'g \R4h\PP\P!VP4R7R7#Ve \R4h\PPRR7#\V\P4'd\VP4\P4'd~\V\ 4'g\\#V44h\PPVUu.uF(p\P!VP4R7NK* upR7#\V\P$4'd\VP4\P4'dVf?\PP\P&PRR7R 7#\V\4'g \R 4h\PP\P&P\P!VP4R7R7R 7#\V\P(4'd!\PPVR 7#\+R \#V4 24huupi) rzExpected len(output) == 1z Expected isinstance(out, IRNode)r) as_tensorr T)as_none) as_tensors)as_optional_tensorz#Expected isinstance(output, IRNode))as_intzUnsupported return type )rrD TensorTypeNoneTyperrrrr export_schemar}rTensorArgumentrrgetElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError)rrr's&& rhandle_single_outputFFallbackKernel.export_extern_kernel_node..handle_single_outputE&s+(8(8%..'IJJftUm446{a',-HII )Ck5+;+;<<%c622,-OPP(1188"/">">CLLN"S9,-CDD(11888FFK88Z**,e.>.>>>"&(33(f66$--44$* #)C&44#,,.I#) 5 K););<<**,e.>.>BB>(1188+8+O+O+V+V$(,W,9 &ff55,-RSS(1188+8+O+O+V+V&3&B&B%+__%6',W,9K77$--44F4CC"%=d;>O=P#QRR9 s;.M%Expected self.op_overload is not None1torch_library_custom_op)rrrM ro )rr))r.r/rr rr rrr rrrrrr{r9aot_moder'serialize_inputsrDrrrrreturnsr-rrr0rrM r+r, custom_ops_maybe_get_opdefrrrrBrextern_kernel_nodesr@)rrrr ordered_kwargsr serializernamed_argumentsrrrM routput_arguments return_schemarro rs& rexport_extern_kernel_node(FallbackKernel.export_extern_kernel_node$&s B MMO    $// d, ,**4;;8J8JK **4899 9  ! !# 0 09  !!ww+T+N+ +*44 $55fFK8 St fe55??MM N NmmDGT!W5==Gnn,,G w<1 **,G!!*..K 4[ JK .1,,-G   .H)M %!++.H      # !HI I t'')>)> ? ?++D,<,<=I25H. /##'',,.&(!   $$T*''''e j s !K&K+c V^8dQhRRRR/#rDr)rs"rrrw&sX2X23X2X2rc  baaaSPpVf \R4hVPR8Xd\V\P P 4'g\\V44h\PP'd4^RI H p\V4V9d\PRV4RSnMVPR8Xd@\V\P P 4'g\\V44hMB\PP'd#V\"P$P&9Sn\PP'Ed\V\P P 4'dSP 'gRV3R lloSP)SP*SP,4wpo\.P0!VVV3R lSP244p\4;QJd<V3R l\7WRP8P:44F 'gK RM1 R M-!V3R l\7WRP8P:444SnSP=V4SP 'dSP?4pSP@f \R 4hSPf \R4hVPCSPE4SP@V3RlSPVSPG44MVPIS4SPKV4\SPL\N4'd6SPQV4SPSV4SPUV4R#R#)r$NzExpected kernel is not NonerLrNzG%s is missing a c-shim implementation, using proxy executor as fallbackT _quantizedc V^8dQhRRRR/#)rrHrrrr)rs"rr,FallbackKernel.codegen..__annotate__&s 7 7] 7t 7rc<\V\P4'dS!VP44#\V\P4#r)rrDrr NumberType)rH is_numbers&rr)FallbackKernel.codegen..is_number&s=a!3!344$Q%5%5%788!!U%5%566rc3J<"TFpSP!V3/SBxK R#5ir)r)rrrrs& rr)FallbackKernel.codegen..&s(?))!6v6?s #c3|<"TF1wr\V\4;'dS!VP4xK3 R#5ir)rcomplexr0)rr rrs& rrr&s5,DDA1g&AA9Q[[+AADs<<Fr{rcH<.SP4OSP4O#r)rrrsrr(FallbackKernel.codegen..&s"F$++-F0C0C0EFr)+r rrUrrDr+r,rr{r9rTrXrOrr.r rrMrcustom_ops_to_c_shimsr rrrPr rrrr-r.rArr,generate_fallback_kernel_with_runtime_lookuprrgenerate_fallback_kernelrrr_r$r*r1) rr8rb rOr args_iter exported_argsrrs f& @@rrGFallbackKernel.codegen&s!! > !>? ?   v %fejj&;&;<<$T&\22ww"""Lv;&;;KKa15D-    -fejj&;&;<<$T&\22= WW f11GGG  % GG   65::#8#899--- 7 7 ..t{{D>+C+CD,, >>+C+CD,)D % W%  $ $ $ ::SP=V4p \?V \AS4\AS4S\CV 4V R7#\D;QJdRS4F 'gK RM RM !RS44oSPGSV4oS'gs\S\PHPJPL4'g)S\PNPPPRJd\PT!R4oRVVVVVVVVV3 R llp\S\P P"4'd|\P&P(P+S4'gM\P PV'g-S'd%\XP[SVSSSSSSSR 7 pVeV#VfS!\]SR 7SSSSSSR 7oMT\V\P,4'd V!V4#S'g \_R 4hS!\aSR 7SSSSSSR 7oRVVVVV3RlloS!V.4p\V\@\b34'd VSn2V#\V\f4'd\cV4Sn2V#V.Sn2V# +'giEL;i)z9Create an instance of FallbackKernel from an _OpOverloadsN)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variant)rrrrrr c38"TFp\V4xK R#5ir)rer s& rr(FallbackKernel.create..'s!K{,s"3"3{rTFrXc V^8dQhRRRR/#)rrr0rr r)rs"rr+FallbackKernel.create..__annotate__'s   . rc 6< S'g \R4hS!SPV4SSSS SS R7p\P'gS'g\ V4'g8\ P PPVP44V#)"Not sure where to find device inforr) rr rM assume_unaligned_fallback_outputrwr{r9rgr-r) rpackedrr8has_unaligned_inputrb rr r rr s & rcreate_direct_output3FallbackKernel.create..create_direct_output's|$%IJJ$$V,"3F777&(00))--foo.?@Mr)rrArr>r=c$V^8dQhRRRRRR/#)rrrrzlist[tuple[Any, int]]rr)rs"rrr;'s"* * C* 2G* C* rc <aa\S\\34'd0\S4!VVV3Rl\ \ S4444#\S\ 4'd@SP4UUu/uF"wr#VS!VS\S4V3.,4bK$ upp#\S\P4'd\SPS4S S4p\P'gS'g\S4'g4\P P"P%VP&4V#\S\(4'dS#\S\P*4'dSP,P.#\S\P0P2\434'g\7S4'd\9\;SR7S SS4#Se\=R\S4 R24hR#uuppi)c3j<"TF(pS!SV,S\S4V3.,4xK* R#5ir)r)rrgenerate_outputrrs& rrAFallbackKernel.create..generate_output..'s7$/$F1Iw4<:K9L/LMM/s03rNzFallbackKernel output type z is not supported)rrrrrrrr rDr rr r rMr?rwr{r9rgr-rrrrr]rrr,r/rr~ r) rrr rfrrr8rGrAr@s ff rrG.FallbackKernel.create..generate_output's&4-00F|$"3v;/$FD))%+LLN$2g$v,9L8M.MNN$2FELL11!((0 ;;;*,V44GG--11#((; FC(( FELL11{{'''..0@A$V,,(f- %(5d6l^CTUGs8(H )4rL*_fused_moving_avg_obs_fq_helper_functionalrr r{r9r r rr r r r rrrDr+r,r0rrg is_builtinr torch._library._out_variantr4r5r6r7r-rr rLr_make_out_variant_kernel_namerrrrrry higher_orderprintr8rTExternKernelMultiOut try_creater~ rr. rrM r)rrb rrfake_incorrect_kernelscontextrr r4r5r6r7out_oprrB out_resultrM r8rGrAr r@r rr sff*l @@@@@@@@rrFallbackKernel.createI's#'"Q"Q!S / /1$79J9JKG!mG W''@@@F..((  00.."44 fejj33 4 4  , ,V Ed Ef E vuzz44 5 5NN((33F;;>5<<88  Ffnn--'/~26:!c*;F*C&D&I--n=&! ,"&"7!'DV'L & "c!K{!Kccc!K{!KKn= vu66@@NN O O//555\\%(F   , vuzz44 5 5NN((33F;;GG'''-88"3$79 J%!!  !&)"3F 5 5'7 7$%IJJ!0"3F* * X"."5 ge} - -$FN   & &"7^FN&YFNaWWs (Q"" Q3 )rrr}r rr rr)rrrrrSr/r rrrrrVrrrMrErrrGr r:r0rrrrs@rr r %s N4HLN4N4`   %  2%%N+$,# E(NX2X2t    00dxxrr cVa]tRtRtRtRRltRRltRRR R/R V3R llltR tV;t #) ComplexViewi(z9View a complex number as two dtyped numbers or vice versacV^8dQhRR/#rsr)rs"rrComplexView.__annotate__ (rqrc R#rrrs&rr' ComplexView.should_allocate (rrcV^8dQhRR/#rCr)rs"rrrZ (s$$m$rc &VP^4.#rrlrs&rrM(ComplexView.get_inputs_that_alias_output (s"##rrNrc8V^8dQhRRRRRRRRR R R R R RRR/#rur)rs"rrrZ(s[    &  &  +  & E   rc  4<\SV`VVVVVVVR7R#)r>N)rr/ rrrb r rvr rrrs &&&&&&$$rr/ComplexView.__init__(s,      /  rr) rrrrrSr' rMr/rrrs@rrXrX(s5C$ )- HL   rrXc&]tRtRtRtRRltRtR#)MemoryCheckKerneli&(z 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 V^8dQhRRRR/#rDr)rs"rrMemoryCheckKernel.__annotate__-(s  3  rc VP4VPwr#p\V4p\V4pV'dVPR4RV RV RV R2pM RV RV R2pVPV4R#)z.Override codegen to write direct function callzW# note: don't currently distinguish between buffers returned and dealloc'd in last stepzcheck_memory_step(allocated=z, freed=z, is_final_step=rRN)r0rrcr)rr8 alive_list dead_list is_final_step alive_repr dead_reprcalls&& rrGMemoryCheckKernel.codegen-(s 224/3/A/A, }*% O    i 2*Xi[P`an`oopqD1*Xi[PQRD$rrN)rrrrrSrGrrrrrere&(s   rrec.]tRtRt$R]R&RRltRtR#)r. i?(rr8cV^8dQhRR/#rr)rs"rrMultiOutputLayout.__annotate__C(rrc VP#rrrs&rrBMultiOutputLayout.get_deviceC(rrrN)rrrrrrBrrrrr. r. ?(s rr. ca]tRtRtRRltRRV3Rlllt]!R4RRRll4tRR ltR R lt R R lt Rt V;t #)rr iG(c V^8dQhRRRR/#rDr)rs"rrMultiOutput.__annotate__H(s44344rc VPV4VP'g%VPV4VPV4R#R#r)codegen_multi_output!skip_size_stride_alignment_checksr$r*rFs&&rrGMultiOutput.codegenH(s=$$T*555  % %g .  * *7 36rc ,V^8dQhRRRRRRRRR R /#) rrrSr rrlist[tuple[Any, ...]]rzrrrr)rs"rrrwN(sC S S S S' S ,0 S  Src <\SV`RW.R4\PP V4Vn\PP V4W0nW@nR#r( ) rr/r{r9r! rr" rrz)rrr rrzrs&&&&&rr/MultiOutput.__init__N(sK vw3GG++D1  ""4( 1R.rc V^8dQhRRRR/#r#r)rs"rrrw\(s>>!> !>rc VP^,p\V\4'g \V4hVP V4#r)rrrrr')rrrs&& rr' MultiOutput.get_free_symbol_uses[(s:[[^ *f-- , ,..}==rcV^8dQhRR/#rsr)rs"rrrwd(s   rc \VP4^8H;'d"\VP^,\4#r)rrrrI rs&rr' MultiOutput.should_allocated(s54;;1$  t{{1~'8 9 rcV^8dQhRR/#rCr)rs"rrrwi(s  m rc VPUu.uFKp\V\4'gK\VP 44^8gK;VP 4NKM up#uupir)rrr rrMrrs& rrM(MultiOutput.get_inputs_that_alias_outputi(s\{{ "#~. C4467!; CLLN"   sA#A# A#cV^8dQhRR/#r r)rs"rrrwq(s! ! !8! rc aa\4pVPFNp\V\4'gKVP \ P !VP444KP VP4oVP4P4oRVV3RllpVP4p\P'*;'g'VRJ;'g\VP4'*p\ P!W0P!4RVR7p\ P"!VVP$\4R7#)c$V^8dQhRRRRRR/#r8 r)rs"rr1MultiOutput.get_read_writes..__annotate__(rH rcv<\V4^8wd \R4h\P!SS!V4R4#rJ rL rM s&&rr> *MultiOutput.get_read_writes..dummy(rO rNr2 r r)rDrrrr-rNr rr>rrBrMr rqrrWr<r r ) rrrr> r8should_normalizewrite_rwrVrs & @@rr MultiOutput.get_read_writesq(s/9l;;C#v&& ,..s||~>?}}//#002 ; ; "11 1 ' '~ ' '&++&&   33 ==?B2B &&??"   r)rrrzr) rrrrrGr/rer'r' rMr rrrs@rrr rr G(sA4 S SM*>+>  ! ! rrr c`a]tRtRtRtRV3Rllt]RRl4tRRltR R lt R t V;t #) ri(zMultiOutput for opaque objects.c ,V^8dQhRRRRRRRRR R /#) rrrSr rrr} opaque_valuerrrr)rs"rrOpaqueMultiOutput.__annotate__(s<1111' 1  1  1rc 8<\SV`WVRR7W@nR#)TrzN)rr/r)rrr rrrs&&&&&rr/OpaqueMultiOutput.__init__(s  SWX$0!rcV^8dQhRR/#)rrrr)rs"rrr(s??u?rc \R4h)zOpaqueMultiOutput has no dtype)rrs&rr7OpaqueMultiOutput.dtype(s=>>rcV^8dQhRR/#)rrrr)rs"rrr(s#4rc V#rrrs&rr#OpaqueMultiOutput.wrap_for_lowering(r rcV^8dQhRR/#r r)rs"rrr(s   !8  rc p\4pVPFNp\V\4'gKVP \ P !VP444KP \\ P !VP44.4p\ P!VV\4R7#)r ) rDrrrr-rNr rr )rrrr s& rr !OpaqueMultiOutput.get_read_writes(s.8l;;C#v&& ,..s||~>?0:  ! !$--/ 2 30 &&"   r)r) rrrrrSr/rYr7rr rrrs@rrr(s3)11??   rrc2]tRtRtRtRRltRRltRtR#) AllocatingMultiOutputi(zMultiOutput 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. cV^8dQhRR/#rsr)rs"rr"AllocatingMultiOutput.__annotate__(r rc R#rRrrs&rr' %AllocatingMultiOutput.should_allocate(rrc V^8dQhRRRR/#rDr)rs"rrr(s44344rc rVP'g%VPV4VPV4R#R#r)rzr$r*rFs&&rrGAllocatingMultiOutput.codegen(s/555  % %g .  * *7 36rrN)rrrrrSr' rGrrrrrr(s 44rrc V^8dQhRRRR/#)rrTrrrr)rs"rrr(s22*?2C2rcVPpVPPPR4^,pVPpRV RV RV 2#)z8Build fully-qualified kernel name for an out-variant op.z::z torch.ops.rM)rUr-rrArV)rTnsrrs& rrMrM(sP   Bnn!!''-a0G##H t1WIQxj 11rcva]tRtRt$RtR]R&R]R&RV3RlltR R lt]R R R R/RRll4t Rt V;t #)rPi(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_nodesc ,V^8dQhRRRRRRRRRR /#) rrrrTrrrrrrr)rs"rr!ExternKernelMultiOut.__annotate__(s< " " "& "! "  "  "rc n<\SV`!V/VBW n.Vn\ V4VnWnR#r)rr/rrrMrr )rrTrrrrs&$$*,rr/ExternKernelMultiOut.__init__(s8 $)&)*(*%"?"G!rc V^8dQhRRRR/#rDr)rs"rrr(s77377rc JVPV4VPV4R#r)rA generate_extern_kernel_multi_outrFs&&rrGExternKernelMultiOut.codegen(s W%006rrNrAFc@V^8dQhRRRRRRRRR R R R R RRRRRRR/ #)rrb rr rr8rr rr r9 r rrr rr rArrz&Sequence[AllocatingMultiOutput] | Noner)rs"rrr(s{JJ%JJ J & J ' J+J&JEJ"J 0Jrc  ^RIHp Hp Hp V !VP4'gR#\ V\ \34'gR#V !V4p V fR#V !V 4p\;QJdRV4F 'dK RM RM !RV44'gR#\V4\V48wdR#V!\VR7VVVVVVV VR7 p.p\V4Fwpp\VPVP.VPO.VP!4OR7p\#VV\%V4V3.R 7p\&P('gV 'g\+V4'g4\,P.P0P3VP44VP7V4K VVnVVn\ V\ 4'd \ V4#\V4#) zGCreate an ExternKernelMultiOut if the op has a matching .out() variant.)r4r5r7Nc3V"TFp\V\P4xK! R#5ir)rrDr r`s& rr2ExternKernelMultiOut.try_create.. )sG1:a..s')FTr)rrrTrr8r7r5r6)rr r)rLr4r5r7r-rrrrrr. rrr8r7r{r6rrrMr?rwr{r9rgr-rr@rrM )rrb r r8r r r rrrAr4r5r7rTrr@rM r tensor_outr multi_outs&&&&&&&&$$ rrQExternKernelMultiOut.try_create(s  fnn--.5$-88' >)&1 sGGsssGGGG ~ #m"4 4 V ,    /'  02&~6MAz !(( &&(z''(-**,- F .~.23I 777&(44))--inn= NN9 %%7(+2'  ne , ,> !G}r)r rrr) rrrrrSrr/rGr:rQrrrs@rrPrP(sS 99 " "7JHLJ%*JJrrPcD]tRtRt$RtR]R&RRltRRltR R ltR R lt R Rlt RRlt RWRRllt RR/RRllt RRltRRltRRltRRltRXR R!lltR"R#ltR$R%ltRXR&R'lltR(R)ltR*R+ltR,R-ltR.R/ltR0R1ltR2R3ltR4R5ltR6R7ltR8R9ltR:R;ltR<R=lt R>R?lt!R@RAlt"]#!R4RXRBRCll4t$RDRElt%RFRGlt&RWRHRIllt'](RJRKl4t)RLRMlt*RNROlt+RPRQlt,](RRRSl4t-RTRUlt.].t/RVt0R#)Yr i;)z; TensorBox / StorageBox allow in-place mutation of Tensors rrcV^8dQhRR/#rsr)rs"rrMutableBox.__annotate__C)rrc 6VPP4#rrrs&rr!MutableBox.has_exceeded_max_readsC)rrcV^8dQhRR/#rr)rs"rrrF)rrc 6VPP4#rrrs&rrBMutableBox.get_deviceF)rrcV^8dQhRR/#rr)rs"rrrI)s''A'rc 6VPP4#rrrs&rrMutableBox.make_loaderI)yy$$&&rcV^8dQhRR/#rr)rs"rrrL)s((>(rc 6VPP4#r)rrrs&rrMutableBox.make_indexerL)yy%%''rcV^8dQhRR/#rxr)rs"rrrO)s&&.&rc 6VPP4#r)rrnrs&rrnMutableBox.get_strideO)rrcV^8dQhRR/#rr)rs"rrrR)rrc 6VPP4#rrrs&rrMutableBox.get_nameR)rrNc V^8dQhRRRR/#rr)rs"rrrU)s77J7$7rc 8VPPV4#r)rrrs&&rrMutableBox.has_large_inner_fnU)syy++I66rrTc$V^8dQhRRRRRR/#rr)rs"rrrX)rrc :VPPWR7#rrrs&&$rrMutableBox.mark_reuseX)rrcV^8dQhRR/#rr)rs"rrr[)s((d(rc 6VPP4#rrrs&rrMutableBox.realize_hint[)rrcV^8dQhRR/#rr)rs"rrr^)s''V'rc 6VPP4#r)rrrs&rrMutableBox.unwrap_view^)rrcV^8dQhRR/#rsr)rs"rrra)s+++rc 6VPP4#r)rrrs&rrMutableBox.is_input_buffera)yy((**rcV^8dQhRR/#rr)rs"rrrd)s))t)rc 6VPP4#r)rrrs&rrMutableBox.freeze_layoutd)syy&&((rc$V^8dQhRRRRRR/#rr)rs"rrrg)s)OO"O37O Orc 8VPPW4#r)rrrs&&&rr*MutableBox.freeze_layout_with_stride_orderg)syy88NNrc V^8dQhRRRR/#rr)rs"rrrl)s>>=>T>rc 8VPPV4#r)rrrs&&rr(MutableBox.freeze_layout_with_fill_orderl)syy66u==rc V^8dQhRRRR/#rr)rs"rrro)s??4F?4?rc 8VPPV4#r)rrrs&&rr(MutableBox.freeze_layout_with_same_ordero)syy66v>>rc$V^8dQhRRRRRR/#rr)rs"rrrr)s)XX/X@DX Xrc 8VPPW4#r)rrrs&&&rr+MutableBox.freeze_layout_with_exact_stridesr)syy99-WWrcV^8dQhRR/#r r)rs"rrrw)s++!8+rc 6VPP4#r)rr rs&rr MutableBox.get_read_writesw)rrcV^8dQhRR/#rr)rs"rrrz)r rc 6VPP4#rr rs&rr MutableBox.get_readsz)rrcV^8dQhRR/#rr)rs"rrr})rrc 6VPP4#rr rs&rrMutableBox.num_reads})rrcV^8dQhRR/#rr)rs"rrr)rrc 6VPP4#rrrs&rr MutableBox.get_storage_numel)rrcV^8dQhRR/#rr)rs"rrr)r rc 6VPP4#rr rs&rr+MutableBox.get_reduction_type)r rcV^8dQhRR/#r.r)rs"rrr)r rc 6VPP4#rr rs&rr1MutableBox.get_reduction_size)r rcV^8dQhRR/#rsr)rs"rrr)rrc 6VPP4#rrrs&rr6MutableBox.is_extern)rrcV^8dQhRR/#rsr)rs"rrr)s$$$$rc 6VPP4#r)rr;rs&rr;MutableBox.is_no_op)rrc V^8dQhRRRR/#r>r)rs"rrr)s444&4rc 8VPPV4#rr rs&&rr@MutableBox.constant_to_device)syy++F33rcV^8dQhRR/#rCr)rs"rrr)s..M.rc 6VPP4#r)rrErs&rrEMutableBox.get_mutation_names)r rcV^8dQhRR/#rr)rs"rrr)s..C.rc 6VPP4#r)rrIrs&rrIMutableBox.get_operation_name)r rcV^8dQhRR/#rCr)rs"rrr)rrc 6VPP4#r)rrMrs&rrM'MutableBox.get_inputs_that_alias_output)syy5577rcV^8dQhRR/#rr)rs"rrr)rrc 6VPP4#rrrs&rrMutableBox.realize)rrc V^8dQhRRRR/#r#r)rs"rrr)s==!= !=rc 8VPPV4#rrmr&s&&rr'MutableBox.get_free_symbol_uses)syy--m<MutableBox.get_layout)rrcV^8dQhRR/#rjr)rs"rrr)s+++rc 6VPP4#rr+rs&rrlMutableBox.get_output_spec)rrcV^8dQhRR/#r.r)rs"rrr)rrc 6VPP4#rrrs&rr<MutableBox.get_size)rrcV^8dQhRR/#rQr)rs"rrr)s{rc .VPP#r)rr7rs&rr7MutableBox.dtype)syyrcV^8dQhRR/#rr)rs"rrr)s   rc \VP\4'dR\V4P R\VP4P R2pRpVPPpM&\V4P R2pVPpRpV\ \ V44V.pRPV4#)rIz))rRr>)rrr rrrKrrL)rline0endlr{rFs& rrMutableBox.__str__)s dii , ,Dz**+1T$))_-E-E,FaHEDIINNEDz**+1-EIIED  3u:    yyrrrr)1rrrrrSrrrBrrrnrrrrrrrrrrrr r rr r+r1r6r;r@rErIrMrrer'r rrrYrr>rlr<r7rr9rrrrr r ;)s L2&'(&$7DDD('+)O >?X +%%-..%$4..8#L)=*= *+3++&+$ "Hrr cl]tRtRt]]RRl44t]]RRl44t]RRl4tRtR #) ri)c V^8dQhRRRR/#)rrrrr)rs"rrTensorBox.__annotate__)sII*I/DIrc R#rrrs&rrTensorBox.create)sFIrc V^8dQhRRRR/#)rrrrrr)rs"rrr@)s..V. .rc R#rrrBs&rrrC)s+.rcV^8dQhRR/#)rrrr)rs"rrr@)s++V+rc Z\V\4'dV#\\V44#r)rrrrrBs&rrrC)s% d1 2 2KD)**rrN)rrrrrrVrrrrrrr)s@ II ..++rc]tRtRtRtRRltRRltRRltR R ltR R lt R Rlt RR/RRllt RR/RRllt RRlt RtR#)ri)z/ StorageBox allow in-place mutation of Tensors cV^8dQhRR/#rsr)rs"rrStorageBox.__annotate__)src \VP\\34'd6VPP 4\ P P9#R#r)rrrrrr{r9rrs&rrStorageBox.is_input_buffer)s= dii+!? @ @99%%'177+?+?? ?rcV^8dQhRR/#rsr)rs"rrrJ)s  $ rc \VP\4;'d6VPP4\P P 9#r)rrrhrr{r9rrs&rrStorageBox.is_module_buffer)s> tyy> 3 : : ""$(9(99 rcV^8dQhRR/#rr)rs"rrrJ)src  \PVP4'dVPP4#\ VP\ \ \\34'g\\VP44hVPP4pVPP4pVPP4pVf \R4h\R\VVPP!4VPP#4RR7VPR7Vn\$P&P)VP4VPn\$P&P-VP4VP.VPnWPnW PnVPPP4VPnVPPP6VPnVPP*#)NrF)r8r7r5rr)rrrrrrDrrr rrrrrBrr?rAr<r{r9r! rr" rrrrr)rrrr8s& rrStorageBox.realize)s  " "499 - -99%%' '$))iD$%GHH dii1 1ii//1 II++- %%' > !>? ?"!ii))+YY'')     00;  ""499- LL  + ' #yy~~88  IINN22 yy~~rcV^8dQhRR/#rr)rs"rrrJ*sdrc \VP\\34'd>VPP 4P ^8dVP 4R#R#R#)z< Called on buffers we expect to be forced to realize later. N)rrrDrrnontrivial_read_countrrs&rrStorageBox.realize_hint*sI tyy9i"8 9 9 **,BBQF LLNG :rc V^8dQhRRRR/#)rrrrrr)rs"rrrJ*s  c d rc L^RIHpVP4Uu.uF2pV!V4'dK\PP V4NK4 ppV'gR#\ V4p\V4p\V4pWQ8;'dWV, ^8;'dWg8H#uupi)r)is_nonfreeable_buffersF) r rYr r{r9get_dep_size_hintrrr{)rrrYr size_of_reads total_sizemax_sizemin_sizes&& r$has_accumulated_enough_reads_by_size/StorageBox.has_accumulated_enough_reads_by_size*s@~~' ')#. +AGG % %c *'  ' }%}%  # % %%* % %$  s B!#B!cV^8dQhRR/#rsr)rs"rrrJ.*s   rc \PpVf>\VP4'd\PpM>\P pM-\ V\4'g\R\V4 24h\ VP\4;'dmVP4V8;'gRVP4;'g:\PRJ;'d VP\P4#)Nz.expected int realize_acc_reads_threshold, got )rMrealize_acc_reads_thresholdrdrrealize_cpu_acc_reads_threshold$_realize_acc_reads_threshold_defaultrrrrrDrr realize_acc_reads_size_thresholdr_)rrcs& rr!StorageBox.has_exceeded_max_reads.*s&,&H&H# & .dii  .4.T.T+??,93??$DTJeEfDgh$))Y/   NN : :  &&(  77tC==;; rrTc$V^8dQhRRRRRR/#)rrrrrrr)rs"rrrJG*s"!!S!$!RV!rc aV^8Edt\VP\\34'EdMVPP 4oRSP 9dR#\ VP4'd.ROp\;QJdV3RlV4F 'gK RM RM!V3RlV44'dR#VPP4p\VP\4'dBV'd:V\P8d%SP\^V^, 48dR#VP4'dR#T;'d"VP4\P8#R#)zR A heuristic to decide if we should realize a tensor that is used multiple times. inline_asm_elementwiseTc3@<"TFqSP9xK R#5ir)used_ops)rropcounts& rr5StorageBox.should_realize_on_reuse..[*s@iG,,,isF)expr.log10log1plog2sigmoidtanh)rrrDrrrlrdrrrMrealize_opusers_thresholdrrrrrealize_reads_threshold)rrr heavy_opsrealize_thresholdrms&&$ @rshould_realize_on_reuse"StorageBox.should_realize_on_reuseG*s 19DII 9/EFFii002G'7+;+;;dii   3@i@333@i@@@$(II$K$K$M!tyy)44# @ @@#a1BQ1F*GG&&((TT4>>#3f6T6T#T Trc$V^8dQhRRRRRR/#rr)rs"rrrJj*s!TTrc ZVPWR7'dVP4R#R#)rN)ryrrs&&$rrStorageBox.mark_reusej*s%  ' ' ' G G LLN HrcV^8dQhRR/#rr)rs"rrrJn*rrc 6VPP4#rr rs&rrStorageBox.num_readsn*rrrBN)rrrrrSrrrrr_rryrrrrrrrr)sH  > & 2!!FD%%rrcH]tRtRt$R]R&R]R&RtR]R&RtR ]R &R tR#) Subgraphir*rrr graph_moduleNr inductor_config_patcheszGraphLowering | Noner9r)rrrrrrr9rrrrrrr*s$ I&&5929"&E &rrc V^8dQhRRRR/#)rbuffersrrrr)rs"rrrz*sLL"2LtLrcVUu.uF+p\V\4'dVP4MTNK- pp\\ RV444\V48#uupi)c38"TFp\V4xK R#5ir)r )rr^s& rr'_has_aliased_buffers..*s;7"V**7r)rrrrrD)rr^s& r_has_aliased_buffersrz*sdF!+6? C CO  z;7;; VP P R/4P R4p \/V RR4VnVPfVP,'d!\0P2!VP,4M \54pT;_uu_4\PP7VP8VVP:R 7Vn\P<!VP4;_uu_4VPP>!V!RRR4RRR4VPP@pRpVF+p \#V \$4'dKV PC4pM Vf \R 4h\EVV\GVR 7R 7oR V3Rllp\!V4UUu.uFwppV!VV4NK pppVSn$V#uup iuup i +'giL;i +'giL;iuuppi)z|For each operand, get a realized input, force it to have the same strides as the subgraph inputs, then use an InvokeSubgraph)constrain_to_fake_tensorNeager_input_valszGExpected current_node.args[1] is torch.ops.higher_order.invoke_subgraphrfrnested_region_configrrrrrr)rrrc$V^8dQhRRRRRR/#)rrrindrrz:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutputr)rs"rr+InvokeSubgraph.create..__annotate__*s$   !$  G rc |<\V\\34'dV#VP4pVf \ R4h\ \ VVP4VP4VP4VP4PVP4PR7S\V3.RR7#)Nrrz Tr)rrrL rBrrr rrAr<rnr>rrr)rrr8invoke_subgraphs&& r create_output,InvokeSubgraph.create..create_output*s:N"OPP **,>()FGG"%$..0#__.%002%00299"("3"3"5"?"? $C[M6:  r)%rrr{r9r rr1rrDryrNrrrrrs rrrrrr@rrrMr!r rrrrr graph_outputsrBrr. rM )rrrrr fake_operandsrr fx_operandsr new_operandsroperand nested_configrrM r8rrroutsrs&&* @rrInvokeSubgraph.create*su 7ww++  +00445GH H  HF""eii&<&<&I&II$$Q'uyy/E/E/U/UU(a,Q/8MF""eii&<&<&I&II$$Q'uyy/E/E/U/UU(a'++FG4K4?@KqVVE]]KM@ AI!I1"3"3A"6!I%' %h/LC/ARS##G,##,WC6HI 0   + + 3(--11(B?CC&M078$0H , >> !333 X==> ]  !"!6!6,,#0"*--"7" ((88NN&& 69 ....Gg'<== ++-  > !>? ?($F3    0;DG:LM:LYQ fa(:LM"& sA "JJ98 dNs7O<O A$O4O! O46P! O1 ,O44 P c V^8dQhRRRR/#rDr)rs"rrr+rrc (VPV4R#r)codegen_invoke_subgraphrFs&&rrGInvokeSubgraph.codegen+rr)rr)rrrrrSrrrrM r/r$r:rrGrrrs@rrr*sZ!%Ho$(,H%,'+G $+ ) )8|||..rrc V^8dQhRRRR/#)rrzint | torch.SymIntrz int | Exprr)rs"rrr+s%*rc^\V\4'dV#VPP#)zLConvert an int or SymInt to a plain int or sympy Expr for use in IR layouts.)rrrr])rs&r _maybe_exprr+s"!S 66;;rca]tRtRt$RtRtR]R&RtR]R&RtR]R &R t R ]R &Rt R ]R&RRV3Rlllt RRlt ] RRRll4tRRltRRltRtV;t#)Switchi&+a IR node representing torch.cond and torch.switch. For torch.cond: ``selector`` is a boolean scalar tensor, branches contains exactly two subgraphs (true branch, false branch), and is_cond=True. For torch.switch: ``selector`` is an integer scalar tensor, branches contains one subgraph per case, and is_cond=False. Attributes: selector: A scalar tensor selecting which branch to execute. branches: Non-empty sequence of subgraphs representing the individual branches. operands: Input tensors passed to the individual subgraphs. is_cond: True when this node represents a torch.cond (boolean 2-branch select). outputs: MultiOutput nodes representing the node's outputs. Nr3selectorzSequence[Subgraph] | NonebranchesrrFris_condSequence[MultiOutput] | NonerM c4V^8dQhRRRRRRRRR R R R R R/#)rrrrzSequence[Subgraph]rrrr. rr rrrrr)rs"rrSwitch.__annotate__>+sQ)))%)# ) " ) E )) )rc <WnW nW0nW`n\ V.VO4wrx\ S V`RVVVR7VeWPn\PPV4Vn \PPV4R#N)rrrr) rrrr_split_by_sym_typerr/rr{r9r! rr" ) rrrrrrrsym_argsr rs &&&&&&& rr/Switch.__init__>+s!    2H3Hx3H I "   (%6 "GG++D1  ""4(rcV^8dQhRR/#r!r)rs"rrrZ+sHH~Hrc LVPe\VP4#.#r)rrrs&rr$Switch.get_subgraphsZ+s&*mm&?tDMM"GRGrc ,V^8dQhRRRRRRRRR R /#) rrrrr"rr~ rrrzlist[MultiOutput]r)rs"rrr^+sClll!l" l  l  lrc VPV4pVUu.uFqPPV4NK pp\PPPR,p\ V\ 4'g\\V44h.pVFNp\ V\4'd%VPVPR,4K=VPV4KP \PPPR,p RRlp VEF_p V PeK\PPV PVV PR7V nV PPP4P^,p \ V \ 4'g\\V 44hV U u.uF.p \ V \4'dV PR,MT NK0 pp \P !V P4;_uu_4V PP"!V!V !V PP$W4V PnRRR4EKb \&;QJdRV4F 'gK RM RM !RV44'd \R 4hVUu.uFqPP$NK ppV'dR MR p\)VV4F8wpp\+V4'gK\R V R VP RV RV 24h V^,pVR,Fp\-V4\-V48wd\VV34h\/\)VV44FwpwppVP14VP148wd\VVV34hVP34VP348wd\VVV34hVP54P6VP54P68wgK\VVV34h K \9RW1.,44p\;\PP<P>\PPPPARR44pVf \R4h\CVVV\EVR7VVR7p\/\)V\PPPR,44UUUUu.uFwpwpp\G\IVP14eVP14MTVP34VPK4Uu.uFp\MV4NK upVPO4Uu.uFp\MV4NK upVP54P6VP54PPR7V\RV3.4NK pppppVVn*^RI+H,p\[4pVF-pV!VP4wp p!p V\[V!4,pK/ \]V4U"u.uF(p"\_VV",P`VV",V4NK* up"Vn1V#uupiuup i +'giEK;iuupiuupiuupiuuppppiuup"i)z:Create a sequence of IRNodes from a switch/cond statement.rfc(V^8dQhRRRRRRRR/#) rrr fake_tensorszSequence[torch.Tensor] branch_fakesz#Sequence[torch.Tensor | int | None]rrr)rs"rr#Switch.create..__annotate__z+s0  + 0 >   rc J.p\WV4FwrEp\V\4'dVPV4K/VP 4p\ V4'dVP 4pVP\ P\V4VRR74K V#)Fr]) rrrr@r6r<r+rJr)rrrretrrK branch_fakers&&& r_require_exact_strides-Switch.create.._require_exact_stridesz+s C-0\.)kf&;<<JJv&"kkmG199"-"4"4"6JJ$::%f-we;.&JrNrc3<"TFqPRJxK R#5ir)r9)rbranchs& rr Switch.create..+s;(||t#(sTFz(Expected all branch.graph to be not Nonez torch.condz torch.switchz7Output aliasing is currently not supported in compiled z. The outputs of the z subgraph of z are aliased: r c3r"TF-p\V\4'dKVP4xK/ R#5ir)rrrB)ros& rrr+s, *a!67 ALLNN*s77rzcannot determine devicer)rrrrrrrz )-check_input_alias_and_mutation_return_outputsrl)2rs r{r9r rrrrrrBr@rrrrrrrrrrrrrrBrAr>rrr?r:r r1rr. rr rr5rr6rrrM torch._higher_order_ops.utilsrrDrr' rr& )#rrrrrrrrfx_op fake_outputsrrbranch_out_argsrrrbranch_outputsr b_outputs ref_outputsrr_ob_or8rswitchr merged_outputrrM rmutated_operand_indicesr&branch_mutatedrs#&&&&& rr Switch.create]+s$$X.2:;(Q%%a((; ! 4 4 9 9" = +x00 k!23 3$&  E%&&$$UZZ%67$$U+ ! ww++007  6!H~~%!"!6!6,,#0"*--"7" #+"7"7"="="I"I"K"P"PQR"S!/8<<(o)>??IX+IXAZ4%8%8AFF5Ma?+((88NN&& 6 4J 44l4HNN0 98!0 3;(;333;(; ; ; !KL LCKL8,,448L"),~!$X~!> FI#I..$MgYW**0++mG9N[dZeg"?%Q' '++I;3y>1$k9%=>>!*3{I+F!G :C>>#s~~'77(!S#77==?cmmo5(!S#77>>#**cnn.>.E.EE(!S#77 "H,  *   6 GG   & & GG % % ) )*=t D  > !:; ;$F3/  6/8K!5!5!:!:5!AB/% $/**FM# ((*6",,. **,4A4F4F4HI4Hb+b/4HI6C6J6J6LM6LKO6LM!,,.55$//1;;   "/%  .! 4><F=f>Q>QR )Q1na $z.'A A # 56# 6 8C=//# G6#  CCSrc \V4'gV#VUu.uF+p\V\4'dVP4MTNK- pp\ 4p.p\ W4FfwrV\ V4V9d'VP\PV44K;VP\ V44VPV4Kh V#uupir) rrrrrDrr r@r+rrr-)rr^unwrapped_buffers seen_buffersroriginal_inputunwrapped_buffers& r_clone_aliased_inputsWhileLoop._clone_aliased_inputs`,s#N33! !) (%/v$G$GF   V S(  )3 +-03N0V ,N"#|3 l55nEF  $4!56 n- 1W ! s1Cc V^8dQhRRRR/#)rr'rrr)rs"rrr{,sLLvL&Lrc \V\4'dV#\V\\34'd \V4#\V\4'd\P V4#\ R\V4 24h)zNYI unsupported output type: )rrrrrr rrr)r's&r_maybe_wrap_as_tensor_box#WhileLoop._maybe_wrap_as_tensor_boxz,se c9 % %J j/: ; ;S> ! [ ) )##C( (!>tCykJK Krc 0V^8dQhRRRRRRRRRRR R /#) rcond_fnrbody_fnrrrrrrrr)rs"rrr,sNmmmm) m , m  m #mrc ^RIHpRRlp\PPP R,p\PPP R,p W,p V U u.uFqP R,NK p p VU u.uFqP R,NK p p V U u.uFqP R,NK pp VU u.uFqPV 4NK pp \PV4pV!W4pVU u.uFqPV 4NK pp V!VV4pVV,pW3EF.pVPeK\V \4'g\\V 44h\PPVPV VP R7Vn\P"!VP4;_uu_4VPP$!V !VVJde\'VPP(4\'V 48wd \R4hV!VPP(V 4VPnRRR4EK1 VP'dVP'g \R4hVPP(pVPP(p\+V4'd\R V 24h\'V4^8wd \V4hV^,p\V\,4'gXVP/4\0P28wd \V4h\'VP544^8wd \V4h\'V4^8:d \R 4hV^,P74pVf \R 4h\'V4\'V48wd\VV34h\9\;VV44FwpwppR R lpV!VP54VP544V!VP=4VP=44VP74VP748wd\VVVV34hVP/4VP/48wgK\VVV34h Vf \R 4h\?\PP@PB\PPP PERR44p\VVVV\GVR7VVR7pVPe?\VPPH\0PJPL4'g \R4hV!VPPHV 4^,p\OV4pVUu.uF pVV,NK p p\QV 4p!.p".Vn).Vn*V'Ed \'V4^8wd \R4h\9\PPP R,4Fwpp#\W\YV#PZV#P\V#P_4U$u.uFp$\aV$4NK up$V#Pc4U%u.uFp%\aV%4NK up%R7V\dV3.4p&VPRPgV&4V"PgV&4K EM%\9V4EFwpp#VV9djV\'V48d \R4h\iV!4p'VPTPg\kV'PlV'V44V"PgV'4Kw\W\YV#P74V#P/4V#P54V#P=4V#Po4PpR7V\dV3.4p&VPRPgV&4V"PgV&4EK \;VV"4Flwp(p)V(Ps4\PPt9gK5\PPvPyV)Ps44Kn V"#uup iuup iuup iuup iuup i +'giEK;iuupiuup$iuup%i)zcreate the while_loop IR node. stack_output controls whether it stack each iterations' output, which is necessary for training. )check_input_alias_and_mutationc$V^8dQhRRRRRR/#)r tensor_boxesrrz'list[int | torch.SymInt | torch.Tensor]rrr)rs"rr&WhileLoop.create..__annotate__,s&  * A  rc r\V4\V48wd \R4h.p\W4Fwr4\V\P 4'dL\ PV4pVP\PWTP4RR74KpVPV4K V#)z/Expected len(tensor_boxes) == len(fake_tensors)Fr]) rrrrrDr rrr@r+rJr6)r rrrfknew_tbs&& rr0WhileLoop.create.._require_exact_strides,s< C $55$%VWWCl9b%,,//'@@DFJJ$::"IIKu; JJrN):*JrrfNrzHExpected len(subgraph.graph.graph_outputs) == len( fake_carried_inputs )z(Expected cond_fn.graph and body_fn.graphzOutput aliasing is currently not supported in compiled torch.while_loop. The outputs of the body_fn subgraph of torch.while_loop are aliased: z9torch.while_loop is assumed to have at least one operand.rc$V^8dQhRRRRRR/#)r lhs_exprszSequence[int | Expr] rhs_exprsrrr)rs"rrr-s& < </ </ < ._guard_list_equals-sIy>S^3()TUU #I 9HCGG$$11#;!:rrr)rrrrrrrz_Expected body_fn.graph is not None and isinstance( body_fn.graph.module, torch.fx.GraphModule )z-NYI: while_loop_stack_output input mutations.rzonly carries can be mutated.)r8r7r5r6rrl)=rr r{r9r rrrs rrrrrrrrrrrrrrrrArDrr<rBrrrnr?r:r r1r. modulefx GraphModulerDrrM r& rr rr8r7r5rr6rr@rr' rr>rrrrr-)*rrr rrrr rfx_carried_inputsfx_additional_inputs fx_all_inputsrfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_ all_inputsr cond_outputs body_outputsr r8rrborr while_loop mutated_idxsmutated_idx_setrr mutated_inputs_iter all_outputsrrrr mutated_inputrr's*&&&&&& rrWhileLoop.create,sf Q <GG0055b9 ww3388<)@ 2?@-Q66%==-@6GH6Gvve}}6GH9M!N9MA&&--9M!N9GHA,,Q/H#99/J0V !>? ?  3|#4 4 #   %S,%GHKAxB < r{{}bkkm < r}}  @}}"--/1$aR%899||~/$aR[11%I( > !>? ?5 GG   & & GG % % ) )*=t D  *0!!$F3/%   MM %7==//1E1EFF q  6 MM /  %\25DE_c*S//_E#>2$&  &( # <?#q($%TUU()=)=)B)B5)IJ V'%}}$ll8> F "k"o F:@--/J/B B/J  C[M  "")))4""9- K )6 V/)c.11,-KLL$()<$=M//66&}';';]JW &&}5 +##)#4#4#6"("2"2"4!'!2#)#4#4#6#)#4#4#6#=#= #  !I&&--i8&&y1/ 72NK8HC||~!5!55 ++// ?9AH!NHO 988\F GJs=/b2b71b<ccBc  c  c%4c* c c V^8dQhRRRR/#rDr)rs"rrru-rrc  VPWP4VPVP4VP\ VR/44R#r)codegen_while_looprrrrM rrFs&&rrGWhileLoop.codegenu-s?""4):):;88 MMOT\\749Lb+Q rcV^8dQhRR/#rr)rs"rrr{-rrc \VRR4;p'dW\\PPP V4pVf \ R4h\VP44#\4#rrrs& rr"WhileLoop.get_unbacked_symbol_defs{-rr)rrrrrrr)rrrrrSrrrrrrM r/r$rVrrr:rrGrrrrs@rrr,,s].2N+215.5%)M?)%)M?),0G )0))>2LLmm^   rrcXa]tRtRtR RR/RV3RlllltRV3RlltRR ltR tV;t#) ri-Nrc8V^8dQhRRRRRRRRR R R R R RRR/#rur)rs"rrEffectfulKernel.__annotate__-s[2222& 2 & 2 + 2&2E2 2rc  <\S V`VVVVVRVR7^RIHpV!V4p V f \ R4hWn\ PPPV R4Vn V\ PPV &R#)Nr>) _get_effectz Expected effect_type is not None) rr/torch._higher_order_ops.effectsr=r effect_typer{r9 effectful_opsr1prev_effect_buffer) rrrb r rvr rrr=r?rs &&&&&&&$ rr/EffectfulKernel.__init__-s      /   @!&)   !CD D&"#''"7"7";";K"N-1k*rcV^8dQhRR/#r r)rs"rrr;-s!8rc <\SV`4pVPeHVPP \ P !VPP444V#r)rr rArr-rNr r)rrrs& rr EffectfulKernel.get_read_writes-sUg-/  " " .    ! !$$T%<%<%E%E%GH rcV^8dQhRR/#rsr)rs"rrr;-rrc R#rRrrs&rr EffectfulKernel.has_side_effects-rr)r?rAr) rrrrr/r rrrrs@rrr-s-2HL22<rrcH]tRtRt]!R4RRRll4tRRltRtR#) ri-c V^8dQhRRRR/#r#r)rs"rrNonTensorObj.__annotate__-r rc \4#rrCr&s&&rr'!NonTensorObj.get_free_symbol_uses-r rcV^8dQhRR/#rr)rs"rrrK-src R#rrrs&rrNonTensorObj.realize-rrrNr)rrrrrer'rrrrrrr-s#N+, rrcl]tRtRt$R]R&R]R&RRltRR R lltR R ltR RltRRlt Rt R#)r i-rr%FakeScriptObject | torch.ScriptObjectrJcV^8dQhRR/#rr)rs"rrTorchBindObject.__annotate__-#rc VP#rrrs&rrTorchBindObject.get_name-rrNc V^8dQhRRRR/#rr)rs"rrrT-(=rc VP#rrrs&&rr!TorchBindObject.codegen_reference-rrcV^8dQhRR/#)rrrRr)rs"rrrT-s@rc VP#rrLrs&rrTorchBindObject.get_value-rVrcV^8dQhRR/#)rrztorch.ScriptObjectr)rs"rrrT-s''0'rc \VP\P4'd VP#VPP#r)rrJrDrreal_objrs&rrTorchBindObject.get_real_obj-s3 djj%"4"4 5 5:: ::&& &rcV^8dQhRR/#rr)rs"rrrT-s==s=rc VP4p\V4'd^#\VR4'g \R4h\ VP 44p\ P!V4^,pVUu.uFJp\V\P4'gK%VP4VP4,NKL pp\P!\P V^4#uupi)r__obj_flatten__z4Expected hasattr(real_script_obj, "__obj_flatten__"))rr/r,rrrer&rrrDr rr4r*rOoperatorr-)rreal_script_obj flat_dict flat_elemsr flat_sizess& r get_buf_bytesTorchBindObject.get_buf_bytes-s++-  / /(9:: !WX X88:; ((3A6   !U\\* )ANN qwwy ( (   j!<<  s ;"C3")C3rr) rrrrrrrrrrkrrrrr r -s+ I 00' ==rr cB]tRtRt$RtR]R&RRltR RR lltR tR#) OpaqueValueTypeConstanti-aIR 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. rrJcV^8dQhRR/#rr)rs"rr$OpaqueValueTypeConstant.__annotate__-s  # rc ,\VP4#r)rcrJrs&rr OpaqueValueTypeConstant.get_name-sDJJrNc V^8dQhRRRR/#rr)rs"rrrp-s(=rc \VP4wr#VP4F"wrEV\PP V&K$ V#r)r.rJr r{r9opaque_value_type_classes)rrobj_repr opaque_typesrrHs&& rr)OpaqueValueTypeConstant.codegen_reference-sA!4TZZ!@ &&(DA34AGG - -a 0)rrr rrrrrSrrrrrrrrnrn-s J rrncH]tRtRt$R]R&R]R&RRltR R R lltR tR#) ri-rrrr8cV^8dQhRR/#rr)rs"rrGeneratorState.__annotate__.rUrc VP#rrrs&rrGeneratorState.get_name.rrNc V^8dQhRRRR/#rr)rs"rrr|.rYrc VP#rrrs&&rr GeneratorState.codegen_reference.rrrr)rrrrrrrrrrrrr-s I rrcL]tRtRt$RtR]R&R]R&RRltR R R lltR tR #)ri.z Represents an opaque object (e.g., ProcessGroup) that is passed through as a graph input. Similar to GeneratorState, this wraps the object with its placeholder name so codegen can reference it properly. rrrrJcV^8dQhRR/#rr)rs"rrOpaqueObjectState.__annotate__.rUrc VP#rrrs&rrOpaqueObjectState.get_name.rrNc V^8dQhRRRR/#rr)rs"rrr.rYrc VP#rrrs&&rr#OpaqueObjectState.codegen_reference.rrrrryrrrrr.s" I Jrrcj]tRtRtRRltRRltRRRllt]R R l4t]R R l4t R t R#)_CollectiveKerneli.cV^8dQhRR/#rsr)rs"rr_CollectiveKernel.__annotate__.rqrc R#rrrs&rr' !_CollectiveKernel.should_allocate.rrcV^8dQhRR/#rsr)rs"rrr.rrc R#rRrrs&rr"_CollectiveKernel.has_side_effects.rrNc V^8dQhRRRR/#rJr)rs"rrr".s   :    rc |\VP4\PPJd \ R4hVPpVeWnMVPPVnVPPUu.uF"q3P'gKVPNK$ upVn R#uupi)z,Setting cpp kernel needs a valid op_overloadN) rr rDr+r,rrr-rr.r/r)rrrb rs&& rr%_CollectiveKernel.set_cpp_kernel_name".s   !)>)> > !OP P!!   #)>>#6#6D #NN44. 4q FAFF4. *. s B9B9c ,V^8dQhRRRRRRRRRR /#) rrb rrzIRNode | list[IRNode]rrrrrr)rs"rrr6.sC/@/@/@&/@ /@  /@  /@rc  \PP;_uu_4VP!W.VO5/VBpRRR4XPpVP pVP pVP'd\V RVP 24hRp VFnp \V \4'dKV P4\PPV P44V eK^V P4p Kp V f\RV R24hV!\V R7VVVV4p \ P"!V4p V P$P'V U u.uFp \)\V R7W4NK up 4V P*P'V Uu.uFqP4NK up4RV9djV P$P-\)\V R7VR,V 44V P*P-VR,P44R#R# +'giEL;iuup iuupi)NrzIn-place collective zE requires at least one tensor argument; got only non-tensor IR nodes.rr')r{r9r rr r r rrrrrr rrBr~ r& tree_leavesr& r r' rr@)rrb rrrrr r r r8 tensor_argr@inpsrrs&&&*, rcreate_inplace _CollectiveKernel.create_inplace5.sWW   ''HHHF((  00..  # # # F81V-E-E,F!GH H%J*l33    GG ' ' (;(;(= >~#..0 & > &vh/:;  f %      !!&)&&OS Tt^Jf5s Ct T !!T"BTc<<>T"BC F?  # # * *z8&-P     % %fUm&<&<&> ? G  < U#CsH0 II 0 I c ,V^8dQhRRRRRRRRRR /#) rrb rrzTensorBox | list[TensorBox]rrrrz%list[MultiOutput] | _CollectiveKernelr)rs"rrr~.s<999,9 9  9 / 9rc  \PP;_uu_4VP!W.VO5/VBpRRR4XPpVP pVP pVPp VP'd\V RVP 24hVF+p \V \4'dKV P4K- \V\4'EdVPWv4p V f \R4hV!\V R7VVVV 4p \!V4U Uu.uF(wr\#VP%V4V \V 3.4NK* upp V n\)V P&V4Fawr\*P,'g\/V4'dK.\PP0P3VP44Kc V P&#V!VP%V4VVVV 4p \*P,'g\/V4'g4\PP0P3V P44V .V nV # +'giEL6;iuupp i)Nrrr)r{r9r rr r r r rrrr rrrr. rrr r rM rrMr?rwrgr-r)rrb rrrrr r r r rr8r@rrrs&&&*, rcreate_out_of_place%_CollectiveKernel.create_out_of_place}.s#WW   ''HHHF..((  00..  # # # F82f.F.F-G!HI I%Jj/::""$& nd + +__[AF~$%BCC!0 F"+>!:  ";IA ((0AYK "; FN #6>>>B :::BSCCGG--11#((; C >> !$$^4 F666>O??))--fkk:$XFNMe  .sI'*.I;' I8 )rrr) rrrrr' rrr:rrrrrrrr.s<  &/@/@N99rrcFa]tRtRtR RR/RV3RlllltRRltRtV;t#) _AllReduce_Kerneli.Nrc8V^8dQhRRRRRRRRR R R R R RRR/#rur)rs"rr_AllReduce_Kernel.__annotate__.ePPPP& P & P + P&PEP Prc  V<\SV`VVVVVRVR7VPR4R#)Nr>+aoti_torch_cpu__c10d_functional_all_reduce_rr/rrbs &&&&&&&$rr/_AllReduce_Kernel.__init__.=      /     !NOrc V^8dQhRRRR/#rDr)rs"rrr.//3//rc VPR4VPV4\VP\4'dVP V4R#R#z+torch/csrc/inductor/aoti_torch/c/shim_cpu.hNinclude_extra_headerrirrr_r$rFs&&rrG_AllReduce_Kernel.codegen.C$$%RS,,T2 dkk6 * *  % %g . +rrrrrs@rrr.s&PHLPP,//rrcFa]tRtRtR RR/RV3RlllltRRltRtV;t#) _AllReduceKerneli.Nrc8V^8dQhRRRRRRRRR R R R R RRR/#rur)rs"rr_AllReduceKernel.__annotate__.seOOOO& O & O + O&OEO Orc  V<\SV`VVVVVRVR7VPR4R#)Nr>*aoti_torch_cpu__c10d_functional_all_reducerrbs &&&&&&&$rr/_AllReduceKernel.__init__.s=      /     !MNrc V^8dQhRRRR/#rDr)rs"rrr.rrc VPR4VPV4\VP\4'dVP V4R#R#rrrFs&&rrG_AllReduceKernel.codegen.rrrrrrs@rrr.s&OHLOO,//rrcza]tRtRtRRR/RV3RlllltRRltRR lt]R R l4tR V3R llt Rt V;t #) _WaitKerneli.Nrc8V^8dQhRRRRRRRRR R R R R RRR/#rur)rs"rr_WaitKernel.__annotate__.rrc  V<\SV`VVVVVRVR7VPR4R#)Nr>+aoti_torch_cpu__c10d_functional_wait_tensorrrbs &&&&&&&$rr/_WaitKernel.__init__.rrc V^8dQhRRRR/#rDr)rs"rrr/rrc VPR4VPV4\VP\4'dVP V4R#R#rrrFs&&rrG_WaitKernel.codegen/rrcV^8dQhRR/#rr)rs"rrr/s$4rc VP^,p\V\4'g \R4h\V\4'dAVP^,p\V\4'g\\ V44hV.#\V\ 4'dUVP^,p\V\4'd*VP^,wrEVPV,.#.#.#r)rrrrrrrr r)rrrcollr&rs& rget_volatile_reads_WaitKernel.get_volatile_reads/skk!n#v&& !CD D c, - - 1 Aa(($T!W--3J [ ) )::a=D$ 122Q C())IIrc$V^8dQhRRRRRR/#)rrb rrrrrr)rs"rrr2/s!   I $ rc \PP;_uu_4VPW4pRRR4XP'd\ V RVP 24hV!\ VP4R7VVPVPVP4pVPP\\ VP4R7W$44R# +'giL;i)Nrr)r{r9r rrrr~ rBr r r r& r@r' )rrb rrr@s&&& r create_wait_WaitKernel.create_wait1/s WW   ''4F  # # # F81V-E-E,F!GH H cnn. /      " "  ! !   && :S^^-=> L  s C** C: cV^8dQhRR/#r r)rs"rrrB/s!8rc <\SV`4pVP4pVF@pVPP \ P !VP444KB V#r)rr rrr-rNr r)rrvolatile_readsvrrs& rr _WaitKernel.get_read_writesB/sSg-/ 002 B    ! !,"6"6r{{}"E F!rrr) rrrrr/rGrr:rr rrrs@rrr.sDPHLPP,/6  rrc V^8dQhRRRR/#rrrrrr)rs"rrrN/s  6 .@ rcl\V\\34'd \V4#\V\\ 34'd>\ \P,!4pVFpV\V4,pK V#\V\P4'd \V4#\ 4#r) rrAr"r:rrrDrr$rBrDr rrrHs& rrBrBN/s!h%&&$Q'' At} % % u|| $ &A ,Q/ /A Au|| $ $$Q''|rc V^8dQhRRRR/#rr)rs"rrr^/s  & %7 rcl\V\\34'd \V4#\V\\ 34'd>\ \P,!4pVFpV\V4,pK V#\V\P4'd \V4#\ 4#r) rrAr"r9rrrDrr$rCrDr rs& rrCrC^/s!h%&&A At} % % u|| $ &A #A& &A Au|| $ $A|rc$V^8dQhRRRRRR/#)rrrrz torch.fx.Noderrr)rs"rrrn/s&TTsT}TTrcx\V\4'Ed#\VP\4'Ed\VPP\4'd)VPPP RV4R#\VPP\ 4'EdVPPP RV4\VPP\4'dg\VPPP\4'd3VPPPP RV4R#\VPP\4'dVPPP'gw\VPPP^,\ 4'd:VPPP^,P RV4R#R#R#R#R#R#R#r ) rrrrrrrrrr rr)rrs&&rassign_origin_nodern/st&)$$FKK)L)L fkk&& . . KK   / / q A  ((& 1 1 KK   / / q A&++**N;;    %%uAA   %%88J6;;++[99 ((000fkk..55a8&AAKK$$++A.AA-QRSB1:2*M$r)rrr)r rrrrrrrR)TFNFN({__conditional_annotations__ __future__rrWrrTr*rPloggingrftextwraprcollections.abcrrrrrr r enumr r numbersr typingrrrrrrrrrrrrtyping_extensionsrrrrrrr unittest.mockr!rr"r#r$torch._export.serde.schema_exportserderrtorch._library.custom_opsrrrrg rtorch._loggingrDtorch.fxtorch.utils._pytree_pytreer&torch._dispatch.pythonr%torch._dynamo.utilsr&torch._export.serde.serializer'*torch._higher_order_ops.auto_functionalizer(torch._inductorr)r r+"torch._library.fake_class_registryr,r-torch._library.opaque_objectr.r/torch._native.ops.sumr0torch._prims_commonr1r2r3r4r5r6%torch.fx.experimental.symbolic_shapesr7r8r9r:r;r<r=r>r?r@rA torch.fx.noderBtorch.utils._ordered_setrDtorch.utils._python_dispatchrEtorch.utils._sympy.functionsrFrGrHrIrJrKtorch.utils._sympy.symbolrLr<rMrNcodegen.commonrOrPrQrRrSrTrUrVrWrXrY loop_bodyrZ ops_handlerr[r\r]r^runtime.benchmarkingr_ runtime.hintsr`rarbrcrdrerfrgrhrirjrkrlrmrnrorprqrrrsrtrurvrwrx virtualizedryrzr{r|r} torch.typesr~codegen.cutlass.templatercodegen.wrapperrr9rkernel.gemm_epiloguerrrrrZ __version__rrr rrrrrrrArr+r,rarr getLoggerrr.rKrLrrrrrrrrrrr r r rr#r)r.rMrTrardrvrrrrrrr[rrBrDrlrrrr_rrr6rPrurrr r=rRrVrrbrerrrrr(rrrrr$ r5 rH r[ r_ rSr_rr?r[ rp rt r~ r rr* rr< rhrL rrm rr r rr rrPrimitiveInfoTyper r r r r r7 rI rT r` r rr r rXr r r r+rLr\r r' rrrrrrrZrrrrrrrrrr-rrrrr rXrer. rr rrrMrPr rrrrrrrrrrrr rnrrrrrrrBrCr)rs@rr s "" MM:     ''22..,,$$;(?M#2RQ1    /?+" NN-:2*)B&'95$7%"(OY'''NJ t_ T] T] T]Dj) EkD()(**//%**2P2PP iP)5 !   8??4 8 yy~~'T k     c4[();(JKdR   i ) d#$$!$H!$      U U     D >;*'7*2%GP';k,k,\ UaaaHW FW W t&% % % P i  F ;| $ ;y ! K # ;y ! K  ;u  ;u  {=) <8 XNXNvO O O h)-1:."8D>8D>"BH"LM YM8S98Svn'nbs 1s ld +d NQ 5Q Q j   a5aaH* ::z$_ v_ _ DbbbJ?(??DH9(H9H9V!(!!Hb;bbJ ShSSl,,,^|H|H~6"K|KK$ S| S S(< 7 7 w Zw w t C&CaHVaHH$Gf$GNT%{%D.]*]*@UNV]NNbU &fi & & & K [ &6( F    d#  $ U~4_~4~4B d#$4X4m+_m+`?>?D%K$&,tC#I4E4L/MMLL^"|" ,0[/ ;Q,.Q,h5N50((<N@{%N{%|55 UT?TTn Z9Zz ;; d# A A$ AUH<HHV"U#l##L / (+" +"\ V Bi@?L?D' ' T - $> \> BAlAH ,)|,)`/,/d<:B5B(8-828-827)l7)t))|))X=3=3@1E1Eh51|51p-L-8<;<<;~U } &} @U .  @  2 K ,K \   F 4K 4 22f>fV TTTn + +M%M%`U'v''LLUX.\X.X.vUt \t t n " "UW W W t *n*Z6#=l#=#=Ll(\ "^^B/)/>/(/>P#Pl     TTYvNJsr88 ss