+ jYN ^RIHt^RIt^RIt^RIt^RIt^RIt^RIt^RIt^RI t ^RI t ^RI t ^RI t ^RI t ^RIHt^RIHtHtHt^RIHt^RIHtHtHtHt^RIt^RIHt^RIt^RIt^RIH uH!t"^RI#H$t$^R I%H&t&H't'^R I(H)t)H*t*^R I+H,t,^R I-H.t.H/t/H0t0H1t1H2t2H3t3^R I4H5t5H6t6H7t7H8t8^RI9H:t:H;t;HH?t?^RI@HAtAHBtBHCtCH t ^RIDHEtE^RIFHGtGHHtHHItIHJtJ^RIKHLtL^RIMHNtN^RIOHPtP^RIQHRtR^RISHTtTHUtUHVtVHWtWHXtXHYtYHZtZH[t[^RI\H]t]H^t^^RI_H`t`HataHbtbHctcHdtd^RIeHftf^RIgHhthHitiHjtj^RI HktkHltlHmtmHntnHotoHptpHqtqHrtrHstsHtttHutuHvtvHwtwHxtxHytyHztzH{t{H|t|^RI}H~tHtHtHt^RIHt^RIHt^R IHtHtHtHtHtHtHtHtHtHtHtHtHtHtHtHt^R!IHtHtHtHtHtHtHtHt^R"IHt^R#IHtHtHtHtHtHtHt^R$IHt]'d-^R%IHt^R&IHt^R'IHt^R(IBHt^R)IHt^R*IHt]!R+4t] EPn!]4t]EPtEPw]R,4t]EPtEPw]R-4t]EPtEPw]R.4t]E!4tDR/tR0R1ltR2R3lt!R4R54t]!R4R6R7l4t!R8R94t]EP!R:R;44t]EP!R<R=44t]EP!R>R?]44t]EP!R@RA]44tRBRCltRDREltRFRGlt!RHRI]4t]!4EPtRJRKltRLRMltRNROltRPRQltRRRSltRTRUltRVRWltRXRYlt!RZR[]4t]EP!R\R]7!R^R_44tR`RaltRRbRcllt!RdRe]4t]EP]EPR\4]EP]EPR\4]EPRf4!RgRh]4t!RiRj4t]EP!RkRl44t!RmRn4t]EP!RoRp44t]]]]3,,t]]Rq3,t]]R,Rq3,t]EP!R\R]7!RrRs44t]EP!R\R]7!RtRu44t]],t!RvRw]]]3,4t]EP!RxRy44t!RzR{]],4t!R|R}]4t!R~R]4tRRltR#)) 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_cpu_backendhas_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 TritonMeta)get_max_y_gridnext_power_of_2)BaseSchedulerNode$FusedExternTritonKernelSchedulerNodeFusedSchedulerNode Scheduler SchedulerNode)get_broadcasted_shape)coor_benchmark_device_idxcoor_device_strget_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 TMA_ALIGNMENT 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_reprDerivedIterationRangesRootIterationRangesIterationRangesEntryIterationRangesRootPartialAccumulate SIMDKernelSIMDScheduling)%tiling_scores_suggest_inner_reduction) 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_hintsschedulefusioncVR9pV'dRMRpVR9d,\P'd VR 9dRMRpV RV RV 2#V RV 2#) anytriton_helperstl_strict.2)rmaxminprodfmax)rrr)rr)r#strict_signed_zero)reduction_type use_helpermodulestricts& _/home/luca/screen-checking/.venv/lib/python3.14/site-packages/torch/_inductor/codegen/triton.pyget_triton_reduction_functionrsn#HHJ!+ F//(((^~-M   >*!F844>*++cV^8dQhRR/#)exprobject)formats"r __annotate__rs   rc\V\P4'gR#\V\P4;'g-VP;'d\ VP 4^8H#)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. F) isinstancesympyExprInteger is_integerlen free_symbolsrs&ris_sympy_integer_likersT dEJJ ' ' dEMM *   77C 1 12a7rc$V^8dQhRRRRRR/#)rr sympy.Exprdtype torch.dtypereturnr)rs"rrrs$&'&' &'(&'&'rcaaVP'dVP\4'gV#VP\Jd\ VP !#RRloRVV3RlloS!V4#)c V^8dQhRRRR/#)rnodez sympy.Basicrboolr)rs"rr6_materialize_trunc_to_float_expr..__annotate__s    rcX\\VRR4;'g\VRR44#) is_BooleanF is_Relational)rgetattr)rs&ris_predicate_expr;_materialize_trunc_to_float_expr..is_predicate_exprs- D, . W W'$QV2W  rc V^8dQhRRRR/#)rrrrr)rs"rrrs$$J$:$rc<VP\4'gV#VP\Jd\VP!#S!V4'gVP 'dV#\ ;QJd".VV3RlVP4FNK 5M!VV3RlVP44pWP8XdV#VP!V!#)c3<"TF=p\V\P4'dS!V4'g S!V4MTxK? R#5iN)rrr).0argrrewrite_float_subexprs& r R_materialize_trunc_to_float_expr..rewrite_float_subexpr..sG !#uzz**3DS3I3I "# & !s 2AA)hasrfuncrargsrtuple)rnew_argsrrs& rr?_materialize_trunc_to_float_expr..rewrite_float_subexprsxx ##K 99 "+ + T " "doooK5 yy  55 yy    yy Kyy(##r)is_floating_pointrrrrr)rrrrs&&@@r _materialize_trunc_to_float_exprrs]  " " "$((:*>*>  yyJTYY'' $$$ ! &&rcN]tRt^t$Rt/tR]R&/tR]R&]RRl4t Rt R #) 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$V^8dQhRRRRRR/#)rrzCallable[..., str]convert_outputrrNoner)rs"rrOpDtypeSupport.__annotate__s#66#56t6PT6rc VPp\\P\P.4VP V&W P V&R#r)__name__rtorchfloat32float64rr)clsrrop_names&&& rregister_upcastOpDtypeSupport.register_upcasts;--(2EMM5==3Q(RW%'5G$rrN) r __module__ __qualname____firstlineno____doc__r__annotations__r classmethodr__static_attributes__rrrrrs1 <>8=')O_)66rrcV^8dQhRR/#rrstrr)rs"rrrsCrc\4'gR#^RIp\VPPR4'dR#R#)zX import AttrsDescriptor if the triton version is new enough to have this class defined. rNAttrsDescriptorz4from triton.compiler.compiler import AttrsDescriptor)rtriton.compiler.compilerhasattrcompiler)tritons rgen_attr_descriptor_importrs3   #v''):;;Erc @]tRtRtRt]!]P]P.4t ]!]P]P]P.] O4t ] Uu/uF+pV\P!\ V, R2RRR7bK- upt] Uu/uF9pV\P!\ V,P%4 R2RRR7bK; upt]RR l4t]R R l4t]R R l4t]RRl4t]RRl4t]RRl4tRtR#uupiuupi) TritonSymbolsizM Stores sympy.Symbol instances and constants associated with triton codegen. offsetTinteger nonnegativeBLOCKrpositivec V^8dQhRRRR/#)rrrrrr)rs"rrTritonSymbols.__annotate__s00:0.0rc  RpVPpVEFp\V\P4'dM\P P PVP,p\VP4pEM*\V\P\P\P\P\P\P 34'dRpM\P P"P%V4pVf\'RVP 24hVP(p\P P+4p R.V ,p VP,fRpM1\/VP1V44WP,&\V 4p\3W&4pEK Vf \'R4hV#)NzUnregistered range symbol: 1zexpr_shape must not be Noner)rr r!TMPrXkernelcse varname_mapnamershape UNBACKED_INTSIZEPRECOMPUTED_SIZEINDEXFLOATUNBACKED_FLOATrange_tree_nodesgetAssertionErrorroottriton_tensor_ndim tensor_dimrget_block_sizer?) rr expr_shape expr_varsvarcse_var var_shapertreendimrs && rget_block_shapeTritonSymbols.get_block_shapes_&( %% Cc488,,((,,22388<!'--0 %%II))JJJJ''    xx0044S9<(+Fsxxj)QRRyyxx224 ??*!#I-01C1CD1I-JE//* %e I/zEJAD   !>? ?rc V^8dQhRRRR/#rrrmrrr)rs"rrrPs&&/&j&rc 6VPP4#r)r  block_sizerrs&&rrTritonSymbols.get_block_sizeOsyy##%%rc V^8dQhRRRR/#rr)rs"rrrTs((O( (rc 6VPP4#r)r  block_offsetrs&&rget_block_offsetTritonSymbols.get_block_offsetSsyy%%''rc$V^8dQhRRRRRR/#)rr TritonKernelsymbol sympy.Symbolr str | Noner)rs"rrrXs$  ! +7   rc VPPV4;peVPP4#VPF&p\ W$4'gK\ V, R2u# R#)Nmask)rr r  mask_name block_typesr r)rrr%rsymts&&& rmask_name_for_symbol"TritonSymbols.mask_name_for_symbolWsd ++//7 7D D99&&( (OODf++$T*+400$rc$V^8dQhRRRRRR/#)rrr$r%r&rrr)rs"rrres$;;!;+7; ;rc VPPV4;peVPP#\ W P 4#r)rr r  is_reductionr reduction_types)rrr%rs&&& ris_reduction_index_symbol'TritonSymbols.is_reduction_index_symboldsA++//7 7D D99)) )f&9&9::rc$V^8dQhRRRRRR/#)rrr$rrrrr)rs"rrrms"   J SW rc aa\;QJd*VV3RlVP4F 'gK R# R#!VV3RlVP44#)zHReturn whether any symbol in a composite block shape is reduction-owned.c3H<"TFpSPSV4xK R#5ir)r3)rr%rrs& rr;TritonSymbols.has_reduction_index_symbol..os( +  ) )&& 9 9+s"TF)rr)rrrsff&rhas_reduction_index_symbol(TritonSymbols.has_reduction_index_symbollsLs ++ ss s s ++   rrN)rrrrrrr!R0_INDEXR1_INDEXr2XBLOCKYBLOCKZBLOCKr+rSymbolr block_offsetsupper block_sizesrrrr!r-r3r9r)r,s0rrrsU!$--!?@Odkk4;; VoVWK D ellj./v6RVWWM  D ell$%%'( .t   K00d&&((  ;;  { s 1D?Drc]tRtRt$R]R&R]R&R]R&R]R &R ]R &R ]R &RRltRRltRRltRRltRRlt ] RRl4t Rt R#)IndexingOptionsiur index_strOrderedSet[str] mask_varsr' expand_strr _has_rindexrindexzSequence[int | str] | None expand_shapecV^8dQhRR/#rrrr)rs"rrIndexingOptions.__annotate__~s$$$$rc ,\VP4#r)rrHselfs&rhas_maskIndexingOptions.has_mask~sDNN##rcV^8dQhRR/#rNr)rs"rrrOs99d9rc J\VP\P4#r)rrKr!rrQs&r has_indirectIndexingOptions.has_indirects"4::txx88rcV^8dQhRR/#rNr)rs"rrrOs  D rc VP#r)rJrQs&r has_rindexIndexingOptions.has_rindexsrcV^8dQhRR/#rNr)rs"rrrOsKKTKrc \;QJd&RVP4F 'gK R# R#!RVP44#)c3V"TFp\V4PR4xK! R#5i)tmpN)r startswith)rr)s& rr.IndexingOptions.has_tmpmask..s"J>43t9''..>')TF)rrHrQs&r has_tmpmaskIndexingOptions.has_tmpmasks3sJ4>>JssJsJsJ4>>JJJrcV^8dQhRR/#rNr)rs"rrrOs  4 rc a\;QJd;V3Rl\PP44F 'gK R# R#!V3Rl\PP444#)c3<"TFUpSPFBp\V\4;'d&VP;'dVP V4xKD KW R#5ir)rHrrr1 owns_mask)rrr)rRs& rr,IndexingOptions.has_rmask..sS 5 tS ! P Pd&7&7 P PDNN4

