^jѯ ddlmZddlZddlZddlZddlZddlZddlZddlZddl 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/m0Z0m1Z1m2Z2m3Z3dd l4m5Z5m6Z6m7Z7ddl8m9Z9m:Z:m;Z;mZ>ddl?m@Z@mAZAmBZBm Z ddlCmDZDddlEmFZFmGZGmHZHmIZIddlJmKZKddlLmMZMddlNmOZOddlPmQZQddlRmSZSmTZTmUZUmVZVmWZWmXZXmYZYddlZm[Z[m\Z\ddl]m^Z^m_Z_m`Z`maZambZbddlcmdZdddlemfZfddl mgZgmhZhmiZimjZjmkZkmlZlmmZmmnZnmoZompZpmqZqmrZrmsZsmtZtmuZumvZvmwZwddlxmyZzm{Z{m|Z|m}Z}dd l~mZd!d"lmZd!d#lmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZmZd!d$lmZmZmZmZmZmZmZd!d%lmZmZmZmZmZmZmZd!d&lmZer,dd'lmZdd(lmZdd)lmZdd*lAmZd!d+lmZd!d,lmZed-Ze j`eZejfjied.Zejfjied/Zejfjied0ZeDZCd1Zd2Zdld3Z dmd4ZGd5d6Zeddnd7ZGd8d9Zej~Gd:d;Zej~Gd<d=Zej~Gd>d?eZej~Gd@dAeZdodBZdpdCZ dqdDZGdEdFeZeǫjZdrdGZdrdHZdsdIZdrdJZdtdKZdudLZdvdMZdwdNZGdOdPeZdxdQZdydzdSZGdTdUeZejeՐjdRejeՐjdReՐjdVGdWdXeիZGdYdZZej~Gd[d\ZGd]d^Zej~Gd_d`ZGdadbeeeeeefzfZej~GdcddZGdedfeeZGdgdheZGdidjeZd{dkZy)|) annotationsN)abstractmethod)CallableIterableSequence) lru_cache)Anycast TYPE_CHECKINGTypeVar) PRECEDENCE)get_interface_for_device)identitypreserve_rng_state)is_integer_dtype type_to_dtype) OrderedSet)CeilDivFloorDivMinModularIndexing 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 get_warp_sizenative_matmul_persistent_rblock ReductionHintTRITON_MAX_BLOCKTRITON_MAX_RSPLIT)get_max_y_gridnext_power_of_2)BaseSchedulerNode$FusedExternTritonKernelSchedulerNodeFusedSchedulerNode Scheduler SchedulerNode)get_broadcasted_shape)get_raw_stream_name)_TMA_SUPPORTED_DTYPES cache_on_selfDelayReplaceLinedevice_supports_fp64get_bounds_index_exprget_fused_kernel_nameget_kernel_metadatais_welford_reduction Placeholderprefix_is_pointwiseprefix_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_indicesis_unaligned_buffer_nameselect_tile_hintshould_unwrap_unspec_argsignature_to_meta/use_uint8_triton_storage_for_cuda_float8_e4m3fn)SymbolicCallArg) ModuleTypeDtypePropagationOpsHandler)ShapeEnv)IRNode)BlockShapeType)SIMDKernelFeatures_T perf_hintsschedulefusionc<|dv}|rdnd}|dvr|d|dS|d|S)N)anymaxminprodtriton_helperstl)rr.2)reduction_type use_helpermodules i/var/www/ramen.bs-engineer-server.com/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}}}}ZeddZeddZedd Ze dd Ze dd Zedd Zy cc}}}}wcc}}}}w) TritonSymbolszU Stores sympy.Symbol instances and constants associated with triton codegen. offsetTinteger nonnegativeBLOCKrpositivec Jd}|j}|D]}t|tjrHtj j j|j}t|j}nt|tjtjtjtjtjtj frd}ntj j"j%|}|Jd|j|j&}tj j)} dg| z} |j*d}n2t-|j/|| |j*<t| }t1||}|J|S)NrzUnregistered range symbol: 1)rr r!TMPrUkernelcse varname_mapnamershape UNBACKED_INTSIZEPRECOMPUTED_SIZEINDEXFLOATUNBACKED_FLOATrange_tree_nodesgetroottriton_tensor_ndim tensor_dimstrget_block_sizer?) rr expr_shape expr_varsvarcse_var var_shapertreendimrs rget_block_shapezTritonSymbols.get_block_shapes\&( %%  FCc488,((,,22388<!'--0 %%II))JJJJ''   xx0044S9'Q+Fsxxj)QQ'yyxx224 ??*!#I-01C1CD1I-JE$//* %e I/z9EJ? FB%%%rc6|jjSr)r block_sizerrs rrzTritonSymbols.get_block_sizeDsyy##%%rc6|jjSr)r block_offsetrs rget_block_offsetzTritonSymbols.get_block_offsetHsyy%%''rc|jj|x}|jjS|jD]}t ||st |dcSy)Nmask)rrr mask_name block_typesr r)rrsymbolrsymts rmask_name_for_symbolz"TritonSymbols.mask_name_for_symbolLsh ++//7 7D D99&&( (OO 1Dfd+$T*+400 1rc|jj|x}|jjSt ||j Sr)rrr is_reductionr reduction_types)rrrrs ris_reduction_index_symbolz'TritonSymbols.is_reduction_index_symbolYsC++//7 7D D99)) )fc&9&9::rcDtfd|jDS)zHReturn whether any symbol in a composite block shape is reduction-owned.c3BK|]}j|ywr)r )rrrrs rrz;TritonSymbols.has_reduction_index_symbol..ds&   ) )&& 9 s)rr)rrrs`` rhas_reduction_index_symbolz(TritonSymbols.has_reduction_index_symbolas% ++   rN)rrrr})rrjrr)r TritonKernelr sympy.Symbolr str | None)rrrrrr)rrrrrr)rrrrrr!R0_INDEXR1_INDEXr XBLOCKYBLOCKZBLOCKrrSymbolr block_offsetsupper block_sizesrrrrrr r )rrrrs0000rrrs!$--!?@Odkk4;; VoVWK   ellj./v6RVWWM   ell$%%'( .t  K..`&&(( ! +7    ;!;+7; ;;  w s *D 8D rceZdZUded<ded<ded<ded<d ed <d ed <dd ZddZddZddZddZe ddZ y)IndexingOptionsr index_strOrderedSet[str] mask_varsr expand_strr _has_rindexrindexzSequence[int | str] | None expand_shapec,t|jSr)rrselfs rhas_maskzIndexingOptions.has_maskssDNN##rcJt|jtjSr)rr!r!rr$s r has_indirectzIndexingOptions.has_indirectvs"4::txx88rc|jSr)r r$s r has_rindexzIndexingOptions.has_rindexysrc:td|jDS)Nc3PK|]}t|jd yw)tmpN)r startswith)rrs rrz.IndexingOptions.has_tmpmask..}sJ43t9''.Js$&)rrr$s r has_tmpmaskzIndexingOptions.has_tmpmask|sJ4>>JJJrcdtfdtjjDS)Nc3K|]F}jD]5}t|txr|jxr|j |7Hywr)rrrr owns_mask)rrrr%s rrz,IndexingOptions.has_rmask..sU   tS ! Pd&7&7 PDNN4

j@g|zz } |jC|tjjD\}}|jDcgc]}j |d}}|j }|j|} tGjHtjjD|j}||tjjjS|||| |||| }|jU|||Scc} wcc} wcc}w#tJ$r3tMtOtQt7|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!rWFrx)stride_sorter_cls shape_env) rArBrDrrIrFrHrKrM)rYzIterable[sympy.Expr]rlist[sympy.Expr])+rUgraphrZrrRrPstatically_known_equalsall remove_dimsrrrno_x_dimprefixpoprrrrrrlowerrKrJaddnum_reduction_dimsinside_reductionrnumelsfeaturesrSOnemaybe_sort_with_stride_order _shape_envr' argsort_symAssertionErrorlistreversedrangerXcompute_boundary_check)rrArB range_treesr get_max_blockr^rMr[dimsingleton_dimsrrIincluded_final_shape_dimsrsymsym_strreduction_ndimrKstriderHrFrDresultrZs @rcreatezBlockDescriptorOptions.creates77## N#6<<0 $V^^4AG@R@R 9((C/ // OOA %/L!&& ;D$ +,,;C!#hnn.G*737J77S155c:; ;44))-.#ahhoo2F2WW!!..0 EGGKK=>9 9K!' C C/177;M;M!D!  GMnn E ?s$L+L0.L5'8L::9M65M6cDtj|}t|||iSzN Replaces instances of {symt}_offset with the new expression. rrrNr%r replacementrroffsets rreplace_offsetz%BlockDescriptorOptions.replace_offset&&  --d3$+ 677rc|tjD](}|j|tjd|}*|SNrrr rrrr%rrs 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)rUrarZrrrrrr9rneeds_yz_grid_overflowrvrrrbrRrrnZeror!rrPrstatically_known_multiple_ofrNrerrL)r%ryrxrZt block_to_maxneeds_overflow_grididxs rrwz-BlockDescriptorOptions.compute_boundary_check4s 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'.ps)   4 4QXXt D s.0)rrPr$s rr*z!BlockDescriptorOptions.has_rindexos" ((   rc"|jSr)r*r$s rr4z BlockDescriptorOptions.has_rmaskus  rcyrrr$s rr/z"BlockDescriptorOptions.has_tmpmaskxrrc4t|jSr)rrr$s rr&zBlockDescriptorOptions.has_mask{sD'')**rc| dfd }|j}|j}|rJ|jj|j}|jj|j}t ||D cgc]#\} } | rt j jn| %} } } |||| }| } |jjs4|s2t| t|k(r|jj| } tjj|xr7t| t|k(xrtfdt | |D} t|jr*| s(d|dtjj!|d}||j}|jjs}|r|jj"n|jj$}d|d|d}|Dcgc]}| c}|r |jn$|jj|j}||||}|fScc} } wcc}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 cTt|||}t|t|k7r||Sr)triton_reshapetriton_shape_dims)value old_shape new_shaper current_shapes r apply_reshapezKBlockDescriptorOptions.codegen_broadcast_and_reshape..apply_reshapes0 $E9i@F +/@/KK ) Mrc3pK|]-\}}j|dxsj||/ywrWN)rb)rpre_dimpost_dimrZs rrzGBlockDescriptorOptions.codegen_broadcast_and_reshape..sH&GX00!<G33GXFGs36tl.broadcast_to(, ) tl.trans(rrr Sequence[sympy.Expr | int | str]rrrr)rFrHrKrevertziprrnro is_identityrrUrarZrcrr index_to_strsort_idxrevert_sort_idx)r%r initial_shaperIallow_implicit for_storerrFrHrzis_broadcastingpre_broadcast_shapeimplicit_broadcast_shapesupports_implicit_broadcastr permute_dimsirrZs @@rcodegen_broadcast_and_reshapez4BlockDescriptorOptions.codegen_broadcast_and_reshape~sy2;H   7 8  .. 22 "00778L8LMO $ 2 2 9 9$:P:P Q  ),O=N(O $_+EGGKK 3  e]4GH#6 ""..'(C ,<< (,'9'9'@'@AT'U $ 77##&4' ( )S-= = *--E{)S $ t%% &/J"5'AHH,A,A/,R+SSTU ,M(( !!--""++''77   wba8E7CD!]1-DM$$''..t/C/CD eY <m##{ hEs (I8 Irr`)rAr@rBrrxlist[IterationRangesRoot]rrryCallable[[str], int]r^z"type[BlockParameters.StrideSorter]rMrrr?rrrrrr!rrrrrr)ryrrxrrr)rrCr;) rrrrErIrErrrrrz,tuple[str, Sequence[sympy.Expr | int | str]])rrrrrrLrMr=rrPrRrUrrrrrwrr(r*r4r/r&rrrrr?r?s  ))!!%%0/(,O%,Hd !!''####i i$ i / i # i,i>ii iiV88-78?C8 8 2 +2 /2   2 h$ !+p$p$,p$* p$  p$  p$ 6p$rr?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)rUrrrBrrRrPr7)r%rrfrs rformatzTensorDescriptorOptions.formats HH ! !''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0  ,DIIdO+>rNTrrrr)rrrrrrrrrs?rrc>eZdZ ddZddZdd dZd dZy) BlockPtrOptionscDtj|}t|||iSrrrs rrzBlockPtrOptions.replace_offsetrrc|tjD](}|j|tjd|}*|Srrrs rrzBlockPtrOptions.remove_roffsetsrrc 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) rUrrrUrrBrrRrPrDr7)r%rrrrUrrs 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. )rrrUrrrnr)r%rrblockradvances radvance_roffsetzBlockPtrOptions.advance_roffset;st**40 ,,  ##FFD9%%feggllDAB   sAA,Nrrrr)rr!rr)rrrrrrrrrrrrs688-78?C8 8 7:rrc|Dcgc]3}t|tr|ntjj |5c}Scc}w)zFormat mixed symbolic and pre-rendered Triton shape dimensions. Nested-reduction codegen builds reshape/broadcast shapes from both sympy expressions and already-rendered block-size strings. )rrrUrrrrzs rrrOs@PU HKz#s#)>)>s)CC  s8Ac>ddjt|dSN[r])r7rrs rtriton_shape_strrZs! tyy*51231 55rct|trt|tsJt|}t|}||k(r|S|Dcgc] }|dk7s | c}|k7rd|ddj|dSd}g}|D]G}|t |kr|||k(r|j d|dz }0|dk(sJ|j d I|t |k(sJ|d dj|d Scc}w) zUG3tyy'?&@CC C F" ]# # c0B(B MM#  1HC3; ; MM& ! " #m$ $$ $WAdii'( **.s  C(C(ceZdZd"dZd"dZed#dZd"dZd"dZd"dZ d"dZ d"dZ d"d Z d"d Z d"d Zd"d Zd"d Zd"dZd"dZd"dZd$dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZd"dZ d"dZ!d"dZ"d"d Z#y!)% TritonPrinterct|jdk(sJd|j|jddtjj dS)NrWlibdevice.trunc(r).to(rrr_printrUr index_dtyper%rs r_print_TruncToIntzTritonPrinter._print_TruncToIntO499~"""t{{499Q<89qxx?S?S>TTU V rct|jdk(sJ|j|jd}d|d|dS)NrWrz(libdevice.trunc(z) + tl.zeros_like())rrr)r%rrs r_print_TruncToFloatz!TritonPrinter._print_TruncToFloatsF499~""" DIIaL)#5');E7"EErcLttjjsyy)zReturn the Triton float type for scalar shape math. Uses tl.float64 by default but falls back to tl.float32 on devices that lack fp64 support (e.g. Intel Arc consumer GPUs). tl.float32 tl.float64)rDrUracurrent_devicerrr_get_scalar_float_typez$TritonPrinter._get_scalar_float_types$AGG$:$:;rc|jrtt|}|Stjrt j jr|}|S|j}d|d|d}|S)N tl.full([], rr) rrintr$ is_fbcoderversionhipr)r%rret float_types r _print_FloatzTritonPrinter._print_Floatsp ??c$i.C     EMM$5$5FC 446J b A6C rct|jdk(sJ|j|jdtddz }|j }|d|dS)NrWrAtom?.to(r)rr parenthesizer r)r%rrr s r_print_ToFloatzTritonPrinter._print_ToFloats]499~"""   diilJv,>,D E002 D A&&rc|j\}}|jr3|jr'|j|jdtddz S|j |}|j |}d|d|dS)N % rrz!triton_helpers.remainder_integer(rr)ris_nonnegative stringifyr rr%rquotdivquot_sdiv_ss r_print_PythonModzTritonPrinter._print_PythonModsrII 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_FloorDivs~II c   3#5#5>>$))VZ5G#5MN NT" C 26(#eWAFFrcP|j|jdtddz S)N / rr)rrr rs r_print_IntTrueDivzTritonPrinter._print_IntTrueDivs#~~dii 60BS0HIIrct|jdk(sJd|j|jddtjj dSNrWlibdevice.floor(rrrrrs r _print_floorzTritonPrinter._print_floorrrct|jdk(sJd|j|jddtjj dSr%rrs r_print_FloorToIntzTritonPrinter._print_FloorToIntrrct|jdk(sJd|j|jddtjj dSNrWlibdevice.ceil(rrrrrs r_print_ceilingzTritonPrinter._print_ceilingK499~""" TYYq\!: ;5AUAU@VVWXXrct|jdk(sJd|j|jddtjj dSr+rrs r_print_CeilToIntzTritonPrinter._print_CeilToIntr.rc,d|j|dS)Nz tl.sqrt_rn(().to(tl.float32)))rrs r _helper_sqrtzTritonPrinter._helper_sqrtsdkk$/00ABBrc |j|jd}|j|jd}|j}d|d|d|d|d S)NrrWzlibdevice.pow((rz), (r)rrr)r%rbaseexpr s r_print_FloatPowzTritonPrinter._print_FloatPows`{{499Q<(kk$))A,'002  eJC vRuA..rc|j|jd}|j|jd}|j|jd}d|d|d|dS)NrrWr# tl.where(rr)doprintr)r%rcpqs r _print_WherezTritonPrinter._print_Wheres_ LL1 & LL1 & LL1 &1#Rs"QCq))rct|jdk(r|j|jdSt|jdz}t|}|j||jd|}|j||j|d}t d||fD\}}|dvs Jd|dd |d |d |d |d |d |d |d |dS)zI Helper for max/min code generation. cmp: > or < rWrr#Nc3(K|] }d|d yw)r9rNrrr]s rrz6TritonPrinter._print_min_max_helper..s.!q1X.s)>Q ;;tyy|, ,$))n!4j KKTYYt_- . KKTYYst_- ..1v..1j C$rErRrUrXr[r]r`rcrergrirkrmrorqrsr{rrrrrs F 'F GJ   Y Y C U/$* 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)rOrQrs rtriton_compute_typerus *51 22rc^|tjk(rtj}t|S)z@Convert torch.dtype to triton type, with fix for storing tl.bool)rrint8rOr~s rtriton_store_typerzs"    u rct|r+|jr|jdkrtjSt |S)z0Implicit upcasts used for Triton reduction types)r is_signeditemsizerint32rQr~s rupcast_acc_dtypers05??u~~7J{{ u %%rc*tt|S)z:Convert torch.dtype to triton type, with reduction upcasts)rrr~s rtriton_acc_typers /6 77rc<|jdkxr |jS)Nr#)rrr~s rlow_precision_fprs >>Q  :5#:#::rct|tsy|j}t|tjr t |SdSr)rr]rrr)rrs rlow_precision_fp_varrs6 c; ' IIE&0 &D E "O%Orct|tr |jSt|tjj rt t|Syr)rr]rr _prims_commonNumberrrMrs rtriton_arg_dtypers>#{#yy#u**112T#Y'' rctjj xr>t|txr,|j t jt jfvSr) r$rcodegen_upcast_to_fp32rr]rrfloat16bfloat16rs rneeds_upcast_to_float32rsD 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)r%rboundsrr __class__s rrzTritonCSEVariable.__init__sH vuE:*4, E"EE  E"EE rc:|D]}t|tr&|jj|j9t|tj sTt jtj|x} ||jj|yr) rrrupdaterrrrrUrri)r%rrkwargsrrs rupdate_on_argsz TritonCSEVariable.update_on_argssw 2C#01%%cmm4C. "/!C!CAHHc!RRINN&&y1 2rr) rrrzValueRanges[Any]r torch.dtyperr}rr)rrrrr __classcell__rs@rrrsH !% F F! F F  F  F 2rrcddlm}|S)Nrry)!torch._inductor.dtype_propagationrzrys rget_dtype_handlerrsL % ''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)rr)r upcast_strings rmaybe_upcast_argz.maybe_upcast_float32..maybe_upcast_args -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(,6rrr) itemsr itertoolschainvaluesrrrrrrO)rrr upcast_argskeyval upcast_kwargsrany_needs_upcast result_dtypeneeds_downcastdowncast_stringrrrs rwrappedz8maybe_upcast_float32..decorator..wrappeds<@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..decorators&&t^< 00rr<)rCallable[..., Any]rrr)rrrs` @rmaybe_upcast_float32rs '< rc~eZdZdZej ej Ze dY dZdZ ed[dZ edZ e dZ edZe d\d Zeed Zed Zed Zed ZeedZeedZeedZeedZedZedZedZedZedZedej>dddddZ eedZ!eedZ"e dZ#e dZ$edZ%eedZ&eed Z'eed!Z(eed"Z)eed#Z*eed$Z+eed%Z,eed&Z-eed'Z.ed]d(Z/ed^d)Z0ed_d*Z1ed+Z2eed,Z3eed-Z4eed.Z5eed/Z6eed0Z7eed1Z8eed2Zed3Z9eed4Z:ed5Z;ed6Zed9Z?ed:Z@ed;ZAed<ZBed=ZCed>ZDed?ZEed@ZFedAZGedBZHedCZIedDZJeedEZKeedFZLeedGZMeedHZNeedIZOedJZPeedKZQe dLZReedMZSeedNOdPZTeedNOdQZUeedRZVeedSZWedTZXedUZYeedVZZedWZ[eedXZ\y)`TritonOverrideszEMap element-wise ops to Triton e.g., ops.to_dtype(x,...) -> x.to(...)NTc>tjtjf d fd }|>t|||tj j tj _|!t|rd|d}tj}|tjk(rd|dS|tjk(r| |js||dS|r t|}n t|}|*|vr&|tjk(s t|r|d|dS|d|dS) Nc||k(ry|vr|vr ||k7rJd|tjk(s|tjk(ry|tjk(s|tjk(ryy)NrzCConversions between float8_e5m2 and float8_e4m3fn is not supported!rr#)r float8_e5m2 float8_e4m3fn) src_dtype dst_dtype fp8_dtypess r_get_min_elements_per_threadz>TritonOverrides.to_dtype.._get_min_elements_per_threadsI%Z'+* UU  U E---e>O>O1OE///9@S@S3Srz$triton_helpers.fp8e4m3fn_to_float32(rr9z != 0)z.to(tl.int16).to(tl.uint8)z.to(tl.float32).to(r)rrrrrr)rrrrrUrmin_elem_per_threadrvrruint8rrrr)r]rruse_compute_typesr out_dtypers @rto_dtypezTritonOverrides.to_dtypesD        " /:   .  ,/,Y>,,,AHH (  %T & 7qc;A I EJJ qc= ekk !  !i&A&AYEV S23 3 +E2I)%0I  !#ejj(,? ?,QCr!A6 6rcZtjjr d|d|dSd|d|dS)Nz tl.maximum(rrztriton_helpers.maximum(rrrs rr zTritonOverrides.maximumrrcd|d|d|dS)Nr@rrr)rPrQrBs rwherezTritonOverrides.where s1#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)NFr\ztl.loadz other=0.0T) rr9rKrrUrrxr_has_constant_maskr_cacher)rreduction_rangekvs ris_where_neededz,TritonOverrides.dot..is_where_needed/ss. >?hh2226O"// //xx**?;  ++113 !18 Q;!3C  !rctjjd|j}tj j Dcgc]}|js|j!}}t|dk(sJdtj|d||}tj jjtj j||j|jScc}w)NdotrWz'don't tile reduction when native matmulrrr)r% Reduction default_valuerrUrrxrrrTritonKernelOverridesrrgeneratecomputer)rdefaultrreduction_mask where_vars r where_condz'TritonOverrides.dot..where_condDsll00 BG-.XX-A-A%)TEVEV N~&!+ V-V V+-33N14EsGTI88<<((  )399CII) s C5C5c |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 xrrWrz, (rrz'right tl.dot operand cannot depend on yrr)rNotImplementedError) rrrIrzshape_2dbroadcast_neededRBLOCKrrrs r#reshape_transpose_broadcast_for_dotz@TritonOverrides.dot..reshape_transpose_broadcast_for_dotbs|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 bmmr"tf32ieeeztl.dot(rz, input_precision=""))rrErIrErr)rUris_native_matmulrdense_size_listrrrr!rrrrrr&r'rtrrrbackendscudamatmulfp32_precision) rPrQorig_aorig_brr+r1input_precisionr0rrrs @@@@rr!zTritonOverrides.dots:xx((((A * 2 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 3O3DBGGrrW) 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\\rKc .ddd|dddd S)Nztl.inline_asm_elementwise(rrz ], dtype=z , is_pure=z, pack=rr)r asm_literalasm_triton_typeconstraints_literalr?r@s rasm_callz8TritonOverrides.inline_asm_elementwise..asm_calls>,[M.s& AD-cU"TF! < s = z$ = triton_helpers.inline_asm_unpack()rOr7splitstripr. enumeraterrrr]rrrrr r reprreplacerUrr'rnewvarr writeline)asmr>rr?r@rAinputs_rCconstraint_partsinput_constraints cast_inputsrrNrBrL first_inputr'rr packed_argsrIrJrKs `` @@@rinline_asm_elementwisez&TritonOverrides.inline_asm_elementwisesN&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_expr!"STTrctd)Nz/ops.value_expr not implemented outside a kernelrerfs r value_exprzTritonOverrides.value_exprrhrctd)Nz+ops.masked not implemented outside a kernelre)rbodyothers rmaskedzTritonOverrides.masked#s!"OPPrcd|dS)Nzlibdevice.lgamma(rrrs rlgammazTritonOverrides.lgamma'#1#Q''rcd|dS)Nzlibdevice.erf(rrrs rerfzTritonOverrides.erf- s!$$rcd|dS)Nzlibdevice.cosh(rrrs rcoshzTritonOverrides.cosh3!1%%rcd|dS)Nzlibdevice.sinh(rrrs rsinhzTritonOverrides.sinh9rwrcd|dS)Nzlibdevice.acos(rrrs racoszTritonOverrides.acos?rwrcd|dS)Nzlibdevice.acosh(rrrs racoshzTritonOverrides.acoshE"!A&&rcd|dS)Nzlibdevice.asin(rrrs rasinzTritonOverrides.asinKrwrcd|dS)Nzlibdevice.asinh(rrrs rasinhzTritonOverrides.asinhQr~rcd|d|dS)Nzlibdevice.atan2(rrrr]rs ratan2zTritonOverrides.atan2W"!Bqc++rcd|ddS)Nztl.full([], 0x08Xz), tl.uint32).to(tl.float32, bitcast=True)r)bitss r_f32_const_from_bitsz$TritonOverrides._f32_const_from_bits]sSz)RSSrc >djdgdgt|zz}djtt|}tj j jtj jd|d|d|dtj|S) Nrz=frztl.inline_asm_elementwise('z', 'z', [z*], dtype=tl.float32, is_pure=True, pack=1)r") r7rr9rrUrrr&r'rr)rWrrr>args_strs r_inline_asm_f32zTritonOverrides._inline_asm_f32asii#T): :; 99Sd^,xx||$$ HH  )#d;-@zC E-- %  rc |jtjfdfd }dfd }tj}|d}|d}|d|d}||d|tj }tj d|g}|d |d |d |d} || d | } ||d | |d } || | |d} || | |d} || | |d} || | |d} || | |d} || | |d} || d | } || | | } |d| }||d|d|}|d |d |d | d} d}||d|tj}|| dtj}||d|tj}||d}||d|tj }|d |d |d | dS)Nctjjjtjj||SNr")rUrrr&r')rrrs rgenz(TritonOverrides._cuda_atanf..genxs688<<((  $e5) rc8tjd|||gS)Nzfma.rn.f32 $0, $1, $2, $3;)rr)rPrQrBrs rfmaz(TritonOverrides._cuda_atanf..fma}s#"22,q!Qi ri?irZrz > zrcp.approx.ftz.f32 $0, $1;r@rrvi ;lOaxi~9#=lz ){i!k=l"|i`L>l'*U}-in?it?z"tl.full([], 0x80000000, tl.uint32)z.to(tl.uint32, bitcast=True) & z.to(tl.uint32, bitcast=True)rz.to(tl.float32, bitcast=True)z <= )rrrrrr)rPr rQr rBr rr )rrrrrrruint32)r]rrrBoneinfabs_xlargercpzz2poly correctionr neg_resultpi_over_two_minus_result sign_masksign result_bits signed_bits signed_resultnot_nanrs @r _cuda_atanfzTritonOverrides._cuda_atanfmsH05      0 0 m ml1#Q'(ugS& 3-- (5'5  )E7"SEE7!4 5 A3c! 1Z="a m44Qz]+4Qz]+4Qz]+4Qz]+4Qz]+4Qz]+B4s4&)* ZA&1VH& #&q}a mZ#P yr*B)C2fXQOP8 aS7 {CU\\RVH$@A5<<P TF#k]3U\\B {m+HIJ tC5)5::6Ywir-6(!DEErct|}t|}tjjst j jjdk(rt|tr|tjk(s|rd}|rb|}t jjjt jj|dtj|j }t"j%|}|$|tjk7r|dt'|dS|S|rdnd}d||d}|r&|$|tjk7r|dt'|dS|S)Nr8rr"rrrzlibdevice.atan()rrrr r rUraget_current_device_or_throwrMrr]rrrr&r'rrrrO)r]r needs_upcastrrupcasts ratanzTritonOverrides.atans5!#.q1  !!335::fD1k*%--'<L$ HHLL))HH$$c)--'' * %003F'LEMM,I k,&?%@BBM&2""1#fXQ/ E-%5==2HXT+e"4!5Q7 7 rcd|dS)Nzlibdevice.atanh(rrrs ratanhzTritonOverrides.atanhr~rcd|d|dS)Nzlibdevice.copysign(rrrrs rrzTritonOverrides.copysigns%QCr!A..rcd|dS)Nzlibdevice.erfc(rrrs rerfczTritonOverrides.erfcrrcd|dS)Nzlibdevice.erfinv(rrrs rerfinvzTritonOverrides.erfinvrqrcd|d|dS)Nzlibdevice.hypot(rrrrs rhypotzTritonOverrides.hypotrrcd|dS)Nzlibdevice.log10(rrrs rlog10zTritonOverrides.log10rrcd|dS)Nzlibdevice.log2(rrrs rlog2zTritonOverrides.log2rrcd|d|dS)Nzlibdevice.ldexp(rz.to(tl.int32))r)r]ns rldexpzTritonOverrides.ldexps"!Bqc88rcd|d|dS)Nzlibdevice.nextafter(rrrrs r nextafterzTritonOverrides.nextafters&aS1#Q//rc|d|SNr6rrs r logical_andzTritonOverrides.logical_andCs|rc |dS)Nz == 0rrPs r logical_notzTritonOverrides.logical_notsE{rc|d|SNrrrs r logical_orzTritonOverrides.logical_orrrcd|d|dS)Nr9 ^ rrrs r logical_xorzTritonOverrides.logical_xors1#S1~rc|d|Srrrs r bitwise_andzTritonOverrides.bitwise_andCs|rc d|S)N~rrs r bitwise_notzTritonOverrides.bitwise_not s 1#wrc|d|Srrrs r bitwise_orzTritonOverrides.bitwise_orrrc|d|S)Nrrrs r bitwise_xorzTritonOverrides.bitwise_xorrrc|d|S)Nz << rrs rbitwise_left_shiftz"TritonOverrides.bitwise_left_shiftD}rc|d|S)Nz >> rrs rbitwise_right_shiftz#TritonOverrides.bitwise_right_shiftrrcd|d}tjr,tjj \}d|d|d|dSd|d|dS)Nr9).to(tl.uint32)ztriton_helpers.rand4x(rrztl.rand(r_can_use_4x_randomrUrr6seedrblocks rrandzTritonOverrides.rand s`VHO,  - - /xx//1HU+D6F82eWAF F$r&++rcttjtxrttjtj j j xratjjjdk(xr4tjjdk(xrtj S)Nr8rW)rrUrrr _inductorselect_algorithmTritonTemplateKernelrarrMrr$align_random_eagerrrrrz"TritonOverrides._can_use_4x_random(s qxx . .%//::OO .335::fD  . ++-2  . --- rc nd|d}d|d|d}d|d|d|d}d|d|d}d |d |d |d |d |d S)Nr9rz((z)*(rz)//(z//4))z)%(z!triton_helpers.rand_eager_kernel(r+z, VEC=rr) r base_offsetthreads_per_roundtidvectid_u32denomrC tid_truncs r rand_eagerzTritonOverrides.rand_eager4scU/*SE./r2 eWCuE 2 UG2. 24&;-q2i[X^_b^ccdeercd|d}tjr,tjj \}d|d|d|dSd|d|dS)Nr9rztriton_helpers.randn4x(rrz tl.randn(rrs rrandnzTritonOverrides.randn>s`VHO,  - - /xx//1HU,TF"VHBugQG G4&6(!,,rc ,d|d}d|d|d|d|d S)Nr9rztriton_helpers.randint64(rrr)rrlowhighs r randint64zTritonOverrides.randint64Fs1VHO,*4&6("SED6KKrctd)Nz.ops.load_seed not implemented outside a kernelre)rrs r load_seedzTritonOverrides.load_seedKs!"RSSrcNtjjrd|dSd|dS)Nz tl.rsqrt(rzlibdevice.rsqrt(rrs rrsqrtzTritonOverrides.rsqrtOs/ ==  qc# #%aS* *rcd|dS)Nzlibdevice.log1p(rrrs rlog1pzTritonOverrides.log1pXrrcd|dS)Nzlibdevice.tan(rrrs rtanzTritonOverrides.tan]rtrcRtjjjj |}|rt |dr |j }nd}tjrBtjjr(tdkDr|tjk7r|d|dSd|dS)Nr)rzlibdevice.fast_tanhf(rzlibdevice.tanh()rUrrrrrrr$rrr r rr)r]rrs rtanhzTritonOverrides.tanhcs((,,**..q1 ww0MMEE   !!"$v-&!+1#Q/ /$QCq) )rcd|dS)Nz tl.sigmoid(rrrs rsigmoidzTritonOverrides.sigmoidysQCq!!rcd|d|d|dS)Nz(libdevice.signbit(z ) != 0) if (z).dtype is tl.float32 else z < 0rrs rsignbitzTritonOverrides.signbit~s# "!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}|t jk(rVttj j"s3t j}|t jk(rt j}|j||}|j||} d|d| d} |6||k7r1t|r|r| dt%|d} | S| dt%|d} | S)Nztriton_helpers.pow_integer(rrr:r)rpowrrrrrrrint64rrrrrDrUrarrO) rrPrQrr5exponentr pow_dtypecast_acast_brs rr zTritonOverrides.pows(*..q!4  #(8(F**1l;Da!4!4!;!;< Q ,s  1b !D D215S9PQR9S U]]EMM: : $L15E ]]   %.B GG " "/  Iu}},$}} ((I6((I6!&F815  # (A -# &xtK ,E+FaHF #84 L(A'B!D rcNtjjrd|dSd|dS)Nzlibdevice.log(rz tl_math.log()r$ruse_pytorch_libdevicers rlogzTritonOverrides.logs3  6 6$A3a( (aS""rF)rcd|dS)Nzlibdevice.isinf( ).to(tl.int1)rrs rrzTritonOverrides.isinf"!M22rcd|dS)Nzlibdevice.isnan(rrrs rrzTritonOverrides.isnanrrcd|dS)Nzlibdevice.nearbyint(rrrs rroundzTritonOverrides.rounds&aS**rcd|dS)Nr&rrrs rfloorzTritonOverrides.floorr~rctjdtj}tjdtj}tj||}tj |||}tj ||}tj |tj|||}tj |tj|||}tj ||}tj |tj||}tj||}tj |tj||}tj |||SNrrW) r rrreqrltrrtruncdiv)rPrQzerorb_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))r rrrrrrr)r]rleftrightrs rrzTritonOverrides.signsw LLEKK (||SVVAq\EJJ7 cffQlUZZ8ggdE"d1#W%%rcd|dS)Nrrrrs rtrunczTritonOverrides.truncr~rc|d|S)Nrrrs rrzTritonOverrides.truncdiv s D}rcd|dS)Nr,rrrs rceilzTritonOverrides.ceil rwrNT)rrrtorch.dtype | None)rrrrrrrr)rrrr)rWrr Sequence[Any]rr}rr )r]r]rr])]rrrrrre_LOG_2_Er|rrrrrrrrrrrrr6rrrrrr rr!rrr_rarcrgrjrnrprsrvryr{r}rrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrrr rrrrrr$rr*rr-rrrrrsVOtyy H)- H&H&&H&H&T DD,<<  33## ,* ) )&&''"""77 77 **vHvHpmm CCJ####UUUUQQ((%%&&&&&&''&&'',,TT    4F4Fl:''//&&((,,''&&9900,,    ff--LLTT++ ''%%**&""  ++&&P##/303/303++''--6&&'' &&rrrceZdZdZfdZeejdZedZ edZ edZ e dZ e dZe d Ze d d ZxZS) r%a 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)rr_setup_libdevice_routing)r%rrrs rrzTritonKernelOverrides.__init__ s# $)&) %%'rcddlmtjjj D]}t ||sJt||}fd}|dk(rDt dsJtj|||}||_ t||t|kd}tj|||}||_ t||t|y)z.decomposition_router1 s977emm+)!,,>7#3X>qAGGGrr)r:r;c^|jtjk(r d|d|dS||S)Nz libdevice.r9r)rrr)r]r:r;s r dtype_routerzDTritonKernelOverrides._setup_libdevice_routing..dtype_router@ s277emm+'z1#Q77)!,,rN) torch._inductor.codegen.commonr8rrr'op_requires_libdevice_fp64rr functoolspartialrsetattrr|)rfn_name original_implr<fnr>r8s @rr6z.TritonKernelOverrides._setup_libdevice_routing& s D,,GG 4G3( ((#C1M H )#/;;;&&(QX& Wl2&67 - ""]WB"BK C,r"2 3; 4rcrtjj}dg|z}|j|||S)NrWr)rUrrr)rrrrrs rrzTritonKernelOverrides.constantL s9 xx**,d ##E5#>>rct||}tjj|dd}t |t sJ|j r |j }ntj|j}tjj}|tjtjfvr|}n|}t|}tjjj!tjj"|j%d|j&d||t)||}|j*|_|S)NF block_ptrtma_compatibility_checkerr9r)rrr)rrUrindexingrrr"rrr!get_index_dtype_as_torch_dtyperrr rQrr&r'rrrEr) rrrrLrr cast_dtype output_dtypers rrgz TritonKernelOverrides.index_exprU s/e<88$$ ET% (O444  ))E!11(..AE hh==? U[[%++. .$JJ +:6 hhll## HH   LL1X//02J ?(. $ !**  rc 2tjj}|tjtj fvr|ntj tj_ |j ||}|tj_||k7s|j|k7r_tjjjtjjd|dt|d||j}|S#|tj_wxYw)z Like :meth:`index_expr`, but honors ``dtype`` by setting the kernel index dtype before emitting, and casting the result if needed. r9rrr") rUr _index_dtyperrr rgrrr&r'rOr)rrrreal_index_dtypers rrjz TritonKernelOverrides.value_expr| s 8800u{{EKK88Eekk  5..u-C$4AHH ! u $ U(:((,,''  C5k%013ii (C  %5AHH !s #C??Dc |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)r"output)opz)graph for body does not contain an outputFloadrWT)rrz.shape, rr&)rrr)rr r rUrrr&r'rrra find_nodesrtargetrt mask_loadsris_boolrir"wraprr rrdiscard) rrlrmnodes need_whererrrnew_maskrr s rrnzTritonKernelOverrides.masked s   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)rUrrinput seed_offset)rrrs rrzTritonKernelOverrides.load_seed sIhhmm!!$'se3qxx}}889KVTUUV W rc0d|d}tjjj|x}r|Stjjj |j |j }tjjj tj|j }tjjj|d|d|dtjjj|||f||fS)Nzfrexp(rr"rz = triton_helpers.frexp() rUrrtry_getrUrrrrr'rVput)r] cache_keycse_valmantissar s rfrexpzTritonKernelOverrides.frexp sQCqM hhll**95 57 5N88<<&&QWWAGG&D88<<&&U[[&H ""j8*$?$+1-$D *U$s  C Cc|jSr)rr$s rrz(BlockParameters.StrideSorter.is_identity) s$$ $rcy)zBCreate a `StrideSorter` that can be used to sort block parameters.Nrrrr_s rrz#BlockParameters.StrideSorter.create- srcb|js|jDcgc]}|| c}S|Scc}wr)rrr%attrrs rsortz!BlockParameters.StrideSorter.sort4 s.##)-7AQ77K8 ,cb|js|jDcgc]}|| c}S|Scc}wr)rrrs rrz#BlockParameters.StrideSorter.revert9 s0##)-)=)=>AQ>>K?rNrzlist[int | sympy.Expr]r_r{rrJ)rrrr dataclassesfieldrrr=rrrrrrrrr StrideSorterzBlockParameters.StrideSorter s##%6[%6%6E%BB   %  %   U#9 UFN U ) U   U    rrc<eZdZfdZe ddZxZS)$BlockParameters.IdentityStrideSorterc"t|yrrrr%rs rrz2BlockParameters.IdentityStrideSorter.__post_init__@  G ! #rc L||ttt|S)Nrr)rtrvrrs rrz+BlockParameters.IdentityStrideSorter.createC s'!1eC(8$9:; rr)rrrrrrrrs@rIdentityStrideSorterr> s4 $  #9 FN  )   rrc@eZdZdZfdZe ddZxZS)+BlockParameters.TensorDecriptorStrideSorterzT Sorts BlockParameters dimensions with strides in descending order. c"t|yrrrs rrz9BlockParameters.TensorDecriptorStrideSorter.__post_init__R rrcttt|} 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)reverser)rtrvrr'rrrs)rrr_ identity_sortrs rrz2BlockParameters.TensorDecriptorStrideSorter.createU sa!s+;'M ) !,,Y8HRVW !1!  " )( )sA AAr)rrrrrrrrrs@rTensorDecriptorStrideSorterrL s9  $  #9 FN  )   rrc t|}td||fD\}}|di|Dcic]}|||||zc}Scc}w)z0 Concatenates block parameters. c3FK|]}tj|ywr)rasdictrHs rrz*BlockParameters.__add__..w sBq[''*Bs!r)rMr)r%rmrrPrQrs r__add__zBlockParameters.__add__r 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. )r_r)rrRr@rrrr)r%r^r_rKrrrAs rrpz,BlockParameters.maybe_sort_with_stride_orderz 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@rrr)r%rrrrs ` rrdzBlockParameters.remove_dims sO   5@5G5G5M5S5S5UVcsK$$V  VsA N)rmr@rr@)r^ztype[StrideSorter]r_r{rz4tuple[BlockParameters, BlockParameters.StrideSorter])rrGrr@)rrrrrrrtrrrPrRrU dataclassrrrrrprdrrrr@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|_yr) r current_loop prior_loop collections defaultdictdequeready_for_reuse loop_count store_count)r%rs rrz+CooperativeReductionWorkspaceCache.__init__ s@ *66{7H7HIrc|jj|}|r|jS|jj |d\}}}|j j |||f||fSr)rrpopleftr workspacerr)r%nbytescachedws_namerY ws_offsets rallocatez+CooperativeReductionWorkspaceCache.allocate sk%%))&1 >># # $ 3 3FE BI   &'9!=>##rc|jD]&\}}}|j|j||f(|j|_g|_|xjdz c_yNrW)rrrrr)r%rrrs r on_loop_endz.CooperativeReductionWorkspaceCache.on_loop_end s_*.// F &FGY   ( / /)0D E F++ 1rcH|j}|xjdz c_|Sr)r)r%priors rincrement_store_countz8CooperativeReductionWorkspaceCache.increment_store_count s#   A rN)rr)rrrrrrrrrrrrr s $rrc$eZdZUded<dZdZy)FixedTritonConfigzdict[str, int]r$c |j|Srr$r%rs rrzFixedTritonConfig.__getitem__ s{{4  rc||jvSrrrs r __contains__zFixedTritonConfig.__contains__ st{{""rN)rrrrrrrrrrr s !#rrceZdZdZddZy) 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)rUr _load_maskr)r%rjrs r augment_keyzTritonCSE.augment_key s,88&& &4 &tyy) ) rN)rjrrzstr | tuple[str, str])rrrrrrrrrr s  rrcxeZdZUdZded<ded<ded<ded<d Zd ed <d Z dd Z d ddZddZ y )TMACompatibilityCheckerzO Checks if the TMA API can be used for load / store triton operations. rrrrrrforceNr buffer_namecd|_y)Nz2Cannot use TMA descriptor for load / store since: )failed_debug_prefixr$s rrz%TMACompatibilityChecker.__post_init__ s #W rc|jrytjjjdk(r4t j jddk\rtjs+tjjjdk(r$tjjr ts!tjd|jy|j r7|j"j$r!tjd|jy|j&t(vr,tjd |j|j&yy) NTr8r 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.z/%s dtype %s has no CUtensorMapDataType mapping.)rrUrarrMrr8get_device_capabilityr$assume_aligned_inputsruse_tensor_descriptorrrdebugrrrrerrAr$s r can_use_tmaz#TMACompatibilityChecker.can_use_tma s ::GG779>>&H 88:1=B4477668==F 33)+ II[((   >>dkk22 IIA((  ::2 2 IIA((   rc  |jr}|jDcgc]+}tjjj |-}}tjjj t j|}n!|j}t j|}tjjj|dt jds"tjd|j|y|jj}|ddD]}tjjjt||zdt jdt jdrbtjd|j||ytjjjt||zdt jdt jds#tjd |j||y|j At#|j r,tjd |j|j y|j$d}tjjj|t jdr,tjd |j|j$yd} d} |j&D])} t(j*D]} t-| | s| } | } )+| r| sJ|d t(j*|j.j0r;|j2s.| t(j4vrt6| } d}|j.j8D],}|j:s|j<| k(s |j>}n|-tjd |j| |j$y|j.jA|}|jC| |i|z}tjjjE|t jds-tjd|j|j$|yy d dd}||zdz }|jGtH|jGt|}tKtMt jN|| d}||j.jQt6| kDr#tjd|j||y|j.jS| } |j.jTrW||j.jT| kDrtjd|j| |j.jT| |ytW||j.jXj[| d|j.jX| <ycc}w#t\$r.tjd|j|j$YywxYw)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. r\rWz=%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 TMA descriptor base must be 16 byte aligned but constant_offset=%s with %d-byte dtype yields a non-16-byte aligned base.z%s TMA descriptor base must be 16 byte aligned but buffer %s is recorded as unaligned (e.g. produced by a sub-pointer view / reinterpret_tensor).z>%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)r]rrrs 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)r]rrrrzsympy.Expr | Nonerr)/rrRrUrarZ!replace_backed_symbols_with_hintsrsympifyrbrrrrrrrrrrrPrrrr rpersistent_reductionrr rrxrrfnumel_get_persistent_RBLOCKsubsstatically_known_geqrTrr9rnsolve max_blockr fixed_configrtma_min_block_sizesrrw)r% block_paramsrBstrRconstant_offset_expr element_sizerinnermost_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$%77#3#3#U#U o.$ #**G#(==#A ww77 U]]STEUV IIO((  zz** crl F77##;; 5q%--:KL a  k,,    $ww77 0<?EMMRTDU V MM!   IIH(($      ',DTEUEU,V II.((    !- 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;;NANGDKK33NCq`  U,, ,,   s&0X""B/X'BX'AX''4YYc|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. )rr$s rrMz TMACompatibilityChecker.can_lift szzrr;)r)rr@rBzsympy.Expr | intrr) rrrrrrrrr rMrrrrr sm  O K"K"X. .f-.c%c*c  cJ rrceZdZUdZeZded<eZded<dZ e Z e Z ded<eZd ed <d Zd ed < dr dsfd ZdtdZdudZdvdZdwdZdxdZedydZedydZdzdZdydZdZdZdZdZ dydZ!dZ"edwdZ#d dd dd dd  d{d!Z$ d| d}d"Z%d|d#Z&d$Z'd%Z( d~d&Z)d'Z*dd(Z+dd)Z, dd*Z-d+Z.d,Z/ed-Z0dd.d/Z1d0Z2dd1Z3 dd2Z4dd3Z5 d dd4Z6dd5Z7d6Z8dd7Z9 d dd8Z:dwd9Z;dd:Z< dd;Z= dd<Z> dd=Z?dd>Z@d?ZAd@ZBdAZCdBZD ddCZE ddDZF ddEZG ddFZH ddGZIe ddHZJ ddIZK ddJZL ddKZMddLZNdMZOdvdNZPddOZQdPZRdQZSeTeUd dwdRZVeTdSZWeTdTZXddUZYeZjdydVZ\ddwdWZ]edXZ^ddYZ_edZZ`e dd[Zad\Zbdd]Zcd^Zd d dd_Zedd`ZfddaZgddbZhdxdcZi dddZjdxdeZkddfZlddgZmddhZndydiZoddjZpeqdvdkZrddlZsddmZteqddnZuddoZvddpZw ddqZxxZyS)rzdA class to represent a triton kernel and helpers to generate triton kernel programmatically rthelper_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|_t3|_t7j8t,|_t-t.t<f|_||_ t7jB|_"d|_#d|_$ttJ|_&d|_'|jPr|jS|jT|jVr|jY|j[|jVr|j]d|_/g|_0y)NrF)1 optimize_maskris_combo_kernelper_subkernel_blocksrrr newvar_prefixsuffixrprologue_cacher_prologuepost_loop_combinepost_loop_storerr outside_loop_varsrrcount block_ptr_iddictrblock_ptr_to_bufferrtrrrpointer_advancementsrr hint_overrideCounter _load_counts_pdl_load_index _pdl_has_waitr1autotune_hints triton_metarkcodegen_reduction_numelsrlcooperative_reductioninit_cooperative_reductioncodegen_range_treeinit_cooperative_reduction_maskhas_load_with_contiguous_rdimstores_with_contiguous_rdim) r%tilingrrrr#rrrrs rrzTritonKernel.__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(rc:td|jDS)Nc3JK|]}t|jduywr)rrrrs rrz2TritonKernel.triton_tensor_ndim..S sQ3td23Qs!#)sumrxr$s rrzTritonKernel.triton_tensor_ndimR sQ@P@PQQQrc^dg|jz}d||<ddj|dS)Nrrrrr)rr7)r%rsizess rindexing_size_strzTritonKernel.indexing_size_strU s942244a499U#$A&&rcdg|jz}|jD],}|j|j||j<.|S)Nr)rr3rblock_size_str)r%r7rs rr6zTritonKernel.dense_size_listZ sW//11++- ?D*)-)<)<)>doo& ? rcHddj|jdSr)r7r6r$s rdense_size_strzTritonKernel.dense_size_stra s$499T11345Q77rc|j$|jd|jdSdg|jz}d||j<dj |}|jd|j d|dS) N = tl.full(z, True, tl.int1)rrrz = tl.full([z], True, tl.int1)[r)rrr<rr7r:)r%entryr7rs rcreate_constant_maskz!TritonKernel.create_constant_maskd s    #??$%[1D1D1F0GGWX 42244"%e5! !",,.//A& L rctjstjjsy|j}|Dcgc]}t |t jr| }}t|dk(ry tjjj|||}td|DScc}w#t$rYywxYw)NFrc3&K|] }|dk( ywrr)rrs rrz4TritonKernel._has_stride1_on_rdim.. s966Q;9)r$ deterministic test_configsforce_filter_reduction_configsrr rr rrUrarZ stride_varsZeroDivisionErrorr)r! support_varsr reduce_varsrGs r_has_stride1_on_rdimz!TritonKernel._has_stride1_on_rdimq 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)ra)rrs rrz>TritonKernel.has_store_with_contiguous_rdim.. s (, d # r)rcr0r$s rhas_store_with_contiguous_rdimz+TritonKernel.has_store_with_contiguous_rdim s' 040P0P    rct|Sr)rO)r%rs r dtype_to_strzTritonKernel.dtype_to_str s 5!!rc|jxretjjtjj |j j|j jSr) rkrUchoices should_use_cooperative_reductionrarrmrreduction_numelr$s rrSz-TritonKernel.should_use_cooperative_reduction sP$$ )S)S GG / / 1 MM   MM ) )*  rcjsJjD]$}|j|xjdz c_&jd}jrt |jd}j j|_tj _ jjdtfdjDrjjdyy)z/One time setup code for cooperative reductions.NrWr]ra  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)rrrrr%s rrz:TritonKernel.init_cooperative_reduction.. s0   ''- - s(+z>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))r+rxgrid_dimrlrrr semaphoressemaphores_namer%cooperative_reduction_workspace_cacherlsplicerrV)r%r sem_counts` rr,z'TritonKernel.init_cooperative_reduction 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))rerlrV_has_constant_xmaskr\)r% rsplit_aranges rr.z,TritonKernel.init_cooperative_reduction_mask 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_loopr4s rrz2TritonKernel.codegen_range_tree.. s=D4<<=sr5Trzrbase = )rxrciteration_ranges_codegen_headerrlrkrVrfiteration_ranges_ranges_coder_get_reduction_symbols_flatten_reduction_indicesr\rcodegen_reduction_indices)r%rrn_basesrbases rr-zTritonKernel.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. Trr$s rneed_numel_argszTritonKernel.need_numel_args src|jxr4tjj|j|j Sr)rkrUrRshould_use_persistent_reductionrmr+r$s rrnz,TritonKernel.should_use_persistent_reduction s5$$ )R)R MM455*  rc|jxrGt|j|jdzk(xr |jxr|jddk(S)NrWr)rrrlrjrr$s r want_no_x_dimzTritonKernel.want_no_x_dim sY  % % 1DKK D$;$;a$?? 1!! 1!!(+q0  rcy)Nztl.device_assertrr$s rassert_functionzTritonKernel.assert_function s!r) copy_shapedense_indexing override_maskrJrKmask_constant_indexc  !"#$jj}d} t!t|t j dD](} t | tjsJ| xstj| } |r=t| tjr?jj| j } !j#| j$t| tj&tj(tj*tj,tj.tj0frtj3| x} !j5| +t6j8j:xs|xsj<duxrdk7} d}d}t}j?D]A}|jA|jBrd}nd}|j5|jEC|r&jFrt6j8jHsrjKrz|sxj<sltM!|z dk(r[jOsJ|rHjPdk(r9 dd" dfd $ d"$fd #d!#fd }|}||Sd}d}jS}fd }tUrstMjWd k(r |\}}nLtYd gtMjWz}t[d gtMjWz}d|d|d}j\sj^s|rtfdj`D!n t!j<r!j5j<tc|!|| |S| r|sԉjdrjfr!ji}j<r|j5j<tdj?D}|jk|sq|jm|}|jo}t |tpsJ|js||j#|j$|jk|sqdgtMjWz}|D]o}t |tXsJj?D]H}|ju|s|jv}t |txsJjW|||<Jqddj{t}tX|zdz}t[|}d|d|d}n/|\}}d|d|d}|!n|sr|\}}d|d|d}|!|| s|r |\} }nd}|r t|g!j<r!j5j<j!tc|!|| |S)zO Compute the index and mask to pass to tl.load() or tl.store() FrrNrTtl.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,). NrrPrRrU)rXmatch_affine_block_exprrr@rrrr)r! range_treers 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]9\}} t9t3| j5|j5| ;c} }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)rstatically_known_power_of_2)rrrrZs rrzETritonKernel.indexing..match_mod_div_block.. sH!==eYOOH$@@GGHs14rrWr{) rrsymbolsrArBWildrrr^rrrrUrarZremove_precomputed_replacementsrrXmatch_mod_div_block_exprget_slice_numelsrrfrrrrrXrrrNrr@)r!r} index_varrmodulonum_dimsr match_resultdimsrRblock_index_exprs slice_numelslinear_block_sizerzrPrrdrrZr%s @@rmatch_mod_div_blockz2TritonKernel.indexing..match_mod_div_blockr sd&'--/ !& "!))%**ykJ! v  (() HY$>?++oi&OPQ &77++ @@G @@AQAQR2KK9eX   ' ! %2CCDI 77++ NN:+<+<= ".  %2$@$@$L!) 88aI1'*,qr*:DH&E#s -x/O/OPU/V!88=  1 (!2 3y-*H*H*TU3 3'HLM18;;A>M +#) %3Ns7>I?(I2I c~fD]5}tj||j}|||}|3|cSy)ze Match a block indexing subexpression involving a single range tree. N)rXfactor_index_exprr)rr} match_funcfactored_index_exprmatchr~rs 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?| t@r8t-t.jC|jD|jFsy|Scc}}wcc}w)Nc3<K|]}|jywr)rr4s rrzBTritonKernel.indexing..match_block_expr..s*QT4;;=*QrWFtemplate_out_shape)rArBrxrryrMr^)rB)$rNrrrr3rXget_subexpr_involving_symbolrrr@rr intersectionrr5 filter_masksr$r use_block_ptrrrrr rrMr)transpose_discontiguous_tensor_descriptorrrrr issubclassrr rArB)rrindex_relative_to_xyr_indexrxrindex_subexprs range_symbolsrsubexprrAr options_classrMr^transpose_contiguousoptionsrsr!rrr%rKs rmatch_block_exprz/TritonKernel.indexing..match_block_exprs}.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/m-DE04/1J1-5TT(/(?(?U $}Q"s I# +I)cr7ttrddfSddjdDzdzfSjt j fS)Nz.shaperrc32K|]}t|ywrr)rrBs rrzATritonKernel.indexing.._get_expand_str..os*Fa3q6*Frr)rrr7r<rr6)rsr%sr_get_expand_strz.TritonKernel.indexing.._get_expand_strjshj#.(\0$66*F:*F!FFLjXX**,eD4H4H4J.KKKrrWrrz , tl.int32)c3xK|]1}|js#j|s|j3ywr)rrrrWs rrz(TritonKernel.indexing..s6',,T5L5LT5RNN$'s7:)r"c3<K|]}|jywrrr4s rrz(TritonKernel.indexing..s()-DNN$(rrrrDrrrr)r!rr}rlrBlockParameters | None)rrr}rlrr)rzBlockDescriptorOptions | None)@prepare_indexingrrr8operator attrgetterrrrrr r r!rrrrrrrrrrrrrrir$rrtrr3rvar_listrallow_block_ptrrrris_indirect_indexingrrrr6rrrrrxrrkr5copyissubset differencergrr]r2rrr7r9r)%r%r!rsrtrurJrKrv index_varsr*rrr need_dense have_densehave_loop_varsdense_mask_varsrrrrr"rr mask_shape axis_masks tmp_masksr- expand_listrrzexpand_shape_strrYrr~rrs%``` ` @@@@rrLzTritonKernel.indexing s%%e,''  %/\ *(*=*=f*EF -Cc5<<0 00#}'N'Nc(JTXX.((..sxx8  !2!23%%II))JJJJ''  "/!C!CD#!NNIMM),7 -< MM ( ( + +d*qj   +5<++- 2D&&t}}5!%"    0 1  2t33 8S8S-1==?"OOI/0A5--e4  J. ! /B ' ,x !x /Bx 'x t   / ( (` ` F'(G" '+ %%e,  L ! 'S!5!5!78A=+:+<( L !s4+?+?+A'B!BC $aS3t/C/C/E+F%FG ":,b ;GI  D$8$8)>.'^^- ??NN4??3'(151H1H1J( %--j9 * 5 5j AI#--/C%c+<===&&s+%%cmm4 %--j9 #ec$*>*>*@&AA &KD%dC000 $ 7 7 9K>>$/"&//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$// * )$     %   rc|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'r\F)rrJtma_descriptorr"rOc3K|]A}tjjj|t j dCywrN)rUrarZrbrr)rrs rrz1TritonKernel.codegen_block_ptr..Cs=!'GG,,DD & a(8sAA z#duplicate advancement for pointer 'z ' at type 'rK)rrrrkrxrcr*rMrrrrhrnextrrnamedvarruint64rir^rrVrlr!rr rrcr")r%rrrLrmcheckblock_descriptorblock_ptr_line block_varblock_descriptor_id named_var line_bodyradvance_offsets advancementss rcodegen_block_ptrzTritonKernel.codegen_block_ptrs2'') 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 &&rcd|d|jd}|j}|j}|jjs6|jj |}|jj |}t t|j|D]8\}\} } tjjj| | s4d||<:|jjs!|jj||_n||_|j||j|jdd\}} tjj|} ||j j"vrU|j j"|} t%| t&s,tjj| j(d} |dt+| d}t%|t,r d |d||dS|j/|||S) NrrrFTrrrr tl.store()rIrFrHrKrrrRrrUrarZrbrrrP get_dtyperinplace_buffersrrd other_namesrrcodegen_descriptor_store_line)r%rrLrJrrmrFrHrrz broadcast_dimrY store_dtypebufs rcodegen_block_ptr_store_linez)TritonKernel.codegen_block_ptr_store_lineTs #5'H,@,@+AC #22$66%%11&44;;OLO ( 6 6 = =>O P *3 $$o 6*  / %C%#}ww77]K).!#&  / %%11)1)?)?)D)DEV)WH &):H &99      : qgg''- 499,, ,))++D1Cc:.gg//0BC '/ <=Q? h 0ykE75'; ;11)XuMMrc`|dtjj|jdS)zFGenerate the descriptor load line. Override for backend customization.z.load(rrUrrrU)r%rrLs rcodegen_descriptor_load_linez)TritonKernel.codegen_descriptor_load_lines-"#6!((*?*?@P@P*Q)RRSTTrcf|dtjj|jd|dS)zGGenerate the descriptor store line. Override for backend customization.z.store(rrr)r%rJrLrs rrz*TritonKernel.codegen_descriptor_store_lines2GAHH$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)NFrI0) assignmentr)rrrrLrrr&r:texprrename_indexingindirect_assertget_load_bufferrr&rr) r%rrrhrrLrr:size_strlinebuffers r check_boundszTritonKernel.check_boundss $ +++==RV=W(O444&& (0(9(9(;8$$8=5--d344## esx %%h/ &$5 Lrc|js|jr |jS|jr5|jdj r|j s |jS|jS)Nr\) r(r/r'rkrxrcr*rlloads)r%rLs rrzTritonKernel.get_load_buffersb  "h&:&:&<<<   ! !  $,,'')99 :: rc|jD]4}|j|s|j|jcSyr)r3r2rr)r%rrs r_range_tree_mask_shapez#TritonKernel._range_tree_mask_shapesE++- BD~~d#t'>'>'@AA Brc\t|tr |jS|j|Sr)rrrr)r%rs r _mask_shapezTritonKernel._mask_shapes( d- .:: **400rc|yt|}|D],}|j|}|yt|t|}.|Sr)rrr?)r%rr result_shaperrs r_broadcast_shape_with_masksz(TritonKernel._broadcast_shape_with_maskssX =U|  RD))$/J!0uZ?PQL  R rztl.extra.cuda.gdc_wait()z%tl.extra.cuda.gdc_launch_dependents()ctjjjjsyt t jtjjjrytjjryt jjjdk(xr$tjj!ddk\S)NFr8rr)rrr$r enable_pdlrrUrrrr r rarrMr8rrrr_enable_pdl_codegenz TritonKernel._enable_pdl_codegens%%,,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..s9|4488HHs=A) read_writeswritesrrreadsr)r prev_depsrr prev_nodes` r matching_depz.matching_depsY( ((!--44I%OOIy7L7L7R7RS " rc3.K|] }|ywrr)rrr s rrz9TritonKernel._handle_pdl_before_access..s,S\!_,Ssr\)rrUrrra scheduler previous_noderrkrxrcrlrVGDC_WAIT)r% wait_bufferr dependencies need_waitrr r s ` @@@r_handle_pdl_before_accessz&TritonKernel._handle_pdl_before_accesss'') 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)NrWr\) r use_countrkrxrcrrVr 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)_linesrMrrrrgrr)r%code new_lineshas_waitprevious_launchls r _filter_pdlzTritonKernel._filter_pdls  AAw#~$--1"4#HAw#~$//Q"6".MM/2"%i.   Q    rcP|jjt|||yr)saved_partial_accumulaterrm)r%rrrros rrpzTritonKernel.partial_accumulate$s$ %%,, dNC 8 rc |jj}|jxxdz cc<t}|j ||}t j j}t||}|j|d|j||dt|dd}t|tr"|j|jrd|_|j#|j%} t'd|j)|j+D} |j-|rd} nd| sd} n_|j.rQ|j0s|j2d j4r,fd } d } t7j8t:d | }nd } | sr=|j=r-|j>rdtA|j>} n |rd} nd} nd } d}tBjDjFr&|jHjK}|vr|dkD} |j-| xr|j. xr| xr| }d }|rd}d}d}tMre|}|tNjPtNjRfvr=tBjDjTrtNjV}n|dtY|dz }d}nt|tZt\frt|j_||| \}} t|tZr d|| | |d}n|ja||}|jc||jd|jfdd\}}nti|rd|d|d}|jj}d}n^d|d|jld|jn| | |d }|jpr |jp}ntrju|j}|tNjPtNjRfvr/tBjDjTr|dz }tNjV}|tNjvk(r/tNjxjz|dz }tNjv}|j}|}|j||jj|||||}|j|||jdkDr xxdzcc<t|tsJ|j|_E|r|}|tNjvk(r#t|dvrdt|d k(rdnd!d"}d#|d$|d}|jj||||jp}|jr|jrd%}n|tNjvk(rd }nd&}|j>rtA|j>n|}d'|jnd$|d$|d}|jj||||j|j|j}|j.r|jss|jj||S)(zc Load from the memory location 'name', offset by some indexing expression 'index'. rWTFtma_load_for_template_epiloguerrrrIc3&K|] }|dk( ywrr)rrs rrz$TritonKernel.load..Vs AF rCz, eviction_policy='evict_last'r\c kDrsryy)N evict_last evict_firstr)expected_countr*indirect_indexing load_countsrsr decide_laterz'TritonKernel.load..decide_laterast$~5"3'$rz, eviction_policy=''zrz, other=z , other=0rz, cache_modifier='.cg'Nrrrrbrrr), rrTr")TrueFalsez tl.full([1], r2rz , tl.int1)rrz0.0rr@)Mrrer%rrrUrarrvrLtma_compatibility_checker_clsrrrrKr!r/r*r/rget_strides_of_loadris_broadcastedrkrrxrcrArBrCr& _load_otherrir$r skip_l1_cachermbuffer_read_countsrtrrrrrrOrrrrrrPrIrrrr:r"rrrr r rrrr&rrrrrrrrr4rri)r%rr!r make_lineoriginal_indexruses_uint8_storagerLr/ is_coalescedepr0rm has_read_depsr9r8cachemodappend_broadcastrrr load_bufferr bcast_operandr  other_valr-r*r.r/s ` @@@@rrWzTritonKernel.load+sriiood#'' DQ=E  55e<!!$'LUTWX== &*&H&Hd$DeL 'I'!   h 0T5N5N NN6 26D .((* **,   44^DKKM     ~ .1B1B  " "  % %)9)9")=)E)E % %).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<<=MxXD&DD((((#'# E e'~6!#d>*:"=#+#6#6 !#d8+=+=*>c(BSBSATUWTXY^X_`h_iijk (($11E)99(..IE%--88MM88))  "u}}'8'8'@& **84  &&{D9XX&& 4U'  ##K<   ! #   " *&7888'11  &M "s::K'K#ZF)BA#J*U&m_B7G6HJD**Th6K6K+J!!** Dejj(!DD7;7G7GM$"2"23T#8#4#4"5R |2i[PQR!XX..::"(((*<*< / $$X-?-?-A*  " " & &z 2rc |jj|}|}tjj |}d}||dk(r*|dk(xs t |dd} |j ||d| |}|j|d|du||dk(} t| tr6|j| jr|jj|||jjv} |j|} | r'| r%|j j#t%|d t| t&t(fr-|j+||| \} }|j-|| | ||}nV|| j.}t1|r~|j2r|d |j5d z }nZ|j6Nt9d |j6Ds2d j;t=t>|j6}|d |d z }d|d|d|d | j@d }n|dk(rd|_!| j.}t1|rZ|j6Nt9d|j6Ds2d j;t=t>|j6}|d |d z }d|d|d|d | j@d }ntEd|tGjH}|jJs7|jLr+|jO|jQ||j |jS|j |d|j j#t%|||jJs|jTjW||jYy)zn store the 'value' to the memory location 'name', offset by some indexing expression 'index'. Ntma tma_storeFTr( atomic_add)rtrJrKrvztl.debug_barrier().broadcast_to(rc38K|]}t|dk(ywrNrrHs rrz%TritonKernel.store..8s9&'CFcM9rrrr1c38K|]}t|dk(ywrKrrHs rrz%TritonKernel.store..D?aCFcM?rLztl.atomic_add(z, sem='relaxed')z store mode=r)-rrUrUrarrr4rLrrrKr!r0rrr6storesrVr^rrrrrrrr<rrcr7r9rr:atomic_add_foundr- contextlib ExitStackrkr+ enter_contextguard_cooperative_storerrriclose)r%rr!rmoderr;rrKrrL is_inplacer6rrmr indexing_str value_shape exit_stacks rstorezTritonKernel.storeseiit$!!$'$(! <45=EMFWT;%FE(,(J(J )K) %== dl&? $ 4 !  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)r'rVrS)r%condmsgs rdevice_assert_asyncz TritonKernel.device_assert_asyncXs( !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):)r[rrVr^indent)r%rrrs rrTz$TritonKernel.guard_cooperative_store[sC 88NNPd.?uK,PQR}}rctd}|D]0}|t|ds| |j}"||jz}2|S)Nr)rr)r% variablesmaskselems r_combine_maskszTritonKernel._combine_masksdsK 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] rrWr#rrryztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(rz, )r")r(rir1ONE_ELEMENT_PER_THREADrrerrrr r-rr'rr&rrrfr)r%r boundariesboundary_indicesindexing_dtyper(sortersorter_indicesboundaries_ptr boundary_sizeboundaries_underlying_numelboundary_stride sorter_ptr sorter_stride triton_dtyperrds r bucketizezTritonKernel.bucketizeps$  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|fSrwrx)r%rrryrzr7rs rreduction_resize_and_shapez'TritonKernel.reduction_resize_and_shapes'') 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. rWNr0r")rjrr6rr&rrr)r%rrr target_ndimr target_shapes rreduction_collapse_dimsz$TritonKernel.reduction_collapse_dimss  " "a 'L--/$2I2II ,,. $\k2hZ? xx   3u:}l C % !  rc KLMNOPQRSTUdKdRdLRfd }tjRfd|}}}|rJtj||}RrtjnRrtjnj sJt djD} j| t| } jr| jjjdjd} jrYj} t| dk\sJ| D cgc] } d| vsd | vs | } } d d j!| d Nt#| Tn)j%Nt#jTj'NTfd |}d}dKdL|KzLz}ddddddd}d}|vrt)|t"r|\}}j+j,z O dMOfd P dNPfd }dOOQfd }KLfd}|f}|j.j0vrj.j0|St3}t5}|vrj6j9nd}t;j}t=OOj,zD]}d||< t d| D}Lvr_|Jj.j?|t#|j.j?|t#|f}|D] }||_ nc|vr2|Jj.j?|t#|}||_ n-j.j?|t#|}||_ dj!| MMfdUjBrtDjFjI}dSdPSUfd }d!k(rt)|t"sd} nwt)|t"r&tK||D!"cgc]\}!}"||!|"} }!}"nAd"k(r3j.jMjN||jP#} n |||} |vrt)| tRsJ|J|r"d$tU|d%jW|d&}#nCtUj.jMjNd'| d(| d)|| jX}#|Q|jN|| |#|nd*k(r8jZrj]||U|}ndj_|}nPd+k(rMt)| t`sJ| \}$}%}&t#fd,jcjN|$|%|&OD}nd!k(rt)|t"rpt)| t`sJj.j?tetR|jX}'jgjN||'| d| d-O}nyji|}net)| tRsJPjN| | jP\}(})}*j.jMjN|(|)|*}nLvrtetj|dn|}+j.jmd.|+|t#j},tDjFjo}j'tp|}t)|t"sd"k(rj} t| dk\sJ| D cgc] } d| vsd | vs | } } t#| |,_,d d j!| d Njrju|,d/Nd |d |d&n5jrju|,d/j%d |d |d&|vrd.|+d0}#|Jjrju|#d/j%d tjv|jxd jW|d&|Q|!d$tU|d%jW|d&n| d}-jNj{d1|,d2|#d3Qd4|,d |#d |d |-d5|,d6U|,d7|,d8|#d6U|#d7|#d8|j|||,|#|ntrj]||U|}nd!k(rd.|d9}.d.|d:}/jrju|.d/j%d;|d&jrju|/d<j%d |d&t)|t"rE|\}0}1jNj{d=|.d2|/d>|.d |/d |0d |1d tjd?n<jNj{d=|.d2|/d@|.d |/d |d tjd? jNj{dA|.d6U|.d7|.dA|/d6U|/d7|/dA |}2j.j?tetj|2jX}'jgj||2|'|.|/O}ntEj}3|3|,|}4d"k(r!jNju|,d6|4n'jNju|,d6U|4|,tjk(rDj.jMj||,dBtj|,jX},|j|tetR||,djZrtDjFjo}tj}5j|jfD]2}6|6judC|5j|6j4|vrmLvr9tet"tRtRf|\}7}8|7}+tU|7}9tU|8}:nlKvr0tetR|}+tU|}9j|dD}:n8tetR|}+j|dE}9j|dD}:|+dF};j|ju|;d6|9j|;|}<j6j9}|+dG}=j|ju|=d6|:j|=|tjv|jx}>|j||<|>|ntrd*k(sJtet"tRtRtRf|\}?}@}Aj|?t5|d}Bj|@t5|d-}Cj|At5|dH}Djj|?|@|A|B|C|DO nd!k(rtet"tRtRf|\}2}'t)|t`sJj|2t5|d}Ej|'t5|d-}Fjgj|2|'|E|FOnMjtetR|t5|}G|jtetR||Gd|5jtetj|j.j0|<t)|t"r|n|f}Hjj|HtdI|HDsJ|retHD]W\}}ILvr|dkDrKLzvr|n|}JIjP|Jk7s+j|juId6|IdJtJd&Y|Scc} wcc}"}!wcc} w)QzS codegen reduction of value to Triton according the reduction_type cH|duxr|jxr|jdkS)Nr)rr)rs r should_upcastz-TritonKernel.reduction..should_upcasts$D=KQ%8%8KQZZ!^ Krct|jr$tj|tjS|Sr)rr rrr)rrs r maybe_upcastz,TritonKernel.reduction..maybe_upcasts4!- UEMM2  rc(|jSrr~)rrs rz(TritonKernel.reduction..smAGG.Drc3<K|]}|jywrrr4s rrz)TritonKernel.reduction..I4>>+Irr\rrXYrrrcxjjjd|dd|jS)Nrrrr")rr&r'r)rr<r%rYs rrz(TritonKernel.reduction..s?dhh'' "1#R'7q9gg! (r)argmaxargmin) argmax_value argmin_value)argmax_with_valueargmin_with_valuerr)rrrrrrNct } 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. r!rz[None,:,:,None]rWz [:,:,None]r9rrr)rrrr6rr}rPr) rr result_typetriton_reduction_fnis_bmmrrrzrrr%s rfinal_reductionz/TritonKernel.reduction..final_reduction9s#@"O 00FE&T11349{{... %wo6F00a0E %wj1F-ekk-1-E $ ? ?*+1UG2cU!.final_reduction_defineZs0*&%EKE1a MMZLE73 4rr!c j|||j} j||t|tr$t t j|jnt jj}|dk(rV|\}}|jd|d|dd|d|dd|d j|d|d j|dy|d k(r=|jd|d |d d|d|dd|d j|d dy|jd|d |d d|d|dd|d j|d dy)Nvalue_and_indexz rz = triton_helpers.z _with_index(z) rO rz_val, z_idx = triton_helpers._val_idx) rrrr]r rrUrrMr\r{) rrrr! result_kind result_value result_indexrzroot_opr%s rfinal_argreducez/TritonKernel.reduction..final_argreducegs00 LE00e[1U[[%++.XX<<> E//-7* l !N"\N2DWI\Z_Y``bchbiiklokpq!N#d&;&;|n&N%OP!N#d&;&;|n&N%OP' Lzl2H Q]^c]ddfglfmmopsotuLD$9$9ZL:M$N#OP Lzl2H Q]^c]ddfglfmmopsotuLD$9$9ZL:M$N#OPrcvryvryy)Nrrr!r)arg_value_reduction_typesarg_with_value_reduction_typesrsrargreduce_result_kindz5TritonKernel.reduction..argreduce_result_kinds!??(!::rrc3>K|]}t|dr|ywr)rKrs rrz)TritonKernel.reduction..s!& (;CF(CC& sr"r6c:s|Stj||Sr)r%r)tvalfvalr]s rr+z*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)rrriinforr)rrrs rupdate_constant_dtypez5TritonKernel.reduction..update_constant_dtypes ..)uzz2I#O )X]#Ou{{95999 ;;y1555Y!7!;!;; ;;y1555#Orc||j}jt|}jj j |||j|j Sr)r_map_tuple_or_scalarrirr&r'r)rr( default_strr%rrr+s r _mask_valuez+TritonKernel.reduction.._mask_valuesb/EKKP"77 wO xx((LLuk2++++ )ronline_softmax_reducer!r~r9rrrzindex, z.shape)welford_reducewelford_combinec3vK|]0\}}jjj||2yw)r"N)rr&r')rrrrr%s rrz)TritonKernel.reduction..s9#$uHH%%dllEe%T#rrWrYr>_index _next, z_next = triton_helpers.z%imum_with_index( # ) rO_next _max_sumz, float('-inf'), = tl.zeros(z zT_next = triton_helpers.online_softmax_combine_with_sum( z3 ) zK_next = triton_helpers.online_softmax_combine( rz .to(tl.int8)zif HAS_RSPLIT:rr_bval_bidxr#c3<K|]}t|tywr)rrrHs rrz)TritonKernel.reduction..4sJ:a!23Jrr)rr/rr)rr]rr])rr]rr/rz.tuple[str, torch.dtype | None, BlockShapeType])rr]rr]rr/rr)r!rr])Spytreetree_anytree_maprrrkrrxrr8rrrfr5r6rr7rr<rrrrjrreduction_cacherrrmselect_index_dtypertrvrUrrr%r#r$rr&r'rr]rrPrr+rwelford_reduce_fallbackr_welfordr %online_softmax_reduce_final_reduction prepare_softmax_twopass_fallbackr rdefault_accumulatorrirlrVrrr\rrHr$rget_reduction_combine_fnrrrQrRrrSrar{*codegen_cooperative_reduction_peer_combinewelford_reduce_final_reductionrUrrrcrRr)Vr%rrrrr do_upcastoriginal_dtypeoriginal_src_dtyperdreduction_range_prefix dense_sizesr xy_sizes_onlyarg_index_reduction_typesarg_reduction_types arg_root_ops logical_indexrrrrjacc_typetorch_acc_typerrrresult_mask_varsrrr(r masked_valuerraccumulator_indexmeanm2weight result_sum_result_dtype_shape result_prefix accumulatorraccumulator_maxaccumulator_sum value_max value_sum result_max combine_fnupdatedrZrrr block_value block_indexblock_value_namepeer_valpeer_idx_inputpeer_idx result_mean result_m2 result_weight peer_meanpeer_m2 peer_weightpeer_maxpeer_sumpeers result_tupler target_dtyperrr]r<rzrrrrrYr+sV```` @@@@@@@@@@@r reductionzTritonKernel.reductions L OO$DeL & OOL%8E)6y)A yI%25%9EMMuE$$$$I8H8HII % u  ?? LL )!%!1!1"!5!A%>QRda Aq 1R R5( $xx00uEKK0X *5': !44!, <<<"... C ./uT5F5F{5S4TTUV&),)) LL./E.Fgl^[bc"-"."4"4 *)%'~6LL %)+  #33--!%!4!4"NE:xQV"J"&!=!=eU!KJ#44!,999%1"r6"#(,  dBU)#  #::eU+%lH===!%#4 Z+H+N+N"1"J"&!K!K ""$Q$Q"J"&!F!Fue!TJ!, <<<*9LL, 0B0B+'"XX..LL'v/ "%CCS*%a(  ((++M?#$D0023,K ll66~yQG// wGGgu-!U*"&"6"6"8K{+q000)4%!%t sd{%M%).m(2B"WIRPXzYZ[II''&-{43F3F3H2IG9TVW_V``ab!44&' f$=!"... ##()T5H5H5J4K2{{;/334Bt7H7H7U6VVWY'~6%0M*+51B1B;1O0PPQR2359  ## W%6$77NwiX M$5#6br)M S{m5,A;!O PQ"#3z5F4Gu2MO`'a&bc  **%)+ &n5!00z8U  #::$%j\"6$%j\"6 ##&'{43F3F3H2IIZ[cZddef ##&'|D4G4G4I3J"XJVWXeU++0(IyLL''()0AB,-R/@9+RPY{Z\]c]q]q\rsLL''()0AB,-R/@5'FL`L`Kab ##$%S6Gu4M)_(`a$%S6Gu4M)_(`a( !XX__tC'<'B'B- "GG**##  88S $[%8!U*LL**k]#gY+GHLL**&-s:g{+K*LM *#'(("3"3..&-|4#jj)// #4#K'**j1   % %ll66~yQG#--/J..0D0DE 7 ./((6 7 !44!%CC15k;672.L,%1M"%l"3K"%l"3K#'@@$(j$AM"%j/K"&"7"7:,d8K"LK$(j$AM"&"7"7:,d8K"LK"&"7"7:,d8K"LK&3_E#: &&004D3ES 1VWJJ$i#mm>>@ $1?%!8&&00N3C3{m1TUJJ"K1M1M (()+ &n5%)99998<+{K?@*95 Y !KK$Y/AJ II$Y/AJ #MM!$Y/AJ 33((!  #::)-+{23Z*& J"'8444JJ 0 ;WQZ JJ 0 ;WQZ::((GGj13CI3NPW'((${J*GPT    .23 .C  +%/ E%Bz  %%l3J\JJJJ &|4  6!%CCA&14RRT'(  <<</**44!(#fXT2El2S1TTUV }Yp SD%s#< A@; A@;2AA4 AAAAc|j|||}|j|||}tdDcgc]'}t|jj |)c}\}}|j d|d|d|d|d|dt jd|d|j|d|d|j|d||fScc}w)Nr#r~ r9 = triton_helpers.online_softmax_reduce( ) rO) rrvrrrUr\r$rr{) r%rrrrzrrYrrs r_online_softmax_reducez#TritonKernel._online_softmax_reduceIs66vPUV66vPUVMRSTX!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)rrrrr%s rrz(TritonKernel._welford..^s(   ( ( > s ztriton_helpers.welford(rrc2t|d|dzdzSr)rrs r reduced_shapez,TritonKernel._welford..reduced_shapeds$qcAgi(889 9rr"rOc3VK|] }j||j"ywr)r}r)rrr%s rrz(TritonKernel._welford..ms*   + +E5;; ? s&))rrUrrVr7rr) r%rrrrrzrwelfordr rwelford_resultsrCs `` `` rrzTritonKernel._welfordZs 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_meanrrr_m2_weightrrrrrrz<_next = triton_helpers.welford_combine( z, rrz;_next = triton_helpers.welford_reduce( z1, roffset == 0 ) z rOrrN)rrjrrr6r"unknownrlrVr<r'r\rr)r%rrrr+rrrzraccumulator_m2accumulator_weightrrrrs rrzTritonKernel.welford_reducers%%'$*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.welfordr"rO)rtrrRrrrUr\r)r%rrrrrrrrzrr result_exprsr result_exprrrs rrz+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)NrrrrrO)rr\r$rr{) r%rrrrrrzrrrs rrz2TritonKernel.online_softmax_reduce_final_reductions66vxO66vxO  L:,' !O#4Bse2f>R>R=ST LD11ZLBCD LD11ZLBCD   :%%rcD|jr|jdStS)NRSPLIT)rr7r$s 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. r]zxindex < xnumelNrz_ws = (rcz).to(tl.pointer_type(z)) tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), rrTrQ_peersrrrz = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, r)rlr_rrr[rrr\rrOcreate_cse_varr"rrrVri) r%rr default_valxnumelrrrrrs rrz7TritonKernel.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)rrrrrIrc38K|]}t|dk(ywrKrrHs rrz/TritonKernel.store_reduction..9rNrLrrIrrrr1)!rkrUrarrLr4rrUrQrRr+rSrTrrrrrrVr^rrrrrrrrcr7r9rr:rU) r%rr!rrrLrrZrXrYs 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  rct|jj|j|j||||S)z=Emit a tl.reshape into a new tile shape, returning a CSE var.r")rr&r' _reshape_expr)r%rrrs r emit_reshapezTritonKernel.emit_reshapeGs=xx  LL   ue , !  rc |t|}|jj|j|d|d|d||S)zDEmit a Triton reduction primitive along `axis`, returning a CSE var.r9rrr")rrr&r')r%rraxisrr reduce_fns r emit_reducezTritonKernel.emit_reduceUsM2.A xx  LLk5'D6 + !  rc|j||}d|dt|d}t|t|t|} |jj |j | ||S)zreshape(broadcast_to(reshape(value, pre_broadcast), broadcast), final). Used for nested-reduction broadcasts that lift a reduced-resolution value (one element per group) to full or half resolution. rrrr")r(rrrtrr&r') r%rrrFrIr out_shapereshaped broadcastedrs remit_broadcast_via_reshapez'TritonKernel.emit_broadcast_via_reshapefs{%%e-@Axj+;O+L*MQ O k4+@${BSTxx  LL  !  rct|dd}|d|dt|dStt|t |t |S)Nrrrr)rrrrrt)rrrs rr(zTritonKernel._reshape_exprsO E7D1   r*:5*A)B!D Dc%j$y/4;GGrc 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)rrYr"N)rr)rrrrrrs rrz,TritonKernel._lift_helper..s@%A~u s1#Qqc]%u{{ Ks47rc32K|]}t|ywrrrHs rrz,TritonKernel._lift_helper..sRc!fRrz def {name}():ryrry)ShapePropagationOpsHandlerc4eZdZ dfd Zy)+TritonKernel._lift_helper..CSEProxyc d|z t||i|}t ||i|}jt ||i|||S)NrYr")rr&) r%rrrrO output_shaper dtype_handlerhelper helper_name overrides shape_handlers r_defaultz4TritonKernel._lift_helper..CSEProxy._defaults4&z)  w!  #" #  w!  #" # ||,GIt,d=f=&& $rN)rrrztuple[Any, ...]rrqrr )rrrrB)rr=r>r?r@rAsrCSEProxyr:s-  '6 @N   rrCc32K|]}t|ywrr)rrUs rrz,TritonKernel._lift_helper..sBF Brreturn rz)r_rVr\rvrrRrr7rr from_iterablerrrz!torch._inductor.shape_propagationr8r.rarUset_ops_handlerrrigetvalue)r%rFrdtypesrr signaturerzr8rCoutputsrr=r>r?r@rAs ` @@@@@@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|]}|jywrrr4s rrz$TritonKernel.scan..rrz(ops.scan not supported inside ops.maskedc32K|]}t|ywrrQrrs rrz$TritonKernel.scan..Fe*51Frrrr"rrrr\rrz float('nan')rr>c2djd|DS)NrLc3&K|] }|d ywrDNrrrs rrz1TritonKernel.scan..csv...csv88.cse_multiple.. LI488$$Y/L%(r"rO rrvrcrrrrUrr'rVrrir)rrrdrJrr cache_keysrjrr result_varsrr\r%s 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((r1cT|jyt|j}d|d<|S)Nrr\)rrt)rrs r_partial_scan_shapez.TritonKernel.scan.._partial_scan_shape"s*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, )"rkr+rrxrr8rrrArBrr&r'rMrrjrrrr<r6rrrrUr7rrlrVrQrrrr) r%rJrrrdbroadcasted_values accumulators cse_computecombine_helper_fnrzrr value_dtyper reduced_sizerreduced_size_strr(rhpartial_scan_varsrkpartial_scan_varpartial_reduce_vars accs_nextfull_scan_vars full_scan partial_scanrgacc_nextpartial_reducerr\s ` @rscanzTritonKernel.scans$$$$--5v5-I8H8HII % u ??N$NN" FvFF''(9(94<<H  --j&&I%%'$*A*AA/ 1LE5((++ '1%89;kk ,K HH%% ";-r$2E2E2G1HJD0023 &E  % %e ,&u-H,,#335 #& R "hhooEoN %&tyy'>&?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#.Trrz(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc32K|]}t|ywrrRrSs rrz$TritonKernel.sort.._rTrrrrr"c2djd|DS)NrLc3&K|] }|d ywrWrrXs rrz1TritonKernel.sort..csv..krYrCrZr[s rr\zTritonKernel.sort..csvjr]rcNt|}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|!ywrr`rbs rrz:TritonKernel.sort..cse_multiple..prcrdr"rOre)rrlrdrJrrrfrjrrrgrr\r%s rrhz'TritonKernel.sort..cse_multiplems&'A;@8DaTF"QCr%1DJDLLLAKLI Y/LL%(0B$C E5e5;;?K LL " "{#$Cv. *-[*)E 4% I+0J( Y 3 4% %ELrir\rrnumelr#ztriton_helpers.sort_with_index(rrWz , stable=z , descending=zUnhandled sort)rkr+rrxrr8rrrArBrr&r'rrjrrrRr<r6rrrsrrr)r%rJrstable descendingrdrnrzrrrlrhrrrgr input_varr\s` @rrzTritonKernel.sortKs@$$$$--5v5-I8H8HII % u ??N$NN"(( A ( ''(9(94<<H %%'$*A*AAFvFF6{c&k)))&f-  5 "5'D,?,?,A+B!DQiD0023    = &$#00002243C3CB3GHh v;! 12DQ2G1HK]^_K`Jab82cU)F8= AO 't-?OK !12 2%(f%= 1 !J #(J ) 0 0J  1[!!] s AHc|jsy|j|j|jj|jjy)z Generate the output from prologue. This should be extracted from the subgraph, which is why this is partitioned from codegen_body. N)rr\clearr)r%rs rcodegen_prologuezTritonKernel.codegen_prologues@ }}  DMM"  !!#rcf|js=|js1|js%|js|js |j sy|j Dcgc]}|js|}}|jr|jsJdi}t|jD]\}}|j}tjj|t j"}|j%t&|}d|}|j(j+|d|d|j,j/|t j"d||<|j(j+d|j(j+d |j(j1d 5|j3sZ|j d } | j4d k(sJ| j4} |j(j+| d| j6d| d|j(j9|j|j(j;dg|j(j9|j|j(j9|j|j(j9|j|j(j9|j t|jD]\}}|j<} d|}tj>|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_BLOCKr"z/split_size = min(RSPLIT_SIZE, xnumel - xoffset)z@for _ in tl.range(0, split_size, XBLOCK, num_stages=NUM_STAGES):rW)rrr]zmask =  < rzxindex += XBLOCKr9z, 0))rr'rOz&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(rrr7z = tl.advance(rz + tl.program_id(1)zR if HAS_RSPLIT: triton_helpers.x_grid_barrier(z) Tr)F indexing_coderrOr'rrrxrcmix_order_reductionrrRr%rr%r#rrr<rrirlrVrrrar_rfrr\ writelinesrrrr&runittestmockpatchobjectrvrrkr+r r rrrdrhrur"rrrrrrrr^r!rUrr invalidater cache_clearrZr[rr) r%r loop_trees accumname2varr partial_accumrr(rr?r]rrtriton_reduction_functionnewvalrrlevelrf loop_startloop_end num_stagesrJ advancement prev_treeprev_advancementsprev_advancement prev_block prev_num_itercurprevsem_ptrs r codegen_bodyzTritonKernel.codegen_bodys   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)Nrr#z!Unsupported numel argument type: )rladd_numel_to_call_argsrrrrrwrUrarZoptimization_hint_with_override inner_exprr#rrrwrM)r%r numel_argsrhints rkernel_benchmark_extra_argsz(TritonKernel.kernel_benchmark_extra_argsQs     !+-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=rrfrrOz = torch.zeros(z*Don't find the buffer or const tensor for rErDN) rzdef call(args):zwith rz = get_raw_stream(z.run(*args, stream=)rrz def benchmark_all_configs(args):z.benchmark_all_configs(*args))rrzif __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)textwrapdedentrrUrarimport_get_raw_stream_asr$s rimports_for_benchmark_kernelz)TritonKernel.imports_for_benchmark_kernels:  F177%%>>?OP Q   rc|jry|jry|jr|jsJy|jryy)Nrr+rr pointwise)rr+rrkr$s 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') ) r_r\triton.language.extra.tlx ImportErrorrrVr$rproton_profilingrI)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_castsrrrrs rtriton_meta_commonzTritonKernel.triton_meta_common s5%+$B$B B113!00<<  rcXidtjjjdtj dtj dtjjdtjdtjdtjdtjd tjd tjd tjjd tjj d tjj"dtj$xstj&dtj&dtj(j*dtjj,dtjj.i}tj0rtj2|d<tj4j6d|d<t j8rd|d<tj:rLtj:|d<tj<|d<tj>|d<tj@|d<tjBr9tjB|d<tjD|d<tjF|d<|S)N backend_hashassert_indirect_indexingautotune_local_cacheautotune_pointwiseautotune_remote_cacheforce_disable_cachesdynamic_scale_rblockincremental_autotune max_autotunemax_autotune_pointwisemin_split_scan_rblockspill_threshold store_cubinrDbatch_invariantrF&mix_order_reduction_allow_multi_stagesdynamic_disable_pipelining$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$rrrrrrrrrrrrrrDrrErFrr*write_are_deterministic_algorithms_enabledrr r rrrrrrrr)r inductor_metas rinductor_meta_commonz!TritonKernel.inductor_meta_common(sv EKK//HHJ &(G(G  #F$?$?  !&--"B"B  $V%A%A  #F$?$?   #F$?$?  #F$?$?  F//  %f&C&C  $V]]%H%H  v}}<<  6==44  V11KV5K5K  v55 -f.A.A.`.`! " 5fmm6j6j# $ )&--*R*R% *  < <::< @ A ==   (&*M( #    )-M+ &  # #171I1IM- .7=7U7UM3 48>8W8WM4 5FF K L  + +00 5 677 < =>> C Drc|j|j|j|j|jt |j d}|jr|j|d<tjstjjr|jxs |j|d<|jr|j|d<|j r|j |d<|j"|j"|d<|j$|j$|d<|j&r|j(|d<|j+x}||d <|j,rd |d <tj.s tj0stj2r|j5d z }|||d <tj.stj2r|j7}|||d<|S)z Used by both standalone codegen_kernel() and ComboKernel.combo_grid_meta() (which calls this on each sub_kernel). )rerPnum_load num_store num_reductionr( RSPLIT_SIZE"has_loadstore_with_contiguous_rdimr tiling_scores min_xblock min_rblockrr4Tadd_persistent_rblockgeA kernel_num_gb kernel_flop)rerPrr r setr(r rsplit_sizer$rDrErFr/rNrr rrr+r$_get_native_matmul_persistent_rblockrbenchmark_kernelrbenchmark_combo_kernelestimate_kernel_num_bytesestimate_flops)r%rrrflopss rinductor_meta_per_kernelz%TritonKernel.inductor_meta_per_kernel[s  $ 5 5 !//!$"5"56   # #!%!1!1C    6#6#6#U#U22766 4 5  # #)-)A)AC% &   #'#5#5C ?? & $C  ?? & $C   % %*.*C*CC& '??A AF N5;C1 2  % %+/C' (  # #'',,335;F!'-O$  " "f&C&C'')E %*M" rcH|jjxr |j }|j}t |j dk(}|r|r|jj |j }|jjjd}|j}|'d|vr#d|vr|dt|ddz }|tk\}n,|jj|tjk(}|j jj"} |jjj"} t$j&j(j+| tt$j&j(j+| dz } | dk\r|rt$j&j(j-|jj.drDt$j&j(j1|jj.dr|d kry y ) Nr#rr]r0_rWg?ii TF)rmrrr rr1 memory_stats persistentmemoryrzcount_per_threadrINNER_REDUCTION_RATIO_THRESHOLDget_reduction_hintr5INNERloopedbytesrUrarZoptimization_hintstatically_known_leqrTstatically_known_gt) r% looped_redr  two_d_redr dim_statsmem_ops_per_threadr_coalesce_ratiocontiguous_red looped_mempersistent_memsaved_bytes_ratios rrz"TritonKernel.add_persistent_rblocks]]//1S$:S:S6S **  $) )==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'",rc !"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}|jjttj| |d|j|j}t"D]}d|d"|j8<||_N||_O|j|j|j|j|jtzjj8|jstj~jst"dg|d<nt"g|d<|jD]$}|jd|j|&|jr2d|jd|jjd|d|d }n|jrF|jj|j}d|jd|d|d|d|d }nMt|"}|d|j8dnd}d|jd|d |d!|d|d"|jd }|j||xsttj}|jd#|d$d jd%!Dd&|j5tj~jr|jd'|d(|j||j$jD]\}} |j|d)| |j|jtj~jr|jd*|d(dddtj r/|j|j|jd+|jS#1swYXxYw),z Convert the TritonKernel from Inductor SIMD IR to triton code, including inductor triton heuristics, imports, metadata, and benchmarking infra. Ncpuz"triton_helpers.set_driver_to_cpu()z"triton_helpers.set_driver_to_gpu()rctrjt|jt|dy)NT) is_constexpr)rPrr[rY)rargdefsrKs radd_constexpr_argz6TritonKernel.codegen_kernel..add_constexpr_args/-/  h!78 NN78$? @rrrr  NUM_STAGES) size_dtyper8ztl.dot)rKrr native_matmulTlaunch_cooperative_grid) grid_type kernel_namemutated_arg_names optimize_memrWrr)pointer_range_overrideconfigsrz$ @triton_heuristics.z( config=zI, filename=__file__, triton_meta=z$, inductor_meta=z; ) @triton.jit z!( size_hints=z%, reduction_hint=ztile_hint=TileHint.rDrzH filename=__file__, triton_meta=z*, min_elem_per_thread=zdef r9c3<K|]}|jywr) full_namerHs rrz.TritonKernel.codegen_kernel..s*JQ1;;=*Jrr7zpl.enter_scope("r4rOzpl.exit_scope("r)jr_rlrrKrkrUrarZr(r9rr\rrrMr$rrrrrRrrer rrrinv_precomputed_replacementsrr mutations input_buffersrirremoved_buffersr` inner_nameoutput_buffersrdrrg zero_moderh ZERO_ON_CALLr8r3rfrrrYrxrrrrr+rrurr2rrrrr<rrlr'r is_inference is_backward_get_grid_typerrIDESCRIPTIVE_NAMErrrqr)rrrr#r r rPemit_pointer_range_32rprrVrrrmr$r rsrrr7rarcodegen_static_numelsaliasesrrrI)#r%rr size_hintsrfr numel_hint size_hint device_typerYrrr mutated_argsmutation mutation_argargnamersizeargr9triton_meta_signaturer)rArarg_numr>heuristics_linereduction_hintr tile_hintr?oldnewr8rKs# @@rcodegen_kernelzTritonKernel.codegen_kernelsP  ![[..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<P>PQN#$$($7$7$9#:; *~.$$2#34!!,0##0"34 O$J :D>B>N-dii[:TVI#$$($7$7$9#:; *~R {;!!,0##0"34))-)A)A(BC O O$:c+"9"9:  ;-q*J'*J!J K2 N [[] B}}--!1+bAB  & &t , II--/ 1S#c#/0 1 KK "}}-- R@A B  " " KK--!%%o6  }}# B Bs Cee#ctjjj|}t |t j tfrt|}t|}|Sd}tjjj||sC|dkDrtd||dz}tjjj||sC|S)NrWi@z!Failed to find static RBLOCK for r#) rUrarZsimplifyrrrrr9r)rw)rrs rrz#TritonKernel._get_persistent_RBLOCKs!!**62 fu}}c2 3f+C!#&C Cgg&&;;FCH?$'H%QRRqgg&&;;FCH Jrc|jr|jr |jry|jDcgc]2}|jr$t |j |j4}}|rtj|SdScc}wr) r5rr+rxrr4rrrr)r%rrblockss rrz1TritonKernel._get_native_matmul_persistent_rblocks%%,,))((    ,D,G,G ,S T  &-tyy!6$6  s7BcN tj|y#t$rYywxYw)NTF)rrrw)rs rhas_persistent_RBLOCKz"TritonKernel.has_persistent_RBLOCKs*   / / 7  s  $$c|jdr d|d<d|d<tj|js|j drJd|d<yy)NrToverride_persistent_reductionFoverride_cooperative_reduction) contains_oprrkrTr)kernel_features kernel_kwargss r apply_feature_required_overridesz-TritonKernel.apply_feature_required_overridessb  & &v .=AM9 :>CM: ;11/2Q2QR$(()HI II=BM9 :Srcd 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_integersdU]]C$89 9rznumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)rzBLOCK: tl.constexpr = r]zXBLOCK: tl.constexpr = 1N)rrrr)rxrrkrUrarZrgrrVrfrrr+rrrr5rrre)r%rrursimplified_tree_numelrrs rrSz"TritonKernel.codegen_static_numelss,$ :$$ ;D$$(=(=()(8(8(A(A$**(M%$%:;NNdkk](3?T;U:V#WX  T%>%>-- JJt';';DJJ'GHEFugMfgC55djjAC,,!#rl$++"3"3"5!66LSERS{{c!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)NrWr#rz"Unsupported number of dimensions: )r5rxrrrr/MixOrderReductionGridr+CooperativeReductionGridGrid1Drr9rGrid2DWithYZOverflowGrid2Dr5BatchMatmulGrid3DGrid3Drw)r%rrs rrPzTritonKernel._get_grid_type s 8H8HI***+I J  # #6M6$:: :  ' '6M6$== = !V$++ + !V3t22D4D4DEF(===$++ + !V$$(:::$++ +=aSABB#JsD c|jD]}t|jtjtj fr |j}n*t jjj||}|jr |js|j||jt|yr)rxrrrrrrUra wrapper_codegenerate_numel_exprrrkrrM)r%rr arg_typesrrs rrz#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)rrr)r) rUrarwrite_triton_header_oncerrrworkspace_argsgenerate_workspace_allocationgenerate_kernel_callr)rdeallocate_workspaces) r%rr deallocate_wswrapperrYrrwss r call_kernelzTritonKernel.call_kernel)s''&&((*%)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()) rUrarrrrrrf cpp_wrapperrV)r%rrYrarg_signaturesr arg_signaturers rcodegen_nan_checkzTritonKernel.codegen_nan_check@s''&&*.))*B*B*D'9na"%i"@ , C-377&&%%TUXTYY\]`\aade)-BCD%%d+(-BCD%%d+ ,rct|i|Sr)r)r%rrs rr!zTritonKernel.create_cse_varOs $1&11rcF|jd|j|j|j}|jj s|j r+|jdk(r|jj|y|jj|y)NrOr]) rrrrrrcrrfrrVrl)r%r?rs rcodegen_iteration_ranges_entryz+TritonKernel.codegen_iteration_ranges_entryRsy**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|j|jd||S) Nrrrryrrz + rsplit_startz tl.arange(0, ) rr8r5rrr+rrrr)r%r?rrrs rrez)TritonKernel.iteration_ranges_ranges_code]s+++%%e&6&67  ! !  A%'')Q.D&& *5*C4 }A&  & &))""x/Ftzz%*:*:*<=>avfXNNrc^|j}|j}dg|z}d|d|d|dS)NrWrrr)rr)r%r?rrrrs riteration_ranges_scalar_codez)TritonKernel.iteration_ranges_scalar_codersC&& &&(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(rr9z + tl.program_id(rWz) * tl.num_programs(rrr#r)rrWr#ryr)rXrr5r pid_cacherr)r%r?rreversed_pid_mappids riteration_ranges_get_pidz%TritonKernel.iteration_ranges_get_pidzs~~)))u~~.a0  & &u -cU+ENNQ,>+??STYTbTbSccefC  T%<%<%>!%C#$q1 "#3ENN#C"DAFCoo!!#s+   z )U$t//02 2 rc |jr |jry|jdk(xr[|j xrL|j xr=t j jj|jt S)NFrW) rrrXhas_zdimr+rUrarZr)rr8)r%r?s rrz#TritonKernel.needs_yz_grid_overflowsw   D$=$= NNa  YNN" Y... YGG$$99%++~GWXX  rc|jr |j|jdSt|jS)Nr)rrr6)r%rfs rrzTritonKernel.max_blocks;   $$ '7u%=> > //rc|jsy|jr5tjjj |j dry|jsy|jrW|jjd|jvr.|j|jjddk(rZy|jsMtjjj|j dr|jr |jsy|jrf|jrZ|j|j }tjj!}t#|}t%|t&r?||kr:y|jdk(r|j(rd}n|j+|j}|jr|j,r||j/z}tjjj1|j |r[|j2dk7xsJ|j4xs<tjjj7|j t9Sy)NFrrrWTr])supports_constant_maskr5rUrarZstatically_known_ltrrrrfrrrbrrrrr3rrrerr+rrrXrr)r8)r%rrr warp_sizes 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%f-I)S)i).C [[C DMMIt{{3I   !;!;!DOO$55I 77   8 8Y O "W==W77##88^EUV  rcd|jd}|jdk(sJ|j|S)Nrr])rxrfr)r%xtrees rr_z TritonKernel._has_constant_xmasks5  #||s"""&&u--rc|jD]3}|j|s|j|j5|jdy)Nr)rxrr]r)r%rrs rrzTritonKernel.filter_maskssJ$$ 4D&&t,!!$.."23 4 &!rcttjd|jDcgc] }t| c}Scc}wr)rtrr rjr)r%rs rget_reduction_prefixesz#TritonKernel.get_reduction_prefixessB]::;U>UV  t    s<c|jDcgc]}|js|}}djtd|D}|j d|j ||jDcgc]#}|jrt j|%}}t|}|j d|j |ycc}wcc}w)z^ Generates code that flattens ND reduction numels, block sizes, etc. into 1D. rvc3:K|]}|jdyw)rN)rfr4s rrz8TritonKernel.codegen_reduction_numels..s"UTdkk]%#8"Usz rnumel = zRBLOCK: tl.constexpr = N) rxrr7r8r\rrrrM)r%rrreduction_treesr rn_blocksrs rr*z%TritonKernel.codegen_reduction_numelss -1,<,<RD@Q@Q4RRF"U_"UUV $**V"4!567 ((     ( ( .  y) / 60B/CDES  sCC9(Cc ||j}|Dcgc]}tj||fi|c}Scc}w)zK Helper to initialize symbols like rn_numel, rn_base, etc. )rrr)r%rr rn_prefixesrfs rrfz#TritonKernel._get_reduction_symbolss=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. rTrrWN)rrfrvrrMrr)r%r rn_numelsrs r_get_reduction_index_coeffsz(TritonKernel._get_reduction_index_coeffs s{113 //PT/U ;@[AQTUAU;V 47M)C!GI. / ]]1    sA0c:|j}t||S)zK Compute linear reduction indices from N dimensional ones. )rrL)r% multi_indscoeffss rrgz'TritonKernel._flatten_reduction_indicess113,,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)rfrgr\r)r%r rn_offsetsrn_indsrrindexs rrhz&TritonKernel.codegen_reduction_indices s 00 d1 --gtQU-V11*= 4#4#4W#=">?@009 $"3"3F";!<=>rc f|jr |jr|jr|jrJd|j D]3\}}}|rdnd}|j ||d|j |5|j |jd|j |jd|j||j |jd|jd|j |jy|j}|jr%|j |jd|d|dn|jD|j |jd|j||j |d n|j|d|j|}n|j!||d }|j"s|j%d nd } |j'|d |j)|d| |jd|g|j+|r!|j |j-|y|dk(r |j"s3|j |jd|jd|dyy)Nz=derived reduction roots do not support cooperative reductionsz: tl.constexprrrOrcrz offset + r5z offset = 0rrr z offset = rvr]r)has_custom_codegen_headerr+rrnamed_constantsrVrrrrerrrfrcrXrrrrrrrr@) r%r?rr}r constexpr annotationr]rrs rrdz,TritonKernel.iteration_ranges_codegen_header0s  * * ,**--&& OO  O ).(=(=(? Q$T91:- #zl#d6G6G6M5NOP Q NN::,c$"3"3E4F4F4H"I!J#44U;<>  NN??$%S C8I8I%++8V7WX   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 NNeoo/0EJJ^FO^ ^@I &15-199C9& 9 $ 9  9.9+9 9v. 9 ( 1<  (s s s & s  5 s  / s j&DO&" 0F P#0 &!  B>>> >@    0                 "  > :  6    !   4HH/H HH>S1>S;R>S >S@"'" " ( " !"BD"'D"(D" D"  D" ! D"L $q%f0wr  t_%"%"N  00d/b11fgR 7 C+ C._replace_argsO 9>>"%553: /0%-- ,;;(*$+GM ,rrrLr)rzast.Callrrrr)split_and_set_rangesrr get_rangescodegenrdeepcopyr kernel_ast*torch._higher_order_ops.triton_kernel_wraprrrrOrNamerr store_noder splitlineslineno col_offsetr get_lines_refr'r7)r%rnew_astrrrnew_store_value_nodersrc_with_store_replaced src_linesstore_line_index indentationsr" load_lines compute_lines new_src_liness rrz$FusedUserDefinedTritonKernel.codegens  C22##22==?J    . . 6 6z B  C-- 7 78 8@ =''(A---"xx(>(>(C(CD , , ," ,    # . .4FO4 4(<(rrc 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^) rrrr^rr>r<debug_triton_codedebug_device_str)r%rrrs rrzTritonScheduling.__init__sM #  GIw$? OO :D$0B CD(9% :rctjjstjjr't g|j t jS|j Sr)r$rcooperative_reductionsforce_cooperative_reductionsrbackend_featuresrZREDUCE_TO_SINGLE_ELEMENT)rrs rget_backend_featuresz%TritonScheduling.get_backend_features sR 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)r:ForeachKernelSchedulerNodec36K|]}t|ywr)r)rrr s rrz3TritonScheduling.codegen_comment..!s>? 189sz Fused node name list: r)rUrarrG make_commentr$ debug_fusiontorch._inductor.schedulerr:r rrrcommentr7r-write_provenance_debug_handle) r% node_scheduler?rorigins_detailed_originsr:r node_names debug_handler s @rcodegen_commentz TritonScheduling.codegen_comments''&&%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_overrider_ async_compileuse_process_poolrrVr\rUrarrM define_kernelrI)r%rrsrc_coder? subs_namer kernel_pathrGcompile_wrapperrmetadata_commentr detailed_originss r_emit_kernel_to_wrapperz(TritonScheduling._emit_kernel_to_wrapper6s  & &         (*  ) ) +  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) NrrrYrtriton_z #pragma CMT#pykernel_metadata)rUrar src_to_kernelr$rdescriptive_namesrFrRcustomize_fused_kernel_namerVr7next_kernel_suffix aot_inductormodel_name_for_generated_filesunique_kernel_namesrTrrIrQrr*r)rQr" rGr&is_metric_table_enabledlog_kernel_metadata) r%r r rrr? fused_namekernel_categoryr _basenamerYr s rr zTritonScheduling.define_kernel`s''&& 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|]}|jywrrrrs rrz9TritonScheduling.benchmark_fused_nodes..s:WA1::<:Wr)r )generate_kernel_code_from_nodesr+rWbenchmark_codegened_moduler)r%r^n_spills_thresholdr rs 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)rB rE mss r store_cachez@TritonScheduling.benchmark_codegened_module..store_caches&(T3r7+rc}tjj|r.t|5}t |j cdddSy#1swYyxYwr)rA rB existsopenr<readrB fdrE s 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 %srrWc4jdSr clone_argsrcallwrapped_jit_functionsrrz=TritonScheduling.benchmark_codegened_module.. D!@!5!@!@$!G!JKrrc"jSrrS rrW srrz=TritonScheduling.benchmark_codegened_module..s ? 4 ? ? Frz+The fused kernel for %s took %.3f ms to run)rrUrarXrrrrrrr@ get_argsrV r$ rT Exceptionr$r.disallow_failing_autotune_kernels_TESTING_ONLYr< launchersrn_spillsr0 benchmarkr@r)r%rr: r device_interfacerI rP r2r_ rrrE rV rH rW s ` @@@@@rr9 z+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 JJcD|jd}|xr td|jD}|j}|rddlm}|}|rd|d<|j ||tjj||||}||i|}|j|||S)Nr|c3<K|]}|jywr) is_split_scan)rrs rrz9TritonScheduling.create_kernel_choices..s( %)D   ( rrW)TritonSplitScanKernelFrn) rorscheduler_nodesrtriton_split_scanrf rrrUrRtriton_kernel_kwargsadd_multi_kernel_choices) r%rp kernel_argsrqis_scanre rrf rs rcreate_kernel_choicesz&TritonScheduling.create_kernel_choicess "--f5 C( -<-L-L-N( % +/*:*:  @/K >CM: ;44_mT 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) NrmrnFi)rnrmrWc|jSr)r)rs rrz;TritonScheduling.add_multi_kernel_choices..Js q'='=rrx)r$r multi_kernelrrr+rrrmrTrUrarZr)rmust_keep_buffersr) r%rrk rqkernelsoptional_persistentoptional_cooperativerrmkernel2s rrj z)TritonScheduling.add_multi_kernel_choicess (.h}}))N$99 -BSBS +C ?  &;; MDUDU ,E A   NN    #38  __44Fww44VUC-T--$'8=E'5+E+ENN((((+rc fdfd}fd}dg}}d}tjj}t|tj_tjj} t| tj_t j dkD} t jdkD} t j} |j|d| | d| } | D]^\}}}|Dcgc]}|j}}|Dcgc]}|D]}|j}}}t|dk(r'||d\}}||z }|dz }|j||J|jtt j"d }t%j&|t(j+d |j,|\'|z }|z }|jj,j/j0j2  j4d j6}t|dk(sJ|dj8dkDrt;d xnUtjj=}t?j@ fd |t?j@ fd|t(j+dtd|D||z }|z }|jj,d xx tBjDjGa|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~jJtjjjddzSr> r? rD srrE z@TritonScheduling.benchmark_combo_kernel..cache_file_path[s6<<+ ++77##CLL1!4~E Erc}tjj|rCt|5}t d|j j DcdddSy#1swYyxYw)Nc32K|]}t|ywr)r<)rr2s rrzNTritonScheduling.benchmark_combo_kernel..load_cache..cs Eaq Err)rA rB rK rL rrM rPrN s rrP z;TritonScheduling.benchmark_combo_kernel..load_cache_s^"$Dww~~d#$ZF2 E2779??3D EEFFFs .A,,A5c\}t|tdztzy)NrLrG )rB rE rH ms_clones rrI z.store_cachefs&"$D s2w}s8}< =rrgT)subkernel_nodescustom_part_algorithmenable_autotune mixed_sizesonly_gen_src_coderrWNr$ rQ rc4jdSrrS rU srrz9TritonScheduling.benchmark_combo_kernel..rX rrY c(jdSrrS r[ srrz9TritonScheduling.benchmark_combo_kernel..s;0;;TB1ErzDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc3<K|]}|jywrr6 r7 s rrz:TritonScheduling.benchmark_combo_kernel..s ) # ww66",-A"B"#''"<"<%/0G%H" 77!;;;a? %GG::%"&+#"!5 ; -=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\\ *15 5D 54.    ) ) +CA ,D#7%<"22G IOs *M/M4)rzScheduler | Nonerr)rz torch.devicer)r)rtuple[float, str])rN)r zOrderedSet[str] | Nonerr )rpr~rk list[Any]rqrqrlist[TritonKernel])rrrk r rqrqrr )rrrrrrrrrZFOREACH BUCKETIZEINPLACE_BUFFERSMASKED_SCATTER_WITH_INDEXSCANSORTTRITON_TEMPLATESTUPLE_REDUCTIONr rrr r r" r r; r9 rm rj rrrs@rrrs?)K)!  " "  $ $  * *  4 4        + +  * *  :$$MB(YT6p  OSW$5KW$ W$rQ+QQ& Q  Q8333& 3  3jt3rrcg}|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_renderrr0torch._inductor.codegen.cuda_combined_schedulingr 3torch._inductor.codegen.xpu.xpu_combined_schedulingr rr get_backendrorMrUraset_current_devicer&generated_kernel_countr8 r rQrra) rlinesmulti_templater r rbackendold_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)rrrrrrr<)rrrrv)rrrrrr0)rrrr)rrrr)rzCSEVariable | Anyrr)rr rr/)rr rr)rrzr)rrrzCallable[[_T], _T])rr:rrv) __future__rrrrQrrrArloggingrrrA rabcrcollections.abcrrrrtypingr r r r rsympy.printing.precedencer rtorch._loggingtorch.utils._pytreer'_pytreertorch._dynamo.device_interfacertorch._dynamo.utilsrrtorch._prims_commonrrtorch.utils._ordered_setrtorch.utils._sympy.functionsrrrrrrtorch.utils._tritonrrrutils._sympy.symbolrrr r!utils._sympy.value_rangesr"rr$r%r&r r( codecacher)r*r+r,rr- ops_handlerr.runtimer/runtime.benchmarkingr0 runtime.hintsr1r2r3r4r5r6r7runtime.runtime_utilsr8r9rr:r;r<r=r>shape_propagationr? stream_utilsr@rArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQ virtualizedrRr rSrTrUwrapper_benchmarkrVblock_analysisrXcommonrYrZr[r\r]r^r_r`rarbrcrdrerfrgrhsimdrirjrkrlrmrnro triton_utilsrprqrrrsrtrurvrrwtypesrxrrz%torch.fx.experimental.symbolic_shapesr{r|r}simd_kernel_featuresr~r getLoggerrr_logginggetArtifactLogger perf_hint_log schedule_log fusion_logr#rrrrrrrrr?rrrrrrrArrrrrrrrrrrrrrrr _initialize_pointwise_overridesr%rtr@rrrrrrrrrr rrrr sL"    8844 0 $$C<?/ YX4))(FF;('.D6.&CBB/$% L>&8 Bg!00<H ~~//*E ^^ - -h A  #$, &' &'(&'&'R 6 6  4$h h V $ $ $ N [$[$[$|  ?4??8 =,==@6+ +/+0+ +BpQMpQf 3 &8 ;P2 28( (VW&kW&x 33T:22D9//9I"OI"X$+$+N N N N b!!H ### %sU38_'<<=  oood V=M:/0V=MrzJ(<J(ZP3~P3f#r