9j> TddlmZddlZddlZddlZddlZddlZddlZddlZddl Z ddl Z ddl Z ddl Z ddl Z ddlmZddlmZmZmZddlmZddlmZmZmZmZddlZddlmZddlZddlZddlm cm!Z"ddl#m$Z$dd l%m&Z&m'Z'dd l(m)Z)m*Z*dd l+m,Z,dd l-m.Z.m/Z/m0Z0m1Z1m2Z2dd l3m4Z4m5Z5m6Z6ddl7m8Z8m9Z9m:Z:m;Z;ddlm?Z?m@Z@mAZAm Z ddlBmCZCddlDmEZEmFZFmGZGmHZHddlImJZJddlKmLZLddlMmNZNddlOmPZPddlQmRZRmSZSmTZTmUZUmVZVddlWmXZXmYZYddlZm[Z[m\Z\m]Z]m^Z^m_Z_ddl`maZaddlbmcZcddl mdZdmeZemfZfmgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZqddlrmsZtmuZumvZvmwZwdd lxmyZyd!d"lzm{Z{d!d#l|m}Z}m~Z~mZmZmZmZmZmZmZmZmZmZmZmZmZmZd!d$lmZmZmZmZmZmZmZd!d%lmZmZmZmZmZd!d&lmZer,dd'lmZdd(lmZdd)lmZdd*l@mZd!d+l|mZd!d,lmZed-Ze jPeZejVjYed.ZejVjYed/ZejVjYed0ZeCZBd1Zd2Zdjd3Z dkd4ZGd5d6Zeddld7ZGd8d9ZejnGd:d;ZejnGd<d=ZejnGd>d?eZejnGd@dAeZ dmdBZGdCdDeZej|ZdndEZdndFZdodGZdndHZdpdIZdqdJZdrdKZdsdLZGdMdNeZdtdOZdudvdQZGdRdSeZejeːjdPeːjdTGdUdVe˫ZGdWdXZejnGdYdZZGd[d\ZejnGd]d^ZGd_d`eeeeeefzfZejnGdadbZGdcddeeZGdedfeثZGdgdheZdwdiZy)x) annotationsN)abstractmethod)CallableIterableSequence) lru_cache)Anycast TYPE_CHECKINGTypeVar) PRECEDENCE)get_interface_for_device)identitypreserve_rng_state)is_integer_dtype type_to_dtype) OrderedSet)CeilDivFloorDivModularIndexing TruncToFloat TruncToInt)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)BaseSchedulerNode$FusedExternTritonKernelSchedulerNodeFusedSchedulerNode Scheduler SchedulerNode)get_broadcasted_shape)get_raw_stream_name) 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_hintsschedulefusionc<|dv}|rdnd}|dvr|d|dS|d|S)N)anymaxminprodtriton_helperstl)r~r.2)reduction_type use_helpermodules ^/media/conek/DATA/Code/OCR/venv/lib/python3.12/site-packages/torch/_inductor/codegen/triton.pyget_triton_reduction_functionrsE#@@J!+ F'>*!,,>*++ct|tjsyt|tjxs&|jxrt |j dk(S)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_likersI dEJJ ' dEMM *  7C 1 12a7rc|jr|jts|S|jturt |j Sdddfd |S)NcLtt|ddxs t|ddS)N is_BooleanF is_Relational)boolgetattr)nodes ris_predicate_exprz;_materialize_trunc_to_float_expr..is_predicate_exprs* D, . W'$QV2W  rc$|jts|S|jturt|jS|s |j r|St fd|jD}||jk(r|S|j|S)Nc3vK|]0}t|tjr|s|n|2ywN)rrr).0argrrewrite_float_subexprs r zR_materialize_trunc_to_float_expr..rewrite_float_subexpr..s@ #uzz*3DS3I "# &  69)hasrfuncrargsrtuple)rnew_argsrrs rrz?_materialize_trunc_to_float_expr..rewrite_float_subexprsxx #K 99 "+ + T "dooK yy    tyy Ktyy(##r)rz sympy.Basicreturnr)r sympy.Exprrr)is_floating_pointrrrrr)rdtyperrs @@r _materialize_trunc_to_float_exprrsR  " "$((:*>  yyJTYY'' $$ ! &&rc@eZdZUdZiZded<iZded<eddZy) 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_outputsc|j}ttjtjg|j |<||j |<yr)__name__rtorchfloat32float64rr)clsrconvert_outputop_names rregister_upcastzOpDtypeSupport.register_upcasts=--(2EMM5==3Q(RW%'5G$rN)rzCallable[..., str]rrrNone) r __module__ __qualname____doc__r__annotations__r classmethodrrrrrrs1 <>8=')O_)66rrcdtsyddl}t|jjdryy)zd 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_importrs-  #v''):;Erc eZdZdZeej ejgZeejejejgeZ e Dcic]%}|tjt|ddd'c}}}}Ze Dcic]3}|tjt|j#ddd5c}}}}Zed dZed dZed d Zy cc}}}}wcc}}}}w) TritonSymbolszU Stores sympy.Symbol instances and constants associated with triton codegen. offsetTinteger nonnegativeBLOCKrpositivec d}|j}|D]}t|tjr?tj j j|j}|j}nht|tjtjtjtjtjtjfrd}n|j Dcgc]}t||s|}}t#|dk(sJd|j|d} tj j%} dg| z} tj j'D cgc]} t(| | j*k(r| } } t#| dk(sJdt-|j/| d| | dj0<t3| }t5||}|J|Scc}wcc} w)NrrQAmbiguous type: r1z# of Match expected to 1)rrr TMPrOkernelcse varname_mapnameshape UNBACKED_INTSIZEPRECOMPUTED_SIZEINDEXFLOATUNBACKED_FLOAT block_typesrtriton_tensor_ndimactive_range_treesrprefixstrget_block_size tensor_dimrr<)rr expr_shape expr_varsvarcse_var var_shapesymtsymbol_matchessymndimrtree tree_matchs rget_block_shapezTritonSymbols.get_block_shape s&( %% % FCc488,((,,22388<#MM %%II))JJJJ''   &)__"!sD8QD"">*a/N3CCHH:1NN/$Q'xx224 !" ; ; =!#$++5  :!+G-GG+25c6H6HTU6W2Xjm../!%L /z9EJK% FN%%%1"s%G+7G+(G0c4|j|jSr) block_sizesrrrs rrzTritonSymbols.get_block_sizeAstyy))rc4|j|jSr) block_offsetsrrs rget_block_offsetzTritonSymbols.get_block_offsetEs  ++rN)rrrru)rrdrz sympy.Symbol)rrrrrr R0_INDEXR1_INDEXreduction_typesXBLOCKYBLOCKZBLOCKrrSymbolrrupperrrrrr)rrrrs0000rrrs!$--!?@Odkk4;; VoVWK   ellj./v6RVWWM   ell$%%'( .t  K44l**,,Q s *C( 8C0 rceZdZUded<ded<ded<ded<d ed <d ed <dd ZddZddZddZddZe ddZ y)IndexingOptionsr index_strOrderedSet[str] mask_varsz str | None expand_strr _has_rindexrindexzSequence[int | str] | None expand_shapec,t|jSr)rr selfs rhas_maskzIndexingOptions.has_maskSsDNN##rcJt|jtjSr)rr r rrs r has_indirectzIndexingOptions.has_indirectVs"4::txx88rc|jSr)r rs r has_rindexzIndexingOptions.has_rindexYsrc:td|jDS)Nc3PK|]}t|jd yw)tmpNr startswithrmasks rrz.IndexingOptions.has_tmpmask..]sJ43t9''.J$&r}r rs r has_tmpmaskzIndexingOptions.has_tmpmask\sJ4>>JJJrc:td|jDS)Nc3PK|]}t|jd yw)rNrrs rrz,IndexingOptions.has_rmask..`sH3t9'',Hrrrs r has_rmaskzIndexingOptions.has_rmask_sHHHHrc|jr2djttt|jSdS)N & r)r joinsortedmaprrs rmask_strzIndexingOptions.mask_strbs4 =ANNEJJvc#t~~67 8 PV rNrrrr) rrrrrrrrr#propertyr)rrrrrJsTN ,,$9 KI  rrceZdZUdZded<ded<ded<ded <d ed <d ed <d ed<ded<dZded<dZded<ed'dZed'dZ ed'dZ ed'dZ e dd d(dZ d)dZd*dZ d+dZd,d Zd-d!Zd-d"Zd-d#Zd-d$Zd-d%Z d.d&Zy)/BlockDescriptorOptionsz 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]orderrr Sequence[sympy.Expr]broadcast_shape list[bool]broadcasting_dims final_shapeBlockParameters.StrideSorter stride_sorterNzlist[int] | None_boundary_checkFrcan_liftc.|jjSr)r0rrs rrzBlockDescriptorOptions.shapes{{   rc.|jjSr)r0 block_shapers rr?z"BlockDescriptorOptions.block_shapes{{&&&rc.|jjSr)r0stridesrs rrAzBlockDescriptorOptions.strides{{"""rc.|jjSr)r0offsetsrs rrDzBlockDescriptorOptions.offsetsrBr)r<c tjjd fd }||j|_||j|_|j D cgc]} j | d} } t| rd| d<|j| }|j|tjj\}} |jD cgc]} j | d} } |j }|j| }|Dcgc]}tj|}}tjjr%|djdk(sJ|j!dtjj"}tjj$st'|jt'tjj(|z k(rItjj*j-r!|t.j0j2g|zz } t5j6tjj|j}||tjjjA|||||| | | }|jC|||Scc} wcc} wcc}w#t8$r3t;t=t?t'|j}YwxYw) z2Helper to create a BlockDescriptorOptions instancecL|Dcgc]}j|c}Scc}wr)lookup_precomputed_size)exprsrsizevarss r lookup_sizez2BlockDescriptorOptions.create..lookup_sizes"GLMtH44T:M MMs!rQF)stride_sorter_cls shape_envrx) r0r1r3r r8r5r7r:r<)rHzIterable[sympy.Expr]rlist[sympy.Expr])"rOgraphrIrrAr?statically_known_equalsall remove_dimsmaybe_sort_with_stride_order _shape_envrrrno_x_dimrpopnum_reduction_dimsinside_reductionrnumelsfeatures is_reductionrSOner& argsort_symAssertionErrorlistreversedrangerGcompute_boundary_check)rr0r1 range_treesr  get_max_blockrLr<rJdimsingleton_dimsr:strider7r5rr8reduction_ndimr3resultrIs @rcreatezBlockDescriptorOptions.creates77## N#6<<0 $V^^4AG@R@R 9((C/ // OOA 44))FNN#s188??';n'LL!!..0 EGGKK=>9 9K ?%%agg&8&8&..IE GG,,DD_U#+/'   %%m[A C $ S$ ?%FNN(;"<=>E ?s$J-J2J7)8J<<9K87K8cDtj|}t|||iSzN Replaces instances of {symt}_offset with the new expression. rrrHrr replacementrroffsets rreplace_offsetz%BlockDescriptorOptions.replace_offset&  --d3$+ 677rc|tjD](}|j|tjd|}*|SNrrrrsrrrrrs rremove_roffsetsz&BlockDescriptorOptions.remove_roffsets;!11 ED&&tU]]1-=tDD E rc tjj}|Dcic]7}tj|j |t |j 9}}tttjj|}tt|jDcgc]%}|j|j|t j"j$s|r:tjt&j(|j*|j,vsb|j/|j||j*|s|j/|j|t1|j*||sMtjj2r1|j*|tjt&j4k(s|(c}|_ycc}wcc}w)z6List of indices to pass to tl.load(boundary_check=...)N)rOrPrIrrrrr}r(rneeds_yz_grid_overflowrcrrrQrArr]Zeror rr?rstatically_known_multiple_ofrHrVrr;)rrfrerIt block_to_maxneeds_overflow_grididxs rrdz-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 r)r}r?rs rrz!BlockDescriptorOptions.has_rindex;s" ((   rc"|jSr)rrs rr#z BlockDescriptorOptions.has_rmaskAs  rcyrrrs rrz"BlockDescriptorOptions.has_tmpmaskDrrc4t|jSr)rrrs rrzBlockDescriptorOptions.has_maskGsD'')**rcT|j}|j}|rJ|jj|j}|jj|j}t ||D cgc]#\}} | rt j jn|%} }} t||| }|jjs4|s2t| t|k(r|jj| } tjj|xr7t| t|k(xrtfdt | |D} t|jr(| s&d|dtj j#|d}|j} |jjsj|r|jj$n|jj&} d|d| d}|r |jn$|jj|j} t|| |}|Scc} }w)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 c3pK|]-\}}j|dxsj||/ywrQN)rQ)rpre_dimpost_dimrIs rrzGBlockDescriptorOptions.codegen_broadcast_and_reshape..sH&GX00!<G33GXFGs36tl.broadcast_to(, ) tl.trans()r5r7r:revertziprr]r^triton_reshape is_identityrrOrPrIrRr}r index_to_strsort_idxrevert_sort_idx)rvalue initial_shaper8allow_implicit for_storer5r7rgis_broadcastingpre_broadcast_shapesupports_implicit_broadcast old_shape permute_dimsrIs @rcodegen_broadcast_and_reshapez4BlockDescriptorOptions.codegen_broadcast_and_reshapeJs2.. 22 "00778L8LMO $ 2 2 9 9$:P:P Q  ),O=N(O $_+EGGKK 3  um5HI""..'(C ,<< #'"4"4";";[[ [[z88-78?C8 8 2 +2 /2   2 h$ !+aa,a* a  a  a arr.ceZdZdddZy)TensorDescriptorOptionsc,tjj}|jdk7r|d||jdn|d||jd||j d||j g}ddj|dS) a 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)rOrrr1rrAr?r&)rrrrfrs rformatzTensorDescriptorOptions.formats HH ! !''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0  ,DIIdO+>rNTrrrr)rrrrrrrrrs?rrc>eZdZ ddZddZdd dZd dZy) BlockPtrOptionscDtj|}t|||iSrnrorps rrszBlockPtrOptions.replace_offsetrtrc|tjD](}|j|tjd|}*|Srvrwrxs rryzBlockPtrOptions.remove_roffsetsrzrc tjj}g|j}|s|Dcgc]}|j |}}|j dk7r|d||j dn|d||j d||jd||jd||jd||g}d d j|dScc}w) a 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) rOrrrDryr1rrAr?r3r&)rrrrrrDrrs rrzBlockPtrOptions.formats HH ! !!DLL/BIJt++F3JGJ''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0Qtzz]O $qzl #  $DIIdO#4A66KsC"c tj|}|jDcgc]A}|j||||j|tj j |z C}}|Scc}w)av 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. )rrrDrsrr]r})rrrblockradvances radvance_roffsetzBlockPtrOptions.advance_roffsetst**40 ,,  ##FFD9%%feggllDAB   sAA,Nrrrr)rr rr)rrrrsryrrrrrrrs688-78?C8 8 7:rrcrt|trt|tsJ|Dcgc]!}tjj |#}}|Dcgc]!}tjj |#}}||k(r|S|Dcgc] }|dk7s | c}|k7rd|ddj |dSd}g}|D]G} |t |kr| ||k(r|jd|dz }0| dk(sJ|jd I|t |k(sJ|d dj |d Scc}wcc}wcc}w) zUG3tyy'?&@CC C F" ]# # c0B(B MM#  1HC3; ; MM& ! " #m$ $$ $WAdii'( **%JI.s&D*&D/ D4D4ceZdZd!dZd!dZd!dZd!dZd!dZd!dZd!dZ d!dZ d!d Z d!d Z d!d Z d!d Zd!d Zd!dZd!dZd"dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZd!dZ d!dZ!y )# TritonPrinterct|jdk(sJd|j|jddtjj dS)NrQlibdevice.trunc(r).to(rrr_printrOr index_dtyperrs r_print_TruncToIntzTritonPrinter._print_TruncToInt,O499~"""t{{499Q<89qxx?S?S>TTU V rct|jdk(sJ|j|jd}d|d|dS)NrQrz(libdevice.trunc(z) + tl.zeros_like())rrr)rrrs r_print_TruncToFloatz!TritonPrinter._print_TruncToFloat3sF499~""" DIIaL)#5');E7"EErc|jrtt|}|Stjrt j jr|}|Sd|d}|S)N tl.full([], , tl.float64))rrintr# is_fbcoderversionhip)rrrets r _print_FloatzTritonPrinter._print_Float<s\ ??c$i.C     EMM$5$5FC !m4C rct|jdk(sJ|j|jdtddz }|dS)NrQrAtom?z.to(tl.float64))rr parenthesizer )rrrs r_print_ToFloatzTritonPrinter._print_ToFloatHsI499~"""   diilJv,>,D EO$$rc|j\}}|jr3|jr'|j|jdtddz S|j |}|j |}d|d|dS)N % rrz!triton_helpers.remainder_integer(rr)ris_nonnegative stringifyr rrrquotdivquot_sdiv_ss r_print_PythonModzTritonPrinter._print_PythonModNsrII c   3#5#5>>$))UJv4F4LM MT" C 26("UG1EErc|jsJ|j\}}|jr3|jr'|j|jdtddz S|j |}|j |}d|d|dS)N // rrz!triton_helpers.div_floor_integer(z, r)rrrrr rrs r_print_FloorDivzTritonPrinter._print_FloorDivXs~II c   3#5#5>>$))VZ5G#5MN NT" C 26(#eWAFFrcP|j|jdtddz S)N / rr)rrr rs r_print_IntTrueDivzTritonPrinter._print_IntTrueDives#~~dii 60BS0HIIrct|jdk(sJd|j|jddtjj dSNrQlibdevice.floor(rrrrrs r _print_floorzTritonPrinter._print_floorjrrct|jdk(sJd|j|jddtjj dSrrrs r_print_FloorToIntzTritonPrinter._print_FloorToIntqrrct|jdk(sJd|j|jddtjj dSNrQlibdevice.ceil(rrrrrs r_print_ceilingzTritonPrinter._print_ceilingxK499~""" TYYq\!: ;5AUAU@VVWXXrct|jdk(sJd|j|jddtjj dSrrrs r_print_CeilToIntzTritonPrinter._print_CeilToInt}rrc,d|j|dS)Nz tl.sqrt_rn(().to(tl.float32)))rrs r _helper_sqrtzTritonPrinter._helper_sqrtsdkk$/00ABBrc|j|jd}|j|jd}d|d|dS)NrrQzlibdevice.pow((z).to(tl.float64), (z).to(tl.float64)))rr)rrbaseexps r_print_FloatPowzTritonPrinter._print_FloatPowsG{{499Q<(kk$))A,'!&9#>OPPrcD|jdjrdt|jdd}n"d|j|jdd}|jd}|jrdt|d}nd|j|d}d|d|d S) Nrrr(z).to(tl.float64)rQlibdevice.pow(rr)r is_Integerfloatr)rrr exp_valr s r_print_PowByNaturalz!TritonPrinter._print_PowByNaturals 99Q< " "!% ! "5!6mDDt{{499Q<011ABD))A,    w 0 >Cdkk'*++; or < rQrr"Nc3(K|] }d|d yw)rrNrrrNs rrz6TritonPrinter._print_min_max_helper..s.!q1X.s)>Q ;;tyy|, ,$))n!4j KKTYYt_- . KKTYYst_- ..1v..1j C$r@rBrDrFrHrPrrrrr+s F % F GJ   Y Y CQ/"* C,55; N O O N O O N O O O   Qrrc*tt|S)zCConvert torch.dtype to triton type and upcast [b]float16 to float32)rIrKrs rtriton_compute_typerSs *51 22rc^|tjk(rtj}t|S)z@Convert torch.dtype to triton type, with fix for storing tl.bool)rrint8rIrRs rtriton_store_typerVs"    u rct|r+|jr|jdkrtjSt |S)z0Implicit upcasts used for Triton reduction types)r is_signeditemsizerint32rKrRs rupcast_acc_dtyper\s05??u~~7J{{ u %%rc*tt|S)z:Convert torch.dtype to triton type, with reduction upcasts)rSr\rRs rtriton_acc_typer^%s /6 77rc<|jdkxr |jS)Nr")rZrrRs rlow_precision_fpr`*s >>Q  :5#:#::rct|tsy|j}t|tjr t |SdSr)rrWrrr`)rrs rlow_precision_fp_varrb.s6 c; ' IIE&0 &D E "O%Orct|tr |jSt|tjj rt t|Syr)rrWrr _prims_commonNumberrr"rs rtriton_arg_dtyperg6s>#{#yy#u**112T#Y'' rctjj xr>t|txr,|j t jt jfvSr) r#rcodegen_upcast_to_fp32rrWrrfloat16bfloat16rfs rneeds_upcast_to_float32rl>sD MM 0 00 9 sK ( 9 II%--8 8rc<eZdZ d dfd ZdZxZS)TritonCSEVariablecnt|||||t|_|Jd|Jdy)Nrz!TritonCSEVariable must have dtypez!TritonCSEVariable must have shape)super__init__rr )rrboundsrr __class__s rrrzTritonCSEVariable.__init__GsH vuE:*4, E"EE  E"EE rcF|D]}t|tr&|jj|j9t|tj sTt jD]6}t||s|jjt|dgy)Nr) rrnr updaterrrrrr)rrrkwargsrrs rupdate_on_argsz TritonCSEVariable.update_on_argsTs C#01%%cmm4C.*55D%c40--*T2B1C4/H.IJ rr) rrrszValueRanges[Any]r torch.dtyperrurr)rrrrrrx __classcell__rts@rrnrnFsH !% F F! F F  F  F rrncddlm}|S)Nrrq)!torch._inductor.dtype_propagationrrrqs rget_dtype_handlerr~bsL % ''rTc dddfd }|S)z Codegen helper to upcast arguments to float32, depending on the config and dtype. This decorates tl.math/libdevice codegen functions. c,t|rdnd}||S)N.to(tl.float32)rrl)r upcast_strings rmaybe_upcast_argz.maybe_upcast_float32..maybe_upcast_argns -DS-I)r }o&&rcFtjdfd }|S)Nc|Dcgc] }| }}|jDcic]\}}||}}} |i|} xr4tdtj||j D}|sdn#t t j|i|} | tjdfv} | r| dt| dnd} || Scc}wcc}}w)Nc32K|]}t|ywrrrrs rrzKmaybe_upcast_float32..decorator..wrapped..}s6(,6.to(rr) itemsr} itertoolschainvaluesrr~rrrrI)rrwr upcast_argskeyval upcast_kwargsrkany_needs_upcast result_dtypeneeds_downcastdowncast_stringrrrs rwrappedz8maybe_upcast_float32..decorator..wrappedvs<@AS+C0AKAHN WHCS"23"77WMW;8-8F- #6$??4A63  (@W.0$--@$Q&Q  *%--1FFN"l&>{<013  Xo./ /)BWs CCr+)rr)rrrrs` r decoratorz'maybe_upcast_float32..decoratorrs&&t^< 00rr+)rCallable[..., Any]rrr)rrrs` @rmaybe_upcast_float32rhs '< rc,eZdZdZej ej Ze dS dTdZ edUdZ edZ e dZ e dVdZeed Zed Zed Zed Zeed ZeedZeedZeedZedZedZedZedZedZedej<dddddZeedZ eedZ!e dZ"edZ#eedZ$eedZ%eedZ&eed Z'eed!Z(eed"Z)eed#Z*eed$Z+eed%Z,eed&Z-eed'Z.eed(Z/eed)Z0eed*Z1eed+Z2eed,Z3eed-Zed.Z4eed/Z5ed0Z6ed1Z7ed2Z8ed3Z9ed4Z:ed5Z;ed6Zed9Z?ed:Z@ed;ZAed<ZBed=ZCed>ZDeed?ZEeed@ZFeedAZGeedBZHeedCZIedDZJeedEZKe dFZLeedGZMeedHIdJZNeedHIdKZOeedLZPeedMZQedNZRedOZSeedPZTedQZUeedRZVy)WTritonOverrideszEMap element-wise ops to Triton e.g., ops.to_dtype(x,...) -> x.to(...)NTcZ dd}|>t|||tjjtj_|tj k(rd|dS|tj k(r| |js||dS|r t|}n t|}|d|dS)Nc||k(rytjtjf}||vr||vr ||k7rJd|tjk(s|tjk(ry|tjk(s|tjk(ryy)NrzCConversions between float8_e5m2 and float8_e4m3fn is not supported!rXr")r float8_e4m3fn float8_e5m2) src_dtype dst_dtype fp8_dtypess r_get_min_elements_per_threadz>TritonOverrides.to_dtype.._get_min_elements_per_threadsI% ##!!J Z'+* UU  U E---e>O>O1OE///9@S@S3Srrz != 0)z.to(tl.int16).to(tl.uint8)rr)rryrryrr) r~rOrmin_elem_per_threadrruint8rrSrV)rNrruse_compute_typesr out_dtypes rto_dtypezTritonOverrides.to_dtypes " /:   6  ,/,Y>,,,AHH ( EJJ qc= ekk !  !i&A&AYEV S23 3 +E2I)%0ID 1%%rc|j|jk(sJ|j|k7r|dt|d}|dt|d}t||k7r|dtt|d}|S)Nrr, bitcast=True))rZrrIrK)rNrrouts rto_dtype_bitcastz TritonOverrides.to_dtype_bitcasts!!U^^333 77i #T+i013A4 E*+?; u % .Ek*=e*DEFaHC rc Ttjj|}t||}t |}|dk(r5t j d|dkr|dk(r d|d|dS|dk(r d|d|dS|dkr#|jsd |d d}d|d |d |d |d Sd|d |d |d S)Nrg?z tl.float32tl.full(z, 0x80000000, tl.uint32).to(rz tl.float64z$, 0x8000000000000000, tl.uint64).to(ztl.rXrrr)rrd dtype_to_typercrSmathcopysignrY)rrrtype_ triton_valrItriton_signed_types r_shaped_constantz TritonOverrides._shaped_constants##11%8"5<0 )%0  A:$--U3a7l*!%(D[MQ`aa ,!%(L[MYhii 19U__#&{12&7!8 eWBzl"5G4Hk]Z[\ \eWBzl"[MC Crc*|j||gS)Nrp)r)rrrs rconstantzTritonOverrides.constants##E5#;;rct|tjjr|j ||St ||k(r|Sd|dt |dS)Nrrr)rrrdrerrgrI)rrrs r_cast_libdevice_argz#TritonOverrides._cast_libdevice_argsX c5..55 6<<U+ + C E )UO3%u[/022rcd|dS)Nr/rrrNs rabszTritonOverrides.abs aS""rct|dd}t|dd}|tjk(r7|tjk(r$tjj r d|d|d}n d|d|d}t |s t |rMtj||}|tjtjfvr|dt|d}|S)Nrtriton.language.div_rn(rrrrr) rrrr#eager_numericsdivision_roundingrbr~truedivrjrIrNyx_dtypey_dtyperrs rrzTritonOverrides.truedivs!Wd+!Wd+ u}} $5==(%%77,A3b15CaSA3a.C  "&:1&=)+33Aq9IU]]EMM::T+i"8!9; rcrt|dd}t|dd}|tjk(r|tjk(r d|d|d}n d|d|d}t|s t|rMt j ||}|tj tjfvr|dt|d}|S) z Division with round-to-nearest rounding mode. Always uses triton.language.div_rn for float32 inputs to match eager CUDA behavior. rNrrrrrr)rrrrbr~rrjrIrs rdiv_rnzTritonOverrides.div_rn)s !Wd+!Wd+ emm #5==(@+A3b15CaSA3a.C  "&:1&=)+33Aq9IU]]EMM::T+i"8!9; rcd|d|d}t|s t|rMtj||}|tjtj fvr|dt |d}|S)Nrrrr)rbr~modrrjrrI)rNrrrs rrzTritonOverrides.mod?sl!Cs!n  "&:1&=)+//15IU]]EMM::T+i"8!9; rc:tjrd|dSd|dS)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 rr zTritonOverrides.expHs+   !!A& &#A3a( (rcd|dS)Nzlibdevice.exp2(rrrs rexp2zTritonOverrides.exp2X!1%%rcd|dS)Nzlibdevice.expm1(rrrs rexpm1zTritonOverrides.expm1]"!A&&rcd|dS)Nz tl.sqrt_rn(rrrs rsqrtzTritonOverrides.sqrtbsQCq!!rctjj}|dk(ry|dk(r d|d|dS|dk(r|dS|8tjtj d t j|Std |) 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_ONLYopsmaximumrrr[r`)rNbugs rreluzTritonOverrides.reluhsmm88 / !# O #8s:YZ[Y\\]^ ^ J S:  [;;s||Au{{;Q? ? LSGT rcZtjjr d|d|dSd|d|dS)Nz tl.minimum(r, tl.PropagateNan.ALL)ztriton_helpers.minimum(rrrrr%r&s rminimumzTritonOverrides.minimumz< ==   2aS(>? ?,QCr!A6 6rcZtjjr d|d|dSd|d|dS)Nz tl.maximum(rrztriton_helpers.maximum(rrrs rrzTritonOverrides.maximumrrcd|d|d|dS)Nrrrr)r%r&rs rwherezTritonOverrides.wheres1#Rs"QCq))rc  tjjsJ||}}d}d}||r||}||r||} d  fd }ttjj dk\sJdt t jtj t t jtj t t jtj t t jtjtjjjtjj||t!|j" g|j$ f}tjjjtjj||t!|j" g|j$ f}t&j(j*j,j.dk(rd}nd}d |d |d |d S)a 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. cxttt|jsytj j d}|jsJtj j|rytj jjjD]\}}||k(s d|vsd|vsyy)NFrKztl.loadz other=0.0T) r}r(rEr rOrrer\_has_constant_maskr_cacher)rreduction_rangekvs ris_where_neededz,TritonOverrides.dot..is_where_neededss. >?hh2226O"// //xx**?;  ++113 !18 Q;!3C  !rctjjd|j}tj j Dcgc]}|jr|jd}}t|dk(sJdtj|d||}tj jjtj j||j|jScc}w)NdotrrQz'don't tile reduction when native matmulrrr)r$ Reduction default_valuerrOrrer\rrTritonKernelOverridesrrgeneratecomputer)rdefaultrreduction_mask where_vars r where_condz'TritonOverrides.dot..where_condsll00 BGHH00$$;;-t$N ~&!+ V-V V+-33N14EsGTI88<<((  )399CII) s"C3c |vr|Dcgc] }| k(rdn| }}|gk(rE|vsJdddg}|vr|d<|vr|d<t|||}|gk7}|r d|ddd}|S|gk(rK|vsJd ddg}|vr|d<|vr|d<t|||}d |d }|gk7}|r d|ddd}|Stcc}w) a  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 xrrQrz, (rrz'right tl.dot operand cannot depend on yrr)rNotImplementedError) rrr8rgshape_2dbroadcast_neededRBLOCKrrrs r#reshape_transpose_broadcast_for_dotz@TritonOverrides.dot..reshape_transpose_broadcast_for_dots|6&JW X3v 3!> X Xvv..]2<2 :]*"(HQK]*"(HQK'umXF$,/?#? #.ugS6("ME4L1 00]2=2 :]*"(HQK]*"(HQK'umXF$E7!,$,/?#? #.ugS6("MEL*)W!YsCrztl.dot can only do mm and bmmrtf32ieeeztl.dot(rz, input_precision=""))rr4r8r4rr)rOris_native_matmulrdense_size_listrrrr rrrrrrrrarrrbackendscudamatmulfp32_precision) r%r&orig_aorig_brrrinput_precisionrrrrs @@@@rrzTritonOverrides.dots:xx((((A * 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 3O3DBGGrrQ) constraintsris_purepack input_dtypesc t||#djdg|Dcgc]}dc}z}|jdDcgc]}|j} }| Dcgc]}|j dr|} }g} t t || dt|D]o\} \} }|dk(rH|Ft| tr6| j|| k7r$| j| dt|| dV| jt| qtjjr7t!|j#d d t!|j#d d n d |d d |d fd }d kr|dj| S|d}t$j&j(}t$j&j*}|j-||j.}djfd| D}|j1|d|||j1|d|d|dd|Scc}wcc}wcc}w)Nrz=rr",=hrr\z\\r c .ddd|dddd S)Nztl.inline_asm_elementwise(rrz ], dtype=z , is_pure=z, pack=rr)r asm_literalasm_triton_typeconstraints_literalrrs rasm_callz8TritonOverrides.inline_asm_elementwise..asm_call{s>,[M.s& AD-cU"TF! < s = z$ = triton_helpers.inline_asm_unpack()rIr&splitstripr enumeraterrrrWrrrrrrreprreplacerOrrrnewvarr writeline)asmrrrrrinputs_rconstraint_partsinput_constraints cast_inputsirrr first_inputrrrk packed_argsrrrs `` @@@rinline_asm_elementwisez&TritonOverrides.inline_asm_elementwiseKsN&e,  ))TF6-Bac-B$BCK 0;/@/@/EF!AGGIFF(8R1 S@QQRR $S1B=S[1Q%RS -KAxQS ,sK0IIa0""cU${<?/K.LA#NO""3s8, - ==  s)++D&9K"&{"3";";D&"I cU!*K$%k]!"4    19DIIk23 3Qi ((""hhll%{/@/@Aii HS   VHC(='>?@h:6("[MQSTXSYYZ [  m.C GRs I I &I=Icd|dS)Nz tl_math.cos(rrrs rcoszTritonOverrides.cosrrcd|dS)Nz tl_math.sin(rrrs rsinzTritonOverrides.sinrrctd)Nz/ops.index_expr not implemented outside a kernelr)rrrs r index_exprzTritonOverrides.index_exprs!"STTrctd)Nz+ops.masked not implemented outside a kernelr5)rbodyothers rmaskedzTritonOverrides.maskeds!"OPPrcd|dS)Nzlibdevice.lgamma(rrrs rlgammazTritonOverrides.lgamma#1#Q''rcd|dS)Nzlibdevice.erf(rrrs rerfzTritonOverrides.erf s!$$rcd|dS)Nzlibdevice.cosh(rrrs rcoshzTritonOverrides.cosh!1%%rcd|dS)Nzlibdevice.sinh(rrrs rsinhzTritonOverrides.sinhrCrcd|dS)Nzlibdevice.acos(rrrs racoszTritonOverrides.acosrCrcd|dS)Nzlibdevice.acosh(rrrs racoshzTritonOverrides.acosh"!A&&rcd|dS)Nzlibdevice.asin(rrrs rasinzTritonOverrides.asinrCrcd|dS)Nzlibdevice.asinh(rrrs rasinhzTritonOverrides.asinhrJrcd|d|dS)Nzlibdevice.atan2(rrrrNrs ratan2zTritonOverrides.atan2"!Bqc++rcd|dS)Nzlibdevice.atan(rrrs ratanzTritonOverrides.atanrCrcd|dS)Nzlibdevice.atanh(rrrs ratanhzTritonOverrides.atanhrJrcd|d|dS)Nzlibdevice.copysign(rrrrPs rrzTritonOverrides.copysigns%QCr!A..rcd|dS)Nzlibdevice.erfc(rrrs rerfczTritonOverrides.erfcrrcd|dS)Nzlibdevice.erfinv(rrrs rerfinvzTritonOverrides.erfinvr=rcd|d|dS)Nzlibdevice.hypot(rrrrPs rhypotzTritonOverrides.hypotrRrcd|dS)Nzlibdevice.log10(rrrs rlog10zTritonOverrides.log10rrcd|dS)Nzlibdevice.log2(rrrs rlog2zTritonOverrides.log2rrcd|d|dS)Nzlibdevice.ldexp(rz.to(tl.int32))r)rNns rldexpzTritonOverrides.ldexps"!Bqc88rcd|d|dS)Nzlibdevice.nextafter(rrrrPs r nextafterzTritonOverrides.nextafter s&aS1#Q//rc|d|SNr%rrs r logical_andzTritonOverrides.logical_andCs|rc |dS)Nz == 0rr%s r logical_notzTritonOverrides.logical_notsE{rc|d|SNz | rrs r logical_orzTritonOverrides.logical_orrjrcd|d|dS)Nr ^ rrrs r logical_xorzTritonOverrides.logical_xor!s1#S1~rc|d|Srhrrs r bitwise_andzTritonOverrides.bitwise_and&Cs|rc d|S)N~rrls r bitwise_notzTritonOverrides.bitwise_not*s 1#wrc|d|Srorrs r bitwise_orzTritonOverrides.bitwise_or.rvrc|d|S)Nrrrrs r bitwise_xorzTritonOverrides.bitwise_xor2rvrc|d|S)Nz << rrs rbitwise_left_shiftz"TritonOverrides.bitwise_left_shift6D}rc|d|S)Nz >> rrs rbitwise_right_shiftz#TritonOverrides.bitwise_right_shift:rrc d|d}d|d|dS)Nr).to(tl.uint32)ztl.rand(rrrseedrs rrandzTritonOverrides.rand>s%VHO,$r&++rc nd|d}d|d|d}d|d|d|d}d|d|d}d |d |d |d |d |d S)Nrrz((z)*(rz)//(z//4))z)%(z!triton_helpers.rand_eager_kernel(r+z, VEC=rr) r base_offsetthreads_per_roundtidvectid_u32denomr" tid_truncs r rand_eagerzTritonOverrides.rand_eagerCscU/*SE./r2 eWCuE 2 UG2. 24&;-q2i[X^_b^ccdeerc d|d}d|d|dS)Nrrz tl.randn(rrrrs rrandnzTritonOverrides.randnMs%VHO,4&6(!,,rc ,d|d}d|d|d|d|d S)Nrrztriton_helpers.randint64(rrr)rrlowhighs r randint64zTritonOverrides.randint64Rs1VHO,*4&6("SED6KKrctd)Nz.ops.load_seed not implemented outside a kernelr5)rrs r load_seedzTritonOverrides.load_seedWs!"RSSrcNtjjrd|dSd|dS)Nz tl.rsqrt(rzlibdevice.rsqrt(rrs rrsqrtzTritonOverrides.rsqrt[s/ ==  qc# #%aS* *rcd|dS)Nzlibdevice.log1p(rrrs rlog1pzTritonOverrides.log1pdrrcd|dS)Nzlibdevice.tan(rrrs rtanzTritonOverrides.tanir@rcRtjjjj |}|rt |dr |j }nd}tjrBtjjr(tdkDr|tjk7r|d|dSd|dS)Nr)rzlibdevice.fast_tanhf(rzlibdevice.tanh()rOrrrgetrrr#rrrrrr)rNrrs rtanhzTritonOverrides.tanhos((,,**..q1 ww0MMEE   !!"$v-&!+1#Q/ /$QCq) )rcd|dS)Nz tl.sigmoid(rrrs rsigmoidzTritonOverrides.sigmoidsQCq!!rcd|d|d|dS)Nz(libdevice.signbit(z ) != 0) if (z).dtype is tl.float32 else z < 0rrs rsignbitzTritonOverrides.signbits# "!L3NqcQU V rcd|d|dS)Nzlibdevice.fmod(rrrrs rfmodzTritonOverrides.fmods!2aS**rctj||}|mt|rb|j||}t |t j jr |j|t jn|}d|d|dSt|xs t|}|}|t jt jfvr-t|s|rt jnt j}|j||}|j||} d|d| d} |6||k7r1t|r|r| dt|d} | S| dt|d} | S)Nztriton_helpers.pow_integer(rrrr)r~powrrrrrdrerint64rlrrr`rI) rr%r&rr exponentr pow_dtypecast_acast_brks rrzTritonOverrides.powsg(*..q!4  #(8(F**1l;Da!4!4!;!;< Q ,s  1b !D D215S9PQR9S U]]EMM: : $L15E ]]  ((I6((I6!&F815  # (A -# &xtK ,E+FaHF #84 L(A'B!D rcd|dS)Nz tl_math.log(rrrs rlogzTritonOverrides.logrrF)rcd|dS)Nzlibdevice.isinf( ).to(tl.int1)rrs risinfzTritonOverrides.isinf"!M22rcd|dS)Nzlibdevice.isnan(rrrs risnanzTritonOverrides.isnanrrcd|dS)Nzlibdevice.nearbyint(rrrs rroundzTritonOverrides.rounds&aS**rcd|dS)Nrrrrs rfloorzTritonOverrides.floorrJrctjdtj}tjdtj}tj||}tj |||}tj ||}tj |tj|||}tj |tj|||}tj ||}tj |tj||}tj||}tj |tj||}tj |||SNrrQ) rrrr[eqrltsubrytruncdiv)r%r&zerooneb_zerob_nega_negrs rfloordivzTritonOverrides.floordivs||Au{{+ll1ekk* 4 IIfc1 %q$ IIeSWWT1-q 1 IIeSWWT1-q 1q$ IIeS__Q/ 3||Aq!yy 5t<yyt,,rcftjdtj}tjtj ||tj }tjtj ||tj }tj||}|d|dS)Nrr.dtype))rrrr[rrrUr)rNzleftrightrs rsignzTritonOverrides.signsw LLEKK (||SVVAq\EJJ7 cffQlUZZ8ggdE"d1#W%%rcd|dS)Nrrrrs rtrunczTritonOverrides.truncrJrc|d|S)Nrrrs rrzTritonOverrides.truncdivs D}rcd|dS)Nrrrrs rceilzTritonOverrides.ceil rCrNT)rryrtorch.dtype | None)rryrryrryrr)Wrrrrrrae_LOG_2_E staticmethodrrrrrrrrrrrr rrrrrrrrrrr/r1r3r6r:r<r?rBrErGrIrLrNrQrTrVrrYr[r]r_rdrfrirmrprsruryr{r}rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrsOtyy H)- :&:&&:&:&x DD,<<33## ,* ) )&&''"""77 77 **xHxHtmm CCJ####UUQQ((%%&&&&&&''&&'',,&&''//&&((,,''&&9900,,ff--LLTT++ ''%%**&""  ++  D##/303/303++''--6&&'' &&rrrceZdZdZfdZeejdZedZ edZ e dZ e dZ e dZe d d ZxZS) ra Map 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. cDt||i||jyr)rqrr_setup_libdevice_routing)rrrwrts rrrzTritonKernelOverrides.__init__ s# $)&) %%'rcddlmtjjj D]}t ||sJt||}fd}|dk(rDt dsJtj|||}||_ t||t|kd}tj|||}||_ t||t|y)z.decomposition_router2s977emm+)!,,>7#3X>qAGGGrr)rrc^|jtjk(r d|d|dS||S)Nz libdevice.rr)rrr)rNrrs r dtype_routerzDTritonKernelOverrides._setup_libdevice_routing..dtype_routerAs277emm+'z1#Q77)!,,rN) torch._inductor.codegen.commonrr _inductorr&op_requires_libdevice_fp64rr functoolspartialrsetattrr)rfn_name original_implrfnrrs @rrz.TritonKernelOverrides._setup_libdevice_routing's D,,GG 4G3( ((#C1M H )#/;;;&&(QX& Wl2&67 - ""]WB"BK C,r"2 3; 4rcrtjj}dg|z}|j|||S)NrQrp)rOrrr)rrrrrs rrzTritonKernelOverrides.constantMs9 xx**,d ##E5#>>rct||}tjj|dd}t |t sJ|j r |j }ntj|j}tjj}|tjtjfvr|n|}tjj } dtj_tjj"j%tjj&|j(t+|||}|tj_|tjtjfvritjj"j%tjj&|j-||t/||j0}n|}|j2D]l}t5|t6j8stj:|tjj"j<|j>j@}n||k7r_tjj"j%tjj&|j-||||j0}|jB|_!|S#|tj_wxYw)NF block_ptrtma_compatibility_checkerrsrrr)"rrOrindexingrrr rrr get_index_dtype_as_torch_dtyperr[rr# test_configsruntime_triton_dtype_assertrrrrr@rrKrrrr r promote_typesrrrr ) rrrrrrorigr index_vars rr6z TritonKernelOverrides.index_exprVs=/e<88$$ ET% (O444  ))E!11(..AEhh==? u{{EKK&@@k"">> C>CF   ;((,,''  "",T2 (C?CF   ; ekk2 2((,,''   S%()%0ii (C E!..  !)TXX6!//qxx||77 GMME   #hhll++HH$$LLk2%)) ,!**  =?CF   ;s A.KK1c |ztjj`tjj j tjj|dtj|j}|jjd}|sJdd}|D]>}|jD]-}|jdk7st|jds+d }>@|rdn|}tjj|| 5}|} ddd|r j j"r t|}tjj j tjjd | d t%|d | dt'j(|| j*| j}t-j.| |} n } | j0j3| S#1swYxYw)N .to(tl.int1)routput)opz)graph for body does not contain an outputFloadrQT)rrz.shape, rrr)rrrrOrrrrrrrP find_nodesrtargetrm mask_loadsrsis_boolrcr!wraprrrr discard) rr8r9nodes need_whererrrnew_maskrkrs rr:zTritonKernelOverrides.maskeds   1 1 =88<<((  & %jjjj )D %%%2AAAu  Dyy ::'+CCHHQK+P!%J   # XX U 3 xVF  }}$$U HHLL))  6((=+?*@6('R"''.llll *E))Hfe4CC h' )  s G88Hctjjj|}d|dtjjj d|dS)Ntl.load( + load_seed_offsetr)rOrrinput seed_offset)rrrs rrzTritonKernelOverrides.load_seedsIhhmm!!$'se3qxx}}889KVTUUV W rc0d|d}tjjj|x}r|Stjjj |j |j }tjjj tj|j }tjjj|d|d|dtjjj|||f||fS)Nzfrexp(rrrz = triton_helpers.frexp() rOrrtry_getr$rrrr[rr%put)rN cache_keycse_valmantissars rfrexpzTritonKernelOverrides.frexpsQCqM hhll**95 57 5N88<<&&QWWAGG&D88<<&&U[[&H ""j8*$?$+1-$D *U$s  C Cc|jSr)r4rs rrz(BlockParameters.StrideSorter.is_identity+ s$$ $rcy)zBCreate a `StrideSorter` that can be used to sort block parameters.Nrrr1rMs rrlz#BlockParameters.StrideSorter.create/ srcb|js|jDcgc]}|| c}S|Scc}wr)rrrattrr,s rsortz!BlockParameters.StrideSorter.sort6 s.##)-7AQ77K8 ,cb|js|jDcgc]}|| c}S|Scc}wr)rrr<s rrz#BlockParameters.StrideSorter.revert; s0##)-)=)=>AQ>>K?r?Nr1zlist[int | sympy.Expr]rMrsrr9)rrrr dataclassesfieldrr7r,rrrrlr>rrrr StrideSorterzBlockParameters.StrideSorter s##%6[%6%6E%BB   %  %   U#9 UFN U ) U   U    rrDc<eZdZfdZe ddZxZS)$BlockParameters.IdentityStrideSorterc"t|yrrqr7rrts rr7z2BlockParameters.IdentityStrideSorter.__post_init__B  G ! #rc L||ttt|S)Nr1r)rarcrr:s rrlz+BlockParameters.IdentityStrideSorter.createE s'!1eC(8$9:; rrA)rrrr7rrlrzr{s@rIdentityStrideSorterrF@ s4 $  #9 FN  )   rrMc@eZdZdZfdZe ddZxZS)+BlockParameters.TensorDecriptorStrideSorterzT Sorts BlockParameters dimensions with strides in descending order. c"t|yrrHrIs rr7z9BlockParameters.TensorDecriptorStrideSorter.__post_init__T rJrcttt|} tj||d}|||S#t $r|}YwxYw)a 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)reverserL)rarcrr&r_r`)rr1rM identity_sortrs rrlz2BlockParameters.TensorDecriptorStrideSorter.createW sa!s+;'M ) !,,Y8HRVW !1!  " )( )sA AArA)rrrrr7rrlrzr{s@rTensorDecriptorStrideSorterrON s9  $  #9 FN  )   rrTc t|}td||fD\}}|di|Dcic]}|||||zc}Scc}w)z0 Concatenates block parameters. c3FK|]}tj|ywr)rBasdictrs rrz*BlockParameters.__add__..y sBq[''*Bs!r)r"r)rr9rr%r&rs r__add__zBlockParameters.__add__t sR4jBT5MBB19a8sc1S6AcF?*8998sA c |j|j|}tditj|j Dcic]\}}||j |c}}}||fScc}}w)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. )rMr)rlrAr/rBrWrr>)rrLrMr:rrr0s rrTz,BlockParameters.maybe_sort_with_stride_order| s~*000S   !, 2 24 8 > > @C]'',,  }$$ s A1c fd}tditj|jDcic]\}}|||c}}Scc}}w)zA Remove dimensions where removable_dims is True. cRt|Dcgc] \}}|s| c}}Scc}}wr)r)ititem is_removableremovable_dimss r filter_dimsz0BlockParameters.remove_dims..filter_dims s3+.b.*A&D,# s#r)r/rBrWr)rr_r`rrs ` rrSzBlockParameters.remove_dims sO   5@5G5G5M5S5S5UVcsK$$V  VsA N)r9r/rr/)rLztype[StrideSorter]rMrsrz4tuple[BlockParameters, BlockParameters.StrideSorter])r_r6rr/)rrrrrBrCrarrr?rArD dataclassrDrMrTrXrTrSrrrr/r/ s0k//EE E$5K$5$5d$KK!K 1 1 1$ GG G 1 1 1$ GG G&&&P |  #l##J:%!3%@H% =%" rr/c*eZdZdZdZddZdZdZy)"CooperativeReductionWorkspaceCachez The scratch space used for cooperative reductions can be reused after two reduction loops. This keeps track of what can be reused. c||_g|_g|_tjtj |_d|_d|_yrv) r current_loop prior_loop collections defaultdictdequeready_for_reuse loop_count store_count)rrs rrrz+CooperativeReductionWorkspaceCache.__init__ s@ *66{7H7HIrc|jj|}|r|jS|jj |d\}}}|j j |||f||fSr)rjrpopleftr workspacerer)rnbytescachedws_namer( ws_offsets rallocatez+CooperativeReductionWorkspaceCache.allocate sk%%))&1 >># # $ 3 3FE BI   &'9!=>##rc|jD]&\}}}|j|j||f(|j|_g|_|xjdz c_yNrQ)rfrjrrerk)rrprrrss r on_loop_endz.CooperativeReductionWorkspaceCache.on_loop_end s_*.// F &FGY   ( / /)0D E F++ 1rcH|j}|xjdz c_|Srv)rl)rpriors rincrement_store_countz8CooperativeReductionWorkspaceCache.increment_store_count s#   A rN)rpr)rrrrrrrtrwrzrrrrcrc s $rrcc$eZdZUded<dZdZy)FixedTritonConfigzdict[str, int]r#c |j|Srr#rr]s rr,zFixedTritonConfig.__getitem__ s{{4  rc||jvSrr~rs r __contains__zFixedTritonConfig.__contains__ st{{""rN)rrrrr,rrrrr|r| s !#rr|ceZdZdZddZy) TritonCSEz Subclasses CSE to apply the current load mask to the cache key to avoid CSEing variables across separate masked blocks. cZtjjx}r||jfS|Sr)rOr _load_maskr)rrrs r augment_keyzTritonCSE.augment_key s,88&& &4 &tyy) ) rN)rrrzstr | tuple[str, str])rrrrrrrrrr s  rrcdeZdZUdZded<ded<ded<ded<d Z dd Z dd Zdd Zy )TMACompatibilityCheckerzO Checks if the TMA API can be used for load / store triton operations. TritonKernelrryrrrforcecd|_y)Nz2Cannot use TMA descriptor for load / store since: )failed_debug_prefixrs rr7z%TMACompatibilityChecker.__post_init__ s #W rcB|jrytjjjdk(r4t j jddk\rtjs+tjjjdk(r$tjjr ts!tjd|jy|j r7|j"j$r!tjd|jyy) NTrr xpuz}%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.)rrOrPget_current_device_or_throwr"rrget_device_capabilityr#assume_aligned_inputsruse_tensor_descriptorrrdebugrrrrVrs r can_use_tmaz#TMACompatibilityChecker.can_use_tma s ::GG779>>&H 88:1=B4477668==F 33)+ II[((   >>dkk22 IIA(( rc  |jrA|jDcgc]+}tjjj |-}}n |j}tjjj |dtjds"tjd|j|y|jj}|ddD]}tjjj t||zdtjdtjdrbtjd|j||y|jd}tjjj |tjdr,tjd |j|jyd}d}|j D])} t"j$D]} t'| | s| }| })+|r|sJ|d t"j$|j(j*r;|j,s.|t"j.vrt0|} d} |j(j2D],} | j4s| j6| k(s | j8} n| -tjd |j| |jy|j(j;| }|j=||i|z}tjjj?|tjds-tjd |j|j|yy d dd }||zdz }|jAtB|jAt|}tEtGtjH||d}||j(jKt0|kDr#tjd|j||y|j(jM|} |j(jNrW||j(jN| kDrtjd|j| |j(jN| |ytQ||j(jRjU| d|j(jR| <ycc}w#tV$r.tjd|j|jYywxYw)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. rKrQz=%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 zN%s could not find reduction range tree for innermost prefix %s Block shape: %szj%s persistent reduction innermost block shape cannot load 16 bytes. Block shape: %s, persistent RBLOCK: %dc||z }|r||z}|Srr)rNrrrs rindexing_div_repzQTMACompatibilityChecker.are_block_parameters_compatible..indexing_div_rep s a%C!AgJrzC%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)rNrrrrzsympy.Expr | Nonerr),rrArOrPrI!replace_backed_symbols_with_hintsrQrrrrrrrZrr?rrrrrpersistent_reductionrrrrer\rnumel_get_persistent_RBLOCKsubsstatically_known_geqr#rr6rnsolve max_blockr fixed_configr~tma_min_block_sizesrrL)r block_paramsstrA element_sizeriinnermost_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_compatiblez7TMACompatibilityChecker.are_block_parameters_compatible s ::'..  BB2FG #**Gww77 U]]STEUV IIO((  zz** crl F77##;; 5q%--:KL a  k,,    "!- 8 8 < 77   3 3 !5==#3  IIP((((  ##3@@ N+77  !.*=+9(+5(   $(< $%%QR_RkRkQl m < KK , ,NN$ (E(EE%//C$D !J[[,, >>ahh2G&G!"J !  d,,) ,,   $ B B: N %**,@BS+TU "77##88%u}}R'8 A,, ,,%  TIB  ,0!!) 3\ABF (2(:(:.)'/+;<&"1 10""DKK$9$934%II]00-&  !!%!9!9:N!O;;++% (@(@(PP < 44* KK44^D*  %GJ& 77;;NANGDKK33NCkZ  U,, ,,   s&0TB/T3BT;AT4T?>T?c|jS)aH 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. )rrs rr<z TMACompatibilityChecker.can_lift szzrNr*)rr/rr) rrrrrr7rrr<rrrrr sS  O KX& &P% B rrc^eZdZUdZeZded<eZded<dZ e Z e Z ded<eZd ed <d Zd ed < d` dafd ZedbdZedbdZdcdZdbdZdZdZdZdZdbdZdZedddZd dd dd d dedZ df dgdZ dfdZ! dhdZ"d Z#d!Z$d"Z%ed#Z&dd$d%Z'd&Z(did'Z) djd(Z*dkd)Z+ dl dmd*Z,dnd+Z-d,Z.dod-Z/ dp dqd.Z0ddd/Z1drd0Z2 dsd1Z3 dtd2Z4 dud3Z5dud4Z6d5Z7d6Z8d7Z9d8Z: dvd9Z; dwd:Z< dxd;Z= dyd<Z> dzd=Z?did>Z@d?ZAd{d@ZBd|dAZCdBZDdCZEeFeGd dddDZHeFdEZIeFdFZJdldddGZKedHZLedIZMdJZNd}dKZOdLZP d~ ddMZQdndNZRddOZSddPZTddQZU ddRZVddSZWddTZXddUZYddVZZdbdWZ[ddXZ\e]d{dYZ^ddZZ_dd[Z`e]dd\Zadd]Zbdd^Zc dd_ZdxZeS)rzdA class to represent a triton kernel and helpers to generate triton kernel programmatically rhelper_functionszCallable[[sympy.Expr], str]kexprTztype[BlockPtrOptions]block_ptr_options_clsztype[TensorDescriptorOptions]tensor_descriptor_options_clsNz bool | None3transpose_discontiguous_tensor_descriptors_overrideFc ||_||_||_t||fi|t |j |j|_i|_ t|_ t|_ t|_ tt|_||_t%j&|_t+t,t,f|_t1|_t5j6t*|_t+t,t:f|_||_t5j@|_!d|_"d|_#ttH|_%d|_&|jNr|jQ|jR|jTr|jW|jY|jTr|j[d|_.g|_/y)NrF)0 optimize_maskris_combo_kernelrqrrr newvar_prefixsuffixrprologue_cacherYprologuepost_loop_combinepost_loop_storerr outside_loop_varsrrcount block_ptr_iddictrblock_ptr_to_bufferrrrgrhpointer_advancementsrr hint_overrideCounter _load_counts_pdl_load_index _pdl_has_waitr0autotune_hints triton_metarYcodegen_reduction_numelsr8cooperative_reductioninit_cooperative_reductioncodegen_range_treeinit_cooperative_reduction_maskhas_load_with_contiguous_rdimstores_with_contiguous_rdim) rtilingrrrrrrwrts rrrzTritonKernel.__init__ s$1(%4 *6*T//=.0(6(8 1?1A/=/?!+C!2#6 %OO-#'S>#3 / 1  # #D ) !$(S>#3 *6A6I6I6K ")6826   ) )$)) 4  % %  + + - !  % %  0 0 2-2*68(rctjstjjsy|j}|Dcgc]}t |t jr| }}t|dk(ry tjjj|||}td|DScc}w#t$rYywxYw)NFrc3&K|] }|dk( ywrr)rris rrz4TritonKernel._has_stride1_on_rdim..> s966Q;9)r# deterministicrforce_filter_reduction_configsrrrrrrOrPrI stride_varsZeroDivisionErrorr})r  support_varsr reduce_varsrs r_has_stride1_on_rdimz!TritonKernel._has_stride1_on_rdim$ s  F$7$7$V$V)) $ c=#@#@A   { q  ''**66uk<XK9[999! !  s#B.1+B33 B?>B?c<td|jD S)Nc32K|]}t|ywr)r[)rrs rrz>TritonKernel.has_store_with_contiguous_rdim..B s (, d # r)rRrrs rhas_store_with_contiguous_rdimz+TritonKernel.has_store_with_contiguous_rdim@ s' 040P0P    rct|Sr)rI)rrs r dtype_to_strzTritonKernel.dtype_to_strF s 5!!rcp|jxr)tjj|jSr)rYrOchoices should_use_cooperative_reductionr[rs rrz-TritonKernel.should_use_cooperative_reductionI s-$$ )S)S MM*  rcjsJjD]$}|j|xjdz c_&jd}jrt |jd}j j|_tj _ jjdtfdjDrjjdyy)z/One time setup code for cooperative reductions.NrQrNra  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) c3ZK|]"}|jrj| $ywr)r\rrrrs rrz:TritonKernel.init_cooperative_reduction..j s0   ''- - s(+z>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))rregrid_dimrZrrr semaphoressemaphores_namerc%cooperative_reduction_workspace_cacher8splicer}r%)rr sem_counts` rrz'TritonKernel.init_cooperative_reductionN s))))$$ #D}}( "  #KK$    4+<+5W II6 2     ((  II  P   rcd}|js|d}|jjd||jr|jj dy|jrJ|jjdy)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))rVr8r%_has_constant_xmaskr)r rsplit_aranges rrz,TritonKernel.init_cooperative_reduction_masks s{> }},oY7M .}o>?  # # % II   }} $$ II  e rc2|jD]q}|js|j||j,|js9|jj |j d|j|s|jrtd|jDrS|jddd}|j|}|jjd|j|y|j|jyy)Nzbase = c34K|]}|jywr)is_looprrs rrz2TritonKernel.codegen_range_tree.. s=D4<<=sr Trzrbase = )reriteration_ranges_codegen_headerr8rYr%riteration_ranges_ranges_coder}_get_reduction_symbols_flatten_reduction_indicesrrcodegen_reduction_indices)rrrn_basesrbases rrzTritonKernel.codegen_range_tree s$$ D<<44T499E&& ##{{m74+L+LT+R*ST   =D,<,<==66Dd777A   8D,=,=e,D+E!FG..tyy9 !rcy)z Indicate whether we need provide numel as arguments for the generated kernel calls in the benchmark. Should be true for pointwise/reduction kernels but false for triton matmul kernels. Trrs rneed_numel_argszTritonKernel.need_numel_args src|jxr4tjj|j|j Sr)rYrOrshould_use_persistent_reductionr[rrs rrz,TritonKernel.should_use_persistent_reduction s5$$ )R)R MM455*  rc|jxrGt|j|jdzk(xr |jxr|jddk(S)NrQr)rrrZrXrrs r want_no_x_dimzTritonKernel.want_no_x_dim sY  % % 1DKK D$;$;a$?? 1!! 1!!(+q0  rcy)Nztl.device_assertrrs rassert_functionzTritonKernel.assert_function s!r) copy_shapedense_indexing override_maskrrc  !"#$jj}d}t!t|t j dD]y} t | tjsJ|xst| tj}|rAt| tjr?jj| j } !j#| j$t| tj&tj(tj*tj,tj.tj0frtj2D cgc]} t| | r t4| } } t7| dk(r t7| dk(sJd| j !j9| dd|t:j<j>xs|xsj@duxrdk7} d }d}t}jCD]@}|jE|jFrd }nd}|j9|jHdB|r&jJrt:j<jLsrjOrz|sxj@slt7!|z dk(r[jQsJ|rHjRd k(r9 dd " dfd $ d"$fd #d!#fd }|}||Sd}d}jU}fd}tWrst7jYdk(r |\}}nLt[dgt7jYz}t]dgt7jYz}d|d|d}j^s j`rtfdjbD!n t!j@r!j9j@te|!|||S| r|sˉjfrjhr!jk}j@r|j9j@tgd}|jm|sq|jo|}|jq}t |trsJ|ju||j#|j$|jm|sqdgt7jYz}|D]y}t |tZsJjCD]R}|jw|jHs|jx}t |tzsJjY|||<T{ddj}ttZ|zdz}t]|}d|d|d}n/|\}}d|d|d}|!n|sr|\}}d|d|d}|!|| s|r |\} }nd}|r t|g!j@r!j9j@j!te|!|||Scc} w) zO Compute the index and mask to pass to tl.load() or tl.store() FrrrrQrrNTtl.int32ctj||j}|yt|jgt j |g|gt j|gS)z Matches expressions of the form: idx = s * xindex This implies stride (s,), and shape (XBLOCK,). Nrr?rArD)rRmatch_affine_block_exprsymbolr/rrrr)r  range_treeris rmatch_affine_blockz1TritonKernel.indexing..match_affine_block sl-DD:,,.>&%++,!.!=!=j!I J#H*;;JGH rc |j}tjdtjtj |g\}}t dt|j|jt|||jt|||z}tjjj|}j|j }t#j$||||}|y|\}} } t#j&|} tjjj)|j*t-fd| Dryt.j1|} t3| j5| dgt7| dd|ddD cgc]C\}} tj8t3| j5|j5| Ec} }z}| Dcgc]#}t;||t.j=|i%}}t?|Dcgc]}j5|c}|| | Scc} }wcc}wcc}w) 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)excluderr"Nc3lK|]+}j| xrj| -ywr)r~statically_known_power_of_2)rrrrIs rrzETritonKernel.indexing..match_mod_div_block..u sH!==eYOOH$@@GGHs14rrQr) rrsymbolsrrWildr~rrrrrrOrPrIremove_precomputed_replacementsrrRmatch_mod_div_block_exprget_slice_numelsrrr}rrrrGrMinrHrr/)r rrrmodulonum_dimsr match_resultdimsrAblock_index_exprs slice_numelslinear_block_sizergr?rrdrrIrs @@rmatch_mod_div_blockz2TritonKernel.indexing..match_mod_div_block( sh&'--/ !& "!))%**ykJ! v  (() HY$>?++oi&OPQ &77++ @@G @@AQAQR2KK9eX   ' ! %2CCDI 77++ NN:+<+<= ".  %2$@$@$L!) 88aI1'*,qr*:DH&E#s II-x/O/OPU/V!88=  1 (!2 3y-*H*H*TU3 3'HLM18;;A>M +#) %3Ns7AI (I%<I*c~fD]5}tj||j}|||}|3|cSy)ze Match a block indexing subexpression involving a single range tree. N)rRfactor_index_exprr)rr match_funcfactored_index_exprmatchrr)s rmatch_block_subexprz2TritonKernel.indexing..match_block_subexpr s\''# %J+>*O*Oj//1+''':JGE($  %rc tjjDcic]\}}||jc}}}j }|Dcgc]&}t j ||j(}}td|D}t}t||D]@\}}t|j|jdkDry||} | y|| z }B|t|z } jt j"j$r j&n j(} t j"j$rd} tj*} nt-t.j1} j2 j2}nt j"j4}t7dr|duz}|rtj8ntj*} | j;|| |j<| | }| t>k(r,t-t.jA|jBsy|Scc}}wcc}w)Nc3<K|]}|jywr)rrs rrzBTritonKernel.indexing..match_block_expr.. s*QT4;;=*QrQFtemplate_out_shape)r0r1rer rfr<rL)"rHrange_tree_nodesrrrrRget_subexpr_involving_symbolrrr/rr intersectionrsum filter_masksr#r use_block_ptrrrrMr rr<r)transpose_discontiguous_tensor_descriptorrrTrlrrrr0)rrindex_relative_to_xyr_indexrerindex_subexprs range_symbolsrsubexprr0r options_classr<rLtranspose_contiguousoptionsr r r r/rrs rmatch_block_exprz/TritonKernel.indexing..match_block_expr sl.8$2G2G2M2M2OP$!QAqvvIP/+#557 !, "(DD3T[[]""!+*Q[*Q Q .0 %(n%E +MD'=55g6J6JKLqP#1$?F~# F*L +5s>7JJ!!),}}22..;;==..$H(7(L(L%04/1J1- 9AACHPP#$!TT- #MMSS-t%9:, $0FF,0(CC,AA& (..'$* +'"&..%&7/!$;;04/1J1-5TT ${Q"s I +Icr7ttrddfSddjdDzdzfSjt j fS)Nz.shaperrc32K|]}t|ywrr)rrs rrzATritonKernel.indexing.._get_expand_str..$ s*Fa3q6*Frr)rrr&dense_size_strrr)r rsr_get_expand_strz.TritonKernel.indexing.._get_expand_str shj#.(\0$66*F:*F!FFLjXX**,eD4H4H4J.KKKrrrz , tl.int32)c3vK|]0}|js"j|s|jd2ywrN)r\rrrs rrz(TritonKernel.indexing..< s:',,T5L5LT5R{{m4('r)r )xmaskymaskzmaskr0_maskrrrrrrr)r rrrfrBlockParameters | None)rrrrfrrN)rzBlockDescriptorOptions | None)Aprepare_indexingrrr'operator attrgetterrrrrrrr rrrrrvr rrrrrrrrrr'r#rr rrr6var_listrallow_block_ptrr9ris_indirect_indexingrrrrrrrrrerrYrcopyissubset differencerWrnrrrrr&r(r8)%rr r r r rr index_varsrrrrprefix_matches need_dense have_densehave_loop_varsdense_mask_varsrrBrAr r rrG mask_shapexyzr tmp_masksr expand_listrrgexpand_shape_strr(r rr/r)s%``` ` @@@@rrzTritonKernel.indexing sk%%e,''  %/\ *(*=*=f*EF! :Cc5<<0 00#~]22(JTXX.((..sxx8  !2!23%%II))JJJJ''   !. 9 9"%c40t$"" ~&!+>*a/N3CCHH:1NN/ !2 3489C! :H MM ( ( + +d*qj   +5<++- 6D&&t}}5!%"   4;;-t 4 5  6t33 8S8S-1==?"OOI/0A5--e4  J. ! /B ' ,x !x /Bx 'x t   / ( (_ _ 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$--d3 * 5 5d ;I#--/C%c+<===&&s+%%cmm4 %--d3 #ec$*>*>*@&AA &KD%dC000 $ 7 7 9K??4;;7"&//C#-c3#77#7/3/C/C/Ec/JK, KK!388C[,A#BBSH $[1 .ykJ * l*9+R8H7IKI'I  Z"1"3<! "M?3I ?? MM$// * )$     %   a "s&Z c|j}t|tr|r&|r$|dk(sJd}n|sd}n|r|dk(sJd|d}nd|}|jr)|jdj r|j s |jr|jr!||jvr|j|}||fS|j|d}|jj|}|r t||fSt|j} t|trd| }nd | }|jj!|t"j$g } |jj'|| t)||d |} |jr+|j*j-| ||j|<n|j.j-| t|trn||j0|<t2j4D]L} |j7| } t9d | Dr'|j:| }||vsJd |d| d| ||<N||fS|j|}||fS)aGenerate 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'rKF)rrrtma_descriptorrrc3K|]A}tjjj|t j dCywrN)rOrPrIrQrr)rrs rrz1TritonKernel.codegen_block_ptr.. s=!'GG,,DD & a(8sAA z#duplicate advancement for pointer 'z ' at type 'r )rrrrYrerrr<rrrrrnextrrnamedvarruint64rrXrr%r8rrrrrRr)rrrrr9checkblock_descriptorblock_ptr_line block_varblock_descriptor_id named_var line_bodyradvance_offsets advancementss rcodegen_block_ptrzTritonKernel.codegen_block_ptr s2'') h 7 8 --- ---+E94KL+E95  ! !  $,,##%     SD,?,?%?$(#6#6s#; h &&e"*e!D HH,,^< y>500&*4+<+<&=#h8)23F2G'H$)78K7L'M$ HH--$ELL.  ^Y7(2B1C3~FV/WX $$MM++I6/?D'',II'' 2h8BFD,,-=>!. = =I*2*B*B4*H+:  %'+'@'@'F /|CABRASS^_c^ddefC:I %56!I& && (s3 &&rc,d|d|jd}tt|j|jD]B\}\}}tj j j||s4d|j|<D|j||j|jdd}tj j|} ||jjvrU|jj|} t| ts,tj j| j d} |dt#| d}t|t$r d |d||dS|d tj&j)|j*d|dS) NrrrFTrrrr tl.store(z.store()r8r!rr5rOrPrIrQr7rr? get_dtyperinplace_buffersrr^ other_namesrVrrrrD) rrrrrr9rrg broadcast_dim store_dtypebufs rcodegen_block_ptr_store_linez)TritonKernel.codegen_block_ptr_store_lines #5'H,@,@+AC*3 $$h&>&> ?*  8 %C%#}ww77]K27**3/  8 66      7 gg''- 499,, ,))++D1Cc:.gg//0BC '/ <=Q? h 0ykE75'; ;GAHH$9$9(:J:J$K#LBugUVWWrc|s|syt|tjsJ|j|dd}t|tsJ|j }|j r |jnd}|rt|j|nd}|j||rdnd||} |j|} |jj| | dtjy)NFr0) assignmentr)rrrrrrrr)texprrename_indexingindirect_assertget_load_bufferrrrr[) rrrlowerrrrr)size_strlinebuffers r check_boundszTritonKernel.check_bounds,s $ +++==RV=W(O444&& (0(9(9(;8$$8=5--d344## esx %%h/ &$5 Lrc|js|jr |jS|jr5|jdj r|j s |jS|jS)NrK) rrrrYrerrr8loads)rrs rrzTritonKernel.get_load_bufferFsb  "h&:&:&<<<   ! !  $,,'')99 :: rztl.extra.cuda.gdc_wait()z%tl.extra.cuda.gdc_launch_dependents()ctjjjjsyt t jtjjjrytjjryt jjjdk(xr$tjj!ddk\S)NFrrr)rrr#r enable_pdlrrOrselect_algorithmTritonTemplateKernelrrrPrr"rrrrr_enable_pdl_codegenz TritonKernel._enable_pdl_codegenXs%%,,77 ahh @ @ U U V ==   GG / / 1 6 6& @ ; 00215: rconsider_readsc|jsytjjtjj $tjj j ndfd|sJduxstfd|D}|sy|jr%|jdjr |j}|j|jy)NcJjj}r*tj|jj}t fd|DS)Nc3K|]7}jj|j|jk(9ywr)mutation_renamesrr)rw current_nodedeps rrzOTritonKernel._handle_pdl_before_access..matching_dep..us9|4488HHs=A) read_writeswritesrrreadsr})r prev_depsrr prev_nodes` r matching_depz.matching_deppsY( ((!--44I%OOIy7L7L7R7RS " rc3.K|] }|ywrr)rr(rs rrz9TritonKernel._handle_pdl_before_access..{s,S\!_,SsrK)rrOrrrP scheduler previous_noder}rYrerr8r%GDC_WAIT)r wait_bufferr dependencies need_waitrrrs ` @@@r_handle_pdl_before_accessz&TritonKernel._handle_pdl_before_accessfs'') xx,, /0ww/@/@/LAGG   + +RV  |%S,Sl,S)S    T%5%5b%9%A%A))Kdmm,rc|jsy|jdkDry|jr%|jdjr |j }|j |j|j |jy)NrQrK) r use_countrYrerrr%r GDC_LAUNCH)r launch_buffer result_vars r_handle_pdl_after_loadz#TritonKernel._handle_pdl_after_loadso'')    ! #   T%5%5b%9%A%A 22M  .0rc$g}d}d}|jD]t}t|tur|j|vr|r%d}t|tur,|j|vr||j |t |}|j|v||_y)NFT)_linesr"rrrrWrr)rcode new_lineshas_waitprevious_launchls r _filter_pdlzTritonKernel._filter_pdls  AAw#~$--1"4#HAw#~$//Q"6".MM/2"%i.   Q    rcP|jjt|||yr)saved_partial_accumulaterrg)rrrrrs rrzTritonKernel.partial_accumulates$ %%,, dNC 8 rc  |jj}|jxxdz cc<t}|j ||}t j j}|j|d|j||dd}t|tr"|j|jrd|_|j|j!}t#d|j%|j'D} |j)|rd} nX| sd} nS|j*rE|j,dj.r,fd } d } t1j2t4d | }nd } |sr8|j7r(|j8rd t;|j8} nd} nd } d} t<j>j@r&|jBjE}|vr|dkD} |j)| xr|j* xr| xr| }d }|rd}d}d}tGre|}|tHjJtHjLfvr=t<j>jNrtHjP}n|dtS|dz }d}nt|tTtVfr|jY||| \}} t|tTr d|| | |d}n/|dt jZj]|j^d}|ja||jb|jddd}|jd}ntg|rd|d|d}|jh}d}n^d|d|jjd|jl| | |d }|jnr |jn}ntpjs|j}|tHjJtHjLfvr/t<j>jNr|dz }tHjP}|tHjtk(r/tHjvjx|dz }tHjt}|j{|}|j}||j~j|||||}|j|||jdkDr xxdzcc<t|tsJ|j|_D|rd|d|d}|j~j||||jn}|jr|jrd}n|tHjtk(rd }nd!}|j8rt;|j8n|}d"|jld|d|d}|j~j||||j}|j*r|jss|jj||S)#zc Load from the memory location 'name', offset by some indexing expression 'index'. rQTFrrrc3&K|] }|dk( ywrr)rr,s rrz$TritonKernel.load..s AF rz, eviction_policy='evict_last'rKc kDrsryy)N evict_last evict_firstr)expected_countrindirect_indexing load_countsrsr decide_laterz'TritonKernel.load..decide_laterst$~5"3'$rz, eviction_policy=''zrz, other=rdz, cache_modifier='.cg'Nrrrr z.load(rwrr), rrrrrz0.0Truerr)Jrr rrrTrOrPryrtma_compatibility_checker_clsrrrr rrrr}get_strides_of_loadris_broadcastedrYrerrrr?r _load_otherrcr#r skip_l1_cacher[buffer_read_countsrmrrjrkrirrIrrrurrrDrr?r8rr rr)r rrrrrrrrrrrrnr rrr#rr')rrr r make_lineoriginal_indexrrr is_coalescedeprr9 has_read_depsrrcachemodappend_broadcastrrrm load_bufferrr other_valrrrrs ` @@@@rrzTritonKernel.loads iiood#'' DQ=E  55e<!!$'== &*&H&H 'I'!   h 0T5N5N NN6 26D .((* **,   44^DKKM     ~ .1B1B  " "t'7'7';'C'C % %).N+B!))*:FLQIB :8+<+<+>"=1A1A#B"CD%E  == & &!%!A!A!C )) 24 81 <  ##N3 3 ))) !!    /H $ #D )D77==77!MMEd;u#5"6a88DE(_6M$NO*.*@*@#x+' %h8%&6%7wrd8*AND./vahh6K6KHL\L\6]5^^_`D==((((#'# >!,,&~6!#d>*:"=#+#6#6 !#d8+=+=*>c(BSBSATUWTXY^X_`h_iijk (($11E)99(..IE%--88MM88))  "u}}'8'8'@& **84  &&{D9XX&& 4U'  ##K<   ! #   " *&7888'11  %j\4D3EQGD**Th6K6K+J!!** Dejj(!DD7;7G7GM$"2"23T#8#4#4"5R |2i[PQR!XX..U*:J:J/ $$X-?-?-A*  " " & &z 2rc |jj|}|}tjj |}d}||dk(r|dk(} |j ||d| }|j |d|du|} t| tr6|j| jr|jj|||jjv} |j|} | r'| r%|jj!t#|dt| t$t&fr-|j)||| \} }|j+|| | ||}nV|| j,}t/|r~|j0r|d|j3dz }nZ|j4Nt7d |j4Ds2d j9t;t<|j4}|d|dz }d |d |d |d | j>d }n|dk(rd|_ | j,}t/|rZ|j4Nt7d|j4Ds2d j9t;t<|j4}|d|dz }d|d |d |d | j>d }ntCd|tEjF}|jHs7|jJr+|jM|jO||j|jQ|j|d|jj!t#|||jHs|jRjU||jWy)zn store the 'value' to the memory location 'name', offset by some indexing expression 'index'. NtmaTr)r rrztl.debug_barrier().broadcast_to(rc38K|]}t|dk(ywrNrErs rrz%TritonKernel.store..s9&'CFcM9rrxrr atomic_addc38K|]}t|dk(ywrrErs rrz%TritonKernel.store..?aCFcM?rztl.atomic_add(z, sem='relaxed')z store mode=r),rrrOrPryrrrrrr rrrzrstoresr%rXrrrurrrrrFrrRr&r(rr)atomic_add_foundr contextlib ExitStackrYr enter_contextguard_cooperative_storerrr'close)rrr rmoderrrrrr is_inplacerrmr9r indexing_str value_shape exit_stacks rstorezTritonKernel.storeisLiit$!!$'$(! <45=EME(,(J(J )K) % == dl&? !  h 0T5N5N NN6   , , 3 3D 9TYY666 ,,^< . KK ! !,t5I"J K h2I J K&*&<& . Q ?? #C5\N#eWBxGXGXFYYijD% D6&:; ;))+ $$)C)C  $ $T%A%A$ %T U &&t{{D&N l467$$  " " & &u -rcZ|jjd|dt|dy)Nztl.device_assert(rr)rr%r")rcondmsgs rdevice_assert_asyncz TritonKernel.device_assert_asyncs( !24&49+QGHrc|jj}|jt|d|d|j S)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):)rrzr%rXindent)rrrrs rrz$TritonKernel.guard_cooperative_storesC 88NNPd.?uK,PQR}}rctd}|D]0}|t|ds| |j}"||jz}2|S)Nr )rr )r variablesmaskselems r_combine_maskszTritonKernel._combine_maskssK 3D|t[)= NNE!DNN2E 3 rcV|jjtj|jj |d}|j |d} |j |d} |j |d} |r|jj |dnd} |r|j |dnd} |tjk(rd}n!|tjk(rd}n td|j|j|dg|r|dgng|jj|jd |d |d | d | d | d |d |d |d | d | d |d ||j }|j!|j||j#|||}||_|S) z3 See [Note: Inductor bucketize op] rrQr"rrrztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(rz, )r)rr'r0ONE_ELEMENT_PER_THREADrr rrr[rrrrrrrrrr )rr boundariesboundary_indicesindexing_dtypersortersorter_indicesboundaries_ptr boundary_sizeboundaries_underlying_numelboundary_stride sorter_ptr sorter_stride triton_dtyperkrs r bucketizezTritonKernel.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##F,>$))E*+1--rc|j}|dk(rd|d|fS|j}dg||z zdg|zz}|g|d||z dg|znd}|ddj|d|fSrr)rrrrr r rs rreduction_resize_and_shapez'TritonKernel.reduction_resize_and_shape#s'') A:6ugQ?F F)))VHw,>>=B=N 9e'uw( 9A3= 9TX $))E*+1-y88rc|jdk(r|S|j|jz }|j}|d|dgz}|jj |t t ||||t|S)zC Reshape to RBLOCK, collapsing all reduction dims. rQNrr)rXrrrrrrr)rrrr target_ndimr target_shapes rreduction_collapse_dimsz$TritonKernel.reduction_collapse_dims/s  " "a 'L--/$2I2II ,,. $\k2hZ? xx   3u:}l C % !  rc8:;<=>?@ABdEd?dF?fd }tj?fd|}}|rJtj||}?rtjn?rtjnj sJt djD}j|t|}jr|jjjdjd} jrYj} t| dk\sJ| D cgc] } d| vsd | vs | } } d d j!| d ;t#| An)j%;t#jAj';Afd |}d} dvrt)|t"r|\}} j+j,z < dG<fd = dH=fd }<>fd}|f}|j.j0vrj.j0|St3}t5}t7j}d|<<j.j9|t#|}t d|D|_dj!|::fdBj<rst>j@jC}d@dI@Bfd }dk(rd}nwt)|t"r&tE||Dcgc]\}}|||}}}nAdk(r3j.jGjH||jJ}n |||}dvrt)|tLsJtNjPjS}| r"dtU| djW|d }nCtUj.jGjHd!| d"|d#||jX}d$d%d>|jH|||||_%nd&k(r8jZrj]||B|}nmj_|}nYd'k(rMt)|t`sJ|\}}}t#fd(jcjH|||<D}ndk(rje|}nt)|tLsJ=jH||jJ\}} }!j.jGjH|| |!}nj.jgd)||t#j}"t>j@ji}j'tj|}t)|t"sdk(rj} t| dk\sJ| D cgc] } d| vsd | vs | } } t#| |"_,d d j!| d ;jljo|"d*;d |d |d n5jljo|"d*j%d |d |d dvrd)|d+}jpjs}#jljo|d*j%d tjt|#jvd jW|#d d$d%d>| !dtU| djW|#d n| d,}$jHjyd-|"d.|d/>d0|"d |d |d |$d1|"d2B|"d3|"d4|d2B|d3|d4|jz||"|n4t}rj]||B|}ndk(r3d)|d5}%d)|d6}&jljo|%d*j%d7|d jljo|&d8j%d |d jHjyd9|%d.|&d:|%d |&d |d t~jd; jHjyd9|%d2B|%d3|%d9|&d2B|&d3|&d9 |}'j.j9|'jX}(jjz|'|(|%|&<}nt?j})|)|"|}*dk(r!jHjo|"d2|*n'jHjo|"d2B|*|"tjk(rDj.jGjz|"d<tj|"jX}"|jz||"djZrt>j@ji}tj}+jzjfD]2},|,jod=|+j|,j4dvrjzjo|d>j|d?j|d@|}-jpjs}#j||#tjt|#jv}.|j||-|.nFt}rd&k(sJ|\}/}0}1j|/t5|d}2j|0t5|dA}3j|1t5|dB}4jj|/|0|1|2|3|4< ndk(rw|\}'}(t)|t`sJj|'t5|d}5j|(t5|dA}6jj|'|(|5|6<n1j|t5|}7|j||7d|+j|j.j0|<t)|t"r|n|f}8jj|8tdC|8DsJ|rD|8D]?}9|9jJ|k7sjzjo|9d2|9dDt|d A|Scc} wcc}}wcc} w)JzS codegen reduction of value to Triton according the reduction_type cH|duxr|jxr|jdkS)NrX)rrZ)r(s r should_upcastz-TritonKernel.reduction..should_upcastNs$D=KQ%8%8KQZZ!^ Krct|jr$tj|tjS|Sr)rrrrr)rrs r maybe_upcastz,TritonKernel.reduction..maybe_upcastQs4!- UEMM2  rc(|jSrrR)rrs rz(TritonKernel.reduction..\smAGG.Drc3:K|]}|jdywrIrrs rrz)TritonKernel.reduction..eMDdkk]$/MrKrrXYrrrcxjjjd|dd|jS)Nrrrr)rrrr)rrFrrs rrz(TritonKernel.reduction..s?dhh'' "1#R'7q9gg! (rN)argminargmaxct } j||} dk(rYt jdk(}|jJ|r|d}dg|jd}n>|d}g|jd}n( j |d|dd|j\}}||d j |d}n |j}|||fS) zK Helper to generate a reduction call, e.g. tl.sum. rrXz[None,:,:,None]rQz [:,:,None]rrrr)rrrrrr rr) rr result_typetriton_reduction_fnis_bmmrkrrgrrrs rfinal_reductionz/TritonKernel.reduction..final_reductions#@"O 00FE&T11349{{... %wo6F00a0E %wj1F-ekk-1-E $ ? ?*+1UG2cU!.final_reduction_defines0*&%EKE1a MMZLE73 4rcj||}j||}|jd|d|dd|d|dd|dj|dd y) N z_val, z_idx = triton_helpers.z _with_index(r) r_idx )rrr )rrrr rgrroot_oprs rfinal_argreducez/TritonKernel.reduction..final_argreduces00FE00FE MM F:,.DWI\Z_Y``bchbiiklokpq C 5 5 D6I JKL rrrc3>K|]}t|dr|ywrh)rErs rrz)TritonKernel.reduction..s!* (;CF(CC* sr%c:s|Stj||Sr)rr)tvalfvalrs rrz*TritonKernel.reduction..where_conds  (..tT4@ @rcd|js|tjk(r|S||k(s|dk(r|S|tj|jk(rtj|jS|tj|j k(rtj|j S|S)z7update reduction constant mask value to match dst_dtyper)rrriinfor~r)rrrs rupdate_constant_dtypez5TritonKernel.reduction..update_constant_dtypes ..)uzz2I#O )X]#Ou{{95999 ;;y1555Y!7!;!;; ;;y1555#Orc||j}jt|}jj j |||j|j S)Nr)r_map_tuple_or_scalarrcrrrr)rr default_strrrr6rs r _mask_valuez+TritonKernel.reduction.._mask_valuesb/EKKP"77 wO xx((LLuk2++++ )ronline_softmax_reducerrR)r!r rrrrzindex, z.shape)r~rwelford_reducewelford_combinec3vK|]0\}}jjj||2yw)rN)rrr)rrrrrs rrz)TritonKernel.reduction..4s9#$uHH%%dllEe%T#rr( = tl.full(_indexr r*_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_bvalrQr"c3<K|]}t|tywr)rrnrs rrz)TritonKernel.reduction..sJ:a!23Jr2r)r(rrr)rrWrrW)rrWr#rrz.tuple[str, torch.dtype | None, BlockShapeType])rrWrrWr#rrrrrW)Rpytreetree_anytree_maprrrYrrer8r'rrrrrrr&rrFr8rrrXrreduction_cacher^r\rar$r rr$rrrrrrrWrOrrrrrrr<welford_reduce_fallbackr_welford prepare_softmax_twopass_fallbackrjdefault_accumulatorrcr8r%r[select_index_dtyper5r~rrrCr#r%online_softmax_reduce_final_reductionget_reduction_combine_fnrrUrrrrrr *codegen_cooperative_reduction_peer_combinewelford_reduce_final_reductionrrrvrRrS)Crrrrrr do_upcastoriginal_dtyperreduction_range_prefix dense_sizesr xy_sizes_only logical_indexr(r/racc_typetorch_acc_type result_shaperrr: masked_valuerr(accumulator_dtypeaccumulator_indexmeanm2weight_result_dtype_shape accumulatorrraccumulator_maxaccumulator_sum result_max result_sum combine_fnupdatedrr~peer_valpeer_idx result_mean result_m2 result_weight peer_meanpeer_m2 peer_weightpeer_maxpeer_sumpeers result_tuplerkrrFrgr&r.rr6rrsC```` @@@@@@@@@r reductionzTritonKernel.reductionCs7 L OO$DeL  OOL%8E)6y)A yI%25%9EMMuE$$$$MDA%>QRda Aq 1R R5( $xx00uEKK0X *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&n5!00z8U  #::$%j\"6$%j\"6 ##&'{43F3F3H2IIZ[cZddef ##&'|D4G4G4I3J"XJVWX ##$%W_,=>()O+>@ JJ U[[-E-I-I  4 4j(HU%n5%)99998B5 Y  KK$Y/AJ II$Y/AJ #MM!$Y/AJ 33((!  #::)3& J!'8444JJ 0 ;WQZ JJ 0 ;WQZ::((GG 0 ;W't';';ZPTU    .8  +%/ E%Bz  %%l3J\JJJJ & <<>1**44!(#fXT2En2U1VVWX  Q Yd SR%s8 t t 't. t<tc|j|||}|j|||}tdDcgc]'}t|jj |)c}\}}|j d|d|d|d|d|dt jd|d|j|d|d|j|d||fScc}w)Nr"rR r9 = triton_helpers.online_softmax_reduce( ) r) rrcrrr$rr#rr ) rrrkrlrgrr(rmrns r_online_softmax_reducez#TritonKernel._online_softmax_reduce!s66vPUV66vPUVMRSTX!V#dhhooEo&B"C!V J  L:,' !O#4Bse2f>R>R=ST LD11ZLBCD LD11ZLBCD   :%%"Ws,Cc fd|||fD\}}}d|d|d|dd }fd}|||fD cgc]/} jj|| j1} } jdj | D cgc] } t | c} d|t fd| DScc} wcc} w) z; Helper to codegen triton_helpers.welford. c3DK|]}j|ywr)r)rrrrrs rrz(TritonKernel._welford..6s(   ( ( > s ztriton_helpers.welford(rrc2t|d|dzdzSr)r)rrgs r reduced_shapez,TritonKernel._welford..reduced_shape<s$qcAgi(889 9rrrc3VK|] }j||j"ywr)r r)rrrs rrz(TritonKernel._welford..Es*   + +E5;; ? s&))rr$rr%r&rr) rrrdrerfrgrwelfordrrwelford_resultsr"s `` `` rrPzTritonKernel._welford2s F+ b&,D6B4r&C5J : F+  HHOO%}U[[/IO J   DII&G!s1v&GHIWIVW (    'Hs 4B:B? c|j|jz }t|dt|j |t j }t|dt|j |t j } t|dt|j |t j } |jj|d|jd|d|jj| d|jd|d|jj| d|jd|d|dk(r>|\} } } |jjd |d | d | d |d| d| d | d| d| d n8|dk(sJ|jjd |d | d | d|d|d| d| d|jjd|d||d|d| d|| d| d| d|| d| d |}|j|j|dd|| | || S)z%Helper to codegen a welford reduction_meanrrrs_m2_weightrFrrr=r*rAz<_next = triton_helpers.welford_combine( z, rBr<z;_next = triton_helpers.welford_reduce( z1, roffset == 0 ) z rrCrN)rrXrnrrr!unknownr8r%rFrrrWr)rrrrrr^rrgrjaccumulator_m2accumulator_weightrdrerfrss rr<zTritonKernel.welford_reduceJs%%'$*A*AA'l% ,,./&&(  +l# ,,./&&(   /l' ",,./&&(   m<(;(;(='>b ! L  l4+>+>+@*AH:Q O  !",t/B/B/D.ERzQR S  . .$ D"f LL   W^$4GRS   MZ;-u(={KLM  C n-=U+C^ TUV  J2D1EU/KM_$`#ab   ! 22  " "          rc t|j|||||| } |||g} tt| | D]E\} \} \}}| "|jj | |} | | | <|j | d|Gt| S)z0Helper to codegen call to triton_helpers.welfordrr)rarPr!rrr$rr)rrrsrtrurdrerfrgrr result_exprsr, result_exprrrs rrWz+TritonKernel.welford_reduce_final_reductionsdmmFD"fc5IJ#Y > 09#lF:S0T 6 ,A, ^eU""hhooEoG "- Q MM[MUG4 5  6 \""rc|j|||}|j|||} |jd|d|d|d| d|dtjd|d|j |d|d|j |d||fS)Nrrrrr)rrr#rr ) rrrmrnryrzrgrrkrls rrTz2TritonKernel.online_softmax_reduce_final_reductions66vxO66vxO  L:,' !O#4Bse2f>R>R=ST LD11ZLBCD LD11ZLBCD   :%%rcD|jr|jdStS)NRSPLIT)rr4rs r max_rsplitzTritonKernel.max_rsplits"   $$X. .  rc|jd}|jsdnd}||jz|jz}|jj |\}}|j jd|d|d|j|dt|d|d |d |d d |j|dddg|tj} |jj| d|dt|d| S)a  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. rNzxindex < xnumelNr-z_ws = (r z).to(tl.pointer_type(z)) tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), rrTr _peersrrrz = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, r)rZrrZrrrtrrrrIcreate_cse_varr!rrr%rc) rrr default_valxnumelrrprrrsr{s rrVz7TritonKernel.codegen_cooperative_reduction_peer_combinesSS!(,(@(@(B %..(4??+<<!GGPPQWX %% GG9C0A0A)0L/MMbcnotcubvw$%J:,VXY]X^_   & ##l& !X&&&( $  &&g[ -eers~eeAAC D  rcH|jsJd|_tjj|}|j |d|j ||dd}d|_|j j|}tj}|jr+|j|j||j|j|j|t|t t"frY|jj%t'||j)|||j+||d|j-nt|t.sJ|j0}t3|rZ|j4Nt7d|j4Ds2dj9t;t<|j4} |d| d z }|jj%t'|d |d |d |d|j>d |jAy) NFT)rrrrrrec38K|]}t|dk(ywrrErs rrz/TritonKernel.store_reduction..rrrrrrxrr)!rYrOrPryrrrrrrrrrrrrrrr%rXrrrrrrrrRr&r(rr)r) rrr rrrrrrrs rstore_reductionzTritonKernel.store_reductions $$$$ %!!$'== &*&H&H 'I'!  !%iit$))+  % %  $ $,,T43G3GH  &&t';';SA h2I J K  * *55  ,+H,C,C,E+HI   h8 88#--L%e,KK+?5;;??"iiC(=> . Q ??  * *uDc%8CTCTBUUVW  rc N t jdt tdDcgc]*t fdt t ||D,}}djdtjj|D} jd|dtdd d l m }d d lm}|| G fd d t } j#5t%j&| 5||} djd| D} jd| dddddd|j(j+ j-Scc}w#1swYAxYw#1swYExYw)Nz @triton.jitr"c3rK|].\}\}}jdd|||j0yw)rr(rN)rjr)rrcrrrr,s rrz,TritonKernel._lift_helper..)s@%A~u s1#Qqc]%u{{ Ks47rc32K|]}t|ywrrErs rrz,TritonKernel._lift_helper../sRc!fRrz def {name}():r!rrq)ShapePropagationOpsHandlerc4eZdZ dfd Zy)+TritonKernel._lift_helper..CSEProxyc d|z t||i|}t ||i|}jt ||i|||S)Nr(r)rr) rrrrw output_dtype output_shaper dtype_handlerhelper helper_name overrides shape_handlers r_defaultz4TritonKernel._lift_helper..CSEProxy._defaultAs4&z)  w!  #" #  w!  #" # ||,GIt,d=f=&& $rN)rrrztuple[Any, ...]rwrrr )rrrr)rrrrrrsrCSEProxyr@s-  '6 @N   rrc32K|]}t|ywrrE)rrs rrz,TritonKernel._lift_helper..ZsBF Brreturn r")rYr%rVrcrr!rr&rr from_iterablerr}rr!torch._inductor.shape_propagationrr-rrOset_ops_handlerrr'getvalue)rrrdtypesr,r signaturerrrroutputsrrrrrrs ` @@@@@@r _lift_helperzTritonKernel._lift_helpersy  !'e1X   )23vv3F)G   IIRioo.K.KD.QRR = 267#% * PP24 24   ~ 0]]_ 2a// ; 2$iGiiB'BBG   wwi0 1 2 2 $$(():k(RRk ` 2 2 2 2s)/F F2FFF FF$c BjsJjrJdtdjD}j |t |}j rJdg}g}td|D}tjjjj}j|||}jjz } t!||D]U\} } jjj| dt#| d| | j$} jjjd| d j'd| tj)} |j+| t-| } j.rljj)}d |d <jj1| |}d d j3|d }| j4rdnd}j6j9|d|d |d | d|j+|Xdfd}|d|d| d |d|||}j.sd}|Dcgc]*}|d|dt;|j<||,}}|t|t|}|t||}t!||Dcgc]+\}}|d|d |d|j<|j$-}}}t!|||D]*\}}}jj9|d|d |d,n|}|D]$}t?|t@sJt||_!&t|Scc}wcc}}w)z: Perform an associative scan on 'values'. TODOc3:K|]}|jdywrIrrs rrz$TritonKernel.scan..lrrz(ops.scan not supported inside ops.maskedc32K|]}t|ywrrKrrs rrz$TritonKernel.scan..tFe*51FrrrrrrrrKrrz float('nan')z-1r?c2djd|DS)Nr!c3&K|] }|d ywrNrrrs rrz1TritonKernel.scan..csv...csv88.cse_multiple.. LI488$$Y/L%(rr rrcrRrrrr$rrr%r rr)rrrrrcr, cache_keysrrr result_varsrrrs r cse_multiplez'TritonKernel.scan..cse_multiplesF A;@8DaTF"QCr%1DJDLLLAKLI Y/LL'*&&&9"UEe5;;?K LL " "{#$Cv. *-[*)E 4% I+0J( Y 3 4% %ELD"D91D!ztl.associative_scan((rcT|jyt|j}d|d<|S)NrrK)rra)rrs r_partial_scan_shapez.TritonKernel.scan.._partial_scan_shapes*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, )"rYrrrer8r'rrrrrrrrrrXrrSrrFrrr^rr$r&rr8r%rKrrrnr ) rrrorrbroadcasted_values accumulators cse_computecombine_helper_fnrgrr value_dtyper^ reduced_sizerjreduced_size_strrrpartial_scan_varsrpartial_scan_varpartial_reduce_vars accs_nextfull_scan_vars full_scan partial_scanracc_nextpartial_reducerrs ` @rscanzTritonKernel.scan_s$$$$--5v5-MD&?q#A ,1,C,C. ##"m;/?.@7)2hZWXY##K09 1< = &$)#C(:$;#.D.DE-.>?# ##5#6>Q8RSI'l(;=NON03>CT/U  ,I| ,YKr,qI&,,&,,K:=<)<: 5+~ &&"m#.rrz(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc32K|]}t|ywrrrs rrz$TritonKernel.sort..rrrrrrc2djd|DS)Nr!c3&K|] }|d ywrrrs rrz1TritonKernel.sort..csv..rrrrs rrzTritonKernel.sort..csvrrcNt|}t|Dcgc] }|d|d|}}t fd|Dr)|Dcgc]} jj |c}St ||D cgc],\}} jj || j.} }} jj | d|t | |D]*\} }|r|| _ jj|| ,t| Scc}wcc}wcc} }w)Nrc3TK|]}jj|!ywrrrs rrz:TritonKernel.sort..cse_multiple..rrrrr)rrrrrcr,rrrrrrrrs rrz'TritonKernel.sort..cse_multiples&'A;@8DaTF"QCr%1DJDLLLAKLI Y/LL%(0B$C E5e5;;?K LL " "{#$Cv. *-[*)E 4% I+0J( Y 3 4% %ELrrKrrnumelr"ztriton_helpers.sort_with_index(rrQz , stable=z , descending=zUnhandled sort)rYrrrer8r'rrrrrrrrrXrrr!rFrr\rr`rr rs)rrrstable descendingrrrgr,rrrrrrr input_varrs` @rr>zTritonKernel.sorts@$$$$--5v5-MD|jt j"} tA|j} |j,jC|j(| d| d| jDd}d dl#}|jHjJjM|d|j(5| |||}ddd|j(j+|d dddtOtQ|jD]$}|j(j+d|d|d&n/|jRrtQ|d kDrt|D]\}}|j(j1| 5|j4}|jTrdnd}|jTrdn|d}t jVjXrt[dkDrd}nd}|j(j+d |d!|d"|d"|j]d#|d$ ddd|j(j1|d z 5|j_||j(ddd|j(j1tQ| 5|ja|j(|j(j9|j|j(j9|j|j(j9|j|j(j9|jdddtcgt|D]u\}}|j(j1|d z 5|jd|jfjiD]\}}|tQ|d z kry||d z}|jd|jf}||vrT||}tjjm|}to|jp|}ts||Dcgc] \}}|||zz }}}|j(j+tu|jv||d%|d"txjzj}|d& ddd|j,j|j|jxn|j(j9|j|j(j9|j|j(j9|j|j(j9|j|j(j9|j|jTrb|js |j rJ|jd'}|j(j9d(|d)d*+|jj|js%|j(j9|j |jj|jj|jj|jj|jj|j jycc}w#1swYxYw#1swYxYw#1swYxYw#1swYxYw#1swYxYwcc}}w#1swYxYw),a 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_BLOCKrz/split_size = min(RSPLIT_SIZE, xnumel - xoffset)z@for _ in tl.range(0, split_size, XBLOCK, num_stages=NUM_STAGES):rQ)rrrNmask =  < rzxindex += XBLOCKrz, 0))rrrz&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(rrrz = tl.advance(rz + tl.program_id(1)zR if HAS_RSPLIT: triton_helpers.x_grid_barrier(r+Tr)F indexing_coderrrrrrermix_order_reductionrr!rrr$rrRrrr8rcr8r%rrjrrrrr writelinesrrUrrrunittestmockpatchobjectrcrrYrrrrrrrrbrrrrrrrrrXrrOrr invalidater cache_clearrrrwr) rr loop_trees accumname2varr partial_accumrrrentryrNrrotriton_reduction_functionnewvalrrplevelr loop_startloop_end num_stagesr advancement prev_treeprev_advancementsprev_advancement prev_block prev_num_itercurprevsem_ptrs r codegen_bodyzTritonKernel.codegen_body3s   zz{{||%%## '+'7'7Ht4<5;;W33M7Kse} ##f3G9-?') ?VS!>!>?@  ##$>/" /"  ++( $ 8 8 C#,+^I;bI^I^_jIkHllm n)<##D$:$:;  "A #D II  T// 0 II  TZZ ( II  T\\ * II  T[[ ) //0  % %  " "d&:&:--..ABG II  33:)<    6 6 B B D'' II  T11 2   "    $$& ""${IrI) ?) ?fJJ . .4/"#swg*g&Hg, g*g,Bg9hB0hB#h&4h Ah&g) $g,,g69h h h h&&h0 cg}|jr3g}|jd|g|D]}t|tr|j t |/t|t rZtjjj|j|j}|j t |t|tjrQtjjj||j}|j t |tdt!||S)Nrrz!Unsupported numel argument type: )radd_numel_to_call_argsrrrrrorOrPrIoptimization_hint_with_override inner_exprrrrrLr")rr numel_argsrhints rkernel_benchmark_extra_argsz(TritonKernel.kernel_benchmark_extra_argss     !+-J  ' 'J ;! Vc3'KKC)_577++KK&*&8&8LDKKD *UZZ077++KK&*&8&8LDKKD *$'Hc %TUU! V" rct}|jj\}}}}|jgd|j 5t j }g}t||D]\} } dt|} tjj| } | rtjjj| j|j} tjjj| j!|j}|j#| d| d|d| j%d| j'd n| tjj(vrtjj(| }tjjj|j+|j} tjjj|j-|j}|j#| d| d|d|j.d|j0d nt3| t4rftjjj7| j8|j}d | j:vrd }|j#| d |nt3| t<rtjj?}tjjj7| j |j}|j#| d |d|d| j0dntAd | |jC| |jE|jG|j#ddjI|dddd|jgdtjj?}|jJ}|j 5|j#dtjjLjO|d|j 5|j#tjjLjQ|tS|}|j#|d|d|j#tUtVjXd|ddddddd|jgd|j 5|j#dtjjLjO|d|j 5|j#tjjLjQ||j#dtUtVjXddddddd|jgd|j 5|j#d|j#d|j#d|j#dtjj?jZd|j#d||j#d|j#d ddd|S#1swYxYw#1swYxYw#1swYxYw#1swY xYw#1swYxYw#1swY|SxYw)!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=rrrrz = torch.zeros(z*Don't find the buffer or const tensor for rrN) r(zdef call(args):zwith rz = get_raw_stream(z.run(*args, stream=)r(r(z def benchmark_all_configs(args):z.benchmark_all_configs(*args))r(r(zif __name__ == '__main__':z5 +!'!$x.!12gg,,X677++LL &*&8&8MDWW--NN(&*&8&8OF$$#*$4TF"VHJs~~O_N``ijmjwjwjyizz{|!2!22#$77#4#4X#>L77++LL$))+&*&8&8MDWW--NN$++-&*&8&8OF$$#*$4TF"VHJ|ObObNcclmymmmAABC 1"#''"2"2"R"R &*&8&8#S#K% 4&' $$z[M%BC6WW@@BFGG,,LL T5G5GME$$#*OE7*VHIV]VcVcUddef#DXJO  *k5 +l   T==? @   wtyy';& ?u: @x 9:<<>$$ ]]_    uQWW%7%7%D%DU%K$LAN O   GG&&11%82%8   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    Q: @: @B            sf N[$A [1B [:[-A [5:A%[([5B%\[ [ [[%([2 -[55[?\ ctjdjtjj j dS)Nzl from torch._dynamo.testing import rand_strided {} import torch get_raw_stream)textwrapdedentrrOrPr3import_get_raw_stream_asrs rimports_for_benchmark_kernelz)TritonKernel.imports_for_benchmark_kernelws:  F177%%>>?OP Q   rc|jry|jry|jr|jsJy|jryy)Nrrrr} pointwise)rrrrYrs r_get_heuristiczTritonKernel._get_heuristicsD   !  ' '*  & &(( (()  " "rcTt}|jd ddl}|jdt x}r|j ||jdt jjr|jd|jS#t$rYtwxYw)NzP import triton import triton.language as tl rzU 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') ) rYrtriton.language.extra.tlx ImportErrorrr%r#rproton_profilingr)rimportsr attr_descs rgen_common_triton_importsz&TritonKernel.gen_common_triton_importss!"   , NN 34 49 4   i (   == ) ) NN !!+   sB B'&B'cxtj |jtjjdS)N)enable_fp_fusion launch_pdl disable_ftz)r#emulate_precision_castsrrrWrs rtriton_meta_commonzTritonKernel.triton_meta_commons5%+$B$B B113!00<<  rc~tjjjtj tj tjjtjtjtjtjtjtjjtjjtjj tj"tj$j&tjj(d}tj*rtj,|d<tj.j0d|d<t j2rd|d<tj4rLtj4|d<tj6|d<tj8|d<tj:|d <tj<r9tj<|d <tj>|d <tj@|d <|S) N) backend_hashassert_indirect_indexingautotune_local_cacheautotune_pointwiseautotune_remote_cacheforce_disable_cachesdynamic_scale_rblock max_autotunemax_autotune_pointwisemin_split_scan_rblockspill_threshold store_cubinrr&mix_order_reduction_allow_multi_stages$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`rarbrcrdrerfrrrrg*write_are_deterministic_algorithms_enabledrhrrrrjrkrlrmrnrorp)r inductor_metas rinductor_meta_commonz!TritonKernel.inductor_meta_commons"KK//HHJ(.(G(G$*$?$?"(--"B"B%+%A%A$*$?$?$*$?$?"//&,&C&C%+]]%H%H%}}<<!==44#11.4.A.A.`.`68W8WM4 5FF K L  + +00 5 677 < =>> C Drc -.t}i}|jjD]_\}}t|r |jst j jj|}tt|}|||<a||j|jt j jj}|dk(r|jdn|jdtj r|j|j#|j$j'\-} .} t).D]\} } t+| t,st/t0j2| j4} | t j jj6vsbt-| j8t j jj6| .| <t;} |j<D]}||j$j>vr(| jA|j$j>|||j$jBvrj|t j jDvrN||jDvr@| jAt/tF|j$jB|jH||j$jJvs|j$jJ|}t+|tLrJ| jA|tO-.D]O\}} t+| tPs| jRtTjVk(s5| jA|j8QtY| } |j[D]Z}t-|j\d|j^}.ja|-jatc|j8\-.fd}|jdD]K}|jfr |jhr|jj)||j\jmdM|jnr|d|jpr|d |d ts.|jt- }|twjxt j jitzj|jj~jxr0d t|jvxsd t|jvd |j}|jnrd|d<t j jxst j j}|jjt|jttj| ||j|j|j|j|jd |j}|jpr|j|d <tjstjjr|jxs |j|d<t jjjgxr |jh }|j}t|jdk(}|r|r|jj|j}|jjjd}|j}|'d|vr#d|vr|dt|ddz }|tk\}n,|jj|tjk(}|jjj} |jjj}!t j jj| tt j jj|!dz }"|"dk\r|rt j jj|jjdrHt j jj|jjdr |dkrd|d<|jr|j|d<|jr|j|d<|jnr|jh|d<d}#tj stjr|jdz }#|#|#|d <tj r|j}$|$|$|d!<t.D]}%d|d".|%j8<||_r||_s|j|j|j|j|jtzjj|jrt.d#$g|d%<nt.g|d%<|jD]$}&|jd&|j|&&|jr2d'|jd(|jjd)|d*|d+ }'n|jrF|jj|j}(d'|jd,|d-|(d)|d*|d+ }'nYd&})t|dk(rtt.d.k(rd/})nd0})d'|jd,|d1|)d2|d*|d3|jd+ }'|j|'|xsttj}*|jd4|*d5d1jd6-Dd7|j5tj~jr|jd8|*d9|j ||j$j D]\}+},|j|+d:|,|j|jtj~jr|jd;|*d9dddtj r!|j|j |#|jS#1swYKxYw).add_constexpr_argFs/-/  h!78 NN78$? @rrr RSPLIT_SIZE NUM_STAGES) size_dtyperzztl.dot)rr0r/ native_matmulTlaunch_cooperative_grid) grid_typer kernel_namemutated_arg_names optimize_memrVrnum_load num_store num_reduction"has_loadstore_with_contiguous_rdimr"rrNr0_rQg?ii add_persistent_rblock tiling_scoresrrgeA kernel_num_gb kernel_flopr/r)pointer_range_overrideconfigsrz$ @triton_heuristics.z( config=zI, filename=__file__, triton_meta=z$, inductor_meta=z; ) @triton.jit z!( size_hints=z%, reduction_hint=rXztile_hint=TileHint.SQUARE,ztile_hint=TileHint.DEFAULT,rzH filename=__file__, triton_meta=z*, min_elem_per_thread=zdef rc3<K|]}|jywr) full_namers rrz.TritonKernel.codegen_kernel..s*JQ1;;=*Jr2rzpl.enter_scope("rrzpl.exit_scope(")rYrZrrErYrOrPrIoptimization_hintr6rrrSrr"r#benchmark_kernelrIrr)r!rr_r rrrinv_precomputed_replacementsrr mutations input_buffersr'rzremoved_buffersrZ inner_nameoutput_buffersr^rra zero_moderb ZERO_ON_CALLr'rrrrrSrer\rrrrrrnrr1rlrrrrrr8rrY is_inference is_backward_get_grid_typersetrrDDESCRIPTIVE_NAMErVrrrrru rsplit_sizerrrrrrr[rrr memory_stats persistentmemoryrgcount_per_threadr~INNER_REDUCTION_RATIO_THRESHOLDget_reduction_hintr2INNERloopedbytesstatically_known_leqreduction_numelstatically_known_gtrrjestimate_kernel_num_bytesestimate_flopsrkrrtrrrrrrjrr%rrLrlrr6r&rrPcodegen_static_numelsaliasesrCr)/rrr size_hintsrr numel_hint size_hint device_typer(r,rr mutated_argsmutation mutation_argargnamersizeargr{triton_meta_signaturerrrt looped_redr two_d_redr dim_statsmem_ops_per_threadr_coalesce_ratiocontiguous_red looped_mempersistent_memsaved_bytes_ratior7flopsarg_numrheuristics_linereduction_hint tile_hintroldnewrzrs/ @@rcodegen_kernelzTritonKernel.codegen_kernels4  ![[..0 +MFE"6*43H3H));;EBJ'J8I!*Jv   + < KK668 9''==?DDKe# @A @A&& D==?@#'99#;#;#= Iq * FAs#w'ellCHH5QWW--JJJ#*!''"2"2"O"OPV"W$IaL  )3  /H499222  !8!8!BCDII555AGG$;$;;D$8$88  )B)B8)LMXX499333#yy77A %lJ???  . /6 3 /LGS3 -MM%6%C%CC  .  /l+ ++- 2D U3TZZ@G   W % NN77<<0 1 2 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 B B: NQT  22>BR! "S("GG$$99MM115GG$$88MM114'",9= 56   -1-?-?M/ *  # #373K3KM/ 0  % %484M4MM0 1  " "f&>&>335;F!17 o.  " "'')E /4 m,+95 BG@AK $Yw%7%<%< = B'* dii(  # ==   (T-B-B&/ RT&U%VK "&/ &:%;K "++ F NN2  KK     #$$($7$7$9#:; --4478!!,0##0"34 O " "!]]==d>P>PQN#$$($7$7$9#:; *~.$$2#34!!,0##0"34 OI:!#/ :;q@ }} B Bs :Cs  sctjjj|}t |t j tfrt|}t|}|Sd}tjjj||sC|dkDrtd||dz}tjjj||sC|S)NrQi@z!Failed to find static RBLOCK for r") rOrPrIsimplifyrrrrr6rrL)rrs rrz#TritonKernel._get_persistent_RBLOCK(s!!**62 fu}}c2 3f+C!#&C Cgg&&;;FCH?$'H%QRRqgg&&;;FCH JrcN tj|y#t$rYywxYw)NTF)rrrL)rs rhas_persistent_RBLOCKz"TritonKernel.has_persistent_RBLOCK9s*   / / 7  s  $$cd d}|jD]e}|jr |jrdtjj j |j}||r)|j|jdt||jr|jr|jr1|j|j|j}d|d}n3|j|j}|j r t#|d}|j|jj%d||jdk(sG|j&sU|jdhy ) a 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. cBt|tjtfSr)rrrrrs ris_static_integerz=TritonKernel.codegen_static_numels..is_static_integerSsdU]]C$89 9rznumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)rzBLOCK: tl.constexpr = rNzXBLOCK: tl.constexpr = 1N)rrrr)rer\rYrOrPrIrrr%rrrrrrrrr~rrV)rrrrsimplified_tree_numelrrs rrz"TritonKernel.codegen_static_numelsAs,$ :$$ ;D$$(=(=()(8(8(A(A$**(M%$%:;NNdkk](3?T;U:V#WX  T%>%>-- JJt';';DJJ'GHEFugMfgC55djjAC,,!#rl$++"3"3"5!66LSERS{{c!dmm9:' ;rcJt|jDcgc]}t|j c}}|jr|dk(sJt j S|jr|dk(sJt jS|dk(rt jS|dk(rItt|j|jrt jSt jS|dk(r,|jrt j St j"St%d|cc}w)NrQr"rz"Unsupported number of dimensions: )r7rerr\rr.MixOrderReductionGridrCooperativeReductionGridGrid1Dr}r(r|Grid2DWithYZOverflowGrid2DrBatchMatmulGrid3DGrid3DrL)rrrcs rrzTritonKernel._get_grid_typeks 8H8HI***+I J  # #6M6$:: :  ' '6M6$== = !V$++ + !V3t22D4D4DEF(===$++ + !V$$(:::$++ +=aSABB#JsD c|jD]}t|jtjtj fr |j}n*t jjj||}|jr |js|j||jt|yr)rerrrrrrOrP wrapper_codegenerate_numel_exprr\rYrr")rrr9 arg_typesrrs rr z#TritonKernel.add_numel_to_call_argss$$ -D$**u}}ell&CDzzww++??dK$$(=(=  &  d, -rctjj}|j|jj \}}}}|j ||||jjD]}|j||j||d||j|j|r|jyy)NT)rrrrt) rOrPrwrite_triton_header_oncerr)r workspace_argsgenerate_workspace_allocationgenerate_kernel_callrrtdeallocate_workspaces) rrr deallocate_wswrapperr(r9rwss r call_kernelzTritonKernel.call_kernels''&&((*%)YY%=%=%?"9a ##D)Y?))** 6B  1 1" 5 6 $$  ((,, %    & & ( rctjj}|jj \}}}}t ||D]w\}}t |tstjjr|jd|d|dJd|d}|j|d|d}|j|yy)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()) rOrPrrr)rrr` cpp_wrapperr%)rrr(r9arg_signaturesr arg_signaturers rcodegen_nan_checkzTritonKernel.codegen_nan_checks''&&*.))*B*B*D'9na"%i"@ , C-377&&%%TUXTYY\]`\aade)-BCD%%d+(-BCD%%d+ ,rct|i|Sr)rn)rrrws rrzTritonKernel.create_cse_vars $1&11rcF|jd|j|j|j}|jj s|j r+|jdk(r|jj|y|jj|y)NrrN) rrrrrootrrrrr%r8)rr rs rcodegen_iteration_ranges_entryz+TritonKernel.codegen_iteration_ranges_entrysy**SD,@,@,L!M NO ::  $":":u||s?R    ( ( . II   %rc||jJ|j|j}|jr$|jdk(r|jdk(rd}|j}|dk7rd|dnd}|j r|j r|jr|d}d|jjd ||S) NrrXrrrrz + rsplit_startz tl.arange(0, zBLOCK)) rindexing_size_strrrrrrr\rr)rr rrrs rrz)TritonKernel.iteration_ranges_ranges_codes+++%%e&6&67  ! !  A%'')Q.D&& *5*C4 }A&  & &))""x/Fu||1134F4&IIrc^|j}|j}dg|z}d|d|d|dS)NrQrrr)rr)rr rrrrs riteration_ranges_scalar_codez)TritonKernel.iteration_ranges_scalar_codesC&& &&(sTz$r%;-q99rc|jJd|jd}|j|r#d|d|jdzd|jd}|jr,|jdk(rd dd d }d||jd}|jj ||}|j d k7r|d |j dS|S)Nztl.program_id(rrz + tl.program_id(rQz) * tl.num_programs(rrXr"r)rrQr"rr)rr|rr pid_cacherr)rr rreversed_pid_mappids riteration_ranges_get_pidz%TritonKernel.iteration_ranges_get_pids~~)))u~~.a0  & &u -cU+ENNQ,>+??STYTbTbSccefC  T%<%<%>!%C#$q1 "#3ENN#C"DAFCoo!!#s+   z )U$t//02 2 rc|jrtjry|jdk(xr[|j xrL|j xr=t jjj|jt S)NFrQ) rr#!combo_kernel_per_subkernel_blocksrhas_zdimrrOrPrIrrr5)rr s rr|z#TritonKernel.needs_yz_grid_overflowsw   F$L$L NNa  YNN" Y... YGG$$99%++~GWXX  rc|jr |j|jdSt|jS)Nr)rrr3)rrs rrzTritonKernel.max_blocks;   $$ '7u%=> > //rc|jr5tjjj |j dry|j sy|jrW|jjd|jvr.|j|jjddk(rZy|jsMtjjj|j dr|jr |jsy|jr~|jrr|j|j }tjj}t!j"|j$xsd}t'|t(r?||kr:y|jdk(r|j*rd}n|j-|j}|jr|j.r||j1z}tjjj3|j |r[|j4dk7xsJ|j6xs<tjjj9|j t;Sy)NrFrrQT rN)rrOrPrIstatically_known_ltrrrrrrrQr\rrrr1rl warp_sizerrrVrrrr~rrrr5)rrrr0rs rrzTritonKernel._has_constant_masks  ww33DJJC!!   DKK$5$5$7#8!>$BSBS!S  DKK$5$5$7#8!>?1D%%ww77 AF))d.G.G   !:!:33DJJ?I WW88:F(//7AAGRI)S)i).C [[C DMMIt{{3I   !;!;!DOO$55I 77   8 8Y O "W==W77##88^EUV  rcd|jd}|jdk(sJ|j|S)NrrN)rerr)rxtrees rrz TritonKernel._has_constant_xmask?s5  #||s"""&&u--rc|jD]2}|j|s|j|jd4|jdy)Nrr)rerrr)rr rs rr8zTritonKernel.filter_masksDsN$$ 8D&&t,!!T[[M"67 8 &!rcttjd|jDcgc] }t| c}Scc}wr)rarrrXr)rrs rget_reduction_prefixesz#TritonKernel.get_reduction_prefixesLsB]::;U>UV  t    s<c|jDcgc]}|js|}}djtd|D}|j d|j ||jDcgc]+}|jrt j|j-}}t|}|j d|j |ycc}wcc}w)z^ Generates code that flattens ND reduction numels, block sizes, etc. into 1D. rKc3:K|]}|jdyw)rNrrs rrz8TritonKernel.codegen_reduction_numels..Ys"UTdkk]%#8"Urz rnumel = zRBLOCK: tl.constexpr = N) rer\r&r'rrrrrrG)rrrreduction_treesr rn_blocksrs rrz%TritonKernel.codegen_reduction_numelsSs -1,<,<RD@Q@Q4RRF"U_"UUV $**V"4!567 ((     % %dii 0  y) / 60B/CDES  sCC90Cc ||j}|Dcgc]}tj||fi|c}Scc}w)zK Helper to initialize symbols like rn_numel, rn_base, etc. )r rr)rrrw rn_prefixesrs rrz#TritonKernel._get_reduction_symbolses=113 JUV xx0;F;VVVs!9c|j}|jddd}tt|dz Dcgc]}t ||dzdc}t j dgzScc}w)z Compute coefficients to convert ND reduction indices to linear indices. For example: rindex = r0_index * r1_numel * ... * rn_numel + ... + rn_index. rTrrQN)r rrcrrGrr)rr rn_numelsrs r_get_reduction_index_coeffsz(TritonKernel._get_reduction_index_coeffsls{113 //PT/U ;@[AQTUAU;V 47M)C!GI. / ]]1    sA0c:|j}t||S)zK Compute linear reduction indices from N dimensional ones. )rrF)r multi_indscoeffss rrz'TritonKernel._flatten_reduction_indicesys113,,rc$|jddd}|jddd}|j|}|jd|j||j|}|jd|j|y)zX Generates code that converts ND reduction indices into linear indices. rTrr z roffset = z rindex = N)rrrr)rr rn_offsetsrn_indsrrrindexs rrz&TritonKernel.codegen_reduction_indicess 00 d1 --gtQU-V11*= 4#4#4W#=">?@009 $"3"3F";!<=>rc|j}|jr%|j|jd|d|dn|jD|j|jd|j ||j|dn|j |d|j |}n|j||d}|js|jdnd}|j|d|j|d ||jd|g|j|r!|j|j|y|d k(r |js%|j|d |jd |d yy)Nrz offset + r z offset = 0rrr|z offset = rKrNrrr)rrr%rrrrrrrrrrcreate_constant_mask)rr rrNr block_sizes rrz,TritonKernel.iteration_ranges_codegen_headers LL == NNejj\QCy4@ A ^^ # NNejj\T-N-Nu-U,VW X NNaS + ,+Id&G&G&N%OP881#V M,0+C+C1779+U#  OOc4#@#@#G"HJ<Xzzl#dV,   " "5 ) NN444U; <s(t77 NNaS |3qc? @ 8r)rTNNF) rdict[str, sympy.Expr]rzFixedTritonConfig | Nonerz int | Nonerrrrr*rr+)r rr zstr | tuple[str] | NonerzTMACompatibilityChecker | None)r)rrrrrz)BlockPtrOptions | TensorDescriptorOptionsrztuple[str, str])rrrrrrrr)rrY)rrrr)rrr rr rrr rrrWrrNrrr-)rCSEVariable | NoneNN)rrWrz.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]rrWrryrrrztuple[str, sympy.Expr] | Nonerr rrW)rztuple[str, BlockShapeType])rrWrryrrW) rryrryrrMr%CSEVariable | tuple[CSEVariable, ...]rr")rryrJ)rrr rrrW)rtuple[CSEVariable, ...]rtuple[torch.dtype, ...]rr)rr$rozUCallable[[tuple[CSEVariable, ...], tuple[CSEVariable, ...]], tuple[CSEVariable, ...]]rr#rr#) rr$rr#rrrrrr#)rr)r7z float | NonerrY)rz type[triton_heuristics.GridExpr]r)rrrz IRNode | Nonerr)rrn)r re)r rfrr)r rfrr rr)r rfrr)rrrr)rrfrr)r rrr)rrYrr)rrrzlist[sympy.Symbol]r)rrOrr)r rfrrYrr)frrrrrrrrrrSrrrrrrrrrrrr,rrrrrrrrrr rrurrrrrrrrrrrrrrrrr r rr}rrPr<rWrTrrVrrrr>rrr%rCrIrLrrrSrYrurrrrrr rrrrrrrr|rrrr8r>r rrrrrrrzr{s@rrr sl&I%%).E &.O$;!3B0B"#@HL7K 15$( %39%39 / 39 " 3939 39j::6  " #J*:0  ""/3DHb b , b $Bb R f'f'f'< f'  f'P#XJMMM M  M4 *H8J    :?-<1  "  :H {|SW\\ *\3>\FO\ \|I &15-199C9& 9 $ 9  9.9+9 9v. 9 ( 1<  (\\\& \ 5 \ / \|&DO&" 0F P#0 &!  B>>> >@>S1>S;R>S >S@"'" " ( " !"BD"'D"(D" D"  D" ! D"L $q%f0wr  t_%"%"N  --^yv  (;TC( -LP)),)DH). ,2 &J*:(:14: :.  0 9v. "  F$W  -? A(A0>A ArrcfeZdZdZ dfd ZddZ d d fd Zd dZxZS) FusedUserDefinedTritonKernelz When fusing a user-defined triton kernel with epilogues, we use this class to generate the modified triton kernel source c ^t|||ddddd||_t|jjj t jsJ|jjj |_|jjsJt|jjjdk(sJtj|jjjdj|_|j j#dd|_y)NrTF)r[rrrrrrQr(r)rqrrscheduler_noder kernel_noderr$UserDefinedTritonKernelir_nodecan_fuse_epiloguer kernel_storesrastunparsestore_value_nodeoriginal_stored_exprr#)rrr[r(rts rrrz%FusedUserDefinedTritonKernel.__init__s   !!  -    + + 0 0"2L2L   483F3F3R3R3W3W ||--///4<<--445:::$'KK LL & & - -a 0 A A% !%)$=$=$E$EdB$O!rct|jjdk(sJ||jjdjk(r|j}|j |}t jj|}|jr |j}ntj|j}|jj|j|||}|St!d|d)NrQrrz!Epilogue attempted to load from 'zE'. Inductor indexing variables are not defined in user kernel scope. )rr+ mutable_argsget_namer1rrOrPryr rrr rrrr`)rrr  loaded_exprrrrrs rrz!FusedUserDefinedTritonKernel.loads4<<,,-222 4<<,,Q/88: :33K}}U+HGG%%d+E$$ --%55hnnE** KuE+J !3D6:UU rct|jjjtj sJ||jjjj k(r||_yt|%||||yr) rr(fused_epiloguerr$ComputedBufferr4new_store_cse_varrqr)rrr rrrts rrz"FusedUserDefinedTritonKernel.storesl$--<<AA2CTCTUUU 4&&55::CCE E&+D " GM$ud 3rc$|5|j|jjj}|jjj |dddt j |jj}ddl m }m }||}t|jdk(sJtj|j j"} dd}||jdj$dd|tj&|}|j)} ||}|jdj$j*dz } d|jdj$j,z} |j.j1D cgc]} | | z } } |j2j1D cgc]} | | z }} | d| | z|z| | dz}dj5|S#1swYxYwcc} wcc} w) Nr)identify_triton_storesidentify_triton_stores_from_astrQct|j|kDr||j|<|jD]}|j|k(s||_yr)rrkeywordsrr) call_noder<positional_indexnew_argkeywords r _replace_argz:FusedUserDefinedTritonKernel.codegen.._replace_arg sO 9>>"%553: /0%-- ,;;(*$+GM ,rrr!r()r?zast.Callr<rr@r)split_and_set_rangesr(r7 get_rangescodegenrUdeepcopyr+ kernel_ast*torch._higher_order_ops.triton_kernel_wrapr;r<rrr.Namer9r store_noder/ splitlineslineno col_offsetr get_lines_refrr&)rrXnew_astr;r<r-new_store_value_noderCsrc_with_store_replaced src_linesstore_line_index indentationsr load_lines compute_lines new_src_liness rrFz$FusedUserDefinedTritonKernel.codegens  C22##22==?J    . . 6 6z B  C-- 7 78 8@ =''(A---"xx(>(>(C(CD , , ," ,    # . . 4FO 4  4<(rr&c leZdZUdZeZded<eejejejejejejejej gZdfd ZeddZddZdZdZddd Z d dd Z dd Z dd Zd ZxZS)TritonSchedulingz5Scheduling backend for Triton kernel code generation.z type[Any] kernel_typect||| t|dsy|jD]$}t |t t fst|_&y)Nr) rqrrrrrr;r9debug_triton_codedebug_device_str)rrrrts rrrzTritonScheduling.__init__FsM #  GIw$? OO :D$0B CD(9% :rctjjstjjr't g|j t jS|j Sr)r#rcooperative_reductionsforce_cooperative_reductionsrbackend_featuresrTREDUCE_TO_SINGLE_ELEMENT)rr0s rget_backend_featuresz%TritonScheduling.get_backend_featuresNsR MM 0 0}}99P#&&P(O(OP ###rc tjj}t||\}}|r|j |t j rtddlm}m t fd|DsX|Dcgc]}t||r|j }}|j |jddj||rt||} |j!|| yycc}w)Nr)r7ForeachKernelSchedulerNodec36K|]}t|ywr)r)rrcrfs rrz3TritonScheduling.codegen_comment..es>? 189sz Fused node name list: r)rOrPrrB make_commentr# debug_fusiontorch._inductor.schedulerr7rfr}rr4commentr&r,write_provenance_debug_handle) r node_schedulerrorigins_detailed_originsr7rc node_names debug_handlerfs @rcodegen_commentz TritonScheduling.codegen_commentYs''&&%8%P""    )     CP+!!%67JJL  $$''>tyy?T>UV BL  1 1+| L s.#C$c |j||||||ryt} tjrtj ||| j d|d| j |dtjj} | j d| jdd|} |||\} } | d | zd z| zz } |j|| j| y) zDEmit kernel to wrapper, with support for external template handlers.Nzasync_compile.triton(z, '''Trz''', device_str='z')z# kernel path: r() emit_kernel_overriderY async_compileuse_process_poolrr%rrOrPrr" define_kernelr)rrrsrc_coder subs_namerm kernel_pathrBcompile_wrapperrAmetadata_commentrndetailed_originss r_emit_kernel_to_wrapperz(TritonScheduling._emit_kernel_to_wrapperzs  & &         (*  ) ) +  H 5!!$9)e"LMxt4<<>!!$5n6I6I5J""MN,[M:$7 w$O!!D7NT14DDDk?+C+C+EGWXrc tjj}||jvr|j|}|Stj j r$t|tj j nd}|r tjj||}t|dd}djd|||jg}tjjrtjjd|}||j|<tj jr|nd}|j!t#t$j&|}|j!t#t$j(|}|j!dd}t+t-|j/d\} } } |j1||||||| t2t5j6d rt5j8|| ||S) Nrrr(rtriton_z #pragma CMT#pykernel_metadata)rOrPr src_to_kernelr#rdescriptive_namesrArcustomize_fused_kernel_namerPr&next_kernel_suffix aot_inductormodel_name_for_generated_filesunique_kernel_namesr#rrDrr6r)r(r r~rBr%is_metric_table_enabledlog_kernel_metadata) rrxrmrrr fused_namekernel_categoryry _basenamer(rzs rrwzTritonScheduling.define_kernels''&& w,, ,!//9Kf_==22&mV]]5T5TU  YYBB:xX @J2ANO((?J8R8R8TUK""AA"(!4!4!S!S TTUVaUbc /:G ! !( +'-}}'H'H iI  ''K,H,H(I;WH''K,C,C(DiPH '' s;H(08>>;K1Ld(S %Iq+  ( (# ../@A++KhOrc|j|d}tj|}|j||t d|DS)z Benchmark fused list of nodes and return the execution time in milliseconds on randomly generated inputs. T)rc3<K|]}|jywrr4rrcs rrz9TritonScheduling.benchmark_fused_nodes..s:WA1::<:Wr2)rp)generate_kernel_code_from_nodesr*rbenchmark_codegened_moduler)rrn_spills_thresholdrxrs rbenchmark_fused_nodesz&TritonScheduling.benchmark_fused_nodessV 77PT7Ux(.. # :WQV:W0W/  rc 8 ttjj}t 5|j tjj 5d fd fd} fd}||n tdg}tjd|j| jfcdddcdddSj j j  j dj(}t+|d k(sJ|dj,|kDr t'd  nytjj } t/j0 fd |  t+j2dkDr' t/j0 fd t5|  z tjd| | jfcdddcdddS#t$rn}t j"j$rtjd||t'd  | jfcYd}~cdddcdddSd}~wwxYw#1swYnxYwdddy#1swYyxYw)z$Benchmark an already compiled moduleNc~jJtjjjddzSNrz .kernel_perf__file__ospathsplitextrsrcache_file_pathzDTritonScheduling.benchmark_codegened_module..cache_file_paths6||///ww'' 5a8>IIrc>}t|tyrr+r)rrmss r store_cachez@TritonScheduling.benchmark_codegened_module..store_caches&(T3r7+rc}tjj|r.t|5}t |j cdddSy#1swYyxYwr)rrexistsopenrreadrfdrs r load_cachez?TritonScheduling.benchmark_codegened_module..load_cachesM&(77>>$'d0r$RWWY/000s AA r%kernel src code for %s written to: %srz*Exception (%s) in compiling fused nodes %sinfrQc4jdSrv clone_argsrcallwrapped_jit_functionsrrz=TritonScheduling.benchmark_codegened_module..+ D!@!5!@!@$!G!JKrr0c"jSrrrrsrrz=TritonScheduling.benchmark_codegened_module..4s ? 4 ? ? Frz+The fused kernel for %s took %.3f ms to run)rrOrPrrr0rrrrrget_argsrrr Exceptionr#r.disallow_failing_autotune_kernels_TESTING_ONLYr launchersrn_spillsr/ benchmarkrr)rrrrpdevice_interfacerrrrr0rrrrrs ` @@@@@rrz+TritonScheduling.benchmark_codegened_modulesf4AGG4G4GH  R $  # #AGG$G$G$I JR $B J , )4 *i[:Q  II7   B~3<<'?R $R $R $B<<>D88D#&;;  (4)44d;A>?-66Iy>Q& &&|$$'995\<<>!**K!+==>Bk33F"6{B II=  Ms||#eR $R $R $N (==OO @ 5\ 3<<''cR $R $R $N (OR $R $R $R $R $sh.J"AI;< J(I;8HCI;. J I8 AI3I8I; J3I88I;;J JJc|jd}|xr td|jD}|j}|rddlm}|}|rd|d<|jdr d|d <d|d<t j|js|jd rJd|d <tjj||||}||i|}|j|||S) Nrc3<K|]}|jywr) is_split_scan)rrs rrz9TritonScheduling.create_kernel_choices..Gs( %)D   ( r2rQ)TritonSplitScanKernelFoverride_cooperative_reductionr>Toverride_persistent_reduction) contains_opr}scheduler_nodesr[triton_split_scanrrrrrrOrtriton_kernel_kwargsadd_multi_kernel_choices) rkernel_features kernel_args kernel_kwargsis_scanrr[rrs rcreate_kernel_choicesz&TritonScheduling.create_kernel_choices@s "--f5 C( -<-L-L-N( % +/*:*:  @/K >CM: ;  & &v .=AM9 :>CM: ;11/2Q2QR$(()HI II=BM9 : 66 +} k;];,,V[-PPrc |g}tjjs|S|jxr|j d }|j xr|j d }|r%|j |j|i|ddi|r|jj}tjjj|dr[|j |j|i|ddix}|r2|jr&|j |j|i|dddt|dkDr.|ddD]} |j| _|j!d|S) NrrFi)rrrQc|jSr)r)rs rrz;TritonScheduling.add_multi_kernel_choices..s q'='=rr)r#r multi_kernelrrrrr[r[rrOrPrIrrmust_keep_buffersr>) rrrrkernelsoptional_persistentoptional_cooperativerr9kernel2s rrz)TritonScheduling.add_multi_kernel_choiceses (.h}}))N$99 -BSBS +C ?  &;; MDUDU ,E A   NN    #38  __44Fww44VUC-T--$'8=E'5+E+ENN((((+rcfdfd}fd}dg}}d}tjj}t|tj_tjj} t| tj_t j dkD} t jdkD} |j|d| | d} | D]:\} }}|Dcgc]}|j}}|Dcgc]}|D]}|j}}}t|dk(r'||d\}}||z }|dz }|j|| J| jttj d } t#j$| t&j)d |j*|\'|z }|z }|jj*j-j.j0j2dj4}t|dk(sJ|dj6dkDrt9d xnUtjj;}t=j>fd | t=j>fd| t&j)dtd|D||z }|z }|jj*=|tj_| tj_|||fScc}wcc}}w)a[ Benchmark combo kernel partitions and return total execution time. Generates kernel code for each partition and benchmarks them. Single-node partitions use benchmark_fused_nodes(), while multi-node partitions use the combo kernel benchmarking path. Returns (total_ms, total_clone_ms, file_list). c~jJtjjjddzSrrrsrrz@TritonScheduling.benchmark_combo_kernel..cache_file_paths6<<+ ++77##CLL1!4~E Erc}tjj|rCt|5}t d|j j DcdddSy#1swYyxYw)Nc32K|]}t|ywr)r)rrs rrzNTritonScheduling.benchmark_combo_kernel..load_cache..s Eaq Err!)rrrrrrrrs rrz;TritonScheduling.benchmark_combo_kernel..load_caches^"$Dww~~d#$ZF2 E2779??3D EEFFFs .A,,A5c\}t|tdztzy)Nr!r)rrrms_clones rrz.store_caches&"$D s2w}s8}< =rrgT)subkernel_nodescustom_part_algorithmenable_autotune mixed_sizesonly_gen_src_coderQrrrc4jdSrvrrsrrz9TritonScheduling.benchmark_combo_kernel..rrrc(jdSrvrrsrrz9TritonScheduling.benchmark_combo_kernel..s;0;;TB1ErzDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc3<K|]}|jywrrrs rrz:TritonScheduling.benchmark_combo_kernel..s ) # ww66",-A"B"#''"<"<%/0G%H" 77!;;;a? ::%"&+#" ; -=? + (Hfj=GHT 0H H/?OeOAQZZ\O\OEO:!# 6z!} E G#!#  &' ''''K,C,C(DiPH""8,C II7   &D88D#&;;  0%00$7: ;,66Iy>Q& &&|$$q( %e ,X<<>!**K!'00E! IIV<<<   M NH h &N   S\\ *? +@#7%<"22C IOs L:7L?)rzScheduler | Nonerr)r0z torch.devicer)r{)rtuple[float, str])r{N)rpzOrderedSet[str] | Nonerr)rrvr list[Any]rrrlist[TritonKernel])rrrrrrrr)rrrrrr[rrrTFOREACH BUCKETIZEINPLACE_BUFFERSMASKED_SCATTER_WITH_INDEXSCANSORTTRITON_TEMPLATESTUPLE_REDUCTIONrbrrrrdrrr~rwrrrrrrzr{s@rrZrZ5s?)K)!  " "  $ $  * *  4 4        + +  * *  :$$MB(YT6p  OSW$5KW$ W$r#Q+#Q#Q& #Q  #QJ333& 3  3jp3rrZcg}|j}|t|tjsJ|r0|j$|j |j d|Sddlm}ddl m }|j}|J|jj|}t|t||fsJdt|t j"j%|5t&j(}|j+|j-j/}|t&_ddd|j |j d|j t1j2d|S#1swYRxYw)Nz" Unfinalized multi template bufferr)CUDACombinedScheduling)XPUCombinedSchedulingz]Scheduling backend should be SIMD or CUDACombined when generating debug Triton strings, got: z Triton code:z )get_template_noderr$MultiTemplateBuffermake_kernel_renderrr40torch._inductor.codegen.cuda_combined_schedulingr 3torch._inductor.codegen.xpu.xpu_combined_schedulingr r.r get_backendrir"rOrPset_current_devicer%generated_kernel_countrrr rFr) rlinesmulti_templater r r0backendold_generated_kernel_count triton_codes rr]r] st E++-N  !Z@V@V%WW W.;;C  ((JKL< L9  "!!!..,,V4 n&<>ST  llpqxlykz {  WW ' ' / H*1)G)G &!AA eg .HG * H   ( 67 X__[&9: L H Hs #A E==F)rr)rrrryrrr+)rrrr4rr4rrr)rryrry)rryrr)rzCSEVariable | Anyrr)rr rr)rr rr)rrrr)rrrzCallable[[_T], _T])rr7rr) __future__rr.rgrrUrBrrloggingrrPrrFabcrcollections.abcrrrrtypingr r r r rsympy.printing.precedencer rtorch._loggingtorch.utils._pytreer&_pytreerKtorch._dynamo.device_interfacertorch._dynamo.utilsrrtorch._prims_commonrrtorch.utils._ordered_setrtorch.utils._sympy.functionsrrrrrtorch.utils._tritonrrrutils._sympy.symbolrrrr utils._sympy.value_rangesr!rr#r$r%rur' codecacher(r)r*r+rr, ops_handlerr-runtimer.runtime.benchmarkingr/ runtime.hintsr0r1r2r3r4runtime.runtime_utilsr5r6rr7r8r9r:r;shape_propagationr< stream_utilsr=r>r?r@rArBrCrDrErFrGrHrIrJrK virtualizedrLrrMrNrOwrapper_benchmarkrPblock_analysisrRcommonrSrTrUrVrWrXrYrZr[r\r]r^r_r`rarbsimdrcrdrerfrgrhri triton_utilsrjrkrlrmrnrrotypesrpr}rr%torch.fx.experimental.symbolic_shapesrsrtrusimd_kernel_featuresrvrw getLoggerrr_logginggetArtifactLogger perf_hint_log schedule_log fusion_logrrrrrrrrarr.rrrrrrrSrVr\r^r`rbrgrlrnr~rrrr_initialize_pointwise_overridesrrr/rcr|rrrrrr&rZr]rrrr< s"    8844 0 $$C<?/ YX4))(FF;('.D6. CBB/$% L>&8 Bg!00<H ~~//*E ^^ - -h A  #$, &' &'(&'&'R 6 6  4$Q,Q,h    B ~~~B  ?4??8 =,==@+ +/+aQMaQH 3 &8 ;P 8( (V| &k| &B22D9//9J"OJ"Z$+$+N N N N b!!H ### %sU38_'<<=  AAAHN7A:/0N7AbnB(<B(JU3~U3p#r