mm mrc  a\PPoRV3RllpV!VP4VnV!VP4VnVP U u.uFp SP V ^4NK p p \V 4'dRV R&VPV 4pVU u.uFp \PV 4NK p p \PP'dJV^,PR8wd!\RV^,P: 24hV P^4\!4p VP Fp\#V\$P&4'gK%VP(FQp\+V4P-4p\/V4'g\1V4'gK@V P3V4KS K \PP4p\PP6'g\9V 4\9\PP:4V, 8XdX\PP<P?4'd*V \$P@PB.V,, p VPEV\PPFR7wppVPUu.uFpSP V^4NK ppVP pVPV4p\HPJ!\PPFVP4pV!V\PPPSV4VVV VVVVR7 pVPUWS4V#uup iuup iuupi \d4\M\O\Q\9TP4444pLi;i) z2Helper to create a BlockDescriptorOptions instancec V^8dQhRRRR/#)rexprszIterable[sympy.Expr]rrr)rs"rr3BlockDescriptorOptions.create..__annotate__s  N N3 N8H NrcP<VUu.uFpSPV4NK up#uupir)lookup_precomputed_size)rrsizevarss& r lookup_size2BlockDescriptorOptions.create..lookup_sizes&GLMutH44T:uM MMs#Fxz0expected first range tree prefix to be 'x', got )r shape_env) rzr{r}rHrrrrr)+rXgraphrrrrstatically_known_equalsall remove_dimsrrrno_x_dimprefixr poprrrrrrlowerrMrLaddnum_reduction_dimsinside_reductionrnumelsfeaturesr1SOnemaybe_sort_with_stride_order _shape_envr& argsort_symlistreversedrangercompute_boundary_check)rrzr{rrHrrrrdimsingleton_dimsrrincluded_final_shape_dimsrsymsym_strreduction_ndimrstriderrr}resultrs&$$$$$$$ @rcreateBlockDescriptorOptions.creates077## N N#6<<0 $V^^4AG@R@R @RH , ,S! 4@R   ~  !&N2 ##N3GRRkd}33D9k R 88   1~$$+$F"1~,,/1 OOA %/L!&&D$ ++,,C!#hnn.G*7337J77S7S155c:-'44)))-.#ahhoo2F2WW!!..00 EGGKK=>9 9K!' C C/177;M;M!D!  GMnn FTFH , ,VQ 7n  !,,##$56 ?%%agg&8&8&..IE GG,,DD_U#+/'   %%mA g SP  ?%FNN(;"<=>E ?s$!N)NN 9N%%;O#"O#c(V^8dQhRRRRRRRR/#rrr replacementr,r!rr)rs"rrr5,888-78?C8 8rc J\PV,p\WV/4#z> Replaces instances of {symt}_offset with the new expression. rrArPrRrrr,roffsets&&&& rreplace_offset%BlockDescriptorOptions.replace_offset5$  --d3$+ 677rc V^8dQhRRRR/#rrrrr)rs"rrr>J:rc \PF*pVPV\P!^4V4pK, V#rrr2rrrrRrr,s&& rremove_roffsets&BlockDescriptorOptions.remove_roffsets>4!11D&&tU]]1-=tDD2 rc$V^8dQhRRRRRR/#)rrrrrrrr)rs"rrrCs(2 2 +2 /2   2 rc  \PPpVUu/uF@p\PVP ,V!\ VP ,4bKB pp\\\PPV44p\\VP44Uu.uEFwpVPVPV,\ P"P$4'dKGV'dF\P\&P(,VP*V,P,9gVP/VPV,VP*V,4'dKVP/VPV,\1VP*V,V44'dEK\PP2'd?VP*V,\P\&P4,8XdEKuVNEKz upVnR#uupiuupi)z6List of indices to pass to tl.load(boundary_check=...)N)rXrrrrCr,rrrsrneeds_yz_grid_overflowrrrrrrrZeror!r>rrstatically_known_multiple_ofrPrr=r)rRrrrt block_to_maxneeds_overflow_grididxs&&& rr-BlockDescriptorOptions.compute_boundary_checkCs 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'? ?###rcV^8dQhRR/#rNr)rs"rrr|sdrc R#FrrQs&rrW#BlockDescriptorOptions.has_indirect|rcV^8dQhRR/#rNr)rs"rrrs  D rc \;QJd&RVP4F 'gK R# R#!RVP44#)c3~"TF3p\P\R\P4V4xK5 R#5ir$N)rr9r rXr)rrs& rr4BlockDescriptorOptions.has_rindex..s; )  4 4^QXX.  )s;=TF)rrrQs&rr[!BlockDescriptorOptions.has_rindexsMs ((  ss s s ((    rcV^8dQhRR/#rNr)rs"rrrs!!4!rc "VP4#r)r[rQs&rrl BlockDescriptorOptions.has_rmasks  rcV^8dQhRR/#rNr)rs"rrrsTrc R#rrrQs&rrd"BlockDescriptorOptions.has_tmpmaskrrcV^8dQhRR/#rNr)rs"rrrs++$+rc 4\VP44#r)rrrQs&rrSBlockDescriptorOptions.has_masksD'')**rc 0V^8dQhRRRRRRRRRRR R /#) rvaluer initial_shaper~rallow_implicitr for_storerz,tuple[str, Sequence[sympy.Expr | int | str]]r)rs"rrrsNp$p$p$,p$* p$  p$  p$ 6p$rc aaVoRV3RllpVPpVPpV'dKVPPVP4pVPPVP4p\ Wx4U U u.uF)wrV 'd\ P PMT NK+ p p p V!WV 4pT p VPP'g=V'g5\V 4\V48XdVPPV 4p \PPoT;'dl\V 4\V48H;'dM\;QJd(V3Rl\ W44F 'dK RM RM!V3Rl\ W444p \VP4'd1V 'g)RV R\PP!V4 R2pVoVPpVPP'gV'dVPP"MVPP$pR V RV R2pVUu.uF pSV,NK upoV'd VPM$VPPVP4pV!WV4pVS3#uup p iuupi) 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 c(V^8dQhRRRRRRRR/#rrr old_shape Sequence[sympy.Expr | int | str] new_shaperr)rs"rrJBlockDescriptorOptions.codegen_broadcast_and_reshape..__annotate__s0   7 8  rcV<\WV4p\V4\V48wdVoV#r)triton_reshapetriton_shape_dims)rr rr current_shapes&&& r apply_reshapeKBlockDescriptorOptions.codegen_broadcast_and_reshape..apply_reshapes. $Ei@F +/@/KK ) Mrc3~<"TF2wrSPV^4;'gSPW4xK4 R#5iN)r)rpre_dimpost_dimrs& rrGBlockDescriptorOptions.codegen_broadcast_and_reshape..sK*T%G00!<GG33GFG)Ss==FTtl.broadcast_to(, ) tl.trans()rrrrevertziprrr is_identityrrXrrrrr index_to_strsort_idxrevert_sort_idx)rRrrrrr rrrris_broadcastingpre_broadcast_shapeimplicit_broadcast_shapesupports_implicit_broadcastr  permute_dimsirrs&&&&&& @@rcodegen_broadcast_and_reshape4BlockDescriptorOptions.codegen_broadcast_and_reshapes2;H   .. 22 "00778L8LMO $ 2 2 9 9$:P:P Q  ),O(O (O$+EGGKK 3(O  e4GH#6 ""...'(C ,<< (,'9'9'@'@AT'U $ 77##&4' ' ( )S-= =  *--E)S*--E)S $ t%% & &/J"5'AHH,A,A/,R+SSTU ,M(( !!---""++''77   wba8E7CD|!]1--|DM$$''..t/C/CD e <m##{ hEs /J9J?)r)rrrrrrrrrvrrrrrrrrrrrWr[rlrdrSr-rrrrrxrxs  ))!!%%0/(,O%,Hd !!''####mmm^8 2 h$  !+p$p$rrxc&]tRtRtRRRlltRtR#)TensorDescriptorOptionsic V^8dQhRRRR/#rrrrr)rs"rr$TensorDescriptorOptions.__annotate__s??3??rc 0\PPpVP^8wdV RV!VP4 R2MTRV!VP4 2RV!VP 4 2RV!VP 4 2.pRRPV4 R2#)z Codegen a call to tl.make_tensor_descriptor() Args: name: variable name for pointer roffset: unused, but kept for compatibility with BlockPtrOptions.format() Returns: "tl.make_tensor_descriptor(...)"  + (rshape=strides= block_shape=ztl.make_tensor_descriptor(r)rXrr$r{rrrrq)rRrrfrs&&& rrTensorDescriptorOptions.formats HH ! !''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0  ,DIIdO+>rrNT)rrrrrrrrrr0r0s ??rr0cJ]tRtRtRRltRRltR RRlltRR ltR tR #) BlockPtrOptionsi c(V^8dQhRRRRRRRR/#rr)rs"rrBlockPtrOptions.__annotate__"rrc J\PV,p\WV/4#rrrs&&&& rrBlockPtrOptions.replace_offset"rrc V^8dQhRRRR/#rr)rs"rrr?+rrc \PF*pVPV\P!^4V4pK, V#rrrs&& rrBlockPtrOptions.remove_roffsets+rrc V^8dQhRRRR/#r2r)rs"rrr?0s77377rc  \PPp.VPOpV'g VUu.uFqPP V4NK ppVP ^8wdV RV!VP 4 R2MTRV!VP 4 2RV!VP4 2RV!VP4 2RV!VP4 2RV!V4 2.pRR PV4 R2#uupi) z Codegen a call to tl.make_block_ptr() Args: name: variable name for pointer roffset: should rn_offset be included in offsets=..., for use with tl.advance() Returns: "tl.make_block_ptr(...)" r5rr6r7r8zorder=zoffsets=ztl.make_block_ptr(r) rXrr$rrr{rrrr}rq)rRrrr9rrrs&&& rrBlockPtrOptions.format0s HH ! !!DLL/BIJ'++F3'GJ''1,&Qt3345Q7Qtzz]O $q' (1T--./ 0Qtzz]O $qzl #  $DIIdO#4A66KsC+c V^8dQhRRRR/#)rr,r!rrr)rs"rrr?MsDZrc  \PV,pVPUu.uFEpVPW2V4VPV\P P V4, NKG ppV#uupi)aF Codegen string to pass to tl.advance(name, ...). Advance is the difference between offsets in each loop iteration. To compute it, we replace rN_offset with multiples of RN_BLOCK. Since we expect rN_offset to vary in range(0, rN_numel, RN_BLOCK), the first iteration has rN_offset=0, while the second has rN_offset=RN_BLOCK. )rrCrrrrr)rRr,rblockradvances&& radvance_roffsetBlockPtrOptions.advance_roffsetMs}**40 ,,  '##FD9%%feggllDABB'   sA A5rNr;) rrrrrrrrLrrrrr=r= s8 7:rr=c V^8dQhRRRR/#)rrrr list[str]r)rs"rrras=)rcVUu.uF:p\V\4'dTM\PP V4NK< up#uupi)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. )rrrXrr$rrs& rrrasGPU OTz#s##)>)>s)CCu  sAAc V^8dQhRRRR/#)rrrrrr)rs"rrrls66<66rc>RRP\V44 R2#[r])rqrrs&rtriton_shape_strrXls! tyy*51231 55rc(V^8dQhRRRRRRRR/#r r)rs"rrrps0++ +/+0+ +rc\V\4'd\V\4'g \R4h\V4p\V4pW48XdV#VUu.uF qUR8wgK VNK upV8wdRV RRP V4 R2#^p.pVFcpV\ V48d*WV,8XdVP R4V^, pK<VR8wd\RV: 24hVP R 4Ke V\ V48wd\R V R \ V4 24hV R RP V4 R 2#uupi)zUG3tyy'?&@CC C F ]# #c0B(B MM#  1HCs{$'>th%GHH MM& ! c-  6se4M@R?S T  WAdii'( **#.s D;'D;c]tRtRtRRltRRlt]RRl4tRR ltR R lt R R lt RRlt RRlt RRlt RRltRRltRRltRRltRRltRRltR R!ltR"R#ltR$R%ltR&R'ltR(R)ltR*R+ltR,R-ltR.R/ltR0R1ltR2R3ltR4R5ltR6R7ltR8R9lt R:R;lt!R<R=lt"R>R?lt#R@RAlt$RBt%RC#)D TritonPrinteric V^8dQhRRRR/#rrrrrr)rs"rrTritonPrinter.__annotate__  j S rc \VP4^8wd"\R\VP4 24hRVPVP^,4 R\P P R2#)rexpected 1 arg, got libdevice.trunc().to(rrrr _printrXr index_dtyperRrs&&r_print_TruncToIntTritonPrinter._print_TruncToInte tyy>Q  #7DII7G!HI It{{499Q<89qxx?S?S>TTU V rc V^8dQhRRRR/#rhr)rs"rrrisFF FsFrc \VP4^8wd"\R\VP4 24hVPVP^,4pRV RV R2#)rrlz(libdevice.trunc(z) + tl.zeros_like())rrr rp)rRrrs&& r_print_TruncToFloat!TritonPrinter._print_TruncToFloats\ tyy>Q  #7DII7G!HI I DIIaL)#5');E7"EErcV^8dQhRR/#rr)rs"rrrisCrc \\\PP4'gR#R#)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)rFrXrcurrent_devicerrr_get_scalar_float_type$TritonPrinter._get_scalar_float_types$AGG$:$:;;rc V^8dQhRRRR/#rhr)rs"rrris    rc VP'd\\V44pV#\P!4'd&\ P P'dT pV#VP4pRV RV R2pV#) tl.full([], rr) rrintr# is_fbcoderversionhipr)rRrret float_types&& r _print_FloatTritonPrinter._print_Floatsx ???c$i.C      EMM$5$5$5FC 446J b A6C rc V^8dQhRRRR/#rhr)rs"rrris'':'#'rc \VP4^8wd"\R\VP4 24hVPVP^,\R,R, 4pVP 4pV RV R2#)rrlAtom?.to(r)rrr  parenthesizer r)rRrrbrs&& r_print_ToFloatTritonPrinter._print_ToFloatss tyy>Q  #7DII7G!HI I   diilJv,>,D E002 D A&&rc V^8dQhRRRR/#rhr)rs"rrrisFFZFCFrc VPwr#VP'dBVP'd0VPVPR\R,R, 4#VP V4pVP V4pRV RV R2#) % rrz!triton_helpers.remainder_integer(rr)ris_nonnegative stringifyr rprRrquotdivquot_sdiv_ss&& r_print_PythonModTritonPrinter._print_PythonModsvII     3#5#5#5>>$))UJv4F4LM MT" C 26("UG1EErc V^8dQhRRRR/#rhr)rs"rrris G GJ G3 Grc TVP'g \R4hVPwr#VP'dBVP'd0VP VPR\ R,R, 4#VP V4pVP V4pRV RV R2#)zexpr must be integer // rrz!triton_helpers.div_floor_integer(z, r)rr rrrr rprs&& r_print_FloorDivTritonPrinter._print_FloorDivs !78 8II     3#5#5#5>>$))VZ5G#5MN NT" C 26(#eWAFFrc V^8dQhRRRR/#rhr)rs"rrrisJJjJSJrc `VPVPR\R,R, 4#) / rr)rrr rrs&&r_print_IntTrueDivTritonPrinter._print_IntTrueDivs#~~dii 60BS0HIIrc V^8dQhRRRR/#rhr)rs"rrris    rc \VP4^8wd"\R\VP4 24hRVPVP^,4 R\P P R2#rrllibdevice.floor(rnrrorrs&&r _print_floorTritonPrinter._print_floorrurc V^8dQhRRRR/#rhr)rs"rrrirjrc \VP4^8wd"\R\VP4 24hRVPVP^,4 R\P P R2#rrorrs&&r_print_FloorToIntTritonPrinter._print_FloorToIntrurc V^8dQhRRRR/#rhr)rs"rrrisYY:Y#Yrc \VP4^8wd"\R\VP4 24hRVPVP^,4 R\P P R2#rrllibdevice.ceil(rnrrorrs&&r_print_ceilingTritonPrinter._print_ceilinga tyy>Q  #7DII7G!HI I TYYq\!: ;5AUAU@VVWXXrc V^8dQhRRRR/#rhr)rs"rrrisYYZYCYrc \VP4^8wd"\R\VP4 24hRVPVP^,4 R\P P R2#rrorrs&&r_print_CeilToIntTritonPrinter._print_CeilToIntrrc V^8dQhRRRR/#rhr)rs"rrrisCCCCrc ,RVPV4 R2#)z tl.sqrt_rn(().to(tl.float32)))rprrs&&r _helper_sqrtTritonPrinter._helper_sqrtsdkk$/00ABBrc V^8dQhRRRR/#rhr)rs"rrris U UJ U3 Urc  VPVP^,4pVPVP^,4pVP4pRV RV RV RV R2 #)rzlibdevice.pow((rnz), (rx)rprr)rRrbaseexprs&& r_print_FloatPowTritonPrinter._print_FloatPows`{{499Q<(kk$))A,'002  eJC vRuA..rc V^8dQhRRRR/#rhr)rs"rrri#s****rc VPVP^,4pVPVP^,4pVPVP^,4pRV RV RV R2#)r tl.where(rr)doprintr)rRrcpqs&& r _print_WhereTritonPrinter._print_Where#s_ LL1 & LL1 & LL1 &1#Rs"QCq))rc$V^8dQhRRRRRR/#)rrrcmprrr)rs"rrri)s&CC*C3C3Crc \VP4^8Xd#VPVP^,4#\VP4^,p\V4pVPV!VPRV!4pVPV!VPVR!4p\;QJd.RWV34FNK 5M!RWV344wrVVR 9d\ RV R24hRV RV RV RV R V RV RV RV R 2#) z1 Helper for max/min code generation. cmp: > or < Nc3."TF pRV R2xK R#5i)rrNrrrs& rr6TritonPrinter._print_min_max_helper..;s.v!q1XvszUnexpected comparator: ''rz * ( z= z) + rx)><)rrrptyperr )rRrrmidrabs&&& r_print_min_max_helper#TritonPrinter._print_min_max_helper)s tyy>Q ;;tyy|, ,$))n!4j KKTYYt_- . KKTYYst_- .u.v.uu.v.. j  #;C5!BC C1#T!AcU"QCtA3d1#Qse1QCrBBrc V^8dQhRRRR/#rhr)rs"rrri@55z5c5rc &VPVR4#)rrrrs&&r _print_MinTritonPrinter._print_Min@))$44rc V^8dQhRRRR/#rhr)rs"rrriCrrc &VPVR4#)rrrrs&&r _print_MaxTritonPrinter._print_MaxCrrc V^8dQhRRRR/#rhr)rs"rrriFs;;z;c;rc \VP4^8wd"\R\VP4 24hRVPVP^,4 R2#)rrl tl_math.abs(rryrrs&&r _print_AbsTritonPrinter._print_AbsFsO tyy>Q  #7DII7G!HI Idkk$))A,78::rc V^8dQhRRRR/#rhr)rs"rrriLNNZNCNrc \VP4^8wd"\R\VP4 24hRVPVP^,4 R2#)rrlzlibdevice.cos((rryrrs&&r_print_OpaqueUnaryFn_cos&TritonPrinter._print_OpaqueUnaryFn_cosLP tyy>Q  #7DII7G!HI I TYYq\!: ;;LMMrc V^8dQhRRRR/#rhr)rs"rrriROOjOSOrc \VP4^8wd"\R\VP4 24hRVPVP^,4 R2#)rrlzlibdevice.cosh((rryrrs&&r_print_OpaqueUnaryFn_cosh'TritonPrinter._print_OpaqueUnaryFn_coshRP tyy>Q  #7DII7G!HI I!$++diil";!<Q  #7DII7G!HI I  DIIaL 9:%@T@T?UUV W rc V^8dQhRRRR/#rhr)rs"rrris Q Q Qs Qrc R\VP4^8wd"\R\VP4 24hVPwr#VP'd%V^8d\RV 24h\ RV R24hVP V\ R,4pRV RV RV) 2#) rzexpected 2 args, got zexpected ndigits < 0, got zOFor integer inputs, only non-negative ndigits are currently supported, but got rMulzlibdevice.nearbyint(1e * z) * 1e)rrr r ValueErrorrr )rRrnumberndigits number_strs&& r_print_RoundDecimal!TritonPrinter._print_RoundDecimals tyy>Q  #8TYY8H!IJ J))    !|$'A'%KLLabiajjkl &&vz%/@A 'yJ>Q  : :5#:#::rc V^8dQhRRRR/#)rrzCSEVariable | Anyrrr)rs"rrrsPP/PDPrc\V\4'gR#VPp\V\P4'd \ V4#R#r)rr_rrrM)rrs& rlow_precision_fp_varrPs< c; ' ' IIE&0 &D&D E "O%Orc V^8dQhRRRR/#)rrr rtorch.dtype | Noner)rs"rrrs#"4rc\V\4'd VP#\V\PP 4'd\ \V44#R#r)rr_rr _prims_commonNumberrrrs&rtriton_arg_dtyperWsD#{##yy#u**1122T#Y'' rc V^8dQhRRRR/#)rrr rrr)rs"rrrsrc\PP'*;'dF\V\4;'d.VP \ P\ P39#r) r#rcodegen_upcast_to_fp32rr_rrfloat16bfloat16rVs&rneeds_upcast_to_float32r]sN MM 0 00 9 9 sK ( 9 9 II%--8 8rc8a]tRtRtRRV3RllltRtRtV;t#)TritonCSEVariableic ,V^8dQhRRRRRRRRR R /#) rrrboundszValueRanges[Any]rrrrrrr)rs"rrTritonCSEVariable.__annotate__sC F F F! F F  F  Frc <\SV`WW4R7\4VnVf \ R4hVf \ R4hR#)rWNz!TritonCSEVariable must have dtypez!TritonCSEVariable must have shape)super__init__rrHr )rRrrarr __class__s&&&&&rreTritonCSEVariable.__init__sG u:*4, = !DE E = !DE E rc vVFp\V\4'd(VPPVP4K@\V\P 4'gKb\ P\R\P4V4;pfKVPPV4K R#r) rr_rHupdaterr@rr-r rXrr)rRrrkwargsrr*s&&&& rupdate_on_args TritonCSEVariable.update_on_argssC#011%%cmm4C.. "/!C!C^QXX6"I  NN&&y1r)rHr)rrrrrerkr __classcell__rfs@rr_r_s F F 2 2rr_T)frozenc@]tRtRt$R]R&R]R&R]R&R]R &R tR #) TritonOpTraceEntryirrtuple[Any, ...]rztuple[tuple[str, Any], ...]rjr rrN)rrrrrrrrrrqrqs I  '' KrrqcV^8dQhRR/#)rrr~r)rs"rrrs((5(rc^RIHpV!4#)rr})!torch._inductor.dtype_propagationr~r}s rget_dtype_handlerrvsL % ''rc V^8dQhRRRR/#)rrrrzCallable[[_T], _T]r)rs"rrrs(((9K(rc*aaRRloRVV3RllpV#)z Codegen helper to upcast arguments to float32, depending on the config and dtype. This decorates tl.math/libdevice codegen functions. cV^8dQhRR/#rr)rs"rr*maybe_upcast_float32..__annotate__ s'''rc8\V4'dRMRpV V 2#).to(tl.float32)rr])r upcast_strings& rmaybe_upcast_arg.maybe_upcast_float32..maybe_upcast_arg s#-DS-I-I)r }o&&rc V^8dQhRRRR/#)rrzCallable[..., Any]rr)rs"rrrzs*/ArcL<a\PSS4RVVV3RllpV#)cV^8dQhRR/#rr)rs"rr=maybe_upcast_float32..decorator..__annotate__s 0 0 0rc|<VUu.uF pS!V4NK ppVP4UUu/uFwrEVS!V4bK pppS !V/VBpS ;'dy\;QJd>R\P!WP 444F 'gK RM3 RM/!R\P!WP 4444pV'gRM#\ \ 4S P4!V/VBp V \PR39p V 'dV eR\V 4 R2MRp V V 2#uupiuuppi)c38"TFp\V4xK R#5irr}rrs& rrKmaybe_upcast_float32..decorator..wrapped..s6AC(,,ATFNrrr) itemsr itertoolschainvaluesrrvrrrrR)rrjr upcast_argskeyval upcast_kwargsrany_needs_upcast result_dtypeneeds_downcastdowncast_stringrrrs*, rwrapped8maybe_upcast_float32..decorator..wrappeds-<@ADS+C0DKAHN WHCS"23"77MW;8-8F-  ##6$??4A6###6$??4A63  (.0$--@$Q&Q  *%--1FFN"l&>{<013  Xo./ /)BWs D3D8)rr)rrrrsf r decorator'maybe_upcast_float32..decorators$&&t^< 0 00rr)rrrsf @rmaybe_upcast_float32rs '< rc ]tRtRtRt]P !]P4t] ReRRll4t ] RRl4t ] R 4t ] R 4t] R 4t] R R l4t] ]!4R44t] R4t] R4t] R4t] ]!4R44t] ]!4R44t] ]!4R44t] ]!4R44t] R4t] R4t] R4t] R4t] R4t] R4t] RRR] PBRRR^R R/R!l4t"] ]!4R"44t#] ]!4R#44t$] R$4t%] R%4t&] R&4t'] ]!4R'44t(] ]!4R(44t)] ]!4R)44t*] ]!4R*44t+] ]!4R+44t,] ]!4R,44t-] ]!4R-44t.] ]!4R.44t/] ]!4R/44t0] R0R1l4t1] R2R3l4t2] R4R5l4t3] R64t4] ]!4R744t5] ]!4R844t6] ]!4R944t7] ]!4R:44t8] ]!4R;44t9] ]!4R<44t:] ]!4R=44t] R>4t;] ]!4R?44t<] R@4t=] RA4t>] RB4t?] RC4t@] RD4tA] RE4tB] RF4tC] RG4tD] RH4tE] RI4tF] RJ4tG] RK4tH] RL4tI] RM4tJ] RN4tK] RO4tL] ]!4RP44tM] ]!4RQ44tN] ]!4RR44tO] ]!4RS44tP] ]!4RT44tQ] RU4tR] ]!4RV44tS] RW4tT] ]!4RX44tU] ]!RYRZ7R[44tV] ]!RYRZ7R\44tW] ]!4R]44tX] ]!4R^44tY] R_4tZ] R`4t[] ]!4Ra44t\] Rb4t]] ]!4Rc44t^Rdt_R#)fTritonOverridesi0zEMap element-wise ops to Triton e.g., ops.to_dtype(x,...) -> x.to(...)NTc V^8dQhRRRR/#)rrr src_dtyperRr)rs"rrTritonOverrides.__annotate__6s"K&K&K&&K&rc a\P\P3oRV3RllpVe>\V!W!4\P P 4\P nVe(\V4'dRV R2p\PpV\P8XdRV R2#V\P8Xd VeVP'gVfV R2#V'd \V4pM \V4pVe6VS9d/V\P8Xg\V4'd V RV R2#V RV R2#) c$V^8dQhRRRRRR/#)rrr dst_dtyperrr)rs"rr.TritonOverrides.to_dtype..__annotate__As$  " /:   rc<W8Xd^#VS9dVS9dW8wd \R4hV\P8XgV\P8Xd^#V\P8XgV\P8Xd^#^#)rzCConversions between float8_e5m2 and float8_e4m3fn is not supported!)r r float8_e5m2 float8_e4m3fn)rr fp8_dtypess&&r_get_min_elements_per_thread>TritonOverrides.to_dtype.._get_min_elements_per_threadAsw%Z'+*$YE---e>O>O1OE///9@S@S3Srz$triton_helpers.fp8e4m3fn_to_float32(rrz != 0)z.to(tl.int16).to(tl.uint8)z.to(tl.float32).to(r)rrrrrXrmin_elem_per_threadrzrruint8rr=rAr)rrruse_compute_typesr out_dtypers&&&& @rto_dtypeTritonOverrides.to_dtype5s:         4  ,/,Y>,,,AHH (  %T & & 7qc;A I EJJ qc= ekk !  !i&A&A&AYEV S23 3 +E2I)%0I  !#ejj(,.is_where_neededoss. >??hh2226O"///$%JKKxx**?;;  ++1138 Q;!3C 4rc\PPRVP4p\P P Uu.uF&q"P'gKVP4NK( pp\V4^8wd \R4h\PV^,W4p\P PP\P PW@PVP R7#uupi)dotz'don't tile reduction when native matmulrr)r$ Reduction default_valuerrXrrr1r*rr TritonKernelOverridesr rgeneratecomputer)rdefaultrreduction_mask where_vars& r where_cond'TritonOverrides.dot..where_condsll00 BG-.XX-A-A-ATEVEV  -A >"a'$%NOO-33N14EsTI88<<((  )99CII) s DDc$V^8dQhRRRRRR/#)rrr~rrrr)rs"rr)TritonOverrides.dot..__annotate__s-G G /G .G  G rc<S V9dVUu.uF q3S 8XdRMTNK ppVSS.8XdXSV9d \R4hRR.pSV9dSV^&SV9dSV^&\WV4pVSS.8gpV'd RV RS RS R2pV#VSS.8Xd^SV9d \R4hRR.pSV9dSV^&SV9dSV^&\WV4pRV R 2pVSS.8gpV'd RV RS RS R2pV#\huupi) aX Generate a reshape, transpose, and broadcast for the tl.dot. tl.dot requires specific shape requirement : (Y,R) x (R,X) but the current triton codegen eagerly broadcast the tl.arange so it needs to be reshaped to meet the requirement. This is done by three steps. 1. remove the empty dimension (dim with size 1) and make it 2d with tl.reshape 2. permute the dimension if needed (e.g., (X,R) -> (R,X)) with tl.trans 3. broadcast if needed with broadcast_to. - This shows up when matmul operand is broadcasted with torch.expand/repeat. - e.g., torch.rand((16,)).expand(16,16) @ B e.g., (Y,1,R), (Y,R) -> tl.reshape(var, (Y,R)) e.g., (1,X,R), (R,X) -> tl.trans(tl.reshape(var, (X,R))) e.g., (1,X,1), (R,X) -> tl.broadcast_to(tl.trans(tl.reshape(var, (X,1))), (R,X)) TODO : eventually we want to remove this function when lazy broadcasting arrives rz&left tl.dot operand cannot depend on xrz, (rrxz'right tl.dot operand cannot depend on yr r)r rNotImplementedError) rrrrshape_2dbroadcast_neededRBLOCKr=r>r?s &&& r#reshape_transpose_broadcast_for_dot@TritonOverrides.dot..reshape_transpose_broadcast_for_dotsk6&JW X-3v 3!>- Xvv..]*()QRR:]*"(HQK]*"(HQK'uXF$,/?#? #.ugS6("ME2L/ 00]*()RSS:]*"(HQK]*"(HQK'uXF$E7!,$,/?#? #.ugS6("MEL*)S!YsC.ztl.dot can only do mm and bmmrtf32ieeeztl.dot(rz, input_precision=""))rXris_native_matmulr rdense_size_listrrrCr!r=r>r?r;rrrrrrrbackendscudamatmulfp32_precision) rrorig_aorig_brrr(input_precisionr'r=r>r?s && @@@@rrTritonOverrides.dotNs:xx((( !@A A , 4 6 " "1 A 6 " "1 AG G R qxx'') *Q . !@A A]..t{{;<]..t{{;<]..t{{;<]..t}}=> HHLL ! ! HH   /4=66BR S''6" "  HHLL ! ! HH   /4=66BR S''6" "  >>   % % 4 4 >$O$O2aS 3O3DBGGr constraintsris_purepack input_dtypesc  aaaaa\V4oVf*RPR.VUu.uFpRNK up,4pVPR4Uu.uFqP4NK p pV Uu.uFqP R4'dKVNK p p.p \ \ WjR\V444Fwp wrVR8XdZVeV\V \4'd@V PW\,8wd)V PV R\W\,4 R24KhV P\V 44K \PP'd8\!V4P#R R 4o\!V4P#R R 4oM R V R 2oR V R 2oVVVVV3R lpS^8:dV!RPV 44#V^,p\$P&P(p\$P&P*pVP-VVP.R 7pRPV3RlV 44pVP1V RV!V4 24VP1V RV RV RS R24V#uupiuupiuupi)Nrz=rr,=hrr\z\\rc .<RS RS RV RS RS RS R2 #)ztl.inline_asm_elementwise(rr\z ], dtype=z , is_pure=z, pack=rr)r asm_literalasm_triton_typeconstraints_literalr8r9s&rasm_call8TritonOverrides.inline_asm_elementwise..asm_call:s>,[M.Gs$ HS-cU"TF! < s = z$ = triton_helpers.inline_asm_unpack()rRrqsplitstripra enumerater"rrr_rr_rrrrreprreplacerXrrrnewvarr writeline)asmr7rr8r9r:inputs_rconstraint_partsinput_constraints cast_inputsr,rHrrE first_inputrrr packed_argsrBrCrDs$$$dd$* @@@rinline_asm_elementwise&TritonOverrides.inline_asm_elementwise sO&e,  ))TF6-B6ac6-B$BCK 0;/@/@/EF/E!GGI/EF(8R(81 S@QQQ(8R $S=S[1Q%RSKAxS ,sK00II0""cU${?@h:6("[MQSTXSYYZ [  m.C GRs I7 I</J Jc RV R2#)z tl_math.cos(rrrs&rcosTritonOverrides.cosPrrc RV R2#)z tl_math.sin(rrrs&rsinTritonOverrides.sinVrrc \R4h)z/ops.index_expr not implemented outside a kernelr$rrrs&&&r index_exprTritonOverrides.index_expr\!"STTrc \R4h)z/ops.value_expr not implemented outside a kernelrcrds&&&r value_exprTritonOverrides.value_expr`rgrc \R4h)z+ops.masked not implemented outside a kernelrc)r)bodyothers&&&rmaskedTritonOverrides.maskedds!"OPPrc RV R2#)zlibdevice.lgamma(rrrs&rlgammaTritonOverrides.lgammah#1#Q''rc RV R2#)zlibdevice.erf(rrrs&rerfTritonOverrides.erfn s!$$rc RV R2#)zlibdevice.cosh(rrrs&rcoshTritonOverrides.cosht!1%%rc RV R2#)zlibdevice.sinh(rrrs&rsinhTritonOverrides.sinhzr{rc RV R2#)zlibdevice.acos(rrrs&racosTritonOverrides.acosr{rc RV R2#)zlibdevice.acosh(rrrs&racoshTritonOverrides.acosh"!A&&rc RV R2#)zlibdevice.asin(rrrs&rasinTritonOverrides.asinr{rc RV R2#)zlibdevice.asinh(rrrs&rasinhTritonOverrides.asinhrrc RV RV R2#)zlibdevice.atan2(rrrrrs&&ratan2TritonOverrides.atan2"!Bqc++rc V^8dQhRRRR/#)rbitsrrrr)rs"rrrsTT3T3Trc RVR R2#)ztl.full([], 0x08Xz), tl.uint32).to(tl.float32, bitcast=True)r)rs&r_f32_const_from_bits$TritonOverrides._f32_const_from_bitssSz)RSSrc(V^8dQhRRRRRRRR/#) rrRrr Sequence[Any]rrrr r)rs"rrrs)   S   n  QT  rc  NRPR.R.\V4,,4pRP\\V44p\P P P\P PRV RV RV R2\PVR7#)rz=fr9ztl.inline_asm_elementwise('z', 'z', [z*], dtype=tl.float32, is_pure=True, pack=1)r) rqrrsrrXrrrrrr)rRrrr7args_strs&&& r_inline_asm_f32TritonOverrides._inline_asm_f32sii#T): :; 99Sd^,xx||$$ HH  )#d;-@zC E-- %  rc V^8dQhRRRR/#)rrr_rr)rs"rrrs4F4F{4F{4Frc  aVPo\P3RV3RlllpRV3Rllp\PpV!R4pV!R4pV!RV R24pV!V RV 2\P 4p\P R V.S4pV!R V R V R V R24p V!V R V 24p V!V!R 4W!R44p V!WV!R44p V!WV!R44p V!WV!R44p V!WV!R44p V!WV!R44p V!WV!R44p V!V R V 24p V!WV 4p V!RV 24pV!V!R4V!R4V4pV!R V R V R V R24p RpV!V RV 2\P4pV!V R2\P4pV!V RV 2\P4pV!V R24pV!V RV 2\P 4pV!R V R V R V R24#)c$V^8dQhRRRRRR/#)rrrrrrr_r)rs"rr1TritonOverrides._cuda_atanf..__annotate__s"  c + BS rc<\PPP\PPWSR7#r)rXrrrr)rrrs&&rgen(TritonOverrides._cuda_atanf..gens488<<((  $5) rc(V^8dQhRRRRRRRR/#)rrr rrrr)rs"rrrs(  3 3 3 3 rc6<\PRWV.S4#)zfma.rn.f32 $0, $1, $2, $3;)rr)rrrrs&&&rfma(TritonOverrides._cuda_atanf..fmas!"22,qQi ri?irrz > zrcp.approx.ftz.f32 $0, $1;rrr1i ;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 <= )rrrrrrruint32)rrrroneinfabs_xlargercpzz2poly correctionr neg_resultpi_over_two_minus_result sign_masksign result_bits signed_bits signed_resultnot_nanrs& @r _cuda_atanfTritonOverrides._cuda_atanfsD05        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(!DEErc \V4p\V4p\PP'Eg\ P P4PR8Xd\V\4'dV\P8Xg V'dRpV'dcTp\ PPP\ PPV R2\PVP R7p\"P%V4pVe'V\P8wdV R\'V4 R2#V#V'dRMRpRV V R2pV'd+Ve'V\P8wdV R\'V4 R2#V#)r0Nr|rrrrzlibdevice.atan()rWr]rrrrXrget_current_device_or_throwrrr_rrrrrrrrrR)rr needs_upcastrrupcasts& ratanTritonOverrides.atans;!#.q1  !!!335::fD1k**%--'<L$ HHLL))HH$$c)--'' * %003F'LEMM,I k,&?%@BBM&2""1#fXQ/ E-%5==2HXT+e"4!5Q7 7 rc RV R2#)zlibdevice.atanh(rrrs&ratanhTritonOverrides.atanhrrc RV RV R2#)zlibdevice.copysign(rrrrs&&rrTritonOverrides.copysign s%QCr!A..rc RV R2#)zlibdevice.erfc(rrrs&rerfcTritonOverrides.erfcrrc RV R2#)zlibdevice.erfinv(rrrs&rerfinvTritonOverrides.erfinvrsrc RV RV R2#)zlibdevice.hypot(rrrrs&&rhypotTritonOverrides.hypotrrc RV R2#)zlibdevice.log10(rrrs&rlog10TritonOverrides.log10!rrc RV R2#)zlibdevice.log2(rrrs&rlog2TritonOverrides.log2&rrc RV RV R2#)zlibdevice.ldexp(rz.to(tl.int32))r)rns&&rldexpTritonOverrides.ldexp+s"!Bqc88rc RV RV R2#)zlibdevice.nextafter(rrrrs&&r nextafterTritonOverrides.nextafter0s&aS1#Q//rc V RV 2#rprrs&&r logical_andTritonOverrides.logical_and6Cs|rc V R2#)z == 0rrs&r logical_notTritonOverrides.logical_not;sE{rc V RV 2#rrrs&&r logical_orTritonOverrides.logical_or?rrc RV RV R2#)r ^ rrrs&&r logical_xorTritonOverrides.logical_xorDs1#S1~rc V RV 2#rrrs&&r bitwise_andTritonOverrides.bitwise_andICs|rc RV 2#)~rrs&r bitwise_notTritonOverrides.bitwise_notMs 1#wrc V RV 2#rrrs&&r bitwise_orTritonOverrides.bitwise_orQrrc V RV 2#)rrrs&&r bitwise_xorTritonOverrides.bitwise_xorUrrc V RV 2#)z << rrs&&rbitwise_left_shift"TritonOverrides.bitwise_left_shiftYD}rc V RV 2#)z >> rrs&&rbitwise_right_shift#TritonOverrides.bitwise_right_shift]rrc RV R2p\P4'd-\PP 4wpRV RV RV R2#RV RV R2#)r).to(tl.uint32)ztriton_helpers.rand4x(rrztl.rand(r_can_use_4x_randomrXrr.seedrblocks&& rrandTritonOverrides.randascVHO,  - - / /xx//1HU+D6F82eWAF F$r&++rc \\P\4;'d\\P\P P P4'*;'dr\PP4PR8H;'d?\PP4^8H;'d\P'*#)r0)rrXrr$r _inductorselect_algorithmTritonTemplateKernelrrrr r#align_random_eagerrrrr"TritonOverrides._can_use_4x_randomis qxx . . .%//::OO . .335::fD  . . ++-2  . . --- rc  nRV R2pRV RV R2pRV RV RV R2pRV RV R2pRV R V R V R V R V R 2 #) rrz((z)*(rxz)//(z//4))z)%(z!triton_helpers.rand_eager_kernel(r+z, VEC=rr) r  base_offsetthreads_per_roundtidvectid_u32denomr< tid_truncs &&&&& r rand_eagerTritonOverrides.rand_eageruscU/*SE./r2 eWCuE 2 UG2. 24&;-q2i[X^_b^ccdeerc RV R2p\P4'd-\PP 4wpRV RV RV R2#RV RV R2#)rrztriton_helpers.randn4x(rrz tl.randn(rr s&& rrandnTritonOverrides.randnscVHO,  - - / /xx//1HU,TF"VHBugQG G4&6(!,,rc  ,RV R2pRV RV RV RV R2 #)rrztriton_helpers.randint64(rrr)r rlowhighs&&&&r randint64TritonOverrides.randint64s1VHO,*4&6("SED6KKrc \R4h)z.ops.load_seed not implemented outside a kernelrc)rrs&&r load_seedTritonOverrides.load_seeds!"RSSrc Z\PP'dRV R2#RV R2#)z tl.rsqrt(rzlibdevice.rsqrt()rrrrs&rrsqrtTritonOverrides.rsqrts2 ==   qc# #%aS* *rc RV R2#)zlibdevice.log1p(rrrs&rlog1pTritonOverrides.log1prrc RV R2#)zlibdevice.tan(rrrs&rtanTritonOverrides.tanrwrc \PPPP V4pV'd \ VR4'dVP pMRp\P'dO\PP'd/\4R8d V\P8wd VeRV R2#RV R2#)rNzlibdevice.fast_tanhf(rzlibdevice.tanh())rXrrrr rrr#rrrrrr)rrrs& rtanhTritonOverrides.tanhs((,,**..q1 ww00MMEE   !!!"$v-&!+1#Q/ /$QCq) )rc RV R2#)z tl.sigmoid(rrrs&rsigmoidTritonOverrides.sigmoidsQCq!!rc  \PP4PR8XdRV RV RV RV RV R2 #RV RV RV R V R2 #) xpurz%).to(tl.int64, bitcast=True) < 0 if (z.).dtype is tl.float64 else (libdevice.signbit(z ) != 0) if (z).dtype is tl.float32 else z < 0z(libdevice.signbit(z).dtype is tl.float32 or (z).dtype is tl.float64 else )rXrrrrs&rsignbitTritonOverrides.signbits 77 . . 0 5 5 >A3c++,#.cs$   "!%#/s33d  rc RV RV R2#)zlibdevice.fmod(rrrrs&&rfmodTritonOverrides.fmods!2aS**rc \4PW4pVey\V4'dhVPW4p\ V\ P P4'd!VPV\ P4MT pRV RV R2#\V4;'g \V4pTpV\ P\ P39d:\V4'g V'd\ PM\ PpV\ P8Xd_\\P P"4'g6\ PpV\ P8Xd\ PpVPW4pVPW'4p RV RV R2p VeDW78wd>\V4'dV'dV R\%V4 R2p V #V R\%V4 R2p V #)Nztriton_helpers.pow_integer(rrrr)rvpowrrrrrTrUrint64r]rrrMrFrXrrrR) rrrrrexponentr pow_dtypecast_acast_brs &&& rrCTritonOverrides.pows(*..q4  #(8(F(F**1;Da!4!4!;!;<< Q ,s  1b !D D215SS9PQR9S U]]EMM: : $L115E ]]   %.B GG " "/ /  Iu}},$}} ((6((6!&F815  # (A --# &xtK ,E+FaHF #84 L(A'B!D rc Z\PP'dRV R2#RV R2#)zlibdevice.log(rz tl_math.log()r#ruse_pytorch_libdevicers&rlogTritonOverrides.log s6  6 6 6$A3a( (aS""rF)rc RV R2#)zlibdevice.isinf( ).to(tl.int1)rrs&rrTritonOverrides.isinf "!M22rc RV R2#)zlibdevice.isnan(rOrrs&rrTritonOverrides.isnan rQrc RV R2#)zlibdevice.nearbyint(rrrs&rroundTritonOverrides.round s&aS**rc RV R2#)rrrrs&rfloorTritonOverrides.floor rrc \P!^\P4p\P!^\P4p\P!W4p\P !WCV4p\P !W4p\P !V\P!W 4V4p\P !V\P!W!4V4p\P !W4p\P !V\P!V4V4p\P!W4p\P !V\P!V4V4p\P !WBV4#r) rrrrFeqr ltrrtruncdiv)rrzerorb_zerob_nega_negrs&& rfloordivTritonOverrides.floordiv% s||Au{{+ll1ekk*  IIf1 %q IIeSWWT-q 1 IIeSWWT-q 1q IIeS__Q/ 3||A!yy 5t<yyt,,rc l\P!^\P4p\P!\P !W4\P 4p\P!\P !W4\P 4p\P!W#4pV RV R2#)rr.dtype))rrrrFrr\r@r)rrleftrightrs& rrTritonOverrides.signA sq LLEKK (||SVVA\EJJ7 cffQlUZZ8ggd"d1#W%%rc RV R2#)rmrrrs&rtruncTritonOverrides.truncJ rrc V RV 2#)rrrs&&rr]TritonOverrides.truncdivP s D}rc RV R2#)rrrrs&rceilTritonOverrides.ceilW r{rrNT)`rrrrrrre_LOG_2_Er8rrrrrrrrrrrrrrrrrrrrr rrrrZr]r`rerirnrqruryr}rrrrrrrrrrrrrrrrrrrrrrrrrrrr rrr r%r(r+r.r1r6r9r=r@rCrLrrrUrXrbrrjr]rorrrrrr0srOyy HK&K&Z&DD,<<  33## ,* ) )&&''"""<<<<''**xHxHtCCmm C  C  CCCJ####UUUUQQ((%%&&&&&&''&&'',,TT    4F4Fl:''//&&((,,''&&9900,,    ff--LLTT++ ''%%**&""  &++&&P##/303/303++''--6&&'' &&rrrca]tRtRtRtV3Rlt]]PR44t ]R4t ]R4t ]R4t ] R4t] R 4t] R 4t] R R l4tR tV;t#)rie aMap 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. c H<\SV`!V/VBVP4R#r)rdre_setup_libdevice_routingrRrrjrfs&*,rreTritonKernelOverrides.__init__m s# $)&) %%'rc a^RIHo\PPP Fp\ W4'g\RV RV 24h\W4pV3RlpVR8XdS\ SR4'g \R4h\P!W2VR7pWn \W\V44KRp\P!WRVR7pWn \W\V44K R #) z.decomposition_router s777emm+)!,,"#3X>qAGGGrr9z$OpDecompositions must define sigmoid)r|r}cbVP\P8Xd RV RV R2#V!V4#)z libdevice.rr)rrr)rr|r}s&&&r dtype_routerDTritonKernelOverrides._setup_libdevice_routing..dtype_router s277emm+'z1#Q77)!,,rN)torch._inductor.codegen.commonrzrrr&op_requires_libdevice_fp64rr r functoolspartialrsetattrr8)rfn_name original_implr~fnrrzs& @rrv.TritonKernelOverrides._setup_libdevice_routingt s D,,GGG3(($witC5%IJJ#C1M H )#/;;()OPP&&(QX& l2&67 - ""WB"K C,r"2 3?Hrc x\PP4p^.V,pVPWVR7#)rrW)rXrr r)rrrrrs&&& rrTritonKernelOverrides.constant s7 xx**,d ##E#>>rc \W4p\PPVRRR7p\ V\ 4'g\ R\V4 24hVP'dVPpM\PVP4p\PP4pV\P\P39dTpMTp\!V4p\PP"P%\PP&VP)RVP* R2V4V\-V4VR7pVP.VnV#)FN block_ptrtma_compatibility_checkerexpected IndexingOptions, got rr)rrar)rrXrindexingrrEr rrLrrrKget_index_dtype_as_torch_dtyperrFrDrTrrrrrFrGrH) rrrrrrq cast_dtype output_dtypers &&& rre TritonKernelOverrides.index_expr s)/<88$$ ET% (O44 #A$x.AQ!RS S  ))E!11(..AE hh==? U[[%++. .$JJ +:6 hhll## HH   LL1X//02J ?(. $ !**  rc  <\PPpV\P\P 39dTM\P \PnVP W4pV\PnW28wgVPV8wd`\PPP\PPRV R\V4 R2VVPR7pV# T\Pni;i)z Like :meth:`index_expr`, but honors ``dtype`` by setting the kernel index dtype before emitting, and casting the result if needed. rrnrr) rXr _index_dtyperrFrDrerrrrrRr)rrrreal_index_dtypers&&& rri TritonKernelOverrides.value_expr s 8800u{{EKK88Eekk  5..-C$4AHH !  $ U(:((,,''  C5k%013ii (C  %5AHH !s &DDc  Ve}\PPea\PP P \PPV R2\PVPR7pVPPRR7pV'g \R4hRpVFpVPFupVPR8wg_\VP^,4'g=VP^,\PP P 9gKrRpK K V'dRMTp\PP#WR 7;_uu_4pV!4p RRR4V'dX P$P&'d \V4p\PP P \PPR V R \)V4 R V R 2\*P,!V4V P.V PR7p\0P2!XW4p MX p V P4P7X4V # +'giL;i)N .to(tl.int1)routput)opz)graph for body does not contain an outputFloadT)rrz.shape, rre)rarr)rrrrXrrrrrrr find_nodesr rtargetrx store_cache mask_loadsrais_boolrkr"wraprrr rHdiscard) r)rlrmnodes need_whererrrnew_maskrrs &&& rrnTritonKernelOverrides.masked s   1 1 =88<<((  & %jjjj )D %%%2 !LM M DyyJJ&(/ << xx{ahhll&>&>>!%J!# XX  3 3xVF4 }}$$$U HHLL))  6((=+?*@6('R"''.llll *E))Hf4CC h' )4 3s 4I++ I; c \PPPV4pRV R\PPP RV4 R2#)tl.load( + load_seed_offsetr)rXrrinput seed_offset)rrrs&& rr(TritonKernelOverrides.load_seed sIhhmm!!$'se3qxx}}889KVTUUV W rc 8RV R2p\PPPV4;p'dV#\PPP VP VP R7p\PPP \PVP R7p\PPPV RV RV R24\PPPWV34W43#)zfrexp(rrrz = triton_helpers.frexp() rXrrtry_getrPrrrrFrrQput)r cache_keycse_valmantissarEs& rfrexpTritonKernelOverrides.frexp# sQCqM hhll**95 57 5N88<<&&QWWAGG&D88<<&&U[[&H ""j8*$)>%??$KCPTP]P]L^K__cdghlh}h}d~cA!%U3t/D/D+E%F G  $ 1B BD ?$?A+--?$D *U$s  DD$c VP#r)rrQs&rr#(BlockParameters.StrideSorter.is_identity s$$ $rc$V^8dQhRRRRRR/#rrzlist[int | sympy.Expr]rrrrr)rs"rr)BlockParameters.StrideSorter.__annotate__ s) U U#9 UFN U ) Urc R#)zBCreate a `StrideSorter` that can be used to sort block parameters.Nrrrrs&&&rr#BlockParameters.StrideSorter.create src xVP'g#VPUu.uF q!V,NK up#V#uupir)r#r%rRattrr,s&& rsort!BlockParameters.StrideSorter.sort s5###)-7AQ77K87c xVP'g#VPUu.uF q!V,NK up#V#uupir)r#r&rs&& rr!#BlockParameters.StrideSorter.revert s9###)-)=)=>)=AQ)=>>K?r)rr&N)rrrrr dataclassesfieldr&rrvr#rrrrr!rrrr StrideSorterBlockParameters.StrideSortero sd##%0%6%6E%BB $  %  %   U   U    rrc>a]tRtRtV3Rlt]RRl4tRtV;t#)$BlockParameters.IdentityStrideSorteri c $<\SV`4R#rrdrrRrfs&rr2BlockParameters.IdentityStrideSorter.__post_init__  G ! #rc$V^8dQhRRRRRR/#rr)rs"rr1BlockParameters.IdentityStrideSorter.__annotate__ s$  #9 FN  ) rc  LV!V\\\V444R7#)rr%)rrrrs&&&rr+BlockParameters.IdentityStrideSorter.create s'!1eC(8$9:; rr) rrrrrrrrrmrns@rIdentityStrideSorterr s $    rrcBa]tRtRtRtV3Rlt]RRl4tRtV;t #)+BlockParameters.TensorDecriptorStrideSorteri zD Sorts BlockParameters dimensions with strides in descending order. c $<\SV`4R#rrrs&rr9BlockParameters.TensorDecriptorStrideSorter.__post_init__ rrc$V^8dQhRRRRRR/#rr)rs"rr8BlockParameters.TensorDecriptorStrideSorter.__annotate__ s$  #9 FN  ) rc \\\V444p\P!W!RR7pV!VVR7# \ dTpLi;i)aV If the strides are not all known constants or if the strides are already sorted in descending order, return identity sort. For example if block_shape @ strides is [ZBLOCK, XBLOCK, YBLOCK] @ [8, 1, 16] The indices to sort the strides in descending order will be [2, 0, 1]. The indices to revert back to the original order will be [1, 2, 0]. T)reverser)rrrr&rr )rrr identity_sortr%s&&& rr2BlockParameters.TensorDecriptorStrideSorter.create s^!s+;'M ) !,,YRVW !1!  " )( )sA AAr) rrrrrrrrrrmrns@rTensorDecriptorStrideSorterr s!  $    rrc V^8dQhRRRR/#)rrmryrr)rs"rrBlockParameters.__annotate__ s::_::rc  \V4p\;QJd.RW34FNK 5M!RW344wr4T!R/VUu/uFqUW5,WE,,bK upB#uupi)z Concatenates block parameters. c3N"TFp\P!V4xK R#5ir)rasdictrs& rr*BlockParameters.__add__.. sBMq[''**Ms#%r)rr)rRrmrrrrs&& r__add__BlockParameters.__add__ s\4juBTMBuuBTMBB9a8as16AF?*a8998sA' c$V^8dQhRRRRRR/#)rrztype[StrideSorter]rrrz4tuple[BlockParameters, BlockParameters.StrideSorter]r)rs"rrr s$%%!3%@H% =%rc  VPVPVR7p\R/\P!V4P 4UUu/uFwrEWCP V4bK uppBpWc3#uuppi)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. )rr)rrryrrrr)rRrrrrrrzs&&& rr,BlockParameters.maybe_sort_with_stride_order s}*000S   !, 2 24 8 > > @ @HC'',, @  $$ s A1c V^8dQhRRRR/#)rremovable_dimsrrryr)rs"rrr s  *  rc  aV3Rlp\R/\P!V4P4UUu/uF wr4W2!V4bK uppB#uuppi)z1 Remove dimensions where removable_dims is True. cd<\VS4UUu.uFwrV'dKVNK upp#uuppir)r")ititem is_removablers& r filter_dims0BlockParameters.remove_dims..filter_dims s9+.b.*A*A&D#*A s ,,r)ryrrr)rRrr$rrs&f rrBlockParameters.remove_dims sS   5@5G5G5M5S5S5UV5UsK$$5UV  VsArN)rrrrrrrrrrrrr dataclassrrrrrrrrrrryryd s*//EE E$/$5$5d$KK!K + 1 1$ GG G + 1 1$ GG G***X |  #l##J:%"  rryc8]tRtRtRtRtRRltRtRtRt R #) "CooperativeReductionWorkspaceCachei z The scratch space used for cooperative reductions can be reused after two reduction loops. This keeps track of what can be reused. c Wn.Vn.Vn\P!\P 4Vn^Vn^VnR#rN) r current_loop prior_loop collections defaultdictdequeready_for_reuse loop_count store_count)rRrs&&rre+CooperativeReductionWorkspaceCache.__init__ s> *66{7H7HIrcV^8dQhRR/#)rnbytesrr)rs"rr/CooperativeReductionWorkspaceCache.__annotate__ s$$z$rc VPPV4pV'dVP4#VPP VR4wr4pVP P WV34W53#r)r1r popleftr workspacer,r_)rRr6cachedws_namerT ws_offsets&& rallocate+CooperativeReductionWorkspaceCache.allocate se%%))&1 >># # $ 3 3FE BI   &9!=>##rc VPF)wrpVPV,PW#34K+ VPVn.VnV;P^, unR#r)r-r1r_r,r2)rRr6r<r=s& r on_loop_end.CooperativeReductionWorkspaceCache.on_loop_end sT*.// &FY   ( / /0D E+:++ 1rc PVPpV;P^, unV#)r)r3)rRpriors& rincrement_store_count8CooperativeReductionWorkspaceCache.increment_store_count s#   A r)rr,r2r-r1r3N) rrrrrrer>rArErrrrr)r) s $rr)c.]tRtRt$R]R&RtRtRtR#)FixedTritonConfigi zdict[str, int]r#c (VPV,#rr#rRr"s&&rrFixedTritonConfig.__getitem__" s{{4  rc WP9#rrJrKs&&r __contains__FixedTritonConfig.__contains__% s{{""rrN)rrrrrrrNrrrrrHrH s !#rrH.c0]tRtRt$RtR]R&R]R&RtR#)_UnresolvedLoadIndexStatei. z=3.4.0, a CUDA device with cc>=9.0, use_tensor_descriptor=True, and assume_aligned_inputs=Truez1%s stores with XBLOCK=1 cannot transfer 16 bytes.z/%s dtype %s has no CUtensorMapDataType mapping.)rkrXrrrr#ruse_tensor_descriptorrrrLdebugrorr0get_device_capabilityassume_aligned_inputsr rrr!has_strict_sum_multirow_reductionrrC)rR device_types& r can_use_tma#TMACompatibilityChecker.can_use_tmah sj :::gg99;@@ %  333*,,-// _,,  6) 88:1=B444%' 333)++ IIN((   >>> KK {{##EEGG IIC((  ::2 2 IIA((   rc$V^8dQhRRRRRR/#)r block_paramsryr{zsympy.Expr | intrrr)rs"rrrr s-RR%R*R  Rrc  \PP4PpVP'dVP Uu.uF,p\PP PV4NK. pp\PP P\P!V44pM"VP p\P!V4p\PP PVR,\P!^44'g$\PRVPV4R#VR8XdR#VPP pVRRFp\PP P\#W,^\P!^44\P!^44'dKn\PRVPVV4R# \PP P\#Wg,^\P!^44\P!^44'g%\PRVPVV4R#VP$eI\'VP$4'd.\PRVPVP$4R#VP(R,p \PP PV \P!^44'd.\PR VPVP(4R#Rp Rp V P*F2p \,P.Fp \1W4'gKT p T p K0 K4 V 'd V 'g\3V R \,P. 24hVP4P6P94^8Xd8V \,P:9d#\PR VP4R#VP4P<'EdRV \,P:9Ed<\>V ,pRpVP4P@F6pVPB'gKVPDV8XgK*VPFpM Vf/\PR VPVVP(4R#VP4PIV4pV PKV V/4V,p\PP PMV\P!^44'g/\PR VPVP(V4R#R#RRRllpW,^, pVPO\PV4PO\"V4p\S\U\PV!VV ^444pVVP4PY\>V ,48d%\PRVPVV4R#VP4PZ'd@V \\P^8Xd+V^8d$\PRVPV4R#VP4P`'d\VP4Pb'g@V \\P^8Xd+V^8d$\PRVPV4R#VP4PeV 4p VP4Pf'dfVVP4PfV ,8dA\PRVPV VP4PfV ,V4R#R#\iVVP4PjPmV ^44VP4PjV &R#uupi \nd0\PRTPTP(4R#i;i)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. zG%s tensor descriptors require innermost stride to be 1. Strides are: %sFrtTNzU%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 zF%s strict sum linear accumulation requires a reduction block size of 1zN%s could not find reduction range tree for innermost prefix %s Block shape: %szn%s persistent reduction innermost block shape cannot transfer 16 bytes. Block shape: %s, persistent RBLOCK: %dc(V^8dQhRRRRRRRR/#)rrrrrzsympy.Expr | Nonerr)rs"rrMTMACompatibilityChecker.are_block_parameters_compatible..__annotate__Y s0!!) rc6W, pV'd W2,pV#rr)rrrrs&&& rindexing_div_repQTMACompatibilityChecker.are_block_parameters_compatible..indexing_div_repY s %C!gJrzC%s the minimum block size to satisfy expression %s is too large: %dzK%s no_x_dim kernel has XBLOCK fixed at 1 but TMA requires min block size %dzI%s combo kernel with shared XBLOCK cannot guarantee TMA min block size %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: %srr)8rXrrrrkrr!replace_backed_symbols_with_hintsrsympifyrrrLrvrorrErrlrvrrrr+r r rrstrict_sum_rblockr2persistent_reductionrrr1rnumel_get_persistent_reduction_blocksubsstatically_known_geqrOrr9rnsolve max_blockrr!r=is_combo_kernelper_subkernel_blocksr$ fixed_configrtma_min_block_sizesr r2)rRr~r{rzstrconstant_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_compatible7TMACompatibilityChecker.are_block_parameters_compatible sgg99;@@ :::'...B  BB2F. $%77#3#3#U#U o.$ #**G#(==#A ww77 U]]STEUVV IIY((   % zz** crlF77##;; 5q%--:KL a  k,,   #$ww77 0?EMMRTDU V MM!    IIH(($      ',DTEUEU,V,V II.((    !- 8 8 < 77   3 3 !5==#3   IIP((((  ##3@@N+77 !.==+9(+5( 8A%)= ())UVcVoVoUpq  KK 2 2 4 9$ (E(EE IIX((  KK , , ,$ (E(EE%//C$D !J[[,,>>>ahh2G&G!"J-!  d,,) ,,   $ K KJ W %**,@BS+TU "77##88%u}}R'8 E,, ,,%  fIb 3ABF (2(:(:.)'/+;<&"1 10""DKK$9$934%II]00-&  ! KK(((, ;&*IIe00& !KK/// KK<<<, ;&*IIc00& !!%!9!9:N!O;;+++% (@(@(PP < 44* KK44^D*  %Q4GJ& 77;;NANGDKK33NCM|  U,, ,,   sK 2` .B8`%(`%=`%`%`%;=`%:6`%1A `%A`%%6aacV^8dQhRR/#rNr)rs"rrrr s  $ rc VP#)a Can you lift the make_tensor_descriptor call to the top of the kernel? This requires being certain that all of the shape, stride, and block_shape information is handled in arguments or top level definitions. Right now we assume this is always possible if you force TMA. )rkrQs&rr TMACompatibilityChecker.can_lift szzrrnr) rrrrrrrlrr{rrrrrrrjrjW sF  O K"K"X?BRh  rrjc |a]tRtRt$Rt]tR]R&]t R]R&Rt ] t ] tR]R &]tR ]R &R tR ]R&RRV3Rlllt]RRl4t]RRl4tRRltRRltRRltRRltRRltR R!ltR"R#ltR$R%lt]R&R'l4tR(R)lt R*R+lt!R,R-lt"RR.R/llt#]R0R1l4t$]RR2R3ll4t%R4R5lt&R6R7lt']R8R9l4t(R:R;lt)R<R=lt*R>t+R?t,R@t-RAt.RBRClt/]0RDREl4t1RFt2]RGRHl4t3RIR RJRRKR RLRRMR RNR/RORPllt4RRQRRllt5RRSlt6RTt7RUt8RVRWlt9RXt:RYRZlt;R[R\ltR`t?]Ra4t@RbR/RcltARdtBReRfltCRgRhltDRiRjltERkRlltFRmRnltGRoRpltHRqRrltIRsRtltJRRuRvlltKRwRxltLRytMRzR{ltNRR|R}lltOR~RltPRRltQRRltRRRltSRRltTRRltURtVRtWRtXRtYRRltZRRlt[RRlt\RRlt]RRlt^]RRl4t_RRlt`RRltaRRltbRRltcRtdRRlteRRltfRtgRth]i]j!R 4RRl44tk]iRRl4tl]iR4tmRRltn]oPRRl4tqRRRlltr]R4tsRRlttRRltu]R4tv]RRl4twRtxRRltyRtzRRRllt{RRlt|RRlt}RRlt~RRltRRltRRltRRltRRltRRltRRltRRlt]0RRl4tRRltRRlt]0RRl4tRRltRRltRRltRtV;t#)r$i z\A class to represent a triton kernel and helpers to generate triton kernel programmatically rhelper_functionszCallable[[sympy.Expr], str]kexprTztype[BlockPtrOptions]block_ptr_options_clsztype[TensorDescriptorOptions]tensor_descriptor_options_clsNz bool | None3transpose_discontiguous_tensor_descriptors_overrideFc 0V^8dQhRRRRRRRRR RR R /#) rtilingdict[str, sympy.Expr]rzFixedTritonConfig | None hint_override int | Nonerrrrrr)rs"rrTritonKernel.__annotate__ sF?9?9%?9 / ?9 " ?9?9#?9 ?9rc <W0nW@nW`nWpn\S V`!V3/VB\ VPVP4Vn /Vn \4Vn \4Vn \4Vn\\ ,!4VnW n\&P(!4Vn\,\.\.3,!4Vn\34Vn\6P8!\,4Vn\,\.\<3,!4Vn/Vn \4Vn!RVn"RVn#WPn$\6PJ!4Vn&^Vn'RVn(.Vn).Vn*/Vn+/Vn,/Vn-\6PJ!4Vn.\\^,!4Vn0RVn1VPd'dVPgVPh4VPj'dVPm4VPo4VPj'dVPq4RVn9.Vn:R#NF); optimize_maskrrrrdrerX newvar_prefixsuffixrprologue_cacheraprologuepost_loop_combinepost_loop_storerr outside_loop_varsrrcount block_ptr_iddictrblock_ptr_to_bufferrrr.r/pointer_advancementsrrhost_tma_descriptor_args_host_tma_non_materializable$_host_tma_non_materializable_buffers_emitted_device_tmarCounter _load_counts_pdl_load_index _pdl_has_waitop_traceop_trace_buffer_arg_names_op_trace_buffer_names_op_trace_cse_names_op_trace_symbol_names_op_trace_symbol_countsr0autotune_hints triton_metarcodegen_reduction_numelsrlcooperative_reductioninit_cooperative_reductioncodegen_range_treeinit_cooperative_reduction_maskhas_load_with_contiguous_rdimstores_with_contiguous_rdim) rRrrrrrrrrjrfs &&&&&&&&,rreTritonKernel.__init__ s$1(%4*>! *6*T//=.0(6(8 1?1A/=/?!+C!2#6 %OO-#'S>#3 / 1  # #D ) !$(S>#3 LN%=G\)LP1#( *6A6I6I6K "24 46&68#35 HJ#ALATATAV$)68.2   ) )$)) 4  % % %  + + - !  % % %  0 0 2-2*68(rcV^8dQhRR/#rNr)rs"rrr sOO$Orc T\VP;'g VP4#r)rrrrQs&ruses_tmaTritonKernel.uses_tma s!D11MMT5M5MNNrcV^8dQhRR/#rNr)rs"rrr s(((rc VP#r)rrQs&ruses_device_tmaTritonKernel.uses_device_tma s'''rcV^8dQhRR/#rrrr)rs"rrr# sRRCRrc :\RVP44#)c3P"TFp\VPRJ4xK R#5ir)rr rrs& rr2TritonKernel.triton_tensor_ndim..$ s#Q@P3td233@Ps$&)sumrrQs&rr TritonKernel.triton_tensor_ndim# sQ@P@PQQQrc V^8dQhRRRR/#)rr,rrrr)rs"rrr& s''3'3'rc dR.VP4,pRW!&RRPV4 R2#)rr]rUrrV)r rq)rRr,sizess&& rindexing_size_strTritonKernel.indexing_size_str& s742244499U#$A&&rcV^8dQhRR/#rrrOr)rs"rrr+ src R.VP4,pVP4F/pVPfKVP4WP&K1 V#)r)r rkr block_size_str)rRrrs& rr.TritonKernel.dense_size_list+ sP//11++-D*)-)<)<)>oo&. rcV^8dQhRR/#rr)rs"rrr2 s888rc HRRPVP44 R2#rT)rqr.rQs&rdense_size_strTritonKernel.dense_size_str2 s$499T11345Q77rc V^8dQhRRRR/#rentryrorrr)rs"rrr5 s   *=  #  rc "VPf%VP4 RVP4 R2#R.VP4,pRW!P&RP V4pVP4 RVP 4 RV R2#) N = tl.full(z, True, tl.int1)rr]rz = tl.full([z], True, tl.int1)[rV)r r*rr rqr)rRrrrs&& rcreate_constant_mask!TritonKernel.create_constant_mask5 s    #??$%[1D1D1F0GGWX 42244"%5! !",,.//A& L rc$V^8dQhRRRRRR/#rrrstorerrr)rs"rrrB s!%%c%T%c%rc V\PP9dV#V'dVPP V4#VPP V4#r)rXrremoved_buffersrrrrRrrs&&$r_trace_kernel_argTritonKernel._trace_kernel_argB sD 177** *K 99##D) )yyt$$rc$V^8dQhRRRRRR/#rr)rs"rrrI s!11#11#1rc WP9dQR\VP4 2VPV&VPPVP WR74VPV,#)bufferr)rrrr_rrs&&$r _trace_bufferTritonKernel._trace_bufferI sn 22 2T889:;  ' ' -  * * 1 1&&t&9 **400rc V^8dQhRRRR/#)rrr_rztuple[str, str, str, Any]r)rs"rrrS s  + 2K rc 0VPVP9d0R\VP4 2VPVP&RVPVP,\VP4VP VP 43#)r`r)rrrrr_normalize_trace_valuer)rRrs&&r_trace_cse_varTritonKernel._trace_cse_varS sw 88433 336s4;S;S7T6U1VD $ $SXX .   $ $SXX .  N  ' ' 2   rc V^8dQhRRRR/#)rr%r&rr'r)rs"rrr^ s  \ j rc \P\P\P\P\P 3F#p\ W4'gK\V,u# R#r)r!r=r>r?r;r<r r)r%r,s& r_trace_symbol_prefix!TritonKernel._trace_symbol_prefix] sP KK KK KK MM MM  Df++!$'' rc V^8dQhRRRR/#)rr%r&rr)rs"rrrj sL\rc "WP9dVPV,#\V\P4'dPVPP VP 4pVe'\P!V4pW0PV&V#VPV4pVfV#VPV,pVPV;;,^, uu&\P!V V 23/VPBpW0PV&V#r) rr r!rrr rrr@rr assumptions0)rRr%r normalizedrrs&& r _trace_symbolTritonKernel._trace_symbolj s 00 0..v6 6 &$(( + +++// M,,V4 $$V,1,\\VHUG"4L8K8KL .8##F+rc V^8dQhRRRR/#rhr)rs"rrr s*rc \VP\P!R4R7Uu/uFpW P V4bK ppV'd \ W4p\ V4#uupi)rr)rrroperator attrgetterrrPtexpr)rRrr% replacementss&& r_normalize_trace_expr"TritonKernel._normalize_trace_expr sh!!2!28K8KF8ST T &&v. .T   d1DT{  sA(c V^8dQhRRRR/#)rrr rr)rs"rrr sCCrc a\V\4'dSPV4#\V\P4'dRSP V43#\V\ P4'dR\V43#\V\\34'd3\;QJd.V3RlV4FNK 5#!V3RlV44#\V\4'da\;QJd..V3Rl\VP444FNK 5#!V3Rl\VP4444#V#)rrc3F<"TFpSPV4xK R#5irr)rrrRs& rr6TritonKernel._normalize_trace_value.. sGA44Q77!c3N<"TFwrVSPV43xK R#5irr()rrrrRs& rrr) s* 5HCd11#67 5"%)rr_rrrr#rrrrrrrrr)rRrsf&rr#TritonKernel._normalize_trace_value s e[ ) )&&u- - eUZZ ( (T77>? ? eU[[ ) )SZ( ( edE] + +5GG5 G5GGG G eT " "5 &u{{} 55 5 &u{{} 5  rc ,V^8dQhRRRRRRRRR R /#) rrrrrrrjrrr rrr)rs"rrr s<% % % %  %  %  % rc  aSP'gR#VR8Xd*VwrVSPVRR7SPV43pMVR8XdKVwrVrSPVRR7SPV4SPV4SPV 43pM2\;QJd.V3RlV4FNK 5M!V3RlV44pSPP \ TT\;QJd..V3Rl\VP444FNK 5M'!V3Rl\VP4444SPV4R74R#) NrFr rTc3F<"TFpSPV4xK R#5irr()rrrRs& rr/TritonKernel.record_op_trace.. s!#UPTD$?$?$D$DPTr*c3N<"TFwrVSPV43xK R#5irr()rrrrRs& rrr1 s*&< $55e<=&!% K"";e"<++E2OW_.2 +K"";d";++E2++E2++D1 O$e#UPT#Uee#UPT#UUO  $u&,V\\^&<uu&,V\\^&<226:  rcV^8dQhRR/#rNr)rs"rrr s::t:rc \P'g#\PP'gR#VPpVUu.uF'p\ V\ P4'gK%VNK) pp\V4^8XdR#\PPPWV4p\;QJdRT4F 'gK R# R#!RT44#uupi \dR#i;i)Fc3*"TF q^8HxK R#5irr)rrs& rr4TritonKernel._has_stride1_on_rdim.. s9[6Q;[T)r# deterministic test_configsforce_filter_reduction_configsrr rr2rrXrr stride_varsZeroDivisionErrorr)rK support_varsr reduce_varsr?s& r_has_stride1_on_rdim!TritonKernel._has_stride1_on_rdim s  F$7$7$V$V$V)) $ #c=#@#@A C#  { q  ''**66u<XKs9[9ss9s9s9[999! !  s "C)0C) *C.. C=<C=c$V^8dQhRRRRRR/#)rrr0rrRrrr)rs"rrr s&)! rc VeV\P8XdR#VPFNp\V\\ P 34'dK+\V\ P4'dKMR# R#)zWhether the host launcher can build a TensorDescriptor for this access (not a TMA eligibility check -- that's TMACompatibilityChecker). FT)rrrrrrrr@)rrrs&& r_is_host_tma_materializable(TritonKernel._is_host_tma_materializable s]  %**!4''C#U]]344#u||,, ( rcV^8dQhRR/#rr)rs"rrr s4H4HT4Hrc j\PP'g\4VnR#\4Vn^RIHp\RVP44pV'gR#^RIH pH p\P!\4pVPVPP 4FpVP"P$Fp\'VR4'g:\'VR4'd&VPP)VP*4KNWWP*,P-VP.VP034K K VP34EFDwrWP9dKV Fwr\V 4p \4;QJd&RV P64F 'gK RM R M!RV P644'dVPP)V4Mn\8P:!V 4FNp \=WV34'gKV P6V ,'gK3VPP)V4M KM WP9dEK\?V 4^8gEK)VPP)V4EKG R#) zxPopulate _host_tma_non_materializable_buffers with buffers that can't be expressed as a single host-side TMA descriptor.N)NodeScheduleMarkerc3@"TFqP4xK R#5irr%rs& rrCTritonKernel._prescan_host_tma_materializability.. s'SBR$ BR)rr var_namesrc3V"TFp\V\P4xK! R#5ir)r r!r)rrbs& rrrNs OL1<''S$BRBR'SS! J  # #D ) '11$--2M2MND''--sK00sF++AAEEchhO#HH-44cii5OP .O 388:OHDDD$) %i03OE A $~~5 77   8 8} U U ))--c2 %%))#t4rcV^8dQhRR/#rNr)rs"rrr@s   rc \;QJd0RVP4F'dK R'*# R'*#!RVP44'*#)c38"TFp\V4xK R#5ir)rc)rrs& rr>TritonKernel.has_store_with_contiguous_rdim..As 0P d # #0PrFT)rrrQs&rhas_store_with_contiguous_rdim+TritonKernel.has_store_with_contiguous_rdim?sX3 040P0P 33  3  3 040P0P    rc V^8dQhRRRR/#r:r)rs"rrrEs""+"#"rc \V4#r)rR)rRrs&&r dtype_to_strTritonKernel.dtype_to_strEs 5!!rcV^8dQhRR/#rNr)rs"rrrHs  $ rc VP4eR#VP;'df\PP \P P 4VPPVPP4#r) _strict_sum_rblockrrXchoices should_use_cooperative_reductionrrrrreduction_numelrQs&rr{-TritonKernel.should_use_cooperative_reductionHsg  " " $ 0$$  )S)S GG / / 1 MM   MM ) )*  rc aSP'g \R4hSPF,pVPfKV;P^, unK. SPR,pSP 'd\ VSP R,4pSPPV4Sn \SP4Sn SPPR4\;QJd)V3RlSP4F 'gK RM RM!V3RlSP44'dSPPR 4R#R#) z/One time setup code for cooperative reductions.zexpected cooperative_reductionNrr=a  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) c3x<"TF/pVP'gKSPV4'*xK1 R#5ir)r1r rrrRs& rr:TritonKernel.init_cooperative_reduction..ns4 (   .''- - -(s::TFz>rsplit_end = tl.where(rsplit_end < rnumel, rsplit_end, rnumel))rr rgrid_dimrrrr semaphoressemaphores_namer)%cooperative_reduction_workspace_cacherlsplicerrQ)rRr sem_countsf rr'TritonKernel.init_cooperative_reductionQs))) !AB B$$D}}( " %KK$     4+<+5W II6 2    3 (( 333 ((   II  P   rc JRpVP'gV R2pVPPRV 24VP4'dVPP R4R#VP'd \ R4hVPPR4R#)z$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 zexpected no_x_dim to be FalsezSrsplit_mask = xmask if RSPLIT_IS_POWER_OF_2 else ((rsplit_arange < RSPLIT) & xmask)N)rrlrQ_has_constant_xmaskrr )rR rsplit_aranges& rr,TritonKernel.init_cooperative_reduction_maskws> }}},oY7M .}o>?  # # % % II   }}}$%DEE II  e rc VPFpVP'gVPWP4K2VP'gKFVPP VP RVPV4 24K VP'd\;QJd&RVP4F 'gK RM RM!RVP44'dUVPRRRR7pVPV4pVPPRVPV4 24R#VPVP4R#R#)zbase = c38"TFqPxK R#5ir)is_looprs& rr2TritonKernel.codegen_range_tree..s=,;'g!T;'gSP@RJ;'dS^8gp RpRp\4pSPC4FHpVPEVPF4'dRpMRpVP9VPI44KJ V'd3SPJ'd!\:P<PL'gS'dSPO4'dV'gSP@'g|\QS!V, 4^8XdeSPSS4'gNV'dFSPTR8Xd5RR lo"R V3R llo$R V"V$3R llo#RVVV!V#VV3RllpV!4pVeV#RpRpSPWS4pVV3Rlp\YS4'Ed#S'g\QSP[44^8Xd V!4wppMT\]^.\QSP[44,4p\_^.\QSP[44,4pRV RV R2pSP`'gSPb'g V'd!\V3RlSPd44o!M \4o!SP@'dS!P9SP@4\gVS!VV SVR7#V 'EdRV'EgISPh'Ed SPj'Ed S!Pm4pSP@'dVP9SP@4\RSPC444pVPoV4'g}VPqV4pVPs4p\ V\t4'g\R\V4 24hVPwV4VP'VP(4KR.\QSP[44,pVFp\ V\\4'g\R\V4 24hSPC4FopVPyV4'gKVPzp\ V\|4'g\R\V4 24hSP[4V,VV&Kq K RRP\\\V44,R,p\_V4pRV RV R2pM;V!4wppRV RV R2pVo!M%V'gS'dV!4wppRV RV R2pVo!VfV 'g V'd V!4wp pMR pV'd \V.4o!SP@'dS!P9SP@4SPS!4\gVS!VV SVR7#)!z? Compute the index and mask to pass to tl.load() or tl.store() Frrzexpected sympy.Symbol, got NTtl.int32c$V^8dQhRRRRRR/#rrKr range_treerorBlockParameters | Noner)rs"rr+TritonKernel.indexing..__annotate__#s$  ! /B ' rc\P!WP44pVfR#\VP.\ P V4.V.\ PV4.R7#)zg Matches expressions of the form: idx = s * xindex This implies stride (s,), and shape (XBLOCK,). Nrrrr)rZmatch_affine_block_exprr%ryrrrr!)rKrrs&& rmatch_affine_block1TritonKernel.indexing..match_affine_block#sj-DD,,.>&%++,!.!=!=j!I J#H*;;JGH rc$V^8dQhRRRRRR/#rr)rs"rrr9s)x x !x /Bx 'x rc Z<aaVP4p\P!R\P!\P V.R7R7wr4\ ^\VP4VP\W#44VP\W#V44,4p\PPoSPV4pSPVP 4p\"P$!WWe4pVfR#Vwpp p \"P&!V4p \PPoSP)VP*4o\,;QJd VV3RlV 4F 'gK RM RM!VV3RlV 44'dR#\.P1V4p \3V SP5V ^,44.\7V R,VR,4UU u.uF9wrm\9\3V SP5V44SP5V 44NK; up p,pV Uu.uF#p\;W\.P=V4/4NK% pp\?VUu.uFpSP5V4NK upVV VR 7#uup piuupiuupi) 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)excluderNc3<"TF:pSPVS4'*;'dSPV4'*xK< R#5ir)rstatically_known_power_of_2)rrrrs& rrETritonKernel.indexing..match_mod_div_block..sM".!==eYOOHH$@@GGH!-s "AATFrNNr) r%rsymbolsrrWildrrrrrrrXrrremove_precomputed_replacementsrrZmatch_mod_div_block_exprget_slice_numelsrrrrrrrr"rrPr!ry)rKr index_varrmodulonum_dimsr match_resultdimsrblock_index_exprs slice_numelslinear_block_sizerrrrAdrrrRs&& @@rmatch_mod_div_block2TritonKernel.indexing..match_mod_div_block9s&'--/ !& "!))%**ykJ!   (() HY$>?++oi&OPQ &77++ @@G @@AQAQR2KKe   ' ! %2CCDI 77++ NN:+<+<= 3".333".  %2$@$@$L!) 88aI1'*,r*:DH&E'F  -x/O/OPU/V!88=  'F 1 (!2 3!2-*H*H*TU!2 3'HLM18;;A>M +#) %3Ns7?J)J#8J(c$V^8dQhRRRRRR/#)rrrrrorrr)rs"rrrs&    / ( rc<SS3F7p\P!WP44pV!W14pVfK5Vu# R#)zE Match a block indexing subexpression involving a single range tree. N)rZfactor_index_exprr%)rr match_funcfactored_index_exprmatchrrs&& rmatch_block_subexpr2TritonKernel.indexing..match_block_subexprsS''#J+>*O*O//1+''':GE($ #rcV^8dQhRR/#)rrzBlockDescriptorOptions | Noner)rs"rrrs` ` &C` rc N<\SSPP4UUu/uFwrWPbK upp4pSP 4pVUu.uF'p\ P !W$P44NK) pp\RV44p\4p\W54FHwrH\VPVP44^8dR#S!W4p V fR#Wy, pKJ V\V4, p SPS4\ P"P$'d SP&M SP(p \ P"P$'dRp \P*p M\-\.S4oSP14p SP2eSP2pM\ P"P4p\7SR4'd VSRJ,pV'd\P8M\P*p V P;VV VSSP<V V R7p\?V\@4'd@\-\.S4oSPCVPDVPFR7'gR#V#uuppiuupi)c3@"TFqP4xK R#5irrMrs& rrBTritonKernel.indexing..match_block_expr..s*Q[T;;==[rONFtemplate_out_shape)rzr{rrHrrr)r{)$rPrrrrkrZget_subexpr_involving_symbolr%rryr"r intersectionrr filter_masksr#r use_block_ptrrrrr rjrr)transpose_discontiguous_tensor_descriptorrrrrrr0rrzr{)rrindex_relative_to_xyr_indexrrindex_subexprs range_symbolsr~subexprrzr options_classrrtranspose_contiguousoptionsrrKrHrrRrs rmatch_block_expr/TritonKernel.indexing..match_block_exprs.8$2G2G2M2M2OP2O$!AvvI2OP/+#557 !, "!,(DD3[[]!, "!+*Q[*Q Q .0 %(%EMD=55g6J6JKLqP#1?F~# *L&F5s>7JJ!!),}}222..;;==...$H(7(L(L%04/1J1- 9AACHPP#$!TT- #MMSS-t%9::, $0FF,0(CC,AA& (..'$* +'"&..%&7/g'>??04/1J1-5TT(/(?(?U $}Q"s J -J"c<S'dF\S\4'dS R2R3#RRPRS44,R,S3#SP4\ SP 443#)z.shapeNrUrc38"TFp\V4xK R#5irr)rrs& rrATritonKernel.indexing.._get_expand_str..6s*F:a3q66:rrV)rrrqrrr.)rrRsr_get_expand_str.TritonKernel.indexing.._get_expand_str1skj#..(\0$66*F:*F!FFLjXX**,eD4H4H4J.KKKrrrz , tl.int32)c3<"TFBpVP'dKSPV4'dK0VP4xKD R#5ir)r1r r*rs& rr(TritonKernel.indexing..Ns@' 0,,%595L5LT5R%DNN$$ 0sA A A )rLc3@"TFqP4xK R#5irr*rs& rrr|s(1JNN$$1JrO expected TritonCSEVariable, got rzexpected str mask, got zexpected int dim, got rUr=rVrrr)Bprepare_indexingrrrrrr rrr@r rrr3r r!rrrrrirHrrrrrrr-rr#rrrfrkrvar_listr*allow_block_ptrrr{ris_indirect_indexingrqr$rr.rrrrrrErr-copyissubset differencerr_rrir rrqrsr)%rRrKrrrrrr index_varsr[rrr* need_dense have_densehave_loop_varsdense_mask_varsrrrrIrLrFr mask_shape axis_masks tmp_masksr` expand_listr)rexpand_shape_strrTrHrrrs%ffd$$$d$ @@@@rrTritonKernel.indexingsa%%e,''  %/\ *(*=*=f*EFFCc5<<00$'B49+%NOO#}'N'Nc(JTXX..((..sxx8  !2!23%%II))JJJJ''   "/!C!CD#!NNIMM),9G> MM ( ( + + + +d*qj   +5<++-D&&t}}55!%"    0 1 .t333 8S8S8S-1==??"OOOI/0A5--e44  J. ,x 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$8)>)>.'^^- ???NN4??3'(151H1H1J( %--j99 * 5 5j AI#--/C%c+<==,>tCykJ&&s+%%cmm4 #ec$*>*>*@&AA &D%dC00,/FtDzl-STT $ 7 7 9>>$//"&//C#-c3#7#7&4& * l*9+R8H7IKI'I  Z"1"3<! "M?3I ??? MM$// * )$     %   rc(V^8dQhRRRRRRRR/#)rrrrrz)BlockPtrOptions | TensorDescriptorOptionsrztuple[str, str]r)rs"rrrs8I'I'I'I'< I'  I'rc DVP4p\V\4'EduV'd!V'dVR8wd\RV: 24hRp\ 4'd\ P P'dVP'gxVP^8XdgW P9dWVPV\PPV44'd"W P9dW0PV&W$3#\ 4'd|\ P P'd\VP'gJVP^8wd9VPP!V4W P9dVPVRVnM7V'gRpM,V'dVR8wd\RV: 24hRV: R2pMRV: 2pVP$'d:VP&R,P('dVP+4'gVP'EdVP'd'W P,9dVP,V,pWd3#VP/VRR7pVP0P3V4pV'd\5V4V3#\7VP84p \V\:4'dR V 2pMR V 2pVP0P=V\>P@.R 7p VP0PCWz4\EW R V 24p VP'd+VPFPIV 4W`P,V&MVPJPIV 4\V\:4'dfWPLV&\NPPFCp VPSV 4p VPTV ,pWn9d\R V RV R24hWV&KE Wd3#VP/V4pWd3#)a%Generate a block pointer or tensor descriptor for Triton kernel operations. This method creates either a block pointer (for regular Triton operations) or a tensor descriptor (for TMA operations) based on the indexing type. It handles caching and reuse of descriptors for performance optimization. Args: name: The name of the buffer/tensor being accessed var: The variable name for the pointer indexing: Block pointer options or tensor descriptor options containing indexing information and boundary check settings other: Additional parameters string (e.g., padding options) Returns: A tuple containing: - block_descriptor: The generated block pointer or tensor descriptor variable name - other: Modified additional parameters string with boundary check options , other=0.0zexpected ', other=0.0', got rT, boundary_check=z, padding_option='zero'F)rrtma_descriptorrrJz#duplicate advancement for pointer 'z ' at type 'rr)+rrr0r rr#renable_host_side_tmarr{rrGrXr get_dtyperrrrrrr[rrrrrnextrr=namedvarruint64rr`rrQrlrrr2rLr)rRrrrrmcheckblock_descriptorblock_ptr_line block_varblock_descriptor_id named_var line_bodyr,advance_offsets advancementss&&&&& rcodegen_block_ptrTritonKernel.codegen_block_ptrs2'') h 7 8 8M)(+Gy)QRR *++MM666 ))),,1@@@44Xqww?P?PQU?VWW;;;9A11#6z!*++MM666 ))),,11155c:77755c:(,D $M)(+Gy)QRR+E94KL+E95  ! ! !  $,,,##%%       S,?,?%?$(#6#6s#; l &&i"*e!D HH,,^< y>500&*4+<+<&=#h88)23F2G'H$)78K7L'M$ HH--$ELL.  ^7(1C3~FV/WX $$$MM++I6/?'',II'' 2h88BF,,-=>!. = =*2*B*B4*H'+'@'@'F +;"0"EFVEWWbcgbhhi j#:I%56!> && (s3 &&rc ZRV RVP R2pVPpVPpVPP'g7VPP V4pVPP V4p\ \VPV44F=wpwr\PPPW4'gK9RWx&K? VPP'g"VPPV4VnMWrnVPVVPVPRRR7wrK\PPV4p WP P"9ddVP P"V,p \%V \&4'g1\PPV P(^,4p V R\+V 4 R2p\%V\,4'd RV RV V R2#VP/W2V4#)rrrFTrr r tl.store()rrrrr#r!rMr"rXrrrrr-rrrinplace_buffersrrf other_namesrAr=codegen_descriptor_store_line)rRrrrrrmrrrr broadcast_dimrT store_dtypergs&&&&&& rcodegen_block_ptr_store_line)TritonKernel.codegen_block_ptr_store_lineEs #5'H,@,@+AC #22$66%%111&44;;OLO ( 6 6 = =>O P *3 $$o 6* %C%#ww77KK).!& * %%111)1)?)?)D)DEV)WH &): &99      : gg''- 99,, ,))++D1Cc:..gg//0BC '/ <=Q? h 0 0ykE75'; ;11)uMMrc `V R\PPVP4 R2#)zFGenerate the descriptor load line. Override for backend customization.z.load(rrXrr$r)rRrrs&&&rcodegen_descriptor_load_line)TritonKernel.codegen_descriptor_load_linexs-"#6!((*?*?@P@P*Q)RRSTTrc fV R\PPVP4 RV R2#)zGGenerate the descriptor store line. Override for backend customization.z.store(rrr0)rRrrrs&&&&rr**TritonKernel.codegen_descriptor_store_line|s2GAHH$9$9(:J:J$K#LBugUVWWrc(V^8dQhRRRRRRRR/#)rrrrdrrrBr)rs"rrrs8MMMM M  Mrc tV'g V'gR#\V\P4'g\R\ V4 24hVP VRRR7p\V\ 4'g\R\ V4 24hVPpVP4'd VPMRpV'd\VPV44MRpTPYc'dRMRW4p VPV4p VPPWR\ P"R7R#)Nzexpected sympy.Expr, got Frr0) assignmentr)rrrr rrrErFrSrtr!rename_indexingindirect_assertget_load_bufferrrrrF) rRrrdrrBrrFrtsize_strliner s &&&&& r check_boundsTritonKernel.check_boundss $ ++ #rc VP4F;pVPV4'gKVPVP44u# R#r)rkrir r )rRr)rs&& r_range_tree_mask_shape#TritonKernel._range_tree_mask_shapesA++-D~~d##t'>'>'@AA.rc V^8dQhRRRR/#)rr)zstr | TritonCSEVariablerrr)rs"rrrs11 71N1rc h\V\4'd VP#VPV4#r)rr_rrE)rRr)s&&r _mask_shapeTritonKernel._mask_shapes+ d- . .:: **400rc$V^8dQhRRRRRR/#)rrrrHz#OrderedSet[str | TritonCSEVariable]rr)rs"rrrs(7  rc VfR#\V4pVF0pVPV4pVfR#\V\V44pK2 V#r)rrIr?)rRrrH result_shaper)r s&&& r_broadcast_shape_with_masks(TritonKernel._broadcast_shape_with_maskssQ =U| D))$/J!0uZ?PQL  rztl.extra.cuda.gdc_wait()z%tl.extra.cuda.gdc_launch_dependents()c \PPPP'gR#\ \ P\PPP4'dR#\PP'dR#\ PP4PR8H;'d)\PP!4^,^ 8#)Fr0)rrr#r enable_pdlrrXrrrrrrrrr0rwrrr_enable_pdl_codegen TritonKernel._enable_pdl_codegens%%,,777 ahh @ @ U U V V ==    GG / / 1 6 6& @ ; ; 00215: rconsider_readsc aaaaVP4'gR#\PPo\PP e%\PP P MRoVVV3RloV'g \R4hSRJ;'g;\;QJdV3RlV4F 'gK RM RM!V3RlV44pV'gR#VP'd0VPR,P'd VPpVPVP4R#)NcL<aSf \R4hSf \R4hSPPpS'd,\P!VSPP 4p\ ;QJd VV3RlV4F 'gK R# R#!VV3RlV44#)Nzprev_node must not be Nonezcurrent_node must not be Nonec3<"TF8pSSPPVPVP48HxK: R#5ir)mutation_renamesr r)rw current_noder[s& rrOTritonKernel._handle_pdl_before_access..matching_dep..s8"A|4488HH"sAATF)r rVwritesrrrWr)r[ prev_depsrTrZ prev_nodesf r matching_dep.matching_deps $%ABB#$%DEE!--44I%OOIy7L7L7R7RS 3"33 3 3" rzdependencies must not be emptyc34<"TF pS!V4xK R#5irr)rrr_s& rr9TritonKernel._handle_pdl_before_access..s,Sl\!__lsTFr)rRrXrrZrrc previous_noder rrrrrlrQGDC_WAIT)rR wait_bufferrT dependencies need_waitrZr_r^s&&d* @@@r_handle_pdl_before_access&TritonKernel._handle_pdl_before_accesss'')) xx,, /0ww/@/@/LAGG   + +RV   !AB B%SS,Sl,S,Sl,S)S    T%5%5b%9%A%A%A))Kdmm,rc JVP4'gR#VP^8dR#VP'd0VPR,P'd VP pVP VP4VP VP4R#)Nr) rR use_countrrrrrQrd GDC_LAUNCH)rR launch_buffer result_vars&&&r_handle_pdl_after_load#TritonKernel._handle_pdl_after_loadsx''))    ! #   T%5%5b%9%A%A%A 22M  .0rcV^8dQhRR/#rcoderar)rs"rrr s   rc N.pRpRpVPFp\V4\JdVPV9d V'dK1Rp\V4\Jd2VPV9d!VeVP V4\ V4pVPV4K W!nR#)FNT)_linesrrrdrlrrr_)rRrs new_lineshas_waitprevious_launchls&& r _filter_pdlTritonKernel._filter_pdl s AAw#~$--1"4#HAw#~$//Q"6".MM/2"%i.   Q  rc$V^8dQhRRRRRR/#)rrKrrrorzLoadIndexBasis | Noner)rs"rrrs$(((':( (rc \PPp.pVPFpVPP V4pVfK$VP VJgK6VPVP\PP4'dKrVPV4K \V4^8:dR#.p\PPpV'dp\V4F^wrVPVPV4'gK)VPV4WP,pVP!V 4Ks R#VPWP"4'gR#\%V4#)aFind split digits that reconstruct a range tree's flat index. Each entry is one digit in a mixed-radix index. A complete basis has at least two non-unit digits, starts at divisor 1, and forms a contiguous divisor chain whose total extent equals the root's numel. Symbolic lengths are supported when size analysis can prove those identities without guards; otherwise this conservatively returns None. N)rXrrrrr r rlengthrrrr_rrMdivisorrrr) rRrKrr remainingr%rbasisrr,s &&& r_load_index_split_basis$TritonKernel._load_index_split_basiss77##02 ((F))--f5E!JJ$& 88uww{{SS  ') y>Q ,.''++%i033EMM7KKLL'||+GMM!$ 1//DDU|rcV^8dQhRR/#rNr)rs"rrrEs   $  rc FVP\J;'dVPP4;'dgVP'*;'dN\ ;QJd&RVP 4F 'gK RM RM!RVP 44'*#)z=Whether this kernel uses a supported shared coordinate scope.c3B"TFp\V\4xK R#5ir)rrlrs& rr@TritonKernel._supports_load_index_basis_reuse..Ks",D4!;<<,TF)rfr$rr1rrrrQs&r _supports_load_index_basis_reuse-TritonKernel._supports_load_index_basis_reuseEs NNl *   **,  ...  C ,,CCC ,,  rc V^8dQhRRRR/#)rstateLoadIndexStaterrUr)rs"rrrQs  ~ . rc aa\S\4'd SP#\;QJd".VV3RlSP4FNK 5#!VV3RlSP44#)z.Ws) HXD ( (d ; ;HXs),)rrTrVrr)rRrsffr_resolve_load_index_bases&TritonKernel._resolve_load_index_basesQs[ e4 5 5;; u HLHXHX u u HLHXHX   rc(V^8dQhRRRRRRRR/#)rrKrrrorLoadIndexBasisrr)rs"rrr[s8JJJ"J J  Jrc \RV4\PP4p\ VUu.uFqUP 4NK up4p/p\ PPpVPV, Fp VPPV 4pVeVPVJdK4VPVP\PP4;'d'VPVP VP"4p V 'gKWGV &K V'd \%W4#T#uupi)zAReplace a tree's flat symbol with its exact split-coordinate sum.c3b"TF%qP4VP,xK' R#5ir)r%r)rrs& rr>TritonKernel._rewrite_full_range_with_basis..cs ?\\^emm + +s-/)rrrrrr%rXrrrrr r rrrr~rrP) rRrKrr split_indexr basis_symbolsr"rr% is_full_ranges &&&& r_rewrite_full_range_with_basis+TritonKernel._rewrite_full_range_with_basis[s ? ? ##FuLLN#FG 79 77##((=88F))--f5E} $ 6$<< uww{{MM225<<L }'2V$93?z%.IEI$GsEc$V^8dQhRRRRRR/#)rrrrKrrr)rs"rrrus!&&C& &z&rc |VP4'gV#\\VP4pVP V4pVP P V4pVf\W"4VP V&V#VPV8Xd VP#VPV4pTp\VPVRR7FwrV fK VPWxV 4pK Wr8wd8\PP P#WpP%44p\'W'V4VP V&V#)aERewrite a full-range load coordinate to an observed split basis. Reduction bodies can use split coordinates while later pointwise work uses one flattened coordinate for the same range root. Looped reductions may reuse a basis within their body; closing the loop invalidates CSE and clears the basis before codegen enters a different scope. TritonKernel subclasses and derived iteration families may use separate codegen scopes. Cooperative reductions also use an untested partitioned execution model. Only the base, non-cooperative kernel with base range trees is accepted here. T)r)rr rXrrgrZr rQrKrrr"rrrXrrsimplify_with_ranges var_rangesrT) rRrrKrrrrVrrrs &&& r_reuse_load_index_basis$TritonKernel._reuse_load_index_basisus4466L9dhh'OOD) &&**95 =1J%0WC " "9 -L ;;% << ..u5t//tDKD <.s MqFMr;z, eviction_policy='evict_last'cH<SS,S8dS'g S'dR#R#) evict_last evict_firstr)expected_countr[indirect_indexing load_countsrsr decide_later'TritonKernel.load..decide_laterst$~5"3'$rz, eviction_policy=''zrz, other=z , other=0rz, cache_modifier='.cg'rrrr&r5rx), r|rrttl.debug_barrier()rrTruez tl.full([1], z , tl.int1)rrz0.0r7rrr)rFalse)[rrrrrrrXrrrzr#rrrr_rrrktma_compatibility_checker_clsrrrrErCrKrr[rdrget_strides_of_loadris_broadcastedrrrrrrrErS _load_otherrk skip_l1_cacherbuffer_read_countsrxrr[r\rZrrRr=r0r#r1r-rrrrrrIrFrtrLrrrrrr;rinvalidated_storesrrrQr`rhrrorkr_r rHrrrNrrlr)!rRrrKr make_lineoriginal_indexruses_uint8_storageskip_tma tma_checkerrrd is_coalescedeprrm has_read_depsrrcachemodappend_broadcastrr=r load_bufferrn bcast_operandr^ other_valrr[rrs!&f& @@@@rrTritonKernel.loads3iiood#'' DQ=E ,,T59 55e<!!$'LUX == - - -88@88:AAGGRH1155c:==55  $:V:V #;   33d$DeL 4 == &1!  h 0 0T5N5N NN6 6 26D .((* **, s  44^DKKM sss  44^DKKM     ~ . .1B1B  " " "  % % %)9)9")=)E)E)E % %).N+B!))*:FLQIB :8+<+<+>+>"=1A1A#B"CD##%E  == & & &!%!A!A!C )) 24 81 <  ##N3 3  )))  !!     /H $ #D ) )D77==777!MMEd;u#5"6a88DE(_6M$NOO*.*@*@#+' %h88%&6%7wrd8*AND<<=MxXD&DD((((#'# E e'~661155c:--11#t<!#d>*:"=#+#6#6 1155c:--11#t<!#d8+=+=*>c(BSBSATUWTXY^X_`h_iijk ((($11E)99(..IE%--88MM888))  "u}}'8'8'@& **84 DHH// /335::eC  ! !,t5I"J K &&{D9XX&& 4U'  ##K<   ! #   " *&788 #CDDTCU!VW W'11   &M "s::K'K#ZF)BA#J*U&m_B7G6HJD**T6K6K+J!!!*** Dejj(!DD7;7G7G7GM$"2"23T#8#4#4"5R |2i[PQR!XX..::"(((*<*< / $$$ tyy(@  " " & &z 2rc ,V^8dQhRRRRRRRRR R /# rrrrKrrr_r4rWrrr)rs"rrrs;bbb *b3>bFOb brc  B VPPV4pTp\PP V4pVP WV4pRp V'g<VeVR8Xd1VR8H;'g\ VRR4p VPVVRV VR7p VPVRVRJV VR8HR7p \V \4'd=VPV P4'dVPPV4WPP9p VP!V4p V 'd.V 'd&VP"P%\'VR 44\V \(\*34'd*VP-WV 4wrVP/WWV4pEMVEf/V P0p\3V4'dVP4'dVR VP74 R 2, pMVP8e\:;QJd&R VP84F 'dK RM RM!R VP844'g7R P=\?\@VP844pVR V R 2, pVPBPEV4VPFPIVR4RV RV RV R V PJ R 2 pMVR8XdRVn&V P0p\3V4'dVP8e\:;QJd&RVP84F 'dK RM RM!RVP844'g7R P=\?\@VP844pVR V R 2, pRV RV RV R V PJ R2 pM\ORV 24h\PPR!4pVPT'g=VPV'd+VPYVP[WP"44VP]VP"VRR7VP"P%\'VV44VPT'gVP^PEV4VPa4R#)z^ store the 'value' to the memory location 'name', offset by some indexing expression 'index'. Ntma tma_storeFTr atomic_add)rrrrr.broadcast_to(rc3>"TFp\V4R8HxK R#5irNrrs& rr%TritonKernel.store..s9+6aCFcM;rr'r5rc3>"TFp\V4R8HxK R#5irrrs& rrr?;aCFcM;rztl.atomic_add(z, sem='relaxed')z store mode=)rT)1rrrXrrrkrrrrrErCrKrr_r(rstoresrQr`r=r0r#r-rFrrrrrrqrsrrrrrrtatomic_add_foundr$ contextlib ExitStackrr enter_contextguard_cooperative_storerhrclose)rRrrKrr4rrrbuffer_misalignedrrkr is_inplacerrrmr= indexing_str value_shape exit_stacks&&&&& rrTritonKernel.storesiit$!!$' 88EJ$(! dldemEMFFWT;%FE(,(J(J )K) %== dl&? $ 4 !  h 0 0T5N5N NN6 6   , , 3 3D 9YY666 ,,^< . KK ! !,t5I"J K h2I J K K&*&<& . Q ?? #C5\N#eWBxGXGXFYYijD% D6&:; ;))+ $$$)C)C)C  $ $T%A%A$ %T U &&t{{D&N l467$$$  " " & &u -rcV^8dQhRR/#rr)rs"rrrsIIIrc \VPPRV R\V4 R24R#)ztl.device_assert(rrN)rrQrN)rRcondmsgs&&&rdevice_assert_async TritonKernel.device_assert_asyncs( !24&49+QGHrc VPP4pVP\VRV R244VP 4#)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):)rrErQr`indent)rRrr rs&&& rr$TritonKernel.guard_cooperative_storesC 88NNPd.?uK,PQR}}rcV^8dQhRR/#)r variablesCSEVariable | Noner)rs"rrrs  ); rc RpVFApVfK \VR4'gKVfVPpK/W#P,pKC V#)NrH)rrH)rRrmaskselems&* r_combine_masksTritonKernel._combine_maskssGD|t[))= NNE!NN2E rc8V^8dQhRRRRRRRRRR R R R R RR/#)rrr_ boundariesz.tuple[str, sympy.Expr, sympy.Expr, sympy.Expr]boundary_indicesindexing_dtyperrgrsorterztuple[str, sympy.Expr] | Nonesorter_indicesrrr)rs"rrr s[999C9& 9 $ 9  9.9+9 9rc VPP\P4VPP V^,4pVP V^,4p VP V^,4p VP V^,4p V'd#VPP V^,4MRp V'dVP V^,4MRp V\P8XdRpM#V\P8XdRpM \R4hVP!VPV^,.V'd V^,.M.O5!VPPVPRV RV RV RV RV RV RV RV RV RV RV R2VVPR7pVP!VPV4VP#WV4pVVnV#) z# See [Note: Inductor bucketize op] rrztl.int64z5Bucketize only supports indexing with int32 and int64z'triton_helpers.bucketize_binary_search(rz, )r)rrr0ONE_ELEMENT_PER_THREADrrr$rrFrDr$rhrrrrrorrH)rRrrrrrgrrboundaries_ptr boundary_sizeboundaries_underlying_numelboundary_stride sorter_ptr sorter_stride triton_dtyperrs&&&&&&&& r bucketizeTritonKernel.bucketize s$  C CDA7))*Q-8 &*&7&7 1 &F#++JqM:39TYY__VAY/v 8>))&)4F U[[ (%L u{{ *%L%G  && LL*Q- :@F1I;b "" LL5fXRbr2M1NbQ`Paac nBgRl"]O2  !,,#   ##DLL&9##FnM  rcV^8dQhRR/#rr)rs"rrrHs...rc VP4pV^8XdRV R2#VPpR.W#, ,R.V,,pV RRPV4 R2#)r!triton_helpers.promote_to_tensor(rr]rrUrrVr rrq)rRrndimsnreducers&& rreduction_resizeTritonKernel.reduction_resizeHsh'') A:6ugQ? ?)))VHw,>>$))E*+1--rcV^8dQhRR/#)rrztuple[str, BlockShapeType]r)rs"rrrQs 9 9:T 9rc VP4pV^8Xd RV R2V3#VPpR.W4, ,R.V,,pVe.VRW4, O^.V,O5MRpV RRPV4 R2V3#) rr rr]rNrUrrVr )rRrrr r rrs&&& rreduction_resize_and_shape'TritonKernel.reduction_resize_and_shapeQs'') A:6ugQ?F F)))VHw,>>=B=N 9e'u( 9A3= 9TX $))E*+1-y88rc$V^8dQhRRRRRR/#)rrr_rrrr)rs"rrr]s$  ( 1<  rc VP^8XdV#VP4VP, pVP4pVRVR.,pVPP V\ \ V4WV4V\V4R7#)z3 Reshape to RBLOCK, collapsing all reduction dims. Nr'r)rr r.rrrrr)rRr rr target_ndimr target_shapes&&&& rreduction_collapse_dims$TritonKernel.reduction_collapse_dims]s  " "a 'L--/$2I2II ,,. $\k2hZ? xx   3u:} C % !  rc ,V^8dQhRRRRRRRRRR/#) rrrrrrVrz%CSEVariable | tuple[CSEVariable, ...]rr)rs"rrrqsCy y y y & y  5 y  / y rc )aaaaaSaTaUaVaWaXaYaZa[a\a]a^RRloZRVZ3Rllp\P!VZ3RlV4pSpSpV'dW\P!WT4pSZ!S4'd\PMSoSZ!S4'd\PMSoSP 'g \ R4h\RSP44p SPV 4\V 4p SP'dV PSP4SPR],P^,p SP4RJ;'dSR 8Ho[S[;'dSP'*p SP 'dzSP#4p \%V 4^8d\ R \%V 4 24hV U u.uFp R V 9g R V 9gKV NK pp R RP'V4 R2oV\)V4o]M)SP+4oV\)SP#44o]SP-VVVV]3RlV4pR^pR_oSR`oTVSS,ST,pRRRRRRRRRRRR/pRpSV9d\/V\(4'dVwppSP14SP2, oWRVWVVVV[3RlloXRVX3RllpRaVWVYV3RllpVSVTV3RlpSSV3pVSP4P69dSP4P6V,#\9S4p\;S4pSV9dSP<P?4MRp\ASP#44p\CSWSWSP2,4FpR VV&K \R!V 44pSST9dmVf \ R"4hSP4PEV\)V4R#7SP4PEV\)V4R#73pVF pVVn#K MqSV9d>Vf \ R"4hSP4PEV\)V4R#7pVVn#M-SP4PEV\)V4R#7pVVn#R$P'V 4oUVU3R%lo^SP'Ed\HPJPMSS4pR&o\R'VVV\V^3R(llp SR)8Xd\/V\(4'gRp!M\/V\(4'd)\OVV4U"U#u.uFwp"p#V !V"V#4NK p!p"p#MBSR*8Xd3SP4PQSPRWDPTR+7p!M V !VV4p!SV9d\/V!\V4'g\ R,\YV!4 24hVf \ R-4hV'd#R.\[V4 R/SP]V4 R02p$MC\[SP4PQSPRR1V R2V! R32VV!P^R#74p$VS,oYV!SPRVV!V$V!44E M/SR48Xd?SP`'dSPcVSVS^VS4pEMSPeSV4pEMSR58Xd\/V!\f4'g\ R6\YV!4 24hV!wp%p&p'\(;QJd6.VV3R7lSPiSPRV%V&V'SWS44FNK 5M/!VV3R7lSPiSPRV%V&V'SWS444pEM>SR)8Xd\/V\(4'd\/V!\f4'g\ R8\YV!4 24hSP4PES\k\VV4P^R#7p(SPmSPRVV(V!^,V!^,SWS4pEMSPoSV4pEMz\/V!\V4'g\ R,\YV!4 24hSX!SPRV!V!PT4wp)p*p+S['dBSP<Pq4'g"\s\k\tV44p,V, R9V) R02p)SP4PQSPRV)V*V+R#7pEMSST9d\k\tV4^,MTp-SP4PwR:V- 2V\)SP#44R#7p.\HPJPySS4pSP-\rV4p\/V\(4'Eg7SR*8XdSP#4p \%V 4^8d\ R \%V 4 24hV U u.uFp R V 9g R V 9gKV NK pp \)V4V.n/R RP'V4 R2oVSPzP}V. R;SV RV RV R024MV 'dN\)V4V.n/R RP'V4 R2p/SPzP}V. R;V/ RV RV R024M5SPzP}V. R;SP+4 RV RV R024SV9EdR:V- R<2p$Vf \ R"4hSPzP}V$ R;SP+4 R\P~!V4P RSP]V4 R024VS,oYVe"R.\[V4 R/SP]V4 R02MV R2p0SPRPR=V. R>V$ R?SY R@V. RV$ RV RV0 RAV. RBS^!V. RC2V.4 RDV$ RBS^!V$ RC2V$4 RD24V!SPVV.V$V!44EM\S4'dSPcVSVS^VS4pEMSR)8XEdR:V RE2p1R:V RF2p2SPzP}V1 R;SP+4 RGV R024SPzP}V2 RHSP+4 RV R024\/V\(4'dWVwp3p4SPRPRIV1 R>V2 RJV1 RV2 RV3 RV4 RK\P R\P RL24MMSPRPRIV1 R>V2 RMV1 RV2 RV RK\P R\P RL24SPRPRNV1 RBS^!V1 RC2V14 RNV2 RBS^!V2 RC2V24 RN2 4Tp5SP4PES\k\tV54P^R#7p(SPmSPV5V(V1V2SWS4pEMV 'd\k\ZV4p,SP4PQSPRS^!VV,4VPTVP^R#7p6SX!SPRV6R4wp7p8p9SP4PQSPRV7V8V9R#7p:SPRP}V. RBV. ROV: 24SPP}V RBV. 24M\HP!SS4p;V;!V.V4pV>P}RQ4V=PV>P44K5 SV9EdySST9d>\k\(\V\V3,V4wp?p@T?p-\[V?4pA\[V@4pBMoSSS9d1\k\VV4p-\[V4pASPV RR24pBM8\k\VV4p-SPV RS24pASPV RR24pBV- RT2pCSPP}VC RBXA 24SPVCSV4pDSP<P?4pV- RU2pESPP}VE RBXB 24SPVEV\P~!V4P4pFV!SPVVDVFV!44EM\S4'dSR48wd\ RVS: 24h\k\(\V\V\V3,V4wpGpHpISPVG\;S4V^,4pJSPVH\;S4V^,4pKSPVI\;S4V^,4pLSPSPVGVHVIVJVKVLSWS4 EM SR)8Xd\k\(\V\V3,V4wp5p(\/V\f4'g\ RW\YV4 24hSPV5\;S4V^,4pMSPV(\;S4V^,4pNSPmSPV5V(VMVNSWS4MMSP\k\VV4\;S4V4pOV!SP\k\VV4VOR4V=P4\k\tV4SP4P6V&\/V\(4'dTMV3pPSPPVP4\;QJdRXXP4F 'dK RYM RZM !RXXP44'g \ R[4hV'dx\XP4FhwppQSST9d V^8dKSSSST,9dTMTpRXQPTVR8wgK;SPP}XQ RBVQ R\\XR4 R024Kj V#uup iuup#p"iuup i)bzC codegen reduction of value to Triton according the reduction_type c V^8dQhRRRR/#)rrrRrrr)rs"rr,TritonKernel.reduction..__annotate__|s L L/ LD Lrc`VRJ;'d$VP;'dVP^8#r)rrE)rs&r should_upcast-TritonKernel.reduction..should_upcast|s*D=KKQ%8%8KKQZZ!^ Krc V^8dQhRRRR/#)rrr_rr)rs"rrr s    rc<S!VP4'd&\P!V\P4#T#r)rrrrr)rr! s&r maybe_upcast,TritonKernel.reduction..maybe_upcasts7!-- UEMM2  rc(<S!VP4#rr<)rr! s&r(TritonKernel.reduction..smAGG.Drexpected inside_reductionc3@"TFqP4xK R#5irrrs& rr)TritonKernel.reduction..I8H>>++8HrONrz%expected at least 3 dense sizes, got XYrUrrVcx<SPPSPRV RS R2VPSR7#)rrrr)rrrr)rrrRrs&rr( r) s=dhh'' "1#R'7q9gg! (rargmaxargmin argmax_value argmin_valueargmax_with_valueargmin_with_valuerrc$V^8dQhRRRRRR/#)rrr_ result_typerRrz.tuple[str, torch.dtype | None, BlockShapeType]r)rs"rrr s(& .& .& .,& .< & .rc <\S 4pS PWS 4pS R8Xdk\S P44^8HpVPf \ R4hV'dV R2p^.VPO^NpMKV R2p.VPO^NpM5S 'dRMRpS P V RV RS V R 2VP4wrVVeV R S PV4 R 2pM VPpWRV3#) z3 Helper to generate a reduction call, e.g. tl.sum. rzvalue.shape must not be Nonez[None,:,:,None]z [:,:,None]zB, reduction_ordering=tl.constexpr(tl.ReductionOrdering.INNER_TREE)rrrrr) rr rr.rr r rur) r rr8 triton_reduction_fnis_bmmrrreduction_orderingrrrrR strict_sums &&& rfinal_reduction/TritonKernel.reduction..final_reductions"#@"O 00FE&T11349;;&()GHH %wo6F00a0E %wj1F-ekk-1-E"Y# !% ? ?*+1UG2cU;M:NaPKK!  &"84(9(9+(F'GqI#kk - -rc(V^8dQhRRRRRRRR/#)rrnr_rr8 rRrrr)rs"rrr s2 5 5# 5 5, 5  5rcN<S!WV4wppVPV RV 24R#)z= Generate a reduction and assign it to an existing variable. rJN)r)r rnrr8 rTr> s&&&& rfinal_reduction_define6TritonKernel.reduction..final_reduction_defines.*&EKE1a MMZLE73 4rrKc<S PWVP4pS PTT\V\4'd%\ \ PVP4M\ PP44pVR8XdWVwrVVPRV RV RS RV RV RS RV RS PV 4 RV RS PV 4 R24R #VR8Xd?VPRV R V R S RV RV RS RV RS PV R 24 R24R #VPRV R V R S RV RV RS RV RS PV R 24 R24R #)value_and_indexz rz = triton_helpers.z _with_index(z) rJ rz_val, z_idx = triton_helpers._val_idxN) r rrr_r rrXrrrr ) r rnrrK result_kind result_value result_indexrroot_oprRs &&&&& rfinal_argreduce/TritonKernel.reduction..final_argreduces00 LE00e[11U[[%++.XX<<> E//-7*  !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#OPrc,<SS9dR#SS9dR#R#)rE rrKr)arg_value_reduction_typesarg_with_value_reduction_typesrsrargreduce_result_kind5TritonKernel.reduction..argreduce_result_kind4s!??(!::rrc3Z"TF!p\V^,4'dKVxK# R#5ir+)rMrs& rrr, Is!&  C(;CF(CCC5s+ +z*index_dtype must be set for arg reductionsrrpcD<S'gV#\PSW4#r)rr )tvalfvalrs&&rr*TritonKernel.reduction..where_condbs (..tT@ @rcVP'gV\P8XdV#W8XgV^8XdV#V\P!V4P8Xd!\P!V4P#V\P!V4P 8Xd!\P!V4P #V#)z7update reduction constant mask value to match dst_dtype)rrriinforr)rrrs&&&rupdate_constant_dtype5TritonKernel.reduction..update_constant_dtypejs ...)uzz2I#O)X]#Ou{{95999 ;;y1555Y!7!;!;; ;;y1555#OrcV^8dQhRR/#rrr_r)rs"rrr |s  { rc<S!VSVP4pSP\V4pSPP SP S!W4VPVP R7#r)r_map_tuple_or_scalarrkrrrr)rr default_strrRrr[ rs&& r _mask_value+TritonKernel.reduction.._mask_value|s`/EKKP"77 wO xx((LLu2++++ )ronline_softmax_reducerr<zexpected CSEVariable, got zexpected index_dtype to be setrrnrrzindex, z.shape)welford_reducewelford_combinezexpected Sequence, got c3z<"TF0wrSPPSPVSVR7xK2 R#5i)rN)rrr)rrrrrRs& rrr, s=#) HH%%dllEe%TT)s8;z,expected masked_value to be a Sequence, got r5rTr_index _next, z_next = triton_helpers.z%imum_with_index( # ) rJ_next _max_sumz, float('-inf'), = tl.zeros(z zT_next = triton_helpers.online_softmax_combine_with_sum( z, z3 ) zK_next = triton_helpers.online_softmax_combine( rF rz .to(tl.int8)zif HAS_RSPLIT:rH rG _bval_bidx.expected reduction_type 'welford_reduce', got zexpected Sequence default, got c3B"TFp\V\4xK R#5ir)rr_rs& rrr, sJ\:a!233\rFTz2all result_tuple entries must be TritonCSEVariablerr)r1 r2 )r3 r4 )r5 r6 )rK)Xpytreetree_anytree_maprrrr rrrrrrfr_rryrr-r.rrqrrr` rr rrreduction_cacherJrGrselect_index_dtyperrrPrHr$rrr"rrrr_rrrurrre welford_reduce_fallbackr_welfordr %online_softmax_reduce_final_reduction prepare_softmax_twopass_fallbackryrkr rdefault_accumulatorrlrQrZ rrrrJr#rrget_reduction_combine_fnrr@rrrrrr *codegen_cooperative_reduction_peer_combinewelford_reduce_final_reductionrrrirrMr=)_rRrrrrr% do_upcastoriginal_dtypeoriginal_src_dtyperreduction_range_prefixstrict_sum_loop dense_sizesrd xy_sizes_onlyarg_index_reduction_typesarg_reduction_types arg_root_ops logical_indexrB rM rR racc_typetorch_acc_typerqrMr,result_mask_varsrnrrrb masked_valuerraccumulator_indexmeanm2weight result_sum_result_dtype_shaper^ result_prefix accumulatorresult_size_strraccumulator_maxaccumulator_sum value_max value_sum result_maxrn chunk_expr chunk_dtype chunk_shapechunk combine_fnupdatedrrgrJ rK 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_dtyperP rQ rrrr> rL r! r= r[ rrs_ffff& @@@@@@@@@@@@r reductionTritonKernel.reductionqs L  OO$DeL & OOL8E)6y)A)A yI%25%9%9EMMuE$$$ !<= =I8H8HII % u  ??? LL )!%!1!1"!5!A%>QR>Qda Aq 1>Q R 5( $xx00uKK0X *5': !44!, <<(4T,5G4HI&()IJJ C ./uT5F5F{5S4TTUV&),)) LL./E.Fgl^[bc"-"."4"4 *)%'~6LL %)+  #33---!%!4!4"NE:xQV"J"&!=!=eU!KJ#44!,99(+B4 CUBV)WXX%1"r6"U#(,  dBU)#UU#(,  dBU)#  #::eU++%lH==,J#L124"&#4 Z+H+N+N"1"J"&!K!K ""$Q$Q"J"&!F!Fue!TJ!, <<(4T,5G4HI+:LL, 0B0B+'dmm&U&U&W&W(c7);2B"WIRPXzYZ[%(-l(;K%()$))L*A)B!&DOII''&-{?2C2gYbQYPZZ[\II''&-{43F3F3H2IG9TVW_V``ab!44&' f$=!&()UVV ##()T5H5H5J4K2{{;/334Bt7H7H7U6VVWY'~6%0M*+51B1B;1O0PPQR2359  ## W%6$77NwiX M$5#6br)M S{m5,A;!O PQ"#3z5F4Gu2MO`'a&bc  **%)+ &n55!00z8U  #::$%j\"6$%j\"6 ##&'{43F3F3H2IIZ[cZddef ##&'|D4G4G4I3J"XJVWXeU+++0(IyLL''()0AB,-R/@9+RPY{[#112"V5N5N4OPLL''()0AB,-R/@5'J#112"V5N5N4OP ##$%S6Gu4M)_(`a$%S6Gu4M)_(`a( !XX__tC'<'B'B- "GG**## !C)**LLud+++++ + 8GLL&$84 K))LL%% *  &&+c+c%'QR&&00J>@ $1?%!8&&00N3C3{m1TUJJ"K1M1M (()+ &n55!%55(HHZ[9=+{K?@*95 Y !KK$Y/AJ II$Y/AJ #MM!$Y/AJ 33((!  #::)-+{23Z*& J"'844(9$w-I JJ 0 ;WQZ JJ 0 ;WQZ::((GGj13CI3NPW'((${J*GPT    .23 .C  +%/ E%B%Bz  %%l3sJ\JsssJ\JJJ !UV V &|4 6!%CCA&14RRT'(  <<</**44!(#fXT2El2S1TTUV5yYB Sh%s# AR<AR<&AS6AS AScV^8dQhRR/#rrrr)rs"rrr,s&&DO&rc VPWV4pVPWV4p\^4Uu.uF(p\VPP VR74NK* upwrxVP RV RV RV RV RV R\ P R\ P RV RVPV 4 RV RVPV 4 R24Wx3#uupi)rr< r9 = triton_helpers.online_softmax_reduce( , ) rJ) r rrrrPrr#rrr ) rRr r r rrrTr r s &&&&&& r_online_softmax_reduce#TritonKernel._online_softmax_reduce,s66vPUV66vPUVMRSTX!VX#dhhooEo&B"CX!V   L:,' !O#4Bse<%%&b)B)B(CD LD11ZLBCD LD11ZLBCD   %%"Ws.CcV^8dQhRR/#r r)rs"rrr>s  [ rc  aaaaVVV3RlW#V34wr#pRV RV RV RS R2 pV3RlpW#V3U u.uF0p SPPSV!V P4R7NK2 p p SPRP V U u.uFp \ V 4NK up 4 RV 24\ ;QJd.V3RlV 4FNK 5#!V3RlV 44#uup iuup i) z+ Helper to codegen triton_helpers.welford. c3J<"TFpSPSVS4xK R#5ir)r )rrr rrRs& rr(TritonKernel._welford..Bs* +  ( ( > >+s #ztriton_helpers.welford(rrcB<\V^SVS^,R,4#r+)rrQs&r reduced_shape,TritonKernel._welford..reduced_shapeHs$qcAgi(889 9rrrJc3Z<"TF pSPWP4xK" R#5ir)r r)rrrRs& rrr Qs* (  + +E;; ? ?((+)rrPrrQrqrr) rRr r r r rrwelfordr rwelford_resultsr<s ff&&&ff rr{ TritonKernel._welford>s F+ &,D6B4r&C5J : F+ + HHOO%}U[[/IO J+   DII&G!s1v&GHIWIVWu ( u u (    'Hs 6CC c VP4VP, p\V R2\VP 44V\ P !4R7p\V R2\VP 44V\ P !4R7p \V R2\VP 44V\ P !4R7p VPPV RVP4 RV R24VPPV RVP4 RV R24VPPV RVP4 RV R24VR8Xd>Vwrp VPPR V R V R V R V RV RV R V RV RV R 24MGVR8wd\RV: 24hVPPR V R V R V RV RV RV RV R24VPPRV RV!V R2V4 RV RV!V R2V 4 RV RV!V R2V 4 R2 4TpVPVPVRRVV V VV4 #)z%Helper to codegen a welford reduction_meanrrra_m2_weightrp rrrf ri rj z<_next = triton_helpers.welford_combine( z, rk re rs z;_next = triton_helpers.welford_reduce( z1, roffset == 0 ) z rJrl r N)r rr_rr.r"unknownrlrQrrrr r r)rRrnrrrr rrr accumulator_m2accumulator_weightr r r r s&&&&&&& rre TritonKernel.welford_reduceVs%%'$*A*AA'l% ,,./&&(  +l# ,,./&&(   /l' ",,./&&(   m<(;(;(='>b ! L  l4+>+>+@*AH:Q O  !",t/B/B/D.ERzQR S  . .$ Df LL   W^$4GRS   MZ;-u(={KLM  C n-=U+C^ TUV  J2D1EU/KM_$`#ab   ! 22  " "          rc  \VPWWgW44p W#V.p \\W44FEwp wp wrV f!VPP WR7p WV &VP V RV 24KG \V 4#)z0Helper to codegen call to triton_helpers.welfordrrJ)rr{ rMr"rrPrr)rRr r r r r r r rrr result_exprsr, result_exprrrs&&&&&&&&&& rr +TritonKernel.welford_reduce_final_reductionsdmmF"cIJ# > 09#l:S0T ,A, ^e""hhooEoG "-Q MM[MUG4 5 1U \""rc .VPWV4pVPWV4p VPRV RV RV RV RV R\P R\P RV RVP V 4 RV RVP V 4 R24W#3#)r rr r r rJ)r rr#rrr ) rRr r r r r rrr r s &&&&&&&& rr| 2TritonKernel.online_softmax_reduce_final_reductions66vO66vO  L:,' !O#4Bse<%%&b)B)B(CD LD11ZLBCD LD11ZLBCD   %%rc XVP'dVPR,#\#)RSPLIT)rr6rQs&r max_rsplitTritonKernel.max_rsplits%    $$X. .  rcV^8dQhRR/#r^ r)rs"rrrs rc 2VPR,pVP4'gRMRpWBP,VP4,pVPP V4wrxVP PRV RV RVPV4 R\V4 RV R V R V R 2R R 7VPV R2RR.V\P!4R7p VPPV RV R\V4 R24V #)z Generate code to save a [XBLOCK, RSPLIT] temporary workspace, where each thread block writes a different column. After the barrier, every thread block loads the completed value so that it can compute the final value independently. rzxindex < xnumelNrm z_ws = (rz).to(tl.pointer_type(z)) tl.store(z%_ws + (xindex * RSPLIT + rsplit_id), rr TrL_peersr=r r z = tl.load(z_ws + (xindex * RSPLIT + rsplit_arange), rsplit_mask, eviction_policy='evict_first', other=triton_helpers.if_mask(rsplit_mask, rx)rrrEr rr>rrr$rRcreate_cse_varr"r rrQrk) rRrnr default_valxnumelr)r6r<r=r s &&&& rr 7TritonKernel.codegen_cooperative_reduction_peer_combinesRS!(,(@(@(B(B ..(4??+<<!GGPPQWX %% GG9C0A0A)0L/MMbcnotcubvw$%J:,VXY]X^_   & ##l& !X&&&( $  &&g[ -eers~eeAAC D  rc$V^8dQhRRRRRR/#)rrrrKrrr_r)rs"rrrs-@@@@ @rc 6VP'g \R4hRVn\PP V4pVP VRVP VVRRR7R7pRVnVPPV4p\P!4pVP'd+VPVPWP44VPVPV4\!V\"\$34'd\VPP'\)VVP+VVVP-V4VRVP/4: 2444EM\!V\04'g\R\3V4 24hVP4p\7V4'dVP8e\:;QJd&RVP84F 'dK RM RM!RVP844'g7R P=\?\@VP844p VR V R 2, pVPP'\)VR V R V RV R VPB R 2 44VPE4R#)r* FT)rrr rkrrrNc3>"TFp\V4R8HxK R#5irrrs& rr/TritonKernel.store_reduction..#rrrrrr'r5r)#rr rXrrrrrrrrrrrrrhrr=r0rQr`r-rrrErrFrrrrqrsrrtr) rRrrKrrrrrrrs &&&& rstore_reductionTritonKernel.store_reductions; $$$ !<= = %!!$'== &*&H&H 'I'!  !%iit$))+  % % %  $ $,,T3G3GH  &&t';';SA h2I J K K  * *55  ,+H,C,C,E+HI   h88$'Ed8nEU%VWW#--L%e,,KK+?5;;??5;;???"iiC(=> . Q ??  * *uDc%8CTCTBUUVW  rc(V^8dQhRRRRRRRR/#)rrr_rrrrrr)rs"rrr1s2     0      rc rVPPVPVPW4VVR7#)z=Emit a tl.reshape into a new tile shape, returning a CSE var.r)rrr _reshape_expr)rRrrrs&&&&r emit_reshapeTritonKernel.emit_reshape1s;xx  LL   u , !  rc 0V^8dQhRRRRRRRRR R R R/#) rrr_rraxisrrrrrrr)rs"rrr?sF            rc  |\V4pVPPVPV RV RV R2VVR7#)zDEmit a Triton reduction primitive along `axis`, returning a CSE var.rrrr)rrrr)rRrrr rr reduce_fns&&&&&& r emit_reduceTritonKernel.emit_reduce?sM2.A xx  LLk5'D6 + !  rc4V^8dQhRRRRRRRRRRR R R R/#) rrr_r(rrrrr out_shaperrr)rs"rrrPsP   > :  6    !   rc VPW4pRV R\V4 R2p\V\V4\V44p VPP VP V VVR7#)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 rXrrrrr) rRrr(rrrr reshaped broadcastedr=s &&&&&&& remit_broadcast_via_reshape'TritonKernel.emit_broadcast_via_reshapePsx%%eAxj+;O+L*MQ O k4+@${BSTxx  LL  !  rc$V^8dQhRRRRRR/#)rrr_rrrrr)rs"rrrks+HHH/H Hrc \VRR4pVfRV R\V4 R2#\\V4\ V4\ V44#)rNr[rr)rrXrrr)rrr s&& rr TritonKernel._reshape_exprjsO E7D1   r*:5*A)B!D Dc%j$y/4;GGrc$V^8dQhRRRRRR/#)rrtuple[CSEVariable, ...]dtypestuple[torch.dtype, ...]rrr)rs"rrrts)>S>S1>S;R>S >Src  aa a a aaa\4o S PR4\4o \^4Uau.uF[o\;QJd*.V V3Rl\ \ W#444FNK 5M#!V V3Rl\ \ W#4444NK] ppRPR\PPV444pS PRV R24\4oRo^RI H p^RIHpV!4oV!4o !V V V VVV3R lR \ 4p S P#4;_uu_4\$P&!V !44;_uu_4V!V!p RPR V 44p S PR V 24R R R 4R R R 4VP(P+S P-4SR7#uupi +'giLH;i +'giLS;i)z @triton.jitc3t<"TF-wpwr#SPRS RV 2W2PR7xK/ R#5i)rrTrN)rr)rrrrrr,s& rr,TritonKernel._lift_helper..~s?)G%A~ s1#Qqc]%{{ KK)Gs58rc38"TFp\V4xK R#5irrrs& rrr sR.Qc!ff.Qrz def {name}():rr})ShapePropagationOpsHandlerc4<]tRtRtRVVVVVV3RlltRtR#)+TritonKernel._lift_helper..CSEProxyic(V^8dQhRRRRRRRR/#) rrrrrrrjrrr r)rs"rr8TritonKernel._lift_helper..CSEProxy.__annotate__s,   '6 @N  rc <S RV 2, o \SV4!V/VBp\S V4!V/VBpSPS\S V4!V/VBVVR7#)rTr)rr) rRrrrjr output_shaper dtype_handlerhelper helper_name overrides shape_handlers &&&& r_default4TritonKernel._lift_helper..CSEProxy._defaults4&z) &!  #" # '!  #" # ||It,d=f=&& $rrN)rrrrr r)rr r r r r srCSEProxyr s   rr c38"TFp\V4xK R#5irr)rrs& rrr sB'F 'rreturn N)r)rarQr^rrrMr"rqrr from_iterablerrur~!torch._inductor.shape_propagationr r-rrXset_ops_handlerrrgetvalue)rRrrr r,r signaturer~r r outputsrr r r r r s&&&&` @@@@@@r _lift_helperTritonKernel._lift_helperts  !'e1X   E)23v3F)GEE)23v3F)G   IIRioo.K.KD.QRR = 267#% * PP24 24   ~ 0]]__a// ;;$iGiiB'BBG   wwi0 1<_ $$(():k(RRk `<;__s0G AG=$G'!3G G' G$ G'' G7 c(V^8dQhRRRRRRRR/#)rr r r zUCallable[[tuple[CSEVariable, ...], tuple[CSEVariable, ...]], tuple[CSEVariable, ...]]rr rr)rs"rrrs8E"E"'E" E" ( E" !E"rc  0aaSP'g \R4hSP'd \R4h\RSP44pSP V4\ V4pSP'd \R4h.p.p\;QJd.RV4FNK 5M !RV44p\P!SPPSP4pSPW#V4pSP4SP , p \#W14EFawrSPPSPV R\%V 4 R2V V P&R7p SPPSPR V R SP)4 R2V \SP+44R7p VP-V 4\/V 4p SP0'dKSP+4pR VR&SPP3WR7pR R P5V4 R 2pV P6'dRMRpSP8P;V RV R V R V R24VP-V4EKd RoVV3RlpV!RS!V4 RV R V R2VVV4pSP0'gRpVUu.uF+pV!RV R2\=VP>4V!V4R7NK- ppV!\V4\V44pV!\V4V4p\#VV4UUu.uF,wppV!RV R V R2VP>VP&R7NK. ppp\#VVV4F+wpppSPP;V RV R V R24K- MTpVF@p\AV\B4'g\R\EV4 24h\V4Vn#KB \V4#uupiuuppi)z* Perform an associative scan on 'values'. r* TODOc3@"TFqP4xK R#5irrrs& rr$TritonKernel.scan..r- rOz(ops.scan not supported inside ops.maskedc38"TFp\V4xK R#5irrTrrs& rrr+ Fve*511vrrrrrrrrUrVz float('nan')r7rc2RPRV44#)rc3*"TF q R2xK R#5ir=Nrrrs& rr1TritonKernel.scan..csv...csv88.cse_multiple..#LI488$$Y//r FTrrJ rrrrr r"rPrrrQrHrr)r=rrr rr, cache_keysrrr result_varsrnr8 rRs&&&& r cse_multiple'TritonKernel.scan..cse_multiples#F A;@8D8aF"QCr%18JDsLLsssLLLLAKLI Y/LL'*&&9&9NUe;;?&9  LL " "{#$Cv. *-[)E% +0( Y3*F% %ELE 2$E'1Eztl.associative_scan((rcZVPfR#\VP4pRVR&V#)Nrr)rr)rrs& r_partial_scan_shape.TritonKernel.scan.._partial_scan_shapes*99$ OE #E"I Lrztriton_helpers.select_one((z1), rbase == (RBLOCK - 1), dim=-1, keep_dims=True)ztl.where(roffset > 0, z = tl.where(roffset > 0, rr)$rr rrrrrrrfrrrrrrr% r rr"r=rrr.r_rJrrPrqrrlrQrTrrr_rrH) rRr r rrbroadcasted_values accumulators cse_computecombine_helper_fnrrr value_dtyper reduced_sizer reduced_size_strrrE partial_scan_varsrI partial_scan_varpartial_reduce_vars accs_nextfull_scan_vars full_scan partial_scanrD acc_nextpartial_reducernr8 s f&&& @rscanTritonKernel.scans.$$$ !<= =  % % % ( (I8H8HII % u  ??? !KL L FvFFvFF''(9(94<<H  --j&I%%'$*A*AA/LE((++ '1%89;kk ,K HH%% ";-r$2E2E2G1HJD0023 &E  % %e ,&u-H,,,#335 #& R "hhooEoN %&tyy'>&?q#A ,1,C,C,C. ##"m;/?.@7)2hZWXY##K090< = &$)#C(:$;#.D.DE-.>? ): ##5#6>Q8RSI'l(;=NON03>CT/U  0V+I| ,YKr,qI&,,&,, 0V :=<)<:5+~ &&"m#.Fr- rOz(ops.sort not supported inside ops.maskedz3ops.sort is only supported in persistent reductionsc38"TFp\V4xK R#5irr- r. s& rrrb Qr/ rz)expected len(dtypes) == len(values), got r^rrrrc2RPRV44#)rc3*"TF q R2xK R#5ir2 rr3 s& rr1TritonKernel.sort..csv..`r5 r;r6 r7 s&rr8 TritonKernel.sort..csv_r: rc <\V4p\V4Uu.uF qP RV RV 2NK pp\;QJdV 3RlV4F 'dK RM RM!V 3RlV44'd,VUu.uFpS PP V4NK up#\ W14UU u.uF+wrS PP WPR7NK- p pp S PPS !V 4 RV 24\ W4F.wrV'dW+n S PPW{4K0 \V 4#uupiuupiuup pi)rc3Z<"TF pSPPV4xK" R#5irr= r? s& rr:TritonKernel.sort..cse_multiple..erA r FTrrJrB )r=rK rr rr,rC rrrrD rnr8 rRs&&&& rrE 'TritonKernel.sort..cse_multiplebs%&'A;@8D8aF"QCr%18JDsLLsssLLLLAKLI Y/LL%($C$CLEe;;?$C  LL " "{#$Cv. *-[)E% +0( Y3*F% %ELrG z#last range tree must be a reductionrrnumelztriton_helpers.sort_with_index(z , stable=z , descending=zUnhandled sortr)rr rrrrrrrfrrrrrrr rrrrMrr.r1r r"rHra)rRr rr^ r_ rrM rr,rrK rE rl r=rD rn input_varr8 sf&&&& @rrTritonKernel.sort;sw$$$ !<= =  % % % ( (I8H8HII % u  ??? !KL L((( !VW W''(9(94<<H %%'$*A*AAFvFFvFF v;#f+ % ;CK=SQW[MZ &f-  . "5'D,?,?,A+B!DiD0023  .   = &$#000 !FG G2243C3CB3GHHh v;! 12DQ2G1HK]^_K`Jab82cU)F8= AO 'tOK !12 2%(%= !J #( ) 0 0J &>[!!_ s;AJ)cV^8dQhRR/#rrr)rs"rrrs $ $^ $rc VP'gR#VPVP4VPP4VPP4R#)z Generate the output from prologue. This should be extracted from the subgraph, which is why this is partitioned from codegen_body. N)rrrar)rRrss&&rcodegen_prologueTritonKernel.codegen_prologuesB }}}  DMM"  !!#rc VP'g]VP'gKVP'g9VP'g'VP'gVP 'gR#VP Uu.uFqP'gKVNK ppVP'Ed@VP'g \R4h/p\VP4FwrEVPp\PP!V\"P$4pVP'\(V4pRV 2pVP*P-V RV R24VP.P1V\"P$R+R7W8&K VP*P-R4VP*P-R4VP*P3^R 7;_uu_4VP54'gxVP ^,p V P6R 8wd\R V P6: 24hV P6p VP*P-V R V P8 R V R24VP*P;VP4VP*P=R.4VP*P;VP4VP*P;VP4VP*P;VP4VP*P;VP 4VP>'dVP*P-R4\VP4EF wrEVP@p RV 2p\PB!VP\"P$4p \EVP4p VP.PGVP*V RV R2V PHR,R7p^RI%pVPLPNPQVRVP*4;_uu_4V !W8,V4pRRR4VP*P-V RX 24EK RRR4\S\UVP44F%pVP*P-RV RV R24K' EMVPV'Ed:\UV4^8Ed)\V4EF+wppVP*P3VR 7;_uu_4VP6pVPX'dRMRpVPX'dRMV R2p\"PZP\'d\_4R-8dRpMRpVP*P-RV RV RV RVPa4 R V R!2 4RRR4VP*P3V^,R 7;_uu_4VPcWP*4RRR4EK. VP*P3\UV4R 7;_uu_4VPeVP*4VP*P;VP4VP*P;VP4VP*P;VP4VP*P;VP4RRR4\g.\V4O4EFwppVP*P3V^,R 7;_uu_4VPhVPj,Pm4EFLwppV\UV4^, 8dVV^,,pVPhVPj,pVV9dcVV,p\nPqV4p\sVPtV4p\wVV4UUu.uFwppVVV,, NK ppp\x;QJdR"V4F 'dK R#M R$M !R"V44'dKVP*P-\{VP|V,V R%V R\~PPV4 R&244EKO RRR4VP.PVP4VP4EK MVP*P;VP4VP*P;VP4VP*P;VP4VP*P;VP4VP*P;VP4VPX'doVP'gVP 'dKVP R'2pVP*P;R(V R)2R$R*7VPP4VP'g&VP*P;VP 4VPP4VPP4VPP4VPP4VPP4VP P4R#uupi +'giEL;i +'giEL;i +'giEL;i +'giEK~;i +'giEL};iuuppi +'giEL;i).z Concat output code from index_code, loads, compute, stores, suffix into self.body. For pointwise kernels, this is called just once at the end. For reduction kernels, this generates a loop over the reduction axis. Nz1Mix order reduction requires persistent reductionaccumz = tl.full([R0_BLOCK], z, tl.float32)[None, :]rz/split_size = min(RSPLIT_SIZE, xnumel - xoffset)z@for _ in tl.range(0, split_size, XBLOCK, num_stages=NUM_STAGES):)rrzexpected entry prefix 'x', got zmask =  < rzxindex += XBLOCKzxoffset += XBLOCKrz, 0)rrJz&tl.store(ws_ptr + (tl.program_id(0) + z3 * tl.num_programs(0)) * r0_numel + r0_index, accumz , r0_mask) rsplit_startr7 rsplit_endz, num_stages = 2rzfor zoffset in tl.range(rrr c3"TFCp\PPPV\P !^44xKE R#5ir+)rXrrrrr)rrs& rr,TritonKernel.codegen_body..*sC+6GG,,DD & a(8+6sA A FTz = tl.advance(rz + tl.program_id(1)zR if HAS_RSPLIT: triton_helpers.x_grid_barrier(z) r )rR0_BLOCK)rz )r4r)I indexing_coderArrrrrrmix_order_reductionrr rMrrr$rr~ rrr` rkrlrQrrrrrrr writelinesrrr rrrunittestmockpatchrrrrrrrrrBrrrrr,rrrrrr"rr`rrXrr$r`r cache_clearrrrAra) rRr loop_trees accumname2varr partial_accumrrrrrrr triton_reduction_functionnewvalr~ r levelr loop_startloop_end num_stagesr advancement prev_treeprev_advancementsprev_advancement prev_block prev_num_itercurprevsem_ptrs & r codegen_bodyTritonKernel.codegen_bodysR     zzz{{{|||%%%### '+'7'7H'7t<5;;W33M7Kse} ##f3G9-+S3,bS!>!>?@ ##$>$>5W!,/"3+6 333+6  % ++( $ 8 8 C#,+^I;bI^I^_jIkHllm n;38N##D$:$:;  "S BV II  T// 0 II  TZZ ( II  T\\ * II  T[[ ) //0  % % %  " " "d&:&:&:--..ABG II  33:)<    6 6 B B D''' II  T11 2   "    $$& ""$[I@POOU,++r433877:990/"877s=m m +E4m" Dm"!m 1-m" 4m6>&m6%A m6%n ;B0n;Cn9<n3 n9&n9:n9A#n9 m m"" m3 6 n n  n0 3n99 o cV^8dQhRR/#rr)rs"rrrWsYrc .pVP4'EdJ.pVPRV.4VEF-p\V\4'dVP \ V44K6\V\ 4'd\\PPPVPVPR7pVP \ V44K\V\P4'dS\PPPVVPR7pVP \ V44EK\R\!V4 24h V#)rrz!Unsupported numel argument type: )radd_numel_to_call_argsrrr_rr{rXrroptimization_hint_with_override inner_exprrrrr2r)rRr numel_argsrhints& rkernel_benchmark_extra_args(TritonKernel.kernel_benchmark_extra_argsWs    ! !+-J  ' 'J ;!c3''KKC)_5577++KK&*&8&8LDKKD *UZZ0077++KK&*&8&8LDKKD *$'Hc %TUU!"" rc V^8dQhRRRR/#)rnum_gbz float | Nonerrar)rs"rrros{{|{{rc \4pVPP4wr4rVVP.RO4VP 4;_uu_4\ P !4p.p\WE4EFOwrR\V4 2p \PPV 4p V 'd\PPPV P4VPR7p \PPPV P!4VPR7pVP#V RV RV R\%V P'44 RV P)4 R2 4EM6V \PP*9d\PP*V ,p\PPPVP-4VPR7p \PPPVP/4VPR7pVP#V RV RV R\%VP04 RVP2 R2 4EM1\5V \64'dj\PPP9V P:VPR7pR V P<9d^pVP#V R V 24M\5V \>4'd\PPA4p\PPP9V P VPR7pVP#V R V R\%V4 RV P2 R24M\CR V 24hVPEV 4EKR VPGVPI44VP#R RPKV4 R24RRR4VP.RO4\PPA4p\MVPN4wppVP 4;_uu_4V'dVP#V4VP#R\PPPPSV4 R24VP 4;_uu_4VP#\PPPPUV44\WVPN4pVP#V RV R24VP#\Y\ZP\4 RV R24RRR4RRR4VP.RO4VP 4;_uu_4V'dVP#V4VP#R\PPPPSV4 R24VP 4;_uu_4VP#\PPPPUV44VP#R \Y\ZP\4 R24RRR4RRR4VP.RO4VP 4;_uu_4VP#R4VP#R4VP#R4VP#R\PPA4P^ R24VP#RV 24VP#R4VP#R4RRR4V# +'giELk;i +'giEL;i +'giEL!;i +'giEL/;i +'giEL;;i +'giT#;i) 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. rarg_r z = rand_strided(rz , device='z ', dtype=rrrJz = torch.zeros(z*Don't find the buffer or const tensor for r r=Nzwith r]z = get_raw_stream(z.run(*args, stream=z.benchmark_all_configs(*args)z!!$x.!12gg,,X677++LL &*&8&8MDWW--NN(&*&8&8OF$$#*$4TF"VHJ_b_m_m_oOpNqqz{~|I|I|K{LLMN!2!22#$77#4#4X#>L77++LL$))+&*&8&8MDWW--NN$++-&*&8&8OF$$#*$4TF"VHJ_k_r_rOsNtt}KQQ~RRST 11"#''"2"2"R"R &*&8&8#S#K% 4&' $$z[M%BC66WW@@BFGG,,LL T5G5GME$$#*OE7*_U[E\D]]fgngtgtfuuvw#DXJO  *k&?l   T==? @   wtyy';& ?ux 9:<<>89M9MN u ]]__  /   uQWW%7%7%D%DU%K$LAN O  GG&&11%82.2F2FG   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   Y__F! __(! ___  sfN:]/6A.^$B^ 7^2A.^? A%^+ ^?B%_/ ^  ^  ^ ^( + ^< 6 ^?? _  _$ c \P!R4P\PP P R44#)zp from torch._dynamo.testing import rand_strided {} import torch get_raw_stream)textwrapdedentrrXrr import_get_raw_stream_asrQs&rimports_for_benchmark_kernel)TritonKernel.imports_for_benchmark_kernels=  &##<<=MN O  Prc VP'dR#VP'dR#VP'd VP'g \ R4hR#VP'dR#R#)rrr* rr pointwise)rrrrr rQs&r_get_heuristicTritonKernel._get_heuristicsY    !  ' ' '*  & & &((($%@AA)  " " "rcV^8dQhRR/#rr)rs"rrrs%"%"#%"rc p\4pVPR4^RIpVPR4\ 4;p'dVP V4VPR4\ PP'dVPR4VP4# \dLi;i)zP import triton import triton.language as tl NzU 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') ) rartriton.language.extra.tlx ImportErrorrrQr#rproton_profilingr" )rimportsr attr_descs& rgen_common_triton_imports&TritonKernel.gen_common_triton_importss!"   , NN 34 49 4   i (   == ) ) ) NN !!+   sB'' B54B5cV^8dQhRR/#)rrr7r)rs"rrr.s  : rc R\P'*RVP4R\PP/#)enable_fp_fusion launch_pdl disable_ftz)r#emulate_precision_castsrRrr rs&rtriton_meta_commonTritonKernel.triton_meta_common-s; F$B$B B #113 600<<  rc /R\PPP4bR\P bR\P bR\PPbR\PbR\PbR\PbR\PbR\PbR \PbR \PPbR \PP bR \PP"bR \P$;'g\P&bR\P&bR\P(P*bR\PP,bR\PP./Cp\P0'd\P2!4VR&\P4P6eRVR&\P8!4'dRVR&\P:'dM\P:VR&\P<VR&\P>VR&\P@VR&\PB'd:\PBVR&\PDVR&\PFVR&V#) backend_hashassert_indirect_indexingautotune_local_cacheautotune_pointwiseautotune_remote_cacheforce_disable_cachesdynamic_scale_rblockincremental_autotune max_autotunemax_autotune_pointwisemin_split_scan_rblockspill_threshold store_cubinr<batch_invariantr>&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#r r rr r r r r r r r r r r<r r=r>r r *write_are_deterministic_algorithms_enabledr rrrr r r r r r r )r inductor_metas& rinductor_meta_common!TritonKernel.inductor_meta_common5s EKK//HHJ &(G(G  #F$?$?  !&--"B"B  $V%A%A  #F$?$?   #F$?$?  #F$?$?  F//  %f&C&C  $V]]%H%H  v}}<<  6==44  V11KKV5K5K  v55 -f.A.A.`.`! " 5fmm6j6j# $ )&--*R*R% *  < < <::< @ A ==   (&*M( #     )-M+ &  # # #171I1IM- .7=7U7UM3 48>8W8WM4 5FF K L  + + +00 5 677 < =>> C DrcV^8dQhRR/#)rrrr)rs"rrrhsbb.brc ^aaRVPRVPRVPRVPRVPR\ VP 4/pVP'dVPVR&\P'g!\PP'd$VP;'g VPVR&VP'dVPVR &VP 'dR VR &VP"'dR VR &VP$'dVP$VR &VP&eVP&VR&VP(eVP(VR&VP*'dVP,VR&VP/4;peW!R&VP0'dR VR&VP34;peW!R&\P4'g-\P6'g\P8'd VP;4R, pVeW1R&\P4'g\P8'dVP=4pVeWAR&VP>'EdlVP@PC4wrVp\ERV44o/oVP,'dVPFF|pVPH'gKVPJPM4 R2pVS9gK=VPOVPP4p VPR'd \UV ^4p V SV&K~ VV3Rlp Rp /p VP>P[4FmwrVP\Uu.uF q!V4NK ppRTRVP^Uu.uF q!V4NK upRVP`Uu.uF q!V4NK up/W&Ko WR&V# \V\X3dEK2i;iuupiuupiuupi)zs Used by both standalone codegen_kernel() and ComboKernel.combo_grid_meta() (which calls this on each sub_kernel). rrnum_load num_store num_reductionr RSPLIT_SIZE"has_loadstore_with_contiguous_rdimrTrrr min_xblock min_rblockrr3add_persistent_rblockrgeA kernel_num_gb kernel_flopc3b"TF%p\VR4'gKVPxK' R#5i)rN)rr)rrs& rr8TritonKernel.inductor_meta_per_kernel..s"'$-Sf1EIs//rc<\V4pVS9dV#VS9d SV,#\V4# \\3dTu#i;ir)rr TypeErrorr2)rbs_str fixed_blocks sig_arg_namess& r_resolve_block_dimATritonKernel.inductor_meta_per_kernel.._resolve_block_dimsSAM) LL('..!q6M!:.! L!s 2AAc^\V4# \\3d\T4u#i;ir)rr r2r)rbs&r_resolve_tensor_dimBTritonKernel.inductor_meta_per_kernel.._resolve_tensor_dims-"q6M!:."q6M"s ,,rrrr)1rrr r r setrr| rsplit_sizer#r<r=r>rrqrrrrr r rr$_get_native_matmul_persistent_rblockr rybenchmark_kernelr benchmark_combo_kernelestimate_kernel_num_bytesestimate_flopsrrr rrr1rrBrrr-rr r2rrrr)rRrrJr flopsrTr# rt block_namerr r resolvedinneroptsrbrr r s& @@rinductor_meta_per_kernel%TritonKernel.inductor_meta_per_kernelhs   5 5   T// c$"5"56   # # #!%!1!1C     6#6#6#U#U#U227766 4 5  # # #)-)A)AC% & ==="C O    %)C! "    #'#5#5C ?? & $C  ?? & $C   % % %*.*C*CC& '??A AF N5;1 2  % % %+/C' (--/ /F <'-# $  # # #''',,,335;F!'-O$  " " "f&C&C&C'')E %*M"  ( ( (!%!9!9!; AqQ&'$-'M,.L(((**B(* (9':%%@ %]:%&*&J&J288&T#'#8#8#8*-c2,C;> Z 8+ ! " H#<<BBD 7;7G7GH7G!*1-7GH!4djjIj1!4jI M 1 3A 6 M# E/7* + ;%.z#:% $%,IJMs*(,PP#P P% 'P*PPcV^8dQhRR/#rNr)rs"rrrs00t0rc VPP4;'dVP'*pVPp\ VP 4^8HpV'EdV'EdVPP VP 4pVPPP^,pVPpVe/RV9d(RV9d!\W PP4pM,VPPV4\P8HpVP PP"pVPPP"p \$P&P(P+V4\-\$P&P(P+V 4^4, p V R8dV'd\$P&P(P/VPPR4'dN\$P&P(P1VPPR4'd V^ 8:dR#R#)rrr0_g?iiTF)rr1rrrr memory_stats persistentmemoryrcount_per_threadrsr|get_reduction_hintr4INNERloopedbytesrXrroptimization_hintrstatically_known_leqstatically_known_gt) rR looped_redr two_d_redr$ dim_statsmem_ops_per_threadcontiguous_red looped_mempersistent_memsaved_bytes_ratios & rr "TritonKernel.add_persistent_rblocks]]//1SS$:S:S6S **  $) :))==55dkkBL$//66::1=I!*!;!; )=(]*!F!==#@#@" MM44]C$**+ &,,3399J)44;;AAN ! 0 0 B B: NQT  22>BR! "S("GG$$99MM115GG$$88MM114'",rcV^8dQhRR/#rr)rs"rrrsUU3Urc pa&a'\4p/pVPP4FkwrE\V4'dVP'gK*\ P PPV4p\\V44pWsV&Km VfVPVP44\ P P4PpVR8XdVPR4MVPR4\P 'd VPVP#44VP$P'4wo&p o'p \)S'4Fwr\+V \,4'gK\/\0P2V P44p V \ P PP69gKl\-V P8\ P PP6V ,4S'V &K \;4p VP<EFBpWP$P>9d-V PAVP$P>V,4WP$PB9dtV\ P PD9dUWPD9dEV PA\/\FVP$PBV,4PH4WP$PJ9gKVP$PJV,p\+V\L4'd \OR4hV PAV4EKE \QS&S'4FZwpp \+V \R4'gKV PT\VPX8XgK?V PAVP84K\ \[V 4p VP]4F[p\-VP^ R2VP`4pS'PcV4S&Pc\eVP844K] V&V'3RlpVPfF\pVPh'dVPj'dK)VPlfK9V!VP^Po4 R24K^ VPp'd V!R 4VPr'dV!R 4V!R 4\uS'VPvS&R 7p^R IpVPPVR4VV9gK(VV,VPV&K@ VP'gVPPRR4VPVP4\|PPeFVP'g!\PP'g\S'R4RR7.VR&M\S'RR7.VR&VPF%pVPR4VPV4K' VP'd6RVP4 RVPP: R V: R!V: R"2 pMVP'dJVPPVP4p RVP4 R#V: R$V R V: R!V: R"2 pMR\VS'4p!V!eR%V!P8 R&2MRp"RVP4 R#V: R'V" R(V: R!V: R)VP R"2 pVPV4T;'g\\P4p#VPR*V# R+R'PR,S&44 R-24VP4;_uu_4\PP'dVPR.V# R/24VPV4VP$P4Fwp$p%VPV$ R0V% 24K VPVP4\PP'dVPR1V# R/24RRR4\P 'd0VPVPVPR2444VP4# +'giLe;i)5z Convert the TritonKernel from Inductor SIMD IR to triton code, including inductor triton heuristics, imports, metadata, and benchmarking infra. Nrtz"triton_helpers.set_driver_to_cpu()z"triton_helpers.set_driver_to_gpu()z%mutation_arg must not be a RemovedArgrc<\4'dSP\V44SP\VRR74R#)T) is_constexprN)rSr_r]r[)r argdefsr# s&radd_constexpr_arg6TritonKernel.codegen_kernel..add_constexpr_argYs2-//  h!78 NN78$? @rrr r NUM_STAGES) size_dtyper<  _coor_enabledr# r r native_matmulztl.dotTlaunch_cooperative_grid grid_type kernel_namemutated_arg_names optimize_memr)pointer_range_override'skip_cpp_wrapper_input_tensor_alignmentconfigs)rJ rz$ @triton_heuristics.z( config=zI, filename=__file__, triton_meta=z$, inductor_meta=z; ) @triton.jit z!( size_hints=z%, reduction_hint=ztile_hint=TileHint.r=rzH filename=__file__, triton_meta=z*, min_elem_per_thread=zdef rc3@"TFqP4xK R#5ir) full_namers& rr.TritonKernel.codegen_kernel..s*J'Q;;=='rOr zpl.enter_scope("r,rJzpl.exit_scope("r )rrrrr)rrarrrMrrXrrr, r9rrr rrr#r r rr rMrrgr rr@rinv_precomputed_replacementsrr mutations input_buffersrr(rrb inner_nameoutput_buffersrfr r"ri zero_moderj ZERO_ON_CALLrrrkrrr_r[rr1rr rBrr| ryrq"torch.fx.experimental.proxy_tensorrB rr r7r1rrrrC rrlrr is_inference is_backward_get_grid_typerrKDESCRIPTIVE_NAMEr r rurr rq r_r rrrzrrremit_pointer_range_32rtrrQrr rr( rrwrr rqrr codegen_static_numelsaliasesr r r" )(rRrrs size_hintsrr numel_hint size_hintrzrTr,rr% mutated_argsmutation mutation_argargnamersizeargr= triton_meta_signaturerB props_devicerrH r arg_num tma_fieldsfinal_kernel_metarr heuristics_linereduction_hintr tile_hintrF oldnewr< r# s(&& @@rcodegen_kernelTritonKernel.codegen_kernels0  ![[..0MF"6**43H3H3H));;EBJ'J8I!*v  1 < KK668 9''==?DDKe# @A @A&&& D==?@#'99#;#;#= Iq *FA#w''ellCHH5QWW--JJJ#*!''"2"2"O"OPV"W$IaL +)3 H99222  !8!8!BCII555AGG$;$;;$8$88  )B)B8)LMXX99333#yy77A lJ77()PQQ  .'8 3LGS3 --MM%6%C%CC  . 4l+ ++-D U3TZZ@G   W % NN77<<0 1. A $$D   T%>%>%>& !2!2!4 5U; <%  % % % h '  # # # m , l + 1 $"2"2G!  Eww::< ??!<< (9(9:L"& 2*11,?ROO**11??VV!S^3TTx3t||CT7T ))+  #   % % %6:K1 2 ww++BBqww/B/B  ,,.77 3{;;<  L  ++-  '')  +95G@AK $Yw%7%<%< =6'* dii( 002  !99;E    " "5$ /)),=e,D""5) }}}    " "#8$ ? # ==   (  ! ! !)L)L)L+-<@&K ")TR&K "++F NN2  KK ,    #$$($7$7$9#:; --4478!!,0##0"34 O " " "!]]==d>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 [[]]}}---!1+bAB  & &t , II--/S#c#/00 KK "}}--- R@A  " " " KK--!%%o6  }}#]s7B;l%3l%% l5 c v\PPPV4p\ V\ P \34'd\V4p\V4pV#^p\PPPW4'g!VR8d\RV 24hV^,pKOV#)rz!Failed to find static RBLOCK for i@) rXrrsimplifyrrrrr9r- r2)rl rs& r_get_persistent_RBLOCK#TritonKernel._get_persistent_RBLOCKs!!**62 fu}}c2 3 3f+C!#&C Cgg&&;;FHH?$'H%QRRqJrcV^8dQhRR/#rr)rs"rrr's 7 7j 7rc ^VP'd%VP'dVP'dR#VPUu.uF;pVP'gK\ VP VP44NK= ppV'd\P!V4#R#uupir) r-rrrr1r3rt rrr)rRrrblockss& rr 1TritonKernel._get_native_matmul_persistent_rblock's%%%,,,)))(( (   U +D,G,G ,S T(  &-tyy!6$6  s B* (B*cV^8dQhRR/#rr)rs"rrr6s333rc VP4;pe=\PPP W!4'dV#\ R4hVP V4#)NzQpersistent strict sum requires its planned reduction block to cover the reduction)ryrXrrrr rt )rRrl rJs&& rr,TritonKernel._get_persistent_reduction_block6s[--/ /F <ww44VDD  ) **622rc V\PV4R# \dR#i;i)TF)r$rt r2)rl s&rr"TritonKernel.has_persistent_RBLOCK@s*   / / 7  s  ((c$V^8dQhRRRRRR/#)rkernel_featuresr kernel_kwargsrrrr)rs"rrrIs)CC+CCM: ;11/2Q2QRR  !@AA$V>CM9 : Src hRRlpVPEFpVP'dVP'dk\PP P VP4pV!V4'd*VPVP R\V4 24VP'dVP'dVP'd2VPVPVP44pRV R2pM9VPVP4pVP 'd \#V^4pVPVPP%4 RV 24VPR8XgEKtVP&'gEKVPR4EK R #) ay We get a small speedup from hard coding numels if they are static. This code stomps on the passed-in values by writing an constant to the top of the kernel. In a kernel like: def KERNEL_NAME(in_ptr0, in_ptr1, out_ptr2, xnumel, r0_numel, XBLOCK : tl.constexpr, R0_BLOCK : tl.constexpr): We would add xnumel = 4096 r0_numel = 768 After the signature, before the kernel code, if we decided to make these static. As its hardcoded, it becomes a better signal to triton on how to unroll and do some static indexing. So, it's not so much that downstream knows that its a static numel, as that you just plop a constant into the kernel. c V^8dQhRRRR/#)rrrrrr)rs"rr8TritonKernel.codegen_static_numels..__annotate__ns : :J :4 :rcB\V\P\34#r)rrrrrs&ris_static_integer=TritonKernel.codegen_static_numels..is_static_integernsdU]]C$89 9rznumel = z*triton_helpers.constexpr_next_power_of_2((z + RSPLIT - 1) // RSPLIT)zBLOCK: tl.constexpr = rzXBLOCK: tl.constexpr = 1N)rr1rrXrrrs rrQrrrrrr9rr-rrBr)rRrsr rsimplified_tree_numelrrs&& rr\ "TritonKernel.codegen_static_numels\s>$ :$$D$$$(=(=(=()(8(8(A(A$**(M%$%:;;NNdkk](3?T;U:V#WX   T%>%>%>--- JJt';';DJJ'GHEFugMfgC>>tzzJC,,,!#rl$++"3"3"5!66LSERS{{c!dmmm9:'%rcV^8dQhRR/#)rrz type[triton_heuristics.GridExpr]r)rs"rrrsCC @Crc \VPUu.uFp\VP'*4NK up4pVP'd&V^8wd\ RV 24h\ P#VP'd&V^8wd\ RV 24h\ P#V^8Xd\ P#V^8XdP\\VPVP44'd\ P#\ P#V^8Xd3VP 'd\ P"#\ P$#\'RV 24huupi)rzexpected n == 1, got z"Unsupported number of dimensions: )rrrr1r| r r.MixOrderReductionGridrCooperativeReductionGridGrid1DrrsrGrid2DWithYZOverflowGrid2Dr-BatchMatmulGrid3DGrid3Dr2)rRrrs& rrY TritonKernel._get_grid_types) 8H8HI8H***+8HI J  # # #Av$'avfXNNrc$V^8dQhRRRRRR/#)rrrorr rrr)rs"rrrs$::(:14: :rc fVPpVP4p^.V,pRV RV RV R2#)rrrr)rqr )rRrrrqrrds&&& riteration_ranges_scalar_code)TritonKernel.iteration_ranges_scalar_codesC&& &&(sTz$r%;-q99rc V^8dQhRRRR/#rr)rs"rrrs.Acrc VPf \R4hRVP R2pVPV4'd(RV RVP^, RVP R2pVP'd4VP 4^8Xd^^^^^^/pRW1P, R2pVP P W"4pVPR8wdV R VP R2#V#) Nzentry.grid_dim must not be Noneztl.program_id(rrz + tl.program_id(z) * tl.num_programs(rxrr)rr rr-r  pid_cacher rq)rRrrreversed_pid_mappids&& riteration_ranges_get_pid%TritonKernel.iteration_ranges_get_pids >> ! !BC Cu~~.a0  & &u - -cU+ENNQ,>+??STYTbTbSccefC  T%<%<%>!%C !1aAq1 "#3NN#C"DAFCoo!!#+   z )U$t//02 2 rc V^8dQhRRRR/#)rrrorrr)rs"rrrs   ,?  D  rc bVP'dVP'dR#VP^8H;'dtVP'*;'d[VP'*;'dB\ P PPVP\44'*#r) rrrhas_zdimrrXrrr- rr8)rRrs&&rr#TritonKernel.needs_yz_grid_overflows    D$=$=$= NNa  Y YNN" Y Y... Y YGG$$99%++~GWXX  rc V^8dQhRRRR/#)rrrrrr)rs"rrrs0000rc VP'd%VPVP4 R2,#\VP4,#)r)rrBr5)rRrs&&rrTritonKernel.max_blocks>    $$ '7u%=> > //rc V^8dQhRRRR/#)rrrorrr)rs"rrr#s<<':<t<rc VP4'gR#VP'd=\PPP VP ^4'dR#VP'gR#VP'dcVPP4 R2VP9d7VPVPP4 R2,^8XdR#MrVP'ga\PPPVP ^4'd'VP'dVP'gR#VP'dvVP'ddVPVP 4p\PP!4p\#V4p\%V\&4'd W$8dR#MAVPR8XdVP('d^pMVP+VP4pVP'd)VP,'dW P/4,p\PPP1VP V4'dhVP2^8g;'gQVP4;'g=\PPP7VP \944#R#)FrTr)supports_constant_maskr-rXrrstatically_known_ltrrrrrBrrr1rrrr2rrrrrr rrr r- r8)rRrrr warp_sizes&& rr TritonKernel._has_constant_mask#s&**,,  ww33DJJCC!!!    DKK$5$5$7#8!>$BSBS!S  DKK$5$5$7#8!>?1DE%%%ww77 AFF)))d.G.G.G    !:!:!:<U>UV V t  V   sAc V^8dQhRRRR/#rr rarrr)rs"rrrusFF~F$Frc VPUu.uFq"P'gKVNK ppRP\RV444pVP RVP V4 24VPUu.uF,pVP'gK\ PV4NK. pp\V4pVP RVP V4 24R#uupiuupi)zN Generates code that flattens ND reduction numels, block sizes, etc. into 1D. r1c3>"TFqP R2xK R#5i)rN)rrs& rr8TritonKernel.codegen_reduction_numels..{s"U_Tkk]%#8_rz rnumel = zRBLOCK: tl.constexpr = N) rr1rqrrrrrrrO)rRr rreduction_treesrl rn_blocksrJs&& rr%TritonKernel.codegen_reduction_numelsus -1,<,<R,c V^8dQhRRRR/#)r multi_indsrrrr)rs"rrrs--5E-*-rc 8VP4p\W!4#)z; Compute linear reduction indices from N dimensional ones. )r rN)rRr coeffss&& rr'TritonKernel._flatten_reduction_indicess113,,rc V^8dQhRRRR/#r r)rs"rrrs???4?rc &VPRRRR7pVPRRRR7pVPV4pVPRVPV4 24VPV4pVPRVPV4 24R#)zH Generates code that converts ND reduction indices into linear indices. rTrrKz roffset = z rindex = N)rrrr$)rRr  rn_offsetsrn_indsrrindexs&& rr&TritonKernel.codegen_reduction_indicess 00 d1 --gtQU-V11*= 4#4#4W#=">?@009 $"3"3F";!<=>rc$V^8dQhRRRRRR/#)rrrorsrarrr)rs"rrrs-4M4M"4M4M  4Mrc  VP4'Ed*VP'd0VP'dVP'd \ R4hVP 4F9wr4pV'dRMRpVP V V RVPV4 24K; VP VP RVPVP44 RVPV4 24VP VP4 RVP RVPVP4 24R#VPpVP'd&VP VP RV RV R24MVPfEVP VP RVPV4 24VP V R 24MVP eV RVPV4 2pMVP#W R 24pVP$'gVP'4 R 2MR p VP)V R VP+V4 RV 2VP RV 2.4VP-V4'd#VP VP/V44R#VR8XdVP$'g5VP VP4 RVP RV R24R#R#)z=derived reduction roots do not support cooperative reductionsz: tl.constexprrrJrru Nz offset + rz offset = 0rrr z offset = r1rr)has_custom_codegen_headerrrr1r named_constantsrQr$rr rr*rrrrr r r| rBr} r r r) rRrrsrr constexpr annotationrr=rs &&& rr,TritonKernel.iteration_ranges_codegen_headers  * * , ,***---&&&$S).(=(=(?$91:- #zl#d6G6G6M5NOP)@ 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%OP88#V M,0+C+C+C1779+U#  OOc4#@#@#G"HJ<Xzzl#dV,   " "5 ) ) NN444U; <s(t777 NNeoo/0EJJr;rErIrNrdrlrRrhrorzrrrrrrrrrrrr r r r r r r{ re r r| r r r r r r r r% r[ rrq r r r r r rrr r r r rcached_propertyr rp rt r rrr r\ rY r r r r r rr r rrr rrr rrr rrrrrmrns@rr$r$ sd&I%%).E &.O$;!3B0B"#@HL7K?9?9BOO((R' 8  %1   * % N::6"4Hl8  " $L,:0  11  ""h /3 h  h  h h EIh "h TI'V1NfUXM8  1 "*H8J     -9> -D1  "(T   J4&P kZbHI 9v. 9 (y v&$ 0I V#0& ! B@D   " 4HH>S@E"NJ"X $A%F0{z P t_%"%"N  00dbH00dUn 73CC&(;TC, -). ,2 &O,:0  0 <|. "  F$W  -? 4M4Mrr$cZa]tRtRtRtRV3RlltRRltR RRlltR R ltR t V;t #) FusedUserDefinedTritonKernelizz When fusing a user-defined triton kernel with epilogues, we use this class to generate the modified triton kernel source c(V^8dQhRRRRRRRR/#) rrrrrscheduler_noder;rrr)rs"rr)FusedUserDefinedTritonKernel.__annotate__s8"P"P%"P%"P= "P  "Prc  b<\SV`VV^RRRRR7W0n\VPPP \ P4'g6\R\VPPP 4 24hVPPP Vn VPP4'g \R4h\VPPP4^8wd6\R\VPPP4 24h\P !VPPP^,P"4VnVP$P'RR 4VnR#) rTNF)rrrrrrz&expected UserDefinedTritonKernel, got z ir_node epilogue must be fusableexpected 1 store, got r r)rdrer r kernel_noderr$UserDefinedTritonKernelr rir_nodecan_fuse_epiloguer kernel_storesrastunparsestore_value_nodeoriginal_stored_exprrO)rRrrr rfs&&&&rre%FusedUserDefinedTritonKernel.__init__sc   !!  -    + + 0 0"2L2L  !8d>Q>Q>]>]>b>b9c8de 483F3F3R3R3W3W ||--// !CD D t||))00 1Q 6 (T\\-G-G-N-N)O(PQ %(KK LL & & - -a 0 A A% !%)$=$=$E$EdB$O!rc$V^8dQhRRRRRR/#)rrrrKrrr_r)rs"rrr s"Z4Erc d\VPP4^8wd,\R\VPP4 24hWPP^,P 48XdVP pVP V4p\PPV4pVP'dVPpM\PVP4pVPPVP W5VR7pV#\RV R24h)rzexpected 1 mutable arg, got rz!Epilogue attempted to load from 'E'. Inductor indexing variables are not defined in user kernel scope. )rr mutable_argsr get_namer rrXrrrLrrrKrrrA)rRrrK loaded_exprrrrrns&&& rr!FusedUserDefinedTritonKernel.loads t||(( )Q . .s4<<3L3L/M.NO  <<,,Q/88: :33K}}U+HGG%%d+E$$$ --%55hnnE** KE+J !3D6:UU rc ,V^8dQhRRRRRRRRR R /#rr)rs"rrr -s4 *3>FO rc \VPPP\P 4'g6\ R\VPPP4 24hWPPPP48XdVPPPPp\V\P4'g\ R\V4 24hVPpVPPVPV R\V4 R2VVP R7VnR#\ RV R24h)zexpected ComputedBuffer, got zexpected ir.Layout, got rrrz Epilogue attempted to store to 'r N)rr fused_epiloguerr$ComputedBufferr rr layoutLayoutrrrrrArnew_store_cse_var)rRrrKrr4r r,s&&&&& rr"FusedUserDefinedTritonKernel.store-s1$--<<AA2CTCTUU /T5H5H5W5W5\5\0]/^_  &&55::CCE E((77<<CCFfbii00$'?V ~%NOO ,,K%)XX%6%6 '/ <=Q?!kk &7&D "!24&9UU rcV^8dQhRR/#rr)rs"rrr Hs=(=(=(rc V;_uu_4VPVPPP44pVPPP V4RRR4\ P !VPP4p^RI H pH pV!V4p\VP4^8wd"\R\VP4 24h\P !VP"P$4pRRlpV!VP^,P&R^V4\P(!V4pVP+4p V!V4pVP^,P&P,^, p RVP^,P&P.,p VP0P34U u.uF qV ,NK p p VP4P34U u.uF qV ,NK pp V RV V ,V,WR,pRP7V4# +'giEL;iuup iuup i)N)identify_triton_storesidentify_triton_stores_from_astr c$V^8dQhRRRRRR/#)r call_nodezast.Callr rpositional_indexrr)rs"rr:FusedUserDefinedTritonKernel.codegen..__annotate__]s& , , , ," ,rc\VP4V8dW0PV&VPFpVPV8XgKW4nK R#r)rrkeywordsrr)r( r r) new_argkeywords&&&& r _replace_arg:FusedUserDefinedTritonKernel.codegen.._replace_arg]sF 9>>"%553:/0%--;;(*$+M.rrrr )split_and_set_rangesr r get_rangescodegenrdeepcopyr kernel_ast*torch._higher_order_ops.triton_kernel_wrapr% r& rrr r Namer! r store_noder splitlineslineno col_offsetrA get_lines_refrrq)rRrnew_astr% r& r new_store_value_noder/ src_with_store_replaced src_linesstore_line_index indentationsry load_lines compute_lines new_src_liness& rr3 $FusedUserDefinedTritonKernel.codegenHs  T22##22==?J    . . 6 6z B -- 7 78 8@ }## $ ) #9#m>R>R:S9T!UV V"xx(>(>(C(CD ,    # . .r<debug_triton_codedebug_device_str)rRrcrrfs&& rreTritonScheduling.__init__sN #  GIw$?$? OOD$0B CDD(9%$rcV^8dQhRR/#)rr z torch.devicer)rs"rrrK s$$,$rc \PP'g!\PP'd(\ .VP O\ PN4#VP #r)r#rcooperative_reductionsforce_cooperative_reductionsrbackend_featuresr\REDUCE_TO_SINGLE_ELEMENT)rr s&&rget_backend_features%TritonScheduling.get_backend_featuressX MM 0 0 0}}999P#&&P(O(OP ###rc ^a \PPp\W4wrEV'dVP V4\ P 'd^RIHpH o \;QJdV 3RlV4F 'gK RM RM!V 3RlV44'gbVUu.uF&p\Wv4'gKVP4NK( ppVP VP RRPV4 24V'd \VV4p VP!W)4R#R#uupi)r)r:ForeachKernelSchedulerNodec3<<"TFp\VS4xK R#5ir)r)rrrY s& rr3TritonScheduling.codegen_comment..sCPa 1899=sTFz Fused node name list: rN)rXrr rI make_commentr# debug_fusiontorch._inductor.schedulerr:rY rrr commentrqr,write_provenance_debug_handle) rRrUrF r origins_detailed_originsr:r node_names debug_handlerY s &&& @rcodegen_comment TritonScheduling.codegen_comments''&&%8%P"    )      3CP333CP+*!!7!AJJL* $$''>tyy?T>UV BL  1 1+ L s $D*<D*c  VPVVVVVV4'dR#\4p \P4'd\P WS4T P RV: R24V P VRR7\PP4p V P RV P R24^RI H p V !4'd3V'd+^R I Hp \P P#W|!44pR V 2p V!Wa4wrV R V,R ,V,, p VP%WIP'4V 4R#) zDEmit kernel to wrapper, with support for external template handlers.Nzasync_compile.triton(z, '''Tr z''', device_str='z')rA ) cache_dirz# kernel path: r )emit_kernel_overridera async_compileuse_process_poolrrQrrXrrrrV rB 'torch._inductor.runtime.cache_dir_utilsrh ospathrelpath define_kernelr" )rRr rsrc_coderF subs_namerU kernel_pathrIcompile_wrapperrrB rh metadata_commentra detailed_originss&&&&&&&&& r_emit_kernel_to_wrapper(TritonScheduling._emit_kernel_to_wrappers'  & &          (*  ) ) + +   5!!$9)e"LMxt4<<>!!$5n6I6I5J""MN E ??{ I''//+y{CK,[M:$7 $O!D7NT14DDDk+C+C+EGWXrc  Z\PPpWP9dVPV,pV#\P P 'd%\V\P P 4MRpV'd \PPWa4p\V4R,pRPRWvVP4.4p\PP'd \PP RV 2pWTPV&\P P'dTMRpVP!\#\$P&4V4pVP!\#\$P(4V4pVP!RR4p\+\-VP/44R4wrp VP1VVVVVVV \24\4P6!R4'd\4P8!W[V4V#) r:Nr4NrTrtriton_z #pragma CMT#pykernel_metadata)rXrr src_to_kernelr#rdescriptive_namesrHrzcustomize_fused_kernel_namerYrqnext_kernel_suffix aot_inductormodel_name_for_generated_filesunique_kernel_namesrOrrKrZ r r)r(rLrw rIr%is_metric_table_enabledlog_kernel_metadata) rRrq rUrr rF fused_namekernel_categoryrr _basenamerTrs s &&&& rrp TritonScheduling.define_kernel s''&& ,, ,!//9Kf_==222&mV]]5T5TU  YYBB:X @J2NO((?8R8R8TUK""AAA"(!4!4!S!S TTUVaUbc /: ! !( +'-}}'H'H'H iI  ''K,H,H(I;WH''K,C,C(DiPH '' s;H(08>>;K1Ld(S %I+  ( (# ../@AA++KhOrcV^8dQhRR/#)rrtuple[float, str]r)rs"rrrK 8 s   DU  rc VPVRR7p\P!V4pVPWB\ RV44R7#)zk Benchmark fused list of nodes and return the execution time in milliseconds on randomly generated inputs. T)r c3@"TFqP4xK R#5irr rrs& rr9TritonScheduling.benchmark_fused_nodes..@ s:WQVA::<.cache_file_pathW s<<<'()HIIww'' 5a8>IIrc@<S!4p\V\S44R#rr+r)rn r mss rr@TritonScheduling.benchmark_codegened_module..store_cache\ s&(T3r7+rc<S!4p\PPV4'd8\V4;_uu_4p\ VP 44uuRRR4#R# +'giR#;ir)rm rn existsopenrreadrn fdr s r load_cache?TritonScheduling.benchmark_codegened_module..load_cache` sL&(77>>$''dr$RWWY/$$s A'' A8 r %kernel src code for %s written to: %sz*Exception (%s) in compiling fused nodes %srexpected 1 launcher, got c><S!SP!S!^,4#r clone_argsrcallwrapped_jit_functionsrr( =TritonScheduling.benchmark_codegened_module.. sD!5!@!@$!G!JKrr c$<SP!S!#rr rr srr( r s 4 ? ? Frz+The fused kernel for %s took %.3f ms to run)rrXrrzrr rrrLrvr get_argsr rz r Exceptionr#r.disallow_failing_autotune_kernels_TESTING_ONLYr launchersrr n_spillsr/ benchmarkrG r)rRrr rc r device_interfacerr rrr r rr r r r s&f&&& @@@@@rr +TritonScheduling.benchmark_codegened_moduleC sl4AGG4G4GH   # #AGG$G$G$I J JB J  , )4 *i[:Q  II7   #\>s||+? K J ! D<<>D88D#&;;  ()44d;A>?-66I9~"$'@Y@P%QRR|$$'995\<<>!**K!+==>Bk33F"6{B II=  Ms||#g K J ! P (==OOO @ 5\ 3<<''c K J ! P (O K J J ! sn6K22AK K  K20(K I6DK ; K2 K AK 2K 3K 7 K2 K K  K$ K22 L c(V^8dQhRRRRRRRR/#) rr r kernel_args list[Any]r rrlist[TritonKernel]r)rs"rrrK s8QQ+QQ& Q  Qrc VPR4pT;'dQ\;QJd*RVP44F 'gK RM RM!RVP444pVPpV'd ^RIHpTpV'dRVR&VP W4\PPWaW#4pV!V/VBpVPWV4#)r[ c3@"TFqP4xK R#5ir) is_split_scan)rrs& rr9TritonScheduling.create_kernel_choices.. s( -NT   -NrOTF)TritonSplitScanKernelr ) rrscheduler_nodesrI triton_split_scanr r rXrztriton_kernel_kwargsadd_multi_kernel_choices) rRr r r is_scanr rI r rs &&&& rcreate_kernel_choices&TritonScheduling.create_kernel_choices s "--f5  CC( -<-L-L-N( CCC( -<-L-L-N( % +/*:*:  @/K >CM: ;44_T 66 + k;];,,V-PPrc(V^8dQhRRRRRRRR/#) rrr$r r r rrr r)rs"rrrK s23333& 3  3rc  FV.p\PP'gV#VP;'dVP R4'*pVP ;'dVP R4'*pV'd&VP VP!V/VBRR/B4V'dVPPp\PPPVR4'diVP VP!V/VBRR/B;p4V'd:VP'd(VP VP!V/VBRRRR/B4\V4^8d5VR,Fp VPV nK VP!RR7V#)r r FircVP#r)r)rs&rr( ;TritonScheduling.add_multi_kernel_choices.. s q'='=rr)r#r multi_kernelrr rr_rI rr|rXrrr- rmust_keep_buffersr) rRrr r kernelsoptional_persistentoptional_cooperativerl rmkernel2s &&&& rr )TritonScheduling.add_multi_kernel_choices s (.h}})))N$99  -BSBS +C ?  &;;  MDUDU ,E A   NN   #38  __44Fww44VUCC!--$'8=E'5+E+E+ENN(((+rc  aaaaaaa V3RloV3RlpVVV3Rlp^.reRp\PPp\PPp \ V4\Pn\ V 4\Pn\ P ^8p \ P^8p \ Pp VPVRV V RV R7p V EFwrpVUu.uFpVP4NK ppVUUu.uFpVFpVP4NK K ppp\V4^8Xd9VV^,,wppVV, pV^, pVPV4KVf \R4hVP\!\"P$4R 4p\&P(!V4o\*P-R VSP.4V!4wooSe1VS, pVS, pVPSP.4EKESP14oSP2oSP4o S!S P6!S!^,4S P8p\V4^8wd\R \V4 24hV^,P:^8d\=R 4;ooMY\PP?4p\@PB!VVV 3R lVR7o\@PB!VV 3RlVR7o\*P-R\ RV44SS4V!4VS, pVS, pVPSP.4R;o;oo \DPFPI4EK V\PnV \PnWWV3#uupiuuppi T\PnT \Pni;i)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<SPf \R4h\PP SP4^,R,#r r r srr @TritonScheduling.benchmark_combo_kernel..cache_file_path!s<||#$%DEE77##CLL1!4~E Erc<S!4p\PPV4'd\V4;_uu_4p\;QJd0.RVP 4P 44FNK 5M)!RVP 4P 444uuRRR4#R# +'giR#;i)c38"TFp\V4xK R#5ir)r)rrrs& rrNTritonScheduling.benchmark_combo_kernel..load_cache..!s E3Daq3DrNr )rm rn r r rr rKr s rr ;TritonScheduling.benchmark_combo_kernel..load_cache !sz"$Dww~~d##$ZZ2 5 E2779??3D E55 E2779??3D EE Z  Z s A$B22 C cn<S!4p\V\S4R,\S4,4R#)rNr )rn r r ms_clones rr.store_cache!s&"$D s2w}s8}< =rgT)subkernel_nodescustom_part_algorithmenable_autotune mixed_sizesonly_gen_src_coderNzsrc_code must not be Nonerz r r rc><S!SP!S!^,4#rr r srr( 9TritonScheduling.benchmark_combo_kernel.._!s%9%D%Dd%KA%N Orr c2<SP!S!^,#rr r srr( r c!s 4 ? ? Fq IrzDThe fused kernel for %s took %.3f ms to run, %.3f ms to clone inputsc3@"TFqP4xK R#5irr r s& rr:TritonScheduling.benchmark_combo_kernel..i!s@Zzz||ZrO)%rXrrinplaced_to_removerr#combo_kernels_autotunecombo_kernel_allow_mixed_sizes!combo_kernel_per_subkernel_blocksgenerate_combo_kernel_code get_nodesr rr_r rOrrKr r*rrLrvr r r rz r r r rrr/r r accelerator empty_cache)!rR node_listnode_benchmark_resultsr rtotal_ms file_listtotal_clone_msremoved_buffers_originplaced_to_remove_origr r rkernel_code_listrq r node_grouprfused_node_listsrrnamesnode_msrn r r rr r rr r r s!&&& @@@@@@@rr 'TritonScheduling.benchmark_combo_kernel su F   > ) # !ww66"#''"<"<",-A"B%/0G%H" 77!;;;a? %GGQ A#>> )&* /'"&%9 ?  1A,*AK#LDNN$4 #L3CS3C%UU3CSz?a'$::a=$IMGT'H"a'N$$T*#()DEE#++C 0G0G,H)T!&&x0 ;LL *| H>NH"h.N$$S\\2||~xx'*{{$)44d;A>?0:: y>Q&(+DS^DT)UVVQ<((1,$)%L0BWW@@BF%..O%B +44I% H  Z@Z@@   B(*  .5999t2!!--/G1AJ';AGG #)@AGG &22M$MSF';AGG #)@AGG &s+$O,OO "O.I4O O,Prr))r NF) rrrrrr$rI rrr\FOREACH BUCKETIZEINPLACE_BUFFERSMASKED_SCATTER_WITH_INDEXSCANSORTTRITON_TEMPLATESTUPLE_REDUCTIONrT rerrV re rw rp r r r r r rrmrns@rrH rH s?)K)!  " "  $ $  * *  4 4        + +  * *  ::$$MB1Yf6p  c$JQ83j}3}3rrH c V^8dQhRRRR/#)rrr:rrOr)rs"rrry!s((-()(rc.pVP4pVe,\V\P4'g \ R4hV'd3VP f%VP VP4 R24V#^RIH p^RI H pVP4pVf \ R4hVPPV4p\V\W434'g\ R\!V4 24h\"P$P'V4;_uu_4\(P*pVP-VP/44P14pV\(nRRR4VP VP4 R24VP \2P4!XR44V# +'giLZ;i) Nz4multi_template must be None or a MultiTemplateBufferz" Unfinalized multi template buffer)CUDACombinedScheduling)XPUCombinedSchedulingzdevice must not be Nonez]Scheduling backend should be SIMD or CUDACombined when generating debug Triton strings, got: z Triton code:z )get_template_noderr$MultiTemplateBufferr make_kernel_renderr_r 0torch._inductor.codegen.cuda_combined_schedulingr 3torch._inductor.codegen.xpu.xpu_combined_schedulingr r rc get_backendrrrrXrset_current_devicer%generated_kernel_countr r rLr r) rlinesmulti_templater r r backendold_generated_kernel_count triton_codes & rrM rM y!sz E++-N*^R=S=S"T"TSTT.;;C  ((JKL@ L=  " > !:; ;..,,V4 n&<T  !optu|p}o~ WW ' ' / /*1)G)G &!AA eg .HG *0   ( 67 X__[&9: L0 /s A F33 G r;) __future__rr r.rrrrrloggingrrrm r abcrcollections.abcrrrrtypingr r r r rsympy.printing.precedencer rtorch._loggingtorch.utils._pytreer&_pytreeru torch._dynamo.device_interfacertorch._dynamo.utilsrrtorch._prims_commonrrtorch.utils._ordered_setrrSrrrrrrtorch.utils._tritonrrrrutils._sympy.symbolrrr r!utils._sympy.value_rangesr"rr#r$r%rj r' codecacher(r)r*r+rvr, ops_handlerr-runtimer.runtime.benchmarkingr/ runtime.hintsr0r1r2r3r4r5r6r7runtime.runtime_utilsr8r9rcr:r;r<r=r>shape_propagationr? stream_utilsr@rArBrCrDrErFrGrHrIrJrKrLrMrNrOrPrQrRrSrT virtualizedrUrrVrWrXwrapper_benchmarkrYblock_analysisrZcommonr[r\r]r^r_r`rarbrcrdrerfrgrhrirjsimdrkrlrmrnrorprqrrrRrs triton_utilsrtrurvrwrxryrzr r{typesr|rur~%torch.fx.experimental.symbolic_shapesrrrrr getLoggerrrL_logginggetArtifactLogger perf_hint_log schedule_log fusion_logrrrrrrr'rErxr0r=rrXrrfrr!r=rArGrJrMrPrWr]r_rqrvrrrrrC_initialize_pointwise_overridesrrryr)rHrrrdrrUrQrTrrXrjr$r rH rM rrrrG s"    8844 0 $$C<?/YX4))(FF;('.   D6 (CBB/$   H% L>&8 B !00<H ~~//*E ^^ - -h A   , &'R 6 6  4$j j Z $ $ $ N b$b$b$J  ?4??8 =,==@6+LGQMGQT 3 &8 ;P2 2@ d#$( (Vk&kk&`!33T:22D9//9U"OU"p$+$+N R R R j!!H ###U38_$ +S01~,c12 d#$ d#$+-DD%|34* oood \HM:/0\HM~P\(<\(~n3~n3b(r