3j %SSKJr SSKrSSKrSSKrSSKrSSKrSSKrSSKrSSK r SSK r SSK J r J r JrJrJr SSKJrJr SSKJr SSKJr SSKJrJrJrJrJrJrJrJrJrJ r J!r! SSK"J#r#J$r$J%r%J&r&J'r'J(r( SS K)J*r* SSK+r+SS K+J,r,J-r-J.r. SSK/J0s J1s J2r3 SSK4J5s J6r7 SSK8r9SSK:r9SSK;J6s JJ?r? SS K@JArA SS KBJCrC SSKDJErE SSKFJGrG SSKHJIrI SSKJJKrKJLrL SSKMJNrNJOrOJPrPJQrQJRrRJSrS SSKTJUrUJVrVJWrWJXrXJYrYJZrZJ[r[J\r\J]r]J^r^J_r_ SSK`Jara SSKbJcrc SSKdJere SSKfJgrgJhrhJiriJjrj SSKkJlrl SSKmJnrnJoro SSKpJqrqJrrrJsrsJtrtJuru SSKoJvrvJwrwJxrxJyryJzrzJ{r{ SSK|J}r} SSK~JrJrJrJr SSKJr SS KJrJr SS!K6JrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJrJr SS"KJrJrJr \(a+SS#KTJr SS$K`Jr SS%KJr SS&KJr SS'KJr SS(KJr SS)K6Jr O\rS*\S+'SSKr\GR`rS,r\&"S.5r\ "S/5r\ "S05r\ "S15r\\,-rS*\S2'\\-\,-rS*\S3'\9GRxGRz\9GRxGR|-rS*\S4'\GR"\5r\R*"\ GRS5S69r\9GR>GRr\!\S7\\S74S8S9\\!\\\S74S7S8S94S-4rS*\S:'GS;S;jr\GR"S,S<9"S=S>55rGSSAjrGS>SBjrGS?SCjr/SDQr/SEQrGS@GSASFjjrGSBSGjrGS@GSASHjjr\GSCGSDSIjj5r\GSCGSESJjj5rGSCGSFSKjjrGSGSLjrGSHSMjrGSISNjrGSISOjrGSJSPjrGSKSQjrGSLSRjrGSMSSjrGSNSTjrSUr"SVS95r\"S-S<9"SWSX55r\"SYSZ\55rGSOS[jr\"S\S]\55r\"S^S_\55r\"S`5\"Sa5\"Sb5\"Sc5\"Sd5\"Sd5\"Se5Sf.rSg\Sh'GSPGSQSijjr\"SjSk\55rS\-"S54GSRSljjr\ \\,\\,/\4rS*\Sm'"SnSo\5r"SpSq\5r"SrSs\5r\"StSu\55r\"SvSw\55r\"SxSy\55rGSSSzjrGSSS{jrGSTGSUS|jjrGSVS}jrGSWS~jr\"SS\55r\"SS\55r\"SS\55r\"SS\55r\"SS\55r\"SS\55r\"SS\55r\"SS\55r"SS\5r\"SS\55r\"SS\55Gr\"SS\55GrGSXSjGrGSYSjGr"SS5Gr\"SSG\55Gr"SSG\5Gr"SSG\5Gr"SSG\5Gr"SS\5Gr "SSG\5Gr \"SSG\55Gr "SSG\5Gr \"S-S<9"SS\\r55Gr \"S-S<9"SSG\ \55Gr"SSG\ 5Gr"SSG\5Gr"SSG\5Gr\"SS\55Gr\"SS\55Gr\"S-S<9"SSG\55Gr\GR"S,S<9"SS55Gr"SSG\5Gr"SSG\5Gr\\-G\-\-G\\\-\-G\--Gr"SS5Gr"SSG\5Gr"SSG\5Gr"SSG\5Gr"SSG\5Gr"SSG\5Gr "SSG\5Gr!GSZSjGr"\"S-S<9"SSG\55Gr#"SSG\#5Gr$"SSG\$5Gr%\"S-S<9"SSG\#55Gr&\"S-S<9"SSG\&55Gr'"SSG\'5Gr("SSG\&5Gr)"SSG\ 5Gr*"SSG\&5Gr+"SSG\+5Gr,"SSG\+5Gr-"SSG\&5Gr."SSG\&5Gr/"SSG\&5Gr0"SSG\&5Gr1"SSG\&5Gr2"SSG\25Gr3"SSG\)5Gr4"SSG\&5Gr5"SSG\&5Gr6"SSG\'5Gr7"SSG\&5Gr8"SSG\&5Gr9"SSG\&5Gr:"GSGSG\&5Gr;\"S-S<9"GSGS55Gr<"GSGSG\)5Gr=\"S-S<9"GSGSG\=55Gr>"GSGS G\=5Gr?\"GS GS G\55Gr@"GS GS G\&5GrA"GSGSG\A5GrB"GSGSG\A5GrCGS[GSjGrD"GSGSG\=5GrE\GR"GSGS\55GrF"GSS7G\F5GrG"GSGSG\F5GrH\"S-S<9"GSGS\55GrIGS\GSjGrJ\"S-S<9"GSGSG\&55GrK\"S-S<9"GSGS G\&55GrLGS]GS!jGrM\"S-S<9"GS"GS#G\&55GrN"GS$GS%G\=5GrO"GS&GS'\5GrP\"GS(GS)G\P55GrQ\"GS*GS+G\P55GrR\"GS,GS-G\P55GrS\"GS.GS/G\P55GrT"GS0GS1G\=5GrU"GS2GS3G\U5GrV"GS4GS5G\U5GrW"GS6GS7G\U5GrXGS^GS8jGrYGS^GS9jGrZGS_GS:jGr[g!\a SrS-rG NJf=f(`) annotationsN)Callable GeneratorIterableIteratorSequence)AbstractContextManager nullcontext)Enum)partial) AnycastClassVarLiteraloverload SupportsFloat SupportsInt TYPE_CHECKING TypeAliasTypeVarUnion) assert_neverNeveroverride ParamSpecSelfTypeIs)patch)ExprIntegerSymbol)identity)GraphModuleSerializer)can_auto_functionalize)metricsget_free_symbols)FakeScriptObject)get_opaque_obj_repris_opaque_value)compute_required_storage_lengthis_boolean_dtype(is_contiguous_for_memory_format_or_falseis_float_dtypemake_channels_last_strides_for StrideType) &_remove_effect_token_unbacked_bindingscompute_unbacked_bindings free_symbolsfree_unbacked_symbolsGuardOnDataDependentSymNodehas_free_unbacked_symbols IterateExprsrebind_unbackedresolve_unbacked_bindingsShapeEnvSymTypes)Node OrderedSet)_disable_current_modes)CleanDivFloorDivModModularIndexing)SymT)config dependencies)BackendFeature CodegenSymbolget_scheduling_for_deviceindex_prevent_reorderingKernel)Depextract_free_symbols#extract_input_node_reduction_rangesextract_read_writesSymbolUsageCollectorOpsHandler var_builder)LoopBody) OpCounterCSE OpCountResult ReductionType StoreMode) benchmarker)DeviceProperties ReductionHint)argsort argsort_sym cache_on_selfcache_on_self_and_argsceildivconvert_shape_to_inductorconvert_shape_to_symintdeveloper_warningdo_bench_using_profilingdtype_from_sizeget_dtype_sizeget_kernel_metadataGPU_ALIGN_BYTES ir_dataclass is_dynamicis_gpu sympy_dotsympy_index_symbolsympy_index_symbol_with_prefix sympy_product sympy_substensor_is_aligned)opsOpsValueV) SympyBoolean)Argument) IntLikeType)CUTLASSTemplate)PythonWrapperCodegen) GraphLowering)IndentedBufferrrwTF_P_T_U_V_IntLike_NumLike _OpOverloadsz prefix TensorBoxr!IRNode _NodeOrNodesc.[U[[45$N) isinstanceintr xs L/home/wildlama/miniconda3/lib/python3.13/site-packages/torch/_inductor/ir.py _is_staticrs a#w (()frozencR\rSrSr%S\S'S\S'S\S'S\S 'S \S 'S \S 'Srg)GraphPartitionSignatureOrderedSet[sympy.Symbol] symbol_inputsz0dict[str, IRNode | sympy.Expr | TorchBindObject] input_nodes list[IRNode] output_nodeszdict[str, bool]input_deallocationboolskip_cudagraph list[str]constant_namesN__name__ __module__ __qualname____firstlineno____annotations____static_attributes__rrrrrs/,+BA('rrc&^SU4SjjmT"U5 g)Nc >Ucg[U[[45(aUH nT"U5 M g[U[5(a UR 5H nT"U5 M g[U[ [ [[[RRR[[[[ ["4 5(dS[%U5S35eg)NzFound zE, which is not a supported top level IR node. See [Note: Inductor IR])rlisttupledictvalues ExpandView DynamicScalar AssertScalarrsympylogicboolalgBooleanrrEffectfulKernelShapeAsConstantBufferOpaqueMultiOutputtype)nodesnode_check_tensorboxs rr%validate_ir.._check_tensorboxs =  e} - - & t $ $  &'! KK''//#)%  e %jk r)r_NodeOrNodes | NonereturnNoner) node_or_nodesrs @r validate_irrs>]#rcb^[T[5(d[T55eSU4SjjnU$)Nc0>[[T5"U0UD6$r)getattrrq)argskwargsnames rfnops_wrapper..fnssD!42622r)robjectrrrrr)rstrr)rrs` r ops_wrapperrs+ dC ,$t*, 3 Irc f^[[U[[U5555mSU4SjjnU$)Nc>[U5[T5:Xde[[U55Vs/sH oTUPM sn$s snfrlenrange)indexi inv_orders rreindex inverse_reorder..reindexsC5zS^+++-23u:->?->il#->???Ar Sequence[_T]rr)rziprr)orderrrs @rinverse_reorderrs*Sc%j 123I@ Nrc^SU4SjjnU$)Nc>[U5[T5:Xde[[U55Vs/sH oTUPM sn$s snfrr)rrrs rrsame_reorder..reindex#sB5zSZ''').s5z):;):AeAh):;;;rrr)rrs` r same_reorderr"s< Nrc^^SUU4SjjnU$)Nc >T"T"U55$rr)rreindex1reindex2s rr fuse_reindexing..reindex.s((r)rrrz Sequence[_V]r)rrrs`` rfuse_reindexingr*s)) Nr)rrE)rrrrEcjUb[SU55(a [U5nU$[X5nU$)z) Convert strides to fill order (argsort) c3b# UH%n[U[[R45v M' g7frrrrr .0ss r !get_fill_order..>s#QS 1sEMM.B C CS-/)allr[r\)seq shape_env sorted_idxs rget_fill_orderr8s= CQSQQQ$+CL  !0 rc[U5VVs0sHupX!_M nnn[[U55Vs/sHoCUPM nnU$s snnfs snf)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_orderrFsR(1'7 8'783ch'7F 8%*3u:%67%6)%6J7 97s AAc[X5n[[U55Vs/sHnSPM nn[U5H upVXTU'M U$s snf)z! Convert strides to stride order r)rrrr)rrr_outrelems rget_stride_orderrRsL !/s >JCHo &o1oC &Z(D ) J 's A cgrrrreplace_symbols_with_hintss rir_node_to_tensorr_sRUrcgrrrs rrrcsrcUcgU(a%[RRRnO[nUR 5Vs/sH o2"U5PM nn[ U5(a0UR5RVs/sH o2"U5PM nnO[RU5nUR5nUR5n[U5n[U5n[RRRR5 [ R""XEXgS9R%5nSSS5 U$s snfs snf!,(df  W$=f)N)sizestridedtypedevice)rsgraphsizevarsoptimization_hintr"get_sizeis_storage_and_layout get_layoutr FlexibleLayoutcontiguous_strides get_dtype get_devicerarsuppress_guardstorch empty_stridedzero_) rrshape_fnrrr r r ts rrris y"77##55!" .AHQKD .Q'(||~'<'<=' 6 5 Hs EE $E E cD[U[5(a U(dS/$U$r)rrvalues rmay_convert_to_optionalrs!%uv Lrc@[U[5(dUcU$[U[R5(a UR$[U[ [ 45(a[UR55$[SUS[ U5RS35 g)Nzget_device_type(: )) rrrr rr OutputSpecget_device_typerrrrs rr$r$sz!SQY Au|| $ $vv A + , ,q||~..#A3ba)9)9(:!<=rc[U5nUS;a[[US35S:XaggUb[U5=ncgSSKJn [ U[5(d[U55e[X#5$)N)cpucudaxpu_backendtritonTFrE)TritonScheduling) r$rrFrJcodegen.tritonr+rr issubclass)rr device_schedulingr+s r is_tritonr/sz Q F'' 6fXX. /8 ;!:6!B B  K0 ' . .G5F0GG . ' ::rc[U5S:H$)Nr&)r$rs ris_cpur1s 1  &&rcj^[U[5(aMUR5b<[UR 55(d[UR 55(ag[ R"U4SjUR 5SS56n[ R"[ R"UR 5SS5[ R"UR 5SS55n[ R"X#5n[RRRU5$)NFc3f># UH&n[R"[UT5S5v M( g7frN)rEqrB)rr alignments rr-is_aligned_realized_tensor..s( F2EQ%((3q)$a ( (2Es.1rE)rrmaybe_get_strider4 get_striderrAndOrr5Lersr r guard_or_false)rr6aligned_stridesaligned_last_dim is_aligneds ` ris_aligned_realized_tensorrBs q& ! !    '  0 0  . .ii F!,,."2E FOxx #Q'!**,r2BA)F?=J 77   * *: 66rc~[U5[U5:Xa[U5[U5:Xde[X U5H{up4n[RRR US5(aM7[RRR [R"XE55(aM{ g g)zH Returns true if the strides are equal, ignoring dimensions of size 1 . rEFT) rrrsr r statically_known_leqr>rr5)strides1strides2shapedims1s2s rsignificant_strides_equalrKs u:X &3x=CM+II I5H5  77   0 0a 8 8 ww..uxx/?@@ 6 rct[U5(dU$[S[XR5555(aU$[ XR5UR 55(dU$[ U5up#/URQn[UR 55H<upV[RRRUS5(dM6XXE'M> [URURUR UUR"UR$5n['[)X'S95$)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# UH0up[RRRX5v M2 g7frrsr r statically_known_equalsrrIrJs rr2try_match_insignificant_strides..s1 7FB 008878:rEdatalayout)rrrr:rKras_storage_and_layoutr rrsr r rD FixedLayoutr r roffset is_pinnedrReinterpretView)tensorstridesstorage old_layout new_striderr new_layouts rtry_match_insignificant_stridesras  ! ( (   '#4#4#67  $W.?.?.A6??CT U U /7G%:$$%J&//+, 77   0 0A 6 6#JJM- J _'E FFrcURRSS9Sn[UR5VVs/sHup#UPM snnURS'SSKJn U"U5 gs snnf)Noutputopruser_visible_output_idxs)record_original_output_strides)r  find_nodesrrmetatorch._inductor.compile_fxrg)gm output_noderrrgs rgm_original_output_stridesrmsd((%%%215K#K$4$454554K/0J"2& 4sA#c[5nUH9nU[UR5SS9-nU[UR5SS9-nM; [ U5$)NF unbacked_only)r>r'rr:r)inputssym_varsinps rget_symbolic_inputsrtsP!+H$S\\^5II$S^^%5UKK >rc[U[5(a URn[U[5(aUR 5n[U[ 5(a URn[U[ 5(aUR5$S$r)rrrTBaseView unwrap_view StorageBoxBufferget_namers r try_get_namer{sc!Y FF!X MMO!Z  FF%a001::<:d:rcJ\rSrSr%Sr\"5rS\S'SrS\S'\ R"SS 9r S \S '\ R"SS 9r S \S '\ R"SS 9r S\S'\ R"SS 9rS\S'\ R"SS 9rS\S'\\R$SRSj55r\\R$SSSj55r\STSj5rSUSjrSVSjrSWSjrSXSjrSYSjrSZSjrS[SjrS\SjrSXSjrS]S^S jjrS_S`S!jjr SaS"jr!SbS#jr"ScS$jr#SdS%jr$SeS&jr%SfS'jr&SgS(jr'ShS)jr(SiS*jr)\*SjS+j5r+SkS,jr,SgS-jr-SlS.jr.SmSnS/jjr/SoS0jr0SpS1jr1SgS2jr2SqS3jr3SrS4jr4SsS5jr5SiS6jr6StS7jr7SlS8jr8SgS9jr9SmSuS:jjr:SvS;jr;SWS<jrjr>SwSxS?jjr?SyS@jr@SzSAjrASwS{SBjjrBS|SCjrCS}SDjrDS~SEjrESSFjrFSwSSGjjrGSlSHjrHShSIjrISgSJjrJSgSKjrKSSLjrLSSMjrMStSNjrNSSOjrO\P(a \*SaSPj5rQSQrRgSQrRg)ri%zBase class for all intermediate representation (IR) nodes in TorchInductor. Note: This is an abstract base class. Most methods raise NotImplementedError and must be overridden by concrete subclasses. zClassVar[OrderedSet[Any]]_current_originsNzClassVar[int | None]_current_stream_idxF)initOrderedSet[Any]originslist[str] | None tracebacktorch.fx.Node | None origin_nodedict[str, Any]r int | None stream_idxc## [RnX-[lSv U[lg!U[lf=f7fr)rr})rolds rcurrent_originsIRNode.current_origins:s4%%"%- * &)F #cF #sA1 A >Ac## [RnU[lSv U[lg!U[lf=f7fr)rr~)rrs rcurrent_stream_idxIRNode.current_stream_idxDs2 ((%/" - ),F &F &s?/ ? <?cL[U[[[[[ 45$r)rComputedBuffer InputsKernel InputBufferrZTemplateBuffer)rs ris_realized_nodeIRNode.is_realized_nodePs&     rc,[RU5$r)rcreateselfs rwrap_for_loweringIRNode.wrap_for_lowering]s%%rc0[RXU5 gr)r __setattr__)rattrrs r_post_init_setattrIRNode._post_init_setattr`s 4u-rcJ[UR5nURSU5 URS[R(a[ R "5OS5 URSS5 URS05 URSUR5 g)Nrrrrr)r>r}rrFdebug_ir_tracebackr format_stackr~)rrs r __post_init__IRNode.__post_init__fsT223  73  V5N5N//1TX   t4  r2  d.F.FGrcB[SUR555$)Nc38# UHoRv M g7frrrdeps rr(IRNode.get_read_names..r?.>s((.>r> get_readsrs rget_read_namesIRNode.get_read_namesq?dnn.>???rcUR$r)rrs r get_tracebackIRNode.get_tracebackts ~~rcUR$rrrs rget_origin_nodeIRNode.get_origin_nodewrcgrrrs rget_defining_opIRNode.get_defining_opzrc/$)z'Return subgraphs contained in this noderrs r get_subgraphsIRNode.get_subgraphs} rc[5nURn[U[5(a-UR 5nUR (a [U/5nUHn[ US5(a.UR(aURUR5 MB[RRRRS05RUR/5n[U[5(dMUHQn[RRR RUS5nU(dM@URU5 MS GM U$)N stack_trace postToPre)r>rr ExternKernelrrhasattrraddr _inductordebug _inductor_post_to_pre_grad_nodesgetrr#_inductor_pre_grad_node_stack_trace)r stack_tracesrrrpre_grad_nodes node_namers rget_stack_tracesIRNode.get_stack_tracess)3 ,, dL ) )..0K$k]3Dt]++0@0@  !1!12OO))JJNN#c$))R( ".$77!/I--QQUU%t #{$((5"0,rc6S[USS53nU(a[U5S:aUSSS3nUR5(dU/$/nUR5H8nURS5 X4R S5- nURS 5 M: U/U-$) Nzorigins=r@=z...zstack_traces = { })rrrappendsplit)rshortenrstack_trace_strrs r common_reprIRNode.common_reprsWT9b9:; s7|b( "c*G$$&&9 002K  " "#5 6 006 6O  " "3 '3y?**rc.[U5[URU55-n[[[U55nU(a5[ SR U55n[ U5RSUS3$[ U5RSUS3$)Nz, z( z )(r")rrmaprindentjoinrr)rlinesr multiline new_liness r str_helperIRNode.str_helpersU d4#3#3G#<==Se_% uzz%01I4j))*#i[< <4j))*!E7!4 4rcUR$rr rs rrIRNode.get_dtype zzrcDUR5$![a gf=fr)rNotImplementedErrorrs rmaybe_get_dtypeIRNode.maybe_get_dtypes& >># #"    c2[S[U5S35e)Nz#get_layout() is not implemented by !rrrs rrIRNode.get_layouts!$GT |ST"UVVrcDUR5$![a gf=fr)rrrs rmaybe_get_layoutIRNode.maybe_get_layout& ??$ $"  rc"UR5$r)rrs rget_output_specIRNode.get_output_specs  rcDUR5$![a gf=fr)rrrs rmaybe_get_output_specIRNode.maybe_get_output_specs( '') )"  rc>[UR5[5$)z4True for single tensor output (excludes MultiOutput))rrLayoutrs rhas_tensor_outputIRNode.has_tensor_outputs$446??rc2[S[U5S35e)Nz!get_size() is not implemented by rrrs rrIRNode.get_sizes!$Ed4j\QR"STTrcDUR5$![a gf=fr)rrrs rmaybe_get_sizeIRNode.maybe_get_size% ==? ""  rc"UR5$rrrs rrG IRNode.shapes}}rc4[UR55$r)rnrrs r get_numelIRNode.get_numelsT]]_--rc[RRR[R "UR 5S55$Nrrsr r statically_known_truerr5rrs ris_zero_elementsIRNode.is_zero_elements0ww55ehht~~?OQR6STTrc0[S[U535e)a If the IRNode refers to data which has not been materialized (e.g., it is a Pointwise/Reduction that could potentially have more compute fused into it), realize the IRNode into physical memory, ending the possibility of fusing into it, but allowing, e.g., multiple users to access the data without having to recompute. Check StorageBox.realize for a particularly notable implementation. TODO(ezyang): I think, in principle, every IRNode should have an implementation of this, and most of the time no-op is OK, but you really do have to audit each IRNode for this, so for now, raise an error if it's not implemented. Note that some code in graph.py will catch this thrown error and suppress it with a warning. zrealize NYI on rrs rrealizeIRNode.realizes "ODJ<"@AArc0[S[U535e)Nzcodegen_reference NYI on rrwriters rcodegen_referenceIRNode.codegen_references!$=d4j\"JKKrcgrrrs rrIRNode.get_devicerrc0UR5nUceU$r)rrr s rget_device_or_errorIRNode.get_device_or_errors"!!! rcgNFrrs rhas_exceeded_max_readsIRNode.has_exceeded_max_reads rc>[[U5R5errrrrs r make_loaderIRNode.make_loader!$t*"5"566rc>[[U5R5err,rs r make_indexerIRNode.make_indexerr/rc>[[U5R5err,rs rr:IRNode.get_strider/rcDUR5$![a gf=fr)r:rrs rr9IRNode.maybe_get_striderrc>[[U5R5err,rs rrzIRNode.get_namer/rcDUR5$![a gf=fr)rzrrs rmaybe_get_nameIRNode.maybe_get_name!r rczUR5[RR;$![a gf=fr')rzrsr  graph_inputsrrs ris_input_bufferIRNode.is_input_buffer's4 ==?agg&:&:: :"  s *- ::cgr'rr thresholds rhas_large_inner_fnIRNode.has_large_inner_fn-r*rcgrrruserss r mark_reuseIRNode.mark_reuse0 rcgrrrs r realize_hintIRNode.realize_hint3rJrc>[[U5R5err,rs rrwIRNode.unwrap_view6r/rc>[[U5R5err,rs r freeze_layoutIRNode.freeze_layout9r/rc>[[U5R5err,rr allow_paddings rfreeze_layout_with_stride_order&IRNode.freeze_layout_with_stride_order<"$t*"5"566rc>[[U5R5err,rrs rfreeze_layout_with_fill_order$IRNode.freeze_layout_with_fill_orderAr/rc>[[U5R5err,rr s rfreeze_layout_with_same_order$IRNode.freeze_layout_with_same_orderDr/rc>[[U5R5err,r exact_stridesrUs r freeze_layout_with_exact_strides'IRNode.freeze_layout_with_exact_stridesGrXrc>[[U5R5err,rs rget_read_writesIRNode.get_read_writesLr/rc6UR5R$rrgreadsrs rrIRNode.get_readsO##%+++rc4[UR55$r)rrrs r num_readsIRNode.num_readsRs4>>#$$rc>[[U5R5err,rs rget_storage_numelIRNode.get_storage_numelUr/rc>[[U5R5err,rrps rget_free_symbol_usesIRNode.get_free_symbol_usesXrXrc>[[U5R5err,rs rget_reduction_typeIRNode.get_reduction_type]r/rc>[[U5R5err,rs rget_reduction_sizeIRNode.get_reduction_size`r/rcgr'rrs r is_externIRNode.is_externcr*rcgr'rrs ris_no_opIRNode.is_no_opfr*rc>[[U5R5err,r#s rconstant_to_deviceIRNode.constant_to_deviceir/rc>[[U5R5err,rs rget_mutation_namesIRNode.get_mutation_nameslr/rc>[[U5R5err,rs rget_operation_nameIRNode.get_operation_nameor/rc>[[U5R5err,rs rget_inputs_that_alias_output#IRNode.get_inputs_that_alias_outputrr/rcgrrrs rr  IRNode.dtypews(+rr)rzOrderedSet[Node]rGenerator[None, None, None])rrrrrrrrrr)rrrr rrrrrOrderedSet[str])rrrrrzOperation | Nonerzlist[Subgraph]T)rrr Sequence[str])TT)rzSequence[object]rrrrrrr torch.dtype)rztorch.dtype | Nonerr)rz Layout | Nonerr#)rzOutputSpec | NonerrrSequence[Expr])rSequence[_IntLike] | None)rz)_IntLike | sympy.Rel | Sequence[_IntLike]rrr str | NonerrzIndentedBuffer | Nonerrrtorch.device | Noner torch.devicer$Callable[[Sequence[Expr]], OpsValue]r Callable[[Sequence[Expr]], Expr]rSequence[_IntLike]rrrBrrrrGrrrFr Sequence[int]rUrrrrrrrr rrrrcrrUrrrrdependencies.ReadWritesrzOrderedSet[Dep]rrrrrprrrr rrrrr)Srrrr__doc__r>r}rr~ dataclassesfieldrrrrr staticmethod contextlibcontextmanagerrrrrrrrrrrrrrrrrrrrrrrrpropertyrGrrrrrr$r(r-r1r:r9rzr:r>rCrHrLrwrQrVr[r_rdrgrrorrrvryr|rrrrrrrr rrrrrr%s3=,/>04-4 +00e.rrrrs rrOperation.get_read_namesrrc6UR5R$rrjrs rrOperation.get_readsrmrc[errrs r get_outputsOperation.get_outputsrrc[5$rr=rs rget_unbacked_symbol_defs"Operation.get_unbacked_symbol_defs |rc[5$)a When unbacked_only=True: Returns the unbacked symbols which are required to be in scope in order to successfully perform codegen for this buffer. For example, a buffer that corresponds to an extern kernel call that takes i0 as an argument would return {i0} here. This is used to generate necessary dependencies that ensure we actually bind i0 in codegen before you try to use it. Note that this is NOT transitive; in particular, if this buffer takes in as input another buffer with dynamic shape (e.g., (i0,)), we will not report it here, because you will already have a dependency on that buffer, which will eventually have a dependency on i0 if necessary. When unbacked_only=False: Similar to `unbacked_only=True` but including all free symbols instead of only free unbacked symbols. r=rus rrvOperation.get_free_symbol_usess ,|rcg)z Gets extra global memory size needed by this buffer. Some algorithms (e.g. group gemm) may require extra global memory in the generated code. rrrs rget_workspace_sizeOperation.get_workspace_sizes r)rrNrrr)rr)rrr)rr)rrrrrr)rrrrrrr list[Buffer]rrrrr)rrrrrrrrrrrrrrrgrrrrrrvrrrrrrr{sr2" #$'"-@,"%*! !0rrc\rSrSr%S\S'S\S'S\S'S\S '\"S5S!S"S jj5rS#S jrS$S jr\r S%S jr S&Sjr S'Sjr S'Sjr \S(Sj5r\\R$4S)Sjj5r\S*Sj5rS+Sjr\S$Sj5rS,S-SjjrS!S.SjjrS/SjrS0SjrS1SjrS2SjrS'SjrS3SjrS4Sjr S r!g)5Loopsirr rr Callable[..., Any]inner_fnrrangesc^[5R"/U4SjUR5QURT5P76$)Nc3<># UHn[UT5v M g7frr&rerps rr-Loops.get_free_symbol_uses..s F+Qq-00+)r>unionr inner_fn_free_symbolsrus `rrvLoops.get_free_symbol_usess>|!! F$++ F  & &} 5  rc  URSURRS3[UR5UR 5/UVs/sHo"S[ X53PM sn-SUR<3/-5$s snf)N'= origin_node=)rr rrr  inner_fn_strrr)rnamesrs r_to_str Loops._to_strsDKK$$%Q'DJJ!!#  rD read_buffersrs rrLoops.get_read_namesHs$//1>>??rcH[UR5R5$r)rrDr]rs rroLoops.num_readsKs4((*7788rc2[S[U5S35e)Nz+get_reduction_size() is not implemented by rrrs rr|Loops.get_reduction_sizeN!9$t*Q G  rc2[S[U5S35e)Nz+get_reduction_type() is not implemented by rrrs rryLoops.get_reduction_typeSrcrc2[S[U5S35e)Nz+constant_to_device() is not implemented by rrr#s rrLoops.constant_to_deviceXrcrrrr)rrrrrrrr)rr rr rr)r rrrDrr)rrUrzSequence[Sequence[_IntLike]]rrrprrOrderedSet[Symbol])rVr!rrrrrrr)"rrrrrr^rvrr!__repr__rrrr- classmethodrrrDINDEXr:r]rDrArrCrrXrrror|ryrrrrrrrs    G$$) ! ! %   )H  # #:>**  ((+  ; W  @9   rrcUR(a [R"[S5U5$[R"SU5$)Nnanr)is_floating_pointrqconstantfloat)rr s r nop_loader_fnrs^s1 ||E%L%00||Au%%rcd\rSrSrS SjrS Sjr\rS SjrS SjrSSjr SSjr Sr g ) PointwiseiectUR5(a[[URS9$UR$Nr)rr rsr r rs rr-Pointwise.make_loadergs,  " "= ; ;}}rc$URS5$rr rs rr!Pointwise.__str__nr#rc/$rrrs rr|Pointwise.get_reduction_sizess rcgrrrs rryPointwise.get_reduction_typevrrc|UR5n[R"U=(d SU"U5U"U55$Nunnamed)r-rqstorer output_nameindexervarsloaders r store_outputPointwise.store_outputys2 !!#yy1 74=&,OOrcUR5n[R"[SU5"U5n[ UUR UUR S9$FMove this to a given device. Requires that all reads are to constants.override_devicer r r r )r-rrConstantBufferrur r rr rs rrPointwise.constant_to_devicesI!!#n.?HP**;;   rrNrr)rzSequence[sympy.Expr]r)rrr!Callable[[Sequence[Expr]], Never]rrrrr) rrrrr-r!rkr|ryrrrrrrruruesR)HPP3P P  P  rrucR\rSrSr%S\S'SrS\S'S SjrS SjrS rg) Scatteriroutput_indexerNrW scatter_modec UR5n[R"[SU5"U5n[ UUR UUR URURS9$)rr)r r r r rr) r-rrrrr r rrrs rrScatter.constant_to_devices[!!#n.?HP**;;..**   rcUR5nUcSn[R"UU"URU55U"U5URS9$)Nr)mode)r-rqrrrrs rrScatter.store_outputsT !!#  #Kyy  D''- . 4L""   rrr)rrrrrrrr ) rrrrrrrrrrrrrrsB44"L)"    3    rr logical_ormaximumminimummulr bitwise_xor)anyrMminprodsumdotxor_sumz"dict[str, Callable[..., OpsValue]]REDUCTION_COMBINE_FNc^^^T[;a [T$TS;aSUUU4SjjnU$TS:Xa SSjnU$[ST35e)Nargmaxargminc>Uup#UupET S:Xa[R"X$5nO[R"X$5n[R"X$5n[ T 5(a[R "X"5n[R "XD5n [R "U[R"X55n[R "U[R"X55nT (a[R"X55O[R"X55n [R "U[R"Xz55n[R"XbU5[R"XcU54$)Nr) rqltgteqr.ner logical_andwhere)aba_valuea_indexb_valueb_indexmaskequala_isnanb_isnantiearg_break_ties_leftr reduction_types rargmax_combine_fn3get_reduction_combine_fn..argmax_combine_fns ! G G)vvg/vvg/FF7,Ee$$&&2&&2~~dCFF7,DEucoog.OP'w(VVG-  >>$(CDD $1 $1 rwelford_combinec\Uup#nUupVnXR- nXG-n Xy- n X(U --X6-X-U-U --U 4$rr) rra_meana_m2a_weightb_meanb_m2b_weightdelta new_weight w2_over_ws rwelford_combine_fn4get_reduction_combine_fn..welford_combine_fns]&' "F(%& "F(OE!,J -I** emh6BB rzunknown reduction_type=)rtuple[object, object]rrrtuple[OpsValue, OpsValue])r#tuple[OpsValue, OpsValue, OpsValue]rrrr)rr)rr rrrs``` rget_reduction_combine_fnrs--#N33 / / $ )>  &  :! , , 2 2 1  "!"$;N;K"LMMrcJ^\rSrSr%S\S'S\S'S\S'S\S 'S!S jr\r\"S5S"S#U4S jjj5rS$S jr S%S jr S&Sjr S'Sjr S(Sjr S"S#SjjrS)Sjr\S*S+Sjj5r\S,Sj5r\\R*S4S-Sjj5r\S.Sj5r\S.Sj5r\S/Sj5r\S0Sj5r\S*S1Sjj5r\S2Sj5r\S3Sj5r\S*S4Sjj5r\S5Sj5rS r U=r!$)6 Reductionirreduction_rangesrVrr src_dtyperZreduction_hintc$URS5$)N)r rrr rs rr!Reduction.__str__s||LMMrc|>^[TU]T5[5R"U4SjUR56-$)Nc3<># UHn[UT5v M g7frr&rs rr1Reduction.get_free_symbol_uses.. s P:OQq-00:Or)superrvr>rrrrp __class__s `rrvReduction.get_free_symbol_usess7w+M:Z\=O=O P$:O:O P>   rcUR$r)rrs rr|Reduction.get_reduction_sizes$$$rcUR$r)rrs rryReduction.get_reduction_type"""rc [R"URURURUR X455n[R "U=(d SU"U5U5 gr)rq reductionr rrr store_reduction)rrrrreduction_varsrs rrReduction.store_reductionsR JJ NN    MM$ /   K49gdmUKrcX[UR5[UR5-$r)rr rrs r index_lengthReduction.index_length#s!4;;#d&;&;"<<.inner_reduction_splitsosrr>zUse 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: %sr8rrrc `^ UR5nUce[S[UUR5UR 5S9US9nUR 5nUR ceUR Vs/sH=n[U[5(dM[U[R5(aM;UPM? nn/nSn[URSS9Hm [U 4SjU55(dMURT R5 T R ["R$R&;dMd["R$R&T R n[)UR*SS5nUR-5 [)UR*SS5U:wdMSnM XV4$s snf) Nr r rrrUrTFcUR$rrrs r@Reduction.num_splits..get_read_indices..saffrkeyc3T># UHoTRR;v M g7fr)rr3)rr3mds rrAReduction.num_splits..get_read_indices..sF:aBHH111:%(r T)rrrrrrg range_varsrrrNumbersortedrkrrrrrsr name_to_bufferrrU decide_layout) r3r cb read_writesr indiceschangedbuforiginal_strider s @rget_read_indices.Reduction.num_splits..get_read_indicessn\\^F% %%%!++- B,,.K))5 55%///Aa&/9!U\\/J/  GG[..4DEF:FFFNN288,ww!''"8"88gg44RWW=*1#**h*M))+"3::x>/Q&*GF# #!s4F+F+2F+rc3B# UHoS:H=(d US:v M g7f)rrENrrs rr'Reduction.num_splits..s9AQ!a%)rrrrrr)r3rrztuple[Sequence[Expr], bool])*rnrsr r all_unbacked_explicitly_hintedrZrr has_featurerHREDUCE_TO_SINGLE_ELEMENTrFsplit_reductionsrYrmulti_processor_count functoolsr choicesreduction_split_factorINNERrrrrrrO!replace_backed_symbols_with_hintslogrrrrGindex_vars_squeezerr|simplify_with_ranges stride_hintsrkeysrOUTER)#r  dst_dtyperr r rrreduction_numel input_nodeexprsrr should_splitpropsnum_smmin_elements_per_threadrouter_reduction_splitsr new_rangesnew_reduction_rangesextracted_numel_hintr3rrrrrranges1 num_outer num_innerrjr\outers# r num_splitsReduction.num_splitsAs!-"78ww>>uEE ((!+ + ww//AA/RWW%%77 f8MN %/ ##FN,S,ST T ( ( ''  U " ((!+ + ''/,,"$ @I@Q@Q 00&$A "AJ@Q@Q 00&%A "  &)    &< " ?*+?LEz$**E11%*Z*K*K\\.2BDI5(00 ! $F,A. )!,JGQ w<1  ((!+ +'3'F'F JJL!..0( $NW  A  55aAAgg&&33>4 #7G 999EQ Q  y  &&(>$) !&&(>$) YJIs Q Qc^^^^^^[RRRT5m[ X#5mSUUU4SjjmUS;a4[ T[ RT55mSUU4SjjmU4Sj$TmT$)z1Convert inner_fn from a reduction to an pointwisec >^[R"TUU4Sj[R"TVs/sHn[ U5PM sn655$s snf)Nc36># UHnT"TU5v M g7frr)rrrvalue_fns rr=Reduction._unroll_reduction_fn..fn..s&#UF++#s)r"reduce itertoolsproductr)rr combine_fnrrCs` rr*Reduction._unroll_reduction_fn..fnsN##"+"3"3,<=,sArrc>UVs/sHn[R"U5PM nnT"X5[R"T"U5[R 54$s snfr)rexpandrq index_exprrint64)rrr flatten_indexr s rrC0Reduction._unroll_reduction_fn..value_fn sO4::6a%,,q/6:U+NN=#8%++F;s Ac>T"U5S$r6r)rrs rr0Reduction._unroll_reduction_fn..s E1r)rrrr )rrrrrr)rsr r  guard_int_seqr_fixed_indexerrr)r rrrrHrOrrCs`` @@@@r_unroll_reduction_fnReduction._unroll_reduction_fns77++99:JK-nH    1 1* 112BCM  ) 3E *  . -HIrc $^^^^^^[RRR[ T55mTS:Xa`SU4Sjjn U "S5U "S5U "S5U "S5S.mTT;d TS35eSUUU4Sjjn [ R UUU [U5S9$TS:Xa-TS;a SU4S jjn O SUU4S jjn [ R UTXS9$[T[5(at[T5[R:aW[ U5S:wd[UR5(a.TS :wa([ R UTURTTTU5US9$UR!UTUTUTTTU 5 upSU4S jjnU"U5nU["R$:XaU nUS :Xa\U ce[&R("[*SS5 [-U 5unnSSS5 WceWceUR/UTUTUTUUTU5 $US:aUR1UTUTUTTUUU 5 nSn[R2R4(a&[U[65(aSSjnU"U5nU(am[UR8[:5(d[UR855eUR8R<SUlTUl UUl!TUl"U$[6R [;UTTUTTUUS95nU$!,(df  GN5=f)za Create a reduction node. May split the reduction to multiple layers to expose more parallelism. rc$>T[R:Xa [U5$TR(a0[U[5(d[ U55e[ U5$[U[5(d[ U55e[U5$r) rrrprrrrrrr)valr-s rpy_cnst!Reduction.create..py_cnst2sm *9$00%c=99D49D9 :%%c;77BcB7s8OrrE)rrrrz* not supported for zero-dimension tensors!c8>[R"TTT5$rrqrq)rr-rrtypes_to_initss rconst_fn"Reduction.create..const_fnHs||ON$CYOOrrrJc2>[R"ST5$rr])rr-s rrReduction.create..fnVs<<955rcr>TVs/sHn[RRPM nnT"X5$s snfrrr7r8)rrreduction_indexr rs rrrb[s2=M&N=Muww||=MO&N#E;;'O$4rcl>[T5(aU$US:a[U[R5$U$r6)rrMrF min_num_split)rr.s r_maybe_increase_split/Reduction.create.._maybe_increase_splits2/** qy5&"6"677 rr8Nr>TcLUR5n[U5S:wag[[U55nU[R R ;ag[R R Un[U[5(dgURR5ceU$r6) rrnextiterrsr rrrrTry)cur_node read_namesbufnamers r_find_split_reduction/Reduction.create.._find_split_reductions"*!8!8!:J:!+#"4 #34Gagg&<&<<#''009C%c>::#88668DDDJrr)rYrrzbool | float | int)rrrrr)rrrr)rnrrzComputedBuffer | None)#rsr r simplifyrnrurrrr rrFunroll_reductions_thresholdrjrrUr>rZrrrrrO!create_multilayer_existing_rangescreate_multilayerr*mix_order_reductionrrTrr _split_size_original_inner_fn_original_ranges_original_reduction_ranges)r1r r-rr r rrrr/rZr_rhintrrir6r7rsplit_reductionrqr.r^s ` ` `` @@rrReduction.creates~$''**33MBR4ST a  $qz"1: qz O"_4 !""LM 4 P P##!F| $  a !5566 <<##Y$   0 0O$v'I'IIv&!+vfkk/B/B%' ##11.  $ nn            &e, ]22 2!N B;) ))n.>E3V40 0F) ))'3 3388 $  QY''  C#O}}00ZY5O5O'*$#8"<"/"6"6 BBO0012B/>.B.B.S.STU.V+5=2390=M:J !!1-#-    kFEs L Lc US;aL[U5(a [S5$[U5(ag[R"U5R $US;aL[U5(a [S5$[U5(ag[R"U5R $[U5(aSOSn[U5(aSOSnUUUUUX"U4X"U4[S5U4S .U$) N)rMrz-infF)rrinfTrrE)rrrrrwelford_reduceronline_softmax_reduce)r.rrr,riinforrM)rr zeroones rdefault_accumulatorReduction.default_accumulators . .e$$V}$!%(({{5)--- . .e$$U|#!%(({{5)---(//uQ&u--d1#40 $D1&+FmT%:     rc:US:Xag[RX5$)Nrr)rrrr s r default_valueReduction.default_value s! - -,,^CCrcUS:XaU$US::a*US::a$U[R:Xa[R$US::a*US::a$U[R:Xa[R$U$)Nr8i)rZr, OUTER_TINY)rrrs r_multilayer_second_step_hint&Reduction._multilayer_second_step_hintsg B;! ! C^^UunUGtmnTU-U-mSUU UU 4SjjnT (ac[T 5n[R"[R"TU5[R"T U55n[R"XST 5$U"5$)Nc$>T"TT"T/55$rr)rr new_indexrsrbodyCReduction._multilayer_wrap_loader..wrapper_fn..bodyasi');<.wrapper_fnZs"1 _*/ 'Y ?2_DG = =-o> vvNN7K8NN?K@zz$g66v r)rSequence[Symbol]rerrrr) rViewdynamic_reshape_indexerrsr r rrr5rB) r1rrr.rrrr/ dense_indexrrrs ` ` `` @@r_multilayer_wrap_loader!Reduction._multilayer_wrap_loaderEs>> Z .. / ((>> HHS%0! 4   # 6F    (rc^^^[ST55(d ST<35e[RU[U5[U5-5mSUUU4SjjnU$)Nc3*# UH oS:Hv M g7f)rENrrr3s rrDReduction._multilayer_wrap_loader_existing_ranges..ys3?a6?z8Only enabled for numel_hint == 1, found original_ranges=c >US[T5nU[T5SnT"UT"[U5[U5-55$r)rr) merged_indexnew_reduction_index original_idxrroriginal_rangesrs rrEReduction._multilayer_wrap_loader_existing_ranges..wrapper_fnsQ((>#o*>?L$S%9%;T"/UQUQ5$rr)rreintermediate_loaders rintermediate_fn;Reduction.create_multilayer_helper..intermediate_fns''A'A'AB BrNr)rrrerrrr)rfloat16bfloat16rrrrrr-rsr r rrnrrr)r1r r-rrrrr6r7rrrintermediate_dtype intermediaterrrs @rcreate_multilayer_helper"Reduction.create_multilayer_helpers$0 ??   !''          *668 C% C8J C  C WW%%77 o8VW 99 ~ -Cs?/C"DDDD (&!+C,@,B!C-#-   rc [U5n [XS- -U5n URXr5n URUUU UU U U 5nUR UUUUUU/UQUPU /UUU 5 $)rrE)rnrArrr)r1r r-rr r rrrrr/r.rrrs rrvReduction.create_multilayers&((89o;UC ##N>00        ++       f e  L     rc jURUUUUU5n URUUUU UU/UQUQUU SU 5 $)rr8)rr) r1r r-rr rrr6r7rrrs rru+Reduction.create_multilayer_existing_rangessc$@@   %   ++      % +o + +     rrrrrirr) rrrrrrrrrrrrzSequence[Sequence[Expr]]rr)r rr-rrrr zCallable[_P, OpsValue]r rrrrzReductionType | Literal['scan']r.rr/ IRNode | Nonertuple[ReductionHint, _IntLike]) r zrUrlrZrrrrrrrrrrvrur __classcell__rs@rrrs((!!!!NHK(  ) %# L L3 L L ) L  L=    %)jjjj) j # j - j8jj"j (jjX)N),)) ) 2 ))V)6(=(=$(NNN N % N  N)N&N&N"N NN`$/ &>DD$/D &DD%(:G  &4A >%)('(-(" (  (  (/("( ((TD($2  &  0  J8= = =  = ' = ( = $2= = ,= &= = &=  = = ~%)+ + +  + % +  + )+ &+ + &+ "+  + + Z$ $ $  $ % $ ( $ $2$ "$ ,$ &$ &$  $ $ rrc ^^^SUUU4SjjnU$)1A closure containing math to read a given elementc>Tb[U5[T5:Xde[U5[T5:XdeTn[UTT5Hup#nUS:wdMXU--nM U$r6)rr)rresultrstszrXrr s rr_fixed_indexer..indexer*sj!c%jCK&???5zSY&&&ufd3KCRQw(*4 r)rrrrr)rr rXrs``` rrTrT#s Nr INNER_FN_TYcz^\rSrSr%S\S'SU4SjjrSSjrSrU=r$) MultiOutputReductioni9r output_indexc >^[T5(aT4m[T5S:XaTSn OSU4Sjjn [T U] UUU UUUUUS9 Xlg)NrErc4>^^[UU4SjT55$)Nc34># UH o"TT5v M g7frr)rrr reduction_idxs rr@MultiOutputReduction.__init__..loader..SsHiR]33i)r)rr inner_fnss``rr-MultiOutputReduction.__init__..loaderPsHiHHHrr)rrrrrztuple[OpsValue, ...])callablerr__init__r) rr r-rr rrrrrrrs ` rrMultiOutputReduction.__init__<s I  " I y>Q q\F I# I4B I% I -))  )rc N[R"URURURUR X455n[ U[[45(d[U55eXPRn[R"U=(d SU"U5U5$r) rqrr rrr rrrrrr)rrrrrrrs rr$MultiOutputReduction.store_reductionas JJ NN    MM$ /   &5$-00>$v,>0(()"";#;)WT]ERRr)r)r rr-rrz#INNER_FN_TY | Sequence[INNER_FN_TY]r rrrrrVrrrrZrr) rrrrrrrrrr ) rrrrrrrrrrs@rrr9s#)#)#)7 #) " #) , #)&#)#)&#)#)JSS3S S ) S SSrrcj\rSrSr\\R S4SSjj5rSrg)OnlineSoftmaxReductionisNc ^^^^^^^[UUUUUUU4Sj[U555n U Hn U R5 M U $)z. Create the reduction disregarding splitting. c3p># UH+n[R[TTTTTSTTU5 5v M- g7f)rN)rrr) r output_idxr r-r r rrrs rr0OnlineSoftmaxReduction.create..sO 0    $$+"   0s36)rrr) r1r r-rr r r num_outputrr/resultsrs `````` ` rrOnlineSoftmaxReduction.createtsB   $J/   A IIKrr)r rr-rrrr r r rrrrrrrZr/rrSequence[TensorBox]) rrrrrlrZrrrrrrrrss)6(=(=$(!!! ! % !  !)!!&!"! !!rrc\rSrSr\\R 4SSjj5r\SSj5r \S Sj5r Sr g) WelfordReductionic^^^^^US;de[RR(aS5e[RR R [T55nS UUU4Sjjn US:XaU "S5n U "S5n U "S5n XU 4$US:XaDS UUUU4SjjmUS:XaT"US5U "S5U "S54$[U4SjU55$[RTTTUSTTUUS 9upU[R:XaU nUS:aURTTUTTUUU5$[S 5Vs/sH)n[R![#TTUTTUTUU5 5PM+ nnUHnUR%5 M U$s snf) N)rrzGwelford reduction usage is explicitly disabled, please check you configcV>^SUU4Sjjn[RTTU[T5S9$)Nc2>[R"TT5$rr])rr rYs rr 8WelfordReduction.create..const..inner_fns||rrrrrrrrurr)rYr r r r s` rconst&WelfordReduction.create..consts7   ##!F| $ rrrEcV>^SUU4Sjjn[RTTU[T5S9$)Ncr>TVs/sHn[RRPM nnT"X5$s snfrrd)rrrerrs rr 7WelfordReduction.create..copy..inner_fns2=M&N=Muww||=MO&N!#77'Orfrrr)rr r r r rs` rcopy%WelfordReduction.create..copys788!''!%< (rrc34># UH nT"U5v M g7frr)rrrs rr*WelfordReduction.create..s: "T"XX r)rr.r)rYrrr)rrrr)rFmtiadisable_welford_reductionrsr r rsrnrrr>rZrrvrrrrr)r1r r rr rrrr.rmeanm2weightr|rrrrrs `` `` @rrWelfordReduction.creates!FFFF;;88 U 8''**33MBR4ST   a 8DqB1XFV# # a  L   !11IaL)58U1X==: :::& **    aL  )++   ]22 2!N 19((   2$Ah '     $""  '  A IIK% s<0F cg)N)rrrrrs rrWelfordReduction.default_value src ^^^^^^^[T5m[RRR [ R "[TT5S55(+n U (aBUS:wa<S U4Sjjn TRUTUS[U SS9[U SS94UTSTUS9$[TTS- -T5m[RUT[UUUUU4SjU55/UQTPT/UU5n U Hn U R5 M S Sjm[RRR[U55n TR!TX5n[RUT[U4S jU 55UT/SU5$) rrrc2>[R"UT5$rr])rrrr s rrq4WelfordReduction.create_multilayer..constant5 s||E511rrrE)r r rr rrrrc 3L># UHnTRUTTTTSS9v M g7f)rrN)r)rrrr1r.rrs rr5WelfordReduction.create_multilayer..M s> (F++$# ,(s!$cU"/UQUQ5$rr)rrers rintermediate_loader_fnBWelfordReduction.create_multilayer..intermediate_loader_fn` s 4E4O45 5rc3T># UHn[TUR5S9v M g7f))rN)r r-)rrrs rrrp s&&A.q}}G&r )rrrrrrrrr)rrrerrrrrr)rnrsr r rrr5rBrvr rArrrrrr)r1r r rr rrrrrrq intermediatesrrrrr.s` ` ` ` @@@rrv"WelfordReduction.create_multilayer s ((89((>> HHS%0! 4   +<< 2# 24B 2KN 2 2 ((aLHA.HA. !10-)  o;UC (//     (   f e  L  # &A IIK 6! 6+ 69 6  6WW%%77 f8MN 99 :  &&   &   G    rrN)r rr rrSequence[Callable[..., Any]]r rrrrrVrrZrrr)r rr rrrr rrrrrVrrrrZrr) rrrrrlrZrrrrrvrrrrrrs)6(=(=yyy0 y  y ( y&y&y yyv$/ & \ \ \ 0 \  \ ( \ &\ \ &\  \ \ rrc^\rSrSr%S\S'S\S'S\S'S\S'S \S 'S \S 'S \S'S\S'\"S5S S!U4Sjjj5rS"U4SjjrS#SjrS$Sjr S%Sjr S%Sjr S%Sjr S&Sjr S'SjrS S!Sjjr\\R$4SS.S(Sjjj5r\S)Sj5rSrU=r$)*Scani| r scan_rangesr=Callable[[tuple[Any, ...], tuple[Any, ...]], tuple[Any, ...]]rHzFCallable[[Sequence[_IntLike], Sequence[_IntLike]], Sequence[_IntLike]]rrZrrrtuple[torch.dtype, ...]dtypestuple[Callable[..., Any], ...]rc>^[TU]T5[5R"U4SjUR56-[5R"U4SjUR 56-$)Nc3<># UHn[UT5v M g7frr&rs rr,Scan.get_free_symbol_uses.. O>N"1m44>Nrc3<># UHn[UT5v M g7frr&rs rrr% Hi"1m44ir)rrvr>rrrrs `rrvScan.get_free_symbol_uses s] G ( 7l  Od>N>NO l  HdiiH   rc>[UR5[UR5-[UR5:Xde[TU]5 gr)rr rrrrrrs rrScan.__post_init__ =4;;#d&6&6"773tyy>III rc"^URX45m[U4SjUR55n[R"UR UR U5n[R"U=(d SU"T5X`R5$)Nc32># UH o"T5v M g7frrrr rs rr'Scan.store_reduction.. D^x}}^r) rrrrqrr!rHrr)rrrr scan_varsrrrs @rrScan.store_reduction skll4+DT^^DD$++t?yy  $9gclF;L;L4M  rcg)Ncustomrrs rryScan.get_reduction_type srcUR$r)rrs rr|Scan.get_reduction_size rrcUR$rrrs rr Scan.get_size yyrcUR$rrrs rr-Scan.get_pointwise_size r'rcX[UR5[UR5-$r)rr rrs rrScan.index_length !4;;#d&6&6"777rcURUR5nURUR[R5nUR X5nU4$r)r:r rrDrrrrrrs rrAScan.inner_fn_args C DKK(T--t}}=ll5)v rcURUR5nURUR[R5nUR X#5n[ URXAS9$rR)r:r rrDrrrNr rrprrrs rrScan.inner_fn_free_symbols M DKK(T--t}}=ll5)#DMM3TTrT)can_fallback_to_atenc ^^^/USTQUTS-SQmUT/m[RRU[R5(dS/[ U5-$[ U5S:aB[RRU[R 5(dS/[ U5-$[RRn U R[T55n [ U5[ U5:XdeU R[R"U S55(a=[[ U55V s/sHn [RUX,X<US9PM sn $UR!UUSUSTTTUU S9up}["nU S:at[$R&R(SL=(d [*=(a [,S:=(a [ U5S:HnU(dU(aS/[ U5-$Sn O[.nSUUU4Sjjn[[ U55V s/sH/n [0RU"S UX,UX<UUTTUUUU S. U D65PM1 nn UHnUR35 M U$s sn fs sn f) NrErr)r r r axispointwise_rangesrrH scan_numelz3.3.0c>[U5[T5:Xde[U5[T5:Xde/USTQUQUTSQ$rr)r scan_indexrNrOrs rrScan.create..reindex Sz?c+&66 66u:%5!66 66>U5D\>J>tu> >r) r r r!r rrr rrHrrr)rrrSrrrr)rsr rrHSCANrTUPLE_REDUCTIONr rsrnrrr=rrurr>rrversionhip has_tritontriton_version SplitScanrr)r1r r!rrrNrHrrLrr rPrr> scan_typesupports_splitrrrrOrs ` @@rr Scan.create s=T%4[<4q +;<Dzl ww""6>+>+>??6CK' ' v;?177#6#6 N22$ $ 6CK' '77##&&}['AB 6{c)n,,,  ) )%((:q*A B B%*#f+$6%7L   ! .&4 ! %7 &)^^)q\-#!!&4 & " > !!T)Wj.V^w=V%v;!# "' 6CK//!"J%  ? ?.!&c&k 2% $!3 #   ! .!&4'+ +)##1!- "!3%  *F NN AP s ;$I3 6I8c L^^SUU4Sjjn [RUUUU UUSUS9$)Nc,>T"/USTQUQUTSQ5$rr)rrrNr s rr#Scan.num_splits..wrapper_fn7 s*Fc%4jF=F3tu:FG Grr)r r-rr r rrr.)rrrrrrr)rr>) r1r r r rNrOrrHrPrs `` rr>Scan.num_splits* sA H H###(!&$  rrrrir) rrrz%Callable[[Sequence[_IntLike]], Never]rrr4rrr rrrrh)r rr!r rz+tuple[Callable[[Sequence[Expr]], Any], ...]rrrNrrHrrrZrLrrr rSequence[TensorBox | None])r rr rr rrNrrOrrrrHrrPrrr)rrrrrr^rvrrryr|rr-rrArrlrZrrr>rrrs@rrr| s MM SS!! ##--F#   $       7    $    8 U )6(=(=`&*``(`? `  `  `R`&`#`` $``D   7    (  # R   (  rrc\rSrSrSrg)r\iG rN)rrrrrrrrr\r\G srr\cn^\rSrSr%S\S'S\S'S\S'S\S'S \S 'S \S 'S \S'S\S'S\S'\"S5SSU4Sjjj5rS U4SjjrS!SjrS"Sjr S#Sjr S#Sjr S#Sjr S$Sjr S%SjrSSSjjr\\R$4S&Sjj5rSrU=r$)'SortiL r sort_rangesrz:Callable[[Sequence[Expr], Sequence[Expr]], Sequence[Expr]]rrZrrrr r!r"rrstable descendingc>^[TU]T5[5R"U4SjUR56-[5R"U4SjUR 56-$)Nc3<># UHn[UT5v M g7frr&rs rr,Sort.get_free_symbol_uses..b r&rc3<># UHn[UT5v M g7frr&rs rrrme r(r)rrvr>rrhrrs `rrvSort.get_free_symbol_uses] s] G ( 7l  Od>N>NO l  HdiiH   rc>[UR5[UR5-[UR5:Xde[TU]5 gr)rr rhrrrr+s rrSort.__post_init__i r-rc6^URX45m[U4SjUR55n[R"UR XPR UR5n[R"U=(d SU"T5X`R5$)Nc32># UH o"T5v M g7frrr0s rr'Sort.store_reduction..u r2r3r) rrrrqsortr!rirjrr)rrrrrrrrs @rrSort.store_reductionm soll40DT^^DD$++v{{DOOLyy  $9gclF;L;L4M  rcg)Nrurrs rrySort.get_reduction_type{ srcUR$r)rhrs rr|Sort.get_reduction_size~ rrcUR$rr<rs rr Sort.get_size r>rcUR$rrrs rr-Sort.get_pointwise_size r'rcX[UR5[UR5-$r)rr rhrs rrSort.index_length rCrcURUR5nURUR[R5nUR X5nU4$r)r:r rhrDrrrEs rrASort.inner_fn_args rGrcURUR5nURUR[R5nUR X#5n[ URXAS9$rR)r:r rhrDrrrNr rIs rrSort.inner_fn_free_symbols rKrc  T^^^/USTQUTS-SQmUT/m[RRU[R5(dS/[ U5-$[RR n U R[T55n [RR(a[RRn OHSn [RR=(a% U R[R"X55n U (dS/[ U5-$[ U5[ U5:XdeU R[R"U S55(a=[![ U55Vs/sHn["R%UX.X>US9PM sn$SUUU4Sjjn[![ U55Vs/sH3n[&R%[)SUX.UX>UUTTUUUUUS. U D65PM5 nnUHnUR+5 M U$s snfs snf)NrErrc>[U5[T5:Xde[U5[T5:Xde/USTQUQUTSQ$rrR)r sort_indexrNrOrhs rrSort.create..reindex rUr) r r r!r rrr rhrrrrirj)rrrrrrr)rsr rrHSORTrr rsrnrFr*decompose_sort_opspersistent_reductionsrrr=rrurrrgr)r1r r!rrrNrirjrrr  sort_numelis_persistent_kernel max_rblockrrrrrOrhs ` @@rr Sort.create s!=T%4[<4q +;<Dzl ww""6>+>+>??6CK' '77##&&}['AB == + +#)==#F#F J 33U22588J3ST !$6CK' '6{c)n,,,  ) )%((:q*A B B%*#f+$6%7L   ! .&4 ! %7  ? ?0!&c&k 2' &!3 %   ! .!&4'+ +##1!-!) $!3'  ,F NN Q s $H  :H%rrrir) rrrrrrrrrr rrrr)r rr!r rz'tuple[Callable[[list[Expr]], Any], ...]rrrNrrirrjrrrZrr rrd)rrrrrr^rvrrryr|rr-rrArrlrZrrrrrs@rrgrgL sC  GG!! ##-- LF#   $       2    '    8 U )6(=(=QQ(Q; Q  Q  QQQ&QQ $QQrrgc:[USS9 g![a gf=f)NFfreezeT)rVrrs rrr s&a. s c[USS9upUR5(aUR5 UR5$![a gf=fNFr)rVshould_pad_strides pad_strides is_contiguousr)r_bufferrUs r is_contiguous_storage_and_layoutr sS/%@  $ $ & &    ##%% sAA AAc [U[5(a[URUUUUUS9$[U[5(a5[URUUUUUS9upgXRR 54$[U[ 5(aU(asU(a6UR5 UR 5R5(deO6UbURX4S9 O#UbURXTS9 OUR5 [ U5UR 54$[U[5(a#[URUS9upXR4$[e)z Try to simplify x into a StorageBox and a Layout. allow_padding only affect how we apply stride_order. When allow_padding is True, we have the freedom to add padding when applying the stride_order. rwant_contiguous stride_orderrUrcrUr)rrrVrTrxrryrQrrVrdrrZrUr) rrrrrUrcrrUbuffers rrVrV sJ!Y$ FF+%''   !Z  ) FF+%''  &&##%%%!V !||~335555)11 2*22!3!!}alln,,!_%%* FF xx rc^[USS9up#URU5$![a gf=fr)rVis_stride_orderedr)rrrrUs r"is_stride_order_storage_and_layoutr; s8/%@'' 55 s  ,,c[U[[45(a[UR5$[U[ 5(a}UR n[RRRUR[UR5-[5(+n[UR5=(d U$[U[5(a+UR!5[RR";$gr')rrrx is_unalignedrTrZrUrsr r statically_known_multiple_ofrXrer rgryrzunaligned_buffers)rrUhas_unaligned_layouts rrrE s$J/00DII&&$((#$77#3#3#P#P MMN6<<8 8/$  DII&>*>>$}}!''";";;; rc\rSrSr%S\S'\"S5SSSjj5rSSjrSSjrSSjr \ S Sj5r S!S jr S"S jr S#S jrS$S jrS%S jrS&SjrS'SjrS(SjrS)SjrS*SjrS'SjrS'SjrS+SjrS,SjrS-SjrS.SjrSrg)/rviW rrTc8URRU5$rrTrvrus rrvBaseView.get_free_symbol_uses[ syy--m<T"T"U55$rrrinnerrs rr&BaseView.make_indexer..indexerf & &r)rrrr)rTr1r)rrrrs @@rr1BaseView.make_indexerb s4 &&(%%' ' 'rcp^^URR5mUR5mSUU4SjjnU$)Nc >T"T"U55$rrrs rr$BaseView.make_loader..loadero rrr)rTr-r)rrrrs @@rr-BaseView.make_loaderk s4 %%'%%' ' ' rc6URR5$r)rTrrs rr BaseView.dtypet syy""$$rc6URR5$rrTrrs rrBaseView.get_layoutx yy##%%rc6URR5$rrTrrs rrBaseView.get_device{ rrcgrrrs rrBaseView.get_origin_node~ rrc6URR5$rrTrzrs rrzBaseView.get_name yy!!##rc"UR5$rr rs rr-BaseView.get_pointwise_size }}rc8URRU5$rrTrHrFs rrHBaseView.mark_reuse yy##E**rc6URR5$rrTr(rs rr(BaseView.has_exceeded_max_reads yy//11rc6URR5$rrTrrs rrBaseView.realize yy  ""rc8URR5 grrTrLrs rrLBaseView.realize_hint s  rc6URR5$rrTrrrs rrrBaseView.get_storage_numel yy**,,rc6URR5$rrTrrs rrBaseView.is_extern yy""$$rc[UR[5(d[UR55eURR 5$r)rrTrvris_module_bufferrs rrBaseView.is_module_buffer s9$))X..?TYY?.yy))++rc6URR5$rrTrrs rrBaseView.get_read_names yy''))rc[R"[SS5 [UR 5UR 55R sSSS5 $!,(df  g=fr=)rrrrPr-rrkrs rrBaseView.get_reads sD \\.*:D A&  " e B A As 2A A'czUn[U[5(a#URn[U[5(aM#U$r)rrvrT)rrs rrwBaseView.unwrap_view s1H%%AH%%rcUR5n[R"[SU5"U5n[ UUR 5UUR 5S9$r)r-rrrrurrrs rrBaseView.constant_to_device sN!!#n.?HP.."==?   rrNrrir*Callable[[Sequence[Expr]], Sequence[Expr]]rrrrrrrrrrrrrrrrr)rrrrrr^rvrr1r-rr rrrrzr-rHr(rrLrrrrrrrwrrrrrrvrvW s LJ'=(=C%%&&$+2#!-%,*  rrvc`\rSrSr%S\S'\S Sj5r\S Sj5rS Sjr S Sjr Sr g )ri rrc[RRnUVs/sHn[R"U5PM nnUR 5nS/[ U5[ U5- -[U5-n[ U5[ U5:Xde[[ U55HnXS:XaXEceXEX'MXEb0[RRRXE5(aMQ[U5(aMc[U5(aMuXnXEnUceUceXg- nURUSS9S:XaMSUR 5SUSU35e U$s snf)zReplace `-1` with correct sizesNr8rfallbackzBroadcast failed in ExpandView(, z) on dimension ) rsr r rrLrrrris_size_one_or_falser6r) rnew_sizer rold_sizerv1v2diffs r_normalize_sizeExpandView._normalize_size sy77##-56XELLOX6::<6S]S]:;d8nL8}H ---s8}%A{b {...&k $(8(8(M(M ))./99 [[~%~~%~w..!"/ 6ajjl^2hZ_`^ab-&<E7s E)cURX5n[U5(Ga [U5up4[U5[UR5- nUS:de[ R R/U-n[URUR5H^upxUR[RRRU5(dUO[ R R5 M` [UR UR"[%U5UUR&UR(5n [+X9S9$[-XS9$)NrrS)rTr)rrrVrrrr7r8rr rrsr r rrWr r rrXrYrZr) r1rrr]r^skipr_r rr`s rrExpandView.create s &&q3  # #"7": Gx=3z#77D19 9'',,$.J #J$5$5z G !!77++@@FF!H %!!  X!!$$ J#C Cq00rcUR$rr<rs rrExpandView.get_size r>rc^^UR5nURR5m[U5[T5- mSUU4SjjnU$)Nc>[UTS5n[U5[T5:Xde[[T55H*nTUS:XdM[RR X'M, U$r6)rrrrr7r8)rractualrs rr*ExpandView.make_reindexer..reindex s_tu&Eu:V, ,,3v;'!9>$ww||EH(Lrrrrr)rrTr)rtargetrrrs @@rrExpandView.make_reindexer sU##%6{S[( !    rrN)rrrrrr)rrrrrrvrr) rrrrrrrrlrrrrrrrrr sA %%N114 3rrc`\rSrSr%S\S'\S Sj5r\S Sj5rS SjrS Sjr Sr g ) PermuteViewi rdimsc URU5n[U5[[[U555:Xde[ U5(a[ U5up4[ URURUVs/sHoTRUPM snUVs/sHoTRUPM snURUR5n[X6S9$[XS9$s snfs snf)NrS)rTr ) _map_neg_dimsr>rrrrVrWr r rr rXrYrZr )r1rr r]r^rr`s rrPermuteView.create s  &$:eCI.>#????  # #"7": G$!!  -12T#T2/34t!""1%t4!!$$ J#C C--34s :C C$cXUVs/sHo"S:aUO [U5U-PM sn$s snfrrR)r1r rHs rrPermuteView._map_neg_dims/ s+@DEaxSY_4EEEs'c[URUR55[[[ UR555:XdeUR R 5nURVs/sHo!UPM sn$s snfr)r>rr rrrTr)rrrs rrPermuteView.get_size3 sq$,,TYY78J #dii. !=    yy!!#!%+AQ+++s8B cZ^[UR5VVs0sHupX!_M snnm[[UR55Vs/sHnTUPM snm[ T5[ [[UR555:XdeSU4SjjnU$s snnfs snf)Nc8>TVs/sHoUPM sn$s snfrr)rrinvs rr+PermuteView.make_reindexer..reindexA s'**c!Hc* **sr)rr rrr>)rrr<rrs @rrPermuteView.make_reindexer: s!*$)) 45 4qt 45$S^454!s1v45#*U3tyy>-B"CCCC +! +  + 65s B" B(rN)rrr rrrv)r rr list[int]rr) rrrrrrlrrrrrrrrr r  sB ..$FF,  3 rr cV\rSrSr\SS.SSjj5r\S Sj5rS SjrSr g) SqueezeViewiI N)rHc [U5(Ga][U5up4/n/nUbF[U[5(d[ U55eSU::aU[ UR 5:de[[UR UR55HunupUcT[RRRU5(d$URU5 URU 5 M]M_Xr:wa$URU5 URU 5 MUS:XaMS5e [UR UR"UUUR$UR&5n [)X:S9$Uce[*R-UUR/5V s/sH5n [RRRU 5(aM3U PM7 sn 5$UR/5US:Xde[*R-U[UR/55VV s/sHup{Xr:wdM U PM sn n5$s sn fs sn nf)NrrEzexpected squeezed size to be 1rS)rrVrrrrrrrr rsr r rrrWr r rXrYrZrrr) r1rrHr]r^rr_rrr r`rs rrSqueezeView.createK s  # #"7": GHJ!#s++6T#Y6+CxC#joo*>$>>>%.s:??JDUDU/V%W!>D;77++@@FF -"))&1Gx -"))&1#qyJ*JJy&X%!!  !!$$ J#C C ;;;ZZ\)77++@@C) ::<$) ));;q1::<1H"U1HAH11H"UV V#Vs2H9 H9  H> +H> c^^UVs/sH oS:wdM UPM nn[U5VVs/sHup1US:wdM UPM snnm[U5mSUU4SjjnX$4$s snfs snnf)NrEc>[U5[T5:Xd UST35e[RR/T-n[ TU5H up#X1U'M [ U5$)N )rrr7r8rr)rrrrlengthnot_ones rr%SqueezeView.squeezer..reindex sbu:W- C%'/C C-%*WW\\NV$;Igu-!"#.# #r)rrrztuple[Expr, ...])rr)rrrrrr"r#s @@rsqueezerSqueezeView.squeezery sc $.t!AvAt.!*4;AF1;T $ $  /;s AA A A c[S5e)Nzuse SqueezeView.create())AssertionError)rrTs rrSqueezeView.__init__ s788rr)rrrHrrr)rrrz>tuple[list[int], Callable[[Sequence[Expr]], tuple[Expr, ...]]])rTr rr) rrrrrlrrr%rrrrrrrI sC48+W+WZ!! G!! 9rrc~\rSrSr%S\S'S\S'S SjrSSjrSSjr\r\ SS j5r SS jr S r g ) GenericViewi rrrrcUR$r)rrs rrGenericView.make_reindexer s||rc[[UR55Vs/sHn[[R U5PM nn[ URU55nSSR[[U55SU3$s snf)Nzlambda rr!) rrrrmrDrmrrrrr)rr9 index_old index_news r reindex_strGenericView.reindex_str svCHTYYCX CXa *4::q 9CX  i01 3sI#6789+FF  s$BczURURSUR3SUR53/5$)Nsize=zreindex=)rrTrr1rs rr!GenericView.__str__ s= YY% {+x8H8H8J7K-L M  rc$U"U[U5US9$)NrTrr)r)r1rrrs rrGenericView.create sX@@rcUR$rr<rs rrGenericView.get_size r>rrNrr)rrrrrrrrvr) rrrrrrr1r!rkrlrrrrrrr+r+ sp  77 3 G HA A!A< A  AArr+c\rSrSrSr\S Sj5r\\S Sj55r \S Sj5r \S SSjj5r \S SSjj5r S r g)ri z This class handles tensor reshaping by computing appropriate index transformations to map the new shape back to the original storage layout. c[R"U5n[R"U5n[RRR R nU"[R"US55(aX-nU$r)rrLrsr r r evaluate_exprLt)rrr=s rhandle_negative_indexView.handle_negative_index s[ll3||D!((22@@ #q) * **C rc^^^^[T[5(d[T55eTRUR 5T5umm[ R RRTT5(aU$[[T55S:=(d [[T55S:n[U5nS SjmS UUUU4SjjnST;aSU4SjjnT"U[T5US9$U(aT"UT[RT55$[U5(dU"U5$[!USS9upxUR"n [ R RR%T5n [ R RR%U 5n [ R RR%T5n SSKJn U "U U U US 9nUbUVs/sH@n[+US 5(aUR,R.O[0R2"U5PMB nn[5UR6UR8TUUR:UR<5n[?UUS 9$U"U5$s snf)Nrc[USS9up4[URURUUURUR 5n[ X5S9$)NT)rrS)rVrWr r rXrYrZ)rsrr_r]r^r`s rcreate_reinterpret_view,View.create..create_reinterpret_view sT#8T"R G$!!  !!$$ J#C Crc>TRTT5nT"U[T5US9$![a5 [R U5nT"UT[ R T55s$f=f)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. r7)rrr5rrequire_contiguousrr)rrr1rCrrs r"handle_unbacked_or_dynamic_reshape7View.create..handle_unbacked_or_dynamic_reshape sm 55hIXHH. !33A6.x!B!B8!L  s"&[S/[T5-5$r)rr)rrs r fake_reindex!View.create..fake_reindex saS3x=011rr7Fr)_compute_stridesize_obliviousrrS)rsrrrr_rrrZrrrr)rr rztuple[int, ...]) rrrresolve_negative_sizerrsr r statically_known_list_equalsrr4rrrrrrVr to_symints_or_intstorch._subclasses.fake_implsrLrrexprrr rWr r rXrYrZ)r1rrunbacked_symbols_in_sizesrrGrJr]r^ old_strideold_size_symintold_stride_symintnew_size_symintrLnew_stride_symintrr_r`rCrs` ` @@rr View.create so(H--=tH~=- 66qzz|XN( 77   8 88 L LH %h/ 01 4 8(23a7 "9;  D D#1 D?M D  D     , = 2ADNLI I*8^>>xH  %Q''5a8 84AeD&& ''**==hGGG,,?? K''**==hG@,   4    (+*A 'q&11 u}}Q7GG*  %!!  !!$$ J# C C2!44!sAI1cPUVs/sH,n[RRRU5PM. nnUVs/sH,n[RRRU5PM. nn[ U5n[ [ U55HInXS:XdM [RRX'[[U5[U55X' O [RRR[U5[U55 X4$s snfs snf)Nr8) rsr r rsrrrrr7Oner@rn check_equals)rrrrs rrPView.resolve_negative_size2 s;CC(QAGG$$--a0(C:BC(QAGG$$--a0(C>s8}%A{b #ggkk &}X'> h@WX  & %%mH&=}X?VW!!DCs 3D3D#NcURXU5nU$![[[4a= [ U5/nURX5nURXR5n[ Xg5nU$f=fr)_dynamic_reshape_indexerr(r5 IndexErrorrnr)r1rr dense_dimrflatrrs rrView.dynamic_reshape_indexerC su :228yQG  ;ZH :!(+,D33HCH33DCH%h9G  :sAA('A(c`^^^[RRRmSU4Sjjn[ [ U55Vs/sHn[ [RU5PM snm[[TU55n[U5nUSL=(a& U[ U5S- :g=(a [ U5S:HnU(a'UceURU5nURU5 /mU(GaU(GaUR5n UR5upU S:Xa=TR[RR5 URX45 GOzU S:XaURU 5 GOaU"X5S:XaTRU 5 GOBU"X5S:arU"X5S:a+UR5upX-U -n X-n U"X5S:aM+TRU 5 [RRR!X5 OU"X5S:a[RR"nU nTR[%XU55 X-nU"X5S:aAUR5nTR[%XU55 X-nX-n U"X5S:aMA[RRR!X5 O[&eU(a U(aGMU(alUR5n [RRR!U S5 TR[RR5 U(aMlU(aEUR5up[RRR!U S5 U(aMEUbB[ U5S:Xa3TR)5 TR5nTR+UU5 OTR)5 [ T5[ U5:XdeSUU4SjjnU$s snf)z7 Perform a reshape entirely by modifying indexing math c>T"[R"X55(agT"[R"X55(agT"[R"X55(ag[R R RX5(ag[R R RX5(ag[[R"X55e)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. rr8rE) rr5r>Gtrsr r rr5)rrr>s r compare_sizes4View._dynamic_reshape_indexer..compare_sizes_ sehhqn--ehhqn--ehhqn--ww<^[U5[T5:Xd[U5[T545e[[TU55m[U4SjT55$)Nc3<># UHn[UT5v M g7fr)ro)rr replacementss rrAView._dynamic_reshape_indexer..reindex.. sHiA|44ir)rrrr)rrmr view_exprs @rr.View._dynamic_reshape_indexer..reindex sOu:T* CSZT,C C*D% 01LHiHH Hr)rrrrrrr)rsr r r>rrrmrDVIEWrrr0rrr7r8r^r]rCr(reverseinsert)rrrcrir stack_new stack_oldreordering_dense_dimold_dimsize_oldvarsize_newvar2 size_new2divisormodulus dense_exprrr>rros @@@rraView._dynamic_reshape_indexerT su))88 >JCHH BV BVQ *499a 8BV T8,- N  T ! #S^a// #H "  ( ((mmI.G   W % I }}H%MMOMC1}  .  #1Q  *x2a7  %x2Q6#H7!;&/mmoOD/C/C'3H$H7!;  %  --hAx2Q6''++"  w!GH!+#H7!;'mmoG$$_S7%KL%/G'1H $H7!;   --hA$$;II> }}H GG   ) )(A 6   UWW\\ *i %MMOMC GG   ) )(A 6i  S]a%7    "J   Y 3    9~X... I! I  I IY s$P+r)rrrrrr)rrrrrr)rrrrrztuple[list[Expr], list[Expr]]r)rrrrrcrr&Callable[[Sequence[_T]], Sequence[_V]])rrrrrcrrr)rrrrrrr?rlrrrPrrarrrrrr s  k5k5Z" ",:" &""  !% $%  0  !%z z zz 4 zzrrc^\rSrSr%SrS\S'SU4SjjrSSjr\rSSjr SSjr SS jr \ SS j5r SS jrSS jrSS jrSSjrSSjrSSjr\"S5SSSjj5rS S!SjjrS"SjrSrU=r$)#rZi z*Pretend our storage has a different layoutrrUc>[TU]5 [UR[5(a0[ R USURR55 gg)NrT)rrrrTrvrrrwr+s rrReinterpretView.__post_init__ sC  dii * *   tVTYY-B-B-D E +rcPURURUR/5$r)rrTrUrs rr!ReinterpretView.__str__ s&     rc6URR5$rrrs rrzReinterpretView.get_name rrc.URR$r)rUr rs rrReinterpretView.get_device s{{!!!rcgrrrs rrReinterpretView.get_origin_node rrc.URR$r)rUr rs rr ReinterpretView.dtype s{{   rc@[URR5$r)rrUrrs rrReinterpretView.get_size sDKK$$%%rc@[URR5$r)rrUr rs rr:ReinterpretView.get_stride sDKK&&''rc^SU4SjjnU$)NcZ>TRR5n[R"TR 5U"U55nTRR TR R :wa6[R"UTR TR R 5$U$r)rUr1rqloadrzr rTto_dtype_bitcast)rr tmp_loaderrs rr+ReinterpretView.make_loader..loader spkk..0G$--/75>BJ{{  DIIOO3++J DIIOOTT!!rrrrrrrrrs` rr-ReinterpretView.make_loader s " rc6URR5$r)rUr1rs rr1ReinterpretView.make_indexer{{''))rcUR$rrUrs rrReinterpretView.get_layoutr'rcgrrrs rrQReinterpretView.freeze_layout rJrc[URRU5[URRU5-[URRU5-$r)r'rUrr rXrus rrv$ReinterpretView.get_free_symbol_usessQ T[[--} =t{{11=A Bt{{11=A B rc p[RRRURUR R UR RUR RUb URO#[RRRUR RS9$rw) rsr  wrapper_codecodegen_reinterpret_viewrTrUrr rX writeliner rs rr!ReinterpretView.codegen_referencesww##<< II KK   KK   KK   & 2F  8L8L8V8V++## =  rcgr6rrs rroReinterpretView.num_reads%rrrrrrrrrrrrrrrr)rrrrrrrr!rkrzrrrr rr:r-r1rrQr^rvrrorrrs@rrZrZ s4 NF  H$"!!&( * -.$) ! ! /   rrZcn\rSrSr%SrS\S'\S Sj5rS Sjr\r \ SSj5r SSjr SS jr S rg ) DtypeViewi)z(Pretend our storage has a different typer target_dtypec[U5(aX[U5up4[URUURUR UR UR5n[X5S9$[XS9$)NrS)rTr) rrVrWr rr rXrYrZr)r1r new_dtyper]r^r`s rrDtypeView.create/sj  # #"7": G$!!!!!!$$ J#C Ca88rcPURURUR/5$r)rrTrrs rr!DtypeView.__str__>s  4+<+<=>>rcUR$r)rrs rr DtypeView.dtypeCs   rc6URR5$rrTrrs rrDtypeView.get_sizeGrrcP^^TRR5mSUU4SjjnU$)Nc|>[R"T"U5TRTRR5$r)rqrrrTr )rrrs rr%DtypeView.make_loader..loaderMs*''c D4E4EtyyW WrrrTr-)rrrs` @rr-DtypeView.make_loaderJs( %%' X X rrN)rrrrrrvrrrr)rrrrrrrlrr!rkrr rr-rrrrrr)sE2 9 9?H !!$rrcv\rSrSrSr\SSj5r\SS Sjj5rSrg) SliceViewiSziView that represents a slice along a single dimension. Corresponds to tensor[..., start:end:step, ...]. c^^^^^ ^ [RRm UR5Um[ SX4T455(a![ R m [ RmOR[ SX4T455(a![ R m [ RmOT Rm T RmSUU U 4SjjmSUUU4SjjnU"USTS5nU"XCTT5nX44$)zb Normalize start and end such that both are in the range [0, x.get_size()[dim]] and start <= end. c38# UHn[U5v M g7fr)r4rrs rr0SliceView.normalize_start_end..dsH1GA$Q''1Grc3# UHJn[U[5(dMUR[R[R 5v ML g7fr)rrhasrMinMaxrs rrrgs; ,!T" (AEE%))UYY ' '+s A5Ac>TRX5(aUOT"X5nTRX25(aUnU$T"X25nU$r)statically_known_geqrD)rlowerupper clamped_lower clamped_fullmax_funcmin_funcr s rclamp,SliceView.normalize_start_end..clampss^221<<(1BT  00FF   m3   rcD>UcU$TRUT5nT"XU5$r)r?)rYrrrrr1dim_sizes r clamp_wrap1SliceView.normalize_start_end..clamp_wrap~s.{++C:CU+ +rr)rrrrrrrr) rYrrrrrr Expr | intrr) rsr r rrrrr evaluate_min evaluate_max) r1rrHstartendrrrrrr s ` @@@@@rnormalize_start_endSliceView.normalize_start_endYs77##::<$ H%h1GH H HyyHyyH  (+   yyHyyH,,H,,H  , ,$' ,03 ,>H ,  , ,5!Xq1Xx8zrc ^^^^ [R"T5m[T[5(d TS:dT5eTS:XaUS:aTS:XaU$[ UR 55m U(aURUTTU5umn[UT- TS- -T5T T'[U5(a[U5upx[ UR5n U TT-U T'[URURT U URURTT--UR 5n [#XzS9$SUU UU4Sjjn [%UT U S9$![a GNf=f)NrlrErSc>[U5[T5:Xd SUST35e[U5nUTT-T-UT'U$)Nz wrong ndim r!)rr)rrHrrsteps rr!SliceView.create..reindexsRu:X. P+eWAhZ0P P.KEsd*U2E#JLrr7r)rrLrr TypeErrorrrrrArrVr rWr r rXrYrZr) r1rrHrrrrr]r^r_r`rrs `` ` @rrSliceView.createsm||D!$%%7471 zcY.419 % 00CDJE3 uq!94@  # #"7": Gj//0J(o4JsO$!!  !!J$5$5c$:U$BB$$ J#C C !    ah@@K   sE EErN) rrrHrrrrrrztuple[int, int])rET)rrrHrrrrrrrrrrr) rrrrrrlrrrrrrrrSs // /),/36/ //b3A 3A3A 3A  3A  3A3A 3A3ArrcR\rSrSr%S\S'S\S'S SjrS SjrSSjrSS jrS r g ) BaseConstantirr rr cgNrrrs rrBaseConstant.get_sizesrcUR$rr%rs rrBaseConstant.get_devicer'rcgrrrs rrBaseConstant.get_origin_noderrc[5$rr=rs rrBaseConstant.get_readsrrrNrrrr) rrrrrrrrrrrrrrrs"  rrcR\rSrSr%S\S'S\S'S\S'S SjrSS jrSS jrS rg )Constantir rrr rr c^SU4SjjnU$)NcZ>[R"TRTR5$r)rqrqrr rrs rr$Constant.make_loader..loaders<< DJJ7 7rrrrs` rr-Constant.make_loaders 8 rcgrrrs rrConstant.realizerJrc@[URURUS9$)N)rr r )rrr r#s rrConstant.constant_to_devicesdjj 6JJrrNrrr) rrrrrr-rrrrrrrrs# J    KrrcH\rSrSr%S\S'S\S'S\S'S SjrS S jrS rg )IndexingConstantir rrr rr c^SU4SjjnU$)NcZ>[R"TRTR5$r)rqrMrr rs rr,IndexingConstant.make_loader..loaders>>$**djj9 9rrrrs` rr-IndexingConstant.make_loaders : rc@[URURUS9$)N)rr r )rrr r#s rr#IndexingConstant.constant_to_devicesdjj 6RRrrNrr)rrrrrr-rrrrrrrs J   Srrc`SnSn[[[X555HupEUS:XaM [RR R XR5(d0[RR R XS5(d gU[R"SU5-nX$-nM gNrEFT) reversedrrrsr r rOrr)r rGexpected_strideexpected_stride_maxrys ris_contiguous_strides_for_shaper sOs5123 6 ww77   ''""::1RRuyyA.4 rc<[RUR-$r)rFpadding_alignment_bytesitemsizers rget_align_for_dtyper  s  ) )U^^ ;;rcB\rSrSrSrSSjrS SjrS S SjjrSrg) r#iztAbstract base for Layout, MultiOutputLayout, NoneLayout. Represents the memory layout of the output of an Operation.c>[[U5R5err,rs rrOutputSpec.get_devicer/rc>[[U5R5err,rs r storage_sizeOutputSpec.storage_sizer/rc>[[U5R5err,rus rrvOutputSpec.get_free_symbol_usesrXrrNrrrr) rrrrrrrrvrrrrr#r#s1C77%*7!7 !77rr#c\rSrSrSrS\"S5S4SSjjr\SSj5r\RS Sj5r\SS j5r \ RS S j5r \S!S j5r \ RS"S j5r S#S jr \ r S$SjrS%SjrS&Sjr\S'Sj5rS&SjrS(SjrS&Sjr\S)Sj5rS*SjrS&SjrS+SjrS,SjrS-SjrS!Sjr\"S5S.S/Sjj5rSrg)0riz_ Layout base class Carries tensor meta-information including offset and whether it is pinned. NrFc^Uc[RU5nXlX l[ U5[ U5:Xd SUSU35e[ SU55(deX0lX@lXPlX`l UR(a"URRS:XdS5egg)Nr4 , stride=c3N# UHn[U[[45v M g7fr)rrrrs rr"Layout.__init__..7s#66t>*4>>*:;MDz""#2dkk&6&6%78H7ITZZLYII;i }VH]O1 N rcUR$rr%rs rrLayout.get_deviceir'rc [R [R"[ UR 5[ UR 5URURURS9sSSS5 $!,(df  g=f)N)r r  pin_memory) rs fake_moderrrarr r r rYrs r get_exampleLayout.get_examplelsN [[&&' 2' 4jj{{>> [[s AA77 BcB[URUR5$r)r r rrs rrLayout.is_contiguousvs.t{{DIIFFrc[U5nUS;d USS:Xag[U[U5U5Hup4nUS:wdMX4:wdM g g)N)rrEFT)rrr/)rGr\ndimleftrightrs ris_channels_last_contiguous"Layout.is_channels_last_contiguousysY5z v qQ!$  *5 1  " DqyT]" rc [UR[[R [ [UR 5555UR 5HupnUS:wdMX:wdM g gr)rr rrrrr)rrArBrs r is_transposedLayout.is_transposedsZ!$ KK ^66tHTYY.sorted_indicess,J=@AS'$$W-SA AAs-r8TrMF)rJrrr) rr rrrsr r rrrr _shape_envr=) rrrrH non_1_indicesr rMstride_orderedrTs rrLayout.is_stride_orderedsl4;;3u:--- $DII. .ww11#1BaG .  +88-Q++a.-82?@-Qa-@ B u%E *s5z"A'-yN8 $#s5zA~&A!$~!e'<.s* 6 q3 . / /6rrONc>Tcg[U[R5(dg[U4SjUR55$)NFc3F># UHnTRU5v M g7fr)is_unbacked_symint)rrrs rrILayout._pad_strides..contains_unbacked_symints..s!R@Q1y33A66@Qs!)rrrrr3)rTrs rcontains_unbacked_symints6Layout._pad_strides..contains_unbacked_symintss7 dEJJ//R@Q@QRR Rrc34># UH nT"U5v M g7frr)rrr_s rrrZsN:a6q99:rrE)rT)rTzsympy.Expr | intrr)r rrFpad_channels_lastrrCrsget_current_noderrirrrFchainpad_dynamic_shapesr rOrrrrrrrrr padding_stride_thresholdrr_r%num_comprehensive_padding) in_stridesrr aligncurrent_fx_noderirrr new_stridespaddedrankrprev_idxr require_paddingr_rs @@r _pad_stridesLayout._pad_stridess-$E* z?a  ''F,N,N - -  ,,. ?F + +0D0D0H0H u1 1   __Z6   ((Z *1!''<*H*HAGG&&d  S N:NNN ' I> ,\: "'J"89"8Qq"8 9&' qM"":ab>;ID!(+H *T^;F6C#78(V<<<(Na'PVUZZ0NV5N5N  & #*6#9E#A  <  ))Q.);:s5 I5c[U[5(d[U55eURceUR URUR UR 5Ulgr)rrrr rprr rs rrLayout.pad_strides sP$//;d;/{{&&&'' TYY K rcP[R=(a [U[5$r)rFcomprehensive_paddingrrrs rrLayout.should_pad_stridess++P 40PPrc[U[5(aU$UR5(aUR5 [URUR UR URURUR5$r) rrWrrr r rr rXrYrs ras_fixedLayout.as_fixedsf dK ( (K  " " $ $     KK JJ II KK KK NN   rc[R(dS[U5RS35eUR 5R 5$)Nzconvert z to FixedLayout first)rr>rrrxr1rs rr1Layout.make_indexer$sG,, tDz**++@ A ,}}++--rc[U[5=(a URUR:H=(a URUR:H=(ay URUR:H=(aY UR UR :H=(a9 UR UR :H=(a URUR:H$r)rrr r rr rXrY)rothers r__eq__ Layout.__eq__*s uf % 2 u||+ 2 ekk) 2 UZZ' 2 u||+  2  u||+  2 %//1 rcX[URURUR5$r)r+rr rXrs rrLayout.storage_size5s.tyy$++t{{SSrc[URU5[URU5-[URU5-$r)r'rr rXrus rrvLayout.get_free_symbol_uses8s= TYY 6t{{M: ;t{{M: ; r)rrrr r rYr )r rr rrrr zSequence[Expr] | NonerXrrYrrrrrrrrrrrrrrr)r torch.Tensorr)rGrr\rrr)rrrr)rhrrrr rrrrrrWr)r}rrrrr) rrrrrr rrrsetterr rXr!rkrr:rrrCrFrrUrprrrxr1r~rr^rvrrrrrrs)-qz     &       2 [[ ]] ]]  HG!,> "!F- L!L)7L@KL LL\L Q  .  TH%$) ! ! & rrc"\rSrSrSrSSjrSrg)rWiCz A Tensor layout we cannot changecX[URURUR5$)r)rTrr rXrs rr1FixedLayout.make_indexerFsdiidkkBBrrNr)rrrrrr1rrrrrWrWCs *CrrWc^\rSrSrSrSrSSjr\SSj5r\SSj5r \SSj5r \SSj5r \SS j5r \ SS j5r\RSS j5r\ SS j5r\RSS j5r\ S Sj5r\RS!Sj5rS"S#SjjrS"S$SjjrS%SjrS&SjrS'SjrS(SjrS)S*U4SjjjrSrU=r$)+riKzp A Tensor layout that we are allowed to change Assumption: layout change should NOT add or remove free symbols FcJ[R"U5R5$)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. )rdeepcopyrxrs r!get_fixed_layout_without_freezing0FlexibleLayout.get_fixed_layout_without_freezingTs}}T"++--rc[U5S:Xa/$[RR/n[ USS5HnUR X!S-5 M [ [ U55$)NrrEr8)rrr7r]rrr)sizesreversed_stridesrs rr!FlexibleLayout.contiguous_strides^s^ u:?I!GGKK=U12Y'D  # #DB+?$? @(H-.//rc[[[U555[U5:XdX45e[RR nS/[U5-nUHnX#U'X U-nM U$)z Create a stride based on the order the dimensions should be filled in. In this format, channels last would be: [1, 3, 2, 0] N)r>rrrr7r])rr next_strider\rs r fill_orderedFlexibleLayout.fill_orderedgsm%E +, 50AAQE>QAggkk &3u:%A$AJ%a0Krc[[[U555[U5:Xde[U5n[R X5$)zz Create a stride based on the sorted order of a permuted range. In this format, channels last would be: [3, 0, 2, 1] )r>rrrrr)rrrs rrQFlexibleLayout.stride_orderedxs@%E +, 50AAAA,U3 **5==rcDU[R:Xa[RU[5$U[R :Xa[RU[ 5$U[R:Xa[RU5$[RSU5 [e)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) r channels_lastrrQNHWC_STRIDE_ORDERchannels_last_3dNHWDC_STRIDE_ORDERcontiguous_formatrr'rr)r memory_formats r stride_ordered_for_memory_format/FlexibleLayout.stride_ordered_for_memory_formats E// /!008IJ J e44 4!008JK K e55 5!44U; ; IIP & %rc[U5[U5:Xde[RRR U5n[ [ [U55URS9n[RX5$)z Create a stride that has the same stride order as given stride For example, if given stride is [1000, 1, 100, 10], the fill order should be [1, 3, 2, 0] r) rrsr r guarding_hints_or_throwrr __getitem__rr)rr rs r same_orderedFlexibleLayout.same_ordereds`5zS[(((!!99&AE#f+.F4F4FG **5==rcUR$rr!rs rrFlexibleLayout.sizer#rc4URSU5 Xlg)Nr)!assert_free_symbol_uses_unchangedrr%s rrrs ..vu= rcUR$rr'rs rr FlexibleLayout.strider)rc4URSU5 Xlg)Nr )rrr%s rr r ..x? rcUR$rr-rs rrXFlexibleLayout.offsetr)rc4URSU5 Xlg)NrX)rrr%s rrXrrrcBURURU5nUR5(a-U(a&URX0RUR5n[ UR URURUURUR5$r) rQrrrpr rWr rXrY)rrrUr_s ras_stride_orderFlexibleLayout.as_stride_ordersx((E:  " " $ $**:yy$**MJ KK JJ II  KK NN   rcUnUR5(a-U(a&URX0RUR5n[ UR URURUUR UR5$r)rrprr rWr rXrY)rrcrUr_s ras_exact_stridesFlexibleLayout.as_exact_stridessg#  " " $ $**:yy$**MJ KK JJ II  KK NN   rc4URURU5nUR5(a&URX RUR5n[ UR URURUURUR5$r) rrrrpr rWr rXrY)rrr_s r as_fill_orderFlexibleLayout.as_fill_orderst$($5$5dii$G  " " $ $**:yy$**MJ KK JJ II  KK NN   rc4URURU5nUR5(a&URX RUR5n[ UR URURUURUR5$r) rrrrpr rWr rXrY)rr r_s r as_same_orderFlexibleLayout.as_same_orderst&&tyy&9  " " $ $**:yy$**MJ KK JJ II  KK NN   rc t0nSH/nSH&nX#4n[[[X5U55X'M( M1 U$)N)rr rXTF)r>r'r)rinitial_free_symbolsrrprs rget_initial_free_symbol_uses+FlexibleLayout.get_initial_free_symbol_usessI!0D!. +,6$WT%8-H-$)"/1$#rcSH9nURX4n[[X#55nXT:XaM.SUSU35e g)Nrz)Expected free symbols unchanged, but got z vs )rr>r')rrrrpold_free_symbolsnew_free_symbolss rr0FlexibleLayout.assert_free_symbol_uses_unchangedsX*M#88$9NO )*:5*PQ #7 ;U(a[RX45nO[RU5n[TU]XX6US9 UR 5Ulg)NrY)rrrrrrr)rr r rrrYr\rs rrFlexibleLayout.__init__sO $11$EG$77=G K%)$E$E$G!r)rrrrr)rrrr)rrrrrr)rrrrrr)rrrztorch.memory_formatrr)rrr rrrrrrrr)rrrUrrrW)rcrrUrrrW)rrrrW)r rrrW)rz$dict[tuple[str, bool], sympy.Symbol])rrrr7rrr') r rr rrrrSequence[int | Integer] | NonerYrrr)rrrrrr>rrrrrQrrrrrr rXrrrrrrrrrrs@rrrKs N.00  > >&&-@& &&4 > >&8 >  > > [[ ]] ]] ;@ " 37  "HM / @D      $8< HHH H 5 H  H HHrrcp^\rSrSrSrSU4SjjrS SjrS Sjr\"S5S S Sjj5r Sr U=r $) NonOwningLayouti)z,Is a view into the storage of another tensorc>UR5n[TU] URURUR UR 5 Xlgr)rrrr r rr view)rrrUrs rrNonOwningLayout.__init__,s?"  MM LL KK MM   rc>UR5R5$r)rxr1rs rr1NonOwningLayout.make_indexer6s}}++--rcURR5RnUS:XagSSKJn [ R RRX5$)NrTrE) ALIGNMENT) rrrXutilsrrsr r r)rrXrs rmaybe_guard_aligned#NonOwningLayout.maybe_guard_aligned9sB%%'.. Q;$ww<[U[5(aUR5OUn[TU]UR UR URURURURS9 X l X0l g)Nr r rr rXrY) rrrxrrr r rr rXrYrr)rrUrrfixedrs rrCommBufferLayout.__init__`sh &0%G%G!V <<++<<<<oo  !1$r)rr)rUzFlexibleLayout | FixedLayoutrrrr) rrrrrrrrrrs@rrrQs; %$O%,%)% %%rrc\rSrSr%S\S'\R "SS9rS\S'\R "SS9rS\S 'SS jr SS jr SS jr S r g) NoneLayoutisrr cS/$rrrrrrNoneLayout.~srdefault_factoryrrcS/$rrrrrrrs1#rr cgrrrs rrNoneLayout.storage_sizerrcU$rrrs rrxNoneLayout.as_fixed rcUR$rr%rs rrNoneLayout.get_devicer'rrNrrr) rrrrrrrrr rrxrrrrrrrssC !'' DD)D#))+FFIFrrc^\rSrSrS U4Sjjr\S Sj5r\RSSj5rSSjrSSjr SSjr \ SSSjj5r SS jr SS jrS rU=r$)MutationLayoutSHOULDREMOVEic>[TU]UR5UR5UR 5S5 XlUR 5R5n[RRU5 gr) rrr$rrr  get_bufferrzrsr mark_buffer_mutated)rr rrs rr#MutationLayoutSHOULDREMOVE.__init__sc   & & (     OO      ))+ ##D)rc6UR5R$r) real_layoutr rs rr !MutationLayoutSHOULDREMOVE.strides!(((rcgrrr%s rr rs rc>UR5R5$r)rrrs rr'MutationLayoutSHOULDREMOVE.storage_sizes!..00rc^SU4SjjmT"UR5n[U[5(d[U55eU$)Nc>[U[5(aT"UR5$[U[5(aT"UR 55$[U[ 5(aT"UR 5$U$r)rrr rvrw MutableBoxrT)r  unwrap_viewss rr;MutationLayoutSHOULDREMOVE.get_buffer..unwrap_viewssb&"<==#FMM22&(++#F$6$6$899&*--#FKK00Mr)r r rr )r rryr)rrrs @rr%MutationLayoutSHOULDREMOVE.get_buffers9 dkk*&&))74<7) rchUR5Rn[U[5(deU$r)rrUrr)rrUs rr&MutationLayoutSHOULDREMOVE.real_layouts,"))&&)))) rc UR5 [RRUR 55 [ U[ 5(a URnUR5 U(d[RUR5UR5UR5[UR5UR55VVs/sH.upE[RR R#XE5PM0 snnS9n[ U[$[&45(deURnUR5 [)US5(dU5e[ URR*[,5(d$[/URR*55e[1U5URlUR$s snnf)NrrT)rrsr rrzrrrTrLrurrrr-rrr check_equals_and_simplifyrvrrrUrrr)r1srcdst unsafe_aliasrrrs r realize_into'MutationLayoutSHOULDREMOVE.realize_intosV  ##CLLN3 c9 % %((C ##~~'mmo*!$CLLNCLLN C CGG$$>>qD C $DdXz$:;; ;;))C sF##(S(##((//>::QDS!jr"S?S"jr#S2S#jr$S$r%U=r&$)@ryirrr#rUcF>[TU]5 URSS5 gr)rrrr+s rrBuffer.__post_init__s   t4rc>UR5R5$r)rr1rs rr1Buffer.make_indexers --//rcJUR(dU5eUR$rrrs rrzBuffer.get_namesyy$yyyrc[UR[5(aURR5$[ [ UR5R 5er)rrUrr:rrrrs rr:Buffer.get_examples@ dkk6 * *;;**, ,!$t{{"3"<"<==rc>UR5R5$r)rrrs rrBuffer.get_devices##%0022rcgrrrs rrBuffer.get_defining_oprrc6UR5R$r)rr rs rr  Buffer.dtypes &&&rc:/UR5RQ$r)rrrs rrBuffer.get_sizes("''((rc:/UR5RQ$r)rr rs rr:Buffer.get_stride s*"))**rc6UR5R$r)rrXrs r get_offsetBuffer.get_offsets '''rc[UR[5(a UR$[[ UR5R 5er)rrUrrrrrs rrBuffer.get_layouts7 dkk6 * *;; !$t{{"3"<"<==rcUR$rrrs rrBuffer.get_output_specr'rc"UR5$r)rrs rrrBuffer.get_storage_numels~~rc6UR5R$r)rrYrs r get_is_pinnedBuffer.get_is_pinneds ***rc[UR[5(a@[UR[5(d URR 5Ulgggr)rrUrrrxrs rrQBuffer.freeze_layoutsF dkk6 * *: KK4 4 ++..0DK4 *rc[UR[5(d[UR55eURR XS9UlgNr)rrUrrrrTs rrV&Buffer.freeze_layout_with_stride_order%sB$++~66IT[[8II6kk11%1U rc[UR[5(d[UR55eURR U5Ulgr)rrUrrrrZs rr[$Buffer.freeze_layout_with_fill_order+s=$++~66IT[[8II6kk//6 rc[UR[5(d[UR55eURR U5Ulgr)rrUrrrr^s rr_$Buffer.freeze_layout_with_same_order/s=$++~66IT[[8II6kk//7 rc[UR[5(d[UR55eURR XS9Ulgr8)rrUrrrrbs rrd'Buffer.freeze_layout_with_exact_strides3sG$++~66IT[[8II6kk22 3  rc[RRR[R "UR 5S55$rrrs rrBuffer.is_zero_elements;rrcz^TR5(a[[TR5S9$SU4SjjnU$)Nrc>TR5n[R"TR=(d SU"U55$r)r1rqrrrrrs rr"Buffer.make_loader..loaderCs/'')G88DII2GENC Crr)rr rsrrs` rr-Buffer.make_loader>s3  " "=0@A A D rc"UR5$rrzrs rrBuffer.codegen_referenceIrrcgrrrs rrBuffer.decide_layoutLrJrc[UR[5(a%URRR 5/$gr)rrUrrrzrs rr#Buffer.get_inputs_that_alias_outputOs2 dkk? 3 3KK$$--/0 0rc[UR[5(a%URRR 5/$gr)rrUrr rzrs rrBuffer.get_mutation_namesTs3 dkk#= > >KK&&//12 2rc6[UR5/5$r)r>rzrs rrBuffer.get_read_namesYs4==?+,,rc[5$rr=rus rrvBuffer.get_free_symbol_uses\ |rc[5$rr=rs rrBuffer.get_unbacked_symbol_defsbrrcgrrrs rrBuffer.realizeerJrcgr'rrs rshould_allocateBuffer.should_allocatehsrrrrr)rztorch.Tensor | torch.SymIntrrrr)rrrrrrrrrr)r rrr)rcrrUrrrrrrrrrrr)'rrrrrrr1rzr:rrrr rr:r*rrrrr3rQrVr[r_rdrr-rrrrrr^rvrrrZrrrs@rryrys$   50> 3'')+(>  +1;@V"V37V V 78 CH * ;?  U    -H%$)! !&  rrycJ\rSrSrSSjrSSjr\RrS SjrSr g) OperationBufferimcU/$rrrs rrOperationBuffer.get_outputsps v rcU$rrrs rrOperationBuffer.get_defining_opsrrcX[RU5 [RU5 gr)ryrrrs rrOperationBuffer.__post_init__ysT"%rrNrrrr) rrrrrrrrrrrrrr]r]ms#55&rr]c\rSrSrSSjrSrg)ri~cgr6rrs rroInputBuffer.num_readsrrrNr)rrrrrorrrrrr~srrc\rSrSrSrSrg) DonatedBufferiaA 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)rrrrrrrrrririsrric8\rSrSr%SrS\S'SSjrS SjrSrg) riNrrc^SU4SjjnU$)Nc>TR5R5n[R"[R R TR5TR5U"U55$r) rr1rqrrsr  constant_namerzrrDs rr*ConstantBuffer.make_loader..loadersPoo'446G88%%dmmot7K7KL rrrrs` rr-ConstantBuffer.make_loaders  rc[[RRUR 5U5UR S9$NrrU)rrsr rmrzrUr#s rr!ConstantBuffer.constant_to_devices/&&t}}?   rrrr) rrrrrrr-rrrrrrrs+/O(/ rrcl\rSrSrS Sjr\"S5S S Sjj5rS S SjjrSSjrSSjr Sr g)NoneAsConstantBufferic[5$rr=rs rrNoneAsConstantBuffer.get_readsrrc[5$rr=rus rrv)NoneAsConstantBuffer.get_free_symbol_usesrTrNcJ[RRR$r)rsr rnone_strrs rr&NoneAsConstantBuffer.codegen_referencesww##,,,rc[SS9$Nr%)rrs rr$NoneAsConstantBuffer.get_output_specs &&rcgr'rrs rr&NoneAsConstantBuffer.has_tensor_outputr*rrrrrrrrr) rrrrrr^rvrrrrrrrrurusC23$)! !4 -'rrucd\rSrSr%S\S'\"S5S S Sjj5rS S SjjrS SjrSr g)rirrTc.[URU5$r)r'rTrus rrv*ShapeAsConstantBuffer.get_free_symbol_usess  =99rNch[RRRUR5$r)rsr rcodegen_sizevarrTrs rr'ShapeAsConstantBuffer.codegen_references!ww##33DII>>rcgr'rrs rr'ShapeAsConstantBuffer.has_tensor_outputr*rrrrrrr) rrrrrr^rvrrrrrrrrs< J34$):!: !:5: ?rrc^\rSrSr%SrS\S'SrS\S'SrS \S 'SrS \S 'Sr S \S'Sr S \S'\ RS'Sj5r \\ RS'Sj55rS(SjrS)SjrS*SjrS+SjrS,Sjr\"S5S-S.Sjj5rS/U4SjjrS0SjrS1SjrS2SjrS3Sjr\S4Sj5rS5S6Sjjr\S7S8Sjj5rS9S jr S9S!jr!S(S"jr"S0S#jr#S0S$jr$S:S%jr%S&r&U=r'$);rizZ Represents a buffer that is computed during kernel execution rather than being an input. rrTFzClassVar[bool]_force_realizeNrrxCallable[..., Any] | Noneryrrzr{c ## URceURceURceURce[ UR [ 5(d[UR 55eUR nURn[ URURURURURURURURS9nX0l[URURUR5UlURR!U5 Sv XlX lg!XlX lf=f7f)Nr)rxryrzr{rrTrrrUr r rrrrWget_default_sizes_body clear_cache)rold_datar^new_datas rwith_original_inner_fn%ComputedBuffer.with_original_inner_fns5+++&&222$$000..:::$))Y//EDO3DE/99[[  % nn00,,!%!@!@'66",,'66 H!I&%%DK  ' ' 3 3D 9  I$K!I$KsBE%B5E E%E""E%c## [RnS[lSv U[lg!U[lf=f7fNT)rr) old_values r force_realizeComputedBuffer.force_realizes2#11  6,0N ) ,5N )IN )s?/ ? <?cURb UR$[URS5(aURR$g)z} Returns self.name if it exists, otherwise returns the name of the data node if that exists. If neither exist, returns None. Nr)rrrTrs rget_computed_buffer_name'ComputedBuffer.get_computed_buffer_names: 99 99  499f % %99>> !rc6URR5$rrTrors rroComputedBuffer.num_reads rrc6URR5$rrTrrs rrComputedBuffer.get_readsrrc6URR5$rrrs rrComputedBuffer.get_read_namesrrct[UR[[[[ 45(d.[ R"[5[5[5S9$[R"[SS5 URR5(aT[UR5URR5URR!55sSSS5 $[UR5URR#55sSSS5 $!,(df  g=f)Nrkwrites index_exprsr>T)rrTrrrgrurG ReadWritesr>rrrryrPget_store_functionr-r|rrs rrgComputedBuffer.get_read_writess$))itY%GHH** l!|&L  \\.*:D Ayy++--*++-II002II002B A+++-II&&(B A As:A*D).1D)) D7cURRU5URRU5-nUR5(a!X R 5RU5-nU$r)rUrvrThas_store_functionrg)rrprs rrv#ComputedBuffer.get_free_symbol_uses+sd$11  II * *= 9:  " " $ $ **,AA-P PF rc>UR5(dgUR[RR;a?UR 5S:Xa+UR (dURR5$[TU]!5$r) ryrrsr mutated_buffersrorrTr-rr+s rr-ComputedBuffer.make_loaderEsc'')) !8!88 A%''99((* *w"$$rcV[UR[[[[ 45$r)rrTrrrgrurs rr!ComputedBuffer.has_store_functionPs$))itY%GHHrcUR5R5R5n[UR[ [ [45(a+[URRURU5$[UR[5(d[UR55e[URRURU5$r)rrxr1rrTrrrgr rrrurr)rrs rr!ComputedBuffer.get_store_functionSs//#,,.;;= dii)T4!8 9 949944diiI Idii33 DT$))_ D349911499gF Frc [UR[5(Ga[R"UR R 5UR R55uupnUR5Rn[SU55(deUVVs/sHkn[U[R5(dM$[URUVs0sH%ofS:wdM U[RR _M' sn5PMm nnnU(a[UR ["[$45(aUR R'X5nOUnUVs/sH,n[(R*R,R/X5PM. n nSSKJn U "XR555$gs snfs snnfs snf)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# UH/n[U[R[R45v M1 g7fr)rrGStarDep MemoryDeprs rr0ComputedBuffer.get_fill_order..js2A1|33\5K5KLMMs79rrEpick_loop_orderN)rrUrrGr(rTr-r|rgrkrrrorrr7r8rrgrrsr r r* schedulerrr) r index_varsrrrkr3vrrTstride_lengthsrs rrComputedBuffer.get_fill_order[sx dkk> 2 2.:.M.M ,,. 0L0L0N/ + (Z!((*00E  Aa!7!78Y 177n$WnUVPV_Q _n$WX  dii$66"ii// KG(GMR"MRTAGG$$11$@U"7&~}}GG#%X"s*%#F? F? F:-F: F?(3G:F?c[UR[5(a:UR5nU(aUR U5 gUR 5 ggr)rrUrrr[rQrZs rrComputedBuffer.decide_layoutsC dkk> 2 2'')E2259""$ 3rcp[R"UR5UR5SS9up[R "[ SUR55 [UR5UR5(aUOUSSU/UQ76nSSS5 /n/n/n/nUR5HfupXS;a-U(aeURU5 URU 5 M:XS;deURU5 URU 5 Mh Xg4WXE44$!,(df  N=f)NqrrrEr) rGr(r-r|rrrrrSrryitemsr) rr var_rangesrr reduce_vars index_size reduce_sizerrs rr%ComputedBuffer.get_default_sizes_bodys-(::  # # %t'>'>'@ \\.*;T__=N O'')0022Ra DP !#   $$&DAG|&&!!!$!!!$G|#|""1%""1%'($0III)P Os 7D'' D5cX^^^TR5uup4nupgU(aU"X44XVU45uup4nupg/URR5QmUb[U[5(a[ U5S:XdeUup[U[ 5(d[U55e[U [5(d[U 55e[SU 55(deURn X:Xd U U45eU V s/sH oT;dM U PM n n TU - m/UR5Qm[RRT[R 5(dTR#UR%55 SUUU4Sjjn Xg-n ['[)T55(+=(d [*R,(+nU "UU UU5unnnU "X}XN5unnn[.R0"UUSS9uunnn[3UU"U5U"U5/UUU5nUU4U4$s sn f)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. rc3B# UHn[U[5v M g7fr)rr)rfs rr6ComputedBuffer.simplify_and_reorder..sH4Gqz!T**4Grc >TRXUT 5upEnTR5S:Xae[U5S:XaV[[ [U555nU"U5SS:wa*UVs/sHoUPM nn[ U5n[ U5nU"U5nU(aD[RRRUU[T X55upIn [Xi5n OUn XKU4$s snf)Nrrr) _apply_loop_reorderingryrrrrrrsr r _simplify_loopsrKr)x_vars support_varsrsimplify_loopsnewsizesreindex0rrrr_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. .! 9s+C*pr) rSequence[sympy.Symbol]rrrrrrrdtuple[list[int], Callable[[Sequence[int]], Sequence[int]], Callable[[Sequence[int]], Sequence[int]]])rindexing_exprsrrrrrrrrrget_write_exprsrsr rrHPREFER_STORE_LOOP_ORDERextendget_read_exprsrjr$rFloop_ordering_after_fusionrGindex_vars_no_squeezerS)rextra_indexing_constraintsrecompute_sizes_body_funcrrrrrextra_indexing_rangesextra_indexing_exprexpected_var_rangesrrrshould_merge_loops iter_ranges iter_reindexr reduce_rangesreduce_reindex iter_varsrrrs` @@rr#ComputedBuffer.simplify_and_reorders4  ' ' )  %Z  %Z % *)4k1J  )) 94..5578 % 15u==23q8 9:T 6 !3T:: WDAV2I. # 1 1N0--/0 ww""4)O)OPP    3 3 5 64 /*4 /04 /!4 /! 4 /  4 /4 /l"/ t,- - VV5V5V1V (<     ( $ \1,@ {, ( ~q 0[T U]SUS9 [RU5UlX0l[ RRU5Ul [ RRU5 0Ul 0Ul 0Ul U(a [U5O0UlX@l/UlUbmURSn[%U[&5(d[)U55eUR+5nUV s/sHn [-[/US9X5PM sn UlU=(d [15UlSUlSUlgs sn f)Nrrrr%)rrrunwrap_storagerqmake_kernel_renderrsr register_bufferrregister_operationrepilogue_fusable_outputs_multi_output_childrenr _named_inputsmutated_inputsmutation_outputsrrrrMutationOutputrr>allowed_prologue_inpsallow_epilogue_fusionallow_prologue_fusion) rrUrqr r r named_inputs first_inputr rrs rrTemplateBuffer.__init__s1 d62"11&9 "4GG++D1  ""4(+-9;%?A# #/D B  -68  %++a.Kk622 ED4E E2 ++-F*%)Cz8#D)%D ! " 1Z\ " 37"26"%s9Ec[UR[5(a [S5eUR 5R $)Nz1Multi-output templates do not have a single dtype)rrUMultiOutputLayoutrrr rs rr TemplateBuffer.dtypes9 dkk#4 5 5%C  &&&rc URSS9$)NT normalize)rPrs rrgTemplateBuffer.get_read_writess''$'77rc ^^[5nURHn[U[[45(d[ U55eUm[TR [5(d[ TR 55eTR R5mSUU4SjjnU[R"UTR5SUS9R-nM U$)z(Build read dependencies from all inputs.c|>[U5S:Xde[R"TR5T"U55$r)rrqrrz)rrrs inp_indexers rdummy4TemplateBuffer._read_deps_from_inputs..dummys16{a'''xx  E0BCCrrr' r Sequence[Any]rr0 rr ) r>rqrrZryrrUrr1rGrPrrk)rr( rkinp_rawr- rsr, s @@r_read_deps_from_inputs%TemplateBuffer._read_deps_from_inputss.8l{{Gg'@AA P4= PA,3Ccjj&11 C4 3C C1**113K D D \55s||~rYe E# rc *^^[UR[5(ar[[R "UR 5[R"S5SSS9/5n[R"URU5U[5SSS9$UR 5mUR5R5mSUU4Sjjn[R"X0R5SUS9nU=RURU5-slU$) aExtract read/write dependencies for this TemplateBuffer. When the layout is MultiOutputLayout (multi-output templates), the buffer itself has no data layout, so we cannot build an indexer. Instead, synthesize a trivial write dep and derive read deps from the named tensor inputs (``_named_inputs``). For single-output templates with a concrete layout, fall through to the standard path. rr) var_namesrN)rkrrr rcb>[U5S:Xde[R"TT"U5S5$Nrfakerrqrrrrrs rr- 1TemplateBuffer.extract_read_writes..dummy,v;!# ##99T75>6: :rr' r/ )rrUr$ r>rGrrzrr rr2 rr1rPrrk)rr( rr- depsrrs @@rrP"TemplateBuffer.extract_read_writess dkk#4 5 53= ** q)9Rb4F **11)<&L  }}//#002 ; ;// ==?B)  d11)<<  rc6[RR$r)rr7r]rs rr|!TemplateBuffer.get_reduction_sizesww{{rcgrrrs rry!TemplateBuffer.get_reduction_typerrcgrrrs rrZTemplateBuffer.should_allocate rrc*UR5/4S4$rr )rrrs rr#TemplateBuffer.simplify_and_reorder s$      rc6[UR[5$)zFWhether this template produces multiple outputs via MultiOutputLayout.)rrUr$ rs ris_multi_outputs_template(TemplateBuffer.is_multi_outputs_templates$++'899rcUR$r)r rs rget_allowed_prologue_inps(TemplateBuffer.get_allowed_prologue_inpss)))rcg)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)r hook_outputss r_finalize_codegen TemplateBuffer._finalize_codegen!srcP[U[5(a+[UR[5(a UR$[R U5n[U[ 5(a URn[UR[5(aUR5 U$)zWRealize a TensorBox, preserving MultiOutput layout (unlike ExternKernel.realize_input).) rrrT MultiOutputr realize_inputrxrUrrQ)r1r2rs rrealize_template_input%TemplateBuffer.realize_template_input/sr b) $ $BGG[)I)I77N++B/ fj ) )[[F fmm^ 4 4  " r)direct_alias_at_leafon_tensor_leafon_non_tensor_leafc~^^^^^^^0m[R"5mSUUUUUUU4Sjjm[T"U/55$)zLWalk a structured output tree, creating MultiOutput nodes for tensor leaves.c v>[U[[45(a?/n[U5H,up4UR T"U/UQ[ U5U4P55 M. U$[ T 5n[U[R5(aT (aUT ;a[RT U5/$[U5nUT ;aT U/$[[RU5TU5nUTRUR!5'T bT "UR!5XqU5 [U5nUT U'U/$T bT "U5 /$r)rrrrrrrlrTensorrridrR FallbackKerneltensor_to_layoutr rz)rcrrritemleaf_idxtidmor2rV leaf_counterrX rW seen_outputs template_bufwalks rrf 0TemplateBuffer.build_multi_outputs..walkJs0&4-00+-(0GANN4.K.K$v,9J.K#LM 1L)H&%,,//'H8L,L%,,-A(-KLMMj,&(-.. "33F;\7FH 33BKKMB!-"2;;="xHr]$& S!t !-"8,Ir)rcrrzlist[tuple[type, int]]rlist[TensorBox])rFcountr) r1re structuredrV rW rX rc rd rf s ` ```@@@rbuild_multi_outputs"TemplateBuffer.build_multi_outputs;s4.0  (   6T*b)**r) r r r r r rr rqr r r r)NNN)rUr#rqSequence[IRNode]r rr Iterable[IRNode] | Noner OrderedSet[str] | Noner zdict[str, IRNode] | Nonerrrr)r( rrzOrderedSet[dependencies.Dep]r)r( rrrrrrr )rr rrrz9tuple[tuple[Sequence[Expr], list[Expr]], LoopBody | None]r)rN zdict[str, str]rzFinalizeCodegenResult | None)r2rrr) re rrj rrV zdict[int, IRNode] | NonerW zFCallable[[str, MultiOutput, list[tuple[type, int]], int], None] | NonerX zCallable[[int], None] | Nonerztuple[TensorBox, ...])rrrrrrrr rgr2 rPr|ryrZrrH rK rO rlrT rk rrrs@rrrsM378<150707!076 07 0 07 6 07/07 0707d''8$$L OS?C  $K  $=   C  :* *  %    :>;?)+$)+)+ 7 )+  )+9)+ )+)+rrc^\rSrSrSSU4Sjjjr\"S5S S U4Sjjj5rS SjrS SjrSr U=r $) TritonTemplateBufferihc>[TU]UUUUUS9 URceURUR0UlSUlSUlg)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 N)rrrr subgraph_inps subgraph_outs)rrUrqr r r rs rrTritonTemplateBuffer.__init__is]"    )"7  yy$$$)-DII(>%FJ9=rc,>[TU]U5nUR(a URO/nUR(a URO/nUHzn[ U[ R 5(aUR[XQ55 M>[ U[5(a"URURU55 MuUcMze UH?n[ U[5(a"URURU55 M:UcM?e U$r) rrvrt rs rrrupdater'r)rrpresrt rs rsrrs rrv)TritonTemplateBuffer.get_free_symbol_usessg*=9.2.@.@**b .2.@.@**b  C#uzz** +C?@C(( 333MBC{"{ !!C#v&& 333MBC{"{ !  rc U/URQ$rr rs rr TritonTemplateBuffer.get_outputs-t,,--rc&SURS3nU$)NzTritonTemplateBuffer(layout=r"r)rrs rr!TritonTemplateBuffer.__str__s,T[[M; r)r rs rt r ) rUrrqrm r zCallable[_P, _T] | Noner rn r ro rrrrrr) rrrrrr^rvrr!rrrs@rrq rq hs 378< >>!>4 > 0 > 6 > >><23$)! !4..rrq c^\rSrSrSrSU4SjjrSSjrSSjrSSjrSSjr SSjr SS jr SS jr SS jr SS jrS rU=r$) 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>[TU]5 XlX0lX lX@lSUlSUl0UlSUl SUl 0Ul 0Ul gr') rrrrUr descriptionfailed_benchmark_with_cudagraphsrrk decompositiondecomposition_kwargsconfig_patches)rrrrUr rs rrChoiceCaller.__init__s`   &'! 05',./38<46!.0rc^^UR5mUR(a[R"UU4Sj5$[R (a[ UU4Sj5$[R"TTSU0SS9$)Nc>T"T6$rralgorsrr(ChoiceCaller.benchmark..s T4[rc>T"T6$rrr srrr s D$Krrr%) to_callabler rXbenchmark_gpu_with_cuda_graphrF/profile_bandwidth_with_do_bench_using_profilingrc benchmark)rrrr s `@rr ChoiceCaller.benchmarks\!  * *<<=PQ Q  A A+,?@ @$$T4%dKKrc[errrs r call_nameChoiceCaller.call_namerrc[errrs rr ChoiceCaller.to_callablerrc"UR5$)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 }}rc[errrs rr ChoiceCaller.hash_keyrrc[errrs rrlChoiceCaller.output_noderrc0$)zRInformation returned here is logged to the autotune log file when that is enabled.rrs r info_dictChoiceCaller.info_dictrrcg)Nunsupported_choicerrs rautoheuristic_idChoiceCaller.autoheuristic_ids#rcSUlg)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 rkrrUr) rrrrrUrr rrr)rr rrrrrr)rr )rr)rz6dict[str, PrimitiveInfoType | list[PrimitiveInfoType]]r)rrrrrrr r r r r rlr r r rrrs@rr r sp11"1 1  1  16L""""$rr c\rSrSrSSjrSrg)TritonTemplateCallerBaseic[errrs rget_make_kernel_render/TritonTemplateCallerBase.get_make_kernel_renderrrrN)rr )rrrrr rrrrr r s"rr c^\rSrSrSrS U4Sjjr\S Sj5r\SSj5rSSSjjr \ RSSj5r SSjr SSS jjrSS jrS rU=r$)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>[TU]UUSUS9 X0l0UlX@lX l[ SU55Ul0Ulg)N)rUrqr r c3# UH]n[U[5=(dA [U[RRR 5=(a UR v M_ g7fr)rr rrselect_algorithmExternKernelCallerhas_out_variant)rchoices rr/MultiTemplateBuffer.__init__..sU% - v7 8 65??#C#C#V#VW+**  -sA%A') rr_choice_timings_fn_choice_timings_choicesoriginal_inputsr_output_plannable_make_kernel_renders)rrUrqchoice_timings_fnunfiltered_choicesr rs rrMultiTemplateBuffer.__init__sd #"7  #4LN,> %!$% - % " <>!rcUR$)zN Are all possible choices TritonTemplates or Extern Kernels with out variants )r rs routput_plannable$MultiTemplateBuffer.output_plannable#s %%%rcUR$r)r rs rr#MultiTemplateBuffer.choices*s }}rczXR;aURU5URU'URU$r)r r )r hint_overrides rchoice_timings"MultiTemplateBuffer.choice_timings.s<  4 4 4262I2I-2XD  /##M22rc#8# [U[RRR5(d[ U55eUR UR :XdeURnUR5UlSv X lg!X lf=f7fr) rrrr TritonTemplateCallerrrUr r )rcallerrenders rswap_as_triton_caller)MultiTemplateBuffer.swap_as_triton_caller5s EOO44II    <  {{fmm+++(("("?"?"A - &, #f #sBBBBBBc\[U[RRR5(d[ U55eUR 5URR:XdeUR5URR:XdeUR5Ul gr) rrrr r rrrUrr:r r r )rr s rfinalize_as_triton_caller-MultiTemplateBuffer.finalize_as_triton_callerCs EOO44II    <  }}&--"4"4444 FMM$8$8888"("?"?"ArcRURUS9n[X"RS9nX2U4$)N)r r)r rr)rr timings min_choices rget_min_choice"MultiTemplateBuffer.get_min_choiceKs3%%M%Bkk2 J/00rcUR5H"up#UR5URU'M$ URSUlg)z;Finalize with multiple callers for different hint overridesN)rr r r )rcallersr r s rfinalize_as_triton_callers.MultiTemplateBuffer.finalize_as_triton_callersRsE&-]]_ !M7=7T7T7VD % %m 4&5#'";";D"Ar)r r r r r r r ) rUrrqrm r z1Callable[[int | None], dict[ChoiceCaller, float]]r list[ChoiceCaller]r rrrr)rr r)r rrzdict[ChoiceCaller, float])r r rr )r r rr)r rrztuple[ChoiceCaller, float])r z*dict[int | None, TritonTemplateCallerBase]rr)rrrrrrrr r#r rrr r r r rrrs@rr r s>>!>M > / > / > >8&& +/3'3 "3 - -B+/1'1 #1BAB BBrr c\^\rSrSrSU4SjjrSSjrSSjrSrU=r$) CUTLASSTemplateBufferi]cJ>[TU]XU5 X@lXPlX`lgr)rrworkspace_sizetemplatesupports_epilogue_fusion)rrUrqr r r r rs rrCUTLASSTemplateBuffer.__init__^s& );<, (@%rc8URb UR$S$rr rs rr(CUTLASSTemplateBuffer.get_workspace_sizems&*&9&9&Et""L1Lrc~UR5H)n[R"UR5SS5 M+ gr)rrqrrz)rrcs remulate_store_fn&CUTLASSTemplateBuffer.emulate_store_fnps,&&(F IIfoo't 4)r)r r r )rUrrqrm r Callable[_P, _T]r rr rwr rrrrr) rrrrrrr rrrs@rr r ]sd A A! A- A  A " A#' A  AM55rr cT^\rSrSrSU4SjjrSU4SjjrSrU=r$)CppTemplateBufferiucL>[TU]XU5 X@lXPlSUlgr)rrr r outputs)rrUrqr r r rs rrCppTemplateBuffer.__init__vs& );<  ,0 rc>[UR[5(a[UR[5(d[ UR55eURSn[U[ 5(d[ U55eURn[U[5(d[ U55eU$[TU]%5$r) rrUr$ r rrryrrr)r first_outputrUrs rrCppTemplateBuffer.get_layouts dkk#4 5 5dllH55 ItDLL7I I5<<?LlF33 GT,5G G3!((Fff-- ;tF| ;-M7%' 'r)r r r ) rUrrqrm r r r rwr r rrr)rrrrrrrrrs@rr r usL 1 1! 1- 1 " 1  1  1 ( (rr cV^\rSrSrSrSSU4SjjjrSSjrSrU=r$) CuteDSLTemplateBufferiz Buffer for CuteDSL (CUTLASS Python DSL) template kernels. Similar to other template buffers but specialized for CuteDSL operations. c >[TU]XU5 X@lXPlU/UlUb[ UR S[5(d[UR S55eUR SR5nU=RUVs/sHn[[US9Xp5PM sn- slggs snf)Nrr%) rrr r r rrqrrrr r) rrUrqr r r r rrs rrCuteDSLTemplateBuffer.__init__s );< ,&*V  %dkk!nf55 KtDKKN7K K5[[^..0F LL))Cz8#D) L &sB>cUR$rr rs rr!CuteDSLTemplateBuffer.get_outputs ||r)r r r r) rUrrqrm r r r r r rn rrr) rrrrrrrrrrs@rr r s[37 !-    0  *rr c^\rSrSrSrSS U4SjjjrS SjrS SjrS S SjjrSr U=r $)NVUniversalGemmBufferiz Buffer for NVIDIA Universal GEMM kernels. Unlike CuteDSL templates which use Jinja templates, this generates simpler Python code that directly calls the cutlass_api library. c >[T U]XSS9 X0lX@lU/UlX`lXPlXplXlXl Xl URRURRS.UlURUlg)N)r ) kernel_namemin_cc)rrkernelaccumulator_typer r variant scale_type_a scale_type_bswizzle_type_aswizzle_type_bmetadatar r kernel_metadata_make_kernel_renderr ) rrUrqr r r r r r r r rs rrNVUniversalGemmBuffer.__init__s DA 0&*V , ((,,"??66oo,,  #'":":rcUR$)z#Return the workspace size in bytes.r rs rr(NVUniversalGemmBuffer.get_workspace_sizes"""rcUR$rr rs rr!NVUniversalGemmBuffer.get_outputsr rc^ SSKJn SSKJn /nURHVn[ U[ 5(a URn[ U[5(a URnURU5 MX [UR5nU"UUUURURURURUR UR"UR$UR&S9 m U 4SjnT U4$)z Create a kernel renderer for code generation. Returns (kernel, render) tuple where: - kernel: NVUniversalGemmKernel object with call_kernel() method - render: function that returns source code string r)NVUniversalGemmKernel) Placeholder) r rrlr r r r r r r r c$>TR5$r)r ) render_kernelsrr 9NVUniversalGemmBuffer._make_kernel_render..renders '') )r)Btorch._inductor.codegen.nv_universal_gemm.nv_universal_gemm_kernelr torch._inductor.utilsr rqrrrTrxrr KERNEL_NAMEr r r r r r r r ) rout_noder r r rrsr r r s @rr )NVUniversalGemmBuffer._make_kernel_renders  6!# ;;C#y))hh#z**hh   s # +112 -##  00!22..LL****....   *f$$r) r r r r r r r r r r r )rNNNN)rUrrqrm r r r r r r r rr Any | Noner r r r r r rrrrr)r r r rrztuple[Any, Any]) rrrrrrrrr rrrs@rr r s #'#'%)%);;!; ;  ;  ;;!;!;#;#; ;;B#:>*%*%,6*% *%*%rr c&[SU55$)Nc3B# UHn[U[5v M g7frrr)rr9s rr#is_node_sequence..s4ez!V$$er)r)rs ris_node_sequencer! s 4e4 44rc\rSrSr%S\S'SSjrSSjrSSjr\SSj5r \ SSj5r SS jr SS jr \"S5SSS jj5rS rg )ri#Sequence[IRNode | Sequence[IRNode]]rqcnURUn[U[5(deUR5$rrqrrrz)rrinputs r input_nameInputsKernel.input_names/ A%((((~~rc^[[R"5n[RmURHqn[ U[ 5(aURU4SjU55 M5[ U[5(aMLURT"UR555 Ms [[R"U4SjUR555n[R"UU[5S9$)Nc3P># UHnT"UR55v M g7frrH)rrrs rr/InputsKernel.get_read_writes.. sBEqWQZZ\22E#&c3P># UHnT"UR55v M g7frrH)rrrs rrr+ 's#. /AGCLLN # #/Ar, r) r>rGrMrrqrrrw rrrzrr)rrkr& rrs @rrgInputsKernel.get_read_writess<++,.&&[[E%** BEBBE#899 '%.."234!L,,-. /3/?/?/A.  &&"   rc6UR5R$rrjrs rrInputsKernel.get_reads1rmrc[U[5(a URn[U[5(a URn[U[5(a*[U[ 5(d[ RU5n[U[5(aURU5$[U[5(aU$[U[[ 45(d[U55eU$r) rrrTrxrvrZrrS unwrap_storage_for_inputTorchBindObjectryrr1rs rr2 %InputsKernel.unwrap_storage_for_input4s a # #A a $ $A a " ":a+I+I**1-A a # # //2 2 a ) )H!fo677@a@7rc/nUHdn[U[5(a&UVs/sHn[RU5PM nnO[RU5nUR U5 Mf U$s snfr)rrrr2 r)rq inputs_newrrs rr InputsKernel.unwrap_storageGsm79 A!X&&GHIq!\::1=qI 99!<   a   JsA/cgrrrs rrInputsKernel.is_externTrrcgr6rrs rroInputsKernel.num_readsWrrc[[R"5nURHIn[ U[ 5(aX#R U5-nM-UHnX$R U5-nM MK U$r)r>rr!rqrrrv)rrpr3rs inner_inps rrv!InputsKernel.get_free_symbol_usesZsg u|| $ &;;C#v&&--m<<!$I77 FFA"%  rrN)rrrrrrrO)rqr# rzlist[IRNode | Sequence[IRNode]]rrrr)rrrrrr' rgrrlr2 rr rror^rvrrrrrrs //  ,,$ 3  (  N+$) !  ! , rrc(\rSrSrSSjrSSjrSrg) NopKernelihcgrrrs rrNopKernel.is_no_opirrc[5$rr=rs rrNopKernel.get_readslrrrNrr)rrrrrrrrrrrA rA hs rrA c\rSrSrSr\S Sj5r\S S Sjj5r\"S5S SSjj5r \SSj5r SSjr S r g) ConcatKernelipzb There isn't actually a real kernel for concat, we just change the storage for the upstream data. c  USR5nUSR5n[USR55nS/nXR/nSUs=::a[ U5:de e[ S[ U55HnXR5n UR XR5 [ U 5[ U5:XdeXR5U:XdeXR5U:Xde[ [ U55HEn X:Xa XZX-XZ'M[RRRXZX5XZ'MG UR XR5 M [RU5n [R(a#[R!XUSR"5n [ [ U55H|nXn [%U 5(dMU R'5n [)U [*5(dM@[R-U R.U R05(dMq[3U5n O [5SU55n[RR6R8SnUSLa7[)U[5(a"[5SU55(a [3U5n [;SU55nUce[=S[+UUUU US9/S 9n[?U5n/n[AU5GHunn[)U[B[D45(d[GU55eURIU[JRMUX&UXxSS 95n[)U[N5(d[GU55e[)URP[5(d[GURP55eURPR U5 [)URR[B5(aURRRU5nO URRn[)U[>5(dGM4URW5(dGMLUR5=ncGMb[YURF5(dGM[[U5(aGMUR UR]55 GM [ U5S:aR[RR_U[`Rb5(a[RReU5 [RRgU5Ul4URkURP5Ul([RRmU5 U$) z& Create the concat kernel from inputs rrEc38# UHn[U5v M g7fr)rrs rr&ConcatKernel.create..s-WPV1.CA.F.FPVrFc3# UHvnSUR;=(a[ URSR[RS9=(d* URSR[RS9v Mx g7f)rYrN)rirrrrrargs rrrJ sx (C!HHUO11@S@S1Txx44&+&<&<5 (sA>Bc3z# UH1n[U5=(a UR5Rv M3 g7fr)rrrYrs rrrJ s- KQa !! $ A)A)A A6s9;N)r r rr rYrrUrqr)7rrrrrrrrsr r r rrrFrurrpr rrrrWrCrr r/r current_noderrrG rxrrvrrrrrryrqrTrwr>rjrirrrHFOREACHregister_operation_listr rr r )r1rqrHr r r offsets_start offsets_endr input_sizer< output_striderrUany_input_is_storage_and_layout fx_node_argsrY concat_kernelr op_namesrsrinput_unwrappeddevs rrConcatKernel.createvs %%'q ##%q **,- }o C'#h-'''''q#f+&A++-J   /z?c(m3 339&&(E1 119'')V3 333x=)8"*+ "=HK"#''"2"2"L"L Z]#HK *   x} -'(6'H'H'R  ' '"//M s6{#A A$Q''K88fmmTT$B8$LM$+.-WPV-W*W'ww++003  ,u 4<.. (   ;8DM KQ  !!!$$#   M*'FAscHj#9:: EDI E:++  Cq!1;>!L lF33 GT,5G G3m22D99 U4 @T@T;U U9  ' ' 5#((H--"%(("6"6"8"%((?J77#3355NN,,S9388$$"<00 ? ? AB1(4 x=1 !4!4V^=S=S!T!T GG + +H 5WW44]C "11-2F2FG  ""=1 rNc2[U[5(aURURU5$[U[[ 45(d[ U55e[UR[5(a[URR[5(aURR(dgUcg[UR55[UR55:wag[S[UR5UR5555$[URS5=(aJ [URR[ 5=(a [UR["5(+$)NFTc3x# UH0up[RRRX5v M2 g7frrNrPs rr=ConcatKernel.can_realize_into_without_copy..s1EFB  88@@ErRrU)rrcan_realize_into_without_copyrTrvrxrr rUrWr rr:rrrrExternKernelAlloc)r1r rs rrc *ConcatKernel.can_realize_into_without_copys! c9 % %44SXXsC C#*566AS A6 chh 3 4 4sxx <<xx00{3>>#$CNN,<(==!#.."2CNN4DE  CHHh ' <388??N; <sxx):;; rc,[RX5$r)rA rvrus rrv!ConcatKernel.get_free_symbol_usess--dBBrc [U[5(d&[U5(a[U5up4[X4S9n[U[5(d[ U55e[U[ 5(aUR URU5$[U[5(aiUR5 [URS5(deURX5(a&[U5URl UR$[RUR!5UR#5UR%5['UR)5UR)55VVs/sH.upV[*R,R.R1XV5PM0 snnS9nUR Xr5$s snnf)NrSrUr)rrZrrVrrrrTrxrrrc rrUrurrrr-rrrsr r r )r1r rr]rUrrpws rrConcatKernel.realize_intosZ #//$S))"7"<%7B#//:c:/ c9 % %##CHHc2 2 c: & & KKM388X.. ..00::"1#"6xx   >>#--/__&   ??DA  ::1@?  (( s75Gcgrrrs rrZConcatKernel.should_allocate:rrr)rqrm rHrrrxr)r rrrrrrr)r rrrrrr) rrrrrrlrrc r^rvrrZrrrrrG rG ps uun/3! ! ,! ! ! FN+$)C!C !C,C ))BrrG c$^\rSrSr%SrSrS\S'\R"\ S9r S\S'S r S \S 'S r S \S 'S r S \S'\R"\S9rS\S'S rS\S'S rS\S'\R"\ S9rS\S'S rS\S'\R"\ S9rS\S'\R"\S9rS\S'SCSDU4SjjjrSESjrSFSjrSGS jrSGS!jrSHSIS"jjrSJS#jrSHSKS$jjrSLS%jrSMS&jrSNS'jr \!SOS(j5r"\#SPS)j5r$\#SQS*j5r%\#SRS+j5r&\#SRS,j5r'\#SSSTS-jj5r(\#SUSVS.jj5r)\#SUSWS/jj5r*\#SRS0j5r+\#SRS1j5r,\#SRS2j5r-\#SRS3j5r.SGS4jr/SXS5jr0SHSYS6jjr1SZS7jr2S[S8jr3SUS\S9jjr4SNS:jr5SJS;jr6SJS<jr7SJS=jr8S]S>jr9S^S?jr:\;"S5SUS_S@jj5rSBr?U=r@$)`ri>z 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. rr0 constant_argsrrrNReinterpretView | None output_viewrpython_kernel_namecpp_kernel_name Iterable[str]ordered_kwargs_for_cpp_kernel_OpOverloads | None op_overloadzlist[dict[str, Any]] | Nonearg_propertieszdict[str, dict[str, Any]]allarg_propertiesz dict[str, dict[str, Any]] | Nonekwarg_propertiesz"dict[sympy.Symbol, pytree.KeyPath]unbacked_bindingszlist[MutationOutput]r c @>[T U]UUUS9 X@lU(aUO0UlX`lXlUR U5 URU5 XlUR5 0Ul /Ul [RRUl0UlgNrP )rrrn rrp rv set_cpp_kernel_nameset_python_kernel_namert collect_arg_kwarg_propertiesrz r rsr rR fx_noder) rrrUrqrn rrp rq rr rt rv rs rrExternKernel.__init__Zs   + &fB &&   1 ##$67-J* ))+!# "ww++ +-rc U/URQ$rr{ rs rrExternKernel.get_outputszr} rc[5$rr=rs rr%ExternKernel.get_unbacked_symbol_defs}rrc[UR[RR5(afURR R Vs/sH:nUR(aMURURURS.PM< snO.[[UR55Vs/sHn0PM snUl[UR[RR5(aSURR R Vs0sH'nURURURS._M) snO0Ul[UR[RR5(aUR (dRURR R Vs/sH!oR(dMURPM# snUlURR R Vs/sHoR(dMUPM snUlg/Ulgs snfs snfs snfs snfs snf)N)rrr)rr)rrv r_ops OpOverload_schema arguments kwarg_onlyr real_typerrrrqrw rx rt schema_kwargs)rrrs rr )ExternKernel.collect_arg_kwarg_propertiess$**EJJ,A,ABB))11;;    >55$($4$4$<$<$F$F6$Fq,,FAFF$F62 ++33=="=a="D "$D ? 7 6"s0I/(I: I#.I III+Ic[UR[5(a!UR5 UR 5 ggr)rrUrapply_constraintrQrs rrExternKernel.decide_layouts0 dkk> 2 2  ! ! #     3rc[X5up4U(aURU5 U(dUR5nU(aSSKJn U"XSS9nUR X&5 gg)NrE)'set_kernel_post_grad_provenance_tracingT)r)rf make_commenttry_get_kernel_namerr write_provenance_debug_handle)rwrapperr origin_str_detailed_origin_strr debug_handles rcodegen_commentExternKernel.codegen_comments],?t+M(    ,224K  FBTL  1 1+ L rc[errrr s rcodegenExternKernel.codegenrrcrXl[RR(a3[ UR [ RR5(dgUR nURcURS:XaURS:XaURRS5SOURRSS5nSSKJn UR!SUR305n[#SU5S S 9nUS :aUS U3nS US 3UlgUR$R&Ulgg)Natenr.rrinductor_fallback_opszaten.c3l# UH*oRS5(dM[USS5v M, g7f)rrEN) startswithr)rrs rr3ExternKernel.set_cpp_kernel_name..s'KAc9JZS12ZZs44rEr_vz at::_ops::z::call)rr rsr  cpp_wrapperrrv rr r namespace _overloadnamerrreplacetorchgen.aoti.fallback_opsr rrMr r)rrr r opnamer version_infolatest_versions rr} ExternKernel.set_cpp_kernel_names.ww""*   ejj33+ +  !!    '6)++y8OO))#.q100c:M4885@Q9RTVW !$KK""A% &xr.)9:F)3F86'B$'-~~':':$7 (rcXlUbgURnUcg[U[RR 5(aSUR 3UlgURRSS5SUR 3Ulg)Nztorch.ops.higher_order.._ops..ops.r ) rq rv rrr HigherOrderOperatorrrr )rrq r s rr~ #ExternKernel.set_python_kernel_names"4  ) !! >   > > ? ?(??P&QD #$$,,Xw?@&//ARS  #rcZSSKJn UR5=n(a URO[R R n[R R(a UR$[R R(a[[R RU5(d([[R R55eURcg[R RRURU5$UR$)NrE) CppWrapperCpu)codegen.cpp_wrapper_cpur rrrsr  device_type fx_wrapperrq r rrrr get_c_shim_func_name)rr dr s rr ExternKernel.try_get_kernel_names:!%!22A29L9L 77  ** * WW agg22MBB D$$E B##+77''<<$$f ** *rc0UR5nUceU$r)r rs rget_kernel_nameExternKernel.get_kernel_names!'') rc [RUR5UR5UR 5UR 5UR 5UR5S9nUR5 U$)N)r r r r rr) rurrrr-rrrr)rri s r copy_inputExternKernel.copy_input sa   <<>++-]]_::<))+oo'    rc ^^X#S.n[R"U5unm/m/n/n/nUGHn U =[S`=bju [RR R RU SS9n TRS5 URU 5 URU 5 Mz =[S`=bu TRS5 URU 5 U RRn U RRS:XaU ceUR[RRU R!55 GM =["S`=bCu TRS5 URU 5 URU R$5 GMl [&S`=b'uTRS5 URU 5 GM TRS5 URU 5 URU 5 GM SUU4Sjjn UV s/sHoR)U 5PM nn UHn [+U 5(dM[-U SS 9 M! /nUGH'n [/U [05(dhU R35[RR4;a<UR[RR4U R355 M[/U [05(dhU R35[RR6;a<UR[RR6U R355 M[/U [85(a"URU R;55 GM5[/U [<5(aURU R>5 GMh[/U [R@RBR5(arU RRn U RRS:XaU ceUR[RRU R!55 GM UR[EU 55 GM* U "X5unnU"U0UD6nSn[RFR =n(a[RHRJRMS 5n[O5n[RHRP[RRRTRVLaUS n[Y[RH5nU [[U[RHU5 SSS5 []UUU5n[/U[^[`45(dU/OUnUHn[/U[Rb5(dM$URd(dM7[fRh(aMNS n[RRHRJRMS S5=n(aUSU3nU[Rl5M UUUU U4$s sn f!,(df  N=f)a]Partition kernel args into tensor and non-tensor, realize tensor inputs, re-run fake tensor propagation with the realized strides, and return (example_output, tensor_args, non_tensor_args, unflatten_args, unbacked_bindings). unflatten_args(new_tensor_args, new_non_tensor_args) reconstructs the original (args, kwargs) tree from replacement lists. )rrrN)r|Fr'Tc6>/n[U5n[U5nTH@nU(aUR[U55 M&UR[U55 MB [R"UT5nUR S/5UR S054$)Nrr)rmrrlpytreetree_unflattenr) new_tensor_argsnew_non_tensor_argsr it_tensorsit_non_tensors is_tensorr3args_flat_is_tensor args_specs runflatten_args3ExternKernel.process_kernel..unflatten_argsUsFo.J!"56N0 MM$z"23MM$~"67 1 %%fi8A55$aeeHb&99 9rrrYrEzEsparsity not handled. Please file issue for sparse inference weights.rz Found from : )r rr rrztuple[list[_T], dict[str, _T]])6r tree_flattenrrsr r rcreate_symintnoderGeneratorStater rrrr'default_generators clone_stateOpaqueObjectStaterrrS rrVrrvrz constantstorchbind_constantsr3 get_valueropaque_example_valuerirrr9rR rirr r _higher_order_opseffects with_effectsr1r8r2rrr[ is_sparserFgraph_partitiondisable_cudagraphs_reason)r1r rr binded_args args_flat tensor_argsnon_tensor_argsreal_non_tensor_argsrN r device_indexr r example_argsnew_args new_kwargsexample_outputrz r node_meta_valctxexample_out_lirmsgrr r s @@rprocess_kernelExternKernel.process_kernels" $6 %22;? 9*,$& (*  CTV77++55GGRVGWD'..u5#**40(//5  &^%'..u5#**3/#&::#3#3L::??f49QQQ(// 55lCOOQ &)&('..u5#**3/(// :) X'..t4&&s+'..u5#**3/(//4;> :) :@L : + : :6AA[((+[ AA$Q''%a5  Aa**qzz|qww?P?P/P##AGG$5$5ajjl$CDq(++JJLAGG$?$??##AGG$?$? $MNA//##AKKM2A011##A$:$:;Au11@@AA xx~~ xx}}.<3KKK##JJ11,?KKM##$5a$89+. .lQ*8Z8GK -- -9 -NN//33E:M0; C~~$$(?(?(G(G(T(TT -a 0 >>J 9>=!  ntUm<<    A1ell++KKK...]"#''"6"6";";"?"? t"TT;T E!2;-@C471         YBjsYY Y&c H[U[5(d[U55e[U[5(aU$UR 5n[ R RUR55nUceUR5nUbSUR;a[U[[[45(a[UR[5(a}[URS[ R"S9(d*[URS[ R$S9(a)UR'[)UR+555 OUR-5 [.R0"UR+5SS9upVUSnUR35"U5n[ R R4R7X5n[ R R4R9X5n [ R R4R;X5n [=Xy5U -n X:wa[>RASU U U5 [Be[URD[GURI5URK5UR+5U U SS9S 9$) 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. rYrL r3rrz@convert_to_reinterpret_view failed: stride=%s offset=%s index=%sFrrS)&rrvrrZrwrsr rrzrriryrrUrr-rrrr_r/rrQrGr(r1r r) stride_vars offset_varrkr'rrrTrWr$r) r1r x_unwrap_viewrx_unwrap_view_fx_node index_argsrr rr\rXexpecteds rconvert_to_reinterpret_view(ExternKernel.convert_to_reinterpret_views"!X&&/Q/& 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 jUc [5$[U[[RR R [45(a [US9$[U[5(am[5 [RR[R"UR UR#5UR%5S95sSSS5 $[U[&5(aU$[U[(5(aUR+UR,5$[U[.5(a1[/UR+UR,5UR15S9$[U[25(a@UR55 [7UR955(aUR;U5$[U[>5(aUR55 U$[U[@[[B45(aU$UREU5$!,(df  GNK=f![<a Nxf=f)N)rT)r r rS)#rurrrrrrrrrr?rsr add_tensor_constantrr[rrrrrrS rTrZrrvrrrwr rrx NonTensorObjrr r4 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*)"+s4AH(H% H"% H21H2c[U5(a@[UR55S:XaU$UR5H nUS:XdM Us $ URU5$rT)rrr:r )r1rr s rrequire_stride1ExternKernel.require_stride1#sR  # #1<<>"a',,.Q;H)~~a  rc  UcUceUR5S;a U(dU$[U5(Ga/[UR5[5(aU(a[ X5=(a( [ UR5R5(+n[USSU(aJ[[RRRUR5R55OUUS9 U$[USSSUUS9 U$[UR5[[45(ayU(a$UR5R!U5(d>U(aG[#X1R5RUR%55(aUb ['X5$U$[UR5=n[(5(a[UR+5=n[5(a [-S5e[U[5(aOU(aUR!U5(d0U(a+[#X7RUR%55(aU$[U[.5(akU(a$UR5R!U5(d>U(a9[#X1R5RUR%55(aU$[U[05(a[UR2[45(a[UR2[65(d[UR95=n5(a[;US5(ao[UR2[<5(dPUR?UR25UlU(aURAXUS9$U(aURCXUS9$Sn UR%5n Ub[RRn [G[IUR%555V s/sHJn U RKX<S 5(dMU RMUR%5U S 5(dMHU PML n n U H.n [NRPRRRUXS S 5nM0 URWU5n[USSUUUS9 U(a[ X5(deU$U (a<U bUce[NRPRRRYX5n['X5$U$![Da GNgf=fs sn f) N)rrETF)rrrrUrzHthe MutationLayoutSHOULDREMOVE's real layout shouldn't be FlexibleLayoutrTrrrrE)-rrrrrrr4r rVrrsr r rrWrrrKrrarrr(rrrTrvrZrwrrd r require_stride_orderrequire_exact_stridesrrrrOrrrloweringslice_r rL)r1rrrcrUuse_current_stride_ordermutation_layoutrrw expanded_dims orig_sizer rrHs rrequire_stridesExternKernel.require_strides-s M$=== ;;=F "=H ! # #!,,..990R00K3ALLN4I4IJJ-*#(- 8 - ! 0 0 H H$%LLN$9$9!" "'&3H*#(-%)&3&3 HALLN[/,JKK1<<>;;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% NN1  !''   5a?? ?? (]-F FF((//=A21D DW'  s* 6T3T3,U &U7U3 UUc URUUVs/sH:n[U[R5(aURR OUPM< snUS9$s snf)N)rcrU)r rrSymIntrrT)r1rrcrUrs rr "ExternKernel.require_exact_stridess]"" KXKXaz!U\\:: A=( #  sAA c"URXUS9$)N)rrU)r )r1rrrUs rr !ExternKernel.require_stride_orders""1"OOrc.URU[5$r)r rr4 s rrequire_channels_last"ExternKernel.require_channels_lasts''+<==rc.URU[5$r)r rr4 s rrequire_channels_last_3d%ExternKernel.require_channels_last_3ds''+=>>rcSSjnU"U5(aU$URU[RUR555$)NcUR5nU[RR ;=(a' [RR UR $![[4a gf=fr')rzAttributeErrorrrsr r is_mkldnn)rrs ris_mkldnn_tensor9ExternKernel.require_contiguous..is_mkldnn_tensors] zz|177,,,R1B1B41H1R1R R#$78  sAA0/A0rrrrr rrr)r1rr s rrFExternKernel.require_contiguoussC S A  H,,>44QZZ\B rchURU[RUR555$rr! r4 s rrequire_contiguous_strides'ExternKernel.require_contiguous_stridess-(( ~00>  rcgrrrs rr ExternKernel.apply_constraintrJrc[U[5(d[U55e[U[5(d [U5nUR(dS5e[ U5n[ UR5nX4:aq[ RSURXC- 5 [X45H?nURUSnURXb;aX&OURUS5 MA U$)Nz/ExternKernel.arg_properties should not be emptyzv%s has %d unprovided positional arguments. Will check if they are in the keyword arguments or will use default values.rr) rrrrrw rr'rrv rr)rrrn_args n_pos_argsrarg_names rfill_non_provided_args#ExternKernel.fill_non_provided_argss$))54:5)$%%:D""U$UU"T,,-    II^  #   6...q1&9 )$,,Q/@/ rc[RR(Ga`/nSnU(aeUR(aT[ UR 5[ U5:XdS5eURVs0sHoDR S5U_M nn[UR 5HupVUb3UceUR X5nU(aUR S5OSnOb[ UR5U-n UR(a7U [ UR5:aURU R S5OSnUR[RRRXh55 M U$UR V s/sH,n [RRRU 5PM. sn $s snfs sn f)NzDnames passed to codegen_const_args does not match self.constant_argsrr) rsr r rw rrn rrrqrrval_to_arg_str) rrrname_to_arg_propertiesrN rrproptype_rrs rcodegen_const_argsExternKernel.codegen_const_args,s 77   F &* ",,4--.#e*<Z<594G4G*4GSGGFOS(4G'*"$"4"45)5 ,,,155eh?D04DHHV,$Edkk*Q.C ..3T=P=P9Q3Q++C044V<!  agg22AA!KL6MDHDVDVWDVqAGG((77:DVW W'*&Xs 4G3G c [RR(aDURb7UR /UR QUR QUR5nSnOUR nSn/n[U5HupE[RR(aUR(aU[UR5:dS5eURURS5nUR[RRRXV55 MUR[RRRU55 M U(aURUR!55 U$)NFTz-Invalid access to ExternKernel.arg_propertiesr)rsr r rv r, rqrn rrrw rrrrr/ rr3 )rrqneed_codegen_constant_argsrrrr2 s r codegen_argsExternKernel.codegen_argsMs* 77  4#3#3#?003$++3 2 23T[[F*/ &[[F)- &f%DAww""**q3t7J7J3K/KCK++A.226: AGG00??IJ AGG00??BC& & KK//1 2 rc X;aURU5$XR;aURRU5$URRU5=nbURS5$[US35e)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)rrrx r()rr+ rrN s rget_kwargs_valueExternKernel.get_kwargs_valuegst  ::h' ' {{ ";;??8, ,))--h7 7C D77?+ +z)GHIIrc [RR(aURb[ UR 5S:Xa/$/nUR HnU(aUS:XaMURU5n[U[5(aURU5 MKURceURRU05RS5nUR[RRRXE55 M U$URR!5VVs/sH3updUS[RRRU53PM5 nnnU$s snnf)Nrrrr)rsr r rv rr rt r: rrrrx rrr/ rr)rskip_outrr+ rr2 ks rcodegen_kwargsExternKernel.codegen_kwargsts8 77  +D4F4F0G10L F >>E 1))(3a&&MM!$11=== 2266xDHHPEMM!''"6"6"E"Ea"OP?" !KK--//DA#Qqww++::1=>?/  s6:E5cURbSURRn[USS5nURSS5nUR SS5SnUSU3nU$SnU$) Nrunknown_namespacer r r rEr unknown_op)r r rr rsplit)rr  op_namespaceop_names r get_op_nameExternKernel.get_op_namesv << #\\((F"6<9LML'//'BL'..sA6q9L%ax0G#Grc [R(a[RR(d[ UR 55S:Xag[RRRUR 55n[RRRUR55nUR5nURSUR5SUSUSU<S3 5 ggg)Nrzassert_size_stride(rr") rF size_assertsrsr r rnrrcodegen_shape_tupler:rG rrz)rr rr rF s rcodegen_size_asserts!ExternKernel.codegen_size_assertss   qww':':T]]_-277'';;DMMOLDWW))==doo>OPF&&(G   %dmmo%6bb7+UVW (; rc j[R(a[RR(d~UR 5nU[RR ;nUR5nU(a!URSUS[SU<S35 gURSUSUS35 ggg)Nzassert_alignment(rr"z # buffer z (op: z) is assumed to be not aligned) rFalignment_assertsrsr r rzrrG rrg)rr ralignedrF s rcodegen_alignment_asserts&ExternKernel.codegen_alignment_assertss  # #AGG,?,?==?D!''";";;G&&(G!!'vR/@7+QO!!vVG94RS-@ #rc[RR(a[RR (agUR 5 UR5nURSUSUS35 g)zS Track outputs of fallback operators if config.test_configs.track_memory_lifecycle Nz track_tensor(z, 'z')) rF test_configstrack_memory_lifecyclersr r "write_memory_track_allocation_oncerzr)rr rs rcodegen_memory_tracking$ExternKernel.codegen_memory_trackingsV""99QWW=P=P 224}}M$s4&;P   - -d B%%C "3' 'A&;;%%'C "3' 'A(rc ,[USS5nSU<3/nU[R"U5Vs/sH'nURS[XR53PM) sn- nUR SUR <35 UR U5$s snf)Nrq zpython_kernel_name=rr)rrfieldsrrrr)rr rrs rr!ExternKernel.__str__sd$8$? !+ 1  $++D1 1zzl!GD**56 71    |D$4$4#789u%%  s.B)rx rrw rn rr r rr rv rt rp rq r rz rNNNNrN)rrrUr#rqr# rn r0 rdict[str, Any] | Nonerp ro rq rrr rrt rs rv ru rrrrrr)r rxr rrrr rxrrrr rrr)rq rrrrr)rrrr)r rrr rr rzftuple[Any, list[Any], list[Any], Callable[[Any, Any], Any], dict[sympy.Symbol, pytree.KeyPath] | None])rrrrZrO)NNF) rrrSequence[int] | NonercrrUrrrr)rrrcrrUrrr)rrrrrUrrr)rr0 rrrr0 )rrrrrr)r+ rrr rr )r= rrr)rz'tuple[list[Sequence[Expr]], list[Expr]])rztuple[Expr, Sequence[Expr]]r)Arrrrrrn rrrrrrp rq rr rrt rv rw rx ry rz r rrrr rr r r} r~ r r rr rlr r rS r r r r r r rFr$ r r, r3 r7 r: r? rG rL rQ rW rZ rb r^rvr!rkrrrs@rrr>s $&M=%(..tDFND*.K'.%) )"&OZ&4?3D3D4!=(,K$+26N/63>3D3D40:>6=->t-T*T(*(,.2)-&*79+/...4 . % . & .,.'.$.(5.). ..@.#$J! HLM+M:DM M""#;J +$   ^ !^ *-^ 9<^  ^ ^ @E E N!!!!F!!'+37# a a$a1 a  a  aaFQV    '9  JN      DIPP,P=AP PP >>??"   "!"+9" "HXB4 J4    =$'BN+$) !  ! ,  &Hrrcx^\rSrSrSSjrSSU4SjjjrS SjrSrU=r$) ExternKernelOutic&URU5 gr)generate_extern_kernel_outr s rr ExternKernelOut.codegen **40rc :>URU5n [U [5(d[U 55e[T U]SUU UU=(d 0SUUUU 5 [ RRU5Ul [ RRU5 gr) r rrrrrrsr r rr ) rrUrqrn rrp rq rr rt rv unwrapped_inputsrs rrExternKernelOut.__init__ s ..v6*H55Mt[R"[R5n[TU][ U[RU/S9/UR URU//SS[RRS9 g)Nrzaten.randint.low_outzat::_ops::randint_low_out::call)rUrqrn rq rr rv ) rrrNrrrWrrMr randintlow_out)rri r limitsrs rrRandomSeeds.__init__.slU[[) kkW !::vzzE7;5> ,,  rr)ri rr rrrrrrrrrrrs@rr~ r~ -s   rr~ c|^\rSrSrSSjrSS U4SjjjrS SjrS SjrSrU=r $) rd iAc&URU5 gr)generate_extern_kernel_allocr s rr ExternKernelAlloc.codegenBs,,T2rc Z>URU5n [SU 55(de[T U] SU[ [ [ U 5UU=(d 0SUUUU5 /Ul[RRU5Ul [RRU5 g)Nc3B# UHn[U[5v M g7frr )rrs rr-ExternKernelAlloc.__init__..QsC2BQ:a((2Br) r rrrrrrr rsr r rr ) rrUrqrn rrq rr rt rv ry rs rrExternKernelAlloc.__init__Es ..v6C2BCCCCC    &!#3 4  Lb    )  ') GG++D1  ""4(rcgr'rrs rrZ!ExternKernelAlloc.should_allocateer*rc[errrs rr "ExternKernelAlloc.apply_constrainthrr)rr rn )rNNNrN)rUr#rqrm rn r0 rrm rq rrr rrt r0 rv ru rrrr) rrrrr rrZr rrrs@rrd rd As3(*(,)-&*79+/))!)% ) & ) ' )$)(5))) ))@""rrd ch^\rSrSrSrS U4SjjrS SjrS SjrS SjrS Sjr Sr U=r $)r ilzH An output buffer that represents the mutation of a pre-existing buffer c>[TU]SUS9 UR5n[RR U5 U/UlX0l[RRU5Ul grq) rrrzrsr rmutation_names mutating_noder r)rrU mutated_noder mutated_node_namers rrMutationOutput.__init__qs` d62(113 ##$5601(5GG++D1 rcUR$r)r rs rrMutationOutput.get_defining_op{s!!!rcUR$r)r rs rr!MutationOutput.get_mutation_names~rrcgr'rrs rrZMutationOutput.should_allocater*rcjUR5nSU5Vs/sH nUcMUPM sn$s snf)Nc3`# UH$n[RRU5v M& g7fr)rsr try_get_buffer)rrs rr6MutationOutput.get_mutation_buffers..s"P..t44s,.)r)rr rs rget_mutation_buffers#MutationOutput.get_mutation_bufferss@002QP P P   s00)r r r)rUr#r rr rrrrdrrrrm ) rrrrrrrrrZr rrrs@rr r lsF2 2062GP2 2"#  rr c^\rSrSr%Sr0rS\S'\S Sj5r\S Sj5r S U4Sjjr S Sjr SS jr S r U=r$) TMADescriptoriaL 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[U5S:XdeUSS:Xa[U/USQ76$USS:Xde[U/USQ76$)Nrr experimentalrEri)rTMADescriptorExperimentalTMADescriptorStable)r1r[tma_metas r _create_implTMADescriptor._create_impls\8}!!! A;. (,VBhqkB BA;(* **&v< < [TU]S[[UUR 5S95[ [ [U5[U5S5 Xl [RRU5Ul [RRU5 g)NrS)rrrrZrrrryrr[rsr r rr )rr[rqrn rs rrTMADescriptor.__init__s   !,,.  &!6 * -   GG++D1  ""4(rc&URU5 gr)generate_tma_descriptorr s rr TMADescriptor.codegen''-rcUR$r)r[rs r get_tensorTMADescriptor.get_tensorr'r)rr[)r[rr ztuple[str, tuple[Any, ...]]rr )r[rrqr0 rn r0 rrrn r)rrrrrr rrlr rrr r rrrs@rr r s(*F $)=='B= =='B ))&3)DQ) )..rr cH^\rSrSrSrSSU4SjjjrSrU=r$)r iz 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^>[U5S;de[U5[U5:XdeUcUR5RnX lX0lX@l[UR5UlU/n/URQURQUR Pn[TU]!UUUS9 g)N)rErr[rqrn ) rrr r  block_dims element_sizermrr)rr[r r r rqrn rs rr"TMADescriptorExperimental.__init__s4yF"""4yC O+++  !++-66L $( N  YY __     '  r)r r r rmr) r[rr list[int | torch.SymInt]r r r rrrrrrrrrrrrs@rr r sG$(   ' -  !     rr c0^\rSrSrSrSU4SjjrSrU=r$)r iz the new host-side TMA descriptor API (the ones obtained via TensorDescriptor.from_tensor). See also TMADescriptorExperimental for the old API. c2>X l[TU] UU/US9 g)Nr ) block_shaperr)rr[r rs rrTMADescriptorStable.__init__s&& 8%  r)r )r[rr r r rs@rr r s  rr cR^\rSrSrSSU4SjjjrSSjrSrU=r$)SubgraphBuffericV>[TU]SX5 X0lX@l[R R U5Ul[R RU5 [R RURXE5Ul [UR5(de[UR5nUHTnXRRUR'URRR!UR5 MV UV s/sHoRPM sn UlSSKJs Jn [R*"UR5 SSSS.n 0U EU=(d 0En U R-U 5 URR."UR6 SSS5 U(a<URR0H"n U R3UR555 M$ SSS5 gs sn f!,(df  N_=f!,(df  g=f)NrFATEN) max_autotunemax_autotune_gemmmax_autotune_gemm_backends)rrrkexample_inputsrsr r rr make_subgraphsubgraphr! rqrtr=graph_input_namesr sym_inputstorch._inductor.configrrFset_graph_handlerrrun operationsrr)rrUrrkr subgraph_namer r sym_inpsym_varinductor_config base_patchesmerged_patchesrers rrSubgraphBuffer.__init__s v3,GG++D1  ""4(--dgg~U  ,,,,(5 !G7>MM & &w|| 4 MM + + 2 27<< @"8BBzG<.CodegenGraphi7c2XlURUlgrr r)rr s rr5SubgraphBuffer.codegen..CodegenGraph.__init__8s" !JJ rr N)r ry)rrrrrrrrr CodegenGraphr 7s 'rr )r! rqr'codegen_subgraph_with_flattened_outputsr r r)rr r r outer_inputss rr SubgraphBuffer.codegen6s| ' '   ,,,,7;{{C{!++-{ C77  ' -doo - - YYK DsB )r rkrr r r) rUrrrrktorch.fx.GraphModuler list[Any]r rr rm rn rrrrrr rrrs@rr r s_15,A,A",A ! ,A " ,A  ,A.,A,A\   rr c^\rSrSrSrSSjrSSjr\SSj5rSSjr SSjr \ "S5SSU4Sjjj5r SS jr SU4S jjr\SU4S jj5rSS jrSS jrSrU=r$)UserDefinedTritonKerneliFz6 A user-defined triton kernel (e.g. via @triton.jit). cz^SSKJn SSKJn UR UR 5m/n/n/n[ TU5(a[TS5(a&URU4SjTR55 O.[TS5(deURTR5 [TS5(a<TRH+nURTRRU5 M- O.[TS5(deURTR5 TR nTRmTX4U4$) Nr) Autotuner)kernel_side_table restore_idxc3V># UHnTRRUv M g7fr)r arg_names)rrr s rrBUserDefinedTritonKernel.get_kernel_and_metadata..Xs$*4FqFII''*4Fs&) restore_value reset_idx reset_to_zero)triton.runtime.autotunerr *torch._higher_order_ops.triton_kernel_wrapr get_kernel kernel_idxrrrr r r rrr r configs)rr r r restore_value_argsreset_to_zero_argsrr s @rget_kernel_and_metadata/UserDefinedTritonKernel.get_kernel_and_metadataKs6P"--doo>(*(* fi ( (v}--"))*4:4F4F*v7777"))&*>*>?v{++))A&--fii.A.A!.DE*v7777"))&*>*>?nnGYYFw4FFFrc[R(dgURR(dg[ UR R 5S:wag[ UR5S:Xde[URS[5(dg[URSR[5(dg[URSRR[5(dg[URSRRR[5(dg[SURSRRRR55(dgg)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. FrErc3*# UH oS:Hv M g7fr4rrs rr.sN#Ma6#MrT)rF*epilogue_fusion_user_defined_triton_kernel arg_accessescan_fuse_epiloguer kernel_storesstores mutable_argsrrrTrxrrurr rs rr )UserDefinedTritonKernel.can_fuse_epilogueks%0@@  22 t!!(( )Q .4$$%***$++A. ::$++A.33Z@@$++A.3388.II$++A.3388==yIIN4#4#4Q#7#<#<#A#A#F#F#M#MNNNrc"URUSS9$)N)epilogue_fusion_codegenr s rr UserDefinedTritonKernel.codegens}}Wd};;rc$URX5$)zA epilogue_fusion: (fused epilogue node, modified kerel src code) r )rr r s rcodegen_with_epilogue_fusion4UserDefinedTritonKernel.codegen_with_epilogue_fusions }}W66rc T^SSKJn UR5unnnnURUUURUUUR U5unn n n UR V s0sHoURU 5_M n n U(a{[URRR5S:Xde[[URRR55RnX;deUunnXU'URVs/sHnURPM snmURVs/sH"nUR (dMUR"PM$ nn[%U4SjU55n/n/n/n/n[&R("U R+5[-[&R."S5U 55GHunnUU;aU"5(aMUR1U5 UR1U5 [3U[45(a@UR1UR755 UR1UR955 M[3U[:[<[>[@RB45(a-UR1U5 UR1[EU55 MUU;a)UR1S5 UR1[:5 GMUcYU"5(a)UR1S5 UR1[:5 GMVURG5 URG5 GMy[IS[EU5S U35e URKX5 URMUUUUUU U S URO5URPRS 9 gs sn fs snfs snf) QOverrides the parent member. See https://github.com/pytorch/pytorch/issues/151692r)triton_version_uses_attrs_dictrEc3.># UH nTUv M g7frr)rrr s rr3UserDefinedTritonKernel._codegen..s$F:aYq\:srr8NzUnsupported arg type: r!T) arg_typesraw_argsraw_keys triton_meta inductor_metar*r original_fxnode_name))r r r !define_user_defined_triton_kernelrgridrt r: rr rrrlrmrparams is_constexprnumr>rFrdrrrepeatrrrrrrrrrrrrr0rr generate_kernel_callrr )rr r r r r r r new_namer r extra_launch_argsr> named_argsmutable_arg_nameepilogue_computed_bufferrr constexprsconstexpr_namesrr raw_keys_filteredraw_args_filteredrrN r s @rr UserDefinedTritonKernel._codegens2 I  ( ( *       5 5   KK   II      261S1S 1SAt$$Q' '1S   t((44;;<A AA#D):):)F)F)M)M$NOTT #1 11*9 ' $a+C' (%+]]3]QVV]3 %+]]E]anneaee] E$$F:$FF! ')')"    I$4$4R$8:K L ID#&+I+K+K  $ $T *  $ $S )#v&& C1134  1C#udEJJ!?@@ C   c+( B  %233KKO$$S)%))+%))+),B49+RPSu*UVVI L W/$$  &&#'??$!%!2!2 % u 4EsNN /N%N%cP>[TU]U5[URU5-$r)rrvr'r rs rrv,UserDefinedTritonKernel.get_free_symbol_usess- w+M:=M II}>   rc[5$rr=rs rr0UserDefinedTritonKernel.get_unbacked_symbol_defs rrc >/n0n/nUR5Hup[U [5(aX[R UR U 55n X;a[ RXU5n URU 5 XU'MrURU 5 XU'M [U5S:wdeUSR5Ul [U[5(d[U55e[TU]=S[!URS9U[#U5U5 XlX lUR)5up n [+U S5(deU R,Vs/sH oU;dM UPM snUlSSKJnJn [U 5S:aU SR6O0nSSKnU R:UlUR?UR<5Ul U"UR<5Ul!X@l"U"U 0UEUEU5Ul#URFRHRJVs/sHBn[URMURN5[5(dM3UURNPMD snUl(URPVs/sH!n[S[!URS9UU5PM# snUl*[VRXR[U5 gs snfs snfs snf)Nrr%r )identify_accessed_tensorsidentify_triton_stores).rrrrr2 rS r rrrrr rrrrrrr r r rr rt r r2 r3 rastr  kernel_srcparse kernel_astr kernel_argsr rrrrr r r rsr r )rr r tma_descriptor_metadatar8 rqrrn r> rrr r rrN r2 r3 autotuned_kwargsr4 rrrs rr UserDefinedTritonKernel.__init__#s "$&&( %%'DA!Y'' 99$:L:LQ:OP/%,,Q0JKA a q $$Q'q (6{aQi**, &(++9T&\9+   dkk *  -   % $ < < >Av{++++!++. +Ck/AC+. * 14G q0@71:,,b!**))DOO43DOOD&5  /{ /. / # ((44;; ;+//#((3Y? "K !; ((! ( :T[[93 E(!  ""4(K. 6 ! s K,K)2KK(K"c f>[R(d[TU] 5$URR R Vs0sH5nURURURR5_M7 nn[URR R5n[X0R5VVs0sH upURUR5_M" nnn[R"[!URR R Vs/sHnUR#U5PM sn5[!URR RVs/sHnUR#U5PM sn5[!5S9nU$s snfs snnfs snfs snfNr)rFr rrgr rrkrr8 rzrrrr rGrr>rename) rformal_arg_dep read_renamesformal_arg_writes mut_output write_renamesrrrs rrg'UserDefinedTritonKernel.get_read_writesss@@7*, , #'"3"3"?"?"E"E "E   !1!1.2E2E!F!O!O!Q Q"E  !!2!2!>!>!E!EF/2!#8#8/ /*   !4!4!6 6/  #-- $00<<BBBJJ|,B  $00<<CCCJJ}-C #  9   sX7 7e %e ;e   e N56$) ! ! 7 N)N) N) "0 N) $ N) N)`##J+rr ch^\rSrSrSrS SjrS SjrS SjrS SjrS U4Sjjr Sr U=r $)InplaceBernoulliFallbacki= This needs to be a custom class to handle mutation properly c  [SUR55(deSUR5un[RR(a\UR UR 5SUSSR[[UR55SUR35 gUR UR 5SUSSR[[UR55SUR35 g)Nc3B# UHn[U[5v M g7frr rrs rr3InplaceBernoulliFallback.codegen..s>+Q:a((+rc3^# UH#n[[U5R5v M% g7fr)rrrrN s rrrO s"I[VQ1133[s+-rrz, NULL)r") rrqrsr r rr rrreprrn ending)rr rs rr InplaceBernoulliFallback.codegens>$++>>>>>IT[[I 77     '')*!A3b3tTEWEW;X1Y0ZZabibpbpaqr    '')*!A3b3tTEWEW;X1Y0ZZ[\c\j\j[kl rcgr'rrs rrZ(InplaceBernoulliFallback.should_allocater*rc&URS5/$rr' rs rr+InplaceBernoulliFallback.get_mutation_names"##rc[5$rr=rs rr1InplaceBernoulliFallback.get_unbacked_symbol_defsrrcR>[TU]S[UR5S9UR U/5UUS9 [ R RUR55 [ R RU5Ul [ R RU5 g)Nr%rv ) rrrrr rsr rrzr rr )rrv rrn rs rr!InplaceBernoulliFallback.__init__s   alln -    $ #   ##AJJL1GG++D1  ""4(rrrn rrr)rv rrrrn r rr rrrrrr rZrrrrrrs@rrJ rJ sF $ )' ),2 )DG )  ) )rrJ c^\rSrSrSrS SjrS SjrS SjrS SjrSU4Sjjr \ SSSjj5r S r U=r $)InplaceCopyFallbackirK cNUR5up#nURX2U5 gr)r7 codegen_device_copy)rr rr  non_blockings rr InplaceCopyFallback.codegens%#'#4#4#6 <##Cl;rcgr'rrs rrZ#InplaceCopyFallback.should_allocater*rc&URS5/$rrW rs rr&InplaceCopyFallback.get_mutation_namesrY rc[5$rr=rs rr,InplaceCopyFallback.get_unbacked_symbol_defsrrc >[TU]SUUUSSS9 [RR USR 55 [RR U5Ul[RRU5 g)Nz aten.copy_aoti_torch_copy_)rq rr r) rrrsr rrzr rr )rrUrqrn rs rrInplaceCopyFallback.__init__sr     +.   ##F1I$6$6$89GG++D1  ""4(rcX4Vs/sHo@RU5PM nnU4n[[UR5S9UU5nU$s snfr~)rS ra rr)r1rr rd rrqrn rs rrInplaceCopyFallback.createsU25 ; 1##A& ;% $ cnn. /     .s 9[!!##[srrr") r! rqrrQ rn rr rrR )rr argrefss rr $MutatingFirstArgExternKernel.codegens ,,,, 9T[[ 9 t)) *  ##%&a '(:';1W^^[U[5(dS5e[TU] S[ UR 5S9UR U/5U4S9 [RRUR55 [RRU5Ul [RRU5 SUlSUl[U[ ["[$45(d['U55e[RR(R+UR,R55 g)NzTODO: dynamic shapesr%)rn z"inductor_ops.resize_storage_bytes_z&torch::inductor::resize_storage_bytes_)rrrrrrr rsr rrzr rr rq rr rvrxrrnever_reuse_buffersrrT)rvariablerrs rrResizeStorageBytes.__init__s(C((@*@@(   h113 4    +#+  ##H$5$5$78GG++D1  ""4("FG(Xz9$EFFVXVF ##'' (>(>(@Ar)rr rrq )r rrrrrr rs@rr r sBBrr c6^\rSrSrSU4SjjrSSjrSrU=r$)SetSourceTensorKerneli-c>UR5 [TU] UR5X/S[R R RRS9 [U[[[45(d[U55e[RR R#UR$R'55 [RR R#UR'55 [RR R#UR'55 UR)5n[+[-US9X5[+[-US9X 5/Ulg)Nz!torch.ops.aten.set_.source_Tensor)rq rv r%)rQrrrrrqr set_ source_Tensorrrvrxrrrsr r rrTrzrr rr )r self_tensorstorage_tensorr rs rrSetSourceTensorKernel.__init__.s$$&   % % '  )B ++99  +*i'HII 4 L  I ##'' (8(8(A(A(CD ##''(?(?(AB ##'' 8**, :V4k H :V4n K! rcFURS5URS5/$rTrW rs rr2SetSourceTensorKernel.get_inputs_that_alias_outputBs"DOOA$677rr{ )r rr rrrr)rrrrrrrrrs@rr r -s (88rr c^\rSrSrSrS SjrS SjrSSjrSSjrSSS .SU4S jjjr S r U=r $)ScatterFallbackiFz 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&URU5 gr)generate_scatter_fallbackr s rr ScatterFallback.codegenMs))$/rcgr'rrs rrZScatterFallback.should_allocatePr*rcpURSn[U[5(deUR5/$rr% rrss rr"ScatterFallback.get_mutation_namesSs1kk!n#v&&&& rc[5$rr=rs rr(ScatterFallback.get_unbacked_symbol_defsXrrNTrE include_selfc d>[U[5UlUR(a&X$U4Vs/sHoRU5PM n nU4n O$X$4Vs/sHoRU5PM n nX54n [T U]S[ UR5S9URU 5U XgS.[U5SS/US9 [RRUR55 [RRU5Ul[RR!U5 gs snfs snf)Nr%r rEr )rq rt rv )rr src_is_tensorrS rrrrr rrsr rrzr rr ) rrv rrHrr rEr rtensorsrn rs rrScatterFallback.__init__[s(Y7   78oFo))!,oGF FM78jAj))!,jGA JM   alln -    (  <";/+3^*D#   ##AJJL1GG++D1  ""4(%GBs D(D-)rr rn rrq r)rv rrrrHrrrr rrErr rrrr_ rs@rr r Fs| 0 "!!)!!) !) !)  !)  !)!)!) !)!)rr cp^\rSrSrSrS SjrS SjrS SjrS SjrS U4Sjjr Sr U=r $)IndexPutFallbackizI This needs to be a custom class to handle mutation and indices properly c&URU5 gr)generate_index_put_fallbackr s rr IndexPutFallback.codegens++D1rcgr'rrs rrZ IndexPutFallback.should_allocater*rc&URS5/$rrW rs rr#IndexPutFallback.get_mutation_namesrY rc[5$rr=rs rr)IndexPutFallback.get_unbacked_symbol_defsrrc >X0lUVs/sH ofcMUPM nnX$/UQVs/sHo RU5PM nnSn [T U] S[ WR 5S9UR U5U4SU US9 [RRURS55 [RRU5Ul [RRU5 gs snfs snf)Naoti_torch_index_put_outr%zaten.index_put_)rq rr rv r)rrS rrrrr rsr rr' r rr ) rrv rrr accumulater valid_indicesr rr rs rrIndexPutFallback.__init__s $+=GqG =342M}2MN2MQ%%a(2MN4   alln -    ( M0+#   ##DOOA$67GG++D1  ""4(>NsC-C-C2)rrrn rrr) rv torch._ops.OpOverloadrrrr rr0 r r rrr_ rs@rr r s`2$)*) ) )  )  ) ))rr c2\rSrSr\SSj5rSSjrSrg) DeviceCopyic UR5nUceUR5(d[U5[RR ;a[ SUR555(a[RR(dn[RR(a>[RRU5 [RRU5 URU5$[RRU5 [RRU5 [S5 U4n[R!U5nSnUR#5(aUR%5n['UR(5=(a UR(S:H=(a UnUR(S:H=(a ['UR(5=(a UnU(a%[+U5(aSUR-5l[1[3UUR55UR#5UUS9UR7U5/U5$)Nc3Z# UH!o[RR;v M# g7fr)rsr r rs rr$DeviceCopy.create..sG4Fq***4Fs)+zDeviceCopy in input programr&Tr)rrr{rsr rrrrF aot_inductoruse_runtime_constant_foldingr add_device_inforrbrrFrr:rjrrrrYr rWrrS ) r1rr rd x_devicern r is_destination_pinnedis_source_pinneds rrDeviceCopy.creates<<>### Qqww'>'>>GA4D4D4FGGG''DDww""''/''1''/ / ' )78%  + +A . ::<<\\^F 8== ! KfkkU&: K|  MMU " Kvfkk': K|   5a 8 8'+ALLN $   /    q ! "   rcUR5n[U5S:XdeUR(a2URUSURR 5US5 gURUSUR 5US5 g)NrrrE)r7 rrp rc r)rr rs rr DeviceCopy.codegens{  "4yA~~     ' 'Q));;=tAw   ' 'Q1G1G1I4PQ7 SrrN)rrr rrd rrrrn )rrrrrlrr rrrrr r s- - ^Trr c^\rSrSrSrS SjrS SjrS U4SjjrS Sjr\ "S5SSSjj5r SSjr S r U=r $)DynamicSelectStorageOffsetia 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. c[5$rr=rs rr$DynamicSelectStorageOffset.get_readsrrcgr'rrs rrZ*DynamicSelectStorageOffset.should_allocater*rc>[TU]S[[R"S5S9/5 XlX lX0lX@lXPl X`l gNr&r%) rrrrr unbacked_offset_symbolr base_offsetbase_dim_striderr)rr rr r rrrs rr#DynamicSelectStorageOffset.__init__sG ze1DErJ'=# &.  rc.[UR/5$r)r>r rs rr3DynamicSelectStorageOffset.get_unbacked_symbol_defs s466788rc.[URU5$r)r'rrus rrv/DynamicSelectStorageOffset.get_free_symbol_uses s  M::rc6URXRS9 g)NrQ )codegen_dynamic_select_indexrr s rr "DynamicSelectStorageOffset.codegen s,,T,Dr)r r rrrr rr)r sympy.Symbolrr r sympy.Symbol | intr r rr rrrrrrrrn rrrrrrrZrrr^rvr rrrs@rr r s ,(  ,  !  &989$);!; !;:; EErr c^\rSrSrSrS SjrS SjrS U4SjjrS Sjr\ "S5SSSjj5r SSjr S r U=r $)DynamicSliceSizei 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. c[5$rr=rs rrDynamicSliceSize.get_reads& rrcgr'rrs rrZ DynamicSliceSize.should_allocate) r*rc>[TU]S[[R"S5S9/5 XlX lX0lX@lXPl gr ) rrrrr unbacked_size_symbolrrrr)rr rrrrrs rrDynamicSliceSize.__init__, s> ze1DErJ$8!   rc.[UR/5$r)r>r rs rr)DynamicSliceSize.get_unbacked_symbol_defs< s444566rct[URU5R[URU55$r)r'rrrrus rrv%DynamicSliceSize.get_free_symbol_uses? s0  M:@@ TXX} 5  rc&URU5 gr)codegen_dynamic_slice_sizer s rr DynamicSliceSize.codegenG rw r)rrrrr rr) r r rr rr rr rr rrrrn r rs@rr r  s *"  !  !  7./$) ! ! 0 11rr ch^\rSrSrSrS SjrS SjrS U4SjjrS SjrS Sjr Sr U=r $)riK z3 The result of a call to aten._local_scalar_dense. c[5$rr=rs rrDynamicScalar.get_readsP rrcgr'rrs rrZDynamicScalar.should_allocateS r*rc>UR5 [TU] S[[R "S5S9UR U/55 XlX lgr ) rrrrrr r symkeypath)rr r rTrs rrDynamicScalar.__init__V sI   *ELL$78$:M:Mtf:U  rc.[UR/5$r)r>r rs rr&DynamicScalar.get_unbacked_symbol_defs` s488*%%rc&URU5 gr)codegen_dynamic_scalarr s rr DynamicScalar.codegenc s&&t,r)r r rr)r r r zpytree.KeyPathrTrrrrrn ) rrrrrrrZrrr rrrs@rrrK sF*8@F &--rrc^\rSrSrSrS SjrS SjrS U4SjjrS Sjr\ "S5S SSjj5r SSjr S r U=r $)rig z- The result of a call to aten._assert_scalar c[5$rr=rs rrAssertScalar.get_readsl rrcgr'rrs rrZAssertScalar.should_allocateo r*rcv>[TU]S[[R"S5S9/5 XlX lgr )rrrrr scalarr )rr r rs rrAssertScalar.__init__r s3   ell51 2    rcgrrrs rr AssertScalar.has_side_effects} rrc.[URU5$r)r'r rus rrv!AssertScalar.get_free_symbol_uses s  ];;rc [R(dg[[UR SS955n[ R R(ag[ R R(a^SUS3n[ R RRURSS9nURSUSURSUS 35 g[ R RRURSS9nURS US 35 URS [UR5S35 URUR!5S 35 g)NFrozstd::to_string(r")rszif (!(z()) { throw std::runtime_error("Expected z but received " + z); }zif not (z):z raise RuntimeError(z = None)rFscalar_assertsrlrmrvrsr r r rcodegen_cpp_sizevarr rr codegen_python_sizevarrQ rz)rr rV symbol_strsizevars rr AssertScalar.codegen s4$$  d44454IJK 77    WW *6(!4Jgg**>> e?G    !J488*Tfgqfrrwx gg**AA eBG    4 5    7TXX7GqI J    19 :r)r r rr)r rtr rrrrrrn )rrrrrrrZrr r^rvr rrrs@rrrg sR N+$)<!< !<,< ;;rrc*\rSrSr%S\S'S\S'Srg)ExternKernelNodei rrzexport_schema.NoderrNrrrrr r s I rr c ^\rSrSrSrSSS.SU4SjjjjrSU4SjjrSSjrSSjrSS jr \ SS j5r SS jr SS jr SS jrSr\SSj5r\ SSj5r\SSj5rSrU=r$)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. Nrz c>^[TT]U[U5[U5US9 STlU=(d 0Tl[ U[ RR[ RR45(dSUS[U5S35eUTl UTl Uc0OUTl TRce[R R#TR5 /Tl/Tl[ TR[ RR5(agSTRR)5;agTRR*n[ R,R.R1TR5(a-TR&R3USR555 gS Sjn UR6(aI[9TR5(d/U "TR5(d[;S TR35eTRTR<TR>5upS U4S jjn [ R,R.RAXU5H upU "X5 M g) Nr] F#Fails to create FallbackKernel for r! not supported_c10d_functionalrc8[RRUR5[RRR 5=(dD [ US5=(a1 [RRR UR;$)N py_kernels)r_C%_dispatch_has_kernel_for_dispatch_keyr DispatchKey Functionalizerr rds rhas_functionalize_impl7FallbackKernel.__init__..has_functionalize_impl sg88AA 588//==L)HHH((66"--G  rz'NYI: Can't generate FallbackKernel for c>>^[TR[R5(a+[U[[ 45(d[U55e[ R"TR5(a[U[ [45(aeUcgTRcgSUU4Sjjn[ R"TR5(aUbUH nU"U5 M gg[ R"TR5(deU"U5 g)Nc$>TRRUR55 TRceTRR(a<TR R[ [UR5S9UT55 ggr~) alias_namesrrz alias_infois_writer r rr)rinfors r add_aliasPFallbackKernel.__init__..handle_aliasing_and_mutation..add_alias!sj  '' 5222??++))00&z'H!TR,r)rrrr) rrrListTyperr library_utilsis_tensor_like_typer is_tensorlist_like_type)r rN r optional_tensor_argrs` rhandle_aliasing_and_mutation=FallbackKernel.__init__..handle_aliasing_and_mutation!s$))U^^44!#e}55@tCy@500;;&cE4=9999{&  44TYY???/2+!"5603#%88CCCC#r)rer rr)r ztorch._C.ArgumentrN r rr)!rrruse_runtime_dispatchrz rrr r r rrv r rrq rsr  warn_fallbackr r rr _libraryrmutates_and_returns_first_argrrz is_mutabler$rrqrn zip_schema)rrUr r nontensor_argsr rrz schemar rr& r rN rs` rrFallbackKernel.__init__ s*   +  . !  %*!!2!8b UZZ**EJJ,J,JK   X 04<. W X ","Nb &&222 d556'))+ d&& (F(F G G  !1!1!6!6!8 8 !!)) >>   = =d>N>N O O    & &{1~'>'>'@ A     *4+;+;<<*4+;+;<<%9$:J:J9KL **4;;8J8JK  >--88vNID ( 3Orc@>[TU]5nUR[RR R La]URHMn[U[5(dMUR[R"UR555nMO U$r)rrgrv r_prims rng_primsgraphsafe_run_with_rng_statern rr with_readrGrrz)rrrN rs rrgFallbackKernel.get_read_writes*!swg-/   u||55RR R))c>22"-"7"7$,,S\\^<#K* rc nURUR5UR[USS55$Nrz )(codegen_unbacked_symbol_defs_for_outputsrzr rr s rcodegen_unbacked_symbol_defs+FallbackKernel.codegen_unbacked_symbol_defs6!s0?? MMOT\\749Ld+S  rc[USS5=n(aL[[RRR U5nUce[ UR55$[ 5$r8 rr9rsr r rr>r+rrz resolveds rr'FallbackKernel.get_unbacked_symbol_defs;!_ '.A4 H H  H0  **,=H' ''hmmo. .< rc~[R"SS55n[UR5(deURVs/sHo!"UR 55PM nnUR X0R 5upE[RR(a[UR[RR5(aURXE5n[!URR"R$U5VVs/sH8upb[RR&R)X&R*5PM: nnnO9UVs/sH,n[RR&R)U5PM. nnUR,R/U5 U$s snfs snnfs snf)Nc*\rSrSr%S\S'SSjrSrg))FallbackKernel.codegen_args..ShimiF!r refcUR$r)rE rs rrk2FallbackKernel.codegen_args..Shim.__repr__J!s xxrrNr)rrrrrrkrrrrShimrD F!s H rrH )r dataclassr! rqrr rn rsr r rrv rr r r, rr r rr/ r rrw )rrH rr rrparams rr7 FallbackKernel.codegen_argsE!s\         ,,,,<@KKHKqtA//12K H**;8J8JK  77  :d.>.> @U@U#V#V..t.q!s)$(A**433'rrE)rr3 rrr[ r r ScriptObjectr(r*rrr>rrjr)r r rnon_torch_bind_tensor_argsrN devices device_setr s rrN FallbackKernel.find_device]!su $ J 1:a+IQ J # & ;3>S;C..BR's~~';GS1:  nell 3 3!(( (  UXX224DE  ^ , ,<<& & ntUm 4 4#$($J -7AJ&&vJGA7|q qz!!!&%,,7777&++&&!M"1: ? K TBs"F0F0F5$F50 F:>F:c2SSKJn U"UR5$)Nr) is_impure)torch._library.utilsrV rv )rrV s rr FallbackKernel.has_side_effects!s2))**rc.[UR[RR[RR 45(d+SURS[ UR5S35e[UR[RR 5(d_SURR5;aAURRR(a[UR5(a/$UR$)Nr r!r r ) rrv rr r r rrr r, r$r rs rr+FallbackKernel.get_inputs_that_alias_output!s   uzz44ejj6T6TU   2$2B2B1C2D$$%&n 6  4++UZZ-K-KLL"$*:*:*?*?*AA  ((33&t'7'788I## #rcP[UR5S::deUR$r6)rr rs rr!FallbackKernel.get_mutation_names!s'4&&'1,,,"""rc  [RSUR5UR5 [ U[ 5(d[ U55eURURUR5upURX5nURVs/sHnUR"U40UD6PM nnURn[RR(d/UQUQ$[!S/5nUR#XQU5nSSjn[ U[$R&R(R*5(a#UR-USUS5R.n OUR0R.n [3U 5S:XaBUR4(a UR4O UR6n U SR8n U"X5/n O:[;XR45V Vs/sHupU"U R8U5PM n n nURce[=UR5[>R@"URRC5UU 0S9S9n[RDRGU5 /UQUQ$s snfs snn f) 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  [U[R[R45(aUn[U[[ 45(a[ U5S:XdeUSn[U[R5(aT[U[5(de[RR[R"UR5S9S9$Ube[RRSS9$[U[R5(a[UR5[R5(as[U[5(d[!U55e[RRUVs/sH%n[R"UR5S9PM' snS9$[U[R"5(a[UR5[R5(aUc8[RR[R$RSS9S9$[U[5(de[RR[R$R[R"UR5S9S9S9$[U[R&5(a[RRUS 9$[)S [!U535es snf) NrErr) as_tensorT)as_none) as_tensors)as_optional_tensor)as_intzUnsupported return type )rr TensorTypeNoneTyperrrr export_schemarurTensorArgumentrzr! getElementTyperr OptionalTypeOptionalTensorArgumentIntType RuntimeError) return_typercrs rhandle_single_outputFFallbackKernel.export_extern_kernel_node..handle_single_output!s+(8(8%..'IJJftUm44v;!+++ )Ck5+;+;<<%c62222(1188"/">">CLLN"S9;&;(11888FFK88Z**,e.>.>>>"&(33AT&\A3$--44$* #)C&44#,,.I#) 5 K););<<**,e.>.>BB>(1188+8+O+O+V+V$(,W,9 &ff5555(1188+8+O+O+V+V&3&B&B%+__%6',W,9K77$--44F4CC"%=d;>O=P#QRR7 s,LrrE)r rqr r )rr)rm z1torch.TensorType | torch.ListType | torch.JitTypercIRNode | Sequence[IRNode]rzexport_schema.Argument)$r'rrzrv rr] rr rqrn r, rt r: rsr aot_moder#serialize_inputsrr torchbind CallTorchBindr/ returnsr rr r r rr rf r<rextern_kernel_nodesr)rrrrordered_kwargsr  serializernamed_argumentsrn ru r rm output_arguments return_schemarcrs rexport_extern_kernel_node(FallbackKernel.export_extern_kernel_node!sl B MMO    $//;d;/**4;;8J8JK **4899 9  ! !# 0 09  !!ww+T+N+ +*44 $55fFK3 SJ3 S-3 S$3 Sj fe55??MM N NmmDGT!W5==Gnn,,G w<1 '+lldll8M8MG!!*..K 4[ JK .1,,-G   .H)M %!++.H  +++##'',,.&(   $$T*''''K ` s J*J c^^^TRnUceURS:Xa[U[RR 5(d[ U55e[RR(a2SSK J n [U5U;a[RSU5 STlOURS:Xa:[U[RR 5(d[ U55eOA[RR(a"U[ R"R$;Tl[RR(a[U[RR 5(aTR(dS U4SjjmTR'TR(TR*5unm[,R."UUU4S jTR055n[3U4S j[5XRR6R8555TlTR;U5 TR(aTR=5nTR>ceTRceURATRC5TR>U4S jTRUTRD(a TRDO TRF5 OcURIT5 [TRJ[L5(a3TROU5 TRQU5 TRSU5 TRUU5 g) r Nr rr zG%s is missing a c-shim implementation, using proxy executor as fallbackT _quantizedc>[U[R5(aT"UR55$[U[R5$r)rrri rh NumberType)r is_numbers rr )FallbackKernel.codegen..is_numberC"s=a!3!344$Q%5%5%788!!U%5%566rc3J># UHnTR"U40TD6v M g7fr)r: )rr> rrs rr)FallbackKernel.codegen..P"s(?))!6v6?s #c3z># UH0up[U[5=(a T"UR5v M2 g7fr)rcomplexr )rrrr s rrr U"s2,DDA1g&A9Q[[+AADs8;cH>/TR5QTR5Q$r)r7 r? rsrr(FallbackKernel.codegen..c"s"F$++-F0C0C0EFr)rz torch.JitTyperr)+rv r rrr r rrsr r r r rr'rr( rFr custom_ops_to_c_shimsr rqrn rFrdrt rrr r r r| rq ,generate_fallback_kernel_with_runtime_lookuprzr r generate_fallback_kernelrUrrL rQ rW r: ) rr r r r args_iter exported_argsr rs ` @@rr FallbackKernel.codegen"sz!!!!!   v %fejj&;&;<< Jd6l J<ww""Lv;&;;KKa15D-    -fejj&;&;<< Jd6l J< WW f11GGG  % GG  65::#8#899-- 7 ..t{{D>+C+CD,)D % W%  $ $ ::R@RB5(d'U[RDRFRHLa[R,"S5m[U[RR5(aF[R RJ(d'T(a [LROUUTUWW UW TS 9 nUbU$UcT"[QTS 9UUWW UW S 9mO!T(dS 5eT"[STS 9UUWW UW S 9mSUUUUU4S jjmT"U/5n[U[6[T45(a UTl+U$[U[X5(a[UU5Tl+U$U/Tl+U$!,(df  GN=f)z9Create an instance of FallbackKernel from an _OpOverloadsNr)_is_functionalget_out_arg_nameslookup_manual_out_variantto_out_variantrEr r rr )rUrqrn rrq rv c38# UHn[U5v M g7fr)rrM s rr(FallbackKernel.create.."s!K{,s"3"3{rr&rz has_unaligned_inputr%rrz z"Not sure where to find device infoc >^^[T[[45(a/[T5"UUU4Sj[ [ T5555$[T[ 5(a<TR5VVs0sHup#UT"UT[T5U4/-5_M snn$[T[R5(a}[TRT5T T5n[R(dT(d[T5(d3[R R"R%UR&5 U$[T[(5(aT$[T[R*5(aTR,R.$[T[R0R2[445(d[7T5(a[9[;TS9T TT5$TbS[T5S35egs snnf)Nc3Z># UH nT"TUT[T5U4/-5v M" g7fr)r)rrgenerate_outputrrcs rrAFallbackKernel.create..generate_output.."s7$/$F1Iw4<:K9L/LMM/s(+r%zFallbackKernel output type z is not supported)rrrrrrrrrr[ rR r^ rF assume_unaligned_fallback_outputrprsr rrrrr rrTr rP r(r*rr) rcrrrYrr1r r r packeds `` rr .FallbackKernel.create..generate_output"s&4-00F|$"3v;/$FD))%+LLN$2g$v,9L8M.MNN$2FELL11!((0 ;;*,V44GG--11#((; FC(( FELL11{{'''..0@A (((f- ~1$v,?PQ~Es6%G3)rcr rzlist[tuple[Any, int]]rr )-r *_fused_moving_avg_obs_fq_helper_functionalrr rsr r9r r rrr r r[ torch._library._out_variantr r r r r rrWr r rGr rs r_make_out_variant_kernel_namerrN r rs rt rq higher_orderprintr ExternKernelMultiOut try_createrr$ rr r)r1r rrfake_incorrect_kernelscontextr r r r rz r r r r out_oprU out_resultr r r r r s` @@@@rrFallbackKernel.create"s#'"Q"Q!S  /1$79J9JKG!mG ""6;D;F;  !  fejj33 4 4 ELL: :   Ffnn--'/~26:!c*;F*C&D&I$)00(..0>//05^2245  '! ,"&"7!'DV'L & "!K{!KKn= vu66@@NN O O//555\\%(F vuzz44 5 5GG''-88"3$79 J%!!  !&)"3F ?? ?6!0"3F) ) V"."5 ge} - -$FN   & &"7^FN&YFNkWs L"" L1)r rr rv rz r r( rrUr#r rr rm r. r0 r r rrm rz )dict[sympy.Symbol, pytree.KeyPath] | Nonerrrrn rrq )r zSequence[torch.Tensor] | Noner r0 rr rr)rcrrrW)r rrr rr rr] )rrrrrrrgr: rr7 rrN r rrr| rr r^ rlrrrrs@rr] r] s)-u4HLu4u4u4& u4 & u4 + u4&u4Eu4 u4u4n   0#2#DQ# ##J+$*#w(rS3S3j    ]]rr] cl^\rSrSrSrS SjrS SjrSSS.S U4SjjjrSrU=r $) ComplexViewi##z9View a complex number as two dtyped numbers or vice versacgr'rrs rrZComplexView.should_allocate'#r*rc&URS5/$rrW rs rr(ComplexView.get_inputs_that_alias_output*#s"##rNr c ,>[TU]UUUUUUUS9 g)Nr )rr rrUr r r. r rrz rs rrComplexView.__init__.#s,      /  rrrrr ) rrrrrrZrrrrrs@rr r ##sqC$)-GK   &  &  +  & E    rr c"\rSrSrSrSSjrSrg)MemoryCheckKerneliD#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. cUR5 URup#n[U5n[U5nU(aURS5 SUSUSUS3nO SUSUS3nURU5 g)z.Override codegen to write direct function callzV# note: dont currently distinguish between buffers returned and dealloc'd in last stepzcheck_memory_step(allocated=z, freed=z, is_final_step=r"N)rV rn rQ r)rr alive_list dead_list is_final_step alive_repr dead_reprcalls rr MemoryCheckKernel.codegenK#s 224/3/A/A, }*% O    h 2*Xi[P`an`oopqD1*Xi[PQRD$rrNrn )rrrrrr rrrrr r D#s   rr c*\rSrSr%S\S'SSjrSrg)r$ i]#rr cUR$rr%rs rrMultiOutputLayout.get_devicea#r'rrNr)rrrrrrrrrrr$ r$ ]#s rr$ c^\rSrSrS SjrS S U4Sjjjr\"S5S S Sjj5rS SjrSSjr SSjr Sr U=r $)rR ie#cURU5 UR(d#URU5 URU5 ggr)codegen_multi_output!skip_size_stride_alignment_checksrL rQ r s rr MultiOutput.codegenf#s:$$T*55  % %g .  * *7 36rc>[TU]SX/S5 [RR U5Ul[RR U5 X0lX@lgr) rrrsr r rr rr )rrUr& rr rs rrMultiOutput.__init__l#sK vw3GG++D1  ""4( 1R.rczURSn[U[5(dU5eURU5$r)rqrrrv)rrpr/s rrv MultiOutput.get_free_symbol_usesy#s:[[^ *f--9z9-..}==rcz[UR5S:H=(a [URS[5$)NrEr)rrqrr rs rrZMultiOutput.should_allocate#s04;;1$ t{{1~'8 9 rcURVs/sHIn[U[5(dM[UR 55S:dM9UR 5PMK sn$s snfr)rqrr] rrrzr s rr(MultiOutput.get_inputs_that_alias_output#s\{{ "#~. C4467!; CLLN"   sA"A" A"c|^^[5nURHMn[U[5(dMUR [ R "UR555 MO UR5mUR5R5mSUU4SjjnUR5n[R(+=(d% USL=(d [UR5(+n[ R"X0R!5SUS9n[ R""UUR$[5S9$)Ncb>[U5S:Xde[R"TT"U5S5$r7 r9 r: s rr- *MultiOutput.get_read_writes..dummy#r< rrr' rr/ )r>rqrrrrGrrzrr1rrFrrjrrPrrr) rrkrsr- r should_normalizewrite_rwrrs @@rrgMultiOutput.get_read_writes#s/9l;;C#v&& ,..s||~>?}}//#002 ; ;"11 1 '~ '&++&&   33 ==?B2B &&??"   r)rrr rn r) rUr#r& rrlist[tuple[Any, ...]]r rrrrrrr) rrrrr rr^rvrZrrgrrrs@rrR rR e#s438 S S S' S ,0 S  S SM*$)>!> !>+>     rrR cl^\rSrSrSrSU4Sjjr\S Sj5rS SjrS Sjr Sr U=r $) ri#zMultiOutput for opaque objects.c0>[TU]XUSS9 X@lg)NTr )rrr )rrUr& r opaque_valuers rrOpaqueMultiOutput.__init__#s  SWX$0!rc[S5e)NzOpaqueMultiOutput has no dtype)r rs rr OpaqueMultiOutput.dtype#s=>>rcU$rrrs rr#OpaqueMultiOutput.wrap_for_lowering#rrch[5nURHMn[U[5(dMUR [ R "UR555 MO [[ R "UR55/5n[ R"UU[5S9$r= ) r>rqrrrrGrrzr)rrkrsrs rrg!OpaqueMultiOutput.get_read_writes#s.8l;;C#v&& ,..s||~>?0:  ! !$--/ 2 30 &&"   r)r ) rUr#r& rrr r r rr)rr)rrr) rrrrrrrr rrgrrrs@rrr#s^)111' 1  1  1??   rrc,\rSrSrSrSSjrSSjrSrg) 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. cgrrrs rrZ%AllocatingMultiOutput.should_allocate#rrclUR(d#URU5 URU5 ggr)r rL rQ r s rr AllocatingMultiOutput.codegen#s,55  % %g .  * *7 36rrNrrn )rrrrrrZr rrrrr r #s 4rr cURnURRRS5SnURnSUSUSU3$)z8Build fully-qualified kernel name for an out-variant op.z::rEz torch.ops.r )r r rrr )r nsrF rs rr r #sP   Bnn!!''-a0G##H t1WIQxj 11rc^\rSrSr%SrS\S'S\S'SU4SjjrSSjr\S S S .SS jj5r S r U=r $)r i#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_nodescl>[TU]"U0UD6 X l/Ul[ U5UlXlgr)rrr r r rq rv )rr r rrrs rrExternKernelMultiOut.__init__#s8 $)&)*(*%"?"G!rcHURU5 URU5 gr)r generate_extern_kernel_multi_outr s rr ExternKernelMultiOut.codegen$s W%006rNFr c XSSKJn Jn Jn U "UR5(dg[ U[ [45(dgU "U5n U cgU "U 5n[SU55(dg[U5[U5:wagU"[US9UUUUUUU US9 n/n[U5Hunn[URUR/URQ/UR!5QS9n[#UU[%U5U4/S9n[&R((dU (d[+U5(d3[,R.R0R3UR45 UR7U5 M UUlUUl[ U[ 5(a [ U5$[U5$) zGCreate an ExternKernelMultiOut if the op has a matching .out() variant.r)r r r Nc3V# UHn[U[R5v M! g7fr)rrr[ rN s rr2ExternKernelMultiOut.try_create..'$sG1:a..s')r%)rrz r r r )rUr& r)r r r r r rrrrrr$ rrWr r rGr r rrFr rprsr rrrrr r )r1r r r r r r rrz r r r r r r r r r tensor_outrU multi_outs rr ExternKernelMultiOut.try_create$s  fnn--.5$-88' >)&1 GGGG ~ #m"4 4 V ,    /'  02&~6MAz !(( &&(z''(-**,- F .~.23I 77&(44))--inn= NN9 %%7(+2'  ne , ,> !G}r)rv r r rq ) rr r r r rrr rrrn )r r r r r rr rm r r0 r r rrm rz r r rrz&Sequence[AllocatingMultiOutput] | None) rrrrrrrr rlr rrrs@rr r #s 99 " "& "! "  "  "7HL$)J%JJ J & J ' J+J&JEJ"J 0JJrr c \rSrSr%SrS\S'S.SjrS/SjrS0SjrS1Sjr S2S jr S3S jr S4S5S jjr S6S jr S7SjrS8SjrS.SjrS7SjrS9S:SjjrS;SjrSSjrS?SjrS@SjrSASjrSBSjrSCSjrS.SjrS.SjrSDSjrSESjrS3S jr SES!jr!SBS"jr"\#"S5S9SFS#jj5r$SGS$jr%SHS%jr&S4SIS&jjr'\(SJS'j5r)SKS(jr*SJS)jr+SCS*jr,\(SLS+j5r-S3S,jr.\.r/S-r0g )MriW$z; TensorBox / StorageBox allow in-place mutation of Tensors rrTc6URR5$rrrs rr(!MutableBox.has_exceeded_max_reads_$rrc6URR5$rrrs rrMutableBox.get_deviceb$rrc6URR5$rrrs rr-MutableBox.make_loadere$yy$$&&rc6URR5$r)rTr1rs rr1MutableBox.make_indexerh$yy%%''rc6URR5$r)rTr:rs rr:MutableBox.get_stridek$rrc6URR5$rrrs rrzMutableBox.get_namen$rrNc8URRU5$r)rTrCrAs rrCMutableBox.has_large_inner_fnq$syy++I66rc8URRU5$rrrFs rrHMutableBox.mark_reuset$rrc6URR5$rrrs rrLMutableBox.realize_hintw$rrc6URR5$r)rTrwrs rrwMutableBox.unwrap_viewz$rrc6URR5$r)rTr>rs rr>MutableBox.is_input_buffer}$yy((**rc6URR5$r)rTrQrs rrQMutableBox.freeze_layout$syy&&((rc8URRX5$r)rTrVrTs rrV*MutableBox.freeze_layout_with_stride_order$syy88NNrc8URRU5$r)rTr[rZs rr[(MutableBox.freeze_layout_with_fill_order$syy66u==rc8URRU5$r)rTr_r^s rr_(MutableBox.freeze_layout_with_same_order$syy66v>>rc8URRX5$r)rTrdrbs rrd+MutableBox.freeze_layout_with_exact_strides$syy99-WWrc6URR5$r)rTrgrs rrgMutableBox.get_read_writes$rrc6URR5$rrrs rrMutableBox.get_reads$rrc6URR5$rrrs rroMutableBox.num_reads$rrc6URR5$rrrs rrrMutableBox.get_storage_numel$rrc6URR5$rrrs rryMutableBox.get_reduction_type$rrc6URR5$rrrs rr|MutableBox.get_reduction_size$rrc6URR5$rrrs rrMutableBox.is_extern$rrc6URR5$r)rTrrs rrMutableBox.is_no_op$rrc8URRU5$rr r#s rrMutableBox.constant_to_device$syy++F33rc6URR5$r)rTrrs rrMutableBox.get_mutation_names$rrc6URR5$r)rTrrs rrMutableBox.get_operation_name$rrc6URR5$r)rTrrs rr'MutableBox.get_inputs_that_alias_output$syy5577rc6URR5$rrrs rrMutableBox.realize$rrc8URRU5$rrrus rrvMutableBox.get_free_symbol_uses$syy--m<rQrVr[r_rdrgrrorrryr|rrrrrrrr^rvrrrrrUrrrr r!rkrrrrrrW$sb L2&'(&$7+('+);@O"O37O O >?HMX/X@DX X +%%-..%$4..8#L)$)=!= !=*= *+3++&+$ "Hrrcd\rSrSr\\SSj55r\\SSj55r\S Sj5rSrg) ri$cgrrrTs rrTensorBox.create$sFIrcgrrrTs rrrU$s+.rcX[U[5(aU$[[U55$r)rrrrxrTs rrrU$s% d1 2 2KD)**rrN)rTrrr)rTrrr)rTr)rrrrrrrrrrrrr$s@ II ..++rcr\rSrSrSrSSjrSSjrSSjrSSjrSSjr SSjr SS jr SS jr SS jr S rg )rxi$z/ StorageBox allow in-place mutation of Tensors c[UR[[45(a5URR 5[ R R;$gr')rrTrrZrzrsr r=rs rr>StorageBox.is_input_buffer%s= dii+!? @ @99%%'177+?+?? ?rc[UR[5=(a5 URR5[R R ;$r)rrTrrzrsr r rs rrStorageBox.is_module_buffer%s9 tyy> 3 : ""$(9(99 rc  [RUR5(aURR5$[ UR[ [ [[45(d[UR55eURR5nURR5nURR5nUce[S[UURR5URR!5SS9URS9Ul["R$R'UR5URl["R$R+UR5 UR,URlXRlX RlURRR2URlURR($)NF)r r rrYr)rrrTrzrrurrrgrrrrrrrrrsr r rr rrrr)rrrr s rrStorageBox.realize%sl  " "499 - -99%%' '$))iD$%GHH $ IIK  Hii//1 II++- %%'!!!"!ii))+YY'')     00;  ""499- LL  + ' #yy~~88 yy~~rc[UR[[45(a:URR 5R S:aUR 5 ggg)z< Called on buffers we expect to be forced to realize later. rEN)rrTrurrDnontrivial_read_countrrs rrLStorageBox.realize_hint,%sI tyy9i"8 9 9 **,BBQF LLNG :rc<SSKJn UR5Vs/sH1nU"U5(aM[RR U5PM3 nnU(dg[ U5n[U5n[U5nXQ:=(a XV- S:=(a Xg:H$s snf)Nr)is_nonfreeable_buffersFr) r rcrrsr get_dep_size_hintrrMr)rrBrcr size_of_reads total_sizemax_sizemin_sizes r$has_accumulated_enough_reads_by_size/StorageBox.has_accumulated_enough_reads_by_size6%s@~~' ')#. +AGG % %c *'  ' }%}%  # %%* %$  s B#Bc2[UR[5=(aw UR5[R :=(dO UR 5=(d8 [RSL=(a UR[R5$r) rrTrurorFrealize_acc_reads_thresholdrC realize_acc_reads_size_thresholdrirs rr(!StorageBox.has_exceeded_max_readsI%sq$))Y/  NN vAA A &&( 77tC==;; rcr^US:a[UR[[45(a[ UR5(a9URR 5mSS/n[ U4SjU55(agUR5[R:=(d UR5$g)zR A heuristic to decide if we should realize a tensor that is used multiple times. rEexpsigmoidc3@># UHoTR;v M g7fr)used_ops)rropcounts rr5StorageBox.should_realize_on_reuse.._%s@iG,,,isTF) rrTrurr1rDrrorFrealize_reads_thresholdrC)rrG heavy_opsrts @rshould_realize_on_reuse"StorageBox.should_realize_on_reuseU%s 19DII 9/EFFdii  ))446"I. @i@@@ 6#A#AA-**, rcRURU5(aUR5 ggr)rxrrFs rrHStorageBox.mark_reuseg%s!  ' ' . . LLN /rc6URR5$rrrs rroStorageBox.num_readsk%rrrTNrrr)rBrrr)rGrrrrr)rrrrrr>rrrLrir(rxrHrorrrrrxrx$s4  < &  $%rrxc8\rSrSr%S\S'S\S'SrS\S'S rg) Subgraphio%rrr graph_moduleNzGraphLowering | Noner r)rrrrrr rrrrrro%s I&&"&E &rrcUVs/sH*n[U[5(aUR5OUPM, nn[[ SU555[U5:$s snf)Nc38# UHn[U5v M g7fr)r\ )rrs rr'_has_aliased_buffers..|%s;7"V**7r)rrZrwrr>)buffersrs r_has_aliased_buffersrv%sdF!+6? C CO  z;7;; [TU]SUUS9 Xl[RR U5Ul[RRU5 gr| )rrr rsr r rr )rr rrUrs rrInvokeSubgraph.__init__%sO   ! GG++D1  ""4(rcBUR(a UR/$/$r)r rs rrInvokeSubgraph.get_subgraphs%s"&-- 7R7rc8^SSKJn [RRnSnUR R S5=n(atSnUR[RRRLa8URS[RRRLdeSnUSUSnOSnUR[RRRLa8URS[RRRLdeSnURUSnUV s/sHoR S PM nn UV s/sHoRU 5PM nn /n [U5HRup[!U ["[$[&45(aU R)U 5 M8U R)U"XU 55 MT U nURcz[RR+UR,UUR.S 9Ul[R0"UR5 URR2"U6 SSS5 URR4n SnUH*n [!U ["5(aMU R75n O Uce[9UU[;US 9S 9mSU4S jjn[U 5VVs/sHunnU"UU5PM nnnUTlU$s sn fs sn f!,(df  N=fs snnf)z|For each operand, get a realized input, force it to have the same strides as the subgraph inputs, then use an InvokeSubgraphrE)constrain_to_fake_tensorNeager_input_valsrrrrrYrkr r r%)r rrUc Z>[U[[45(aU$UR5nUce[ [ UUR 5UR5UR5UR5RUR5RS9T[U4/SS9$)NrTr ) rrrurrR rWrrr:rrXrYr)rcindr invoke_subgraphs r create_output,InvokeSubgraph.create..create_output%s:N"OPP **,)))"%$..0#__.%002%00299"("3"3"5"?"? $C[M6:  r)rcrrrrz:ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput)r rrsr rR rirr rrqr r rrrS rrrr r rr rrr r graph_outputsrrr$ r )r1r rrrR fake_operandsrrX fx_operandsr new_operandsroperandr r rrrcoutsrs @rrInvokeSubgraph.create%s 7ww++  +00445GH H  HF""eii&<&<&I&II#((+uyy/E/E/U/UUUU,Q/8MF""eii&<&<&I&II#((+uyy/E/E/U/UUUU'++FG4K4?@KqVVE]KM@ AI!I1"3"3A"6!I%' %h/LC/ARS##G,##,WC6HI 0  >> !WW22((,&mm3HN $$X^^4""M25....Gg'<== ++- !!!($F3    !$  G .;DG:LM:LYQ fa(:LM"& SA "J.54TNs8K;L<LL Lc&URU5 gr)codegen_invoke_subgraphr s rr InvokeSubgraph.codegen&r r)rr )r rrrm rUr$ rrr)r rrrrz@list[ShapeAsConstantBuffer | NoneAsConstantBuffer | MultiOutput]rn )rrrrrr rrr rrrlrr rrrs@rrr%s!%Ho$(,H%,'+G $+ )  ),< )FW )  )8ff,2f IffP..rrc^\rSrSr%SrSrS\S'SrS\S'SrS\S 'Sr S\S 'Sr S \S 'SU4S jjr SSjr \ SSj5r\SSj5rSSjrSSjrSrU=r$) Conditionali&a IR node representing torch.cond Attributes: predicate: A boolean scalar tensor determining which branch to execute. operands: Input tensors passed to both true and false subgraphs. true_subgraph: Subgraph executed when predicate is True. false_subgraph: Subgraph executed when predicate is False. outputs: MultiOutput nodes representing the conditional's outputs. Nr predicaterrr true_subgraphfalse_subgraphSequence[MultiOutput] | Noner c>XlX lX0lX@l[ U/UQ5upx[ T U]SUUUS9 UbX`l[RRU5Ul [RRU5 gN)rrUrqrn ) rrrr_split_by_sym_typerrrz rsr r rr ) rrrrrrUrz sym_argsr rs rrConditional.__init__&s# *, 2I3I3I J "   (%6 "GG++D1  ""4(rc/nUR(aURUR5 UR(aURUR5 U$r)rrrr subgraphss rrConditional.get_subgraphs4&sG      T// 0      T00 1rc\[U[5(aU$URR$r)rrrrT)rs r _maybe_exprConditional._maybe_expr<&s" a  Hvv{{rcB URU5nUVs/sHoPRU5PM nn[RRRSn[ U[ 5(d[U55e/nUHIn[ U[5(a URURS5 M8URU5 MK [RRRSn SSjn X#4Hn U RbM[RRU RUU RS9U l[R"U R5 U RR"U6 U "U RR U 5U RlSSS5 M URceURceURR n URR n SU 4SU 44H&up[#U 5(dM[%SUS U35e ['U 5['U 5:XdX45e[)[+X55HununnUR-5UR-5:Xd UUU45eUR/5UR/5:Xd UUU45eUR15R2UR15R2:XaMUUU45e [5S XA/-55n[7[RR8R:[RRRR=S S55nUcS 5e[?UUUU[AUS 9US9n[)[+U [RRRS55VVVVs/sHununn[C[EUR-5bUR-5OUUR/5URG5Vs/sHn[>RIU5PM snURK5Vs/sHn[>RIU5PM snUR15R2UR15RLS9U[NU4/5PM nnnnnUUl(SSK)J*n U"UR5u nnnU"UR5u nnn[WU5[WU5-n[YU5Vs/sH n[[UUR\UUU5PM" snUl/U$s snf!,(df  GM=fs snfs snfs snnnnfs snf)zNCreate a Sequence of IRNodes from a conditional statement (see .lowering.cond)r8rYc /n[X5Hhup4[U[5(aURU5 M-UR[R [ U5UR5SS95 Mj U$NFr)rrrrrr rr )r fake_tensorsretrcr8 s r_require_exact_strides2Conditional.create.._require_exact_strides^&smC #M @ f&;<<JJv&JJ$::%f-t{{}E; !AJrNrtrue_fnfalse_fnzVOutput aliasing is currently not supported in compiled torch.cond. The outputs of the z% subgraph of torch.cond are aliased: c3p# UH,n[U[5(aMUR5v M. g7fr)rrr)ros rr%Conditional.create..&s, +a!67 ALLNN+s66rz zcannot determine devicer%)rrrrrUrz rr)-check_input_alias_and_mutation_return_outputs)rrm rzSequence[torch.Tensor]rr)0rS rsr rR rrrrr<rrir rrr r rrr(rrrrrrrXrlr9r rrrr$ rR rWrrr rYrr torch._higher_order_ops.utilsrr>rr rUr )r1rrrrrrrfx_op fake_outputsrr true_outputs false_outputsrr rt_of_or rz conditionalrc merged_outputrrrtrue_mutated_inputsfalse_mutated_inputsmutated_operand_indicesrs rrConditional.createB&sP%%i0 2:;(Q%%a((; ! 4 4 9 9" = +x00C${2CC0$&  E%&&$$UZZ%67$$U+ ! ww++007  + 0  $!+H~~%!"!6!6,,#0"*--"7" ((8NN&& 64J 44l4HNN0 98,"}}(((~~)))}}22  44 (,7*m9TUMD#L11$**./TU\T]_V< C $66U8UU6&s<'GHMAzS>>#s~~'77 F!S# F7==?cmmo5 D3} D5>>#**cnn.>.E.EE T3PS} TEI  +   6 GG   & & GG % % ) )*=t D !<#<<!!!#$F3/  :/8L!''"6"6";";E"BC/) (/**FM' ((*6",,. **,@M@R@R@TU@T"+11"5@TU>K>R>R>T>T //3>T",,.55$//1;;   &/)  2&  :':N:N O *Aq%q :(:O:O P +Aq&#--@"AJ E # 56( 6 8C=//# L6( $ [U2.AU7(AV V  V 4VAV 'V7 V V c URU5 URUR5UR[ US055 gr8 )codegen_conditionalr9 rzr rr s rr Conditional.codegen&s9##D)88 MMOT\\749Lb+Q rc[USS5=n(aL[[RRR U5nUce[ UR55$[ 5$r8 r= r> s rr$Conditional.get_unbacked_symbol_defs&rA r)rrrrrrz )rrrrm rrrrrUr$ rz r rrr)rzint | torch.SymIntrzint | sympy.Expr) rrrrrrrrh rzlist[MultiOutput]rn r)rrrrrrrrrrr rrrrrlrr rrrrs@rrr&s  $I}#(,H%,%)M?)&*NO*,0G )0))#) ) ! ) " )E) )8 XXX X " X  XXt   rrc/n/nUHFn[U[5(aURUR5 M5URU5 MH X!4$r)rrrrT)r non_sym_argsrrN s rrr&sOLH c0 1 1 OOCHH %    $   !!rc ^\rSrSr%SrSrS\S'SrS\S'SrS\S'Sr S\S 'Sr S \S 'SU4S jjr SS jr \ SSj5r\ SSj5r\SSj5rSSjrSSjrSrU=r$) WhileLoopi&zSThe IR node for while_loop and while_loop_stack_output. It supports input mutation.Nrcarried_inputsadditional_inputsr cond_subgraph body_subgraphrr c>XlX lX0lX@l[ /UQUQ5up[ T U]SUU US9 UbX`lXpl[RRU5Ul [RRU5 gr)rrrrrrrrz stack_outputrsr r rr ) rrrrrrUrz rrr rs rrWhileLoop.__init__'s-!2** 2 1n 10 1!  "   (%6 "(GG++D1  ""4(rc/nUR(aURUR5 UR(aURUR5 U$r)rrrrs rrWhileLoop.get_subgraphs%'sG      T// 0      T// 0rc[U5(dU$UVs/sH*n[U[5(aUR5OUPM, nn[ 5n/n[ X5HeupV[ U5U;a&UR[RU55 M:UR[ U55 URU5 Mg U$s snfr) rrrZrwr>rr\ rrr r)rrunwrapped_buffers seen_buffersroriginal_inputunwrapped_buffers r_clone_aliased_inputsWhileLoop._clone_aliased_inputs0's#N33! !) (%/v$G$GF   V S(  )3 +-03N0V ,N"#|3 l55nEF  $4!56 n- 1W ! s1C c[U[5(aU$[U[[45(a [U5$[U[5(a[R U5$[ S[U535e)NzNYI unsupported output type: )rrrxrZrR rrl r)rs r_maybe_wrap_as_tensor_box#WhileLoop._maybe_wrap_as_tensor_boxJ'se c9 % %J j/: ; ;S> ! [ ) )##C( (!>tCykJK KrcSSKJn SSjn[RRR Sn[RRR Sn X-n U V s/sHoR SPM n n UV s/sHoR SPM n n U V s/sHoR SPM nn UV s/sHoRU 5PM nn [RU5nU"X5nUV s/sHoRU 5PM nn U"UU5nUU-nX4GHnURbM[U [5(d[U 55e[RRURU URS9Ul[R "UR5 URR""U 6 UULaZ[%URR&5[%U 5:XdeU"URR&U 5URlSSS5 GM UR(aUR(deURR&nURR&n[)U5(a[+S U35e[%U5S :XdU5eUSn[U[,5(dMUR/5[0R2:XdU5e[%UR555S:XdU5e[%U5S:dS 5eUSR75nUce[%U5[%U5:Xd UU45e[9[;UU55HununnSS jnU"UR55UR555 U"UR=5UR=55 UR75UR75:Xd UUUU45eUR/5UR/5:XaMUUU45e Uce[?[RR@RB[RRR RES S55n[UUUU[GUS9UUS9nURb=[URRH[0RJRL5(deU"URRHU 5Sn[OU5nUVs/sHnUUPM n n[QU 5n!/n"/Ul)/Ul*U(Ga[%U5S:XdS5e[9[RRR S5Hunn#[W[YU#RZU#R\U#R_5V$s/sHn$[`RcU$5PM sn$U#Re5V%s/sHn%[`RcU%5PM sn%S9U[fU4/5n&URRRiU&5 U"RiU&5 M GO[9U5GH unn#UU;adU[%U5:dS5e[kU!5n'URTRi[mU'RnU'U55 U"RiU'5 Mq[W[YU#R75U#R/5U#R55U#R=5U#Rq5RrS9U[fU4/5n&URRRiU&5 U"RiU&5 GM [;UU"5Hkun(n)U(Ru5[RRv;dM4[RRxR{U)Ru55 Mm U"$s sn fs sn fs sn fs sn fs sn f!,(df  GM =fs snfs sn$fs sn%f)zcreate the while_loop IR node. stack_output controls whether it stack each iterations' output, which is necessary for training. r)check_input_alias_and_mutationc V[U5[U5:Xde/n[X5H}up4[U[R5(aH[ R U5nUR[RXTR5SS95 MlURU5 M U$r) rrrrr[ rrrrr r ) tensor_boxesrrr2fknew_tbs rr0WhileLoop.create.._require_exact_stridesc's|$L(99 99Cl9b%,,//'@@DFJJ$::"IIKu; JJrN):*Jrr8rYNrzOutput aliasing is currently not supported in compiled torch.while_loop. The outputs of the body_fn subgraph of torch.while_loop are aliased: rEz9torch.while_loop is assumed to have at least one operand.c[U5[U5:Xde[X5H.up#[RRR X#5 M0 gr)rrrsr r r^) lhs_exprs rhs_exprslhsrhss r_guard_list_equals,WhileLoop.create.._guard_list_equals'sC9~Y777 #I 9HCGG$$11#;!:rrz r%)rrrrrUrz rrz-NYI: while_loop_stack_output input mutations.r zonly carries can be mutated.)r r rr rX)rrm rz'list[int | torch.SymInt | torch.Tensor]rr)rSequence[int | sympy.Expr]rrrr)>rrrsr rR rrirS rrrrrr rrr r rrrr(rrrrrrrrr:r9r rrr$ modulefx GraphModuler>rmr r rR rWr r rrrr rrrlr rUrrXrzr=r r)*r1cond_fnbody_fnrrrrrfx_carried_inputsfx_additional_inputs fx_all_inputsrfake_all_inputsfake_carried_inputsfake_additional_inputscarried_inputs_additional_inputs_ all_inputsr cond_outputs body_outputsrr rreborrz while_loop mutated_idxsmutated_idx_setrr mutated_inputs_iter all_outputsrcrrr mutated_inputrsrs* rrWhileLoop.createU's Q * A  :GG0055b9 ww3388<)@ 2?@-Q66%=-@6GH6Gvve}6GH9M!N9MA&&-9M!N9GHA,,Q/H#99/J0V>#?#?@C/E  8N$NN88/8498+4}}..}}22 }}22  - - XXdWeg  < A%3|3% O!233;;=EJJ. 1 1.qzz|$) ,1 ,):" G "A))+!!!?#s<'88  ;  8%S,%GHKAxB <5 <5 < < r{{}bkkm < r}}  @==?bmmo5 J2r67J J5<<>R\\^3 @aR[ @3I"!!!5 GG   & & GG % % ) )*=t D  *0!!$F3/%  }}(Z MM %(("6"6. .   6 MM /  %\25DE_c*S/_E#>2$&  &( # '1, ? , ))=)=)B)B5)IJ V'%}}$llDJKKMRMbk55b9MRFLmmoVo 7 7 ;oV  C[M  "")))4""9- K )6 V/)^!44T6TT4$()<$=M//66&}';';]JW &&}5 +##)#4#4#6"("2"2"4!'!2#)#4#4#6#)#4#4#6#=#= #  !I&&--i8&&y1- 70NK8HC||~!5!55 ++// ?9UAH!NHO98tF"SVs=(`#`($`-`2?`7 s rr"WhileLoop.get_unbacked_symbol_defs5(rA r)rrrrrrrz )rrm rrm rrrrrUr$ rz r rrrrr)rrm rrm )rrrr) rrrrrrm rrm rrrrp rn r)rrrrrrrrrrr rrrrrrlrr rrrrs@rrr&s%].2N+215.5%)M?)%)M?),0G )0)(),) )  ) " )E)) )>2LLWWW) W , W  W #WWr   rrcr^\rSrSrSSS.S U4SjjjjrS U4SjjrS SjrSrU=r$) ri@(Nr c >[T U]UUUUUSUS9 SSKJn U"U5n U ceXl[ R RRU S5Ul U[ R RU 'g)Nr r) _get_effect) rrtorch._higher_order_ops.effectsr effect_typersr  effectful_opsrprev_effect_buffer) rrUr r r. r rrz rrrs rrEffectfulKernel.__init__A(s      /   @!&) &&&&"#''"7"7";";K"N-1k*rc>[TU]5nURbGURR [ R "URR555 U$r)rrgrrkrrGrrz)rrrs rrgEffectfulKernel.get_read_writes^(sUg-/  " " .    ! !$$T%<%<%E%E%GH rcgrrrs rr EffectfulKernel.has_side_effectsh(rr)rrrr rr) rrrrrrgr rrrs@rrr@(s|)-2HL222& 2 & 2 + 2&2E2 22:rrc@\rSrSr\"S5SSSjj5rSrg)r il(c[5$rr=rus rrv!NonTensorObj.get_free_symbol_usesm(rTrrNrr)rrrrr^rvrrrrr r l(s,N+$)! !,rr c`\rSrSr%S\S'S\S'S SjrSSSjjrSS jrSS jrSS jr S r g)r3 it(rr%FakeScriptObject | torch.ScriptObjectrcUR$rrrs rrzTorchBindObject.get_namey(r>rNcUR$rrrs rr!TorchBindObject.codegen_reference|(r>rcUR$rrrs rr TorchBindObject.get_value(rrc[UR[R5(a UR$URR$r)rrrrP real_objrs r get_real_objTorchBindObject.get_real_obj(s3 djj%"4"4 5 5:: ::&& &rcUR5n[U5(ag[US5(de[UR 55n[ R "U5SnUVs/sHEn[U[R5(dM$UR5UR5-PMG nn[R"[RUS5$s snf)Nr__obj_flatten__)r0r*rrr3r r rrr[ r numelr"rEoperatorr)rreal_script_obj flat_dict flat_elemsr flat_sizess r get_buf_bytesTorchBindObject.get_buf_bytes(s++- ? + +(9::::88:; ((3A6   !U\\* )ANN qwwy (   j!<<  s +#C%Crrrr)rr')rztorch.ScriptObjectr) rrrrrrzrr r0r:rrrrr3 r3 t(s& I 00' =rr3 c<\rSrSr%SrS\S'S SjrS S SjjrSrg) 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. r rc,[UR5$r)rQ rrs rrz OpaqueValueTypeConstant.get_name(sDJJrNc[UR5up#UR5H"upEU[RR U'M$ U$r)r)rrrsr opaque_value_type_classes)rrobj_repr opaque_typesr9rs rr)OpaqueValueTypeConstant.codegen_reference(sA!4TZZ!@ &&(DA34AGG - -a 0)rrrrr rrrrrrrzrrrrrr=r=(s J rr=cB\rSrSr%S\S'S\S'S SjrS S SjjrS rg) r i(rrrr cUR$rrrs rrzGeneratorState.get_name(r>rNcUR$rrrs rr GeneratorState.codegen_reference(r>rrrrr)rrrrrrzrrrrrr r (s I rr cF\rSrSr%SrS\S'S\S'S SjrS S S jjrS rg)r i(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. rrr rcUR$rrrs rrzOpaqueObjectState.get_name(r>rNcUR$rrrs rr#OpaqueObjectState.codegen_reference(r>rrrrrrErrrr r (s" I Jrr c\rSrSrS SjrS SjrS S Sjjr\S Sj5r\S Sj5r Sr g)_CollectiveKerneli(cgr'rrs rrZ!_CollectiveKernel.should_allocate(r*rcgrrrs rr "_CollectiveKernel.has_side_effects(rrNcn[UR5[RRLdS5eURnUbXlOUR RUlUR RVs/sH!o3R(dMURPM# snUl gs snf)Nz,Setting cpp kernel needs a valid op_overload) rrv rr r rr r rr r rt )rrr r rs rr} %_CollectiveKernel.set_cpp_kernel_name(sD$$%)>)>> : >!!   #)>>#6#6D #NN44. 4q FAFF4. *. s B2B2c [RR UR"X/UQ70UD6unnnnn SSS5 W (a USU 35eWH@n U R 5 [RR U R 55 MB USR5n U"[U S9UUWW5n [R"U5n U RRU Vs/sHn[[U S9X5PM sn5 U RRU Vs/sHoR 5PM sn5 SU;a]U RR[[U S9USU 55 U RRUSR 55 gg!,(df  GN~=fs snfs snf)Nr!rr%r)rsr r9r rrrzrrr tree_leavesr rr r r)r1r rqrr_example_outputr r r rz tensor_argr r inpsrrss rcreate_inplace _CollectiveKernel.create_inplace(sWW  ""6CDCFC  ! %E2C1D&EE$%J    GG ' ' (;(;(= >&Q**, f %      !!&)&&OS Tt^Jf5s Ct T !!T"BTc<<>T"BC F?  # # * *z8&-P     % %fUm&<&<&> ? ; 0 U#CsF/.G+G/ F>c B[RR UR"X/UQ70UD6unnnnn SSS5 W (a USU 35eWH*n [ U [ 5(aMU R 5 M, [ W[5(aURXe5n U ceU"[U S9UUWW5n [U5V Vs/sH(up[URU5U [U 4/5PM* snn U l [U RU5H_up[R (d[#U5(aM,[RR$R'UR(5 Ma U R$U"URU5UUWW5n [R (d[#U5(d3[RR$R'U R(5 U /U l U $!,(df  GN=fs snn f)Nrr%)rsr r9r rr3 rrrN r$ rrR r^ r rrFr rprrr)r1r rqrrr r r r rz r[r r rr[rs rcreate_out_of_place%_CollectiveKernel.create_out_of_place,)sWW  ""6CDCFC  ! %F3D2E&FF$%Jj/::""$& nd + +__[AF% %%!0 F"+>!:  ";IA ((0AYK "; FN #6>>>B ::BSCCGG--11#((; C >> !$$^4 F66>O??))--fkk:$XFNMe .sH /H H)rr rt rrro ) r rrqzIRNode | list[IRNode]rr rr rr) r rrqzTensorBox | list[TensorBox]rr rr rz%list[MultiOutput] | _CollectiveKernel) rrrrrZr r} rlr]r`rrrrrQrQ(s  ()@)@&)@ )@  )@  )@)@B99,9 9  9 / 99rrQcb^\rSrSrSSS.SU4SjjjjrS SjrSrU=r$) _AllReduce_Kernelii)Nr c N>[TU]UUUUUSUS9 URS5 g)Nr +aoti_torch_cpu__c10d_functional_all_reduce_rrr} r s rr_AllReduce_Kernel.__init__j)=      /     !NOrcURS5 URU5 [UR[5(aUR U5 ggNz+torch/csrc/inductor/aoti_torch/c/shim_cpu.hinclude_extra_headerr rrUrrL r s rr _AllReduce_Kernel.codegen)C$$%RS,,T2 dkk6 * *  % %g . +rrrr rn r rs@rrcrci)s)-PHLPPP& P & P + P&PEP PP,//rrccb^\rSrSrSSS.SU4SjjjjrS SjrSrU=r$) _AllReduceKerneli)Nr c N>[TU]UUUUUSUS9 URS5 g)Nr *aoti_torch_cpu__c10d_functional_all_reducerfr s rr_AllReduceKernel.__init__)s=      /     !MNrcURS5 URU5 [UR[5(aUR U5 ggrjrkr s rr _AllReduceKernel.codegen)rnrrrr rn r rs@rrprp)s)-OHLOOO& O & O + O&OEO OO,//rrpc^\rSrSrS SS.S U4SjjjjrS SjrS Sjr\SSj5rSU4Sjjr S r U=r $) _WaitKerneli)Nr c N>[TU]UUUUUSUS9 URS5 g)Nr +aoti_torch_cpu__c10d_functional_wait_tensorrfr s rr_WaitKernel.__init__)rhrcURS5 URU5 [UR[5(aUR U5 ggrjrkr s rr _WaitKernel.codegen)rnrcURSn[U[5(de[U[5(a7URSn[U[5(d[ U55eU/$[U[ 5(aGURSn[U[5(a!UR SupEURU/$/$/$r)rqrrrQrrR r)rrsrcollrrs rget_volatile_reads_WaitKernel.get_volatile_reads)skk!n#v&&&& c, - - 1 Aa(( 1$q' 1(3J [ ) )::a=D$ 122Q C())IIrcv[RR URX5unnnnnSSS5 W(a USU35eU"[ UR 5S9UWWW5nUR R[[ UR 5S9X(55 g!,(df  N}=f)Nr!r%) rsr r9r rrr rr ) r1r rsrZr r r rz r s r create_wait_WaitKernel.create_wait)s WW  ""6/  ! %E2C1D&EE$ cnn. /       && :S^^-=> L ! s B** B8c>[TU]5nUR5nUH@nURR [ R "UR555 MB U$r)rrgrrkrrGrrz)rrvolatile_readsvrrs rrg_WaitKernel.get_read_writes)sSg-/ 002 B    ! !,"6"6r{{}"E F!rrrr rn r )r rrsrrrr) rrrrrr rrlrrgrrrs@rrwrw)s)-PHLPPP& P & P + P&PEP PP,/2  *rrwcV[U[[45(a [U5$[U[[ 45(a5[ [R"5nUHnU[U5-nM U$[U[R5(a [U5$[ 5$r) rr;rr4rrr>rr!re rr[ rr3rs rre re *s!h%&&$Q'' At} % % u|| $ &A ,Q/ /A Au|| $ $$Q''|rcV[U[[45(a [U5$[U[[ 45(a5[ [R"5nUHnU[U5-nM U$[U[R5(a [U5$[ 5$r) rr;rr3rrr>rr!rf rr[ rs rrf rf *s!h%&&A At} % % u|| $ &A #A& &A Au|| $ $A|rcD[U[5(Ga [UR[5(Ga[URR[5(a'URRR SU5 g[URR[ 5(GanURRR SU5 [URR[5(ad[URRR[5(a1URRRR SU5 g[URR[5(aURRR(dk[URRRS[ 5(a4URRRSR SU5 ggggggg)Nrr) rrrTrxrrryrrR rrq)rr9s rassign_origin_noder *sq&)$$FKK)L)L fkk&& . . KK   / / q A  ((& 1 1 KK   / / q A&++**N;;    %%uAA   %%88J6;;++[99 ((00fkk..55a8&AAKK$$++A.AA-QRSB1:2*M$r)rrrzTypeIs[int | Integer])rrrr)rrrr)rrrz&Callable[[Sequence[_T]], Sequence[_T]])rz&Callable[[Sequence[_U]], Sequence[_V]]rz&Callable[[Sequence[_T]], Sequence[_U]]rrr)rz#Sequence[int | torch.SymInt | Expr]rzShapeEnv | Nonerr)rSequence[int | Integer]rrr)rrrrrr)rrrrrr)rrrrrztorch.Tensor | None)rzSequence[_T] | NonerzSequence[_T | None] | None)rz/IRNode | OutputSpec | torch.device | None | strrr)rz"IRNode | torch.device | None | strrr)rzBuffer | TensorBoxr6rrr)rErrFrrGrrr)r[rr\zSequence[int | torch.SymInt]rr)rkr rr)rqrm rr)rzExpr | Sequence[Expr]r rrrrr)rrr rrrrr)rrr rp rXrrrr )TFNFN)rrrrrrrrrUrrcrrztuple[StorageBox, Layout])rrrrrrr)r rrGrrr)r rrr)rr# rzTypeIs[Sequence[IRNode]])r r rr)rrm rr)rr rz-tuple[list[ShapeAsConstantBuffer], list[Any]])rrrrj)rr r9z torch.fx.Noderr(\ __future__rrrrr"rFloggingr5textwraprcollections.abcrrrrrr r enumr r typingr rrrrrrrrrrtyping_extensionsrrrrrr unittest.mockrrrr r!torch._export.serde.schema_exportserder/ rf rW r* rr" torch._loggingrtorch.fxtorch.utils._pytree_pytreer torch._dynamo.utilsr"torch._export.serde.serializer#*torch._higher_order_ops.auto_functionalizer$torch._inductorr%r r'"torch._library.fake_class_registryr(torch._library.opaque_objectr)r*torch._prims_commonr+r,r-r.r/r0%torch.fx.experimental.symbolic_shapesr1r2r3r4r5r6r7r8r9r:r; torch.fx.noder<torch.utils._ordered_setr>torch.utils._python_dispatchr?torch.utils._sympy.functionsr@rArBrCtorch.utils._sympy.symbolrDrrFrGcodegen.commonrHrIrJrKrLrMrNrOrPrQrR loop_bodyrS ops_handlerrTrUrVrWruntime.benchmarkingrX runtime.hintsrYrZr[r\r]r^r_r`rarbrcrdrerfrgrhrirjrkrlrmrnrorp virtualizedrqrrrsrtru torch.typesrvcodegen.cutlass.templaterwcodegen.wrapperrxr ryrzrrr* __version__r[rZ ImportErrorr{r|r}r~rrrrrr r r r getLoggerrr'rr rrrrrI rrrrrrrrrrrrrr$r/r1rBrKrarmrtr{rrrrsrurrrrrTrrrrrr\rgrrrVrrrvrr rr+rrZrrrrrr r r#rrWrrrrrrryr]rrirrurrr rrq rrPrimitiveInfoTyper r r r r r r r! rrA rG rrs r~ rd r r r r r r rJ ra rr r r r r r r r rrr r] r r r$ rR rr r r rrrxrrrrrrrr r3 r=r r rQrcrprwre rf rrrrrsM" MM:    UT ''22,,$$(?M#2?M    /?QQ*" NN-:0*)B&'95$%"(OY'''NJ t_ T] T] T]Dj) EkD()(**//%**2P2PP iP !   8??4 8 yy~~'T k     c4[();(JKdR   i ) d#$$ $F44,!$LP  , 9H   LP  , 9H   U U 27 +/  :?  26  D  >6 > >;('7*    $$G $G )$G $GN';S,S,l UUUUpF FF F R&% % % P i  F | $ y ! y !   u  u =)<8JN<N<N +<NBF<N<N~d  d  d  R$(1:   & &"8D>8D>"BH"LM YM7S97St#1#L` +` FF 5F F T   [5[[| !3748 9 9 991 9  9 2 99x 4 $^ v^ ^ B[[[|-(--`A9(A9A9H!(!!H[;[[|ShSSl&&&RlAlA^6"K|KK$ S| S S '9 (< 7 7 ` Z` ` F C&C[HV[H|!Gf!GHT%{%D.V*V*rUEV]EEPU &fi & & & K [ &6( F   Ut4_t4t4n  d#$]+_]+@<><~%K$&,tC#I4E4L/MMLL^"|" \B.\B~5N50((6N>Z%NZ%z5 .55 UR?RRj K9K\UE<EEPU"l""J / ((" ("V V B=L=@% % P - $< \< ~WlWt ))|))Z/,/d<8B5B$8-826)l6)r))|))X9T9Tx,E,E^21|21j-L-8<;<<;~U t &t nU .  @  2 I ,I X   F 4K 4 2f>fV TTTn + +n%n%bU'v'' LUB.\B.B.JUe ,e e P "  "2 "U@ @ @ F )n)X6"=l"="=Jl(\ "YYx/)/>/(/>S#Sr   TTiLNJs&l66 mm