ȅi SSKJr SSKrSSKrSSKrSSKrSSKrSSKrSSKrSSK r SSK r SSK r SSK J r SSKJrJrJr SSKJr SSKJrJrJrJrJrJr SSKrSSKJr SSKrSSKrSSKJ s J!r" SSK#J$r$ SS K%J&r&J'r' SS K(J)r) SS K*J+r+ SS K,J-r-J.r.J/r/ SS K0J1r1J2r2J3r3 SSK4J5r5J6r6J7r7J8r8 SSK9J:r: SSK;Jr>J r SSK?J@r@ SSKAJBrBJCrCJDrDJErE SSKFJGrG SSKHJIrI SSKJJKrK SSKLJMrM SSKNJOrOJPrPJQrQJRrRJSrS SSKTJUrUJVrV SSKWJXrXJYrYJZrZJ[r[ SSK\J]r] SSK J^r^J_r_J`r`JaraJbrbJcrcJdrdJereJfrfJgrgJhrhJiriJjrjJkrk SSKlJmrnJoroJprpJqrq SSKrJsrs S S!KtJuru S S"KvJwrwJxrxJyryJzrzJ{r{J|r|J}r}J~r~JrJrJrJrJrJrJrJr S S#KJrJrJrJrJrJrJr S S$KJrJrJrJrJr S S%KJr \(a,SS&KJr SS'KJr SS(KJr SS)K=Jr S S*KvJr S S+KJr \"S,5r\GRD"\5r\GRJGRM\S-5r\GRJGRM\S.5r\GRJGRM\S/5r\@"5r?S0rS1rSdS2jr"S3S45r\"S5SeS5j5r\"S5SeS6j5r"S7S85r\GRb"S9S:55r\GRb"S;S<55r\GRb"S=S>\55r\GRb"S?S@\55rSfSAjr"SBSC\5r\"5GRprSgSDjrSgSEjrShSFjrSgSGjrSiSHjrSjSIjr"SJSK\{5rSkSLjrSlSmSMjjr"SNSO\5r\GRSP5 "SQSR\5r"SSST5r\GRb"SUSV55r"SWSX5r\GRb"SYSZ55r"S[S\\z\\\\\\4445r\GRb"S]S^55r"S_S`\\5r"SaSb\5rSnScjrg)o) annotationsN)abstractmethod)CallableIterableSequence) lru_cache)AnycastOptional TYPE_CHECKINGTypeVarUnion) PRECEDENCE)get_interface_for_device)identitypreserve_rng_state)is_integer_dtype) OrderedSet)CeilDivFloorDivModularIndexing)get_triton_versionhas_triton_packagehas_triton_stable_tma_api)free_symbol_is_type prefix_strsymbol_is_typeSymT) ValueRanges)configirmetricsutils) AsyncCompile) code_hashget_path PyCodeCache write_atomic)'set_kernel_post_grad_provenance_tracing)DefaultHandler)triton_heuristics) benchmarker) AutotuneHintDeviceProperties ReductionHintTRITON_MAX_BLOCKTRITON_MAX_RSPLIT)get_max_y_gridnext_power_of_2)BaseSchedulerNodeFusedSchedulerNode Scheduler SchedulerNode)get_broadcasted_shape) cache_on_selfDelayReplaceLineget_bounds_index_exprget_fused_kernel_nameget_kernel_metadatais_welford_reduction Placeholderprefix_is_reduction sympy_dot sympy_product sympy_subs triton_typetriton_version_uses_attrs_dictupcast_compute_type)_ops ReductionType StoreModeV)"get_kernel_category_by_source_code)BlockPatternMatcher)ArgNameBackendFeature ConstexprArgCSE CSEVariable DeferredLineIndentedBufferInplacedBufferis_buffer_removed OpOverrides PythonPrinter RemovedArgSizeArg TensorArg WorkspaceArgWorkspaceZeroMode) constant_reprIterationRangesIterationRangesEntryIterationRangesRootPartialAccumulate SIMDKernelSIMDScheduling) config_ofequal_1_arg_indicesnon_constexpr_signatureshould_unwrap_unspec_argsignature_to_meta)SymbolicCallArg) ModuleTypeDtypePropagationOpsHandler)ShapeEnv)IRNode)BlockShapeType)SIMDKernelFeatures_T perf_hintsschedulefusioncLUS;nU(aSOSnUS;aUSUS3$USU3$)N)anymaxminprodtriton_helperstl)r{r|.2)reduction_type use_helpermodules X/home/james-whalen/.local/lib/python3.13/site-packages/torch/_inductor/codegen/triton.pyget_triton_reduction_functionrsE#@@J!+ F'>*!,,>*++c[U[R5(dg[U[R5=(d+ UR=(a [ UR 5S:H$)z" Is this expression a Sympy Integer or is it an integer sympy Expr containing no free symbols. The latter case can happen with Identity expr. Fr) isinstancesympyExprInteger is_integerlen free_symbolsexprs ris_sympy_integer_likersL dEJJ ' ' dEMM *  7C 1 12a7rcJ\rSrSr%Sr0rS\S'0rS\S'\S Sj5r Sr g ) OpDtypeSupportz Some Triton ops such as libdevice and tl.math only support float32 and float64. This class records which dtypes are supported by specific IR ops. z"dict[str, OrderedSet[torch.dtype]]supported_dtypeszdict[str, bool]convert_outputscURn[[R[R/5UR U'X R U'gN)__name__rtorchfloat32float64rr)clsfuncconvert_outputop_names rregister_upcastOpDtypeSupport.register_upcasts;--(2EMM5==3Q(RW%'5G$rrN)rzCallable[..., str]rboolreturnNone) r __module__ __qualname____firstlineno____doc__r__annotations__r classmethodr__static_attributes__rrrrrs1 <>8=')O_)66rrcx[5(dgSSKn[URRS5(agg)zX import AttrsDescriptor if the triton version is new enough to have this class defined. rNAttrsDescriptorz4from triton.compiler.compiler import AttrsDescriptor)rtriton.compiler.compilerhasattrcompiler)tritons rgen_attr_descriptor_importrs3   #v''):;;Ercj[5nURS5 SSKnURS5 [ 5=n(aUR U5 URS5 [ RR(aURS5 UR5$![a Nf=f)NzD import triton import triton.language as tl rzM import triton.language.extra.tlx as tlx # noqa: F401 a from torch._inductor.runtime import triton_helpers, triton_heuristics from torch._inductor.runtime.triton_helpers import libdevice, math as tl_math from torch._inductor.runtime.hints import AutotuneHint, ReductionHint, TileHint, DeviceProperties z import triton.profiler as proton import triton.profiler.language as pl pl.enable_semantic('triton') ) rVsplicetriton.language.extra.tlx ImportErrorr writeliner"rproton_profilinggetvalue)importsr attr_descs rgen_common_triton_importsrsG NN   (  /00y0)$ NN }}%%      +    sB%% B21B2c \rSrSrSr\"\R\R/5r \"\R\R\R/\ Q5r \ VVVVs0sH$nU[R"[ US3SSS9_M& snnnnr\ VVVVs0sH2nU[R"[ UR%5S3SSS9_M4 snnnnr\S Sj5r\SS j5r\SS j5rS rg s snnnnfs snnnnf) TritonSymbolszM Stores sympy.Symbol instances and constants associated with triton codegen. offsetTinteger nonnegativeBLOCKrpositivec "SnURnUGHn[U[R5(a?[R R RURnURnGO}[U[R[R[R[R[R[R45(aSnGOUR Vs/sHn[XG5(dMUPM nn[#U5S:XdSUR35eUSn [R R%5n S/U -n [R R'5V s/sHn [(U U R*:XdMU PM n n [#U 5S:XdS5e[-UR/U S55XSR0'[3U 5n[5X&5nGM UceU$s snfs sn f)NrrNAmbiguous type: r1z# of Match expected to 1)rrrTMPrLkernelcse varname_mapnameshape UNBACKED_INTSIZEPRECOMPUTED_SIZEINDEXFLOATUNBACKED_FLOAT block_typesrtriton_tensor_ndimactive_range_treesrprefixstrget_block_size tensor_dimtupler:)rr expr_shape expr_varsvarcse_var var_shapesymtsymbol_matchessymndimrtree tree_matchs rget_block_shapeTritonSymbols.get_block_shapes&( %% Cc488,,((,,22388<#MM %%II))JJJJ''    &)__"%4Ts8QD_">*a/N3CCHH:1NN/$Q'xx224 !" ; ; = =!#$++5 = :!+G-GG+25c6H6HTU6W2Xm../!%L /zEJKN%%%1"s0HH<H H c4URUR$r) block_sizesrrrs rrTritonSymbols.get_block_size2styy))rc4URUR$r) block_offsetsrrs rget_block_offsetTritonSymbols.get_block_offset6s  ++rrN)r sympy.Exprrrr)rrarz sympy.Symbol)rrrrrrrR0_INDEXR1_INDEXreduction_typesXBLOCKYBLOCKZBLOCKrrSymbolrrupperrrrrrr).0rrrs0000rrrs!$--!?@Odkk4;; VoVWK D ellj./v6RVWWM  D ell$%%'( .t   K44l**,,Q s +C1 9C9 rc\rSrSr%S\S'S\S'S\S'S\S 'S \S 'S \S 'SSjrSSjrSSjrSSjrSSjr \ SSj5r Sr g)IndexingOptionsi;r index_strOrderedSet[str] mask_varsz Optional[str] expand_strr _has_rindexrindexz#Optional[Sequence[Union[int, str]]] expand_shapec,[UR5$r)rrselfs rhas_maskIndexingOptions.has_maskDsDNN##rcJ[UR[R5$r)rrrrr s r has_indirectIndexingOptions.has_indirectGs"4::txx88rcUR$r)rr s r has_rindexIndexingOptions.has_rindexJsrc:[SUR55$)Nc3V# UHn[U5RS5v M! g7f)tmpNr startswithrmasks r .IndexingOptions.has_tmpmask..Ns"J>43t9''..>')rzrr s r has_tmpmaskIndexingOptions.has_tmpmaskMsJ4>>JJJrc:[SUR55$)Nc3V# UHn[U5RS5v M! g7f)rNrrs rr,IndexingOptions.has_rmask..Qs"H3t9'',,rrr s r has_rmaskIndexingOptions.has_rmaskPsHHHHrcUR(a2SR[[[UR555$S$)N & r)rjoinsortedmaprr s rmask_strIndexingOptions.mask_strSs4 =ANNEJJvc#t~~67 8 PV rrNrrrr) rrrrrr rrr r&propertyr-rrrrrr;sTN 55$9 KI  rrc\rSrSr%SrS\S'S\S'S\S'S \S 'S \S 'S \S'S \S'S\S'SrS\S'SrS\S'\S)Sj5r \S)Sj5r \S)Sj5r \S)Sj5r \ SS.S*Sjj5rS+SjrS,SjrS-S jrS.S!jrS/S"jrS/S#jrS/S$jrS/S%jrS/S&jrS0S'jrS(rg)1BlockDescriptorOptionsi]z This is a base class that describes a block descriptor used in Triton kernels. It can be used to create either a tensor descriptor (with TensorDescriptorOptions) or a block pointer (with BlockPtrOptions). BlockParametersparamsrconstant_offset list[int]orderrrSequence[sympy.Expr]broadcast_shape list[bool]broadcasting_dims final_shapeBlockParameters.StrideSorter stride_sorterNzOptional[list[int]]_boundary_checkFrcan_liftc.URR$r)r5rr s rrBlockDescriptorOptions.shapeus{{   rc.URR$r)r5 block_shaper s rrE"BlockDescriptorOptions.block_shapeys{{&&&rc.URR$r)r5stridesr s rrHBlockDescriptorOptions.strides}{{"""rc.URR$r)r5offsetsr s rrLBlockDescriptorOptions.offsetsrJr)rAc ,^[RRmS U4SjjnU"UR5UlU"UR5UlUR V s/sHn TR U S5PM n n [U 5(aSU S'URU 5nURU[RRS9upURV s/sHn TR U S5PM n n UR nURU 5nUVs/sHn[RU5PM nn[RR(a&USRS:XdeUR!S5 [RR"n[RR$(d['UR5['[RR(5U- :XaN[RR*R-5(a!U[.R0R2/U-- n[4R6"[RRUR5nU"U[RRRAU5UUUUU U US9 nURCXS5 U$s sn fs sn fs snf![8a3 [;[=[?['UR5555nNf=f) z2Helper to create a BlockDescriptorOptions instancecR>UVs/sHnTRU5PM sn$s snfr)lookup_precomputed_size)exprsrsizevarss r lookup_size2BlockDescriptorOptions.create..lookup_sizes&GLMutH44T:uM MMs$rNF)stride_sorter_cls shape_envrx) r5r6r8rr=r:r<r?rA)rQzIterable[sympy.Expr]rlist[sympy.Expr])"rLgraphrRrrHrEstatically_known_equalsall remove_dimsmaybe_sort_with_stride_order _shape_envrrrno_x_dimrpopnum_reduction_dimsinside_reductionrnumelsfeatures is_reductionrSOner% argsort_symAssertionErrorlistreversedrangerPcompute_boundary_check)rr5r6 range_treesr get_max_blockrVrArSdimsingleton_dimsr?strider<r:rr=reduction_ndimr8resultrRs @rcreateBlockDescriptorOptions.creates77## N#6<<0 $V^^4AG@R@R @RH , ,S! 4@R   ~  !&N2 ##N3!' C C/177;M;M!D! GMnn FTFH , ,VQ 7n  !,,##$56GRRkd}33D9k R 88  q>((C/ // OOA 44))FNN#s188??';n'LL!!..00 EGGKK=>9 9K ?%%agg&8&8&..IE GG,,DD_U#+/'   %%mA C $ S$ ?%FNN(;"<=>E ?s$ KK K9K:LLcB[RUn[XU05$z> Replaces instances of {symt}_offset with the new expression. rrrEr r replacementrroffsets rreplace_offset%BlockDescriptorOptions.replace_offset$  --d3$+ 677rc[RH*nURU[R"S5U5nM, U$Nrrrr~rrr rrs rremove_roffsets&BlockDescriptorOptions.remove_roffsets4!11D&&tU]]1-=tDD2 rc &[RRnUVs0sH8n[RUR U"[ UR 5_M: nn[[[RRU55n[[UR55Vs/sGHKnURURU[ R"R$5(aMBU(a<[R[&R(UR*UR,;dpUR/URUUR*U5(aMUR/URU[1UR*UU55(aM[RR2(a5UR*U[R[&R4:XaGMIUPGMN snUlgs snfs snf)z6List of indices to pass to tl.load(boundary_check=...)N)rLrZrRrrrrrzr,rneeds_yz_grid_overflowrmrrr[rHrrgZerorrrErstatically_known_multiple_ofrEr`rr@)r rprorRt block_to_maxneeds_overflow_grididxs rrn-BlockDescriptorOptions.compute_boundary_checks 77## !/    % %aff -}Z=O/P P  / "#ahh&E&E{"STS_- -44T\\#5F U  ,)55dkkB++C0==>%AA JJsOT-=-=c-B !) E E JJsO&t'7'7'.-s' ( m&C&C D D(r)rzrEr s rr!BlockDescriptorOptions.has_rindex,s" ((   rc"UR5$r)rr s rr& BlockDescriptorOptions.has_rmask2s  rcgrrr s rr "BlockDescriptorOptions.has_tmpmask5rrc4[UR55$r)rrr s rr BlockDescriptorOptions.has_mask8sD'')**rc^URnURnU(aJURRUR5nURRUR5n[ Xg5VV s/sH(upU (a[ R ROUPM* n nn [XU 5nURR(d:U(d3[U 5[U5:XaURRU 5n [RRmU=(a< [U 5[U5:H=(a [U4Sj[ X555n [UR5(a-U (d&SUS[R R#U5S3nURn URR(dtU(aURR$OURR&n SUSU S3nU(a URO$URRUR5n [XU5nU$s sn nf)a> Generate a broadcast and a reshape for the block descriptor. This restores stride-0 dimensions which were removed from the block descriptor. Transposes are also applied to the input using self.stride_sorter: if for_store is True: - First Broadcast the value. Since self.broadcast_shape is stored in descending stride order, it must be reverted to the original order since the input value does not have dims with descending strides - After, transpose the broadcasted value so that dimensions are in descending stride order - Finally reshape to the block shape else (for load): - First broadcast the value to self.broadcast_shape (strides are descending) - Then transpose the value so that dimensions no longer have descending strides - Finally reshape the block to the final kernel tile shape c3|># UH1upTRUS5=(d TRX5v M3 g7frNN)r[)rpre_dimpost_dimrRs rrGBlockDescriptorOptions.codegen_broadcast_and_reshape..ysE*O%G00!<G33GFG)Ns9<tl.broadcast_to(, ) tl.trans()r:r<r?revertziprrgrhtriton_reshape is_identityrrLrZrRr\rzr index_to_strsort_idxrevert_sort_idx)r value initial_shaper=allow_implicit for_storer:r<rqis_broadcastingpre_broadcast_shapesupports_implicit_broadcast old_shape permute_dimsrRs @rcodegen_broadcast_and_reshape4BlockDescriptorOptions.codegen_broadcast_and_reshape;s2.. 22 "00778L8LMO $ 2 2 9 9$:P:P Q  ),O(O (O$+EGGKK 3(O  u5HI""..'(C ,<< #'"4"4";";[[ [[z88-78?C8 8 2 +2 /2   2 h$ !+aa,a* a  a  a arr3c"\rSrSrSSSjjrSrg)TensorDescriptorOptionsic.[RRnURS:waUSU"UR5S3OUSU"UR53SU"UR 53SU"UR 53/nSSRU5S3$) z Codegen a call to tl.make_tensor_descriptor() Args: name: variable name for pointer roffset: unused, but kept for compatibility with BlockPtrOptions.format() Returns: "tl.make_tensor_descriptor(...)" r + (rshape=strides= block_shape=ztl.make_tensor_descriptor(r)rLrrr6rrHrEr*)r rr}fargss rformatTensorDescriptorOptions.formats HH ! !''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0  ,DIIdO+>rrNTrrrr)rrrrrrrrrrrs ??rrcP\rSrSrSSjrS SjrS S SjjrS SjrSrg) BlockPtrOptionsicB[RUn[XU05$ryrzr{s rr~BlockPtrOptions.replace_offsetrrc[RH*nURU[R"S5U5nM, U$rrrs rrBlockPtrOptions.remove_roffsetsrrc [RRn/URQnU(d UVs/sHoPR U5PM nnUR S:waUSU"UR 5S3OUSU"UR 53SU"UR53SU"UR53SU"UR53SU"U53/nS S RU5S3$s snf) z Codegen a call to tl.make_block_ptr() Args: name: variable name for pointer roffset: should rn_offset be included in offsets=..., for use with tl.advance() Returns: "tl.make_block_ptr(...)" rrrrrrzorder=zoffsets=ztl.make_block_ptr(r) rLrrrLrr6rrHrEr8r*)r rr}rrLrrs rrBlockPtrOptions.formats HH ! !!DLL/BIJ'++F3'GJ''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0Qtzz]O $qzl #  $DIIdO#4A66KsC*c [RUnURVs/sHAnURX2U5URU[R R U5- PMC nnU$s snf)aF Codegen string to pass to tl.advance(name, ...). Advance is the difference between offsets in each loop iteration. To compute it, we replace rN_offset with multiples of RN_BLOCK. Since we expect rN_offset to vary in range(0, rN_numel, RN_BLOCK), the first iteration has rN_offset=0, while the second has rN_offset=RN_BLOCK. )rrrLr~rrgr)r rrblockradvances radvance_roffsetBlockPtrOptions.advance_roffsetsw**40 ,,  '##FD9%%feggllDAB'   sAA.rNrrrr)rrrr) rrrrr~rrrrrrrrrs688-78?C8 8 7:rrc[U[5(a[U[5(deUVs/sH"n[RR U5PM$ nnUVs/sH"n[RR U5PM$ nnXE:XaU$UVs/sH ofS:wdM UPM snU:waSUSSR U5S3$Sn/nUHKn U[ U5:a XU:XaURS5 US- nM2U S:XdeURS 5 MM U[ U5:XdeUS SR U5S 3$s snfs snfs snf) zUG3tyy'?&@CC C F ]# #c0B(B MM#  1HC3; ; MM& !  #m$ $$ $WAdii'( **%JI.s)E )E  E#Ec@\rSrSrS"SjrS"SjrS"SjrS"SjrS"SjrS"Sjr S"Sjr S"S jr S"S jr S"S jr S"S jrS"S jrS"SjrS"SjrS#SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"SjrS"Sjr S"Sjr!S r"g!)$ TritonPrinteric[UR5S:XdeSURURS5S[RR S3$)NrNlibdevice.trunc(r).to(rrr_printrLr index_dtyper rs r_print_TruncToIntTritonPrinter._print_TruncToIntM499~"""t{{499Q<89qxx?S?S>TTU V rcUR(a[[U55nU$[R"5(a$[ R R(aUnU$SUS3nU$)Nz tl.full([], z , tl.float64))rrintr" is_fbcoderversionhip)r rrets r _print_FloatTritonPrinter._print_Float#s_ ??c$i.C      EMM$5$5FC !m4C rc[UR5S:XdeURURS[SS- 5nUS3$)NrNrAtom?z.to(tl.float64))rr parenthesizer)r rrs r_print_ToFloatTritonPrinter._print_ToFloat/sI499~"""   diilJv,>,D EO$$rcURup#UR(a8UR(a'URURS[SS- 5$UR U5nUR U5nSUSUS3$)N % rrz!triton_helpers.remainder_integer(rr)ris_nonnegative stringifyrrr rquotdivquot_sdiv_ss r_print_PythonModTritonPrinter._print_PythonMod5spII    3#5#5>>$))UJv4F4LM MT" C 26("UG1EErc,UR(deURup#UR(a8UR(a'URURS[SS- 5$UR U5nUR U5nSUSUS3$)N // rrz!triton_helpers.div_floor_integer(z, r)rrrrrrrs r_print_FloorDivTritonPrinter._print_FloorDiv=s|II    3#5#5>>$))VZ5G#5MN NT" C 26(#eWAFFrcPURURS[SS- 5$)N / rr)rrrrs r_print_IntTrueDivTritonPrinter._print_IntTrueDivHs#~~dii 60BS0HIIrc[UR5S:XdeSURURS5S[RR S3$NrNlibdevice.floor(rrrrrs r _print_floorTritonPrinter._print_floorMrrc[UR5S:XdeSURURS5S[RR S3$rrrs r_print_FloorToIntTritonPrinter._print_FloorToIntSrrc[UR5S:XdeSURURS5S[RR S3$NrNlibdevice.ceil(rrrrrs r_print_ceilingTritonPrinter._print_ceilingYK499~""" TYYq\!: ;5AUAU@VVWXXrc[UR5S:XdeSURURS5S[RR S3$rrrs r_print_CeilToIntTritonPrinter._print_CeilToInt]r#rc,SURU5S3$)Nz tl.sqrt_rn(().to(tl.float32)))rrs r _helper_sqrtTritonPrinter._helper_sqrtasdkk$/00ABBrcSURURS5SURURS5S3$)Nlibdevice.pow(rrrNr)rrrs r_print_FloatPowTritonPrinter._print_FloatPowds?T[[167r$++diiPQl:S9TTU V rc6URSR(a;S[URS5SURURS5S3$SURURS5SURURS5S3$)Nrr,rrNr)r is_Integerfloatrrs r_print_PowByNatural!TritonPrinter._print_PowByNaturalis 99Q< " "#E$))A,$7#84;;tyyQR|;T:UUVW WT[[167r$++diiPQl:S9TTU V rcURURS5nURURS5nURURS5nSUSUSUS3$)NrrNr! tl.where(rr)doprintr)r rcpqs r _print_WhereTritonPrinter._print_Whereps_ LL1 & LL1 & LL1 &1#Rs"QCq))rc[UR5S:XaURURS5$[UR5S-n[U5nURU"URSU65nURU"URUS65n[ SXV455upVUS;d SUS35eS US US US US US US US US3$)z1 Helper for max/min code generation. cmp: > or < rNrr!Nc3.# UH nSUS3v M g7f)(rNrrrXs rr6TritonPrinter._print_min_max_helper..s.v!q1Xvs)>z * ( z= z) + )))rrrtyper)r rcmpmidrabs r_print_min_max_helper#TritonPrinter._print_min_max_helpervs tyy>Q ;;tyy|, ,$))n!4j KKTYYt_- . KKTYYst_- ..v..j C$>Q  :5#:#::rc[U[5(dgURn[U[R5(a [ U5$S$r)rrTrrr)rrs rlow_precision_fp_varrs< c; ' ' IIE&0 &D&D E "O%OrcJ^\rSrSrSSU4SjjjrSrSrU=r$)TritonCSEVariableich>[TU]XX4S9 [5UlUcS5eUcS5eg)Nrz!TritonCSEVariable must have dtypez!TritonCSEVariable must have shape)super__init__rr)r rboundsrr __class__s rrTritonCSEVariable.__init__sD u:*4, E"EE  E"EE rcpUHn[U[5(a'URRUR5 M?[U[R 5(dM`[ RH<n[XE5(dMURR[US3/5 M M g)Nr) rrrupdaterrrrrr)r rrkwargsargrs rupdate_on_args TritonCSEVariable.update_on_argssC#011%%cmm4C..*55D%c00--*T2B1C4/H.IJ6r)rr) rrrzValueRanges[Any]r torch.dtyperrrrr)rrrrrrr __classcell__rs@rrrsS !% F F! F F  F  F F  rrcSSKJn U"5$)Nrrn)!torch._inductor.dtype_propagationrorns rget_dtype_handlerrsL % ''rc:^^^SSjmSU4SjjmSUUU4SjjnU$)z Codegen helper to upcast arguments to float32, depending on the config and dtype. This decorates tl.math/libdevice codegen functions. c[RR(+=(aD [U[5=(a- UR [ R[ R4;$r) r"rcodegen_upcast_to_fp32rrTrrfloat16bfloat16)rs r needs_upcast*maybe_upcast_float32..needs_upcastsD 44 4 =3 , = emmU^^<< rc2>T"U5(aSOSnUU3$)N.to(tl.float32)rr)r upcast_stringrs rmaybe_upcast_arg.maybe_upcast_float32..maybe_upcast_arg%s$-9#->->)B }o&&rcL>^[RTT5 SUUUU4SjjnU$)Nc>UVs/sH nT"U5PM nnUR5VVs0sHupEUT"U5_M nnnT "U0UD6nT =(a7 [U4Sj[R"XR 5555nU(dSO#[ [ 5T R5"U0UD6n U [RS4;n U (aU bS[U 5S3OSn UU 3$s snfs snnf)Nc34># UH nT"U5v M g7frr)rrrs rrKmaybe_upcast_float32..decorator..wrapped..4s6-Sc S!!-S.to(rr) itemsrz itertoolschainvaluesgetattrrrrrrF)rrr upcast_argskeyval upcast_kwargsruany_needs_upcast result_dtypeneeds_downcastdowncast_stringrrrrs rwrapped8maybe_upcast_float32..decorator..wrapped-s<@ADS+C0DKAHN WHCS"23"77MW;8-8F- #6-6__T==?-S63  (.0$--@$Q&Q  *%--1FFN"l&>{<013  Xo./ /'BWs C0C5r0)rr)rrrrrs` r decorator'maybe_upcast_float32..decorator)s$&&t^< 0 0.rr/r0)rCallable[..., Any]rrr)rrrrs` @@rmaybe_upcast_float32rs  ': rc\rSrSrSr\R "\R5r\ SRSSSjj5r \ STSj5r \ S5r \ S5r\ \"5S 55r\ S 5r\ S 5r\ \"5S 55r\ \"5S 55r\ \"5S55r\ \"5S55r\ S5r\ S5r\ S5r\ S5r\ S5r\ S\R:SSS.Sj5r\ \"5S55r\ \"5S55r \ S5r!\ S5r"\ \"5S55r#\ \"5S55r$\ \"5S55r%\ \"5S55r&\ \"5S 55r'\ \"5S!55r(\ \"5S"55r)\ \"5S#55r*\ \"5S$55r+\ \"5S%55r,\ \"5S&55r-\ \"5S'55r.\ \"5S(55r/\ \"5S)55r0\ \"5S*55r1\ \"5S+55r2\ \"5S,55r\ S-5r3\ \"5S.55r4\ S/5r5\ S05r6\ S15r7\ S25r8\ S35r9\ S45r:\ S55r;\ S65r<\ S75r=\ S85r>\ S95r?\ S:5r@\ S;5rA\ S<5rB\ \"5S=55rC\ \"5S>55rD\ \"5S?55rE\ \"5S@55rF\ \"5SA55rG\ SB5rH\ \"5SC55rI\ \"5SD55rJ\ \"5SE55rK\ \"SFSG9SH55rL\ \"SFSG9SI55rM\ \"5SJ55rN\ \"5SK55rO\ SL5rP\ SM5rQ\ \"5SN55rR\ SO5rS\ \"5SP55rTSQrUg)UTritonOverridesiIzEMap element-wise ops to Triton e.g., ops.to_dtype(x,...) -> x.to(...)NTcxSSjnUb=[U"X!5[RR5[RlU[R :XaSUS3$U[R :XaUbUR(dUcUS3$U(a [U5nO [U5nUSUS3$)Nc X:Xag[R[R4nX;aX;a X:waS5eU[R:XdU[R:XagU[R:XdU[R:Xagg)NrzCConversions between float8_e5m2 and float8_e4m3fn is not supported!rr!)r float8_e4m3fn float8_e5m2) src_dtype dst_dtype fp8_dtypess r_get_min_elements_per_thread>TritonOverrides.to_dtype.._get_min_elements_per_threadUs% ##!!J '+* UU  U E---e>O>O1OE///9@S@S3Srr>z != 0)z.to(tl.int16).to(tl.uint8)rr)rrrrrr) r{rLrmin_elem_per_threadrruint8rrr)rXrruse_compute_typesr out_dtypes rto_dtypeTritonOverrides.to_dtypeNs " /:   6  ,/,Y>,,,AHH ( EJJ qc= ekk !  !i&A&AYEV S23 3 +E2I)%0ID 1%%rcURUR:XdeURU:waUS[U5S3nUS[U5S3n[U5U:waUS[[U55S3nU$)Nrrz, bitcast=True))rrrFrH)rXrrouts rto_dtype_bitcast TritonOverrides.to_dtype_bitcasts!!U^^333 77i #T+i013A4 E*+?; u % .Ek*=e*DEFaHC rc [RRU5n[U"U55n[ U5nUS:XaU$US:a(UR (dSUSS3nSUSUSUSUS3 $SUSUSUS3$) Nz tl.float32rztl.rtl.full(rrr)r _prims_common dtype_to_typer`rr)rrrtype_ triton_valrFtriton_signed_types r_shaped_constant TritonOverrides._shaped_constants##11%8"5<0 )%0 , &  19U__#&{12&7!8 eWBzl"5G4Hk]Z[\ \eWBzl"[MC Crc"URX/S9$)Nr)r)rrrs rconstantTritonOverrides.constants##E#;;rcSUS3$)NrVrrrXs rabsTritonOverrides.absaS""rc[USS5n[USS5nU[R:Xa3U[R:Xa[R(a SUSUS3nO SUSUS3n[ U5(d[ U5(aN[ 5RX5nU[R[R4;aUS[U5S3nU$)Nrztriton.language.div_rn(rrr>rr) rrrr"emulate_divison_roundingrrtruedivrrF)rXyx_dtypey_dtyperrs rrTritonOverrides.truedivs!Wd+!Wd+ u}} $5==(//,A3b15CaSA3a.C  " "&:1&=&=)+33A9IU]]EMM::T+i"8!9; rcSUSUS3n[U5(d[U5(aN[5RX5nU[R[R 4;aUS[ U5S3nU$)Nr>rrr)rrmodrrrrF)rXrrrs rrTritonOverrides.modsp!Cs!n  " "&:1&=&=)+//5IU]]EMM::T+i"8!9; rcD[R(aSUS3$SUS3$)z When use_fast_math, use the ftz (flushing to zero) variant of exponent computation. Check https://github.com/triton-lang/triton/issues/5735 for more details. z tl_math.exp(rzlibdevice.exp()r" use_fast_mathrs rexpTritonOverrides.exps+   !!A& &#A3a( (rcSUS3$)Nzlibdevice.exp2(rrrs rexp2TritonOverrides.exp2!1%%rcSUS3$)Nzlibdevice.expm1(rrrs rexpm1TritonOverrides.expm1"!A&&rcSUS3$)Nz tl.sqrt_rn(rrrs rsqrtTritonOverrides.sqrtQCq!!rc[RRnUS:XagUS:Xa SUSUS3$US:XaUS3$Uc:[R"[R "S [ R5U5$[S U<35e) N compile_errorzcompile error! runtime_errorz"triton_helpers.device_assert_then(z == 0, "injected assert fail", raccuracyz + 1rz:unrecognized config triton.inject_relu_bug_TESTING_ONLY = ) r"rinject_relu_bug_TESTING_ONLYopsmaximumrrrrj)rXbugs rreluTritonOverrides.relusmm88 / !# O #8s:YZ[Y\\]^ ^ J S:  [;;s||Au{{;Q? ? LSGT rcd[RR(a SUSUS3$SUSUS3$)Nz tl.minimum(r, tl.PropagateNan.ALL)ztriton_helpers.minimum(rrrrrIrJs rminimumTritonOverrides.minimum : ==   2aS(>? ?,QCr!A6 6rcd[RR(a SUSUS3$SUSUS3$)Nz tl.maximum(rrztriton_helpers.maximum(rr r!s rrTritonOverrides.maximumr$rcSUSUSUS3$)Nr5rrr)rIrJr7s rwhereTritonOverrides.wheres1#Rs"QCq))rc ^^ ^ ^ [RR(deXp2SnSnU"U5(aU"U5nU"U5(aU"U5nS UU U U 4Sjjn[[RR 55S:dS5e[ [ R[R5m [ [ R[R5m [ [ R[R5m [ [ R[R5m[RRR[RRU"U[!UR"5T T/5UR$T T4S9n[RRR[RRU"U[!UR"5TT /5UR$TT 4S9n[&R(R*R,R.S:XaSnOSnS US US US 3$)at Triton code generation for lowering ops.dot to tl.dot. The logic is as follows: 1. Downcasting for performance If the data was previously upcasted to fp32, we downcast back to the original dtype (e.g., fp16 or bf16) for better performance. While surrounding operations may run in fp32, matmul itself is executed at the original precision to optimize throughput. 2. Handling non-constant reduction masks If the reduction mask is not constant and there was any operation between tl.load and tl.dot, we zero out regions outside the mask using tl.where(r0_mask, val, 0). This ensures that values outside the mask do not contribute to the dot product, preventing incorrect results. 3. Shape alignment for tl.dot We massage shapes to match the tl.dot requirement of (Y, R) x (R, X). Current codegen eagerly broadcasts tl.arange to create unique axes. We reshape, transpose, or broadcast to align with the (Y, R) x (R, X) shape. We avoid using 3D dot ((Z, Y, R) x (Z, R, X)) because 3D tl.dot has poor performance. During batched matmul (bmm), we keep ZBLOCK=1 and call the 2D dot kernel instead. c[[[UR55(dg[R R SnUR(de[R RU5(ag[R RRR5Hup#X0:XdM SU;dMSU;dM g g)NFrUztl.loadz other=0.0T) rzr,rBrrLrrorf_has_constant_maskr_cacher)rreduction_rangekvs ris_where_needed,TritonOverrides.dot..is_where_needed;ss. >??hh2226O"// //xx**?;;  ++1138 Q;!3C 4rc[RRSUR5n[R R Vs/sH%nUR(dMURS3PM' nn[U5S:XdS5e[RUSX5n[R RR[R RX@RURS9$s snf)NdotrrNz'don't tile reduction when native matmulrrr)r# Reduction default_valuerrLrrorfrrTritonKernelOverridesr(rgeneratecomputer)rdefaultrreduction_mask where_vars r where_cond'TritonOverrides.dot..where_condPsll00 BGHH000D$$%4;;-t$0  ~&!+ V-V V+-33N14EsTI88<<((  )99CII) s C9 C9c>T U;aUVs/sH o3T :XaSOUPM nnUTT/:XaOTU;dS5eSS/nTU;aTUS'TU;aTUS'[XU5nUTT/:gnU(a SUSTSTS3nU$UTT/:XaUTU;dS 5eSS/nTU;aTUS'TU;aTUS'[XU5nS US 3nUTT/:gnU(a SUSTSTS3nU$[es snf) aX Generate a reshape, transpose, and broadcast for the tl.dot. tl.dot requires specific shape requirement : (Y,R) x (R,X) but the current triton codegen eagerly broadcast the tl.arange so it needs to be reshaped to meet the requirement. This is done by three steps. 1. remove the empty dimension (dim with size 1) and make it 2d with tl.reshape 2. permute the dimension if needed (e.g., (X,R) -> (R,X)) with tl.trans 3. broadcast if needed with broadcast_to. - This shows up when matmul operand is broadcasted with torch.expand/repeat. - e.g., torch.rand((16,)).expand(16,16) @ B e.g., (Y,1,R), (Y,R) -> tl.reshape(var, (Y,R)) e.g., (1,X,R), (R,X) -> tl.trans(tl.reshape(var, (X,R))) e.g., (1,X,1), (R,X) -> tl.broadcast_to(tl.trans(tl.reshape(var, (X,1))), (R,X)) TODO : eventually we want to remove this function when lazy broadcasting arrives rz&left tl.dot operand cannot depend on xrrNrz, (rrEz'right tl.dot operand cannot depend on yrr)rNotImplementedError) rrr=rqshape_2dbroadcast_neededRBLOCKrrrs r#reshape_transpose_broadcast_for_dot@TritonOverrides.dot..reshape_transpose_broadcast_for_dotpsw6&JW X-3v 3!>- Xvv..]2<2 :]*"(HQK]*"(HQK'uXF$,/?#? #.ugS6("ME4L1 00]2=2 :]*"(HQK]*"(HQK'uXF$E7!,$,/?#? #.ugS6("MEL*)W!YsCrztl.dot can only do mm and bmmr5tf32ieeeztl.dot(rz, input_precision=""))rr9r=r9rr)rLris_native_matmulrdense_size_listrrrrrrrrrr9r:rkrrrbackendscudamatmulfp32_precision) rIrJorig_aorig_br1r>rEinput_precisionrDrrrs @@@@rr4TritonOverrides.dots8xx(((( * 6 6 " "1 A 6 " "1 AI /I .I  I I V188++-.!3T5TT3]..t{{;<]..t{{;<]..t{{;<]..t}}=> HHLL ! ! HH   /4=66BR S''6" "  HHLL ! ! HH   /4=66BR S''6" "  >>   % % 4 4 >$O$O2aS 3O3DBGGrrN) constraintsris_purepackc [U5nSRUVs/sHn[U5PM sn5nUc&SRS/UV s/sHn SPM sn -5nSUSUSUSUSUS US 3 $s snfs sn f) Nrz=rr$ztl.inline_asm_elementwise('z', 'z', [z ], dtype=z , is_pure=z, pack=r)rr*r) asmrTrrUrVinputsrFi input_refs_s rinline_asm_elementwise&TritonOverrides.inline_asm_elementwises*%0 YY71A78  ))TF6-B6ac6-B$BCK,SEk]$zlR[\g[hhrszr{|CDHCIIJK K 8-Bs A5 A: cSUS3$)Nz tl_math.cos(rrrs rcosTritonOverrides.cosrrcSUS3$)Nz tl_math.sin(rrrs rsinTritonOverrides.sinrrc[S5e)Nz/ops.index_expr not implemented outside a kernelrA)rrrs r index_exprTritonOverrides.index_exprs!"STTrc[S5e)Nz+ops.masked not implemented outside a kernelrf)rbodyothers rmaskedTritonOverrides.maskeds!"OPPrcSUS3$)Nzlibdevice.lgamma(rrrs rlgammaTritonOverrides.lgamma#1#Q''rcSUS3$)Nzlibdevice.erf(rrrs rerfTritonOverrides.erf s!$$rcSUS3$)Nzlibdevice.cosh(rrrs rcoshTritonOverrides.coshr rcSUS3$)Nzlibdevice.sinh(rrrs rsinhTritonOverrides.sinhr rcSUS3$)Nzlibdevice.acos(rrrs racosTritonOverrides.acosr rcSUS3$)Nzlibdevice.acosh(rrrs racoshTritonOverrides.acosh rrcSUS3$)Nzlibdevice.asin(rrrs rasinTritonOverrides.asinr rcSUS3$)Nzlibdevice.asinh(rrrs rasinhTritonOverrides.asinhrrcSUSUS3$)Nzlibdevice.atan2(rrrrXrs ratan2TritonOverrides.atan2"!Bqc++rcSUS3$)Nzlibdevice.atan(rrrs ratanTritonOverrides.atan r rcSUS3$)Nzlibdevice.atanh(rrrs ratanhTritonOverrides.atanh%rrcSUSUS3$)Nzlibdevice.copysign(rrrrs rcopysignTritonOverrides.copysign*s%QCr!A..rcSUS3$)Nzlibdevice.erfc(rrrs rerfcTritonOverrides.erfc/r rcSUS3$)Nzlibdevice.erfinv(rrrs rerfinvTritonOverrides.erfinv4rqrcSUSUS3$)Nzlibdevice.hypot(rrrrs rhypotTritonOverrides.hypot9rrcSUS3$)Nzlibdevice.log10(rrrs rlog10TritonOverrides.log10>rrcSUS3$)Nzlibdevice.log2(rrrs rlog2TritonOverrides.log2Cr rcSUSUS3$)Nzlibdevice.ldexp(rz.to(tl.int32))r)rXns rldexpTritonOverrides.ldexpHs!!Bqc88rcSUSUS3$)Nzlibdevice.nextafter(rrrrs r nextafterTritonOverrides.nextafterLs&aS1#Q//rcUSU3$Nr)rr!s r logical_andTritonOverrides.logical_andQCs|rc US3$)Nz == 0rrIs r logical_notTritonOverrides.logical_notUsE{rcUSU3$Nz | rr!s r logical_orTritonOverrides.logical_orYrrcSUSUS3$)Nr> ^ rrr!s r logical_xorTritonOverrides.logical_xor]s1#S1~rcUSU3$rrr!s r bitwise_andTritonOverrides.bitwise_andarrc SU3$)N~rrs r bitwise_notTritonOverrides.bitwise_notes 1#wrcUSU3$rrr!s r bitwise_orTritonOverrides.bitwise_orirrcUSU3$)Nrrr!s r bitwise_xorTritonOverrides.bitwise_xormrrcUSU3$)Nz << rr!s rbitwise_left_shift"TritonOverrides.bitwise_left_shiftqD}rcUSU3$)Nz >> rr!s rbitwise_right_shift#TritonOverrides.bitwise_right_shifturrc SUS3nSUSUS3$)Nr>).to(tl.uint32)ztl.rand(rrrseedrs rrandTritonOverrides.randys%VHO,$r&++rc SUS3nSUSUS3$)Nr>rz tl.randn(rrrrs rrandnTritonOverrides.randn~s%VHO,4&6(!,,rc ,SUS3nSUSUSUSUS3 $)Nr>rztriton_helpers.randint64(rrr)rrlowhighs r randint64TritonOverrides.randint64s1VHO,*4&6("SED6KKrc[S5e)Nz.ops.load_seed not implemented outside a kernelrf)rrs r load_seedTritonOverrides.load_seeds!"RSSrcX[RR(aSUS3$SUS3$)Nz tl.rsqrt(rzlibdevice.rsqrt(r rs rrsqrtTritonOverrides.rsqrts/ ==  qc# #%aS* *rcSUS3$)Nzlibdevice.log1p(rrrs rlog1pTritonOverrides.log1prrcSUS3$)Nzlibdevice.tan(rrrs rtanTritonOverrides.tanrurc[RRRR U5nU(a[ US5(a UR nOSn[R(aJ[RR(a+[5S:aU[R:wa UbSUS3$SUS3$)Nr)rzlibdevice.fast_tanhf(rzlibdevice.tanh()rLrrrgetrrr"rrrrrr)rXrrs rtanhTritonOverrides.tanhs((,,**..q1 ww00MMEE   !!"$v-&!+1#Q/ /$QCq) )rcSUS3$)Nz tl.sigmoid(rrrs rsigmoidTritonOverrides.sigmoidrrcSUSUSUS3$)Nz(libdevice.signbit(z ) != 0) if (z).dtype is tl.float32 else z < 0rrs rsignbitTritonOverrides.signbits#"!L3NqcQU V rcSUSUS3$)Nzlibdevice.fmod(rrrr!s rfmodTritonOverrides.fmods!2aS**rcSUSUS3$)Nr,rrrr!s rpowTritonOverrides.pows s"QCq))rcSUS3$)Nz tl_math.log(rrrs rlogTritonOverrides.logrrF)rcSUS3$)Nzlibdevice.isinf( ).to(tl.int1)rrs risinfTritonOverrides.isinf"!M22rcSUS3$)Nzlibdevice.isnan(rrrs risnanTritonOverrides.isnanrrcSUS3$)Nzlibdevice.nearbyint(rrrs rroundTritonOverrides.rounds&aS**rcSUS3$)Nrrrrs rfloorTritonOverrides.floorrrc HUSU3nUSU3nSUSUSUSUSUSUS 3 $) Nrrz tl.where((z < 0) != (z < 0), tl.where(z != 0, z - 1, ), rr)rIrJrrems rfloordivTritonOverrides.floordivsV D}3qclA3j+;C5vVTXSYY\]a\bbcddrcl[R"S[R5n[R"[R "X5[R 5n[R"[R "X5[R 5n[R"X#5nUSUS3$)Nrr.dtype))rrrrrltrsub)rXzleftrightrs rsignTritonOverrides.signso LLEKK (||SVVA\EJJ7 cffQlUZZ8ggd"d1#W%%rcSUS3$)Nrrrrs rtruncTritonOverrides.truncrrcUSU3$)Nrrr!s rtruncdivTritonOverrides.truncdivsD}rcSUS3$)Nr rrrs rceilTritonOverrides.ceilr rrNT)rrrOptional[torch.dtype])rrrr)Vrrrrrmathre_LOG_2_E staticmethodrrrrrrrrrrr r rrr"rr(r4rrr]r`rcrgrlrorsrwrzr}rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrrrr"rrrrrrIsOyy H,0 :&:&):&:&x DD"<<## , ) )&&''"""77 77 **xHxHt"&emmTPQKK####UUQQ((%%&&&&&&''&&'',,&&''//&&((,,''&&9900,,--LLTT++ ''%%**&""  ++**##/303/303++''ee&&'' &&rrrc^\rSrSrSrU4Sjr\\RS55r \S5r \S5r \ S5r \ S5r\ S 5r\ S S j5rS rU=r$) r8iaMap element-wise ops to Triton within a TritonKernel Unlike TritonOverrides, these assume the code is going to be inserted into the body of the main triton kernel and so it may use indexing and mask variables which are assumed to already be defined in the current scope. cF>[TU]"U0UD6 UR5 gr)rr_setup_libdevice_routing)r rrrs rrTritonKernelOverrides.__init__s# $)&) %%'rc^SSKJm [RRR Hn[ X5(de[X5nU4SjnUS:XaE[ TS5(de[R"X2US9nXl [X[U55 MqSn[R"XRUS9nXl [X[U55 M g)zUR[R:waU"U5$[TU5"U5R$r)rrrrr)rX_original_impl_fn_namer/s rdecomposition_routerLTritonKernelOverrides._setup_libdevice_routing..decomposition_router!s777emm+)!,,"#3X>qAGGGrr)r1r2c`UR[R:Xa SUSUS3$U"U5$)Nz libdevice.r>r)rrr)rXr1r2s r dtype_routerDTritonKernelOverrides._setup_libdevice_routing..dtype_router0s277emm+'z1#Q77)!,,rN) torch._inductor.codegen.commonr/r _inductorr%op_requires_libdevice_fp64rr functoolspartialrsetattrr))rfn_name original_implr3fnr6r/s @rr,.TritonKernelOverrides._setup_libdevice_routings D,,GGG3(( ((#C1M H )#/;;;;&&(QX& l2&67 - ""WB"K C,r"2 3;Hrcj[RR5nS/U-nURXUS9$)NrNr)rLrrr)rrrrrs rrTritonKernelOverrides.constant<s7 xx**,d ##E#>>rc[RRUSSS9n[U[5(deUR (a UR nO[ RUR5n[RR5nU[R[R4;aUOUn[RRnS[Rl[RR R#[RR$UR&[)U5UUS9nU[RlU[R[R4;ae[RR R#[RR$UR+Xr5[-U5UR.S9nOUnUR0Htn[3U[4R65(dM$[R8"U[RR R:UR<R>5nMv X%:wa[[RR R#[RR$UR+Xu5UUR.S9nUR@Ul U$!U[Rlf=f)NF block_ptrtma_compatibility_checkerrrrr5)!rLrindexingrrr rrrget_index_dtype_as_torch_dtyperrint64r" test_configsruntime_triton_dtype_assertrr9r:rr=rrHrrrrr promote_typesrrrr) rrrrIrrorigr index_vars rrg TritonKernelOverrides.index_exprEs,88$$ ET% (O4444  ))E!11(..AEhh==? u{{EKK&@@k"">> C>CF   ;((,,''  "",T2 (C?CF   ; ekk2 2((,,''   S()%0ii (C E!.. !)TXX66!//qxx||77 GMME/ #hhll++HH$$LL2%)) ,!**  =?CF   ;s A+KK.c <Ubx[RRb][RR R [RRUS3[RURS9nURRSS9nU(dS5eSnUHHnURH5nURS:wd[URS5(dM2S n MF MJ U(aSOUn[RRXS 9nU"5n SSS5 U(aW R R"(a [U5n[RR R [RRS U S [%U5S U S3[&R("U5U R*U RS9n[,R."WX5n OW n U R0R3W5 U $!,(df  N=f)N .to(tl.int1)r5output)opz)graph for body does not contain an outputFloadrNTrrz.shape, rrrH)rrrrLrrr9r:rrrZ find_nodesrtargetrj mask_loadsris_boolr`r wraprrr(rdiscard) rrjrknodes need_wherenoderrnew_maskrurs rrlTritonKernelOverrides.maskeds   1 1 =88<<((  & %jjjj )D %%%2AAAu Dyy::'+CCHHQK+P+P!%J! # XX  3xVF4 }}$$U HHLL))  6((=+?*@6('R"''.llll *E))Hf4CC h' )4 3s H  Hc[RRRU5nSUS[RRR SU5S3$)Ntl.load( + load_seed_offsetr)rLrrinput seed_offset)rrrs rrTritonKernelOverrides.load_seedsIhhmm!!$'se3qxx}}889KVTUUV W rc*SUS3n[RRRU5=n(aU$[RRR UR UR S9n[RRR [RUR S9n[RRRUSUSUS35 [RRRXU45 X44$)Nzfrexp(rr5rz = triton_helpers.frexp() rLrrtry_getnewvarrrrrr:rput)rX cache_keycse_valmantissaexponents rfrexpTritonKernelOverrides.frexpsQCqM hhll**95 57 5N88<<&&QWWAGG&D88<<&&U[[&H ""j8*$ir7original_stridesrF)initrc[UR5S:de[UR5[UR5:Xde[[ [UR555nURU:HUl[ UR5VVs0sHup#X2_M nnn[ [U55Vs/sHnXBPM snUlgs snnfs snfr)rrrrkrm _is_identity enumerater)r identity_sort_idxrZr/sorted_dims_by_strides_maps r __post_init__*BlockParameters.StrideSorter.__post_init__ st,,-1 11t}}%T-B-B)CC CC $U3t/D/D+E%F G  $ 1B BD ?$?A+-?$D *U$s CC!cUR$r)rr s rr(BlockParameters.StrideSorter.is_identitys$$ $rcg)zBCreate a `StrideSorter` that can be used to sort block parameters.NrrrrWs rrv#BlockParameters.StrideSorter.createsrcpUR(dURVs/sHo!UPM sn$U$s snfrrrr attrrZs rsort!BlockParameters.StrideSorter.sort#0##)-7AQ77K83cpUR(dURVs/sHo!UPM sn$U$s snfrrrs rr#BlockParameters.StrideSorter.revert(rr)rrNrzlist[Union[int, sympy.Expr]]rWrprr>)rrrrr dataclassesfieldrrr1rrrrvrrrrrr StrideSorterBlockParameters.StrideSorters##%0%6%6E%BB   %  %   U#? ULT U ) U   U    rrcH^\rSrSrU4Sjr\SSj5rSrU=r$)$BlockParameters.IdentityStrideSorteri-c">[TU]5 grrrr rs rr2BlockParameters.IdentityStrideSorter.__post_init__/ G ! #rc FU"U[[[U555S9$)Nrr)rkrmrrs rrv+BlockParameters.IdentityStrideSorter.create2s'!1eC(8$9:; rrr) rrrrrrrvrrrs@rIdentityStrideSorterr-s4 $  #? LT  )   rrcL^\rSrSrSrU4Sjr\SSj5rSrU=r $)+BlockParameters.TensorDecriptorStrideSorteri;zD Sorts BlockParameters dimensions with strides in descending order. c">[TU]5 grrrs rr9BlockParameters.TensorDecriptorStrideSorter.__post_init__Arrc[[[U555n[R"X!SS9nU"UUS9$![ a UnNf=f)aV If the strides are not all known constants or if the strides are already sorted in descending order, return identity sort. For example if block_shape @ strides is [ZBLOCK, XBLOCK, YBLOCK] @ [8, 1, 16] The indices to sort the strides in descending order will be [2, 0, 1]. The indices to revert back to the original order will be [1, 2, 0]. T)reverser)rkrmrr%rirj)rrrW identity_sortrs rrv2BlockParameters.TensorDecriptorStrideSorter.createDs^!s+;'M ) !,,YRVW !1!  " )( )s; A  A rr) rrrrrrrrvrrrs@rTensorDecriptorStrideSorterr;s9  $  #? LT  )   rrc [U5n[SX455up4U"S0UVs0sH oUX5XE-_M snD6$s snf)z Concatenates block parameters. c3N# UHn[R"U5v M g7fr)rasdictr?s rr*BlockParameters.__add__..fsBMq[''**Ms#%r)rFr)r rkrrIrJrs r__add__BlockParameters.__add__asL4jBTMBB9a8as16AF?*a8998sA c URURUS9n[S0[R"U5R 5VVs0sHupEXCR U5_M snnD6nXc4$s snnf)z Sort `BlockParameter` with stride_sorter_cls. Returns block parameters as well as a `StrideSorter` which contains information on how the sort can be reverted. )rWr)rvrHr4rrrr)r rVrWr?rrr5s rr^,BlockParameters.maybe_sort_with_stride_orderis}*000S   !, 2 24 8 > > @ @HC'',, @  $$ s A/c ^U4Sjn[S0[R"U5R5VVs0sH up4X2"U5_M snnD6$s snnf)z1 Remove dimensions where removable_dims is True. cd>[UT5VVs/sHupU(aMUPM snn$s snnfr)r)ititem is_removableremovable_dimss r filter_dims0BlockParameters.remove_dims..filter_dimss7+.b.*A*A&D#*A s ,,r)r4rrr)r rrrrs ` rr]BlockParameters.remove_dimszsS   5@5G5G5M5S5S5UV5UsK$$5UV  VsArN)rkr4rr4)rVztype[StrideSorter]rWrprz4tuple[BlockParameters, BlockParameters.StrideSorter])rr;rr4)rrrrrrrrkrrrErHrL dataclassrrrrr^r]rrrrr4r4s*//EE E$/$5$5d$KK!K + 1 1$ GG G + 1 1$ GG G&&&P |  #l##J:%!3%@H% =%" rr4c4\rSrSrSrSrS SjrSrSrSr g) "CooperativeReductionWorkspaceCacheiz The scratch space used for cooperative reductions can be reused after two reduction loops. This keeps track of what can be reused. cXl/Ul/Ul[R"[R 5UlSUlSUlgr) r current_loop prior_loop collections defaultdictdequeready_for_reuse loop_count store_count)r rs rr+CooperativeReductionWorkspaceCache.__init__s> *66{7H7HIrcURRU5nU(aUR5$URR US5up4nUR R XU45 X54$r)rrpopleftr workspacerr)r nbytescachedws_namer\ ws_offsets rallocate+CooperativeReductionWorkspaceCache.allocatese%%))&1 >># # $ 3 3FE BI   &9!=>##rcURH%upnURURX#45 M' URUl/UlU=RS- slgNrN)rrrrr)r rrrs r on_loop_end.CooperativeReductionWorkspaceCache.on_loop_endsT*.// &FY   ( / /0D E+:++ 1rcHURnU=RS- slU$r)r)r priors rincrement_store_count8CooperativeReductionWorkspaceCache.increment_store_counts#   A r)rrrrrrN)rr) rrrrrrrrrrrrrrrs $rrc,\rSrSr%S\S'SrSrSrg)FixedTritonConfigizdict[str, int]r"c URU$rr"r rs rrFixedTritonConfig.__getitem__s{{4  rcXR;$rrrs r __contains__FixedTritonConfig.__contains__s{{""rrN)rrrrrrrrrrrrrs !#rrc"\rSrSrSrSSjrSrg) TritonCSEizy Subclasses CSE to apply the current load mask to the cache key to avoid CSEing variables across separate masked blocks. cb[RR=n(a XR4$U$r)rLr _load_maskr)r rnrs r augment_keyTritonCSE.augment_keys*88&& &4 &yy) ) rrN)rnrrzUnion[str, tuple[str, str]])rrrrrrrrrrrrs  rrcr\rSrSr%SrS\S'S\S'S\S'S\S 'S rSS jrSS jrSS jr Sr g)TMACompatibilityCheckerizG Checks if the TMA API can be used for load / store triton operations. TritonKernelrrrrrforcecSUlg)Nz2Cannot use TMA descriptor for load / store since: failed_debug_prefixr s rr%TMACompatibilityChecker.__post_init__s #W rc,UR(ag[RR5RS:Xah[ R R5SS:aC[RR(a$[R(a[5(d![RSUR5 gUR (a<UR"R$(a![RSUR5 gg)NTrMr z}%s Requires triton>=3.4.0, a CUDA device with cc>=9.0 and `use_tensor_descriptor` and `assume_aligned_inputs` options enabledFz/%s stores with `no_x_dim` cannot load 16 bytes.)rrLrZget_current_device_or_throwrFrrMget_device_capabilityr"ruse_tensor_descriptorassume_aligned_inputsrrdebugr rrr`r s r can_use_tma#TMACompatibilityChecker.can_use_tmas :: GG / / 1 6 6& @ 00215: 33,,)++ II[((   >>dkk22 IIA(( rc  UR(aDURVs/sH,n[RRR U5PM. nnO URn[RRR US[R"S55(d"[RSURU5 gURRnUSSHn[RRR [XT-S[R"S55[R"S55(aMi[RSURUU5 g URSn[RRR U[R"S55(a,[RS URUR5 gSnSnUR H1n ["R$Hn ['X5(dMU nU n M/ M3 U(aU(dUS ["R$35eUR(R*(GaUR,(Gd[.Un Sn UR(R0H4n U R2(dMU R4U :XdM(U R6n O U ceUR(R9U 5nUR;X~05U-n[RRR=U[R"S55(d-[RS URURU5 ggSSS jjnXd-S- nUR?[@U5R?[U5n[C[E[RF"UUS555nUUR(RI[.U5:a#[RS URUU5 gUR(RKU5n UR(RL(aZUUR(RLU :a;[RSURU UR(RLU U5 gg[OUUR(RPRSU S55UR(RPU 'gs snf![Ta. [RSURUR5 gf=f)z Check if the block parameters are valid for TMA. If force, we allow relying on symbolic hints equivalent to what we check for Triton templates. rUrNz=%s TMA API requires innermost stride to be 1. Strides are: %sFNrzU%s TMA API requires outer strides to be 16 byte aligned. Dtype bytes: %d, strides: %sz>%s innermost block shape cannot load 16 bytes. Block shape: %sz, expr must contain a single block type from zj%s persistent reduction innermost block shape cannot load 16 bytes. Block shape: %s, persistent RBLOCK: %dc$X- nU(aX2-nU$rr)rXrrr s rindexing_div_repQTMACompatibilityChecker.are_block_parameters_compatible..indexing_div_repd s %C!gJrzC%s the minimum block size to satisfy expression %s is too large: %dzT%s For block %s, fixed config block size %d is smaller than the minimum required: %dz?%s innermost block shape cannot load 16 bytes. Block params: %sTr)rXrrrrzOptional[sympy.Expr]rr)+rrHrLrZrR symbolic_hintr[rrrrr rrrrErrrrrpersistent_reductionrrrorfrnumel_get_persistent_RBLOCKsubsstatically_known_geqreplacerr5rnsolve max_blockr fixed_configr{tma_min_block_sizesrr)r  block_paramsstrH element_sizersinnermost_block_shapeinnermost_block_typeinnermost_block_symtblock_type_str block_symtinnermost_tree_prefix tree_numelrpersistent_rblockinnermost_block_bytesr solve_exprsolve_expr_simplifiedmin_block_sizes rare_block_parameters_compatible7TMACompatibilityChecker.are_block_parameters_compatibles ::=I=Q=Q=Qr  ..r2=Q G#**Gww77 U]]STEUVV IIO((  zz** crlF77##;; 5q%--:KL a  k,,   #"!- 8 8 < 77   3 3 !5==#3   IIP((((  ##3@@N+77 !.==+9(+5( 8A$(< $%%QR_RkRkQl m < ;; + + +DNNN %//C$D !J[[,,>>>xx#88%&WW  - ) )) $ B B: N %**,@+TU "77##88%u}}R'8 A,, ,,%  fIB  /3!!, 3ABF (2(:(:.)'/+;<&"1 10""DKK$9$934%II]00-&  !!%!9!9:N!O;;++% (@(@(PP < 44* KK44^D*  %Q4GJ& 77;;NANGDKK33NCGv  U,, ,,   s&3TB2TB TAT5UUcUR$)a Can you lift the make_tensor_descriptor call to the top of the kernel? This requires being certain that all of the shape, stride, and block_shape information is handled in arguments or top level definitions. Right now we assume this is always possible if you force TMA. )rr s rrA TMACompatibilityChecker.can_lift szzrr Nr/)r%r4rr) rrrrrrrrr4rArrrrrrsS  O KX! !Fm%m m^ rrc^\rSrSr%Sr\rS\S'\r S\S'Sr \ r Sr S \S 'S\S]U4S jjjr\S^S j5r\S^Sj5rS_SjrS^SjrSrSrSrSrS^SjrSr\S`Sj5rSS SS SS.SaSjjrSbScSjjrSbSjrSdSjrSr Sr!Sr"S r#S S!.S"jr$S#r%SeS$jr&SfS%jr'SgS&jr(ShSiS'jjr)SjS(jr*S)r+SkS*jr,SlSmS+jjr-S`S,jr.SnS-jr/SoS.jr0SpS/jr1SqS0jr2SqS1jr3S2r4S3r5S4r6S5r7SrS6jr8SsS7jr9StS8jr:SuS9jr;SvS:jrSwS=jr?SxS>jr@S?rAS@rB\SA5rCShS`SBjjrD\SC5rE\SD5rFSErGSySFjrHSGrISzS{SHjjrJSjSIjrKS|SJjrLS}SKjrMS~SLjrNSSMjrOS~SNjrPSSOjrQSSPjrRSSQjrSS^SRjrTSSSjrU\VSwSTj5rWSSUjrXSSVjrY\VSSWj5rZSSXjr[SSYjr\SSZjr]S[r^U=r_$)ri z\A class to represent a triton kernel and helpers to generate triton kernel programmatically r{helper_functionszCallable[[sympy.Expr], str]kexprTNzOptional[bool]3transpose_discontiguous_tensor_descriptors_overrideFc >X0lX@lX`l[TU]"U40UD6 [ UR UR5Ul0Ul [5Ul [5Ul [5Ul [["5UlX l[$R&"5Ul[*[,[,4"5Ul[15Ul[4R6"[*5Ul[*[,[:4"5UlXPl[4R@"5Ul!SUl"SUl#[[H"5Ul%SUl&URN(aURQURR5 URT(aURW5 URY5 URT(aUR[5 SUl./Ul/g)NrF)0 optimize_maskr#is_combo_kernelrrr newvar_prefixsuffixrprologue_cacherVprologuepost_loop_combinepost_loop_storerr outside_loop_varsrrcount block_ptr_iddictrblock_ptr_to_bufferr{r9rrpointer_advancementsrr$ hint_overrideCounter _load_counts_pdl_load_index _pdl_has_waitr/autotune_hints triton_metarccodegen_reduction_numelsrjcooperative_reductioninit_cooperative_reductioncodegen_range_treeinit_cooperative_reduction_maskhas_load_with_contiguous_rdimstores_with_contiguous_rdim) r tilingrr=r#rKr>rrs rrTritonKernel.__init__ sx$1(%4 *6*T//=.0(6(8 1?1A/=/?!+C!2#6 %OO-#'S>#3 / 1  # #D ) !$(S>#3 *6A6I6I6K ")6859   ) )$)) 4  % %  + + - !  % %  0 0 2-2*68(rc[R(d [RR(dgURnUVs/sH&n[ U[ R5(dM$UPM( nn[U5S:Xag[RRRXU5n[SU55$s snf![a gf=f)NFrc3*# UH oS:Hv M g7frr)rrss rr4TritonKernel._has_stride1_on_rdim.. s9[6Q;[)r" deterministicrLforce_filter_reduction_configsrrrrrrLrZrR stride_varsZeroDivisionErrorrz)r support_varsr reduce_varsras r_has_stride1_on_rdim!TritonKernel._has_stride1_on_rdim s  F$7$7$V$V)) $ #c=#@#@A #  { q  ''**66u<XK9[999! !  s#C-C*C CCcD[SUR55(+$)Nc38# UHn[U5v M g7fr)rX)rrs rr>TritonKernel.has_store_with_contiguous_rdim.. s 0P d # #0P)r\rXr s rhas_store_with_contiguous_rdim+TritonKernel.has_store_with_contiguous_rdim s' 040P0P    rc[U5$r)rF)r rs r dtype_to_strTritonKernel.dtype_to_str s 5!!rczUR=(a) [RRUR5$r)rcrLchoices should_use_cooperative_reductionrer s rrr-TritonKernel.should_use_cooperative_reduction s-$$ )S)S MM*  rc6^TR(deTRH'nURcMU=RS- slM) TRSnTR(a[ UTRS5nTR RU5Tl[TR 5Tl TRRS5 [U4SjTR55(aTRRS5 gg)z/One time setup code for cooperative reductions.NrNrXra  RSPLIT_NEXT_POWER_OF_2: tl.constexpr = triton_helpers.constexpr_next_power_of_2(RSPLIT) RSPLIT_IS_POWER_OF_2: tl.constexpr = RSPLIT == RSPLIT_NEXT_POWER_OF_2 HAS_RSPLIT: tl.constexpr = RSPLIT > 1 rsplit_id = tl.program_id(0) num_rblocks = (rnumel + RBLOCK - 1) // RBLOCK rsplit_chunk = (num_rblocks + RSPLIT - 1) // RSPLIT * RBLOCK rsplit_start = rsplit_chunk * rsplit_id rsplit_end = rsplit_chunk * (rsplit_id + 1) c3v># UH.nUR(dMTRU5(+v M0 g7fr)rfr,rrr s rr:TritonKernel.init_cooperative_reduction..< s4 (   .''- - -(s99z>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))rSrogrid_dimrdr#rr semaphoressemaphores_namer%cooperative_reduction_workspace_cacherjrrzr)r r sem_counts` rrT'TritonKernel.init_cooperative_reduction s))))$$D}}( " %KK$    4+<+5W II6 2     ((   II  P   rc.SnUR(dUS3nURRSU35 UR5(aURR S5 gUR(aeURRS5 g)Nz$tl.arange(0, RSPLIT_NEXT_POWER_OF_2)z [None, :]zrsplit_arange = z if RSPLIT_IS_POWER_OF_2: rsplit_mask: tl.constexpr = None else: rsplit_mask = rsplit_arange < RSPLIT zSrsplit_mask = xmask if RSPLIT_IS_POWER_OF_2 else ((rsplit_arange < RSPLIT) & xmask))r`rjr_has_constant_xmaskr)r  rsplit_aranges rrV,TritonKernel.init_cooperative_reduction_maskE s~> }},oY7M .}o>?  # # % % II   }} $$ II  e rcZURH}nUR(dURXR5 M1UR(dMDURR UR SURU535 M UR(a[SUR55(aPURSSSS9nURU5nURRSURU535 gURUR5 gg)Nzbase = c38# UHoRv M g7fr)is_looprrs rr2TritonKernel.codegen_range_tree..g s=,=(d U=(d TR@SL=(a TS:gn S nSn[5nTRC5HFnUREURF5(aS nOSnUR9URHS35 MH U(a0TRJ(a[:R<RL(dT(aTRO5(aU(dTR@(d[7T!U- 5S:XanTRQT5(dXU(aQTRRS :XaASS jm"SU4S jjm$SU"U$4S jjm#SUUU!U#UU4SjjnU"5nUbU$SnSnTRUT5nUU4Sjn[WT5(Ga T(d[7TRY55S:Xa U"5unnOL[[S/[7TRY55-5n[]S/[7TRY55-5nSUSUS3nTR^(dTR`(a [U4SjTRb55m!O [5m!TR@(aT!R9TR@5 [eUT!UUTUS9$U (GaU(GdTRf(GaTRh(GaT!Rk5nTR@(aUR9TR@5 [/SQ5nURmU5(d|URoU5nURq5n[ U5(deURuU5 UR#UR$5 URmU5(dM|S/[7TRY55-nUHn[ U[Z5(deTRC5H^nURwURH5(dM%URxn[ U[z5(deTRY5UUU'M` M SSR}[[ZU55-S-n[]U5nSUSUS3nO9U"5unnSUSUS3nUm!O#U(dT(aU"5unnSUSUS3nUm!UcU (dU(a U"5un nOSnU(a [U/5m!TR@(aT!R9TR@5 TRT!5 [eUT!UUTUS9$s sn f) z? Compute the index and mask to pass to tl.load() or tl.store() FrrrrNrrNTtl.int32c[R"XR55nUcg[UR/[ R U5/U/[ RU5/S9$)zg Matches expressions of the form: idx = s * xindex This implies stride (s,), and shape (XBLOCK,). NrrErHrL)rOmatch_affine_block_exprsymbolr4rrrr)r range_treerss rmatch_affine_block1TritonKernel.indexing..match_affine_block sj-DD,,.>&%++,!.!=!=j!I J#H*;;JGH rc >^^UR5n[R"S[R"[R U/S9S9up4[ S[UR5UR[X#55UR[X#U55-5n[R"XURU5nUcgUunnn [R"U5n [ R"R$mTR'UR(5m[+UU4SjU 55(ag[,R/U5n [1XS5/[3U SSUSS5V V s/sH%up[R4"[1X5U 5PM' sn n -nU Vs/sH#n[7X[,R9U505PM% nn[;UUUUS 9$s sn n fs snf) a  Matches higher-dimensional blocks coming from FloorDiv and ModularIndexing. Example expression to match: sN * ((rindex//(d1 * ... * d(N-1)))) + s1 * ModularIndexing(rindex, 1, d1) + ... + s(N-1) * ModularIndexing(rindex, d1 * ... * d(N-2), d(N-1)) This iterates over a block of shape (dN, ..., d1) and stride (sN, ..., s1). (d1,...,d(N-1)) and (s1,...,sN) are wildcards that we match. Note that dN does not appear in the expression, but we solve for it using range tree numels and the other dims. z denom modulo)exclude)rr!Nc3># UH9nTRUT5(+=(a TRU5(+v M; g7fr)rstatically_known_power_of_2)rrr"rRs rrETritonKernel.indexing..match_mod_div_block..; sG".!==eYOOH$@@GGH!-sAArrNr)rrsymbolsr;r<Wildr{rr^rFrrrOmatch_mod_div_block_exprrget_slice_numelsrLrZrRr"rrzrrrrMinrErr4)rrrPdenommodulonum_dims match_resultdimsrHblock_index_exprs slice_numelslinear_block_sizerrqrErrr"rRr s @@rmatch_mod_div_block2TritonKernel.indexing..match_mod_div_block s('--/ !& "!))%**ykJ!  (() HY$>?++oi&OPQ  3KKj&6&6   ' ! %2CCDI 77++ NN:+<+<= ".  %2$@$@$L!-A?1'*,qr*:DH&E&E IIg&7?E&E1 !2 3!2-*H*H*TU!2 3' +#)  3s ,,G"*G c:>TT4HnU"X5nUcMUs $ g)zE Match a block indexing subexpression involving a single range tree. Nr)rr match_funcmatchrrs rmatch_block_subexpr2TritonKernel.indexing..match_block_subexpr^ s3''#J't8E($ #rc >[TTRR5VVs0sHupXR_M snn5nTR 5nUVs/sH'n[ R "X$R55PM) nn[SU55n[5n[X55H@upH[URUR55S:a gT"X5n U c gXy- nMB U[U5- n TRT5 [ R"R$(a[&O[(n [ R"R$(aSn [R*n O[-[.T5mTR15n TR2b TR2nO[ R"R4n[7TS5(aUTSL-nU(a[R8O[R*n U R;UU UTTR<U U S9nU [(:Xa1[-[.T5mTR?UR@5(dgU$s snnfs snf)Nc3@# UHoR5v M g7fr)rrs rrBTritonKernel.indexing..match_block_expr.. s*Q[T;;==[rNFtemplate_out_shape)r5r6rorrprArV)!rErange_tree_nodesrrrrOget_subexpr_involving_symbolrrr4rr intersectionrsum filter_masksr"r use_block_ptrrrrr rrAr;)transpose_discontiguous_tensor_descriptorrrrvr"r4r5)r0rindex_relative_to_xyr_indexrorindex_subexprs range_symbolsr%subexprr5r options_classrArVtranspose_contiguousoptionsrrrrr rGs rmatch_block_expr/TritonKernel.indexing..match_block_exprn sa.8$2G2G2M2M2OP2O$!AvvI2OP/+#557 !, "!,(DD3[[]!, "!+*Q[*Q Q .0 %(%EMD=55g6J6JKLqP#1?F~# *L&F5s>7JJ!!),}}22$0==..$H(7(L(L%04/1J1- 9AACHPP#$!TT- #MMSS-t%9::, $0FF,0(CC,AA& (..'$* +'"&..%&7/!$;;04/1J1-5TT ${Q"s I .I&c>T(a<[T[5(aTS3S4$SSRST55-S-T4$TR5[ TR 554$)Nz.shaperrc38# UHn[U5v M g7frr)rr7s rrATritonKernel.indexing.._get_expand_str.. s*F:a3q66:rjr)rrr*dense_size_strrrK)rr sr_get_expand_str.TritonKernel.indexing.._get_expand_str skj#..(\0$66*F:*F!FFLjXX**,eD4H4H4J.KKKrrrz , tl.int32)c3># UH?nUR(aMTRU5(aM.URS3v MA g7frN)rfr,rrvs rr(TritonKernel.indexing.. sA' 0,,)595L5LT5R)t{{m4( 0sA A A )r )xmaskymaskzmaskr0_maskrr,rrrr)rrrrcrOptional[BlockParameters])rrrrcrr)rz Optional[BlockDescriptorOptions])Aprepare_indexingrrr+operator attrgetterrrrrrrrrrrrrrrrrrrrrrrrr"rrrrrvar_listrallow_block_ptrrris_indirect_indexingrrrrKrrr#r>rorrcrJcopyissubset differencerarr]rrrr*r,r)%r rrrrrFrG index_varsrrrrprefix_matches need_dense have_densehave_loop_varsdense_mask_varsrrrrr rr mask_shapexyzr tmp_masksr expand_listrrqexpand_shape_strr\rrrrs%``` ` @@@@rrITritonKernel.indexing s%%e,''  %/\ *(*=*=f*EFCc5<<00 00#~]22(JTXX..((..sxx8  !2!23%%II))JJJJ''    !. 9 9" 9%c0%Jt$ 9" ~&!+>*a/N3CCHH:1NN/ !2 3489CGH MM ( ( + +d*qj   +5<++-D&&t}}55!%"   4;;-t 4 5 .t33 8S8S-1==??"OOI/0A5--e44  J. ! /B * ,b !b /Bb *b H   .A *   _ _ D'(G" '+ %%e,  L ! ' 'S!5!5!78A=+:+<( L !s4+?+?+A'B!BC $aS3t/C/C/E+F%FG ":,b ;GI  D$8$8&' $ 0 0' 'L  doo.")   j$$$)>)>)>.'^^- ??NN4??3!"HI$--d33 * 5 5d ;I#--/C%c+<====&&s+%%cmm4 %--d33 #ec$*>*>*@&AA &D%dC0000 $ 7 7 9??4;;77"&//C#-c3#7#77#7/3/C/C/Ec/JK, !:'!388C[,A#BBSH $[1 .ykJ * l*9+R8H7IKI'I  Z"1"3<! "M?3I ?? MM$// * )$     %   m "s ]1 ]1cZUR5n[U[5(aU(aU(a US:XdeSnO'U(dSnOU(aUS:XdeSU<S3nOSU<3nUR(a3URSR (aUR 5(dUR(GaUR(a!X R;aURUnXd4$URUSS9nURRU5nU(a [U5U4$[UR5n [U[5(aSU 3nOS U 3nURR!U["R$/S 9n URR'Xz5 [)XS U35n UR(a*UR*R-U 5 X`RU'OUR.R-U 5 [U[5(auXR0U'[2R4HSn UR7U 5n [9S U 55(aM-UR:U nXn;dS USU S35eXU'MU Xd4$URU5nXd4$)a%Generate a block pointer or tensor descriptor for Triton kernel operations. This method creates either a block pointer (for regular Triton operations) or a tensor descriptor (for TMA operations) based on the indexing type. It handles caching and reuse of descriptors for performance optimization. Args: name: The name of the buffer/tensor being accessed var: The variable name for the pointer indexing: Block pointer options or tensor descriptor options containing indexing information and boundary check settings other: Additional parameters string (e.g., padding options) Returns: A tuple containing: - block_descriptor: The generated block pointer or tensor descriptor variable name - other: Modified additional parameters string with boundary check options , other=0.0r, boundary_check=z, padding_option='zero'rUF)r}rFtma_descriptorr5 = c3# UHCn[RRRU[R "S55v ME g7frN)rLrZrRr[rr)rrs rr1TritonKernel.codegen_block_ptr.. sC+:GG,,DD & a(8+:sA A z#duplicate advancement for pointer 'z ' at type 'rC)rrrrcrorrrArArrrkrnextrGrnamedvarruint64rmrUrBrrjrIrrrr\rJ)r rrrIrkcheckblock_descriptorblock_ptr_line block_varblock_descriptor_id named_var line_bodyradvance_offsets advancementss rcodegen_block_ptrTritonKernel.codegen_block_ptrW s2'') h 7 8 8 --- ---+E94KL+E95  ! !  $,,##%%      S,?,?%?$(#6#6s#; h &&e"*e!D HH,,^< y>500&*4+<+<&=#h88)23F2G'H$)78K7L'M$ HH--$ELL.  ^7(1C3~FV/WX $$MM++I6/?'',II'' 2h88BF,,-=>!. = =*2*B*B4*H+:  %'+'@'@'F /CABRASS^_c^ddefC:I%56!!>& && (s3 &&rc^SUSURS3n[[URUR55HGunupx[R R RXx5(dM8SURU'MI URUURURSSS9nUS[[R RU55S3n[U[5(a SUSUUS3$US [RR!UR"5SUS3$) NrrrFTrrr tl.store(z.store()r=rrr:rLrZrRr[r<rrEr get_dtyperrrrrL) r rrIrFrrkrrq broadcast_dims rcodegen_block_ptr_store_line)TritonKernel.codegen_block_ptr_store_line s- #5'H,@,@+AC*3 $$h&>&> ?* %C%#ww77KK27**3/ * 66      7 '/0A0A$0GHIK h 0 0ykE75'; ;GAHH$9$9(:J:J$K#LBugUVWWrcU(dU(dg[U[R5(deURUSSS9n[U[5(deUR nUR 5(a UROSnU(a[URU55OSnURXc(aSOSX5n URU5n URRXS[RS9 g)NFrE0) assignmentr)rrrrIrrr r-texprrename_indexingindirect_assertget_load_bufferrr9rr) r rrlowerrrIrr-size_strlinebuffers r check_boundsTritonKernel.check_bounds s $ ++++==RV=W(O4444&& (0(9(9(;(;8$$8=5--d344## esx %%h/ &5 Lrc&UR5(dUR5(a UR$UR(a?URSR (a!UR 5(d UR$UR$)NrU) rr r:rcrorrrjloads)r rIs rrTritonKernel.get_load_buffer sk  " "h&:&:&<&<<<   ! !  $,,''))99 :: rztl.extra.cuda.gdc_wait()z%tl.extra.cuda.gdc_launch_dependents()c[RRRR(dg[ [ R[RRR5(ag[RR(ag[ RR5RS:H=(a$ [RR!5SS:$)NFrMrr )rr9r"r enable_pdlrrLrselect_algorithmTritonTemplateKernelrrrZrrFrMrr s r_enable_pdl_codegen TritonKernel._enable_pdl_codegen s%%,,77 ahh @ @ U U V V ==   GG / / 1 6 6& @ ; 00215: rconsider_readsc ^^^^UR5(dg[RRm[RR b$[RR R OSmUUU4SjmU(deTSL=(d [U4SjU55nU(dgUR(a*URSR(a URnURUR5 g)Nc>^TceTRRnT(a+[R"UTRR5n[ UU4SjU55$)Nc3># UH8nTTRRURUR5:Hv M: g7fr)mutation_renamesrr)rw current_nodedeps rrOTritonKernel._handle_pdl_before_access..matching_dep..$ s8"A|4488HH"sAA) read_writeswritesrrreadsrz)r7 prev_depsr0r6 prev_nodes` r matching_dep.matching_dep sY( ((!--44I%OOIy7L7L7R7RS " rc34># UH nT"U5v M g7frr)rdr>s rr9TritonKernel._handle_pdl_before_access..* s,Sl\!__lrrU)r-rLrr6rZ scheduler previous_noderzrcrorrjrGDC_WAIT)r  wait_bufferr0 dependencies need_waitr6r>r=s ` @@@r_handle_pdl_before_access&TritonKernel._handle_pdl_before_access s'')) xx,, /0ww/@/@/LAGG   + +RV  |%S,Sl,S)S    T%5%5b%9%A%A))Kdmm,rc4UR5(dgURS:agUR(a*URSR(a UR nUR UR5 UR UR5 g)NrNrU) r- use_countrcrorrCrrE GDC_LAUNCH)r  launch_buffer result_vars r_handle_pdl_after_load#TritonKernel._handle_pdl_after_load3 sr''))    ! #   T%5%5b%9%A%A 22M  .0rc@/nSnSnURHn[U5[LaURU;a U(aM.Sn[U5[La/URU;aUbUR U5 [ U5nURU5 M X!lg)NFT)_linesrFrrErMrarr)r code new_lineshas_waitprevious_launchls r _filter_pdlTritonKernel._filter_pdlC s AAw#~$--1"4#HAw#~$//Q"6".MM/2"%i.   Q  rcNURR[XU55 gr)saved_partial_accumulaterrd)r rrrrus rrvTritonKernel.partial_accumulateT s" %%,, dC 8 rc B ^^^^^URRT5nURmTT==S- ss'[nUR U5mUn[ R RT5nURUSURUUSSS9S9n[U[5(a'URUR5(aSUlUR5mUR!5n[#SUR%U5R'555n UR)U5(aSn OiU (dSn O_UR*(aLUR,SR.(a.UUUUU4S jn TTmS n [0R2"[4S U 5nOS n U(dT(aBUR75(a-UR8(aS [;UR853n OSn OS n Sn [<R>R@(a"URBRE5nUTS:n UR)U5(+=(a( UR*(+=(a U (+=(a U nS nU(aSnSnSn[GT5(alUnU[HRJ[HRL4;aB[<R>RN(a[HRPnOUS[SU5S3- nSnGO[U[T[V45(aURYTX7U 5unn [U[T5(a SUU U US3nO/US[ RZR]UR^5S3nURaUURbURdSSS9nURdnO[gU5(aSUSUS3nURhnSnOcSUSURjSURlU U US3 nURn(a URnnO[pRsUR5nU[HRJ[HRL4;a4[<R>RN(aUS- n[HRPnU[HRt:Xa0[HRvRxcUS- n[HRtnUR{U5nUR}UT5 UR~RUU"U5UUS9nURUU5 URS:a TT==S-ss'[U[5(deURUlDU(aSUSUS3nUR~RUUXgRnS9nUR(aUR(aSnOU[HRt:XaS nOS!nUR8(a[;UR85OUnS"URlSUSUS3nUR~RUUUURS9nUR*(aUR5(d"T(dURRU5 U$)#zS Load from the memory location 'name', offset by some indexing expression 'index'. rNTFrrrEc3*# UH oS:Hv M g7frr)rrZs rr$TritonKernel.load.. s MqFMr^z, eviction_policy='evict_last'rUc6>TTT:aT(dT(agg)N evict_last evict_firstr)expected_countrindirect_indexing load_countsrsr decide_later'TritonKernel.load..decide_later st$~5"3'$rz, eviction_policy=''zrz, other=rz, cache_modifier='.cg'Nrrrrdz.load(rrrErrrSr5rrz0.0Truerr5)JrrgrMrrrLrZrrItma_compatibility_checker_clsrrrerrWrr rzget_strides_of_loadris_broadcastedrcrorr;r<r<r  _load_otherr`r"r skip_l1_cacherebuffer_read_countsrjrrrrrrFrrrrrrLrrEr=rrrr-r rrrrrrrIrr9rPrLrrrrr&rEr)r rrr make_lineoriginal_indexrrIr  is_coalescedeprhrk has_read_depsrprocachemodappend_broadcastrr"r load_bufferrOzero other_valrerrfrgs ` @@@@rrVTritonKernel.load[ siiood#'' DQCK  55e<!!$'== &*&H&H 'I'!   h 0 0T5N5N NN6 6 26D .((* **,   44^DKKM     ~ . .1B1B  " "t'7'7';'C'C % %).N+B!))*:FLQIB :8+<+<+>+>"=1A1A#B"CD%E  == & &!%!A!A!C .t4q8M ##N3 3 ))) !!    /H $ #D ) )D77==77!MMEd;u#5"6a88DE(_6M$NOO*.*@*@#+' %h88%&6%7wrd8*AND./vahh6K6KHL\L\6]5^^_`D==((((#'# >!,,&~66!#d>*:"=#+#6#6 !#d8+=+=*>c(BSBSATUWTXY^X_`h_iijk (($11E)99(..IE%--88MM88))  "u}}'8'8'@& **84  &&{D9XX&& 4U'  ##K<   ! #   " *&78888'11  %j\4D3EQGD**T6K6K+J!!** Dejj(!DD7;7G7GM$"2"23T#8#4#4"5R |2i[PQR!XX..U*:J:J/ $$X-?-?-A-A*  " " & &z 2rc URRU5nUn[RR U5nSnUbUS:XaUS:Hn UR UUSU S9nUR USUSLUS9n [U [5(a;URU R5(aURRU5 XRR;n URU5n U (a,U (a%URR![#US55 [U [$[&45(a)UR)XU 5upUR+XXU5nGOJUcU R,n[/U5(a`UR0bS[3SUR055(d2SR5[7[8UR055nUS US 3- nS US US USU R:S 3 nOUS:XaSUlU R,n[/U5(a`UR0bS[3SUR055(d2SR5[7[8UR055nUS US 3- nSUS US USU R:S3 nO[?SU35e[@RB"5nURD(d;URF(a*URIURKXR55 URMURUSS9 URR![#X55 URD(dURNRQU5 URS5 g)z^ store the 'value' to the memory location 'name', offset by some indexing expression 'index'. NtmaTr_)rrFrGztl.debug_barrier()c3># UHn[U5S:Hv M g7frNrr?s rr%TritonKernel.store..R?;aCFcM;r.broadcast_to(rrrr atomic_addc3># UHn[U5S:Hv M g7frrr?s rrr]rrztl.atomic_add(z, sem='relaxed')z store mode=r/)*rrTrLrZrrkrIrrrerrXrinplace_buffersrmstoresrrUrrrrrrrr\r*r,rr-atomic_add_foundrA contextlib ExitStackrcrS enter_contextguard_cooperative_storerIrErclose)r rrrmoderrrrrGrrI is_inplacermrrkr" indexing_str value_shape exit_stacks rstoreTritonKernel.stores'iit$!!$'$(! <45=EME(,(J(J )K) % == dl&? !  h 0 0T5N5N NN6 6   , , 3 3D 9YY666 ,,^< . KK ! !,t5I"J K h2I J K K&*&<& . Q ?? se4 ~Sr(BSBSATTUVD \ !$(D !#--L%e,,KK+?5;;???"iiC(=> . Q ?? #C5\N#eWBxGXGXFYYijD% D6&:; ;))+ $$)C)C  $ $T%A%A$ %T U &&t{{D&N l467$$  " " & &u -rcZURRSUS[U5S35 g)Nztl.device_assert(rr)r:rrepr)r condmsgs rdevice_assert_async TritonKernel.device_assert_asyncqs( !24&49+QGHrcURR5nUR[USUS355 UR 5$)z For cooperative reductions only one thread block should write out the result. We rotate which thread block does each write for better parallelism zif rsplit_id == (z % RSPLIT):)r{rrrUindent)r rr#rs rr$TritonKernel.guard_cooperative_storetsC 88NNPd.?uK,PQR}}rcSnUH:nUcM[US5(dMUcURnM,X#R-nM< U$)Nr)rr)r  variablesmaskselems r_combine_masksTritonKernel._combine_masks}sGD|t[))= NNE!NN2E rcpURR[R5 URR US5nUR US5n UR US5n UR US5n U(aURR US5OSn U(aUR US5OSn U[R:XaSnO"U[R:XaSnO [S5eUR"URUS/U(aUS/O/Q76 URRURS US US U S U S U S US US US U S U S US 3UURS 9nUR!URU5 UR#XU5nUUlU$) z# See [Note: Inductor bucketize op] rrNr!rrrztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(rz, )r5)rPrr/ONE_ELEMENT_PER_THREADrrgrrrrKrArIr:rr9rrPrr)r r boundariesboundary_indicesindexing_dtypersortersorter_indicesboundaries_ptr boundary_sizeboundaries_underlying_numelboundary_stride sorter_ptr sorter_stride triton_dtyperurs r bucketizeTritonKernel.bucketizes$  C CDA7))*Q-8 &*&7&7 1 &F#++JqM:39TYY__VAY/v 8>))&)4F U[[ (%L u{{ *%L%G  && LL*Q- :@F1I;b "" LL5fXRbr2M1NbQ`Paac nBgRl"]O2  !,,#   ##DLL&9##FnM  rcUR5nUS:XaSUS3$URnS/X#- -S/U--nUSSRU5S3$ NrNz!triton_helpers.promote_to_tensor(rrrrrrrrbr*)r rndimsnreducesizess rreduction_resizeTritonKernel.reduction_resizesh'') A:6ugQ? ?)))VHw,>>$))E*+1--rcUR5nUS:XaSUS3U4$URnS/X4- -S/U--nUb/USX4- QS/U-Q7OSnUSSRU5S3U4$rr)r rrrrrrs rreduction_resize_and_shape'TritonKernel.reduction_resize_and_shapes'') A:6ugQ?F F)))VHw,>>=B=N 9e'u( 9A3= 9TX $))E*+1-y88rcURS:XaU$UR5UR- nUR5nUSUS/-nURR U[ [ U5XV5U[U5S9$)z3 Reshape to RBLOCK, collapsing all reduction dims. rNNrDr5)rbrrKrr9rrr)r r#rr target_ndimr target_shapes rreduction_collapse_dims$TritonKernel.reduction_collapse_dimss  " "a 'L--/$2I2II ,,. $\k2hZ? xx   3u:} C % !  rcJ^^^^^:^;^<^=^>^?^@^ASESjn[R"U5Vs/sHofRPM nn[R"XT5n[ SU55(aJ[ R "T[ R5m[ R "T[ R5mTR(de[STR55nTRU5 [U5nTR(aURTR5 TRSRSn TR (abTR#5n [%U 5S:deU V s/sHn SU ;dSU ;dMU PM n n S S R'U 5S 3m;[)U 5m@O)TR+5m;[)TR#55m@TR-U;UU@4S jU5nS n TS;a[/U[(5(aUupMTR15TR2- m<SFUU4SjnTTU4nUTR4R6;aTR4R6U$[9T5n[;T5n[=TR#55nSUT<'TR4R?U[)U5S9n[SU55Ul SR'U5m:U:4SjmATRB(Ga[DRFRITT5nSm?SHUUU?UA4SjjnTS:XaS nO|[/U[(5(a)[KUU5VVs/sHunnU"UU5PM nnnO>TS:Xa/TR4RMTRNXDRS9nO U"UU5nTS;a[/U[P5(de[RRTRW5nU (a"S[YU 5STR[U5S3nO@[YTR4RMTRNS U S!US"3UUR\S95nS#S$S.Tm>U"TRNUUU5 UUlGOTS%:Xa=TR^(aTRaUTUTAUT5nGOTRcTU5nGOvTS&:XaS[/U[d5(deUunnn[)UU4S'jTRgTRNUUUTU b!S[YU 5STR[U#5S3OU S+3n$TRNR}S,U"S-US.T>S/U"S US US U$S0U"S1TA"U"S23U"5S3US1TA"US23U5S335 U"TR~UU"U5 GO7[T5(aTRaUTUTAUT5nGOTS:XGa0S(US43n%S(US53n&TRpRsU%S)TR+5S6US35 TRpRsU&S7TR+5S US35 TRNR}S8U%S-U&S9U%S U&S US [RS:3 5 TRNR}S8U%S1TA"U%S23U%5S8U&S1TA"U&S23U&5S83 5 Un'TR4R?TU'R\S9n(TRTR~U'U(U%U&T3535 TRUS?3TU5n-TRtRw5n#TRUU#[ Rx"U#5Rz5n.U"TRUU-U.5 GOR[T5(aTS%:XdeUun/n0n1TRU/[;T5US5n2TRU0[;T5US@5n3TRU1[;T5USA5n4TRTRU/U0U1U2U3U4T.maybe_upcastsF ;;MMNN UEMM2  rc3f# UH'o[R[R4;v M) g7fr)rrrr?s rr)TritonKernel.reduction..s M_U]]ENN33_s/1c3># UHoRS3v M g7frrrs rrrMTRRTRSUSTS3URTS9$)Nrrrr5)rr9r:r)r0rr rs r(TritonKernel.reduction...s=dhh'' "1#R'7q9gg! (rN)argminargmaxc>[T 5nT RXT5nT S:Xa_[T R55S:HnURceU(aUS3nS/URQSPnO=US3n/URQSPnO'T R USUSTS3UR5upVUbUS T R U5S3nO URnXRU4$) z3 Helper to generate a reduction call, e.g. tl.sum. r4rz[None,:,:,None]rNz [:,:,None]r>rrr)rrrrKrrrnr) r#r result_typetriton_reduction_fnis_bmmrurrqrrr s rfinal_reduction/TritonKernel.reduction..final_reduction?s#@"O 00FE&T11349{{... %wo6F00a0E %wj1F-ekk-1-E $ ? ?*+1UG2cU!T"XU5un nURUSU35 g)z= Generate a reduction and assign it to an existing variable. rN)r)r#rOrrr\rs rfinal_reduction_define6TritonKernel.reduction..final_reduction_define`s.*&EKE1a MMZLE73 4rc>TRXT5nTRXT5nURSUSUSTSUSUSTSUSTRUS35S 35 g) N z_val, z_idx = triton_helpers.z _with_index(r) r_idx )rrr)r#rOrrrqrroot_opr s rfinal_argreduce/TritonKernel.reduction..final_argreducems00FE00FE MM F:,.DWI\Z_Y``bchbiiklokpq C 5 5 D6I JKL rrr5c3P# UHn[US5(aMUv M g7fr)rB)rrs rrrs!*  C(;CF(CCC5s& &r)cB>T(dU$[RTX5$r)r8r()tvalfvalrs rr>*TritonKernel.reduction..where_conds (..tT@ @rc~UR(dU[R:XaU$X:XdUS:XaU$U[R"U5R:Xa [R"U5R$U[R"U5R :Xa [R"U5R $U$)z7update reduction constant mask value to match dst_dtyper)rrriinfor{r|)rrrs rupdate_constant_dtype5TritonKernel.reduction..update_constant_dtypes ..)uzz2I#O)X]#Ou{{95999 ;;y1555Y!7!;!;; ;;y1555#Orc>T"UTUR5nTR[U5nTRR TR T"X5URUR S9$)Nr5)r_map_tuple_or_scalarr`rr9r:r)rr; default_strr rrr>s r _mask_value+TritonKernel.reduction.._mask_values`/EKKP"77 wO xx((LLu2++++ )ronline_softmax_reducer4r)rrr>rrrzindex, z.shape)r{r|welford_reducewelford_combinec3t># UH-upTRRTRUTUS9v M/ g7f)r5N)rr9r:)rrrrr s rrrs:#) HH%%dllEe%T)s58r\ = tl.full(_indexrr_next, z_next = triton_helpers.z%imum_with_index( # ) r_nextr_max_sumz, float('-inf'), = tl.zeros(z zG_next = triton_helpers.online_softmax_combine( z+ ) z .to(tl.int8)zif HAS_RSPLIT:z_bval = _val_bvalrNr!c3B# UHn[U[5v M g7fr)rrr?s rrrsLAz!%677s)rrr)rrTrrT)rrTrr%rz1tuple[str, Optional[torch.dtype], BlockShapeType])rOrTrrTrr%rrrrT)Vpytree tree_leavesrtree_maprzrrNrrcrrorr+rrrrJrKrr*rrrrrrbrreduction_cacherrrkrlrrr#r6r7rr9r:rTrLrrJrrnrrSrwelford_reduce_fallbackr_welford prepare_softmax_twopass_fallbackrdefault_accumulatorr`rjrreselect_index_dtyperr{rrCr@r"r%online_softmax_reduce_final_reductionget_reduction_combine_fnrrrrrDrrr*codegen_cooperative_reduction_peer_combinewelford_reduce_final_reductionrr\rErrrr)Br rrrrrroriginal_dtypesrreduction_range_prefix dense_sizesr xy_sizes_only logical_indexrrrnacc_typetorch_acc_type result_shaperOr;r masked_valuer0rAaccumulator_dtypeaccumulator_indexmeanm2weight_result_dtype_shape accumulatorrrPaccumulator_maxaccumulator_sum result_max result_sum combine_fnupdatedrbufpeer_valpeer_idx result_mean result_m2 result_weight peer_meanpeer_m2 peer_weightpeer_maxpeer_sumpeersr orig_dtyperrrqrrrrr>sB```` @@@@@@@@r reductionTritonKernel.reductions\ 170B0B50IJ0I990IJ 4 M_M M M++Iu}}EI''u}}=E$$$$MDA%>QR>Qda Aq 1>Q R 5( $xx00uKK0X *5': !55!, <<<<$%HH$K$K$M! *+C ,>+?uTEVEVWhEiDjjk(l%(+)) LL./E.Fgl^[bc"3"."4"4 *)%&+e<^LLL*l2B"WIRPXzYZ[II''&-{43F3F3H2IG9TVW_V``ab!55&' |6$:!"mm>>@  ##()T5H5H5J4K2{{;/334Bt7H7H7U6VVWY&+e<^L%0M*+51B1B;1O0PPQR2359  ## W%6$77NwiX M$5#6br)M S{m5,A;!O PQ"#3z5F4Gu2MO`'a&bc  **J EV&n55!00z8U  #::$%j\"6$%j\"6 ##&'{43F3F3H2IIZ[cZddef ##&'|D4G4G4I3J"XJVWX ##$%W_,=>()O+>@ JJ U[[-E-I-I  4 4j(HU%n55%)99998B5 Y  KK$Y/AJ II$Y/AJ #MM!$Y/AJ 33((!  #::)3& J!'84444JJ 0 ;WQZ JJ 0 ;WQZ::((GG 0 ;W't';';ZPTU    .8  + j% ( (LLLL LL  " " ) )* 5!LL?+q000"%j/O"Cz?c/&:: ::#&z?#CZ!---99 ***44%s3%t,? ,K+LAN$D"j*;<< <<  " " & &z 2?1#55&q)555&&00!l#j\6I/Z[J\6]5^^_`aK4Yd SR%s#zzzz5z  z cURXU5nURXU5n[S5Vs/sH%n[URR US95PM' snupxUR SUSUSUSUSUS[ RSUSURU5SUSURU5S35 Xx4$s snf)Nr!r r9 = triton_helpers.online_softmax_reduce( ) r) rrmrrrlrr"rr) r r#r"r#rqrr\r$r%s r_online_softmax_reduce#TritonKernel._online_softmax_reduces66vPUV66vPUVMRSTX!VX#dhhooEo&B"CX!V   L:,' !O#4Bse2f>R>R=ST LD11ZLBCD LD11ZLBCD   %%"Ws,Cc ^^^^UUU4SjX#U45up#nSUSUSUSTS3 nU4SjnX#U4V s/sH-n TRRTU"U R5S9PM/ n n TRSR U V s/sHn [ U 5PM sn 5SU35 [ U4SjU 55$s sn fs sn f) z+ Helper to codegen triton_helpers.welford. c3J># UHnTRTUT5v M g7fr)r)rrr#rr s rr(TritonKernel._welford..s* +  ( ( > >+s #ztriton_helpers.welford(rrc2>[USTUTS-S-5$)NrrN)r)rrqs r reduced_shape,TritonKernel._welford..reduced_shapes$qcAgi(889 9rr5rc3Z># UH nTRXR5v M" g7fr)rr)rrr s rrr?s* (  + +E;; ? ?((+)rrlrrr*rr) r r#rrrrqrwelfordrArwelford_resultsr$s `` `` rrTritonKernel._welfords F+ &,D6B4r&C5J : F+ + HHOO%}U[[/IO J+   DII&G!s1v&GHIWIVW (    'Hs 4B=C cUR5UR- n[US3[UR 55U[ R "5S9n[US3[UR 55U[ R "5S9n [US3[UR 55U[ R "5S9n URRUSUR5SUS35 URRU SUR5SUS35 URRU SUR5SUS35 US:Xa=Uupn URRS US U S U S USU SU S U SU SU S 35 O9US:XdeURRS US U S U SUSUSU SU S35 URRSUSU"US3U5SU SU"U S3U 5SU SU"U S3U 5S3 5 UnURURUSSUU U UU5 $)z%Helper to codegen a welford reduction_meanrrr_m2_weightrrrrrrz<_next = triton_helpers.welford_combine( z, rrz;_next = triton_helpers.welford_reduce( z1, roffset == 0 ) z rrr8N)rrbrrrKr unknownrjrrr:rrrC)r rOrrr>rrrqr!accumulator_m2accumulator_weightrrrr+s rrTritonKernel.welford_reduces%%'$*A*AA'l% ,,./&&(  +l# ,,./&&(   /l' ",,./&&(   m<(;(;(='>b ! L  l4+>+>+@*AH:Q O  !",t/B/B/D.ERzQR S  . .$ Df LL   W^$4GRS   MZ;-u(={KLM  C n-=U+C^ TUV  J2D1EU/KM_$`#ab   ! 22  " "          rc  [URXXgX55n X#U/n [[X55HAun un upU cURR XS9n XU 'UR U SU35 MC [U 5$)z0Helper to codegen call to triton_helpers.welfordr5r)rkrrrrrlrr)r r#r+r,r-rrrrqrr result_exprsrZ result_exprrrs rr+TritonKernel.welford_reduce_final_reductionIsdmmF"cIJ# > 09#l:S0T ,A, ^e""hhooEoG "-Q MM[MUG4 5 1U \""rc URXU5nURXU5n URSUSUSUSU SUS[RSUSUR U5SUSUR U5S35 X#4$)Nr8rr9r:r)rrr"rr) r r#r$r%r1r2rqrr"r#s rr 2TritonKernel.online_softmax_reduce_final_reductionas66vO66vO  L:,' !O#4Bse2f>R>R=ST LD11ZLBCD LD11ZLBCD   %%rcNUR(aURS$[$)NRSPLIT)r#r3r s r max_rsplitTritonKernel.max_rsplitps"   $$X. .  rc URSnUR5(dSOSnXBR-UR5-nURR U5upxUR RSUSUSURU5S[U5SUS US US 3S S 9 URUS3SS/U[R"5S9n URRU SUS[U5S35 U $)z Generate code to save a [XBLOCK, RSPLIT] temporary workspace, where each thread block writes a different column. After the barrier, every thread block loads the completed value so that it can compute the final value independently. rXzxindex < xnumelNrz_ws = (rez).to(tl.pointer_type(z)) tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), rr:Tstrip_peersrrXrJz = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, rE)rdrrrYr{rrCrrrFcreate_cse_varr rMrDrr`) r rOr default_valxnumelrrrrr3s rr7TritonKernel.codegen_cooperative_reduction_peer_combineusRS!(,(@(@(B(B ..(4??+<<!GGPPQWX %% GG9C0A0A)0L/MMbcnotcubvw$%J:,VXY]X^_   & ##l& !X&&&( $  &&g[ -eers~eeAAC D  rc|UR(deSUl[RRU5nUR USUR UUSSS9S9nSUlUR RU5n[R"5nUR(a*URURXR55 URURU5 [U[ ["45(aZURR%['UUR)UUUR+U5USUR-5<3555 O[U[.5(deUR0n[3U5(a`UR4bS[7SUR455(d2SR9[;[<UR455n USU S 3- nURR%['US US US USUR>S 3 55 URA5 g) NFT)rrrrrErc3># UHn[U5S:Hv M g7frrr?s rr/TritonKernel.store_reduction..rrrrrrrr)!rcrLrZrrIrkrrTrrrSrrrDrIrrrrrUrrrrrrrr\r*r,rr-r) r rrrrrIrrrrs rstore_reductionTritonKernel.store_reductions  $$$$ %!!$'== &*&H&H 'I'!  !%iit$))+  % %  $ $,,T3G3GH  &&t';';SA h2I J K K  * *55  ,+H,C,C,E+HI   h88 88#--L%e,,KK+?5;;???"iiC(=> . Q ??  * *uDc%8CTCTBUUVW  rc f^^ ^ ^ ^^^[5m T RS5 [5m [S5V^s/sH+m[ U U4Sj[ [ X#5555PM- nnSRS[RRU555nT RSUS35 [5mSmS S K J n S S KJn U"5mU"5m "U U U UUU4S jS [ 5n T R#5 [$R&"U "55 U"U6n SRSU 55n T RSU 35 SSS5 SSS5 UR(R+T R-5TS9$s snf!,(df  NC=f!,(df  NL=f)Nz @triton.jitr!c3n># UH*unup#TRSTSU3X2RS9v M, g7f)rr\r5N)rr)rrrrrrZs rr,TritonKernel._lift_helper..s<)G%A~ s1#Qqc]%{{ K)Gs25rc38# UHn[U5v M g7frrr?s rrrjsR.Qc!ff.Qrjz def {name}():rrrn)ShapePropagationOpsHandlerc@>\rSrSrSUUUUUU4SjjrSrg)+TritonKernel._lift_helper..CSEProxyic>T SU3- m [TU5"U0UD6n[T U5"U0UD6nTRT[T U5"U0UD6UUS9$)Nr\r5)rr9) r rrr output_dtype output_shaper dtype_handlerhelper helper_name overrides shape_handlers r_default4TritonKernel._lift_helper..CSEProxy._defaults4&z) &!  #" # '!  #" # ||It,d=f=&& $rrN)rrrztuple[Any, ...]rrxrr )rrrrrxr)rrsrtrurvrwsrCSEProxyros-  '6 @N   rrzc38# UHn[U5v M g7frr)rrTs rrrjsB'F 'rjreturn r)rVrrSrmrrrr*rr from_iterablerrro!torch._inductor.shape_propagationrmr,rrLset_ops_handlerr9rr)r r@rdtypesrZr signaturerormrzoutputsrrsrtrurvrws ` @@@@@@r _lift_helperTritonKernel._lift_helpersg  !'e1X   )23v3F)G   IIRioo.K.KD.QRR = 267#% * PP24 24   ~ 0]]_a// ;$iGiiB'BBG   wwi0 1<_ $$(():k(RRk `<;__s)2F F""2FF" F F"" F0c ^^TR(deTR(aS5e[STR55nTR U5 [ U5nTR (aS5e/n/n[SU55n[R"TRRTR5nTRX#U5nTR5TR- n [!X15GHVupTRRTRU S[#U 5S3U U R$S9n TRRTRSU S TR'5S3U [TR)55S9n UR+U 5 [-U 5n TR.(aMTR)5nS US 'TRR1XS9nS S R3U5S 3nU R4(aSOSnTR6R9USUS US U S35 UR+U5 GMY SmUU4SjnU"ST"U5SU S US3UUU5nTR.(dSnUVs/sH(nU"SUS3[;UR<5U"U5S9PM* nnU"[U5[U55nU"[U5U5n[!UU5VVs/sH)unnU"SUS US3UR<UR$S9PM+ nnn[!UUU5H+unnnTRR9USUS US35 M- OUnUH*n[?U[@5(de[U5Ul!M, [U5$s snfs snnf)z* Perform an associative scan on 'values'. TODOc3># UHoRS3v M g7frrrs rr$TritonKernel.scan..#rrz(ops.scan not supported inside ops.maskedc38# UHn[U5v M g7frrHrrs rrr+Fve*511vrjrrr5rrrrUrrz float('nan')z-1rc2SRSU55$)NrDc3*# UH oS3v M g7frNrrrs rr1TritonKernel.scan..csv..O.csvN88[U5n[U5Vs/sH oPSUSU3PM nn[U 4SjU55(a,UVs/sHnT RR U5PM sn$[ X15VV s/sH(upT RR XRS9PM* n nn T RRT "U 5SU35 [ X5H-upU(aX+l T RRX{5 M/ [U 5$s snfs snfs sn nf)Nrc3Z># UH nTRRU5v M" g7frrcontainsrrnr s rr:TritonKernel.scan..cse_multiple..T#LI488$$Y//rDr5r rrmr\rrrrlrr:rrrmr)r"rrrrrZ cache_keysrnrr result_varsrOrr s r cse_multiple'TritonKernel.scan..cse_multipleQsF A;@8D8aF"QCr%18JDLLLLAKLI Y/LL'*&&9&9NUe;;?&9  LL " "{#$Cv. *-[)E% +0( Y3*F% %ELD" %D'/D,ztl.associative_scan((rcVURcg[UR5nSUS'U$)NrrU)rrk)rrs r_partial_scan_shape.TritonKernel.scan.._partial_scan_shapens*99$ OE #E"I Lrztriton_helpers.select_one((z1), rbase == (RBLOCK - 1), dim=-1, keep_dims=True)ztl.where(roffset > 0, z = tl.where(roffset > 0, )"rcrSrrorr+rrr;r<rr9r:rrrbrrrrrKrrrrlr*rrjrrHrrrr) r rr&rrbroadcasted_values accumulators cse_computecombine_helper_fnrqrr value_dtyper reduced_sizer!reduced_size_strr;rpartial_scan_varsrpartial_scan_varpartial_reduce_vars accs_nextfull_scan_vars full_scan partial_scanracc_nextpartial_reducerOrs ` @rscanTritonKernel.scans$$$$--5v5-MD&?q#A ,1,C,C. ##"m;/?.@7)2hZWXY##K090< = &$)#C(:$;#.D.DE-.>? ): ##5#6>Q8RSI'l(;=NON03>CT/U  0V+I| ,YKr,qI&,,&,, 0V :=<)<:5+~ &&"m## UHoRS3v M g7frrrs rr$TritonKernel.sort..rrz(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc38# UHn[U5v M g7frrrs rrrrrjrrrr5c2SRSU55$)NrDc3*# UH oS3v M g7frrrs rr1TritonKernel.sort..csv..rr^rrs rrTritonKernel.sort..csvrrc d>[U5n[U5Vs/sH oPSUSU3PM nn[U 4SjU55(a,UVs/sHnT RR U5PM sn$[ X15VV s/sH(upT RR XRS9PM* n nn T RRT "U 5SU35 [ X5H-upU(aX+l T RRX{5 M/ [U 5$s snfs snfs sn nf)Nrc3Z># UH nTRRU5v M" g7frrrs rr:TritonKernel.sort..cse_multiple..rrDr5rr)r"rrrrrZrrnrrrrOrr s rr'TritonKernel.sort..cse_multiples&'A;@8D8aF"QCr%18JDLLLLAKLI Y/LL%($C$CLEe;;?$C  LL " "{#$Cv. *-[)E% +0( Y3*F% %ELrrUrrnumelr!ztriton_helpers.sort_with_index(rrNz , stable=z , descending=zUnhandled sort)rcrSrrorr+rrr;r<rr9r:rrbrrrrrKrfr,rjrrr)r rrstable descendingrrrqrZrrrrr"rrO input_varrs` @rrTritonKernel.sorts5$$$$--5v5-MD[!!] s"AH.cUR(dgURUR5 URR5 URR5 g)z Generate the output from prologue. This should be extracted from the subgraph, which is why this is partitioned from codegen_body. N)rBrclearrA)r rTs rcodegen_prologueTritonKernel.codegen_prologues@ }}  DMM"  !!#rc<UR(dVUR(dEUR(d4UR(d#UR(dUR (dgUR Vs/sHoR(dMUPM nnUR(GaUR(dS5e0n[UR5HupEURn[RRU[ R"5nUR%[&U5nSU3nUR(R+USUS35 UR,R/U[ R"SS9X8'M UR(R+S5 UR(R+S 5 UR(R1S S 9 UR35(d[UR S n U R4S :XdeU R4n UR(R+U SU R6SU S35 UR(R9UR5 UR(R;S/5 UR(R9UR5 UR(R9UR5 UR(R9UR5 UR(R9UR 5 [UR5HupEUR<n SU3n[R>"UR[ R"5n [AUR5n UR,RCUR(U SU S3U RDSS9nS SK#nURHRJRMUSUR(5 U "X8U5nSSS5 UR(R+USW35 M SSS5 [O[QUR55H%nUR(R+SUSUS35 M' GO;URR(Ga[QU5S :Ga[U5GH unnUR(R1US 9 UR4nURT(aSOSnURT(aSOUS3n[ RVRX(a[[5S:aSnOSnUR(R+S US!US"US"UR]5S#US$3 5 SSS5 UR(R1US -S 9 UR_XR(5 SSS5 GM UR(R1[QU5S 9 URaUR(5 UR(R9UR5 UR(R9UR5 UR(R9UR5 UR(R9UR5 SSS5 [c/[U5Q5GHrunnUR(R1US -S 9 URdURfRi5HunnU[QU5S - :avUUS -nURdURfUn[jRmU5n[oURpU5nUU5VVs/sHunnUUU-- PM nnnUR(R+[uURvUUS%US"[xRzR}U5S&355 M SSS5 UR,RUR5 UR5 GMu OUR(R9UR5 UR(R9UR5 UR(R9UR5 UR(R9UR5 UR(R9UR5 URT(aiUR(dUR (aGURS'3nUR(R9S(US)3S*S+9 URR5 UR(d%UR(R9UR 5 URR5 URR5 URR5 URR5 URR5 UR R5 gs snf!,(df  GN6=f!,(df  GN=f!,(df  GN=f!,(df  GM=f!,(df  GN=fs snnf!,(df  GN=f),z Concat output code from index_code, loads, compute, stores, suffix into self.body. For pointwise kernels, this is called just once at the end. For reduction kernels, this generates a loop over the reduction axis. Nz1Mix order reduction requires persistent reductionaccumz = tl.full([R0_BLOCK], z, tl.float32)[None, :])rR0_BLOCKr5z/split_size = min(RSPLIT_SIZE, xnumel - xoffset)z@for _ in tl.range(0, split_size, XBLOCK, num_stages=NUM_STAGES):rN)rrrXmask =  < rzxindex += XBLOCKr>z, 0))rr:rz&tl.store(ws_ptr + (tl.program_id(0) + z3 * tl.num_programs(0)) * r0_numel + r0_index, accumz , r0_mask) rsplit_startr rsplit_end)rr!z, num_stages = 2rzfor zoffset in tl.range(rrrlz = tl.advance(rz + tl.program_id(1)zR if HAS_RSPLIT: triton_helpers.x_grid_barrier(rTr\)F indexing_coder'rr:rCrDrormix_order_reductionrrr\rr#r6r rr1rr`rjrrrrrrrr writelinesrr rr9runittestmockpatchobjectrmrrcrSrrrrrrrlrJrrrrrrrrUrIrLrr invalidaterE cache_clearrzr{rr)r r loop_trees accumname2varr partial_accumrr;rentryrXrr&triton_reduction_functionnewvalrr'levelr loop_startloop_end num_stagesrF advancement prev_treeprev_advancement prev_block prev_num_itercurprevsem_ptrs r codegen_bodyTritonKernel.codegen_bodys   zz{{||%%## '+'7'7H'7t<5;;W33M7Kse} ##f3G9-+S',VS!>!>?@ ##Aj + h: 5h== i  i ! i1 4 jj  j c/nUR5(Ga^/nURSU/5 UGHBn[U[5(aUR [ U55 M5[U[ 5(ag[RRRURUR[RS9nUR [ U55 M[U[R 5(a^[RRRUUR[RS9nUR [ U55 GM.[#S[%U535e U$)NrrKfallbackz!Unsupported numel argument type: )radd_numel_to_call_argsrrrrrlrLrZrR size_hint inner_exprrKr"unbacked_symint_fallbackrrrrF)r r numel_argsrhints rkernel_benchmark_extra_args(TritonKernel.kernel_benchmark_extra_argss     ! !+-J  ' 'J ;!c3''KKC)_5577++55&*&8&8!'!@!@6D KKD *UZZ0077++55&*&8&8!'!@!@6D KKD *$'Hc %TUU%"& rc[5nURR5up4pVUR/SQ5 UR 5 [ R "5n/n[XE5GHdupS[U53n [RRU 5n U (a[RRRU R5UR[ R"S9n [RRRU R%5UR[ R"S9nUR'U SU SUSU R)5SU R+5S3 5 GO=U [RR,;a[RR,U n[RRRUR/5UR[ R"S9n [RRRUR15UR[ R"S9nUR'U SU SUSUR2SUR4S3 5 GO.[7U [85(at[RRR;U R<UR[ R"S9nS U R>;aS nUR'U S U35 O[7U [@5(a[RRC5n[RRR;U R URS 9nUR'U S USUSU R4S35 O[ESU 35eURGU 5 GMg URIURK55 UR'SSRMU5S35 SSS5 UR/SQ5 [RRC5nURNnUR 5 UR'S[RRPRSU5S35 UR 5 UR'[RRPRUU55 SU3nUR'USUS35 UR'[W[XRZ5SUS35 SSS5 SSS5 UR/SQ5 UR 5 UR'S[RRPRSU5S35 UR 5 UR'[RRPRUU55 UR'S[W[XRZ5S35 SSS5 SSS5 UR/SQ5 UR 5 UR'S5 UR'S5 UR'S5 UR'S[RRC5R\S35 UR'S U35 UR'S!5 UR'S"5 SSS5 U$!,(df  GN=f!,(df  GN=f!,(df  GN=f!,(df  GN=f!,(df  GN)=f!,(df  U$=f)#a Generates Python code for benchmarking this Triton kernel. - Creates example inputs (random tensors, constants, sizes). - Runs the kernel on the current GPU/stream. - Prints runtime (ms) and throughput (GB/s) using `num_gb`. Args: num_gb (float): The number of gigabytes to use for throughput calculation. Returns: IndentedBuffer: A buffer containing the generated Python benchmark code. )rrzdef get_args():arg_rz = rand_strided(rz , device='z ', dtype=rrhrr)rKz = torch.zeros(z*Don't find the buffer or const tensor for r|rN) rzdef call(args):zwith rstreamz = get_raw_stream(z.run(*args, stream=)rrz def benchmark_all_configs(args):z.benchmark_all_configs(*args))rrzif __name__ == '__main__':z!!$x.!12gg,,X677++66 &*&8&8!'!@!@7D WW--88(&*&8&8!'!@!@9F $$#*$4TF"VHJs~~O_N``ijmjwjwjyizz{|!2!22#$77#4#4X#>L77++66$))+&*&8&8!'!@!@7D WW--88$++-&*&8&8!'!@!@9F $$#*$4TF"VHJ|ObObNcclmymmmAABC 11"#''"2"2"<"< &*&8&8!'!@!@#=#K% 4&' $$z[M%BC66WW@@BFGG,,66 T5G5G7E$$#*OE7*VHIV]VcVcUddef#DXJO  *u&?v   T==? @   wtyy';& ?B 9:<<>$$ ]]_   uQWW%7%7%D%DU%K$LAN O  GG&&11%8!'ug.   K=0B5'!KL  ;22344G }TUV ! JK ]]_   uQWW%7%7%D%DU%K$LAN O  GG&&11%8  c+"9"9:;;XY ! DE ]]_   N    R   0 1   I!''JmJmJoJtJtIuuA    y1 2   = >   N   [_L!_ !__  sf O\/A \*>?OP Q   rcUR(agUR(agUR(aUR(degUR(agg)Nr#rSrr5 pointwise)r#rSrrcr s r_get_heuristicTritonKernel._get_heuristic:sD   !  ' '*  & &(( (()  " "rcz[RRR5[R [R [RR[R[R[R[R[R[RR[RR[RR [R"[R$R&S.n[R((a[R*"5US'[R,R.bSUS'[R0"5(aSUS'[R2(aL[R2US'[R4US'[R6US'[R8US '[R:(a9[R:US '[R<US '[R>US 'U$) N) backend_hashassert_indirect_indexingautotune_local_cacheautotune_pointwiseautotune_remote_cacheforce_disable_cachesdynamic_scale_rblock max_autotunemax_autotune_pointwisemin_split_scan_rblockspill_threshold store_cubinr_r`$are_deterministic_algorithms_enabledTis_hiprprofile_bandwidthprofile_bandwidth_regexprofile_bandwidth_output/profile_bandwidth_with_do_bench_using_profilingcoordinate_descent_tuning coordinate_descent_search_radius'coordinate_descent_check_all_directions) rr%_tritontriton_hash_with_backendr"r+r,rr-r.r/r0r1r2r3r4r5r_rLr`*write_are_deterministic_algorithms_enabledr6rrrr8r9r:r;r<r=r>) inductor_metas rinductor_meta_common!TritonKernel.inductor_meta_commonFs"KK//HHJ(.(G(G$*$?$?"(--"B"B%+%A%A$*$?$?$*$?$?"//&,&C&C%+]]%H%H%}}<<!==44#11.4.A.A.`.` "  < <::< @ A ==   (&*M( #     )-M+ &  # #171I1IM- .7=7U7UM3 48>8W8WM4 5FF K L  + +00 5 677 < =>> C Drc ^-^.[5n0nURR5HupE[U5(aUR(dM([ R RRU5n[U[[R45(dSnO[[U55nXsU'M UcUR[55 [ R R!5R"nUS:XaURS5 OURS5 [$R&(aURUR)55 UR*R-5um-n m.n [/T.5Hup[U [05(dM[3[R4U R65n U [ R RR8;dMj[1U R:[ R RR8U 5T.U 'M [=5n UR>GH'nXR*R@;a(U RCUR*R@U5 XR*RD;amU[ R RF;aOXRF;a@U RC[3[HUR*RDU5RJ5 XR*RL;dMUR*RLUn[U[N5(aeU RCU5 GM* [QT-T.5HXunn [U [R5(dMU RT[VRX:XdM=U RCUR:5 MZ [[U 5n UR]5H[n[1UR^S3UR`5nT.RcU5 T-Rc[eUR:55 M] U-U.4SjnURfHYnURh(aURj(aM'URlcM6U"UR^Ro5S35 M[ URp(aU"S 5 URr(aU"S 5 U"S 5 [uT.URvT-S 9nU[xRz"[ R R!550[|R~R$RR=(a7 S [UR5;=(d S [UR5;S.n[ R R=(d [ R RnUR5R[UR5[[R5U UURURURURURS. UR5EnURr(aURUS '[$R(d[$RR(a"UR=(d URUS'[ RRRi5=(a URj(+nURn[UR5S:HnU(Ga U(GaURRUR5nURRRSnURnUb+SU;a%SU;aUS[USS5- nU[:nO,URRU5[R:HnURRRn URRRn![ R RRU [$RS9[[ R RRU![$RS9S5- n"U"S:aU(a[ R RRURRS5(aN[ R RRURRS5(a US::aSUS'UR(aURUS'UR(aURUS'URp(aURjUS'Sn#[$R&(d[$R(aUR5S - n#U#bU#US!'[$R&(aUR5n$U$bU$US"'[T.5/US#'UR5US$'[T.5Hn%SUS%T.U%R:'M [$R(+US&'UUlwURUR5 UR5 URUR5 URH%n&URS'5 URU&5 M' UR(a5S(UR5S)URR$<S*U<S+U<S,3 n'OUR(aIURRUR5n(S(UR5S-U<S.U(S*U<S+U<S,3 n'O_S'n)[U5S:Xa[[T.55S/:XaS0n)OS1n)S(UR5S-U<S2U)S3U<S+U<S4UGRS,3 n'URU'5 U=(d [[GR5n*URS5U*S6S2GRS7T-55S835 UGR5 [$RGR(aURS9U*S:35 UGR U5 UR*GR 5Hun+n,URU+S;U,35 M URUR5 [$RGR(aURS[5(aTR[U55 TR[USS95 g)NT) is_constexpr)rGrrRrP)rargdefsrs radd_constexpr_arg6TritonKernel.codegen_kernel..add_constexpr_args2-//  h!78 NN78$? @rrrX RSPLIT_SIZE NUM_STAGES) size_dtyperIztl.dot)rr r native_matmul) grid_typerP kernel_namemutated_arg_names optimize_memr`rnum_load num_store num_reduction"has_loadstore_with_contiguous_rdimr!rrXr0_rN)rg?ii Tadd_persistent_rblock tiling_scoresr$rgeA kernel_num_gb kernel_flopconfigs launch_pdlrenable_fp_fusionrz$ @triton_heuristics.z( config=zI, filename=__file__, triton_meta=z$, inductor_meta=z; ) @triton.jit z!( size_hints=z%, reduction_hint=rztile_hint=TileHint.SQUARE,ztile_hint=TileHint.DEFAULT,rzH filename=__file__, triton_meta=z*, min_elem_per_thread=zdef r>c3@# UHoR5v M g7fr) full_namer?s rr.TritonKernel.codegen_kernel..s*J'Q;;=='rrlzpl.enter_scope("rIrzpl.exit_scope(")rVrdrrBrcrLrZrRrrrrrr5rrrrFr"benchmark_kernelr#rrrr\r rrinv_precomputed_replacementsrr mutations input_buffersrrremoved_buffersrW inner_nameoutput_buffersr[rr^ zero_moder_ ZERO_ON_CALLr+rrrrrProrfrrrrSrrkrr0rvrr9rrOrrjr: is_inference is_backward_get_grid_typersetrPrADESCRIPTIVE_NAMEr`rrTrUrVrC rsplit_sizer_rLr`rWrkrrer[rrY memory_stats persistentmemoryrqcount_per_threadr{INNER_REDUCTION_RATIO_THRESHOLDget_reduction_hintr1INNERloopedbytesrrstatically_known_leqreduction_numelstatically_known_gtr$r8estimate_kernel_num_bytesestimate_flopsrgr-rhemulate_precision_castsrQrrrYr9rr#r'rirrr*rrcodegen_static_numelsaliasesrr)/r rrTrrr numel_hintr device_typer\rZrr mutated_argsmutation mutation_argargnamersizeargrJtriton_meta_signaturerQrSrB looped_redr[ two_d_redrs dim_statsmem_ops_per_threadr_coalesce_ratiocontiguous_red looped_mempersistent_memsaved_bytes_ratiorflopsarg_numrtheuristics_linereduction_hint tile_hintrQoldnewrIrs/ @@rcodegen_kernelTritonKernel.codegen_kernelus  ![[..0MF"6**43H3H))77>Jj3 *>??! +C O< !*v )1, < KK13 4''==?DDKe# @A @A&& D==?@#'99#;#;#= Iq *FA#w''ellCHH5QWW--JJJ#*!''"2"2"O"OPV"W$IaL +)3 H99222  !8!8!BCII555AGG$;$;;$8$88  )B)B8)LMXX99333#yy77A %lJ????  .'6 3LGS3 --MM%6%C%CC  . 4l+ ++-D U3TZZ@G   W % NN77<<0 1. A $$D  T%>%>& !2!2!4 5U; <%  % % h '  # # m , l + 1 $"2"2G! /&--agg.Q.Q.ST&&--;;RTYY/P8s4< ? XX&&335Wd>W>W:W **  $) )==55dkkBL$//66::1=I!*!;!; )=(]*$1#7#mC>PRS:T#T !15T!TMM44]C$**+ &,,3399J)44;;AAN ! 0 0 : :V%D%D!;!  **"V-L-L+ ! "S("GG$$99MM115GG$$88MM114'",9= 56   -1-?-?M/ *  # #373K3KM/ 0  % %484M4MM0 1  " "f&>&>335;F!17 o.  " "'')E /4 m,"+I"6!7 I$($<$<$> L!+95G@AK $Yw%7%<%< =6.4.L.L*L &'& dii(  #++F NN2  KK ,   #$$($7$7$9#:; --4478!!,0##0"34 O " "!]]==d>P>PQN#$$($7$7$9#:; *~.$$2#34!!,0##0"34 OI:!#/ :;q@ }}]s Cv>> w c[RRRU5n[ U[ R [45(a[U5n[U5nU$Sn[RRRX5(dIUS:a[SU35eUS-n[RRRX5(dMIU$)Nr!i@z!Failed to find static RBLOCK for ) rLrZrRsimplifyrrrrr5r|r)rrs rr#TritonKernel._get_persistent_RBLOCKs!!**62 fu}}c2 3 3f+C!#&C Cgg&&;;FHH?$'H%QRRqgg&&;;FHH JrcP[RU5 g![a gf=f)NTF)rrr)rs rhas_persistent_RBLOCK"TritonKernel.has_persistent_RBLOCKs*   / / 7  s  %%cRS SjnURGHnUR(aUR(ai[RR R UR5nU"U5(a)URURS[U535 UR(aUR(aUR(a1URURUR55nSUS3nO8URUR5nUR (a [#US5nURURR%5SU35 URS:XdGMlUR&(dGMURS5 GM g ) ay We get a small speedup from hard coding numels if they are static. This code stomps on the passed-in values by writing an constant to the top of the kernel. In a kernel like: def KERNEL_NAME(in_ptr0, in_ptr1, out_ptr2, xnumel, r0_numel, XBLOCK : tl.constexpr, R0_BLOCK : tl.constexpr): We would add xnumel = 4096 r0_numel = 768 After the signature, before the kernel code, if we decided to make these static. As its hardcoded, it becomes a better signal to triton on how to unroll and do some static indexing. So, it's not so much that downstream knows that its a static numel, as that you just plop a constant into the kernel. cB[U[R[45$r)rrrrrs ris_static_integer=TritonKernel.codegen_static_numels..is_static_integersdU]]C$89 9rznumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)rzBLOCK: tl.constexpr = rXzXBLOCK: tl.constexpr = 1N)rrrr)rorfrcrLrZrRrrrrrrrSr:rrrJr{rr`)r rTrrsimplified_tree_numelrrs rr"TritonKernel.codegen_static_numelss,$ :$$D$$(=(=()(8(8(A(A$**(M%$%:;;NNdkk](3?T;U:V#WX  T%>%>-- JJt';';DJJ'GHEFugMfgC55djjAC,,!#rl$++"3"3"5!66LSERS{{c!dmmm9:'%rcH[URVs/sHn[UR(+5PM sn5nUR(aUS:Xde[ R $UR(aUS:Xde[ R$US:Xa[ R$US:XaN[[URUR55(a[ R$[ R$US:Xa[ R$[!SU35es snf)NrNr!rz"Unsupported number of dimensions: )rrorrfrr-MixOrderReductionGridrSCooperativeReductionGridGrid1Drzr,rGrid2DWithYZOverflowGrid2DGrid3Dr)r rrs rroTritonKernel._get_grid_types 8H8HI8H***+8HI J  # #6M6$:: :  ' '6M6$== = !V$++ + !V3t22D4D4DEFF(===$++ + !V$++ +=aSABBJs$DcURHn[UR[R[R 45(a URnO)[ RRRX5nUR(aUR(dMURU5 UR[U55 M gr)rorrrrrrLrZ wrapper_codegenerate_numel_exprrfrcrrF)r rr arg_typesrrs rr#TritonKernel.add_numel_to_call_args s$$D$**u}}ell&CDDzzww++??K$$(=(=(=  &  d,%rc|[RRnUR5 URR 5upVpWUR XU5 URRHnURU5 M URUUSUURS9 U(aUR5 gg)NT)rrrQ) rLrZrwrite_triton_header_oncerrrworkspace_argsgenerate_workspace_allocationgenerate_kernel_callrQdeallocate_workspaces) r rr` deallocate_wswrapperr\rrwss r call_kernelTritonKernel.call_kernels''&&((*%)YY%=%=%?"a ##DY?))**B  1 1" 5+ $$  (( %    & & ( rc[RRnURR 5up#pB[ X45HupV[ U[5(dM[RR(aURSUSUS35 MUSUS3nURU5 SUS3nURU5 M g)Nz:AOTI_TORCH_ERROR_CODE_CHECK(aoti_torch_check_inf_and_nan("z", z));z assert not z.isnan().any().item()z.isinf().any().item()) rLrZrrrrrr] cpp_wrapperr)r rr\rarg_signaturesr arg_signaturer"s rcodegen_nan_checkTritonKernel.codegen_nan_check/s''&&*.))*B*B*D'n"%i"@ C-3377&&%%TUXTYY\]`\aade)-BCD%%d+(-BCD%%d+#Arc[U0UD6$r)r)r rrs rr_TritonKernel.create_cse_var>s $1&11rc\URSURURUR553nURR (d!UR (a,URS:XaURRU5 gURRU5 g)NrrX) rr:rrrootrrrrrrj)r rr"s rcodegen_iteration_ranges_entry+TritonKernel.codegen_iteration_ranges_entryAsy**SD,@,@,L!M NO ::  $":":u||s?R    ( ( . II   %rc>URceURUR5nURnUS:waSUS3OSnUR(a'UR(aUR (aUS3nSUR R5SUU3$)Nrrrrz + rsplit_startz tl.arange(0, zBLOCK))rindexing_size_strrrSrrfrr)r rrrr@s rr)TritonKernel.iteration_ranges_ranges_codeLs+++%%e&6&67&& *5*C4 }A&  & &))""x/Fu||1134F4&IIrc^URnUR5nS/U-nSUSUSUS3$)NrNrrr)rr)r rrrrrs riteration_ranges_scalar_code)TritonKernel.iteration_ranges_scalar_codeYsC&& &&(sTz$r%;-q99rc0URceSURS3nURU5(a#SUSURS-SURS3nURRX"5nURS:waUS URS3$U$) Nztl.program_id(rr>z + tl.program_id(rNz) * tl.num_programs(rErr)rxr pid_cacherr)r rrpids riteration_ranges_get_pid%TritonKernel.iteration_ranges_get_pidas~~)))u~~.a0  & &u - -cU+ENNQ,>+??STYTbTbSccefCoo!!#+   z )U$t//02 2 rcURS:H=(aq UR(+=(aY UR(+=(aA [RR R UR[55(+$r) rxhas_zdimrSrLrZrRr|rr4)r rs rr#TritonKernel.needs_yz_grid_overflowpsa NNa  YNN" Y... YGG$$99%++~GWXX  rcUR(a URUR5S3$[UR5$)Nr)r#rr2)r rs rr"TritonKernel.max_blockxs;   $$ '7u%=> > //rcUR(a:[RRR UR S5(agUR (dgUR(a[URR5S3UR;a0URURR5S3S:XagOKUR(d:[RRRUR S5(agUR(a-UR(aURUR 5nO?URS:XaUR(aSnOUR!UR5nUR(a#UR"(aX R%5-n[RRR'UR U5(aeUR(S:g=(dO UR*=(d< [RRR-UR [/55$g)NrFrrNTrX)rJrLrZrRstatically_known_ltrr=r#rrr>r[rfrrr`r"rSrYrrxrr|r4)r rr"s rr,TritonKernel._has_constant_mask}s  ww33DJJCC!!   DKK$5$5$7#8!>$BSBS!S  DKK$5$5$7#8!>?1DE%%ww77 AFF   !:!:33DJJ?I [[C DMMIt{{3I   !;!;!OO$55I 77   8 8Y O O "W==W77##88^EUV  rcfURSnURS:XdeURU5$)NrrX)rorr,)r xtrees rr TritonKernel._has_constant_xmasks5  #||s"""&&u--rcURH9nURU5(dMURURS35 M; URS5 g)Nrr)ror,r]r)r rrs rrTritonKernel.filter_maskssL$$D&&t,,!!T[[M"67% &!rc[[R5SURVs/sH n[UPM sn$s snfr)rkrrrbr)r rs rget_reduction_prefixes#TritonKernel.get_reduction_prefixessG]::;U>UV V t V   s?cURVs/sHo"R(dMUPM nnSR[SU555nUR SUR U535 URVs/sH3nUR(dM[ RURPM5 nn[U5nUR SUR U535 gs snfs snf)zN Generates code that flattens ND reduction numels, block sizes, etc. into 1D. rc3># UHoRS3v M g7f)rNrrs rr8TritonKernel.codegen_reduction_numels..s"U_Tkk]%#8_rz rnumel = zRBLOCK: tl.constexpr = N) rorfr*r+rr:rrrrD)r r#rreduction_treesr rn_blocksrs rrR%TritonKernel.codegen_reduction_numelss -1,<,<R,rc>URnUR(a%URURSUSUS35 OURcDURURSUR U535 URUS35 OUR bUSUR U53nOURXS35nUR(dUR5S3OSnURUSURU5S U3URSU3/5 URU5(a!URURU55 gUS :XaUR(d%URUS URS US 35 gg)Nrz offset + rz offset = 0rrrLz offset = rrXrrr)rrrrrxrrrrrrrr,create_constant_mask)r rrTrXr" block_sizes rr,TritonKernel.iteration_ranges_codegen_headers LL == NNejj\QCy4@ A ^^ # NNejj\T-N-Nu-U,VW X NNaS + ,+Id&G&G&N%OP88#V M,0+C+C1779+U#  OOc4#@#@#G"HJ<Xzzl#dV,   " "5 ) ) NN444U; <s(t77 NNaS |3qc? @ 8r)rMrOrNrrPrGrIr{rr#rWr9rKrcr>rr=rErJrCrDrBrArzrXr$rQ)rTNNF) rYzdict[str, sympy.Expr]r#zOptional[FixedTritonConfig]rKz Optional[int]r>rrrr/rrrrr0)rrrz Optional[Union[str, tuple[str]]]rGz!Optional[TMACompatibilityChecker])r)rrrrrIz/Union[BlockPtrOptions, TensorDescriptorOptions]rztuple[str, str])rrrrr rrr)rTrV)rrrurx)rrrrr) rrrrrrTrrKrrr)rOptional[CSEVariable]NN)rrTrz.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]rrTrrrrrz Optional[tuple[str, sympy.Expr]]rr rrT)rztuple[str, BlockShapeType])rrTrrrrT) rrrrrrJr+Union[CSEVariable, tuple[CSEVariable, ...]]rr )rrr)rrrrrrT)rtuple[CSEVariable, ...]rtuple[torch.dtype, ...]rr)rr r&zUCallable[[tuple[CSEVariable, ...], tuple[CSEVariable, ...]], tuple[CSEVariable, ...]]rr rr ) rr rr rrrrrr )rr})rzOptional[float]rrV)rz type[triton_heuristics.GridExpr]r$)rrr`zOptional[IRNode]rr)rr)rrb)rrcrr)rrcrr rr)rrcrr)rrrr)rrcrr)rrrr)r#rVrr)r@rrzlist[sympy.Symbol]r)rrYrr)rrcrTrVrr)`rrrrrr8rvrrr:rrrkr;rr)rer1rkrnrrrTrVrUrrrrrIrrr$rrErMr-rIrPrYrvrVrrrrrrrrr5r;rrrr rYrrfrrrrrrrr#r'rCrrrrrorrrr_rrrrrr"r,rrr;rrRrrrrrrrrs@rrr s&I%%).E &.O$;!JN7N 48'+ %39%39 2 39 % 3939 3939j::6  " #J*:0  ""8<GKH H 5 H $EH ^ f'f'f'B f'  f'PX:MMM M  M4 *H8J  :?-<1  "  :H xvSWZZ *Z3>ZFOZ ZxI &480499C9& 9 $ 9  919.9 9v. 9 ( 1<  (iii& i ; i 5 iV&DO&" 0F P#0 &!  B>>> >@>S1>S;R>S >S@"'" " ( " !"BD"'D"(D" D"  D" ! D"L $m%^4||  ,,\AF  (;TC$ -OS))/)GK), ,2 & J:(:14: :  0 /b. "  F$W  -? A(A0>A AArrc ~^\rSrSr%\rS\S'\"\R\R\R\R\R\R\R\R /5rSU4Sjjr\SSj5rSSjrSrSSSjjrSSS jjrSS jrSS jrS rS rU=r$)TritonSchedulingiz type[Any] kernel_typec>[TU]U5 Ub[US5(dgURH+n[ U[ [ 45(dM [UlM- g)Nr^) rrrr^rr9r7debug_triton_codedebug_device_str)r rCr`rs rrTritonScheduling.__init__.sN #  GIw$?$? OOD$0B CDD(9%$rc[RR(d[RR(a'[ /UR Q[ RP5$UR $r)r"rcooperative_reductionsforce_cooperative_reductionsrbackend_featuresrQREDUCE_TO_SINGLE_ELEMENT)rr s rget_backend_features%TritonScheduling.get_backend_features6sR MM 0 0}}99P#&&P(O(OP ###rc^ [RRn[X5upEU(aUR U5 [ R (aSSKJnJ m [U 4SjU55(daUVs/sH%n[Xv5(dMUR5PM' nnUR URSSRU535 U(a[UU5n UR!X)5 ggs snf)Nr)r6ForeachKernelSchedulerNodec3<># UHn[UT5v M g7fr)r)rrr s rr3TritonScheduling.codegen_comment..MsCPa 1899=sz Fused node name list: r)rLrZrr? make_commentr" debug_fusiontorch._inductor.schedulerr6r rzrget_namecommentr*r+write_provenance_debug_handle) r  node_schedulerQrorigins_detailed_originsr6r node_names debug_handler s @rcodegen_comment TritonScheduling.codegen_commentAs''&&%8%P"    )     CP+*!!7!AJJL* $$''>tyy?T>UV BL  1 1+ L s <C?C?cr[RRnXR;aURUnU$[R R (a$[U[R R 5OSn[U5SSnSRSXvUR5/5n[RR(a[RRSU3nXTRU'[R R(aUOSnUR[[ R"5U5nUR[[ R$5U5nURSS5n['[)UR+55S5upn [-5n [.R15(a[.R X5 U R3S U<S 35 U R5US S 9 [RR75n U R3S U R8S35 SU 3n[;X$5unnUSU-S-U-- nUR=X\R?5U5 [@RB"S5(a[@RD"X[U5 U$)Nrrr\rtriton_z #pragma CMT#pyzasync_compile.triton(z, '''Tr\z''', device_str='z')z# kernel path: rkernel_metadata)#rLrZr src_to_kernelr"rdescriptive_namesr>rMr*next_kernel_suffix aot_inductormodel_name_for_generated_filesunique_kernel_namesr rrArqrr(r'r]rV async_compileuse_process_poolrrrrFr? define_kernelrr$is_metric_table_enabledlog_kernel_metadata)r src_coder! rrrQ fused_namekernel_category subs_name _basenamer\ kernel_pathcompile_wrapperrmetadata_commentr" detailed_originss rr5 TritonScheduling.define_kernelbs]''&& ,, ,!//9Ktm==22&mV]]5T5TU  AJ2ANO((?8R8R8TUK""AA"(!4!4!S!S TTUVaUbc /: ! !( +'-}}'H'H iI  ''K,H,H(I;WH''K,C,C(DiPH '' s;H(08>>;K1Ld(S %I+,.O--//$$Y9  % %(=i]%&P Q  " "84 " 8WW@@BN  % %(9.:M:M9Nb&Q R!0 > (;M(S %G% w 58H H H   ! !5579I ../@AA++KhOrcURUSS9n[R"U5nURXB[ SU55S9$)zk Benchmark fused list of nodes and return the execution time in milliseconds on randomly generated inputs. T)rdc3@# UHoR5v M g7frr rrs rr9TritonScheduling.benchmark_fused_nodes..s:WQVA::<TRce[RRTR5SS-$Nrz .kernel_perf__file__ospathsplitextrsrcache_file_pathDTritonScheduling.benchmark_codegened_module..cache_file_paths6||///ww'' 5a8>IIrc>>T"5n[U[T55 grr*r)rR rU mss r store_cache@TritonScheduling.benchmark_codegened_module..store_caches&(T3r7+rc>T"5n[RRU5(a.[U5n[ UR 55sSSS5 $g!,(df  g=fr)rQ rR existsopenr1readrR fdrU s r load_cache?TritonScheduling.benchmark_codegened_module..load_cachesJ&(77>>$''dr$RWWY/$$s A A*rM%kernel src code for %s written to: %srz*Exception (%s) in compiling fused nodes %sinfrNc4>T"TR"T6S5$r clone_argsrcallwrapped_jit_functionsrr=TritonScheduling.benchmark_codegened_module..D!5!@!@$!G!JKrr c">TR"T6$rrg rrk srrrl s 4 ? ? Frz+The fused kernel for %s took %.3f ms to run)rrLrZrrr rrrrrP get_argsrj r) rh Exceptionr"r.disallow_failing_autotune_kernels_TESTING_ONLYr1 launchersrn_spillsr. benchmarkrRr)r rrI r$ device_interfacerZ rb r'rt r rrU rj rY rk s ` @@@@@rrH +TritonScheduling.benchmark_codegened_modulesK4AGG4G4GH   # #AGG$G$G$I JB J , )4 *i[:Q  II7   B~3<<'; K J ! @<<>D88D#&;;  ()44d;A>?-66Iy>Q& &&|$$'995\<<>!**K!+==>Bk33F"6{B II=  Ms||#a K J ! L (==OO @ 5\ 3<<''_ K J ! L (K K J J ! sh.J""AJ J"(J=HCJ4 J" JAI?'J(J, J"?JJ J J"" J0cURS5nU=(a [SUR555nURnU(aSSKJn UnU(aSUS'URS5(a SUS 'SUS'[ RUR5(dURS 5(aeSUS '[RRXaX#5nU"U0UD6nURXU5$) Nrc3@# UHoR5v M g7fr) is_split_scan)rr`s rr9TritonScheduling.create_kernel_choices.. s( -NT   -NrrN)TritonSplitScanKernelFoverride_cooperative_reductionrToverride_persistent_reduction) contains_oprzscheduler_nodesr triton_split_scanr} rrr}rrLrqtriton_kernel_kwargsadd_multi_kernel_choices) r kernel_features kernel_args kernel_kwargsis_scanr{ r r} rs rcreate_kernel_choices&TritonScheduling.create_kernel_choicess "--f5 C( -<-L-L-N( % +/*:*:  @/K >CM: ;  & &v . .=AM9 :>CM: ;11/2Q2QRR$(()HII II=BM9 : 66 + k;];,,V-PPrc $U/n[RR(dU$UR=(a UR S5(+nUR =(a UR S5(+nU(a%UR UR"U0UDSS0D65 U(aURRn[RRRUS5(aeUR UR"U0UDSS0D6=n5 U(a7UR(a&UR UR"U0UDSSS.D65 [U5S:a-USSHn URU lM UR!SS9 U$) Nr r~ Fi)r~ r rNcUR$r)r)r/s rr;TritonScheduling.add_multi_kernel_choices..\s q'='=rr)r"r multi_kernelrrrSrr rer}rLrZrRr|rmust_keep_buffersr) r rr r kernelsoptional_persistentoptional_cooperativerrkkernel2s rr )TritonScheduling.add_multi_kernel_choices*s (.h}}))N$99 -BSBS +C ?  &;; MDUDU ,E A   NN   #38  __44Fww44VUCC!--$'8=E'5+E+ENN(((+rc*^^^^^^^U4SjmU4SjnUUU4SjnS/pTSn[RRn[U5[Rl[RRn[U5[Rl[ R S:n [ RS:n URUSU U SS9n U GH upnUVs/sHoR5PM nnUVVs/sHnUHnUR5PM M nnnU R[[R5S5n [R "U 5m["R%S UTR&5 U"5ummTb'UT- nUT- nUR)TR&5 MTR+5mTR,mTR.mT"TR0"T6S5 TR2n[5U5S :XdeUSR6S:a[9S 5=mmOS[RR;5n[<R>"UUU4S jUS 9m[<R>"UU4SjUS 9m["R%S[SU55TT5 U"5 UT- nUT- nUR)TR&5 GM U[RlU[RlXFU4$s snfs snnf)Nc>TRce[RRTR5SS-$rN rO rT srrU @TritonScheduling.benchmark_combo_kernel..cache_file_pathds6<<+ ++77##CLL1!4~E Erc>T"5n[RRU5(aC[U5n[ SUR 5R 555sSSS5 $g!,(df  g=f)Nc38# UHn[U5v M g7fr)r1)rr's rrNTritonScheduling.benchmark_combo_kernel..load_cache..ls E3Daq3Drjr )rQ rR r] r^ rr_ splitr` s rrb ;TritonScheduling.benchmark_combo_kernel..load_cachehsW"$Dww~~d##$Z2 E2779??3D EE Z Zs .A11 A?c\>T"5n[U[T5S-[T5-5 g)NrDrX )rR rU rY ms_clones rrZ .store_cacheos&"$D s2w}s8}< =rrgT)subkernel_nodescustom_part_algorithmenable_autotune mixed_sizesonly_gen_src_coder) rd rNre c4>T"TR"T6S5$rrg ri srr9TritonScheduling.benchmark_combo_kernel..rm rrn c(>TR"T6S$rrg rp srrr s0;;TB1ErzDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc3@# UHoR5v M g7frrD rE s rr:TritonScheduling.benchmark_combo_kernel..s<A::<<r) rLrZrhrinplaced_to_remover"combo_kernels_autotunecombo_kernel_allow_mixed_sizesgenerate_combo_kernel_code get_nodesr r rrArr)rVrrrP rrq rj r) rh rt rru r1rr.rv )r  node_listrb rZ total_ms file_listtotal_clone_msremoved_buffers_originplaced_to_remove_origr r kernel_code_listr8 r\ node_groupr`fused_node_listsr^rnamesrt r rrU rj rrY r rk s @@@@@@@rbenchmark_combo_kernel'TritonScheduling.benchmark_combo_kernel_s  F  > ) # ww66",-A"B"#''"<"<%/0G%H" 77!;;;a? ::%"&+#" ; (8 #H=GHZT 0Z H/?O/?eAQZZ\\/?EO''K,C,C(DiPH""8,C II7   &D88D#&;;  %00$7: ;,66Iy>Q& &&|$$q( %e ,X<<>!**K!'00E! IIV<<<   M NH h &N   S\\ *k(8l#7%<"22o IOs L 7$Lr)rCzOptional[Scheduler]rr)r z torch.devicer)rx)rtuple[float, str])rxN)r$ zOptional[OrderedSet[str]]rr )r rsr list[Any]r rxrlist[TritonKernel])rrr r r rxrr )rrrrrr rrrQFOREACH BUCKETIZEINPLACE_BUFFERSMASKED_SCATTER_WITH_INDEXSCANSORTTRITON_TEMPLATESTUPLE_REDUCTIONr rrr r& r5 rJ rH r r r rrrs@rr r s)K)!  " "  $ $  * *  4 4        + +  * *  :$$MB=~  RVW$5NW$ W$r#Q+#Q#Q& #Q  #QJ333& 3  3j]3]3rr c4/nUR5nUb![U[R5(deU(a1URc$UR UR 5S35 U$SSKJn UR5nUceURRU5n[U[U45(dS[U535e[RR!U5 ["R$nUR'UR)55R+5nU["lSSS5 UR UR 5S35 UR [,R."WS55 U$!,(df  NX=f)Nz" Unfinalized multi template bufferr)CUDACombinedSchedulingz]Scheduling backend should be SIMD or CUDACombined when generating debug Triton strings, got: z Triton code:z )get_template_noderr#MultiTemplateBuffermake_kernel_renderrr 0torch._inductor.codegen.cuda_combined_schedulingr rrC get_backendrfrFrLrZset_current_devicer$generated_kernel_countrG r r]r r)r`linesmulti_templater r backendold_generated_kernel_count triton_codes rr r s_ E++-N  !Z@V@V%W%WW W.;;C  ((JKL2 L/ "!!!..,,V4'N4J#KLL klpqxlykz { LWW ' ' /*1)G)G &!AA eg .HG *0   ( 67 X__[&9: L0 /s .A F  F)rrr0)rrrr9rr9rrr )rrrr)rrrr)rzUnion[CSEVariable, Any]rr)rror)rrrzCallable[[_T], _T])r`r6rr}) __future__rrrrr;rloggingr&rrQ r abcrcollections.abcrrrrtypingr r r r r rrsympy.printing.precedencerrtorch._loggingtorch.utils._pytreer%_pytreertorch._dynamo.device_interfacertorch._dynamo.utilsrrtorch._prims_commonrtorch.utils._ordered_setrtorch.utils._sympy.functionsrrrtorch.utils._tritonrrrutils._sympy.symbolrrrrutils._sympy.value_rangesr rr"r#r$r3 r& codecacher'r(r)r*rr+ ops_handlerr,runtimer-runtime.benchmarkingr. runtime.hintsr/r0r1r2r3runtime.runtime_utilsr4r5rCr6r7r8r9shape_propagationr:r;r<r=r>r?r@rArBrCrDrErFrGrH virtualizedrIrrJrKrLwrapper_benchmarkrMblock_analysisrOcommonrPrQrRrSrTrUrVrWrXrYrZr[r\r]r^r_simdr`rarbrcrdrerf triton_utilsrgrhrirjrkrrltypesrmrro%torch.fx.experimental.symbolic_shapesrprqrrsimd_kernel_featuresrsrt getLoggerrr_logginggetArtifactLogger perf_hint_log schedule_log fusion_logrwrrrrrrrrr3rrrrr6rrrrrrrrrrr_initialize_pointwise_overridesr8r{r4rrrrrrrr r rrrr s"  88EE 0 $$C<0/KK YX4))(FF;('.DWW5 CBB/$% L>&8 B !00<H ~~//*E ^^ - -h A  #$,  6 6  4$ 4%%PQ,Q,h    B ~~~B  ?4??8 =,==@+ +/+pQMpQf 3 &8 ;P 8( .bx &kx &v//9H"OH"V$+$+N N N N b!!H ### %uS%S/-A'BBC  jjjZg5A:/0g5ATk]3~]3@ r